找回密码
 立即注册

QQ登录

只需一步,快速开始

查看: 84|回复: 2

DAY81:阅读Texture Fetching

[复制链接]
发表于 2018-8-29 14:07:30 | 显示全部楼层 |阅读模式
G. Texture Fetching
This appendix gives the formula used to compute the value returned by the texture functions of Texture Functions depending on the various attributes of the texture reference (see Texture and Surface Memory).
The texture bound to the texture reference is represented as an array T of
  • N texels for a one-dimensional texture,
  • N x M texels for a two-dimensional texture,
  • N x M x L texels for a three-dimensional texture.
It is fetched using non-normalized texture coordinates x, y, and z, or the normalized texture coordinates x/N, y/M, and z/L as described in Texture Memory. In this appendix, the coordinates are assumed to be in the valid range. Texture Memory explained how out-of-range coordinates are remapped to the valid range based on the addressing mode.

G.1. Nearest-Point Sampling
In this filtering mode, the value returned by the texture fetch is
  • tex(x)=T for a one-dimensional texture,
  • tex(x,y)=T[i,j] for a two-dimensional texture,
  • tex(x,y,z)=T[i,j,k] for a three-dimensional texture,

where i=floor(x), j=floor(y), and k=floor(z).
Figure 13 illustrates nearest-point sampling for a one-dimensional texture with N=4.
For integer textures, the value returned by the texture fetch can be optionally remapped to [0.0, 1.0] (see Texture Memory).
Figure 13. Nearest-Point Sampling Filtering Mode. Nearest-point sampling of a one-dimensional texture of four texels.


                               
登录/注册后可看大图


回复

使用道具 举报

 楼主| 发表于 2018-8-29 14:26:18 | 显示全部楼层
G.2. Linear Filtering
In this filtering mode, which is only available for floating-point textures, the value returned by the texture fetch is
  • tex(x)=(1−α)T+αT[i+1] for a one-dimensional texture,
  • tex(x,y)=(1−α)(1−β)T[i,j]+α(1−β)T[i+1,j]+(1−α)βT[i,j+1]+αβT[i+1,j+1] for a two-dimensional texture,
  • tex(x,y,z) =
    (1−α)(1−β)(1−γ)T[i,j,k]+α(1−β)(1−γ)T[i+1,j,k]+
    (1−α)β(1−γ)T[i,j+1,k]+αβ(1−γ)T[i+1,j+1,k]+
    (1−α)(1−β)γT[i,j,k+1]+α(1−β)γT[i+1,j,k+1]+
    (1−α)βγT[i,j+1,k+1]+αβγT[i+1,j+1,k+1]
    for a three-dimensional texture,

where:
  • i=floor(xB), α=frac(xB), xB=x-0.5,
  • j=floor(yB), β=frac(yB), yB=y-0.5,
  • k=floor(zB), γ=frac(zB), zB= z-0.5,
α, β, and γ are stored in 9-bit fixed point format with 8 bits of fractional value (so 1.0 is exactly represented).
Figure 14 illustrates linear filtering of a one-dimensional texture with N=4.
Figure 14. Linear Filtering Mode. Linear filtering of a one-dimensional texture of four texels in clamp addressing mode.


                               
登录/注册后可看大图





G.3. Table Lookup
A table lookup TL(x) where x spans the interval [0,R] can be implemented as TL(x)=tex((N-1)/R)x+0.5) in order to ensure that TL(0)=T[0] and TL(R)=T[N-1].
Figure 15 illustrates the use of texture filtering to implement a table lookup with R=4 or R=1 from a one-dimensional texture with N=4.
Figure 15. One-Dimensional Table Lookup Using Linear Filtering


                               
登录/注册后可看大图



回复 支持 反对

使用道具 举报

 楼主| 发表于 2018-8-29 15:03:56 | 显示全部楼层
sisiy 发表于 2018-8-29 14:26
G.2. Linear FilteringIn this filtering mode, which is only available for floating-point textures, th ...

这个附录章节主要说明了, 在纹理读取/采样的时候, 内置的常用的最近点采样(也就是不插值), 和普通的1D/2D/3D方向上线性插值时候的, 坐标变换和值变换的关系.还讲述了在线性插值的情况下, 将一种大小的纹理, 做为另外一种大小的查找表时候的坐标变换情况.注意这里首先说了最近点采样, 也就是不插值.可以从配图种看到, 采样的坐标(纹理读取使用的坐标)和采样得到的值(读取得到的值, 以下均使用读取)之间,具有尖锐的边缘变化.所以一般情况下, 可能并不适合图像处理之类的用途使用, 除非你总是使用整数的坐标.而常规的, 我们之前的章节看到的例子, 将1个图片总是归一化到[0.0, 1.0)这种坐标的时候(X/Y/Z)每个方向的时候,几乎总是需要进行某种插值, 来规避这种尖锐的变化, 让读取到的图片看起来点和点之间的过渡平滑一些(还记得之前的图像旋转的例子吗?)
但是再说线性插值之前, 我们需要注意一下, 读取时候指定的坐标,和读取到的纹理的点(纹理元素, 纹元)之间的变换关系,这个关系之前的章节基本没有说.本章节额外的指出了. 有0.5个坐标值的差异.这个问题非常重要. 因为论坛上天天有人挂在这里。
例如: http://bbs.gpuworld.cn/thread-10385-1-1.html
这个人就是忘记了0.5的偏差(注意这个人是2D的. X/Y每个方向上都有这个偏差),然后一旦像这个人还混合了边界模式(此人使用了钳位模式), 最终读取到的结果可能变得很诡异. 难以找到原因,论坛也不只这一个人的问题.历年来大量的人都挂在这上面了. 特别是前些年深度学习还不流行的时候, 大家很多做图像处理的, 喜欢使用纹理, 这里就会出现问题,然后需要补充的是,当读取时候采用了归一化坐标的时候(从[0, N)到[0.0, 1.0f)的变换),这里的偏差就不是0.5个了, 而是半个点.这里本章节里面的例子中, 图片的下方可以看出是大约0.125的坐标偏差.(大约等于0.5/4),这里也需要注意一下.然后再说一下第二段落的线性插值.这种插值模式是最基本的, 能由硬件的tex单元(SM里的一个部分)免费提供的(但可能运算精度不高).
除了这种插值外的其他所有的插值方式都不能免费得到,需要用户自行按照上一段落的最近点读取到附近的一些点的值后, 手工进行计算.然后就算是这种简单的免费提供的插值运算,本段落里面的式子也看起来颇为复杂.例如3D时候的: (1−α)(1−β)(1−γ)T[i,j,k]+α(1−β)(1−γ)T[i+1,j,k]+
(1−α)β(1−γ)T[i,j+1,k]+αβ(1−γ)T[i+1,j+1,k]+
(1−α)(1−β)γT[i,j,k+1]+α(1−β)γT[i+1,j,k+1]+
(1−α)βγT[i,j+1,k+1]+αβγT[i+1,j+1,k+1]s

这样的一个线性变换关系.但实际上没有必要看这个. 有个简化版本的的:首先将坐标去掉0.5个值的偏移.然后直接就可以当作普通的理解了. 我用1D的举个例子:
假设你是坐标0.5, 实际上存在的两个点[0]的值是1.2, 点[1]的值是8.0,那么坐标[0.5]的值就是正好1.2和8.0的平均值, 也就是4.6,如果点继续往[1]处偏移, 例如[0.6]坐标的点, 则值是[0] * 0.4 + [1] * 0.6这样的平均.直到改点完全移动到了[1]处, 则值就整好是[1]的值(注意坐标是纠正偏移后的), 可以直接看成[0] * 0.0 + [1] * 1.0 = [1]的值了.这样的理解很形象.
2D和3D的类似, 只是参与计算的扩大到了附近的4个或者8个点而已.然后这种线性插值在这种理解下, 读取到的值只是最近的点按照靠近自己的临进程度进行线性计算了. 很容易看.
然后回到第三段落.这里是将一种不同的后备存储的纹理大小, 做为另外一种大小的查找表来使用.此时往往需要进行线性插值.本段落给出了坐标的变换关系, 用户可以自行看一下.需要说明的是, 如果是同样大小的当作查找表用, 而且只需要读取整数位置的点的时候, 则无需这样麻烦,直接用普通的非归一化坐标, 然后别忘记有加上的0.5个坐标偏移即可。
很多人以前将texture当作查找表用, 是因为之前的SM内部并没有能缓冲显存的cache存在(计算能力1.x时候的老问题).而同时可能查找表较小, 和/或此时用满了shared memory, 没有其他选择.在现在的硬件上, 如果不考虑缩放的因素(例如本章节最后段落里面对于两种不同大小的说明),那么可以直接使用普通的缓冲区进行查表. 当然, 能考虑shared memory的时候还是应当考虑的.需要说明的是, 在Maxwell和Pascal上, 查找表很小的时候, 一定要启用L1(Unified cache),只走L2的时候往往查表性能很悲剧(因为表很小, 只能访问L2的一个部分).可以考虑使用const type * __restrict__ p的形式的指针,或者考虑在Maxwell 2代和Pascal上, 使用-dlcm=ca的ptx参数设定.或者考虑__ldg()之类的函数.此时将会利用上这两代卡上的Unified cache.在这种情况下, unified cache(你也可以叫它L1或者texture cache),类似一个后备覆盖了整个显存的巨大隐形纹理.
只有这样, 进行查表访存的时候才能不使用纹理, 而具有纹理查表的效果.大致需要说明的就是这些.到下一个章节就是具体的计算能力说明了.后续是一个很重要的章节.
回复 支持 反对

使用道具 举报

您需要登录后才可以回帖 登录 | 立即注册

本版积分规则

关闭

站长推荐上一条 /1 下一条

快速回复 返回顶部 返回列表