CUDA编程之CUDA Sample-5_Domain_Specific-simpleGL

翊桐公羽

于 2024-08-27 17:07:43 发布

阅读量603

点赞数 17

文章标签： c++

本文链接：https://blog.csdn.net/zhujingwen_tongji/article/details/141469890

版权

OpenGL

OpenGL，全称为"Open Graphics Library"，是一种跨语言和跨平台的 API,用于显示 2D 和 3D 计算机图形。

OpenGL 可以在硬件级别实现, NVIDIA 公司提供的 OpenGL 实现是随 CUDA 驱动程序一起提供的。

GLUT 是 OpenGL 实用程序工具包,一种独立于窗口系统的工具包,用于编写 OpenGL 程序。它实现了一个简单的窗口应用程序编程接口(API)来使用 OpenGL。GLUT 大大简化了学习和探索OpenGL 编程的过程。

该工具包主要包括以下功能:

窗口任务:创建、销毁、改变大小等。
回调函数:处理用户事件,如鼠标和键盘输入。
鼠标和键盘:处理鼠标和键盘输入事件。

SimpleGL

SimpleGL 演示了 CUDA 和 OpenGL 之间的互操作性。该程序使用 CUDA 修改顶点位置,并使用 OpenGL 渲染几何图形。

具体来说:

该程序展示了如何在 CUDA 和 OpenGL 之间进行数据交互和共享。
CUDA 负责计算和修改几何体的顶点位置数据。
OpenGL 则负责使用这些更新后的顶点数据进行图形渲染,在屏幕上显示出几何图形。
这种 CUDA 与 OpenGL 的无缝集成,使得开发者可以利用 CUDA 的高性能计算能力,同时又可以使用 OpenGL 的图形渲染功能,实现各种复杂的图形应用。

具体实现过程如下：

首先创建一个 OpenGL 中的顶点缓冲对象(Vertex Buffer Object, VBO),用于存储顶点数据。这是一个空的 VBO,还没有填充任何数据。
将这个 VBO 在 CUDA 中进行注册,这样 CUDA 就可以访问和操作这个 VBO 中的数据。
接下来将 VBO 映射到 CUDA 的内存空间中,这样就可以直接在 CUDA 中修改 VBO 中的顶点数据。
在 CUDA 内核中,对这些顶点数据进行计算和修改,比如改变顶点的位置坐标。
修改完成后,取消 VBO 与 CUDA 内存的映射关系。
最后,使用 OpenGL 的渲染流程,根据修改后的 VBO 数据来绘制图形,显示在屏幕上。

InitGL代码详解

bool initGL(int *argc, char **argv)
{
    glutInit(argc, argv);
    glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE);
    glutInitWindowSize(window_width, window_height);
    glutCreateWindow("Cuda GL Interop (VBO)");
    glutDisplayFunc(display);
    glutKeyboardFunc(keyboard);
    glutMotionFunc(motion);
    glutTimerFunc(REFRESH_DELAY, timerEvent,0);

    // initialize necessary OpenGL extensions
    if (! isGLVersionSupported(2,0))
    {
        fprintf(stderr, "ERROR: Support for necessary OpenGL extensions missing.");
        fflush(stderr);
        return false;
    }

    // default initialization
    glClearColor(0.0, 0.0, 0.0, 1.0);
    glDisable(GL_DEPTH_TEST);

    // viewport
    glViewport(0, 0, window_width, window_height);

    // projection
    glMatrixMode(GL_PROJECTION);
    glLoadIdentity();
    gluPerspective(60.0, (GLfloat)window_width / (GLfloat) window_height, 0.1, 10.0);

    SDK_CHECK_ERROR_GL();

    return true;
}

初始化 GLUT 库并设置显示模式为 RGBA 颜色和双缓冲。
创建一个指定大小和标题的 OpenGL 窗口。
注册各种回调函数,如渲染、键盘、鼠标等事件的处理函数。
检查是否支持所需的 OpenGL 扩展。
设置默认的 OpenGL 状态,如背景颜色、深度测试等。
设置视口和投影矩阵。
检查 OpenGL 错误并返回初始化结果。

这个函数为后续使用 CUDA 和 OpenGL 进行交互奠定了基础,为应用程序提供了一个可用的 OpenGL 渲染环境。

createVBO

void createVBO(GLuint *vbo, struct cudaGraphicsResource **vbo_res,
               unsigned int vbo_res_flags)
{
    assert(vbo);

    // create buffer object
    glGenBuffers(1, vbo);
    glBindBuffer(GL_ARRAY_BUFFER, *vbo);

    // initialize buffer object
    unsigned int size = mesh_width * mesh_height * 4 * sizeof(float);
    glBufferData(GL_ARRAY_BUFFER, size, 0, GL_DYNAMIC_DRAW);

    glBindBuffer(GL_ARRAY_BUFFER, 0);

    // register this buffer object with CUDA
    checkCudaErrors(cudaGraphicsGLRegisterBuffer(vbo_res, *vbo, vbo_res_flags));

    SDK_CHECK_ERROR_GL();
}

该杉树创建了一个 OpenGL 的顶点缓冲对象(VBO)并将其与 CUDA 进行注册,以便 CUDA 可以访问和操作这个 VBO 中的数据。具体的步骤:

首先断言传入的 vbo 指针不能为空。
使用 glGenBuffers 函数创建一个新的 VBO,并将其绑定到 GL_ARRAY_BUFFER 目标。
使用 glBufferData 函数初始化这个 VBO,分配了 mesh_width * mesh_height * 4 * sizeof(float) 字节的内存空间,并将其标记为 GL_DYNAMIC_DRAW(表示这个缓冲区的内容会经常被修改)。
解绑 GL_ARRAY_BUFFER 目标。
接下来,使用 cudaGraphicsGLRegisterBuffer 函数将这个 VBO 在 CUDA 中进行注册,并将注册标志设置为 vbo_res_flags。这样 CUDA 就可以访问和操作这个 VBO 中的数据了。
最后,检查是否有 OpenGL 错误发生。

这段代码创建了一个 OpenGL VBO,并将其与 CUDA 进行了互操作,为后续使用 CUDA 计算和更新顶点数据,以及使用 OpenGL 进行渲染做好了准备工作。这种 CUDA 和 OpenGL 的无缝集成,是实现高性能图形应用的关键技术之一。

runCuda

oid runCuda(struct cudaGraphicsResource **vbo_resource)
{
    // map OpenGL buffer object for writing from CUDA
    float4 *dptr;
    checkCudaErrors(cudaGraphicsMapResources(1, vbo_resource, 0));
    size_t num_bytes;
    checkCudaErrors(cudaGraphicsResourceGetMappedPointer((void **)&dptr, &num_bytes,
                                                         *vbo_resource));
    //printf("CUDA mapped VBO: May access %ld bytes\n", num_bytes);

    // execute the kernel
    //    dim3 block(8, 8, 1);
    //    dim3 grid(mesh_width / block.x, mesh_height / block.y, 1);
    //    kernel<<< grid, block>>>(dptr, mesh_width, mesh_height, g_fAnim);

    launch_kernel(dptr, mesh_width, mesh_height, g_fAnim);

    // unmap buffer object
    checkCudaErrors(cudaGraphicsUnmapResources(1, vbo_resource, 0));
}

这段代码使用 CUDA 来修改 OpenGL 顶点缓冲对象(VBO)中的顶点数据。具体步骤如下:

首先使用 cudaGraphicsMapResources 函数将 VBO 映射到 CUDA 的内存地址空间中。这样 CUDA 就可以直接访问和修改 VBO 中的数据。
获取映射后的 CUDA 内存指针 dptr 以及映射的字节数 num_bytes。
调用 launch_kernel 函数,传入 dptr 指针、mesh_width、mesh_height 和 g_fAnim 等参数,在 CUDA 中执行核函数来修改顶点数据。它会对 VBO 中的顶点位置进行一些计算和更新。
最后,使用 cudaGraphicsUnmapResources 函数取消 VBO 与 CUDA 内存的映射关系,这样 CUDA 就无法再访问 VBO 中的数据了。

这个过程展示了 CUDA 和 OpenGL 之间的无缝集成。CUDA 可以直接访问并修改 OpenGL 的 VBO 数据,这样可以充分利用 CUDA 的并行计算能力来加速图形渲染任务。之后 OpenGL 就可以使用更新后的 VBO 数据进行渲染,从而实现高性能的图形应用。

Kernel详解

void launch_kernel(float4 *pos, unsigned int mesh_width,
                   unsigned int mesh_height, float time)
{
    // execute the kernel
    dim3 block(8, 8, 1);
    dim3 grid(mesh_width / block.x, mesh_height / block.y, 1);
    simple_vbo_kernel<<< grid, block>>>(pos, mesh_width, mesh_height, time);
}

launch_kernel 函数的作用是在 CUDA 上启动一个内核函数,该内核函数会对传入的 OpenGL VBO 中的顶点数据进行计算和更新。具体步骤如下:

首先定义了 CUDA 内核函数的块大小和网格大小:
- 块大小为 (8, 8, 1)
- 网格大小为 (mesh_width / 8, mesh_height / 8, 1)
这样可以确保内核函数在 CUDA 上并行执行,并且每个块处理的顶点数据量适中。
然后调用了一个名为 simple_vbo_kernel 的 CUDA 内核函数,并传入了以下参数:
- pos: 指向要更新的 VBO 顶点数据的指针
- mesh_width: 网格的宽度
- mesh_height: 网格的高度
- time: 一个用于计算顶点位置的时间变量

launch_kernel 函数负责设置 CUDA 内核的执行配置,并启动内核函数来更新 VBO 中的顶点数据。这种结合 CUDA 计算能力和 OpenGL 渲染能力的方式,可以实现高性能的图形应用程序。

__global__ void simple_vbo_kernel(float4 *pos, unsigned int width, unsigned int height, float time)
{
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;

    // calculate uv coordinates
    float u = x / (float) width;
    float v = y / (float) height;
    u = u*2.0f - 1.0f;
    v = v*2.0f - 1.0f;

    // calculate simple sine wave pattern
    float freq = 4.0f;
    float w = sinf(u*freq + time) * cosf(v*freq + time) * 0.5f;

    // write output vertex
    pos[y*width+x] = make_float4(u, w, v, 1.0f);
}

simple_vbo_kernel 内核函数,它的作用是计算并更新 OpenGL VBO 中顶点的位置。下面是具体的实现步骤:

首先根据当前线程的 blockIdx 和 threadIdx 计算出当前线程在整个网格中的坐标 (x, y)。
然后根据网格大小,计算出当前顶点的 UV 坐标 (u, v)。这里将 UV 坐标从 [0, 1] 范围映射到 [-1, 1] 范围。
接下来计算一个简单的正弦波模式,频率为 4.0f。这个计算结果 w 会用于修改顶点的 Y 坐标,从而实现一个简单的波浪动画效果。
最后,使用 make_float4 函数构造一个新的顶点位置 (u, w, v, 1.0f),并将其写入到 pos 数组中对应的位置。这个 pos 数组就是之前映射到 CUDA 的 OpenGL VBO 的数据。

总的来说,这个内核函数会遍历整个网格,为每个顶点计算出新的位置,从而实现一个简单的波浪动画效果。这种结合 CUDA 并行计算和 OpenGL 渲染的方式,可以非常高效地实现各种复杂的图形效果。

display函数详解

void display()
{
    sdkStartTimer(&timer);

    // run CUDA kernel to generate vertex positions
    runCuda(&cuda_vbo_resource);

    glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);

    // set view matrix
    glMatrixMode(GL_MODELVIEW);
    glLoadIdentity();
    glTranslatef(0.0, 0.0, translate_z);
    glRotatef(rotate_x, 1.0, 0.0, 0.0);
    glRotatef(rotate_y, 0.0, 1.0, 0.0);

    // render from the vbo
    glBindBuffer(GL_ARRAY_BUFFER, vbo);
    glVertexPointer(4, GL_FLOAT, 0, 0);

    glEnableClientState(GL_VERTEX_ARRAY);
    glColor3f(1.0, 0.0, 0.0);
    glDrawArrays(GL_POINTS, 0, mesh_width * mesh_height);
    glDisableClientState(GL_VERTEX_ARRAY);

    glutSwapBuffers();

    g_fAnim += 0.01f;

    sdkStopTimer(&timer);
    computeFPS();
}

display() 函数是OpenGL 的渲染函数,它的主要作用是:

启动 CUDA 计时器 sdkStartTimer(&timer);
调用 runCuda(&cuda_vbo_resource) 函数,在 CUDA 中计算并更新 VBO 中的顶点数据;
清除颜色和深度缓冲区 glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
设置模型视图矩阵,包括平移、旋转等变换;
绑定 VBO 并设置顶点指针 glBindBuffer(GL_ARRAY_BUFFER, vbo); glVertexPointer(4, GL_FLOAT, 0, 0);
启用顶点数组 glEnableClientState(GL_VERTEX_ARRAY);
使用 glDrawArrays(GL_POINTS, 0, mesh_width * mesh_height) 绘制所有顶点;
禁用顶点数组 glDisableClientState(GL_VERTEX_ARRAY);
交换前后缓冲区 glutSwapBuffers();
递增时间变量 g_fAnim += 0.01f;
停止 CUDA 计时器并计算帧率 sdkStopTimer(&timer); computeFPS();

这个函数的主要流程是:

首先更新 CUDA 中的顶点数据;
然后设置好 OpenGL 的渲染状态;
绑定 VBO 并使用 glDrawArrays 绘制顶点;
最后交换缓冲区,更新时间变量,并计算帧率。

这样可以实现一个基于 CUDA 和 OpenGL 的动画渲染效果。CUDA 负责高性能的顶点计算,而 OpenGL 负责将计算结果渲染到屏幕上,两者相结合可以实现非常流畅的图形效果。