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

评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值