GLSL与CUDA的互操作（2）—— 读取与使用纹理数据

adagio_chen 2014-06-24 02:29:29

加精

当我们使用render to texture技术把数据绘制到纹理以后，如何在CUDA中读取opengl的纹理并用于计算呢？这里需要用到CUDA resource 、 CUDA array 和 CUDA texture 这三样东西。

CUDA resource 用来注册一个 opengl的资源（这里是纹理数据）； CUDA array 用来获取资源的地址；最后把 CUDA array绑定到 CUDA texture上，就可以在kernel中使用纹理读取函数来读取数据了。



// 加入互操作相关的头文件

#include <cuda_gl_interop.h>



// 声明二维纹理，4个通道，每个通道8位，用uchar4表示

texture<uchar4, 2, cudaReadModeElementType> cuda_tex;



__global__

void Kernel ( ...... )

{

    uint j = blockIdx.x;

    uint i = threadIdx.x;



    // 读取纹理

    uchar4 texel = tex2D(cuda_tex, i, j);

}







extern "C"

void Launch_Kernel(int width, int height, GLuint gl_tex, ......)

{

    // 注册和绑定cuda resource

    cudaGraphicsResource *resource;

    checkCudaErrors(cudaGraphicsGLRegisterImage(&resource, gl_tex, GL_TEXTURE_2D, cudaGraphicsMapFlagsNone));

    checkCudaErrors(cudaGraphicsMapResources(1, &resource, NULL));



    // 获得此资源地址

    cudaArray_t array_d;

    checkCudaErrors(cudaGraphicsSubResourceGetMappedArray(&array_d, resource, 0, 0));



    // 绑定到CUDA纹理

    checkCudaErrors(cudaBindTextureToArray(cuda_tex, array_d));



    // 调用kernel函数

    Kernel<<<height, witdh>>>( ...... )

        

    // 释放资源

    checkCudaErrors(cudaGraphicsUnmapResources(1, &resource, NULL));

    checkCudaErrors(cudaGraphicsUnregisterResource(resource));

    checkCudaErrors(cudaUnbindTexture(cuda_tex));

}