1、CUDA 介紹與案例,1,介紹,應(yīng)用領(lǐng)域： 游戲、圖形動畫、科學計算可視化、 地質(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ù)學運算與存儲器運算的比率）、高度并行化的計算而設(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ù)學庫，即 CUFFT 和 CUBLAS,CUDA安裝后,編程模型,變量和函數(shù),案例,介紹,CUDA程序編譯,CUDA的源文件被nvcc編譯 NVCC編譯器會分離源碼中設(shè)備代碼和主機代碼，主機代碼交由一般的C/C+編譯器（gcc等）編譯，設(shè)備代碼由NVCC編譯； 內(nèi)核必須使用nvcc編譯成二進制代碼才能在設(shè)備端執(zhí)行。,v_2011,編程模型,變量和函數(shù),案例,介紹,內(nèi)核函數(shù),編程模型,變量和函數(shù),案例,介紹,串行代碼在主機上執(zhí)行，而并行代碼在設(shè)備上執(zhí)行 當調(diào)用kernel的時候，則CUDA的每個線程并行地執(zhí)行一段指令，這與通常C函數(shù)只執(zhí)行一次不一

5、樣。,執(zhí)行模型,說明： 在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

6、,2.1 線程模型（體系結(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)存； G

7、rid之間通過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 ,情況類似： 對于一維塊來說，兩者是相同的； 對

8、于大小為 (Dx，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ù)組下標求和； ！用

9、內(nèi)置變量確定數(shù)組的下標,或 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 N) Ci = Ai + Bi; /參見例 ：vec

10、torAdd!,？哪個線程對哪個數(shù)組下標求和； 即第幾個Block（blockIdx）的第幾個線程，對相應(yīng)的數(shù)組下標求和；,2.2,說明： 所有線程都可訪問全局存儲器； 所有的線程都可以訪問只讀存儲區(qū)域：“常量存儲區(qū)”和“紋理存儲區(qū)”； 紋理存儲器提供不同的尋址模式，為某些特定的數(shù)據(jù)格式提供數(shù)據(jù)過濾的能力； Block中所有thread都在同一個處理核心上運行。因此每個block的thread數(shù)目受到處理核心所擁有存儲器資源的限制，在NVIDIA Tesla架構(gòu)中，一個線程block最多可包含512個線程;,每個線程塊都有一個共享存儲器(Shared Memory)，該存儲器對于塊內(nèi)的所有線程

11、都是可見的，并且與塊具有相同的生命周期。 每個線程都有一個私有的本地存儲器（LocalMemory）。 在同一個block中的thread通過共享存儲器共享數(shù)據(jù)，相互協(xié)作，實現(xiàn)同步： _syncthreads( );,register 訪問延遲極低； 每個線程占有的register有限，編程時不要為其分配過多 私有變量； local memory 寄存器被使用完畢，數(shù)據(jù)將被存儲在局部存儲器中； 此外，存儲大型結(jié)構(gòu)體或者數(shù)組（無法確定大小的數(shù)組），線程的輸入和中間變量；,shared memory 訪問速度與寄存器相似； 實現(xiàn)線程間通信的延遲最小； 保存公用的計數(shù)器,或block的公用結(jié)果； 聲

12、明時使用關(guān)鍵字 _ shared_; constant memory 只讀地址空間，位于顯存，用于存儲需頻繁訪問的只讀參數(shù)； 由于其在設(shè)備上有片上緩存，比全局存儲器讀取效率高很多； 使用_ constant_ 關(guān)鍵字; 定義在所有函數(shù)之外；,globalmemory 存在于顯存中，也稱為線性內(nèi)存; cudaMalloc（）函數(shù)分配; cudaFree（）函數(shù)釋放; cudaMemcpy（）進行主機端與設(shè)備端數(shù)據(jù)傳輸; 對于二維和三維數(shù)組： cudaMallocPitch（）和cudaMalloc3D（）分配線性存儲空間； cudaMemcpy2D（），cudaMemcpy3D（）向設(shè)備存儲器進

13、行拷貝；,對象數(shù)組指針 用于指向 global memory的存儲區(qū)域, 由Host端分配，例： float* d_A， d_B ，d_C; int N = 50000; size_t size = N * sizeof(float); cudaMalloc(void*) cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice) ； cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice) ；,texture memory：只讀； 不是一塊專門的存儲器，涉及顯存、紋理緩存、紋理拾取等紋理流水線操作； 在ke

14、rnel中訪問紋理存儲器的操作, 稱為紋理拾?。╰exture fetching） 紋理拾取使用的坐標與數(shù)據(jù)在顯存中的位置可以不同，通過紋理參考約定二者映射方式； 將顯存中數(shù)據(jù)與紋理參考關(guān)聯(lián)的操作，稱為紋理綁定（texture binding）；,3,變量和函數(shù),3.1 變量類型申明 3.2 函數(shù)類型申明 3.3 內(nèi)置變量 3.4 內(nèi)置向量 3.5 數(shù)學函數(shù) 3.6 與Opengl互操作,3.1,3.1,3.2 函數(shù),gridDim（有多少Block） 變量的類型為 dim3，包含網(wǎng)格的維度。 blockDim (有多少Thread） 變量的類型為 dim3，包含塊的維度。 blockIdx(

15、第幾個Block) 變量的類型為 uint3，包含網(wǎng)格內(nèi)的塊索引 threadIdx(第幾個Thread) 變量的類型為 uint3，包含塊內(nèi)的線程索引。 備注：dim3 是整型的向量類型，未指定的任何組件都將初始化為1,如 dim3 dimBlock(4, 2, 2); uint3 是無符號整型的向量類型；,3.3 內(nèi)置變量,Type name : char1、uchar1、char2、uchar2、char3、uchar3、char4、 uchar4、short1、ushort1、short2、ushort2、short3、ushort3、short4、ushort4、int1、uint1

16、、int2、uint2、int3、uint3、int4、uint4、long1、ulong1、long2、ulong2、long3、ulong3、long4、ulong4、float1、float2、float3、float4、double2 均為結(jié)構(gòu)體，第 1、2、3、4 個組件分別可通過字段 x、y、z 和 w 訪問； 均附帶形式為 make_ 的構(gòu)造函數(shù) 例如： int2 v=make_int2(6, 9);,3.4 內(nèi)置向量,CUDA支持標量運算的數(shù)學函數(shù),包括： pow、sqrt、cbrt、hypot、exp、exp2、espm1、log、log2、log10、log1p、sin、c

17、os、tan、asin、acos、atan、atan2、sinh、cosh、tanh、asinh、acosh、atanh、ceil、floor、trunc、round等; 當函數(shù)有前綴“_”時，運算速度較快。,3.5 數(shù)學函數(shù),3.6 與Opengl的互操作 將一個Opengl緩沖對象注冊到 CUDA，通過 cudaGLRegisterBufferObject( ) ； 例如： GLuint bufferObj; cuGLRegisterBufferObject(bufferObj);,注冊完成后，內(nèi)核即可使用 cuGLMapBufferObject( ); /映射緩沖對象 返回設(shè)備存儲器地址

18、，讀取或?qū)懭刖彌_對象； 即將CUDA內(nèi)存的指針指向OpenGL的緩沖區(qū)； 例如： CUdeviceptr devPtr; int size; cuGLMapBufferObject(,用完之后 cuGLUnmapBufferObject( )； /解除映射 cuGLUnregisterBufferObject( )； /取消注冊；,編程模型,變量和函數(shù),案例,介紹,4,案例,4.1. 向量加法； 4.2. 矩陣乘積; 4.3. 矩陣分塊乘積;,編程模型,變量和函數(shù),案例,介紹,在設(shè)計并行程序時，需要找到被分解問題關(guān)鍵參數(shù)和線程下標的對應(yīng)關(guān)系，并建立映射機制，以使GPU的并行計算能力得到充分發(fā)揮

19、； 此外，考慮GPU的軟硬件資源，CUDA編程模型有一些參數(shù)上的限制，主要包括對所分配線程和線程塊數(shù)量的規(guī)定； 如GeForces系列GPU在每個線程塊支持的最大線程數(shù)是512，而每個網(wǎng)格所具有最大線程塊的數(shù)量是65535,4.1.向量和矩陣加法,編程模型,變量和函數(shù),案例,介紹,/方法一：_global_ void VecAdd(float* A, float* B, float* C) int i = threadIdx.x;Ci = Ai + Bi; int main() / Kernel invocationVecAdd(A, B, C); 參見例：matrix_src.doc,編程模

20、型,變量和函數(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 N) Ci = Ai

21、 + Bi; ,編程模型,變量和函數(shù),案例,介紹,擴展： 矩陣加法 _global_ void matAdd(float *Ad, float *Bd, float *Cd) int i = threadIdx.x; int j = threadIdx.y; Cdij = Adij + Bdij; int main( ) / Kernel invocation dim3 dimBlock(N, N); matAdd(Ad, Bd, Cd); . ,編程模型,變量和函數(shù),案例,介紹,說明內(nèi)存和線程管理的基本特性 本地存儲器、寄存器的用法 線程ID的用法 主機和設(shè)備之間數(shù)據(jù)傳輸?shù)腁PI 為了方便，

22、以方形矩陣說明,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 Width; +i) for (int j = 0; j Width; +j) double sum = 0; for (int k = 0; k Width;

23、 +k) double a = Mi * width + k; double b = Nk * width + j; sum += a * b; Pi*Width + j = sum; ,編程模型,變量和函數(shù),案例,介紹,向GPU傳輸矩陣數(shù)據(jù) void MatrixMulOnDevice(float* M, float* N, float* P, int Width) int size = Width * Width * sizeof(float); float* Md, Nd, Pd; /設(shè)置調(diào)用內(nèi)核函數(shù)時的線程數(shù)目 dim3 dimBlock(Width, Width); dim3 dimGrid(1, 1); /在設(shè)備存儲器上給M和N矩陣分配空間，并將數(shù)據(jù)復(fù)制到設(shè)備存儲器中 cudaMalloc(,GPU上的矩陣乘法,編程模型,變量和函數(shù),案例,介紹,計算結(jié)果