【CUDA】__shfl_xorによるbutterfly exchange

warp shuffleの関数の中で一番直観的でないのが,
このint __shfl_xor(int value, int laneMask)です.

f:id:yusuke_ujitoko:20160209153114p:plain

https://people.maths.ox.ac.uk/gilesm/cuda/lecs/lec4.pdfより引用

がわかりやすくて,指定されたlaneMaskが,どのlaneIDのビットをflipさせるかを決めています.

__shfl_xorを使うと,
warp内で,一定間隔ごとに値群を交換できるbutterfly exchangeが可能です.

__global__ void test_shfl_xor(int *out_dev, int *in_dev, int const mask){
	int value = in_dev[threadIdx.x];
	value = __shfl_xor(value, mask);
	out_dev[threadIdx.x] = value;
}

__shfl_xorを使ったreduction

この__shfl_xorが効力を発揮するのが、reductionです。
通常はshared memoryを使ってreductionを最適化していましたが、
warp shuffleを使うと速度が上がるようです。