姓名：王一帆

學(xué)號(hào)：20011110225

轉(zhuǎn)載自https://bbs.csdn.net/topics/390798229

【嵌牛導(dǎo)讀】

對(duì)CUDA入門知識(shí)進(jìn)行簡單介紹

【嵌牛正文】

CUDA是什么

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

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

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

64位Ubuntu12.04安裝CUDA5.5

具體步驟請(qǐng)點(diǎn)擊此處http://bookc.github.io/2014/05/08/my-summery-the-book-cuda-by-example-an-introduction-to-general-purpose-gpu-programming/。

對(duì)CUDA?C的個(gè)人懵懂感覺:

????????如果粗暴的認(rèn)為C語言工作的對(duì)象是CPU和內(nèi)存條（接下來,稱為主機(jī)內(nèi)存），那么CUDA?C工作的的對(duì)象就是GPU及GPU上的內(nèi)存（接下來,稱為設(shè)備內(nèi)存），且充分利用了GPU多核的優(yōu)勢(shì)及降低了并行編程的難度。一般通過C語言把數(shù)據(jù)從外界讀入，再分配數(shù)據(jù)，給CUDA?C，以便在GPU上計(jì)算，然后再把計(jì)算結(jié)果返回給C語言，以便進(jìn)一步工作，如進(jìn)一步處理及顯示，或重復(fù)此過程。

主要概念與名稱

主機(jī)

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

設(shè)備

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

線程(Thread)

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

線程塊(Block)

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

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

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

線程格(Grid)

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

線程束

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

核函數(shù)（Kernel）

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

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

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

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

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

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

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

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

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

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

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

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

????????4.1.?threadIdx，顧名思義獲取線程thread的ID索引；如果線程是一維的那么就取threadIdx.x，二維的還可以多取到一個(gè)值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.?對(duì)于一維的block，線程的threadID=threadIdx.x。

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

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

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

函數(shù)修飾符

????????1.?__global__，表明被修飾的函數(shù)在設(shè)備上執(zhí)行，但在主機(jī)上調(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.?注意事項(xiàng)：

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

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

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

????????3.4.?不可以在主機(jī)代碼中使用cudaMalloc()分配的指針進(jìn)行主機(jī)內(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ù)可以在主機(jī)內(nèi)存和GPU內(nèi)存之間互相拷貝數(shù)據(jù)。

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

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

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

cudaFree()

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

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

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

GPU內(nèi)存分類

全局內(nèi)存

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

共享內(nèi)存

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

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

????????3.?目的：對(duì)于GPU上啟動(dòng)的每個(gè)線程塊，CUDA?C編譯器都將創(chuàng)建該共享變量的一個(gè)副本。線程塊中的每個(gè)線程都共享這塊內(nèi)存，但線程卻無法看到也不能修改其他線程塊的變量副本。這樣使得一個(gè)線程塊中的多個(gè)線程能夠在計(jì)算上通信和協(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.?特點(diǎn)：常量內(nèi)存用于保存在核函數(shù)執(zhí)行期間不會(huì)發(fā)生變化的數(shù)據(jù)。變量的訪問限制為只讀。NVIDIA硬件提供了64KB的常量內(nèi)存。不再需要cudaMalloc()或者cudaFree(),而是在編譯時(shí)，靜態(tài)地分配空間。

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

????????6.?性能提升的原因：

????????6.1.?對(duì)常量內(nèi)存的單次讀操作可以廣播到其他的“鄰近”線程。這將節(jié)約15次讀取操作。（為什么是15，因?yàn)椤班徑敝赴雮€(gè)線程束，一個(gè)線程束包含32個(gè)線程的集合。）

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

紋理內(nèi)存

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

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

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

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

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

????????4.1.1.?用texture<類型>類型聲明，如texture<float>?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<float，2>?texIn。

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

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

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

固定內(nèi)存

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

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

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

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

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

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

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

原子性

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

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

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

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

使用事件來測量性能

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

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

流

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

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

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

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

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

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

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

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

????????9.?針對(duì)多個(gè)流：

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

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

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

技巧

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

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