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上发帖

原文发布于微信公众号 - 吉浦迅科技(gpusolution)

原文发表时间:2018-06-25

本文参与腾讯云自媒体分享计划,欢迎正在阅读的你也加入,一起分享。

发表于

我来说两句

0 条评论
登录 后参与评论

相关文章

来自专栏数说戏聊

Python3分析Excel数据

使用xlrd和xlwt扩展包,确定工作簿中工作表的数量、名称和每个工作表中行列的数量。 1excel_introspect_workbook.py

13120
来自专栏PPV课数据科学社区

【学习】《R实战》读书笔记(第二章)

“读书会是一种在于拓展视野、宏观思维、知识交流、提升生活的活动。PPV课R语言读书会以“学习、分享、进步”为宗旨,通过成员协作完成R语言专业书籍的精读和分享,达...

36890
来自专栏图形学与OpenGL

OpenGL开发库的详细介绍zz

开发基于OpenGL的应用程序,必须先了解OpenGL的库函数。它采用C语言风格,提供大量的函数来进行图形的处理和显示。OpenGL库函数的命名方式非常有规律...

34330
来自专栏杨建荣的学习笔记

Python之Numpy初识

今天翻了下计划,要学习Numpy了,所以得调动脑细胞的积极性,看看能有什么收获。 首先得了解下什么是Numpy,从我的印象中,一般提到这个工具都会和机器学习关...

367110
来自专栏Golang语言社区

麻将游戏数据结构和AI算法

用休息时间零零散散写完了网络麻将游戏,感觉其中有不少值得记录的东西。 基础数据结构     数据结构确定决定了程序的开发难易程度,就像是游戏的骨架,对于电脑AI...

1.1K20
来自专栏代码永生,思想不朽

utf8中文字符串的多模式匹配算法的优化

上个月接触到了我组的一个关于在海量文本中匹配字符串业务。读源代码时发现一些问题,并针对这些问题做了优化工作,效果非常明显。

52730
来自专栏SimpleAI

令人困惑的TensorFlow【1】

我叫 Jacob,是 Google AI Resident 项目的研究学者。我是在 2017 年夏天加入该项目的,尽管已经拥有了丰富的编程经验,并且对机器学习的...

11720
来自专栏CreateAMind

pytorch初体验

一部分的内容在2017年1月18日Facebook发行的PyTorch相比TensorFlow、MXNet有何优势? - 罗若天的回答 - 知乎 已有。

13510
来自专栏Play & Scala 技术分享

Scala Macro 现状介绍

39650
来自专栏华章科技

令人困惑的TensorFlow!谷歌大脑工程师帮你解决麻烦

导读:虽然对于大多数人来说 TensorFlow 的开发语言是 Python,但它并不是一个标准的 Python 库。这个神经网络框架通过构建「计算图」来运行,...

17830

扫码关注云+社区

领取腾讯云代金券