CUDA与OpenGL交互开发

最近在学习OpenGL，过程中需要使用CUDA进行并行计算。因此，需要解决OpenGL与CUDA的交互问题。学习记录如下：

• Step1. 共享数据区

想到交互，不难想到通信，数据共享等词语。这里使用的是共享数据的方式来完成OpenGL与CUDA的交互。而OpenGL与CUDA都有着自己独特的类型定义。因此，对于共享的数据区，我们需要给它起两个不同的名字，分别为OpenGL以及CUDA服务

OpenGL：

[cpp]  view plain copy

1. GLuint bufferObj; 

[cpp]  view plain copy

1. GLuint bufferObj;  

CUDA：

[java]  view plain copy

1. cudaGraphicsResource *resource; 

[java]  view plain copy

1. cudaGraphicsResource *resource;  

• Step2.将显卡设备与OpenGL关联（已废除）

注：在CUDA5.0版本以及以后的版本，不需要此操作。参考NVIDIA官方文档如下：

cudaError_t  cudaGLSetGLDevice ( int  device  )

Deprecated This function is deprecated as of CUDA 5.0. This function is deprecated and should no longer be used. It is no longer necessary to associate a CUDA device with an OpenGL context in order to achieve maximum interoperability performance.

具体的设置代码为：

[cpp]  view plain copy

1. cudaDeviceProp prop; 
2. int dev; 
3. memset(&prop, 0, sizeof(cudaDeviceProp)); 
4. prop.major = 1; 
5. prop.minor = 0; 
6. cudaChooseDevice(&dev, &prop); 
7. cudaGLSetGLDevice(dev); 

[cpp]  view plain copy

1. cudaDeviceProp prop;  
2. int dev;  
3. memset(&prop, 0, sizeof(cudaDeviceProp));  
4. prop.major = 1;  
5. prop.minor = 0;  
6. cudaChooseDevice(&dev, &prop);  
7. cudaGLSetGLDevice(dev);  

• Step3. 初始化OpenGL

[cpp]  view plain copy

1. #define DIM 512 
2. glutInit(argc, argv); 
3. glutInitDisplayMode(GLUT_DOUBLE | GLUT_RGBA); 
4. glutInitWindowSize(DIM, DIM); 
5. glutCreateWindow("bitmap"); 
6. glewInit(); 

[cpp]  view plain copy

1. #define DIM 512  
2. glutInit(argc, argv);  
3. glutInitDisplayMode(GLUT_DOUBLE | GLUT_RGBA);  
4. glutInitWindowSize(DIM, DIM);  
5. glutCreateWindow("bitmap");  
6. glewInit();  

这里需要注意的是：

1. 需要使用opengl扩展库：glew32.dll。若您没有glew32扩展库，点此下载  (其安装方式与glut，freeglut等相同)
2. （重要）需要在opengl初始化代码最后加上：glewInit()，否则会在后面执行到glGenBuffers报运行时错误：0xC0000005: Access violation.
3. 使用glew库需要： #include "gl/glew.h"，且其声明的位置尽量放在代码最顶端，否则编译报错。
4. 具体示例代码，点此下载

——————————————————   华丽的分割线 ————————————————————

到此为止，基本的准备工作就完成了。下面开始实际的工作。

共享数据缓冲区是在CUDA C核函数 和 OpenGL渲染操作之间 实现互操作的关键部分。为了实现两者之间的数据传递，我们首先需要创建一个缓冲区。

• Step4. 使用OpenGL API创建数据缓冲区

[cpp]  view plain copy

1. const GLubyte* a; 
2. a = glGetString(GL_EXTENSIONS); 
3.  
4. glGenBuffers(1, &bufferObj);//生成一个缓冲区句柄 
5. glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, bufferObj);//将句柄绑定到像素缓冲区（即缓冲区存放的数据类型为：PBO） 
6. glBufferData(GL_PIXEL_UNPACK_BUFFER_ARB, DIM*DIM*4, NULL, GL_DYNAMIC_DRAW_ARB);//申请内存空间并设置相关属性以及初始值 

[cpp]  view plain copy

1. const GLubyte* a;  
2. a = glGetString(GL_EXTENSIONS);  
3.   
4. glGenBuffers(1, &bufferObj);//生成一个缓冲区句柄  
5. glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, bufferObj);//将句柄绑定到像素缓冲区（即缓冲区存放的数据类型为：PBO）  
6. glBufferData(GL_PIXEL_UNPACK_BUFFER_ARB, DIM*DIM*4, NULL, GL_DYNAMIC_DRAW_ARB);//申请内存空间并设置相关属性以及初始值  

这里，我们使用如下枚举值：

[cpp]  view plain copy

1. GL_PIXEL_UNPACK_BUFFER_ARB 

[cpp]  view plain copy

1. GL_PIXEL_UNPACK_BUFFER_ARB  

表示指定缓冲区存储的内容是一个Pixel Buffer Object（PBO）

[cpp]  view plain copy

1. GL_DYNAMIC_DRAW_ARB 

[cpp]  view plain copy

1. GL_DYNAMIC_DRAW_ARB  

表示应用程序将会对缓冲区进行修改

这时，可能你会疑问，前两句代码是干嘛的？

因为，GL_PIXEL_UNPACK_BUFFER_ARB对glew扩展库的版本有要求，所以最好检查一下当前环境是否支持GL_PIXEL_UNPACK_BUFFER_ARB枚举值。

可以看到， 我的环境是支持GL_ARB_pixel_buffer_object的，如果您的环境不支持该枚举值，可能需要您更新glew扩展库版本。

• Step5. 把缓冲区分享给CUDA

由于我们的目的是要使用CUDA的并行计算能力，所以CUDA必须要有权利访问共享数据缓冲区。

要实现该操作，需要将缓冲区句柄注册为一个图形资源，即Graphics Resource；然后“分享给CUDA”

[cpp]  view plain copy

1. cudaGraphicsGLRegisterBuffer(&resource, bufferObj, cudaGraphicsMapFlagsNone) 

[cpp]  view plain copy

1. cudaGraphicsGLRegisterBuffer(&resource, bufferObj, cudaGraphicsMapFlagsNone)  

代码中的resource即之前定义的：

[java]  view plain copy

1. cudaGraphicsResource *resource; 

[java]  view plain copy

1. cudaGraphicsResource *resource;  

方法cudaGraphicsGLRegisterBuffer的参数3表示缓冲区属性，它有以下三种可选值：

1.  cudaGraphicsRegisterFlagsNone : Specifies no hints about how this resource will be used. It is therefore assumed that this resource will be read from and written to by CUDA. This is the default value.

2.  cudaGraphicsRegisterFlagsReadOnly : Specifies that CUDA will not write to this resource.（只读）

3.  cudaGraphicsRegisterFlagsWriteDiscard : Specifies that CUDA will not read from this resource and will write over the entire contents of the resource, so none of the data previously stored in the resource will be preserved.（只写）

• Step6. 让CUDA映射共享资源，并获取相对于显卡而言的设备指针

[cpp]  view plain copy

1. uchar4* devPtr; 
2. size_t size; 
3. cudaGraphicsMapResources(1, &resource, NULL); 
4. cudaGraphicsResourceGetMappedPointer((void**)&devPtr, &size, resource); 

[cpp]  view plain copy

1. uchar4* devPtr;  
2. size_t size;  
3. cudaGraphicsMapResources(1, &resource, NULL);  
4. cudaGraphicsResourceGetMappedPointer((void**)&devPtr, &size, resource);  

CUDA官方文档中这样描述：CUDA在访问图形接口（比如openGL）的共享资源之前，需要首先对其进行映射(map)，然后才可以访问共享数据区，CUDA对资源的访问过程中，OpenGL不能对该数据区其进行任何操作，直到CUDA对数据区解除映射（unmap）为止。

Nvidia的原文描述如下：

Map graphics resources for access by CUDA. Maps the count graphics resources in resources for access by CUDA.

The resources in resources may be accessed by CUDA until they are unmapped. The graphics API from whichresources were registered should not access any resources while they are mapped by CUDA. If an application does so, the results are undefined.

映射完成后，我们需要获得缓冲区对于显卡（设备）而言的指针，即代码中的 devPtr。没有设备指针，我们怎么进行并行计算呢。

• Step7. 执行CUDA核函数

[cpp]  view plain copy

1. dim3 grids(DIM/16, DIM/16); 
2. dim3 threads(16, 16); 
3.  
4. kernel_opengl<<<grids, threads>>>(devPtr); 

[cpp]  view plain copy

1. dim3 grids(DIM/16, DIM/16);  
2. dim3 threads(16, 16);  
3.   
4. kernel_opengl<<<grids, threads>>>(devPtr);  

一个简单的核函数kernel_opengl的定义如下：

[cpp]  view plain copy

1. __global__ void kernel_opengl(uchar4* ptr){ 
2.     int x = threadIdx.x + blockIdx.x * blockDim.x; 
3.     int y = threadIdx.y + blockIdx.y * blockDim.y; 
4.     int offset = x + y * blockDim.x * gridDim.x; 
5.  
6.     float fx = x/(float)DIM - 0.5f; 
7.     float fy = y/(float)DIM - 0.5f; 
8.  
9.     unsigned char green = 128 + 127 * sin(abs(fx*100) - abs(fy*100)); 
10.     ptr[offset].x = 0; 
11.     ptr[offset].y = green; 
12.     ptr[offset].z = 0; 
13.     ptr[offset].w = 255; 
14. } 

[cpp]  view plain copy

1. __global__ void kernel_opengl(uchar4* ptr){  
2.     int x = threadIdx.x + blockIdx.x * blockDim.x;  
3.     int y = threadIdx.y + blockIdx.y * blockDim.y;  
4.     int offset = x + y * blockDim.x * gridDim.x;  
5.   
6.     float fx = x/(float)DIM - 0.5f;  
7.     float fy = y/(float)DIM - 0.5f;  
8.   
9.     unsigned char green = 128 + 127 * sin(abs(fx*100) - abs(fy*100));  
10.     ptr[offset].x = 0;  
11.     ptr[offset].y = green;  
12.     ptr[offset].z = 0;  
13.     ptr[offset].w = 255;  
14. }  

此时，执行完核函数，CUDA的使命也就完成了。它的产出就是： 缓冲区的数据已经被更新了~~

• Step8. 解除CUDA对共享缓冲区的映射

[cpp]  view plain copy

1. cudaGraphicsUnmapResources(1, &resource, NULL) 

[cpp]  view plain copy

1. cudaGraphicsUnmapResources(1, &resource, NULL)  

如果不解除映射，那么OpenGL将没有权限访问共享数据区，因此也就没有办法完成图像的渲染显示了。

• Step9. 调用OpenGL API显示

[cpp]  view plain copy

1. glutKeyboardFunc(key_func); 
2. glutDisplayFunc(draw_func); 
3. glutMainLoop(); 

[cpp]  view plain copy

1. glutKeyboardFunc(key_func);  
2. glutDisplayFunc(draw_func);  
3. glutMainLoop();  

其中，显示回调函数为：

[cpp]  view plain copy

1. static void draw_func(void){ 
2.         glDrawPixels(DIM, DIM, GL_RGBA, GL_UNSIGNED_BYTE, 0); 
3.         glutSwapBuffers(); 
4. } 

[cpp]  view plain copy

1. static void draw_func(void){  
2.         glDrawPixels(DIM, DIM, GL_RGBA, GL_UNSIGNED_BYTE, 0);  
3.         glutSwapBuffers();  
4. }  

乍一看，可能感觉会比较奇怪。因为draw_func里面没有使用到缓冲区句柄bufferObj，那么数据如何会显示呢？

因为，之前的代码：

[cpp]  view plain copy

1. glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, bufferObj); 

[cpp]  view plain copy

1. glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, bufferObj);  

该调用将共享缓冲区指定为一个像素源，OpenGL驱动程序随后会在所有对glDrawPixels()的调用中使用这个像素源。这也就说明了我们为什么使用glDrawPixels来绘制图形了。

通过查看glDrawPixels的文档：

[cpp]  view plain copy

1. void glDrawPixels( 
2.   GLsizei width, 
3.   GLsizei height, 
4.   GLenum format, 
5.   GLenum type, 
6.   const GLvoid *pixels 
7. ); 

[cpp]  view plain copy

1. void glDrawPixels(  
2.   GLsizei width,  
3.   GLsizei height,  
4.   GLenum format,  
5.   GLenum type,  
6.   const GLvoid *pixels  
7. );  

不难发现其最后一个参数为一个缓冲区指针。如果没有任何缓冲区被指定为GL_PIXEL_UNPACK_BUFFER_ARB源，那么OpenGL将从这个参数指定的缓冲区进行数据复制并显示。但在本例中，我们已经将共享数据缓冲区指定为GL_PIXEL_UNPACK_BUFFER_ARB。此时，该参数含义将变为：已绑定缓冲区内的偏移量，由于我们要绘制整个缓冲区，因此这便宜量就是0.

—————————————————————— 华丽的分割线 ————————————————

最终，运行程序，得到指定的结果。

转自：http://blog.csdn.net/lingling_y/article/details/8915163