DAY33：阅读Share Memory和Constant Memory

Shared Memory

Because it is on-chip, shared memory has much higher bandwidth and much lower latency than local or global memory.

To achieve high bandwidth, shared memory is divided into equally-sized memory modules, called banks, which can be accessed simultaneously. Any memory read or write request made of naddresses that fall in n distinct memory banks can therefore be serviced simultaneously, yielding an overall bandwidth that is n times as high as the bandwidth of a single module.

However, if two addresses of a memory request fall in the same memory bank, there is a bank conflict and the access has to be serialized. The hardware splits a memory request with bank conflicts into as many separate conflict-free requests as necessary, decreasing throughput by a factor equal to the number of separate memory requests. If the number of separate memory requests is n, the initial memory request is said to cause n-way bank conflicts.

To get maximum performance, it is therefore important to understand how memory addresses map to memory banks in order to schedule the memory requests so as to minimize bank conflicts. This is described in Compute Capability 3.x, Compute Capability 5.x, Compute Capability 6.x, and Compute Capability 7.x for devices of compute capability 3.x, 5.x, 6.x and 7.x, respectively.

Constant Memory

The constant memory space resides in device memory and is cached in the constant cache.

A request is then split into as many separate requests as there are different memory addresses in the initial request, decreasing throughput by a factor equal to the number of separate requests.

The resulting requests are then serviced at the throughput of the constant cache in case of a cache hit, or at the throughput of device memory otherwise.

Texture and Surface Memory

The texture and surface memory spaces reside in device memory and are cached in texture cache, so a texture fetch or surface read costs one memory read from device memory only on a cache miss, otherwise it just costs one read from texture cache. The texture cache is optimized for 2D spatial locality, so threads of the same warp that read texture or surface addresses that are close together in 2D will achieve best performance. Also, it is designed for streaming fetches with a constant latency; a cache hit reduces DRAM bandwidth demand but not fetch latency.

Reading device memory through texture or surface fetching present some benefits that can make it an advantageous alternative to reading device memory from global or constant memory:

· If the memory reads do not follow the access patterns that global or constant memory reads must follow to get good performance, higher bandwidth can be achieved providing that there is locality in the texture fetches or surface reads;

· Addressing calculations are performed outside the kernel by dedicated units;

· Packed data may be broadcast to separate variables in a single operation;

· 8-bit and 16-bit integer input data may be optionally converted to 32 bit floating-point values in the range [0.0, 1.0] or [-1.0, 1.0] (see Texture Memory).

本文备注/经验分享：

今天的主要内容是除了昨日的Local memory之外, 继续提到其他存储器: Constant memory, Shared memory, 以及. Texture和Surface的后备存储器和对应的缓存。 首先说, 昨日没有说的一点是, 目前部分的架构的A卡支持寄存器索引. 可以很大程度的减轻显存读写传输. 考虑到从Maxwell开始, N卡的结构正在大幅度向A卡靠拢(例如从Maxwell起, 像A卡一样的分组的寄存器/SP; 和A卡一样的Shared memory (LDS)), 还有昨日所说的使用了不能在编译时刻确定的下标在N卡上将不再是导致编译器使用local memory的因素——没错. 你现在用的其实是一种支持CUDA的越来越像A卡的N卡，这对用户来说是一个好事. 在保持软件易用性, 和广泛的生态环境的同时, 硬件的性能越来越好。

其次, 关于constant memory, 这个可能是用户最常见和关心的, 例如用户使用的__constant__标记的数组. 它能提到多少性能? 比普通的global memory上的数组又有什么好处?一般来说, 如果没有使用过大过度多的constant的情况下, 它的代价基本上等于0. 如同手册本章所说, constant memory实际上是普通的显存, 外加特别的constant cache构成,你常用的__constant__和编译器自动生成或者你手写在kernel里面的常数(例如a = b * 1234.567f + 890f), 和你的kernel的参数, 都落入此类.这三种是最常见的三种手工使用或者自动使用(对于后两种情况)constant memory的情况.而后面的两种往往却容易被人忽略.constant cache是一种比较特别的cache, 也叫uniform cache,适合warp里面的所有线程, 没有分歧(一致)的访问某些数据的情况.(来自Raspberry Pi的树莓派的用户, 可以将它看成是树莓派的GPU里面的uniforms访问, 有点类似) 在正常使用的情况下(warp里面一致, 没有爆cache---例如你没有使用过多的constant), 它的性能非常好, 甚至往往可以看成是0等待, 数据立等可取.因为这种很低的延迟的特性, 往往对constant cache里面的数据读取, 可以写成指令的一部分, 而不需要单独的一步Load-then-use。(来自x86的用户可以看成这个是作为RISC的GPU的指令的特例, 正常GPU指令都是RISC风格的, 需要首先将数据载入寄存器(访存指令), 然后再在寄存器中计算(计算指令); 而不像x86那样可以: add eax, [8888]这样的一条指令可以同时载入存储单元[8888]的数据, 并进行加法.但constant memory中的数据除外, 它可以直接作为指令的一部分) 这点也从一定的角度说明在正常假定能constant cache总是hit的情况下, 计划中的延迟非常低.需要说明的是, 手册本章节说了, 如果对constant memory的访问不一致, warp内部产生了N个不同的地址, 则实际的性能会下降到1 / N。 "decreasing throughput by a factor equal to the number of separate requests..." 这也是为何它的另外一个名字叫uniform cache, 不一致的访问的确会影响性能的,但在实际使用中, 特别是maxwell+,并未能感觉到性能的剧烈下降. 有的时候甚至在地址几乎每个线程都不同的情况下.这点是和手册矛盾的. 原因暂未知道.不过我建议总是使用手册建议的行为.

来自A卡的用户可以将这个看成是A卡的scalar cache, 或者一定情况下的LDS的uniform load.这是说的constant memory. (请注意以上的说法只是一种大致的说法, 精确的说, 你的kernel中的常数也可能被编译成立即数. 或者constant cache也可以通过其他方式, 例如load uniform指令使用. 等等)。 以及, 需要补充的是, constant memory目前被构造成N * 64KB的形式, 但实际的constant cache较小, 例如只有4-5KB，不建议用户在短期内在kernel里大范围的读取它. 以避免爆cache.(具体的一个较小的working set应该有多小, 暂时无直接能提供的数据. 建议以实际应用为准)。 然后, 关于shared memory, 这个是非常重要的一个存储器.往往具有多种用途: (1) 作为用户手工管理的缓存. 它在性能上堪比CPU的L1 cache, 但在灵活性上可以比喻成一段用户随心所欲使用的小缓冲区. 非常灵活.如果一段数据需要反复被使用, 而且用户知道这点, 应当考虑使用shared memory,因为它往往比自动管理的cache, 在用户手工安排下, 具有更好的性能. 而且特别是maxwell+的Shared memory, 采用了类似GCN的A卡的LDS的结构, 支持所谓的"远程原子操作"，在很多数据使用上具有优势.(例如: 直方图统计、 Compact一个List) 比上几代的(1.X/2.X/3.X)的shared memory在原子操作这种数据访问上, 具有飞跃式的性能提升。 从用户的角度来看, 应当只认为目前有两种shared memory, 一种是Maxwell之前的所有架构(1.X/2.X/3.X), 另外一种则是(5.X/6.X/7.0)，这也"可能是"为何NV越过了计算能力4.X的原因(你会看到计算能力并不连续, 中间并没有4.X这一代)， 因为shared memory有了飞跃式的改进.(另外的一个没有计算能力4的说法是因为当年微软在4.0的DOS上所遭遇的滑铁卢)，同时, 你会看到手册本章这里说, N-way bank conflict会导致N-way的性能下降(例如吞吐率在Fermi上降低到1/N, 访问延迟提升到N倍), 但在Maxwell+开始, 因为改进的shared memory, 就算你不是在用shared上的原子操作，普通的有bank conflict的读写, 实际上性能并没有下降N倍, 很多第三方的评测能表明这一点(例如arxiv上有很多关于maxwell+的shared memory的奇特性能表现的测试文章), 但目前为何maxwell+的shared memory性能这么好, 尚未知道原因.回到常规的. 你应当知道shared memory分成32个bank, 每个bank只有4B(除了Kepler)或者最多8B(Kepler).只要同时读写的地址, 哪怕是不连续, 不合并的, 只要能落入不同的bank中, 那么shared memory就能全性能的提供数据.这就构成了shared memory的第二个用途, (2) 将原本不适合其他存储器类型(例如global memory)上的低效访存模型, 变成shared memory上的高效访问模型.这个第二点shared memory用途是非常重要的用途，例如手册后面有一个矩阵转置的例子, 在正常安排了shared memory的数据存储形状后, 可以高效的进行纵向读写, 而以往在所有的存储器中, 这种都是低效的.请注意这种例子(现在还没到, 因为这几篇文章都是简略介绍性质的), 往往需要你用一种巧妙的方式来规避bank conflict, 例如我原本有float A[16][16]，我想让这个A能高效的被纵向和横向读取, 往往可以实际的定义成A[16][16 + 1]，也就是浪费掉1/17的空间, 但可以让任何的横向或者纵向(或者可能是其他的访问形式)都变成无bank conflict的.用户可以想想一下为何会这样.这里给出一个常见的(其实手册后面有解释, 但没有太详细)的1个warp的访问模型:对于4B的元素类型(例如float或者int),有__shared__ int B[N]，那么只要是B[id * S + T]。 只要S是奇数, 就不会产生bank conflict,因为容易从数学上证明, 任何这种写法的S, 只要S和banks的数量32互质, 那么warp的32个跨步为S个元素的访问, 就会均匀的落入32个banks中(易证, 这里不证)，而和32互致的显然是奇数.所以S为1,3,5,7,9...即可.例如刚才的__shared__ float A[16][17]，横向的访问第4行: A[3][tid]等于写成: A[3][tid * 1 + 0],根据刚才的说法, 无bank conflict，而纵向的访问第5列: A[tid][5]等于写成: float *A' = (float *)A然后再A'的[tid * 17 + t]访问(1D化计算总偏移量), 显然17也无bank conflict，这就是为何你看到的很多代码, 总是喜欢在写成[16][16 + 1]这种定义的原因. shared memory的这个特性很重要, 很多只使用1次的数据也往往通过这种形式, 得到了显著的加速. 而并非惯性中人们容易认为的, shared memory需要有重复的数据使用, 才能得到加速. 不是的.此外, shared memory还适合那种下标随机的访问, 根据信息学的说法, 完全随机的下标数列, 将会基本上在各个区间上的取值概率相同,也就是说, 如果我有__shared__ int C[N]，对C[random_id]的访问, 往往会均匀的落入32个banks中(的不同深度)，此时使用shared memory作为这种随机的下标, 或者不可能预先计算出来的, 或者和实际的具体运行时刻的具体数据有关的下标的, 查找表之类的应用, 非常有效.我们可以常见到AES的GPU上的实现中, 对4个表格的查找, 就往往是在shared memory中进行的。

这是shared memory的第二个重要用途。实际中, shared memory中往往有第三个用户, 在block内部或者warp内部间的数据交换.诚然, block中的多个线程往往可以通过在global memory完成交换, 但在shared memory上的交换提供了高性能的保证(global memory需要走L1/L2 cache, 而自动管理的cache是一个不确定性的东西)，需要说明的是, 在目前所有计算能力的卡中(除了不支持的3.0的卡), 更小范围的交换(warp内部)应当考虑warp shuffle, 未来的CUDA或者N卡可能会引入专用的warp shuffle上的额外加速. 但目前, warp shuffle等于不使用shared memory任何空间的shared memory上的特殊数据交换. 应当使用的.这三点用途, 是shared memory的最常见的使用. 非常重要.大家参加过的所有培训, 无论是否明确的指出这3点, 实际上的讲述的内容都会落入到这三点中, 但你明确的知道了这3点, 有利于你的理解. 剩下的还有texture/surface，它们访问的时候所使用的存储器, 和存储器所对应的cache，这个实际上在之前的章节说texture, surface的时候说过了。 这里再重新强调一下: (1)如果使用了CUDA Array, 那么CUDA Array本身在显存里就构成了一种特殊的存储方式(NV没有公布, 可以去看AMD的文档说明. 或者直接去查询google patent, NV将这个作为一个专利申请公布了), 这种特殊的存储方式提供了2D/3D以及1D(废话)上的纹理元素之间的临近存储, 有利于提高性能. (2)texture cache(可能会合并在其他cache中), 提供了额外的一层缓冲, 由利于提高性能.，然逅关于本章节的另外2点, 我不能理解, 因此暂时无法评价, 但这里给出:关于texture cache(或者合并后的cache): a cache hit reduces DRAM bandwidth demand but not fetch latency。 无法理解的原因是因为, 例如计算能力5.X/6.X上, 它和L1 Cache合并为unified cache, 无法理解在这种情况下, unified cache为何无法降低延迟. 可能是手册这里指的是老式的以前的texture cache, 具有较大延迟?另外一条是这里: Packed data may be broadcast to separate variables in a single operation; 这里不知道手册具体指的是什么.其他本章节的所有内容都给予了解释. 读者可以自行思考这两条.

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

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