基于纹理缓冲实现OpenGL和CUDA的交互

2023-10-30

基于缓冲对象的交互

  要实现OpenGL和CUDA交互,最常用便捷的方式就是,在OpenGL中创建缓冲对象,将其注册并绑定到一个内存指针,将这个指针传入CUDA核函数中进行读写。关于这点,可以参考笔者之前的文章------基于C++与CUDA的N卡GPU并行程序——OpenGL图形互操作性

// 创建窗口缓冲
int c=1;
char *dummy;
glutInit( &c, &dummy );
glutInitDisplayMode( GLUT_SINGLE | GLUT_RGBA );
glutInitWindowSize( width, height );
glutCreateWindow( "bitmap" );

// 创建缓冲对象
GLuint  bufferObj;
cudaGraphicsResource *resource;
glGenBuffers( 1, &bufferObj );
glBindBuffer( GL_PIXEL_UNPACK_BUFFER_ARB, bufferObj );
glBufferData( GL_PIXEL_UNPACK_BUFFER_ARB, width * height * 4,
              NULL, GL_DYNAMIC_DRAW_ARB );
cudaGraphicsGLRegisterBuffer( &resource, bufferObj, cudaGraphicsMapFlagsWriteDiscard );

// 每次循环,调用idle_func()函数,使用CUDA核函数计算
// static method used for glut callbacks
static void idle_func( void ) {
        GPUAnimBitmap*  bitmap = *(get_bitmap_ptr());
        bitmap->ticks++;
        int time = glutGet(GLUT_ELAPSED_TIME);
        if(time - bitmap->timebase > 1000){
            printf("帧率为: %4.2f\n", float(bitmap->ticks) * 1000.0 / float(time - bitmap->timebase));
            bitmap->timebase = time;
            bitmap->ticks = 1;
        }
        uchar4* devPtr;
        size_t size;
        // 将缓冲对象绑定到内存指针
        cudaGraphicsMapResources( 1, &(bitmap->resource), NULL );
        cudaGraphicsResourceGetMappedPointer( (void**)&devPtr, &size, bitmap->resource);
        // 调用CUDA核函数
        bitmap->fAnim( devPtr, bitmap->dataBlock, bitmap->render);//-------------------------
        // 解除缓冲对象
        cudaGraphicsUnmapResources( 1, &(bitmap->resource), NULL );
        glutPostRedisplay();
    }

// 每次循环,调用Draw()函数,绘制画面
// static method used for glut callbacks
static void Draw( void ) {
    GPUAnimBitmap*   bitmap = *(get_bitmap_ptr());
    glDrawPixels( bitmap->width, bitmap->height, GL_RGBA,
                  GL_UNSIGNED_BYTE, 0 );
    glFlush();
}

// 绘制主函数
void anim_and_exit( void (*f)(uchar4*,void*,void*), void(*e)(void*) ) {//----------------
        GPUAnimBitmap**   bitmap = get_bitmap_ptr();
        *bitmap = this;
        fAnim = f;
        animExit = e;
        glutDisplayFunc( Draw );
        glutIdleFunc( idle_func );
        glutMainLoop();
    }

在这个链接的文章中,CUDA内核计算的结果会写到缓冲对象中,再使用glDrawPixels()绘制到屏幕上。但是这个glDrawPixels()函数有点太慢了。

基于纹理缓冲的交互(主机内存传输)

  glDrawPixels()函数太慢了,一般都不太会用了,可以使用速度更快的glTextureSubImage2D()函数。这里不再是创建缓冲对象,而是创建一个纹理缓冲,然后在着色器shader中,把这个纹理缓冲的内容,贴到屏幕上进行绘制。首先在主机内存上开辟空间,将纹理贴图的图片数据,保存到这个内存上,然后把这个纹理数据传到OpenGL的纹理缓冲上

// 主机内存上保存纹理贴图
uchar4 *textureImage = (uchar4*)malloc(width*height*sizeof(uchar4));
for(int i = 0; i < height; i++){
    for(int j = 0; j < width/2; j++){
        textureImage[i*width+j].x = (unsigned char)(100);
        textureImage[i*width+j].y = (unsigned char)(1);
        textureImage[i*width+j].z = (unsigned char)(0);
    }
}

// 创建纹理缓冲,绑定纹理
glGenTextures(1, &texture);
glActiveTexture(GL_TEXTURE0);
glBindTexture(GL_TEXTURE_2D, texture);
GLsizei imageWidth = width, imageHeight = height;
glTexImage2D(GL_TEXTURE_2D, 0, 4, imageWidth, imageHeight, 0, GL_RGBA, GL_UNSIGNED_BYTE, textureImage);
glEnable(GL_TEXTURE_2D);//启用纹理贴图
glGenerateTextureMipmap(texture);

// 每次循环,调用Draw()函数,绘制画面
// static method used for glut callbacks
static void Draw( void ) {
    GPUAnimBitmap*   bitmap = *(get_bitmap_ptr());
    glTextureSubImage2D(bitmap->texture, 0, 0, 0, bitmap->width, bitmap->height, GL_RGBA8, GL_UNSIGNED_BYTE, bitmap->textureImage);
    // 绘制一个四边形片段,四边形正好就是这个这个屏幕上的边框
    // 使用glTexCoord2f()定义纹理贴图的坐标
    glBegin(GL_QUADS);
    glTexCoord2f(0.0f,0.0f);
    glVertex2f(-1.0f,-1.0f);
    glTexCoord2f(1.0f,0.0f);
    glVertex2f(1.0f,-1.0f);
    glTexCoord2f(1.0f,1.0f);
    glVertex2f(1.0f,1.0f);
    glTexCoord2f(0.0f,1.0f);
    glVertex2f(-1.0f,1.0f);
    glEnd();
    glFlush();
}

在Draw()中绘制一个正好覆盖到屏幕边框的四边形,使用glTexCoord2f()定义纹理贴图的坐标,这样就可以把这张纹理贴图绘制到窗口。

基于纹理缓冲的交互(显存交互)

  以上介绍的是使用glTextureSubImage2D()函数,将主机内存的图像数据,传输到纹理缓冲。但是从CPU主机内存传输数据到GPU显存上,需要使用DMA引擎经过PCIE总线进行访问,这个速度显然时比较慢的。而且CPU主机内存上的数据,无法直接用CUDA核函数进行计算和交互。
  所以真正实用的一种方式,是将上面所述的两种方式结合起来,实现OpenGL和CUDA的快速交互。先使用OpenGL在显存上创建一个缓冲对象,将这个缓冲对象注册绑定到内存指针,之后CUDA核函数就可以对这个缓冲对象的内存进行读写交互。最后创建绑定一个纹理缓冲,使用glTextureSubImage2D()函数,将显存上的缓冲对象数据传输到纹理缓冲中,再显示到屏幕上。因为需要传输的数据本身就位于显卡上,不需要从CPU传过去,速度就更快了,而且可以直接使用CUDA核函数进行计算。

// 创建缓冲对象
GLuint  bufferObj;
cudaGraphicsResource *resource;
glGenBuffers( 1, &bufferObj );
glBindBuffer( GL_PIXEL_UNPACK_BUFFER_ARB, bufferObj );
glBufferData( GL_PIXEL_UNPACK_BUFFER_ARB, width * height * 4,
              NULL, GL_DYNAMIC_DRAW_ARB );
cudaGraphicsGLRegisterBuffer( &resource, bufferObj, cudaGraphicsMapFlagsWriteDiscard );

// 每次循环,调用Draw()函数,绘制画面
// static method used for glut callbacks
static void Draw( void ) {
    GPUAnimBitmap*   bitmap = *(get_bitmap_ptr());
    glTextureSubImage2D(bitmap->texture, 0, 0, 0, bitmap->width, bitmap->height, GL_RGBA8, GL_UNSIGNED_BYTE, NULL);
    // 绘制一个四边形片段,四边形正好就是这个这个屏幕上的边框
    // 使用glTexCoord2f()定义纹理贴图的坐标
    glBegin(GL_QUADS);
    glTexCoord2f(0.0f,0.0f);
    glVertex2f(-1.0f,-1.0f);
    glTexCoord2f(1.0f,0.0f);
    glVertex2f(1.0f,-1.0f);
    glTexCoord2f(1.0f,1.0f);
    glVertex2f(1.0f,1.0f);
    glTexCoord2f(0.0f,1.0f);
    glVertex2f(-1.0f,1.0f);
    glEnd();
    glFlush();
}
从缓存加载纹理

  glTextureSubImage2D()的data参数可以通过两种方式来解释。第一种方式是通过用户程序中存储的自然数据指针解释。第二种data数据的解释方式,则是通过绑定到GL_PIXEL_UNPACK_BUFFER目标的缓存对象来完成,作为缓存对象的偏移位置。用户程序此时可以将数据存储到缓存对象当中,然后再传递到纹理对象里。如果GL_PIXEL_UNPACK_BUFFER目标没有绑定任何缓存对象,那么data会被解释成一个本地指针,如果绑定了缓存,那么data会被解释成缓存中的一个偏移位置。
  使用缓存对象来存储纹理数据的一个主要的好处在于:数据不是立即从缓存对象向纹理进行传输的,而是在着色器请求数据的时候才会执行这一操作。因此应用程序的运行和数据的传输操作可以并行进行。如果数据在应用程序本地内存中,那么glTextureSubImage2D()需要先对数据进行拷贝,然后函数才会返回,这是不可能并行完成的。不过这种方法的好处在于:应用程序在函数返回之后依然可以自由修改之前传输的data数据。

FBO帧缓存对象

  除了上述的方法之外,另一种方式是使用帧缓存对象FBO的一个缓存附件——渲染缓存,使用glBlitFramebuffer()来完成。
相关资料有
网址1
网址2
网址3
网址4
网址5

本文内容由网友自发贡献,版权归原作者所有,本站不承担相应法律责任。如您发现有涉嫌抄袭侵权的内容,请联系:hwhale#tublm.com(使用前将#替换为@)

基于纹理缓冲实现OpenGL和CUDA的交互 的相关文章

  • 将 GPUJPEG 项目移植到 Windows

    我目前正在尝试移植 GPUJPEG 在 Sourceforge 上 http sourceforge net projects gpujpeg 库 基于 CUDA 从 Unix 到 Windows 现在我被卡住了 我不知道发生了什么或为什么
  • 如何在 Windows 上的 nvidia GPU 的 Visual Studio 2010 中配置 OpenCL?

    我在华硕笔记本电脑上的 Wwindows 7 操作系统上使用 NVIDIA GeForce GTX 480 GPU 我已经为 CUDA 4 2 配置了 Visual Studio 2010 如何在 Visual Studio 2010 上为
  • 使用内置显卡,没有NVIDIA显卡,可以使用CUDA和Caffe库吗?

    使用内置显卡 没有 NVIDIA 显卡 可以使用 CUDA 和 Caffe 库吗 我的操作系统是 ubuntu 15 CPU为 Intel i5 4670 3 40GHz 4核 内存为12 0GB 我想开始学习深度学习 CUDA 适用于 N
  • 某些子网格未使用 CUDA 动态并行执行

    我正在尝试 CUDA 5 0 GTK 110 中的新动态并行功能 我遇到了一个奇怪的行为 即我的程序没有返回某些配置的预期结果 不仅是意外的 而且每次启动都会出现不同的结果 现在我想我找到了问题的根源 似乎当生成太多子网格时 某些子网格 由
  • 如何用Go语言的cgo编译Cuda源码?

    我用 cuda c 编写了一个简单的程序 它可以在 eclipse nsight 上运行 这是源代码 include
  • CUDA:如何在设备上填充动态大小的向量并将其内容返回到另一个设备函数?

    我想知道哪种技术可以填充设备上的动态大小数组 int row 在下面的代码中 然后返回其内容 以供另一个设备函数使用 为了将问题置于上下文中 下面的代码尝试使用在 GPU 上运行的高斯 勒让德求积来跨越勒让德多项式基组中的任意函数 incl
  • 如何将CUDA时钟周期转换为毫秒?

    我想用一些代码来测量时间within我的内核需要 我已经关注了这个问题 https stackoverflow com questions 11209228 timing different sections in cuda kernel连
  • cudaMemcpyToSymbol 的问题

    我正在尝试复制到恒定内存 但我不能 因为我对 cudaMemcpyToSymbol 函数的用法有误解 我正在努力追随this http developer download nvidia com compute cuda 4 1 rel t
  • 寻找 CUDA 中的最大值

    我正在尝试在 CUDA 中编写代码来查找最大值 对于给定的一组数字 假设您有 20 个数字 并且内核在 2 个块 每块 5 个线程 上运行 现在假设 10 个线程同时比较前 10 个值 并且thread 2找到最大值 因此线程 2 正在更新
  • 同时使用 2 个 GPU 调用 cudaMalloc 时性能较差

    我有一个应用程序 可以在用户系统上的 GPU 之间分配处理负载 基本上 每个 GPU 都有一个 CPU 线程来启动一个GPU处理间隔当由主应用程序线程定期触发时 考虑以下图像 使用 NVIDIA 的 CUDA 分析器工具生成 作为示例GPU
  • CUDA程序导致nvidia驱动程序崩溃

    当我超过大约 500 次试验和 256 个完整块时 我的 monte carlo pi 计算 CUDA 程序导致我的 nvidia 驱动程序崩溃 这似乎发生在 monteCarlo 内核函数中 任何帮助都会受到赞赏 include
  • 传递给 CUDA 的结构中的指针

    我已经搞砸了一段时间了 但似乎无法正确处理 我正在尝试将包含数组的对象复制到 CUDA 设备内存中 然后再复制回来 但当我遇到它时我会跨过那座桥 struct MyData float data int dataLen void copyT
  • 运行时 API 应用程序中的 cuda 上下文创建和资源关联

    我想了解如何在 cuda 运行时 API 应用程序中创建 cuda 上下文并与内核关联 我知道这是由驱动程序 API 在幕后完成的 但我想了解一下创作的时间线 首先 我知道 cudaRegisterFatBinary 是第一个 cuda a
  • CUDA 估计 2D 网格数据的每块线程数和块数

    首先我要说的是 我已经仔细阅读了所有类似的问题 确定每个块的线程和每个网格的块 https stackoverflow com questions 4391162 cuda determining threads per block blo
  • cuda中有模板化的数学函数吗? [复制]

    这个问题在这里已经有答案了 我一直在寻找 cuda 中的模板化数学函数 但似乎找不到 在普通的 C 中 如果我调用std sqrt它是模板化的 并且将根据参数是浮点数还是双精度数执行不同的版本 我想要这样的 CUDA 设备代码 我的内核将真
  • CUDA 矩阵加法时序,按行与按行比较按栏目

    我目前正在学习 CUDA 并正在做一些练习 其中之一是实现以 3 种不同方式添加矩阵的内核 每个元素 1 个线程 每行 1 个线程和每列 1 个线程 矩阵是方阵 并被实现为一维向量 我只需用以下命令对其进行索引 A N row col 直觉
  • 在 cudaFree() 之前需要 cudaDeviceSynchronize() 吗?

    CUDA 版本 10 1 帕斯卡 GPU 所有命令都发送到默认流 void ptr cudaMalloc ptr launch kernel lt lt lt gt gt gt ptr cudaDeviceSynchronize Is th
  • 从 CUDA 设备写入输出文件

    我是 CUDA 编程的新手 正在将 C 代码重写为并行 CUDA 新代码 有没有一种方法可以直接从设备写入输出数据文件 而无需将数组从设备复制到主机 我假设如果cuPrintf存在 一定有地方可以写一个cuFprintf 抱歉 如果答案已经
  • 最小化 MC 模拟期间存储的 cuRAND 状态数量

    我目前正在 CUDA 中编写蒙特卡罗模拟 因此 我需要生成lots使用随机数cuRAND图书馆 每个线程处理一个巨大的元素floatarray 示例中省略 并在每次内核调用时生成 1 或 2 个随机数 通常的方法 参见下面的示例 似乎是为每
  • 有条件减少 CUDA

    我需要总结一下100000值存储在数组中 但带有条件 有没有办法在 CUDA 中做到这一点以快速产生结果 任何人都可以发布一个小代码来做到这一点吗 我认为 要执行条件约简 您可以直接将条件引入为乘法0 假 或1 真 加数 换句话说 假设您希

随机推荐