关于CUDA中__shared__变量的理解误区及代码问题咨询

阅读NVIDIA手册后编写了如下并行规约代码，尝试对每个线程块内的整数求和并打印结果：

__global__ void kernel(int *devData)
{
    __shared__ int sum;
    int i = blockDim.x  * blockIdx.x + threadIdx.x;
    if (threadIdx.x == 0)
        sum = 0;
    __syncthreads();
    sum += devData[i];
    __syncthreads();
    if (threadIdx.x == 0)
        printf("sum of block %d is %d\n", blockIdx.x, sum);
}

int main(void)
{
    // init device
    int devIdx = 0;
    cudaError_t err = cudaSuccess;
    gpuDeviceInit(devIdx);
    int i;
    int data[100];
    int *devData;
    for (i = 0; i < 100; i++)
        data[i] = 1;
    err = cudaMalloc(&devData, 100 * sizeof(int));
    checkCudaErrors(err);

    // copy data to device
    err = cudaMemcpy(devData, data, 100 * sizeof(int), cudaMemcpyHostToDevice);
    checkCudaErrors(err);

    int blocksPerGrid = 10;
    int threadsPerBlock = 10;
    
    // call kernel function
    kernel <<<blocksPerGrid, threadsPerBlock>>> (devData);
    checkCudaErrors(cudaGetLastError());
    cudaDeviceReset();
    return 0;
}

运行结果为：

sum of block 0 is 1
sum of block 6 is 1
sum of block 2 is 1
sum of block 8 is 1
sum of block 1 is 1
sum of block 7 is 1
sum of block 4 is 1
sum of block 3 is 1
sum of block 9 is 1
sum of block 5 is 1

预期每个线程块的求和结果应为10，疑问：

• __shared__变量是否被线程块内所有线程共享？
• 对CUDA中__shared__变量的理解存在哪些错误？

首先明确：__shared__变量确实是线程块内所有线程共享的，你在这一点上的理解没有错误，问题出在共享变量的写入逻辑上。

你的代码中，sum += devData[i]这一步存在严重的数据竞争：线程块内的10个线程会同时对同一个共享变量sum执行读-修改-写操作。由于GPU线程的并行执行特性，这些操作不会按顺序执行，最终只有一个线程的写入结果会保留，其他线程的累加操作被覆盖，所以sum最终值为1（单个线程的贡献）。

要解决这个问题，有两种常见方案：

方案1：使用原子操作 #

对共享变量的累加操作改用原子函数，确保每次只有一个线程访问sum：

atomicAdd(&sum, devData[i]);

替换原代码中的sum += devData[i]即可，这种方式实现简单但性能较低，适合小线程块场景。

方案2：标准并行规约实现 #

这是更高效的方式，利用共享内存分阶段求和：

1. 每个线程先将对应数据加载到共享内存的独立位置，避免竞争
2. 按步长逐步合并相邻线程的计算结果，最终得到块内总和

修改后的核心代码示例：

__global__ void kernel(int *devData)
{
    __shared__ int sum[10]; // 每个线程对应一个共享内存位置
    int tid = threadIdx.x;
    int i = blockDim.x * blockIdx.x + tid;
    sum[tid] = devData[i]; // 先加载到独立位置，无竞争
    __syncthreads();

    // 并行规约：逐步合并
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) {
            sum[tid] += sum[tid + s];
        }
        __syncthreads();
    }

    if (tid == 0) {
        printf("sum of block %d is %d\n", blockIdx.x, sum[0]);
    }
}

这种方式通过分阶段的无竞争累加，大幅提升性能，是CUDA并行规约的标准实现方式。

内容的提问来源于stack exchange，提问作者shyHelios