CUDA与OpenGL互操作之纹理映射

作者: 爬上墙头deng红杏 | 来源:发表于2020-06-13 08:33 被阅读0次

CUDA与OpenGL互操作之纹理映射
CUDA与OpenGL互操作之顶点缓存映射
OpenGL之纹理映射Mipmap
OpenGL ES学习之纹理贴图
OpenGL纹理坐标&&隧道案例（二）
OpenGL笔记八：2D纹理坐标解析
第七课：纹理滤波、光源和键盘控制
OpenGL ES学习笔记（三）：纹理相关概念
简析OpenGL ES 3.0 Sobel Filter实现边缘
OpenGL坐标概念

引言

在《CUDA与OpenGL互操作之顶点缓存映射》中讲解了如何使用CUDA操作OpenGL中的顶点缓存数据，从而实现CUDA与OpenGL互操作。有时候我们常常使用CUDA对图像进行后处理，比如降噪增强。然后把结果显示在窗口上。这个时候使用VBO进行互操作显然是不合适的，CUDA为我们提供了另外一种手段来实现该目的，即使用纹理映射方式。

取消注册函数

__host__ cudaError_t cudaGraphicsUnregisterResource ( cudaGraphicsResource_t resource )

映射显卡资源

__host__ cudaError_t cudaGraphicsMapResources ( int  count, cudaGraphicsResource_t* resources, cudaStream_t stream = 0 )

映射已经注册的资源到CUDA流中，参数说明：

count：映射的资源数
resources：映射的资源
stream：映射的CUDA同步流
该函数具备默认同步动作

取消映射函数

cudaError_t cudaGraphicsUnmapResources(int count, cudaGraphicsResource_t *resources, cudaStream_t stream __dv(0));

获取纹理数组

__host__ cudaError_t cudaGraphicsSubResourceGetMappedArray ( cudaArray_t* array, cudaGraphicsResource_t resource, unsigned int  arrayIndex, unsigned int  mipLevel )

获取一个数组，通过该数组可以访问已映射图形资源的子资源，参数说明：

array：返回的数组，通过该数组可以访问资源的子资源
resource：用于访问的已映射资源
arrayIndex：数组索引，数组纹理或cubemap面的索引，通过cudaGraphicsCubeFace对cubemap纹理进行子资源访问
mipLevel：用于子资源访问的Mipmap级别

使用示例

参照CUDA simpleCUDA2GL示例，使用Qt进行改写：

#pragma once
#include <QOpenGLWidget>
#include <QOpenGLFunctions>
#include <QOpenGLBuffer>
#include <QOpenGLVertexArrayObject>

#include <cuda_runtime.h>
#include <cuda_gl_interop.h>

class Shader;
class QOpenGLTexture;
class CuTexImg : public QOpenGLWidget, protected QOpenGLFunctions
{
public:
    CuTexImg(QWidget *parent = nullptr);
    ~CuTexImg();

protected:
    virtual void initializeGL();
    virtual void resizeGL(int w, int h);
    virtual void paintGL();

private:
    void makeObject();
    void runCuda();
    void cleanup();

private:
    Shader* m_shader;
    QOpenGLTexture* m_texture;
    QOpenGLBuffer m_vbo, m_ibo;
    QOpenGLVertexArrayObject m_vao;

    uchar4* m_cuPtr;
    cudaGraphicsResource_t m_resource;
};

#include "CuTexImg.h"
#include "Shader.h"
#include <QOpenGLTexture>
#include "helper_cuda.h"

static const int DIM = 512;

extern "C" void
launch_cudaProcess(dim3 grid, dim3 block, uchar4 *g_odata, int imgw);

CuTexImg::CuTexImg(QWidget *parent) : QOpenGLWidget(parent), m_vbo(QOpenGLBuffer::VertexBuffer), m_ibo(QOpenGLBuffer::IndexBuffer)
{
    resize(DIM, DIM);
}

CuTexImg::~CuTexImg() 
{

}

void CuTexImg::initializeGL() 
{
    initializeOpenGLFunctions();

    makeObject();

    checkCudaErrors(cudaGraphicsGLRegisterImage(&m_resource, m_texture->textureId(), GL_TEXTURE_2D, cudaGraphicsMapFlagsWriteDiscard));
    checkCudaErrors(cudaMalloc((void**)&m_cuPtr, DIM*DIM * sizeof(uchar4)));
    runCuda();

    connect(context(), &QOpenGLContext::aboutToBeDestroyed, this, &CuTexImg::cleanup);
}

void CuTexImg::resizeGL(int w, int h) 
{
    glViewport(0, 0, w, h);
}

void CuTexImg::paintGL() 
{
    glClearColor(0.2f, 0.3f, 0.3f, 1.0f);
    glClear(GL_COLOR_BUFFER_BIT);

    m_shader->bind();
    m_texture->bind();
    m_vao.bind();
    glDrawElements(GL_TRIANGLES, 6, GL_UNSIGNED_INT, 0);    
}

void CuTexImg::makeObject()
{
    m_shader = new Shader("shaders/cuTexImg.vert", "shaders/cuTexImg.frag", this);

    float vertices[] = {
        // positions        // texture coords
         1.0f,  1.0f, 0.0f, 1.0f, 1.0f, // top right
         1.0f, -1.0f, 0.0f, 1.0f, 0.0f, // bottom right
        -1.0f, -1.0f, 0.0f, 0.0f, 0.0f, // bottom left
        -1.0f,  1.0f, 0.0f, 0.0f, 1.0f  // top left 
    };
    unsigned int indices[] = {  // note that we start from 0!
        0, 1, 3,  // first Triangle
        1, 2, 3   // second Triangle
    };

    m_vao.create();
    m_vao.bind();

    m_vbo.create();
    m_vbo.bind();
    m_vbo.allocate(vertices, sizeof(vertices));

    m_ibo.create();
    m_ibo.bind();
    m_ibo.allocate(indices, sizeof(indices));

    m_shader->setAttributeBuffer("aPos", GL_FLOAT, 0, 3, 5 * sizeof(float));
    m_shader->enableAttributeArray("aPos");
    m_shader->setAttributeBuffer("aTexCoord", GL_FLOAT, 3 * sizeof(float), 2, 5 * sizeof(float));
    m_shader->enableAttributeArray("aTexCoord");
    
    m_texture = new QOpenGLTexture(QOpenGLTexture::Target2D);
    m_texture->setMagnificationFilter(QOpenGLTexture::Linear);
    m_texture->setMinificationFilter(QOpenGLTexture::Nearest);
    m_texture->setWrapMode(QOpenGLTexture::Repeat);
    m_texture->setFormat(QOpenGLTexture::RGBA8U);       // 格式！！！m_cuPtr的数据类型必须与该格式匹配,采样器的类型也必须一致
    m_texture->setSize(DIM, DIM);                       // m_cuPtr的内存尺寸必须与纹理尺寸一致
    m_texture->allocateStorage();
}

void CuTexImg::runCuda()
{
    cudaArray_t glArray;
    checkCudaErrors(cudaGraphicsMapResources(1, &m_resource, NULL));
    checkCudaErrors(cudaGraphicsSubResourceGetMappedArray(&glArray, m_resource, 0, 0));
    
    dim3 blocks(16, 16);
    dim3 grids((DIM - 1) / 16 + 1, (DIM - 1) / 16 + 1);
    launch_cudaProcess(grids, blocks, m_cuPtr, DIM);

    checkCudaErrors(cudaMemcpyToArray(glArray, 0, 0, m_cuPtr, DIM * DIM * sizeof(uchar4), cudaMemcpyDeviceToDevice));
    
    checkCudaErrors(cudaGraphicsUnmapResources(1, &m_resource, nullptr));
}

void CuTexImg::cleanup()
{
    makeCurrent();

    // 释放显存
    checkCudaErrors(cudaFree(m_cuPtr));
    // 取消资源注册   
    checkCudaErrors(cudaGraphicsUnregisterResource(m_resource));

    // 销毁VAO/VBO
    m_vao.destroy();
    m_vbo.destroy();

    // 释放纹理
    delete m_texture;

    doneCurrent();
}

接下来是cuda代码，cuda核函数基本与示例代码保持一致

#include "helper_cuda.h"

// clamp x to range [a, b]
__device__ float clamp(float x, float a, float b)
{
    return max(a, min(b, x));
}

// convert floating point rgb color to 8-bit integer
__device__ uchar4 rgbToInt(float r, float g, float b)
{
    r = clamp(r, 0.0f, 255.0f);
    g = clamp(g, 0.0f, 255.0f);
    b = clamp(b, 0.0f, 255.0f);
    return make_uchar4(r, g, b, 0);
}

static __global__ void cudaProcess(uchar4* g_odata, int imgw)
{
    int tx = threadIdx.x;
    int ty = threadIdx.y;
    int bw = blockDim.x;
    int bh = blockDim.y;
    int x = blockIdx.x*bw + tx;
    int y = blockIdx.y*bh + ty;

    uchar4 c4 = make_uchar4((x & 0x20) ? 100 : 0, 0, (y & 0x20) ? 100 : 0, 0);
    g_odata[y*imgw + x] = rgbToInt(c4.z, c4.y, c4.x);
}

extern "C" void
launch_cudaProcess(dim3 grid, dim3 block, uchar4 *g_odata, int imgw)
{
    cudaProcess << < grid, block >> > (g_odata, imgw);
}

最后是顶点着色器代码和片段着色器代码

#version 330 core
layout (location = 0) in vec3 aPos;
layout (location = 1) in vec2 aTexCoord;

out vec2 TexCoord;

void main()
{
    gl_Position = vec4(aPos, 1.0);
    TexCoord = aTexCoord;
}

#version 330

in vec2 TexCoord;
uniform usampler2D texImage;

void main()
{
    vec4 c = texture(texImage, TexCoord);
    gl_FragColor = c / 255.0;
}

运行结果：

Cuda&GL.jpg

参考文章：
https://www.khronos.org/registry/OpenGL-Refpages/gl4/html/glTexImage2D.xhtml

CUDA与OpenGL互操作之纹理映射
引言在《CUDA与OpenGL互操作之顶点缓存映射》中讲解了如何使用CUDA操作OpenGL中的顶点缓存数据，从...
CUDA与OpenGL互操作之顶点缓存映射
引言 CUDA相比GLSL更加灵活，逐渐成为主流的GPGPU编程技术。CUDA加速计算后的结果往往需要在显示器中绘...
OpenGL之纹理映射Mipmap
一、Mipmapping 1、Mipmapping是一个功能强大的纹理技术，它可以提高渲染的性能以及提升场景的视觉...
OpenGL ES学习之纹理贴图
OpenGL ES学习之纹理贴图基本原理启用纹理映射功能后，如果想把一副纹理应用到相应的几何图元，就必须告知渲...
OpenGL纹理坐标&&隧道案例（二）
一.隧道案例效果二.纹理坐标纹理坐标主要是把图形映射到OpenGL中去，图形中每个顶点关联到对应的纹理坐标下，...
OpenGL笔记八：2D纹理坐标解析
前言期待您移步上篇：OpenGL笔记七：纹理常用API(二) 纹理采样为了能够把纹理映射(Map)到三角形上，...
第七课：纹理滤波、光源和键盘控制
纹理滤波在纹理映射的过程中，如果图元的大小不等于纹理的大小，OpenGL便会对纹理进行缩放以适应图元的尺寸。我们...
OpenGL ES学习笔记（三）：纹理相关概念
纹理映射到转换后的顶点纹理概念：纹理是一个用来保存图像颜色元素值的OpenGL ES缓存。应该尽量使用最小的...
简析OpenGL ES 3.0 Sobel Filter实现边缘
文档iOS OpenGL ES 3.0 数据可视化 4：纹理映射实现2维图像与视频渲染简介使用索贝尔算子（Sobe...
OpenGL坐标概念
openGL 顶点,坐标系,纹理坐标Android OpenGL es 纹理坐标设定与贴图规则对Android o...