用户
 找回密码
 立即注册
daxiaoyuyu 该用户已被删除
发表于 2013-11-26 17:49:15
73102
660ti 显卡:(有7个SM)
基本限制:
1. 1个block最多有1024个threads。 (1个block只能够加载在1个SM中。)
2. 1个SM最多有2048个threads。(maxThreadsPerMultiProcessor)
3. 在最好的情况下,1个SM有2048个线程,按 warp(32个threads) 为执行单位,分为 2048/32 = 64组。
4. (我猜的,不知道是否对?)一个SM中有192个core,按照执行单位,192个core也会按照warp为执行单元分为 192/32 = 6组,对应执行64组warp。
5.我的grid为 313 * 313 * 1,block 为 16 * 16 * 1,排除register等其因素的限制,我的1个SM一次可以加载8个block,其中有的block执行完后才能够加载其它的block。


我有以下几点疑惑:

当我用Nsight调试得时候,执行到断点的时候,如下面代码,发现 CUDA Info (Warps) 中有 7个block,序号分别为[0,1,2,3,4,5,6],当序号分别为[0,1,2,3,4,5,6]的block全部执行完后,CUDA Info (Warps) 再加载序号为[7,8,9,10,11,12,13]的block,也是7个。

问1:是不是因为我的显卡有7个SM,SM0分配所得的block序号是[0,7,14,21,28,35,42,49],SM1 的为 [1,8,15,22,29,36,43,50],SM2 的为 [2,9,16,23,30,37,44,51],SM3 的为 [3,10,17,24,31,38,45,52] ,SM4 的为 [4,11,18,25,32,39,46,53],SM5 的为 [5,12,19,26,33,40,47,54],SM6 的为 [6,13,20,27,34,41,48,55],所以 CUDA Info (Warps) 显示每个SM首先执行的那个block?

问2:
到代码12行断点的时候(如下图),发现所有的线程都会停在这里,Status都为Breakpoint。我很疑惑问什么所有线程都能够同时到达断点,因为1个SM中只有192个core,不是应该首先执行前面的192个线程吗?怎么会同时执行1个block中的所有256个线程?

问3:
当往下面执行,到下面其它断点的时候,发现block 0的所有线程基本上都能够同时到达断点,但是其他block的线程基本上都变成绿色,而且一直停在 Source Line 为 12 的地方,直到 block0 执行完后才开始执行block1,以此类推。我彻底晕了,block 0在SM 0,block 1在 SM 1 的话,不是应该和block0和block1同时执行的吗?


本帖子中包含更多资源

您需要 登录 才可以下载或查看,没有帐号?立即注册

x
使用道具 举报 回复
发表于 2013-12-2 14:55:35
LZ您好:

您总结的5点基本限制基本上是正确的,其中第四点,根据手册说法,是scheduler每周期将一个warp的一条指令或者可以双发射的两条指令发射到SP和其他的执行单元上,其中并无说明是否将SP按照32个一组安排。同时,一个warp也并非是从开始到结束总在相同的32个sp上执行,请注意scheduler是针对指令发射的,而不是把warp绑定到某组SP上。

关于scheduler 的发射讨论,请您参阅一下以前的帖子。

关于您的3个问题,现在分别答复如下:

1:CUDA info中显示的是active blocks/active warps的内容,是当前正在执行的blocks/warps。以及在给SM分配block的时候,是按照顺序先给每个SM一个block,然后如果还有剩下的blocks并且同时SM还能resident更多的block的时候,再分配一轮,以此类推。

但如果不同的block执行结束的早晚不同的时候,后面补充block的顺序可能与此不同。

2:第一,您的断点在kernel一开始的地方,此时所有的active线程进度都相同,而且他们都满足断点条件。同时您可以灵活选定您需要观察的线程,通过设置不同的focus来实现这一点。

第二,192个core并不代表需要先执行前 192个线程,事实上每次发射到CUDA core上执行的只是某几个warp的当前一两条指令而已,并不是某个warp的线程绑定在一组CUDA core上执行结束才退出。而且,一直都有SM中的SP数量要远远小于一个block的最大允许线程数,如果是一组线程绑定一组CUDA core,执行结束才释放,那么block内的__syncthreads()将无法实现。

3:如果您需要观察后面线程在后一个断点的执行状态,请您将focus设定到需要的线程位置,并F5执行到所需断点处观察即可。

nsight这只是为您提供了一种调试机制而已,这并不代表实际执行中block1要等到block0执行结束之后才开始执行。CUDA中block之间的顺序是不保证的,不管是在同一个SM上还是在不同的SM上。

大致如此,祝您好运~
使用道具 举报 回复 支持 反对
发表于 2013-12-3 18:06:07
我的感受是,nsight调试时,为了方便用户查看变量,把有些本来并行的工作串行执行了。
使用道具 举报 回复 支持 反对
发新帖
您需要登录后才可以回帖 登录 | 立即注册