姓名：王一帆

學(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

具體步驟請點(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)勢及降低了并行編程的難度。一般通過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ù)詳解請看下面的流一節(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)存的請求并提供高效的內(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è)請求，表示在流中執(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ù)伺糠。