DAY38：阅读存储器修饰符

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

本文共计468字，阅读时间15分钟

前情回顾：

DAY36：阅读”执行空间"扩展修饰符

DAY37：阅读不同存储器的修饰符

B.2.4. __managed__

The __managed__ memory space specifier, optionally used together with __device__, declares a variable that:

· Can be referenced from both device and host code, e.g., its address can be taken or it can be read or written directly from a device or host function.、

· Has the lifetime of an application.

See __managed__ Memory Space Specifier for more details.

B.2.5. __restrict__

nvcc supports restricted pointers via the __restrict__ keyword.

Restricted pointers were introduced in C99 to alleviate the aliasing problem that exists in C-type languages, and which inhibits all kind of optimization from code re-ordering to common sub-expression elimination.

Here is an example subject to the aliasing issue, where use of restricted pointer can help the compiler to reduce the number of instructions:

In C-type languages, the pointers a, b, and c may be aliased, so any write through c could modify elements of a or b. This means that to guarantee functional correctness, the compiler cannot load a[0] and b[0] into registers, multiply them, and store the result to both c[0] and c[1], because the results would differ from the abstract execution model if, say, a[0] is really the same location as c[0]. So the compiler cannot take advantage of the common sub-expression. Likewise, the compiler cannot just reorder the computation of c[4] into the proximity of the computation of c[0] and c[1] because the preceding write to c[3] could change the inputs to the computation of c[4].

By making a, b, and c restricted pointers, the programmer asserts to the compiler that the pointers are in fact not aliased, which in this case means writes through c would never overwrite elements of a or b. This changes the function prototype as follows:

Note that all pointer arguments need to be made restricted for the compiler optimizer to derive any benefit. With the __restrict__ keywords added, the compiler can now reorder and do common sub-expression elimination at will, while retaining functionality identical with the abstract execution model:

The effects here are a reduced number of memory accesses and reduced number of computations. This is balanced by an increase in register pressure due to "cached" loads and common sub-expressions.

Since register pressure is a critical issue in many CUDA codes, use of restricted pointers can have negative performance impact on CUDA code, due to reduced occupancy.

本文备注/经验分享：

如同昨天说过的所有静态/动态分配的shared memory/global memory一样, unified memory也有两种分配方式. 本章节说的__managed__即是静态分配的. 例如说: __managed__ int a; 只要你的kernel能运行, 它就可以直接使用a, 而不需要考虑a的空间是何时分配的之类的问题, 比较简便. 此外, 如同所有的unified memory特性一样, a能被CPU和GPU都访问到(还记得我们之前的章节说过, unified memory是升级版本的zero-copy memory吗?) 在很多场合用起来非常简单. 特别的是, 当kernel产生了一个较大的结果, 例如填充了128MB的结果缓冲区, 而你只需要根据结果缓冲区中的情况, 不可在写代码的时刻预测的, 只使用里面的, 例如16MB的内容，那么应当考虑使用unified memory, 较新的GPU硬件能按需的为你回传你需要访问的内容。 感兴趣的人可以看一下unified memory里面的如何CPU或者GPU按需的page-fault, 然后自动传输的，此时的性能将可能会超过手工的传输(全部)。但这里不详细说明unified memory, 后面将有章节单独描述它. 这是一个很给力的特性, 但一般不推荐在Windows上使用(Windows上一定情况下退化成普通的zero-copy memory, 而且性能很惨)。

这是说的静态分配. 和__managed__对应的是cudaMallocManaged, 它将动态的分配一块managed/unified memory(就如同普通的__device__对应的是cudaMalloc*()一样, 后者将动态的分配普通显存)。其实静态分配有很多好处, 有些要求严格的项目是很多时候不允许动态分配的，因为静态分配只要运行起来了, 存储器的使用情况不会发生变化, 程序要么一切正常的运行了, 要么运行不起来(资源不足)，而动态的则可能随着运行中的数据变化(例如你的分配情况依赖于一个具体的数据), 有不可预测的后果，例如运行正常15小时后, 突然挂掉.这对类似雷达测量之类使用GPU的场合有时候是不可接受的。

等等了. 所以本章节(和昨天的那部分)中的静态和动态分配都有各自的用途* 然后本章节的另外一部分则告诉你, 如何更有效的使用指针。 CUDA引入了一个__restrict__的扩展来设定C99的Restricted Pointer (VC和GCC有各自的关键字, 无非是前后的下划线和restrict的位置的区别不同, 而CUDA C的版本则是前后都有两个下划线) 例如VC版本的是__restrict 使用restrict解决了一个重要的C语言里的问题.就是指针不再像以前那样能乱指了(被restrict了么)，如同本章节的说法, 有的时候, 编译器无法进行一些优化, 因为它会按照最坏情况进行估计, 例如本章代码的: void foo(const float* a, const float* b, float* c) { c[0] = a[0] * b[0]; c[1] = a[0] * b[0]; c[2] = a[0] * b[0] * a[1]; c[3] = a[0] * a[1]; c[4] = a[0] * b[0]; c[5] = b[0]; ... } 这里面使用了3个指针a,b,c，无法确定这三个指针是否有任何重叠, 例如实际上可能只有1个缓冲区, 而a,b,c只是它们的别名而已。 例如昨天的章节有人问, 使用了多个extern __shared__ 会如何?此时将会产生重叠/重名的指针(alias) 通过一个指针写入, 很可能改变另外一个指针指向的内容的状态.而__restrict__的出现改变了这一点， type * __restrict__ a type * __restrict__ b type * __restrict__ c 有这样的3个指针. 通过__restrict__修饰后, 你暗示编译器各自指向的内容只能通过各自的a,b,c指针进行访问. 暗示它通过一个指针的写入改变了另外一个指针的值.因此编译器可以放心的进行一些优化, 例如通过b写入后, 不必担心a里面之前读取到的值是否已经改变, 是否需要重新读取, 而可以安心的使用老值。 此时有助于减少无辜的生成的指令. 提高性能.请注意, 这只是一个暗示, 如果你暗示了编译器是一套, 但是做的是另外一套做法，例如, 有的读者比较调皮, 想尝试一下使用重叠的3个指针, 却告诉了编译器是__restrict__的, 结果会如何？那么编译器编译出来的代码很可能运行出错, 请不要这样做(编译器无法在编译时刻检测到你的指针有重叠, 也无法在运行时刻检测到你有重叠, CUDA C和C均不是具有完备的Runtime的语言, 这样做将导致未定义的结果)，所以如果一旦要使用__restrict__来暗示CUDA C编译器, 就一定要做到你的指针使用行为和你所暗示过的一致. 否则将导致未知后果(例如kernel挂掉)，类似的, 本章节还提到, 含有数组元素访问, 例如a[0] * b[1] + c[2] - d[3]这种代码的公用表达式, 在使用了__restrict__的指针a,b,c,d后, 编译器可以安心的做公用表达式消除优化，(这是一种常见的优化, 也是很多新人常见的问题: 例如有人问: 我有3行代码: a * b + c + e a * b + c - g a * b + c - f 它们都含有a * b + c的部分, 我感觉这样编译器会生成冗余指令, 我是不是应当手工提取表达式出来, 只计算一下, 像这样: t = a * b + c t + e t - g 以及, 用t - f 这样能提高性能吗? 答案是你不需要这样做, 因为现代的编译器都具有公用表达式消除能力, 通常情况下的重复的代码部分均将被自动提取出来, 只计算一次的. 因此无需手工处理。手工处理还降低了代码的可读性. 却得不到想象中的性能提升的。而本章节则说了, 如果是使用指针, 必须是__restrict__的, 否则享受不到性能的提升.因为编译器只有在这种情况下, 才能安全的提出公用的表达式, 进行优化。

此外, 在一定的计算能力下, 配合const + __restrict__一起使用, 可以使用SM里面的类似L1的Read Only Cache或者Unified Cache之类的东西, 此时应当考虑一起使用它们. 有助性能提升, 这个当年NV在Kepler的时代, 特意强调了很多次.在多次的GTC的演讲中提到这个问题.因为Kepler当年是一个很难发挥全部性能的卡, 如同前几章说过, 需要TLP + ILP都手段一起上, 才有可能多少发挥出来性能(甚至一起上各种手段都发挥不出来性能)，而一起使用了const + __restrict__后, 有助于大量使用read only cache, 还有助于编译器自动进行ILP。在实际的Kepler卡上, 这样做后, 编译器能在生成的指令中, 将你的代码打乱顺序, 你可能在行3, 行80处, 行90处都有1处访存读取, 对于有这2个修饰的指针, 编译器可以以增加寄存器使用量的代价, 将后续的很远位置的访存, 自动重新调整顺序, 提到前面, (例如等效的在程序开头连续进行了3次访存, 读取你行3, 行80, 行90需要的数据)，此时等于进行了全自动的ILP, 在Kepler这种卡上, 意义重大.可是K80上应当尽量使用。 虽然本章节说了, 这样做有可能增加寄存器使用量(你知道使用过多寄存器有可能会反而降低性能的),但是K80是一张好卡。Kepler里面的唯一良心.它的一个SM是两个SM拼凑起来的, SP(计算单元)数量不变的情况下,其他资源基本都翻倍了(例如寄存器翻倍了)，应当黑用. 性能往往都是正面的.大致如此吧. 总之的一点是, 能有明确的指针(或者数组的名字)的使用, 就应当直接就地使用,尽量使用下标/偏移量变换,而不应当多使用指针变换.前后虽然是等价的。但后者很可能有效的迷惑编译器, 生成较低质量的代码。毕竟代码的生成是你(使用CUDA C描述)和编译器(翻译官)的共同工作.

此外, 再重复一点, 不使用指针变换, 而总是使用下标或者偏移量变换,是维护代码可维护性的一个很关键性的因素.已经见过无数后来人在维护前人的代码(例如前面的同事离职了), 陷入大量指针推导/指针变换的陷阱中无法自拔.

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

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