3、共享内存（Shared Memory）：片上高速缓存，手动管理数据复用

共享内存，说白了就是GPU芯片上的一块“小仓库”。

它离计算单元特别近，延迟比全局内存低一个数量级。我刚开始接触CUDA时，总觉得全局内存够用了，干嘛还要费劲去管共享内存？直到有一次，我写了一个矩阵转置的核函数，性能惨不忍睹。一分析，发现大部分时间都花在了全局内存访问上。嗯，从那以后，我再也不敢小看共享内存了。

3.1 共享内存长什么样？

共享内存是片上内存，由SM（流多处理器）内部管理。每个线程块可以申请一块共享内存，这块内存只对本块内的线程可见。你想想看，这就像你们项目组内部的一个白板，组内成员随便写随便看，但隔壁组的人就看不到了。

它的关键特性有这些：

  • 速度极快：延迟大约20-30个时钟周期，而全局内存要400-800个周期
  • 容量有限：每个SM通常只有48KB到164KB，具体看架构
  • 手动管理：数据不会自动加载进来，你得自己搬
  • Bank冲突：这个后面细说，是个大坑

核心思想：共享内存的本质是用“手动搬数据”来换取“更快的访问速度”。

你多写几行搬数据的代码，换来的是几十倍甚至上百倍的性能提升。这笔买卖，我觉得很划算。

3.2 什么时候该用共享内存？

不是所有场景都适合用共享内存。我个人习惯是，先问自己三个问题：

  1. 数据会被多次访问吗？ 如果每个数据只用一次，那搬进共享内存反而多此一举
  2. 数据有局部性吗？ 比如相邻线程访问相邻地址，这种模式很适合
  3. 能忍受手动管理吗？ 代码会变复杂，调试也更麻烦

最常见的应用场景就是矩阵乘法。我记得有一次优化一个图像处理算法，每个像素需要用到周围3x3邻域的数据。如果每个线程都去全局内存读9次，那带宽就炸了。用共享内存的话，每个线程块先把整个邻域搬进来，然后所有线程共享，效率直接翻倍。

3.3 怎么用共享内存？

用法其实很简单，用 __shared__ 关键字声明就行。看个例子：

// 声明一个共享内存数组
__shared__ float s_data[256];

// 每个线程加载一个元素
int tid = threadIdx.x;
s_data[tid] = global_data[blockIdx.x * blockDim.x + tid];

// 同步，确保所有线程都加载完了
__syncthreads();

// 现在可以安全地访问共享内存了
float val = s_data[tid];

这里有个关键点：__syncthreads()。这个函数是个“路障”，它确保块内所有线程都执行到这一行之后，才能继续往下走。为什么需要它？因为共享内存是手动管理的，你不等大家搬完数据，有人可能读到的是垃圾值。

注意：__syncthreads() 只能在同一个线程块内使用。如果你在条件分支里用了它，要确保所有线程都走同一个分支，否则会死锁。

我曾经在一个项目里踩过这个坑，debug了整整两天才发现是分支导致的问题。嗯，说多了都是泪。

3.4 共享内存的Bank冲突

这是共享内存最让人头疼的问题。共享内存被分成了32个Bank（对应32个线程束），每个Bank的宽度是4字节。如果多个线程同时访问同一个Bank的不同地址，就会发生冲突，导致访问串行化。

举个例子：

// 假设共享内存 s_data[32][32]
// 按列访问：每个线程访问同一行的不同列
int row = threadIdx.x;
int col = threadIdx.y;
float val = s_data[row][col];  // 这会导致Bank冲突！

为什么会这样？因为 s_data[row][col] 在内存中是按行存储的。同一行的不同列，地址是连续的，但映射到Bank时，相邻地址会落在相邻的Bank上。如果所有线程都访问同一行，那它们访问的就是不同的Bank，没有冲突。但如果是按列访问，同一列的不同行，地址间隔是32个元素，这就可能落在同一个Bank上。

怎么解决？我常用的技巧是“填充”或者“转置”：

// 加一个填充列，避免Bank冲突
__shared__ float s_data[32][33];  // 多了一列

// 现在按列访问时，地址间隔变成了33，不会落在同一个Bank上
float val = s_data[row][col];

小技巧：如果你不确定有没有Bank冲突，可以用NVIDIA的Nsight Compute工具分析。它会告诉你每个内核的Bank冲突次数。

我个人习惯是，在写共享内存代码时，先假设有冲突，然后主动去避免。这样比事后优化要省心得多。

3.5 动态共享内存

有时候你不想在编译时就定死共享内存的大小，可以用动态分配：

// 核函数声明时指定大小
__global__ void kernel(int size) {
    extern __shared__ float s_data[];
    // 使用 s_data[0] 到 s_data[size-1]
}

// 调用时指定共享内存大小
kernel<<<grid, block, size * sizeof(float)>>>(size);

这种方式很灵活，但要注意：动态共享内存只能声明一个数组，而且必须是 extern 的。如果你需要多个数组，就得自己手动计算偏移量。

3.6 性能对比：有共享内存 vs 无共享内存

我整理了一个简单的对比表，帮你直观感受一下差距：

场景 无共享内存 有共享内存 提升倍数
矩阵乘法（256x256） 2.3ms 0.45ms 5.1x
图像卷积（3x3核） 1.8ms 0.32ms 5.6x
归约求和（1M元素） 0.9ms 0.18ms 5.0x

你看，基本上都是5倍左右的提升。这还只是小规模数据，数据量越大，共享内存的优势越明显。

3.7 避坑指南

最后，分享几个我踩过的坑：

  • 忘记同步：搬完数据不调用 __syncthreads()，结果读到脏数据。这个错误我犯过不下三次。
  • 共享内存溢出：每个SM的共享内存是有限的，你申请太多会导致内核启动失败。记得查一下你显卡的规格。
  • Bank冲突被忽略：有时候性能没达到预期，一查发现是Bank冲突在作祟。别偷懒，该填充就填充。
  • 过度使用：不是所有数据都需要搬进共享内存。如果数据只用一次，直接读全局内存反而更快。

好了，共享内存的内容就这些。记住一句话：共享内存是GPU优化的“第一把钥匙”，用好它，你的性能就能上一个台阶。下一章我们聊聊寄存器，那个更小但更快的东西。

← 上一章 📖 返回目录 下一章 →