CUDA 共享变量 - C++

在 CUDA 中，共享变量是指在一个工作组（block）内所有线程（thread）可见的变量。共享变量存储在共享内存（shared memory）中，这是位于每个工作组内的一块高速存储器，可用于线程间共享数据。

在 CUDA 中，可以使用如下语句在共享存储器中定义共享变量：

__shared__ int sharedVar;

共享变量仅能在同一工作组内的线程中共享和访问：

if(threadIdx.x == 0) {
    sharedVar = 42; // 线程 0 初始化共享变量
}
__syncthreads();
printf("thread %d: sharedVar=%d\n", threadIdx.x, sharedVar); // 所有线程输出共享变量的值

需要注意的是，在初始化共享变量后，需要调用 __syncthreads() 函数来确保所有线程已更新共享变量：

__syncthreads();

共享变量的使用通常能够提高访问速度，但也需要注意避免数据竞争（data race）等问题。可以使用锁（mutex）等同步机制来解决这些问题。

下面的示例演示了如何使用共享变量来并行计算数组的和：

__global__ void sum(int* array, int n, int* result) {
    __shared__ int sharedArray[256]; // 声明一个大小为 256 的共享数组
    int tid = threadIdx.x; // 线程 ID
    int i = blockIdx.x * blockDim.x + tid; // 数组索引
    int sum = 0;
    // 将数组元素拷贝到共享数组中
    while(i < n) {
        sum += array[i];
        i += blockDim.x * gridDim.x;
    }
    sharedArray[tid] = sum;
    __syncthreads();
    // 递归地将共享数组中元素求和
    for(int i = blockDim.x/2; i > 0; i >>= 1) {
        if(tid < i) {
            sharedArray[tid] += sharedArray[tid + i];
        }
        __syncthreads();
    }
    // 线程 0 将共享数组中的结果写入全局内存
    if(tid == 0) {
        *result = sharedArray[0];
    }
}

以上是一种常见的共享内存操作方式，可在开发具有更高效性能的 CUDA 应用中发挥重要的作用。