簡單的 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 後,會在迴圈裡,把 dataAdataB 的每一項各自相加,並把結果記錄在 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 則是透過 cudaMemcpya, 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.cppVectorAdd.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

2 thoughts on “簡單的 CUDA 程式:VectorAdd”

  1. 錯詞:另外要注意的一點,”隻鱷”份程式為了簡單化,

發佈留言

發佈留言必須填寫的電子郵件地址不會公開。 必填欄位標示為 *