纹理对象是CUDA针对纹理参考缺点而提出的升级版,其作用和纹理参考完全一致,但是使用方法更加灵活。与纹理参考相比,CUDA对其进行各方面的升级,一方面是可以再代码中申请和销毁,另一方面则可以作为设备函数的参数进行传入;可以满足一些特殊的需求。
一个纹理对象是用cudaCreateTextureObject()产生的。cudaCreateTextureObject()有4个参数,常用的前三个是
- cudaTextureObject_t *texObj:需要生产的纹理对象;
- cudaResourceDesc *resDesc:资源描述符,用来获取述纹理数据;
- 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端才能正常使用。