OS XのOpenCL - その26

OS XでのOpenCLの続き。今日で最後にする。もともとはBayer配列のカメラデータをRGBAに変換したかった。最後にOpenCLを使えばこんなに速く軽くなったぜ、みたいな結論になればカッコよかったんだけど残念ながらそうはならなかった。一応GPU版はできてCPUの負荷を減らすことは（ちょっとだけ）できたので、一応の目的は達したけど、ちっとも速くない。やっぱり地道な最適化しないとダメなんだけどそれはそれ非常に難しい、ということがわかった。

17.1  OpenCLのサンプルコード

AppleのデベロッパライブラリにはOpenCLのサンプルコードがいくつかある。しかし見てみるとどうもOpenCL規格のAPIだけを使ったものばかりに見える。これはどういうことだ。

具体的に言うと今日現在でOS X developer LibraryにあるSample CodeのうちタイトルにOpenCLの言葉が含まれているものをあげてみると

1. OpenCL N-Body Simulation
2. DeviceSelectCLGL
3. DeviceSelectCL
4. Offline Compilation Using the OpenCL Compiler
5. OpenCL_FFT
6. OpenCL RayTraced Quaternion Julia-Set Example
7. OpenCL Hello World Example
8. OpenCL Procedural Grass and Terrain Example
9. GPU Histogram
10. GPU Histogram
11. OpenCL Procedural Noise Example
12. OpenCL Parallel Reduction Example
13. Trajectories
14. OpenCL Procedural Geometric Displacement Example
15. OpenCL Parallel Prefix Sum (aka Scan) Example
16. OpenCL Matrix Transpose Example

これだけあったけど、このガイドに出てくるGCDのキューを使ったものはひとつもなかった。なんじゃこりゃ。しかも10.10のsdkのもとではコンパイルできないものがいくつもある（10.10が古いのではなくて、サンプルコードの方がリバイズされていない）。

こういう状態なのでAppleが独自の実装を使わせたいのか、もうやめたいのか全然わからない。ひどい状況である。

17.2  自前のサンプルコード

ということで、参考になるサンプルコードは存在しないということがわかった。まあ、FireWire（IEEE1394）のsdkでも似たようなこと（ガイドにしか参考にできるコードが存在しない）があったので驚かない。

実はここで僕が書いたサンプルコードをいくつかあげようと思っていた。しかし、CPUと比べてもちっとも速くない。最初の目標だったBayer配列の画像をRGBAに変換するのをOS X拡張OpenCLで書いたけど、CPUよりちょっと速いぐらいで、ガイドのガウシアンブラーのように一桁以上速くなるなんていうことはなかった。もちろんガイドがやってるような最適化はほとんどやってないので当然ではある。

もっと単純なモノクロ画像の対数圧縮はnVIDIAを積んだiMacではCPUよりもずっと遅くなった。対数圧縮は単なるピクセルごとのテーブルルックアップなので、簡単すぎてデータアクセスが律速してしまったということだろう。

Appleが使いやすいように簡単にしてくれてるけど、OpenCL、というかGPUの性能を十分に引き出すためには、ただGPUを使えばいい、というわけにはいかなくて、最適化のプロセス、特にメモリアクセスの最適化をまじめにやらないと無意味だということがわかった。

17.3  実行して気がついたこと

ところで書いていてひとつ、問題が発生した。個数やバイト長さを与える引数（例えばsizeXやiStrideなど）をsize_tで宣言すると実行時できない。Xcodeのデバグモードで見るとカーネルコンパイルが失敗してるらしい。

これはsize_tの長さがホスト側（OS Xではunsigned longにtypedefされている）とカーネル側（陽にはわからないがunsigned intらしい）で違っていることが原因らしい。カーネル内部での、例えばget_global_id()関数などがsize_tを返すようになっていてそれと合わせたつもりだった。

OpenCLの標準規格の中であれば整数の精度（ビット幅）はプログラマが意識して合わせることになるが、OS XではXcodeがかってにヘッダを書いてくれて、そこの宣言とあっていればホスト側は問題ないのでビット幅に無頓着になってしまう。

ヘッダの自動生成も良し悪し、ということである。しかしいずれにしても、大した問題ではないけど。

18  まとめ

OS X専用のOpenCLをざっくりみてみた。もとのOpenCLのコードに比べて簡単になる。具体的には

• cl_device_idやcl_contextなどお膳立てに必要なものはGCDのキューが隠蔽するので意識する必要がない
• カーネルを呼ぶときは普通の関数を呼ぶのと同じでいい
• カーネル引数も普通の関数に渡すのと同じでいい
• キューへの投入はdispatch_async、dispatch_syncなどGCDの普通の関数が使える
• キューへの投入はブロックが使えるので、その外の自動変数なんかがそのまま書ける
• カーネルが書き換えたデータへのアクセスもブロック内部で書けばいい（キューイングの必要はない）
• ARCをオンにしていれば、キューやブロックのメモリ管理も気にしなくていい

などなど、とにかく簡単に、簡単に、とできている。まさしく「キッシュイータ向け」と言っていい。

そのためおそらく微妙な調整はOpenCLの本来の規格に従った作業をしないといけない場合があるだろう。実際にそういう必要性を感じたことはないのでよくわからない。例えばcl_contextにエラー時のコールバックを渡したい、というような場合はどうすればいいのかよくわからない。

OS XのOpenCLは、GCDもブロックもオープンソースになっていて少なくともBSDにはポートされているので、「キッシュイータ用のOpenCL」として規格に追加してもいいような気もするんだけど、今の時点では単なるAppleの独自拡張の形になっている。

でも、Appleの独自拡張って、危ないんだよなあ。Javaなんて大宣伝した挙句にあっさり打ち切りだし、これからはガベージコレクションだretain releaseを書かなくてもいいんだぜ、と言った直後にObsoleteだ、使っちゃダメ、だもんなあ。サンプルコードが一つもない、というのも引っかかるよなあ。どうなることやら。

コメント 1

沒有醫生的處方
cialis diario compra http://cialisvonline.com/ Cialis purchasing
by Cialis generico online (2018-04-15 04:08)