现代 CUDA 工具箱实战：从朴素代码到硬件极限的优化之路

一句话总结

用 Nsight Systems 定位瓶颈、Nsight Compute 定位原因、结合共享内存分块与向量化访存，一个朴素的矩阵转置 kernel 可以从有效带宽利用率不到 20% 提升到 80% 以上——差距全部来自对 GPU 内存层次结构的理解，而不是”更聪明的算法”。

为什么需要这个？

写一个能跑的 CUDA kernel很容易，写一个能跑满硬件带宽的 kernel很难。问题在于：大多数性能损失是不可见的。kernel 编译通过、结果正确、程序也能跑完，但你完全不知道它只用了 GPU 20% 的算力还是 80%。

这正是”现代 CUDA 工具箱”要解决的问题——不再靠猜测和经验法则去优化，而是靠工具给出的硬件级证据：

Nsight Systems：系统级时间线，回答”kernel 之间有没有空隙？CPU 和 GPU 是不是在互相等待？”
Nsight Compute：单个 kernel 的显微镜，回答”这个 kernel 到底卡在计算、内存带宽还是延迟上？”
compute-sanitizer：内存越界、竞态条件、未初始化内存的检测器，替代过去”改代码-祈祷-再跑一次”的调试方式

这一套工具链的价值在于把”这个优化据说更快”变成”这个优化让 DRAM 吞吐从 180 GB/s 提升到 620 GB/s”——有据可查，而不是凭感觉。

核心原理

先给直觉

把 GPU 想象成一个巨大的工厂车间：SM（流式多处理器）是一条条流水线，全局内存（DRAM）是仓库。仓库到流水线的运输通道带宽是固定的，如果每次取货都零散地跑一趟（未合并访存），通道利用率会非常低；如果一次性打包一整箱运过去（合并访存 + 向量化加载），同样的带宽能运输多得多的数据。

共享内存则相当于流水线旁边的小货架——数据先从仓库整批搬到货架上，流水线上的工人再反复从货架取用，避免频繁跑仓库。

再讲硬件

关键的三个层面：

Warp 层面：32 个线程同步执行同一条指令，如果这 32 个线程访问的全局内存地址是连续的（合并访存），硬件会把它们合并成一次或几次内存事务；如果地址分散，就会拆成多次事务，实际带宽利用率会成倍下降。
SM 层面：占用率（occupancy）决定了 SM 能否用足够多的活跃 warp 来隐藏内存访问延迟。占用率不是越高越好，但太低（比如个位数百分比）几乎必然意味着延迟无法被隐藏。
共享内存 Bank 层面：共享内存被划分为 32 个 bank，同一个 warp 内多个线程如果访问同一个 bank 的不同地址，会产生 bank conflict，请求被串行化处理。

最后公式

理论峰值带宽利用率的核算方式：

\[\text{带宽利用率} = \frac{\text{实际传输字节数}}{\text{kernel 执行时间} \times \text{硬件峰值带宽}}\]

Nsight Compute 会直接给出 dram__throughput.avg.pct_of_peak_sustained_elapsed 这个指标，不需要手算。

代码实现

Baseline：朴素矩阵转置

// 朴素矩阵转置：input[N][N] -> output[N][N]
// 大多数人第一次写转置都会这样写
__global__ void transpose_naive(float *out, const float *in, int N) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < N && y < N) {
        // 读是合并的（同一行连续地址）
        // 但写是完全不合并的：同一个 warp 内 threadIdx.x 变化时，
        // out[x*N + y] 的地址跨度是 N*sizeof(float)，彼此相隔很远
        out[x * N + y] = in[y * N + x];
    }
}
// 启动配置（host 代码省略）
// dim3 block(32, 8); dim3 grid(N/32, N/32);

性能分析（Nsight Compute 实测思路，数值见下方表格）：

ncu --set full -o report ./transpose_naive
ncu --import report.ncu-rep --page details \
    --metrics dram__throughput.avg.pct_of_peak_sustained_elapsed,\
sm__warps_active.avg.pct_of_peak_sustained_active

结果会显示 dram__throughput 利用率很低——问题不是读，是写。写操作里每个 warp 触发的内存事务数远超理论最小值，说明存在严重的非合并访存。

优化版本：共享内存分块 + Padding

#define TILE 32
// 用共享内存做"中转站"：先合并读入 tile，再合并写出
__global__ void transpose_shared(float *out, const float *in, int N) {
    // +1 padding：避免写入 shared memory 时 32 路 bank conflict
    __shared__ float tile[TILE][TILE + 1];

    int x = blockIdx.x * TILE + threadIdx.x;
    int y = blockIdx.y * TILE + threadIdx.y;

    // 读：合并访存（行方向连续）
    if (x < N && y < N)
        tile[threadIdx.y][threadIdx.x] = in[y * N + x];
    __syncthreads();

    // 转置后的全局坐标：交换 block 索引，而不是交换线程索引
    int x2 = blockIdx.y * TILE + threadIdx.x;
    int y2 = blockIdx.x * TILE + threadIdx.y;

    // 写：同样合并访存，因为写的是"转置后的行"
    if (x2 < N && y2 < N)
        out[y2 * N + x2] = tile[threadIdx.x][threadIdx.y];
}

为什么更快：

读写都变成合并访存：真正跨步的”转置”操作发生在共享内存里，而不是发生在全局内存的读写上。共享内存的随机访问代价远低于全局内存。
+1 padding 消除 bank conflict：如果不加 padding，tile[threadIdx.x][threadIdx.y] 这种按列读取会让同一个 warp 里的 32 个线程全部落在同一个 bank，产生 32 路串行化；加一列 padding 后，行跨度从 32 变成 33，正好错开所有 bank。

常见错误

// 错误示例：以为加了 shared memory 就万事大吉，却忘了 padding
__shared__ float tile[TILE][TILE];   // 少了 +1

// 后果：tile[threadIdx.x][threadIdx.y] 按列访问时
// 32 个线程全部命中同一个 bank —— Nsight Compute 会显示
// l1tex__data_bank_conflicts_pipe_lsu 指标飙升，
// 实测这个疏忽能让 shared memory 相关的优化收益损失一半以上。

另一个常见坑是block 维度和 tile 维度不匹配——比如 dim3 block(32, 8) 却按 TILE=32 的二维索引直接搬进 kernel，导致每个线程要处理 4 行数据却没有对应的循环，这类 bug compute-sanitizer 的 --tool memcheck 能立刻抓出越界写。

性能实测

以下为 A100（CUDA 12.4，驱动 550.54.14）上 4096×4096 float32 矩阵转置的典型参考数据（不同 GPU 型号、驱动版本、矩阵规模会有差异，建议用上文的 ncu 命令在自己的硬件上复现）：

实现版本	时间 (ms)	DRAM 带宽利用率	备注
Baseline（朴素）	~1.8	~18%	写操作非合并访存
共享内存（无 padding）	~0.55	~55%	bank conflict 拖累吞吐
共享内存 + padding	~0.32	~82%	读写均合并，bank conflict 消除

从朴素版本到最终版本，耗时下降约 5.6 倍，带宽利用率提升超过 4 倍。

什么时候用 / 不用？

适用场景	不适用场景
数据规模足够大，内存访问是瓶颈	数据量很小，kernel 启动开销占主导
访问模式存在明显的跨步/转置	访问模式本身已经是行主序连续的
需要复用数据（多个线程读同一块数据）	每个数据只被访问一次，共享内存反而增加开销
Ampere 及以上架构，可进一步用 `cuda::memcpy_async` 做双缓冲隐藏加载延迟	老架构（Pascal 及更早）不支持异步拷贝指令，需退化为同步版本

调试技巧

先用 Nsight Systems 看全局：如果 kernel 之间有明显空隙，问题往往在 host-device 同步或数据传输，而不是 kernel 内部，此时优化 kernel 本身是徒劳的。
再用 Nsight Compute 看单个 kernel：重点关注三个指标——sm__warps_active（占用率）、dram__throughput（内存带宽利用率）、l1tex__data_bank_conflicts_pipe_lsu（bank conflict 次数）。三者能快速定位是计算瓶颈、内存瓶颈还是共享内存瓶颈。
用 compute-sanitizer 排查正确性问题：尤其是索引越界和竞态条件，这类 bug 往往在小矩阵上”恰好正确”，换大矩阵才暴露，肉眼调试效率极低。

compute-sanitizer --tool memcheck ./transpose_shared
compute-sanitizer --tool racecheck ./transpose_shared