Roofline 与瓶颈判断 · 速查表
CUDA 性能调优 · 快速参考 · 配合 Lesson 01
访存合并(Coalescing)速记
事务(transaction): 显存按固定大小块搬运的最小单位。A100 上 1 个 sector = 32B,常 4 个凑成 128B 一起处理。不管你要几字节,都按整块搬——非合并 = 开很多箱却每箱只用一点。
规则: warp(32 线程)的访存按 128B 对齐事务进行。让 warp 内相邻线程访问相邻地址,即 threadIdx.x 乘在数据的最低维上。
| 索引模式(warp 内相邻线程) | 地址差 | 是否合并 |
|---|
x[i], i=tid | 1 元素 | ✅ 合并 |
x[i*stride] | stride 元素 | ❌ 非合并 |
行主序 a[row*N+col], col=tid.x | 1 元素 | ✅ 合并 |
行主序 a[row*N+col], row=tid.x | N 元素 | ❌ 非合并 |
| 术语 | 含义 |
|---|
| request | 一个 warp 的一条访存指令 |
| sector | 32B 货箱;sectors/request 理想≈4,越大越不合并 |
- Nsight 信号:Memory Workload 段 "Uncoalesced global accesses";看 sectors per request(理想≈4,越大越糟)。
- 修复:转置数据布局、调整线程↔数据映射、或向量化(
float4)。
列访问(stride=N)正解: 不要硬改访问模式 —— 先用一次按行的合并访问把 tile 搬进 shared,再在片上按列读。shared 不怕 stride(只提防 bank conflict)。这正是 tiling 的动机。
对齐(alignment): 128B 事务只能从 128 的倍数地址起。数据再连续,若起点偏移(如 data+1、未对齐字段)→ 跨格 → 事务翻倍。cudaMalloc 保证 ≥256B 对齐;float4 天然对齐;必要时 __align__(128)。
| 缓存层级(A100) | 作用 |
|---|
| L1(每 SM,与 shared mem 共用) | 最快,容量小 |
| L2(全 GPU 共享,40MB) | 能摊销部分非合并/重复访问 |
| HBM 显存 | 真正的"远路",带宽即峰值带宽 |
- Memory% 高 + DRAM% 高 → 真打 HBM,缓存没兜住。
- Memory% 高 + DRAM% 低 → 多被 L2 接住,瓶颈在数据复用 → 考虑 shared memory。
- ⚠️ 别依赖缓存救场——能否命中不可控。仍应一开始就写合并访问。
Shared Memory 速记
何时用: 仅当数据被反复读(block 内多线程或单线程多次)。无复用(逐元素)→ 无益。有复用 → 全局读 1 次、shared 读 N 次,全局访存量降为 1/N,算术强度升高(roofline 右移)。
三步: ① __shared__ float tile[N]; ② 协作加载(合并访问)→ __syncthreads(); ③ 复用 tile。
| 层级(A100) | 延迟约 | 范围 |
|---|
| 寄存器 | ~1 cyc | 线程私有 |
| Shared(片上 SRAM,与 L1 共用) | ~20-30 cyc | block 内共享,手动管理 |
| L2 | ~200 cyc | 全 GPU,自动 |
| 全局 HBM | ~400-800 cyc | 全部,自动 |
- ⚠️
__syncthreads() 必须全 block 线程都到达;放进 if 分支只让部分线程执行 → 死锁。
- ⚠️ 写 shared 与读他人写的位置之间,必须有
__syncthreads(),否则读到垃圾值(且时隐时现)。
Bank Conflict 速记
32 个 bank: shared memory 分 32 条通道。连续 4B 字依次落 bank 0..31。bank = (地址/4) % 32。一个 warp 内若多线程访问不同地址但同一 bank → 串行化,N 路冲突慢 N 倍。
| warp 访问模式 | 结果 |
|---|
s[tid] 连续字 | 32 不同 bank → ✅ 1 拍 |
s[0] 全员同址 | broadcast → ✅ 不算冲突 |
s[tid*32] / 按列 tile[tid][k] | 全落 bank 0 → ❌ 32 路冲突 |
- 修复:padding ——
tile[32][33](列宽+1,与 32 互质,打乱 bank 对齐)。
- Nsight 信号:Memory Workload 段 "Shared Memory Bank Conflicts" 非零。
Tiled 矩阵乘要点
朴素 GEMM:A 行 B 列反复从全局读、零复用 → 访存瓶颈。Tiling:把 T×T 小块协作加载进 shared(合并)→ __syncthreads() → 在 shared 复用 T 次 → __syncthreads() → 下一块。全局访存降为约 1/T,算术强度升高,roofline 右移逼近峰值。padding As/Bs 防 bank conflict。
Occupancy 速记
定义: 活跃 warp 数 ÷ SM 最大 warp 数(A100:64 warp = 2048 线程/SM)。衡量调度器有多少候选 warp 可切换来藏延迟。
铁律: 高不一定快(ILP 也能藏延迟,提它有代价如 register spilling);但过低一定伤性能。→ 保下限、不强求上限。只在「延迟瓶颈+occupancy 低」时才提。
| 限制 occupancy 的资源 | 怎么卡住 |
|---|
| 寄存器/线程 | 太多 → 装不下更多线程 |
| shared mem/block | 太多 → 同驻 block 变少 |
| block 大小 | 太小 → 撞每 SM 最大 block 数 |
- block 大小取 32 倍数(常 128/256);grid ≥ SM 数。
cudaOccupancyMaxPotentialBlockSize() 让运行时推荐 block 大小。- 查寄存器用量:
nvcc -Xptxas -v。Nsight 看理论 vs 实测 occupancy(差距大=负载不均/尾部效应)。
核心公式
| 量 | 公式 | 单位 |
|---|
| 算术强度 AI | 总 FLOPs ÷ 总 Bytes | FLOP/Byte |
| 访存屋顶 | 带宽 × AI | FLOP/s |
| 计算屋顶 | 峰值算力(水平封顶) | FLOP/s |
| 脊点(ridge) | 峰值算力 ÷ 带宽 | FLOP/Byte |
判断法则:
把 kernel 的 AI 与脊点比较 ——
AI < 脊点 → 访存瓶颈(优化访存);
AI > 脊点 → 计算瓶颈(优化计算)。
瓶颈 → 优化方向
| 瓶颈类型 | 症状(Nsight) | 优化手段 |
|---|
访存瓶颈
memory-bound | Memory Throughput 高、SM 低 |
访存合并、用 shared memory 复用、减少冗余读写、提高带宽利用、向量化访存(float4) |
计算瓶颈
compute-bound | SM Throughput 高、Memory 低 |
用更快指令、tensor core、减少计算量、提高 ILP、避免低吞吐指令 |
两者都低
latency-bound | 两个利用率都不高 |
提高 occupancy / 并行度,藏延迟(更多 warp 或更多 ILP) |
典型 AI 算子的算术强度
| 算子 | 强度 | 通常瓶颈 |
|---|
| 逐元素(add/relu/scale) | 极低 ≈0.1 | 访存 |
| 归一化 / softmax | 低 | 访存 |
| reduction(求和等) | 低 | 访存 |
| 小矩阵乘 / GEMV | 中低 | 访存 |
| 大矩阵乘 GEMM(已 tiling) | 高 数十+ | 计算 |
常见数据中心卡脊点(FP32,约值)
| 卡 | 峰值 FP32 | 带宽 | 脊点 |
|---|
| V100 | ~15.7 TF | ~0.9 TB/s | ~17 |
| A100 | ~19.5 TF | ~1.5 TB/s | ~13 |
| H100 (SXM) | ~67 TF | ~3.35 TB/s | ~20 |
注:tensor core / 低精度(FP16/BF16/TF32)峰值更高,脊点会显著右移,更易成为访存瓶颈。确认你的卡型号后核对官方 datasheet。
Nsight 速用
- 抓瓶颈:
ncu --section SpeedOfLight ./prog
- 看
Compute (SM) Throughput vs Memory Throughput,谁高谁是瓶颈。
- 详细访存分析:
ncu --section MemoryWorkloadAnalysis ./prog
调优铁律
1. 永远先测量、定位瓶颈,再优化。
2. 只优化瓶颈那一边 —— 优化非瓶颈边,总耗时不变。
3. 优化后再测一遍,用数据确认提升。
来源:CUDA C++ Best Practices Guide ·
Nsight Compute Profiling Guide