CUDA是什么

CUDA，Compute?Unified?Device?Architecture的簡稱，是由NVIDIA公司創(chuàng)立的基于他們公司生產(chǎn)的圖形處理器GPUs（Graphics?Processing?Units,可以通俗的理解為顯卡）的一個并行計算平臺和編程模型。

通過CUDA，GPUs可以很方便地被用來進(jìn)行通用計算（有點像在CPU中進(jìn)行的數(shù)值計算等等）。在沒有CUDA之前，GPUs一般只用來進(jìn)行圖形渲染（如通過OpenGL，DirectX）。

開發(fā)人員可以通過調(diào)用CUDA的API，來進(jìn)行并行編程，達(dá)到高性能計算目的。NVIDIA公司為了吸引更多的開發(fā)人員，對CUDA進(jìn)行了編程語言擴展，如CUDA?C/C++,CUDA?Fortran語言。注意CUDA?C/C++可以看作一個新的編程語言，因為NVIDIA配置了相應(yīng)的編譯器nvcc,CUDA?Fortran一樣。更多信息可以參考文獻(xiàn)。

主要概念與名稱

主機

將CPU及系統(tǒng)的內(nèi)存（內(nèi)存條）稱為主機。

設(shè)備

將GPU及GPU本身的顯示內(nèi)存稱為設(shè)備。

線程(Thread)

一般通過GPU的一個核進(jìn)行處理。（可以表示成一維，二維，三維，具體下面再細(xì)說）。

線程塊(Block)

1.?由多個線程組成（可以表示成一維，二維，三維，具體下面再細(xì)說）。

2.?各block是并行執(zhí)行的，block間無法通信，也沒有執(zhí)行順序。

3.?注意線程塊的數(shù)量限制為不超過65535（硬件限制）。

線程格(Grid)

由多個線程塊組成（可以表示成一維，二維，三維，具體下面再細(xì)說）。

在CUDA架構(gòu)中，線程束是指一個包含32個線程的集合，這個線程集合被“編織在一起”并且“步調(diào)一致”的形式執(zhí)行。在程序中的每一行，線程束中的每個線程都將在不同數(shù)據(jù)上執(zhí)行相同的命令。

核函數(shù)（Kernel）

1.?在GPU上執(zhí)行的函數(shù)通常稱為核函數(shù)。

2.?一般通過標(biāo)識符__global__修飾，調(diào)用通過<<<參數(shù)1,參數(shù)2>>>，用于說明內(nèi)核函數(shù)中的線程數(shù)量，以及線程是如何組織的。

3.?以線程格（Grid）的形式組織，每個線程格由若干個線程塊（block）組成，而每個線程塊又由若干個線程（thread）組成。

4.?是以block為單位執(zhí)行的。

5.?叧能在主機端代碼中調(diào)用。

6.?調(diào)用時必須聲明內(nèi)核函數(shù)的執(zhí)行參數(shù)。

7.?在編程時，必須先為kernel函數(shù)中用到的數(shù)組或變量分配好足夠的空間，再調(diào)用kernel函數(shù)，否則在GPU計算時會發(fā)生錯誤，例如越界或報錯，甚至導(dǎo)致藍(lán)屏和死機。

dim3結(jié)構(gòu)類型

1.?dim3是基亍uint3定義的矢量類型，相當(dāng)亍由3個unsigned?int型組成的結(jié)構(gòu)體。uint3類型有三個數(shù)據(jù)成員unsigned?int?x;?unsigned?int?y;?unsigned?int?z;

2.?可使用亍一維、二維或三維的索引來標(biāo)識線程，構(gòu)成一維、二維或三維線程塊。

3.?dim3結(jié)構(gòu)類型變量用在核函數(shù)調(diào)用的<<<,>>>中。

4.?相關(guān)的幾個內(nèi)置變量

4.1.?threadIdx，顧名思義獲取線程thread的ID索引；如果線程是一維的那么就取threadIdx.x，二維的還可以多取到一個值threadIdx.y，以此類推到三維threadIdx.z。

4.2.?blockIdx，線程塊的ID索引；同樣有blockIdx.x，blockIdx.y，blockIdx.z。

4.3.?blockDim，線程塊的維度，同樣有blockDim.x，blockDim.y，blockDim.z。

4.4.?gridDim，線程格的維度，同樣有g(shù)ridDim.x，gridDim.y，gridDim.z。

5.?對于一維的block，線程的threadID=threadIdx.x。

6.?對于大小為（blockDim.x,?blockDim.y）的?二維?block，線程的threadID=threadIdx.x+threadIdx.y*blockDim.x。

7.?對于大小為（blockDim.x,?blockDim.y,?blockDim.z）的?三維?block，線程的threadID=threadIdx.x+threadIdx.y*blockDim.x+threadIdx.z*blockDim.x*blockDim.y。

8.?對于計算線程索引偏移增量為已啟動線程的總數(shù)。如stride?=?blockDim.x?*?gridDim.x;?threadId?+=?stride。

函數(shù)修飾符

1.?__global__，表明被修飾的函數(shù)在設(shè)備上執(zhí)行，但在主機上調(diào)用。

2.?__device__，表明被修飾的函數(shù)在設(shè)備上執(zhí)行，但只能在其他__device__函數(shù)或者_(dá)_global__函數(shù)中調(diào)用。

常用的GPU內(nèi)存函數(shù)

cudaMalloc()

1.?函數(shù)原型：?cudaError_t?cudaMalloc?(void?**devPtr,?size_t?size)。

2.?函數(shù)用處：與C語言中的malloc函數(shù)一樣，只是此函數(shù)在GPU的內(nèi)存你分配內(nèi)存。

3.?注意事項：

3.1.?可以將cudaMalloc()分配的指針傳遞給在設(shè)備上執(zhí)行的函數(shù)；

3.2.?可以在設(shè)備代碼中使用cudaMalloc()分配的指針進(jìn)行設(shè)備內(nèi)存讀寫操作；

3.3.?可以將cudaMalloc()分配的指針傳遞給在主機上執(zhí)行的函數(shù)；

3.4.?不可以在主機代碼中使用cudaMalloc()分配的指針進(jìn)行主機內(nèi)存讀寫操作（即不能進(jìn)行解引用）。

cudaMemcpy()

1.?函數(shù)原型：cudaError_t?cudaMemcpy?(void?*dst,?const?void?*src,?size_t?count,?cudaMemcpyKind?kind)。

2.?函數(shù)作用：與c語言中的memcpy函數(shù)一樣，只是此函數(shù)可以在主機內(nèi)存和GPU內(nèi)存之間互相拷貝數(shù)據(jù)。

3.?函數(shù)參數(shù)：cudaMemcpyKind?kind表示數(shù)據(jù)拷貝方向，如果kind賦值為cudaMemcpyDeviceToHost表示數(shù)據(jù)從設(shè)備內(nèi)存拷貝到主機內(nèi)存。

4.?與C中的memcpy()一樣，以同步方式執(zhí)行，即當(dāng)函數(shù)返回時，復(fù)制操作就已經(jīng)完成了，并且在輸出緩沖區(qū)中包含了復(fù)制進(jìn)去的內(nèi)容。

5.?相應(yīng)的有個異步方式執(zhí)行的函數(shù)cudaMemcpyAsync()，這個函數(shù)詳解請看下面的流一節(jié)有關(guān)內(nèi)容。

cudaFree()

1.?函數(shù)原型：cudaError_t?cudaFree?(?void*?devPtr?)。

2.?函數(shù)作用：與c語言中的free()函數(shù)一樣，只是此函數(shù)釋放的是cudaMalloc()分配的內(nèi)存。

下面實例用于解釋上面三個函數(shù)

GPU內(nèi)存分類

全局內(nèi)存

通俗意義上的設(shè)備內(nèi)存。

共享內(nèi)存

1.?位置：設(shè)備內(nèi)存。

2.?形式：關(guān)鍵字__shared__添加到變量聲明中。如__shared__?float?cache[10]。

3.?目的：對于GPU上啟動的每個線程塊，CUDA?C編譯器都將創(chuàng)建該共享變量的一個副本。線程塊中的每個線程都共享這塊內(nèi)存，但線程卻無法看到也不能修改其他線程塊的變量副本。這樣使得一個線程塊中的多個線程能夠在計算上通信和協(xié)作。

常量內(nèi)存

1.?位置：設(shè)備內(nèi)存

2.?形式：關(guān)鍵字__constant__添加到變量聲明中。如__constant__?float?s[10];。

3.?目的：為了提升性能。常量內(nèi)存采取了不同于標(biāo)準(zhǔn)全局內(nèi)存的處理方式。在某些情況下，用常量內(nèi)存替換全局內(nèi)存能有效地減少內(nèi)存帶寬。

4.?特點：常量內(nèi)存用于保存在核函數(shù)執(zhí)行期間不會發(fā)生變化的數(shù)據(jù)。變量的訪問限制為只讀。NVIDIA硬件提供了64KB的常量內(nèi)存。不再需要cudaMalloc()或者cudaFree(),而是在編譯時，靜態(tài)地分配空間。

5.?要求：當(dāng)我們需要拷貝數(shù)據(jù)到常量內(nèi)存中應(yīng)該使用cudaMemcpyToSymbol()，而cudaMemcpy()會復(fù)制到全局內(nèi)存。

6.?性能提升的原因：

6.1.?對常量內(nèi)存的單次讀操作可以廣播到其他的“鄰近”線程。這將節(jié)約15次讀取操作。（為什么是15，因為“鄰近”指半個線程束，一個線程束包含32個線程的集合。）

6.2.?常量內(nèi)存的數(shù)據(jù)將緩存起來，因此對相同地址的連續(xù)讀操作將不會產(chǎn)生額外的內(nèi)存通信量。

紋理內(nèi)存

1.?位置：設(shè)備內(nèi)存

2.?目的：能夠減少對內(nèi)存的請求并提供高效的內(nèi)存帶寬。是專門為那些在內(nèi)存訪問模式中存在大量空間局部性的圖形應(yīng)用程序設(shè)計，意味著一個線程讀取的位置可能與鄰近線程讀取的位置“非常接近”。如下圖：

3. 紋理變量（引用）必須聲明為文件作用域內(nèi)的全局變量。

4.?形式：分為一維紋理內(nèi)存?和?二維紋理內(nèi)存。

4.1.?一維紋理內(nèi)存

4.1.1.?用texture<類型>類型聲明，如texture?texIn。

4.1.2.?通過cudaBindTexture()綁定到紋理內(nèi)存中。

4.1.3.?通過tex1Dfetch()來讀取紋理內(nèi)存中的數(shù)據(jù)。

4.1.4.?通過cudaUnbindTexture()取消綁定紋理內(nèi)存。

4.2.?二維紋理內(nèi)存

4.2.1.?用texture<類型,數(shù)字>類型聲明，如texture?texIn。

4.2.2.?通過cudaBindTexture2D()綁定到紋理內(nèi)存中。

4.2.3.?通過tex2D()來讀取紋理內(nèi)存中的數(shù)據(jù)。

4.2.4.?通過cudaUnbindTexture()取消綁定紋理內(nèi)存。

固定內(nèi)存

1.?位置：主機內(nèi)存。

2.?概念：也稱為頁鎖定內(nèi)存或者不可分頁內(nèi)存，操作系統(tǒng)將不會對這塊內(nèi)存分頁并交換到磁盤上，從而確保了該內(nèi)存始終駐留在物理內(nèi)存中。因此操作系統(tǒng)能夠安全地使某個應(yīng)用程序訪問該內(nèi)存的物理地址，因為這塊內(nèi)存將不會破壞或者重新定位。

3.?目的：提高訪問速度。由于GPU知道主機內(nèi)存的物理地址，因此可以通過“直接內(nèi)存訪問DMA（Direct?Memory?Access)技術(shù)來在GPU和主機之間復(fù)制數(shù)據(jù)。由于DMA在執(zhí)行復(fù)制時無需CPU介入。因此DMA復(fù)制過程中使用固定內(nèi)存是非常重要的。

4.?缺點：使用固定內(nèi)存，將失去虛擬內(nèi)存的所有功能；系統(tǒng)將更快的耗盡內(nèi)存。

5.?建議：對cudaMemcpy()函數(shù)調(diào)用中的源內(nèi)存或者目標(biāo)內(nèi)存，才使用固定內(nèi)存，并且在不再需要使用它們時立即釋放。

6.?形式：通過cudaHostAlloc()函數(shù)來分配；通過cudaFreeHost()釋放。

7.?只能以異步方式對固定內(nèi)存進(jìn)行復(fù)制操作。

原子性

1.?概念：如果操作的執(zhí)行過程不能分解為更小的部分，我們將滿足這種條件限制的操作稱為原子操作。

2.?形式：函數(shù)調(diào)用，如atomicAdd（addr,y)將生成一個原子的操作序列，這個操作序列包括讀取地址addr處的值，將y增加到這個值，以及將結(jié)果保存回地址addr。

常用線程操作函數(shù)

1.?同步方法__syncthreads()，這個函數(shù)的調(diào)用，將確保線程塊中的每個線程都執(zhí)行完__syscthreads()前面的語句后，才會執(zhí)行下一條語句。

使用事件來測量性能

1.?用途：為了測量GPU在某個任務(wù)上花費的時間。CUDA中的事件本質(zhì)上是一個GPU時間戳。由于事件是直接在GPU上實現(xiàn)的。因此不適用于對同時包含設(shè)備代碼和主機代碼的混合代碼設(shè)計。

2.?形式：首先創(chuàng)建一個事件，然后記錄事件，再計算兩個事件之差，最后銷毀事件。如：

流

1.?扯一扯：并發(fā)重點在于一個極短時間段內(nèi)運行多個不同的任務(wù)；并行重點在于同時運行一個任務(wù)。

2.?任務(wù)并行性：是指并行執(zhí)行兩個或多個不同的任務(wù)，而不是在大量數(shù)據(jù)上執(zhí)行同一個任務(wù)。

3.?概念：CUDA流表示一個GPU操作隊列，并且該隊列中的操作將以指定的順序執(zhí)行。我們可以在流中添加一些操作，如核函數(shù)啟動，內(nèi)存復(fù)制以及事件的啟動和結(jié)束等。這些操作的添加到流的順序也是它們的執(zhí)行順序?？梢詫⒚總€流視為GPU上的一個任務(wù)，并且這些任務(wù)可以并行執(zhí)行。

4.?硬件前提：必須是支持設(shè)備重疊功能的GPU。支持設(shè)備重疊功能，即在執(zhí)行一個核函數(shù)的同時，還能在設(shè)備與主機之間執(zhí)行復(fù)制操作。

5.?聲明與創(chuàng)建：聲明cudaStream_t?stream;，創(chuàng)建cudaSteamCreate(&stream);。

6.?cudaMemcpyAsync()：前面在cudaMemcpy()中提到過，這是一個以異步方式執(zhí)行的函數(shù)。在調(diào)用cudaMemcpyAsync()時，只是放置一個請求，表示在流中執(zhí)行一次內(nèi)存復(fù)制操作，這個流是通過參數(shù)stream來指定的。當(dāng)函數(shù)返回時，我們無法確保復(fù)制操作是否已經(jīng)啟動，更無法保證它是否已經(jīng)結(jié)束。我們能夠得到的保證是，復(fù)制操作肯定會當(dāng)下一個被放入流中的操作之前執(zhí)行。傳遞給此函數(shù)的主機內(nèi)存指針必須是通過cudaHostAlloc()分配好的內(nèi)存。（流中要求固定內(nèi)存）

7.?流同步：通過cudaStreamSynchronize()來協(xié)調(diào)。

8.?流銷毀：在退出應(yīng)用程序之前，需要銷毀對GPU操作進(jìn)行排隊的流，調(diào)用cudaStreamDestroy()。

9.?針對多個流：

9.1.?記得對流進(jìn)行同步操作。

9.2.?將操作放入流的隊列時，應(yīng)采用寬度優(yōu)先方式，而非深度優(yōu)先的方式，換句話說，不是首先添加第0個流的所有操作，再依次添加后面的第1，2,…個流。而是交替進(jìn)行添加，比如將a的復(fù)制操作添加到第0個流中，接著把a的復(fù)制操作添加到第1個流中，再繼續(xù)其他的類似交替添加的行為。

9.3.?要牢牢記住操作放入流中的隊列中的順序影響到CUDA驅(qū)動程序調(diào)度這些操作和流以及執(zhí)行的方式。

技巧

1.?當(dāng)線程塊的數(shù)量為GPU中處理數(shù)量的2倍時，將達(dá)到最優(yōu)性能。

2.?核函數(shù)執(zhí)行的第一個計算就是計算輸入數(shù)據(jù)的偏移。每個線程的起始偏移都是0到線程數(shù)量減1之間的某個值。然后，對偏移的增量為已啟動線程的總數(shù)。