[CUDA]纹理对象 Texture Object

╰+哭是因爲堅強的太久メ 2022-02-15 03:35 329阅读 1赞

纹理对象是CUDA针对纹理参考缺点而提出的升级版，其作用和纹理参考完全一致，但是使用方法更加灵活。与纹理参考相比，CUDA对其进行各方面的升级，一方面是可以再代码中申请和销毁，另一方面则可以作为设备函数的参数进行传入；可以满足一些特殊的需求。

一个纹理对象是用cudaCreateTextureObject()产生的。cudaCreateTextureObject()有4个参数，常用的前三个是

1.  cudaTextureObject\_t \*texObj：需要生产的纹理对象；
2.  cudaResourceDesc \*resDesc：资源描述符，用来获取述纹理数据；
3.  cudaTextureDesc \*texDesc：纹理描述符，用来描述纹理参数；

*  其中cudaTextureDesc定义如下：

struct cudaTextureDesc
    {
        enum cudaTextureAddressMode addressMode[3];
        enum cudaTextureFilterMode filterMode;
        enum cudaTextureReadMode readMode;
        int sRGB;
        int normalizedCoords;
        unsigned int maxAnisotropy;
        enum cudaTextureFilterMode mipmapFilterMode;
        float mipmapLevelBias;
        float minMipmapLevelClamp;
        float maxMipmapLevelClamp;
    };

使用纹理对象主要包括纹理对象创建、纹理访问和纹理对象销毁。

*  纹理对象创建

纹理对象创建之前首先要分别对纹理资源和纹理对象属性进行确定，分别对应cudaResoruceDesc和cudaTextureDesc；然后即可利用cudaCreateTextureObject来创建纹理对象。

*  纹理访问

纹理对象的纹理访问也和纹理参考一样，也是使用tex1D或tex2D等函数进行操作。

*  纹理对象销毁

纹理对象的销毁直接使用cudaDestroyTextureObject即可。值得注意的是应该先销毁对象，然后再释放对应的设备内存。

*  一个使用纹理对象的简单例子：

// 简单纹理变换函数
    __global__ void transformKernel(float* output,cudaTextureObject_t texObj,
    int width, int height, float theta){
        // 计算纹理坐标
        unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
        unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
        float u = x / (float)width;
        float v = y / (float)height;
        // 坐标转换
        u -= 0.5f;
        v -= 0.5f;
        float tu = u * cosf(theta) - v * sinf(theta) + 0.5f;
        float tv = v * cosf(theta) + u * sinf(theta) + 0.5f;
        // 从纹理中读取并写入全局存储
        output[y * width + x] = tex2D<float>(texObj, tu, tv);
    }
     
    int main(){
        // 定义CUDA array
        cudaChannelFormatDesc channelDesc =
        cudaCreateChannelDesc(32, 0, 0, 0,
        cudaChannelFormatKindFloat);
        cudaArray* cuArray;
        cudaMallocArray(&cuArray, &channelDesc, width, height);
        // 拷贝数据到CUDA array
        cudaMemcpyToArray(cuArray, 0, 0, h_data, size,
        cudaMemcpyHostToDevice);
        // 定义资源描述符
        struct cudaResourceDesc resDesc;
        memset(&resDesc, 0, sizeof(resDesc));
        resDesc.resType = cudaResourceTypeArray;
        resDesc.res.array.array = cuArray;
        // 定义纹理对象参数
        struct cudaTextureDesc texDesc;
        memset(&texDesc, 0, sizeof(texDesc));
        texDesc.addressMode[0] = cudaAddressModeWrap;
        texDesc.addressMode[1] = cudaAddressModeWrap;
        texDesc.filterMode = cudaFilterModeLinear;
        texDesc.readMode = cudaReadModeElementType;
        texDesc.normalizedCoords = 1;
        // 生产纹理对象
        cudaTextureObject_t texObj = 0;
        cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL);
        // 分配用于保持结果的内存
        float* output;
        cudaMalloc(&output, width * height * sizeof(float));
        // 调用Kernel
        dim3 dimBlock(16, 16);
        dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x,
        (height + dimBlock.y - 1) / dimBlock.y);
        transformKernel<<<dimGrid, dimBlock>>>(output, texObj, width, height, angle);
        // 销毁纹理对象
        cudaDestroyTextureObject(texObj);
        // 释放内存
        cudaFreeArray(cuArray);
    }

在使用纹理对象时还有一些特殊注意。通常而言，使用纹理对象的目的就是使用多个纹理对象，因此会申请一个cudaTextureObject\_t 数组。注意在使用时必须将其拷贝到device端才能正常使用。