CUDA4.1 RC

いつのまにやらCUDA 4.1のRC版が出てます。
http://developer.nvidia.com/cuda-toolkit-41
http://developer.nvidia.com/content/new-cuda-now-available

さらっと見た感じ以下の3点推し。

  • 新しいLLVMベースのコンパイラによるパフォーマンスのアップ(最大10%)
  • NPPへの1000以上の関数の追加
  • Visual Profilerの強化

新しいコンパイラはコンパイル速度も早くなるみたい。NPPは使ってないし知りません。

Visual Profilerはボトルネックっぽい部分を解析してパフォーマンスアップのためのガイドを表示してくれそうな感じになってる。(ビデオのスライドだとMemcpyのスループットが低いからPinned Memory使えみたいな)

Parallel Nsightもバージョンアップ http://www.nvidia.com/object/parallel-nsight.html
Warp Watch機能とか追加されてる。シングルGPU環境だと結局Profilerしか使えないんだけど・・・

Visual ProfilerとかのスクリーンショットIntroduction to CUDA 4.1 で見れます。

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);
}

サンプル

CUDA SDKのサンプルとしてはSimple OpenGLというサンプル*1が名前の通りシンプルで理解しやすい。

はてなブログはじめました。

深夜3時のこと暇を持て余して辟易としていら id:r7kamura から、ブログ書け捗るぞと言われたのでブログ書きます。

意識高いと毎日ネタ探さなくてもブログ書けちゃうらしいので、その辺の意識高めていきたい。


技術系に縛ると書けなくなりそうだしなんか書けばいいかな程度で。betaとか左上に書いてるし。