CUDA并行計算架構及編程

CUDA并行架構及編程 摘要CUDA是一種由NVIDIA推出的并行計算架構，非常適合大規(guī)模數據密集型計算。CUDA使GPU的超高計算性能在數據處理和并行計算等通用計算領域發(fā)揮優(yōu)勢，本文討論了CUDA的計算架構和基于GPU的CUDAC編程語言,CUDA使GPU流處理器陣列的性能得到充分發(fā)揮。極大地提高了并行計算程序的效率。 關鍵詞并行計算，GPU通用計算，CUDAAbstractCUDAisaparallelcomputingarchitectureintroducedbyNVIDIA,itmainlyusedforlargescaledata-intensivecomputing.CUDAmakesGPUahighperformanceinparallelcomputing,dataprocessingandothergeneralcomputing.thispaperdiscussestheCUDAcomputingarchitectureandCUDACprogramminglanguagebasedonGPU,CUDAmakesGPUstreamprocessorarraysfullyusedandGreatlyimprovedtheefficiencyofparallelcomputingprogram. Keywordsparallelcomputing，GPUgeneralpurposecomputation，CUDA1引言 并行計算是指同時使用多種計算資源解決計算問題的過程。并行計算科學中主要研究的是空間上的并行問題。從程序和算法設計的角度來看，并行計算又可分為數據并行和任務并行。一般來說，GPU更注重于數據并行計算，主要是將一個大任務化解成相同的各個子任務。早期的GPU研究是通過可編程計算單元為屏幕上的每個像素計算出一個顏色值即渲染問題。自CUDAC出現后，基于GPU的通用計算已經成為一個新的研究領域。通常，像素著色器對各種顏色值進行合成并計算出最終的顏色值。實際上，輸入值可以為任意數據，這樣不一定非要使用GPU來處理圖形，還可以實現某些通用計算。由于GPU有著很高的計算吞吐量，從而給大規(guī)模的數據計算應用提供了一種比CPU更加強大的計算能力。 CUDA是一種由NVIDIA推出的并行計算架構，該架構使GPU能夠解決復雜的計算問題。它包含了CUDA指令集架構以及GPU內部的并行計算引擎⑴。隨著顯卡的發(fā)展，GPU越來越強大，在計算上已經超越了通用的CPU。如此強大的芯片如果只是作為顯卡會造成計算能力的浪費，因此NVIDIA推出CUDA，讓顯卡可以用于圖像渲染以外的目的。CUDA的GPU編程語言基于標準的C語言，通過在標準C語言的基礎上增加一小部分關鍵字，任何有C語言基礎的用戶都很容易地開發(fā)CUDA的應用程序。數以千計的軟件開發(fā)人員正在使用免費的CUDA軟件開發(fā)工具來解決各種專業(yè)中的問題［©。這些解決方案涵蓋了石油天然氣勘探、產品設計、醫(yī)學成像以及科學研究等領域。 2CUDA架構 CUDA程序架構分為兩部分：主機和設備。一般而言，主機指的是CPU及其內存，設備指的是GPUM］。在CUDA程序架構中，主程序由CPU來執(zhí)行，而當遇到數據并行處理的部分，CUDA就會將程序編譯成GPU能執(zhí)行的程序，并傳送到GPU。這種函數在CUDA里叫做核函數。在GPU中要執(zhí)行的線程，根據最有效的數據共享來創(chuàng)建線程塊，其類型不止一維。在同一個塊里的線程，使用同一個共享內存。每個線程由線程ID標識。這是線程塊中的線程號。為了幫助基于線程ID的復雜尋址，還可以將線程塊指定為任意大小的一維、二維或高維線程陣列，并使用多個索引分量來標識每個線程⑵。線程塊的大小是也有限制的,所以不能把所有的線程都放到同一個塊里?？梢杂猛瑯泳S度和大小的塊來組成一個網格做批處理。執(zhí)行內核的線程被組織成線程塊。而線程塊又組成了線程格，如圖1。線程塊內一起協作的線程通過一些快速的共享內存有效地共享數據并同步執(zhí)行，以協調內存訪問。用戶可以在核函數中指定同步點。線程塊中的線程在到達此同步點時掛起。 圖1線程被組織為線程格與線程塊 在GPU內部，SM代表流處理器，即計算核心。每個SM中又包含8個標量流處理器（SP）以及少量的其他計算單元⑶，如圖2。實際上SP只是執(zhí)行單元，并不是完整的處理核心。處理核心必須包含取指、解碼、分發(fā)邏輯和執(zhí)行單元。隸屬同一SM的8個SP共用同一套取指和發(fā)射單元，也共用一塊共享存儲器。一個線程塊必須被分配到一個SM中，但是一個SM中同一時刻可以有多個活動線程塊等待執(zhí)行。這可以更好地利用執(zhí)行單元的資源，當一個線程塊進行同步或者訪問顯存等高延遲操作時，另一個線程塊就可以占用GPU執(zhí)行資源。每個線程有一個私有的本地存儲器，即每個標量處理器私有的緩存。每個線程塊有一個共享存儲器，塊內所有的線程都可以訪問。每個線程網格都可以訪問設備內存，即GPU片上的全局存儲器。每個線程均可以讀寫全局存儲器。本地存儲器用于儲存設備代碼中無法放進寄存器的變量⑶。 CUDA架構的所有這些功能都是為了使GPU不僅能執(zhí)行傳統(tǒng)的圖形計算，還能高效的執(zhí)行通用計算。 3CUDAC編程3.1CUDA并行編程 運行在GPU上的CUDA并行計算函數稱為核函數(kernel)。一個kernel函數并不是一個完整的程序，而是整個程序中的一個可以被并行執(zhí)行的步驟⑵。內核函數必須通過_global_函數類型限定符定義，如_global_voidkernel(void)。并且只能在主機端代碼中調用。在調用時，必須聲明內核函數的執(zhí)行參數?，F考慮給一個N位矢量的每一位加上一個常數。為完成此操作設計一個add函數，分別用C代碼和CUDAC表示為： #defineN10#defineN10voidadd(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運行的程序中，通過while循環(huán)串行的給N位矢量每一位加上常數。值得注意的是右邊的CUDA核函數，在實際運行時，CUDA會產生許多在GPU上執(zhí)行的線程，每一個線程都會去執(zhí)行內核這個程序，雖然程序是同一份，但是因為設置了變量blockIdx，這是一個內置變量，在CUDA運行中已經預先定義了這個變量，變量的值是當前執(zhí)行設備代碼線程塊的索引，從而取得不同的數據來進行并行計算。這里考慮的實例是10位矢量的增量運算，當矢量的位數N遠遠大于10且進行的并不是簡單的增量運算時，使用CUDAC就會有明顯的差別。 intmain(void){ inthost_a[N]; int*device_a; cudaMalloc((void**)&device_a,N*sizeof(int));〃在設備上分配內存 for(inti=0;i<N;i++)host_a[i]=i;〃為數組賦值cudaMemcpy(device_a,host_a,N*sizeof(int),cudaMemcpyHostToDevice);addvvvN,l>»(device_a,5); cudaMemcpy(host_a,device_a,N*sizeof(int),cudaMemcpyDeviceToHost); cudaFree(device_a);//釋放內存 return0;} 在定義了核函數后，可以在主函數中調用它。addvvvN,l>>>表示啟動了N個線程塊，每個線程塊中含一個線程。CPU串行代碼完成的工作包括在核函數啟動前進行數據準備和設備初始化的工作。其中cudaMalloc()函數用來在GPU設備上分配內存，第一個參數是一個指針，指向用于保存新分配內存地址的變量,第二個參數指明分配內存的大小，需要注意的是由這個內存分配函數得到的指針是不能在主機內存上使用的。cudaMemcpy()通過設置參數在CPU和GPU之間傳遞數據。主函數先完成數組初始化的工作，再將數據傳入GPU，并行計算完成后再將結果傳回CPU。 3.2線程的同步與通信 上一節(jié)矢量增值運算的例子，通過創(chuàng)建N個含有一個線程的進程快給矢量的各位各加上一個常數。上述運算并不需要考慮線程塊內線程的同步與通信，但實際中的大多數應用都需要線程之間傳遞數據。這也是CUDA所提供的最重要的創(chuàng)新，它使得工作在GPU上的線程可以協作解決問題，允許應用程序更加高效的執(zhí)行［2】。 在同一個block中的線程通過同步函數_syncthreads()來保證塊內的線程同步即線程塊中的所有線程都執(zhí)行到同一位置。為了在硬件上用很小的代價就能實現_syncthreads()函數，一個block中所有線程的數據都必須交由同一處理核心進行處理⑶。所以，這導致每個線程塊中的線程數量受到處理核心硬件資源的限制。目前規(guī)定每個塊里最多只能有512個線程［2］。 在同一個塊中的線程通過共享存儲器交換數據，為了保證線程塊中的各個線程能夠有效協作，訪問共享存儲器的延遲必須很小。所以在GPU中，共享存儲器與執(zhí)行單元的物理距離必須很小。在CUDAC中同一個塊中的線程通過共享變量來實現通信，使用_shared_來定義一個共享變量?，F利用一個大規(guī)模矩陣相乘的例子來說明線程塊內線程同步與通信及線程之間的互斥。矩陣A(MxN)與矩陣B(NxK)的乘法公式為［51： C=*axb,i=0,1,...m-1;j=0,1,...k-1(1) ijittit=0為簡單起見，這里僅考慮1xN矩陣與Nx1矩陣之積，于是公式簡化為： a1abb..bx.=£ab(2) 12Nii.i=1aN 可以先讓每個進程計算aibi，由于要計算累加和每個進程需要保存一個臨時變量，每計算完一個積，索引平移blockDim.x*gridDim.x個單位。這樣線程塊內的每個線程都保存著若干個aibi，還需要在每個線程塊內設置一個共享數組變量_shared__intvar［threadsperblock］，這個數組的大小是線程塊里線程的數量。這樣線程塊里每個線程都將自己保存的累加值賦給var［threadsperblock］。在給var賦值完畢后將var數組累加起來儲存在var［0］中，這時需要調用同步函數_syncthreads()確保進程塊內所有進程都已完成計算方可進行累加。 至此，進程格中的每個進程塊都有了一個累加和var［0］，需要再將這些屬于各進程塊的var［0］再累加起來儲存在一個變量c中，這個操作可由線程塊中任意一個線程來完成，不妨設為線程0。最終將結果c送回CPU。在第二步的累加中，為保證變量c被互斥的訪問，需要一個原子鎖lock()。由上面的討論可得核函數的主要部分： __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<N){temp+=a[index]*b[index];index+=blockDim.x*gridDim.x;〃索引偏移} var[varnum]=temp; —syncthreads();〃同步等待其他進程 inti=blockDim.x/2; while(i!=0){ if(varnum<i)var[varnum]+=var[varnum+i];__syncthreads();i/=2;}訐(varnum==0){〃互斥訪問lock.lock(); *c+=var[0];lock.unlock();} } 在上面的核函數最后一部分使用了一個Lock結構，這個結構中包含了一個atomicCAS函數，它將判斷和賦值的過程原子化[2]。使用Lock結構的好處是連加的過程可以在GPU上進行，CPU得到的是最后結果。核函數定義好后就可在主函數中調用，main函數的流程與上節(jié)相似，主要是同時啟動的核函數數量不同，matrixvvvblockspergrid,threadsperblock>>>表示啟動了blockspergrid個線程塊，每個線程塊中有threadsperblock個線程。 intmain(void){//CPU完成一些數據準備和設備初始化的工作 Locklock; matrixvvvblockspergrid,threadsperblock»>(lock,device_a,device_b,device_c);cudaMemcpy(&c,device_c,sizeof(float),cudaMemcpyDeviceToHost));、、、、、、、、、、、、、、/將結果傳回CPU后釋放申請的內存}4結語 GPU在特定的科學計算方面比CPU有更強大的運算能力，如氣象和醫(yī)學圖像方面的相關算法，數據量都是非常的大，運用CUDA的計算能力，可以很顯著的提高計算速度，具有廣闊的應用前景。CUDA新技術為科學領域進行大規(guī)模運算提供了新的研究方法，由于GPU的特點是處理密集型數據和并行數據計算,因此CUDA非常適合需要大規(guī)模并行計算的領域。目前CUDA除了可以用C語言開發(fā)，也已經提供FORTRAN的應用接口，未來可以預計CUDA會支持C++、Java、Python等各類語言。但是想要充分利用GPU的計算能力，需要將計算問題進行合理的分解以適應GPU的計算模式，發(fā)揮多級內存和大規(guī)模并行計算的優(yōu)勢。 參考文獻 [1] ShaneRyoo,ChristopherI.Rodrigues,SaraS.Baghsorkhi,SamS.Stone,DavidB.Kirk,Wen-mei.OptimizationprinciplesandapplicationperformanceevaluationofamultithreadedGPUusingCUDA[J].Proceedingsofthe13thACMSIGPLANSymposiumonPrinciplesandpracticeofparallelprogramming,2008. [2] JasonSanders,EdwardKandrot.CUDAbyExample[M].AddisonWesleyProfes--sional,2010. ⑶錢悅，圖形處理器CUDA編程模型的應用研究[J].計算機與數字工程V36No12,2008. [4]JohnNickolls,IanBuck,MichaelGarland,KevinSkadron.ScalableParallelProgrammingwithCUDA[J].GPUComputingQueueHomepagearchiveVolume6Issue2,2008. ⑸肖江，胡柯良，鄧元勇.矩陣乘法和FFT性能測試[J].計算機工程35卷第10期2009. ⑹劉勇，使用GPU加速通用科學計算[J]，科技信息，2008年24期.