项目打包下载
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 <GLglut.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 }