8.3_基于图形互操作性的GPU波纹示例

8.3 基于图形互操作性的GPU波纹示例

在8.2节中，我们多次提到了第5章的GPU波纹示例。在第5章的示例程序中创建了一个CPUAnimBitmap并将其传递给一个函数，每当需要生成一帧图像时都会调用这个函数。

int main(void) {  
    DataBlock data;  
    CPUAnimBitmap bitmap(DIM, DIM, &data);  
    data.bitmap = &bitmap;  
    HANDLE_ERROR(cudaMalloc((void**)&data.dev_bitmap, bitmap(image_size())));  
    bitmap anim_and_exit((void*)(void*, int)) generate_frame, (void*)(void*))cleanup());  
}

根据在前面章节中学到的技术，我们希望创建一个GPUAnimBitmap结构。这个结构的作用与CPUAnimBitmap相同，但在GPUAnimBitmap中，CUDA和OpenGL组件能够不需要CPU的介入而实现相互协作。当在应用程序中使用GPUAnimBitmap()时，main()函数将变得更加简单，如下所示：

int main(void) {
    GPUAnimBitmap bitmap(DIM, DIM, NULL);
    bitmap anim_and_exit(
        (void (*)(uchar4*, void*, int)) generate_frame, NULL);
}

GPUAnimBitmap使用了在8.2一节中介绍的相同函数调用。然而，现在这些调用将被放入GPUAnimBitmap结构中，从而使示例（或者包括你自己的应用程序）变得更为简洁。

8.3.1 GPUAnimBitmap结构

在GPUAnimBitmpa结构中有几个数据成员与8.2节中介绍的内容十分相似。

struct GPUAnimBitmap {
    GLuint bufferObj;
    CUDAGraphicsResource *resource;
    int width, height;
    void *dataBlock;
    void (*fAnim)(uchar4*, void*, int);
    void (*animExit)(void*);
    void (*clickDrag)(void*, int, int, int, int);
    int dragStartX, dragStartY;

我们知道OpenGL和CUDA运行时对于GPU缓冲区分别有着不同的名字，并且还知道在调用OpenGL或CUDA C时需要使用不同的名字。因此，在结构中同时保存了OpenGL的名字bufferObj和CUDA运行时名字的resource。此外，由于代码计算的是一张将要显示的位图图像，因此这张图像需要包含宽度和高度。

为了在使用GPUAnimBitmap时可以注册特定的回调事件，我们还保存了一个void*指针dataBlock指向用户的数据。然而，在GPUAnimBitmap中将不会访问这个数据，而只是将其传递给已注册的回调函数。用户注册的回调函数将保存在fAnim、animExit以及clickDrag中。在每次调用glutIdleFunc()时都将调用fAnim()，这个函数将负责生成在动画中绘制的图像数据。在动画退出时，将调用函数animExit()一次。用户应该在这个函数中实现当动画结束时需要执行的清理代码。最后还有一个可选函数clickDrag()，这个函数将响应用户的鼠标点击/拖曳等事件。如果用户注册了这个函数，那么每当按下、拖曳或者释放鼠标时，都将调用该函数。鼠标点击的初始位置保存在（dragStartX, dragStartY）中，当释放鼠标时，点击或拖曳等事件的起始位置和结束位置就会传递给用户。这可以用于实现令人印象深刻的交互式动画。

在初始化GPUAnimBitmap之后紧接着是在前面示例中相同的代码。在将参数保存到相应的结构成员中后，首先向CUDA运行时查询CUDA设备。

GPUAnimBitmap( int w, int h, void *d ) { width = w; height = h; dataBlock = d; clickDrag = NULL; // 首先找到一个CUDA设备并将其设置到图形互操作中udaDeviceProp prop; int dev; memset( &prop, 0, sizeof(udaDeviceProp )); prop major = 1; prop.minor = 0; HANDLE_ERROR(udaChooseDevice( &dev, &prop ));

在找到一个兼容的CUDA设备后，调用CUDA运行时中的CUDASetGLDevice()，并通知运行时将dev作为与OpenGL互操作的设备：

CUDASetGLDevice( dev );

由于接下来将使用GLUT来创建一个窗口绘制环境，因此需要初始化GLUT。然而，在这个过程中存在一个问题，因为glutInit()希望将命令行参数传递给窗口系统。由于我们不需要传递任何参数，因此可以指定0个命令行参数。然而，在某些版本的GLUT中存在一个错误，即当传递0个参数时应用程序会崩溃。因此，我们要欺骗GLUT，使其认为正在传递一个参数，这样就不会使程序崩溃。

int c=1;  
char *foo = "name";  
glutInit(&c, &foo);

接下来，继续像前面示例中一样初始化GLUT。我们创建了一个要绘制的窗口，并将窗口

的标题指定为字符串“ bitmap”。当然，你也可以选择其他的名字作为窗口标题。

glutInitDisplayMode(GLUT_DOUBLE | GLUT_RGBA);  
glutInitWindowSize(width, height);  
glutCreateWindow(" bitmap");

接下来，我们请求OpenGL驱动程序分配一个缓冲区句柄，并将其绑定到GL_PIXL_UNPACK_buffer_ARB，从而确保之后在调用glDrawPixels()时将绘制互操作缓冲区：

glGenBuffers(1, &bufferObj);  
glBindBuffer(GL_PIXL_UNPACK_buffer_arB, bufferObj);

最后，我们请求OpenGL驱动程序分配一块GPU内存。在完成分配操作后，将这个缓冲区通知给CUDA运行时，并通过cudaGraphicsGLRegisterBuffer()来注册bufferObj从而获得一个CUDA C名字。

glBufferData(GL_PIXUNPACK_BUFFERARB, width \* height \* 4, NULL, GL_DYNAMIC_DRAWARB);   
HANDLE_ERROR(udaGraphicsGLRegisterBuffer( &resource, bufferObj,udaGraphicsMapFlagsNone)）;

在设置好GPUAnimBitmap后，还剩下一个问题就是如何执行绘制操作。绘制操作的主要工作是在函数glutIdleFunction()中完成的。这个函数要做三件事情。首先，它将映射共享缓冲区并获得指向该缓冲区的GPU指针。

// 作为GLUT回调函数的静态方法

static void idle_func(void) {
    static int ticks = 1;
    GPUAnimBitmap* bitmap = *(get Bitmap_ptr());
    uchar4* devPtr;
    size_t size;
    HANDLE_ERROR(
        deductaGraphicsMapResources(1, & ( bitmap->resource), NULL);
    );
    HANDLE_ERROR(
        deductaGraphicsResourceGetMappedPointer((void**)&devPtr,
            &size,
            bitmap->resource);
    );

其次，它调用用户指定的函数fAnim()，这个函数可能启动一个CUDA C核函数并用图像数据来填充devPtr指向缓冲区。

bitmap->fAnim( devPtr, bitmap->dataBlock, ticks++);

最后，它将取消GPU指针的映射，这会释放OpenGL驱动程序使用的缓冲区。绘制图像的操作将在调用glutPostRedisplay()时触发。

HANDLE_ERROR(udaGraphicsUnmapResources(1, & ( bitmap->resource), NULL)); glutPostRedisplay();

GPUAnimBitmap结构中还包含了一些重要但却有些不太相关的基础代码。如果你对这些代码感兴趣，可以对其进行分析。然而，即使没有时间或者兴趣来理解GPUAnimBitmap的剩余代码，也没关系，你仍然能够理解接下来的内容。

8.3.2 重新实现基于GPU的波纹动画示例

现在，我们已经实现了一个GPU版本的CPUAnimBitmap，接下来可以继续将基于GPU的波纹应用程序修改为完全在GPU上执行动画。首先要包含gpu anim.h头文件，其中包含了GPUAnimBitmap的实现。我们还包含了与第5章中几乎相同的核函数。

include"../common/book.h" #include“../common/gpu_anim.h"   
#define DIM1024   
__global__void kernel(uchar4 \*ptr,int ticks）{ //将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  $\mathrm{fx} = \mathrm{x} - \mathrm{DIM} / 2;$  float  $\mathrm{fy} = \mathrm{y} - \mathrm{DIM} / 2;$  floatd  $=$  sqrtf(fx\*fx+fy\*fy); unsigned char grey  $=$  (unsigned char)(128.0f  $^+$  127.0f \* cos(d/10.0f-ticks/7.0f)/
(d/10.0f + 1.0f));  
ptr[offset].x = grey;  
ptr[offset].y = grey;  
ptr[offset].z = grey;  
ptr[offset].w = 255;

上面代码中的粗体部分是我们唯一进行的修改。这个修改的原因是因为OpenGL互操作要求共享内存满足图形操作的需求。由于在绘制过程中，每个元素通常包含4部分信息（红/绿/蓝/Alpha），因此目标缓冲区将不再像以前那样只是一个unsigned char数组。现在要求是一个uchar4类型的数组。在第5章中，我们将缓冲区的每四个字节作为一个数据，因此通过ptr[offset*4+k]对进行索引，其中k的取值范围是0到3。但现在，uchar4类型很好表达了在数据中包含四部分信息的语义。

由于CUDA函数kernel()将生成图像数据，因此剩下的工作就是编写一个主机函数，并将其作为回调函数保存在GPUAnimBitmap的成员idle_func()中。对于这里的应用程序，该函数执行的操作就是启动CUDA C核函数：

void generate_frame( uchar4 *pixels, void*, int ticks ) {
    dim3 grids(DIM/16, DIM/16);
    dim3 threads(16, 16);
    kernel<<<grids, threads>>>(pixels, ticks);
}

由于所有的辅助工作都是在GPUAnimBitmap结构中完成的，因此这个函数的代码很简洁。然后，我们只需创建一个GPUAnimBitmap，并且注册动画回调函数generate_frame()。

int main(void) {
    GPUAnimBitmap bitmap(DIM, DIM, NULL);
    bitmap anim_and_exit(
        (void*) (uchar4*, void*, int)) generate_frame, NULL);
}
前の章8.2_图形互操作
次の章8.4_基于图形互操作性的热传导