c++ - 了解CUDA shfl指令

我已经阅读了Shuffle Tips and Tricks论文，但是我不确定如何将其应用于继承的一些易懂的代码:

extern __shared__ unsigned int lpSharedMem[];
int tid = threadIdx.x;
lpSharedMem[tid] = startValue;
volatile unsigned int *srt = lpSharedMem;

// ...various stuff
srt[tid] = min( srt[tid], srt[tid+32] );
srt[tid] = min( srt[tid], srt[tid+16] );
srt[tid] = min( srt[tid], srt[tid+8] );
srt[tid] = min( srt[tid], srt[tid+4] );
srt[tid] = min( srt[tid], srt[tid+2] );
srt[tid] = min( srt[tid], srt[tid+1] );
__syncthreads();

即使没有CUDA，此代码也相当狡猾，但查看this implementation，我会看到:

__device__ inline int min_warp(int val) {
    val = min(val, __shfl_xor(val, 16));
    val = min(val, __shfl_xor(val, 8));
    val = min(val, __shfl_xor(val, 4));
    val = min(val, __shfl_xor(val, 2));
    val = min(val, __shfl_xor(val, 1));
    return __shfl(val, 0);
}

此代码可通过以下方式调用:

int minVal = min_warp(startValue);

因此，我可以用上面的代码替换我比较狡猾的volatile。但是，我无法真正理解发生了什么。有人可以解释我是否正确，以及min_warp()函数中到底发生了什么。

最佳答案

根据int __shfl_xor(int var, int laneMask, int width=warpSize);的描述:

通道ID是线程在扭曲中的索引，范围是0到31。因此，硬件会对每个线程执行按位XOR:sourceLaneId XOR laneMask => destinationLaneId

例如，使用线程0和:

__shfl_xor(val, 16)

然后线程0获得线程16的值。

现在使用线程4:

因此，线程4获得了线程20的值。

如果我们回到实际的算法，我们会看到这是应用min运算符的并行约简。在步骤:

32个线程将其值累加到较低的16个线程中。

16个线程累加到较低的8个线程中。 (其他线程与实际算法无关紧要)

8个线程累加到较低的4个线程中。

4个线程累加到较低的2个线程中...

PD:请注意，两个代码并不完全相同。此偏移量“32”告诉我们您的共享内存阵列的长度为2 * WARP。 (您正在将2 * WARP值减少为1)

srt[tid] = min( srt[tid], srt[tid+32] );

洗牌一将WARP值减小为1。

Warp

c++ - 了解CUDA shfl指令