CUDAのGraphics Interop
CUDAでOpenGLと相互利用するときの話
CUDA C Programming Guideの3.2.11あたりをまとめただけな感じ。
初期化
まずcudaGLSetGLDevice()でデバイスの指定をする。
これはどのCUDA APIの呼び出しよりも先にやる必要あり。
cudaMemcpy()とかやってると実行時にエラー出ます。
登録
次に相互利用のためにCUDAにVBOを登録する。
glGenBuffers()で生成したVBOのIDをcudaGraphicsGLRegisterBuffer()で登録。
GLuint vbo; cudaGraphicsResource* cudaRes; int main() { // Explicitly set device cudaGLSetGLDevice(0); // Initialize OpenGL and GLUT ... // Create buffer object and register it with CUDA glGenBuffers(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(&cudaRes, vbo, cudaGraphicsMapFlagsWriteDiscard); glutMainLoop() }
cudaGraphicsGLRegisterBuffer()の第3引数はVBOのデータアクセスを指定する。
指定できるのは以下の3種類
- cudaGraphicsMapFlagsNone
- cudaGraphicsMapFlagsReadOnly
- cudaGraphicsMapFlagsWriteDiscard
NoneはRead/Write、WriteDiscardはWrite Only。ReadOnlyはそのまま。
マッピング
これでVBOをCUDAからいじる準備が完了。
次にcudaGraphicsMapResources()、cudaGraphicsResourceGetMappedPointer()でカーネル関数から操作できるポインタを得る。
操作後はcudaGraphicsUnmapResources()でマッピングを解除。
void display() { // Map buffer object for writing from CUDA float4* positions; cudaGraphicsMapResources(1, &cudaRes, 0); size_t num_bytes; cudaGraphicsResourceGetMappedPointer((void**)&positions, &num_bytes, cudaRes)); // Execute kernel dim3 dimBlock(16, 16, 1); dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1); createVertices<<<dimGrid, dimBlock>>>(positions, time, width, height); // Unmap buffer object cudaGraphicsUnmapResources(1, &cudaRes, 0); ... }
リソースの解放
アプリの終了時など適当なときにリソースの開放しましょう。
void deleteVBO() { cudaGraphicsUnregisterResource(cudaRes); glDeleteBuffers(1, &vbo); }