了解CUDA shfl指令

Question

我已经阅读了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,这段代码也很狡猾,但看看这个实现,我看到:

__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的按位异或来计算源行ID:返回由结果通道ID保存的var的值.(......)

通道ID是warp中线程的索引,从0到31.因此硬件为每个线程执行按位异或:sourceLaneId XOR laneMask => destinationLaneId

例如,使用线程0和:

__shfl_xor(val, 16)

laneMask = 0b00000000000000000000000000010000 = 16(十进制)

srclaneID = 0b00000000000000000000000000000000 = 0(十进制)

XOR ------------------------------------------------- ---------

dstLaneID = 0b00000000000000000000000000010000 = 16(十进制)

然后线程0获取线程16的值.

现在用线程4:

laneMask = 0b00000000000000000000000000010000 = 16(十进制)

srclaneID = 0b00000000000000000000000000000100 = 4(十进制)

XOR ------------------------------------------------- ---------

dstLaneID = 0b00000000000000000000000000010100 = 20(十进制)

所以线程4获取线程20的值.等等...

如果我们回到实际算法,我们会发现它min是应用运算符的并行减少.在步骤:

1. 32个线程将它们的值累积到较低的16个线程中.
2. 16个螺纹累积到较低的8个螺纹中.(其他线程对于实际算法无关紧要)
3. 8个螺纹累积到较低的4个螺纹中.
4. 4个螺纹累积到下面2个螺纹中......

PD:请注意这两个代码并不完全相同.这个'32'的偏移告诉我们你的共享内存数组是2*WARP长.(您将2*WARP值减少为1)

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

随机播放将WARP值降低为1.