OpenCLの簡単なプログラムを書いてみて速度比較してみるなど

OpenCLの簡単なプログラムを書いてみてCPUとGPUの速度比較してみるなど。
基本的な感触をつかんでみた。

ソースコードなどはこちらに置きました。
https://github.com/dycoon/dycoon_simple/tree/master/OpenCL/SimpleTest

SimpleTest : 整数配列から奇数と偶数がいくつあるか数える
VMMultTest : 3次元ベクトルのアフィン変換の総和を返す
VertexBlendingTest : 3次元ベクトルを4つのアフィン変換で変換し均等にブレンディングしたものの総和を返す

それぞれの計算では、OpenCLによる処理と、C++で書かれたCPU上での処理の結果とかかった時間を返します。
処理内容はだいたい同じですが、OpenCLによる処理では並列処理を行うことを考慮した書き方になっているためその点で違いがあります。
CPUの処理ではシングルスレッドで処理しています。
計測は、
CPU : Core i7-2600 3.40GHz
GPU : NVIDIA Quadro 600
でおこなった。
OpenCLSDKNVIDIA GPU Computing SDK 4.0
Windows7で64bitの実行ファイルで、コンパイラはVC2008

プログラム GPUバッファ転送時間(ms) GPU計算時間(ms) GPU処理時間(GPUバッファ転送時間 + GPU計算時間)(ms) CPU処理時間(ms)
SimpleTest 177 127 304 84
VMMultTest 55 46 101 37
VertexBlendingTest 55 55 110 147

SimpleTestとVMMultTestではメモリーの量に対して処理が簡単すぎるためかメモリー転送を除いた
単純に計算している時間でみてもOpenCLによる計算のほうが遅くなっている。
VertexBlendingTestだと、メモリーの量に対して処理がそこそこ複雑なためか
OpenCLの処理のほうが勝っている。
もっとメモリーの量に対して処理が多ければどうなるのかというのは興味深いが、
現在は試していない。

OpenCLの処理を記述するに当たっては速度を上げるためにいろいろ試行錯誤した。

  • 一度に処理する要素数(呼び出すOpenCL関数内のループの回数でこれを書いたときは5000)
  • clEnqueueNDRangeKernelのlocal_work_sizeの値(これを書いたときは16)
  • __localを使用する(SimpleTestなどではかえって遅くなることも)

あと、引っかかったポイントとしては

  • clBuildProgramでCL_INVALID_BINARYを返してコンパイルが失敗する

以下のようにして最適化を無効にすることで動作するようにはなるようだ、

clBuildProgram(program, 0, NULL, "-cl-opt-disable", NULL, NULL);

ただ、これだと速度が下がってしまう場合があったので
変数を__localにしてみるなど、あてずっぽうに処理を書き換えて今回は対処した。
また、x64では起こらないがWin32では起こるという傾向があった。

  • 構造体の代入は要素ごとにおこなう

構造体を代入すると、コンパイルエラーが返るわけでもエラーコードが返るわけでもないが、
計算結果が正しくない。
こう書いていたのを

r[gid] = sumR;

以下のようにするなど

r[gid].x = sumR.x;
r[gid].y = sumR.y;
r[gid].z = sumR.z;

レイトレースの処理や当たり判定などの処理に使ってみたいという気がします。
それより前にいろいろやらねばならないことがあるので
私がそこら辺をやるのはかなり先の話になりそうですが。