DAY44：阅读Texture Reference API

B.8.2. Texture Reference API

B.8.2.1. tex1Dfetch()

fetches from the region of linear memory bound to the one-dimensional texture reference texRef using integer texture coordinate x. tex1Dfetch() only works with non-normalized coordinates, so only the border and clamp addressing modes are supported. It does not perform any texture filtering. For integer types, it may optionally promote the integer to single-precision floating point.

Besides the functions shown above, 2-, and 4-tuples are supported; for example:

fetches from the region of linear memory bound to texture reference texRef using texture coordinate x.

B.8.2.2. tex1D()

fetches from the CUDA array bound to the one-dimensional texture reference texRef using texture coordinate x. Type is equal to DataType except when readMode is equal to cudaReadModeNormalizedFloat (see Texture Reference API), in which case Type is equal to the matching floating-point type.

B.8.2.3. tex1DLod()

template<class DataType, enumcudaTextureReadMode readMode>
Type tex1DLod(texture<DataType, cudaTextureType1D, readMode> texRef, float x,
              float level);

fetches from the CUDA array bound to the one-dimensional texture reference texRef using texture coordinate x. The level-of-detail is given by level. Type is the same as DataTypeexcept when readMode is cudaReadModeNormalizedFloat , in which case Type is the corresponding floating-point type.

B.8.2.4. tex1DGrad()

template<class DataType, enumcudaTextureReadMode readMode>
Type tex1DGrad(texture<DataType, cudaTextureType1D, readMode> texRef, float x,						 float dx, float dy);

fetches from the CUDA array bound to the one-dimensional texture reference texRef using texture coordinate x. The level-of-detail is derived from the dx and dy X- and Y-gradients.Type is the same as DataType except when readMode is cudaReadModeNormalizedFloat (see Texture Reference API), in which case Type is the corresponding floating-point type.

B.8.2.5. tex2D()

template<class DataType, enum cudaTextureReadMode readMode>
Type tex2D(texture<DataType, cudaTextureType2D, readMode> texRef, 
          float x, float y);

fetches from the CUDA array or the region of linear memory bound to the two-dimensional texture reference texRef using texture coordinates x and y. Type is equal to DataType except when readMode is equal to cudaReadModeNormalizedFloat (see Texture Reference API), in which case Type is equal to the matching floating-point type.

B.8.2.6. tex2DLod()

template<class DataType, enumcudaTextureReadMode readMode>
Type tex2DLod(texture<DataType, cudaTextureType2D, readMode> texRef,  
            float x, float y, float level);

fetches from the CUDA array bound to the two-dimensional texture reference texRef using texture coordinate (x,y). The level-of-detail is given by level. Type is the same as DataTypeexcept when readMode is cudaReadModeNormalizedFloat (see Texture Reference API), in which case Type is the corresponding floating-point type.

B.8.2.7. tex2DGrad()

template<class DataType, enumcudaTextureReadMode readMode>
Type tex2DGrad(texture<DataType, cudaTextureType2D, readMode> texRef,               float x, float y, float2 dx, float2 dy);

fetches from the CUDA array bound to the two-dimensional texture reference texRef using texture coordinate (x,y). The level-of-detail is derived from the dx and dy X- and Y-gradients.Type is the same as DataType except when readMode is cudaReadModeNormalizedFloat (see Texture Reference API), in which case Type is the corresponding floating-point type.

B.8.2.8. tex3D()

template<class DataType, enum cudaTextureReadMode readMode>
Type tex3D(texture<DataType, cudaTextureType3D, readMode> texRef,  
         float x, float y, float z);

fetches from the CUDA array bound to the three-dimensional texture reference texRef using texture coordinates x, y, and z. Type is equal to DataType except when readMode is equal to cudaReadModeNormalizedFloat (see Texture Reference API), in which case Type is equal to the matching floating-point type.

B.8.2.9. tex3DLod()

template<class DataType, enumcudaTextureReadMode readMode>
Type tex3DLod(texture<DataType, cudaTextureType3D, readMode> texRef, 
             float x, float y, float z, float level);

fetches from the CUDA array bound to the two-dimensional texture reference texRef using texture coordinate (x,y,z). The level-of-detail is given by level. Type is the same as DataType except when readMode is cudaReadModeNormalizedFloat (see Texture Reference API), in which case Type is the corresponding floating-point type.

B.8.2.10. tex3DGrad()

template<class DataType, enumcudaTextureReadMode readMode>
Type tex3DGrad(texture<DataType, cudaTextureType3D, readMode> texRef,               float x, float y, float z, float4 dx, float4 dy);

fetches from the CUDA array bound to the two-dimensional texture reference texRef using texture coordinate (x,y,z). The level-of-detail is derived from the dx and dy X- and Y-gradients. Type is the same as DataType except when readMode is cudaReadModeNormalizedFloat (see Texture Reference API), in which case Type is the corresponding floating-point type.

B.8.2.11. tex1DLayered()

template<class DataType, enum cudaTextureReadMode readMode>
Type tex1DLayered( texture<DataType, cudaTextureType1DLayered, readMode> texRef, 
    float x, int layer);

fetches from the CUDA array bound to the one-dimensional layered texture reference texRef using texture coordinate x and index layer, as described in Layered Textures. Type is equal to DataType except when readMode is equal to cudaReadModeNormalizedFloat (see Texture Reference API), in which case Type is equal to the matching floating-point type.

B.8.2.12. tex1DLayeredLod()

template<class DataType, enum cudaTextureReadMode readMode>
Type tex1DLayeredLod(texture<DataType, cudaTextureType1D, readMode> texRef,                     float x, int layer, float level);

fetches from the CUDA array bound to the one-dimensional texture reference texRef using texture coordinate x and index layer as described in Layered Textures. The level-of-detail is given by level. Type is the same as DataType except when readMode is cudaReadModeNormalizedFloat (see Texture Reference API), in which case Type is the corresponding floating-point type.

B.8.2.13. tex1DLayeredGrad()

template<class DataType, enum cudaTextureReadMode readMode>
Type tex1DLayeredGrad(texture<DataType, cudaTextureType1D, readMode> texRef,                      float x, int layer, float dx, float dy);

fetches from the CUDA array bound to the one-dimensional texture reference texRef using texture coordinate x and index layer as described in Layered Textures. The level-of-detail is derived from the dx and dy X- and Y-gradients. Type is the same as DataType except when readMode is cudaReadModeNormalizedFloat (see Texture Reference API), in which case Typeis the corresponding floating-point type.

B.8.2.14. tex2DLayered()

template<class DataType, enum cudaTextureReadMode readMode>
Type tex2DLayered( texture<DataType, cudaTextureType2DLayered, readMode> texRef,
     float x, float y, int layer);

fetches from the CUDA array bound to the two-dimensional layered texture reference texRef using texture coordinates x and y, and index layer, as described in Texture Memory. Type is equal to DataType except when readMode is equal to cudaReadModeNormalizedFloat (see Texture Reference API), in which case Type is equal to the matching floating-point type.

B.8.2.15. tex2DLayeredLod()

template<class DataType, enum cudaTextureReadMode readMode>
Type tex2DLayeredLod(texture<DataType, cudaTextureType2D, readMode> texRef,                     float x, float y, int layer, float level);

fetches from the CUDA array bound to the two-dimensional texture reference texRef using texture coordinate (x,y) and index layer as described in Layered Textures. The level-of-detail is given by level. Type is the same as DataType except when readMode is cudaReadModeNormalizedFloat (see Texture Reference API), in which case Type is the corresponding floating-point type.

B.8.2.16. tex2DLayeredGrad()

template<class DataType, enum cudaTextureReadMode readMode>
Type tex2DLayeredGrad(texture<DataType, cudaTextureType2D, readMode> texRef,                      float x, float y, int layer, float2 dx, float2 dy);

fetches from the CUDA array bound to the two-dimensional texture reference texRef using texture coordinate (x,y) and index layer as described in Layered Textures. The level-of-detail is derived from the dx and dy X- and Y-gradients. Type is the same as DataType except when readMode is cudaReadModeNormalizedFloat (see Texture Reference API), in which case Type is the corresponding floating-point type.

B.8.2.17. texCubemap()

template<class DataType, enum cudaTextureReadMode readMode>
Type texCubemap( texture<DataType, cudaTextureTypeCubemap, readMode> texRef, 
    float x, float y, float z);

fetches from the CUDA array bound to the cubemap texture reference texRef using texture coordinates x, y, and z, as described in Cubemap Textures. Type is equal to DataType except when readMode is equal to cudaReadModeNormalizedFloat (see Texture Reference API), in which case Type is equal to the matching floating-point type.

B.8.2.18. texCubemapLod()

template<class DataType, enum cudaTextureReadMode readMode>
Type texCubemapLod(texture<DataType, cudaTextureType3D, readMode> texRef,                   float x, float y, float z, float level);

fetches from the CUDA array bound to the two-dimensional texture reference texRef using texture coordinate (x,y,z). The level-of-detail is given by level. Type is the same as DataType except when readMode is cudaReadModeNormalizedFloat (see Texture Reference API), in which case Type is the corresponding floating-point type.

B.8.2.19. texCubemapLayered()

template<class DataType, enum cudaTextureReadMode readMode>
Type texCubemapLayered(texture<DataType, cudaTextureTypeCubemapLayered, readMode> texRef,
float x, float y, float z, int layer);

fetches from the CUDA array bound to the cubemap layered texture reference texRef using texture coordinates x, y, and z, and index layer, as described in Cubemap Layered Textures. Type is equal to DataType except when readMode is equal to cudaReadModeNormalizedFloat (see Texture Reference API), in which case Type is equal to the matching floating-point type.

B.8.2.20. texCubemapLayeredLod()

template<class DataType, enum cudaTextureReadMode readMode>
Type texCubemapLayeredLod(texture<DataType, cudaTextureType3D, readMode> texRef,                          float x, float y, float z, int layer, float level);

fetches from the CUDA array bound to the two-dimensional texture reference texRef using texture coordinate (x,y,z) and index layer as described in Layered Textures. The level-of-detail is given by level. Type is the same as DataType except when readMode is cudaReadModeNormalizedFloat (see Texture Reference API), in which case Type is the corresponding floating-point type.

B.8.2.21. tex2Dgather()

template<class DataType, enum cudaTextureReadMode readMode>
Type tex2Dgather( texture<DataType, cudaTextureType2D, readMode> texRef,    
 float x, float y, int comp = 0);

fetches from the CUDA array bound to the 2D texture reference texRef using texture coordinates x and y and the comp parameter as described in Texture Gather. Type is a 4-component vector type. It is based on the base type of DataType except when readMode is equal to cudaReadModeNormalizedFloat (see Texture Reference API), in which case it is always float4.

本文备注/经验分享：

昨天的章节DAY43:阅读纹理之Texture Object API，加上今天的章节, 主要是在device端(也就是kernel代码里)所能使用的纹理读取(识取)函数.还记得纹理么. 之前的章节介绍过好多. 主要因为历史原因有两种, 一种叫纹理引用, 一种叫纹理对象. 后者是新出的. 新写的代码应当总是使用后者.后者有很多好处

(请参考之的地章节：

DAY16：阅读纹理内存之纹理对象API

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

DAY18：阅读纹理内存之Layered Textures

DAY19：阅读纹理内存之Texture Gather) 所以今天的内容基本上也分成两部分. 一部分是针对纹理引用的, 一部分是针对纹理对象的.只不过手册上先说的是新出的纹理对象, 而老的纹理引用则放到了后头.这也多少反映了NV对CUDA用户使用这两种的偏向性的意见吧. 没有太多能说的, 主要需要注意的是两点: 一点是纹理必须要有后背存储空间, 分成普通的线性显存(常规我们最常用的就是. 能用基地址+指针偏移量直接访问的). 和CUDA Array。后者是一种NV私有类型, 不公开内部的存储格式, 等于是一个黑箱子, 只能通过本章节的这些函数访问. 之前也说过, 里面对元素的存放为特殊优化过的, 有的时候可能会有更好的读取性能.前者则你可以通过本章节的函数访问, 也可以直接通过普通指针访问, 当然, 通过普通指针访问就没有纹理读取的好处了. 本章节主要针对纹理读取的.请一定要注意这两种纹理的后备存储的区别.本章节还特定说明了, tex1Dfetch(), 只能用于普通的线性内存.而tex1D()只能用于cuda array的情况.请注意这里说明了前者只能使用两种边界模式。 另外两种边界模式用不了的.这两种模式是越界返回0---一个很有用的特性.以及, 越界返回边界值.所谓的border和clamp边界模式. 还记得4大边界模式么? 直接sisiy的章节里面, 有过一幅画，里面有4种效果.除了这里说的两种外, 还有越界后镜像的，以及越界后重复的.. 没错, 就像你设定桌面时候的模式选择一样.忘记了章节不要紧, 右键设定一下你的桌面吧.另外两种这里没说的模式是mirror, 以及wrap模式,也叫镜像和圈绕.不过这两种模式需要归一化的坐标(也就是说, 普通常规坐标能用2种越界处理模式; 归一话话坐标能用全部4种).还记得这两种坐标模式么???分别是类似[0,N)和[0.0, 1.0)的.之前的章节还有用归一化的坐标进行图片旋转的，欢迎查看一下之前的代码例子(旋转那个). 然后应当没有太多需要说的了. tex2D()是统一的, 可以处理普通的线性内存后备的纹理,也可以处理CUDA Array的.请注意用户还应当知道如何分配这两种空间.最普通的线性存储可以直接使用cudaMalloc()分配.而CUDA Array分配则在基本的cudaMalloc的名字的基础上, 增加了维度和Array字样.例如cudaMalloc3DArray()。 注意这个函数的名字发生了变化, 多了3D和Array字样. 这代表是用来进行3D的CUDA Array分配的. 注意2D的线性内存分配比较奇特, 是cudaMallocPitch(), 而不是你想象中的cudaMalloc2D()。 为何叫pitch前后章节都有说明. 欢迎阅读. 请注意别忘记普通内存只能使用1D和2D的普通texture,而cuda array可以处理所有的各种类型的texture. 这里需要注意的是

ayered的texture在分配的时候, 例如多层的2D的texture(还记得它么? 之前的章节说过, 你可以直接看成它是2D纹理构成的数组),需要使用cudaMalloc3DArray这种分配.注意是2D的用3D的分配.1D的用2D的分配。 分配的时候有单独的一个参数用来说明, 是进行分层的.不要弄混.此外, 本章节的LOD, MIPMAP什么的, 和图像处理有关, 我不懂它们. (类似之前的章节, 我们也越过了它们)，感兴趣的可以用户自己阅读. 总之, 纹理的2大存储, 3大功能(坐标变换, 值变换, 边界处理), 和单独特有的缓存(这个和计算能力有关, 不一定），构成了用户主要使用它的目的. 例如值变换可以选择返回归一化的浮点数纹元值.例如值变换+坐标变换可以直接进行普通的线性插值.例如边界处理可以等效的自动将你的图像或者矩阵围绕上一圈0(很多人喜欢手工围绕一圈0, 用texture读取可以自动免费加上).等等. 除了本章节手册说的这些读取函数(以及, 里面的说明和注意事项外), 除了刚才我们说的各种(包括分配各种后备存储和哪种后备存储能用哪种类型的纹理), 具体的纹理读取(拾取)时候还有很多细节. 例如手册最后面的附录里面提到的半个点的坐标偏移.这些都需要注意.

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

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