基于CUDA的GPU高性能計算

基于CUDA的GPU高性能計算

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

NVIDIA、Intel這3個生產廠商。

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

級應用程序。應用程序如何利用GPU實現(xiàn)加速？

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

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

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

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

InstructionMultipleInstructionSingleDataSISDMISDMultipleDataSIMDMIMD1.3GPU高性能計算技術并行計算模型：SIMD

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

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

nvccC語言編譯器：

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

分析器：

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

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

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

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

intCol=blockIdx.x*16+threadIdx.x; intPvalue=0; //每個線程負責計算塊子矩陣的一個元素

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