DAY32：阅读local Memory

GPUS Lady

发布于 2018-06-25 16:34:39

5540

发布于 2018-06-25 16:34:39

文章被收录于专栏：GPUS开发者GPUS开发者

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

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

Two-Dimensional Arrays

A common global memory access pattern is when each thread of index (tx,ty) uses the following address to access one element of a 2D array of width width, located at address BaseAddress of typetype* (where type meets the requirement described in Maximize Utilization):

BaseAddress + width * ty + tx

For these accesses to be fully coalesced, both the width of the thread block and the width of the array must be a multiple of the warp size.

In particular, this means that an array whose width is not a multiple of this size will be accessed much more efficiently if it is actually allocated with a width rounded up to the closest multiple of this size and its rows padded accordingly. The cudaMallocPitch() and cuMemAllocPitch() functions and associated memory copy functions described in the reference manual enable programmers to write non-hardware-dependent code to allocate arrays that conform to these constraints.

Local Memory

Local memory accesses only occur for some automatic variables as mentioned in Variable Memory Space Specifiers. Automatic variables that the compiler is likely to place in local memory are:

· Arrays for which it cannot determine that they are indexed with constant quantities,

· Large structures or arrays that would consume too much register space,

· Any variable if the kernel uses more registers than available (this is also known as register spilling).

Inspection of the PTX assembly code (obtained by compiling with the -ptx or-keep option) will tell if a variable has been placed in local memory during the first compilation phases as it will be declared using the .local mnemonic and accessed using the ld.local and st.local mnemonics. Even if it has not, subsequent compilation phases might still decide otherwise though if they find it consumes too much register space for the targeted architecture: Inspection of the cubin object using cuobjdump will tell if this is the case. Also, the compiler reports total local memory usage per kernel (lmem) when compiling with the --ptxas-options=-v option. Note that some mathematical functions have implementation paths that might access local memory.

The local memory space resides in device memory, so local memory accesses have same high latency and low bandwidth as global memory accesses and are subject to the same requirements for memory coalescing as described in Device Memory Accesses. Local memory is however organized such that consecutive 32-bit words are accessed by consecutive thread IDs. Accesses are therefore fully coalesced as long as all threads in a warp access the same relative address (e.g., same index in an array variable, same member in a structure variable).

On some devices of compute capability 3.x local memory accesses are always cached in L1 and L2 in the same way as global memory accesses (see Compute Capability 3.x).

On devices of compute capability 5.x and 6.x, local memory accesses are always cached in L2 in the same way as global memory accesses (see Compute Capability 5.x and Compute Capability 6.x).

本文备注/经验分享：

昨日主要说了如何恰当的使用设备上的多种存储器中的global memory, 但还有一部分没有说完. 今天的章节首先承接前一天的说完global memory, 然后会谈到local memory, constant memory, shared memory等等. 从昨天的内容你应当知道, 直接从cudaMalloc分配出来的global memory会对齐到256B的边界, 这几乎对所有的访问来说, 对是一个足够的对齐位置了.但在实际应用中, 我们往往需要使用, 类似2维数组这种结构,如果2维数组的行宽(字节单位)是一个比较奇特的数字, 例如一个777 * 555的8-bit灰度图片,此时如果此图片的首行的开头地址能对齐到256B, 则他的下一行肯定不能(因为第一行的地址, 例如是0或者2560这样的, 能被256整除, 则下一行的首地址是777或者2560 + 777, 并不能被256整除），这样就造成了一个尴尬的场面, 如果首行对齐了, 那么下一行就无法对齐.如果要保证下一行对齐, 那么首行和再下一行就不能对齐了.于是CUDA提供了一个专门的分配函数, 叫cudaMallocPitch(),该函数会要求你输入行宽, 和高度信息, 然后自动对行宽进行填充, 例如说, 你实际要求了777的行宽, 但该函数会最终按照1024(假设的. 后面会给解释)进行分配,这样按照1024B的实际行宽进行了分配后, 实际上下一行的开头也将对齐到256B的边界, 如此类推再下一行.因此这种可以保证从每行的开头都是恰当的对齐了的.而实际的宽度, 在这里叫pitch(实际的间距).请注意这种分配方式可以保证每行的开头都会对齐到一个很好很恰当的位置(具体对齐到什么程度和计算能力有关), 但需要付出, 可能会有每行尾部的浪费的一些存储空间, 例如原本我只需要分配777字节即可, 现在我一行分配了1024字节(依然是假设的), 则我浪费了24%的存储器容量.虽然造成了浪费, 但因为每行的对齐效果, 会可能造成性能上的提升(也可能不会). 因此有的时候还是值得的. 但是手册本章节说的, "significantly"的性能提升, 那是不可能的. 这个说法来自手册的早期版本(计算能力1.X), 只有在那个老版本的卡上才会如此, 而目前1.X和2.X的老卡均不再被支持了, 只是这里的说法没有修改而已. 类似的这里还需要注意一点, 则是本章节对2维数组开头的说法有点不恰当. 特别是这里: For these accesses to be fully coalesced, both the width of the thread block and the width of the array must be a multiple of the warp size. 在实际中, 一个block如果是W线程宽 * H高, 数组的行宽只需要是W * sizeof(元素)的倍数即可. 只需要完全合并才有性能的说法, 只见于1.X (甚至可以上面这个要求都可以不要, 因为2.0+的计算能力的cache提供了一定的buffering和缓冲效果) 然后在实际的今天, 对显存的使用量越来越大, 而各级缓存已经成为GPU的标配,在没有特殊的情况下(例如某个代码卡显存或者L2传输带宽, 每行不对齐会影响性能), 基本上现在这个cudaMallocPitch已经用途不大了,正常大部分情况下的使用, 它基本不会带来性能上的提升(特别是对于卡计算的kernel), 却可能带来不小的显存上的浪费.因此现在的用户应当考虑通过profiler分析这样带来的浪费和可能的性能提升之间的权衡. 并非一定要这样做(虽然手册说要这样做).特别是刚才说了, 有些应用基本会将显存用满, 此时可能用了pitch分配后程序就跑不起来了(显存不够了) 此外, 长期的另外一种对齐每行的方式是不使用该函数,而是手工的将每行进行对齐.例如宽度N(字节单位, 下同. 如果以后遇到元素为单位的, 我会特别说明, 不说明就是字节单位), 我想对齐到32B的边界(例如我知道我某kernel总是会bypass L1的),则可以进行(N + 31) / 32 * 32。

类似这种的方式进行向上取整(或者其他任何你喜欢的向上取整方式, 包括用if判断是否整除之类的, 均可).这种方式有两个好处: (1)是用户能自己控制每行的对齐程度, 例如32B这样的较小的值. (2)因为cudaMallocPitch这种是自动针对特定的计算能力选择一个较优的对齐程度, 例如在一些计算能力上为128B或者256B边界, 使用手工对齐可以不必担心过度的浪费, 和可能的造成显存不够用的忧虑. 此外, 本章节的最开头有一部分说的"Type需要满足Maximize Ultilization(之前的一个章节)"的说法, 可能是编写手册的人喝醉了, 这里不能明白这里的含义.（最开头的确有问题的) 在实际中, NV就自带一个例子, 使用了一个16 * 16的block形状(16,16), 进行矩阵的运算(乘法或者转置)，而这并不满足手册开头的说法的, 至少blockIdx.x并非是手册说的是warpSize的倍数.因此可见我们看看就好. 不必较真.该例子应当在后面就有. 此外, 除了昨天和今天前一部分的global memory外(这是我们能用的最大的, 最常接触的存储器, 因此说了这么多), CUDA常用的存储器还包括local memory, shared memory, constant memory等等. 今天的下一部分就是要说local memory. 在这些能涉及到的存储器中, 实际上往往global memory对应的是显存(也有其他情况), 而shared memory则对应的是SM内部的一段快速的SRAM(容量小不过). 而local memory和constant memory则没有独立的实体存在,并不存在特殊的这两种专用的存储单元的。

而是他们同样的使用显存(或者内存, 看情况), 只不过在缓存策略上和global memory略微不同罢了. 所以实际上, 今天的local memory更多的是一个虚拟上的概念, 用户将它看成是特殊的global memory就好. 这样很容易理解，它的特殊性在于, local memory的地址只对本线程有效. 每个线程的同1个local memory地址(例如都是int *p), 会被自动在最终访存的时候变换成不同的显存位置. 手册这里实际上说到了这点. 如果warp中的32个线程都访问了同一个local memory中的地址, 例如说, 我有一个变量叫float a, 而a在local memory中, 或者换句话说是&a(注意&符号)指向local memory中.则在读写a的时候(例如写入值1.2345f), 则硬件会自动将该地址映射到一个32个float副本(4B副本)的位置, 实际上将自动完成一次完全合并的32个4B的写入. local memory的这个完成地址变换到, 显存上的能完全合并的地址的过程是对用户来说是全自动的. 用户不需要担心.一定情况下可以将local memory理解成能做自动在warp的多个线程间(lanes间)自动自动合并的显存即可. 但是这种特性也带来了额外的代价, 如果一旦warp的每个线程的没有访问同一local memory的地址, 则这种自动变换位置写入32个临近的值的效果不会发生, 甚至可能会更糟糕:例如有一个较大的数组(per thread的数组): float my_big_array[666]; 第一个线程访问了自己的my_big_array[0], 第二个线程访问了自己的my_big_array[1], 如此类推, warp中的第31个线程访问了自己的my_big_array[31]，那么最终在local memory中的访问非但不能合并, 在最后映射成了显存的实际地址的时候,会实际的访问float *s; s[tid * 33]类似这种的。那么最终在local memory中的访问非但不能合并, 在最后映射成了显存的实际地址的时候,会实际的访问float *s; s[tid * 33]类似这种的。这种会形成上一章节中的非常Scattered的global memory访问效果的.此时访问效果不会很好.(实际上只有同一个地址, 例如my_big_array[3], 会被映射成s[3 + tid]的, 这种才可以, 而且是合并的. 刚才说过了，后面会说到, 类似这种每个线程的坐标都不同的访问, 适合的是shared memory, 但这是下一章节的内容了) 本章节还说了, 往往有很多情况会导致编译器使用local memory, 而不是更快的寄存器的. 其中的一个很重要的原因是: N卡不支持寄存器索引, 也就是说, 指令可以明确的说, 我需要访问寄存器R3.但是指令不能说, 我需要访问寄存器R0后面第三个寄存器.如果你以前有x86的编程经验,你可以说我需要访问rax或者ebx，或者XMM3。但是你不能说我需要访问XMM2后面那个寄存器.因此一旦需要进行索引操作, 编译器除非能发现这个索引操作的目标下标, 能在编译时刻确定下来(例如编译器发现, 虽然下标是个变量, 但值我能知道就是3), 此时编译器才会使用寄存器, 例如刚才说的XMM3，否则它在不能在编译器时刻确定下标的情况下, 它只能使用能被索引的local memory, 而不是寄存器.例如编译器可能会有类似这种安排(实际上不是):float my_shadowed_XMM[3];现在用户可以直接访问my_shadowed_XMM其中i是3了.虽然i哪怕不能在编译时刻确定.类似这种效果.而正常情况下为何local memory总是会将同一个位置扩充32倍, 例如:你在每个线程中定义的是float my_shadowed_XMM[3];而实际显存中对应的位置则是float my_shadowed_XMM[3][32];为何要这样安排是正常的CUDA的warp内部是大部分情况没有分支的,编译器认为你很可能每个线程在都使用同一个下标的情况下, 会映射成访存:my_shadowed_XMM[2][tid]; 这里面的2是每个线程的统一下标.而tid是线程在warp中的编号(精确的说是lane id, 但是手册坚持不引入这个概念, 所以这里也不说).这样每次访问都能自动合并的效果就从这里出来了。

请注意本章节还说了, local memory在不同计算能力上的使用缓存情况, 例如使用使用L1之类的. 感兴趣的可以看一下手册的附录, 里面有更详细的描述.还感兴趣的可以看PTX手册, 里面有不同的cache策略描述, 其中对local memory部分的说明很有特色.此外, 手册还说了几种会被可能放入local memory而不是寄存器中的情况: Arrays for which it cannot determine that they are indexed with constant quantities, Large structures or arrays that would consume too much register space, Any variable if the kernel uses more registers than available (this is also known as register spilling). 第一行的这种就是说, 下标不能在编译时刻确定. 第二种说的是巨大的数组(例如float abc[666]), 你要考虑到目前一个线程最多只能使用255(256)个寄存器的. 666的float abc[]需要的数量超过了寄存器的数目. 第三种则是说, 虽然不是一个巨大的数组, 但零散的各种变量累加起来很多. 也超过了数量. 这些均有可能会导致编译器选择将他们放入local memory, 而不是寄存器.但手册这里说的有点问题, 或者说不是问题, 而是为了照顾初学者的心灵, 避免造成困惑: 并非一个巨大的数组总是全部元素都会被放入local memory的, 如果某些数组中的元素被频繁的使用, 或者下标能被确定(例如, 你总是使用abc[666]里面的abc[0]和abc[233], 就是按照这里的直接给出的0和233这种下标写的, 而不是推导出来等于0或者233), 那么编译器很可能选择将abc[0]和abc[233]放入寄存器.而剩下的其他元素(例如后一半元素之类的)放入local memory. (是不是理解起来好难，这是CUDA手册说的不好. OpenCL规范有另外一种描述方法. 哪种方法好理解的多. 一会我给加上) 另外的, 类似情况还有, 虽然有很多变量, 例如100个变量, 你通过某种方式要求kernel每个线程只能使用64个, 那么并非这100个变量里必须是64个在寄存器中, 36个必须在local memory中的？这不一定.实际的情况可能是, kernel代码的前一部分使用某64个变量非常频繁,后一部分代码使用另外64个变量非常频繁, 例如:我的kernel的1-50行使用变量a0, b, C, X, ....这64个变量非常频繁, 而中间夹杂了偶尔使用剩下的Z, WW, QQ, YY这些变量. 此时前一部分的最终生成代码可能是将非常频繁的这一部分放入寄存器,不频繁的部分放入local memory，而后一部分的使用相反了, 原本前一部分被频繁使用的部分不经常使用了, 而原本偶尔使用的被频繁使用.则local memory中的存储情况可能发生了变化, 原本在寄存器中的那些现在在local memory中了, 而原本在local memory中的不在了.所以说, 这实际上是一个动态的过程, 随着代码的进展不同, 而有不同的变化.实际上, 另外一个类似的规范OpenCL, 是另外一种说法, 我感觉这个说法更容易理解一些:kernel并没有寄存器, 也没有local memory, 它只有一种per thread的存储(OpenCL的线程private memory), 用户所有定义的局部变量之类的, 都在这里面. 但是编译器可以选择, 随时的将寄存器作为这个private memory的私有线程存储的缓存,可以随意安排它想被缓冲(而且能够被缓冲, 例如下标确定的数组中的某个元素)进入寄存器,或者从寄存器移出.这样OpenCL的private memory等于CUDA的寄存器 + Local memory,而用户无需担心具体的存储位置，用户只需要知道编译器尽量为你安排使用快速的寄存器, 而不是慢速的local/private memory即可.不过本章节说的, 什么情况会影响编译器作出存储位置的安排, 还是需要考虑的. 本章节还提到了cuobjdump以及, 编译时刻的--ptxas-options=v参数, 这些都很有用. 欢迎实际的使用一下, 看看效果.用户应当想办法尽量能告诉或者暗示编译器尽量使用寄存器的(例如本章节之前的那3条), 但是无法100%的控制它. cublas之类的能完全的控制, 是因为它们不是编译出来的(手写的汇编)。感兴趣的可以看一下前几天提到的maxas, 看下里面如何用sass汇编手工控制寄存器的使用.但是作为CUDA C或者PTX用户, 我们暂时无法控制这点.实际上不能控制这点是个好事, 手工的寄存器安排很累人的。

此外, 再强调一点, local memory中的东西(包括一个普通的kernel中的局部变量)，它的内容和地址只对本线程有效.