前回は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の真の実力

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

__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億回になります。
これを実行させます。

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の特性

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

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 なるオブジェクトがいます。

__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しないと一致しないはず)

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にしてみる。

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/GPUPU/CUDA/CUDA2

A certain engineer "COMPLEX"

.NETでGPUPUを試してみる CUDA編第2回

Introduction

Explanation

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

CUDAの真の実力

global

device

host

GPUの特性

Conclusion

Source Code

A certain engineer "COMPLEX"

.NETでGPUPUを試してみる CUDA編 第2回

Introduction

Explanation

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

CUDAの真の実力

global

device

host

GPUの特性

Conclusion

Source Code

.NETでGPUPUを試してみる CUDA編第2回