详解cuda by example中第八章图形互操作性代码

这两天学到第八章了，发现第八章内容有部分是关于图形学的内容于是就把我对第八章关于图形互操作的代码理解也写一下，仅供学习，如果大家有发现不对的地方欢迎指正。（文章代码顺序按照官方示例代码顺序给出，可按顺序食用）

1.引入头文件

#include "../common/book.h"
#include "../common/cpu_bitmap.h"

#include "cuda.h"
#include "cuda_gl_interop.h"

以下设置了一些指向OpenGL函数的函数指针类型。这些函数指针用于动态加载OpenGL扩展函数。（就是用来加载一些函数，可以理解为加载了4个函数）

PFNGLBINDBUFFERARBPROC    glBindBuffer = NULL;
PFNGLDELETEBUFFERSARBPROC glDeleteBuffers = NULL;
PFNGLGENBUFFERSARBPROC    glGenBuffers = NULL;
PFNGLBUFFERDATAARBPROC    glBufferData = NULL;

定义了维度大小（512）以及两个独立变量（他们指向的是同一个缓冲区，但是OpenGL与cuda对对这个缓冲区有不同的名字所以用两个变量）

#define     DIM    512

GLuint  bufferObj;
cudaGraphicsResource* resource;

2.核函数

关于offset的解释这里不在赘述，有需要了解的可以参考这篇文章详解cuda by example中第七章纹理内存（热量图）代码-CSDN博客
剩下的代码是对传入的数组进行赋值（实际上是对每个像素点的颜色进行赋值分别有人r,g,b透明度，这四个值进行赋值）

__global__ void kernel(uchar4* ptr) {
    // map from threadIdx/BlockIdx to pixel position
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;

    // now calculate the value at that position
    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;
}

3. key_func函数

这个可以理解为是一个退出函数，key参数表示按下的键的 ASCII 码值。当按下esc键时key=27，此时执行退出程序

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

4.draw_func函数

用来画像素点

static void draw_func(void) {
    // we pass zero as the last parameter, because out bufferObj is now
    // the source, and the field switches from being a pointer to a
    // bitmap to now mean an offset into a bitmap object
    glDrawPixels(DIM, DIM, GL_RGBA, GL_UNSIGNED_BYTE, 0);
    glutSwapBuffers();
}

5.main函数

1.指定cuda设备

cudaDeviceProp  prop;
int dev;

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

2.初始化显示窗口

    glutInit(&argc, argv);
    glutInitDisplayMode(GLUT_DOUBLE | GLUT_RGBA);
    glutInitWindowSize(DIM, DIM);
    glutCreateWindow("bitmap");

3.动态加载函数

四个函数的作用：

glBindBuffer(GLenum target, GLuint buffer)：它的作用是将指定的缓冲区对象绑定到目标缓冲区上，使得后续的缓冲区操作都针对这个绑定的缓冲区对象

glDeleteBuffers(GLsizei n, const GLuint *buffers)：n: 要删除的缓冲区对象的数量，buffers: 指向一个数组，其中包含要删除的缓冲区对象的名称（ID）

glGenBuffers(GLsizei n, GLuint *buffers)：n: 指定要生成的缓冲区对象的数量，buffers: 指向一个数组，该数组将被填充为缓冲区对象的 ID（名称）。这些 ID 是由 OpenGL 生成的

glBufferData(GLenum target, GLsizeiptr size, const void *data, GLenum usage)：size: 以字节为单位的缓冲区大小，即要分配的内存大小；data: 指向数据的指针；usage: 表示应用程序将如何使用这个数据的预期模式，target: 指定缓冲区对象的目标

    glBindBuffer = (PFNGLBINDBUFFERARBPROC)GET_PROC_ADDRESS("glBindBuffer");
    glDeleteBuffers = (PFNGLDELETEBUFFERSARBPROC)GET_PROC_ADDRESS("glDeleteBuffers");
    glGenBuffers = (PFNGLGENBUFFERSARBPROC)GET_PROC_ADDRESS("glGenBuffers");
    glBufferData = (PFNGLBUFFERDATAARBPROC)GET_PROC_ADDRESS("glBufferData");

4.调用函数实列

以下是函数的使用时列

表示分配1个缓冲区对象，并将它的 ID 存储在 bufferObj中

glGenBuffers(1, &bufferObj);

绑定buffer数组，将 bufferObj 对象绑定到 GL_PIXEL_UNPACK_BUFFER_ARB 目标上，L_PIXEL_UNPACK_BUFFER_ARB: 这个枚举值指定了缓冲区对象的目标类型。在这个例子中，它表示像素解包缓冲区（Pixel Unpack Buffer），这类缓冲区用于从 CPU 向 GPU 传输像素数据

glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, bufferObj);

这行代码是在 OpenGL 中用于为当前绑定的缓冲区对象分配内存空间并指定其用法，GL_PIXEL_UNPACK_BUFFER_ARB表示从cpu向gpu传输数据，大小为DIM * DIM * 4（4表示4 个字节），GL_DYNAMIC_DRAW_ARB 表示数据将被频繁更新（通过写操作），但很少被读取。

glBufferData(GL_PIXEL_UNPACK_BUFFER_ARB, DIM * DIM * 4,NULL, GL_DYNAMIC_DRAW_ARB)

5.注册函数

这是 CUDA 提供的一个函数，用于注册 OpenGL 缓冲区对象以进行 CUDA 访问。它的作用是将一个 OpenGL 缓冲区与 CUDA 图形资源关联，从而可以在 CUDA 代码中使用这个缓冲区

HANDLE_ERROR(cudaGraphicsGLRegisterBuffer(&resource,bufferObj,cudaGraphicsMapFlagsNone));

cudaGraphicsMapResources: 这是 CUDA 提供的一个函数，用于映射一个或多个图形资源，以便 CUDA 内核访问这些资源。1: 表示映射的资源数量。这里是 1 个资源。&resource: 指向需要映射的 cudaGraphicsResource 资源的指针。调用 cudaGraphicsMapResources 后，CUDA 将可以访问这个 OpenGL 资源。

cudaGraphicsMapResources(1, &resource, NULL)

用于获取一个映射到 CUDA 的图形资源的设备指针和大小。这允许 CUDA 直接访问图形资源的数据。&size: 这是传递给函数的第二个参数，用于存储资源的大小，resource: 这是之前注册的 cudaGraphicsResource 资源句柄。cudaGraphicsResourceGetMappedPointer 获取到指向资源数据的指针（devPtr）和资源的大小（size）。 有了 devPtr 后，可以在 CUDA 内核中直接对OpenGL 资源中的数据进行读取和修改操作。

uchar4* devPtr;
    size_t  size;
    HANDLE_ERROR(cudaGraphicsResourceGetMappedPointer((void**)&devPtr,&size,resource));

6.执行函数，释放资源

实例化线程块与线程，执行核函数执行完后解除cuda与opengl资源的映射


后面的代码就不详说了跟前几章的代码都差不多，这篇文章大体也差不多了，如果有错误还请大家指正。