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の仕様自体は、Direct Xに依存しているわけではない。少し試してみた範囲では、OpenCLより遙かにお手軽に使用できる(まだパフォーマンスは測定していない)。IntelもC++ AMPからOpenCLへのトランスレータをClang/LLVMで実装しているらしい。C++のラムダ式でカーネルへの引数をキャプチャできるので、わざわざsetKernelArgのようなことをする必要もない。利点/欠点は、ここにもまとめられている。Visual C++ 2012にはプロファイラがある、とのことだが、どうやらExpressには付属していない模様。GPUベンダ提供のプロファイラでいけるのか？
- C言語だと、どうしてもOpenCLのようなAPIになってしまうのも仕方ないと思うが、C++のように強力な言語機能をうまく使えば、ユーザにとって非常に見通しの良いGPGPU実装も可能である、という良い手本になっていると思う。