簡單的 CUDA 程式：VectorAdd

在上一篇的《nVidia CUDA API》中，已經大概介紹了一些寫 CUDA 程式時，必須知道的 API；接下來，就先開始簡單的 CUDA 程式吧！

這邊舉一個簡單的向量相加的程式來當例子。首先，先定義測試用的資料：

// initial data 
int   data_size = 100; 
float *dataA = new float[data_size],
       *dataB = new float[data_size],
       *dataC = new float[data_size];

 
for( int i = 0; i < data_size;    i )
{
     dataA[i] = i;     dataB[i] = -1 * i; 
}  

上面的程式，會宣告 dataA, dataB 兩個 float 的一維 array，同時內部的值都差一個負號；而 dataC 的部分，則是預備用來存放 dataA, dataB 相加後的結果。

一般版本

在向量加法的部分，如果以一般 C/C 的寫法，可以寫成：

void add_vector_cpu( float* a, float* b, float *c, int size )
{
     for( int i = 0; i < size;    i )
         c[i] = a[i]   b[i]; 
}

而 main 的部分，則只需要呼叫函式就可以了：

add_vector_cpu( dataA, dataB, dataC, data_size );

在呼叫 add_vector_cpu 後，會在迴圈裡，把 dataA 和 dataB 的每一項各自相加，並把結果記錄在 dataC 中。由於 dataA, dataB 的值都差一個負號，所以 dataC 會是一個所有值都是 0 的 array。

CUDA 版

而如果要寫成 CUDA 的，向量相加的函式會變成：

__global__
void VectorAdd( float* arrayA, float* arrayB, float* output )
{
     int	idx = threadIdx.x;
     output[idx] = arrayA[idx]   arrayB[idx]; 
}
 
void add_vector_gpu( float* a, float* b, float *c, int size )
{
     int	data_size = size * sizeof(float);
 
     // part1, allocate data on device
     float	*dev_A,	*dev_B,	*dev_C;
     cudaMalloc( (void**)&dev_A, data_size );
     cudaMalloc( (void**)&dev_B, data_size );
     cudaMalloc( (void**)&dev_C, data_size );
 
      // part2, copy memory to device
     cudaMemcpy( dev_A, a, data_size, cudaMemcpyHostToDevice );
     cudaMemcpy( dev_B, b, data_size, cudaMemcpyHostToDevice );
 
      // part3, run kernel
     VectorAdd<<< 1, size >>>( dev_A, dev_B, dev_C );
 
      // part4, copy data from device
     cudaMemcpy( c, dev_C, data_size, cudaMemcpyDeviceToHost );
 
      // part5, release data
     cudaFree(dev_A);
     cudaFree(dev_B);
     cudaFree(dev_C);
} 

而在 main function 的呼叫，和 CPU 的版本一樣，只要寫：

add_vector_gpu( dataA, dataB, dataC, data_size ); 

就可以了∼

從上面的程式可以發現，實際在做加法的函式是「VectorAdd」這個函式；在寫的時候，他前面有加上「__global__」來宣告成 CUDA 的 kernel function。而實際用來呼叫的「add_vector_gpu」這個函式，則算是用來把 CUDA 的相關程式封包起來；在裡面除了要讓 device 執行「VectorAdd」這個 GPU 的 kernel 程式外，還要把記憶體由 host（系統記憶體）複製到 device（顯示卡的記憶體）上，也就是 part1, part2 的部分。此外，在「VectorAdd」執行完後，還要再把結果由 device 複製回 host（part4），並且把 device 上的資源釋放掉。所以其實和 CPU 程式比起來，還是稍微繁瑣了些。

下面大概講一下整個流程：

1. part1 的部分，是用來宣告並透過 cudaMalloc 指派記憶體空間在 device 上
2. part2 則是透過 cudaMemcpy 把 a, b 的資料從 host 複製到 device
3. part3 執行 kernel 程式：VectorAdd。而「<<< 1, size >>>」的部分，就是指定用 1*1 的 grid，裡面則是 size*1*1 的 thread block；也就是說，總共會有一個 grid，grid 裡只有一個 block，block 裡有 size 個 thread。
   而執行的結果，就相當於 VectorAdd 會被呼叫 size 次，每次的 threadIdx.x 都會是不同的值。
4. part4 則是透過 cudaMemcpy 把計算完的 dev_C 的資料從 device 複製回 host
5. part5 會將 dev_A, dev_B, dev_C 這些已經用不到的記憶體空間，透過 cudaFree 來釋放掉。

另外要注意的一點，隻鱷份程式為了簡單化，在kernel 程式的設定，採用了把所有 thread 放在同一個 block 的設定；在這個情況下，如果 size 比 device 中每個 block 的最大 thread 數還多的話，會沒辦法正確執行。這個時候，就須要把把它拆成幾個 thread block 來執行了∼

這份程式，應該就算是 Heresy 對於 CUDA 的 Hellow World 吧∼原始碼的話，有放一份在 Microsoft 的 SkyDrive 上（下載：VectorAdd.zip），有興趣的人可以抓下來試試看。

而由於 Heresy 個人是希望可以把 CUDA 的程式和一般的 C/C 程式分開來，所以在 Heresy 的測試專案裡，是分成了 VectorAdd.cpp 和 VectorAdd.cu 兩個檔案。main 是寫在 VectorAdd.cpp 中，而為了要讓 main 裡能夠呼叫 VectorAdd.cu 裡的 add_vector_gpu，所以必須要在兩個檔案都加入

extern "C" void add_vector_gpu( float* a, float* b, float *c, int size );

把 add_vector_gpu 這個函式特別標註成是用 C 的方式來編譯。

在這份程式裡，CPU 的程式和 CUDA 的程式都有；在最後的部分，有驗證兩者的結果，理論上應該會是都一樣的∼

原始發表： 簡單的 CUDA 程式：VectorAdd