cuda的Shuffle技术以及自定义双精度版本

还是数组求和问题引起的，发现之前那个版本http://blog.youkuaiyun.com/lingerlanlan/article/details/24630511

对于数组的维度是有要求的。因为归约每次变为一半，所以对于线程块的数量和每个线程块线程的数量都要是2的倍数。

今天看到这篇文章https://devblogs.nvidia.com/parallelforall/faster-parallel-reductions-kepler/。

对并行归约进行了讨论。目前还没完全读懂，读懂了翻译一下。

现在对刚了解的shuffle技术写一下体会。

这玩意就是使得线程束内的线程可以共享寄存器变量。

比如函数

int __shfl_down(int var, unsigned int delta, int width=warpSize);

有点像在线程间左移变量。

下面用具体例子来说明，

int i = threadIdx.x % 32;
int j = __shfl_down(i, 2, 8);

这里32指一个线程束的线程数量是32

第一句：

int i = threadIdx.x % 32;

每个线程都有一个变量i，即是线程在所在线程束的id。

第二句：

int j = __shfl_down(i, 2, 8);

首先8指明了范围，就是0-7，8-15，16-23，24-31。

2指明了步长。比如i=5的线程，把i值赋值给了i=3的线程中的j变量。本质上就是在一定范围内线程间按照一定的步长来访问另一格线程的寄存器变量。

这幅图很好的说明了

测试例子：

#include <stdio.h>

__global__ void kernel()
{
	int i = threadIdx.x % 32;
	int j = __shfl_down(i, 2, 8);
	printf("%d:%d\n",i,j);
}

int main()
{

	kernel<<<1,32>>>();
	cudaDeviceSynchronize();

return 0;
}

输出结果：

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

注意红色的部分，因为参数8指明了执行范围。

因为库指提供了int和float的shuffle版本，http://docs.nvidia.com/cuda/cuda-c-programming-guide/#warp-shuffle-functions。

双精度的需要自己实现

__device__ inline
double __shfl_down(double var, unsigned int srcLane, int width=32) {
  int2 a = *reinterpret_cast<int2*>(&var);
  a.x = __shfl_down(a.x, srcLane, width);
  a.y = __shfl_down(a.y, srcLane, width);
  return *reinterpret_cast<double*>(&a);
}

这个很巧妙的。用两个32位的int来跟64位的double转换。

其实理解这个，关键是要彻底明白计算机存储数据就是若干个0和1。

而这里巧妙的另外一个地方是用到了

reinterpret_cast函数来强制转换。

这让我想起了曾经面试qq后台开发经历，貌似就是实现两个很大整数数的相加，具体多少位忘了，反正超过32位。

应该就是这种思路。

参考资料：

https://devblogs.nvidia.com/parallelforall/faster-parallel-reductions-kepler/