DAY17：阅读纹理内存之纹理引用API

3.2.11.1.2. Texture Reference API

Some of the attributes of a texture reference are immutable and must be known at compile time; they are specified when declaring the texture reference. A texture reference is declared at file scope as a variable of type texture:

texture<DataType, Type, ReadMode> texRef;

where:

· DataType specifies the type of the texel;

· Type specifies the type of the texture reference and is equal to[【等同于】cudaTextureType1D, cudaTextureType2D, or cudaTextureType3D, for a one-dimensional, two-dimensional, or three-dimensional texture, respectively, or cudaTextureType1DLayered or cudaTextureType2DLayered for a one-dimensional or two-dimensional layered texture respectively【各自地】 optional argument which defaults to cudaTextureType1D;

· ReadMode specifies the read mode; it is an optional argument which defaults to cudaReadModeElementType.

A texture reference can only be declared as a static global variable and cannot be passed as an argument to a function.

The other attributes of a texture reference are mutable and can be changed at runtime through the host runtime. As explained in the reference manual, the runtime API has a low-levelC-style interface and a high-level C++-style interface. The texture type is defined in the high-level API as a structure publicly derived from the textureReference type defined in the low-level API as such:

· normalized specifies whether texture coordinates are normalized or not;

· filterMode specifies the filtering mode;

· addressMode specifies the addressing mode;

· channelDesc describes the format of the texel; it must match the DataType argument of the texture reference declaration; channelDesc is of the following type:

where x, y, z, and w are equal to the number of bits of each component of the returned value and f is:

o cudaChannelFormatKindSigned if these components are of signed integer type,

o cudaChannelFormatKindUnsigned if they are of unsigned integer type,

o cudaChannelFormatKindFloat if they are of floating point type.

· See reference manual

for sRGB, maxAnisotropy, mipmapFilterMode, mipmapLevelBias, minMipmapLevelClamp, and maxMipmapLevelClamp.

normalized, addressMode, and filterMode may be directly modified in host code.

Before a kernel can use a texture reference to read from texture memory, the texture reference must be bound to a texture using cudaBindTexture() or cudaBindTexture2D() for linear memory, or cudaBindTextureToArray() for CUDA arrays. cudaUnbindTexture() is used to unbind a texture reference. Once a texture reference has been unbound, it can be safely rebound to another array, even if kernels that use the previously bound texture have not completed. It is recommended to allocate two-dimensional textures in linear memory using cudaMallocPitch() and use the pitch returned by cudaMallocPitch() as input parameter to cudaBindTexture2D().

The following code samples bind a 2D texture reference to linear memory pointed to by devPtr:

· Using the low-level API:

· Using the high-level API:

The following code samples bind a 2D texture reference to a CUDA array cuArray:

· Using the low-level API:

· Using the high-level API:

The format specified when binding a texture to a texture reference must match the parameters specified when declaring the texture reference; otherwise, the results of texture fetches are undefined.

There is a limit to the number of textures that can be bound to a kernel as specified in Table 14.

（扫描二维码可以看到table 14)

The following code sample applies some simple transformation kernel to a texture.

本文备注/经验分享：

注意手册有个特点, 如果单纯说texture字样, 是指的后备存储，例如这里:Texture objects are created at runtime and the texture is specified when creating the texture object。纹理对象是在运行的时候被创建的, 而创建的时候指定了后备的存储。(这是对比texture reference来说的,后者是像变量一样的被定义出来的, 后备存储是绑定上去的)。所以你看到了这种:texture<DataType, Type, ReadMode> texRef 因为它是全局变量， 所以不能作为参数传递(建议对照texture object和refernece版本的两个图片旋转的例子查看区别)。

其次, 很多卡上这种定义有总数128个限制.而texture object几乎可以任意多个。

您可能会问：什么情况下用纹理对象API，什么情况下用纹理引用API？

新代码请总是使用纹理对象.引用是给老代码用的。正常请不要使用它，使用纹理对象也有助于你迁移到OpenCL。OpenCL提供的是类似CUDA纹理对象的对应版本.

以及，你看texture reference除了形式不同,访问略微有区别.所有的概念(包括那天你看到的各种归一化操作)都是一样的.

cudaChannelFormatKindSigned , cudaChannelFormatKindUnsigned , cudaChannelFormatKindFloat 这是用来表述通道的，x/y/z/w这些所谓的通道，用普通的说法是"分量"，通道是图形学的说法，例如说, 一个4通道, RGBA的图像，实际上等于说它每个元素有4个分量, 分别是R,G,B,A(红值, 绿值, 蓝值, 透明值)。你也可以将它看成是普通的4个数。因为纹理不一定必须用在图像上。你如果用在普通的矩阵上, 那么等于是普通的4个分量而已，例如可能是射影几何里面的坐标和W分量（缩放的意思），具体分量什么意义和算法有关。这些结构实质上和纹理对象的是一样的。

文中出现的low-level API和high-level API:这个是两种用法而已.哪一种都可以的.注意这里的bind操作是texture reference必须要做的。因为texture reference本身被你写成了全局变量.这里再将这个变量和后备的存储之类的, 以及, 一些信息, 绑定在一起.绑定后才能使用.注意这里给出两个版本, 分别是绑定到普通内存和绑定到不透明的CUDA Array。而每种当给出了两种方式. 建议总是使用简化方式(所谓的高级API)，所以一共是4段代码。而有了这个绑定过程后, 下面的texture reference版本的图片旋转kernel，就可以使用了.

需要注意的是, texRef变量这里的特殊用法.有些是在定义的时候固定的, 例如2D的, 读取是否为原始元素值或者归一值,而有些是可以在运行的时候改变的.(也就是有一些不灵活性. 因为有些东西开头就固定住了)，一些老代码的尖括号内的不是cudaTextureType2D, 而是直接是数字2(表示2D的)，这种维护老代码的时候也需要注意了。后面你还会看到surface。surface是等于没有采样器功能的texture(之前的那些免费的高级功能)，只能进行基本的读写。很多时候如果只需要利用CUDA Array或者其他缓存上的特性, 而不考虑坐标变化, 插值, 值归一化, 边界处理之类的采样器(sampler---这也是个图形学叫法), 可以只使用简化版本的surface。

最后的代码，我建议结合前面的texture object看. 注意一下区别. 功能是一样的.

有不明白的地方，请在本文后留言

或者在我们的技术论坛bbs.gpuworld.cn上发帖