Kepler架构新的纹理内存API——Texture Objects

_梦魇花葬 2014-06-07 11:52:39

加精

小伙伴们对CUDA的Texture内存不会感到陌生，下面的代码展示了Texture Memory的使用方式：



#define N 1024

texture<float, 1, cudaReadModeElementType> tex;



// texture reference name must be known at compile time

__global__ void kernel() {

  int i = blockIdx.x *blockDim.x + threadIdx.x;

  float x = tex1Dfetch(tex, i);

  // do some work using x...

}



void call_kernel(float *buffer) {

  // bind texture to buffer

  cudaBindTexture(0, tex, buffer, N*sizeof(float));



  dim3 block(128,1,1);

  dim3 grid(N/block.x,1,1);

  kernel <<<grid, block>>>();



  // unbind texture from buffer

  cudaUnbindTexture(tex);

}



int main() {

  // declare and allocate memory

  float *buffer;

  cudaMalloc(&buffer, N*sizeof(float));

  call_kernel(buffer);

  cudaFree(buffer);

}

基本过程是这样子的：首先声明一个全局的纹理，然后在Kernel函数使用前绑定纹理，最后在Kernel进行拾取纹理操作。使用Texture Memroy可以加速数组的随机访问，并且提供了一些纹理的操作（对于Fermi和Kepler架构的GPU来说，提供的L1和L2缓存也可以加速数组的随机访问，不用在代码中做任何改变便可以加速，因此要综合考虑是否使用texture memory）。
这种纹理操作方式有两个缺点：
1. 使用的时候需要程序员手动绑定和解绑定内存地址，操作繁琐；
2. 纹理不能声明为全局的静态变量，而且不能作为函数参数传递。
从CUDA 5.0版本开始，对于Kepler架构的GPU来说，提供了一种新的操作Texture Memory API，即Texture Objects. 下面的例子展示了使用方式：



#define N 1024



// texture object is a kernel argument

__global__ void kernel(cudaTextureObject_t tex) {

  int i = blockIdx.x *blockDim.x + threadIdx.x;

  float x = tex1Dfetch<float>(tex, i);

  // do some work using x ...

}



void call_kernel(cudaTextureObject_t tex) {

  dim3 block(128,1,1);

  dim3 grid(N/block.x,1,1);

  kernel <<<grid, block>>>(tex);

}



int main() {

  // declare and allocate memory

  float *buffer;

  cudaMalloc(&buffer, N*sizeof(float));



  // create texture object

  cudaResourceDesc resDesc;

  memset(&resDesc, 0, sizeof(resDesc));

  resDesc.resType = cudaResourceTypeLinear;

  resDesc.res.linear.devPtr = buffer;

  resDesc.res.linear.desc.f = cudaChannelFormatKindFloat;

  resDesc.res.linear.desc.x = 32; // bits per channel

  resDesc.res.linear.sizeInBytes = N*sizeof(float);



  cudaTextureDesc texDesc;

  memset(&texDesc, 0, sizeof(texDesc));

  texDesc.readMode = cudaReadModeElementType;



  // create texture object: we only have to do this once!

  cudaTextureObject_t tex=0;

  cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL);



  call_kernel(tex); // pass texture as argument



  // destroy texture object

  cudaDestroyTextureObject(tex);



  cudaFree(buffer);

}

使用Texture Objects时，编译器在编译时期不需要知道在运行时将要使用哪个纹理，因此在编程方面更加的自由。此外，Texture Objects只需要被实例化一次，因此使用的时候不用不断的绑定、解绑定，消除了绑定与解绑定的运行时间，降低了Kernel函数运行的时间，运行效率有所提高，特别是在那些对延迟敏感的HPC程序中。
综合来看，Texture Objects不仅简化了操作，消除了原始Texture API的绑定与解绑定操作，而且加速了程序的运行效率，如果GPU满足Kepler架构的条件，应该尝试着去使用Texture Objects去代替原始的API~~

...全文