Cuda min 翘曲减少产生竞争条件

问题描述

我有点困惑，我已经使用在线教程中概述的扭曲减少很长一段时间了，它从未引起过问题。这些是片段：

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

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

解决方法

您所指的 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

cuda cuda reduction