CUDA与OpenGL混合编程图形渲染

---

前言

一、 通过OpenGL映射GPU内存

1. 从CUDA编程者角度看，OpenGL在GPU上创建并管理的通过缓冲的内存区域称为缓冲对象。
2. CUDA Kernel将一段缓冲映射如CUDA内存空间，可实现CUDA和OpenGL的互操作。当释放该缓冲或解除映射时，其控制权重回OpenGL。因为不需要进行内存拷贝，所以映射是一种处理速度很快的低开销操作，实现OpenGL和CUDA之间的高速互操作能力。
3. 与OpenGL进行互操作，需要在所有其他Runtime调用之前用cudaGLSetGLDevice()函数来指定CUDA设备。注意，cudaSetDevice()和cudaGLSetGLDevice()是互相排斥的。一旦对某个资源对CUDA进行了注册，便可以根据需要通过cudaGraphicsMapResources()和sourceSetMapFlags()方法可用于设备提示标识（例如，只读，只写等），以便CUDA驱动程序进行优化资源管理。

二、CUDA使用的两种OpenGL内存对象

1、像素缓冲对象（PBO）：OpenGL中用于存储像素的一段内存。2D图像是由多个像素和颜色点组成的。CUDA程序映射PBO，并逐个像素生成或修改图像，然后利用OpenGL进行显示。

2、 向量缓冲对象（VBO）：OpenGL中用于存储3D向量的一段内存。CUDA程序映射VBO，生成或修改3D位置信息，之后OpenGL将这些网格渲染成彩色表面、3D线框图像或3D点集。

glBindBuffer()

glBufferData()/glBufferSubData()/glGetBufferSubData()

glMapBuffer()/glUnmapBuffer()

三、CUDA与OpenGL互操作的关键cudaGraphicsResourceGetMappedPointer函数，用于获取映射后的设备指针

典型调用流程‌

1. 注册OpenGL缓冲对象：cudaGraphicsGLRegisterBuffer(&resource, bufferObj, flags)
2. 映射资源：cudaGraphicsMapResources(1, &resource, 0)
3. 获取设备指针：cudaGraphicsResourceGetMappedPointer((void**)&devPtr, &size, resource)
4. 核函数处理：kernel<<<…>>>(devPtr, …)（直接操作显存数据）
5. 解除映射：cudaGraphicsUnmapResources(1, &resource, 0)

四、 CUDA和OpenGL互操作的混合编程步骤

1、 创建窗口及OpenGL运行环境

glutInit(argc, argv);
glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE);
glutInitWindowSize(window_width, window_height);
glutCreateWindow("Cuda GL Interop (VBO)");
glutDisplayFunc(display);
glutKeyboardFunc(keyboard);
glutMotionFunc(motion);
glutTimerFunc(REFRESH_DELAY, timerEvent, 0);

2、设置OpenGL视口和坐标系。要根据绘制的图形是2D还是3D等具体情况设置。(1)和(2)是所有OpenGL程序必需的，这里也没什么特殊之处，需要注意的是，后面的一些功能需要OpenGL 2.0及以上版本支持，所以在这里需要进行版本检查。

 // 版本检查 if (!isGLVersionSupported(2, 0)){fprintf(stderr, "ERROR: Support for necessary OpenGL extensions missing.");fflush(stderr);return false;}// default initializationglClearColor(0.0, 0.0, 0.0, 1.0);glDisable(GL_DEPTH_TEST);// viewport   设置视口glViewport(0, 0, window_width, window_height);// projection 坐标系glMatrixMode(GL_PROJECTION);glLoadIdentity();gluPerspective(60.0, (GLfloat)window_width / (GLfloat)window_height, 0.1, 10.0);

3、创建CUDA环境。可以使用cuGLCtxCreate或cudaGLSetGLDevice来设置CUDA环境。该设置一定要放在其他CUDA的API调用之前。

 // 3. 创建CUDA环境 cudaGLSetGLDevice(gpuGetMaxGflopsDeviceId());

4、产生一个或多个OpenGL缓冲区用以和CUDA共享。使用PBO和使用VBO差不多，只是有些函数调用参数不同。以下是具体过程

 void createVBO(GLuint* vbo, struct cudaGraphicsResource** vbo_res,unsigned int vbo_res_flags)
{assert(vbo);// create buffer object//产生一个buffer IDglGenBuffers(1, vbo);//将其设置为当前非压缩缓冲区，// 如果是PBO方式，parameter1设置为GL_PIXEL_UNPACK_BUFFER，// 如果是VBO方式，parameter1设置为GL_ARRAY_BUFFERglBindBuffer(GL_ARRAY_BUFFER, *vbo);// initialize buffer objectunsigned int size = mesh_width * mesh_height * 4 * sizeof(float);//给该缓冲区分配数据，// PBO方式下，parameter1设置为GL_PIXEL_UNPACK_BUFFER，parameter2设置为图像的长度*宽度*4。// VBO方式下，parameter1设置为GL_ARRAY_BUFFER，parameter2设置为顶点数*16，因为每个顶点包含3个浮点坐标(x，y，z)和4个颜色字节(RGBA)，这样一个顶点包含16BglBufferData(GL_ARRAY_BUFFER, size, 0, GL_DYNAMIC_DRAW);glBindBuffer(GL_ARRAY_BUFFER, 0);// register this buffer object with CUDAcheckCudaErrors(cudaGraphicsGLRegisterBuffer(vbo_res, *vbo, vbo_res_flags));SDK_CHECK_ERROR_GL();
}

5、用CUDA登记缓冲区。登记可以使用cuGLRegisterBufferObject或cudaGLRegisterBufferObject，该命令告诉OpenGL和CUDA 驱动程序该缓冲区为二者共同使用。

 // register this buffer object with CUDAcheckCudaErrors(cudaGraphicsGLRegisterBuffer(vbo_res, *vbo, vbo_res_flags));

6、将OpenGL缓冲区映射到CUDA内存。可以使用cuGLMapBufferObject或cudaGLMapBufferObject，它实际是将CUDA内存的指针指向OpenGL的缓冲区，这样如果只有一个GPU，就不需要数据传递。当映射完成后，OpenGL不能再使用该缓冲区。

7、使用CUDA往该映射的内存写图像数据。前面的准备工作在这里真正发挥作用了，此时可以调用CUDA的kernel，像使用全局内存一样使用映射了的缓冲区，向其中写数据。

8、取消OpenGL缓冲区映射。要等前面CUDA的活动完成以后，使用cuGLUnmapBufferObject或cudaGLUnmapBufferObject函数取消映射

void runCuda(struct cudaGraphicsResource** vbo_resource)
{// map OpenGL buffer object for writing from CUDAfloat4* dptr;checkCudaErrors(cudaGraphicsMapResources(1, vbo_resource, 0));size_t num_bytes;// 是CUDA与OpenGL互操作的关键函数，用于获取映射后的设备指针// 该函数仅适用于通过cudaGraphicsGLRegisterBuffer注册的缓冲对象（如PBO/VBO），// 纹理对象需使用cudaGraphicsSubResourceGetMappedArraycheckCudaErrors(cudaGraphicsResourceGetMappedPointer((void**)&dptr, &num_bytes,*vbo_resource));//printf("CUDA mapped VBO: May access %ld bytes\n", num_bytes);// execute the kernel//    dim3 block(8, 8, 1);//    dim3 grid(mesh_width / block.x, mesh_height / block.y, 1);//    kernel<<< grid, block>>>(dptr, mesh_width, mesh_height, g_fAnim);launch_kernel(dptr, mesh_width, mesh_height, g_fAnim);// unmap buffer objectcheckCudaErrors(cudaGraphicsUnmapResources(1, vbo_resource, 0));
}

9、前面的步骤完成以后就可以真正开始绘图了， OpenGL的PBO和VBO的绘图方式不同，分别为以下两个过程

①如果只是绘制平面图形，需要使用OpenGL的PBO及纹理

//使纹理可用
glEnable(GL_TEXTURE_2D)； 
//生成一个textureID
glGenTextures(1，&textureID)； 
//使该纹理成为当前可用纹理
glBindTexture(GL_TEXTURE_2D，textureID)；
//分配纹理内存。最后的参数设置数据来源，这里设置为NULL，表示数据来自PBO，不是来自主机内存
glTexImage2D(GL_TEXTURE_2D，0，GL_RGBA8，Width， Height，0，GL_BGRA，GL_UNSIGNED_BYTE，NULL)；
//必须设置滤波模式，GL_LINEAR允许图形伸缩时线性差值。如果不需要线性差值，可以用GL_TEXTURE_RECTANGLE_ARB代替GL_TEXTURE_2D以提高性能，同时在glTexParameteri()调用里使用GL_NEAREST替换GL_LINEAR 然后就可以指定4个角的纹理坐标，绘制长方形了
glTexParameteri(GL_TEXTURE_2D，GL_TEXTURE_MIN_FILTER，GL_LINEAR)；
glTexParameteri(GL_TEXTURE_2D，GL_TEXTURE_MAG_FILTER，GL_LINEAR)；

② 绘制3D场景，需要使用VBO

//使顶点和颜色数组可用
glEnableClientState(GL_VERTEX_ARRAY)；
//设置顶点和颜色指针
glEnableClientState(GL_COLOR_ARRAY)；
glVertexPointer(3，GL_FLOAT，16，0)；
//根据顶点数据绘图，参数可以使用GL_LINES， GL_LINE_STRIP， GL_LINE_LOOP， GL_TRIANGLES，GL_TRIANGLE_STRIP， GL_TRIANGLE_FAN， GL_QUADS，GL_QUAD_STRIP，GL_POLYGONglColorPointer(4，GL_UNSIGNED_BYTE，16，12)；
glDrawArrays(GL_POINTS，0，numVerticies)；

10、前后缓存区来回切换，实现动画显示效果。调用SwapBuffers()，缓冲区切换通常会在垂直刷新间隙来处理，因此，可以在控制面板上关掉垂直同步，使得缓冲区切换立刻进行。

效果图

总结

CUDA项目源码地址：https://github.com/chensongpoixs/ccuda_sample/tree/master/02_cuda_texture