CUDA/OpenGL Interop:写入表面对象不会删除以前的内容

Question

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

可能的问题/误解

我是CUDA和OpenGL的新手.在这里,我将尝试解释导致我使用此代码的思维过程:

• 我正在使用a cudaArray来访问纹理(而不是像浮点数组),因为我读到在读/写纹理时最好是缓存局部性.
• 我正在使用曲面因为我在某处读到它是修改a的唯一方法 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);
    }
}

Answer 1

如果我理解正确,您最初注册纹理,映射一次,为表示映射纹理的数组创建表面对象,然后取消映射纹理.然后,每个帧再次映射资源,请求表示映射纹理的数组,然后完全忽略该数组,并使用为第一次映射资源时返回的数组创建的表面对象.从文档:

[...] array每次resource映射时,设置的值都可能会更改.

每次映射资源时都必须创建一个新的表面对象,因为每次都可能得到一个不同的数组.而且,根据我的经验,你实际上会经常得到一个不同的.只有在数组实际更改时才创建新的表面对象可能是有效的.文档似乎允许这样做,但我从未尝试过,所以我无法确定这是否有效...

除此之外:您为纹理生成mipmap.您只覆盖mip级别0.然后使用带有三线性插值的mipmapping渲染纹理.所以我的猜测是你恰好以一个与mip级别0的分辨率不完全匹配的分辨率渲染纹理,因此,你将最终在0级(你写的)和1级(这之间)进行插值.是从原始纹理生成的)...