CUDAカーネルを使用してOpenGLテクスチャを変更しようとしsurf2Dwrite()ていますが、下の画像に示すように、呼び出しがテクスチャの以前のコンテンツと混ざり合っているように見えるという奇妙な問題が発生しています。背面の木製のテクスチャは、CUDAカーネルで変更する前のテクスチャに含まれています。期待される出力には、色のグラデーションのみが含まれ、背後にある木のテクスチャは含まれません。なぜこのブレンドが起こっているのかわかりません。

考えられる問題/誤解

私はCUDAとOpenGLの両方に不慣れです。ここでは、私をこのコードに導いた思考プロセスを説明しようと思います。

• cudaArrayテクスチャの読み取り/書き込み時のキャッシュの局所性に適していると読んだため、（floatの配列などではなく）テクスチャにアクセスするためにを使用しています。
• サーフェスを使用しているのは、それが変更する唯一の方法であるとどこかで読んだからです。 cudaArray
• サーフェスオブジェクトを使用したかったのですが、これは新しい方法であると理解しています。古い方法は、サーフェス参照を使用することです。

チェック/テストの方法がわからない、コードに関するいくつかの考えられる問題：

• 画像形式に一貫性がありませんか？たぶん私はどこかで正しいビット/チャネル数を指定しなかったのですか？多分私はfloatsの代わりにunsigned charsを使うべきですか？

コードの概要

このGitHubGistで、完全に最小限の作業例を見つけることができます。可動部分があるのでかなり長いですが、まとめてみます。MWEを短縮する方法についての提案を歓迎します。全体の構造は次のとおりです。

1. ローカルに保存されたファイルからOpenGLテクスチャを作成する
2. を使用してテクスチャをCUDAに登録します cudaGraphicsGLRegisterImage()
3. テクスチャを表すcudaGraphicsSubResourceGetMappedArray()を取得するための呼び出しcudaArray
4. 作成しcudaSurfaceObject_t、私はへの書き込みに使用することができますcudaArray
5. でテクスチャに書き込むカーネルにサーフェスオブジェクトを渡します。 surf2Dwrite()
6. テクスチャを使用して画面上に長方形を描画します

OpenGLテクスチャの作成

私はOpenGLを初めて使用するため、LearnOpenGLチュートリアルの「テクスチャ」セクションを出発点として使用しています。これが私がテクスチャを設定する方法です（画像ライブラリを使用してstb_image.h）

GLuint initTexturesGL(){
    // load texture from file
    int numChannels;
    unsigned char *data = stbi_load("img/container.jpg", &g_imageWidth, &g_imageHeight, &numChannels, 4);
    if(!data){
        std::cerr << "Error:  Failed to load texture image!" << std::endl;
        exit(1);
    }

    // opengl texture
    GLuint textureId;
    glGenTextures(1, &textureId);
    glBindTexture(GL_TEXTURE_2D, textureId);

    // wrapping
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_MIRRORED_REPEAT);
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_MIRRORED_REPEAT);

    // filtering
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR_MIPMAP_LINEAR);
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR);

    // set texture image
    glTexImage2D(
        GL_TEXTURE_2D,    // target
        0,                // mipmap level
        GL_RGBA8,         // internal format (#channels, #bits/channel, ...)
        g_imageWidth,     // width
        g_imageHeight,    // height
        0,                // border (must be zero)
        GL_RGBA,          // format of input image
        GL_UNSIGNED_BYTE, // type
        data              // data
    );
    glGenerateMipmap(GL_TEXTURE_2D);

    // unbind and free image
    glBindTexture(GL_TEXTURE_2D, 0);
    stbi_image_free(data);

    return textureId;
}

CUDAグラフィックス相互運用

上記の関数を呼び出した後、テクスチャをCUDAに登録します。

void initTexturesCuda(GLuint textureId){
    // register texture
    HANDLE(cudaGraphicsGLRegisterImage(
        &g_textureResource,                       // resource
        textureId,                                // image
        GL_TEXTURE_2D,                            // target
        cudaGraphicsRegisterFlagsSurfaceLoadStore // flags
    ));

    // resource description for surface
    memset(&g_resourceDesc, 0, sizeof(g_resourceDesc));
    g_resourceDesc.resType = cudaResourceTypeArray;
}

レンダリングループ

フレームごとに、次のコマンドを実行してテクスチャを変更し、画像をレンダリングします。

while(!glfwWindowShouldClose(window)){
        // -- CUDA --

        // map
        HANDLE(cudaGraphicsMapResources(1, &g_textureResource));


        HANDLE(cudaGraphicsSubResourceGetMappedArray(
            &g_textureArray,   // array through which to access subresource
            g_textureResource, // mapped resource to access
            0,                 // array index
            0                  // mipLevel
        ));

        // create surface object (compute >= 3.0)
        g_resourceDesc.res.array.array = g_textureArray;
        HANDLE(cudaCreateSurfaceObject(&g_surfaceObj, &g_resourceDesc));

        // run kernel
        kernel<<<gridDim, blockDim>>>(g_surfaceObj, g_imageWidth, g_imageHeight);

        // unmap
        HANDLE(cudaGraphicsUnmapResources(1, &g_textureResource));

        // --- OpenGL ---

        // clear
        glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);

        // use program
        shader.use();

        // triangle
        glBindVertexArray(vao);
        glBindTexture(GL_TEXTURE_2D, textureId);
        glDrawElements(GL_TRIANGLES, 6, GL_UNSIGNED_INT, 0);
        glBindVertexArray(0);

        // glfw:  swap buffers and poll i/o events
        glfwSwapBuffers(window);
        glfwPollEvents();
    }

CUDAカーネル

実際のCUDAカーネルは次のとおりです。

__global__ void kernel(cudaSurfaceObject_t surface, int nx, int ny){
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if(x < nx && y < ny){
        uchar4 data = make_uchar4(x % 255, 
                                  y % 255, 
                                  0, 255);
        surf2Dwrite(data, surface, x * sizeof(uchar4), y);
    }
}