您的当前位置:首页正文

基于GPU的数学形态学运算并行加速研究

2022-01-01 来源:步旅网
第19卷 第l9期 电子设计工程 2011年10月 Vo1.19 No.19 Electronic Design Engineering Oct.2011 基于GPU的数学形态学运算并行加速研究 张聪,邢同举,罗颖,张静,孙强 (电子科技大学光电信息学院,四川成都610054) 摘要:数学形态学运算是一种高度并行的运算,其计算量大而又如此广泛地应用于对实时性要求较高的诸多重要领 域 为了提高数学形态学运算的速度,提出了一种基于CUDA架构的GPU并行数学形态学运算。文章详细描述了 GPU硬件架构和CUDA编程模型.并给出了GPU腐蚀并行运算的详细实现过程以及编程过程中为充分利用GPU资 源所需要注意的具体问题。实验结果表明,GPU并行数学形态学运算速度可达到几个数量级的提高。 关键词:数学形态学;腐蚀;GPU;CUDA;加速比 中圈分类号:TP391.7 文献标识码:A 文章编号:1674—6236(2011)19—0141—03 Parallel accelerating research about mathematical morphology based on GPU ZHANG Cong,XING Tong-ju,LUO Ying,ZHANG Jing,SUN Qiang (School ofopto—Electronic Information,University foElectronic Science and Technology foChina, Chengdu 610054,China) AbstracU Mathematical morphology operation is a highly parallel processing,it involves a large amount of computation and is widely used for so many important filed that have high requirements on real-time.In order to improve the speed of mathematical morphological operations,a GPU Parallel mathematical morphology operations based on CUDA is proposed in this paper.We have made a description of the GPU hardware architecture and the CUDA programming mode1.The speciifc GPU parallel implementation process of erosion operation and some problems that involves make full use of GPU resources are  ̄iven in this paper too.Experimentla results show that,GPU parallel computing mathematical morphology is several orders of magnitude faster than normal morphology operation. Key words:mathematical morphology;erosion;GPU;CUDA;speedup 数学形态学的应用几乎遍及计算机图形图像处理的所 有方面.包括图像滤波、图像分割分类、图像测量、模式识别 1膨胀腐蚀的基本理论 以及纹理分析与合成,其应用还涉及遥感监测、工业自动化 正文内容腐蚀运算:构造一个结构元素,结构元素的原 检测测量、生物医学影像、图像压缩、军事、航空航天等众多 点定位于待处理的目标元素上,结构元素完全包含在目标区 领域。在运用数学形态学处理气象预报、工程计算和模拟等 域中时,则结构元素的原点所对应的目标像素点被保留.负 问题时,需要对巨大数据量进行多次重复计算,这些计算又 责被腐蚀掉,置为背景点。膨胀运算:构造一个结构元素,结 必须在有限的时间内完成,比如在几秒或几分钟内完成。这 构元素的原点定位在待处理的目标元素上。当有目标点被结 就对数学形态学在处理这些问题时的运算速度提出了较高 构元素覆盖时该点被膨胀为目标点。 的要求。在数学形态学中,腐蚀和膨胀是最基本的两个运算, 膨胀腐蚀运算中存在一些重要的运算性质,在设计算法 这两种对偶运算好比形态学字母表中的字母,其他运算都可 时使用这些性质往往能简化处理过程和提高运算速度。对偶 以根据这两种运算表述Ⅲ,因此形态学处理速度决定于膨胀、 性便是其中的一个重要性质,对偶性可用下面公式来表示: 腐蚀处理速度。 (XOB)C=Xq ̄)B;(X ̄B)C=XCOB。这两个公式表明,图像膨胀 人们通过在腐蚀、膨胀处理中采用改进算法来提高其运 等价于该图像的补被相同结构元素腐蚀所得结果的补.反之 算速度嘲。但其速度也只能提高几倍.这在某些大数据量、实 亦然 。在本文中讨论主要集中在腐蚀运算的并行加速.而膨 时性要求较高的应用中仍不能满足需要。针对这种情况,文 胀运算可以通过待处理图像补的腐蚀运算求得。 章中提出了通过图形处理器(GPU)并行处理的方式对数学 形态学运算进行加速。实验结果表明.通过GPU对数学形态 2 GPU简介 学运算进行加速,其运算速度可达到1—2个数量级的提高。 2.1 GPIJ硬件架构 收稿日期:2011-07—29 稿件编号:201107127 与CPU相比,GPU有着很高的计算吞吐量,这主要在于 作者简介:张聪(1986一),男,重庆人,硕士。研究方向:信号与信号处理。 .141—— 《电子设计工程))2011年第19期 GPU和CPU不同的设计理念。GPU最初设计就是为了图形 处理,由于图像渲染的高度并行性,GPU设计者通过增加数 学逻辑单元(ALU)和控制单元(Contro1)的方式提高处理能力 和存储器带宽。这样就在GPU硬件架构中形成了众多的流 处理器,Nvidia GTX 280的架构 如图l所示。 图1中:Nvidia GTX 280中有3O个完整前端流多处理器 图1 NvidiaGTX 280架构 Fig. 1 Nvidia GTX 280 architecture (SM),每个流多处理器包含8个流处理器(SP)。每个SP又可 看成一个SIMD处理器。每个SM拥有独立的完整前端,包括 取指、译码、发射和执行单元等,但每两个SM共享一条存储 器流水线。SP不具有独立的处理器核,它们有独立的寄存器 和指令指针,但没有取指和调度单元构成的完整前端。 2.2 CUDA软件开发环境 由于GPU具有很高的计算吞吐量,利用GPU进行通用 计算已经引起了人们的关注 。CUDA架构正是专门为了 GPU通用计算而设计的一种全新模块,它使得开发人员无需 再过多地考虑严格的资源限制和编程限制 CUDA C是在 CUDA GPU上的编程语言.CUDA C语言本质上是对C的简 单扩展,它的出现使得开发人员不需要学习计算机图形学和 着色语言,就可以编写出高效的GPU并行程序。 CUDA编程模型[91将CPU作为主机(Host)。GPU作为设 备(Device)[10J。在这个模型中GPU和CPU协同工作,CPU上 负责逻辑性较强的事务处理和串行计算,GPU则专注于执行 高度并行化处理任务。运行在GPU上的CUDA函数称为内 核函数(kerne1),一个完整的CUDA程序是由Host串行代码 和Device上的kernel函数共同组成的,如图2所示。 运行在设备上的kernel函数定义和普通C语言函数基 图2 CUDA编程模型 Fig.2 CUDA programming structure 模板中心对准当前待处理像素点,查看被模板覆盖的所有像 素点,如果全为目标点则当前待处理像素点保留为目标点, 否则腐蚀掉。腐蚀可进一步转化为邻域像素点灰度值求和的 问题,每一个像素点都由被模板覆盖所有像素点灰度值的和 决定。由此,确定每一点是否被腐蚀掉,所要执行的操作是相 同的,只是所处理的数据不同而已。这是典型的单指令多数 据问题,具有很好的并行性。 3.1 实现时需要注意问题 本一致,当需要被主机调用时,只需在函数定义前面加上 lobal_修饰符___global void Copy char src,char dst)。在主机调用kernel函数时,需要传递额外的参数,如: Copy<<<blocksPerGrid,threadsPerBlock>>>(ptr_sre,ptr dst)。 kernel函数是以线程网格(Grid)的形式组织的,每个Grid由 在GPU上编码实现腐蚀运算时,为了正确地执行和尽可 若干个线程块(block)组成,每个block由若干线程(thread)组 成 调用kernel函数所需传递的额外参数blocksPerGrid和 threadsPelrBlock分别是Grid和block的维度设计。 能地提高执行效率需要注意几个问题:1)避免当前处理的像 素点被已处理像素点所干扰。在具体实现时,生成一幅和原 图同样大小的新图,每个线程将处理结果存放在这幅新图 3基于GPU的腐蚀运算加速实现 图像腐蚀采用(2n+1)x(2n+1)且系数全为1的模板,即 中,而运算数据则从原图中获取:2)充分利用GPU硬件资源, 尽可能为每个像素点都开辟一个线程.目前GPU每个block 允许开辟的线程数量已达到1 024个.每个Grid允许每一维 进行n尺寸腐蚀。要确定某一像素点是否被腐蚀掉,只需将 一开辟的block数量为65 535。所以对几乎所有的二值图像,都 142一 张聪.等基于GPU的数学形态学运算并行加速研究 能实现为每一个像:泰点开辟一个线程;3)腐蚀运算访问内存 模式具有很强的空间局部性,GPU纹理缓存是专门为了加速 这种访问模式而设计的,在这种内存访问模式下,使用纹理 内存能减少内存流量是性能得到提高。4)block维度设计时, 每个block的线程数量应该为64的倍数f ,这样GPU的性能 才能最大限度地提高。 3.2实现过程 上面已经提到CUDA编程模型是GPU和CPU协同工作 的.所以GPU腐蚀:匿算的实现过程是CPU和GPU共同完成 的。CPU负责内存分配、Grid和block的维度设定、启动在 GPU上运行的kern el函数并向其传递参数、最后从GPU获 取处理结果。GPU负责大量并行运算,并将处理数据保存下 来。其具体显现步骤如下: 1)读取待处理图像,根据图像大小在device上开辟显存 devSrc,并将图像数据复制到这块显存上,并将其绑定二维纹 理内存: 2)在device开辟和devSrc同样大小的显存devDst,用于 GPU保存处理结果。在host上开辟和devDst同样大小的内 存hostDst.用于接收GPU处理的结果; 3)根据待处理图像大小,设计 d和block的维度。为了 提高性能将block维度固定为dim3 blockDim(16,16),则grid的 维度则根据图像宽度width、高度height和block的维度设 定,可设定为dim3 gridDim((width+15)/16,(height+15)/16); 4)启动kernel函数,并将设定好的参数传递到kernel函 数中: 5)kernel函数:执行完毕,将处理结果传回host。host保存 并显示结果。 4实验结果及分析 文章所采用实验平台为:Interl(R)Xeon(R)CPU 5160, 主频2.99 GHz.Windows7操作系统。显卡为GForce GTX 570。文章中所处理的为1 024xl 024、7 350x5 700的图像。表 格l为在对大小为1 024xl 024的图像进行不同尺寸腐蚀运 算时,CPU串行腐蚀和GPU并行腐蚀所耗用时间对比。 表1在: 同腐蚀尺寸下CPU和GPU腐蚀时间对比 Tab.1 The comlmrison of the erosion time between CPU and G PU in diferent erosion sizes 从表格1中可看出,随着腐蚀尺寸的增大,即随着运算 时间复杂度的增加,CPU串行腐蚀运算所耗用时间大幅度增 加,GPU并行腐蚀运算所耗用时间没有大幅增加.GPU并行 运算加速比越来越大。 在图3中,可以看到在对图像进行同一腐蚀尺寸的腐蚀 处理时,GPU腐蚀处理大小为7 350 ̄5 700的图像比处理 l 024x1 024的图像加速比要高。当腐蚀尺寸增加到6时,对 于1 024xl 024的图像加速比可达到80左右,对于7 350x 5 700的图像加速比可达150左右。GPU并行运算之所以有 这么高的加速效果,主要因为它有如此多的数学逻辑单元, 所以能同时运行较多的线程。它通过大量线程的切换来隐藏 访存延迟,当某个线程被阻塞了,马上切换到其他线程进行 处理.所以当待处理数据规模越大时,GPU资源就越能被充 分利用,加速效果就更好。 腐蚀尺寸 图3对不同大小图片.GPU加速比随着腐蚀尺寸的变化而变化 Fig.3 In different size of images,the speedup of GPU changes when the erosion size changes 5 结 论 通过GPU并行处理的方式,使形态学运算速度大幅提 高。实验结果表明,随着待处理数据规模的增加和处理复杂 度的增加,GPU并行腐蚀运算的加速效果显著增加。在本实 验平台上GPU并行腐蚀运算可达到150倍的性能提升.目 前,还未见任何优化算法达到如此高的加速效果的报道。本 文还在支持CUDA的较低档的型号为GrI22O的显卡上进行 了测试,在其上进行并行腐蚀运算也可达到10倍以上的性 能提升。可见,在数学形态学运算时,同时考虑性能和成本两 方面,GPU相对于CPU都有绝对的优势旧。所以,在涉及数学 形态学运算的领域,GPU并行运算是一个必然趋势。 参考文献: 【1】Solille P.形态学图像分析原理与应用【M】.北京:清华大学 出版社.2008:47—48. 【2]陆宗骐,朱煜.数学形态学腐蚀膨胀运算的快速算法【C】,,第 十三届全国图象图形学学术会议,2006一l1—06,2006: 306—31 1. 【3】孙即祥.图像分析【M】.北京:科学出版社,2005:183—184. 【4]Garland M,Grand S L,Niekolls J,et a1.Parallel computing experiences with CUDA[JI.Micro,IEEE,2008,28(4):13—27. [5 Ni5]ckolls J,Dally W J.1’Ile GPU computing era[J1.IEEE Micro,2010,30(2):56—69. [6 Hawi6]ck K A,Leist A,Playne D P.Parallel graph component labelling with GPUs and CUDA[J].Parlalel Computing, 2010,36(12):655—678. (下转第146页) 一】43— 《电子设计工程}2011年第19期 为高电平,并且在甲方的数据报刚传到乙方网卡的数据接收 器的那一时刻就开始进行计数,当计数器计到192b(IP报头) +192b(TCP段头)+88b(正常文件名8个字节及其后缀名3 参考文献: 【l】龚兵.计算机维护技术In】.2版.广州:华南理工大学出版 社.2003. 个字节合计l1个字节)=472b时就使CNT高电平立刻转换 为低电平,然后将这个低电平一直保持到数据报传送完毕为 止,那么,所有寄生在正常程序中的进程后缀名“.exe”都将毫 无保留地被扫地出门了。 [2】Comer D E.Internetworking with TCP ̄P principles,protocol and architecture【M】.America:Prentice Hal1.2000. 【3】FOROUZAN B A.TCPhP protocol suite【M】.2nd ed.New York,NY,USA:McGraw—Hill,Inc.,2002. 3结束语 分析上面的逻辑电路可以知道,这种用来检测并清除病 毒和黑客程序的安全大门具有许多优点,一是安全性能非常 【4】SEBASTIAN M S.Application—speciifc integrated circuits[M]. America:Addison—Wesley Professiona1.2008:25—64. [5】于文强.计算机网络实用技术【M】.北京:清华大学出版社. 2006:l9—28. 好,它对于正常程序一点也不造成任何损伤,只对嵌入到正 常程序中的寄生代码予以清除;二是执行速度极快,整个电 路只占3个时钟周期,以现在常见的微机系统时钟4 GHz进 [6]兰高志.网卡中硬盘保护功能的增设[J],计算机工程,2004, 30(6):133—134. L~N Gao—zhi.Increasing hard—disk protection function on 行计算,则延时约为0.75 ns,也就是说一个纳秒都不到.对于 正常人的肉眼来说,在程序的下载过程中压根儿感觉不到有 netcards[J].Computer Engineering,2004,30(6):133一l34. [7]中国集成电路大全.TrL集成电路【M】.北京:国防工业出版 社,1985:30—78. 任何时间上的滞后;三是效果特别显著,比起传统的软件查 杀技术来可谓是简单方便,新颖独特,易于实现,成本低廉, 为净化Internet环境、保障网络安全、反病毒反黑客开辟了一 条崭新的途径。 【8】何立民.单片机应用系统设计【M】.北京:北京航空航天大学 出版社.1995:69—78. [9】杨全胜,胡友彬.现代微机原理与接口技术【M】.北京:电子 工业出版社.2002:243—291. (上接第143页) 【71 Ufimtsev J S,Schuhen K.GPU—accelerated molecular f101NVIDIA Corporation,NVIDIA C 3.1,NVIDIA CUDAC Programming Guid Version 3.1【S】.2010. [11】左颢睿,张启衡,徐勇,等.基if-GPU的快速Sobel ̄缘检测 modeling coming of age[J].Journal of Molecular and Modelling, 2010,29(2):116—125. [8】吴恩华,柳有权.基于图形处理器(GPU)的通用计算fJ】.计 算法[J】.光电工程,2009,36(1):8-12. 算机辅助设计与图形学学报,2004,16(5):601—612. WU En—hua.LIU You・quan.General purpose computation on ZUO Hao—rui,ZHANG Qi—heng,XU Yong,et a1.Fast sobel edge detection algorithm based on GPU[J].Opto—electronic Engineering,2009,36(1):8-12. GPU[J].Journal of Computer—aided Design&Computer Graphics,2004,16(5):601—612. 【12]Seung I P,Ponce S P,HUANG Jing,et a1.Low—cost,high— speed computer vision using NVIDIA's CUDA architecture 【9】Harish P,Narayanan P J.Accelerating large graph algrithms Oil the GPU using CUDA[J].Computer Science,2007(4873): 1 97-208. 【C1#2008 37th IEEE Applied Imagery Pattern Recognition Workshop.2008:1-7. Bluestar Silicones推出新一代SILCOLEASE(R)Optima概念:更冷、更快、更低 Optima Concept是可以解决与标签和离型材料生产相关的所有成本、利润或生产力问题的SILCOLEASE(R)解决方案。现 有的SILCOLEASE(R1系列产品可与新的Optima部件配合使用,从而应对以下挑战: 降低铂金用量 低温固化 快速运行一快速固化&减少气雾 Optima概念可通过更换零部件或在适当的时候推出完整的新系统来调整涂层配方,从而优化离型材料的生产。过去两年 内开发的最新Optima II产品已经通过测试可应用于各种生产规模.并可在更加广泛的底层和机器操作条件下运行。 咨询编号:201 1 191O10 —146— 

因篇幅问题不能全部显示,请点此查看更多更全内容