CUDA共享内存的用法

• 1. 什么是共享内存？
• 2. 为什么使用共享内存？
• 3. 如何使用共享内存？
  • 3.1. 静态共享内存
  • 3.2. 动态共享内存
  • 3.3. 区别比较
• 4. 注意事项
  • 4.1. 线程同步
  • 4.2. Bank 冲突
  • 4.3. 容量限制
• 5. 典型应用场景
• 6. Demo演示
  • 6.1. 需求说明
  • 6.2. 代码实现
  • 6.3. 源码说明

1. 什么是共享内存？

共享内存是 NVIDIA GPU 上的一个关键的内存层次。你可以把它理解为一个软件管理的缓存或片上内存。

• 位置： 它在物理上位于每个 流多处理器（SM） 的芯片上，而不是像全局内存（用 cudaMalloc 分配）那样在显存的 DRAM 上。
• 速度： 它的访问速度极快，堪比寄存器（L1缓存级别），比全局内存快上百倍。
• 作用域： 共享内存由同一个线程块（Block）内的所有线程共享。
• 生命周期： 它的生命周期与其所属的线程块一致。当线程块开始执行时被分配，当该线程块的所有线程都执行完毕时被释放。

2. 为什么使用共享内存？

核心思想：减少对全局内存的慢速访问。

1. 作为可编程的缓存（Programmable Cache）： 你可以手动将全局内存中的数据“搬运”到共享内存中。线程块内的所有线程可以高速地、反复地访问这块数据，从而避免多次直接访问低速的全局内存。
2. 线程块内部的协作（Inter-Thread Communication）： 因为同一个块内的线程可以共享这块内存，所以线程之间可以通过它来交换数据和协调工作。这是实现很多并行算法（如归约、扫描、矩阵乘法优化等）的基础。

3. 如何使用共享内存？

共享内存使用__shared__限定符来声明一个内存变量，主要有两种声明方式：静态声明和动态声明。

3.1. 静态共享内存

在内核中使用 __shared__ 关键字直接声明一个固定大小的数组。

__global__ void myKernel(float* input, float* output) {
    // 声明一个静态的共享内存数组，大小为512个float
    __shared__ float s_data[512];

    int tid = threadIdx.x;
    int gid = threadIdx.x + blockIdx.x * blockDim.x;

    // 示例：每个线程将全局内存数据加载到共享内存
    s_data[tid] = input[gid];

    // ！！！非常重要：确保所有线程都已完成数据加载
    __syncthreads();

    // 现在，所有线程都可以安全地访问共享内存中的任何数据
    // 例如，进行一个块内归约求和
    for (int s = 1; s < blockDim.x; s *= 2) {
        if (tid % (2 * s) == 0) {
            s_data[tid] += s_data[tid + s];
        }
        __syncthreads(); // 每一轮计算后都需要同步
    }

    // 将结果写回全局内存
    if (tid == 0) {
        output[blockIdx.x] = s_data[0];
    }
}

关键点：

• 大小必须在编译时确定（如 [512]）。
• 同一个内核中的多个静态声明会共享同一块内存区域，总大小不能超过限制（例如 48KB）。

3.2. 动态共享内存

当你需要在启动内核时才决定共享内存的大小时，使用此方式。

步骤：

1. 在内核中：使用 extern __shared__ 声明一个未定大小的数组。一个内核中只能有一个这样的声明。
2. 在主机端：在 <<<grid, block, smem_size>>> 的执行配置中，第三个参数指定每个线程块需要动态分配的共享内存字节数。

// Kernel 定义
__global__ void myDynamicKernel(float* input, float* output) {
    // 声明一个动态的共享内存数组，大小在启动内核时决定
    extern __shared__ float s_dynamic[]; // 通常是一个字节数组

    int tid = threadIdx.x;
    int gid = threadIdx.x + blockIdx.x * blockDim.x;

    // 使用：将字节数组转换为所需类型
    float* s_data = (float*)s_dynamic;

    s_data[tid] = input[gid];
    __syncthreads();

    // ... 后续操作与静态声明相同 ...
}

// 主机端代码调用
int main() {
    // ...
    int block_size = 256;
    int shared_mem_size = block_size * sizeof(float); // 计算需要的字节数

    // 启动内核，第三个参数就是动态共享内存的大小（字节）
    myDynamicKernel<<<num_blocks, block_size, shared_mem_size>>>(d_input, d_output);
    // ...
}

关键点：

• 更灵活，大小可以在主机代码中计算并传入。
• 内核中通过 extern __shared__ 声明的是一个 char 数组（字节数组），通常需要将其类型转换为你需要的类型（如 float*）后再使用。

3.3. 区别比较

特性	静态共享内存	动态共享内存
声明	__shared__ float arr[128];	extern __shared__ char arr[];
大小	编译时固定	运行时通过 <<<>>> 的第三个参数指定（字节）
灵活性	较低	高
数量限制	一个内核可声明多个，但总大小受限	一个内核只能声明一个 extern 数组

4. 注意事项

4.1. 线程同步

• __syncthreads()是 CUDA 提供的一个线程块内的屏障同步函数。

• 必须在访问共享内存之前，确保所有线程都已经完成了对共享内存的写入操作。在上面的例子中，加载数据后和归约计算的每一步之后都需要调用 __syncthreads()。

• 警告：在分支语句（如 if）中使用 __syncthreads() 必须非常小心。它要求同一个线程块中的所有线程都必须遇到这个函数，否则会导致死锁。如下面的错误示例：

  if (threadIdx.x < 32) {
      s_data[threadIdx.x] = ...;
      __syncthreads(); // 危险！只有前32个线程会执行到此，其他线程不会，导致死锁。
  }

4.2. Bank 冲突

• 为了实现高带宽，共享内存被划分为若干个（通常是32个）大小相同的内存，也叫Banks。
• 如果多个线程同时访问同一个 Bank 中的不同地址，这些访问会串行化，导致性能下降。这称为Bank冲突(Bank Conflict)。
• 最佳实践：尽量让线程访问的地址使得线程 ID 与 Bank ID 是交错（Strided） 的，或者让同一个 warp 中的线程访问同一个地址（广播，没有冲突）。例如，在矩阵转置优化中，通过填充（Padding）技巧可以避免 Bank Conflict。

4.3. 容量限制

• 每个 SM 的共享内存总量是有限的（例如，安培架构的 A100 每 SM 为 164KB，但可配置为 100KB 共享内存 + 64KB L1 缓存，或反之）。
• 你为每个线程块分配的共享内存大小会影响 SM 上可同时驻留的线程块数量，从而影响 Occupancy（占用率）。需要根据算法需求和硬件限制进行权衡。

5. 典型应用场景

• 矩阵乘法（GEMM）： 将矩阵 A 和 B 的子块从全局内存加载到共享内存，然后进行大量乘加运算，最后将结果写回全局内存。这是最经典的优化案例。
• 归约（Reduction）： 如求和、求最大值等。每个线程块在共享内存中完成一部分数据的归约，最后将每个块的结果汇总。
• 卷积（Convolution）： 将图像块和滤波器权重加载到共享内存，利用其高速特性进行卷积计算。
• 扫描（Scan / Prefix Sum）： 类似归约，也是一种常见的并行算法模式。
• 查找表（Look-up Tables）： 如果需要一个小型的、被所有线程频繁读取的表，可以将其放入共享内存。

6. Demo演示

6.1. 需求说明

数组平方和： 我们将一个数组中的每个元素平方，然后计算所有这些平方值的和。

• CPU 版本：顺序执行，一个接一个地计算。
• GPU 版本（无共享内存）：每个线程计算一个元素的平方，然后使用原子操作求和。这种方法性能较差，但易于理解，用于对比。
• GPU 版本（使用共享内存）：每个线程块先在共享内存内部进行部分求和，大大减少了对全局内存的原子操作次数。这是高效的做法。

6.2. 代码实现

#include <stdio.h>
#include <stdlib.h>
#include "timecounter.cuh"


// 求数组的平方和：CPU版本
void cpu_quadratic_sum(int *array, int count, int *result)
{
    // 统计运行时间
    SteadyTimeCounter counter("cpu_quadratic_sum");

    int sum = 0.0f;
    for (int i = 0; i < count; i++)
    {
        sum += array[i] * array[i]; // 计算平方并累加
    }
    *result = sum;
}

// 核函数：每个线程计算一个平方，并使用原子操作加到全局变量
__global__ void gpu_naive_sum(int* array, int count, int* result) {
    int global_idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (global_idx < count) {
        int square = array[global_idx] * array[global_idx];
        atomicAdd(result, square); // 慢！原子操作是性能瓶颈
    }
}

// 求数组的平方和：GPU版本(无共享内存)
void gpu_quadratic_sum_no_sharedmemory(int *array, int count, int *result)
{
    // 统计运行时间
    SteadyTimeCounter counter("gpu_quadratic_sum_no_sharedmemory");

    // 创建和分配设备内存
    int *d_array, *d_result;
    cudaMalloc(&d_array, count * sizeof(int));
    cudaMalloc(&d_result, sizeof(int));
    // 将数据拷贝到设备内存
    cudaMemcpy(d_array, array, count * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_result, result, sizeof(int), cudaMemcpyHostToDevice);

    // 设置线程块和网格大小
    int threadsPerBlock = 256;
    int blocksPerGrid = (count + threadsPerBlock - 1) / threadsPerBlock;

    // 启动核函数
    gpu_naive_sum<<<blocksPerGrid, threadsPerBlock>>>(d_array, count, d_result);

    // 等待所有线程完成计算
    cudaDeviceSynchronize();

    // 将结果拷贝回主机
    cudaMemcpy(result, d_result, sizeof(int), cudaMemcpyDeviceToHost);

    // 释放设备内存
    cudaFree(d_array);
    cudaFree(d_result);
}

// 线程块内使用共享内存求平方和
__global__ void gpu_shared_sum(int *array, int count, int *result)
{
    // 声明静态共享内存，大小等于一个线程块的线程数
    extern __shared__ int s_data[]; // 通常是一个字节数组
    // 使用：将字节数组转换为所需类型
    // int* s_data = (int*)s_dynamic;

    // 计算线程索引
    int tid = threadIdx.x;
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    // 1. 每个线程将它的平方值加载到共享内存中
    s_data[tid] = (idx < count) ? array[idx] * array[idx] : 0.0f;

    // 2. 等待块内所有线程都完成加载
    __syncthreads();

    // 3. 在共享内存上进行归约求和
    // 归约算法：迭代地将一半的线程的数据相加
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) {
            s_data[tid] += s_data[tid + s];
        }
        // 每次迭代后都必须同步，确保下一步计算前数据已准备好
        __syncthreads();
    }

    // 4. 现在，共享内存中的s_data[0]就是这个线程块的部分和
    // 由第一个线程（tid=0）使用原子操作将部分和加到全局结果上
    if (tid == 0) {
        atomicAdd(result, s_data[0]);
    }
}

// 求数组的平方和：GPU版本(使用共享内存)
void gpu_quadratic_sum_with_sharedmemory(int *array, int count, int *result)
{    
    // 统计运行时间
    SteadyTimeCounter counter("gpu_quadratic_sum_with_sharedmemory");

    // 创建和分配设备内存
    int *d_array, *d_result;
    cudaMalloc(&d_array, count * sizeof(int));
    cudaMalloc(&d_result, sizeof(int));
    // 将数据拷贝到设备内存
    cudaMemcpy(d_array, array, count * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_result, result, sizeof(int), cudaMemcpyHostToDevice);

    // 设置线程块和网格大小
    int threadsPerBlock = 256;
    int blocksPerGrid = (count + threadsPerBlock - 1) / threadsPerBlock;

    // 启动核函数
    int shared_memory_size = threadsPerBlock * sizeof(int);
    gpu_shared_sum<<<blocksPerGrid, threadsPerBlock, shared_memory_size>>>(d_array, count, d_result);

    // 将结果拷贝回主机
    cudaMemcpy(result, d_result, sizeof(int), cudaMemcpyDeviceToHost);

    // 释放设备内存
    cudaFree(d_array);
    cudaFree(d_result);
}

int main()
{
    constexpr int N = 1024 * 1024; // 1M 个元素
    // constexpr int N = 256;
    int *h_array = new int[N];
    int h_result1 = 0.0f;
    int h_result2 = 0.0f;
    int h_result3 = 0.0f;

    // 初始化主机内存
    for (int i = 0; i < N; i++)
    {
        h_array[i] = (int)(i % 10); // 用一些简单的数据填充
    }

    // CPU串行版本
    cpu_quadratic_sum(h_array, N, &h_result1);
    printf("CPU Sum: %d\n", h_result1);

    // GPU并行版本(无共享内存)
    gpu_quadratic_sum_no_sharedmemory(h_array, N, &h_result2);
    printf("GPU Sum(No SharedMemory): %d\n", h_result2);

    // GPU并行版本(使用共享内存)
    gpu_quadratic_sum_with_sharedmemory(h_array, N, &h_result3);
    printf("GPU Sum(With SharedMemory): %d\n", h_result3);

    delete[] h_array;
    return 0;
}

# 编译
nvcc ./shared_memory.cu

# 运行
./a.out
[steady timecount]  cpu_quadratic_sum used 2522us
CPU Sum: 29884300
[steady timecount]  gpu_quadratic_sum_no_sharedmemory used 268260us
GPU Sum(No SharedMemory): 29884300
[steady timecount]  gpu_quadratic_sum_with_sharedmemory used 2341us
GPU Sum(With SharedMemory): 29884300

6.3. 源码说明

在上面的Demo中，实现了三种“求数组平方和”的方法：

方法	关键代码	说明
CPU串行版本	cpu_quadratic_sum()	这个版本逻辑非常简单，循环遍历每个元素，计算平方并累加。
GPU并行版本(无共享内存)	gpu_quadratic_sum_no_sharedmemory()	这个版本启动了 N 个线程，每个线程负责计算一个元素的平方，并直接使用原子操作 atomicAdd 将结果累加到全局内存的一个变量中。原子操作保证了在多个线程同时进行加法时结果的正确性，但它的速度很慢，因为它需要防止访问冲突。
GPU并行版本(有共享内存)	gpu_quadratic_sum_with_sharedmemory()	1. 每个线程块分配一块共享内存，用于计算该线程块内所有线程数据的部分和。 2. 每个线程计算一个平方值，先存到共享内存中。 3. 在线程块内部，使用一种叫做归约（Reduction） 的算法将共享内存中的所有值求和，得到一个该线程块的部分和。 4. 每个线程块只将其最终的部分和（一个数字）使用原子操作加到全局内存中。