整数演算をCUDAでやるサンプルを読む

x68kも使おうと思えばここまでできるってことがすばらしい。
　http://www.gizmodo.jp/2009/05/x680x054.html
ぺらぺらなフロッピーもどうかと思うのに、カセットテープかよっっ。


Macの一般誌を読んでいたらOpenCLのことが書いてあった。それってCPUかGPUか気にせずに使えるらしい。でも、明示的にGPUを使うかどうかは選べるんでしょうね。でも、SSEみたいに明示的な使い方をしないと上手い事使えない気はします。というか、基本的にCPUを使うようなやり方じゃないと、いらん処理までGPUでやってしまうんじゃないでしょうか。

実はOpenCLはNVIDIAであれ、ATIであれ、どっちでもいいように、インターフェイスは同じにするってことだと思っていたんだけど、どうもそれだけじゃないらしいけど、どうもそれ以上に読み取る事ができない。
　http://journal.mycom.co.jp/column/osx/301/index.html
CUDAに似ているメモリ管理をするらしいので、このままCUDAを信奉しても良いのかな。というか、ATI Streamを使う気になれん、分かりづらくて。少なくともCUDAほどにはお手軽にやれないと思う。



ええと、浮動小数点以外を扱うものとして、サンプルソースのCppIntegrationを見てみる事にしました。

関数が __global__ か __device__ を宣言をしていると、GPU上でやれることを書くので、その中でcutilXxx()とかの普通の関数は呼べなくなっている。そりゃ少しでもCPUで処理するようなものは排除した方が良いに決まっているわな。というわけで、GPUに載せる関数は部分的に限るのがよろしいかと。普通の関数呼び出そうとして、エラーが出まくって困った。

__global__の関数の直前（下だとhoge()）でタイマーを動かし、終わったら切るという単純な仕組みになっています。前回のdctのところから引っ張ってきました。

int timerCUDA = 0;
cutilCheckError(cutCreateTimer(&timerCUDA));
cutilCheckError(cutResetTimer(timerCUDA));
cutilCheckError(cutStartTimer(timerCUDA));

hoge <<>>(nantoka,kantoka);

cutilCheckError(cutStopTimer(timerCUDA));
int result_time = cutGetAverageTimerValue(timerCUDA);
printf("Processing time (gold)      : %f ms \n", result_time);
cutilCheckError(cutDeleteTimer(timerCUDA));

んで、結果は下記のようでkernelとkernel2って方がGPGPUの結果。もともとはgoldとgold2の出力はなかったので、出力を足した。goldの方はCPUでの処理。書かなかったけど、前回のも同様。

Processing time (kernel)    : 0.040535 ms
Processing time (kernel2)   : 0.030229 ms
Processing time (gold)      : 0.000638 ms
Processing time (gold2)     : 0.000692 ms
Test PASSED
Hello World.
Hello World.


Press ENTER to exit...

実行速度が二桁違いますね。何か根本的に間違った事をしている気もする。処理範囲の指定が違うかもしれないのだけれども、あまりに単純な整数演算をしようとするとGPGPUだとオーバーヘッドが大き過ぎてダメダメだっていう事は分かった。少なくともメモリの転送時間は入れてないので、これ以上に遅くなる可能性はない事もないけど、バスが太いからあんまり気にならないかもですね。


処理の流れを書いておきましょう。

処理概要	関数とか
コマンドライン設定	cutilDeviceInit()
GPU側のメモリの確保	cudaMalloc()
GPU側に処理データを転送	cudaMemcpy( , cudaMemcpyHostToDevice)
整数演算。
CPU側に処理データを転送	cudaMemcpy( ,cudaMemcpyDeviceToHost)
GPU側のメモリの解放	cudaFree()

考える事として、SSEとGPGPUの得意分野がかぶっているという事。CPUとGPUの違いはあれ、どっちもやっている事はSIMDでの処理だから、手分けの仕方が難しいかもなぁと思った。だからピンポイントな処理で一緒に同様な処理をするのが良いっぽい。でも大抵の場合、特別なライブラリを使わない限り、SSEはアセンブラで最適化していると思われるので、C言語に戻してからじゃないと使えない。なんか二度手間ですね。HandBreakでは大抵のライブラリがSSEに対応するために、要所要所でアセンブラになっていて、アセンブラを読むのはかなりめんどくさい。アセンブラにする前のソースを手に入れないとダメかも。

タグ：GPGPU