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

原创 2017年04月05日 11:20:12

两种复杂度

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

complexity

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


Reduce 归约

归约的输入如下:

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

什么是可归约的运算符呢?需要满足两个要求:二元性和结合律。
所谓二元性,是运算符是二元运算符,对两个元素进行操作。加减乘除,逻辑与、逻辑或,比较大小都属于二元运算符。结合律是满足:(a+b)+c=a+(b+c)的运算符,加减乘是,除法不是。因为(8/2)/4=18/(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是指用该元素作为输入,与任何值进行给定的运算符操作,得到的仍是该值。
即:[I op a =a]
对于加法来说,0是标识元素,因为0加上任何值不改变大小;
对于乘法来说,1是标识元素;
对于取小值来说,在unsigned char型中,0xFF是标识元素。

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

Input Inclusive output Exclusive output
a0 a0 I
a1 a0+a1 a0
an a0+a1+...+an a0+a1+...+an1

Hills Steele Scan

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

Step n0 n1 n2 n3 n4
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

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

版权声明:本文为博主原创文章,转载请标注出处。

CUDA中并行规约(Parallel Reduction)的优化

Parallel Reduction是NVIDIA-CUDA自带的例子,也几乎是所有CUDA学习者的的必看算法。在这个算法的优化中,Mark Harris为我们实现了7种不同的优化版本,将Bandwi...
  • redline2005
  • redline2005
  • 2013年12月20日 16:49
  • 4300

CUDA学习(归约算法)

parallel reduction Nvidia官网归约ppt 中文ppt 可以理解为将数组中所有数求和的过程并行化 CUDA本身并不支持全局同步,将每一层归约作为一个kernel重复递归调用...
  • tian_ciomp
  • tian_ciomp
  • 2016年12月28日 21:12
  • 852

深入理解CUDA点积运算

本博客主要讲述了《GPU高性能编程CUDA实战》这本书中关于点积运算中难懂的部分。...
  • github_30605157
  • github_30605157
  • 2015年09月03日 19:42
  • 1213

OpenCL的学习---计算直方图的理解

看到《OpenCL编程指南》第14章---计算直方图,有点难理解,我对内存中抽象的东西。所以kernel函数那里看了很久。感谢北邮的大神 http://www.mrobotit.cn/~shanxin...
  • wd1603926823
  • wd1603926823
  • 2017年04月14日 14:24
  • 1257

[菜鸟每天来段CUDA_C] GPU上实现直方图计算

本文通过在GPU上计算直方图说明GPU计算中的原子操作。
  • jonny_super
  • jonny_super
  • 2014年02月23日 18:38
  • 1963

我的CUDA学习笔记

研究生阶段开始接触CUDA。这里要感谢我的老板,感谢他的眼光和对我的信任,把这副担子交给我,我不想让他失望。在学习中慢慢陶醉于CUDA的神奇,被它强大的power所折服,这个过程实在太美妙,让我有很强...
  • houjian2015
  • houjian2015
  • 2015年04月07日 18:47
  • 180

《GPU高性能计算之CUDA》书中实例源代码

  • 2010年06月29日 00:38
  • 10.5MB
  • 下载

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

两种复杂度 Step complexity 即步骤复杂度,完成一个工作需要多少步。 Work complexity 即工作复杂度,完成工作一共需要的工作量。 对于并行计算,由于可以采取多线程的运算...
  • Yan_Joy
  • Yan_Joy
  • 2017年04月05日 11:20
  • 761

CUDA实现的直方图均衡化算法

CUDA实现的直方图均衡化算法因做毕设接触到CUDA这个东西,于是就开始了一段漫长的学习CUDA的过程!关于直方图均衡化算法,也是一时兴起想实现一下。最开始是看CUDA的samples里面有一个his...
  • u010950959
  • u010950959
  • 2016年12月28日 20:23
  • 559

《GPU高性能编程 CUDA实战》(CUDA By Example)读书笔记

入门书籍,十分简单,看完就可以编写一些简单的CUDA程序了
  • FishSeeker
  • FishSeeker
  • 2017年07月13日 21:38
  • 1370
内容举报
返回顶部
收藏助手
不良信息举报
您举报文章:cuda 学习 | GPU的归约、扫描、直方图算法
举报原因:
原因补充:

(最多只允许输入30个字)