CUDA并行計(jì)算架構(gòu)及編程

《CUDA并行計(jì)算架構(gòu)及編程》由會(huì)員分享，可在線閱讀，更多相關(guān)《CUDA并行計(jì)算架構(gòu)及編程（5頁(yè)珍藏版）》請(qǐng)?jiān)谘b配圖網(wǎng)上搜索。

1、CUDA并行架構(gòu)及編程 摘要CUDA是一種由NVIDIA推出的并行計(jì)算架構(gòu)，非常適合大規(guī)模數(shù)據(jù)密集型計(jì)算。CUDA使GPU的超高計(jì)算性能在數(shù)據(jù)處理和并行計(jì)算等通用計(jì)算領(lǐng)域發(fā)揮優(yōu)勢(shì)，本文討論了CUDA的計(jì)算架構(gòu)和基于GPU的CUDAC編程語(yǔ)言,CUDA使GPU流處理器陣列的性能得到充分發(fā)揮。極大地提高了并行計(jì)算程序的效率。 關(guān)鍵詞并行計(jì)算，GPU通用計(jì)算，CUDAAbstractCUDAisaparallelcomputingarchitectureintroducedbyNVIDIA,itmainlyusedforlargescaledata-intensivecomputing.CUDA

2、makesGPUahighperformanceinparallelcomputing,dataprocessingandothergeneralcomputing.thispaperdiscussestheCUDAcomputingarchitectureandCUDACprogramminglanguagebasedonGPU,CUDAmakesGPUstreamprocessorarraysfullyusedandGreatlyimprovedtheefficiencyofparallelcomputingprogram. Keywordsparallelcomputing，GPUge

3、neralpurposecomputation，CUDA1引言 并行計(jì)算是指同時(shí)使用多種計(jì)算資源解決計(jì)算問題的過程。并行計(jì)算科學(xué)中主要研究的是空間上的并行問題。從程序和算法設(shè)計(jì)的角度來(lái)看，并行計(jì)算又可分為數(shù)據(jù)并行和任務(wù)并行。一般來(lái)說(shuō)，GPU更注重于數(shù)據(jù)并行計(jì)算，主要是將一個(gè)大任務(wù)化解成相同的各個(gè)子任務(wù)。早期的GPU研究是通過可編程計(jì)算單元為屏幕上的每個(gè)像素計(jì)算出一個(gè)顏色值即渲染問題。自CUDAC出現(xiàn)后，基于GPU的通用計(jì)算已經(jīng)成為一個(gè)新的研究領(lǐng)域。通常，像素著色器對(duì)各種顏色值進(jìn)行合成并計(jì)算出最終的顏色值。實(shí)際上，輸入值可以為任意數(shù)據(jù)，這樣不一定非要使用GPU來(lái)處理圖形，還可以實(shí)現(xiàn)某些通用計(jì)算

4、。由于GPU有著很高的計(jì)算吞吐量，從而給大規(guī)模的數(shù)據(jù)計(jì)算應(yīng)用提供了一種比CPU更加強(qiáng)大的計(jì)算能力。 CUDA是一種由NVIDIA推出的并行計(jì)算架構(gòu)，該架構(gòu)使GPU能夠解決復(fù)雜的計(jì)算問題。它包含了CUDA指令集架構(gòu)以及GPU內(nèi)部的并行計(jì)算引擎⑴。隨著顯卡的發(fā)展，GPU越來(lái)越強(qiáng)大，在計(jì)算上已經(jīng)超越了通用的CPU。如此強(qiáng)大的芯片如果只是作為顯卡會(huì)造成計(jì)算能力的浪費(fèi)，因此NVIDIA推出CUDA，讓顯卡可以用于圖像渲染以外的目的。CUDA的GPU編程語(yǔ)言基于標(biāo)準(zhǔn)的C語(yǔ)言，通過在標(biāo)準(zhǔn)C語(yǔ)言的基礎(chǔ)上增加一小部分關(guān)鍵字，任何有C語(yǔ)言基礎(chǔ)的用戶都很容易地開發(fā)CUDA的應(yīng)用程序。數(shù)以千計(jì)的軟件開發(fā)人員正在使用

5、免費(fèi)的CUDA軟件開發(fā)工具來(lái)解決各種專業(yè)中的問題［?。這些解決方案涵蓋了石油天然氣勘探、產(chǎn)品設(shè)計(jì)、醫(yī)學(xué)成像以及科學(xué)研究等領(lǐng)域。 2CUDA架構(gòu) CUDA程序架構(gòu)分為兩部分：主機(jī)和設(shè)備。一般而言，主機(jī)指的是CPU及其內(nèi)存，設(shè)備指的是GPUM］。在CUDA程序架構(gòu)中，主程序由CPU來(lái)執(zhí)行，而當(dāng)遇到數(shù)據(jù)并行處理的部分，CUDA就會(huì)將程序編譯成GPU能執(zhí)行的程序，并傳送到GPU。這種函數(shù)在CUDA里叫做核函數(shù)。在GPU中要執(zhí)行的線程，根據(jù)最有效的數(shù)據(jù)共享來(lái)創(chuàng)建線程塊，其類型不止一維。在同一個(gè)塊里的線程，使用同一個(gè)共享內(nèi)存。每個(gè)線程由線程ID標(biāo)識(shí)。這是線程塊中的線程號(hào)。為了幫助基于線程ID的復(fù)雜尋址

6、，還可以將線程塊指定為任意大小的一維、二維或高維線程陣列，并使用多個(gè)索引分量來(lái)標(biāo)識(shí)每個(gè)線程⑵。線程塊的大小是也有限制的,所以不能把所有的線程都放到同一個(gè)塊里?？梢杂猛瑯泳S度和大小的塊來(lái)組成一個(gè)網(wǎng)格做批處理。執(zhí)行內(nèi)核的線程被組織成線程塊。而線程塊又組成了線程格，如圖1。線程塊內(nèi)一起協(xié)作的線程通過一些快速的共享內(nèi)存有效地共享數(shù)據(jù)并同步執(zhí)行，以協(xié)調(diào)內(nèi)存訪問。用戶可以在核函數(shù)中指定同步點(diǎn)。線程塊中的線程在到達(dá)此同步點(diǎn)時(shí)掛起。 圖1線程被組織為線程格與線程塊 在GPU內(nèi)部，SM代表流處理器，即計(jì)算核心。每個(gè)SM中又包含8個(gè)標(biāo)量流處理器（SP）以及少量的其他計(jì)算單元⑶，如圖2。實(shí)際上SP只是執(zhí)行單

7、元，并不是完整的處理核心。處理核心必須包含取指、解碼、分發(fā)邏輯和執(zhí)行單元。隸屬同一SM的8個(gè)SP共用同一套取指和發(fā)射單元，也共用一塊共享存儲(chǔ)器。一個(gè)線程塊必須被分配到一個(gè)SM中，但是一個(gè)SM中同一時(shí)刻可以有多個(gè)活動(dòng)線程塊等待執(zhí)行。這可以更好地利用執(zhí)行單元的資源，當(dāng)一個(gè)線程塊進(jìn)行同步或者訪問顯存等高延遲操作時(shí)，另一個(gè)線程塊就可以占用GPU執(zhí)行資源。每個(gè)線程有一個(gè)私有的本地存儲(chǔ)器，即每個(gè)標(biāo)量處理器私有的緩存。每個(gè)線程塊有一個(gè)共享存儲(chǔ)器，塊內(nèi)所有的線程都可以訪問。每個(gè)線程網(wǎng)格都可以訪問設(shè)備內(nèi)存，即GPU片上的全局存儲(chǔ)器。每個(gè)線程均可以讀寫全局存儲(chǔ)器。本地存儲(chǔ)器用于儲(chǔ)存設(shè)備代碼中無(wú)法放進(jìn)寄存器的變量⑶

8、。 CUDA架構(gòu)的所有這些功能都是為了使GPU不僅能執(zhí)行傳統(tǒng)的圖形計(jì)算，還能高效的執(zhí)行通用計(jì)算。 3CUDAC編程3.1CUDA并行編程 運(yùn)行在GPU上的CUDA并行計(jì)算函數(shù)稱為核函數(shù)(kernel)。一個(gè)kernel函數(shù)并不是一個(gè)完整的程序，而是整個(gè)程序中的一個(gè)可以被并行執(zhí)行的步驟⑵。內(nèi)核函數(shù)必須通過_global_函數(shù)類型限定符定義，如_global_voidkernel(void)。并且只能在主機(jī)端代碼中調(diào)用。在調(diào)用時(shí)，必須聲明內(nèi)核函數(shù)的執(zhí)行參數(shù)?，F(xiàn)考慮給一個(gè)N位矢量的每一位加上一個(gè)常數(shù)。為完成此操作設(shè)計(jì)一個(gè)add函數(shù)，分別用C代碼和CUDAC表示為： #defineN10#de

9、fineN10voidadd(int*a,intb){intindex=0;while(indexvN){a[index]=a[index]+b;index+=1; _global_voidadd(int*a,intb){ intindex=blockIdx.x;if(indexvN)a[index]=a[index]+b;} }在CPU運(yùn)行的程序中，通過while循環(huán)串行的給N位矢量每一位加上常數(shù)。值得注意的是右邊的CUDA核函數(shù)，在實(shí)際運(yùn)行時(shí)，CUDA會(huì)產(chǎn)生許多在GPU上執(zhí)行的線程，每一個(gè)線程都會(huì)去執(zhí)行內(nèi)核這個(gè)程序，雖然程序是同一份，但是因?yàn)樵O(shè)置了變量blockIdx，這是一個(gè)內(nèi)置變

10、量，在CUDA運(yùn)行中已經(jīng)預(yù)先定義了這個(gè)變量，變量的值是當(dāng)前執(zhí)行設(shè)備代碼線程塊的索引，從而取得不同的數(shù)據(jù)來(lái)進(jìn)行并行計(jì)算。這里考慮的實(shí)例是10位矢量的增量運(yùn)算，當(dāng)矢量的位數(shù)N遠(yuǎn)遠(yuǎn)大于10且進(jìn)行的并不是簡(jiǎn)單的增量運(yùn)算時(shí)，使用CUDAC就會(huì)有明顯的差別。 intmain(void){ inthost_a[N]; int*device_a; cudaMalloc((void**)&device_a,N*sizeof(int));〃在設(shè)備上分配內(nèi)存 for(inti=0;i

11、nt),cudaMemcpyHostToDevice);addvvvN,l>?(device_a,5); cudaMemcpy(host_a,device_a,N*sizeof(int),cudaMemcpyDeviceToHost); cudaFree(device_a);//釋放內(nèi)存 return0;} 在定義了核函數(shù)后，可以在主函數(shù)中調(diào)用它。addvvvN,l>>>表示啟動(dòng)了N個(gè)線程塊，每個(gè)線程塊中含一個(gè)線程。CPU串行代碼完成的工作包括在核函數(shù)啟動(dòng)前進(jìn)行數(shù)據(jù)準(zhǔn)備和設(shè)備初始化的工作。其中cudaMalloc()函數(shù)用來(lái)在GPU設(shè)備上分配內(nèi)存，第一個(gè)參數(shù)是一個(gè)指針，指向用于保存新分

12、配內(nèi)存地址的變量,第二個(gè)參數(shù)指明分配內(nèi)存的大小，需要注意的是由這個(gè)內(nèi)存分配函數(shù)得到的指針是不能在主機(jī)內(nèi)存上使用的。cudaMemcpy()通過設(shè)置參數(shù)在CPU和GPU之間傳遞數(shù)據(jù)。主函數(shù)先完成數(shù)組初始化的工作，再將數(shù)據(jù)傳入GPU，并行計(jì)算完成后再將結(jié)果傳回CPU。 3.2線程的同步與通信 上一節(jié)矢量增值運(yùn)算的例子，通過創(chuàng)建N個(gè)含有一個(gè)線程的進(jìn)程快給矢量的各位各加上一個(gè)常數(shù)。上述運(yùn)算并不需要考慮線程塊內(nèi)線程的同步與通信，但實(shí)際中的大多數(shù)應(yīng)用都需要線程之間傳遞數(shù)據(jù)。這也是CUDA所提供的最重要的創(chuàng)新，它使得工作在GPU上的線程可以協(xié)作解決問題，允許應(yīng)用程序更加高效的執(zhí)行［2】。 在同一個(gè)bl

13、ock中的線程通過同步函數(shù)_syncthreads()來(lái)保證塊內(nèi)的線程同步即線程塊中的所有線程都執(zhí)行到同一位置。為了在硬件上用很小的代價(jià)就能實(shí)現(xiàn)_syncthreads()函數(shù)，一個(gè)block中所有線程的數(shù)據(jù)都必須交由同一處理核心進(jìn)行處理⑶。所以，這導(dǎo)致每個(gè)線程塊中的線程數(shù)量受到處理核心硬件資源的限制。目前規(guī)定每個(gè)塊里最多只能有512個(gè)線程［2］。 在同一個(gè)塊中的線程通過共享存儲(chǔ)器交換數(shù)據(jù)，為了保證線程塊中的各個(gè)線程能夠有效協(xié)作，訪問共享存儲(chǔ)器的延遲必須很小。所以在GPU中，共享存儲(chǔ)器與執(zhí)行單元的物理距離必須很小。在CUDAC中同一個(gè)塊中的線程通過共享變量來(lái)實(shí)現(xiàn)通信，使用_shared_來(lái)定

14、義一個(gè)共享變量?，F(xiàn)利用一個(gè)大規(guī)模矩陣相乘的例子來(lái)說(shuō)明線程塊內(nèi)線程同步與通信及線程之間的互斥。矩陣A(MxN)與矩陣B(NxK)的乘法公式為［51： C=*axb,i=0,1,...m-1;j=0,1,...k-1(1) ijittit=0為簡(jiǎn)單起見，這里僅考慮1xN矩陣與Nx1矩陣之積，于是公式簡(jiǎn)化為： a1abb..bx.=￡ab(2) 12Nii.i=1aN 可以先讓每個(gè)進(jìn)程計(jì)算aibi，由于要計(jì)算累加和每個(gè)進(jìn)程需要保存一個(gè)臨時(shí)變量，每計(jì)算完一個(gè)積，索引平移blockDim.x*gridDim.x個(gè)單位。這樣線程塊內(nèi)的每個(gè)線程都保存著若干個(gè)aibi，還需要在每個(gè)線程塊內(nèi)設(shè)置一個(gè)共

15、享數(shù)組變量_shared__intvar［threadsperblock］，這個(gè)數(shù)組的大小是線程塊里線程的數(shù)量。這樣線程塊里每個(gè)線程都將自己保存的累加值賦給var［threadsperblock］。在給var賦值完畢后將var數(shù)組累加起來(lái)儲(chǔ)存在var［0］中，這時(shí)需要調(diào)用同步函數(shù)_syncthreads()確保進(jìn)程塊內(nèi)所有進(jìn)程都已完成計(jì)算方可進(jìn)行累加。 至此，進(jìn)程格中的每個(gè)進(jìn)程塊都有了一個(gè)累加和var［0］，需要再將這些屬于各進(jìn)程塊的var［0］再累加起來(lái)儲(chǔ)存在一個(gè)變量c中，這個(gè)操作可由線程塊中任意一個(gè)線程來(lái)完成，不妨設(shè)為線程0。最終將結(jié)果c送回CPU。在第二步的累加中，為保證變量c被互斥的

16、訪問，需要一個(gè)原子鎖lock()。由上面的討論可得核函數(shù)的主要部分： __global__voidmatrix(Locklock,float*a,float*b,float*c){__shared__floatvar［threadsperblock］;intindex=threadIdx.x+blockIdx.x*blockDim.x;intvarnum=threadldx.x;〃初始化索引 floattemp=0; while(index

17、=temp; —syncthreads();〃同步等待其他進(jìn)程 inti=blockDim.x/2; while(i!=0){ if(varnum

18、后就可在主函數(shù)中調(diào)用，main函數(shù)的流程與上節(jié)相似，主要是同時(shí)啟動(dòng)的核函數(shù)數(shù)量不同，matrixvvvblockspergrid,threadsperblock>>>表示啟動(dòng)了blockspergrid個(gè)線程塊，每個(gè)線程塊中有threadsperblock個(gè)線程。 intmain(void){//CPU完成一些數(shù)據(jù)準(zhǔn)備和設(shè)備初始化的工作 Locklock; matrixvvvblockspergrid,threadsperblock?>(lock,device_a,device_b,device_c);cudaMemcpy(&c,device_c,sizeof(float),cudaMe

19、mcpyDeviceToHost));、、、、、、、、、、、、、、/將結(jié)果傳回CPU后釋放申請(qǐng)的內(nèi)存}4結(jié)語(yǔ) GPU在特定的科學(xué)計(jì)算方面比CPU有更強(qiáng)大的運(yùn)算能力，如氣象和醫(yī)學(xué)圖像方面的相關(guān)算法，數(shù)據(jù)量都是非常的大，運(yùn)用CUDA的計(jì)算能力，可以很顯著的提高計(jì)算速度，具有廣闊的應(yīng)用前景。CUDA新技術(shù)為科學(xué)領(lǐng)域進(jìn)行大規(guī)模運(yùn)算提供了新的研究方法，由于GPU的特點(diǎn)是處理密集型數(shù)據(jù)和并行數(shù)據(jù)計(jì)算,因此CUDA非常適合需要大規(guī)模并行計(jì)算的領(lǐng)域。目前CUDA除了可以用C語(yǔ)言開發(fā)，也已經(jīng)提供FORTRAN的應(yīng)用接口，未來(lái)可以預(yù)計(jì)CUDA會(huì)支持C++、Java、Python等各類語(yǔ)言。但是想要充分利用GPU

20、的計(jì)算能力，需要將計(jì)算問題進(jìn)行合理的分解以適應(yīng)GPU的計(jì)算模式，發(fā)揮多級(jí)內(nèi)存和大規(guī)模并行計(jì)算的優(yōu)勢(shì)。 參考文獻(xiàn) [1] ShaneRyoo,ChristopherI.Rodrigues,SaraS.Baghsorkhi,SamS.Stone,DavidB.Kirk,Wen-mei.OptimizationprinciplesandapplicationperformanceevaluationofamultithreadedGPUusingCUDA[J].Proceedingsofthe13thACMSIGPLANSymposiumonPrinciplesandpracticeofpara

21、llelprogramming,2008. [2] JasonSanders,EdwardKandrot.CUDAbyExample[M].AddisonWesleyProfes--sional,2010. ⑶錢悅，圖形處理器CUDA編程模型的應(yīng)用研究[J].計(jì)算機(jī)與數(shù)字工程V36No12,2008. [4]JohnNickolls,IanBuck,MichaelGarland,KevinSkadron.ScalableParallelProgrammingwithCUDA[J].GPUComputingQueueHomepagearchiveVolume6Issue2,2008. ⑸肖江，胡柯良，鄧元勇.矩陣乘法和FFT性能測(cè)試[J].計(jì)算機(jī)工程35卷第10期2009. ⑹劉勇，使用GPU加速通用科學(xué)計(jì)算[J]，科技信息，2008年24期.