CUDA介紹與案例.ppt

1、CUDA 介紹與案例,1,介紹,應(yīng)用領(lǐng)域： 游戲、圖形動(dòng)畫、科學(xué)計(jì)算可視化、 地質(zhì)、生物、物理模擬等；,編程模型,變量和函數(shù),案例,介紹,2008年SIGGRAPH年會(huì)上，NVIDIA公司推出CUDA (Compute Unified Device Architecture) NVIDIA2008; CUDA是NVIDIA為自己的GPU編寫的一套編譯器及相關(guān)的庫(kù)文件； 將GPU 視作數(shù)據(jù)并行計(jì)算設(shè)備,是一種新的處理和管理GPU 計(jì)算的硬件和軟件架構(gòu)；,編程模型,變量和函數(shù),案例,介紹,GPU到CUDA,編程模型,變量和函數(shù),案例,介紹,CPU 和 GPU 的 每秒浮點(diǎn)運(yùn)算次數(shù) 和存儲(chǔ)器帶寬比較

2、,具有強(qiáng)大浮點(diǎn) 運(yùn)算能力,GPU采用大量的執(zhí)行單元，這些執(zhí)行單元可以輕松的加載并行處理線程，而不像CPU那樣的單線程處理； 主機(jī)和設(shè)備均維護(hù)自己的 DRAM，分別稱為主機(jī)存儲(chǔ)器和設(shè)備存儲(chǔ)器,CPU 和 GPU 之間浮點(diǎn)功能的差異，原因在于 GPU 專為高計(jì)算密集型（數(shù)學(xué)運(yùn)算與存儲(chǔ)器運(yùn)算的比率）、高度并行化的計(jì)算而設(shè)計(jì)； GPU 的設(shè)計(jì)能使更多晶體管用于數(shù)據(jù)處理，而非數(shù)據(jù)緩存和流控制 ；,編程模型,變量和函數(shù),案例,介紹,高度并行化、多線程、多核處理器，操作 系統(tǒng)的多任務(wù)機(jī)制負(fù)責(zé)管理多個(gè)并發(fā)運(yùn)行 的CUDA和應(yīng)用程序?qū)PU的訪問； 基于標(biāo)準(zhǔn)C語(yǔ)言, 可自由地調(diào)用GPU的并行 處理架構(gòu); 同時(shí)適

3、用于圖形和通用并行計(jì)算應(yīng)用程序， 成為圖形處理器的主要發(fā)展趨勢(shì)。,編程模型,變量和函數(shù),案例,介紹,依次安裝 Cuda Driver Cuda Toolkit Cuda SDK 在安裝目錄中包括： bin工具程序及動(dòng)態(tài)鏈接庫(kù) doc文件 include 頭文件 lib 鏈接庫(kù) open64基于open64的CUDA compiler src 一些自帶例子，如 C: CUDA_SDKCsrc 安裝程序會(huì)設(shè)定一些環(huán)境變量： CUDA_BIN_PATH CUDA_INC_PATH CUDA_LIB_PATH,編程模型,變量和函數(shù),案例,介紹,CUDA 軟件棧包含多個(gè)層，如圖所示：設(shè)備驅(qū)動(dòng)程序、應(yīng)用程

4、序編程接口（API）及其運(yùn)行時(shí)兩個(gè)較高級(jí)別的通用數(shù)學(xué)庫(kù)，即 CUFFT 和 CUBLAS,CUDA安裝后,編程模型,變量和函數(shù),案例,介紹,CUDA程序編譯,CUDA的源文件被nvcc編譯 NVCC編譯器會(huì)分離源碼中設(shè)備代碼和主機(jī)代碼，主機(jī)代碼交由一般的C/C+編譯器（gcc等）編譯，設(shè)備代碼由NVCC編譯； 內(nèi)核必須使用nvcc編譯成二進(jìn)制代碼才能在設(shè)備端執(zhí)行。,編程模型,變量和函數(shù),案例,介紹,內(nèi)核函數(shù),編程模型,變量和函數(shù),案例,介紹,串行代碼在主機(jī)上執(zhí)行，而并行代碼在設(shè)備上執(zhí)行 當(dāng)調(diào)用kernel的時(shí)候，則CUDA的每個(gè)線程并行地執(zhí)行一段指令，這與通常C函數(shù)只執(zhí)行一次不一樣。,執(zhí)行模型

5、,說明： 在CUDA架構(gòu)下，一個(gè)程序分為兩部分，即host 端和 device端； Host 端 是在 CPU上執(zhí)行的部分； Device端 是在顯卡芯片上執(zhí)行的部分，該端程序又稱為內(nèi)核 函數(shù)（ kernel function） 通常Host端程序復(fù)制內(nèi)核函數(shù)到顯卡內(nèi)存中，再由顯卡芯片執(zhí)行device端程序，完成后再由host端程序?qū)⒔Y(jié)果從顯卡內(nèi)存中取回；,CUDA 允許程序員定義內(nèi)核函數(shù)，實(shí)現(xiàn)配置； 塊組織為一個(gè)一維或二維線程塊網(wǎng)格，維度由 語(yǔ)法的第一個(gè)參數(shù)指定； 執(zhí)行內(nèi)核的每個(gè)線程都會(huì)被分配一個(gè)獨(dú)特的線程 ID，可通過內(nèi)置的 threadIdx 變量在內(nèi)核中訪問,編程模型,2,2.1 線程

6、模型（體系結(jié)構(gòu)） 2.2 存儲(chǔ)器模型（體系結(jié)構(gòu)）,2.1 線程模型,并行線程結(jié)構(gòu)： Thread: 并行的基本單位 索引：threadIdx（內(nèi)置變量） Block (Thread block): 線程塊 允許彼此同步，通過快速共享內(nèi)存交換數(shù)據(jù) 每個(gè)塊的線程數(shù)應(yīng)是 warp （調(diào)度任務(wù)的最小單位）塊大小的倍數(shù)，每個(gè)塊的線程數(shù)受限 索引：blockIdx（內(nèi)置變量） Grid: 一組thread block 共享全局內(nèi)存 Kernel: 在GPU上執(zhí)行的核心程序 One kernel One grid,說明： Grid 是一組Block, 以1維、2維或3維組織，共享全局內(nèi)存； Grid之間通過

7、global memory交換數(shù)據(jù) Block 是互相協(xié)作的線程組，通過global memory共享數(shù)據(jù)，允許彼此同步；以1維、2維或3維組織； 同一block內(nèi)的thread可以通過shared memory和同步實(shí)現(xiàn)通信;,一個(gè)內(nèi)核可能由多個(gè)大小相同的線程塊執(zhí)行，因而線程總數(shù)應(yīng)等于每個(gè)塊的線程數(shù)乘以塊的數(shù); 線程索引(Index)及ID(IDentity)，索引為(x,y)的線程ID為(x+yDx);對(duì)于大小為(Dx,Dy,Dz)的三維塊，索引為(x,y,z)的線程ID為(x+yDx+zDxDy)； 示意圖,線程塊索引及ID ,情況類似： 對(duì)于一維塊來說，兩者是相同的； 對(duì)于大小為 (D

8、x，Dy) 的二維塊來說，索引為 (x,y) 的線程塊的ID 是 （x + yDx）； 對(duì)于大小為 （Dx，Dy， Dz） 的三維塊來說，索引 為（x, y, z） 的線程的ID 是 （x + yDx +zDxDy）; /+圖示說明,例: 向量的和 _global_ void vecAdd(float* A, float* B, float* C) / 內(nèi)核函數(shù) int i = threadIdx.x; Ci = Ai + Bi； int main( ) vecAdd(dA, dB,dC); / +配置 ,？哪個(gè)線程對(duì)哪個(gè)數(shù)組下標(biāo)求和； ! 第i個(gè)線程對(duì)相應(yīng)的數(shù)組下標(biāo)求和； ！用內(nèi)置變量確定數(shù)

9、組的下標(biāo),或 int main( ) int threadsPerBlock = 256; int blocksPerGrid = (N + threadsPerBlock - 1) /threadsPerBlock; VecAdd(d_A, d_B, d_C, N); ,_global_ void VecAdd(const float* A, const float* B, float* C, int N) / 內(nèi)核函數(shù) int i = blockDim.x * blockIdx.x + threadIdx.x; if (i (A, B, C); 參見例：matrix_src.doc,編程模

10、型,變量和函數(shù),案例,介紹,方法二： int main( ) int threadsPerBlock = 256; int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock; VecAdd(d_A, d_B, d_C, N); ,編程模型,變量和函數(shù),案例,介紹,_global_ void VecAdd(const float* A, const float* B, float* C, int N) int i = blockDim.x * blockIdx.x + threadIdx.x; if (i (Ad, Bd, C

11、d); . ,編程模型,變量和函數(shù),案例,介紹,說明內(nèi)存和線程管理的基本特性 本地存儲(chǔ)器、寄存器的用法 線程ID的用法 主機(jī)和設(shè)備之間數(shù)據(jù)傳輸?shù)腁PI 為了方便，以方形矩陣說明,4.2. 矩陣乘法,編程模型,變量和函數(shù),案例,介紹,矩陣大小為 WIDTH x WIDTH 在沒有采用分片優(yōu)化算法的情況下： 一個(gè)線程計(jì)算P矩陣中的一個(gè) 元素 需要從全局存儲(chǔ)器載入WIDTH次,方塊矩陣乘法,編程模型,變量和函數(shù),案例,介紹,CPU上的矩陣乘法 void MatrixMulOnHost(float* M, float* N, float* P, int Width) for (int i = 0; i

12、 Width; +i) for (int j = 0; j (Md, Nd, Pd，Width); / 從設(shè)備中讀取P矩陣的數(shù)據(jù) cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost); / 釋放設(shè)備存儲(chǔ)器中的空間 cudaFree(Md); cudaFree(Nd); cudaFree (Pd); ,編程模型,變量和函數(shù),案例,介紹,/ 矩陣乘法的內(nèi)核函數(shù)每個(gè)線程都要執(zhí)行的代碼 _global_ void MatrixMulKernel(float* Md, float* Nd, float* Pd, int Width) / 2維的線程ID號(hào) int i = threadIdx.x; int j = threadIdx.y; /Pvalue用來保存被每個(gè)線程計(jì)算完成后的矩陣的元素 float Pvalue = 0;,編程模型,變量和函數(shù),案例,介紹,/每個(gè)線程計(jì)算一個(gè)元素 for (int k = 0; k Width; +k) float Melement = Mdj* Width+k; float Nelement = Ndk * Width+i; Pvalue += Melement *