CUDA介紹與案例.ppt

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

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

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

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

5、,說明： 在CUDA架構(gòu)下，一個程序分為兩部分，即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ù)，實現(xiàn)配置； 塊組織為一個一維或二維線程塊網(wǎng)格，維度由 語法的第一個參數(shù)指定； 執(zhí)行內(nèi)核的每個線程都會被分配一個獨特的線程 ID，可通過內(nèi)置的 threadIdx 變量在內(nèi)核中訪問,編程模型,2,2.1 線程

6、模型（體系結(jié)構(gòu)） 2.2 存儲器模型（體系結(jié)構(gòu)）,2.1 線程模型,并行線程結(jié)構(gòu)： Thread: 并行的基本單位 索引：threadIdx（內(nèi)置變量） Block (Thread block): 線程塊 允許彼此同步，通過快速共享內(nèi)存交換數(shù)據(jù) 每個塊的線程數(shù)應(yīng)是 warp （調(diào)度任務(wù)的最小單位）塊大小的倍數(shù)，每個塊的線程數(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和同步實現(xiàn)通信;,一個內(nèi)核可能由多個大小相同的線程塊執(zhí)行，因而線程總數(shù)應(yīng)等于每個塊的線程數(shù)乘以塊的數(shù); 線程索引(Index)及ID(IDentity)，索引為(x,y)的線程ID為(x+yDx);對于大小為(Dx,Dy,Dz)的三維塊，索引為(x,y,z)的線程ID為(x+yDx+zDxDy)； 示意圖,線程塊索引及ID ,情況類似： 對于一維塊來說，兩者是相同的； 對于大小為 (D

8、x，Dy) 的二維塊來說，索引為 (x,y) 的線程塊的ID 是 （x + yDx）； 對于大小為 （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); / +配置 ,？哪個線程對哪個數(shù)組下標求和； ! 第i個線程對相應(yīng)的數(shù)組下標求和； ！用內(nèi)置變量確定數(shù)

9、組的下標,或 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)存和線程管理的基本特性 本地存儲器、寄存器的用法 線程ID的用法 主機和設(shè)備之間數(shù)據(jù)傳輸?shù)腁PI 為了方便，以方形矩陣說明,4.2. 矩陣乘法,編程模型,變量和函數(shù),案例,介紹,矩陣大小為 WIDTH x WIDTH 在沒有采用分片優(yōu)化算法的情況下： 一個線程計算P矩陣中的一個 元素 需要從全局存儲器載入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è)備存儲器中的空間 cudaFree(Md); cudaFree(Nd); cudaFree (Pd); ,編程模型,變量和函數(shù),案例,介紹,/ 矩陣乘法的內(nèi)核函數(shù)每個線程都要執(zhí)行的代碼 _global_ void MatrixMulKernel(float* Md, float* Nd, float* Pd, int Width) / 2維的線程ID號 int i = threadIdx.x; int j = threadIdx.y; /Pvalue用來保存被每個線程計算完成后的矩陣的元素 float Pvalue = 0;,編程模型,變量和函數(shù),案例,介紹,/每個線程計算一個元素 for (int k = 0; k Width; +k) float Melement = Mdj* Width+k; float Nelement = Ndk * Width+i; Pvalue += Melement *