图形互操作源码分析

/*
* Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
*
* NVIDIA Corporation and its licensors retain all intellectual property and
* proprietary rights in and to this software and related documentation.
* Any use, reproduction, disclosure, or distribution of this software
* and related documentation without an express license agreement from
* NVIDIA Corporation is strictly prohibited.
*
* Please refer to the applicable NVIDIA end user license agreement (EULA)
* associated with this source code for terms and conditions that govern
* your use of this NVIDIA software.
*
*/
/*
图形互操作实验
*/
#include <GL\glut.h>
#include "../common/book.h"
#include "../common/cpu_bitmap.h"
#include "device_launch_parameters.h"
#include "cuda.h"
#include "cuda_gl_interop.h"
#include <math.h>
#include <cuda_runtime_api.h>
PFNGLBINDBUFFERARBPROC    glBindBuffer = NULL;
PFNGLDELETEBUFFERSARBPROC glDeleteBuffers = NULL;
PFNGLGENBUFFERSARBPROC    glGenBuffers = NULL;
PFNGLBUFFERDATAARBPROC    glBufferData = NULL;

#define     DIM    512

/*
定义数据缓冲区的两个句柄
bufferObj是OpenGL对这个数据的名字
resource是cuda对这个变量的名字
*/
GLuint  bufferObj;
cudaGraphicsResource *resource;

/*

*/
__global__ void kernel(uchar4 *ptr) {
//计算像素位置
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
int offset = x + y * blockDim.x * gridDim.x;

//计算相应位置上的值
float fx = x / (float)DIM - 0.5f;
float fy = y / (float)DIM - 0.5f;
unsigned char   green = 128 + 127 * sin(abs(fx * 100) - abs(fy * 100));

// accessing uchar4 vs unsigned char*
ptr[offset].x = 0;
ptr[offset].y = green;
ptr[offset].z = 0;
ptr[offset].w = 255;
}

//退出
static void key_func(unsigned char key, int x, int y) {
switch (key) {
case 27:
// clean up OpenGL and CUDA
HANDLE_ERROR(cudaGraphicsUnregisterResource(resource));
glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, 0);
glDeleteBuffers(1, &bufferObj);
exit(0);
}
}

/*
如果没有任何缓冲区绑定为GL_PIXEL_UNPACK_BUFFER_ARB源，那么OpenGL驱动程序将从这个缓冲区进行复制
由于数据位于GPU上，并且我们已经将共享数据缓冲区绑定为GL_PIXEL_UNPACK_BUFFER_ARB源
因此最后一个参数将变成绑定缓冲区的一个偏移
由于是要赋值整个缓冲区，所以偏移值就是0
*/
static void draw_func(void) {
glDrawPixels(DIM, DIM, GL_RGBA, GL_UNSIGNED_BYTE, 0);
glutSwapBuffers();
}

int main(int argc, char **argv) {
cudaDeviceProp  prop;
int dev;

memset(&prop, 0, sizeof(cudaDeviceProp));
prop.major = 1;
prop.minor = 0;
HANDLE_ERROR(cudaChooseDevice(&dev, &prop));

/*
dev中保存的是符合要求的设备ID
互操作性要求，在其他任何运行时调用它之前要通过cudaGLSetGLDevice()指定
将获得的设备ID dev传递进去
为cuda运行时使用OpenGL驱动程序做好准备
*/
HANDLE_ERROR(cudaGLSetGLDevice(dev));

/*
在执行其他任何操作前需要先执行这些GLUT调用
通过GLUT创建名为bitmap的窗口
并在这个窗口中绘制结果
*/
glutInit(&argc, argv);
glutInitDisplayMode(GLUT_DOUBLE | GLUT_RGBA);
glutInitWindowSize(DIM, DIM);
glutCreateWindow("图形互操作演示");

/*
在OpenGL中创建一个数据缓冲区对象
将句柄存放在全局变量bufferObj中
*/
glBindBuffer = (PFNGLBINDBUFFERARBPROC)GET_PROC_ADDRESS("glBindBuffer");
glDeleteBuffers = (PFNGLDELETEBUFFERSARBPROC)GET_PROC_ADDRESS("glDeleteBuffers");
glGenBuffers = (PFNGLGENBUFFERSARBPROC)GET_PROC_ADDRESS("glGenBuffers");
glBufferData = (PFNGLBUFFERDATAARBPROC)GET_PROC_ADDRESS("glBufferData");

// the first three are standard OpenGL, the 4th is the CUDA reg
// of the bitmap these calls exist starting in OpenGL 1.5
//生成缓冲区句柄
glGenBuffers(1, &bufferObj);
//将句柄绑定到缓冲区
glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, bufferObj);
/*
请求OpenGL驱动程序来分配一个缓冲区
GL_DYNAMIC_DRAW_ARB标志表示缓冲区将被应用程序反复修改
刚开始没有初始值，所以倒数第二个参数为null
保存图像大小为DIM*DIM个32位的值
*/
glBufferData(GL_PIXEL_UNPACK_BUFFER_ARB, DIM * DIM * 4,
NULL, GL_DYNAMIC_DRAW_ARB);

/*
通过调用cudaGraphicsGLRegisterBuffer告诉运行时cuda运行时希望在OpenGL和CUDA间使用像素缓冲区数据OpenGL PBO bufferObj
CUDA运行时将在resource中返回一个句柄指向缓冲区
在随后的CUDA运行时调用中，将通过这个句柄来访问缓冲区bufferObj
cudaGraphicsMapFlagsNone标志表示不需要为缓冲区指定特殊的行为
*/
HANDLE_ERROR(
cudaGraphicsGLRegisterBuffer(&resource,
bufferObj,
cudaGraphicsMapFlagsNone));

/*
告诉CUDA运行时映射共享资源
*/
HANDLE_ERROR(cudaGraphicsMapResources(1, &resource, NULL));
uchar4* devPtr;
size_t  size;
/*
请求一个被映射资源的指针
可以把devPtr作为设备指针来使用
*/
HANDLE_ERROR(
cudaGraphicsResourceGetMappedPointer((void**)&devPtr,
&size,
resource));

dim3    grids(DIM / 16, DIM / 16);
dim3    threads(16, 16);
//devPtr为指向共享缓冲区的指针
kernel << <grids, threads >> >(devPtr);
/*
在执行绘制任务之前取消资源的映射以确保应用程序的CUDA部分和图形部分之间实现同步
*/
HANDLE_ERROR(cudaGraphicsUnmapResources(1, &resource, NULL));

// set up GLUT and kick off main loop
glutKeyboardFunc(key_func);
glutDisplayFunc(draw_func);
glutMainLoop();
}