哪些情况会造成GPU比CPU还慢。

加速计算 CUDA
21

LZ您好：

时值周末，未能在第一时间回帖，请见谅。

下面将逐条与您讨论：

1：线程不等于CUDA CORE，以及前者是软件概念，后者是硬件概念，这一点您的理解是正确的。

关于两者的关系，这里还不是很到位。实际上每个线程都要跑kernel中的代码，您可以将其看做一个指令流，指令流中的具体指令将被发射到GPU的SM内部的不同单元去处理，大部分计算指令被发射到SP/CUDA CORE上执行，一部分计算指令被发射到SFU上执行，访存指令被发射到LSU上等。

SM上一般会有大量的resident threads，比SP数量多很多，他们按warp划分好。SM中的scheduler 以warp为单位调度这些threads。如果某个warp的指令流中的当前指令或者当前连续两条指令是就绪可以执行的状态（比如操作数什么已经就绪），那么scheduler会在某个时刻选中这个warp，发射到SM内的CUDA CORE和SFU，LSU上执行。如果有连续两条指令是就绪的，可能会双发射，将这两条指令发射到SM内不同的执行单元上执行，请注意，这里说的是一个warp的行为，是warp内的每个线程都有同样指令执行。——从这里可以看出，一个warp的线程，或者一个线程是没有一直固定在某个SP上执行的，同时SP功能单一，他基本上只有计算能力，而不像CPU CORE那样是一个大而全的CORE。一个线程中的所有指令，并不都在SP上完成，还需要SM其他部分硬件一同完成（再进一步，需要隶属于整个GPU的硬件协同完成，如L2 cache，如显存控制器等）。

综上，容易理解，为什么一个线程不是始终在一个SP上执行；以及为什么SP不具备完整的功能，而线程所需的功能实际上是由SM各部分分别提供的；以及为什么我前面说您可以粗略地将一个SM看作是一个具备超线程能力的CPU核。

此外，这些计算单元都是流水线化的，很多都是每周期可以接纳一次新的计算，但是这一次的计算等结果出来，可能需要多个周期。考虑到当前一个就绪的warp被scheduler发布到SP上计算，这个计算吞吐量假定是每周期可以执行一次，那么执行这个warp这次计算的这些SP们在下周期又可以进行计算，而此时当前的这个warp的计算结果还没有返回，我们的GPU并不会选择让这个warp等在这里，而是将这个warp切换到等待的队列中，从就绪的队列中再找一个warp进行下一周期的计算，等多个周期之后，这个warp的结果返回了，可以往下计算了，那么它会被转移到就绪的队列中，等待scheduler调度。这个过程中，自动地通过线程切换的方式，实现了计算和指令延迟的互相掩盖，这也是GPU执行的一个基本方式。NV表示，这种线程切换是硬件完成的，是没有额外开销的，因而一般鼓励上大量的线程，通过这种自动的机制，自动用计算掩盖指令延迟，尽量使计算单元时刻处于忙碌状态。

这就是为什么一般建议上大量的线程，以及上极少量的线程效率不可能高的原因。

以上是对您第一个点的讨论。

22

2：前两句基本正确，“SM由sp组成”可进一步改为“SM由sp和其他硬件组成”。（批判神马的就免了，现在不流行这个）

“warp中的32个线程使用SM中的32个sp来执行”——由前面回答1的时候所述，可以是SP也可以是其他单元，以及线程不是自始至终在SP上的，是被切换来切换去的。

当一个block的线程数不足32的时候，（或者不是32的倍数，考虑最后一个warp的时候），因为执行的最小单位是32，所以实际上还会执行一个warp的，会通过某种内部机制，自动补齐到32线程，但是后面补的那些是无效线程，跟着空跑，不保留结果，不会产生意料之外的访存越界，以及从逻辑正确性上讲，一般无需担心这种情况；但是从执行效率讲，会有一些损失。

以上是对您第二点的讨论。

23

3：这段内容在回答1：中做了较为详细的叙述，现在仅作简要补充:

1）一个SM上的resident threads一般很多会比SP多很多，以及切换过程前面有回答。

2）执行单元是SM内的SP，SFU，LSU等，不仅仅是SP。

3）这里说的“灌入”指的是scheduler将一个就绪的warp的指令发射给合适的执行单元（比如32个SP），假如这里执行的是吞吐量为每周期每SP一条指令的计算指令（如浮点乘法），那么到下个周期的时候，这组SP又可以接纳新的计算指令了，虽然此时前一周期的计算结果还没有输出。这是流水线化执行单元的一个特性。
以及，这里指的“灌入”依然指的是“指令”而不是“线程（的整个指令流）”，“灌入”是一种形象的说法。

4）SM上所有resident threads以warp为最小单位，划分为就绪的一组和不就绪的一组。
若某就绪warp当前指令被发射到执行单元，或当前连续两条指令被双发射到执行单元之后，视后面的指令是否依赖于前面的结果，会被再次划分到上述两组中。

如果该warp处于非就绪状态，那么其依然是resident 在该SM上，但当前不会被scheduler选择并发射到执行单元上并利用新的一个周期的执行单元的计算机会（执行单元是流水线的）。“切换出来”就指这种情况。以及，当结果成功返回以后，该组warp会被分配到就绪的那一组，等待scheduler选择。

如果该warp仍处于就绪状态，那么原则上可能直接被scheduler继续发射下面的指令，也可能放入就绪warp的那一组/队列，等待scheduler下一次选择发射指令。具体的调度方式是不公开的。

以上是对您第三点的补充讨论。

24

4：我将分几点讨论本问题：

1）数据量巨大，这个是正常的，远远超过CUDA CORE的数量这个也是正常的，但这两个事情没有特别直接的关系，CUDA CORE是可以反复使用的:)，而数据处理过去就过去了。换句话说，你要是用一个CPU core来处理，不也得处理这么多数据么？反复循环处理就可以了。

2）但是“因此每个BLOCK可以安排512…”这里的因果关系不成立。数据量巨大或者一般大，GPU里面的SP非常多或者一般多，都不影响block的线程规模安排的。

线程规模安排其实一般推荐选择256,192,512这样的适中大小的典型值，具体选择多少还和算法逻辑有关，考虑优化的话，还和GPU计算能力版本有些关系。
如果您需要使用线程编号辅助寻址，那么您维护好线程编号和地址位置的关系即可，并不一定要一个线程处理一个数据，以及这并非是一个限制block线程规模的本质因素。

3）如果某个特定的任务只需要很少的线程，比如说1个warp，那么此时，虽然为了这一个warp启动一个kernel挺亏的，但是如果这样能避免一次“数据传回host——CPU处理——结果传回device”，那么还是值得的。毕竟是在device端本地折腾么。

如果是一些密集计算的大kernel之间夹杂一些这样的小的kernel做一些扫尾工作等，可以无需纠结，直接这样使用即可。

如果您的主要计算算法都是这种只需要少量线程计算的，但是如果需要计算非常多组，而且各组之间互不相关，那么还是可以良好GPU加速的。比如您可以使用一个warp来计算一组，或者使用一个较小的block来计算这样一组，然后上大量的线程来并行化计算很多组数据，那么是OK的。

如果您计算都是这种使用很少线程的，并且也只计算很少组或者只计算一组，以及也不是为了避免host-device通信，那么这样小的计算直接在CPU端搞定就行。

您的第四点讨论于此。

25

5：
1）关于“线程切换”和“流水线化”的描述，请参看今天我详细叙述的内容，本处不再赘述。

2）以及不是按照一个block里面的线程数量少于一个SM中的SP（和其他执行单元）这样计算的。
一个block只能resident在一个SM上，但是一个SM上可以resident多个blocks。对于SM3.X的硬件，一个SMX上可以resident最多16个blocks（但同时受到最多不能超过2048线程的限制）。

3）以及一般需要resident threads远远多于SP数目。

4）以及，切换都是以warp为单位的，而不是整个block切换的。多个resident 在同一个SM上的block中的warp都可以参与切换。

最后回到您的直接问题上来，您的问题实际上是单个结果的计算，计算量有限，但是要计算很多组。
按照您当前的想法，每个block128线程，计算一组的结果，然后通过上多个block来实现并行，这个是可行的。

进一步的优化建议是，您可以一个block 128线程多计算几组数据，比如说5组，一般来说会有所改进。即，适当增加一个线程的工作量，有利于提高总效率。

以及，您的算法可能是某种折半规约的算法，如果您能修改该算法实现，依然是一个block计算多组数据，但是无需在计算一组数据的时候，每次工作线程数量折半，可能会更好一些。

以上是您第五点的讨论。

26

6：这里稍微需要说明一下的是，您的比较实际上是您的CUDA实现在您的GPU上的结果VS您的C语言实现在您的CPU上的结果。

而不是单纯的方法VS方法，或者硬件VS硬件，或者某实现VS某语言。

讨论速度比较或者加速比，都应该把比较类型，比较意图等等说清楚，而不是笼统地一句“比C语言快XX倍”这个说法无意义的，C语言也有实现好坏，也有CPU好坏等，一个具体实现VS一种语言，严格说这么比没什么意义。

这也是我之前从不评价您的加速比的原因。
如果您是修改了算法的某个环节获得了同设备上的加速，那么可以讨论算法加速比，可以讨论减少了多少次计算等。
如果您是将部分环节变成了并行实现，以此获得了加速，那么您可以讨论并行加速比，并且这种加速比的极限取决于不可并行的部分。
如果您的某算法充分挖掘了CPU的效能，并且一个同功能同算法的并行版本充分发掘了GPU的效能，那么您可以讨论一下设备VS设备的加速比，并可分析与理论峰值的关系，是计算密集还是访存密集，瓶颈在哪里，适合哪个结构使用。

如果只是您的初步实现，那么只能说是您的实现@您的CPU VS 您的实现@您的GPU，这里面变数非常多。

其次，debug模式下编译出来的代码含有大量的辅助调试的代码，以及不做优化，实际执行速度和release模式相比要慢得多，所以请无视debug模式下的速度。

最后，前面已经详细讨论了线程规模的安排，不在赘述每block16线程的问题。

这是对您的第六点的讨论。

您在20#给出的6点已经全部讨论完毕，供您参考。

祝您周末愉快~

27

额外补充一楼：

前面21#提到“此外，这些计算单元都是流水线化的，很多都是每周期可以接纳一次新的计算”

这个参数在宏观上体现为各个指令的吞吐率，CUDA C Programming Guide中有表格说明（以及不同版本的手册表格略有出入，官方无解释，请以最新版本手册为准，如有错误NV负责）。

手册中的表格是每SM每周期的指令吞吐量，这个吞吐量不仅受真正的执行单元的流水线设计影响，也受具备执行某指令能力的执行单元的数量的影响，是一个综合的结果，也是一个可以不管细节直接参考的结果。

此外前面讨论各楼中，CUDA CORE==SP，SM==SMX，等等一些同义或者十分近义的术语一般予以混用，并不特别区分，请注意理解。（严格说，只有kepler的SM才叫SMX，但习惯上混用或者仅将SMX称为SM而不反过来。）

28

您的讲解中概念繁多，原理复杂，我要多读几遍才能进行新一轮的提问。
打这么多字，辛苦了。

29

LZ您好：

俗话说“不明觉厉”，但其实“明则不厉”。

欢迎您经常来论坛交流和讨论各项CUDA技术，我和横扫斑竹，玫瑰斑竹将竭诚为您服务。

祝您在论坛学的愉快~

30

1、（21#）“如果某个warp的指令流中的当前指令或者当前连续两条指令是就绪可以执行的状态（比如操作数什么已经就绪），那么scheduler会在某个时刻选中这个warp，发射到SM内的CUDA CORE和SFU，LSU上执行。”一个线程当成一个指令流，则一个warp有32个指令流。您这里是指warp中的“所有线程”还是“只要有一个线程”的当前指令或当前连续指令就绪就可以被SM调走？
2、（21#）“一个线程不是始终在一个SP上执行的”。我现在明白了一个线程要根据指令来决定是在CUDA CORE、SFU还是LSU执行。因此一个线程不是始终在一个SP上执行。当一个线程随同warp被调走之后，在等待结果的过程中被切换出来，这个等待的过程中该线程会不会再次被调走继续执行下一条就绪的指令？
根据您（23#）第4小点的回答，如果指令2与正在执行的指令1无依赖性，那该线程会被分配到就绪的一组，也就是说将被再次调走，执行指令2。这么一来的话，我为什么要把这个线程切换出来呢，我把指令1灌入执行单元后，直接灌指令2就可以了。
3、（21#）“这些计算单元都是流水线化的，很多都是每周期可以接纳一次新的计算。”那就是说一个计算单元里可以同时计算多条指令。我可不可以这样理解，把计算单元比作一条环形铁轨，线程中的指令比作火车，吞吐量比作发车间隔，当一辆火车上轨后，达到一个发车间隔时，另一辆火车接着上轨，而不是在那等第一辆火车回来再上轨。
4、（21#）“SM上一般会有大量的resident threads”，为什么您总是不把resident翻译出来？resident threads跟我们在核函数中分配的线程是不是一回事，有什么关系？例如我的分配格式为<<<1024,512>>>,我这512个线程是在CUDA CORE、SFU或LSU上跑的，跟resident threads有什么关系？是不是说我在核函数里分配的线程对应就是SM上的resident threads？
5、（24#）“数据量巨大或者一般大，GPU里面的SP非常多或者一般多，都不影响block的线程规模安排的”。（25#）“您可以一个block 128线程多计算几组数据，比如说5组，一般来说会有所改进。”我想问一下，假如我一组数据只有128个元素，一个线程对应一个元素，那只需128线程。如果我想安排多一些线程，那就要多组放在一个BLOCK里，假设为两组，我分配256线程。那这两组如何并行呢？我只能想到用一个if语句，if(tid<128{执行第一组｝，if(tid>=128){执行第二组｝。这就成了串行了。您能否解释一下怎么实现？

31

LZ您好：

1：一个warp总是同时被调度的，以及scheduler是以warp为单位调度线程的。关于指令流，实际上因为从大的角度讲，一个grid中所有的线程都跑的是相同的kernel，所以他们需要执行的指令流都是一样的（暂不考虑分支，后面说）。
因为一个warp总是同时被调度，同时被执行，所以这个warp里面的所有事情都是同时的，或者说保证进度一致的。他们总是同时就绪的，同时被scheduler发射，并同时在执行单元上执行。

具体说，如果warp内没有分支，那么32个线程的行为总是一致的，上面一句话很好理解。

如果warp内有分支，那么不同的线程需要执行不同的内容，此时，部分线程会被插入等待指令以等待另外的线程运行结束；或者通过假执行的方式，把不同的分支都执行一遍，每次执行只有对应有效的线程保存结果，其他线程丢弃结果。
总之，有分支的一个warp最终还是会保证所有线程都执行妥当，虽然中途比较折腾，并且站在整个分支的角度上讲，warp仍然是同步的。

32

2：一条指令需要根据其类型在不同的执行单元上执行。

“当一个线程随同warp被调走之后，在等待结果的过程中被切换出来，这个等待的过程中该线程会不会再次被调走继续执行下一条就绪的指令？”首先，总是以warp为单位调度的。其次，某个warp当前的一条或者两条指令被scheduler发射到执行单元之后，下面的指令可能依赖于当前的结果，也可能不依赖于当前指令的结果。

如果依赖，那么该warp就处于非就绪状态，那么必须被切换到等待队列中，等待当前执行的指令的结果返回。
如果不依赖，那么该warp仍然处于就绪状态，此时理论上可以被scheduler继续发射剩下的指令，也可以被scheduler安排到就绪warp的队列中等待下次再发射，并同时发射其他就绪的warp。
具体的调度逻辑涉及公平性等原则，并且该调度逻辑是不公开的，所以并无确切答案告诉您。

以及，宏观上讲，无论是对该warp继续发射，还是换个warp发射，并无区别。
同时这里也 回答了您2：中后半段的问题。

33

3：是的，您的理解是正确的，以及以工厂流水线的类比可能更为直观一些。（他们都叫“流水线”么 :slight_smile: ）

另外需要说明的是“一个计算单元里可以同时计算多条指令”需要理解为，这些指令都出在不同的执行程度/阶段。

34

4：之所以总是写成resident threads是因为这个是英文官方手册上的写法，而且我不知道如何写成中文能够不引起误导，因为似乎并没有确切的中文译法，于是就保留了手册的说法，同时写成这样也方便您对照手册查找。

resident threads和grid中分配的threads是不同的，后者是总的线程数量/形状安排等，而前者指的是，同时能够加载在（通常指一个）SM上的线程。我们知道，如果grid中的线程数量非常多的话，那么他们是以block为单位，分批加载到SM上进行计算的，也就是一批计算完毕了，换下一批。

一个SM上能加载的最大线程数量，在fermi上是1536个线程，在kepler上是2048个线程。
同时一个SM上所能加载的线程数量还受到最大加载block数量的限制。比如fermi上最大可以加载8个blocks。
那么对于fermi而言当一个block的线程少于1536/8=192threads的时候，您SM上的resident threads数量将无法达到上限。

具体的各项限值，请您参考Programming Guide的附录。

35

5：前面4：已经说明选择block内线程数量的一个重要参考因素，那就是尽量达到resident threads数量的上限。（以及这样可以通过自动的线程切换达到计算掩盖延迟的效果）

所以说“数据量巨大或者一般大，GPU里面的SP非常多或者一般多，都不影响block的线程规模安排的”。
如果您的计算量很小，只能使用较小的线程规模，那么一般是不能取得良好的效果的，您有多组这样的小计算量的任务加以“任务级别”的并行。

如果您“一组数据只有128个元素，一个线程对应一个元素，那只需128线程”，但是同时有10000组数据需要处理，那么您安排一个block128线程，开10000个block，可以达到初步的效果。
此时，如果您每个block计算5组（只需要安排好访存即可，很容易），开2000个block，可以取得更好一些的效果。
按照您说的方法，一个block开256个线程，用if划分一下，这个也可以的（虽然一般习惯上不这样），这样您一个block将计算2组数据，以及他们是并行的，而非您想的是串行的，以及这个分支是warp对齐的，不会有明显的效能影响。

36

以及，对5：补充一个简单的例子

假定您一个block有128个线程，需要处理一个数组，int yijun[5][128]，那么：

global void yijun_kerenl(int yijun[128],…)
{
int tid=threadIdx.x; //0~127
for(int n=0;n<5;++n)
{
int input=yijun[n][tid];
… //你的算法
… //保存结果

}
return;
}

这样就就可以只用一个block处理5组数据，当然以上写法仅为示意，实现方法不唯一。

祝您编码顺利~

祝您晚安~

37

继续补充一下5：

“我只能想到用一个if语句，if(tid<128{执行第一组｝，if(tid>=128){执行第二组｝。这就成了串行了。”
您可能认为，当执行第一个分支的时候，后面一半的128个线程是等待着不干活的，而执行到后面分支的时候，前面的线程是不干活的。

其实并非是这样，warp之间的执行顺序是没有保证的（block内的warps可以使用__syncthreads()强制同步）。在一开始的时候，前面128个线程就会走第一个分支，而后面128个线程就会走第二个分支，并不存在谁等谁的问题。
（这个和warp内的分支不一样，warp内的分支是有开销的。而按warp对齐的分支，开销可以忽略。）

当然您也可以使用一些书写技巧，用一个通式表达这256个线程的不同访存位置，从而不使用if判断，这样就没有分支了，但是会有一些其他的计算开销。

总之这里线程和计算任务的安排，线程和访存的安排是很灵活的，只要解决问题，怎么写都可以。

祝您编码顺利~

38

噢，我明白了，这就是为什么每个SM上会有Maximum numbers of threads per multiprocessor,这个数量就是SM上的resident threads, 而且我在GRID安排的线程要比这个数量多才能充分发挥SM的效率，是这样吗？

39

是的，我一直认为if语句是串行的，因为我以前写过这样的程序，发现时间反而变长了。您这里用FOR语句难道不是串行的吗？
既然是可行的，我现在就着手对程序进行优化。

40

LZ您好：

1：按照programming guide 附录上的原文说法是“Maximum number of resident threads per multiprocessor”，这个是SM上的resident threads的最大值（硬件限值），以及，这只是一个制约因素，还有您block规模的大小，一个block内threads数量太少则会导致“Maximum number of resident blocks per multiprocessor”这个限值先达到，而无法继续resident更多的线程。
以及，这个和您的每个block的资源使用量也有关，如果一个block占据了SM上的所有shared memory，那么即便这个block的thread数量很少，也不会有第二个block继续resident在该SM上。

2：grid中的线程数量要足够多，这是一个必要条件，还需要block规模合适，资源使用情况合适等，如前所述。

此外，一般而言需要多上threads，某些特别构造的例子中，也可以使用较少的threads实现较好的效率。但常规编码，请优先考虑多上threads。

另外，多上threads并不是万能的，实际程序运行效果还受到代码撰写情况的影响。

大致如此，供您参考。