CUDA最小warp归约代码出现竞争条件问题排查求助

问题根因 #

你使用的warp归约代码是基于CUDA 9之前、Volta架构以前的旧假设编写的：旧架构中warp内32个线程是锁步执行的，因此不需要显式的warp内同步逻辑，仅靠volatile关键字保证内存读写不被编译器优化即可正常运行。但从Volta架构（计算能力7.0+）开始，CUDA支持warp内线程独立调度，没有显式同步的情况下，归约的后一步读操作可能读到前一步尚未写入完成的共享内存值，就会被竞争条件检测工具识别为数据冲突。
另外你的代码还存在一处隐含风险：当前你调用完min_warp_reduce后使用__syncthreads()做块同步没问题，但如果后续有warp内其他线程直接读取归约结果的场景，仅靠块同步无法保证warp内的内存可见性。

修复方案 #

在min_warp_reduce函数的每一步归约操作后添加__syncwarp()显式同步warp内所有线程，保证前一步的写操作对所有warp内线程可见即可，修改后的函数实现如下：

template <class T, unsigned int blockSize>
__device__
void min_warp_reduce(volatile T * sdata, int tid){
    if (blockSize >= 64) {
        sdata[tid] = min(sdata[tid],sdata[tid + 32]);
        __syncwarp();
    }
    if (blockSize >= 32) {
        sdata[tid] = min(sdata[tid],sdata[tid + 16]);
        __syncwarp();
    }
    if (blockSize >= 16) {
        sdata[tid] = min(sdata[tid],sdata[tid +  8]);
        __syncwarp();
    }
    if (blockSize >=  8) {
        sdata[tid] = min(sdata[tid],sdata[tid +  4]);
        __syncwarp();
    }
    if (blockSize >=  4) {
        sdata[tid] = min(sdata[tid],sdata[tid +  2]);
        __syncwarp();
    }
    if (blockSize >=  2) {
        sdata[tid] = min(sdata[tid],sdata[tid +  1]);
        __syncwarp();
    }
}

如果你的业务代码只需要兼容Kepler/Maxwell/Pascal等旧架构，也可以在编译时添加-arch=sm_60及更低的架构编译选项，也能规避检测报错，但兼容性不如显式加__syncwarp()的方案。

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