GPUのまとめ3 - doridoridoriand’s diary

いい加減コードを示していくよｗ

とりあえずコードを書きます先にCUDAのドライバーなどがインストールされていること前提で進めています。そのうちインストール方法とかまとめます(ほんとか？？)

#include <iostream>
#include <malloc.h>
#include <helper_cuda.h>

# define N           (8192 * 8192)
# define block_size  196608
# define thread_size 512

__global__ void add( int *a, int *b, int *c) {
  int thread_id = blockIdx.x * blockDim.x + threadIdx.x;
  while (thread_id < N) {
    c[thread_id] = a[thread_id] + b[thread_id];
    thread_id += gridDim.x;
  }
}

int main(int argc, char** argv) {
  int *a, *b, *c;
  ind *dev_a, *dev_b, *dev_c;

  a = (int*)malloc(N * sizeof(int));
  b = (int*)malloc(N * sizeof(int));
  c = (int*)malloc(N * sizeof(int));

  cudaMalloc((void**)&dev_a, N * sizeof(int));
  cudaMalloc((void**)&dev_b, N * sizeof(int));
  cudaMalloc((void**)&dev_c, N * sizeof(int));

  for (int i = 0; i < N; i++) {
    a[i] = i;
    b[i] = 2 * i;
  }

  cudaMemcpy(dev_a, a, N * siezeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(dev_b, b, N * siezeof(int), cudaMemcpyHostToDevice);

  add<<<block_size, thread_size>>>(dev_a, deb_b, dev_c);

  cudamemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost);

  cudaFree(dev_a);
  cudaFree(dev_b);
  cudaFree(dev_c);

  free(a);
  free(b);
  free(c);

  return 0;
}

上から説明していきます。CUDA独自の書き方とかまあ色々説明しなければならないことはあるのですが、まあ適宜最小限の説明で行きます。最初に色々言われても実際のコード見ないとわからないこと多いもんね

# define N           (8192 * 8192)
# define block_size  196608
# define thread_size 512

まずここの定数宣言なんだって感じですよね

# define N (8192 * 8192)

今回はCUDAの性能を純粋に測りたかったので単純なベクトル和を計算するのを書いています。よってこれは計算するベクトルの総和。つまり計算対象は 8192 * 8192 あるってことです

# define block_size  196608

使用するブロック数です。これテキトー(というか、さっきまで使ってたGPUに合わせただけ)なので、適宜変えて大丈夫です。これを変えて処理時間とか見みると結構面白い

define thread_size 512

使用するスレッド数です。これも変えておｋ

__global__ void add( int *a, int *b, int *c) {
  int thread_id = blockIdx.x * blockDim.x + threadIdx.x;
  while (thread_id < N) {
    c[thread_id] = a[thread_id] + b[thread_id];
    thread_id += gridDim.x;
  }
}

これが今回のメインの話題となるやつですね

__global__ってついたやつはコンパイル時にGPU用のコードとしてコンパイルされます。もっと正確にいうと、コンパイルはnvccというnVIDIAのコンパイラを使うのですが、コンパイルの流れとしては

コンパイル実行 nvcc -o hoge hoge.cu
nvccがCPUコードとGPUコードを判別。CPU用はgccに渡し、GPU用はnvccが担当する
最後CPU側がGPU バイトコードを呼ぶような感じでリンクして実行ファイル生成(めっちゃ適当)

ってなっていますGPU側はPTXコードというのが生成されたり色々するのですが、 ~~僕があんま詳しくないので~~もう少し詳しく取り扱いたいのでまた今度

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

これがGPUの特徴をよく表していると個人的に思うのですが、並列計算をするときに、SP同士でメモリの参照が狂ったりしないようにそれぞれの順番をuniqになるように定義できるのです。ちなみに使用するスレッドが１つだけの場合、つまり# define thread_size 1のときはthreadIdxはいらなくて良かったりします

以下はとくに変ではない、ふつうのCに登場してくるwhile文なので説明は省略します。
ではメイン関数をば

int *a, *b, *c;
ind *dev_a, *dev_b, *dev_c;

初期化はポインタで指定してね。でないと参照渡しとか出来ないので

cudaMalloc((void**)&dev_a, N * sizeof(int));
cudaMalloc((void**)&dev_b, N * sizeof(int));
cudaMalloc((void**)&dev_c, N * sizeof(int));

これでcuda側のメモリ確保を行います。

for (int i = 0; i < N; i++) {
  a[i] = i;
  b[i] = 2 * i;
}

ベクトル和の値の初期化をしてるだけです

cudaMemcpy(dev_a, a, N * siezeof(int), cudaMemcpyHostToDevice);

これでCPU側で初期化した配列群をGPU側のグローバルメモリにコピーします

add<<<block_size, thread_size>>>(dev_a, deb_b, dev_c);

ここで先ほどのGPUのコードを呼び出しています。 <<<>>>で囲まれた部分でこの関数で使用するブロックとスレッドの数を指定できます

cudamemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost);

ここで計算し終わったデータをGPU側からCPU側にコピーしています

cudaFree(dev_a);

みんなだいすきメモリ解放。Cのfree()のCUDA版です

すげーどんどん説明がテキトーになっているのがわかるかと思われますが、夏バテなので許してください

最後に__global__の関数、つまりCUDAで処理される関数はvoid型しか取れないのでみなさん注意してください

次回はAWSで実行環境を作る話をしたいです。僕はもうお家帰って寝たいです