CUDA 与 OpenGL 的互操作

 

　　cuda 与 opengl 的互操作一般是使用cuda生成数据，然后在opengl中渲染数据对应的图形。这两者的结合有两种方式：

　　　　1、使用opengl中的pbo（像素缓冲区对象）。cuda生成像素数据，opengl直接渲染即可。

　　　　2、使用opengl中的fbo（顶点缓冲区对象）。cuda生成顶点数据，opengl渲染。

　　这两种方法的核心都是将opengl中的缓冲区对象映射到cuda内存空间中（让cuda的内存指针指向opengl的缓冲区），这样就不需要将缓冲区中的数据传输至cuda内存中，然后利用cuda的高并行计算性能加速计算，最后直接使用opengl渲染。

　　

　　一个例子，使用cuda根据时间动态生成16个点，在屏幕上显示。

　　步骤：

　　1、设置与opengl互操作的设备

status=cudaglsetgldevice(0);

　　2、在cuda中注册缓冲区对象

status = cudaglregisterbufferobject(this->vbo);

　　3、映射缓冲区对象：让cuda内存指针指向缓冲区对象对应的空间

// 映射缓冲对象
    float4* position;
    status=cudaglmapbufferobject((void**)&position, this->vbo);

　　4、运行核函数

// 运行核函数
    dim3 dimblock(4, 4, 1);
    dim3 dimgrid(1);
    kernelfunc<<<dimgrid, dimblock>>>(position, clock(), 4, 4);
    cudathreadsynchronize(); //同步线程

　　5、取消映射

status=cudaglunmapbufferobject(this->vbo);

　效果图：

　　

　　注意：当cuda的kernel函数修改cuda指针指向的空间超出opengl缓冲对象大小时，会导致后面出现取消映射失败。（这里所说的cuda指针是映射到opengl缓冲对象的）

　　完整代码如下：

　　

#include "genvertex.cuh"
#include <time.h>


__global__ void kernelfunc(float4* position, float time, unsigned int width, unsigned int height)
{
    unsigned int x = blockidx.x*blockdim.x + threadidx.x;
    unsigned int y = blockidx.y*blockdim.y + threadidx.y;
    float u = x / (float)width;
    float v = y / (float)height;
    u = u*2.0f - 1.0f;
    v = v*2.0f - 1.0f;
    float freq = 4.0f;
    float w = sinf(u*freq + time*0.001f)*cosf(v*freq + time*0.001f)*0.5f;
    position[y*width + x] = make_float4(u*10, w*10, v*10, 1.0f);
}

genvertex::genvertex()
{
    this->setup();
}

genvertex::~genvertex()
{
}

void genvertex::setup() {
    cudaerror_t status;
    //设备设置
    status=cudaglsetgldevice(0);
    if (status != cudasuccess) {
        puts("setup device failed!");
    }
}

void genvertex::setvbo(unsigned int vbo) {
    this->vbo = vbo;
    cudaerror_t status;
    status = cudaglregisterbufferobject(this->vbo);
    if (status != cudasuccess) {
        puts("register buffer object failed!");
    }
}


void genvertex::createvtxwithcuda()
{
    cudaerror_t status;
    // 映射缓冲对象
    float4* position;
    status=cudaglmapbufferobject((void**)&position, this->vbo);
    if (status != cudasuccess) {
        puts("map buffer object failed!");
    }
    // 运行核函数
    dim3 dimblock(4, 4, 1);
    dim3 dimgrid(1);
    kernelfunc<<<dimgrid, dimblock>>>(position, clock(), 4, 4);
    cudathreadsynchronize(); //同步线程
    status=cudaglunmapbufferobject(this->vbo);
    if (status != cudasuccess) {
        puts("unmap buffer object failed!");
    }
}