Nsight 调试中的不解

加速计算 CUDA
1

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同时执行的吗？
[attach]3472[/attach]

2

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上。

大致如此，祝您好运~

3

我的感受是，nsight调试时，为了方便用户查看变量，把有些本来并行的工作串行执行了。