thrustというライブラリがある。
直接使ったことがなくても、CUDAをインストールしたら付いてくるので知らずにインストールしている人は多いと思う。
これはCUDAをC++で使う上で最高レベルに便利なライブラリで、もはやこれ無しでCUDAプログラミングをするのは考えられないくらいだ。CUDA以外にもintel TBB、OpenMP、C++標準のthreadサポートのバックエンドが入っており、たまにこれ無しでOpenMPのコードを書くと少しだるく感じてしまう。
上に貼った公式サイトでは最新リリースが1.8.1 (2015)になっていて、開発が止まっているのかな？ としばらく思っていたが、GitHubのレポジトリのリリースを見たらちゃんと最近も活発に開発されている。単に公式サイトの更新をサボっているだけっぽい。
見ていて驚いたのが、最新の一つ前のリリースで、CUDA 10.1に同梱されているもののリリースノートだ。タイトル通りの変更が入っている。
Release Thrust 1.9.4 (CUDA 10.1) · thrust/thrust · GitHub
背景を少し説明しておこう。
まず、GPUはCPUから独立したアクセラレータなので、GPU上での計算をしている間CPUが待つ必要は多くの場合ない。なので、CPUはGPUに命令を発行したのち、終わるのを待たずに次の処理を始めることができる。もちろんデータの依存関係などはこのとき注意を払う必要がある。
これまで、Thrustの処理は一部だけがasyncだった。一部のthrust関数は処理が終わるまでCPUをブロックするが、一部は非同期に実行することで呼ばれた後すぐにCPUを自由にしていた。だがどれが非同期だったかを常に意識するのはしんどい作業で、データの依存関係などを考える際に少しばかりの不安があった。
このリリースで、thrust::future<T>とthrust::eventが実装され、thrust::async名前空間に複数のアルゴリズムが実装された。future<T>は簡単に言うと「まだ処理が終わっていないかも知れないが、終わるとT型の結果が格納されるハンドル」で、処理が終わっているかどうかのフラグと、いざ必要になった時に計算が終わるまで待って中身を取り出すという操作をさせてくれる。thrust::async::*以下にある関数を呼び出した場合このようなハンドルが返ってきて、即座にCPUは自由になる。その後はfutureを持って回ることで非同期に処理を進めることができ、いざ同期するとなったときにそのハンドル経由で計算が終わるのを待てばよいということになる。
これに伴って、今まで非同期に動いていたthrustのアルゴリズムが必ず同期的に動くようになった。この機能によって、当然、既存のコードは若干の速度低下を被るかも知れない旨が注意喚起されている。だが私としては今まで通りのコードが確実に同期されるという保証と、非同期的に動かしたいときに明示的に指定できるようになったことのメリットが圧倒的に上回っているように感じる。不安が払拭され、実際にバグが減るだろう。
もうひとつ大きいと思った変更は、thrust::optionalがサポートされたことだ。GPUでoptionalが使える！ 代数的データ型の有用性は今更語る必要も無いだろう。オイオイ最高だな。
あと、std::pmrライクなメモリリソースもサポートされている。CPUとGPUのどちらからでも触れるunified memoryを扱うuniversal_memory_resourceなどもサポートされており、かなり凄まじい機能追加に思える。C++最新規格にかなり追いついているので、数年前と比べるとCUDAプログラミングはかなりの速度で簡単になってきていると言っていいだろう。ちょっと前はCUDAのインストールはXを落としてからしたほうが良いよとか言われてたのに、隔世の感がある。
というわけでもっとCUDAプログラミングをしましょう。
