第3章：共享内存优化——从入门到精通

共享内存，说白了就是GPU里的一块“高速缓存”，由你手动控制。它比全局内存快得多，但容量有限。我刚开始接触CUDA时，总觉得全局内存够用就行，直到被一个矩阵乘法的性能瓶颈狠狠教育了一顿。嗯，从那以后，我再也不敢忽视共享内存了。

3.1 共享内存的声明与使用

共享内存在核函数内部声明，用 __shared__ 关键字。它的生命周期是整个线程块。每个线程块都有自己的副本，线程块之间不可见。

我个人习惯把共享内存声明在核函数的开头，这样一目了然。举个例子：

__global__ void sharedMemoryExample(float *input, float *output) {
    __shared__ float sharedData[256];  // 静态声明，大小固定

    int tid = threadIdx.x;
    sharedData[tid] = input[tid];      // 每个线程加载一个元素

    __syncthreads();                   // 确保所有线程加载完毕

    // 做一些计算...
    output[tid] = sharedData[tid] * 2.0f;
}

这里要注意：__shared__ 声明的数组大小必须是编译期常量。你想想看，如果大小不确定怎么办？那就得用动态共享内存了。

3.2 动态共享内存

动态共享内存允许你在运行时指定大小。声明时用 extern __shared__，然后在核函数调用时通过第三个参数传入大小。

我在项目中遇到过这样一个场景：处理不同尺寸的图像，共享内存大小需要动态调整。静态声明显然不行，动态共享内存就派上了用场。

__global__ void dynamicSharedExample(float *input, float *output, int size) {
    extern __shared__ float sharedData[];  // 动态声明，不指定大小

    int tid = threadIdx.x;
    if (tid < size) {
        sharedData[tid] = input[tid];
    }

    __syncthreads();

    if (tid < size) {
        output[tid] = sharedData[tid] + 1.0f;
    }
}

// 调用时指定共享内存大小（字节）
dynamicSharedExample<<<1, 256, 256 * sizeof(float)>>>(d_input, d_output, 256);

3.3 共享内存与同步

共享内存是线程块内所有线程共享的。这就带来了一个经典问题：数据竞争。多个线程同时读写同一个共享内存地址，结果不可预测。

解决办法就是 __syncthreads()。这个函数会阻塞线程块内的所有线程，直到它们都到达这个同步点。说白了，就是“大家等一下，等所有人都准备好了再继续”。

我建议你在以下场景必须使用 __syncthreads()：

• 一个线程写共享内存，另一个线程读同一个位置
• 多个线程协作完成一个计算任务（比如归约、扫描）
• 使用共享内存作为中间缓存，进行数据重排

__global__ void reductionKernel(float *input, float *output) {
    __shared__ float sharedData[256];

    int tid = threadIdx.x;
    sharedData[tid] = input[blockIdx.x * blockDim.x + tid];

    __syncthreads();  // 必须同步，否则下面的归约会读到未初始化的数据

    // 树形归约
    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride) {
            sharedData[tid] += sharedData[tid + stride];
        }
        __syncthreads();  // 每次迭代后都要同步
    }

    if (tid == 0) {
        output[blockIdx.x] = sharedData[0];
    }
}

3.4 共享内存的Padding技巧

共享内存在硬件上被分成多个bank（通常32个）。每个bank每周期可以服务一个地址。如果多个线程同时访问同一个bank的不同地址，就会发生bank冲突，导致串行化。

你想想看，本来可以并行访问的，结果变成了排队，性能自然就下来了。解决办法就是Padding——在共享内存数组的每一行末尾加几个无用元素，让相邻的线程访问不同的bank。

举个例子，一个2D共享内存数组 shared[32][32]，如果按列访问，同一列的32个线程会访问同一个bank，冲突严重。加上Padding变成 shared[32][33]，就能避免这个问题。

// 没有Padding，按列访问时bank冲突严重
__shared__ float noPad[32][32];

// 加上Padding，避免bank冲突
__shared__ float withPad[32][33];  // 每行多一个元素

int row = threadIdx.y;
int col = threadIdx.x;

// 按列访问时，withPad[row][col] 的地址会错开
float val = withPad[row][col];  // 无bank冲突

嗯，这里要注意：Padding的大小不是随便选的。你需要了解目标GPU的bank数量（通常是32）和bank宽度（通常是4字节）。一般来说，每行加1个元素就足够了。

最后，我画了一张图来总结本章的知识体系：

这张图把共享内存优化的四个核心方向串起来了。声明与使用是基础，动态共享内存解决灵活性，同步保证正确性，Padding提升性能。四者缺一不可。

好了，共享内存这块就讲到这里。记住：共享内存是GPU编程的“瑞士军刀”，用好了事半功倍。多写多练，你也能成为高手。