【CUDA】Warp shuffleの__shfl()_down, __shfl()_upサンプル

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()も同様。