cuda 学习 | GPU的归约、扫描、直方图算法

---

两种复杂度

• Step complexity
  即步骤复杂度，完成一个工作需要多少步。
• Work complexity
  即工作复杂度，完成工作一共需要的工作量。

对于并行计算，由于可以采取多线程的运算，可以对每一步的运算时间进行很大的缩减。但对于整个程序，有时需要分很多步骤，后续步骤需要等待前面的步骤处理完得到结果才能继续执行。因此有时步骤的复杂度反而决定了整个程序运行的时间。

---

Reduce 归约

归约的输入如下：

1. 一组元素集合
2. 可归约的运算符

什么是可归约的运算符呢？需要满足两个要求：二元性和结合律。
所谓二元性，是运算符是二元运算符，对两个元素进行操作。加减乘除，逻辑与、逻辑或，比较大小都属于二元运算符。结合律是满足：$(a+b)+c=a+(b+c)$的运算符，加减乘是，除法不是。因为$(8/2)/4 = 1 \neq 8/(2/4)=16$。

用简单的函数表示为：reduce[(1,2,3),+]=6。
在上图中，如果是传统的串行计算，那么总的计算量为7，耗时为7；当我们使用并行计算时，可以先把邻近的相加，并重复操作，直到得到最后的输出。这样总的计算量为7，而耗时仅为3，缩短了一大半的时间！

shared memory调用

之前说到的调用内核函数，除了平时串行程序的输入参数，还需要告诉GPU启动的block与thread数。其实还可以告诉GPU shared memory大小，调用时将其作为第三个参数。

shmem_kernel<<<blocks, threads, Mem>>>

在线程内，需要使用

extern __shared__ float sdata[];

来获取相应的空间。

---

Scan 扫描

扫描和归约很像，简单来说扫描是多次的归约。
扫描需要有三个输入：

1. 一组元素集合
2. 可扫描的运算符
3. 标识元素（identity element）

可扫描的运算符与归约的要求相同，也是二元性和结合律。
第三点identity element是指用该元素作为输入，与任何值进行给定的运算符操作，得到的仍是该值。
即：$\text{[I op a =a]}$
对于加法来说，0是标识元素，因为0加上任何值不改变大小；
对于乘法来说，1是标识元素；
对于取小值来说，在unsigned char型中，0xFF是标识元素。

标识元素是第一点扫描与归约不同的地方，第二点在于输出。
归约输出是一个值，而扫描输出是与输入相同大小的一个数组。
具体还可以分为exclusive和inclusive，区别在于输出第一个元素是否为标识元素。
Step complexity 为$O(log(n))$，Work complexity 为$O(n^2)$。

Input	Inclusive output	Exclusive output
$a_0$	$a_0$	$I$
$a_1$	$a_0+a_1$	$a_0$
…	…	…
$a_n$	$a_0+a_1+...+a_n$	$a_0+a_1+...+a_{n-1}$

Hills Steele Scan

Hills Steele Scan 是一种优化Step complexity的算法，第一步是相邻求和放入下一栏，接着间隔1位求和，然后间隔两位求和……以此类推。Step complexity 为$O(log(n))$，Work complexity 为$O(n^2)$。

Step	$n_0$	$n_1$	$n_2$	$n_3$	$n_4$
1	1	2	3	4	5
2	1	3	5	7	9
3	1	3	6	10	14
4	1	3	6	10	15

Blelloch Scan

这是一种优化Work complexity的算法，比上面的要复杂一些。主要分为reduce和downsweep两步。Step complexity 为$O(2log(n))$，Work complexity 为$O(n)$。

downsweep的操作方法如图右下角，L、R为上一行两个输入，R镜像到左下角值不变，右下角为L、R之和。

---

Histogram

直方图在图像中也是经常用到，比如灰度直方图反映了灰度的分布情况，能从整体把握图像的亮暗、对比度信息。
传统的串行直方图统计，需要遍历每一个像素，然后对应的统计灰度值加1，效率较低，而并行计算中有三种算法。

atomic

上节中讲到过这个方法。如果每个线程负责一个像素，独自相加，最大的问题就是内存访问与修改。通过atomicadd可以很好的解决这个问题，但是缺点是速度较慢。

local histogram

每个线程负责一部分的图像区域，无需使用atomic。之后再进行归约操作变成全局直方图。

sort

这种方法目前还没有具体说明，只是了解一下。