Cuda最小偏差减少会产生竞争条件

Question

我有点困惑，我已经使用了在线教程中概述的翘曲减少方法很长一段时间了，它从来没有引起过问题。以下是代码片段：

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]);
}

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

Answer 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