第2章：CUDA 性能优化基础

写出能跑的 CUDA 代码只是起点,写出跑得快的代码才是 AI Infra 工程师的核心能力。本章建立 CUDA 性能优化的核心方法论:从 Warp 执行模型出发,讲合并访存、Bank Conflict、Occupancy、同步与原子操作,以及向量化加载这些 4-5 倍性能差距的关键开关。

📑 目录

• 1. SIMT 执行模型与 Warp Divergence
• 2. 全局内存:合并访问的力量
• 3. 共享内存:Bank Conflict
• 4. 向量化加载:float4/int4
• 5. Occupancy 与资源平衡
• 6. 同步与原子操作
• 7. Warp Shuffle:跳过 Shared Memory
• 自我检验清单
• 参考资料

---

1. SIMT 执行模型与 Warp Divergence

1.1 Warp 是 GPU 的真实”指令单位”

GPU 一次发射的不是一个线程的指令,而是 整个 Warp(32 线程) 的同一条指令。这就是 SIMT(Single Instruction, Multiple Thread)。

1.2 Warp Divergence 的代价

__global__ void diverge(int* x) {
    int tid = threadIdx.x;
    if (tid % 2 == 0) {
        x[tid] = compute_a();    // 偶数线程做 A,奇数线程闲等
    } else {
        x[tid] = compute_b();    // 然后奇数做 B,偶数闲等
    }
}

Divergence 时性能损失最高 50%(只在两条路径长度相等时,实际可能更糟)。

优化原则:让 Warp 内的线程尽量走同一路径。

// 反例:按 tid 奇偶分支
if (tid % 2 == 0) ... else ...

// 正例:按 Warp ID 分支(同 Warp 内 tid 都同奇偶)
int warp_id = tid / 32;
if (warp_id % 2 == 0) ... else ...

1.3 不可避免的 Divergence:边界处理

__global__ void vec_add(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];   // 最后一个 Warp 可能 divergence
}

边界判断不可避免,但只在最后一个 Block 影响,可以接受。

---

2. 全局内存:合并访问的力量

2.1 什么是合并访问

GPU 一次内存事务最多搬 128 字节(一条 cache line)。Warp 内 32 个线程的访问地址如果对齐到 128B 且连续,只需一次 transaction;否则会拆成多次。

✅ 合并访问(1 次 transaction):
线程 0  访问 a[0]   ┐
线程 1  访问 a[1]   │
...                │  连续地址,128B 对齐
线程 31 访问 a[31]  ┘

❌ 跨步访问(32 次 transaction):
线程 0  访问 a[0]
线程 1  访问 a[1024]   ← 跨度 4096 字节
...

2.2 实测对比

// 合并访问
__global__ void coalesced(float* a, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    a[i] += 1.0f;
}

// 跨步访问(stride = 32)
__global__ void strided(float* a, int n) {
    int i = (blockIdx.x * blockDim.x + threadIdx.x) * 32;
    if (i < n) a[i] += 1.0f;
}

A100 上,合并访问吞吐 ~1500 GB/s,跨步访问 ~50 GB/s——30 倍差距。

2.3 设计准则

让相邻线程访问相邻地址,越简单越好。

如果数据布局天生不合并(比如 column-major 访问 row-major 数组),就用 Shared Memory 中转(参考第 1 章的矩阵转置)。

---

3. 共享内存:Bank Conflict

3.1 什么是 Bank

Shared Memory 被切成 32 个 Bank,每个 Bank 一次能服务一个线程的访问。Warp 内 32 个线程同时访问 32 个不同 Bank → 1 cycle 完成;访问同一个 Bank → 串行,N-way conflict 慢 N 倍。

3.2 Bank 索引计算

每 4 字节(一个 float)轮流分配到一个 Bank:

地址 0  → Bank 0
地址 4  → Bank 1
...
地址 124 → Bank 31
地址 128 → Bank 0   (回到 Bank 0)

3.3 经典冲突:32×32 矩阵转置

__shared__ float tile[32][32];        // ❌ 32-way conflict
tile[threadIdx.y][threadIdx.x] = ...; // 写没冲突(行连续)
... = tile[threadIdx.x][threadIdx.y]; // 读冲突!所有线程访问同一列

读取 tile[*][threadIdx.y] 时,Warp 内 32 线程(threadIdx.x = 0..31)都访问列 threadIdx.y,这一列的所有元素地址间隔正好 32 floats = 128 字节 → 都落在同一个 Bank → 32-way conflict。

3.4 Padding 技巧

__shared__ float tile[32][33];        // ✅ 加一列 padding

把列宽改成 33,每行多一个空 float,列方向访问的步长不再是 32 整数倍 → Bank 错开 → 无冲突。

实测加 padding 后,矩阵转置性能提升 30-40%。

3.5 用 Nsight Compute 验证

ncu --section MemoryWorkloadAnalysis ./my_kernel

报告中的 Shared Memory Bank Conflicts 一栏直接告诉你冲突次数,优化前后对比一目了然。

---

4. 向量化加载:float4/int4

4.1 一次搬 128 位

GPU 支持 16 字节(128 位)的单条 load/store 指令。把 4 个 float 打包成 float4:

// 标量版本:每线程 1 个 float
__global__ void scalar(const float* a, float* b, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) b[i] = a[i] * 2.0f;
}

// 向量化:每线程 4 个 float
__global__ void vectorized(const float4* a, float4* b, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        float4 v = a[i];
        v.x *= 2.0f; v.y *= 2.0f; v.z *= 2.0f; v.w *= 2.0f;
        b[i] = v;
    }
}

4.2 性能收益

向量化的好处不只是”一次搬更多”:

• 减少 instruction count(指令变 1/4)
• 提高指令级并行度
• 更好地利用 cache line

实测 A100,向量加法用 float4 比标量快 1.5-2 倍。

4.3 注意事项

• 数据起始地址必须 16 字节对齐
• 长度不是 4 的倍数时,需要处理边界(剩余的几个用标量补齐)

---

5. Occupancy 与资源平衡

5.1 Occupancy 定义

$$\text{Occupancy} = \frac{\text{每 SM 实际活跃 Warp 数}}{\text{每 SM 最大 Warp 数(64)}}$$

每个 SM 同时能跑多少个 Warp,取决于三个资源谁先耗尽:

1. 寄存器数:每 SM 64K 个,Block 总寄存器数不能超
2. Shared Memory:每 SM ~228 KB
3. Block / Warp 数上限:64 Warp / SM,32 Block / SM

5.2 Latency Hiding

GPU 通过 Warp 切换隐藏长延迟操作(等 HBM、等 Tensor Core)。Occupancy 越高,可切换的 Warp 越多,延迟隐藏越好。

但Occupancy 不是越高越好——有些 Kernel(尤其用 Tensor Core 的 GEMM)用满寄存器,占用 50% Occupancy 反而比 100% 快,因为减少了寄存器溢出。

5.3 实战调优流程

1. ncu --query-metrics-mode all 看当前 Occupancy
2. 如果 Occupancy 低,看 Limiter:Register / Shared Mem / Block Size?
3. 调整:
   • 减少寄存器:__launch_bounds__(threads, blocks) 提示编译器
   • 减少 Shared Mem:用更小的 tile
   • 调 Block Size:128/256/512 试一遍

__global__ __launch_bounds__(256, 4)    // 256 线程/Block,期望 4 个 Block/SM
void my_kernel(...) { ... }

---

6. 同步与原子操作

6.1 同步层次

同步	范围	用途
__syncthreads()	Block 内	Shared Memory 一致性
__syncwarp()	Warp 内	比 __syncthreads 快
cudaDeviceSynchronize()	全 GPU	主机等设备
__threadfence()	全 GPU	内存可见性,不阻塞

6.2 原子操作

atomicAdd(&counter, 1);          // int 原子加
atomicMax(&max_val, x);          // 原子取最大
atomicCAS(&lock, 0, 1);          // 比较并交换

原子操作很慢——所有冲突的线程要串行执行。优先用 reduction 算法替代(比如先 Block 内归约,Block 间用一次 atomic)。

6.3 Reduction 替代 Atomic 的例子

// 反例:每个线程都 atomicAdd
__global__ void bad_sum(const float* x, float* total, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) atomicAdd(total, x[i]);   // N 次 atomic,极慢
}

// 正例:Block 内规约,Block 间 atomic
__global__ void good_sum(const float* x, float* total, int n) {
    __shared__ float s[256];
    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + tid;
    s[tid] = (i < n) ? x[i] : 0;
    __syncthreads();
    for (int k = 128; k > 0; k >>= 1) {
        if (tid < k) s[tid] += s[tid + k];
        __syncthreads();
    }
    if (tid == 0) atomicAdd(total, s[0]);   // 只 atomic 一次/Block
}

第 3 章会讲更优的 Warp Shuffle 版本。

---

7. Warp Shuffle:跳过 Shared Memory

Warp Shuffle 让同一 Warp 内的 32 线程直接交换寄存器值,不经过 Shared Memory——更快、不占 Shared Memory。

// __shfl_xor_sync:树形归约
float val = my_value;
for (int offset = 16; offset > 0; offset /= 2)
    val += __shfl_xor_sync(0xffffffff, val, offset);
// 此时 Warp 内每个线程的 val 都是 32 个原始值的总和

指令	含义
__shfl_sync(mask, var, src)	从 src lane 取值
__shfl_up_sync(mask, var, n)	从上 n lane 取值
__shfl_down_sync(mask, var, n)	从下 n lane 取值
__shfl_xor_sync(mask, var, n)	与 lane^n 交换值

Warp Shuffle 是高性能 Reduce / Scan 算子的核心——下一章 Reduce 算子会大量使用。

---

✅ 自我检验清单

• Warp Divergence:能解释 SIMT 模型,以及为什么 if/else 在 Warp 内会串行
• 合并访问:能写一对正反例(coalesced vs strided),并预测性能差距
• Bank Conflict:能解释 32×32 矩阵转置为什么有 32-way conflict,加 +1 padding 为什么解决问题
• Float4 向量化:能改造一个标量 Kernel 为 float4 版本,知道收益和注意事项
• Occupancy 调优:看 Nsight 报告能识别 limiter,知道怎么用 __launch_bounds__ 调节
• Atomic 与 Reduction:能解释为什么 atomic 慢,并能写出正确的 Block-level 归约版本
• Warp Shuffle:能用 __shfl_xor_sync 写一个 Warp 内的 sum reduction
• Profile 实操:能用 ncu 抓一次 kernel,识别它是 memory bound 还是 compute bound

---

📚 参考资料

• CUDA C Best Practices Guide:https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/
• CUDA Occupancy Calculator:NVIDIA 官方工具
• PeakCrosser:CUDA Reduce 算子优化 —— 知乎专栏,详尽实战
• NVIDIA Blog:Using Shared Memory in CUDA C/C++
• Faster Parallel Reductions on Kepler (NVIDIA Blog) —— Warp Shuffle 经典案例