第四章：GPU编程进阶（CUDA）——Bank Conflict与Shared Memory优化、Warp Divergence与Occupancy、原子操作与同步原语

好，咱们进入第四章。这一章，说白了就是CUDA性能调优的硬核部分。很多同学写CUDA代码，能跑，但跑不快。为什么？瓶颈往往就卡在Shared Memory、Warp调度和原子操作这几个点上。我当年刚接触GPU时，也在这上面栽过跟头，今天咱们就把这些坑一个个填平。

4.1 Bank Conflict与Shared Memory优化

Shared Memory是GPU上的一块高速缓存，比全局内存快两个数量级。但它有个脾气——它被分成了32个Bank（在计算能力5.x及以上是32个）。每个Bank每个周期只能响应一次访问请求。

什么是Bank Conflict？

想象一下，你有32个银行柜台（Bank），每个Warp里的32个线程同时去取钱。如果每个线程去不同的柜台，完美，一次搞定。但如果两个线程去了同一个柜台，那就得排队了——这就是Bank Conflict。

核心结论：Bank Conflict会导致Shared Memory访问串行化，严重降低带宽利用率。

常见的冲突模式：

  • 无冲突：每个线程访问不同的Bank。例如 shared[tid]，其中tid是线程ID。
  • 2路冲突：两个线程访问同一个Bank。性能减半。
  • 32路冲突：所有线程访问同一个Bank。性能降到1/32。

我遇到过最典型的场景是做矩阵转置时，没注意Bank映射，结果性能惨不忍睹。后来加了个padding就解决了。

如何避免Bank Conflict？

嗯，这里有几个实用技巧：

  1. Padding技巧：在声明Shared Memory时，多分配一个元素。比如 __shared__ float s[32][33]; 而不是 s[32][32]。多出来的那一列会让Bank映射错开，避免冲突。
  2. 改变访问模式：尽量让连续的线程访问连续的地址。比如 s[threadIdx.x] 就是无冲突的。
  3. 使用向量化加载：用 float4 一次读4个float，减少访问次数。

我的经验：在写卷积或矩阵乘的Shared Memory版本时，我习惯先画个Bank映射图。画完你就知道哪里会冲突了。别偷懒，这步很值。

代码示例：

// 有Bank Conflict的版本
__shared__ float s[32][32];
float val = s[threadIdx.x][threadIdx.y]; // 如果threadIdx.x相同，就会冲突

// 无Bank Conflict的版本（加Padding）
__shared__ float s[32][33];
float val = s[threadIdx.x][threadIdx.y]; // Padding让Bank错开

4.2 Warp Divergence与Occupancy

Warp是GPU执行的基本单位，32个线程一起执行同一条指令。如果这32个线程走了不同的分支（比如if-else），会发生什么？

Warp Divergence（束内分歧）

说白了，就是同一个Warp里的线程，有的走if，有的走else。GPU没办法同时执行两条路径，只能先执行if路径（其他线程等待），再执行else路径（其他线程等待）。这样性能就减半了。

我见过最夸张的例子，有人写了个switch-case，32个线程走了32个不同的分支。结果一个Warp的执行时间变成了原来的32倍。嗯，这代码要是上线，老板得找你喝茶。

如何避免Warp Divergence？

  • 让分支条件基于线程ID的某几位：比如 if (threadIdx.x < 16)，这样前16个线程走一个分支，后16个走另一个，不会在同一个Warp内产生分歧。
  • 用三元运算符代替分支：有些情况下 a = (cond) ? b : c; 比if-else更友好。
  • 数据重排：把相同条件的数据集中到一起处理。

Occupancy（占用率）

Occupancy是指每个SM上活跃的Warp数量与最大支持Warp数量的比值。很多人以为Occupancy越高越好，其实不一定。

Occupancy 优点 缺点
高（>80%） 能很好地隐藏内存延迟 每个线程可用资源少（寄存器、Shared Memory受限）
低（<50%） 每个线程可用资源多 延迟隐藏能力差

我个人的习惯是：对于计算密集型任务（比如矩阵乘），我会适当降低Occupancy，给每个线程多分配寄存器，减少寄存器溢出。对于内存密集型任务（比如数据拷贝），我会追求高Occupancy来隐藏延迟。

注意：不要盲目追求100% Occupancy。我曾经优化一个kernel，把Occupancy从100%降到67%，性能反而提升了30%。因为每个线程有了更多寄存器，避免了spill到local memory。

4.3 原子操作与同步原语

原子操作是并发编程的基石。在GPU上，多个线程同时读写同一个地址时，必须用原子操作来保证数据一致性。

常用的原子操作：

  • atomicAdd()：原子加
  • atomicSub()：原子减
  • atomicExch()：原子交换
  • atomicCAS()：比较并交换（Compare And Swap）

原子操作的性能陷阱：

原子操作本质上是串行化的。如果32个线程同时对一个地址做atomicAdd，那它们就得排队。性能会急剧下降。

我曾经优化过一个直方图统计的kernel，一开始直接用atomicAdd，结果性能惨不忍睹。后来我用了Shared Memory做局部统计，最后再合并到全局，性能提升了10倍。

优化策略：

  1. 使用Shared Memory做局部归约：每个Block先在Shared Memory里用原子操作统计局部结果，最后再一次性写到全局。
  2. Warp级归约：利用Warp Shuffle指令（__shfl_down_sync()）在Warp内部做归约，避免原子操作。
  3. 使用更细粒度的原子操作：比如用 atomicCAS 实现自定义的原子操作。

同步原语：

除了原子操作，同步也是关键。CUDA提供了：

  • __syncthreads()：Block内所有线程的同步点。注意，它只同步同一个Block内的线程。
  • __threadfence()：保证当前线程的写操作对其他线程可见。
  • __threadfence_block()：Block级别的内存栅栏。

避坑指南：我曾经在if-else分支里用了__syncthreads()，结果死锁了。因为有些线程走了if，有些走了else，它们永远到不了同一个同步点。记住：__syncthreads()必须被Warp里所有线程执行到，否则就死锁。

代码示例：

// 使用Shared Memory减少全局原子操作
__global__ void histogram_kernel(int *data, int *hist, int n) {
    __shared__ int s_hist[256]; // 每个Block的局部直方图
    
    // 初始化Shared Memory
    if (threadIdx.x < 256) s_hist[threadIdx.x] = 0;
    __syncthreads();
    
    // 局部统计
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        int val = data[tid] & 0xFF;
        atomicAdd(&s_hist[val], 1);
    }
    __syncthreads();
    
    // 合并到全局
    if (threadIdx.x < 256) {
        atomicAdd(&hist[threadIdx.x], s_hist[threadIdx.x]);
    }
}

好了，这一章的内容就这些。Bank Conflict、Warp Divergence、原子操作——这三个是CUDA性能优化的核心。你想想看，如果能把这三点吃透，大部分kernel的性能问题你都能找到症结所在。下一章咱们会讲更高级的优化技巧，包括寄存器优化和指令级并行。到时候见。

← 上一章 📖 返回目录 下一章 →