引言
CUDA相比GLSL更加灵活,逐渐成为主流的GPGPU编程技术。CUDA加速计算后的结果往往需要在显示器中绘制出来,如果我们将在显存中的计算结果拷贝到主存,然后使用QPainter绘制到窗口中,这中间的显存-主存,主存-显存的数据拷贝会显得冗余。既然CUDA计算的结果和图形绘制的内容都在显存里,那么可以通过一些技术手段直接将CUDA计算结果绘制到显示器上,即CUDA与Graphics的互操作。
相关接口介绍
注册显卡资源
__host__ cudaError_t cudaGraphicsGLRegisterBuffer(struct cudaGraphicsResource **resource, GLuint buffer, unsigned int flags);
注册一个OpenGL缓冲对象,参数说明:
- resource - 指向返回的对象句柄
*buffer - 被注册的缓冲对象名 - flags - 注册标识
- cudaGraphicsRegisterFlagsNone: 注册的资源在CUDA程序可以进行读写,默认为该枚举
- cudaGraphicsRegisterFlagsReadOnly: 指定CUDA程序只能读该资源
- cudaGraphicsRegisterFlagsWriteDiscard: 指定CUDA程序不会读该资源对象,每次写资源时资源的内容完全被覆盖,不会保留之前的任何数据
取消注册函数
cudaError_t cudaGraphicsUnregisterResource(cudaGraphicsResource_t resource);
映射显卡资源
cudaError_t cudaGraphicsMapResources(int count, cudaGraphicsResource_t *resources, cudaStream_t stream __dv(0));
映射已经注册的资源到CUDA流中,参数说明:
- count:映射的资源数
- resources:映射的资源
- stream:映射的CUDA同步流
该函数具备默认同步动作。
取消映射函数
cudaError_t cudaGraphicsUnmapResources(int count, cudaGraphicsResource_t *resources, cudaStream_t stream __dv(0));
获取设备指针
cudaError_t cudaGraphicsResourceGetMappedPointer(void **devPtr, size_t *size, cudaGraphicsResource_t resource);
获取已映射资源的设备指针,参数说明:
- devPtr:设备指针,CUDA程序可以使用该指针进行数据更新
- size:缓冲数据尺寸
- resources:已映射的资源
使用示例
参照CUDA simpleGL示例,该示例是使用CUDA对顶点数据进行操作后,再将顶点绘制到窗口上。使用Qt进行改写:
class CuGl : public QOpenGLWidget, protected QOpenGLFunctions
{
public:
CuGl(QWidget* parent = nullptr);
~CuGl();
protected:
virtual void initializeGL();
virtual void resizeGL(int w, int h);
virtual void paintGL();
private:
void runCuda();
void cleanup();
private:
float m_fAnim;
Shader *m_shader;
QOpenGLBuffer m_vbo;
QOpenGLVertexArrayObject m_vao;
cudaGraphicsResource_t m_vboResource;
};
static const int DIM = 512;
static __global__ void kernel(float3* 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_float3(u, w, v);
}
CuGl::CuGl(QWidget* parent) : QOpenGLWidget(parent), m_fAnim(0.0f)
{
resize(DIM, DIM);
}
CuGl::~CuGl()
{
}
void CuGl::initializeGL()
{
initializeOpenGLFunctions();
m_shader = new Shader("shaders/cuda.vert", "shaders/cuda.frag", this);
// 创建VAO/VBO
if (m_vao.create())
{
m_vao.bind();
m_vbo.create();
m_vbo.bind();
unsigned int nbytes = DIM * DIM * 3 * sizeof(GLfloat);
m_vbo.allocate(nbytes);
m_shader->setAttributeBuffer("aPos", GL_FLOAT, 0, 3, 3 * sizeof(float));
m_shader->enableAttributeArray("aPos");
}
// 注册cuda资源
checkCudaErrors(cudaGraphicsGLRegisterBuffer(&m_vboResource, m_vbo.bufferId(), cudaGraphicsMapFlagsWriteDiscard));
// 设置清屏颜色
glClearColor(0.1f, 0.1f, 0.1f, 1.0f);
connect(context(), &QOpenGLContext::aboutToBeDestroyed, this, &CuGl::cleanup);
}
void CuGl::resizeGL(int w, int h)
{
glViewport(0, 0, w, h);
QMatrix4x4 model, view, projection;
view.lookAt(QVector3D(0, 0, 3), QVector3D(0, 0, 0), QVector3D(0, 1, 0));
projection.perspective(45.0f, (float)w / h, 0.1f, 100.0f);
m_shader->bind();
m_shader->setUniformValue("model", model);
m_shader->setUniformValue("view", view);
m_shader->setUniformValue("projection", projection);
}
void CuGl::paintGL()
{
// 使用CUDA进行计算
runCuda();
// 使用OPenGL进行绘制
glClear(GL_COLOR_BUFFER_BIT);
m_shader->bind();
m_vao.bind();
glDrawArrays(GL_POINTS, 0, DIM*DIM);
update();
}
void CuGl::runCuda()
{
float3* glPtr;
size_t size;
// 映射已注册的资源
checkCudaErrors(cudaGraphicsMapResources(1, &m_vboResource, 0));
// 获取映射后的资源的设备指针
checkCudaErrors(cudaGraphicsResourceGetMappedPointer((void**)&glPtr, &size, m_vboResource));
// 利用设备指针对数据做计算
dim3 blocks(16, 16);
dim3 grids((DIM - 1) / 16 + 1, (DIM - 1) / 16 + 1);
kernel << <grids, blocks >> > (glPtr, DIM, DIM, m_fAnim);
m_fAnim += 0.01f;
// 取消资源映射
checkCudaErrors(cudaGraphicsUnmapResources(1, &m_vboResource, 0));
}
void CuGl::cleanup()
{
// 取消资源注册
checkCudaErrors(cudaGraphicsUnregisterResource(m_vboResource));
// 销毁VAO/VBO
m_vao.destroy();
m_vbo.destroy();
}
顶点着色器和片段着色器程序:
#version 330 core
layout(location = 0) in vec3 aPos;
uniform mat4 model;
uniform mat4 view;
uniform mat4 projection;
void main()
{
gl_Position = projection * view * model * vec4(aPos,1.0f);
}
#version 330 core
out vec4 fragColor;
void main()
{
fragColor = vec4(0.4, 0.7, 0.6, 1.0);
}
运行效果:
Cuda&OpenGL.png
网友评论