Cuda min warp reduction 产生竞争条件答案

【问题标题】：Cuda min warp reduction produces race conditionCuda min warp reduction 产生竞争条件
【发布时间】：2021-06-09 16:00:57
【问题描述】：

我有点困惑，我已经使用在线教程中概述的减少翘曲已经有一段时间了，它从来没有引起任何问题。这些是 sn-ps：

while (r < total_rotations){
        rot_index(d_refinements, h_num_refinements, abg,&rot_linear_index, r,s);
        concat[threadIdx.x] = min(concat[threadIdx.x],score_offset[rot_linear_index]);
        r += blockDim.x;
    }
    __syncthreads();
    if (BLOCKSIZE >= 1024){if (tid < 512) { concat[tid] = min(concat[tid],concat[tid + 512]);} __syncthreads();}
    if (BLOCKSIZE >= 512){if (tid < 256) { concat[tid] = min(concat[tid],concat[tid + 256]);} __syncthreads();}
    if (BLOCKSIZE >= 256){if (tid < 128) { concat[tid] = min(concat[tid],concat[tid + 128]);} __syncthreads();}
    if (BLOCKSIZE >= 128){if (tid <  64) { concat[tid] = min(concat[tid],concat[tid + 64]);} __syncthreads();}
    if (tid < 32) min_warp_reduce<float,BLOCKSIZE>(concat,tid); __syncthreads();
    if (tid==0){
        min_offset[0] = concat[0];
    }

还有__device__ 代码。

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]);
    if (blockSize >= 32) sdata[tid] = min(sdata[tid],sdata[tid + 16]);
    if (blockSize >= 16) sdata[tid] = min(sdata[tid],sdata[tid +  8]);
    if (blockSize >=  8) sdata[tid] = min(sdata[tid],sdata[tid +  4]);
    if (blockSize >=  4) sdata[tid] = min(sdata[tid],sdata[tid +  2]);
    if (blockSize >=  2) sdata[tid] = min(sdata[tid],sdata[tid +  1]);
}

对我来说，我已经忠实地复制了教程代码，但是竞态条件检查告诉我有几个冲突。我错过了什么？

【问题讨论】：

标签： cuda reduction

【解决方案1】：

您所指的tutorial 已经很老了，没有考虑到Volta 执行模型。它假设扭曲将保持同步。

volta 执行模型不保证这一点，特别是在存在条件代码的情况下。

您应该能够通过添加__syncwarp() 来解决此问题（消除比赛检查错误）：

__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();}
}

__syncwarp() implies 一个内存屏障，所以如果你愿意，你可以选择移除 volatile 装饰器；但这不是这里讨论的正确性/竞赛所必需的。

如果这不能解决问题，您需要提供mcve

【讨论】：