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


由于表格画不出来,这一篇没有加入表格。完整版

两种复杂度

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

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大小,调用时将其作为第三个参数。

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

在线程内,需要使用

1
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)$。

Hills Steele Scan

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

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

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