CUDA メモ その0

cuda コーディングを学習する際のメモです。

Hello, world!

まずは、プログラム学習の定番である、”Hello, Wrold !”の出力を通して、GPUでの実行を練習します。

#include <cuda_runtime.h>
#include <stdio.h>


__global__ void print_hello(){
  printf("Hello, World!\n");
}


int main(void){

  print_hello<<<1,1,0>>>();

  cudaDeviceSynchronize();
  cudaDeviceReset();

}

上のコードは, Hello, world!を出力する関数をGPU上で実行するというものです。 print_helloという関数の前の__global__という接頭文字は、GPU上で実行する関数につけるものであり、その関数の戻り値はvoidである必要があります。 また、このようにGPU上で実行される関数をカーネル(kernel)と読んだりもします。

今後、CPU側をホスト, GPU側をデバイスと呼ぶことにします。 ホストからデバイス上で, print_helloという関数を実行するには、C言語系統の通常の関数呼び出しと比べて、関数名と引数の間に<<<1,1,0>>>を挟むことで呼び出すことできます(数字の意味は下で)。 先頭のcuda特有の関数を使用するには#include <cuda_runtime.h>を導入する必要があります。

cudaDeviceSynchronize()はデバイスとの同期をとる関数で、通常はデイバス関数の終了をまってからホスト側の操作が実行されますが、 デバイス関数がコードの末尾にある場合は、その終了を待つ必要があり、上の場合では呼び出す必要があります。

グリッド, ブロック, スレッド

cudaでは、デバイス関数を多数並列で実行する機構が備えられていて、ユーザーが自在に設定することが可能です。 Kernelはスレッドごとに実行される仕様になっています。 このスレッドの集合をブロック、さらにブロックの集まりをグリッドと呼びます(ここの図見ると想像しやすいです)。

ブロック内のスレッド数の設定は、以下のように行います。

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void get_thread_id(){
  int i = threadIdx.x;
  printf("thread Id:%d \n", i);
}


int main(void){

  dim3 block(32);
  get_thread_id<<<1,block,0>>>();

  cudaDeviceSynchronize();
  cudaDeviceReset();

}

上のコードはスレッドを32つ持つブロックを１つもつグリッドを使用して、カーネルを実行するコードです。 同様に、グリッドも複数作成して計算することも可能です。

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void get_thread_id(){
  int id_thread = threadIdx.x;
  int id_brock  = blockIdx.x;
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  printf("thread id:%d block id:%d i:%d \n", id_thread, id_brock, i);
}


int main(void){

  dim3 block(32);
  dim3 grid(8);
  get_thread_id<<<grid,block,0>>>();

  cudaDeviceSynchronize();
  cudaDeviceReset();

}

ここでは、32のスレッドを持つブロックを8つ生成するコードとなっています。 kernel内部では、threadIdx.xはスレッドがブロック内部でのID, blockIdx.xはグリッド内部でのIDを表します。 上の計算では、計256個のスレッドが立ち上がっているので、各々のIDは各ブロックが持っているスレッドの総数を示すblockDim.xを用いて、 コードにあるようにblockIdx.x*blockDim.x + threadIdx.x。 似た方法で、ブロック、グリッドともに２次元・３次元構造として設定することも可能です。