DAY36:阅读”执行空间&扩展修饰符

B. C Language Extensions

B.1. Function Execution Space Specifiers

Function execution space specifiers denote whether a function executes on the host or on the device and whether it is callable from the host or from the device.

B.1.1. __device__

The __device__ execution space specifier declares a function that is:

· Executed on the device,

· Callable from the device only.

The __global__ and __device__ execution space specifiers cannot be used together.

B.1.2. __global__

The __global__ exection space specifier declares a function as being a kernel. Such a function is:

· Executed on the device,

· Callable from the host,

· Callable from the device for devices of compute capability 3.2 or higher (see CUDA Dynamic Parallelism for more details).

A __global__ function must have void return type, and cannot be a member of a class.

Any call to a __global__ function must specify its execution configuration as described in Execution Configuration.

A call to a __global__ function is asynchronous, meaning it returns before the device has completed its execution.

B.1.3. __host__

The __host__ execution space specifier declares a function that is:

· Executed on the host,

· Callable from the host only.

It is equivalent to declare a function with only the __host__ execution space specifier or to declare it without any of the __host__, __device__, or __global__ execution space specifier; in either case the function is compiled for the host only.

The __global__ and __host__ execution space specifiers cannot be used together.

The __device__ and __host__ execution space specifiers can be used together however, in which case the function is compiled for both the host and the device. The __CUDA_ARCH__ macro introduced in Application Compatibility can be used to differentiate code paths between host and device:

B.1.4. __noinline__ and __forceinline__

The compiler inlines any __device__ function when deemed appropriate.

The __noinline__ function qualifier can be used as a hint for the compiler not to inline the function if possible.

The __forceinline__ function qualifier can be used to force the compiler to inline the function.

The __noinline__ and __forceinline__ function qualifiers cannot be used together, and neither function qualifier can be applied to an inline function.

本文备注/经验分享:

今天这节很简单, 主要是说了"CUDA C的执行空间"扩展修饰符. 也就是本章节说的__global__和__device__, 以及,不常用的__host__ 你应当知道, CUDA C是对C的扩展, 这使得熟悉普通CPU上C开发的用户(例如, 来自VC的用户), 可以通过一种简单的扩展, 能够用C来写GPU上的函数. 因为GPU上需要执行的函数, 或者说被加速的函数, 经常是一个程序的繁重计算或者繁重处理的核心部分, 例如你有一个原有的程序, 通过某种方式分析(例如, 通过你的大脑假设), 某段代码(或者某段函数), 占据了90%+的CPU上的执行时间,此时就应当将此函数考虑是否能单独改写成CUDA C版本(从你的老C版本---Fortran用户请考虑CUDA Fortran)。 而CUDA C默认的Runtime API风格的编译, 允许这个函数依然保留在普通的源代码文件中, 只需要加上一点点处理, 就可以在GPU上运行了。 这一点点处理就是: (1)原本的老CPU函数, 需要加上__global__修饰符前缀(也就是本章节说的specifier, 应当叫qualifier? 好像现在这两个词已经不分了); (2)原本直接对老CPU函数的调用, 需要改成<<<>>>语法; (3)老版本以前能直接返回结果的, 现在必须是void了, 不能返回了. 但可以通过写入global memory, 然后通过cudaMemcpy回来。 而今天本章节说的, 则只是第一点中的__global__这种扩展修饰符的用法, 而暂时不涉及其他。 刚才说了, 主要的execution space修饰符有两个, __global__和__device__ 它们实际上不仅仅指定了有这两个前缀的函数将在GPU上执行,也同时指定了CUDA C编译器遇到这两个前缀后, 会将有这些特殊前缀的函数, 生成GPU代码, 而其他源文件中的剩余部分, 没有这两个前缀的函数, CUDA C编译器自动跳过, 调用你本机上的原来的CPU编译器, 继续编译剩下的部分。 所以这是为何你总是看到, CUDA C需要有一个配套的CPU编译器(Host Compiler手册中叫)的原因。 这是基于CUDA Runtime API应用程序开发的特有现象(Driver API是另外一回事. 但现在不说)。 通过这种方式, 用户看来, 它只需要将源文件改成.cu扩展名, 将一些函数添加上__global__之类的扩展, 就自动能在GPU上执行了.不需要考虑如何将代码传递给GPU, 也不需要考虑如果通知GPU开始执行.很方便的。

本章节说, __global__前缀的函数, 编译完成后, 可以从Host端调用它(通过某种<<<>>>语法), 它从Device上运行.这就是我们常说的"启动一个kernel"的过程.请注意是CPU(host)代码中, 你要求启动Kernel, 而Kernel是在GPU(device)上运行的.也就是通过__global__前缀 + <<<>>>语法, 此时你将在异种设备上开始了代码执行, 或者说计算.俗称"异构计算". 现在你已经知道如何开始了。 但是这样还没完, 如果你的CPU代码有子函数一样, 将所有的需要的代码放置在一个__global__开头的函数中, 可能会过于庞大复杂了,此时引入了__device__前缀, 你可以用它来写一些只能在GPU上运行的子函数,然后常规的能被调用的kernel(__global__开头的函数)可以调用这些小片段(以__device__开头)嵌入到自己内部, 这样被反复使用的一些代码可以被抽取出来, 做成小片段. 很方便的.因为被设计成__global__是被能跨越CPU/GPU边界调用的函数, 而__device__是被设计成只能在GPU上调用运行的函数,__device__虽然失去了从Host上调用它的能力,但却多了可以直接返回函数值的功能: (1)__device__前缀的函数只能从GPU上运行, 但可以更像正常的C函数一样的返回结果; (2)__global__前缀的能从Host上调用, 然后从GPU上运行; 但不能直接返回任何结果(可以通过其他变通手段); 这样一对比, 你就知道改写成有这两个前缀的函数, 在执行上的区别了. 但是还没完. 这两句是经典CUDA C的说法. 从计算能力3.5(二代Kepler)开始, NV增强了GPU的能力,允许从GPU上启动kernel(__global__前缀修饰的函数),从GPU上启动kernel和从CPU上启动kernel相比, 具有很多好处, 同时能进一步的降低通讯延迟(因为这个是GPU给GPU自己发布任务, 而不是遥远的跨PCI-E那端的CPU老大),具体的好处和用法, 以及细节, 后面有一个专门的章节.这个特性叫"动态并行"(dynamic parallism),然后还有另外一种, 叫__host__修饰, 这种不常用.单一的__host__修饰等于没有修饰(常规的CPU函数)。

但因为你之前知道了, __device__往往是抽取出来的一些代码片段(被频繁使用的小片段), 可以给__global__的kernel嵌入使用。但有些代码, 往往可以同时只写一次, 想同时给CPU上的普通函数, 和GPU上的__global__的kernel用,此时可以指定__host__前缀和__device__前缀同时存在。这种写法等效于:编译器自动当成2个函数看, 一个只有__host__, 另外一个只有__device__,名字一样而已,所以这种能同时从CPU和GPU上调用.(因为编译了两次, 生成了两种代码)。 本章节还说了, __CUDA_ARCH__宏, 这个很重要,可以让你的代码中, 有选择的根据计算能力不同, 来编译出来不同的代码.一个常见的实用例子是(现在, 过几年这个例子就木有了),在有warp shuffle的卡上(例如你手头的Maxwell),和在没有warp shuffle的卡上,编译出来两种等价效果的, 但实现方式不同的代码. 例如前者的warp内部的数据交换可以通过__shlf*()内置函数,而后者可以慢慢的通过shared memory甚至global memory进行交换。 例如很多挖矿软件就是这样, 来自适应多种计算能力的老卡和新卡的,同时还尽量在这些计算能力的新卡和老卡上, 取得它们的计算能力下,可能的最佳实现途径。

此外, 本章节的最后还提出了2个inline控制的修饰符.__forceinline__和__noinline__,这是用来控制这些小代码片段(__device__)如何被大kernel函数(__global__)调用的,inline就是将小函数的代码嵌入到大函数中, 当成整体编译。而非inline则是常规的当成2个独立函数编译, 然后一个函数调用另外一个函数(后者涉及, 参数传递, 流程转移, 结果返回等等)。 (而如果直接inline了, 就没有这些过程---但可能会生成较大的代码) 请注意使用了inline"可能"会生成较好的代码(因为取消了函数间的流程转移, 参数传递, 结果返回之类的开销),但也可能会造成性能下降(因为都将小函数内嵌到大函数中了, 可能会有多个副本存在, 增大了代价, 而增大的代码会增加SM里的L1 I-Cache一级指令缓存的工作压力. 目前Maxwell/Pascal(6.1)(6.0我不知道)上该缓存只有8KB的. 爆缓存需要从global memory载入代码, 增加访存压力), 反而可能会降低性能。 因此是否inline和具体的实际代码有关.不是inline内嵌后就一定会加快速度.也不一定不inline就一定是缓慢的.(Maxwell和6.1的Pascal上的L1 I-Cache大小由Scott Gray进行微架构测试所得. 并非NV官方数据. NV保密这个); 不过手册依然给出了这2个控制inline与否的修饰符, 可以控制__device__函数的被嵌入与否,请注意的是, 这里的__noline__只是一个hint提示, 最终决定是否hornor你的noinline请求, 依然由编译器的做最终决定。 如果它任何完全没有道理你将某个函数要求不要inline, 它会无视你的请求的.如果不指定这2个inline控制修饰符的任何一个.则完全自动控制, 最终生成代码如何看编译器心情。也就是说, 除非你经过测试, 使用__forceinline__之类的, 的确增加了性能(例如你跑了profiler, 发现的确时间变快了10%),否则不要使用它。

请注意现在的文档很多已经qualifier/modifier/specifier/keyword这4个词混用了。就是常见的C语言里面出现的各种特殊单词(保留的, 你不能用)的叫法,现在你都当成修饰符/关键字之类的来看即可。我们又不是在教学, 不用太较真,用鲁迅的一句话说, 知道它们的意思即可(起到特定的修饰作用, 如同本章节说的),而不需要"知道茴香豆的茴有几个写法".

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

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

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

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

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

发表于

我来说两句

0 条评论
登录 后参与评论

相关文章

来自专栏海天一树

Neo4j学习(1):Neo4j是什么

  最简单的图是单节点的,一个记录,记录了一些属性。一个节点可以从单属性开始,成长为成千上亿,虽然会有一点点麻烦。从某种意义上讲,将数据用关系连接起来分布到不同...

781
来自专栏北京马哥教育

做到这二十条,Python程序性能轻松翻倍!

算法的时间复杂度对程序的执行效率影响最大,在Python中可以通过选择合适的数据结构来优化时间复杂度,如list和set查找某一个元素的时间复杂度分别是O(n)...

1005
来自专栏大内老A

ASP.NET Core中如影随形的”依赖注入”[上]: 从两个不同的ServiceProvider说起

我们一致在说 ASP.NET Core广泛地使用到了依赖注入,通过前面两个系列的介绍,相信读者朋友已经体会到了这一点。由于前面两章已经涵盖了依赖注入在管道构建过...

3157
来自专栏.NET技术

简单工厂模式和策略模式的区别与结合

    简单工厂模式和策略模式是大部分程序员,在学习设计模式时接触得最早,或在工作实践中也是用得相对比较多的两个设计模式。

1735
来自专栏灯塔大数据

干货 | Python 性能优化的20条招数

使用python时,你是不是需要性能优化?今天灯塔给你带来python性能优化的20条招数,记得收藏哟!

832
来自专栏Java学习网

Java中UUID的2种创建方法——有代码实例

UUID(Universally Unique Identifier)全局唯一标识符,是指在一台机器上生成的字符串,它保证对在同一时空中的所有机器都是唯一的。按...

26510
来自专栏贾老师の博客

UML类图关系整理

741
来自专栏醒者呆

基础大扫荡——背包,栈,队列,链表一口气全弄懂

提到数据结构,不得不说数据类型,有人将他们比作分子和原子的关系,我们都知道大自然最小的构成单位是原子,数据类型描述的是原子的内部,如质子、中子的情况,而数据结构...

34515
来自专栏大内老A

深入理解string和如何高效地使用string

无论你所使用的是哪种编程语言,我们都不得不承认这样一个共识:string是我们使用最为频繁的一种对象。但是string的常用性并不意味着它的简单性,而且我认为,...

23210
来自专栏木子昭的博客

<技巧>python模块性能测试以python列表的内置函数append和insert为例以python列表insert方法和append方法快速创建1至1000的列表为例:

算法是程序的灵魂,优秀的算法能给程序的效率带来极大的提升,而算法的优劣,往往要经过大量的测试. 在硬件环境基本不变的前提下,对算法实验的次数越多,测试算法运...

2936

扫码关注云+社区