# CUDA系列学习（五）GPU基础算法: Reduce, Scan, Histogram

———-

## 1. Task complexity

task complexity包括step complexity（可以并行成几个操作） & work complexity（总共有多少个工作要做）。
e.g. 下面的tree-structure图中每个节点表示一个操作数，每条边表示一个操作，同层edge表示相同操作，问该图表示的task的step complexity & work complexity分别是多少。

Ans:
step complexity: 3；
work complexity: 6。

## 2. Reduce

1) 最简单的顺序执行顺序组织为((1+2)+3)+4…
2) 由于operation之间没有依赖关系，我们可以用Reduce简化操作，它可以减少serial implementation的步数。

### 2.1 what is reduce?

Reduce input:

1. set of elements
2. reduction operation
1. binary: 两个输入一个输出
2. 操作满足结合律： (a@b)@c = a@(b@c), 其中@表示operator
e.g +, 按位与 都符合；a^b(expotentiation)和减法都不是

### 2.1.1 Serial implementation of Reduce:

reduce的每一步操作都依赖于其前一个操作的结果。比如对于前面那个例子，n个数相加，work complexity 和 step complexity都是O(n)（原因不言自明吧~）我们的目标就是并行化操作，降下来step complexity. e.g add serial reduce -> parallel reduce。

### 2.1.2 Parallel implementation of Reduce:

__global__ void parallel_reduce_kernel(float *d_out, float* d_in){
int myID = threadIdx.x + blockIdx.x * blockDim.x;
int tid = threadIdx.x;

//divide threads into two parts according to threadID, and add the right part to the left one, lead to reducing half elements, called an iteration; iterate until left only one element
for(unsigned int s = blockDim.x / 2 ; s>0; s>>=1){
if(tid<s){
d_in[myID] += d_in[myID + s];
}
__syncthreads(); //ensure all adds at one iteration are done
}
if (tid == 0){
d_out[blockIdx.x] = d_in[myId];
}
}

Quiz: 看一下上面的code可以从哪里进行优化？

Ans：我们在上一讲中提到了global，shared & local memory的速度，那么这里对于global memory的操作可以更改为shared memory，从而进行提速：

__global__ void parallel_shared_reduce_kernel(float *d_out, float* d_in){
int myID = threadIdx.x + blockIdx.x * blockDim.x;
int tid = threadIdx.x;
extern __shared__ float sdata[];
sdata[tid] = d_in[myID];
__syncthreads();

//divide threads into two parts according to threadID, and add the right part to the left one, lead to reducing half elements, called an iteration; iterate until left only one element
for(unsigned int s = blockDim.x / 2 ; s>0; s>>=1){
if(tid<s){
sdata[tid] += sdata[tid + s];
}
__syncthreads(); //ensure all adds at one iteration are done
}
if (tid == 0){
d_out[blockIdx.x] = sdata[myId];
}
}

kernel<<<grid of blocks, block of threads, shmem>>>

parallel_reduce_kernel<<<blocks, threads, threads*sizeof(float)>>>(data_out, data_in);

parallel_reduce_kernel
 Times Read Ops Write Ops 1 1024 512 2 512 256 3 256 128 … n 1 1
parallel_shared_reduce_kernel
 Times Read Ops Write Ops 1 1024 1

## 3. Scan

### 3.1 what is scan?

• Example:

• input: 1,2,3,4
• operation: Add
• ouput: 1,3,6,10（out[i]=sum(in[0:i])）
• 目的：解决难以并行的问题

Inputs to scan:

1. input array
2. 操作：binary & 满足结合律（和reduce一样）
3. identity element [I op a = a], 其中I 是identity element
quiz: what is the identity for 加法，乘法，逻辑与，逻辑或？
Ans：
 op Identity 加法 0 乘法 1 逻辑或|| False 逻辑与&& True

### 3.2 what scan does?

I/O content
input [a0$a_0$ a1$a_1$ a2$a_2$ an$a_n$]
output [I$I$ a0$a_0$ a0a1$a_0\bigotimes a_1$ a0a1$a_0\bigotimes a_1\bigotimes$an$\bigotimes a_n$]

#### 3.2.1 Serial implementation of Scan

int acc = identity;
for(i=0;i<elements.length();i++){
acc = acc op elements[i];
out[i] = acc;
}

work complexity: O(n)$O(n)$
step complexity: O(n)$O(n)$

#### 3.2.1 Parallel implementation of Scan

Q: 那么问题的work complexity和step complexity分别变为多少了呢？
Ans:

• step complexity:
取决于n个reduction中耗时最长的，即O(log2n)$O(log_2n)$
• work complexity:
对于每个output元素进行计算，总计算量为0+1+2+…+(n-1)，所以复杂度为O(n2)$O(n^2)$.

 more step efficiency more work efficiency hillis + steele （1986） √ blelloch （1990） √

1. Hillis + Steele

对于Scan加法问题，hillis+steele算法的解决方案如下：

step 0: out[i] = in[i] + in[i-1];
step 1: out[i] = in[i] + in[i-2];
step 2: out[i] = in[i] + in[i-4];

Q: hillis + steele算法的work complexity 和 step complexity分别为多少？

 log(n)$log(n)$ O(n‾‾√)$O(\sqrt n)$ O(n)$O(n)$ O(nlogn)$O(nlogn)$ O(n^2) work complexity √ step complexity √

1. step complexity：
因为第i个step的结果为上一步输出作为in, out[idx] = in[idx] + in[idx - 2^i], 所以step complexity = O(log(n))$O(log(n))$
2. work complexity:
workload = (n1)+(n2)+(n4)+...$(n-1) + (n-2)+ (n-4)+ ...$ ，共有log(n)$log(n)$项元素相加，所以可以近似看做一个矩阵，对应上图，长log(n)$log(n)$, 宽n，所以复杂度为 nlog(n)$nlog(n)$

2 .Blelloch

1. reduce部分：
每个step对相邻两个元素进行求和，但是每个元素在input中只出现一次，即window size=2, step = 2的求和。
Q: reduce部分的step complexity 和 work complexity？
Ans：

 log(n)$log(n)$ O(n‾‾√)$O(\sqrt n)$ O(n)$O(n)$ O(nlogn)$O(nlogn)$ O(n^2) work complexity √ step complexity √

我们依然将答案用白色标出，请选中看答案。

2. downsweep部分：
简单地说，downsweep部分的输入元素是reduce部分镜面反射的结果，对于每一组输入in1 & in2有两个输出，左边输出out1 = in2，右边输出out2 = in1 op in2 （这里的op就是reduce部分的op），如图：

 36 10 3, 11 1, 3, 5, 7

Q: downsweep部分的step complexity 和 work complexity？
And：downsweep是reduce部分的mirror，所以当然和reduce部分的complexity都一样啦。

ANS：这取决于所用的GPU，问题规模，以及实现时的优化方法。这一边是一个不断变化的问题：一开始我们有很多data（work > processor）, 更适合用work efficient parallel algorithm (e.g Blelloch), 随着程序运行，工作量被减少了（processor > work），适合改用step efficient parallel algorithm，这样而后数据又多起来啦，于是我们又适合用work efficient parallel algorithm…

 serial Hillis + Steele Blelloch work O(n) O(nlogn) O(n) step n log(n) 2*log(n) 512个元素的vector512个processor √ 一百万的vector512个processor √ 128k的vector1个processor √

## 4. Histogram

### 4.2. Histogram 的 Serial 实现：

for(i = 0; i < bin.count; i++)
res[i] = 0;
for(i = 0; i<nElements; i++)
res[computeBin(i)] ++;

### 4.3. Histogram 的 Parallel 实现：

1. 直接实现：

kernel:

__global__ void naive_histo(int* d_bins, const int* d_in, const in BIN_COUNT){
int myID = threadIdx.x + blockDim.x * blockIdx.x;
int myItem = d_in[myID];
int myBin = myItem % BIN_COUNT;
d_bins[myBin]++;
}

atomicAdd(&(d_bins[myBin]), 1);

 key 2 1 1 2 1 0 2 2 value 1 1 1 1 1 1 1 1

 key 0 1 1 1 2 2 2 2 value 1 1 1 1 1 1 1 1

1. atomics
2. per_thread histogram, then reduce
3. sort, then reduce by key

## 5. 总结：

02-15 3万+

05-07 5万+

10-10 4万+

06-04 2万+

12-30 850

03-01 21万+

12-15 7388

04-12 7795

#### CUDA编程入门----Thrust库简介

©️2020 CSDN 皮肤主题: 大白 设计师: CSDN官方博客

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、C币套餐、付费专栏及课程。