前回はCUDAの性能がとんでもなく悪かったです。

Introduction

明らかに悪い、とは思っていたけどここまで悪いとは思っていませんでした。
いや本当。

今回は、CUDAの仕組みから。

Explanation

CUDAってどうやって動いている?

今さらですが。
CUDAはGPUを使って計算します。
が、ここで理解してほしいのは、そういう計算処理は本来、CPUがやることで、GPUに計算させているのはあくまでおまけ。
で、当然ながら、現在のPCの仕組み上、GPUとCPUには扱いに大きな差があります。

まず、CPUとメモリの間はレジスタとキャッシュ (1次とか2次とか)が備わっています。
これによりデータやプログラムが可能な限り再利用され、データの移動が最小限になります。

が、GPUとCPU側のメモリの間にそんなものはありません。
たしかに、GPUにはメモリが乗っていますけど、そこにプログラムが使っているデータはありません。
そもそもメモリ (要するにDRAM) はCPUの側に比べれば格段に遅い。

なので、プログラムからGPUに処理を任せるときは、データとプログラムを渡す必要があります。
これ大事。

で、先に話したように、GPUとプログラムの距離は非常に遠い、つまり遅いです。
そもそもDRAMの演算装置間の転送速度は遅いのに、CPU側のメモリとGPUではさらに遅い。
以上を踏まえると、命令やデータをCPU側メモリからGPUに何度も転送するのは最悪、ということです。

で、ここまで書いてあれですが、きちんと用語があって、CPU側のメモリ、とかいう表現はしません。

Term Meaning
ホスト 呼び出し側。CPU・プログラムと思ってくれれば
デバイス GPU側
カーネル デバイス上で実行される処理。ホスト上のソースコードに書かれたデバイス用の処理

と定義されています。
もう一度言いますと、ホストとデバイス間のカーネルやデータの転送は最小限にする、ってことです。

CUDAの真の実力

前段を踏まえて前回のプログラムの問題点を検証します。
まずデバイス側のカーネルは

1
2
3
4
5
__global__ void addKernel(int *c, const int *a, const int *b)
{
int i = threadIdx.x;
c[i] = a[i] + b[i];
}

になります。
CUDAプログラミングにおいて、関数がホスト側、デバイス側で実行されるかどうかの識別は下記の3つの関数修飾子で区別します。

global

デバイス側で実行さ
れる、ホスト側から呼び出される関数であることを示します。戻り値は必ずvoidです。

device

デバイス側で実行され、デバイス側から呼び出される関数であることを示します。要するに、デバイス側のprivateな関数。

host

ホスト側で実行され、ホスト側から呼び出される関数であることを示します。要するに、ホスト側のprivateな関数。

それで、上の addKernel 関数 は、globalなので、デバイス側で実行され、ホスト側でコールされます。
この関数は、別に定義された cudaError_t addWithCuda(int *, const int *, const int *, unsigned int) 関数 で呼ばれます。
前のコードでは、この addWithCuda 関数を10000回呼び出していました。すなわち10000回、カーネルとデータをデバイス間に転送していたことになります。
それは遅いに決まっています。
理想的なのは、デバイス側で10000回処理が実行されることです。


  • × デバイス側に処理を10000回依頼する
  • ○ デバイス側で処理を10000回実行する

ということです。
以上を踏まえコードを修正します。
サンプルコードはページの末尾を参照。
変更点は

  • ループ回数を10000から100000000回に変更
  • addKernel 関数内でループを回すように変更
  • デバイス側の計測は addWithCuda 関数内で実行するよう変更
  • 変数a,b,cはconstを外して、可変長に変更

これだと、ループ回数は配列長*ループ回数なので、5億回になります。
これを実行させます。

1
2
3
4
5
6
CUDA is
{1,2,3,4,5} + {10,20,30,40,50} = {11,22,33,44,55}
time = {5556}
No CUDA is
{1,2,3,4,5} + {10,20,30,40,50} = {11,22,33,44,55}
time = {0}

一気に改善しています。ですが、まだダメです。
というか、CPUが速すぎる気がします。

そこで、デバイス側をコメントアウトして、ホスト側だけの計測をしました。
その結果、ループ回数10億、配列長50で8000-9000msという結果が出ました。
少なくとも、計測処理は正しいことがわかりました。常に0になるとか、そういうのはない。
なので、ループ回数を1億にして、配列長50なら800-900msになるはず。
ループ回数が変化なし、配列長が10倍なら、CUDAはおよそ55000msになるはず。

GPUの特性

はい。
結果はそうなりませんでした。

1
2
3
4
5
6
CUDA is
{1,2,3,4,5} + {10,20,30,40,50} = {11,22,33,44,55}
time = {5555}
No CUDA is
{1,2,3,4,5} + {10,20,30,40,50} = {11,22,33,44,55}
time = {856}

CPUの計算量推定は正しかったのですが、GPUはそうなりませんでした。
結論を言えば、並列化が関係してます。
addKernel 関数 をみると、threadIdx なるオブジェクトがいます。

1
2
3
4
5
6
7
8
__global__ void addKernel(int *c, const int *a, const int *b, const int loop)
{
int i = threadIdx.x;
for (unsigned int t = 0; t < loop; ++t)
{
c[i] = a[i] + b[i];
}
}

これなんでしょう。
これ、NVIDIA GPU Computing Toolkit\CUDA\v7.5\include\device_launch_parameters.h に定義されています。
字面からみればスレッドIDです。

CUDA_C_Programming_Guide.pdf がありますのでそれを見ます。

B.4.4. threadIdx This variable is of type uint3 (see char, short, int, long, longlong, float, double ) and contains the thread index within the block.

1
訳: B.4.4. threadIdx この変数はuint3 (char, short, int, long, longlong, float, double) 型であり、ブロック内のスレッドインデックスです。

とあります。
字面の通りということでしょう。

だとするならば、CUDAはスレッド処理を実行しているはずです。ゆえに、CPUのような計算量推測ができなかったのでしょう。
では、どこでスレッドの起動を指定しているのか、というと

1
2
// Launch a kernel on the GPU with one thread for each element.
addKernel << <1, size >> >(dev_c, dev_a, dev_b, g_loop);

<1, size > とのこと。
これは、1ブロックにつき最大sizeスレッドで関数を実行せよ、という意味らしい。
第6回 CUDAプログラミングモデル①によれば、指定できるスレッド数は512が限界とのこと。

なので、現在sizeは配列長さと等しい。
つまり、配列の 1 要素に対して、1スレッドで動作するため、配列長さを10倍にしても、処理時間が変化しなかったのである。
ともすれば、GPUはCPUのスピードにはかなわないが、大量のスレッドでその分作業をこなすことで、高速化を図るしか道がないことになる。

なので、CPUは単純に配列長に比例し、GPUは配列長の影響を受けないので、配列帳が7-8倍になれば、ホストとデバイスの処理結果が近似または逆転するはずである。
試しに配列長を400にしてみた。
(倍数+1しないと一致しないはず)

1
2
3
4
5
6
CUDA is
{1,2,3,4,5} + {10,20,30,40,50} = {11,22,33,44,55}
time = {5778}
No CUDA is
{1,2,3,4,5} + {10,20,30,40,50} = {11,22,33,44,55}
time = {5295}

思い切って512にしてみる。

1
2
3
4
5
6
CUDA is
{1,2,3,4,5} + {10,20,30,40,50} = {11,22,33,44,55}
time = {7489}
No CUDA is
{1,2,3,4,5} + {10,20,30,40,50} = {11,22,33,44,55}
time = {6686}

追いつかない…

Conclusion

なんで、CUDAがとてつもなく遅かったのかは原因がつかめました。
が、期待した速度が出ていないこともまた事実。
次回は、もう少し原因を探っていきたいです。

Source Code

https://github.com/takuya-takeuchi/Demo/tree/master/CUDA/CUDA2