CUDA与OpenGL互操作之顶点缓存映射

引言

CUDA相比GLSL更加灵活，逐渐成为主流的GPGPU编程技术。CUDA加速计算后的结果往往需要在显示器中绘制出来，如果我们将在显存中的计算结果拷贝到主存，然后使用QPainter绘制到窗口中，这中间的显存-主存，主存-显存的数据拷贝会显得冗余。既然CUDA计算的结果和图形绘制的内容都在显存里，那么可以通过一些技术手段直接将CUDA计算结果绘制到显示器上，即CUDA与Graphics的互操作。

相关接口介绍

注册显卡资源

__host__ cudaError_t cudaGraphicsGLRegisterBuffer(struct cudaGraphicsResource **resource, GLuint buffer, unsigned int flags);

注册一个OpenGL缓冲对象，参数说明：

resource - 指向返回的对象句柄
*buffer - 被注册的缓冲对象名
flags - 注册标识
- cudaGraphicsRegisterFlagsNone: 注册的资源在CUDA程序可以进行读写，默认为该枚举
- cudaGraphicsRegisterFlagsReadOnly: 指定CUDA程序只能读该资源
- cudaGraphicsRegisterFlagsWriteDiscard: 指定CUDA程序不会读该资源对象，每次写资源时资源的内容完全被覆盖，不会保留之前的任何数据

取消注册函数

cudaError_t cudaGraphicsUnregisterResource(cudaGraphicsResource_t resource);

映射显卡资源

cudaError_t cudaGraphicsMapResources(int count, cudaGraphicsResource_t *resources, cudaStream_t stream __dv(0));

映射已经注册的资源到CUDA流中，参数说明：

count：映射的资源数
resources：映射的资源
stream：映射的CUDA同步流
该函数具备默认同步动作。

取消映射函数

cudaError_t cudaGraphicsUnmapResources(int count, cudaGraphicsResource_t *resources, cudaStream_t stream __dv(0));

获取设备指针

cudaError_t cudaGraphicsResourceGetMappedPointer(void **devPtr, size_t *size, cudaGraphicsResource_t resource);

获取已映射资源的设备指针，参数说明：

devPtr：设备指针，CUDA程序可以使用该指针进行数据更新
size：缓冲数据尺寸
resources：已映射的资源

使用示例

参照CUDA simpleGL示例，该示例是使用CUDA对顶点数据进行操作后，再将顶点绘制到窗口上。使用Qt进行改写：

class CuGl : public QOpenGLWidget, protected QOpenGLFunctions
{
public:
    CuGl(QWidget* parent = nullptr);
    ~CuGl();

protected:
    virtual void initializeGL();
    virtual void resizeGL(int w, int h);
    virtual void paintGL();

private:
    void runCuda();
    void cleanup();

private:
    float m_fAnim;
    Shader *m_shader;
    QOpenGLBuffer m_vbo;
    QOpenGLVertexArrayObject m_vao;
    cudaGraphicsResource_t m_vboResource;   
};

static const int DIM = 512;
static __global__ void kernel(float3* pos, unsigned int width, unsigned int height, float time)
{
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;

    // calculate uv coordinates
    float u = x / (float)width;
    float v = y / (float)height;
    u = u * 2.0f - 1.0f;
    v = v * 2.0f - 1.0f;

    // calculate simple sine wave pattern
    float freq = 4.0f;
    float w = sinf(u*freq + time) * cosf(v*freq + time) * 0.5f;

    // write output vertex
    pos[y*width + x] = make_float3(u, w, v);
}

CuGl::CuGl(QWidget* parent) : QOpenGLWidget(parent), m_fAnim(0.0f)
{
    resize(DIM, DIM);
}

CuGl::~CuGl()
{
    
}

void CuGl::initializeGL()
{   
    initializeOpenGLFunctions();

    m_shader = new Shader("shaders/cuda.vert", "shaders/cuda.frag", this);
    
    // 创建VAO/VBO
    if (m_vao.create())
    {
        m_vao.bind();
        m_vbo.create();
        m_vbo.bind();

        unsigned int nbytes = DIM * DIM * 3 * sizeof(GLfloat);
        m_vbo.allocate(nbytes);

        m_shader->setAttributeBuffer("aPos", GL_FLOAT, 0, 3, 3 * sizeof(float));
        m_shader->enableAttributeArray("aPos");
    }

    // 注册cuda资源
    checkCudaErrors(cudaGraphicsGLRegisterBuffer(&m_vboResource, m_vbo.bufferId(), cudaGraphicsMapFlagsWriteDiscard));
        
    // 设置清屏颜色
    glClearColor(0.1f, 0.1f, 0.1f, 1.0f);

    connect(context(), &QOpenGLContext::aboutToBeDestroyed, this, &CuGl::cleanup);
}


void CuGl::resizeGL(int w, int h)
{
    glViewport(0, 0, w, h);

    QMatrix4x4 model, view, projection;
    view.lookAt(QVector3D(0, 0, 3), QVector3D(0, 0, 0), QVector3D(0, 1, 0));
    projection.perspective(45.0f, (float)w / h, 0.1f, 100.0f);

    m_shader->bind();
    m_shader->setUniformValue("model", model);
    m_shader->setUniformValue("view", view);
    m_shader->setUniformValue("projection", projection);
}


void CuGl::paintGL()
{
    // 使用CUDA进行计算
    runCuda();

    // 使用OPenGL进行绘制
    glClear(GL_COLOR_BUFFER_BIT);
    m_shader->bind();
    m_vao.bind();
    glDrawArrays(GL_POINTS, 0, DIM*DIM);

    update();
}


void CuGl::runCuda()
{
    float3* glPtr;
    size_t size;
    // 映射已注册的资源
    checkCudaErrors(cudaGraphicsMapResources(1, &m_vboResource, 0));

    // 获取映射后的资源的设备指针
    checkCudaErrors(cudaGraphicsResourceGetMappedPointer((void**)&glPtr, &size, m_vboResource));

    // 利用设备指针对数据做计算
    dim3 blocks(16, 16);
    dim3 grids((DIM - 1) / 16 + 1, (DIM - 1) / 16 + 1);
    kernel << <grids, blocks >> > (glPtr, DIM, DIM, m_fAnim);
    m_fAnim += 0.01f;

    // 取消资源映射
    checkCudaErrors(cudaGraphicsUnmapResources(1, &m_vboResource, 0));
}

void CuGl::cleanup()
{
    // 取消资源注册
    checkCudaErrors(cudaGraphicsUnregisterResource(m_vboResource));

    // 销毁VAO/VBO
    m_vao.destroy();
    m_vbo.destroy();
}

顶点着色器和片段着色器程序：

#version 330 core
layout(location = 0) in vec3 aPos;

uniform mat4 model;
uniform mat4 view;
uniform mat4 projection;

void main() 
{
    gl_Position = projection * view * model * vec4(aPos,1.0f);
}

#version 330 core
out vec4 fragColor;
void main() 
{ 
    fragColor = vec4(0.4, 0.7, 0.6, 1.0);
}

运行效果：

Cuda&OpenGL.png