第6章：CUDA 性能优化深入

Ch5 给了工具链——怎么看到瓶颈。这一章给方法论——怎么解决瓶颈。CUDA 单 kernel 调优有几十年的工程经验，但核心其实就六件事：让访存合并、避免 bank conflict、平衡 occupancy 和 ILP、用好 shared memory、用 Hopper / Blackwell 的 TMA / WGMMA、做 kernel fusion。这一章把这六件事讲清楚，并用”朴素 GEMM → cuBLAS 级”的 7 步演进作为贯穿全章的实战例子。

📑 目录

1. Memory coalescing：访存合并是性价比最高的优化
2. Bank conflict 与 vector load
3. Occupancy vs ILP 的取舍
4. Shared memory tiling 与 swizzling
5. Hopper TMA 与 WGMMA
6. Kernel fusion 的几种范式
7. 朴素 GEMM → cuBLAS 级的 7 步演进
自我检验清单
参考资料

1. Memory coalescing：访存合并是性价比最高的优化

1.1 为什么 coalescing 这么重要

GPU 的 DRAM（HBM）访问粒度是 128 字节——一个 transaction 至少传 128 字节。warp 里 32 个线程同时访存：

如果 32 个线程访问的地址恰好覆盖连续 128 字节，硬件合并成 1 个 transaction
如果 32 个线程的地址分散在 32 个不同 cache line，变成 32 个 transaction——访存量虚增 32 倍

这就是 “memory coalescing”——一个完全靠数据布局决定的事，但性能差距能到 30 倍。

1.2 怎么写出 coalesced 代码

正确模式（连续 thread → 连续地址）：

// 假设 A 是 [N][M] 的 float 数组，每行长 M
__global__ void copy_row(float* A, float* B, int M) {
    int tid = threadIdx.x;
    int row = blockIdx.x;
    // thread tid 访问 A[row][tid]：连续地址 → coalesced
    B[row * M + tid] = A[row * M + tid];
}

错误模式（连续 thread → 跨步地址）：

__global__ void copy_col(float* A, float* B, int N, int M) {
    int tid = threadIdx.x;
    int col = blockIdx.x;
    // thread tid 访问 A[tid][col]：跨步 M → 不 coalesced
    B[tid * M + col] = A[tid * M + col];
}

1.3 ncu 怎么诊断

在 Nsight Compute 里，看：

L1 Cache Hit Rate：高 hit rate 说明数据复用好（不直接关 coalescing，但相关）
DRAM Throughput / Compute Throughput：memory bound 时优先看 coalescing
Source Counters → Memory Accesses：找到具体哪一行的”sectors per request”高（理想 = 1，差 = 8 / 16 / 32）

1.4 一个常见错误：在 outer loop 上分线程

// 错误：thread x 跑外层 i，访存跨 N
for (int i = threadIdx.x; i < M; i += blockDim.x) {
    for (int j = 0; j < N; j++) {
        sum += A[i * N + j];  // 不 coalesced (相邻 thread 跨 N 远)
    }
}

// 正确：把外层换成 j，让相邻 thread 访问相邻地址
for (int j = threadIdx.x; j < N; j += blockDim.x) {
    for (int i = 0; i < M; i++) {
        sum += A[i * N + j];  // coalesced
    }
}

🌟 经验法则：写循环时先考虑 thread x 跑哪个 index——总让相邻 thread x 访问相邻地址。

2. Bank conflict 与 vector load

2.1 Bank conflict 是什么

Shared memory 物理上分成 32 个 bank（在 4-byte 边界上轮询）。一个 warp 的 32 个线程同时访问 shared memory：

如果 32 个线程访问 32 个不同 bank → 1 个 cycle 完成
如果 N 个线程访问同一个 bank（不同地址）→ 串行 N 次

这是 shared memory 的”内部 coalescing”问题——比 DRAM coalescing 更隐蔽。

2.2 一个经典 bank conflict

__shared__ float tile[32][32];
int tid = threadIdx.x;
// 32 个 thread 同时访问 tile[0..31][0]
// 列方向：tile[0][0], tile[1][0], ..., tile[31][0]
// 这些地址 stride 是 32（不是 1）→ 全在同一 bank → 32-way conflict
float val = tile[tid][0];  // 慢 32 倍

2.3 修复：padding 或 swizzle

Padding 是最简单的修复：

__shared__ float tile[32][33];  // 多一列 padding
// 现在 tile[0][0], tile[1][0], ... 的 stride 是 33（与 32 互质）
// 访存被均匀分布到所有 bank

Swizzle 是更高级的方案——按公式重新映射 row × col 到 bank：

int swizzled_col = (col ^ (row & 0x1f));  // XOR swizzle
float val = tile[row][swizzled_col];

CUTLASS 库里有大量 swizzle 模板可参考。

2.4 Vector load (float4 / int4 / __ldg)

让一条 load 指令同时取 128 bit（4 个 float）：

// 标量 load：4 条指令，4 次访存
for (int i = 0; i < 4; i++) sum += A[i];

// 向量 load：1 条指令，1 次访存（128 bit 对齐时）
float4 v = *(float4*)(A);
sum += v.x + v.y + v.z + v.w;

vector load 减少：

指令数（cleaner pipeline）
访存指令数（更接近 coalesced ideal）

⭐ 经验：能用 float4 就用——通常能拿到 30-50% 的访存效率提升。

3. Occupancy vs ILP 的取舍

3.1 Occupancy 是什么

Occupancy = (实际活跃 warp 数) / (理论最大 warp 数)。理论最大由：

SM 上的 warp slot 数（Hopper 是 64）
Block 占的 register 数 × 32（每 warp 寄存器需求）
Block 占的 shared memory 数

理论上 occupancy 越高，SM 在 stall 时有越多其他 warp 可调度——但这不总是越高越好。

3.2 Low occupancy 也能跑得好：ILP 视角

ILP（Instruction-Level Parallelism）的核心想法是：在一个 warp 内通过 unroll 增加并行度——不依赖 occupancy 来填 stall。

// Low occupancy 但高 ILP 的写法
__global__ void kernel(...) {
    float a0, a1, a2, a3;  // 多个独立寄存器
    a0 = ...;  // 几条独立指令并发执行
    a1 = ...;
    a2 = ...;
    a3 = ...;
    
    // 编译器会把这些视为独立 chain，发到不同 pipeline
}

3.3 为什么 low occupancy 有时更快

例子：GEMM kernel

高 occupancy 写法：每个 thread 算 1 个 output——register 用得少，occupancy 90%
低 occupancy 写法：每个 thread 算 8x8=64 个 output——register 用很多，occupancy 25%

低 occupancy 版本通常反而更快：

单 thread 算 64 output → register 文件复用好 → 减少 shared memory 访问
单 thread 内部 ILP 充分 → 不靠 occupancy 也能掩盖 stall

3.4 ncu 怎么看

Achieved Occupancy：实际 occupancy
Stall Reasons：如果 stall 主要是 long_scoreboard（DRAM stall），高 occupancy 有用
如果 stall 主要是 wait（短 stall）：低 occupancy + 高 ILP 通常更划算

4. Shared memory tiling 与 swizzling

4.1 为什么需要 tiling

Memory bound kernel 的本质是 每次计算从 DRAM 拉数据 → 算一次 → 扔掉。如果同样的数据要被多次复用，每次都从 DRAM 拉就是浪费。

Tiling 的核心：把一块数据先搬进 shared memory，所有相关计算都从 shared memory 读——把 N 次 DRAM 访问变成 1 次。

4.2 GEMM 的经典 tiling

__global__ void gemm_tiled(float* A, float* B, float* C, int M, int N, int K) {
    __shared__ float As[TILE][TILE];
    __shared__ float Bs[TILE][TILE];
    
    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    
    float c = 0;
    for (int t = 0; t < K / TILE; t++) {
        // 1. 协作把 A 和 B 的一个 tile 装进 shared memory
        As[ty][tx] = A[(by*TILE+ty) * K + t*TILE+tx];
        Bs[ty][tx] = B[(t*TILE+ty) * N + bx*TILE+tx];
        __syncthreads();
        
        // 2. 用 shared memory 算这个 tile 的部分和
        for (int k = 0; k < TILE; k++) {
            c += As[ty][k] * Bs[k][tx];  // shared mem 访问
        }
        __syncthreads();
    }
    C[(by*TILE+ty) * N + bx*TILE+tx] = c;
}

每个 tile 元素被复用 TILE 次（沿 k 维）——访存效率 TILE 倍提升。

4.3 Swizzle：避开 bank conflict

As[ty][tx] 的访问模式如果按行读写连续，按列读时会 conflict。常见做法：

Padding（如 §2.3）
Swizzle index：tile[ty][tx ^ (ty & 0x07)]
用 ldmatrix / TMA 这些硬件指令自动处理（H100+）

4.4 Double buffering

把 tile 装载和 tile 计算重叠：

__shared__ float As[2][TILE][TILE];  // 双 buffer
__shared__ float Bs[2][TILE][TILE];

// 装第 0 个 tile
load_tile(As[0], Bs[0], 0);
__syncthreads();

for (int t = 1; t < num_tiles; t++) {
    int curr = (t-1) % 2;
    int next = t % 2;
    
    // 用 curr buffer 算的同时，把 next tile 装进来
    compute_tile(As[curr], Bs[curr]);
    load_tile_async(As[next], Bs[next], t);
    
    __syncthreads();
}

这是 GEMM 性能再上一个台阶的关键技术——计算和访存并行。

5. Hopper TMA 与 WGMMA

5.1 TMA：Tensor Memory Accelerator

TMA 是 H100 引入的异步 DMA 引擎——专门搬大块连续数据，不占 SM 计算单元。

传统做法：

每个 thread 算自己的地址 → 32 个 thread 协作搬一个 tile
占用 32 个 thread 的算力

TMA 做法：

一条 TMA 指令搬整个 tile（直接从 DRAM 到 shared memory）
不占 thread 算力——SM 全部 thread 都可以做计算

优势：

自动处理多维 tile 的访存模式（包括跨步）
自动 swizzle（避开 bank conflict）
异步 → 计算和搬运重叠

5.2 WGMMA：Warpgroup MMA

WGMMA 是 H100 的新 Tensor Core 指令——一个**warpgroup（4 个 warp = 128 threads）**协作做矩阵乘。

传统做法：

一个 warp 32 个 thread 协作做 16×8×16 的矩阵乘 (mma.sync.aligned)
WGMMA 让 4 个 warp 一起做 64×N×16

优势：

单条指令计算量更大 → 指令开销摊薄
输入 A 可以直接来自 DRAM（异步）→ 不需要先搬到 shared memory
高吞吐：H100 上 WGMMA bf16 可以达 1000+ TFLOPS

5.3 用 CUTLASS 写 H100 kernel

直接写 PTX 太底层——推荐用 NVIDIA CUTLASS：

#include <cutlass/gemm/device/gemm.h>

using Gemm = cutlass::gemm::device::Gemm<
    cutlass::bfloat16_t,                   // A type
    cutlass::layout::RowMajor,
    cutlass::bfloat16_t,                   // B type
    cutlass::layout::ColumnMajor,
    cutlass::bfloat16_t,                   // C type
    cutlass::layout::RowMajor,
    float,                                  // accumulator
    cutlass::arch::OpClassTensorOp,
    cutlass::arch::Sm90                    // Hopper
>;

Gemm gemm;
gemm({M, N, K, A, lda, B, ldb, C, ldc, alpha, beta});

CUTLASS 自动用 TMA + WGMMA + 各种 swizzle / pipeline。

6. Kernel fusion 的几种范式

6.1 为什么 fuse

每个 kernel 启动都有：

~10 µs 启动开销
DRAM ↔ register 一次往返

把多个小 kernel 融合成一个大 kernel：

中间结果留在 register / shared memory（不写回 DRAM）
启动开销从 N × 10µs 降到 1 × 10µs

6.2 三种 fusion 范式

1. Epilogue fusion

GEMM 后接 bias add + ReLU：

GEMM → bias → ReLU

不要分 3 个 kernel——把 bias 和 ReLU 合在 GEMM 的 epilogue 里：

template<typename Epilogue>
__global__ void gemm_with_epilogue(...) {
    float c = ... compute matmul ...
    c = Epilogue::apply(c);  // bias + relu in registers
    C[idx] = c;
}

2. Producer-Consumer fusion

Softmax 内部需要先求 max，再求 sum，再 divide：

naïve: max kernel → sum kernel → divide kernel  (3 kernels)
fused: 一个 kernel 用 online algorithm 一次扫

Online softmax 是经典案例（FlashAttention 用的就是这个）。

3. Vertical fusion

把”前一层输出 → 后一层输入”的链路接起来：

attention(q,k,v) → linear_proj → layernorm → next_layer_attention

通常只能在低层（CUDA 库或 Triton kernel）做——PyTorch eager mode 难以做到。torch.compile 的 Inductor 后端会自动做这种 fusion。

6.3 fusion 的代价

可读性下降：一个 kernel 几百行
register pressure：fused kernel 用更多 register → occupancy 可能下降
不可复用：每个上下游组合都要写一份

工程经验：只 fuse 性能 hot path 的 top 5 kernel——剩下的不值得。

7. 朴素 GEMM → cuBLAS 级的 7 步演进

把上面所有概念串成实战路径——以 4096×4096×4096 fp16 GEMM 为例，每一步的预期性能（A100 上）：

步	优化点	性能 (TFLOPS)	相对峰值
0	朴素 3 层循环	~0.5	<0.5%
1	+ 全局 coalesced 访存	~3	1%
2	+ Shared memory tiling (32×32)	~30	10%
3	+ 每 thread 算 8×8 register tile	~100	30%
4	+ Bank conflict 消除（padding 或 swizzle）	~150	45%
5	+ Double buffering	~190	60%
6	+ 用 mma.sync (Tensor Core)	~250	80%
7	+ 用 CUTLASS / cuBLAS（自动 TMA / WGMMA）	~310	95%+

每一步的性能提升不是”魔法”——它和我们前面讲的概念一一对应。这就是为什么把这一章命名为”深入”：性能 = 一步步把硬件特性榨出来的累积效果。

7.1 一个真实经验

很多团队 self-implementation 的 GEMM 到第 4 步就停了——觉得 45% 已经”很不错”。但 cuBLAS 在 95% 才停——剩下的 50% 性能差就是 TMA / WGMMA / 各种 swizzle / 多级 pipeline 的累积。

所以工程经验是：

✅ 教学和短期验证用 self-impl（学得透）
✅ 生产直接用 cuBLAS / CUTLASS
✅ 特殊算子（FlashAttention 这种）才值得自己写

🎯 自我检验清单

Memory coalescing 的 128 字节边界从哪来？为什么相邻 thread 访问相邻地址才能 coalesce？
Shared memory bank conflict 在什么访问模式下发生？padding 和 swizzle 各自怎么避开它？
解释为什么 low occupancy + high ILP 有时比 high occupancy 更快——这件事在什么类型的 kernel 上特别明显？
Hopper TMA 和传统 thread-cooperative load 的核心区别是什么？为什么 TMA 让 SM 算力被释放？
朴素 GEMM 到 cuBLAS 级的 7 步演进，每一步对应本章哪个概念？

📚 参考资料

CUDA C++ Programming Guide：docs.nvidia.com/cuda/cuda-c-programming-guide/
CUTLASS GitHub：github.com/NVIDIA/cutlass
Hopper TMA Programming Guide：docs.nvidia.com/cuda/parallel-thread-execution/
Simon Boehm GEMM 7-step blog：siboehm.com/articles/22/CUDA-MMM
Maxas / mma.sync 文档：docs.nvidia.com/cuda/parallel-thread-execution/#warp-level-matrix-instructions
本系列模块二：CUDA 编程与算子优化（每个算子的详细教程）

下一章预告：Ch7 把视角从 CUDA 拉回到 PyTorch 框架层——torch.compile / TorchInductor / CUDA Graph / AMP / Profiler——给一份完整的 PyTorch 性能调优决策树。

搜索