CUDA介紹與案例

1、CUDA 介紹與案例 介紹介紹編程模型編程模型變量和函數(shù)變量和函數(shù)案例案例1介紹編程模型變量和函數(shù)案例介紹p 應(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)的庫文件；將GPU 視作數(shù)據(jù)并行計(jì)算設(shè)備,是一種新的處理和管理GPU 計(jì)算的硬件和軟件架構(gòu)；編程模型變量和函數(shù)案例介紹GPU到CUDAGPU多多“核核”：SM真正意義的變革真

2、正意義的變革GPGPU通用計(jì)算通用計(jì)算重要突破重要突破GPU通用計(jì)通用計(jì)算開始出現(xiàn)算開始出現(xiàn)世界上第一個(gè)世界上第一個(gè)GPUGeForce 8800GeForce 6800GeForce 3GeForce 256編程模型變量和函數(shù)案例介紹pCPU 和 GPU 的每秒浮點(diǎn)運(yùn)算次數(shù)和存儲(chǔ)器帶寬比較具有強(qiáng)大浮點(diǎn)具有強(qiáng)大浮點(diǎn)運(yùn)算能力運(yùn)算能力n GPU采用大量的執(zhí)行單元，這些執(zhí)行單元可以輕松的加載并行處理線程，而不像CPU那樣的單線程處理；n 主機(jī)和設(shè)備均維護(hù)自己的主機(jī)和設(shè)備均維護(hù)自己的 DRAM，分別稱為主機(jī)存儲(chǔ)器和設(shè)備存，分別稱為主機(jī)存儲(chǔ)器和設(shè)備存儲(chǔ)器儲(chǔ)器 CPU 和 GPU 之間浮點(diǎn)功能的差異，原因

3、在于 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語言, 可自由地調(diào)用GPU的并行 處理架構(gòu); 同時(shí)適用于圖形和通用并行計(jì)算應(yīng)用程序， 成為圖形處理器的主要發(fā)展趨勢。編程模型變量和函數(shù)案例介紹 依次安裝 Cuda Driver Cuda Toolkit Cuda SDK 在安裝目錄中包括： bin工具程序及動(dòng)態(tài)鏈接庫 doc文件 includ

4、e 頭文件 lib 鏈接庫 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)用程序編程接口（API）及其運(yùn)行時(shí)兩個(gè)較高級(jí)別的通用數(shù)學(xué)庫，即 CUFFT 和 CUBLAS CUDA安裝后編程模型變量和函數(shù)案例介紹 CUDA程序編譯p CUDA的源文件被nvcc編譯 NVCC編譯器會(huì)分離源碼中設(shè)備代碼和主機(jī)代碼，主機(jī)代碼交由一般的C/C+

5、編譯器（gcc等）編譯，設(shè)備代碼由NVCC編譯； 內(nèi)核必須使用nvcc編譯成二進(jìn)制代碼才能在設(shè)備端執(zhí)行。v_2011編程模型變量和函數(shù)案例介紹標(biāo)注字體使用 行距背景圖片出處聲明英文 Century Gothic中文 微軟雅黑正文 互聯(lián)網(wǎng)是一個(gè)開放共享的平臺(tái)OfficePLUS 部分設(shè)計(jì)靈感與元素來源于網(wǎng)絡(luò)如有建議請(qǐng)聯(lián)系 OfficePLUS內(nèi)核函數(shù)編程模型變量和函數(shù)案例介紹編程模型變量和函數(shù)案例介紹n 串行代碼在主機(jī)上串行代碼在主機(jī)上執(zhí)行，而并行代碼執(zhí)行，而并行代碼在設(shè)備上執(zhí)行在設(shè)備上執(zhí)行n 當(dāng)調(diào)用當(dāng)調(diào)用kernelkernel的時(shí)的時(shí)候，則候，則CUDACUDA的每個(gè)的每個(gè)線程并行地執(zhí)行一線

6、程并行地執(zhí)行一段指令，這與通常段指令，這與通常C C函數(shù)只執(zhí)行一次不函數(shù)只執(zhí)行一次不一樣。一樣。執(zhí)行模型執(zhí)行模型編程模型變量和函數(shù)案例介紹說明：n在CUDA架構(gòu)下，一個(gè)程序分為兩部分，即host 端和 device端； Host 端 是在 CPU上執(zhí)行的部分； Device端 是在顯卡芯片上執(zhí)行的部分，該端程序又稱為內(nèi)核 函數(shù)（ kernel function） n 通常Host端程序復(fù)制內(nèi)核函數(shù)到顯卡內(nèi)存中，再由顯卡芯片執(zhí)行device端程序，完成后再由host端程序?qū)⒔Y(jié)果從顯卡內(nèi)存中取回；編程模型變量和函數(shù)案例介紹n CUDA 允許程序員定義內(nèi)核函數(shù)，實(shí)現(xiàn)配置；n 塊組織為一個(gè)一維或二維線

7、程塊網(wǎng)格，維度由 語法的第一個(gè)參數(shù)指定；n 執(zhí)行內(nèi)核的每個(gè)線程都會(huì)被分配一個(gè)獨(dú)特的線程 ID，可通過內(nèi)置的 threadIdx 變量在內(nèi)核中訪問編程模型22.1 線程模型（體系結(jié)構(gòu)）2.2 存儲(chǔ)器模型（體系結(jié)構(gòu)）編程模型變量和函數(shù)案例介紹2.1 線程模型編程模型變量和函數(shù)案例介紹并行線程結(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 b

8、lock共享全局內(nèi)存Kernel: 在GPU上執(zhí)行的核心程序One kernel One grid編程模型變量和函數(shù)案例介紹說明： Grid 是一組Block, 以1維、2維或3維組織，共享全局內(nèi)存； Grid之間通過global memory交換數(shù)據(jù) Block 是互相協(xié)作的線程組，通過global memory共享數(shù)據(jù)，允許彼此同步；以1維、2維或3維組織； 同一block內(nèi)的thread可以通過shared memory和同步實(shí)現(xiàn)通信;編程模型變量和函數(shù)案例介紹 一個(gè)內(nèi)核可能由多個(gè)大小相同的線程塊執(zhí)行，因而線程總數(shù)線程總數(shù)應(yīng)等于每個(gè)塊的線程數(shù)乘以塊的數(shù); 線程索引線程索引(Index)(

9、Index)及及ID(ID(IDentity) )，索引為(x,y)的線程ID為(x+yDx);對(duì)于大小為(Dx,Dy,Dz)的三維塊，索引為(x,y,z)的線程ID為(x+yDx+zDxDy)； 示意圖編程模型變量和函數(shù)案例介紹 線程塊索引及線程塊索引及ID ,ID ,情況類似： 對(duì)于一維塊來說，兩者是相同的； 對(duì)于大小為 (Dx，Dy) 的二維塊來說，索引為 (x,y) 的線程塊的ID 是 （x + yDx）； 對(duì)于大小為 （Dx，Dy， Dz） 的三維塊來說，索引 為（x, y, z） 的線程的ID 是 （x + yDx +zDxDy）; /+圖示說明編程模型變量和函數(shù)案例介紹 例: 向

10、量的和 _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ù)組的下標(biāo)編程模型變量和函數(shù)案例介紹或int main( ) int threadsPerBlock = 256; int blocksPerGrid = (N + threadsPerBlock - 1) /threadsPerBlock; Vec

11、Add(d_A, d_B, d_C, N); 編程模型變量和函數(shù)案例介紹 _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;/參見例 ：vectorAdd!？哪個(gè)線程對(duì)哪個(gè)數(shù)組下標(biāo)求和； 即第幾個(gè)Block（blockIdx）的第幾個(gè)線程，對(duì)相應(yīng)的數(shù)組下標(biāo)求和；編程模型變量和函數(shù)案例介紹2.2 編程模型變量和函數(shù)案例介紹編程模型變量和函數(shù)案例介紹說明：p

12、所有線程都可訪問全局存儲(chǔ)器；p 所有的線程都可以訪問只讀存儲(chǔ)區(qū)域：“常量存儲(chǔ)區(qū)”和“紋理存儲(chǔ)區(qū)”； 紋理存儲(chǔ)器提供不同的尋址模式，為某些特定的數(shù)據(jù)格式提供數(shù)據(jù)過濾的能力；p Block中所有thread都在同一個(gè)處理核心上運(yùn)行。因此每個(gè)block的thread數(shù)目受到處理核心所擁有存儲(chǔ)器資源的限制，在NVIDIA Tesla架構(gòu)中，一個(gè)線程block最多可包含512個(gè)線程;編程模型變量和函數(shù)案例介紹p 每個(gè)線程塊都有一個(gè)共享存儲(chǔ)器共享存儲(chǔ)器(Shared MemoryShared Memory)，該存儲(chǔ)器對(duì)于塊內(nèi)的所有線程塊內(nèi)的所有線程都是可見可見的，并且與塊具有相同的生命周期。 每個(gè)線程都有

13、一個(gè)私有的本地存儲(chǔ)器（本地存儲(chǔ)器（LocalMemoryLocalMemory）。p 在同一個(gè)blockblock中的中的threadthread通過共享存儲(chǔ)器共享數(shù)據(jù)，相互協(xié)作，實(shí)現(xiàn)同步同步： _syncthreads( ); 編程模型變量和函數(shù)案例介紹p registerregister 訪問延遲極低； 每個(gè)線程占有的register有限，編程時(shí)不要為其分配過多 私有變量；p local memorylocal memory 寄存器被使用完畢，數(shù)據(jù)將被存儲(chǔ)在局部存儲(chǔ)器中； 此外，存儲(chǔ)大型結(jié)構(gòu)體或者數(shù)組（無法確定大小的數(shù)組），線程的輸入和中間變量；編程模型變量和函數(shù)案例介紹p shared

14、memoryshared memory 訪問速度與寄存器相似； 實(shí)現(xiàn)線程間通信的延遲最??； 保存公用的計(jì)數(shù)器,或block的公用結(jié)果； 聲明時(shí)使用關(guān)鍵字 _ shared_; p constant memoryconstant memory 只讀地址空間，位于顯存，用于存儲(chǔ)需頻繁訪問的只讀參數(shù)； 由于其在設(shè)備上有片上緩存，比全局存儲(chǔ)器讀取效率高很多； 使用_ constant_ 關(guān)鍵字; 定義在所有函數(shù)之外； 編程模型變量和函數(shù)案例介紹p globalmemoryglobalmemory 存在存在于顯存中，也稱為線性內(nèi)存; cudaMalloc（）函數(shù)分配; cudaFree（）函數(shù)釋放; c

15、udaMemcpy（）進(jìn)行主機(jī)端與設(shè)備端數(shù)據(jù)傳輸; 對(duì)于二維和三維數(shù)組： cudaMallocPitch（）和cudaMalloc3D（）分配線性存儲(chǔ)空間； cudaMemcpy2D（），cudaMemcpy3D（）向設(shè)備存儲(chǔ)器進(jìn)行拷貝；編程模型變量和函數(shù)案例介紹編程模型變量和函數(shù)案例介紹 對(duì)象數(shù)組指針對(duì)象數(shù)組指針 用于指向 global memory的存儲(chǔ)區(qū)域, 由Host端分配，例： float* d_A， d_B ，d_C; int N = 50000; size_t size = N * sizeof(float); cudaMalloc(void*)&d_A, size) ;

16、 cudaMalloc(void*)&d_B, size) ; cudaMalloc(void*)&d_C, size) ; cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice) ； cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice) ；編程模型變量和函數(shù)案例介紹p texture memorytexture memory：只讀； 不是一塊專門的存儲(chǔ)器，涉及顯存、紋理緩存、紋理拾取等紋理流水線操作； 在kernel中訪問紋理存儲(chǔ)器的操作, 稱為紋理拾取紋理拾取（texture fe

17、tching） 紋理拾取使用的坐標(biāo)與數(shù)據(jù)在顯存中的位置可以不同，通過紋理參考約定二者映射方式； 將顯存中數(shù)據(jù)與紋理參考關(guān)聯(lián)的操作，稱為紋理綁定（texture binding）；3變量和函數(shù)3.1 變量類型申明變量類型申明3.2 3.2 函數(shù)類型申明函數(shù)類型申明3.3 3.3 內(nèi)置變量內(nèi)置變量3.4 3.4 內(nèi)置向量內(nèi)置向量3.5 3.5 數(shù)學(xué)數(shù)學(xué)函數(shù)函數(shù)3.6 3.6 與與OpenglOpengl互操作互操作編程模型變量和函數(shù)案例介紹3.13.1編程模型變量和函數(shù)案例介紹編程模型變量和函數(shù)案例介紹3.2 函數(shù)函數(shù)編程模型變量和函數(shù)案例介紹n gridDimgridDim（有多少（有多少Blo

18、ckBlock） 變量的類型為 dim3，包含網(wǎng)格的維度。n blockDimblockDim ( (有多少有多少ThreadThread） 變量的類型為 dim3，包含塊的維度。n blockIdxblockIdx( (第幾個(gè)第幾個(gè)Block)Block) 變量的類型為 uint3，包含網(wǎng)格內(nèi)的塊索引 n threadIdxthreadIdx( (第幾個(gè)第幾個(gè)Thread)Thread) 變量的類型為 uint3，包含塊內(nèi)的線程索引。 備注：dim3 是整型的向量類型，未指定的任何組件都將初始化為1,如 dim3 dimBlock(4, 2, 2); uint3 是無符號(hào)整型的向量類型；3.

19、3 3.3 內(nèi)置變量內(nèi)置變量編程模型變量和函數(shù)案例介紹 Type name : char1、uchar1、char2、uchar2、char3、uchar3、char4、 uchar4、short1、ushort1、short2、ushort2、short3、ushort3、short4、ushort4、int1、uint1、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

20、 個(gè)組件分別可通過字段 x、y、z 和 w 訪問； 均附帶形式為 make_ 的構(gòu)造函數(shù) 例如： int2 v=make_int2(6, 9); 3.4 3.4 內(nèi)置向量內(nèi)置向量編程模型變量和函數(shù)案例介紹 CUDA支持標(biāo)量運(yùn)算的數(shù)學(xué)函數(shù),包括：pow、sqrt、cbrt、hypot、exp、exp2、espm1、log、log2、log10、log1p、sin、cos、tan、asin、acos、atan、atan2、sinh、cosh、tanh、asinh、acosh、atanh、ceil、floor、trunc、round等; 當(dāng)函數(shù)有前綴“_”時(shí)，運(yùn)算速度較快。3.5 3.5 數(shù)學(xué)函數(shù)數(shù)

21、學(xué)函數(shù)編程模型變量和函數(shù)案例介紹3.6 3.6 與與OpenglOpengl的互的互操作操作將一個(gè)Opengl緩沖對(duì)象注冊(cè)到 CUDA，通過通過 cudaGLRegisterBufferObjeccudaGLRegisterBufferObject( ) ；例如： GLuint bufferObj; cuGLRegisterBufferObject(bufferObj);編程模型變量和函數(shù)案例介紹 注冊(cè)完成后，內(nèi)核即可使用 cuGLMapBufferObjectcuGLMapBufferObject( ( ); / ); /映射緩沖對(duì)象 返回設(shè)備存儲(chǔ)器地址，讀取或?qū)懭刖彌_對(duì)象； 即將CUDA內(nèi)

22、存的指針指向OpenGL的緩沖區(qū)； 例如： CUdeviceptr devPtr; int size; cuGLMapBufferObject(&devPtr, &size, bufferObj); 用完之后 cuGLUnmapBufferObjectcuGLUnmapBufferObject( ( ) )； /解除映射 cuGLUnregisterBufferObjectcuGLUnregisterBufferObject( ( ) )； /取消注冊(cè)；編程模型變量和函數(shù)案例介紹4案例 4.1. 4.1. 向量加法；向量加法； 4.2. 4.2. 矩陣乘積矩陣乘積; ; 4.3

23、.4.3. 矩陣分塊乘積矩陣分塊乘積; ; 編程模型變量和函數(shù)案例介紹在設(shè)計(jì)并行程序時(shí)，需要找到被分解問題關(guān)鍵參數(shù)和需要找到被分解問題關(guān)鍵參數(shù)和線程下標(biāo)的對(duì)應(yīng)關(guān)系，并建立映射機(jī)制線程下標(biāo)的對(duì)應(yīng)關(guān)系，并建立映射機(jī)制，以使GPU的并行計(jì)算能力得到充分發(fā)揮；此外，考慮GPU的軟硬件資源，CUDA編程模型有一些參數(shù)上的限制，主要包括對(duì)所分配線程和線程塊數(shù)量分配線程和線程塊數(shù)量的規(guī)定； 如GeForces系列GPU在每個(gè)線程塊支持的最大線程數(shù)是512，而每個(gè)網(wǎng)格所具有最大線程塊的數(shù)量是655354.1.向量和矩陣加法編程模型變量和函數(shù)案例介紹/方法一：_global_ void VecAdd(float

24、* A, float* B, float* C) int i = threadIdx.x;Ci = Ai + Bi;int main() / Kernel invocationVecAdd(A, B, C); 參見例：matrix_src.doc編程模型變量和函數(shù)案例介紹方法二：int main( ) int threadsPerBlock = 256; int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;VecAdd(d_A, d_B, d_C, N); 編程模型變量和函數(shù)案例介紹_global_ void VecA

25、dd(const float* A, const float* B, float* C, int N) int i = blockDim.x * blockIdx.x + threadIdx.x; if (i N) Ci = Ai + Bi;編程模型變量和函數(shù)案例介紹擴(kuò)展： 矩陣加法 _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 dimB

26、lock(N, N); matAdd(Ad, Bd, Cd); . 編程模型變量和函數(shù)案例介紹p 說明內(nèi)存和線程管理的基本特性說明內(nèi)存和線程管理的基本特性 本地存儲(chǔ)器、寄存器的用法本地存儲(chǔ)器、寄存器的用法 線程線程IDID的用法的用法 主機(jī)和設(shè)備之間數(shù)據(jù)傳輸?shù)闹鳈C(jī)和設(shè)備之間數(shù)據(jù)傳輸?shù)腁PIAPI 為了方便，以方形矩陣說明為了方便，以方形矩陣說明4.2. 矩陣乘法編程模型變量和函數(shù)案例介紹矩陣大小為矩陣大小為 WIDTH x WIDTHWIDTH x WIDTH在沒有采用分片優(yōu)化算法的情況下：在沒有采用分片優(yōu)化算法的情況下： 一個(gè)線程計(jì)算一個(gè)線程計(jì)算P P矩陣中的一個(gè)矩陣中的一個(gè) 元素元素 需要

27、從全局存儲(chǔ)器載入需要從全局存儲(chǔ)器載入WIDTHWIDTH次次方塊矩陣乘法方塊矩陣乘法編程模型變量和函數(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; +k) double a = Mi * width + k; double b = Nk * width + j; sum += a * b; Pi*

28、Width + j = sum; 編程模型變量和函數(shù)案例介紹向向GPUGPU傳輸矩陣數(shù)據(jù)傳輸矩陣數(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í)的線程數(shù)目設(shè)置調(diào)用內(nèi)核函數(shù)時(shí)的線程數(shù)目 dim3 dimBlock(Width, Width); dim3 dimGrid(1, 1); /在設(shè)備存儲(chǔ)器上給在設(shè)備存儲(chǔ)器上給M M和和N N矩陣分配空間，并將數(shù)據(jù)復(fù)制到設(shè)備存儲(chǔ)器中矩陣分配空間，并將數(shù)據(jù)復(fù)制到設(shè)備存儲(chǔ)器中 cudaMalloc(&Md, size); / Md線性存儲(chǔ)矩陣元素；線性存儲(chǔ)矩陣元素； cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice); cudaMalloc(&Nd, size); / Nd線性存儲(chǔ)矩陣元素；線性存儲(chǔ)矩陣元素； cudaMemcpy(Nd, N, size, CudaMemcpyHostToDevice); /在設(shè)備存儲(chǔ)器上給在設(shè)備存儲(chǔ)器上給P P矩陣分配空間矩陣分配空間 cudaMalloc(&