增大网格尺寸时CUDA kernel返回错误计算结果的问题咨询

问题根因 #

你遇到的异常是myadd内核存在线程竞态条件导致的：

• 原内核中同一线程块内的所有线程都会执行sum[blockIdx.x] = 0写入操作，该操作和后续的atomicAdd_block累加操作没有任何同步约束。
• 当同块内部分线程已经完成原子累加后，还未执行赋值操作的线程一旦运行到sum[blockIdx.x] = 0，会直接将已累加的结果清零，最终求和结果就会小于正确值。
• 网格尺寸较小时，SM资源充足，线程块能被连贯调度，大概率所有线程先完成赋值再执行累加，竞态触发概率极低；当网格尺寸提升到65535时，SM上的线程上下文切换变得频繁，同块线程的执行时序差变大，竞态极易触发。

修复方案 #

修改myadd内核，仅让块内的0号线程完成对应sum元素的初始化，初始化完成后增加块内同步，保证所有线程累加前看到的是初始化后的0值：

__global__ void myadd(const int *in, int *sum) {
    // 仅块内0号线程执行初始化操作
    if (threadIdx.x == 0) {
        sum[blockIdx.x] = 0;
    }
    // 块内同步，确保所有线程感知到初始化完成
    __syncthreads();
    atomicAdd_block(&sum[blockIdx.x], in[threadIdx.x]);
}

修改后重新编译运行即可解决大网格下的结果错误问题。

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