Shared Memory:少读,而非读得对

Lesson 04 · CUDA 性能调优 · 访存瓶颈的第二大武器

前三节都在教你把全局内存读得对(合并、对齐)。这一节是质变: 有些数据会被反复读很多次。与其每次都跑去显存这条"远路",不如搬进片上的 shared memory 读一次、复用多次——直接减少总访存量,把算术强度抬上去。

1. 内存金字塔:远近差百倍

上一节出现过缓存层级。把它按"离计算单元多远 / 多快"排成金字塔:[1]

关键差距:shared memory 比全局内存快约一个数量级。而且它和 L1 共用同一块物理 SRAM—— 区别在于 shared memory 是你手动管理的(你决定放什么、何时放),不像 L1/L2 是硬件自动的、不可控的。

上一节说"别依赖缓存救场",因为缓存命中不可控。shared memory 就是那个可控版的缓存: 你亲手把热数据钉在片上,保证后续访问都是快的。

2. 什么时候用:看「数据复用」

shared memory 不是万能药。它只在一种情况下有用:同一份数据被一个 block 内多个线程、或同一线程多次读取。

3. 三步用法 + 那个致命的同步点

shared memory 的标准套路就三步,但第二步漏了会出隐蔽的错误结果:

__global__ void kernel(float* g) {
    // ① 声明片上数组(整个 block 共享)
    __shared__ float tile[256];

    // ② 协作加载:每个线程搬一份全局数据进 shared(合并访问!)
    tile[threadIdx.x] = g[blockIdx.x * 256 + threadIdx.x];

    __syncthreads();   // ★ 屏障:等全 block 都加载完,再往下走

    // ③ 现在随便复用 tile[...],都是片上高速访问
    float v = tile[threadIdx.x] + tile[(threadIdx.x + 1) % 256];
    // ...
}

__syncthreads() 是一道 block 内屏障:所有线程都到齐了才一起继续。 为什么非要它?因为线程是分 warp 异步执行的——

没有屏障,线程 5 可能在线程 200 还没把 tile[200] 写进去时,就抢先去读 tile[200], 读到一堆未初始化的垃圾。而且这种 bug 时隐时现(取决于调度),极难调试。

4. 为什么这是在"抬高算术强度"

回到 第 1 节的 roofline。shared memory 的本质是减少全局访存量(分母变小), 于是算术强度 = FLOPs / Bytes 上升——kernel 在 roofline 图上向右移动, 从访存瓶颈一侧挪向计算瓶颈一侧。

这就是矩阵乘 tiling(下一节的主题)的核心思想:不用 shared memory 时,每个输出元素都要从全局反复读整行整列; 用了之后,一块数据加载一次、被整个 block 复用,全局访存量降一个数量级,GEMM 才可能从访存瓶颈变成计算瓶颈、逼近峰值算力。[2]

5. 练习:该不该用 / 用对了吗

凭前面的口诀与铁律判断。点选项看反馈。

主源推荐(本节精读)

📘 CUDA C++ Best Practices Guide — Shared Memory ——讲清楚了 shared memory 的用途、协作加载模式与同步。配合 GPU MODE 的矩阵乘讲解更佳。