共享内存优化：从原理到实战

共享内存，说白了就是GPU上的一块“手搓缓存”。它位于每个线程块内部，所有线程都能访问。我刚开始接触CUDA时，总觉得全局内存够用就行，直到被一次性能瓶颈狠狠教育了一回——那次我写了个矩阵转置，全局内存版本跑得比乌龟还慢。后来换成共享内存，速度直接翻了十几倍。嗯，从那以后我再也不敢小看它了。

共享内存的工作原理

共享内存是片上内存，延迟极低——大概比全局内存快100倍。但它容量有限，每个SM（流式多处理器）通常只有几十到一百多KB。你想想看，这就像你办公桌上的小抽屉，放的是最常用的工具，而不是仓库里的大箱子。

共享内存的生命周期和线程块绑定。线程块启动时分配，线程块结束时释放。同一个块内的线程可以互相读写共享内存，实现数据共享和协作。我个人习惯把共享内存当作“手动缓存”来用——把全局内存中会被多次访问的数据先搬进来，再慢慢处理。

Bank Conflict：共享内存的“隐形杀手”

共享内存被划分成32个Bank（存储体），每个Bank宽度为4字节。当多个线程同时访问同一个Bank的不同地址时，就会发生Bank Conflict。这会导致访问串行化，性能骤降。

为什么会这样？因为硬件设计上，每个Bank每个周期只能响应一个请求。如果两个线程撞到同一个Bank，就得排队。我曾在项目中遇到过这样一个案例：一个简单的直方图计算，因为Bank Conflict，性能比预期差了3倍。排查了半天才发现是地址映射的问题。

访问模式	是否冲突	性能影响
所有线程访问同一个地址	不冲突（广播）	极快
每个线程访问不同Bank	不冲突	理想
多个线程访问同一Bank的不同地址	冲突	串行化，性能下降
多个线程访问同一Bank的同一地址	不冲突（多播）	较快

共享内存的分配方式

共享内存有两种分配方式：静态分配和动态分配。静态分配在编译时就确定大小，动态分配则在运行时指定。

静态共享内存：

__global__ void kernel() {
    __shared__ float s_data[256];  // 编译时确定大小
    // ...
}

动态共享内存：

__global__ void kernel() {
    extern __shared__ float s_data[];  // 大小由启动参数指定
    // ...
}

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

我个人更推荐动态分配，尤其是在写通用库函数时。因为不同场景下需要的共享内存大小可能不同，硬编码会限制灵活性。但要注意，动态共享内存只能声明一个未指定大小的数组，而且所有动态共享内存变量共享同一块内存区域。

动态共享内存的实战技巧

动态共享内存的一个常见用法是处理可变大小的数据块。比如在矩阵乘法中，每个线程块处理的子矩阵大小可能不同。这时候动态分配就派上用场了。

__global__ void matmul(float* A, float* B, float* C, int N) {
    extern __shared__ float s_data[];
    float* s_A = s_data;                    // 前半部分存A
    float* s_B = &s_data[BLOCK_SIZE * BLOCK_SIZE];  // 后半部分存B

    int tx = threadIdx.x, ty = threadIdx.y;
    int row = blockIdx.y * BLOCK_SIZE + ty;
    int col = blockIdx.x * BLOCK_SIZE + tx;

    float sum = 0.0f;
    for (int t = 0; t < N / BLOCK_SIZE; ++t) {
        // 加载数据到共享内存
        s_A[ty * BLOCK_SIZE + tx] = A[row * N + t * BLOCK_SIZE + tx];
        s_B[ty * BLOCK_SIZE + tx] = B[(t * BLOCK_SIZE + ty) * N + col];
        __syncthreads();

        for (int k = 0; k < BLOCK_SIZE; ++k) {
            sum += s_A[ty * BLOCK_SIZE + k] * s_B[k * BLOCK_SIZE + tx];
        }
        __syncthreads();
    }
    C[row * N + col] = sum;
}

// 启动内核
int shared_mem_size = 2 * BLOCK_SIZE * BLOCK_SIZE * sizeof(float);
matmul<<<grid, block, shared_mem_size>>>(d_A, d_B, d_C, N);

这段代码里，我用一个 extern __shared__ 数组同时存放了A和B两个子矩阵。通过指针偏移来划分内存区域。这样做的好处是，你只需要修改 BLOCK_SIZE 就能适应不同的硬件配置，而不需要重新编译。

知识体系结构图

下面这张图展示了共享内存优化的核心逻辑和关键路径：

总结与个人心得

共享内存优化，说白了就是两件事：一是把数据从全局内存搬进来，二是确保搬进来的数据访问时没有Bank Conflict。第一件事靠的是对数据复用模式的理解，第二件事靠的是对地址映射的精细控制。

我个人的经验是，写共享内存代码时，先画个数据流图，搞清楚每个线程需要访问哪些数据，再考虑怎么布局共享内存。别一上来就写代码，那样很容易踩坑。另外，__syncthreads() 这个同步函数一定要用对——少了它数据不一致，多了它性能下降。我见过太多人因为同步放错位置而debug到崩溃。

最后说一句：共享内存虽好，但别贪多。每个SM的资源是有限的，共享内存用多了，能同时运行的线程块就少了。这是一个trade-off，需要根据具体场景来权衡。