配色: 字号:
063821_CUDA并行计算的应用研究[1]
2012-08-08 | 阅:  转:  |  分享 
  
%黪琵√傍彩%鼢绍彩易殇驴

中图分类号: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一

万方数据

献花(0)
+1
(本文系yangshiquan...首藏)