基于CUDA的GPU高性能計(jì)算

基于CUDA的GPU高性能計(jì)算

鄒佳林 1420501011主要內(nèi)容GPU以及GPU高性能計(jì)算技術(shù)介紹CUDA編程簡(jiǎn)介基于CUDA的GPU高性能計(jì)算的一個(gè)實(shí)例2一、GPU及GPU高性能計(jì)算技術(shù)1.1GPU簡(jiǎn)介：GPU英文全稱GraphicProcessingUnit，中文翻譯為“圖形處理器。GPU有非常多的廠商都生產(chǎn)，和CPU一樣，生產(chǎn)的廠商比較多，我們大家熟悉的有3個(gè)，分別是AMD、

NVIDIA、Intel這3個(gè)生產(chǎn)廠商。

Intel集成GPU英偉達(dá)TEGRA4GPUAMDRadeonR5M23031.1GPU簡(jiǎn)介GPU的發(fā)展大致可以分成四個(gè)階段：第一個(gè)時(shí)代是1991年以前，CPU作為系統(tǒng)內(nèi)唯一的通用處理器，包攬了圖形處理在內(nèi)的所有計(jì)算任務(wù)。第二個(gè)時(shí)代是1991-2001年。微軟的Windows操作系統(tǒng)極大刺激了圖形硬件的發(fā)展。第三個(gè)時(shí)代是2001-2006年。各種硬件加速計(jì)算的出現(xiàn)使顯卡的性能突飛猛進(jìn)。標(biāo)志性的事件就是可編程的圖形處理器的出現(xiàn)。第四個(gè)時(shí)代是2006年至今，這一時(shí)期的CPU從硬件設(shè)計(jì)之初就開始考慮到了GPGPU的應(yīng)用。2006年NVDIA公布了統(tǒng)一著色器模型（UnifiedShaderModel）和GPU發(fā)展階段GeForce8系列GPU，GPU從此進(jìn)入了通用計(jì)算時(shí)代。4時(shí)間GPU特點(diǎn)1991年以前顯示功能在CPU上實(shí)現(xiàn)1991年~2001年多為二維圖形運(yùn)算，功能單一2001年~2006年可編程圖形處理器2006年至今統(tǒng)一著色器模型，通用計(jì)算GPU1.2GPU與CPU的區(qū)別51.2GPU與CPU的區(qū)別CPU由專為順序串行處理而優(yōu)化的幾個(gè)核心組成。而GPU則由數(shù)以千計(jì)的更小、更高效的核心組成，這些核心專為同時(shí)處理多任務(wù)而設(shè)計(jì)，可高效地處理并行任務(wù)。61.2GPU與CPU的區(qū)別現(xiàn)如今GPU已經(jīng)不再局限于3D圖形處理了，在浮點(diǎn)運(yùn)算、并行計(jì)算等部分計(jì)算方面，GPU可以提供數(shù)十倍乃至于上百倍于CPU的性能。71.2GPU與CPU的區(qū)別1.CPU是計(jì)算機(jī)的運(yùn)算和控制核心，GPU主要用來做圖形處理。2.由于其設(shè)計(jì)目標(biāo)的不同，CPU需要很強(qiáng)的通用性來處理各種不同的數(shù)據(jù)類型，同時(shí)邏輯判斷又會(huì)引入大量的分支跳轉(zhuǎn)和中斷的處理。而GPU面對(duì)的則是類型高度統(tǒng)一的、相互無依賴的大規(guī)模數(shù)據(jù)和不需要被打斷的純凈的計(jì)算環(huán)境。3.CPU與GPU的區(qū)別還存在于片內(nèi)的緩存體系和數(shù)字邏輯運(yùn)算單元的結(jié)構(gòu)差異：CPU雖然有多核，但核心總數(shù)沒有超過16，每個(gè)核都有足夠大的緩存和足夠多的數(shù)字和邏輯運(yùn)算單元，并輔助有很多加速分支判斷甚至更復(fù)雜的邏輯判斷的硬件；GPU的核數(shù)遠(yuǎn)超CPU，如NVIDIAFermi就有512個(gè)核。81.3GPU高性能計(jì)算技術(shù)什么是GPU加速的計(jì)算?GPU加速的計(jì)算是利用一顆圖形處理器(GPU)以及一顆CPU來加速科學(xué)、工程以及企業(yè)

級(jí)應(yīng)用程序。應(yīng)用程序如何利用GPU實(shí)現(xiàn)加速？

密集計(jì)算代碼（約占5%的代碼量）由GPU負(fù)責(zé)完成，剩余串行代碼由CPU負(fù)責(zé)執(zhí)行。91.3GPU高性能計(jì)算技術(shù)GPGPU(GeneralPurposecomputingongraphicsprocessingunits，基于GPU的通用計(jì)算)。GPGPU并不是單純的使用GPU進(jìn)行通用計(jì)算，而是一種利用異構(gòu)計(jì)算資源的大規(guī)模并行計(jì)算。異構(gòu)計(jì)算：CPU+GPU是一個(gè)強(qiáng)大的組合，因?yàn)镃PU包含幾個(gè)專為串行處理而優(yōu)化的核心，而GPU則由數(shù)以千計(jì)更小、更節(jié)能的核心組成，這些核心專為提供強(qiáng)勁的并行性能而設(shè)計(jì)。程序的串行部分在CPU上運(yùn)行，而并行部分則在GPU上運(yùn)行。GPU已經(jīng)發(fā)展到成熟階段，可輕松執(zhí)行現(xiàn)實(shí)生活中的各種應(yīng)用程序，而且程序運(yùn)行速度已遠(yuǎn)遠(yuǎn)超過使用多核系統(tǒng)時(shí)的情形。未來計(jì)算架構(gòu)將是并行核心GPU與多核CPU共同運(yùn)行的混合型系統(tǒng)。101.3GPU高性能計(jì)算技術(shù)計(jì)算模型的分類

1.單指令單數(shù)據(jù)流（SISD）是非并行計(jì)算模型。

2.單指令多數(shù)據(jù)流（SIMD）是GPU的計(jì)算模型。

3.多指令單數(shù)據(jù)流（MISD）指在同一個(gè)數(shù)據(jù)流上執(zhí)行不同的指令。4.多指令多數(shù)據(jù)流（MIMD）是多核CPU的計(jì)算模型。圖：并行處理的費(fèi)林分類法11Single

InstructionMultipleInstructionSingleDataSISDMISDMultipleDataSIMDMIMD1.3GPU高性能計(jì)算技術(shù)并行計(jì)算模型：SIMD

并行計(jì)算指的是在同一時(shí)刻存在多于一個(gè)計(jì)算任務(wù)被執(zhí)行。SIMD的并行思路是讓不同的線程處理它所對(duì)應(yīng)的那部分?jǐn)?shù)據(jù)。當(dāng)線程數(shù)大于或者等于數(shù)據(jù)個(gè)數(shù)時(shí)，理論計(jì)算時(shí)間相當(dāng)于處理一個(gè)數(shù)據(jù)的時(shí)間；如果線程數(shù)少于數(shù)據(jù)個(gè)數(shù)，則某些線程處理的數(shù)量會(huì)增加。

CUDA提出的SIMT（SingleInstructionMultipleThreads）屬于SIMD的范疇，因?yàn)樗彩窃诙鄠€(gè)數(shù)據(jù)上執(zhí)行相同的指令，SIMT允許由用戶來分配線程，具體來說就是CUDA為每個(gè)線程指定了標(biāo)識(shí)符（編號(hào)）。121.3GPU高性能計(jì)算技術(shù)SIMD的兩大特點(diǎn)（即要使用GPU做并行計(jì)算，算法必須滿足以下兩點(diǎn)）：1）每個(gè)線程的任務(wù)互不相關(guān)。2）每個(gè)線程執(zhí)行相同的指令。具有以下特點(diǎn)的算法能夠在GPU上達(dá)到最高的執(zhí)行效率：1）每個(gè)數(shù)據(jù)（數(shù)據(jù)包）都需要經(jīng)過相同的流程來處理。2）數(shù)據(jù)之間沒有相干性，即某些數(shù)據(jù)的計(jì)算不依賴于另外一些數(shù)據(jù)的計(jì)算結(jié)果。3）數(shù)據(jù)量龐大。13二、基于CUDA的GPU編程2.1CUDA簡(jiǎn)介：CUDA(ComputeUnifiedDeviceArchitecture)，是顯卡廠商N(yùn)VIDIA推出的運(yùn)算平臺(tái)。CUDA是一種由NVIDIA推出的通用并行計(jì)算架構(gòu)，該架構(gòu)使GPU能夠解決復(fù)雜的計(jì)算問題。它包含了CUDA指令集架構(gòu)（ISA）以及GPU內(nèi)部的并行計(jì)算引擎。開發(fā)人員現(xiàn)在可以使用C語言來為CUDA架構(gòu)編寫程序，C語言是應(yīng)用最廣泛的一種高級(jí)編程語言。所編寫出的程序于是就可以在支持CUDA的處理器上以超高性能運(yùn)行。142.1CUDA簡(jiǎn)介計(jì)算行業(yè)正在從只使用CPU的“中央處理”向CPU與GPU并用的“協(xié)同處理”發(fā)展。為打造這一全新的計(jì)算典范，NVIDIA發(fā)明了CUDA（ComputeUnifiedDeviceArchitecture，統(tǒng)一計(jì)算設(shè)備架構(gòu)）這一編程模型，是想在應(yīng)用程序中充分利用CPU和GPU各自的優(yōu)點(diǎn)?，F(xiàn)在，該架構(gòu)現(xiàn)已應(yīng)用于GeForce（精視）、ION（翼揚(yáng)）、Quadro以及TeslaGPU（圖形處理器）上，對(duì)應(yīng)用程序開發(fā)人員來說，這是一個(gè)巨大的市場(chǎng)。15CPU+GPU協(xié)同處理2.1CUDA簡(jiǎn)介CUDA體系結(jié)構(gòu)CUDA體系結(jié)構(gòu)開發(fā)庫運(yùn)行環(huán)境驅(qū)動(dòng)16CUDA開發(fā)環(huán)境CUDA開發(fā)環(huán)境nvccC語言編譯器分析器CUDA編程手冊(cè)2.1CUDA簡(jiǎn)介CUDA開發(fā)環(huán)境：

nvccC語言編譯器：

適用于GPU（圖形處理器）的CUDAFFT和BLAS庫。

分析器：

適用于GPU（圖形處理器）的gdb調(diào)試器（在2008年3月推出alpha版）

CUDA運(yùn)行時(shí)（CUDAruntime）驅(qū)動(dòng)程序（目前在標(biāo)準(zhǔn)的NVIDIAGPU驅(qū)動(dòng)中也提供）。CUDA編程手冊(cè)：CUDA開發(fā)者軟件開發(fā)包（SDK）提供了一些范例（附有源代碼），以幫助使用者開始CUDA編程。172.2CUDA執(zhí)行過程CUDA模型的計(jì)算流程大致可以分為四個(gè)部分1）將待處理數(shù)據(jù)從內(nèi)存?zhèn)魉偷斤@存中。2）將程序指令從CPU轉(zhuǎn)移到GPU中。3）在設(shè)備端GPU上執(zhí)行相關(guān)指令，完成對(duì)顯存數(shù)據(jù)的操作。在計(jì)算過程中可能涉及到與內(nèi)存數(shù)據(jù)的頻繁交換，最后將結(jié)果暫存在顯存中。4）將計(jì)算結(jié)果從顯存重新送回內(nèi)存中。182.2CUDA執(zhí)行過程CUDA程序的構(gòu)成一個(gè)CUDAC程序通常由兩部分構(gòu)成：一部分在主機(jī)（CPU）上順序執(zhí)行，另外一部分則在設(shè)備（GPU）上啟動(dòng)成千上萬個(gè)線程并行執(zhí)行，它們?cè)诰幾g過程中由NVIDIA公司的C編譯器（NVCC）區(qū)分開。NVCC簡(jiǎn)化了C語言或PTX的編譯流程：它提供了簡(jiǎn)單的命令行選項(xiàng)，調(diào)用一系列的編譯工具來執(zhí)行它們。NVCC可同時(shí)編譯由主機(jī)代碼（在CPU上執(zhí)行的代碼）和設(shè)備代碼（在GPU上執(zhí)行的代碼）組成的源文件。192.2CUDA執(zhí)行過程CUDA編譯流程分離主機(jī)端代碼和設(shè)備端代碼。將主機(jī)端代碼輸出為C語言代碼供其他工具編譯，或者NVCC在編譯的最后階段調(diào)用主機(jī)編譯器將主機(jī)端代碼輸出為目標(biāo)代碼。編譯設(shè)備端代碼得到其二進(jìn)制形式（cubin對(duì)象）或/和匯編形式（PTX）。將PTX源碼或cubin對(duì)象利用CUDA驅(qū)動(dòng)API裝載并執(zhí)行，或者使用鏈接將PTX源碼或cubin對(duì)象加載到主機(jī)代碼中并將其作為已初始化的全局?jǐn)?shù)據(jù)數(shù)組導(dǎo)入并執(zhí)行。202.2CUDA執(zhí)行過程一般來說，控制從主機(jī)端向設(shè)備端的轉(zhuǎn)移是通過在主機(jī)端調(diào)用由關(guān)鍵字__global__標(biāo)示的且只能從主機(jī)端被調(diào)用的kernel函數(shù)。CUDA程序用類kernelFun<<<M,N>>>(d_a,d_b,d_c);的語句來啟動(dòng)kernel函數(shù)，其中<<<>>>運(yùn)算符中的N和M是主機(jī)端設(shè)置設(shè)備端要啟動(dòng)kernel函數(shù)的參數(shù)，M表示線程塊的數(shù)量（維度），N表示線程塊的大小，(d_a,d_b,d_c)則為kernel函數(shù)的形參，與一般C函數(shù)沒有區(qū)別，該函數(shù)體定義的語句即為后續(xù)每個(gè)線程要執(zhí)行的代碼。212.2CUDA執(zhí)行過程當(dāng)kernel函數(shù)啟動(dòng)運(yùn)行后，執(zhí)行過程轉(zhuǎn)移到設(shè)備端（GPU），然后生成指定的線程數(shù)目，這些線程被組織成塊（block）不同架構(gòu)的支持CUDA的GPU一個(gè)block所能容納的最大線程數(shù)目也不同，有512個(gè)（Tesla架構(gòu)）也有1024個(gè)（Fermi架構(gòu)、Kepler架構(gòu)），塊最后被組織成一個(gè)線程網(wǎng)格（grid）。每調(diào)用一次kernel函數(shù)會(huì)生成一個(gè)線程網(wǎng)格，當(dāng)kernel函數(shù)中的所有線程都完成他們的執(zhí)行任務(wù)后，相應(yīng)的網(wǎng)格也會(huì)終止，并且在調(diào)用下一個(gè)kernel函數(shù)前，程序會(huì)轉(zhuǎn)到主機(jī)端繼續(xù)執(zhí)行，即kernel函數(shù)和主機(jī)端代碼是異步執(zhí)行的。222.3CUDA的線程組織一般而言，在啟動(dòng)kernel函數(shù)時(shí)我們會(huì)把網(wǎng)格中的線程塊組織成二維數(shù)組形式，將線程塊中的線程組織成三維數(shù)組的形式。通過CUDAC拓展定義的一個(gè)類似C結(jié)構(gòu)的數(shù)據(jù)類型dim3，它有3個(gè)無符號(hào)整數(shù)型字段，分別是x——標(biāo)示x方向上線程或線程塊的索引、y——標(biāo)示y方向上線程或線程塊的索引和z——標(biāo)示z方向上線程或線程塊的索引。同一個(gè)kernel中所有的線程都會(huì)執(zhí)行kernel函數(shù)體定義的語句，唯一不同的是每個(gè)線程都要操作數(shù)據(jù)，所以基于CUDA的GPU計(jì)算模型為單指令多數(shù)據(jù)流（SIMD）。232.3CUDA的線程組織CUDA通過內(nèi)置的一些預(yù)初始化變量——threadIdx.x、threadIdx.y、threadIdx.z標(biāo)示當(dāng)前線程所處的線程塊的位置，以及blockIdx.x、blockIdx.y、blockIdx.z標(biāo)示當(dāng)前線程所處的線程塊在整個(gè)網(wǎng)格中所處的位置，gridDim.x、gridDim.y、gridDim.z標(biāo)示網(wǎng)格的維度和blockDim.x、blockDim.y、blockDim.z標(biāo)示每個(gè)塊的維度；使用上面這些CUDA提供的內(nèi)置變量可以在同一個(gè)kernel函數(shù)中將各個(gè)塊中的線程彼此區(qū)分開來，然后決定哪個(gè)線程要處理哪些數(shù)據(jù)。各個(gè)線程塊間布局是相互獨(dú)立的，所以不同線程塊中的相對(duì)應(yīng)位置的線程具有相同的threadIdx.x、threadIdx.y和threadIdx.z，所有的線程塊擁有相同的線程數(shù)目。242.3CUDA的線程組織對(duì)網(wǎng)格中任意點(diǎn)(i,j)，在CUDA代碼中表征一個(gè)線程，該線程對(duì)應(yīng)的網(wǎng)格中的索引可以使用下面公式進(jìn)行表示：i=threadIdx.x+blockIdx.x*blockDim.xj=threadIdx.y+blockIdx.y*blockDim.y252.3CUDA的并行優(yōu)化存儲(chǔ)器訪問優(yōu)化因?yàn)橹鳈C(jī)內(nèi)存和設(shè)備端顯存間傳遞數(shù)據(jù)的過程中耗時(shí)更長(zhǎng)，所以在進(jìn)行CUDA運(yùn)算過程中，應(yīng)盡量減少CPU和GPU間的數(shù)據(jù)傳輸，充分發(fā)揮GPU多線程并行計(jì)算的優(yōu)勢(shì)將更多的操作放置在設(shè)備端完成；盡量減少對(duì)全局存儲(chǔ)器的訪問，更多的使用共享存儲(chǔ)器。指令優(yōu)化1.在kernel函數(shù)中如果不是必要的盡量不要用__syncthreads()語句進(jìn)行同步，另外，同一個(gè)warp中的線程不需要同步。2.遇到迭代累加的時(shí)候，如果給定累加的次數(shù)，可以通過將迭代過程展開的方式消除循環(huán)計(jì)數(shù)更新指令和分支指令。3.任何控制指令（if、else、switch、do、while、for）都會(huì)致使同一個(gè)warp中的線程產(chǎn)生分支。所以我們應(yīng)盡量減少一個(gè)warp中線程產(chǎn)生分支的情況。26三、基于CUDA的GPU高性能計(jì)算的一個(gè)實(shí)例從矩陣相乘開始初始矩陣M,N,我們從矩陣M中取一行，矩陣N中取一列進(jìn)行點(diǎn)積運(yùn)算，從而得到結(jié)果矩陣P中的每個(gè)元素。從右圖可以看到，P中不同的元素的點(diǎn)積運(yùn)算是可以同時(shí)進(jìn)行的。也就是說，這些點(diǎn)積運(yùn)算之間互相不影響。27在C語言中實(shí)現(xiàn)矩陣乘法矩陣乘法中一個(gè)簡(jiǎn)單的主函數(shù)intmain(void){ 1.//分配和初始化矩陣M、N、P //執(zhí)行I/O讀取輸入矩陣M、N 2.//在設(shè)備上執(zhí)行M*N 3.//執(zhí)行I/O寫入輸出矩陣P //釋放矩陣M、N、P的存儲(chǔ)空間 ... return0；}我們?cè)趺磳?shí)現(xiàn)矩陣M、N相乘呢？for(introw=0;row<row1;row++) { for(intcol=0;col<col2;col++) { intnum=col1;// for(intk=0;k<num;k++) result[row][col]+=p1[row][k]*p2[k][col]; } }2829

在CUDA上實(shí)現(xiàn)矩陣乘法

對(duì)于大型矩陣的乘法，點(diǎn)積的個(gè)數(shù)可能會(huì)非常大，如兩個(gè)1000*1000的矩陣乘法就有1000000個(gè)獨(dú)立的點(diǎn)積，每個(gè)點(diǎn)積涉及1000次乘法和1000次累加運(yùn)算。因此，高維的矩陣乘法擁有大量的數(shù)據(jù)并行性。因此，我們可以用GPU進(jìn)行高性能計(jì)算。voidMatrixMulKernel(float*M,float*N,float*P,intwidth){1.//把M和N傳遞到設(shè)備存儲(chǔ)器中 cudaMalloc((void**)&Md,size); cudaMemcpy(Md,M,size,cudaMemcpyHostToDevice); cudaMalloc((void**)&Nd,size); cudaMemcpy(Nd,N,size,cudaMemcpyHostToDevice);//在設(shè)備上分配P cudaMalloc((void**)&Pd,size)2.//調(diào)用kernel的代碼...3.//把P從設(shè)備上傳遞到主機(jī)上 cudaMemcpy(P,Pd,size,cudaMemcpyDeviceToHost);//釋放設(shè)備上的矩陣 cudaFree(Md); cudaFree(Nd); cudaFree(Pd);}30Kernel函數(shù)：__global__voidMatrixMulKernel(int*dev_M,int*dev_N,int*dev_P,introw){ //計(jì)算P和M中元素的行索引

intRow=blockIdx.y*16+threadIdx.y; //計(jì)算P和N中元素的列索引

intCol=blockIdx.x*16+threadIdx.x; intPvalue=0; //每個(gè)線程負(fù)責(zé)計(jì)算塊子矩陣的一個(gè)元素

if(Row<row&&Col<row) { for(intk=