cuda与opengl互操作之VBO

opengl的缓冲区可以映射到CUDA的地址空间，当做global memory被访问。

这样做可以使计算得到的数据直接可视化，提升速度。

因为数据存储在设备端，没有设备端到主机端的传输耗费，不论计算还是可是化都相当的快。

具体使用步骤：
1、创建VBO

[cpp]  view plain  copy

1. glGenBuffers(1, vbo);  
2. glBindBuffer(GL_ARRAY_BUFFER, *vbo);  
3. glBufferData(GL_ARRAY_BUFFER, size, 0, GL_DYNAMIC_DRAW);  
4. glBindBuffer(GL_ARRAY_BUFFER, 0);  

2、注册VBO

[cpp]  view plain  copy

1. struct cudaGraphicsResource *cuda_vbo_resource;  
2. cudaGraphicsGLRegisterBuffer(&cuda_vbo_resource, *vbo, cudaGraphicsMapFlagsWriteDiscard);  

3、映射VBO

[cpp]  view plain  copy

1. cudaGraphicsMapResources(1, &cuda_vbo_resource, 0);  
2. cudaGraphicsResourceGetMappedPointer((void**)&dptr, &num_bytes, cuda_vbo_resource);  

4、使用

[cpp]  view plain  copy

1. launch_kernel(dptr, mesh_width, mesh_height, animTime);  

5、解除映射

[cpp]  view plain  copy

1. cudaGraphicsUnmapResources(1, &cuda_vbo_resource, 0);  

6、解除注册

[cpp]  view plain  copy

1. cudaGraphicsUnregisterResource(cuda_vbo_resource);  

7、删除VBO

[cpp]  view plain  copy

1. glBindBuffer(GL_ARRAY_BUFFER, *vbo);  
2. glDeleteBuffers(1, vbo);  

代码：

[cpp]  view plain  copy

1. //myVBO.cpp  
2. #include <gl/glew.h>  
3. #include <cuda_runtime.h>  
4. #include <cutil_inline.h>  
5. #include <cutil_gl_inline.h>  
6. #include <cutil_gl_error.h>  
7. #include <rendercheck_gl.h>  
8.   
9.   
10. unsigned int window_width = 512;  
11. unsigned int window_height = 512;  
12. unsigned int mesh_width = 256;  
13. unsigned int mesh_height= 256;  
14.   
15. unsigned int timer = 0;  
16.   
17. int animFlag = 1;  
18. float animTime = 0.0f;  
19. float animInc  = 0.01f;  
20.   
21. GLuint vbo = NULL;  
22.   
23. float rotate_x = 0.0, rotate_y = 0.0;  
24. float translate_z = -3.0;  
25.   
26. struct cudaGraphicsResource *cuda_vbo_resource;  
27.   
28. extern "C" void launch_kernel(float4 *pos, unsigned int mesh_width, unsigned int mesh_height, float time);  
29.   
30. void createVBO(GLuint *vbo)  
31. {  
32.     if (vbo)  
33.     {  
34.         glGenBuffers(1, vbo);  
35.         glBindBuffer(GL_ARRAY_BUFFER, *vbo);  
36.   
37.         unsigned int size = mesh_width * mesh_height * 4 * sizeof(float);  
38.         glBufferData(GL_ARRAY_BUFFER, size, 0, GL_DYNAMIC_DRAW);  
39.   
40.         glBindBuffer(GL_ARRAY_BUFFER, 0);     
41.   
42.         cudaGraphicsGLRegisterBuffer(&cuda_vbo_resource, *vbo, cudaGraphicsMapFlagsWriteDiscard);  
43.     }  
44. }  
45.   
46. void deleteVBO(GLuint *vbo)  
47. {  
48.     if (vbo)  
49.     {  
50.         cudaGraphicsUnregisterResource(cuda_vbo_resource);  
51.   
52.         glBindBuffer(GL_ARRAY_BUFFER, *vbo);  
53.         glDeleteBuffers(1, vbo);  
54.   
55.         *vbo = NULL;  
56.     }  
57. }  
58.   
59. void cleanupCuda()  
60. {  
61.     if(vbo) deleteVBO(&vbo);  
62. }  
63.   
64. void runCuda()  
65. {  
66.     float4 *dptr = NULL;  
67.     size_t num_bytes;  
68.   
69.     cudaGraphicsMapResources(1, &cuda_vbo_resource, 0);  
70.     cudaGraphicsResourceGetMappedPointer((void**)&dptr, &num_bytes, cuda_vbo_resource);  
71.   
72.     launch_kernel(dptr, mesh_width, mesh_height, animTime);  
73.   
74.     cudaGraphicsUnmapResources(1, &cuda_vbo_resource, 0);  
75. }  
76.   
77. void initCuda(int argc, char **argv)  
78. {  
79.     if(cutCheckCmdLineFlag(argc, (const char**)argv, "device"))  
80.         cutilGLDeviceInit(argc, argv);  
81.     else  
82.         cudaGLSetGLDevice(cutGetMaxGflopsDeviceId());  
83.   
84.     createVBO(&vbo);  
85.   
86.     atexit(cleanupCuda);  
87.   
88.     runCuda();  
89. }  
90.   
91. void computeFPS()  
92. {  
93.     static int fpsCount = 0;  
94.     static int fpsLimit = 100;  
95.   
96.     fpsCount++;  
97.   
98.     if(fpsCount == fpsLimit)  
99.     {  
100.         char fps[256];  
101.         float ifps = 1.0f / (cutGetAverageTimerValue(timer) / 1000.0f);  
102.         sprintf(fps, "Cuda GL Interop Wrapper: %3.1f fps", ifps);  
103.   
104.         glutSetWindowTitle(fps);  
105.         fpsCount = 0;  
106.   
107.         cutilCheckError(cutResetTimer(timer));    
108.     }  
109. }  
110.   
111. void display()  
112. {  
113.     runCuda();  
114.   
115.     glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);  
116.   
117.     glMatrixMode(GL_MODELVIEW);  
118.     glLoadIdentity();  
119.     glTranslatef(0.0, 0.0, translate_z);  
120.     glRotatef(rotate_x, 1.0, 0.0, 0.0);  
121.     glRotatef(rotate_y, 0.0, 1.0, 0.0);  
122.   
123.     glBindBuffer(GL_ARRAY_BUFFER, vbo);  
124.     glVertexPointer(4, GL_FLOAT, 0, 0);  
125.   
126.     glEnableClientState(GL_VERTEX_ARRAY);  
127.     glColor3f(1.0, 0.0, 0.0);  
128.     glDrawArrays(GL_POINTS, 0, mesh_width*mesh_height);  
129.     glDisableClientState(GL_VERTEX_ARRAY);  
130.   
131.     glutSwapBuffers();  
132.   
133.     if (animFlag)  
134.     {  
135.         glutPostRedisplay();  
136.         animTime += animInc;  
137.     }  
138.   
139.   
140. }  
141.   
142. void fpsDisplay()  
143. {  
144.     cutilCheckError(cutStartTimer(timer));  
145.   
146.     display();  
147.   
148.     cutilCheckError(cutStopTimer(timer));  
149.   
150.     computeFPS();  
151. }  
152.   
153.   
154.   
155. CUTBoolean initGL(int argc, char **argv)  
156. {  
157.     glutInit(&argc, argv);  
158.     glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE);  
159.     glutInitWindowSize(window_width, window_height);  
160.     glutCreateWindow("Cuda GL Interop Demo (adapted from NVDIA's simpleGL)");  
161.   
162.     glutDisplayFunc(fpsDisplay);  
163.   
164.     glewInit();  
165.     if(!glewIsSupported("GL_VERSION_2_0"))  
166.     {  
167.         fprintf(stderr, "ERROR: Support for necessary OpengGL extensions missing.");  
168.         return CUTFalse;  
169.     }  
170.   
171.     glClearColor(0.0, 0.0, 0.0, 1.0);  
172.     glDisable(GL_DEPTH_TEST);  
173.   
174.     glViewport(0, 0, window_width, window_height);  
175.   
176.     glMatrixMode(GL_PROJECTION);  
177.     glLoadIdentity();  
178.     gluPerspective(60.0, (GLfloat)window_width / (GLfloat)window_height, 0.1, 10.0);  
179.   
180.     return CUTTrue;  
181. }  
182.   
183.   
184. int main(int argc, char **argv)  
185. {  
186.     cutilCheckError(cutCreateTimer(&timer));  
187.   
188.     if(CUTFalse == initGL(argc, argv))  
189.         return CUTFalse;  
190.   
191.     initCuda(argc, argv);  
192.     CUT_CHECK_ERROR_GL();  
193.   
194.     glutDisplayFunc(fpsDisplay);  
195.       
196.   
197.     glutMainLoop();  
198.   
199.     cudaThreadExit();           ///  
200.     cutilExit(argc, argv);      //  
201. }  

[cpp]  view plain  copy

1. //kernelVBO.cu  
2.   
3. #include <cuda.h>  
4.   
5.   
6. __global__ void kernel(float4 *pos, unsigned int width, unsigned int height, float time)  
7. {  
8.     unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;  
9.     unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;  
10.   
11.     float u = x / (float)width;  
12.     float v = y / (float)height;  
13.   
14.     u = u * 2.0f - 1.0f;  
15.     v = v * 2.0f - 1.0f;  
16.   
17.     float freq = 4.0f;  
18.   
19.     float w = sinf(u*freq + time) * cosf(v*freq + time) * 0.5f;  
20.   
21.     pos[y*width+x] = make_float4(u, w, v, 1.0f);  
22.   
23. }  
24.   
25. extern "C" void launch_kernel(float4 *pos, unsigned int mesh_width, unsigned int mesh_height, float time)  
26. {  
27.     dim3 block(8, 8, 1);  
28.     dim3 grid(mesh_width/block.x, mesh_height/block.y, 1);  
29.     kernel<<<grid, block>>>(pos, mesh_width, mesh_height, time);  
30.   
31.     cudaThreadSynchronize();  
32. }  

运行结果：

参考自CUDA SDK