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 =http://www.mamicode.com/ 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 = http://www.mamicode.com/(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.5124     //生成缓冲区句柄125     glGenBuffers(1, &bufferObj);126     //将句柄绑定到缓冲区127     glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, bufferObj);128     /*129     请求OpenGL驱动程序来分配一个缓冲区130     GL_DYNAMIC_DRAW_ARB标志表示缓冲区将被应用程序反复修改131     刚开始没有初始值，所以倒数第二个参数为null132     保存图像大小为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 bufferObj139     CUDA运行时将在resource中返回一个句柄指向缓冲区140     在随后的CUDA运行时调用中，将通过这个句柄来访问缓冲区bufferObj141     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 loop173     glutKeyboardFunc(key_func);174     glutDisplayFunc(draw_func);175     glutMainLoop();176 }