图形互操作源码分析

  1 /*
  2 * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
  3 *
  4 * NVIDIA Corporation and its licensors retain all intellectual property and
  5 * proprietary rights in and to this software and related documentation.
  6 * Any use, reproduction, disclosure, or distribution of this software
  7 * and related documentation without an express license agreement from
  8 * NVIDIA Corporation is strictly prohibited.
  9 *
 10 * Please refer to the applicable NVIDIA end user license agreement (EULA)
 11 * associated with this source code for terms and conditions that govern
 12 * your use of this NVIDIA software.
 13 *
 14 */
 15 /*
 16 图形互操作实验
 17 */
 18 #include <GL\glut.h>
 19 #include "../common/book.h"
 20 #include "../common/cpu_bitmap.h"
 21 #include "device_launch_parameters.h"
 22 #include "cuda.h"
 23 #include "cuda_gl_interop.h"
 24 #include <math.h>
 25 #include <cuda_runtime_api.h>
 26 PFNGLBINDBUFFERARBPROC    glBindBuffer = NULL;
 27 PFNGLDELETEBUFFERSARBPROC glDeleteBuffers = NULL;
 28 PFNGLGENBUFFERSARBPROC    glGenBuffers = NULL;
 29 PFNGLBUFFERDATAARBPROC    glBufferData = NULL;
 30 
 31 #define     DIM    512
 32 
 33 /*
 34 定义数据缓冲区的两个句柄
 35 bufferObj是OpenGL对这个数据的名字
 36 resource是cuda对这个变量的名字
 37 */
 38 GLuint  bufferObj;
 39 cudaGraphicsResource *resource;
 40 
 41 /*
 42 
 43 */
 44 __global__ void kernel(uchar4 *ptr) {
 45     //计算像素位置
 46     int x = threadIdx.x + blockIdx.x * blockDim.x;
 47     int y = threadIdx.y + blockIdx.y * blockDim.y;
 48     int offset = x + y * blockDim.x * gridDim.x;
 49 
 50     //计算相应位置上的值
 51     float fx = x / (float)DIM - 0.5f;
 52     float fy = y / (float)DIM - 0.5f;
 53     unsigned char   green = 128 + 127 * sin(abs(fx * 100) - abs(fy * 100));
 54 
 55     // accessing uchar4 vs unsigned char*
 56     ptr[offset].x = 0;
 57     ptr[offset].y = green;
 58     ptr[offset].z = 0;
 59     ptr[offset].w = 255;
 60 }
 61 
 62 //退出
 63 static void key_func(unsigned char key, int x, int y) {
 64     switch (key) {
 65     case 27:
 66         // clean up OpenGL and CUDA
 67         HANDLE_ERROR(cudaGraphicsUnregisterResource(resource));
 68         glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, 0);
 69         glDeleteBuffers(1, &bufferObj);
 70         exit(0);
 71     }
 72 }
 73 
 74 /*
 75 如果没有任何缓冲区绑定为GL_PIXEL_UNPACK_BUFFER_ARB源，那么OpenGL驱动程序将从这个缓冲区进行复制
 76 由于数据位于GPU上，并且我们已经将共享数据缓冲区绑定为GL_PIXEL_UNPACK_BUFFER_ARB源
 77 因此最后一个参数将变成绑定缓冲区的一个偏移
 78 由于是要赋值整个缓冲区，所以偏移值就是0
 79 */
 80 static void draw_func(void) {
 81     glDrawPixels(DIM, DIM, GL_RGBA, GL_UNSIGNED_BYTE, 0);
 82     glutSwapBuffers();
 83 }
 84 
 85 
 86 int main(int argc, char **argv) {
 87     cudaDeviceProp  prop;
 88     int dev;
 89 
 90     memset(&prop, 0, sizeof(cudaDeviceProp));
 91     prop.major = 1;
 92     prop.minor = 0;
 93     HANDLE_ERROR(cudaChooseDevice(&dev, &prop));
 94 
 95     /*
 96     dev中保存的是符合要求的设备ID
 97     互操作性要求，在其他任何运行时调用它之前要通过cudaGLSetGLDevice()指定
 98     将获得的设备ID dev传递进去
 99     为cuda运行时使用OpenGL驱动程序做好准备
100     */
101     HANDLE_ERROR(cudaGLSetGLDevice(dev));
102 
103     /*
104     在执行其他任何操作前需要先执行这些GLUT调用
105     通过GLUT创建名为bitmap的窗口
106     并在这个窗口中绘制结果
107     */
108     glutInit(&argc, argv);
109     glutInitDisplayMode(GLUT_DOUBLE | GLUT_RGBA);
110     glutInitWindowSize(DIM, DIM);
111     glutCreateWindow("图形互操作演示");
112 
113     /*
114     在OpenGL中创建一个数据缓冲区对象
115     将句柄存放在全局变量bufferObj中
116     */
117     glBindBuffer = (PFNGLBINDBUFFERARBPROC)GET_PROC_ADDRESS("glBindBuffer");
118     glDeleteBuffers = (PFNGLDELETEBUFFERSARBPROC)GET_PROC_ADDRESS("glDeleteBuffers");
119     glGenBuffers = (PFNGLGENBUFFERSARBPROC)GET_PROC_ADDRESS("glGenBuffers");
120     glBufferData = (PFNGLBUFFERDATAARBPROC)GET_PROC_ADDRESS("glBufferData");
121 
122     // the first three are standard OpenGL, the 4th is the CUDA reg 
123     // of the bitmap these calls exist starting in OpenGL 1.5
124     //生成缓冲区句柄
125     glGenBuffers(1, &bufferObj);
126     //将句柄绑定到缓冲区
127     glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, bufferObj);
128     /*
129     请求OpenGL驱动程序来分配一个缓冲区
130     GL_DYNAMIC_DRAW_ARB标志表示缓冲区将被应用程序反复修改
131     刚开始没有初始值，所以倒数第二个参数为null
132     保存图像大小为DIM*DIM个32位的值
133     */
134     glBufferData(GL_PIXEL_UNPACK_BUFFER_ARB, DIM * DIM * 4,
135         NULL, GL_DYNAMIC_DRAW_ARB);
136 
137     /*
138     通过调用cudaGraphicsGLRegisterBuffer告诉运行时cuda运行时希望在OpenGL和CUDA间使用像素缓冲区数据OpenGL PBO bufferObj
139     CUDA运行时将在resource中返回一个句柄指向缓冲区
140     在随后的CUDA运行时调用中，将通过这个句柄来访问缓冲区bufferObj
141     cudaGraphicsMapFlagsNone标志表示不需要为缓冲区指定特殊的行为
142     */
143     HANDLE_ERROR(
144         cudaGraphicsGLRegisterBuffer(&resource,
145         bufferObj,
146         cudaGraphicsMapFlagsNone));
147 
148     /*
149     告诉CUDA运行时映射共享资源
150     */
151     HANDLE_ERROR(cudaGraphicsMapResources(1, &resource, NULL));
152     uchar4* devPtr;
153     size_t  size;
154     /*
155     请求一个被映射资源的指针
156     可以把devPtr作为设备指针来使用
157     */
158     HANDLE_ERROR(
159         cudaGraphicsResourceGetMappedPointer((void**)&devPtr,
160         &size,
161         resource));
162 
163     dim3    grids(DIM / 16, DIM / 16);
164     dim3    threads(16, 16);
165     //devPtr为指向共享缓冲区的指针
166     kernel << <grids, threads >> >(devPtr);
167     /*
168     在执行绘制任务之前取消资源的映射以确保应用程序的CUDA部分和图形部分之间实现同步
169     */
170     HANDLE_ERROR(cudaGraphicsUnmapResources(1, &resource, NULL));
171 
172     // set up GLUT and kick off main loop
173     glutKeyboardFunc(key_func);
174     glutDisplayFunc(draw_func);
175     glutMainLoop();
176 }