阶段 2 · 执行模型与内存05 / 15约 20 分钟
内存层级：决定性能上限的那张表
寄存器、共享内存、L1/L2、HBM 的延迟与带宽，以及计算强度这个判断工具
学完这一课你会
- 记住各级存储的延迟与带宽数量级差异
- 会用计算强度（arithmetic intensity）判断 kernel 是访存瓶颈还是算力瓶颈
- 知道寄存器溢出（register spilling）的症状和后果
- 理解为什么绝大多数真实 kernel 都是 memory-bound
如果只允许你记住 CUDA 优化的一件事，那应该是这张表。GPU 的算力早已过剩，真正稀缺的是把数据喂进去的能力。 一张 A100 有 19.5 TFLOPS 的 FP32 算力，但显存带宽只有 1.55 TB/s——也就是说，每从显存读一个 float（4 字节），你必须做够 50 次浮点运算才能让算力单元不闲着。绝大多数 kernel 远远达不到这个比例。
| 存储层级 | 作用域 | 延迟（周期） | 带宽量级 | 容量 |
|---|---|---|---|---|
| 寄存器 | 单线程私有 | ~1 | ~20 TB/s | 每 SM 256 KB |
| 共享内存 / L1 | block 内共享 | ~20–30 | ~10 TB/s | 每 SM 最多 164 KB |
| L2 缓存 | 全设备共享 | ~200 | ~4 TB/s | 40–50 MB |
| 全局内存（HBM） | 全设备 + host | ~400–800 | 1–3 TB/s | 16–80 GB |
| 常量内存 | 全设备只读 | ~1（命中缓存） | 广播优化 | 64 KB |
| 本地内存 | 单线程私有（实为显存） | ~400–800 | 同全局内存 | 受限于显存 |
shell
1nvcc -O3 -arch=sm_80 -Xptxas -v -c kernel.cu2 3# 典型输出：4# ptxas info : Used 38 registers, 8192 bytes smem, 380 bytes cmem[0]5# ↑ 每线程寄存器数 ↑ 每 block 共享内存6#7# 出现下面这行就说明溢出了，必须处理：8# ptxas info : 24 bytes spill stores, 24 bytes spill loads计算强度：先判断瓶颈，再谈优化
计算强度 = 浮点运算次数 / 访问的字节数（单位 FLOP/Byte）。把它和硬件的「拐点」比一比，就能知道你的 kernel 卡在哪。拐点 = 峰值算力 / 峰值带宽，A100 FP32 约为 19500 / 1555 ≈ 12.5 FLOP/Byte。
| 运算 | 计算强度 | 瓶颈 | 优化方向 |
|---|---|---|---|
向量加 c = a + b | 1 FLOP / 12 B ≈ 0.08 | 严重访存受限 | 只能优化访存模式，算力毫无意义 |
SAXPY y = a*x + y | 2 FLOP / 12 B ≈ 0.17 | 严重访存受限 | 同上，目标是打满带宽 |
| 朴素矩阵乘（无 tiling） | 2 FLOP / 8 B = 0.25 | 访存受限 | 用 shared memory 做数据复用 |
| Tiled 矩阵乘（tile=32） | ≈ 8 | 接近平衡 | 继续加大 tile、用寄存器分块 |
| 大矩阵乘（寄存器分块） | > 50 | 算力受限 | 上 Tensor Core，榨 FLOPS |
各级存储的声明方式
CUDA C++
1__constant__ float c_filter[256]; // 常量内存：host 用 cudaMemcpyToSymbol 写入2 3__global__ void demo(const float* __restrict__ g_in, // 全局内存4 float* g_out) {5 __shared__ float s_tile[32][33]; // 静态共享内存（33 是防 bank conflict 的 padding）6 extern __shared__ float s_dyn[]; // 动态共享内存，大小由 <<<,,bytes>>> 给出7 8 float acc = 0.0f; // 寄存器9 float buf[8]; // 索引若能编译期展开 → 寄存器；否则 → 本地内存（慢）10 11 #pragma unroll // 强制展开，保证 buf 留在寄存器里12 for (int i = 0; i < 8; ++i) buf[i] = g_in[i];13 14 for (int i = 0; i < 8; ++i) acc += buf[i] * c_filter[i];15 g_out[threadIdx.x] = acc;16}共享内存和 L1 在物理上是同一块 SRAM，可以按比例切分。默认由驱动决定，但你可以给出偏好——重度使用 shared memory 的 kernel（比如 tiled matmul）应该主动多要一些。
CUDA C++
1// 偏好：尽量多给共享内存2cudaFuncSetCacheConfig(myKernel, cudaFuncCachePreferShared);3 4// Volta 及以后可以精确指定每个 block 的共享内存上限（单位字节）5cudaFuncSetAttribute(myKernel,6 cudaFuncAttributeMaxDynamicSharedMemorySize,7 96 * 1024);自测
先自己在心里回答一遍，再展开对照。答不上来的说明这一段值得重读。