轉自：http://blog.csdn.net/zyazky/article/details/52456503

 首先看一下CPU和GPU的微架構和計算能力對比。例如我的筆記本lenovo Y480，4核CPU,NVIDIA GT650M顯示卡384個CUDA核。

計算能力對比：

CPU:       4 * 2.5=10GFLOPS

GPU:    384 * 0.88=  337.92GFLOPS

顯示卡計算效能是4核i5 CPU的33.792倍，因此我們可以充分利用這一資源來對一些耗時的應用進行加速。

二

GPU的設計初衷是為了加速應用程式中的圖形繪製運算，用於在顯示器上渲染計算機圖形。因此開發人員需要通過

OpenGL或者DirectX等API來訪問GPU，這不僅要求開發人員掌握一定的圖形程式設計知識，而且要想方設法將通用計算問題轉換為圖形計算問題。其次，GPU與多核CPU在計算架構上有著很大不同，GPU更注重於資料平行計算，即在不同的資料上並行執行相同的計算，而對平行計算中的互斥性、同步性以及原子性等方面支援不足。這些因素都限制了GPU在通用平行計算中的應用範圍。

CUDA架構的出現解決了上述問題。CUDA架構專門為GPU計算設計了一種全新的結構，目的正是為了減輕GPU計算模型中的這些限制。在CUDA架構下，開發人員可以通過CUDA C對GPU程式設計。CUDA C/C++是對標準C/C++

的一種簡單擴充套件，學習和使用起來都非常容易，並且其最大的優勢在於不需要開發人員具備計算機圖形學知識。CUDA架構用於實現GPU裝置的通用計算應用,具有高並行性、高頻寬、價效比高、易於程式設計等優勢。CUDA具有如上所述的諸多優勢,但是,其仍然有許多侷限性,例如,CUDA對多分支和並行化低的應用程式執行效率不高,不適合整型資料運算,對處理的任務的細分程度及方式直接影響到CUDA應用程式的效能,。

CUDA是一種專門為提高並行程式開發效率而設計的計算架構與程式設計模型，是一個完整的GPGPU解決方案，提供了硬體的直接訪問介面，而不必像傳統方式一樣必須依賴圖形API介面來實現GPU的訪問。在構建高效能應用程式時，

CUDA架構能充分發揮GPU的強大計算功能，在NVIDIA GPU上編寫程式來完成通用計算任務。

CUDA包含了CUDA指令集架構（ISA）以及GPU內部的平行計算引擎。CUDA 的核心有三個重要抽象概念：執行緒組層次結構、共享儲存器、遮蔽同步，可輕鬆將其作為C 語言的最小擴充套件級公開給程式設計師。CUDA 軟體堆疊由幾層組成，一個硬體驅動程式，一個應用程式程式設計介面(API)和它的Runtime， 還有二個高階的通用數學庫，CUFFT 和CUBLAS。硬體被設計成支援輕量級的驅動和Runtime 層面，因而提高效能。從CUDA體系結構的組成來說，包含了三個部分：開發庫、執行期環境和驅動。

開發庫是基於CUDA技術所提供的應用開發庫。目前CUDA提供了兩個標準的數學運算庫——CUFFT（離散快速傅立葉變換）和CUBLAS（離散基本線性計算）的實現。這兩個數學運算庫所解決的是典型的大規模的平行計算問題，也是在密集資料計算中非常常見的計算型別。開發人員在開發庫的基礎上可以快速、方便的建立起自己的計算應用。

執行期環境提供了應用開發介面和執行期元件，包括基本資料型別的定義和各類計算、型別轉換、記憶體管理、裝置訪問和執行排程等函式。基於CUDA開發的程式程式碼在實際執行中分為兩種，一種是執行在CPU上的宿主程式碼（Host Code），一種是執行在GPU上的裝置程式碼（Device Code）。不同型別的程式碼由於其執行的物理位置不同，能夠訪問到的資源不同，因此對應的執行期元件也分為公共元件、宿主元件和裝置元件三個部分，基本上囊括了所有在GPGPU開發中所需要的功能和能夠使用到的資源介面，開發人員可以通過執行期環境的程式設計介面實現各種型別的計算。

由於目前存在著多種GPU版本的NVidia顯示卡，不同版本的GPU之間都有不同的差異，因此驅動部分基本上可以理解為是CUDA-enable的GPU的裝置抽象層，提供硬體裝置的抽象訪問介面。CUDA提供執行期環境也是通過這一層來實現各種功能的。目前基於CUDA開發的應用必須有NVIDIA CUDA-enable的硬體支援。

三

主要概念與名稱
主機
        將CPU及系統的記憶體（記憶體條）稱為主機。
裝置
        將GPU及GPU本身的顯示記憶體稱為裝置。
執行緒(Thread)
        一般通過GPU的一個核進行處理。（可以表示成一維，二維，三維）。
執行緒塊(Block)
        1. 由多個執行緒組成（可以表示成一維，二維，三維）。
        2. 各block是並行執行的，block間無法通訊，也沒有執行順序。
        3. 注意執行緒塊的數量限制為不超過65535（硬體限制）。
執行緒格(Grid)
        由多個執行緒塊組成（可以表示成一維，二維，三維）。


執行緒束
        在CUDA架構中，執行緒束是指一個包含32個執行緒的集合，這個執行緒集合被“編織在一起”並且“步調一致”的形式執行。在程式中的每一行，執行緒束中的每個執行緒都將在不同資料上執行相同的命令。
核函式（Kernel）
        1. 在GPU上執行的函式通常稱為核函式。
        2. 一般通過識別符號__global__修飾，呼叫通過<<<引數1,引數2>>>，用於說明核心函式中的執行緒數量，以及執行緒是如何組織的。
        3. 以執行緒格（Grid）的形式組織，每個執行緒格由若干個執行緒塊（block）組成，而每個執行緒塊又由若干個執行緒（thread）組成。
        4. 是以block為單位執行的。
        5. 只能在主機端程式碼中呼叫。
        6. 呼叫時必須宣告核心函式的執行引數。
        7. 在程式設計時，必須先為kernel函式中用到的陣列或變數分配好足夠的空間，再呼叫kernel函式，否則在GPU計算時會發生錯誤，例如越界或報錯，甚至導致藍屏和宕機。
dim3結構型別
        1. dim3是基亍uint3定義的向量型別，相當亍由3個unsigned int型組成的結構體。uint3型別有三個資料成員unsigned int x; unsigned int y; unsigned int z;
        2. 可使用亍一維、二維或三維的索引來標識執行緒，構成一維、二維或三維執行緒塊。
        3. dim3結構型別變數用在核函式呼叫的<<<,>>>中。
        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. 對於計算執行緒索引偏移增量為已啟動執行緒的總數。如stride = blockDim.x * gridDim.x; threadId += stride。
函式修飾符
        1. __global__，表明被修飾的函式在裝置上執行，但在主機上呼叫。
        2. __device__，表明被修飾的函式在裝置上執行，但只能在其他__device__函式或者__global__函式中呼叫。
常用的GPU記憶體函式
cudaMalloc()
        1. 函式原型： cudaError_t cudaMalloc (void **devPtr, size_t size)。
        2. 函式用處：與C語言中的malloc函式一樣，只是此函式在GPU的記憶體你分配記憶體。
        3. 注意事項：
        3.1. 可以將cudaMalloc()分配的指標傳遞給在裝置上執行的函式；
        3.2. 可以在裝置程式碼中使用cudaMalloc()分配的指標進行裝置記憶體讀寫操作；
        3.3. 可以將cudaMalloc()分配的指標傳遞給在主機上執行的函式；
        3.4. 不可以在主機程式碼中使用cudaMalloc()分配的指標進行主機記憶體讀寫操作（即不能進行解引用）。
cudaMemcpy()
        1. 函式原型：cudaError_t cudaMemcpy (void *dst, const void *src, size_t count, cudaMemcpyKind kind)。
        2. 函式作用：與c語言中的memcpy函式一樣，只是此函式可以在主機記憶體和GPU記憶體之間互相拷貝資料。
        3. 函式引數：cudaMemcpyKind kind表示資料拷貝方向，如果kind賦值為cudaMemcpyDeviceToHost表示資料從裝置記憶體拷貝到主機記憶體。
        4. 與C中的memcpy()一樣，以同步方式執行，即當函式返回時，複製操作就已經完成了，並且在輸出緩衝區中包含了複製進去的內容。
        5. 相應的有個非同步方式執行的函式cudaMemcpyAsync()，這個函式詳解請看下面的流一節有關內容。
cudaFree()
        1. 函式原型：cudaError_t cudaFree ( void* devPtr )。
        2. 函式作用：與c語言中的free()函式一樣，只是此函式釋放的是cudaMalloc()分配的記憶體。
GPU記憶體分類

全域性記憶體
        通俗意義上的裝置記憶體。
共享記憶體
        1. 位置：裝置記憶體。
        2. 形式：關鍵字__shared__新增到變數宣告中。如__shared__ float cache[10]。
        3. 目的：對於GPU上啟動的每個執行緒塊，CUDA C編譯器都將建立該共享變數的一個副本。執行緒塊中的每個執行緒都共享這塊記憶體，但執行緒卻無法看到也不能修改其他執行緒塊的變數副本。這樣使得一個執行緒塊中的多個執行緒能夠在計算上通訊和協作。
常量記憶體
        1. 位置：裝置記憶體
        2. 形式：關鍵字__constant__新增到變數宣告中。如__constant__ float s[10];。
        3. 目的：為了提升效能。常量記憶體採取了不同於標準全域性記憶體的處理方式。在某些情況下，用常量記憶體替換全域性記憶體能有效地減少記憶體頻寬。
        4. 特點：常量記憶體用於儲存在核函式執行期間不會發生變化的資料。變數的訪問限制為只讀。NVIDIA硬體提供了64KB的常量記憶體。不再需要cudaMalloc()或者cudaFree(),而是在編譯時，靜態地分配空間。
        5. 要求：當我們需要拷貝資料到常量記憶體中應該使用cudaMemcpyToSymbol()，而cudaMemcpy()會複製到全域性記憶體。
        6. 效能提升的原因：
        6.1. 對常量記憶體的單次讀操作可以廣播到其他的“鄰近”執行緒。這將節約15次讀取操作。（為什麼是15，因為“鄰近”指半個執行緒束，一個執行緒束包含32個執行緒的集合。）
        6.2. 常量記憶體的資料將快取起來，因此對相同地址的連續讀操作將不會產生額外的記憶體通訊量。
紋理記憶體
        1. 位置：裝置記憶體
        2. 目的：能夠減少對記憶體的請求並提供高效的記憶體頻寬。是專門為那些在記憶體訪問模式中存在大量空間區域性性的圖形應用程式設計，意味著一個執行緒讀取的位置可能與鄰近執行緒讀取的位置“非常接近”。如下圖：

        3. 紋理變數（引用）必須宣告為檔案作用域內的全域性變數。
        4. 形式：分為一維紋理記憶體和二維紋理記憶體。
        4.1. 一維紋理記憶體
        4.1.1. 用texture<型別>型別宣告，如texture<float> texIn。
        4.1.2. 通過cudaBindTexture()繫結到紋理記憶體中。
        4.1.3. 通過tex1Dfetch()來讀取紋理記憶體中的資料。
        4.1.4. 通過cudaUnbindTexture()取消繫結紋理記憶體。
        4.2. 二維紋理記憶體
        4.2.1. 用texture<型別,數字>型別宣告，如texture<float，2> texIn。
        4.2.2. 通過cudaBindTexture2D()繫結到紋理記憶體中。
        4.2.3. 通過tex2D()來讀取紋理記憶體中的資料。
        4.2.4. 通過cudaUnbindTexture()取消繫結紋理記憶體。
固定記憶體
        1. 位置：主機記憶體。
        2. 概念：也稱為頁鎖定記憶體或者不可分頁記憶體，作業系統將不會對這塊記憶體分頁並交換到磁碟上，從而確保了該記憶體始終駐留在實體記憶體中。因此作業系統能夠安全地使某個應用程式訪問該記憶體的實體地址，因為這塊記憶體將不會破壞或者重新定位。
        3. 目的：提高訪問速度。由於GPU知道主機記憶體的實體地址，因此可以通過“直接記憶體訪問DMA（Direct Memory Access)技術來在GPU和主機之間複製資料。由於DMA在執行復制時無需CPU介入。因此DMA複製過程中使用固定記憶體是非常重要的。
        4. 缺點：使用固定記憶體，將失去虛擬記憶體的所有功能；系統將更快的耗盡記憶體。
        5. 建議：對cudaMemcpy()函式呼叫中的源記憶體或者目標記憶體，才使用固定記憶體，並且在不再需要使用它們時立即釋放。
        6. 形式：通過cudaHostAlloc()函式來分配；通過cudaFreeHost()釋放。
        7. 只能以非同步方式對固定記憶體進行復制操作。
原子性
        1. 概念：如果操作的執行過程不能分解為更小的部分，我們將滿足這種條件限制的操作稱為原子操作。
        2. 形式：函式呼叫，如atomicAdd（addr,y)將生成一個原子的操作序列，這個操作序列包括讀取地址addr處的值，將y增加到這個值，以及將結果儲存回地址addr。
常用執行緒操作函式
        1. 同步方法__syncthreads()，這個函式的呼叫，將確保執行緒塊中的每個執行緒都執行完__syscthreads()前面的語句後，才會執行下一條語句。
使用事件來測量效能
        1. 用途：為了測量GPU在某個任務上花費的時間。CUDA中的事件本質上是一個GPU時間戳。由於事件是直接在GPU上實現的。因此不適用於對同時包含裝置程式碼和主機程式碼的混合程式碼設計。
        2. 形式：首先建立一個事件，然後記錄事件，再計算兩個事件之差，最後銷燬事件。
流
        1. 併發重點在於一個極短時間段內執行多個不同的任務；並行重點在於同時執行一個任務。
        2. 任務並行性：是指並行執行兩個或多個不同的任務，而不是在大量資料上執行同一個任務。
        3. 概念：CUDA流表示一個GPU操作佇列，並且該佇列中的操作將以指定的順序執行。我們可以在流中新增一些操作，如核函式啟動，記憶體複製以及事件的啟動和結束等。這些操作的新增到流的順序也是它們的執行順序。可以將每個流視為GPU上的一個任務，並且這些任務可以並行執行。
        4. 硬體前提：必須是支援裝置重疊功能的GPU。支援裝置重疊功能，即在執行一個核函式的同時，還能在裝置與主機之間執行復制操作。
        5. 宣告與建立：宣告cudaStream_t stream;，建立cudaSteamCreate(&stream);。
        6. cudaMemcpyAsync()：前面在cudaMemcpy()中提到過，這是一個以非同步方式執行的函式。在呼叫cudaMemcpyAsync()時，只是放置一個請求，表示在流中執行一次記憶體複製操作，這個流是通過引數stream來指定的。當函式返回時，我們無法確保複製操作是否已經啟動，更無法保證它是否已經結束。我們能夠得到的保證是，複製操作肯定會當下一個被放入流中的操作之前執行。傳遞給此函式的主機記憶體指標必須是通過cudaHostAlloc()分配好的記憶體。（流中要求固定記憶體）
        7. 流同步：通過cudaStreamSynchronize()來協調。
        8. 流銷燬：在退出應用程式之前，需要銷燬對GPU操作進行排隊的流，呼叫cudaStreamDestroy()。
        9. 針對多個流：
        9.1. 記得對流進行同步操作。
        9.2. 將操作放入流的佇列時，應採用寬度優先方式，而非深度優先的方式，換句話說，不是首先新增第0個流的所有操作，再依次新增後面的第1，2,…個流。而是交替進行新增，比如將a的複製操作新增到第0個流中，接著把a的複製操作新增到第1個流中，再繼續其他的類似交替新增的行為。
        9.3. 要牢牢記住操作放入流中的佇列中的順序影響到CUDA驅動程式排程這些操作和流以及執行的方式。
技巧
        1. 當執行緒塊的數量為GPU中處理數量的2倍時，將達到最優效能。
        2. 核函式執行的第一個計算就是計算輸入資料的偏移。每個執行緒的起始偏移都是0到執行緒數量減1之間的某個值。然後，對偏移的增量為已啟動執行緒的總數。

四

執行緒並行將執行緒的概念引申到CUDA程式設計中，我們可以認為執行緒就是執行CUDA程式的最小單元，在GPU上每個執行緒都會執行一次該核函式。但GPU上的執行緒排程方式與CPU有很大不同。CPU上會有優先順序分配，從高到低，同樣優先順序的可以採用時間片輪轉法實現執行緒排程。GPU上執行緒沒有優先順序概念，所有執行緒機會均等，執行緒狀態只有等待資源和執行兩種狀態，如果資源未就緒，那麼就等待；一旦就緒，立即執行。當GPU資源很充裕時，所有執行緒都是併發執行的，這樣加速效果很接近理論加速比；而GPU資源少於匯流排程個數時，有一部分執行緒就會等待前面執行的執行緒釋放資源，從而變為序列化執行。

塊並行 塊並行相當於作業系統中多程序的情況，CUDA有執行緒組（執行緒塊）的概念，將一組執行緒組織到一起，共同分配一部分資源，然後內部排程執行。執行緒塊與執行緒塊之間，毫無瓜葛。這有利於做更粗粒度的並行。我們的任務有時可以採用分治法，將一個大問題分解為幾個小規模問題，將這些小規模問題分別用一個執行緒塊實現，執行緒塊內可以採用細粒度的執行緒並行，而塊之間為粗粒度並行，這樣可以充分利用硬體資源，降低執行緒並行的計算複雜度。

多個執行緒塊組織成了一個Grid，稱為執行緒格（經歷了從一位執行緒，二維執行緒塊到三維執行緒格的過程)

流並行流可以實現在一個裝置上執行多個核函式。前面的塊並行也好，執行緒並行也好，執行的核函式都是相同的（程式碼一樣，傳遞引數也一樣）。而流並行，可以執行不同的核函式，也可以實現對同一個核函式傳遞不同的引數，實現任務級別的並行。

CUDA中的流用cudaStream_t型別實現，用到的API有以下幾個：cudaStreamCreate(cudaStream_t * s)用於建立流，cudaStreamDestroy(cudaStream_t s)用於銷燬流，cudaStreamSynchronize()用於單個流同步，cudaDeviceSynchronize()用於整個裝置上的所有流同步，cudaStreamQuery()用於查詢一個流的任務是否已經完成。

前面介紹了三種利用GPU實現並行處理的方式：執行緒並行，塊並行和流並行。在這些方法中，各個執行緒所進行的處理是互不相關的，即兩個執行緒不回產生交集，每個執行緒都只關注自己的一畝三分地，對其他執行緒毫無興趣，就當不存在。。。。

當然，實際應用中，這樣的例子太少了，也就是遇到向量相加、向量對應點乘這類才會有如此高的並行度，而其他一些應用，如一組數求和，求最大（小）值，各個執行緒不再是相互獨立的，而是產生一定關聯，執行緒2可能會用到執行緒1的結果，這時就需要利用執行緒通訊技術了。

執行緒通訊在CUDA中有三種實現方式：

1. 共享儲存器；

2. 執行緒同步；

3. 原子操作；

最常用的是前兩種方式，共享儲存器，術語Shared Memory，是位於SM中的特殊儲存器。還記得SM嗎，就是流多處理器，相當於大核。一個SM中不僅包含若干個SP（流處理器，小核），還包括一部分高速Cache，暫存器組，共享記憶體等，結構如圖所示：

從圖中可看出，一個SM內有M個SP，Shared Memory由這M個SP共同佔有。另外指令單元也被這M個SP共享，即SIMT架構（單指令多執行緒架構），一個SM中所有SP在同一時間執行同一程式碼。

為了實現執行緒通訊，僅僅靠共享記憶體還不夠，需要有同步機制才能使執行緒之間實現有序處理。通常情況是這樣：當執行緒A需要執行緒B計算的結果作為輸入時，需要確保執行緒B已經將結果寫入共享記憶體中，然後執行緒A再從共享記憶體中讀出。同步必不可少，否則，執行緒A可能讀到的是無效的結果，造成計算錯誤。同步機制可以用CUDA內建函式：__syncthreads()；當某個執行緒執行到該函式時，進入等待狀態，直到同一執行緒塊（Block）中所有執行緒都執行到這個函式為止，即一個__syncthreads()相當於一個執行緒同步點，確保一個Block中所有執行緒都達到同步，然後執行緒進入執行狀態。

注意的是，位於同一個Block中的執行緒才能實現通訊，不同Block中的執行緒不能通過共享記憶體、同步進行通訊，而應採用原子操作或主機介入。

五

sp: 最基本的處理單元，streaming processor  最後具體的指令和任務都是在sp上處理。GPU進行平行計算，也就是很多個sp同時做處理。

sm:多個sp加上其他的一些資源組成一個sm,  streaming multiprocessor. 其他資源也就是儲存資源，共享記憶體，寄儲器等。

warp:GPU執行程式時的排程單位，目前cuda的warp的大小為32，同在一個warp的執行緒，以不同資料資源執行相同的指令。

grid、block、thread：在利用cuda進行程式設計時，一個grid分為多個block，而一個block分為多個thread.其中任務劃分到是否影響最後的執行效果。劃分的依據是任務特性和GPU本身的硬體特性。

在CUDA架構下，執行緒的最小單元是thread，多個thread組成一個block，多個block再組成一個grid，不同block之間的thread不能讀寫同一shared memory共享記憶體，因此，block裡面的thread之間的通訊和同步所帶來的開銷是比較大的。SM以 32 個 Thread 為一組的 Warp 來執行 Thread。Warp內的執行緒是靜態的，即在屬於同一個warp內的thread之間進行通訊，不需要進行柵欄同步(barrier)。Fermi的設計根據G80和GT200的架構作出的很多缺陷來改變。在Fermi中，每個SM中的數量不再是GT200的8個SP，而是變成了32個SP，NVIDIA現在又稱之為CUDA Core，總共具有16個SM，所以總共有512個SP。而在GT200中，是30個SM，240個SP。

相關推薦