並列"Hello World"から始めるCUDA入門

はじめに

CUDAって、OpenMPなんかに比べると「とりあえず並列に動くコード」を書くのが難しい気がします。
私も何年か前に初めてCUDAを勉強したとき、ブロックとかスレッドとかストリームとかメモリのコピーとか、コード書く前に覚えること多すぎる...と苦しみました。コードが動くとこまでいかないと、モチベーション維持が大変ですよね。

そこで、とりあえず並列で動くHello Worldの書き方を紹介したいと思います！参考になれば幸いです。

並列処理させるための関数を作る

今回は"Hello World"を出力する関数を作り、それをCUDAで並列処理させるために書き換えていきます！
まず、C言語でベースとなるコードを書いていきましょう。

#include <stdio.h>

void hello(){
    printf("Hello World !!\n");
}

int main() {
    hello();
    return 0;
}

出来ました。hello関数を呼び出すと"Hello World !!\n"が表示されます。
このhello関数を並列処理させるには、関数の先頭に__global__を付けるだけでOKです！

#include <stdio.h>

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

int main() {
    hello();
    return 0;
}

これでhello関数はGPUで並列処理される関数に生まれ変わりました。このように先頭に__global__をつけた関数をカーネル関数と呼びます。関数の書き換えはこれだけです。

じゃあ早速コンパイル&実行してみましょう！

: error: a __global__ function call must be configured

なんかエラーが出ましたね。まあ当たり前ですね、並列数書いてないんですから。

並列数を記述する

並列数は関数を呼び出すときに付け加えてやればOKです！

#include <stdio.h>

__global__ void hello(){
    printf("Hello CUDA World !!\n");
}

int main() {
    hello<<< 2, 4 >>>();
    return 0;
}

ここで見慣れない記法が現れました。<<< 2, 4 >>>って何...?

これは、関数を呼び出す際に使用するブロック数とスレッド数を表しています。書き方は、<<< [ブロック数], [スレッド数] >>>です。CUDAではスレッドの上位階層にブロックが用意されていて、「何ブロック目の何番スレッド」と住所のように管理されています。このあたりの話は少しややこしくなるので詳しい説明はまた別途行うこととします。

この記事では、[ブロック数]×[スレッド数]だけ関数が並列処理される！とだけ覚えてください。

ブロック数とスレッド数を追記したコードでコンパイル＆実行してみます。

...何も表示されません。なぜでしょうか？

カーネル関数の同期処理

実は、カーネル関数は関数内の処理の終了を待たずにCPU側（関数の外）の処理が進んでしまうんです！今回の場合だとカーネル関数呼び出し後すぐにプログラムが終了するため、Hello Worldが出力される前にプロセスが終了してしまいました。

この問題は、カーネル関数を呼び出した後にcudaDeviceSynchronize();を記述してやれば解決します。この関数の呼び出しは、カーネル関数でGPUに渡した処理が全部終わるまでCPUは待っててくださいね。という意味になります。これを反映したコードが以下です。

#include <stdio.h>

__global__ void hello(){
    printf("Hello CUDA World !!\n");
}

int main() {
    hello<<< 2, 4 >>>();
    cudaDeviceSynchronize();
    return 0;
}

実行

今度こそ完成です。コンパイル＆実行してみましょう。

Hello CUDA World !!
Hello CUDA World !!
Hello CUDA World !!
Hello CUDA World !!
Hello CUDA World !!
Hello CUDA World !!
Hello CUDA World !!
Hello CUDA World !!

ちゃんと2×4=8並列されてますね。紆余曲折ありましたが意外とすんなり並列処理できたのではないでしょうか。

まとめ

今回はできるだけ難しいことは置いといて、とりあえず並列に動くコードを書くことに焦点を当てて解説してみました。さすがにこの知識だけだと実際にCUDAを活用するには不十分化かと思いますが、少しでも勉強のお手伝いができたなら幸いです。