cuda __shfl_xor

__shfl是cuda提供的wrap(线程束)级别的方法,一般线程束为32。相似的方法有几个。

__shfl()

Direct copy from indexed lane

__shfl_up()

Copy from a lane with lower ID relative to caller

__shfl_down()

Copy from a lane with higher ID relative to caller

__shfl_xor()

Copy from a lane based on bitwise XOR of own lane ID

__shfl_up和__shfl_down均有var和delta参数

var是将该值与当前的lane id(lane id与tid有些关系，但不一定相同)进行绑定,

delta,在__shfl_up是将当前的lane id - delta的result lane id中的var值返回回来。

delta,在__shfl_down是将当前的lane id + delta的result lane id中的var值返回回来。

另注明：__shfl()默认是线程束级别的同步操作，因此不需要使用__syncthreads(),块级别的同步操作。

__shfl_xor中有var和laneMask,var同样是绑定当前的lane id,而laneMask则是与当前的laneId进行异或操作。

以__shfl_xor进行例子举例。

var+=__shfl_xor(var,16);

var+=__shfl_xor(var,8);

var+=__shfl_xor(var,4);

var+=__shfl_xor(var,2);

var+=__shfl_xor(var,1);

作用可以用下图说明

0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31

1 1 1 1 1 1 1 1 1 1 1    1   1   1   1   1   1   1   1   1   1   1   1   1   1   1   1   1   1   1   1   1

2 2 2 2 2 2 2 2 2 2 2    2   2   2   2   2   2                                                                                  var += __shfl_xor(var, 16)

4 4 4 4 4 4 4 4 4                                                                                                                       var += __shfl_xor(var, 8)

8 8 8 8                                                                                                                                      var += __shfl_xor(var, 4)

16 16                                                                                                                                        var += __shfl_xor(var, 2)

32                                                                                                                                             var += __shfl_xor(var, 1)

因此可以作为统计wrap内中线程数该数据的结果和。