CUDA列求和核函数结果异常，求线程执行逻辑解析

问题分析与解决 #

核心问题：数据竞争（Race Condition） #

你的代码中，多个线程同时对d_sum的同一个元素执行+=操作，而这个操作并非原子性的。+=会拆分为读取当前值→计算新值→写回新值三个步骤，当多个线程并行执行这三步时，会出现线程A读取值后，线程B也读取到相同的旧值，两者分别计算后写回，最终只保留了其中一个线程的修改，导致累加结果缺失，且因每次执行的线程时序不同，结果完全随机。

实际执行逻辑 #

<<<2, 1024>>>表示启动2个线程块，每个线程块包含1024个线程，总共有2048个线程并行执行：

• 第一个线程块（blockIdx.x=0）的1024个线程（threadIdx.x从0到1023），都会执行d_sum[0] += d_a[threadIdx.x*2 + 0]，即对矩阵第一列的所有元素做累加；
• 第二个线程块（blockIdx.x=1）的1024个线程，都会执行d_sum[1] += d_a[threadIdx.x*2 + 1]，即对矩阵第二列的所有元素做累加；
• GPU线程的执行是高度并行的，同一时间可能有数百个线程同时读写d_sum[0]或d_sum[1]，没有同步机制的情况下，必然出现数据竞争，导致结果错误且不稳定。

修复方案 #

方案1：原子操作（简单直接，适合小规模数据）

将普通的+=替换为CUDA提供的atomicAdd函数，它能保证整个累加过程是原子性的，不会被其他线程打断：

__global__ void sumColumns(float* d_a, float* d_sum, int n) {
    if (blockIdx.x < 2 && threadIdx.x < n) {
        atomicAdd(&d_sum[blockIdx.x], d_a[threadIdx.x * 2 + blockIdx.x]);
    }
}

方案2：归约（Reduction，适合大规模数据，效率更高）

如果矩阵规模更大，原子操作的性能会下降，这时可以先让每个线程块内部通过共享内存完成局部求和，再将块内结果汇总到全局内存：

__global__ void sumColumns(float* d_a, float* d_sum, int n) {
    __shared__ float s_sum[2];
    int col = blockIdx.x;
    int idx = threadIdx.x;

    // 初始化共享内存
    if (idx == 0) {
        s_sum[col] = 0.0f;
    }
    __syncthreads();

    // 每个线程累加一个元素到共享内存
    if (idx < n) {
        atomicAdd(&s_sum[col], d_a[idx * 2 + col]);
    }
    __syncthreads();

    // 把块内总和写入全局内存
    if (idx == 0) {
        d_sum[col] = s_sum[col];
    }
}

验证结果 #

修改后，第一列全1的累加结果会稳定为1024，第二列全-1的累加结果会稳定为-1024，符合预期。

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