DAY38：阅读春存储器修饰符

sisiy

246

DAY38：阅读春存储器修饰符

« 于: 六月 22, 2018, 02:42:00 pm »

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:

程序代码: [选择]

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];
    ...
}

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:

程序代码: [选择]

void foo(const float* __restrict__ a,
         const float* __restrict__ b,
         float* __restrict__ c);

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:

程序代码: [选择]

void foo(const float* __restrict__ a,
         const float* __restrict__ b,
         float* __restrict__ c)
{
    float t0 = a[0];
    float t1 = b[0];
    float t2 = t0 * t2;
    float t3 = a[1];
    c[0] = t2;
    c[1] = t2;
    c[4] = t2;
    c[2] = t2 * t3;
    c[3] = t0 * t3;
    c[5] = t1;
    ...
}

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.

已记录

sisiy

246

(无标题)

« 回复 #1 于: 六月 22, 2018, 03:15:37 pm »

如同昨天说过的所有静态/动态分配的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后, 编译器可以安心的做公用表达式消除优化，(这是一种常见的优化, 也是很多新人常见的问题:

已记录

sisiy

246

(无标题)

« 回复 #2 于: 六月 22, 2018, 03:18:16 pm »

引用自: sisiy 于六月 22, 2018, 03:15:37 pm

如同昨天说过的所有静态/动态分配的shared memory/global memory一样, unified memory也有两种分配方式.

...

例如有人问: 我有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__的, 否则享受不到性能的提升.因为编译器只有在这种情况下, 才能安全的提出公用的表达式, 进行优化.

已记录

sisiy

246

(无标题)

« 回复 #3 于: 六月 22, 2018, 03:28:01 pm »

此外, 在一定的计算能力下, 配合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描述)和编译器(翻译官)的共同工作.

已记录