メモ: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での話