思いつきメモ OpenCL
OpenCLの参考書籍
最近自分で技術的な文章を書く気力が減退気味。googleで"OpenCL 入門"とかで検索してみても、和書に関連するページがヒットするものの、OpenCLの入門的なページはほとんどヒットしない。なので、参考までに：OpenCL in Actionのリンク。ざっと眺めてみた感じでは、OpenCLの入門にふさわしい内容になっていると思う。データ型、カーネル記述方法、実例など、まんべんなく扱っている。ちなみに、リンク先サイトはすごい充実している。
気が向いたらもう少しOpenCLについて取り上げたいと思うけれど、いつになることやら。はっきり言って、ここまでCPUが早くなってしまった現在、映像処理やシミュレーション、データマイニングといった分野でしかGPGPUは必要ないのではなかろうか？
more…
OpenCL & GPGPUメモ
さらにOpenCLやGPGPUの実験をしているが、そのメモ。
- PCIeを介して、Device to Host, Host to Deviceの双方向転送が同時に実行されない。NVIDIAのGPUだと、Compute Capabilityが1.Xのものだと双方向の同時通信は無理のようなので、AMDもデバイスの制限があるのかも知れない。 これは、複数のコマンドキューを使用して、できるだけコマンドのオーバーラップを試みても、まだ未解決。 カーネルの実行中にデバイスメモリへの転送はできていることは確認できた。
- カーネルが使用するメモリを転送し終わってから、実際にカーネルが動作するまでに100ms以上のディレイが発生することがある。clEnqueueNDRangeKernelを呼び出して、直ちにclFlushを実行しているのだが...(このディレイはAMDの場合のみ。NVIDIAではそれほど遅くはない)
- OpenCLは、単にカーネルを呼び出すにも、カーネル引数をsetKernelArgで一つずつ積み上げなければいけない。しかも、引数のサイズまで指定する必要がある。さらに調査してみたら、GPGPUとしては、C++ AMPという方法もある(ちなみに、リンク先は機械翻訳なので日本語が少し変)らしい。しかも、Visual C++ 2012だと標準で実装されており、
をインクルードするだけで使用できる。Microsoftの実装だと、Direct X 11のシェーダとしてカーネルが実行されるようだ。だが、C++ AMPの仕様自体は ...
OpenCLメモ
OpenCLの実装メモ。
- カーネルのコンパイルエラーを見つける方法
カーネルの記述が間違っていると、clBuildProgramが失敗する。clGetProgramBuildInfoを使うと、コンパイルエラーを取得できるが、もっと便利なのは、AMD APP KernelAnalyzerを使う方法。Source Codeのところに問題のソースを貼り付け、Compileボタンを押すと、エラーがある場合はCompiler Outputのところに、エラーが出力される。
- プロファイリングの方法
Code XLを使用する。新しいプロジェクトを作成し、プロファイルを取得したいプログラムを指定する。ProfileメニューからStart Profilingを実行。GPU プロファイルにはApplication TraceとPerformance Countersの2種類がある。
これはApplicatoin Traceの実行結果：
プログラム実行中に使用したOpenCL関数や、それに掛かった時間などがグラフ表示される。
これがPerformance Countersの実行結果：
使用したカーネルや、カーネルを実行したときのworkgroup sizeもわかる。
- 雑感
やっと一通りカーネルを実行できるようになった。実行時間だけをみると、CPUの5倍程度高速であるように思われるが、前後のメモリ転送が遅いので(とりわけCPU->GPU)、結局CPUで計算した方が3倍程度高速な状態。ちなみに、CPUのプログラムはVisual C++ Express 2012で64bitバイナリでO2最適化/AVX有効の場合 ... more…