DAY35：阅读流程控制语句

我们正带领大家开始阅读英文的《CUDA C Programming Guide》,今天是第35天，我们正在讲解性能，希望在接下来的65天里，您可以学习到原汁原味的CUDA，同时能养成英文阅读的习惯。

本文共计334字，阅读时间15分钟

5.4.2. Control Flow Instructions

Any flow control instruction (if, switch, do, for, while) can significantly impact the effective instruction throughput by causing threads of the same warp to diverge (i.e., to follow different execution paths). If this happens, the different executions paths have to be serialized, increasing the total number of instructions executed for this warp.

To obtain best performance in cases where the control flow depends on the thread ID, the controlling condition should be written so as to minimize the number of divergent warps. This is possible because the distribution of the warps across the block is deterministic as mentioned in SIMT Architecture. A trivial example is when the controlling condition only depends on (threadIdx / warpSize) where warpSize is the warp size. In this case, no warp diverges since the controlling condition is perfectly aligned with the warps.

Sometimes, the compiler may unroll loops or it may optimize out short if or switch blocks by using branch predication instead, as detailed below. In these cases, no warp can ever diverge. The programmer can also control loop unrolling using the #pragma unroll directive (see #pragma unroll).

When using branch predication none of the instructions whose execution depends on the controlling condition gets skipped. Instead, each of them is associated with a per-thread condition code or predicate that is set to true or false based on the controlling condition and although each of these instructions gets scheduled for execution, only the instructions with a true predicate are actually executed. Instructions with a false predicate do not write results, and also do not evaluate addresses or read operands.

5.4.3. Synchronization Instruction

Throughput for __syncthreads() is 128 operations per clock cycle for devices of compute capability 3.x, 32 operations per clock cycle for devices of compute capability 6.0 and 7.0 and 64 operations per clock cycle for devices of compute capability 5.x, 6.1 and 6.2.

Note that __syncthreads() can impact performance by forcing the multiprocessor to idle as detailed in Device Memory Accesses.

本文备注/经验分享：

这章节主要说了流程控制语句, 和这些语句在GPU上的编译和执行效果.以及, 如何人为的更好的使用这些语句, 以便取得GPU这种SIMT架构上的较好执行效果.以及, control-flow一般翻译成流程控制。 本章节主要谈论到了GPU上流程控制语句的3大影响, 我将分别说一下, 注意, 手册这里写的不是很明确, 但我明确的直接说一下. (手册不说的明确有各种原因), 在GPU上, 源代码中使用任何流程控制语句(if, while, for)这种, 如果使用的不当, 可能会影响性能. 这主要体现在3个方面:我按照重要性排序一下 (1)流程控制语句导致最终运行的时候, 产生的分支和跳转发生在warp内部, 这种情况是最严重的。因为在以前的阅读章节中, 你知道SIMT结构只是构造成了线程可以自由执行的假象, 而实际上它们是按照warp一组了执行的,任何在warp内部的分支都将严重的影响性能.(Volta可能稍微好一点,手册这章节也没有说Volta) 例如根据之前的章节的说法, 你知道如果在一个32线程的warp中, 正好分支在16线程的边界, 将导致只有50%的执行效率.甚至如果这种分支, 所执行的路径还不同, 例如一个非常长的分支, 被只有warp中的非常少的数量的线程执行(例如只有1个线程在执行一个超级长的分支), 效率甚至会下降到大约1/32, 也就是只有3%左右了.所以这个是非常可怕的.也是本章节说的, 能significantly降低性能的原因. 而本章节也说了, 要尽量避免这种, 例如可以考虑分支在warp的边界上，例如这里说的用threadIdx / warpSize(32)，请注意这里的这种写法是不能实际编译执行的, 因为这里的threadIdx只是一个综合性的说法,在实际中, 你能用的只有threadIdx.x/y/z这些, 而不是没有分量的总名字.例如, 一个block形状是(8,8,8), 也就是512个线程, 我们知道这是16个warps(512/32 = 16). 但是哪些线程是这16个warp中的同1个warp的?这里就需要涉及硬件是如何组成warp的问题, 目前已知的方式按照先x方向, 再y方向, 再z方向组合的.例如threadIdx为(0,0,0)和(1,0,0)的线程在1个warp里，而threadIdx为(7,7,0)和(7,7,1)的这两个线程, 不在同一个warp中.所以在实际的非1D的形状的block要按照本章节手册的说法进行warp边界处分支的话, 必须考虑到warp是如何组合的.否则你会弄巧成拙。

这是第一点, 尽量规避warp内部分支, 而要分支在warp边界.这一点是非常重要的, 否则GPU会严重降低性能的.这也带来了实际应用中, 一些kernel可能在快结束的时候, 后半部分的代码, block中只需要保留少量线程了, 你应当该怎么保留的问题.例如还是刚才的(8,8,8)的block, 最后扫尾, block只需要12个线程了,你选哪12个的问题.还是应当尽量选择按照x方向优先的.如果你很不幸的选择了, 例如(0,0,0), (0,0,1), (0,1,1), (0,0,2), (0,0,3)...这样的12个出来，那么会执行一堆里面只有1个活跃线程的warps, 性能降低很惨。 (2)第二点是流程控制语句本身需要执行上的代码，例如说, 那么一个for(A; 判断条件; B), 即使你刚才看了第一条, 所有的判断条件都成立或者不成立在warp的边界，甚至更夸张的, 连warp边界都没有发生分支跳转(例如对block中的所有线程都成立)，此时依然有执行上的代价的，为何?因为循环的控制语句本身(语句A, 和每次循环都要执行的判断语句, 和B), 它们会编译出来具体的硬件能执行的指令,而硬件执行这些指令本身, 需要时间/周期上的代价. 例如一个for (int i = 0; i < 10; i++)，虽然所有的warp中的线程都同样的执行了10次该循环体,并未发生(1)说的warp内部的分支(显然的, 因为大家都是一样的执行了i < 10一共10次这种循环, 没有人执行更多或者更少, 很一致的)，但是你要考虑每次循环的时候, 都有对循环控制变量的++的运算(增加1), 和判断它是不是小于10这种运算的,这些都需要浪费周期和时间.所以这就是第二点， 哪怕没有warp内部的分支,这些流程控制语句的本身的执行需要代价. 而手册也给出了建议(写手册的人在这里没有分清问题, 和下面的第三点混杂了, 可以先无视它), 例如对于循环可以考虑展开, 例如刚才的循环: for (int i = 0; i < 10; i++) {do_something();}可以被展开为: do_something(); do_something(); do_something(); do_something(); do_something(); ... //一共10句

这样每句do_something之间, 通过unroll, 取消了循环控制本身的代码(循环控制变量的+1, 和判断<10, 以及, 后续的跳转语句等等)。这种unroll展开, 虽然我们将它归结到第二点, 看起来不是很重要，但是对于非常小的循环体之类的(例如do_something很小, 假设等价于10条编译出来的指令)， 此时循环控制本身(例如i++和i<10和后续跳转之类)编译出来了8条指令,你依然会有可能的效率损失的(例如80%)，当然, 具体的编译结果不是这里的10和8, 这是一个夸张了点的演示, 但是对于一些特定的代码, 控制语句本身的代价不能无视.这是第(2)点需要考虑的. 第(3)点则是执行上的代价, 要知道往往在我们这种SIMT的架构上, 想完成真正自由执行的假象, 需要使用类似一种叫Fork-Join的warp内部执行流程, 例如感兴趣的用户可以参考在分支时候的, 设定汇聚点(SSY), 使用Fork-Join Stack, 进行汇聚(SYNC)重新warp里面的线程在一起, 甚至还可能跳转带来的I-Cache的miss, 等等, 所以实际上N卡引入了predication,注意手册这里又写的不好了, 写成了branch predication, 但这个和单独一个词的predication(predicated execution)是两回事.前者(两个词的branch predication)往往是指的现代的硬件的分支预测, 而后者(predication或者predicated execution)往往是指的带有掩码或者条件的指令执行,例如ARM中的条件执行就类似后者。手册这里的原始写法容易造成误解. 正因为刚才说的有这些诸多分支上, 能伪造成自由warp内部的执行的假象, 而造成的SIMT结构上的代价,所以N卡引入了predication, 或者叫predicated execution,这种执行带有一个指令前缀掩码,对warp内部的32个线程, 每个线程可以选择真执行或者假执行, 例如: if (lady_is_beautiful) { money += 8; //女生漂亮就有钱 } else { money += 0.1; } 像这种小分支的时候, 编译器目前会生成predicatied的执行语句, 即: p = is lady beautiful? @p money += 8 @!p money += 0.1 这样的三条指令.而中间并没有刚才说的(3)点中的任何其他指令.warp内的32个线程都将全部执行money += 8和money += 0.1 无论它们的条件是什么, 但是那些不美丽的线程, 将屏蔽掉@p的结果 而那些美丽的线程, 将屏蔽掉@!p的结果.通过这种条件执行, 假执行的那些线程, 规避掉了昂贵的传统分支. 所以也能节省性能.但是需要说明的是, 什么样的分支会生成这种简单的前缀执行,什么样的分支会生成正常的常规分支指令,要看编译器的心情, 这并不是一个用户能控制的方面.这是第(3)点 需要说的是, 刚才的@p这种前缀, CUDA的GPU中, 目前一种有7个可用寄存器, 也就是p0, p1, ...p6 每个这种寄存器叫predicate寄存器(P-Reg), 只有1-bit, 可以存放一个真值或者假值. (如果你还记得之前的章节, 你会记得编译的时候, CUDA告诉你可以通过-maxrregcount来控制寄存器数量)

(但是为何是max r reg? 而不是max reg?) 这是因为你控制的是常规的寄存器(Regular Register)，这种p寄存器是固定分配的, 每个线程有7个. 不能选择更多, 也不能选择更少. 常规的可以改.精确的说, 每个线程有8个这种p寄存器, 你能用的有7个.另外一个叫PT(T代表true)，所有的不需要这种条件执行前缀的指令, 都实际上总是使用PT寄存器,也就是@pt inst;而pt总是真. 所以等效只有inst。

cuobjdump也不会反汇编的时候显示这个@pt恒真前缀——这个PT寄存器其实还有一个用途, 因为它恒真, 你可以理解成是只读的(一些需要写入P寄存器, 而又不想破坏掉目前的p0-p7的一个的时候, 可以写入pt, pt因为只读, 变成了一个写入黑洞,往里面写入的任何值都将被丢弃, 下次pt读取的时候还是真)。 其实这个也是之前手册为何说, 一些老卡(3.0的Kepler)可以使用63个寄存器(64个)，新卡(第二代Kepler/Maxwell/Pascal+), 可以使用最多255个寄存器/线程(256个)，刚才你知道, 这些都是常规寄存器, 用户可以选的,但是为啥括号里面有1个的差别?因为常规寄存器里面也有一个寄存器, 类似PT，叫做RZ，一般叫做0寄存器, 该寄存器的值恒定为0, 可以快速提供常数0.类似的, 它也是只读的, 一个指令的输出结果如果不要的话, 可以直接写入RZ, 等于也写入了黑洞. 下次读取还是0. 以上3点, 就是常见的, if, for, while等等流程控制语句在实际编译出来指令后, 在执行的时候所可能具有的代价, 以及, 如何能尽量规避它们. 一个能正确应用了这三点的(或者其他点, 以后说)的CUDA C的Kernel,还是可以能在代码中大量的出现if, for, while的, 而不会影响性能.所以很多人认为我不能在GPU上使用if, for, while, 是错误的. 该用就用.用的好了没啥问题. 注意手册本章节将原始CUDA C代码的语句(for这种), 和编译出来的指令部分, 都直接叫指令, 而说它们的执行会如何如何,这是不对的. 必须将两者分清, 否则就是对CUDA的编译器团队的不尊重. 请注意虽然__syncthreads*()家族非常常用, 但是这里需要明确的说一下: (1) 它们用来进行block内部的同步 (2) 除了(1)点, 还可以用来进行少量的信息交换. 例如带有_and, _or, _count后缀的高级版本的__syncthreads()，可以快速统计一个block内部, 满足某种条件的线程有多少等等。

而本章节说的它们的吞吐率, 实际上并不重要, 你真的需要频繁的快速执行这么多__syncthreads()么? 不需要的.哪怕是你用来交换数据, 过多的数据应当考虑走shared memory.所以你看到了maxwell+将这个吞吐率降低到了32而不是128条(从线程的角度)每cycle了.不过手册本章节说的, 同步会带来的代价依然要注意. 过多的同步可能会影响性能. 比如所有其他线程都在等某个warp里的线程完成某个长延迟的访存操作的时候, 而不得不暂停执行 , 甚至整个SM都得暂停下来 等等.

有不明白的地方，请在本文后留言

或者在我们的技术论坛bbs.gpuworld.cn上发帖