필사 모드: 亲手改一个 GPU kernel 到底意味着什么 —— 把一个转置 kernel 提速 5 倍

• 引言 —— 当分析器指向的 kernel 不在任何库里
• kernel 是什么,以及为什么会去改它
• 执行模型 —— 线程、warp、block、grid
• 内存层级 —— 合并访问与存储体冲突
  • 合并访问(coalescing)
  • 存储体冲突(bank conflict)
• 占用率与 roofline —— 该看什么,该忽略什么
  • 占用率是症状,不是目标
  • roofline —— 大多数 kernel 受限于带宽
• 实操 —— 把转置 kernel 分四个阶段改好
  • 完整代码
  • 测量方法中要注意的地方
  • 结果的形态
  • 常见的踩坑方式
• 用 Nsight Compute 看什么
  • 什么时候该收手
• 结语 —— 改 kernel,改的是数据移动的顺序
• 参考资料

引言 —— 当分析器指向的 kernel 不在任何库里

一个训练 step 耗时 420ms，其中 170ms 花在一个不知名的 kernel 上——遇到这种情况，选择一下子就变窄了。如果是 cuBLAS 调用，你无从下手；如果是 PyTorch 的算子，换个算子试试就行。可如果那个 kernel 是有人为了你模型里某个特殊的 mask 逻辑而临时写出来的，能改它的就只有你自己了。

本文就从这一点讲起：亲手"改"一个 GPU kernel 到底是怎样的工作，需要知道什么、需要测什么。先说结论：kernel 优化里八成的工作不是削减计算量，而是改变搬运内存的顺序。本文实操部分要改的那个 kernel,连一个算术运算都没有,但依然能提速 5 倍以上。

参考环境是 CUDA Toolkit 13.3 Update 1、Nsight Compute 2026.2 系列。概念不随硬件世代过时,但工具的 flag 名称会变,命令行请对照你自己版本的 Nsight Compute CLI 文档。

kernel 是什么,以及为什么会去改它

kernel 是一个在 GPU 上运行的函数。和 CPU 函数不同的是,一次调用会同时启动数万个实例。这一个个实例就是线程(thread),线程通过内置变量得知自己是第几个,只处理与之对应的那份数据。

实务中值得亲自动手改 kernel 的场景,比想象中要窄。按顺序排查大致如下。

关键在最后两行。能用库替代的话,替代永远更划算。库里的 kernel 所积累的调优,远超你自己能投入的时间,而且新架构一出,库会替你更新。自己写的 kernel,则要你自己养它一辈子。

执行模型 —— 线程、warp、block、grid

要改 kernel,就得知道硬件是怎么把线程组织起来的。层级分四级。

grid           一次 kernel 调用的全部线程集合
  └ block      放在同一个 SM 上,共享 shared memory 与 __syncthreads()
      └ warp   32 个线程。调度和指令发出的真正单位
          └ thread   拥有自己寄存器的一条执行流

这里实务上最重要的一点是:warp 才是真正的单位。程序员按线程为单位写代码,但硬件是把 32 个线程打包成一束一起发出的。由此带来两个后果。

第一,分支发散(divergence)。如果一个 warp 内的线程走了不同的分支,硬件会依次执行两条路径,把不属于当前路径的线程禁用掉。在 warp 内部分叉的条件语句会增加执行时间;沿 warp 边界分叉的条件语句则是免费的。

第二,内存访问也是按 warp 为单位合并的。这就是 coalescing(合并访问),下一节的主题。

确定 block 大小时,默认做法是取 32 的倍数。如果 block 大小是 33,硬件会发出两个 warp,其中一个 warp 里有 31 个 lane 会闲置。

内存层级 —— 合并访问与存储体冲突

GPU 内存是分层的,每一层的延迟和带宽都相差数量级。

合并访问(coalescing)

全局内存访问是按 32 字节为单位的事务(transaction)来处理的。如果一个 warp 里的 32 个线程读取连续的 32 个 float,即 128 字节,只需 4 次事务就能完成。反过来,如果同样这 32 个线程读取相隔 4096 的 32 个 float,就需要 32 次事务,每次都取回 32 字节却只用掉 4 字节,其余全部丢弃——相当于只用到了带宽的八分之一。

这是造成 kernel 性能差异的最大单一因素。在实操环节你会亲眼看到恰好 8 倍的浪费。

存储体冲突(bank conflict)

shared memory 按 32 个 bank 交织排布。以 4 字节为一个 word 计,地址除以 4 再对 32 取余就是 bank 编号。一个 warp 里的线程若各自命中不同的 bank,一个周期就能处理完;若命中同一个 bank 里的不同地址,就会按命中次数被串行化。

典型的翻车场景是对一个正方形 shared 数组做按列访问。在 tile[32][32] 里让 tile[i][0] 随 i 遍历,所有元素都会落在 bank 0 上——这是 32-way 冲突。解法是把数组多开一列。声明成 tile[32][33] 后,每一行的起始 bank 都会错开一格,列访问就均匀分散到全部 32 个 bank 上。这是用每个 block 多花 32 * 4 字节,也就是 128 字节的 shared memory,换取串行化的消除。

占用率与 roofline —— 该看什么,该忽略什么

占用率是症状,不是目标

占用率(occupancy)是 SM 实际驻留的 warp 数,除以它能驻留的最大 warp 数。初学者常见的误解,是把这个数值当成要最大化的目标——其实不是。

占用率只做一件事:掩盖延迟。当某个 warp 在等待 HBM 的响应时,如果能执行另一个 warp,SM 就不会闲置。所以占用率是"是否有足够多的 warp 来掩盖延迟"这个问题的代理指标,它本身并不等于性能。

有两种情况下,更低的占用率反而更快。

第一,单线程用大量寄存器换取指令级并行度。如果一个线程同时发出四个独立的 load,即便 warp 数量只有四分之一,飞行中的内存请求总量也是一样的。Vasily Volkov 的 Better Performance at Lower Occupancy 正面论证了这一点。虽然是 2010 年的资料,但论点至今依然成立。

第二,kernel 本身已经贴到了带宽上限。内存管道已经饱和时,再塞更多 warp 也无处可去。

所以实务规则是:只在占用率低的时候才去关注它。如果跌到 25% 以下且 kernel 受限于延迟,就该怀疑寄存器用量或 shared memory 分配。如果占用率有 60% 但 kernel 依然慢,占用率就不是元凶,得看别处。

roofline —— 大多数 kernel 受限于带宽

决定该看哪里的工具是 roofline。它的横轴是算术强度(arithmetic intensity),即每搬运一字节所执行的运算次数。

实际性能(FLOP/s)
   ^
   |            ______________  计算上限
   |           /
   |          /   斜率 = 内存带宽
   |         /
   +--------+-------------------> 算术强度 (FLOP/Byte)
          拐点(ridge point)

拐点 = (计算上限 FLOP/s) / (内存带宽 Byte/s)

拐点是硬件的固有属性。在最新的数据中心 GPU 上,这个值在几十到几百 FLOP/Byte 的区间。但如果实际算一算我们常用 kernel 的算术强度,大多数都是个位数。

读法很简单:如果算术强度远低于拐点,这个 kernel 的性能上限就已经定死了。再怎么削减计算指令都无济于事,唯一有意义的改进是减少搬运的字节数,或者改进搬运的方式。

所以 kernel 优化的第一个问题永远是:"这个 kernel 的理论最小流量是多少字节,现在实际搬运了多少字节。"

实操 —— 把转置 kernel 分四个阶段改好

现在动手改一个。目标是 4096 x 4096 的 float 矩阵转置。计算次数为 0,所以只剩内存这一个变量,是个理想的教学案例。

理论最小流量很明确。读一次、写一次,即 2 * 4096 * 4096 * 4 字节,约 134MB。不可能比这更少。因此性能上限就是"不转置、只做拷贝的 kernel",我们先测这个,作为基准线。

完整代码

// transpose.cu
// 编译: nvcc -O3 -arch=sm_80 transpose.cu -o transpose
#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>

static const int TILE = 32;
static const int BLOCK_ROWS = 8;   // 每个 block 32x8 = 256 个线程
static const int N = 4096;

#define CHECK(x) do { cudaError_t e_ = (x); if (e_ != cudaSuccess) { \
    printf("CUDA error: %s (line %d)\n", cudaGetErrorString(e_), __LINE__); \
    exit(1); } } while (0)

// 阶段 0。上限线: 不转置,只做拷贝。
__global__ void copyKernel(float *out, const float *in) {
  int x = blockIdx.x * TILE + threadIdx.x;
  int y = blockIdx.y * TILE + threadIdx.y;
  for (int j = 0; j < TILE; j += BLOCK_ROWS)
    out[(y + j) * N + x] = in[(y + j) * N + x];
}

// 阶段 1。naive: 读是合并访问的,但写是间隔 N 的跨步访问。
__global__ void transposeNaive(float *out, const float *in) {
  int x = blockIdx.x * TILE + threadIdx.x;
  int y = blockIdx.y * TILE + threadIdx.y;
  for (int j = 0; j < TILE; j += BLOCK_ROWS)
    out[x * N + (y + j)] = in[(y + j) * N + x];
}

// 阶段 2。shared memory 分块: 在 SRAM 内部完成转置,
//        让全局内存的读和写都变成合并访问。
__global__ void transposeShared(float *out, const float *in) {
  __shared__ float tile[TILE][TILE];

  int x = blockIdx.x * TILE + threadIdx.x;
  int y = blockIdx.y * TILE + threadIdx.y;
  for (int j = 0; j < TILE; j += BLOCK_ROWS)
    tile[threadIdx.y + j][threadIdx.x] = in[(y + j) * N + x];

  __syncthreads();

  // 把 block 坐标互换,让写操作也落在连续地址上。
  x = blockIdx.y * TILE + threadIdx.x;
  y = blockIdx.x * TILE + threadIdx.y;
  for (int j = 0; j < TILE; j += BLOCK_ROWS)
    out[(y + j) * N + x] = tile[threadIdx.x][threadIdx.y + j];
}

// 阶段 3。用一列 padding 消除 shared memory 的存储体冲突。
__global__ void transposePadded(float *out, const float *in) {
  __shared__ float tile[TILE][TILE + 1];   // 唯一的区别

  int x = blockIdx.x * TILE + threadIdx.x;
  int y = blockIdx.y * TILE + threadIdx.y;
  for (int j = 0; j < TILE; j += BLOCK_ROWS)
    tile[threadIdx.y + j][threadIdx.x] = in[(y + j) * N + x];

  __syncthreads();

  x = blockIdx.y * TILE + threadIdx.x;
  y = blockIdx.x * TILE + threadIdx.y;
  for (int j = 0; j < TILE; j += BLOCK_ROWS)
    out[(y + j) * N + x] = tile[threadIdx.x][threadIdx.y + j];
}

typedef void (*Kern)(float *, const float *);

static void bench(const char *name, Kern k, float *d_out, const float *d_in,
                  const float *h_ref, float *h_out, bool checkTranspose) {
  dim3 grid(N / TILE, N / TILE), block(TILE, BLOCK_ROWS);
  const int WARMUP = 5, ITERS = 50;
  const double bytes = 2.0 * N * N * sizeof(float);

  for (int i = 0; i < WARMUP; i++) k<<<grid, block>>>(d_out, d_in);
  CHECK(cudaDeviceSynchronize());

  cudaEvent_t t0, t1;
  CHECK(cudaEventCreate(&t0));
  CHECK(cudaEventCreate(&t1));
  CHECK(cudaEventRecord(t0));
  for (int i = 0; i < ITERS; i++) k<<<grid, block>>>(d_out, d_in);
  CHECK(cudaEventRecord(t1));
  CHECK(cudaEventSynchronize(t1));

  float ms = 0.f;
  CHECK(cudaEventElapsedTime(&ms, t0, t1));
  double perIter = ms / ITERS;
  double gbs = bytes / (perIter * 1.0e-3) / 1.0e9;

  // 没有正确性验证的性能数字毫无意义。
  CHECK(cudaMemcpy(h_out, d_out, (size_t)N * N * sizeof(float),
                   cudaMemcpyDeviceToHost));
  long bad = 0;
  for (long r = 0; r < N && bad == 0; r++)
    for (long c = 0; c < N; c++) {
      float want = checkTranspose ? h_ref[c * N + r] : h_ref[r * N + c];
      if (h_out[r * N + c] != want) { bad++; break; }
    }

  printf("%-18s %8.3f ms   %8.1f GB/s   %s\n", name, perIter, gbs,
         bad ? "FAIL" : "ok");
  CHECK(cudaEventDestroy(t0));
  CHECK(cudaEventDestroy(t1));
}

int main() {
  size_t bytes = (size_t)N * N * sizeof(float);
  float *h_in = (float *)malloc(bytes), *h_out = (float *)malloc(bytes);
  for (long i = 0; i < (long)N * N; i++) h_in[i] = (float)(i % 1000);

  float *d_in, *d_out;
  CHECK(cudaMalloc(&d_in, bytes));
  CHECK(cudaMalloc(&d_out, bytes));
  CHECK(cudaMemcpy(d_in, h_in, bytes, cudaMemcpyHostToDevice));

  cudaDeviceProp p;
  CHECK(cudaGetDeviceProperties(&p, 0));
  printf("%s  peak HBM = %.1f GB/s\n\n", p.name,
         2.0 * p.memoryClockRate * (p.memoryBusWidth / 8) / 1.0e6);

  bench("copy (upper bound)", copyKernel,      d_out, d_in, h_in, h_out, false);
  bench("naive",              transposeNaive,  d_out, d_in, h_in, h_out, true);
  bench("shared tile",        transposeShared, d_out, d_in, h_in, h_out, true);
  bench("shared + padding",   transposePadded, d_out, d_in, h_in, h_out, true);

  cudaFree(d_in); cudaFree(d_out); free(h_in); free(h_out);
  return 0;
}

测量方法中要注意的地方

要让数字可信,首先测试工具本身要诚实。上面的代码遵守了五条规则。

• 丢弃 warmup。 第一次调用会混入 context 创建和模块加载的开销。
• 重复测量后取平均。 如果单个 kernel 耗时在 1ms 量级,光是时钟波动就能造成 10% 的抖动。
• 用 cudaEvent,而不是 CPU 计时器。 kernel 执行是异步的,CPU 时间只能测到"发出调用"所花的时间。
• 验证正确性。 索引写错了,结果往往反而更快。没有验证的 GB/s 只是数字游戏。
• 换算成有效带宽,而不是看时间。 绝对时间会随规模和硬件变化,但相对理论带宽的百分比在任何地方都可比。

有效带宽的公式很简单:必须搬运的最小字节数,除以实际花费的时间。这里"必须"二字很关键——用的是算法本身所需的字节数,而不是实际(被浪费地)搬运的字节数。只有这样,浪费才会体现在数字上。

结果的形态

下面是在 A100 80GB(sm_80)级别的机器上跑这套测试工具时得到的典型形态。绝对值会因硬件、驱动、时钟状态而有很大差异,请不要直接引用这些数字,而是在自己的 GPU 上跑一遍上面的代码,得出自己的基准线。 有意义的是各阶段之间的相对比例。

有三点值得读出来。

第一,naive 只有上限的五分之一左右。计算次数为 0,搬运的数据量也一样,却慢了 5 倍。唯一的差别是顺序。写操作按 N 的间隔分散开,每次 32 字节的事务只写入 4 字节,丢弃 28 字节。8 倍的浪费和其他效应混在一起,表现为 5 倍的差距。

第二,shared tile 恢复了大部分差距,但没有走到头。全局访问两边都修好了,但瓶颈转移到了 SRAM 内部。tile[threadIdx.x][threadIdx.y + j] 是按列访问,而在一个 32x32 的正方形数组里,列访问全部落在同一个 bank 上。

第三,最后一步的代码差异只是数组声明里的一个字符。把 [TILE] 改成 [TILE + 1] 就是全部改动。kernel 优化常常就是这个样子——不是改算法,而是把数据布局挪动一格。

常见的踩坑方式

这个实操里实际经常踩到的雷。

• 漏掉 __syncthreads()。 在填完 shared tile 到读取之间如果没有同步,结果会非确定性地出错。小规模输入下反而经常凑巧对,这更危险。
• 把 __syncthreads() 放在分支里。 放在 block 内只有部分线程能到达的位置,是未定义行为。
• 索引计算时忘了互换 block 坐标。 如果只加了 shared tile 却不改输出索引,写操作又变回跨步访问,第二阶段的收益就消失了。结果是对的,但没有变快——这是最难察觉的一种形态。
• 不加 -O3 就测。 主机代码优化缺失时,验证循环会主导测量时间,把结论完全颠倒。
• N 太小。 kernel 启动开销是几微秒级别的,如果总耗时只有几十微秒,你测到的其实是开销本身。

用 Nsight Compute 看什么

数字不好看时,告诉你原因的是分析器(profiler)。先看命令行。

# 收集全部 section。单个 kernel 可能耗时数百 ms,所以必须缩小目标范围。
ncu --set full \
    --kernel-name regex:transpose \
    --launch-skip 5 --launch-count 1 \
    -o transpose_report \
    ./transpose

# 只抓取用于定位原因的指标(快得多)
ncu --metrics \
  sm__throughput.avg.pct_of_peak_sustained_elapsed,\
gpu__dram_throughput.avg.pct_of_peak_sustained_elapsed,\
l1tex__data_bank_conflicts_pipe_lsu_mem_shared.sum,\
l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio \
  --kernel-name regex:transpose --launch-count 1 ./transpose

# 用 GUI 打开
ncu-ui transpose_report.ncu-rep

用 --launch-skip 跳过 warmup 执行很关键。分析器如实展示的是第一次执行时缓存冷启动的状态,不跳过的话,分析的就不是稳态。

打开报告后会看到好几个 section,实务上有一个固定的查看顺序。

1. Speed of Light。 把计算吞吐量和内存吞吐量分别显示为占硬件理论峰值的百分比。方向在这里就能定下来。如果内存达到 80% 以上,说明贴住了带宽;如果两者都低于 30%,那就是延迟或占用率的问题。我们这个转置 kernel 的 naive 版本在这个画面里两者都偏低——不是因为管道饱和,而是因为在浪费。

2. Memory Workload Analysis。 这是核心所在。它显示每个请求实际取回了多少个 sector;一次完美合并访问的 32 线程 float load,每个请求应为 4 个 sector。naive kernel 的写操作每个请求会显示 32 个 sector。8 倍的浪费原原本本地体现在这一行里。仅凭这一个指标就能直接回答"是不是合并访问的问题"。

3. Shared Memory 相关指标。 会显示存储体冲突的次数。第二阶段的 kernel 在这里会明显飙高,第三阶段则接近于 0。这是确认 padding 是否真的起作用的地方。

4. Warp State Statistics。 按类别显示 warp 停滞的原因。如果 Stall Long Scoreboard 占绝对主导,说明在等待全局内存的响应;如果是 Stall MIO Throttle,说明是 shared memory 或特殊功能单元那边拥堵。它能帮你分清原因出在内存还是指令。

5. Occupancy。 放到最后看。前四项都干净,却依然慢的时候,才检查是不是驻留 warp 数量不够。不按这个顺序、先看占用率的话,大多会导向错误的优化方向。

Nsight Compute 还有一个 roofline section,用图形展示你的 kernel 是落在斜坡上还是平地上。转置 kernel 的算术强度为 0,所以会打在最左端,光凭这一点就能得出"计算层面没有可优化的空间"这一结论。

什么时候该收手

如果转置 kernel 已经提到了上限的 93%,那么再去追剩下的 7% 几乎没有意义。最好提前定好判断标准。

• 超过理论最小流量的 90% 就收手。 在 memory-bound 的 kernel 上,再往上主要是 tile 边界和 TLB 效应,投入与回报会急剧恶化。
• 重新测量这个 kernel 在总执行时间中的占比。 如果把 170ms 降到了 40ms,现在瓶颈就在别处了。Amdahl 定律在 kernel 优化里同样成立。
• 把维护成本算进去。 手写的 kernel,每次新架构发布都要重新验证一遍。今天比库快 20% 的 kernel,两年后可能反而慢了 30%。
• 先看上一层能不能解决。 如果下一篇要讲的 Triton 版本用 20 行就能写出同一个 kernel,还能拿到相近的性能,那维护这份 CUDA C++ 版本的理由就会减少。

结语 —— 改 kernel,改的是数据移动的顺序

本文实操部分从头到尾没有出现过一次浮点运算,可第一版和最后一版之间依然差了 5 倍。变化的只有一件事:同样的数据,以什么顺序读取、暂存在哪里、以什么顺序写回。

这个事实决定了 GPU kernel 工作的本质。削减计算量的算法层面改进,大多要么库里已经做过了,要么在我们的问题里根本无法更改。留给我们的杠杆,是内存层级内部的数据布局与移动顺序——而幸运的是,那才是更大的那根杠杆。

把工作顺序浓缩成一句话就是:算出理论最小流量,用测试工具测出现在用到了百分之多少,用 Nsight Compute 定位浪费出现在哪里,修改布局,再测一遍。凭感觉改完就说"变快了",相当于在这套流程里漏掉了两次测量,这样得出的结论,换一台机器就会被推翻。

参考资料

• CUDA C++ Programming Guide: https://docs.nvidia.com/cuda/cuda-c-programming-guide/
• CUDA C++ Best Practices Guide(含合并访问、存储体冲突章节): https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/
• An Efficient Matrix Transpose in CUDA C/C++(本实操的原始出处): https://developer.nvidia.com/blog/efficient-matrix-transpose-cuda-cc/
• Nsight Compute CLI 文档: https://docs.nvidia.com/nsight-compute/NsightComputeCli/index.html
• Nsight Compute kernel 性能分析指南(各 section 说明): https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html
• Volkov, Better Performance at Lower Occupancy (GTC 2010): https://www.nvidia.com/content/gtc-2010/pdfs/2238_gtc2010.pdf
• Williams et al., Roofline: An Insightful Visual Performance Model: https://dl.acm.org/doi/10.1145/1498765.1498785