Shared Memory 优化：从入门到精通

Shared Memory 这东西，说白了就是 GPU 上的一块“手写缓存”。

我刚开始接触 CUDA 时，总觉得 Global Memory 够快了啊，干嘛还要折腾 Shared Memory？直到有一次我写了个矩阵乘法，性能惨不忍睹。一分析，发现大部分时间都花在了反复读取 Global Memory 上。嗯，从那以后，我再也不敢小看 Shared Memory 了。

Shared Memory 的分配与使用

Shared Memory 是每个线程块私有的。它的分配方式有两种：静态分配和动态分配。

静态分配

在核函数里直接声明，大小在编译时就定死了。

__global__ void kernel() {
    __shared__ float cache[256];  // 每个线程块分配 256 个 float
    // ...
}

这种方式简单直接。我个人习惯在核函数内部用静态分配，因为代码可读性好，调试也方便。

动态分配

大小在运行时才确定。你需要在外部的核函数调用参数中指定大小。

__global__ void kernel() {
    extern __shared__ float cache[];  // 注意：没有指定大小
    // ...
}

// 调用时指定大小
kernel<<<grid, block, 256 * sizeof(float)>>>();

动态分配的好处是灵活。我在项目中遇到过这样的情况：同一个核函数，处理不同大小的数据块。用动态分配，一个核函数就能搞定所有场景。

小技巧：动态分配的 Shared Memory 大小，建议用 4 字节对齐。比如你只需要 100 个 float，那就分配 100 * sizeof(float) 而不是 100 * 4。这样代码更通用，也更容易维护。

Padding 技术：避免 Bank Conflict

Shared Memory 被分成了 32 个 Bank。每个 Bank 每周期只能响应一次访问。如果多个线程同时访问同一个 Bank，就会发生 Bank Conflict，导致访问串行化。

你想想看，本来一次就能读完的数据，现在要分好几次。性能自然就下来了。

解决思路很简单：让每个线程访问不同的 Bank。Padding 就是最常用的手段。

什么是 Padding？

在数组的每一行末尾，额外加几个“无用”的元素。这样原本对齐到同一个 Bank 的列，就被错开了。

// 没有 Padding：每行 32 个 float
__shared__ float data[32][32];

// 有 Padding：每行 33 个 float，最后一个不用
__shared__ float data[32][33];

为什么加一个就够了？因为 33 不是 32 的倍数。这样每行的起始地址就错开了，Bank 冲突自然就消失了。

  核心要点：Padding 的本质是“破坏对齐”。让数组的 stride 不等于 Bank 数量的整数倍。

我曾经在一个图像处理项目里，用 Padding 把 Shared Memory 的访问效率提升了 40%。当时排查了好久，发现就是 Bank Conflict 在作祟。加上一个 Padding，性能直接起飞。

什么时候需要 Padding？

不是所有场景都需要。我总结了几个判断标准：

按列访问二维数组时，几乎必加
按步长访问一维数组，且步长是 32 的倍数时，建议加
随机访问模式，看具体情况

注意：Padding 会浪费一点 Shared Memory 空间。但通常只浪费几个 float，换来的是几倍的性能提升。这笔买卖，划算。

动态 Shared Memory 的实战技巧

动态 Shared Memory 的用法，上面已经提过了。这里我重点说说实战中的坑。

大小计算要小心

动态 Shared Memory 的大小，是在核函数调用时通过第三个参数指定的。这个参数的单位是字节。

// 错误示范：忘了乘 sizeof
kernel<<<grid, block, 256>>>();  // 只有 256 字节，不是 256 个 float

// 正确写法
kernel<<<grid, block, 256 * sizeof(float)>>>();

我曾经犯过这个错。调试了半天，数据总是对不上。后来发现是大小没算对。嗯，这种低级错误，犯一次就够了。

多个动态数组怎么处理？

一个核函数只能有一个动态 Shared Memory 数组。如果你需要多个，可以用指针偏移来模拟。

__global__ void kernel() {
    extern __shared__ char shared_mem[];
    
    float* data1 = (float*)shared_mem;
    int* data2 = (int*)(shared_mem + N * sizeof(float));
    
    // 现在 data1 和 data2 分别指向不同的区域
}

注意对齐问题。float 和 int 都是 4 字节对齐，所以这样用没问题。如果你混用了 double（8 字节对齐），就要小心了。

知识体系总览

下面这张图，是我梳理的 Shared Memory 优化知识结构。你可以把它当作一个快速索引。

避坑指南

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

忘记同步：Shared Memory 的写入和读取之间，一定要加 __syncthreads()。否则数据还没写完，别的线程就读了。
过度使用：Shared Memory 总共就几十 KB。别一股脑全塞进去，留点给寄存器和其他用途。
忽略对齐：动态 Shared Memory 的起始地址是 16 字节对齐的。如果你在里面放了个 8 字节的数据结构，要小心偏移计算。

Shared Memory 优化，说白了就是“空间换时间”。多花几个字节的存储，换来几倍的性能提升。这笔账，怎么算都划算。

想要更多资料，联系可可

📚 课程目录

Shared Memory 优化：从入门到精通

Shared Memory 的分配与使用

静态分配

动态分配

Padding 技术：避免 Bank Conflict

什么是 Padding？

什么时候需要 Padding？

动态 Shared Memory 的实战技巧

大小计算要小心

多个动态数组怎么处理？

知识体系总览

避坑指南