AkiiNote CUDA 筆記

CUDA

CUDA的寫法大概就像下面那樣 幾乎每個CUDA都很類似：

開一個足夠大的空間然後給初始值 -> 複製記憶體到GPU -> 呼叫kernel function執行GPU -> 把最後的結果複製回來

在寫CUDA的程式 最需要注意的部分：

l   GPU的記憶體不大(1G、2G……)

在資料很龐大的時候

就必須要想怎麼樣才能把資料縮小

所以在資料的處理上要費不少心思

如果需要用到二維的話 建議使用一維的方式來表達y*N+x

(個人認為在呼叫一維的kernel比較簡單

廢話講完了
以下是矩陣相加的例子：
C = A + B

在global 的地方宣告變數
初始值先給NULL 待會再用malloc 動態宣告記憶體

int* h_A = NULL;    // host

int* h_B = NULL;

int* h_C = NULL;

int* d_A = NULL;    // device

int* d_B = NULL;

int* d_C = NULL;

int N = 10; 

一開始讀取指令的值
N 是矩陣的邊長
size 是先計算需要開多大的空間來存放這些值

  assert(argc == 2); 

  N = atoi(argv[1]); 

  int size = N*N*sizeof(int);

  printf("N:%d, size:%d\n", N, size);

   

然後用malloc 開記憶體的空間
以下是host 的A B C

  h_A = (int*)malloc(size); 

  h_B = (int*)malloc(size); 

  h_C = (int*)malloc(size);

  assert(h_A); 

  assert(h_B); 

  assert(h_C); 

 

  // initial array A & B
在來就是初始h_A 和h_B 的值

 在來是宣告cuda 記憶體的空間，因為要用GPU作計算前，必須先把記憶體上的h_A跟h_B複製到GPU上的記憶體

  cudaMalloc((void**)&d_A, size);

  cudaMalloc((void**)&d_B, size);

  cudaMalloc((void**)&d_C, size);

  printf("d_A: %p\n", d_A);

  printf("d_B: %p\n", d_B);

  printf("d_C: %p\n", d_C);

然後
把h_A複製到GPU上的d_A
把h_B複製到GPU上的d_B

  cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);

  cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

接下來就可以執行囉
這邊就是設定
每個Grid 有幾個Blocks
每個Block 有幾個threads
一個thread 就能執行矩陣上的一個點

  int threadsPerBlock = 256;

  int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

然後在呼叫Kernel function 就可以執行囉
後面括弧的地方是參數值(這個應該沒問題啦)

  kernel<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

 最後由GPU執行完，在把加過的值複製回到local memory

  cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

以下這個部分沒有算好的話 可是非常容易爆掉的

threadsPerBlock：每個block有幾個threads。這個就要看你到底是什麼的GPU了，如果會爆掉的話就改小一點吧。

blocksPerGrid：每個grid有幾個block。如果原本需要執行N次，那現在就只需要執行(N + threadsPerBlock - 1) / threadsPerBlock次。

在GPU 裡面

kernel在執行的時候盡量讓資料獨立

而且在GPU做運算得時候
呼叫kernel function上面的參數都是屬於GPU global值
像是以下程式碼的A B 或 C 
(N 也算啦 可是值有一個值 比較不會影響)

所以要加快運算的話可以先把
global 的資料複製到local來執行
執行完後再存回global 減少跟global存取的次數 
比較複雜的kernel function 能加快不少的時間
(矩陣相加不明顯)

__global__ void kernel(int* A, int* B, int* C, int N) {

  int n = blockIdx.x * blockDim.x + threadIdx.x;

// 這邊使用if 是用來判斷n 是不是在矩陣的範圍內

  if ( n < N){

    int local_a = A[n];
    int local_b = A[n];
    // exec……

    C[n] = local_a + local_b;

 }

}

最後執行完後
別忘了把記憶體給釋放掉喔

cudaFree(d_A);

cudaFree(d_B);

cudaFree(d_C);

free(h_A);

free(h_B);

free(h_C);