我正在尝试使用CUDA内核来修改OpenGL纹理，但是遇到一个奇怪的问题，即我的调用surf2Dwrite()似乎与纹理的先前内容融合在一起，如下图所示。背面的木质纹理是使用我的CUDA内核修改之前的纹理。预期的输出将仅包括颜色渐变，而不包括其后面的木材纹理。我不明白为什么会这样混合。

可能的问题/误解

我是CUDA和OpenGL的新手。在这里，我将解释导致我编写此代码的思考过程：

• 我使用acudaArray来访问纹理（而不是例如浮点数数组），因为我读到它在读/写纹理时对于缓存局部性更好。
• 我使用表面是因为我在某处读到这是修改表面的唯一方法 cudaArray
• 我想使用表面对象，我知道这是更新的工作方式。旧的方法是使用表面参考。

我不知道如何检查/测试的代码可能存在的一些问题：

• 我与图像格式不一致吗？也许我没有在某处指定正确的位数/通道？也许我应该使用floats而不是unsigned chars？

代码摘要

您可以在此GitHub Gist中找到完整的最低限度的工作示例。由于所有的活动部件，时间很长，但是我将尝试进行总结。我欢迎有关如何缩短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);
    }
}