メモ:OpenGL Interoperability の話
長いので Interop にして話す。OpenGLとの相互運用のこと。Direct3Dとも可能だが、Direct3Dはいかんせん初期化が面倒なので使いたくないので無視。
CUDA C Programming Guideの3.2.11辺りにGraphics Interoperabilityのページがある*1のでそれと、SDKのサンプルを見ると分かりやすい。後CUDA by Exampleとか。大学図書館に英語の原書があったので、借りて英語がよく分からないながらも読んでます。
正直、SDKのサンプルってDirectXとかもそうだけどなんかコード自体に読ませる気がない長さを持っている気がする。どうしてあんなに長いのか。サンプル名に simple などかかれていても全くシンプルに見えないので困る。
まあ、ちょっと長いほうがコードを読む甲斐があるような気もしますが、そう思わないとやってられない。
で、OpenGLの初期化とか描画周り差し引けばそこまで難しくない。むしろ描画周りのパラメータの設定とかのほうがしんどいと思う。
まず、CUDAの関数やOpenGLの初期化よりも初めに、cudaGLSetGLDeviceで使用するデバイスを決定する。
cudaGLSetGLDevice(0);
CUDA by Exampleを読むと、compute capabilityを指定して対応するデバイスのIDを取得したりしていたが、まあとりあえずそこまでする必要はないと思うので省略。0 で多分デフォルトのデバイスになるはず。
その後、VBO(Vertex Buffer Object)を作って、それをカーネルで使えるように登録する。
int width = ...; // VBO のサイズ int height = ...; GLuint vbo; struct cudaGraphicsResource* vbo_cuda; glGenBuffer(1, &vbo); glBindBuffer(GL_ARRAY_BUFFER, &vbo); unsigned int size = width * height * 4 * sizeof(float); glBufferData(GL_ARRAY_BUFFER, size, 0, GL_DYNAMIC_DRAW); glBindBuffer(GL_ARRAY_BUFFER, 0); // カーネルで使えるように登録 cudaGraphicsGLRegisterBuffer( &vbo_cuda, vbo , cudaGraphicsMapFlagsWriteDiscard );
で、カーネルへ渡すときはリソースのビューを作成する。
float4* v; cudaGraphicsMapResources(1, &vbo_cuda, 0); size_t bytes; cudaGraphicsResourceGetMappedPointer((void**)&v, &bytes, vbo_cuda);
この v がVBOへの先頭のポインタとなるので、これに対してカーネルで頂点情報を書き込めば良い。
書き込んだ後はUnmapしないと描画にも使えないので必ず行う。
kernel<<< ... , ... >>>(v, width, height, ...); cudaGraphicsUnmapResources(1, &vbo_cuda, 0);
後は、VBOを使って正しく描画するだけ。アプリケーション終了時に必ずリソースを破棄するようにstd::atexitを使って終了関数を登録しておく。
void cleanup() { // 登録の解除 cudaGraphicsUnregisterResource(vbo_cuda); glBindBuffer(GL_ARRAY_BUFFER, vbo); glDeleteBuffers(1, &vbo); } std::atexit(cleanup);
という感じでいいらしい。Interop よりも綺麗に描画しようとするとシェーダとか使わないといけなくなるんでしょうねえ。simpleGLしか読んでなかったけど、oceanFFTとか読んでみますか。
*1:version 4.0での話