CUDAの概要

GPUを用いたプログラミングを行う際の選択肢の1つにCUDAがある。ここではCUDAの書き方をざっと説明する[1]。

CUDAとは

  • NVIDEA製のGPGPUを使って計算処理を行うための言語
  • C/C++をベースにGPUデバイスやメモリを扱えるように拡張
  • 多数のドキュメント・サンプルプログラム[2]、ユーザがいる

これからまとめることは基本的なC/C++の経験を前提とするが、GPUを使った経験や並列計算の経験がない人に対しても理解できるように書いて行くつもりである。

ホストとデバイス

まずはじめに、GPUプログラミングを行う上で必要な概念である、ホストとデバイスについて述べていく。といってもそんなに難しくない。

  • ホスト : CPU、またそのメモリ
  • デバイス: GPU、またそのメモリ

そして、CUDAのコードのうち、CPU側の処理について書かれている部分をホストコード、GPU側の処理について書かれている部分をデバイスコードと言う。

以下、典型的なベクトルの加算を表現したGPUのコードを記す。

上のコードを例に取ると、10-18行目がデバイスコード、対して20-49行目がホストコードである。ホストコードの部分はほぼC/C++と遜色ないが、30-34行目のように並列処理を呼ぶ部分や36, 37行目のようにデバイスからホストへ通信をする部分があるのが純粋なC++と異なる。

デバイスコード

カーネル

デバイス上でプログラムを走らせる関数のうち、ホスト側で呼び出されるものをカーネルと言う。先程の

 __global__ void addVec(float *vec1_dev, float *vec2_dev){
  unsigned int i_global = blockIdx.x * blockDim.x + threadIdx.x;             
  vec1_dev[i_global] += vec2_dev[i_global];
} 
...
int main(){
  ...
   //ホスト関数内でカーネルは呼ばれる
   addVec<<<NB,NT>>>(vec1_dev, vec2_dev);
  ...
}

はまさにカーネルだ。カーネルには

  • __global__ 修飾語をつけることでホスト関数と区別する
  • ホスト関数内で呼ばれる
  • デバイスメモリのポインタを操作する
  • ホストメモリにアクセス出来ない←重要
  • 返り値がない
  • スレッドが持つID等の情報に参照できる

といった特徴がある。

上のカーネルを例に取ると、引数の float *vec1_dev 等はデバイスメモリを前提にしている。また、 blockIdx.x 、 threadIdx.x が実際に動くスレッドの情報を教えてくれているのだ。

カーネル上で動く関数

複数のカーネル上で似たような処理を書くことになった場合、それを関数という形でまとめたくなる。ただし、カーネル内でホストコード内の関数を呼び出すことは出来ない。そこでデバイス上で呼び出されてデバイス上で動作する関数が必要で、CUDAはそれをちゃんとサポートしている。

引数を自乗する関数とその使用例を以下に記す。

__device__ float square_dev(float x){
  return x * x;
}
__global__ void kernel(...){
  ...
  x = square_dev(x);
  ...
}
  • __device__ 修飾語をつける
  • 「カーネル内でのみ」呼び出し可能

のような特徴がある。__device__ で修飾された関数は通常カーネル内でインライン展開される。

ホストコード

先程も述べたとおり、ホスト側のコードは基本的にC/C++と遜色ないが、GPUを扱う上で異なる部分がいくつかある。上の例のうち、C/C++と異なる部分を抜き出してきた。

 int main(){
  ...
  //デバイスメモリのポインタの宣言
  float *vec1_dev, *vec2_dev;
  ...
  //デバイスメモリの割当
  cudaMalloc((void**)vec1_dev, NB * NT * sizeof(float));
  ...
  //カーネルの呼び出し
  addVec<<<NB,NT>>>(vec1_dev, vec2_dev);
  ...
  //デバイスからホストへデータの転送
  cudaMemcpy(vec, vec1_dev, NB * NT * sizeof(float), cudaMemcpyDeviceToHost);
  ... 
  //デバイスメモリの解放
  cudaFree(vec1_dev);
  ...
}

まず、カーネルでデバイスメモリを扱うためには、メモリを確保する必要がある。そしてホストからカーネルを呼び出すことで実際にデバイス上のメモリを操作する。もちろん操作後のデータをホスト側から見ることは出来ないので、デバイスからホストへ転送する処理も書く必要がある(逆もしかり)。最後にデバイス側のメモリの解放も記載する必要がある。

References

[1] Cyril Zeller, “CUDA C/C++ Basics Supercomputing 2011 Tutorial”

[2] CUDA Toolkit Documentation

Follow me!

コメントを残す

メールアドレスが公開されることはありません。 * が付いている欄は必須項目です