多次元Block sizeのWarp idについて

概要

const dim3 block_size(4, 4, 4);
kernel<<<1, block_size>>>(...);

調べ方

結果

dim(4, 4, 4);

$ ./a.out | sort
[ 0, 0, 0] : warpid =  0
[ 0, 0, 1] : warpid =  0
[ 0, 0, 2] : warpid =  1
[ 0, 0, 3] : warpid =  1
[ 0, 1, 0] : warpid =  0
[ 0, 1, 1] : warpid =  0
[ 0, 1, 2] : warpid =  1
[ 0, 1, 3] : warpid =  1
[ 0, 2, 0] : warpid =  0
[ 0, 2, 1] : warpid =  0
[ 0, 2, 2] : warpid =  1
[ 0, 2, 3] : warpid =  1
[ 0, 3, 0] : warpid =  0
[ 0, 3, 1] : warpid =  0
[ 0, 3, 2] : warpid =  1
[ 0, 3, 3] : warpid =  1
[ 1, 0, 0] : warpid =  0
[ 1, 0, 1] : warpid =  0
[ 1, 0, 2] : warpid =  1
[ 1, 0, 3] : warpid =  1
[ 1, 1, 0] : warpid =  0
[ 1, 1, 1] : warpid =  0
[ 1, 1, 2] : warpid =  1
[ 1, 1, 3] : warpid =  1
[ 1, 2, 0] : warpid =  0
[ 1, 2, 1] : warpid =  0
[ 1, 2, 2] : warpid =  1
[ 1, 2, 3] : warpid =  1
[ 1, 3, 0] : warpid =  0
[ 1, 3, 1] : warpid =  0
[ 1, 3, 2] : warpid =  1
[ 1, 3, 3] : warpid =  1
[ 2, 0, 0] : warpid =  0
[ 2, 0, 1] : warpid =  0
[ 2, 0, 2] : warpid =  1
[ 2, 0, 3] : warpid =  1
[ 2, 1, 0] : warpid =  0
[ 2, 1, 1] : warpid =  0
[ 2, 1, 2] : warpid =  1
[ 2, 1, 3] : warpid =  1
[ 2, 2, 0] : warpid =  0
[ 2, 2, 1] : warpid =  0
[ 2, 2, 2] : warpid =  1
[ 2, 2, 3] : warpid =  1
[ 2, 3, 0] : warpid =  0
[ 2, 3, 1] : warpid =  0
[ 2, 3, 2] : warpid =  1
[ 2, 3, 3] : warpid =  1
[ 3, 0, 0] : warpid =  0
[ 3, 0, 1] : warpid =  0
[ 3, 0, 2] : warpid =  1
[ 3, 0, 3] : warpid =  1
[ 3, 1, 0] : warpid =  0
[ 3, 1, 1] : warpid =  0
[ 3, 1, 2] : warpid =  1
[ 3, 1, 3] : warpid =  1
[ 3, 2, 0] : warpid =  0
[ 3, 2, 1] : warpid =  0
[ 3, 2, 2] : warpid =  1
[ 3, 2, 3] : warpid =  1
[ 3, 3, 0] : warpid =  0
[ 3, 3, 1] : warpid =  0
[ 3, 3, 2] : warpid =  1
[ 3, 3, 3] : warpid =  1

dim(8, 4, 4);

少し気持ち悪目に

dim(3, 3, 5);

$ ./a.out | sort 
[ 0, 0, 0] : warpid =  0
[ 0, 0, 1] : warpid =  0
[ 0, 0, 2] : warpid =  0
[ 0, 0, 3] : warpid =  1
[ 0, 0, 4] : warpid =  1
[ 0, 1, 0] : warpid =  0
[ 0, 1, 1] : warpid =  0
[ 0, 1, 2] : warpid =  1
[ 0, 1, 3] : warpid =  1
[ 0, 1, 4] : warpid =  1
[ 0, 2, 0] : warpid =  0
[ 0, 2, 1] : warpid =  0
[ 0, 2, 2] : warpid =  1
[ 0, 2, 3] : warpid =  1
[ 0, 2, 4] : warpid =  2
[ 0, 3, 0] : warpid =  0
[ 0, 3, 1] : warpid =  0
...