劉曉剛

摘 要：可編程圖形處理器GPU已經(jīng)演化成高并行度、多線程、擁有強(qiáng)大計(jì)算能力和極高存儲(chǔ)器帶寬的多核處理器，圖形處理器通用計(jì)算技術(shù)GPGPU把個(gè)人計(jì)算機(jī)上的顯卡用于通用計(jì)算，代替CPU完成計(jì)算工作，可以大大提升計(jì)算效率。采用CUDA技術(shù)編程，利用GPU運(yùn)行教學(xué)軟件，效果良好，較大地提高了軟件的計(jì)算能力和運(yùn)行效率，在CPU價(jià)格昂貴，大規(guī)模并行計(jì)算需求旺盛的今天，CUDA技術(shù)可以提高軟件運(yùn)行效率，提高計(jì)算能力，同時(shí)可減少硬件購置成本，為學(xué)校或科研單位節(jié)約預(yù)算。

關(guān)鍵詞：GPU；GPGPU；CUDA；教學(xué)軟件開發(fā)

DOIDOI：10.11907/rjdk.162098

中圖分類號(hào)：TP319

文獻(xiàn)標(biāo)識(shí)碼：A文章編號(hào)：1672-7800（2016）012-0124-03

0 引言

自2004年以來，CPU的發(fā)展已告別主頻時(shí)代，單核CPU性能的提升日趨平緩，CPU制造公司開始用增加計(jì)算核心的方法增強(qiáng)計(jì)算能力，多核CPU層出不窮。盡管如此，依然不能滿足日益增長(zhǎng)的計(jì)算需要，一些計(jì)算密集型的應(yīng)用和實(shí)時(shí)處理程序需要依靠超級(jí)計(jì)算機(jī)的計(jì)算性能來完成。同時(shí)，在市場(chǎng)對(duì)實(shí)時(shí)、高清晰度三維圖形存在極高需求的推動(dòng)下，可編程圖形處理器（Graphics Processing Unit，簡(jiǎn)稱GPU）已經(jīng)演化成高并行度、多線程、擁有強(qiáng)大計(jì)算能力和極高存儲(chǔ)器帶寬的多核處理器，一種新興的加速技術(shù)應(yīng)運(yùn)而生，即圖形處理器通用計(jì)算技術(shù)（General Purpose Computing on Graphics Processing Unit，簡(jiǎn)稱GPGPU），它把個(gè)人計(jì)算機(jī)上的顯卡用于通用計(jì)算，代替CPU完成計(jì)算工作，可以提升計(jì)算效率。相比CPU，GPGPU具有體積小、低功耗、低成本的特點(diǎn)，已經(jīng)成為計(jì)算密集型應(yīng)用的高效解決方案，并在事實(shí)上作為第二個(gè)通用處理器被大量應(yīng)用軟件使用。而CUDA就是著名的顯卡公司NVIDIA對(duì)于GPGPU的完整解決方案。

1 在教學(xué)軟件開發(fā)中使用CUDA的意義

CUDA的全稱是Compute Unified Device Architecture（計(jì)算機(jī)統(tǒng)一設(shè)備架構(gòu)），是NVIDIA公司在2007年推向市場(chǎng)的并行計(jì)算架構(gòu)。CUDA是作為圖形處理器的通用計(jì)算引擎，作為全套工具提供給用戶利用NVIDIA產(chǎn)品進(jìn)行GPGPU開發(fā)。CUDA提供了一種簡(jiǎn)便的方式編寫應(yīng)用于CUDA架構(gòu)上的GPGPU代碼，它含有NVIDIA擴(kuò)展和限制的類C語言，支持大多數(shù)C語言指令和語法，并加入了使程序能在GPU上進(jìn)行多線程計(jì)算的語言擴(kuò)展，使用CUDA可以方便地編寫在GPU上并行運(yùn)行的程序[1]。

眾所周知，學(xué)校、科研單位以及培訓(xùn)機(jī)構(gòu)在資金、預(yù)算上受財(cái)政的限制，缺少高性能的計(jì)算機(jī)和服務(wù)器等硬件設(shè)備。同時(shí)，在高速發(fā)展的互聯(lián)網(wǎng)時(shí)代，也需要各種教學(xué)軟件的高效運(yùn)行，以保證教學(xué)、科研活動(dòng)的正常進(jìn)行。在教學(xué)軟件開發(fā)中使用CUDA技術(shù)有具有較大意義。一方面，帶多處理器的顯示圖形卡越來越便宜，購置成本低，其計(jì)算性能卻可以代替昂貴的高端CPU[2]；另一方面，采用CUDA技術(shù)開發(fā)或改寫教學(xué)軟件，可以大大提高其運(yùn)行效率，進(jìn)而延長(zhǎng)軟件的使用壽命。

由此可見，CUDA技術(shù)可以減少硬件的購置成本，提高軟件的運(yùn)行效率，提高計(jì)算能力，并為學(xué)?；蚩蒲袉挝还?jié)約預(yù)算，值得推廣。同時(shí)，CUDA的軟件開發(fā)類似于C語言開發(fā)，易于被開發(fā)人員所學(xué)習(xí)掌握。下文將通過一個(gè)實(shí)例，介紹基于CUDA技術(shù)軟件的開發(fā)方法，并對(duì)其運(yùn)行性能進(jìn)行對(duì)比分析。

2 開發(fā)實(shí)例

2.1 基礎(chǔ)知識(shí)

在CUDA中，與一個(gè)線程組相關(guān)聯(lián)的硬件實(shí)現(xiàn)稱為一個(gè)流多處理器（Streaming Multiprocessor），或稱多處理器；一個(gè)流多處理器包含若干個(gè)標(biāo)量處理器（scalar processor），即單處理器，或稱CUDA核（CUDA core）。單處理器是具體計(jì)算指令的執(zhí)行單位，多處理器是一套完整的計(jì)算資源的最小單位?？刂破鲗⒁粋€(gè)線程組分配給一個(gè)多處理器，多處理器中的核協(xié)同工作，并行處理所有的線程。一個(gè)支持CUDA的NVIDIA圖形處理器中至少包含一個(gè)多處理器，如GeForce GTX 295，它有60個(gè)多處理器，線程組就被平均分配給這60個(gè)多處理器并行處理。而且在每個(gè)多處理器中，標(biāo)量處理器也是并行工作的。通過線程組的方式，CUDA程序可以適應(yīng)于不同的硬件規(guī)格[3]，從高端顯卡到普通顯卡，這是一個(gè)具有高度可擴(kuò)展性的編程模型。

CUDA的程序?qū)⒑瘮?shù)根據(jù)調(diào)用和執(zhí)行地方的不同分為不同的類型：在CPU上調(diào)用、CPU上執(zhí)行的的函數(shù)稱為主機(jī)函數(shù)；在CPU上調(diào)用、在GPU上執(zhí)行的函數(shù)稱為全局函數(shù)，常稱為內(nèi)核，簡(jiǎn)稱核，它是GPU函數(shù)執(zhí)行時(shí)的基本單位。在調(diào)用此類函數(shù)時(shí)，它將由 N 個(gè)不同的 CUDA 線程并行執(zhí)行 N 次，這與普通的C語言函數(shù)只執(zhí)行一次的方式不同，所以，雖然這里沒有看到循環(huán)的語法，卻類似執(zhí)行了循環(huán)語句。主機(jī)函數(shù)在聲明時(shí)帶限定符__host__，全局函數(shù)在聲明時(shí)必須帶限定符__global__。

CUDA將一個(gè)線程組稱為一個(gè)block，每個(gè)block由若干線程組成，完成一次函數(shù)調(diào)用的所有block組成了一個(gè)grid。block和grid的尺寸都可以用三元向量表示。線程索引（thread index）是線程在每個(gè)block里的索引。由于block的尺寸是三維的，線程索引也是一個(gè)三元向量threadIdx。訪問它的每個(gè)分量需要加上分量名稱如threadIdx.x，threadIdx.y和threadIdx.z。執(zhí)行內(nèi)核的每個(gè)線程都會(huì)被分配一個(gè)獨(dú)特的線程 ID，可通過內(nèi)置的 threadIdx 變量在內(nèi)核中訪問此ID。

2.2 具體實(shí)例

本次實(shí)例中的顯卡是NVIDIA 的NVS 4200M，計(jì)算能力是2.1，只有一個(gè)多處理器，有48個(gè)核，GPU時(shí)鐘頻率為1.62GHz，現(xiàn)在只能算是運(yùn)算能力一般的GPU。CPU是INTEL的Core i5-2410M，主頻2.3GHz，屬于性能中等的CPU。

先通過計(jì)算分析一下GPU和CPU的計(jì)算能力。NVS 4200M顯卡的計(jì)算能力是48×1.62= 77.76 GFLOPS，Core i5-2410M CPU的計(jì)算能力是2.4×4=9.6GFLOPS。因此得出前者的浮點(diǎn)數(shù)計(jì)算性能是后者的8.1倍，下面再通過具體實(shí)例的計(jì)算結(jié)果來驗(yàn)證它們的計(jì)算能力。

程序開始處根據(jù)GPU的硬件參數(shù)定義數(shù)據(jù)常量DATA_SIZE為1048576，線程常量THREAD_NUM為256，線程組常量BLOCK_NUM為32，GPU主頻FREQUENCY是 1620000kHz。

首先通過子函數(shù)GenerateNumbers（ ）隨機(jī)產(chǎn)生DATA_SIZE個(gè)整數(shù)，保存在整型數(shù)組pnData中，用來分別提供給CPU和GPU計(jì)算求和。因?yàn)槭褂猛瑯拥臄?shù)據(jù)和同樣的計(jì)算公式，所以計(jì)算結(jié)果應(yīng)該一致。

CPU計(jì)算的函數(shù)利用循環(huán)來實(shí)現(xiàn)，通過一個(gè)自定義記時(shí)類CTimer來記時(shí)，單位是ms，直接用printf語句輸出時(shí)間，代碼分析如下：

void sumOfCpu（int *pnData）

{ CTimer cc； long nSum=0； //nSum用來保存計(jì)算的和

long tt=cc.getTime （）； //tt用來保存時(shí)間，先獲取當(dāng)前的時(shí)間

for （int i=0；i

{ nSum+=（pnData[i]*pnData[i]*pnData[i]-pnData[i]*pnData[i]-pnData[i]； }

tt=cc.getTime （）-tt； //計(jì)算后的時(shí)間減去計(jì)算前的時(shí)間，得到CPU計(jì)算耗時(shí)

printf（"CPU sum：%d time： %d ms＼＼n"，nSum，tt）； }

GPU計(jì)算的核函數(shù)名為sumOfGpu，采用縮減樹算法來提高存儲(chǔ)器的訪問效率。具體代碼及分析如下，注意除了循環(huán)語句，執(zhí)行 SumOfGpu（ ）的每個(gè)線程都會(huì)執(zhí)行一次加法運(yùn)算。

__global__ static void sumOfGpu（int *pnNumber，int * pnResult，clock_t* pclock_tTime）

{ extern __shared__ int nShared[]； //聲明一個(gè)動(dòng)態(tài)分配的共享存儲(chǔ)器

const int tid=threadIdx.x； //tid中保存線程號(hào)

const int bid=blockIdx.x； //bid中保存線程組號(hào)

int i； int nOffset=1； //nOffset記錄縮減樹算法中每輪增倍的步距

if （tid==0） pclock_tTime[bid]=clock（）； //獲取當(dāng)前時(shí)間并按線程組保存

nShared[tid]=0； //每一個(gè)線程組的和都存儲(chǔ)在對(duì)應(yīng)下標(biāo)的共享存儲(chǔ)器中

for （i=bid*THREAD_NUM+tid；i

__syncthreads（）； //線程同步

nOffset=THREAD_NUM/2； //下面用縮減樹算法計(jì)算共享存儲(chǔ)器中元素的和

while （nOffset>0）{

if （tid

nOffset>>=1； __syncthreads（）； }

if （tid==0）{ //最后用一個(gè)線程來相加共享存儲(chǔ)器nShared中的數(shù)據(jù)

pnResult[bid]=nShared[0]；

pclock_tTime[bid+BLOCK_NUM]=clock（）； }}

主函數(shù)main（）在CPU中運(yùn)行，主要完成數(shù)據(jù)的初始化，調(diào)用子函數(shù)，并輸出結(jié)果。CUDA在運(yùn)行前要分配內(nèi)存空間，運(yùn)行結(jié)束后要釋放這些空間。

int main（）

{ …… //聲明變量，分配內(nèi)存

GenerateNumbers（pnData，DATA_SIZE）； //生成原始數(shù)據(jù)

……

sumOfGpu<<>>

（pnGpuData，pnResult，pclock_tTime）； //調(diào)用GPU核函數(shù)

……//整理匯總GPU的計(jì)算時(shí)間pclock_tTime并輸出

sumOfCpu（pnData）； //調(diào)用CPU函數(shù)求和并輸出

……//釋放內(nèi)存 }

首先計(jì)算一個(gè)比較復(fù)雜的公式求和，即x的立方減x的平方減x的和，如式（1），其中有立方有平方有減法，代碼如上所示。程序運(yùn)行結(jié)果如圖1所示，GPU只需要1ms，而CPU需要15ms。

接著將程序進(jìn)行少量修改，計(jì)算第二個(gè)公式，即（X%10）的三次方和，即式（2），這里有求模和立方，是比較耗時(shí)的計(jì)算。程序運(yùn)行結(jié)果如圖2所示，GPU的計(jì)算時(shí)間需要3ms，CPU計(jì)算時(shí)間則需要16ms。

從上述實(shí)例來看，計(jì)算同樣的公式，NVS 4200M的計(jì)算時(shí)間都比Core i5-2410M大大減少，其計(jì)算能力遠(yuǎn)超這顆主流的CPU。因此，將CUDA技術(shù)用于教學(xué)和科研軟件的開發(fā)，從而提高其計(jì)算能力的建議是可行的。

3 應(yīng)用分析

GPU特別為計(jì)算密集、高并行度計(jì)算（如圖像渲染）而設(shè)計(jì)，將更多的晶體管用于數(shù)據(jù)處理而不是數(shù)據(jù)緩存和流控，因此浮點(diǎn)計(jì)算能力很強(qiáng)。具有以下特點(diǎn)的算法能夠在GPU上達(dá)到很高的執(zhí)行效率：①每個(gè)數(shù)據(jù)或數(shù)據(jù)包都需要經(jīng)過相同的流程來處理；②數(shù)據(jù)之間沒有相干性，即某些數(shù)據(jù)的計(jì)算不依賴于另外一些數(shù)據(jù)的計(jì)算結(jié)果；③數(shù)據(jù)量龐大。

然而，對(duì)于程序中邏輯分支較多、數(shù)據(jù)計(jì)算中需要大量緩存，以及并行程度較低、線程較少的計(jì)算，并不能體現(xiàn)出GPGPU的優(yōu)勢(shì)，不適合用CUDA技術(shù)。

4 結(jié)語

在3D領(lǐng)域，GPU的用途很簡(jiǎn)單，就是為了更好地渲染3D場(chǎng)景，減輕CPU在圖形運(yùn)算方面的負(fù)擔(dān)，而GPGPU則將GPU的應(yīng)用范圍擴(kuò)展到了圖形之外。采用CUDA技術(shù)編程，利用顯示圖形卡的GPU運(yùn)行教學(xué)科研軟件，可以大大提高軟件的計(jì)算能力和運(yùn)行效率。在CPU價(jià)格昂貴、大規(guī)模并行計(jì)算需求旺盛的今天，CUDA是一種好的解決方案，特別適合于投入有限、資金緊張的學(xué)校和科研單位。

參考文獻(xiàn)：

[1] 仇德元. GPGPU編程技術(shù)——從GLSL、CUDA到OpenCL[M].北京： 機(jī)械工業(yè)出版社，2012.

[2] 續(xù)士強(qiáng)，祝永志. 基于GPU加速的快速字符串匹配算法[J].軟件導(dǎo)刊，2015（2）：51-53.

[3] 郭轉(zhuǎn)轉(zhuǎn)，尹延慶，王佩璐.淺談CUP并行技術(shù)CUDA[J].信息通信，2014（5）：103.

（責(zé)任編輯：孫 娟）