理解CUDA shfl教学

Question

我读过洗牌小贴士的文章，但我不知道如何将它应用到我继承的一些不可靠的代码中：

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

即使没有数据自动化系统，这个代码也很不可靠，但是我看到了这一实现：

__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()函数到底发生了什么。

Answer 1

从int __shfl_xor(int var, int laneMask, int width=warpSize);的描述

__shfl_xor()通过使用laneMask按位执行调用方的车道ID的XOR来计算源行ID :返回由此产生的车道ID所持有的var值。(...)

车道ID是从0到31的翘曲范围内线程的索引。因此，硬件对每个线程执行位异或: sourceLaneId XOR laneMask => destinationLaneId

例如，使用线程0和：

__shfl_xor(val, 16)

laneMask =0b0000000000000000000000000000010000= 16 (小数) srclaneID = 0b00000000000000000000000000000000 =0(小数) dstLaneID =0b0000000000000000000000000000010000= 16 (小数)

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

现在使用线程4：

laneMask =0b0000000000000000000000000000010000= 16 (小数) srclaneID =0b000000000000000000000000000000000000100100=4(小数) dstLaneID =0b00000000000000000000000000000000000010100= 20 (小数)

所以线程4得到线程20的值。

如果我们回到实际的算法中，我们会发现它是一个并行约简，其中应用了一个min算子。按步骤：

1. 32个线程将它们的值累加到下面的16个线程中。
2. 16个线程累加成下面的8个线程。(其他线程与实际算法无关)
3. 8个线程累积成较低的4个线程。
4. 4条线累积成下面的2条线.

PD:请注意，这两种代码并不完全相同。这个偏移量'32‘告诉我们，您的共享内存数组是2*翘曲长。(您正在将2*翘曲值降为1)

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

而洗牌一则将翘曲值降为1。

Answer 2

在这些幻灯片上有一个直观的图表，第27页：

http://on-demand.gputechconf.com/gtc/2017/presentation/s7622-Kyrylo-perelygin-robust-and-scalable-cuda.pdf

xor 1第一次交换奇数/偶数输入

xor 2然后交换以前交换的输入的奇偶对。

xor 4然后交换奇数/偶数四元

等。

另一种方法是计算完整的xor掩码和原始输入值。有许多方法可以用不同的经纱洗牌指令来完成这一点。这种xor模式也用于FFT/DFT变换，也称为“蝶形图”。

图解