OpenGL与CUDA内存交互
cudaGraphicsGLRegisterImage | cudaGraphicsMapResources | cudaGraphicsSubResourceGetMappedArray | cudaMemcpy2DToArray | cudaGraphicsUnmapResources
OpenGL和Direct3D中的 一些资源可以映射到CUDA的地址空间
CUDA能够读取OpenGL或Direct3D编写的数据,或者使CUDA能够编写数据供OpenGL或Direct3D使用。
在使用OpenGL互操作和Direct3D互操作功能 映射资源 之前,必须 将资源注册到CUDA 。
- 这些函数返回一个指向类型为 struct cudaGraphicsResource 的CUDA图形资源的指针。
- 注册资源是潜在的高开销,因此通常每个资源只调用一次。
- 使用cudaGraphicsUnRegierResource()注销CUDA图形资源。
- 每个打算使用该资源的CUDA上下文都需要单独注册。
- 一旦将资源注册到CUDA,就可以使用cudaGraphicsMapResources()和cudaGraphicsUnmapResources()将其映射和取消映射的必要次数。
- 可以调用cudaGraphicsResourceSetMapFlag()来指定使用提示(只写、只读),CUDA驱动程序可以使用这些提示来优化资源管理。
- 在映射资源时,通过OpenGL、Direct3D或其他CUDA上下文访问资源会产生未定义的结果。OpenGL互操作和Direct3D互操作给出了每个图形API和一些代码示例的具体说明。
OpenGL与CUDA互操作可以分成两种,一种是OpenGL将Buffer对象注册到CUDA中去,供CUDA读写操作,然后再在OpenGL中使用。一般这种情况下注册的是VBO和PBO,VBO一般用于存储顶点坐标、索引等数据;PBO则一般用于存储图像数据,因此称作Pixel Buffer Object。另一种是OpenGL将Texture对象注册到CUDA中去,经CUDA处理后得到纹理内容,然后在OpenGL中渲染出来。不过不管是哪一种互操作类型,其操作流程是一致的:
- 在OpenGL里面初始化Buffer Object
- 在CUDA中注册OpenGL中的Buffer Object
- CUDA锁定资源,获取操作资源的指针,在CUDA核函数中进行处理
- CUDA释放资源,在OpenGL中使用Buffer Object
cudaGraphicsGLRegisterImage 注册资源
// 注册 OpenGL 纹理或渲染缓冲区对象。
__host__cudaError_t cudaGraphicsGLRegisterImage ( cudaGraphicsResource** resource, GLuint image, GLenum target, unsigned int flags )
// 调用方式
checkCudaErrors(cudaGraphicsGLRegisterImage(&m_cudaRes[0], m_personTexture[0], GL_TEXTURE_2D, cudaGraphicsRegisterFlagsNone));
Register an OpenGL texture or renderbuffer object.
Parameters
resource
- 指向返回对象句柄的指针
image
- 要注册的纹理或渲染缓冲区对象的名称
target
- 识别图像指定的对象类型
flags
- 注册标志
return
cudaSuccess,cudaErrorInvalidDevice,cudaErrorInvalidValue,cudaErrorInvalidResourceHandle,cudaErrorUnknown
描述
注册由图像指定的纹理或渲染缓冲区对象以供 CUDA 访问。注册对象的句柄作为资源返回。
target必须与对象的类型匹配,并且必须是 GL_TEXTURE_2D、GL_TEXTURE_RECTANGLE、GL_TEXTURE_CUBE_MAP、GL_TEXTURE_3D、GL_TEXTURE_2D_ARRAY 或 GL_RENDERBUFFER 之一。
寄存器标志标志指定预期用途,如下所示:
cudaGraphicsRegisterFlagsNone:不指定有关如何使用此资源的提示。因此,假定该资源将由 CUDA 读取和写入。这是默认值。
cudaGraphicsRegisterFlagsReadOnly:指定 CUDA 不会写入此资源。
cudaGraphicsRegisterFlagsWriteDiscard:指定 CUDA 不会从该资源中读取,而是会覆盖该资源的全部内容,因此不会保留之前存储在该资源中的任何数据。
cudaGraphicsRegisterFlagsSurfaceLoadStore:指定 CUDA 将此资源绑定到表面引用。
cudaGraphicsRegisterFlagsTextureGather:指定 CUDA 将对此资源执行纹理收集操作。
支持以下图像格式。为简洁起见,该列表是缩写的。例如,{GL_R, GL_RG} X {8, 16} 将扩展为以下 4 种格式 {GL_R8, GL_R16, GL_RG8, GL_RG16} :
GL_RED、GL_RG、GL_RGBA、GL_LUMINANCE、GL_ALPHA、GL_LUMINANCE_ALPHA、GL_INTENSITY
{GL_R, GL_RG, GL_RGBA} X {8, 16, 16F, 32F, 8UI, 16UI, 32UI, 8I, 16I, 32I}
{GL_LUMINANCE, GL_ALPHA, GL_LUMINANCE_ALPHA, GL_INTENSITY} X {8, 16, 16F_ARB, 32F_ARB, 8UI_EXT, 16UI_EXT, 32UI_EXT, 8I_EXT, 16I_EXT, 32I_EXT}
当前不允许使用以下图像类:
带边框的纹理
多重采样渲染缓冲区
笔记:
请注意,此函数还可能从以前的异步启动返回错误代码。
cudaGraphicsMapResources 映射资源
// 映射图形资源以供 CUDA 访问。
__host__cudaError_t cudaGraphicsMapResources ( int count, cudaGraphicsResource_t* resources, cudaStream_t stream = 0 )
// 调用方式
checkCudaErrors(cudaGraphicsMapResources(2, m_cudaRes, 0)); // 为提高效率,一次map 2个资源
- count:- 要映射的资源数量
- resources:- 为 CUDA 映射资源
- stream:- 同步流
描述:
映射资源中的count个图形资源以供CUDA 访问。
CUDA 可以访问resources中的资源,直到它们被取消映射。注册resources的图形 API 在由 CUDA 映射时不应访问任何资源。如果应用程序这样做,则结果是未定义的。
此函数提供同步保证,即在cudaGraphicsMapResources()之前发出的任何图形调用将在流中发出的任何后续 CUDA 工作开始之前完成。
如果资源包含任何重复条目,则返回 cudaErrorInvalidResourceHandle。如果当前映射任何资源以供 CUDA 访问,则返回 cudaErrorUnknown。
cudaGraphicsSubResourceGetMappedArray
cudaGraphicsSubResourceGetMappedArray
// 获取一个cuda数组,通过该数组访问映射的图形资源的子资源。
__host__cudaError_t cudaGraphicsSubResourceGetMappedArray ( cudaArray_t* array, cudaGraphicsResource_t resource, unsigned int arrayIndex, unsigned int mipLevel )
// 调用方式
checkCudaErrors(cudaGraphicsSubResourceGetMappedArray(&m_pCudatextPtr, m_cudaRes[0], 0, 0));
cudaMemcpy2DToArray
cudaError_t cudaMemcpy2DToArray ( struct cudaArray * dst,
size_t wOffset,
size_t hOffset,
const void * src,
size_t spitch,
size_t width,
size_t height,
enum cudaMemcpyKind kind
)
// 调用方式
checkCudaErrors(cudaMemcpy2DToArray(m_pCudatextPtr, 0, 0, (void*)m_pdRgba,
m_renderWidth * 4 * sizeof(uchar),
m_renderWidth * 4 * sizeof(uchar), m_renderHeight,
cudaMemcpyDeviceToDevice));
从 src 指向的内存区域复制一个矩阵(每行宽度字节的高度行)到从左上角开始的 CUDA 数组 dst (wOffset, hOffset),其中 kind 是 cudaMemcpyHostToHost、cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost 或 cudaMemcpyDeviceToDevice 之一,以及 指定复制的方向。 spitch 是 src 指向的二维数组在内存中的宽度(以字节为单位),包括添加到每行末尾的任何填充。 wOffset + width 不得超过 CUDA 数组 dst 的宽度。 宽度不得超过间距。 如果 spitch 超过允许的最大值,cudaMemcpy2DToArray() 将返回错误。
Parameters:
- dst - Destination memory address
- wOffset - Destination starting X offset
- hOffset - Destination starting Y offset
- src - Source memory address
- spitch - Pitch of source memory
- width - Width of matrix transfer (columns in bytes)
- height - Height of matrix transfer (rows)
- kind - Type of transfer
cudaGraphicsUnmapResources
// 取消映射图形资源。
__host__cudaError_t cudaGraphicsUnmapResources ( int count, cudaGraphicsResource_t* resources, cudaStream_t stream = 0 )
// 调用方式
checkCudaErrors(cudaGraphicsUnmapResources(2, m_cudaRes, 0));
参数
count
- 取消映射的资源数量
resources - 取消映射的资源
stream - 同步流
描述
取消映射资源中的计数图形资源。
一旦未映射,资源中的资源可能无法被 CUDA 访问,直到它们再次被映射。
此函数提供同步保证,即在cudaGraphicsUnmapResources()之前在流中发布的任何 CUDA 工作将在任何后续发布的图形工作开始之前完成。
如果资源包含任何重复条目,则返回 cudaErrorInvalidResourceHandle。如果当前没有映射任何资源以供 CUDA 访问,则返回 cudaErrorUnknown。
调用参考
void InitCudaGLFbo(void) {
// cuda 注册 gl 的纹理
checkCudaErrors(cudaGraphicsGLRegisterImage(&m_cudaRes[0], m_personTexture[0], GL_TEXTURE_2D, cudaGraphicsRegisterFlagsNone));
checkCudaErrors(cudaGraphicsGLRegisterImage(&m_cudaRes[1], m_personTexture[1], GL_TEXTURE_2D, cudaGraphicsRegisterFlagsNone));
AllocCudaGlMem();
}
void BindTextureFromCuda(cReconstruction::Ptr p_reconstruction) {
checkCudaErrors(cudaGraphicsMapResources(2, m_cudaRes, 0)); // 为提高效率,一次map 2个资源
// color 得到的结果是3通道的,先转为4通道再拷贝到纹理
// 在CUDA中锁定资源,获得操作Texture的指针,这里是CudaArray*类型
checkCudaErrors(cudaGraphicsSubResourceGetMappedArray(&m_pCudatextPtr, m_cudaRes[0], 0, 0));
p_reconstruction->GetRGBAToGPU(m_pdRgba);
// 数据拷贝至CudaArray。这里因为得到的是CudaArray,处理时不方便操作,于是先在设备内存中
// 分配缓冲区处理,处理完后再把结果存到CudaArray中,仅仅是GPU内存中的操作。
checkCudaErrors(cudaMemcpy2DToArray(m_pCudatextPtr, 0, 0, (void*)m_pdRgba,
m_renderWidth * 4 * sizeof(uchar),
m_renderWidth * 4 * sizeof(uchar), m_renderHeight,
cudaMemcpyDeviceToDevice));
// depth
checkCudaErrors(cudaGraphicsSubResourceGetMappedArray(&m_pCudatextPtr, m_cudaRes[1], 0, 0));
//ChronoGraphs::start("GetDepthRGBAToGPU", TIMER_DEBUG_INFO);
// p_reconstruction->GetDepthRGBAToGPU(pd_depth_rgba);
//ChronoGraphs::during("GetDepthRGBAToGPU", TIMER_DEBUG_INFO);
p_reconstruction->GetDepthRGBAToGPU(m_pdDepthRgba);
checkCudaErrors(cudaMemcpy2DToArray(m_pCudatextPtr, 0, 0, (void*)m_pdDepthRgba,
m_renderWidth * 4 * sizeof(uchar),
m_renderWidth * 4 * sizeof(uchar), m_renderHeight,
cudaMemcpyDeviceToDevice));
checkCudaErrors(cudaDeviceSynchronize());
// 处理完后即解除资源锁定,OpenGL可以利用得到的Texture对象进行纹理贴图操作了。
checkCudaErrors(cudaGraphicsUnmapResources(2, m_cudaRes, 0));
}