轉載一篇介紹CUDA

鑑於本身的畢設須要使用GPU CUDA這項技術，想找一本入門的教材，選擇了Jason Sanders等所著的書《CUDA By Example an Introduction to General Purpose GPU Programming》。這本書做爲入門教材，寫的很不錯。本身以爲從理解與記憶的角度的出發，書中不少內容均可以被省略掉，因而就有了這篇博文。此博文記錄與總結此書的筆記和理解。注意本文並無按照書中章節的順序來寫。書中第8章圖像互操做性和第11章多GPU系統上的CUDA C，這兩章沒有看。等有時間了再看吧，趕忙碼字。
CUDA是什麼
        CUDA，Compute Unified Device Architecture的簡稱，是由NVIDIA公司創立的基於他們公司生產的圖形處理器GPUs（Graphics Processing Units,能夠通俗的理解爲顯卡）的一個並行計算平臺和編程模型。
        經過CUDA，GPUs能夠很方便地被用來進行通用計算（有點像在CPU中進行的數值計算等等）。在沒有CUDA以前，GPUs通常只用來進行圖形渲染（如經過OpenGL，DirectX）。
        開發人員能夠經過調用CUDA的API，來進行並行編程，達到高性能計算目的。NVIDIA公司爲了吸引更多的開發人員，對CUDA進行了編程語言擴展，如CUDA C/C++,CUDA Fortran語言。注意CUDA C/C++能夠看做一個新的編程語言，由於NVIDIA配置了相應的編譯器nvcc,CUDA Fortran同樣。更多信息能夠參考文獻。
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]
        若是粗暴的認爲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）
        1. 在GPU上執行的函數一般稱爲核函數。
        2. 通常經過標識符__global__修飾，調用經過<<<參數1,參數2>>>，用於說明內核函數中的線程數量，以及線程是如何組織的。
        3. 以線程格（Grid）的形式組織，每一個線程格由若干個線程塊（block）組成，而每一個線程塊又由若干個線程（thread）組成。
        4. 是以block爲單位執行的。
        5. 叧能在主機端代碼中調用。
        6. 調用時必須聲明內核函數的執行參數。
        7. 在編程時，必須先爲kernel函數中用到的數組或變量分配好足夠的空間，再調用kernel函數，不然在GPU計算時會發生錯誤，例如越界或報錯，甚至致使藍屏和死機。

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>  //頭文件   //核函數聲明，前面的關鍵字__global__ __global__  void  kernel(  void  ) { }   int  main(  void  ) {      //核函數的調用，注意<<<1,1>>>，第一個1，表明線程格里只有一個線程塊；第二個1，表明一個線程塊裏只有一個線程。      kernel<<<1,1>>>();      printf (  "Hello, World!\n"  );      return  0; }

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()分配的內存。
        下面實例用於解釋上面三個函數

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 ) );      //核函數執行      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內存分類
全局內存
        通俗意義上的設備內存。
共享內存
        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 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. 扯一扯：併發重點在於一個極短期段內運行多個不一樣的任務；並行重點在於同時運行一個任務。
        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之間的某個值。而後，對偏移的增量爲已啓動線程的總數。
實例程序
感興趣的讀者能夠下載本書附帶的示例代碼點擊此處下載https://developer.nvidia.com/sites/default/files/akamai/cuda/files/cuda_by_example.zip。

轉載地址：https://www.cnblogs.com/welhzh/p/8316989.html