CUDA 基础算法之reduce、scan、histogram

前言

之前对于CUDA的学习基本上就是不会就查，拿来就用的状况，对一些基础算法的了解不是特别深，之前在面试的时候还有被问到scan扫描算法计算数组的前缀和，表示还没有详细的了解以致只能尴尬地说不清楚，是真的贼尴尬啊，后来去学了些视频课，才逐渐有了一些些基础。（说起来之前还不知道有step complexity 和 work complexity这两种复杂度呢）

1、step and work complexity

step complexity 指可以将一个操作并行成为几个步骤，work complexity 指最终有多少步骤的工作要做；
如下图中所示：

求一数组所有元素之和（所有元素个数为8），其step complexity为3；而work complexity为7.

2、reduce（规约）

规约的大致过程同样如下图所示：

图中所示为规约求和，但是同样求最大最小值之类的同样可以应用规约的思想。如图中所示，如果为串行规约（数组有n个元素），那么step complexity 和work complexity的大小都是O(n-1)；利用并行规约的思想，可以达到如下效果，step complexity：O(logn)；work complexity：O(n-1)；
下面的代码样例是reduce求和的一个基本版本：

__global__ void reduce(int* dev_i, int* dev_o)
{
   
	extern __shared__ int sdata[];
	unsigned int tid = threadIdx.x;
	unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
	sdata[tid] = dev_i[i];
	__syncthreads();
	for (unsigned int s = 1; s < blockDim.x; s *= 2)
	{
   
		int index = 2 * s * tid;
		if (index < blockDim.x)
		{
   
			sdata[index] += sdata[index + s];
		}
		__syncthreads();
	}