目录
背景
通过减少循环的开销（分支指令和循环计数器维护）以及增大单次处理的数据量，来榨干 GPU 的最后一滴性能。
分支
在CUDA/GPU语境下，"分支(branch)"指的是程序执行路径出现选择的地方，本质上对应到SASS（GPU的汇编指令）层面，就是各种条件跳转指令。常见会产生分支的C++/CUDA源码结构包括：
if (condition) { ... } // 最常见 if (condition) { ... } else { ... } for (...) { ... } // 循环本身也是分支（每次判断是否继续循环） while (...) { ... } switch (...) { ... } condition ? a : b; // 三元运算符，编译后也可能是分支（也可能被优化成无分支的select指令）理解divergence的关键，是理解GPU的**SIMT（Single Instruction, Multiple Threads）**执行模型，这和CPU的多线程完全不同。
1. Warp是最小调度单位，32个线程"绑在一起"执行同一条指令
一个warp里的32个线程，物理上共享同一个指令计数器(Program Counter)，意味着这32个线程永远在执行同一条指令，不存在"warp内有的线程在执行第5条指令，有的在执行第10条指令"这种情况。
2. 遇到分支时，硬件用"谓词掩码(predicate mask)"来决定哪些线程真正生效
当warp执行到一条
if (condition)对应的分支指令时：
硬件先让warp内所有32个线程都计算一遍
condition这个判断表达式，得到每个线程各自的true/false结果，存进一个32位的活跃掩码(active mask)，比如11110000111100001111000011110000这种位图，1表示该位置的线程这次该走if分支，0表示不该走如果这32位全部相同（要么全是1，要么全是0）——这就是uniform分支，硬记数说"线程束没有分化"。硬件只需要按掩码结果，让整个warp要么集体执行if内的代码，要么集体跳过，只需要一次指令发射，效率最高
如果这32位有1也有0——这就是divergent分支，硬件必须做下面这件事：
这是divergence最核心的执行机制，分两步（在较老架构上是串行的两阶段，Volta之后引入了Independent Thread Scheduling，机制更灵活，但本质开销逻辑类似）：
第一步：用当前的active mask（比如掩码里值为1的那些线程），只让这些线程真正执行if分支内的指令，而掩码为0的线程，这一轮指令周期里什么都不做（不是不存在，而是被"禁用"，物理上依然占着位置，只是这条指令对它们不产生任何效果，相当于NOP）
第二步：等第一组（掩码为1的）线程把if分支内的所有指令都执行完，硬件翻转掩码，让原本是0的那些线程（也就是该走else分支，或者该跳过if直接往后走的那些）开始执行，而原本执行过if的那些线程这次变成"禁用"状态
这两步是串行的，不是并行的！ 也就是说，一个divergent的warp，原本可能一条指令就能完成的工作，现在要花两倍（或者更多，如果分支更复杂）的执行时间，因为两条路径上的指令分别都要走一遍，只是每次只有一部分线程在真正干活，另一部分在"陪跑"等待。
分支有可能会造成线程束分化，从而指令吞吐浪费，执行时间变长，运算资源被浪费，没有跑满，会增加开销
展开循环
for (itn i=0;i<tid;i++) { // to do something }对于这种循环是最垃圾的，一个循环造成了线程束分化，当tid=0时没有进入循环，tid=1循环一次，tid=2循环2次……tid=31循环31次
循环次数不一致，会导致严重的线程束分化：跑得快的线程要停下来等跑得慢的线程，计算资源被浪费。
那么就需要去优化，如何优化？？？------->>>循环展开
循环展开的核心思想是：减少循环次数，从而减少分支指令的执行次数。
比如展开前
for (int i=0; i<32; i++) { a[i] = b[i] + c[i]; }
- 比较
i < 32- 递增
i++- 条件跳转回循环头部
展开后
for (int i=0; i<32; i+=4) { a[i+0] = b[i+0] + c[i+0]; a[i+1] = b[i+1] + c[i+1]; a[i+2] = b[i+2] + c[i+2]; a[i+3] = b[i+3] + c[i+3]; }直觉上很明显，循环控制相关的指令（比较、递增、跳转）总数减少到原来的1/4，这就是展开带来的直接收益。
在CPU上，这部分开销不是免费的——尤其是跳转指令对CPU的指令流水线不友好，每次跳转都可能引发分支预测失败的风险，一旦预测错误，CPU要清空流水线重新取指，代价相当可观。展开后跳转次数减少，这部分代价也跟着降低。
GPU上有一个CPU不那么强调的额外因素：GPU没有CPU那种复杂的分支预测器，每次循环跳转对GPU来说是相对"诚实"的开销（不像CPU能靠预测器提前猜测、隐藏跳转代价），所以减少跳转次数这件事在GPU上的收益占比，相对会比CPU上更明显一些。
SM内部的指令流水线依然可以让独立指令的load操作重叠发出：
- 原始版本：发出
b[i]的load → 等待数据回来 → 做加法 → store到a[i]→ 进入下一轮循环，重复- 展开版本：可以连续发出
b[i+0]、b[i+1]、b[i+2]、b[i+3]、c[i+0..3]这8条load请求，这些请求没有依赖关系，可以同时在"飞行"中，线程不需要等第一条数据回来才发第二条，等所有数据都陆续返回后再统一做4次加法访存合并
发出来的这几条指令，同时load，如果地址相近，可以合并成一条访存指令
对于现代编译器，这种简单的代码，结构一致性的代码，就算你不是手动展开，编译器也会自动帮你展开
展开归约
基于以上的前置知识，回到上一节，我们深度聊了一些归约的问题
那么对于之前的归约，是否可以展开呢？？？完全可以，如果我们把要算的算一轮呢？那这样我们的for循环是不是可以少一次，block的数量是不是可以少一半
对于之前，我们的数据是1<<24，每个block1024，那么就需要16384个block
我们的线程是一一对应数据量的，如果我们每一个线程对应两个数据量呢？那block是不是可以少一半，那调度器的压力是不是可以少一点
__global__ void reduceUnroll2(int * g_idata,int * g_odata,unsigned int n){ unsigned int tid = threadIdx.x; unsigned int idx = blockDim.x*blockIdx.x*2+threadIdx.x; if (tid >= n) return; int *idata = g_idata + blockIdx.x*blockDim.x*2; if(idx+blockDim.x<n){ g_idata[idx]+=g_idata[idx+blockDim.x]; } __syncthreads(); for (int stride = blockDim.x/2; stride>0 ; stride >>=1){ if (tid <stride){ idata[tid] += idata[tid + stride]; } __syncthreads(); } if (tid == 0) g_odata[blockIdx.x] = idata[0]; }blockDim.x*blockIdx.x*2这里的意思就是原先一个block算一个block的数据量，但是现在一个block算两个的量，也就是数据总量就是那么多，原本一个block该负责
blockDim.x个元素，现在每个block要负责blockDim.x*2个元素。，后面还是交错分配
reduceUnroll2比reduceInterleaved快了将近2倍，这个提升幅度非常显著
就是因为一行简单的提前归约，导致性能快了接近2倍，为什么？？？// reduceInterleaved：一个block处理blockDim.x个元素，然后开始归约 // grid.x = n / blockDim.x = 16384 // reduceUnroll2：一个block先把两份数据合并成一份，然后再用同样的归约逻辑 // grid.x = n / blockDim.x / 2 = 8192 unsigned int idx = blockDim.x * blockIdx.x * 2 + threadIdx.x; if(idx + blockDim.x < n) g_idata[idx] += g_idata[idx + blockDim.x]; // 这一行是关键 __syncthreads(); // 之后的归约循环和reduceInterleaved完全一样回看
reduceInterleaved每一轮归约之后都有一个
__syncthreads()，这个同步点把所有访存的并发机会强制截断了。等所有线程都完成当前轮次的写入，才能进入下一轮的读取。
每一轮都是一个完整的"读-算-写-等待"的串行周期，上一轮和下一轮之间完全没有重叠。每个block里面都有一个__syncthreads()，这个才是性能大杀器
grid.x = 16384个block 每个block处理1024个元素 block内部归约轮数 = log2(1024) = 10轮 每轮：load+load+加法+store+__syncthreads() 总循环轮次 = 10轮 每轮有效处理的数据范围 = 1024个元素（但每轮活跃线程数在递减）grid.x = 8192个block（减半！） 每个block处理2048个元素 预归约1轮：load+load+加法+store+__syncthreads() block内部归约轮数 = log2(1024) = 10轮（和之前一样！！） 总循环轮次 = 1 + 10 = 11轮 但每轮有效处理的数据范围 = 2048个元素16384个block × 10轮 = 163840次block内同步 8192个block × 11轮 = 90112次block内同步每个block处理数据量翻倍 ↓ 需要的block总数减半(16384→8192) ↓ 全局总同步轮次减半(163840→90112) ↓ __syncthreads()的总累积开销大幅降低 ↓ 整体耗时接近减半CPU发起kernel launch ↓ GigaThread Engine接收这次launch的所有block ↓ 按照SM的资源情况（寄存器/shared memory/warp slot是否够用） 把block一批一批地分发给各个SM ↓ SM执行完一个block后，GigaThread Engine再补充新的block进来
reduceUnroll2把grid.x从16384减到8192，GigaThread Engine需要追踪、分发、回收的block总数减少了一半。具体体现在：
追踪开销减少：GigaThread Engine内部维护着一张"哪些block已完成、哪些还在跑、哪些等待分发"的状态表，block数量减半，这张表的维护开销跟着减少。
分发轮次减少：RTX 4060（AD107核心）有24个SM，每个SM能同时驻留的block数有限（受寄存器和shared memory约束）。GigaThread Engine需要反复把block"喂"给SM，block总数减半意味着需要喂的轮次也跟着减少，调度器的工作量直接降低。
block完成通知减少：每个block执行完毕后，SM要通知GigaThread Engine"这个block跑完了，来新的"，这个通知本身也有开销，16384次通知变成8192次通知，开销减半。
- block数量减半 → 同等SM资源下调度压力减半
- 总同步轮次从163840降到90112，降幅约45%
- 加上block减少带来的launch和调度开销降低
__syncthreads()减少（总同步轮次降低） ↓ kernel总执行时间缩短（实测从0.00113201→0.00065589，快了约1.7倍） ↓ 搬运的数据总量基本不变（归约的数据量没变，sector/request=3.77≈3.58，访存模式基本一致） ↓ 相同数据量 ÷ 更短时间 = 单位时间内搬运更多数据 ↓ 带宽利用率从42.17%升到71.50%那么如果我们是提前一个block算四份数据呢？如果算八份数据呢？效率会不会提升？？？
可以理论分析一下，如果是四份的话，那就是block减少到原来的1/4，数量是4096
那__syncthreads()=11*4096=45056，相较于一开始的降72.5%，那性能时间应该会提升个四倍左右？
算8份数据的__syncthreads()=11*2048=22528，降86.25%，但是应该不会提升这么大，因为你还需要并行啊，需要多个block去填满SM，否则如果计算资源有空闲，你让每个block的计算很重，达不到轻量并行效果，所以性能的提升不是线性的，而是一个先突增然后缓慢增加，越来越慢
reduceInterleaved: 0.00113392 展开2: 0.00062298 （相比Interleaved快了约1.82倍） 展开4: 0.00050401 （相比展开2快了约1.24倍） 展开8: 0.00045585 （相比展开4快了约1.11倍）可以看到性能的提升并不是线性的，边际收益在快速衰减。
展开倍数继续增大 ├── 预归约阶段的额外开销越来越重（7行→15行→31行...指数增长） ├── 每个线程需要的寄存器数量增多（要暂存更多中间值）→ 寄存器溢出到local memory → 访存开销暴增 ├── grid.x减小到接近SM数量级别 → SM开始"吃不饱" → 延迟无法被掩盖 └── 带宽利用率已经接近硬件上限 → 继续优化访存没有空间接下来我们展开大一点看一下
实际的实验结果要看你当时的环境，这个有点像一个效率上升到最大，后面反而降下来了，主要还是得看的你数据量啊数据计算的难易程度啊之类的，要结合实际去分析
这里博主猜想会不会是因为单个block过于重，所以导致性能曲线是先增大后减小呢？？？
博主实力还不够，可能这块知识只是猜想，后期如果能力上来在回来验证
完全展开的归约
既然可以把前面的展开，那后面是不是也可以展开，当一个block只剩下64个数据块的时候，这部分是不是可以展开？
__global__ void reduceUnrollWarp8(int * g_idata,int * g_odata,unsigned int n){ unsigned int tid = threadIdx.x; unsigned int idx = blockDim.x*blockIdx.x*8+threadIdx.x; if (tid >= n) return; int *idata = g_idata + blockIdx.x*blockDim.x*8; if(idx + 7*blockDim.x < n){ g_idata[idx] = g_idata[idx] + g_idata[idx + blockDim.x] + g_idata[idx + 2*blockDim.x] + g_idata[idx + 3*blockDim.x] + g_idata[idx + 4*blockDim.x] + g_idata[idx + 5*blockDim.x] + g_idata[idx + 6*blockDim.x] + g_idata[idx + 7*blockDim.x]; } __syncthreads(); for (int stride = blockDim.x/2; stride>32 ; stride >>=1){ if (tid <stride){ idata[tid] += idata[tid + stride]; } __syncthreads(); } //做最后64数据块的归约 if(tid<32){ volatile int *vmem = idata; vmem[tid]+=vmem[tid+32]; vmem[tid]+=vmem[tid+16]; vmem[tid]+=vmem[tid+8]; vmem[tid]+=vmem[tid+4]; vmem[tid]+=vmem[tid+2]; vmem[tid]+=vmem[tid+1]; } if (tid == 0) g_odata[blockIdx.x] = idata[0]; }注意不一样的地方是for循环的结束条件是stride>32，那么就会剩下64个数据块没有被归约
volatile
在c语言当中我们讲过volatile关键字，这个关键字就是可见性，强制要求编译器每次读写某个数据时，不能从寄存器当中读写，而是必须从内存当中读写
因为如果编译器认为你这个线程反复读写同一个地址
vmem[tid]，优化在寄存器当中，导致数据没有刷新到内存，那么后续别的线程在用到上一步的中间结果的时候，会导致内存当中的数据是旧的，因为寄存器是私有的，不加绝对会出错tid=16写vmem[16] ↓ volatile强制写到L1，不留在寄存器里 L1 Cache（SM内所有线程共享） ↑ volatile强制从L1读，不用寄存器里的缓存值 tid=0读vmem[16] → 拿到最新值 ✓有
__syncthreads()，它本身就是一个内存屏障(memory barrier)，会强制所有线程在继续执行之前把所有写操作刷新到内存，并且让所有后续读操作能看到最新的值。所以有__syncthreads()的轮次，编译器的寄存器缓存优化不会造成跨线程的可见性问题，不需要volatile。很明显最后几轮是可以减少__syncthreads()的数量的，每个block减少了6轮，那总共有2048个block，又减少了12288个__syncthreads()
__syncthreads()为什么开销大：因为 Block 内的 Warp 可能分布在同一个 SM 的不同子分区。__syncthreads()需要跨子分区协调，让不同子分区的所有 Warp 都到达同步点，这涉及全局的通信和等待，比子分区内部的 Warp 切换开销大得多。分析到这里，我们似乎没有看过指标，通过网络上的查询
smsp__warp_issue_stalled_barrier_per_warp_active.pct==PROF== Disconnected from process 5419 [5419] main@127.0.0.1 warmUp(int *, int *, unsigned int) (16384, 1, 1)x(1024, 1, 1), Context 1, Stream 7, Device 0, CC 8.9 Section: Command line profiler metrics ---------------------------------------------------- ----------- ------------ Metric Name Metric Unit Metric Value ---------------------------------------------------- ----------- ------------ smsp__warp_issue_stalled_barrier_per_warp_active.pct % 18.52 ---------------------------------------------------- ----------- ------------ reduceNeighbored(int *, int *, unsigned int) (16384, 1, 1)x(1024, 1, 1), Context 1, Stream 7, Device 0, CC 8.9 Section: Command line profiler metrics ---------------------------------------------------- ----------- ------------ Metric Name Metric Unit Metric Value ---------------------------------------------------- ----------- ------------ smsp__warp_issue_stalled_barrier_per_warp_active.pct % 18.54 ---------------------------------------------------- ----------- ------------ reduceNeighboredLess(int *, int *, unsigned int) (16384, 1, 1)x(1024, 1, 1), Context 1, Stream 7, Device 0, CC 8.9 Section: Command line profiler metrics ---------------------------------------------------- ----------- ------------ Metric Name Metric Unit Metric Value ---------------------------------------------------- ----------- ------------ smsp__warp_issue_stalled_barrier_per_warp_active.pct % 52.98 ---------------------------------------------------- ----------- ------------ reduceInterleaved(int *, int *, unsigned int) (16384, 1, 1)x(1024, 1, 1), Context 1, Stream 7, Device 0, CC 8.9 Section: Command line profiler metrics ---------------------------------------------------- ----------- ------------ Metric Name Metric Unit Metric Value ---------------------------------------------------- ----------- ------------ smsp__warp_issue_stalled_barrier_per_warp_active.pct % 45.47 ---------------------------------------------------- ----------- ------------ reduceUnroll2(int *, int *, unsigned int) (8192, 1, 1)x(1024, 1, 1), Context 1, Stream 7, Device 0, CC 8.9 Section: Command line profiler metrics ---------------------------------------------------- ----------- ------------ Metric Name Metric Unit Metric Value ---------------------------------------------------- ----------- ------------ smsp__warp_issue_stalled_barrier_per_warp_active.pct % 32.71 ---------------------------------------------------- ----------- ------------ reduceUnroll4(int *, int *, unsigned int) (4096, 1, 1)x(1024, 1, 1), Context 1, Stream 7, Device 0, CC 8.9 Section: Command line profiler metrics ---------------------------------------------------- ----------- ------------ Metric Name Metric Unit Metric Value ---------------------------------------------------- ----------- ------------ smsp__warp_issue_stalled_barrier_per_warp_active.pct % 26.98 ---------------------------------------------------- ----------- ------------ reduceUnroll8(int *, int *, unsigned int) (2048, 1, 1)x(1024, 1, 1), Context 1, Stream 7, Device 0, CC 8.9 Section: Command line profiler metrics ---------------------------------------------------- ----------- ------------ Metric Name Metric Unit Metric Value ---------------------------------------------------- ----------- ------------ smsp__warp_issue_stalled_barrier_per_warp_active.pct % 14.12 ---------------------------------------------------- ----------- ------------ reduceUnrollWarp8(int *, int *, unsigned int) (2048, 1, 1)x(1024, 1, 1), Context 1, Stream 7, Device 0, CC 8.9 Section: Command line profiler metrics ---------------------------------------------------- ----------- ------------ Metric Name Metric Unit Metric Value ---------------------------------------------------- ----------- ------------ smsp__warp_issue_stalled_barrier_per_warp_active.pct % 11.33 ---------------------------------------------------- ----------- ------------reduceNeighboredLess这个为什么最高，消除divergence，但barrier暴涨，而且工作的线程聚集在前面
warp0~15：在认真做归约计算，执行时间相对较长 warp16~31：if(index<blockDim.x)不满足，直接跳到__syncthreads()等待 结果：warp16~31早早到达barrier， warp0~15还在慢慢算， 所有warp都要等warp0~15最慢的那个完成， 等待时间比所有warp均匀工作时更长工作分配不均匀，导致"快的线程等慢的线程"的时间变长，每次
__syncthreads()的实际等待时间比reduceNeighbored更久，即使调用次数一样，stall占比反而更高。
reduceInterleaved（45.47%）比reduceNeighboredLess（52.98%）好是因为你采用了交错分配，更快的完成了任务，导致等待的时间就少了，所以占比肯定上去了这印证了一个重要结论：barrier stall高低，不只取决于
__syncthreads()调用次数，更取决于block内各warp到达barrier时的时间差异——差异越大，先到的warp等待时间越长，stall越高。但是接下来的趋势大家的类似初始条件相同，但一直因为展开，所以减少了__syncthreads()
这样完成的更快，展开系列通过减少block总数、降低总同步轮次 barrier stall持续下降（45.47%→32.71%→26.98%→14.12%）
reduceUnrollWarp8进一步消除最后6轮的__syncthreads() barrier stall降至最低（11.33%）
更进一步的，我们将中间的循环展开
注意之前的最后一轮展开，是在一个线程束当中，所以不用__syncthreads()，而我们把for循环展开，一定要加__syncthreads()，不然肯定出错
__global__ void reduceCompleteUnrollWarp8(int * g_idata,int * g_odata,unsigned int n){ unsigned int tid = threadIdx.x; unsigned int idx = blockDim.x*blockIdx.x*8+threadIdx.x; if (idx >= n) return; int *idata = g_idata + blockIdx.x*blockDim.x*8; if(idx + 7*blockDim.x < n){ g_idata[idx] = g_idata[idx] + g_idata[idx + blockDim.x] + g_idata[idx + 2*blockDim.x] + g_idata[idx + 3*blockDim.x] + g_idata[idx + 4*blockDim.x] + g_idata[idx + 5*blockDim.x] + g_idata[idx + 6*blockDim.x] + g_idata[idx + 7*blockDim.x]; } __syncthreads(); if(blockDim.x>=1024 && tid <512) idata[tid]+=idata[tid+512]; __syncthreads(); if(blockDim.x>=512 && tid <256) idata[tid]+=idata[tid+256]; __syncthreads(); if(blockDim.x>=256 && tid <128) idata[tid]+=idata[tid+128]; __syncthreads(); if(blockDim.x>=128 && tid <64) idata[tid]+=idata[tid+64]; __syncthreads(); //做最后64数据块的归约 if(tid<32){ volatile int *vmem = idata; //int *vmem = idata; vmem[tid]+=vmem[tid+32]; vmem[tid]+=vmem[tid+16]; vmem[tid]+=vmem[tid+8]; vmem[tid]+=vmem[tid+4]; vmem[tid]+=vmem[tid+2]; vmem[tid]+=vmem[tid+1]; } if (tid == 0) g_odata[blockIdx.x] = idata[0]; }实验结果
这里的优化在于把中间的for循环展开，写成了一个一个的if语句
对于if语句：if(blockDim.x>=1024 && tid <512)
这种会在编译器完全确定，
blockDim.x是编译期已知的常量（kernel launch时确定，编译器可以在编译时就知道这个值），所以这些if条件在编译阶段就会被判断，不会出现在运行时的指令流里。对于blockDim.x=1024的时候，直接优化
idata[tid]+=idata[tid+512]; __syncthreads(); idata[tid]+=idata[tid+256]; __syncthreads(); idata[tid]+=idata[tid+128]; __syncthreads(); idata[tid]+=idata[tid+64]; __syncthreads();当你blockDIm.x=512的时候，直接优化
idata[tid]+=idata[tid+256]; __syncthreads(); idata[tid]+=idata[tid+128]; __syncthreads(); idata[tid]+=idata[tid+64]; __syncthreads();没有任何循环跳转、没有循环计数器、没有循环条件判断，是完全展开的线性指令序列。
总结
注意有时候性能提升很少，绝大多数是因为编译器进化的问题，编译器帮你把绝大多数能够优化的提前优化了，而不需要你显示的去优化，所以有时候实验的结果不太和理论分析一样，这取决于硬件，取决于实验的环境，取决于编译器等等
版本 新增的优化点 reduceInterleaved 解决divergence，访存合并改善 reduceUnroll2/4/8 预归约减少block总数，总同步__syncthreads()轮次降低 reduceUnrollWarp8 最后6轮去掉 __syncthreads()，用volatile代替reduceCompleteUnrollWarp8 block内归约循环彻底展开为编译期常量判断的线性序列，消除循环控制开销+让编译器做更激进的指令调度
扫一扫
529
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
