程式前沿 幫助程式設計師解決問題，增加專業技能，提升個人能力與未來世界競爭力。

CUDA程式設計指南閱讀筆記

CUDA程式設計指南閱讀筆記
目錄

CUDA程式設計指南閱讀筆記（一）

原創 2013年08月01日 20:36:00

隨著多核CPU和眾核GPU的到來，並行程式設計已經得到了業界越來越多的重視，CPU-GPU異構程式能夠極大提高現有計算機系統的運算效能，對於科學計算等運算密集型程式有著非常重要的意義。這一系列文章是根據《CUDA C語言程式設計指南》來整理的，該指南是NVIDIA公司提供的CUDA學習資料，介紹了CUDA程式設計最基本最核心的概念，是學習CUDA必不可少的閱讀材料。

初學CUDA，筆記錯誤之處在所難免，還請發現問題的諸位讀者不吝賜教。

1. 什麼是CUDA？

      CUDA全稱是Compute Unified Device Architecture，中文名稱即統一計算裝置架構，它是NVIDIA公司提出了一種通用的平行計算平臺和程式設計模型。使用CUDA，我們可以開發出同時在CPU和GPU上執行的通用計算程式，更加高效地利用現有硬體進行計算。為了簡化平行計算學習，CUDA為程式設計師提供了一個類C語言的開發環境以及一些其它的如FORTRAN、DirectCOmpute、OpenACC的高階語言/程式設計介面來開發CUDA程式。

2. CUDA程式設計模型如何擴充套件？

     我們知道，不同的GPU擁有不同的核心數目，在核心較多的系統上CUDA程式執行的時間較短，而在核心較少的系統上CUDA程式的執行時間較多。那麼，CUDA是如何做到的呢？
      並行程式設計的中心思想是分而治之：將大問題劃分為一些小問題，再把這些小問題交給相應的處理單元並行地進行處理。在CUDA中，這一思想便體現在它的具有兩個層次的問題劃分模型。一個問題可以首先被粗粒度地劃分為若干較小的子問題，CUDA使用被稱為塊（Block）的單元來處理它們，每個塊都由一些CUDA執行緒組成，執行緒是CUDA中最小的處理單元，將這些較小的子問題進一步劃分為若干更小的細粒度的問題，我們便可以使用執行緒來解決這些問題了。對於一個普通的NVIDIA GPU，其CUDA執行緒數目通常能達到數千個甚至更多，因此，這樣的問題劃分模型便可以成倍地提升計算機的運算效能。
      GPU是由多個流水多處理器構成的，流水處理器以塊（Block）為基本排程單元，因此，對於流水處理器較多的GPU，它一次可以處理的塊（Block）更多，從而運算速度更快，時間更短。而反之對於流水處理器較少的GPU，其運算速度便會較慢。這一原理可以通過下圖形象地看出來：

3. CUDA基本概念（上）

      本節將介紹CUDA的一些基本的程式設計概念，該節用到的例子來自於CUDA Sample中的VectorAdd專案。

3.1 核心（Kernels）

      CUDA C是C語言的一個擴充套件，它允許程式設計師定義一種被稱為核心函式（Kernel Functions）的C函式，核心函式執行在GPU上，一旦啟動，CUDA中的每一個執行緒都將會同時並行地執行核心函式中的程式碼。

      核心函式使用關鍵字__global__來宣告，執行該函式的CUDA執行緒數則通過<<<…>>>執行配置語法來設定。（參見章節”C語言擴充套件”），每一個執行核心函式的執行緒都由一個唯一的執行緒ID，這一ID可以通過在核心函式中訪問threadIdx變數來得到。

      下面通過一些示例程式碼來展示剛剛提到的這些概念該如何應用在程式設計中：

[cpp] view plain copy

  1. // Kernel definition  
  2. __global__ void VecAdd(float* A, float* B, float* C) {  
  3.     int i = threadIdx.x;  
  4.     C[i] = A[i]   B[i];  
  5. }  
  6.   
  7. int main() {  
  8.     …  
  9.     // Kernel invocation with N threads  
  10.     VecAdd<<<1, N>>>(A, B, C);  
  11.    …  
  12. }  

在上面的程式碼中，N個執行緒將會並行地同時執行加法運算。

3.2 執行緒層次（Thread Hierarchy）

      CUDA的每一個執行緒都有其執行緒ID，執行緒的ID資訊由變數threadIdx給出。threadIdx是CUDA C語言的內建變數，通常它用一個三維陣列來表示。使用三維陣列的方便之處在於可以很方便地表示一維、二維和三維執行緒索引，進而方便地表示一維、二維和三維執行緒塊（thread block）。這樣，無論是陣列、矩陣還是體積的計算，都可以很容易地使用CUDA進行運算。
      執行緒的索引與執行緒ID之間存在著直接的換算關係，對於一個索引為(x, y, z)的執行緒來說：
      1、如果執行緒塊（block）是一維的，則執行緒ID = x
      2、如果執行緒塊是二維的，假設塊尺寸為（Dx，Dy），那麼執行緒ID = x y * Dx
      3、如果執行緒塊是三維的，設其尺寸為（Dx，Dy，Dz），那麼執行緒ID = x y * Dx z * Dx * Dy
      下面的例子展示了兩個NxN矩陣相加的CUDA實現：
[cpp] view plain copy

  1. // Kernel definition  
  2. __global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) {  
  3.     int i = threadIdx.x;  
  4.     int j = threadIdx.y;  
  5.     C[i][j] = A[i][j]   B[i][j];  
  6. }  
  7.   
  8. int main() {  
  9.     …  
  10.     // Kernel invocation with one block of N * N * 1 threads  
  11.     int numBlocks = 1;  
  12.     dim3 threadsPerBlock(N, N);  
  13.     MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);  
  14.     …  
  15. }  

     每個執行緒塊（block）中的執行緒數量是有限制的，因為依據前面所說，同一執行緒塊（block）中的所有執行緒都會被分配到同一個處理器核上執行，共享有限的儲存資源，因此對於當前的GPU，執行緒塊所能包含的最大執行緒數目為1024。

      上面的例子中numBlocks代表執行緒塊的數量，這裡的值為1。在一般的CUDA程式中，這個值通常大於1，也就是說將會有多個執行緒塊被分配到多個處理器核中同時進行處理，這樣就大大提高了程式的並行性。
      在CUDA中，執行緒塊包含線上程格（grid）當中，執行緒格可以是一維、二維或者三維的，執行緒格的尺寸一般根據待處理資料的規模或者處理器的數量來指定。執行緒格中所包含的執行緒塊數目通常遠遠大於GPU處理器核心的數目。下圖展示了執行緒格（grid）、執行緒塊（block）以及執行緒（thread）之間的關係：
      核心函式的呼叫可以簡化為kernel<<<A，B>>>(parameters)，在尖括號中，A代表執行緒格（grid）的尺寸，它可以是三維的，用型別dim3表示，也可以是一維的，用int型別表示。B代表執行緒塊（block）的尺寸，它與A類似，也可分別用dim3或int型別表示。
      在核心函式內部，CUDA為我們內建了一些變數用於訪問執行緒格、執行緒塊的尺寸和索引等資訊，它們是：
      1. gridDim：代表執行緒格（grid）的尺寸，gridDim.x為x軸尺寸，gridDim.y、gridDim.z類似。拿上圖來說，它的gridDim.x = 3，gridDim.y = 2，gridDim.z = 1。
      2. blockIdx：代表執行緒塊（block）線上程格（grid）中的索引值，拿上圖來說，Block（1,1）的索引值為：blockIdx.x = 1，blockIdx.y = 1。
      3. blockDim：代表執行緒塊（block）的尺寸，blockDIm.x為x軸尺寸，其它依此類推。拿上圖來說，注意到Block（1,1）包含了4 * 3個執行緒，因此blockDim.x = 4, blockDim.y = 3。
      4. threadIdx：執行緒索引，前面章節已經詳細探討過了，這裡不再贅述。
      明白了這些變數的含義，那麼下面的矩陣加法程式便不難理解了：
[cpp] view plain copy

  1. // Kernel definition  
  2. __global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) {  
  3.     int i = blockIdx.x * blockDim.x   threadIdx.x;  
  4.     int j = blockIdx.y * blockDim.y   threadIdx.y;  
  5.     if (i < N && j < N)  
  6.         C[i][j] = A[i][j]   B[i][j];  
  7. }  
  8.   
  9. int main() {  
  10.     …  
  11.     // Kernel invocation  
  12.     dim3 threadsPerBlock(16, 16);  
  13.     dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);  
  14.     MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);  
  15.     …  
  16. }  

      在上面的程式中，執行緒塊（block）的尺寸是16×16，這是CUDA程式設計中一個非常普遍的選擇。執行緒格（grid）包含了足夠多的執行緒塊（block）來進行計算。

      執行緒塊（block）是獨立執行的，在執行的過程中執行緒塊之間互不干擾，因此它們的執行順序是隨機的。
      同一執行緒塊中的執行緒可以通過訪問共享記憶體（shared memory）或者通過同步函式__syncthreads()來協調合作。這些概念將在以後的章節中詳細解釋。