8.2_图形互操作

8.2 图形互操作

为了说明在图形库与CUDA C之间的互操作机制，我们将编写一个包含两步骤的应用程序。第一个步骤是使用CUDA C核函数来生成图像数据。在第二个步骤中，应用程序将这个数据传递给OpenGL驱动程序并进行渲染。要实现这个功能，我们将使用在前面章节中介绍的大部分CUDA C，以及一些OpenGL或者GLUT函数调用。

首先，我们要包含GLUT和CUDA的头文件从而确保定义了正确的函数和枚举类型。我们还定义了应用程序渲染窗口的大小。窗口为 512×512512 \times 512 个像素，这是一个相对较小的绘制量。

define GL_GLEXT_PROTOTYPES #include "GL/glut.h" #include "CUDA.h" #include "cuda_gl_interop.h" #include "/common/book.h" #include "/common/cpu bitmap.h" #define DIM 512

此外，我们声明了两个全局变量来保存句柄，这些句柄指向将要在OpenGL和CUDA C之间共享的数据。我们马上会看到如何使用这两个变量，它们将保存指向同一个缓冲区的不同句柄。之所以需要两个独立的变量，是因为OpenGL和CUDA对于这个缓冲区各自有着不同的“名字”。变量bufferObj是OpenGL对这个数据的命名，而变量resource则是CUDA C对这个变量的命名。

GLuint bufferObj;  
cudaGraphicsResource *resource;

现在，让我们来看看实际的应用程序。要做的第一件事情就是选择运行应用程序的CUDA设备。在许多系统上，这并不是一个复杂的过程，因为这些系统通常只包含一个支持CUDA的GPU。然而，随着越来越多的系统包含了多个支持CUDA的GPU，就需要通过某种方法从中进行选择。幸运的是，CUDA运行时提供了这种功能。

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));

你或许还记得在第3章中看到的cudaChooseDevice()，我们现在再次使用它。基本上，这段代码告诉运行时选择一个拥有1.0或者更高版本计算功能集的GPU。代码的原理是，首先创建一个cudaDeviceProg结构并将其初始化为空，然后将major版本设置为1，minor版本设置为0。接下来，将这个结构传递给cudaChooseDevice()，这个函数将告诉运行时选择系统中的某个满足cudaDeviceProg结构指定条件的GPU。在第9章中，我们将看到GPU计算功能集的更多含义，但就目前来说，我们只需要知道它基本上表示GPU支持的各种功能。所有支持CUDA的GPU都至少包含1.0版本的计算功能集，因此这个函数调用的结果就是运行时可以选择任何一个支持CUDA的设备，并且在变量dev中返回这个设备的标识符。然而，我们无法确保这个设备是最好的或是最快的GPU，也不能确保不同版本的GPU运行时会选择同一个设备。

如果设备选择的结果看上去没有太多的作用，那么为什么还要费力填充一个CUDADeviceProp结构，并调用CUDAChooseDevice()来获得一个有效的设备ID？而且，我们之前从来没有这么做，那么为什么现在需要这么做？这些都是很好的问题。事实证明，我们需要知道CUDA设备的ID，这样才可以告诉CUDA运行时应该使用哪个设备来执行CUDA和OpenGL。我们通过调用CUDASetGLDevice()来实现这个功能，并把在CUDAChooseDevice()中获得的设备ID dev传递进去。

HANDLE_ERROR(udaGLSetGLDevice(dev));

在CUDA运行时初始化之后，就可以继续调用GL工具箱（GL Utility Toolkit，GLUT）的设置函数来初始化OpenGL驱动程序。如果你之前使用过GLUT，那么下面这些函数调用看上去就很熟悉：

// 在执行其他的GL调用之前，需要首先执行这些GLUT调用。glutInit( &argc, argv );glutInitDisplayMode(GLUT_DOUBLE | GLUT_RGB);glutInitWindowSize(DIM, DIM);glutCreateWindow(" bitmap");

在main()的这个位置上，我们通过调用CUDASetGLDevice()为CUDA运行时使用OpenGL驱动程序做好准备。然后，我们初始化GLUT并且创建一个名为“ bitmap”的窗口，并将在这个窗口中绘制结果。现在，我们可以开始执行实际的OpenGL互操作！

共享数据缓冲区是在CUDA C核函数和OpenGL渲染操作之间实现互操作的关键部分。要在

OpenGL和CUDA之间传递数据，我们首先要创建一个缓冲区在这两组API之间使用。首先在OpenGL中创建一个像素缓冲区对象，并将句柄保存在全局变量GLuint bufferObj中：

glGenBuffers(1, &bufferObj);  
glBindBuffer(GL_PIXL_UNPACK_bufferARB, bufferObj);  
glBufferData(GL_PIXL_UNPACK_bufferARB, DIM * DIM * 4, NULL, GL_DYNAMIC_DRAWARB);

如果你从来没有用过OpenGL中的像素缓冲区对象（Pixel Buffer Object，PBO），那么可以通过以下三个步骤来创建一个：首先，通过glGenBuffers()生成一个缓冲区句柄。然后，通过glBindBuffer()将句柄绑定到像素缓冲区。最后，通过glBufferData()请求OpenGL驱动程序来分配一个缓冲区。在这个示例中请求分配一个缓冲区来保存DIM×DIM个32位的值，并且使用权举值GL_DYNAMIC_DRAW_ARB来表示这个缓冲区将被应用程序反复修改。由于没有任何数据预先加载到缓冲区，因此将glBufferData()的倒数第二个参数设置为NULL。

在设置图形互操作性中，剩下的工作就是通知CUDA运行时，缓冲区bufferObj将在CUDA与OpenGL之间共享。要实现这个操作，需要将bufferObj注册为一个图形资源（Graphics Resource）。

HANDLE_ERROR(udaGraphicsGLRegisterBuffer( &resource,bufferObj,udaGraphicsMapFlagsNone)）;

通过调用cudaGraphicsGLRegisterBuffer(), 我们告诉CUDA运行时希望在OpenGL和CUDA中使用OpenGL PBO bufferObj。CUDA运行时将在变量resource中返回一个句柄指向缓冲区。在随后对CUDA运行时的调用中，将通过这个句柄来访问bufferObj。

标志cusad GraphicsMapFlagsNone表示不需要为缓冲区指定特定的行为，当然我们也可以通过标志cusad GraphicsMapFlagsReadOnly将缓冲区指定为只读的。我们还可以通过标志cusad GraphicsMapFlagsWriteDiscard来指定缓冲区中之前的内容应该抛弃，从而使缓冲区变成只写的。这些标志使得CUDA和OpenGL驱动程序根据缓冲区的访问模式对硬件配置进行优化，当然这些标志并不一定必须设置。

对glBufferData()的调用需要OpenGL驱动程序分配一个足够大的缓冲区来保存DIM×DIM个32位的值。在随后的OpenGL调用中，我们通过bufferObj来引用这个缓冲区，而在CUDA运行时调用中，则通过指针resource来引用这个缓冲区。由于我们将在CUDA C核函数中对这个缓冲区进行读写，因此需要多个指向该对象的句柄。我们需要设备内存中的一个实际地址并传递给核函数。首先告诉CUDA运行时映射这个共享资源，然后请求一个指向被映射资源的指针。

uchar4* devPtr;
size_t size;
HANDLE_ERROR(udaGraphicsMapResources(1, &resource, NULL));
HANDLE_ERROR(udaGraphicsResourceGetMappedPointer((void**)&devPtr, &size, resource));

然后，可以把devPtr作为设备指针来使用，此外这个数据还可以作为一个像素源由OpenGL使用。在完成这些设置步骤后，main()剩余工作的执行流程为：首先，启动核函数并将指向共享缓冲区的指针传递给它。尽管我们还没有看到这个核函数的代码，但可以提前告诉你该核函数的作用是生成将要显示的图像数据。接下来，取消对共享资源的映射。一定要在执行绘制任务之前执行取消映射的调用，这是为了确保在应用程序的CUDA部分和图形部分之间实现同步。特别是，取消映射的调用将使得在cudaGraphicsUnmapResources()之前的所有CUDA操作完成之后，才会开始执行图形调用。

最后，我们通过GLUT注册键盘回调函数和显示回调函数（key_func 和 draw_func），并通过glutMainLoop()将执行控制交给GLUT绘制循环。

dim3 grids(DIM/16, DIM/16);  
dim3 threads(16, 16);  
kernel<<<grids, threads>>>(devPtr);  
HANDLE_ERROR(udaGraphicsUnmapResources(1, &resource, NULL));  
//设置好GLUT并启动循环  
glutKeyboardFunc(key_func);  
glutDisplayFunc draw_func);  
glutMainLoop();

这个应用程序的剩余部分包括三个函数，kernel()、key_func()、draw_func()。现在我们就来看这些函数。

核函数的参数包括一个设备指针，函数的任务是生成图像数据。在下面的示例中，我们将使用从第5章波纹示例中修改而来的核函数：

// 根据波纹代码修改而来，其中使用了uchar4类型，  
// 这是图形交互使用的数据类型  
__global__void kernel(uchar4 *ptr) {  
// 将threadIdx/BlockIdx映射到像素位置  
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));  
// 此时访问的uchar4类型而不是unsigned char*类型  
ptr[offset].x = 0;  
ptr[offset].y = green;  
ptr[offset].z = 0;  
ptr[offset].w = 255;

这里用到了许多熟悉的知识。例如，将线程索引和块索引转换为x坐标和y坐标的方法，以及线性化偏移的方法等，在前面已经分析过多次了。然后，我们执行一些计算来判断位于（x，y）位置上像素的颜色，并将这些值保存到内存中。我们再次使用CUDA C在GPU上生成一张图像。需要注意的是，这张图像随后将在不需要CPU介入的情况下直接交给OpenGL。另一方面，在第5章的波纹示例中，在GPU上生成图像的方式与这里的方式非常相似，但应用程序随后需要将缓冲区复制回CPU以便显示。

那么，如何通过OpenGL来绘制CUDA生成的缓冲区？好的，回顾在main()中执行的设置过程，你会发现以下函数：

glBindBuffer(GL_PIXEL_UNPACK_bufferARB, bufferObj);

这个调用将共享缓冲区绑定为一个像素源，OpenGL驱动程序随后会在所有对glDrawPixels()的调用中使用这个像素源。这意味着，我们需要调用glDrawPixels()来绘制CUDA C核函数生成的图像数据。因此，下面就是draw_func()需要执行的工作：

static void draw_func(void) {
glDrawPixels(DIM, DIM, GL_RGBA, GL_UNSIGNED_BYTE, 0);
glutSwapBuffers();
}

你可能已经发现glDrawPixels()的最后参数为一个缓冲区指针。如果没有任何缓冲区绑定为GL_PIXL_UNPACK_buffer_ARB源，那么OpenGL驱动程序将从这个缓冲区中进行复制。然而，由于数据已经位于GPU上，并且我们已经将共享缓冲区绑定为GL_PIXL_UNPACK_buffer_ARB源，因此最后一个参数将变成绑定缓冲区内的一个偏移。由于我们要绘制整个缓冲区，因此这个偏移值就是0。

示例代码的最后一部分看上去有些奇怪，这是因为我们决定为用户提供一种方法来退出应用程序。回调函数key_func()将响应Esc键，并将这个键作为释放内存并退出的信号：

static void key_func(unsigned char key, int x, int y) {
switch (key) {
    case 27:
        // 释放OpenGL和CUDA
        HANDLE_ERROR(udaGraphicsUnregisterResource( resource));
        glBindBuffer(GL_PIXL_UNPACKBUFFERARB, 0);
        glDeleteBuffers(1, &bufferObj);
        exit(0);
}

当运行这个示例时，将用“NVIDIA绿色（NVIDIA Green）”和黑色绘制一个具有催眠效果的图片，如图8.1所示。你可以尝试用这张图片对你的朋友（或者敌人）进行催眠。


图8.1 基于图像互操作性的催眠图片示例

上一章8.1_本章目标
下一章8.3_基于图形互操作性的GPU波纹示例