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上发帖

原文发布于微信公众号 - 吉浦迅科技(gpusolution)

原文发表时间:2018-05-23

本文参与腾讯云自媒体分享计划,欢迎正在阅读的你也加入,一起分享。

发表于

我来说两句

0 条评论
登录 后参与评论

相关文章

来自专栏黑白安全

用紧缩创建Wordlists

很多时候,在渗透测试活动中,您将发现需要绕过的身份验证表单,以便访问应用程序或远程系统。拥有大而好的单词列表总是有帮助,但作为渗透测试人员,您必须能够根据具体情...

861
来自专栏信安之路

通过实例学习ROP技术

这两天闲着无聊就在 exploit-db 上找个小软件练练手,但是 exploit-db 上面给的 poc 并能正常运行,无奈只好自己写了一个,顺便写篇文章把自...

920
来自专栏技术博文

PHPExcel所遇到问题的知识点总结

工作中进行excel的时候遇到了两个问题, 1.excel表中列值过大,由于没有进行特殊处理,程序没法正常运行; 2.列值中含有日期格式的文本,不能正确读取; ...

2645
来自专栏州的先生

对Django模型对象进行序列化和反序列化

1253
来自专栏Golang语言社区

厚土Go学习笔记 | 04. 导入和导出的不同 用math.Pi来举例

go语言代码中的import是导入包。 导入单个的包可以写成 import "fmt" 如果导入多个包的话,可以用圆括号进行组合导入,写成下面这个样子。 imp...

2666
来自专栏Java 技术分享

Struts2 转换器

2537
来自专栏CDA数据分析师

实用小工具,教你轻松转化Python通用数据格式

已独立成项目在github上面 dataformat, 涉及模块 os, getopt, sys。 1 需求 在进行hadoop测试时,需要造大量数据,例如某个...

1865
来自专栏java一日一条

Java管理扩展指南之MBean简介

MBean是一个被管理的Java对象,就像Javabean组件一样,但是它遵从JMX规范的设计模式。MBean可以表示设备、应用或者任何需要被管理的资源。MBe...

831
来自专栏程序员的SOD蜜

PDF.NET数据开发框架 之SQL-MAP使用存储过程

有关SQL-MAP的规范性介绍,请看下面的文章: PDF.NET(PWMIS数据开发框架)之SQL-MAP目标和规范 在SQL-MAP中使用存储过程 1...

23410
来自专栏xingoo, 一个梦想做发明家的程序员

C/C++ 遇到0xcccccccc访问冲突

最近一直在纠结这个问题. ? 最近写代码,总是遇到这个问题,一旦遇到这个问题,以前好使的代码也就不好使了。很费解,上网搜集了下资料.... 这个0xcccccc...

1817

扫码关注云+社区