warp shuffleの__shfl_up()や__shfl_down()は、指定オフセット分だけずれたスレッドから値を読み出します。
#include <stdio.h> __global__ void shfl_test(int arg){ int laneId = threadIdx.x & 0x1f; int value; if(laneId == 3){ value = arg; }else{ value = 0; } value = __shfl_up(value, 1); printf("thread:%d, value: %d\n", threadIdx.x, value); } int main(){ shfl_test<<<1, 32>>>(1234); return 0; }
__shfl_down()も同様。