第4章 CUDA核函数优化

第四章：CUDA核函数优化——从理论到实战

各位同学，欢迎来到第四章。前面我们学会了怎么写CUDA代码，怎么启动核函数。但说实话，能跑和跑得快，那是两码事。这一章，我们就来聊聊怎么让GPU真正“跑起来”。

我个人习惯把GPU优化分成几个层次：线程束级别、内存级别、计算单元级别。今天咱们一个一个来啃。

4.1 线程束（Warp）与分支发散

先问个问题：GPU里最小的执行单位是什么？很多人会说是线程。其实不对，真正的最小执行单位是线程束（Warp）。

一个Warp包含32个线程，它们一起执行同一条指令。你想想看，如果这32个线程走的是不同的分支路径，会发生什么？

核心概念：分支发散（Branch Divergence）——当Warp内的线程进入不同分支时，所有分支都会被串行执行，没走那个分支的线程就被“屏蔽”了。

我在项目中遇到过这样一个案例：一个简单的if-else判断，让性能直接掉了40%。排查后发现，就是分支发散搞的鬼。

// ❌ 糟糕的分支发散
__global__ void bad_kernel(float* data, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx % 2 == 0) {
        data[idx] = data[idx] * 2.0f;
    } else {
        data[idx] = data[idx] + 1.0f;
    }
}

// ✅ 优化后：用条件运算代替分支
__global__ void good_kernel(float* data, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    float val = data[idx];
    val = (idx % 2 == 0) ? (val * 2.0f) : (val + 1.0f);
    data[idx] = val;
}

避坑指南：我曾经为了代码可读性，在核函数里写了一大堆switch-case。结果性能惨不忍睹。后来我学乖了：能用三元运算符就别用if-else，能用查表法就别用分支判断。

4.2 内存合并访问

说到内存访问，很多新手会忽略一个关键点：全局内存的访问模式。

GPU从全局内存读数据时，是以128字节为单位的。如果同一个Warp里的32个线程访问的是连续的内存地址，那一次就能搞定。如果访问的是分散的地址，那就得多次访问——这就是内存合并访问的核心思想。

// ❌ 非合并访问：按列访问
__global__ void bad_access(float* A, int width) {
    int row = threadIdx.x;
    int col = blockIdx.x;
    float val = A[col * width + row];  // 跨步访问
}

// ✅ 合并访问：按行访问
__global__ void good_access(float* A, int width) {
    int row = blockIdx.x;
    int col = threadIdx.x;
    float val = A[row * width + col];  // 连续访问
}

说白了，让线程ID和内存地址保持线性关系，这就是合并访问的精髓。我刚开始做矩阵乘法时，就是没注意这个细节，性能只有理论峰值的10%。

4.3 Bank Conflict

共享内存（Shared Memory）是GPU里的一块“宝地”，但用不好也会踩坑。共享内存被分成32个Bank，每个Bank每周期只能响应一次访问。

如果同一个Warp里的多个线程访问同一个Bank的不同地址，就会发生Bank Conflict。这会导致访问被串行化，性能下降。

注意：Bank Conflict只发生在共享内存的访问上。全局内存没有这个问题，但全局内存有合并访问的问题。

// ❌ 有Bank Conflict：步长为2
__shared__ float shared[32][32];
float val = shared[threadIdx.x][threadIdx.y];  // 同一列访问同一个Bank

// ✅ 无Bank Conflict：加个padding
__shared__ float shared[32][33];  // 多一列padding
float val = shared[threadIdx.x][threadIdx.y];  // 错开Bank

嗯，这里要注意：加padding是最常用的解决Bank Conflict的方法。我做过一个FFT的优化，加了一行padding后，性能提升了30%。

4.4 Occupancy（占用率）计算与调优

Occupancy，说白了就是SM上活跃的Warp数量占最大Warp数量的比例。很多人以为Occupancy越高越好，其实不一定。

我举个例子：一个核函数用了很多寄存器，每个SM只能同时跑很少的Warp。这时候Occupancy低，但每个Warp跑得快。另一个核函数寄存器用得少，Occupancy高，但每个Warp可能因为内存延迟而等待。

关键是要找到平衡点。

资源限制	影响	调优方向
寄存器数量	每个线程用的寄存器越多，Occupancy越低	用__launch_bounds__限制寄存器数
共享内存	每个Block用的共享内存越多，Occupancy越低	尽量复用共享内存，减少申请量
Block大小	Block太小或太大都会影响Occupancy	通常128-256线程/Block比较合适

Occupancy计算公式：

Occupancy = (活跃Warp数) / (SM最大Warp数) × 100%

活跃Warp数 = min(寄存器限制Warp数, 共享内存限制Warp数, Block限制Warp数)

我个人习惯用NVIDIA的Occupancy Calculator（Excel表格）来快速估算。不过现在Nsight已经集成了这个功能，更方便了。

4.5 使用NVIDIA Nsight进行性能分析

光靠猜是不行的，得用工具说话。Nsight就是我们的“照妖镜”。

我常用的几个功能：

Kernel Profiling：看每个核函数的执行时间、Occupancy、内存带宽利用率
Warp State：看Warp在干什么——是在计算、在等待内存、还是在stall
Memory Access Pattern：可视化内存访问模式，一眼看出是不是合并访问
Source Correlation：把性能数据和源代码关联起来，精确定位瓶颈

我的调试流程：

先跑一次基线，记录性能数据
看Occupancy，如果低于50%，先调资源限制
看Warp Stall原因，如果是内存等待，优化内存访问模式
看分支发散情况，重构代码减少分支
重复1-4，直到性能达标

我曾经接手过一个项目，核函数跑得很慢。用Nsight一看，发现Warp有60%的时间在stall等待内存。优化了内存合并访问后，性能直接翻倍。所以说，工具用得好，下班走得早。

本章知识体系

下面这张图总结了本章的核心逻辑，我建议你保存下来，做优化时对照着看：

好了，这一章的内容就到这里。记住：优化不是玄学，是有章可循的工程实践。多动手、多跑Nsight、多对比数据，你也能成为GPU优化高手。

公众号:蓝海资料掘金营，微信deep3321

📚 课程目录

第四章：CUDA核函数优化——从理论到实战

4.1 线程束（Warp）与分支发散

4.2 内存合并访问

4.3 Bank Conflict

4.4 Occupancy（占用率）计算与调优

4.5 使用NVIDIA Nsight进行性能分析

本章知识体系