CUDA / HPC 面试复习：从执行模型到性能分析

这块我以前最容易答成名词表。block、warp、shared memory、occupancy 都知道，但一旦面试官追一句“所以你平时怎么判断程序慢在哪”，就开始散。

后来发现还是得先把顺序定下来：线程怎么铺，数据怎么走，最后瓶颈怎么看。

先说执行模型。

CUDA 这套东西如果只背 grid / block / thread 其实不太够，还是要把“哪些东西能同步，哪些不能”一起记。block 内线程能同步，也能共享 shared memory，出了 block 默认就别想当然地认为大家能随便对齐进度。这个约束会直接影响 kernel 怎么拆。

warp 也一样，平时很容易把它当定义题记过去，但真正在性能上常出事的是 divergence。同一个 warp 里线程分支走乱了，硬件就只能把不同路径拆开跑，这时候“线程很多”也没什么用。

然后是内存。

很多 kernel 跑不快，本质上不是算得不够多，而是数据喂不进去。这个时候再去看 register、shared memory、global memory 这些层次就比较自然了。register 最快，但预算很紧；shared memory 适合块内复用；global memory 最大，但延迟也高，很多时候瓶颈就卡在这里。

我现在会特别提醒自己别把 occupancy 当成万能指标。它有用，但它说明的是一个 SM 上活跃 warp 的多少，不是性能本身。occupancy 很高，照样可能因为 global memory 没合并、访存模式太差、或者指令依赖太重而跑不快。之前我老盯这个数，后来发现很容易被它带偏。

再往后就是 profiler。

比起一上来调 block size，我现在更愿意先看热点到底在哪。先确认是不是这个 kernel 占了主要时间，再看它更像算力瓶颈还是带宽瓶颈。nsys 看整条时间线比较方便，CPU 和 GPU 有没有串住、流有没有并起来，这种东西它更直观。ncu 更像拿来盯一个 kernel 细看，访存、占用率、指令统计这些都会更清楚。

如果面试里被问“你一般怎么优化一个 kernel”，我觉得最差的答法就是直接开始报术语，什么 unroll、向量化、调块大小。更像样一点的顺序应该还是先判断瓶颈，再决定是改访存、改并行划分，还是改 launch 配置。

有几个点我感觉很容易被追问：

• 为什么 occupancy 高不一定快。
• shared memory 为什么不是总能带来收益。
• branch divergence 到底会把什么东西拖慢。

最后还是拿矩阵乘法举例最顺。这个题之所以老被拿来讲，不是因为它有多基础，而是它正好能把“共享内存到底解决了什么”说清楚。

矩阵乘法里同一块输入数据会被很多线程重复使用，所以它很适合把子块搬进 shared memory：

__shared__ float tileA[32][32];
__shared__ float tileB[32][32];

这里真正该记住的不是“shared memory 很快”，而是原本重复的 global memory 访问能不能变成一次加载、多次复用。只要把这个点想明白，后面再讲 tile 大小、寄存器压力、bank conflict 这些细节，就不会像在背答案。