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)。
64位Ubuntu12.04安裝CUDA5.5
????????具體步驟請點擊此處 http://bookc.github.io/2014/05/08/my-summery-the-book-cuda-by-example-an-introduction-to-general-purpose-gpu-programming/。
[b ]對CUDA?C的個人懵懂感覺[/b]
????????如果粗暴的認(rèn)為C語言工作的對象是CPU和內(nèi)存條（接下來,稱為主機內(nèi)存），那么CUDA?C工作的的對象就是GPU及GPU上的內(nèi)存（接下來,稱為設(shè)備內(nèi)存），且充分利用了GPU多核的優(yōu)勢及降低了并行編程的難度。一般通過C語言把數(shù)據(jù)從外界讀入，再分配數(shù)據(jù)，給CUDA?C，以便在GPU上計算，然后再把計算結(jié)果返回給C語言，以便進(jìn)一步工作，如進(jìn)一步處理及顯示，或重復(fù)此過程。
?主要概念與名稱
主機
????????將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)屏和死機。
C/C++ code?

1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17

/* ?*?@file_name?HelloWorld.cu??后綴名稱.cu ?*/ #include?<stdio.h> #include?<cuda_runtime.h>??//頭文件 //核函數(shù)聲明，前面的關(guān)鍵字__global__ __global__?void?kernel(?void?)?{ } int?main(?void?)?{ ????//核函數(shù)的調(diào)用，注意<<<1,1>>>，第一個1，代表線程格里只有一個線程塊；第二個1，代表一個線程塊里只有一個線程。 ????kernel<<<1,1>>>(); ????printf(?"Hello,?World!\n"?); ????return?0; }

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ù)
C/C++ code?

1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20

#include?<stdio.h> #include?<cuda_runtime.h> __global__?void?add(?int?a,?int?b,?int?*c?)?{ ????*c?=?a?+?b; } int?main(?void?)?{ ????int?c; ????int?*dev_c; ????//cudaMalloc() ????cudaMalloc(?(void**)&dev_c,?sizeof(int)?); ????//核函數(shù)執(zhí)行 ????add<<<1,1>>>(?2,?7,?dev_c?);??? ????//cudaMemcpy() ????cudaMemcpy(?&c,?dev_c,?sizeof(int),cudaMemcpyDeviceToHost?)?; ????printf(?"2?+?7?=?%d\n",?c?); ????//cudaFree() ????cudaFree(?dev_c?); ????return?0; }

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<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.?位置：主機內(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)建一個事件，然后記錄事件，再計算兩個事件之差，最后銷毀事件。如：
C/C++ code?

1 2 3 4 5 6 7 8 9 10

cudaEvent_t?start,?stop; cudaEventCreate(?&start?); cudaEventCreate(?&stop?); cudaEventRecord(?start,?0?); //do?something cudaEventRecord(?stop,?0?); float???elapsedTime; cudaEventElapsedTime(?&elapsedTime,start,?stop?); cudaEventDestroy(?start?); cudaEventDestroy(?stop?)；

流
????????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ù)。
實例程序
感興趣的讀者可以下載本書附帶的示例代碼點擊此處下載 https://developer.nvidia.com/sites/default/files/akamai/cuda/files/cuda_by_example.zip。

總結(jié)

以上是生活随笔為你收集整理的cuda资料的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯，歡迎將生活随笔推薦給好友。