CUDA优化的冷知识18| texture和surface
CUDA优化的冷知识18| texture和surface
大家可以访问：
https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html 来阅读原文。
这是一本很经典的手册。
CUDA优化的冷知识13 |从Global memory到Shared memory
CUDA优化的冷知识14|local memory你可能不知道的好处
上一篇说的3大特性, 都等于在访存的同时, 还附加上一定的固定功能的运算/变换处理. 这种特性, 叫采样器特性(sampler). 而我们都知道, 采样器是在只读路径上的. 而去掉了采样器的texture在CUDA里叫做surface.
因为本优化实践手册编写的年代较早, 这里没有怎么提到surface. 我们简单的说法一下surface.
surface不具有刚才说的texture的采样器只读路径上的这些优势,但是surface具有额外的特性, 它可以写入, 而texture不能.除此之外, surface和texture还具有非采样器的另外的一个重要特性.这个重要特性是在多年前, 也包括最近一些年出的而没有动脑更新的书的经常重点强调的地方,即texture本身的缓存效果. 这是很多人至今还在坚持使用texture的重要因素. 我们来简单看下. 这个因素等于是在说存储本身, 而不是在对该存储的读取路径上的优势. 主要优势有两点, 一个是cache效果. 在某些卡上, 普通的读取不具有较好的缓存效果, 而texture读取有. 例如哪怕是到现在依然被CUDA 11.1所支持5.X硬件, 也是如此.
例如5.0的maxwell的卡, 对于普通的读取不能使用L1/read-only cache, 而texture和另外一种只读的读取方式(不维持一致性(NC)只读读取, 或者常见的__ldg()之类), 却可以充分利用. 此时使用texture或者surface读取, 就能获取此缓存效果上的优势了. 否则你的任何读取可能在此卡上都要走L2. 很亏. 这也是手册本章节说的, 具有带宽上的放大效果(注意, 本章节的其他内容这里不赞同, 因为手册可能很久没更改了, 例如手册说, 使用纹理和DRAM的直接读取具有一样的延迟啥的).另外的一种存储上的优势则是, 例如在使用cuda array的时候, 数据在显存中的排列本身, 可能是被重新排布过的.
也就是说, 一个元素的临近的几个元素可能会被重新排列在一起, 从而带来了空间上的局部性, 你访问了某个元素后, 剩下的某些元素可能就已经在cache中了. 下次读取就带来了更好的访存效果.这是纹理在存储上的两点优势.
注意我们一共说了3+2=5点, 其中前3点因为涉及到一个只读的采样器(SM内部的一个专门的硬件)带来的功能, 所以只能通过纹理来得到此方面的优势/优化;而后面2点, 则是存储方面的, 则是对于texture和surface都适用的.当然, 在现在的卡的架构上, 以前重点的后两点(存储方面)的优势已经不那么明显了, 以前曾经是最重要的用途. 不过现在的代码依然可以考虑后两点, 实际的测试一下看看能否带来性能提升.
此外, 本手册还提醒(不是优化)了读者, 因为texture和surface, 只是对普通的显存的特殊化, 并不真正的存在一种叫纹理之类的存储器,所以实际上他们依然是显存(也可能是内存映射的).等于是在显存的基础上, 套了一层东西. 所以如果你用某种手段, 直接改写了套在纹理/表面背后的实际的显存中的数据存储, 那么这种改变,必须等到下次kernel启动才能生效, 在本次kernel中, 试图读取对后背内容的改变, 将导致未定义的结果. (论坛已经有N个这样的例子了. 我发现我们的用户非常有个性,如果手册说了什么, 那么他做的肯定是这个的反面. 很有意思) 不是你用纹理(例如今天的5点带来的优势)就一定性能会高, 直接使用指针普通访存就性能一定渣.这不一定的. (因为如果一定是纹理或者普通访存就某个一定好, 那么就没有必要提供另外一种访存的可能性/方式了)所以本章节很多地方我们用了"潜在的", "可能的", 而不是说一定的. 这点需要注意. 因为总是有人在论坛问: "为何我用了纹理, 性能没有提升?"
好了. 这是今天的主要内容, 关于纹理存储的优势/优化方面的. 说完纹理, 基本上重要的访存方面的优化就基本说完了. 其实也没有太多方面,