了解CUDA SHFL指令

我读过Shuffle Tips and Tricks纸，但我不知道究竟是如何将其应用到一些狡猾的代码，我继承：了解CUDA SHFL指令

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()函数中究竟发生了什么。

来源

2016-09-06 Ken Y-N

看看这个https://devblogs.nvidia.com/parallelforall/faster-平行削减-开普勒/ – Hopobcn

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

__shfl_xor通过与laneMask执行呼叫者的车道ID的按位XOR来计算源极线ID：返回通过将得到的车道ID保持var值。（...）

车道ID是线程的索引的经内，从0到31因此，硬件执行用于每个线程一个按位XOR：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运算符。在步骤：

32个线程将它们的值累加到较低的16个线程中。
16个线程累积到较低的8个线程中。（其他线程对于实际算法无关紧要）
8个线程累积到较低的4个线程中。
4线程acumulate进入下2个线程...

PD：请注意，这两个代码是不完全一样的。这个'32'的偏移告诉我们你的共享内存数组是2 * WARP长。（你正在减少2个* WARP值到1）

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

而洗牌一个降低WARP值到1

来源

2016-09-06 14:53:29 Hopobcn

了解CUDA SHFL指令

回答

相关问题