CUDA是什么
CUDA来庭，Compute?Unified?Device?Architecture的簡稱达吞，是由NVIDIA公司創(chuàng)立的基于他們公司生產的圖形處理器GPUs（Graphics?Processing?Units,可以通俗的理解為顯卡）的一個并行計算平臺和編程模型吟温。
通過CUDA宽菜，GPUs可以很方便地被用來進行通用計算（有點像在CPU中進行的數(shù)值計算等等）葵姥。在沒有CUDA之前紊浩，GPUs一般只用來進行圖形渲染（如通過OpenGL旭等，DirectX）酌呆。
開發(fā)人員可以通過調用CUDA的API，來進行并行編程搔耕，達到高性能計算目的隙袁。NVIDIA公司為了吸引更多的開發(fā)人員，對CUDA進行了編程語言擴展弃榨，如CUDA?C/C++,CUDA?Fortran語言菩收。注意CUDA?C/C++可以看作一個新的編程語言，因為NVIDIA配置了相應的編譯器nvcc,CUDA?Fortran一樣鲸睛。更多信息可以參考文獻娜饵。
主要概念與名稱
主機
將CPU及系統(tǒng)的內存（內存條）稱為主機。
設備
將GPU及GPU本身的顯示內存稱為設備官辈。
線程(Thread)
一般通過GPU的一個核進行處理箱舞。（可以表示成一維，二維拳亿，三維晴股，具體下面再細說）。
線程塊(Block)
1.?由多個線程組成（可以表示成一維风瘦，二維队魏，三維，具體下面再細說）万搔。
2.?各block是并行執(zhí)行的胡桨，block間無法通信，也沒有執(zhí)行順序瞬雹。
3.?注意線程塊的數(shù)量限制為不超過65535（硬件限制）昧谊。
線程格(Grid)
由多個線程塊組成（可以表示成一維，二維酗捌，三維呢诬，具體下面再細說）涌哲。
線程束
在CUDA架構中，線程束是指一個包含32個線程的集合尚镰，這個線程集合被“編織在一起”并且“步調一致”的形式執(zhí)行阀圾。在程序中的每一行，線程束中的每個線程都將在不同數(shù)據(jù)上執(zhí)行相同的命令狗唉。
核函數(shù)（Kernel）
1.?在GPU上執(zhí)行的函數(shù)通常稱為核函數(shù)初烘。
2.?一般通過標識符__global__修飾，調用通過<<<參數(shù)1,參數(shù)2>>>分俯，用于說明內核函數(shù)中的線程數(shù)量肾筐，以及線程是如何組織的。
3.?以線程格（Grid）的形式組織缸剪，每個線程格由若干個線程塊（block）組成吗铐，而每個線程塊又由若干個線程（thread）組成。
4.?是以block為單位執(zhí)行的杏节。
5.?叧能在主機端代碼中調用唬渗。
6.?調用時必須聲明內核函數(shù)的執(zhí)行參數(shù)。
7.?在編程時奋渔，必須先為kernel函數(shù)中用到的數(shù)組或變量分配好足夠的空間谣妻，再調用kernel函數(shù)，否則在GPU計算時會發(fā)生錯誤卒稳，例如越界或報錯蹋半，甚至導致藍屏和死機。
dim3結構類型
1.?dim3是基亍uint3定義的矢量類型充坑，相當亍由3個unsigned?int型組成的結構體减江。uint3類型有三個數(shù)據(jù)成員unsigned?int?x;?unsigned?int?y;?unsigned?int?z;
2.?可使用亍一維、二維或三維的索引來標識線程捻爷，構成一維辈灼、二維或三維線程塊。
3.?dim3結構類型變量用在核函數(shù)調用的<<<,>>>中也榄。
4.?相關的幾個內置變量
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球拦，線程格的維度冰啃，同樣有gridDim.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ù)在設備上執(zhí)行莲镣，但在主機上調用福稳。
2.?__device__，表明被修飾的函數(shù)在設備上執(zhí)行瑞侮，但只能在其他__device__函數(shù)或者__global__函數(shù)中調用的圆。
常用的GPU內存函數(shù)
cudaMalloc()
1.?函數(shù)原型：?cudaError_t?cudaMalloc?(void?**devPtr,?size_t?size)。
2.?函數(shù)用處：與C語言中的malloc函數(shù)一樣半火，只是此函數(shù)在GPU的內存你分配內存越妈。
3.?注意事項：
3.1.?可以將cudaMalloc()分配的指針傳遞給在設備上執(zhí)行的函數(shù)；
3.2.?可以在設備代碼中使用cudaMalloc()分配的指針進行設備內存讀寫操作钮糖；
3.3.?可以將cudaMalloc()分配的指針傳遞給在主機上執(zhí)行的函數(shù)梅掠；
3.4.?不可以在主機代碼中使用cudaMalloc()分配的指針進行主機內存讀寫操作（即不能進行解引用）。
cudaMemcpy()
1.?函數(shù)原型：cudaError_t?cudaMemcpy?(void?*dst,?const?void?*src,?size_t?count,?cudaMemcpyKind?kind)店归。
2.?函數(shù)作用：與c語言中的memcpy函數(shù)一樣阎抒，只是此函數(shù)可以在主機內存和GPU內存之間互相拷貝數(shù)據(jù)。
3.?函數(shù)參數(shù)：cudaMemcpyKind?kind表示數(shù)據(jù)拷貝方向消痛，如果kind賦值為cudaMemcpyDeviceToHost表示數(shù)據(jù)從設備內存拷貝到主機內存挠蛉。
4.?與C中的memcpy()一樣，以同步方式執(zhí)行肄满，即當函數(shù)返回時谴古，復制操作就已經完成了质涛，并且在輸出緩沖區(qū)中包含了復制進去的內容。
5.?相應的有個異步方式執(zhí)行的函數(shù)cudaMemcpyAsync()掰担，這個函數(shù)詳解請看下面的流一節(jié)有關內容汇陆。
cudaFree()
1.?函數(shù)原型：cudaError_t?cudaFree?(?void*?devPtr?)。
2.?函數(shù)作用：與c語言中的free()函數(shù)一樣带饱，只是此函數(shù)釋放的是cudaMalloc()分配的內存毡代。
下面實例用于解釋上面三個函數(shù)
GPU內存分類
全局內存
通俗意義上的設備內存。
共享內存
1.?位置：設備內存勺疼。
2.?形式：關鍵字__shared__添加到變量聲明中教寂。如__shared__?float?cache[10]。
3.?目的：對于GPU上啟動的每個線程塊执庐，CUDA?C編譯器都將創(chuàng)建該共享變量的一個副本酪耕。線程塊中的每個線程都共享這塊內存，但線程卻無法看到也不能修改其他線程塊的變量副本轨淌。這樣使得一個線程塊中的多個線程能夠在計算上通信和協(xié)作迂烁。
常量內存
1.?位置：設備內存
2.?形式：關鍵字__constant__添加到變量聲明中。如__constant__?float?s[10];递鹉。
3.?目的：為了提升性能盟步。常量內存采取了不同于標準全局內存的處理方式。在某些情況下躏结，用常量內存替換全局內存能有效地減少內存帶寬却盘。
4.?特點：常量內存用于保存在核函數(shù)執(zhí)行期間不會發(fā)生變化的數(shù)據(jù)。變量的訪問限制為只讀媳拴。NVIDIA硬件提供了64KB的常量內存谷炸。不再需要cudaMalloc()或者cudaFree(),而是在編譯時，靜態(tài)地分配空間禀挫。
5.?要求：當我們需要拷貝數(shù)據(jù)到常量內存中應該使用cudaMemcpyToSymbol()旬陡，而cudaMemcpy()會復制到全局內存。
6.?性能提升的原因：
6.1.?對常量內存的單次讀操作可以廣播到其他的“鄰近”線程语婴。這將節(jié)約15次讀取操作描孟。（為什么是15，因為“鄰近”指半個線程束砰左，一個線程束包含32個線程的集合匿醒。）
6.2.?常量內存的數(shù)據(jù)將緩存起來，因此對相同地址的連續(xù)讀操作將不會產生額外的內存通信量缠导。
紋理內存
1.?位置：設備內存
2.?目的：能夠減少對內存的請求并提供高效的內存帶寬廉羔。是專門為那些在內存訪問模式中存在大量空間局部性的圖形應用程序設計，意味著一個線程讀取的位置可能與鄰近線程讀取的位置“非常接近”僻造。如下圖：
3. 紋理變量（引用）必須聲明為文件作用域內的全局變量憋他。
4.?形式：分為一維紋理內存?和?二維紋理內存孩饼。
4.1.?一維紋理內存
4.1.1.?用texture<類型>類型聲明，如texture?texIn竹挡。
4.1.2.?通過cudaBindTexture()綁定到紋理內存中镀娶。
4.1.3.?通過tex1Dfetch()來讀取紋理內存中的數(shù)據(jù)。
4.1.4.?通過cudaUnbindTexture()取消綁定紋理內存揪罕。
4.2.?二維紋理內存
4.2.1.?用texture<類型,數(shù)字>類型聲明梯码，如texture?texIn。
4.2.2.?通過cudaBindTexture2D()綁定到紋理內存中好啰。
4.2.3.?通過tex2D()來讀取紋理內存中的數(shù)據(jù)轩娶。
4.2.4.?通過cudaUnbindTexture()取消綁定紋理內存。
固定內存
1.?位置：主機內存框往。
2.?概念：也稱為頁鎖定內存或者不可分頁內存鳄抒，操作系統(tǒng)將不會對這塊內存分頁并交換到磁盤上，從而確保了該內存始終駐留在物理內存中搅窿。因此操作系統(tǒng)能夠安全地使某個應用程序訪問該內存的物理地址，因為這塊內存將不會破壞或者重新定位隙券。
3.?目的：提高訪問速度男应。由于GPU知道主機內存的物理地址，因此可以通過“直接內存訪問DMA（Direct?Memory?Access)技術來在GPU和主機之間復制數(shù)據(jù)娱仔。由于DMA在執(zhí)行復制時無需CPU介入沐飘。因此DMA復制過程中使用固定內存是非常重要的。
4.?缺點：使用固定內存牲迫，將失去虛擬內存的所有功能耐朴；系統(tǒng)將更快的耗盡內存。
5.?建議：對cudaMemcpy()函數(shù)調用中的源內存或者目標內存盹憎，才使用固定內存筛峭，并且在不再需要使用它們時立即釋放。
6.?形式：通過cudaHostAlloc()函數(shù)來分配陪每；通過cudaFreeHost()釋放影晓。
7.?只能以異步方式對固定內存進行復制操作。
原子性
1.?概念：如果操作的執(zhí)行過程不能分解為更小的部分檩禾，我們將滿足這種條件限制的操作稱為原子操作挂签。
2.?形式：函數(shù)調用，如atomicAdd（addr,y)將生成一個原子的操作序列盼产，這個操作序列包括讀取地址addr處的值饵婆，將y增加到這個值，以及將結果保存回地址addr戏售。
常用線程操作函數(shù)
1.?同步方法__syncthreads()侨核，這個函數(shù)的調用草穆，將確保線程塊中的每個線程都執(zhí)行完__syscthreads()前面的語句后，才會執(zhí)行下一條語句芹关。
使用事件來測量性能
1.?用途：為了測量GPU在某個任務上花費的時間续挟。CUDA中的事件本質上是一個GPU時間戳。由于事件是直接在GPU上實現(xiàn)的侥衬。因此不適用于對同時包含設備代碼和主機代碼的混合代碼設計诗祸。
2.?形式：首先創(chuàng)建一個事件，然后記錄事件轴总，再計算兩個事件之差直颅，最后銷毀事件。如：
流
1.?扯一扯：并發(fā)重點在于一個極短時間段內運行多個不同的任務怀樟；并行重點在于同時運行一個任務功偿。
2.?任務并行性：是指并行執(zhí)行兩個或多個不同的任務，而不是在大量數(shù)據(jù)上執(zhí)行同一個任務往堡。
3.?概念：CUDA流表示一個GPU操作隊列械荷，并且該隊列中的操作將以指定的順序執(zhí)行。我們可以在流中添加一些操作虑灰，如核函數(shù)啟動吨瞎，內存復制以及事件的啟動和結束等。這些操作的添加到流的順序也是它們的執(zhí)行順序穆咐〔鳎可以將每個流視為GPU上的一個任務，并且這些任務可以并行執(zhí)行对湃。
4.?硬件前提：必須是支持設備重疊功能的GPU崖叫。支持設備重疊功能，即在執(zhí)行一個核函數(shù)的同時拍柒，還能在設備與主機之間執(zhí)行復制操作心傀。
5.?聲明與創(chuàng)建：聲明cudaStream_t?stream;，創(chuàng)建cudaSteamCreate(&stream);拆讯。
6.?cudaMemcpyAsync()：前面在cudaMemcpy()中提到過剧包，這是一個以異步方式執(zhí)行的函數(shù)。在調用cudaMemcpyAsync()時往果，只是放置一個請求疆液，表示在流中執(zhí)行一次內存復制操作，這個流是通過參數(shù)stream來指定的陕贮。當函數(shù)返回時堕油，我們無法確保復制操作是否已經啟動，更無法保證它是否已經結束。我們能夠得到的保證是掉缺，復制操作肯定會當下一個被放入流中的操作之前執(zhí)行卜录。傳遞給此函數(shù)的主機內存指針必須是通過cudaHostAlloc()分配好的內存。（流中要求固定內存）
7.?流同步：通過cudaStreamSynchronize()來協(xié)調眶明。
8.?流銷毀：在退出應用程序之前艰毒，需要銷毀對GPU操作進行排隊的流，調用cudaStreamDestroy()搜囱。
9.?針對多個流：
9.1.?記得對流進行同步操作丑瞧。
9.2.?將操作放入流的隊列時，應采用寬度優(yōu)先方式蜀肘，而非深度優(yōu)先的方式绊汹，換句話說，不是首先添加第0個流的所有操作扮宠，再依次添加后面的第1西乖，2,…個流。而是交替進行添加坛增，比如將a的復制操作添加到第0個流中获雕，接著把a的復制操作添加到第1個流中，再繼續(xù)其他的類似交替添加的行為收捣。
9.3.?要牢牢記住操作放入流中的隊列中的順序影響到CUDA驅動程序調度這些操作和流以及執(zhí)行的方式届案。
技巧
1.?當線程塊的數(shù)量為GPU中處理數(shù)量的2倍時，將達到最優(yōu)性能坏晦。
2.?核函數(shù)執(zhí)行的第一個計算就是計算輸入數(shù)據(jù)的偏移萝玷。每個線程的起始偏移都是0到線程數(shù)量減1之間的某個值嫁乘。然后昆婿，對偏移的增量為已啟動線程的總數(shù)。
