第四章：共享内存与同步机制

共享内存，说白了就是GPU里的一块“手递手”区域。每个线程块独享一份，速度比全局内存快两个数量级。我刚开始接触CUDA时，总觉得全局内存够用就行，直到被一次性能瓶颈狠狠教育了一回——嗯，从那以后我再也不敢忽视共享内存了。

4.1 共享内存的声明与生命周期

声明共享内存有两种方式：静态声明和动态声明。我个人习惯用静态声明，代码更清晰。

// 静态声明：编译时确定大小
__shared__ float cache[256];

// 动态声明：运行时确定大小
extern __shared__ float cache[];
// 调用时：kernel<<<grid, block, shared_mem_size>>>(...);

生命周期很简单：从线程块创建开始，到线程块销毁结束。你想想看，这就像每个线程块自己带了个小仓库，用完了就拆掉，互不干扰。

4.2 __syncthreads() 的正确使用场景

为什么需要同步？因为共享内存是线程块内所有线程共享的。如果一个线程写，另一个线程读，你不加同步，读到的可能是垃圾数据。

我在项目中遇到过这样一个坑：做矩阵转置时，线程A把数据写进共享内存，线程B立刻去读。结果呢？数据还没写完，B就读到了旧值。整个结果全乱套了。

__global__ void example_kernel() {
    __shared__ float sdata[256];
    
    int tid = threadIdx.x;
    sdata[tid] = input[tid];  // 写操作
    
    __syncthreads();  // 等所有线程写完
    
    float val = sdata[255 - tid];  // 读操作，安全了
    output[tid] = val;
}

4.3 共享内存的bank冲突

共享内存被分成32个bank（对应warp的32个线程）。如果多个线程同时访问同一个bank的不同地址，就会发生bank冲突，性能骤降。

访问模式	bank冲突	性能影响
所有线程访问同一地址	无冲突（广播）	最优
连续地址访问	无冲突	良好
步长为2的访问	2路冲突	性能减半
随机访问	可能32路冲突	最差

解决bank冲突的方法很简单：加padding。比如声明共享内存时多留几个元素，让地址错开。

// 避免bank冲突的技巧
// 原本：__shared__ float data[32][32];
// 改成：
__shared__ float data[32][33];  // 多一列padding

4.4 利用共享内存减少全局内存访问

这是共享内存最核心的应用。说白了就是：把数据从慢速的全局内存搬到快速的共享内存，然后反复使用。

我做过一个矩阵乘法的优化案例，效果很明显：

__global__ void matmul_shared(float* A, float* B, float* C, int N) {
    __shared__ float Asub[TILE_SIZE][TILE_SIZE];
    __shared__ float Bsub[TILE_SIZE][TILE_SIZE];
    
    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    
    float sum = 0.0f;
    
    for (int t = 0; t < N / TILE_SIZE; t++) {
        // 加载数据到共享内存
        Asub[ty][tx] = A[by * TILE_SIZE + ty][t * TILE_SIZE + tx];
        Bsub[ty][tx] = B[t * TILE_SIZE + ty][bx * TILE_SIZE + tx];
        
        __syncthreads();  // 确保加载完成
        
        // 从共享内存计算
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += Asub[ty][k] * Bsub[k][tx];
        }
        
        __syncthreads();  // 防止覆盖数据
    }
    
    C[by * TILE_SIZE + ty][bx * TILE_SIZE + tx] = sum;
}

4.5 知识体系总览

下面这张图是我自己整理的共享内存知识框架，你可以照着这个思路去理解：

4.6 实战中的避坑指南

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

• 共享内存大小超限：我曾经在代码里声明了64KB的共享内存，结果在计算能力3.5的设备上直接报错。每个SM的共享内存上限是固定的，查一下设备规格再动手。
• 忘记同步：这是最常见的bug。写共享内存后不加__syncthreads()，读到的数据就是错的。调试起来特别痛苦，因为不是每次都会出错。
• 过度使用同步：同步是有开销的。我见过有人每个操作后面都加同步，性能直接腰斩。只在必要的地方加，比如写后读的边界处。
• bank冲突忽视：用NVIDIA Nsight Compute工具跑一下，看看有没有bank冲突。我优化过一个项目，解决bank冲突后性能提升了30%。

共享内存用好了，性能提升是立竿见影的。但也要记住：它不是万能药。数据量太大时，该用全局内存还是得用。关键是在合适的地方用合适的工具。