轉自：https://blog.csdn.net/fangjin_kl/article/details/53906874

一、CPU和GPU的區別 CPU (Central Processing Unit) 即中央處理器 GPU (Graphics Processing Unit) 即圖形處理器 GPGPU全稱General Purpose GPU，即通用計算圖形處理器。其中第一個“GP”通用目的（GeneralPurpose）而第二個“GP”則表示圖形處理（GraphicProcess） CPU雖然有多核，但總數沒有超過兩位數，每個核都有足夠大的快取和足夠多的數字和邏輯運算單元，並輔助有很多加速分支判斷甚至更復雜的邏輯判斷的硬體；

GPU的核數遠超CPU，被稱為眾核（NVIDIA Fermi有512個核）。每個核擁有的快取大小相對小，數字邏輯運算單元也少而簡單（GPU初始時在浮點計算上一直弱於CPU）。

從結果上導致CPU擅長處理具有複雜計算步驟和複雜資料依賴的計算任務，如分散式計算，資料壓縮，人工智慧，物理模擬，以及其他很多很多計算任務等。

GPU由於歷史原因，是為了視訊遊戲而產生的（至今其主要驅動力還是不斷增長的視訊遊戲市場），在三維遊戲中常常出現的一類操作是對海量資料進行相同的操作，如：對每一個頂點進行同樣的座標變換，對每一個頂點按照同樣的光照模型計算顏色值。GPU的眾核架構非常適合把同樣的指令流並行傳送到眾核上，採用不同的輸入資料執行

當程式設計師為CPU編寫程式時，他們傾向於利用複雜的邏輯結構優化演算法從而減少計算任務的執行時間，即Latency。  當程式設計師為GPU編寫程式時，則利用其處理海量資料的優勢，通過提高總的資料吞吐量（Throughput）來掩蓋Lantency

其中綠色的是計算單元，橙紅色的是儲存單元，橙黃色的是控制單元。

GPU採用了數量眾多的計算單元和超長的流水線，但只有非常簡單的控制邏輯並省去了Cache。而CPU不僅被Cache佔據了大量空間，而且還有有複雜的控制邏輯和諸多優化電路，相比之下計算能力只是CPU很小的一部分

二、CUDA CUDA(Compute Unified Device Architecture)，是英偉達公司推出的一種基於新的並行程式設計模型和指令集架構的通用計算架構，它能利用英偉達GPU的平行計算引擎，比CPU更高效的解決許多複雜計算任務。

使用CUDA的好處就是透明。根據摩爾定律GPU的電晶體數量不斷增多，硬體結構必然是不斷的在發展變化，沒有必要每次都為不同的硬體結構重新編碼，而CUDA就是提供了一種可擴充套件的程式設計模型，使得已經寫好的CUDA程式碼可以在任意數量核心的GPU上執行。如下圖所示，只有執行時，系統才知道物理處理器的數量。 

三、CuDNN NVIDIA cuDNN是用於深度神經網路的GPU加速庫。它強調效能、易用性和低記憶體開銷。NVIDIA cuDNN可以整合到更高級別的機器學習框架中，如加州大學伯克利分校的流行CAFFE軟體。簡單的，插入式設計可以讓開發人員專注於設計和實現神經網路模型，而不是調整效能，同時還可以在GPU上實現高效能現代平行計算。

cuDNN 使用者手冊（英文）

CuDNN支援的演算法

卷積操作、相關操作的前向和後向過程。 pooling的前向後向過程 softmax的前向後向過程 啟用函式的前向後向過程  ReLU sigmoid TANH Tensor轉換函式，其中一個Tensor就是一個四維的向量。   Baseline Caffe與用NVIDIA Titan Z 加速cuDNN的Caffe做比較

四、CUDA程式設計 參考自 一篇不錯的CUDA入門部落格

開發人員可以通過呼叫CUDA的API，來進行並行程式設計，達到高效能運算目的。NVIDIA公司為了吸引更多的開發人員，對CUDA進行了程式語言擴充套件，如CUDA C/C++,CUDA Fortran語言。注意CUDA C/C++可以看作一個新的程式語言，因為NVIDIA配置了相應的編譯器nvcc,CUDA Fortran一樣。

如果粗暴的認為C語言工作的物件是CPU和記憶體條（接下來,稱為主機記憶體），那麼CUDA C工作的的物件就是GPU及GPU上的記憶體（接下來,稱為裝置記憶體），且充分利用了GPU多核的優勢及降低了並行程式設計的難度。一般通過C語言把資料從外界讀入，再分配資料，給CUDA C，以便在GPU上計算，然後再把計算結果返回給C語言，以便進一步工作，如進一步處理及顯示，或重複此過程。

主要概念與名稱 主機  將CPU及系統的記憶體（記憶體條）稱為主機。 裝置  將GPU及GPU本身的顯示記憶體稱為裝置。 執行緒(Thread)  一般通過GPU的一個核進行處理。（可以表示成一維，二維，三維，具體下面再細說）。 執行緒塊(Block)  1. 由多個執行緒組成（可以表示成一維，二維，三維，具體下面再細說）。  2. 各block是並行執行的，block間無法通訊，也沒有執行順序。  3. 注意執行緒塊的數量限制為不超過65535（硬體限制）。 執行緒格(Grid)  由多個執行緒塊組成（可以表示成一維，二維，三維，具體下面再細說）。 執行緒束  在CUDA架構中，執行緒束是指一個包含32個執行緒的集合，這個執行緒集合被“編織在一起”並且“步調一致”的形式執行。在程式中的每一行，執行緒束中的每個執行緒都將在不同資料上執行相同的命令。 核函式（Kernel）  在GPU上執行的函式通常稱為核函式。 一般通過識別符號__global__修飾，呼叫通過<<<引數1,引數2>>>，用於說明核心函式中的執行緒數量，以及執行緒是如何組織的。 以執行緒格（Grid）的形式組織，每個執行緒格由若干個執行緒塊（block）組成，而每個執行緒塊又由若干個執行緒（thread）組成。 是以block為單位執行的。 只能在主機端程式碼中呼叫。 呼叫時必須宣告核心函式的執行引數。 在程式設計時，必須先為kernel函式中用到的陣列或變數分配好足夠的空間，再呼叫kernel函式，否則在GPU計算時會發生錯誤，例如越界或報錯，甚至導致藍屏和宕機。

/*  * @file_name HelloWorld.cu  字尾名稱.cu  */

#include <stdio.h> #include <cuda_runtime.h>  //標頭檔案

//核函式宣告，前面的關鍵字__global__ __global__ void kernel( void ) { }

int main( void ) {     //核函式的呼叫，注意<<<1,1>>>，第一個1，代表執行緒格里只有一個執行緒塊；第二個1，代表一個執行緒塊裡只有一個執行緒。     kernel<<<1,1>>>();     printf( "Hello, World!\n" );     return 0; } dim3結構型別 dim3是基於 uint3 定義的向量型別，相當亍由3個unsigned int型組成的結構體。uint3型別有三個資料成員unsigned int x; unsigned int y; unsigned int z; 可使用一維、二維或三維的索引來標識執行緒，構成一維、二維或三維執行緒塊。 dim3結構型別變數用在核函式呼叫的<<<,>>>中。 相關的幾個內建變數  threadIdx，顧名思義獲取執行緒 thread 的ID索引；如果執行緒是一維的那麼就取 threadIdx.x，二維的還可以多取到一個值threadIdx.y，以此類推到三維threadIdx.z。 blockIdx，執行緒塊的ID索引；同樣有blockIdx.x，blockIdx.y，blockIdx.z。 blockDim，執行緒塊的維度，同樣有blockDim.x，blockDim.y，blockDim.z。 gridDim，執行緒格的維度，同樣有gridDim.x，gridDim.y，gridDim.z。 對於一維的block，執行緒的threadID=threadIdx.x。 對於大小為（blockDim.x, blockDim.y）的 二維 block，執行緒的threadID=threadIdx.x+threadIdx.y*blockDim.x。 對於大小為（blockDim.x, blockDim.y, blockDim.z）的 三維 block，執行緒的threadID=threadIdx.x+threadIdx.y*blockDim.x+threadIdx.z*blockDim.x*blockDim.y。 對於計算執行緒索引偏移增量為已啟動執行緒的總數。如stride = blockDim.x * gridDim.x; threadId += stride。 9. 函式修飾符 __global__，表明被修飾的函式在裝置上執行，但在主機上呼叫。 __device__，表明被修飾的函式在裝置上執行，但只能在其他device函式或者global函式中呼叫。 常用的GPU記憶體函式 cudaMalloc() 函式原型： cudaError_t cudaMalloc (void **devPtr, size_t size)。 函式用處：與C語言中的malloc函式一樣，只是此函式在GPU的記憶體你分配記憶體。 注意事項：  可以將cudaMalloc()分配的指標傳遞給在裝置上執行的函式； 可以在裝置程式碼中使用cudaMalloc()分配的指標進行裝置記憶體讀寫操作； 可以將cudaMalloc()分配的指標傳遞給在主機上執行的函式； 不可以在主機程式碼中使用cudaMalloc()分配的指標進行主機記憶體讀寫操作（即不能進行解引用）。 cudaMemcpy() 函式原型：cudaError_t cudaMemcpy (void *dst, const void *src, size_t count, cudaMemcpyKind kind)。 函式作用：與c語言中的memcpy函式一樣，只是此函式可以在主機記憶體和GPU記憶體之間互相拷貝資料。 函式引數：cudaMemcpyKind kind表示資料拷貝方向，如果kind賦值為cudaMemcpyDeviceToHost表示資料從裝置記憶體拷貝到主機記憶體。 與C中的memcpy()一樣，以同步方式執行，即當函式返回時，複製操作就已經完成了，並且在輸出緩衝區中包含了複製進去的內容。 相應的有個非同步方式執行的函式cudaMemcpyAsync()，這個函式詳解請看下面的流一節有關內容。 cudaFree() 函式原型：cudaError_t cudaFree ( void* devPtr )。 函式作用：與c語言中的free()函式一樣，只是此函式釋放的是cudaMalloc()分配的記憶體。 下面例項用於解釋上面三個函式

#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) );     //核函式執行     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記憶體分類 全域性記憶體 通俗意義上的裝置記憶體。

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

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