Lesson 07 · CUDA 性能调优 · 把 roofline 钉到你的 A100 与 3060 上
前六节你在纸面上建好了调优框架。现在你手上有两块能跑的卡—— 是时候把抽象的 roofline 钉到真实硬件上了。本节你将亲手测出自己两块卡的屋顶, 并发现一件反直觉的事:同一个 kernel,在 A100 上是计算瓶颈,在 3060 上却变成访存瓶颈。 这是消费卡和数据中心卡最值钱的一条直觉。
第 1 节我们用了 A100 的「官方」数字。但真机调优有一条铁律:规格表是营销数字,不是你能达到的数字。
你的 3060 是 笔记本版(RTX 3060 Laptop GPU, 6GB)。它的 FP32 峰值不是固定值—— 笔记本厂商把功耗墙(TGP)配在 60W 到 130W 之间,同名的卡实际性能能差 20–30%。[3] 你抄到的「12.7 TFLOP/s」可能根本不是你这台的数字。
推论:roofline 的两道屋顶,都该测,不该抄。带宽用微基准测,算力用峰值微基准测。 规格表只用来对照「我测出来的离理论上限差多远」。
访存屋顶的斜率就是显存带宽。CUDA 自带一个现成的微基准 bandwidthTest,
它在显存里搬一大块数据,直接报告你能达到的 device-to-device 带宽。先在两块卡上都跑一遍。
# CUDA 11.7 自带的编译好版本,路径通常在:
/usr/local/cuda/extras/demo_suite/bandwidthTest --dtod
# 若不在,从 cuda-samples 编译 bandwidthTest 后:
./bandwidthTest --memory=pinned --mode=shmooDevice to Device Bandwidth 那一行——那才是 roofline 斜屋顶的真实斜率。
忽略 H2D/D2H(那是 PCIe 带宽,不是显存带宽)。
若找不到 bandwidthTest,用这个拷贝核自测(c[i]=a[i]:读 4B + 写 4B = 每元素搬 8 字节):
// bw.cu — 编译: nvcc -O3 -arch=sm_80 bw.cu -o bw (3060 改 sm_86!)
__global__ void copy(float* c, const float* a, size_t n){
size_t i = blockIdx.x*blockDim.x + threadIdx.x;
if(i < n) c[i] = a[i];
}
// 主机端要点:N=64M floats(远大于 L2);用 cudaEvent 计时;
// 带宽(GB/s) = 8*N(字节) / 时间(s) / 1e9 ;预热一次再正式测十次取中位数把实测带宽填进脊点公式(回忆第 1 节):脊点 = 峰值算力 ÷ 带宽。下面是规格表估值, 等你跑完微基准,用实测值替换——这才是你这两台机器的真屋顶。
† 笔记本 3060 算力随 TGP 浮动,务必实测替换。[2][3]
注意这个差异的方向:3060 的脊点(≈35)远在 A100(≈12.6)右边。 原因不是 3060 算力强,而是它的带宽相对算力太弱: A100 的带宽是 3060 的 ~5 倍,算力只有 ~1.7 倍。脊点 = 算力/带宽,带宽掉得更狠,脊点就被推向右。
本节核心洞察。设想一个算术强度 AI = 20 的 kernel(比如一个复用度中等的小 tile GEMM):
| 卡 | 脊点 | AI=20 落在哪 | 瓶颈 | 该往哪优化 |
|---|---|---|---|---|
| A100 | 12.6 | 20 > 12.6,右侧 | 计算瓶颈 | 提计算(tensor core、ILP) |
| 3060 | 34.5 | 20 < 34.5,左侧 | 访存瓶颈 | 提访存(合并、shared) |
这条直觉直接服务你的 mission:在 3060 上日常迭代会更早撞访存墙,逼你把访存优化(第 3–5 节)做扎实; 搬到 A100 跑,同一 kernel 可能翻成计算瓶颈,这时第 6 节往后的计算侧优化才登场。 两块卡正好覆盖 roofline 的两侧,是绝佳的双视角训练场。
不要看上文,凭记忆判断。用本节脊点:A100=12.6,3060=34.5。
bandwidthTest 找不到?ncu 在 WSL 报权限错?
实测带宽只有理论的 70% 正常吗?算力峰值怎么单独测?随时追问,把不清楚的挖透。
📘 Nsight Compute Profiling Guide — Roofline / Memory Workload ——本节把 roofline 落到真机,这份指南讲 ncu 如何自动画出你 kernel 在 roofline 上的点。 读「GPU Speed of Light Roofline Chart」一节,正好对接你刚手算的脊点。