技术开发 频道

CUDA纹理存储器的特性及其使用

  3.2.4.1 CUDA数组

  在显存中可以分配的空间有两种:CUDA 数组和线性内存。此外,常数存储器中通过缓存加速读取的数据实际也存在于显存中。CUDA数组和线性内存都可以与纹理参照系绑定,但CUDA数组对纹理拾取访问进行了优化,在设备端也只能通过纹理拾取访问。

  声明CUDA数组之前,必须先以结构体channelDesc描述CUDA数组中的数据类型。

struct cudaChannelFormatDesc {

    
int x, y, z, w;

    
enum cudaChannelFormatKind f;

};

  其中,x, y, z和w分别是每个返回值成员的位数,而f是一个枚举变量,可以取一下几个值:

  • cudaChannelFormatKindSigned,如果这些成员是有符号整型;
  • cudaChannelFormatKindUnsigned,如果这些成员是无符号整型;
  • cudaChannelFormatKindFloat,如果这些成员是浮点型;

  然后,我们要确定CUDA数组的维度和尺寸。CUDA数组可以通过cudaMalloc3DArray()或cudaMallocArray()函数分配。用cudaMalloc3DArray可以分配一维、二维或者三维的CUDA数组,而cudaMallocArray()一般用于分配二维CUDA数组。在使用完CUDA数组后,要使用cudaFreeArray函数释放显存。

  由cudaMalloc3DArray分配的CUDA数组使用cudaMemcpy3D()完成与其他CUDA数组或者线性内存的数据传输。CUDA API中使用结构体cudaExtent描述3D Array和3D线性内存在三个维度上的尺寸,在描述一维、二维和三维数组分别用以下的形式:

  cudaextent extent = make_cudaextent([1,8192],0,0);

  cudaextent extent = make_cudaextent([1,65535],[1,32768],0);

  cudaextent extent = make_cudaextent([1,2048],[1,2048],[1,2048]);

  其中方括号[]内为允许的寻址范围。注意到二维CUDA数组的第一个维度的寻址范围大于一维CUDA数组的寻址范围,因此在一维CUDA数组的尺寸不够用时,将二维CUDA数组的第二个维度设为1代替一维CUDA数组,获得更大的寻址范围。

  下面是声明一个数据类型为char2型,宽×高×深为64×32×16的CUDA 3D数组,对其初始化,最后释放数组的示例代码:

cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 8, 0, 0,cudaChannelFormatKindunsigned); //每个像元由两个char构成

cudaExtent extent
= make_cudaextent(64,32,16);//建立cudaExtent结构体,描述CUDA数组的维度和尺寸

cudaArray
* cuArray;

cudaMalloc3DArray(
&cuArray, &channelDesc, extent); //为cuArray开辟空间



cudaFreeArray(cuArray);

  下面则是使用cudaMallocArray声明一个由float型构成,尺寸为64×32的CUDA数组,对其赋值,并最后释放的示例代码:

cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0,cudaChannelFormatKindunsigned); //每个像元由一个float构成

cudaArray
* cuArray;

cudaMallocArray(
&cuArray, &channelDesc, 64, 32); //为cuArray开辟空间

cudaMemcpyToArray(cuArray,
0, 0, h_data, &channelDesc);//第二和第三个参数分别表示在宽度和高度上的偏移量,假设h_data中的数据已经初始化

cudaFreeArray(cuArray);

  用于在CUDA数组和主机端或者设备端线性内存,以及在CUDA数组间传输数据的函数还有很多,这些还是还有一些异步调用版本,关于这些函数的具体使用方法请参考CUDA Reference mannual。

  3.2.4.1 声明纹理参照系

  纹理参照系中的一些属性必须在编译时之前被显示声明。纹理参照系通过一个作用范围为全文件的texture型变量声明:

texture<Type, Dim, ReadMode> texRef;

  其中,

  • Type 确定了由纹理拾取返回的数据类型;Type可以是B3.1节中描述的任意一种由基本整型或者单精度浮点型组成能的1-,2-或者4-元组向量类型。
  •  Dim 确定了纹理参照系的维度,默认为1。
  •  ReadMode 可以是cudaReadModeNormalizedFloat或者cudaReadModeElementType。如果ReadMode是cudaReadModeNomalizedFloat,并且Type是16-或者8-bit整型,那么返回的值将是一个浮点数。此时,原来整形的值域会被映射到[0.0,1.0](对无符号整型),或者[-1.0,1.0](对有符号整型)。例如,一个值为0xff的8-bit无符号整型会被映射为1.0f。如果使用cudaReadModeElementType,那么就不会对输出进行转换。ReadMode是一个可选参数,如果不写,那么默认就是cudaReadModeElementType。

  例如,下面的代码声明了一个二维,像元数据为unsigned char型,但将返回值转换为float型的纹理参照系:

texture<unsigned char, 2, cudaReadModeNormalizedFloat> texRef;

  3.2.4.2 设置运行时纹理参照系属性

  纹理参照系中的其它属性可以不必声明,并在运行时进行修改。这些参数规定了纹理的寻址模式,是否进行归一化,以及纹理滤波。runtimeAPI拥有底层的C风格和高层的C++风格两种接口。高层API中的texture类型是从底层的textureReference中派生而来的。TextureReference是一个下面的代码描述的结构体。

struct textureReference {

    
int normalized;

    
enum cudaTextureFilterMode filterMode;

    
enum cudaTextureAddressMode addressMode[3];

     struct cudaChannelFormatDesc channelDesc;

}
  • normalized 设置是否对纹理坐标是否进行归一化。如果normalized是一个非零值,那么就会使用归一化到[0,1)的坐标进行寻址,否则对尺寸为width, height, depth的纹理使用坐标[0,width-1], [0,height-1], [0,depth-1]寻址。例如,一个尺寸为64×32的纹理可以通过x维度范围为[0,63],y维度范围[0,31]的坐标寻址。如果采用归一化方式对尺寸为64×32的纹理进行寻址,在x和y维度上的坐标就都是[0.0,1.0)。这样就可以保证纹理的坐标与纹理的尺寸无关。
  •  filterMode用于设置纹理的滤波模式,即如何根据坐标计算返回的纹理值。滤波模式可以是cudaFilterModePoint或者cudaFilterModeLinear。滤波模式为CudaFilterModePoint时,返回值是与坐标最接近的像元的值。CudaFilterModeLinear模式只能对返回值为浮点型的纹理使用,启用这一种模式时将拾取纹理坐标周围的像元,然后根据坐标与这些像元之间的距离进行插值计算。对一维纹理可以使用线性滤波,对二维纹理可以使用双线性滤波。返回值会是对最接近纹理坐标的两个像元(对一维纹理),四个像元(对二维纹理)或者八个像元(对三维纹理)进行插值后得到的值。
  • addressmode说明了寻址模式,即如何处理超出寻址范围的纹理坐标;addressmode是一个大小为3的数组,三个元素分别说明对第一、二、三个纹理坐标的取址模式;取址模式可以是cudaAddressModeClamp或cudaAddressModeWrap中的一种,前者将超出寻址范围的纹理坐标”钳位”到寻址范围内的最大或最小值,后者将超出寻址范围的纹理坐标“折叠”进合理范围。cudaAddressModeWrap只支持归一化的纹理坐标。                                                                                                         对非归一化的坐标,如果寻址的坐标超过了范围[0,N],大于N的坐标将被钳位,设为N-1。对归一化的坐标,有钳位和循环两种处理方式,在钳位方式下,超过[0.0,1.0)范围的坐标将被钳位到[0.0,1.0);循环方式一般用于周期循环纹理,它只使用了纹理坐标中有用的小数部分,例如1.25会被当作0.25处理,而-1.25则会被当成0.75处理。
  • channelDesc描述纹理获取返回值类型,我们已经在3.2.4.1小节讲解CUDA array时介绍过这个结构体。纹理参照系的返回值类型描述必须和与之绑定的CUDA array的数据类型描述相同,或者和与之绑定的线性内存中的元素类型相同。normalized, addressMode和filterMode可以直接在主机端代码中修改。它们只适用于与CUDA数组绑定的纹理参照系。

  附录D中列出了关于纹理拾取的更多信息。

2
相关文章