%黪琵√傍彩%鼢绍彩易殇驴
中图分类号:TP391
2010年第4期
文献标识码:A文章编号:1009-2552{2010)04-0011—05
CUDA并行计算的应用研究
董荦,葛万成,陈康力
(同济大学电子与信息工程学院,上海200092)
摘要:统一设备架构(CUDA)是NVIDIA公司提出的一个基于GPU通用计算的开发环境,它
针对GPU多处理单元的特性,通过并行计算提高大规模运算的速度。根据CUDA技术的特点,
提出了基于CUDA的并行图像锐化、中值滤波和字符搜索算法,并论述其关键技术和基本执行
流程。试验结果表明,提出的方法相对于CPU方法在运算速度上有不同程度的提高和下降。这
同时体现了CUDA的优势和局限性,为其更复杂的应用提供了参考和依据。
关键词:图形处理器;统一计算设备构架;图像锐化;中值滤波;并行计算
StudyonapplicationofparallelcomputationonCUDA
DONGLuo。GEWan.cheng.CHENKang—li
(DepartmentofInformation&CommunicationEn西neering,TongjiUniversity,Shanghad200092,China)
Abstract:ComputeUnifiedDeviceArchitecture(CUDA),whichwasproposedbyNVIDIA
Corporation.isadevelopmentenvironmentbasedonGPUgeneral—purposecomputation.Incognizanceof
thefeaturesofmulticoreofGPU.itcouldincreasethespeedofthemassivedatacomputingthrough
parallelcomputation.Inthispaper,theCUDAbasedmethodsforparallelimagesharpening,median
filteringandcharactersearchingareproposedbasedontheCUDAtechnologyanditsfeatures,andsome
keytechnologiesandbasicexecutionprocessesarediscussed.Theexperimentalresultsshowthatthe
computingspeedoftheproposedmethodsinmostcasesismuchfasterthanthatoftheCPUbased
me出ods,butin$omecasesslower.ItindicatesthattheCUDAtechnologyhasbothmeritsanddemerits
andtheexperimentalresultsprovidereferencesforcomplexapplications.
Keywords:GPU;CUDA;imagesharpening;medianfiltering;parallelcomputing
0引言
近年来GPU已经具备了实现大规模快速计算
的编程能力,NVIDIA公司提出的计算统一设备架
构(ComputerUnifiedDeviceArchitecture,CUDA)技
术就是这方面的杰出代表。CUDA编程给人们一种
新的理念,人们可以将GPU高速并行处理能力用于
数字图像处理算法、离散模拟等通用计算领域,大幅
度提高程序运算速度。目前国内外关于CUDA的
研究主要集中于算法的设计和CUDA优势的阐明。
本文将对CUDA并行计算技术的应用进行研究,重
点考查其在具体使用中的优势和局限性。
1CUDA工作原理
1.1GPU简介
GPU(GraphicProcessingUnit)就是图形处理
器,早期的计算机图形处理相对简单,所以都交由
CPU来完成。从NVIDIA公司在1999年推出了第
一个真正意义上的GPU——GeForce256至今,GPU
已经历八代的发展。如图l所示,在过去lO年间
收稿日期:2009一09一14
基金项目:获德国罗德施瓦茨公司资助
作者简介:董荦(1985一),男,同济大学电子与信息工程学院硕士
研究生,研究方向为信号与信息处理。
一1l一
万方数据
GPU的计算能力有了飞速的发展,基本上是平均每可以通过sharedmemory通信,也可以同步。
6个月就有性能翻倍的产品面世。同时,GPU计算
性能的发展速度也大大快于CPU‘1|。尤其在较大
规模的并行运算上的优势更为明显,并且相对传统
的大型计算机又具有成本的优势。
m,,,……,.,…一……………….口佟哆——
图1CPU和GPU的每秒浮点运算次数比较
GPU在速度上之所以会有如此巨大的变革主
要是因为它本身的设计原因。从微构架上看,GPU
使用更多的晶体管来用于数据处理,其数量远大于
数据缓存和流控制。所以,GPU的优势是无逻辑关
系数据的并行运算。这样一来,它对流控制的要求
就不是很高,而且运算将代替大数据缓存来隐藏内
存存取的延迟。
然而,GPU要实现通用运算还存在许多困难,
比如:通过API编程复杂,难于学习;由于DRAM内
存带宽有限,一些程序会遇到瓶颈;GPU编程缺乏
灵活性,大大限制了GPU性能的发挥。这些缺点随
着NVIDIA公司CUDA的推出而有了解决方案。
1.2CUDA概况
CUDA(ComputeUnifiedDeviceArchitecture,统
一计算设备架构)是一种新的处理和管理GPU计算
的硬件和软件架构,它将GPU视作一个数据并行计
算设备,并且无需把这些计算映射到图形API。CU-
DA程序的开发语言以C语言为基础,并对C语言
进行扩展。
基于CUDA开发的程序代码在实际执行中分
为两类,一类是运行在CPU上的宿主(host)代码,
一类是运行在GPU上的设备(device)代码。其中
运行在GPU上的并行程序称为kernel。
CUDA并行处理结构即线程组织结构如图2所
示,执行内核的线程(thread)被组织成线程块
(block),线程块又组成栅格(面d)。这样同一个
kemel程序可以并行运行在一个栅格所包括的所有
线程块中的线程上旧J。而同一个线程块中的线程
一12一
图2线程组织结构不薏图
1.3GPU硬件构架与CUDA程序执行原理
NVIDIA的GPU中,最基本的处理单元是SP
(StreamingProcessor),而一个NVIDIA的GPU中有
很多的SP。而几个sP加上registor,sharememory,
texturecache和constantcache等一起组成一个SM
(StreamingMultiprocessor)o几个SM又组成一个
TPC(TextureProcessingClusters)o如图3所示。
图3GPU的硬件构架
在GS0/G92核心中,共有128个SP,8个sP为
一组,组成16个SM,再以2个SM为一组,组成一
个8个TPC。而在新一代的GT200中,sP增加到
240,仍然以8个sP组成一个SM,以3个SM组成
一个11PC,共lO个TPC。
GPU执行CUDA程序的模式称为SIMT(单指
令多线程)。每个SP对应一个线程,每个SM对应
一个或几个线程块。但SM实际执行时并不是以线
程块为单位的,而是以warp块为单位。warp块由线
程块分割而成,每个warp块一般由32个线程组成。
万方数据
warp的分组是由SM自动进行,以连续的方式来分
组。线程块中的线程数量不是32倍数时,剩下的会
被独立分成一组。比如线程数为70,就会被分为3
组warp:0—3l、32—63、64-70。
执行时,每发出一条指令,SIMT单元都会选择
一个已准备好warp块执行,并将下一条指令发送到
该warp块的活动线程。如果某行指令需要等待时,
SM会自动切换到下一个warp块来执行。以此来隐
藏线程的延迟和等待来达到大量平行化的目的。
在CUDA中,GPU不能直接存取主内存,只能
存取显卡上的显卡内存。因此,需要将数据从主内
存先复制到显卡内存中,进行运算后,再将结果从显
卡内存中复制到主内存中。这些数据传递的速度受
限于PCIExpress的带宽。如PCIExpress1.0可以
提供双向各4GB/s的带宽,而PCIExpress2.0则可
提供8GB/s的带宽。
1.4CUDA计算的特点分析
基于CUDA计算的优点:
(1)GPU通常具有更大的内存带宽。例如Ge—
Force9800GT具有超过57.6GB/s的内存带宽,而目
前高阶CPU的内存带宽则在10GB/s左右。
(2)GPU具有更多的执行单元。例如GeForce
9800GT具有112个SP。
(3)GPU较同等级的CPU具有价格优势。
基于CUDA计算的缺点:
(1)对并行化程度不高的程序运行效率较低。
(2)GPu目前通常只支持32bits浮点数,且很
多都不能完全支持IEEE754规格,有些运算的精确
度可能较低。目前许多GPU并没有另外的整数运
算单元,因此整数运算的效率较差。
(3)GPU通常不具有分支预测等复杂的流程控
制单元,因此对于具有高度分支的复杂程序,效率
较差。
(4)由于PCI—E接口带宽限制,若频繁于内存
和显存间交换数据会降低程序的运行效率。
2基于CUDA的算法设计与实现
针对CUDA计算的特点选取了三种复杂度不
同的算法进行并行化设计。
2.1拉普拉斯算子图像锐化原理
图像锐化是以对图像的微分运算或差分运算为
基础的。微分运算是求像素点灰度值的变化率,而
图像内不同物体边缘处的像素点的灰度值往往变化
比较明显,因此微分运算或差分运算可以起到增强
边缘信息的作用。图像处理中最常用的微分方法是
求梯度。本文采用的拉普拉斯算子为二阶偏导数,
其定义为p】:
v2厂:套+套
。如‘砂‘
对于数字图像,在某个像素点(x,y)处的拉普
拉斯算子可采用如下差分形式近似:
V0^(i√)=厂(i—l√一1)+以i一1√)+
火i—lJ+1)+八i,-『一1)+八iJ+1)+
以i+lJ一1)+厂(i+1√)+
“i+lJ+1)一毪“i√)
本文采用文献[4]中的推广形式,考虑进对角
线方向,这样就形成了一个8领域算子,其模板为:
lll
1—81
11l
2.2中值滤波原理
中值滤波是一种非线性图像局部平滑技术。它
对脉冲干扰和椒盐噪声的抑制效果较好,在抑制噪
声的同时能使边沿减少模糊。
具体方法是对一个滑动窗口内的诸像素灰度值
排序,用其中值代替窗口中心像素(石,Y)原来的灰
度值。正确选取窗口尺寸的大小是决定中值滤波效
果的重要环节,当窗口内噪声点的个数大于窗口宽
度一半时,滤波效果不好。
2.3CUDA并行图像锐化和中值滤波算法设计
从拉普拉斯算子的原理可以看出,每个像素的值
只与其8领域的像索值有关,并且每个像素值的算法
是完全相同的。于是使一个像素对应一个线程。由
此首先设计kemel函数LinearSharpen,伪代码如下:
一global—voidLinearSharpen(输入图像指针,输出图
像指针,图像宽度,图像高度)
{
(1)intnWeight[3][3]={一1;·1,一1,·1,8,-1,·
1,-1,-1};
(2)intindex=width堆BLOCK—SIZE母block-
Idx.y+BLOCK—SIZE%blockldx.x+
width宰threadldx.Y+hreadldx.x:
(3)intnTmp[O][0]=f(index);
(4)for(YY=O;YY<3;YY++)
for(X_X=0;xx<3;XX++)
{
dGrad+=nTrap[yy][xx]nWeight[yy][殛];
}
(5)editlplmage[index]=f(dGrad);
}
函数定义中一global~声明函数为kernel函数。
伪代码中(1)为设置8领域模板权值。(2)计算当
一13—
万方数据
前像素点的索引值,其中width为图像宽度,BLOCK
—SIZE为一个线程块包含线程的行数或列数(由于
每个线程块最多能包含512个线程∞J,本文定义线
程块大小为16×16),blockldx.x和blockIdx.Y分别
表示当前线程块的茗和Y方向索引值,threadldx.x
和threadldx.Y分别表示当前线程的茗和Y方向索引
值。(3)根据index计算8领域点的索引值并将其
像素灰度值付给nTmp[3][3]矩阵。(4)计算梯度。
(5)由梯度计算当前像素点锐化后的像素灰度值。
主程序伪代码:
main()
{
读入图像文件并初始化host;
将数据传到device;
(1)dim3threads(BLOCK—SIZE,BLOCK—
SIZE);
(2)dim3grid(wWidtlCthreads.x,hHeight/threads.
Y);
(3)clock—tcosttime=0;
costtime一=clock();
(4)LinearSharpen“<鲥d,threads>>>(d—im-
age,d—editimage,wWidth,hHeight);
同步;
costtime+=clock();
把结果传回host;
写出图像文件;
}
其中(1)设定线程块大小为16×16(每个线程
块内包含16×16=256个线程)。(2)设定栅格大
小,wWidth,hHeight为图像像素宽度和高度。以256×
256的图像为例,将图像分为16×16=256个子图
像,每个子图像对应一个线程块,则栅格的大小为
16×160(3)定义时间变量记录程序运行时间。
(4)调用kernel函数让其在设定好的栅格和线程块
上并行运行。
3×3中值滤波的算法设计与此类似,只是kemd
不同。
2.4并行字符搜索的实现
字符的搜索也是常用的算法之一。具体实现方
法为:
(1)将txt文本读人host,并按行分组;
(2)将文本数据传到dedce;
(3)运行kernel处理数据,其中一个线程对应
一行字符;
(4)将产生的结果(目标字符所在的行数、列
数和编号)传回host。
3实验结果与分析
实验平台:NVIDIAGeforce9400GT,Pentium
(R)43.00GHz,DDR25331GB。
实验采用的原始图像如图4所示,锐化和中值
滤波的结果分别如图5和图6所示。
图4原始图像图5锐化后图像图6中值滤波后图像
为了有更好地对比,将Lena图像调整为多种分
辨率。在CUDA方法和传统的CPU方法下分别对6
种不同分辨率的原始图像进行锐化和中值滤波,记
录程序运行时间。运行结果如表1,表2所示。
总体来看两种算法的运行速度较CPU方法都
有一定提高。由表I和表2可以看出Device与host
之间数据传递时间和图片像素数成线性。证明程序
设计合理没有频繁传递数据。两种算法加速比基本
随像素增加而增大,线性锐化和3×3中值滤波在大
于1280×1280像素后加速比略有下降。像素较低
时加速比小于1的原因是线程数量不足,不能充分
一14~
隐藏数据传输和读取延迟。
纵向对比,两种算法kernel程序的复杂度分别
为O(20)和0(47)(其中中值滤波的复杂度主要来
自多数值排序),两算法的加速比平均值分别为1.6
和18.6。分析原因是由于图像锐化的kernel程序
复杂度较低不足以发挥GPU并行处理的优势以隐
藏数据传输读取的延迟。
字符搜索的速度对比情况如表3所示,加速比
小于1,随文本行数增加加速比逐渐增加并趋于稳
定。分析原因,是由于SM共享内存的容量限制,必
万方数据
须将文本按行分组,这样并行处理的结果也是分组
的。传回host之后还要将其合并。而分组、传递和
合并的总延迟是并行运算无法完全隐藏的。
综上所述,CUDA程序设计应该注意以下几点:
(1)控制kernel程序复杂度,不应太低。
(2)一个线程块中的线程数最好是32的倍数,
在上限允许下尽量增加线程的数量。
(3)避免选择并行程度不高的程序、需要频繁
传递数据的程序和需要频繁同步的程序。
裹1锐化速度对比
图像分辨率256x256512x512768x7681024×10241280×12801536x1536
CPU耗时(瑚)11.043.8100.0182.9273.3392.1
CUDA耗时(∞)26.745.356.3108.2109.4203.2
数据传递耗时(啪)20.429.728.136.335.943.7
加速比0.410.971.781.692.501.93
表23x3中值滤波速度对比
图像分辨率256x256512x512768×7681024×10241280x12801536×1536
4结束语
随着大规模计算需求的增长,GPU通用计算技术
迅速发展,而其编程复杂性又制约了其应用。CUDA
正是在这种情况下应运而生的。本文对CUDA关键
技术特点进行了分析,设计了三种不同复杂度的
CUDA算法,最后通过对比实验得出两个结论:一方
面,合理的CUDA并行程序设计使运行速度大幅提
高。本次实验中用到的显卡,是支持CUDA的低档
显卡,只有16个sP。而像GTX280这样的高级显卡
SP已经达到240个。因此,CUDA在GPU通用计算
领域的应用显示出极大的潜力。另一方面,全面考
虑CUDA的优势和局限性,选择合适的程序以及合
理的设计才能将CUDA的性能发挥到极致。
参考文献:
[I]NVIDIA.NVIDIACUDAProgr删ngGuide[EB/OL].http://devd-
older.download.nvidia.conV∞mputeJeuda/2—2/toolkit/dote/
NVIDIA—CUDA—Programming—Guide一2.2.1.pdf,2009-5—26.
[2]多相复杂系统国家重点实验室多尺度离散模拟项目组.基于GPU
的多尺度离散模拟并行运算Em3.科学出版社,2009:7—16.
[3]陈天华.数字图像处理[M].清华大学出版社,2007:162—164.
[4]杨枝灵,王开,等.Visual''C++数字图像获取处理及应用[M].
人民邮电出版社,2003:173—180.
[5]NVIDIA.NVIDIACUDAReferenceManual[EB/OL].http://de-
veloper.download.nvidia.com/computefcuda/2—2/toolkit/does/
CUDA—Reference—Manual一2.2.pdf。2009-4.
责任编辑:李光辉
●P舢●m●●—_’●-●巾●巾●-_●-“--峄●“-m◆舢●-◆●◆_H■-H●_H●_H●哪●_H●一_●●●_HHH-_‘●_◆●_H-_h●巾●_●-H●-。●_H卜●H-●峄_◆●H-_-●一h●^H●-HH峥-H●--
(上接第10页)
【2]Pa]ehaudhuris。SahaAK。Johnson加.AdaptiveClockSyn-chroni.
sationinSensorNetworksProceedingsoftheThirdInternational
SymposiumonInformationProcessinginSensorNetworks[C],
2004:340—348.
[3]C,enedwals,Ku吨R,S时vsstavaM.1ilnillg.sy啦hot0。0l细Sensor
Net哪[c].ACMsensys,LesAngeles,CA,200B(11):1睨一1∞.
[4】PingS.【嘲唧M嘲轴柏n朗t髓r瞻syr虻bd脚i∞forWirelessSensor
№M幽[z].IlltdFamemmhCemer:IRB-TR.妇8-013,June。2008.
[5]
[6】
[7】
№∞tim,Ku8yb,SimonG,eta1.TheFlooding甄nle跏℃I啪niz出m酬[c].h∞出够0f妇2rid
In船nmiomlC0l蛔∞Era-
beddedNetworked蜘Sy咖m(Se啷y删),2004:39-49.
PailFan,lndnmeelcIlak啊b0啊,NmmyI和d1.Clock却ncllmdz面帆
forw抵Net’mrks[c].OPODIS2004。LNCS3544,20晒:400—
414.
胁118DL]ntemettimesynchxliz甜鲫:ThenetworktimeplI勘col
【J].m髓Tr=um''fiommCompute口,1991,39(10):1482—1493.
责任编辑:李光辉
一15一
万方数据
|
|