DAY31：阅读global memory

我们正带领大家开始阅读英文的《CUDA C Programming Guide》,今天是第31天，我们正在讲解性能，希望在接下来的69天里，您可以学习到原汁原味的CUDA，同时能养成英文阅读的习惯。

本文共计652字，阅读时间20分钟

今天开始我们讲解GPU存储访问，需要花费2天。

今天我们的笔记内容是到目前为止最长的一次（可能后面会越来越长）

5.3.2. Device Memory Accesses

An instruction that accesses addressable memory (i.e., global, local, shared, constant, or texture memory) might need to be re-issued multiple times depending on the distribution of the memory addresses across the threads within the warp. How the distribution affects the instruction throughput this way is specific to each type of memory and described in the following sections. For example, for global memory, as a general rule, the more scattered the addresses are, the more reduced the throughput is.

Global Memory

Global memory resides in device memory and device memory is accessed via 32-, 64-, or 128-byte memory transactions. These memory transactions must be naturally aligned: Only the 32-, 64-, or 128-byte segments of device memory that are aligned to their size (i.e., whose first address is a multiple of their size) can be read or written by memory transactions.

When a warp executes an instruction that accesses global memory, it coalesces the memory accesses of the threads within the warp into one or more of these memory transactions depending on the size of the word accessed by each thread and the distribution of the memory addresses across the threads. In general, the more transactions are necessary, the more unused words are transferred in addition to the words accessed by the threads, reducing the instruction throughput accordingly. For example, if a 32-byte memory transaction is generated for each thread's 4-byte access, throughput is divided by 8.

How many transactions are necessary and how much throughput is ultimately affected varies with the compute capability of the device. Compute Capability 3.x, Compute Capability 5.x, Compute Capability 6.x and Compute Capability 7.x give more details on how global memory accesses are handled for various compute capabilities.

To maximize global memory throughput, it is therefore important to maximize coalescing by:

· Following the most optimal access patterns based on Compute Capability 3.x, Compute Capability 5.x, Compute Capability 6.x and Compute Capability 7.x,

· Using data types that meet the size and alignment requirement detailed in Device Memory Accesses,

· Padding data in some cases, for example, when accessing a two-dimensional array as described in Device Memory Accesses.

Size and Alignment Requirement

Global memory instructions support reading or writing words of size equal to 1, 2, 4, 8, or 16 bytes. Any access (via a variable or a pointer) to data residing in global memory compiles to a single global memory instruction if and only if the size of the data type is 1, 2, 4, 8, or 16 bytes and the data is naturally aligned (i.e., its address is a multiple of that size).

If this size and alignment requirement is not fulfilled, the access compiles to multiple instructions with interleaved access patterns that prevent these instructions from fully coalescing. It is therefore recommended to use types that meet this requirement for data that resides in global memory.

The alignment requirement is automatically fulfilled for the built-in types of char, short, int, long, longlong, float, double like float2 or float4.

For structures, the size and alignment requirements can be enforced by the compiler using the alignment specifiers __align__(8) or __align__(16), such as

Any address of a variable residing in global memory or returned by one of the memory allocation routines from the driver or runtime API is always aligned to at least 256 bytes.

Reading non-naturally aligned 8-byte or 16-byte words produces incorrect results (off by a few words), so special care must be taken to maintain alignment of the starting address of any value or array of values of these types. A typical case where this might be easily overlooked is when using some custom global memory allocation scheme, whereby the allocations of multiple arrays (with multiple calls to cudaMalloc() or cuMemAlloc()) is replaced by the allocation of a single large block of memory partitioned into multiple arrays, in which case the starting address of each array is offset from the block's starting address.

本文备注/经验分享：

本次主要说了设备上访存的事项, 这个章节主要讨论了global memory上的访存情况, 什么时候是优化的, 什么时候是不好的. 请注意这里的说法对所有的计算能力都适用(2.0+). 关于每个计算能力的特殊情况, 手册后面还有单独说明. 我这里也主要说一下global memory. 请注意因为之前说过, global memory其实可能会有多种, 一种是用的显存, 一种是用的内存映射成的, 还可能是全自动管理的(unified memory), 或者是其他显卡的显存(P2P Access) ，但你会看到所有的CUDA书籍(几乎全部), 包括NV的手册, 都将这些都看成global memory, 而无论底下的实际存储器的特性.这是因为, 在2.0+的计算能力上, SM总是这样的方式访问global memory:SM ---> L2 --> 后备存储(显存, 内存等等) 所以这里需要说明的是, 和很多人第一次读到这里的理解不同. 这里说的所有特性, 实质上是指的L2的特性, 而不是指的显存.(你在本章节看到的device memory) 正因为L2 cache无法被越过, 所以正常情况下NV会将cache的使用分成两种:有cache的访存, 和无cache的访存，和大部分人的想法相反, 前者是指使用L1(或者变种) + L2，后者是指使用L2。这里的无cache的字样很迷惑很多人。 本章节提到, 一次访存将进行1次或者多次32-Byte的, 或64-byte的或128-byte的传输, 这里同样是指的L2 cache，L2到SM的通道是32B宽度的, (每个SM都有1个32B的到L2的通道),而L1在SM内部的读写宽度是128B的.这就实际上导致了一个问题,如果总是启用L1 cache, 而L1 cache在SM内部发生了一次cache line / block (一个临近的小区域) miss (内容不在L1里), 则L1会连续的从L2传输过来4次32B的内容, 以满足自己的128B连续存储要求。这样的话, 一个程序根据实际的不同, 一次访存可能会导致一次或者多次的L2的传输的.所以这也是手册为何说(后面), 为何有的kernel, 禁用调一级缓存后, 只使用二级缓存(L2强制总是启用的, 禁用不掉), 反而性能会提升.而有的kernel需要使用L1才有较好的性能.因为有的程序一旦读取的太分散, 只使用L2的话, 只会有单次的不间隔的32B传输,而一旦启用L1, 一次miss将导致连续的多次4x32B传输的, 而程序只分散的需要里面的少量内容,用L1会导致额外的过量读取/传输, 浪费性能.为何要提到L2的传输大小是32B?因为NV的L1/L2传输的128B/32B大小不同, profiler所报告的L2 transactions数目是指的32B计算的,不清楚这个,在使用profiler分析访存方面的问题的时候, 会迷惑(如果你认为是128B或者是warpSize * 每个线程请求的大小的话).

这是前言. 本章节实际上主要是说, N卡的访存有几个需要注意的问题: 对齐方面的问题. 对齐分为两种, 一种是访存读写的元素来说的. 另外一种是从warp整体所形成的读写范围说的. 首先说一下前一个. 任何一个warp中的线程, 如果要读取1B(字节), 那么该线程给出的这个字节的地址必须能被被1整除; 如果要读取2B, 那么该线程给出的访问该2B的地址, 必须能被2整除; 如果要读取4B, 那么该线程给出的访问该4B的地址, 必须能被4整除. 类似的对8B, 16B也适用.举个例子来说, 如果有一个float值, 放置在global memory中,那么这个float值(4B大小)的地址, 必须对齐到4的边界(能被4整除)。如果有一个人, 给出的地址是不对齐的, kernel将直接挂掉(而不是这里说的读取到错误的值).类似的, double必须对齐到8B的边界(地址能被8整除)，而uint8_t或者char这种, 没有对齐要求(因为能被1整除是任何整数都具有的特性)。请注意, 这种元素地址(从每个线程的角度)只要不对齐, 访问global memory必须会挂掉kernel. 而后一个, 则是从warp整体来说的,如果warp中的每个线程所读取的元素都对齐好了, 那么实际上的访存请求被SM里的Load/Store Unit(LSU)会被先整理一下, 再发出请求的.这里的整理包括对warp中的32个线程的请求进行合并, 选出多组合并后连续的地址, 再从这些地址的最小的位置 + 长度信息, 进行发出请求. 举个例子来说: __globa__ void your_lady(.... float *p ....) { int id = threadIdx.x; p[id] = 8888; } 如果有这么一个代码, 我只启动了1个只有32个线程的block,那么你会看到, 代码实际上将写入32次8888(32个线程么)，分别这些8888的位置是: p, p + 1, p + 2, ...p + 31 因为这里的float *p来自cudaMalloc, 自带元素对齐属性(这个后面说)，所以这里只需要考虑前面说的方面2, warp整体即可.首先SM里的LSU会整理一下这32次访存, LSU会发现所有的32个访存地址上是连续的,于是它会整理成:从p地址开始的连续32*4B=128B, 均需要将内容设定成8888，这样最终会形成4次L2传输, 每次都是32B. 类似的, 读取也是这样: __globa__ void your_lady(.... float *p ....) { int id = threadIdx.x; .... = p[id]; } 同样LSU会整理一下, 然后连续的从L2载入32*4B(以4次32B的形式, 因为之前你知道L2到SM的通道是32B的),或者如果L1命中, 只有1次从L1的载入, 1 * 128B.此时SM不对外发出请求. 这是最完美和最简单的情况. warp整体的请求是全部能连续合并的, 而且warp请求的最低地址(p这里. 从p开始的128B么, [p, p + 128)也就是), 是满足一定的对齐要求的. 如果你看这个代码: __globa__ void your_lady(.... float *p ....) { float *p2 = p + 1; int id = threadIdx.x; .... = p2[id]; } 首先说, 单个线程本身是对齐到元素的(4B), kernel不会挂掉.其次, SM将该warp中的32个线程的请求整理一下后, 发现地址是从:[p2, p2 + 128)的连续128B的 因为P2是P + 1(+1是指的元素+1, 地址是+4, 因为一个float是4B)，所以实际上的请求是[p + 4, p + 132)，因为p来自cudaMalloc, 有一定的最低对齐方面的特性. 所以这里的最终整理后的范围, 在发出请求之前, LSU会发现无法形成连续的32B传输的。最终它会形成5次L2的传输(多了一次): 传输[p, p + 32), 不要里面的[p, p + 4)部分 传输[p + 32, p + 64), 里面的内容都要 传输[p + 64, p + 96), 里面的内容都要 传输[p + 96, p + 128), 里面的内容都要 传输[p + 128, p + 160), 不要里面的[p + 132, p + 160)部分 你会看到如果warp整体不对齐, 但是warp里面的分量线程对齐到元素, 会正常执行kernel(不会挂), 但会导致额外的传输. 例如如果这次走L2, 会导致5次而不是4次传输, 有用的数据只有80%. 这就是手册本章节说, warp整体导致的传输, 例如这里的32B的L2传输, 必须需要对齐到这种传输的大小的倍数边界的原因, 不对齐就会被LSU自动整理对齐, 但有额外的浪费的传输. 所以实际上: (1)从线程的角度, 读取的元素不对齐: kernel直接挂掉 (2)从warp的角度, 合并后的范围不对齐: 会自动拆分组合成对齐的传输, 浪费掉一些内容, 性能下降. 但kernel不会挂掉. 就是说, SIMT(还记得它吗)会给人一种有很多线程, 每个线程你都能独立运行的假象, 但实际上是warp为整体执行的.既然是同时存在线程执行(假象)和warp执行(实际)的两种.那么任何访存都要从这两种的角度来说. 嗯嗯. 所以就有了刚才的(1)(2)点。

什么是对齐？

对齐是对地址来说的, 而地址就相当于一排房子, 你家的门牌是88号, 邻居的门牌是87号, 而你另外的一个邻居是89号，这里的号, 就是地址.而地址对齐, 就是这个号本身能否被某个数整除.例如你家的88号地址很不错, 能对齐到4，就是说, 你家的地址是4的倍数.而你的两个邻居都没有对齐.

而家里要住人的,有的人比较高贵, 需要住特种的编号的房子,例如你的88号的房子,而不能住87号的.这就是有的人对对齐性有要求.一个人叫float, 这个人有4B, 也就是4个房子那么大的体积,他要求入住的时候, 自己的头必须在88号这种位置.所以他能住在你家.然后上身在你的邻居89号房子，下身在你的邻居的邻居90号房子，脚在91号房子.这就是float这种元素(巨人)的对齐性要求.类似的, 这种巨人还能住在0号房子, 4号房子(均是他的头部所在的位置)，但不能从7号房子这种不对齐的地址开始住.因为一旦你试图这样, 他会愤怒的拆掉你的房子的(kernel会挂掉).所以你知道了他有对齐性要求, 安排房间的时候就不会给他不对齐的地址了.这就是从线程的角度来说, 每个元素的要求.类似的, double是一种更大的巨人(8B),它必须要求住在对齐到8的倍数的这种房间的.他需要连续8个房子住自己的头, 脖子, 上身, 下身, 膝盖, 脚腕, 脚指等等.而头必须在0, 8, 16, 24, 32, 40, ...这种房子，当然也包括你家的88号的房子.类似的还有half, half是一种比较小的人,只需要2B, 也就是两套房子就能住下,

回到前面, 每个访存的线程, 必须满足他要读写的元素(巨人)的基本对齐要求, 不满足巨人就会拆房子, kernel挂掉.但实际上, 你是知道的, N卡用了一种叫SIMT的东西, 来假象的支持海量线程,实际上它执行的是warp, 32个线程一组.所以任何实际点的讨论均不能离开warp.而如同你的线程有对齐性要求, warp他也有.只不过warp的要求比较宽泛, 对齐不好, kernel不能挂掉, 只是会顶多损失点性能吧了.(不在乎性能的可以不看下面的)。

如果台下的诸位只需要能CUDA入门, 可以无视warp，有了warp的概念和知道针对它的优化的细节, 性能会更好,但不知道, 程序不会出错, 顶多可能会慢一点，就如同这里的线程的元素对齐性你必须知道(不知道就kernel挂了)，而warp你最好知道(性能会更好), 实在不知道一样可以用CUDA.完全无障碍的.

但是手册既然说了warp, 这里我们还是要阅读一下的.否则就浪费手册的好意了。

回到warp的对齐性要求来说. warp的对齐性要求体现在整体上.这个整体是指的warp中的32个线程的访存操作被合并整理后的结果. 手册本章节给出了多种情况，每种情况都有自己的特色. 这个马上就说.

而warp在进行访存的时候, 是直接和L1或者L2打交道的.(L2再和后备存储器, 例如显存打交道, 那个是另外一个事了)，而每次warp从SM和L1/L2打交道的时候,如果本次打交道的是L1, 则传输大小总是128B, 无论这128B是否满足warp整体整理出来的访存范围.而如果本次warp打交道的是L2, 则传输大小总是32B, 无论这32B是否满足warp整体整理出来的访存范围.因为手册本章节说了, L1总是从128字节的对齐地址连续传输128B的。而L2总是从32B的对齐地址(能被32整除的地址)连续传输32B的，如果warp整体的访存范围和L1/L2的传输范围不符合,那么SM的LSU会自动请求最接近的一些范围, 舍弃掉其中不要的,这样来满足warp整体的访存完成, 这些不要的, 不对齐的等等,不会导致kernel挂掉,但可能因为warp整体访问的不是最优范围, 有一定的性能损失.而刚才的3段代码, 演示了warp进行L2访问的情况下, 如果范围不对齐到32B(L2自身的传输范围)的边界,会如何使用了5次传输, 而不是4次, 来自动拼凑成满足的范围.来完成warp整体的访问的.这就是刚才一共说的(1)(2)点中的(2)点, warp整体可能会导致不优化的访存,但只要单个线程的元素是满足对齐要求的, 则整体整理出来的范围无论是否满足对齐要求,均可以进行.无非是性能上的好坏变化而已.

甚至这种变化可能会导致性能非常糟糕:

举个例子来说, 每个线程都只间隔17B访问1B, 例如:

char *p

p[tid * 17 + 99999]

tid是线程编号, warp中的每个线程都隔离了17个字节访问1次.这样会warp整体整理出来一大堆需要访问的范围的.每个范围至少都需要1次L2的32B传输(先不考虑L1),而里面有用的只有大约1-2个字节，传输效率只有1/16-1/32(为何是1-2个? 有可能有2个线程能访问在同1个32B里, 也有可能不在)，也就是6%-3%左右的访存效率.此时性能损失的很严重.然而, kernel依然可以正常的执行下去. 只是会变慢.类似这种的.

所以本章节说, In general, the more transactions are necessary, the more unused words are transferred in addition to the words accessed by the threads, reducing the instruction throughput accordingly.

一般的, (对于warp整体整理出来的访存请求)需要越多次传输, 则会传输更多的需要的数据(words)之外的无用数据.而指令的吞吐率也会相应的越来越下降.(为何这里是指令. 因为访存也是指令. 只不过不是计算指令. 是访存指令罢了. 根据硬件的不同, 一条warp整体访存可能会随着里面的范围约零散, 被replay的次数越多, 执行的越慢, 有效的传输比例也越低)

For example, if a 32-byte memory transaction is generated for each thread's 4-byte access, throughput is divided by 8.

这里是说, 如果每个32B传输里面, 只有4B有用的, 则吞吐率下降到1/8

例如这种:

float *p

p[tid * 8]

每个线程越过8个元素(乘以8么. 0, 8, 16, 24, 32...)访问1个元素.每个元素是4B的(占据4个房子的巨人)，注意元素地址, 和字节地址的区别, 这里有4倍的关系，那么每个32B传输中, 有用的只有4B，其他跟随传输的28B都浪费了.效率只有1/8的.(具体情况如何, 看手册后面的计算能力具体介绍)

总之你应当记住

尽量让warp整体整理出来的范围尽量在一起，不要分散，整理出来的最低地址范围也最好是32B或者128B的倍数这种. 一个特例是: 范围完全在一起(中间没有无用的不要的元素), 首地址也在32B或者128B的倍数,这种叫"充分合并的访存"，这种是性能最好的。这些都是SIMT带来的问题.

他和CPU上真正的线程自由执行的区别就在这里.SIMT看上去能像CPU那样的自由执行海量线程,但是他不能的.不考虑这些, 性能就会下降.但是假象总是存在的---你的kernel还是能正常执行的.只是快慢而已.所以为何本章节是优化global memory的访问.

后面还说了其他的, 我简单说一下.

(1)自带的基本元素类型, 例如float, double, float2, float4, int, short都自带了基本的对齐性要求，这种要求很好记. 都是元素多大, 要求就多大.例如float是32-bit(4B)的, 那么需要地址能被4整除, 对齐到4B边界.double需要64-bit(8B),那么地址需要被8整除, 对齐到8B边界.

支持的最大基本元素是16B的(例如double2), 这种需要对齐到16B边界.而如果用户自己定义了一种结构, 例如这里给出的例子:

struct __align__(8) {

float x;

float y;

};

默认是没有对齐要求的(或者说, 对齐到里面的基本元素)。

如果用户想模仿一下类似自带的float2的对齐(对齐有更好性能)，可以手工通过__align__(XXX)指定。其中XXX是你想要对齐的最低要求，通过这种指令, 可以尽量减少访问你自定义的结构体所需要的基本访存指令的数目(基本访存是指, 1B, 2B, 4B, 8B, 16B这种指令)

完成同样的访问, 指令一般情况下, 数量越少越少.性能会越高.但是你需要注意的是, 虽然手工指定了__align__提高了性能，但也提高了要求(对齐性要求的更高了), 也更容易造成你写代码时候出错.这是一种取舍的问题.程序员应当在书写的时候作出易用性和高性能之间的取舍.

手册还给出了第二个例子:

struct __align__(16) {

float x;

float y;

float z;

};

这个例子手册没有说任何解释.但这里是一个著名的CUDA的坑.用户这里自己定义了一个有3个float的结构体,但却要求对齐到16B，这个有3个结构体的float看上去像是float3，但实际上它不是. 真正的float3在CUDA中是要求对齐到4B的,也就是说, CUDA自带的真正的float3只要求对齐要里面的1个元素, 这样自带的float3的任何一次访问都会被拆分成3次4B访问(因为整体不满足对齐要求)，而这个例子给出的伪float3, 性能更好,因为他要求对齐到16B的边界.每次访问只会生成1条16B的指令即可.但是需要注意, 它比自带的float3有更高的要求, 同时还浪费了隐形的最后4B(因为12B是16B中有效的3/4)，很多老人也容易栽倒在这上面——注意，这段代码是对的. 但容易让人出错，用好了, 不会出错. 性能还更好，用不好. 挂了。

然后最后一段还说了一个重要问题.就是对于看了昨天的内容的人来说,他往往会试图拼凑多个小缓冲区, 拼凑成一个大缓冲区.(因为昨天说了, 多次小传输不如一次拼凑起来的大传输)，多次小缓冲区分配往往使用多次cudaMalloc，而一次大缓冲区拼凑分配往往提前计算需要的小缓冲区累加起来的大小, 然后总共一次cudaMalloc分配完，这里实际上是会导致一个问题的.因为cudaMalloc自带提供一个很高的返回的基地址对齐性,手册里说是对齐到256B的边界.所以实际上你的多个小缓冲区都默认对齐到足够适合你用的超级宽的要求上了.而你自己拼凑的往往没有注意到这点.这样以往不小心利用了这点的, 分配了多个小缓冲区的代码, 在合并后可能就会挂掉.因为现在是你自己手工推算指针偏移量进行合并的，一不小心就会挂掉的.所以手册提醒你千万要注意这点.

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

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