warp shuffleの関数の中で一番直観的でないのが,
このint __shfl_xor(int value, int laneMask)です.
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を使うと速度が上がるようです。