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 条评论
登录 后参与评论

相关文章

来自专栏玩转JavaEE

MongoDB管道操作符(二)

上篇文章中我们已经学习了MongoDB中几个基本的管道操作符,本文我们再来看看其他的管道操作符。 ---- $group 基本操作 $group可以用来对文档进...

2886
来自专栏软件开发 -- 分享 互助 成长

C++ STL之priority_queue

    STL中的priority_queue(优先队列)是一种会按照自定义的一种方式(数据的优先级)来对队列中的数据进行动态的排序的容器,不同优先级的情况下,...

2118
来自专栏xingoo, 一个梦想做发明家的程序员

Elasticsearch——multi termvectors的用法

前一篇已经翻译过termvectors的使用方法了,这对于学习如何使用tf-idf来说是很有帮助的了。 更多内容参考我整理的ELK教程 什么是TF-ID...

20610
来自专栏灯塔大数据

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

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

1662
来自专栏xingoo, 一个梦想做发明家的程序员

Elasticsearch聚合初探——metric篇

Elasticsearch是一款提供检索以及相关度排序的开源框架,同时,也支持对存储的文档进行复杂的统计——聚合。 前言 ES中的聚合被分为两大类:Met...

22410
来自专栏码匠的流水账

分布式id生成方案概述

对于每个标识,都需要有一个命名空间(namespace),来保证其相对唯一性。 分布式的ID生成,以Twitter Snowflake为代表的, Flake 系...

1132
来自专栏水击三千

UML学习-状态图

1.状态图概述 状态图(Statechart Diagram)主要用于描述一个对象在其生存期间的动态行为,表现为一个对象所经历的状态序列,引起状态转移的事件(E...

26710
来自专栏牛客网

考点总结:互联网校招技术岗都考些什么?数据结构算法游戏 + 场景c++面向对象javaJVMSpringandroid数据库计网线程安全linux前端询问面试官

数据结构 红黑树 pk 平衡二叉树 hash表处理冲突的方法 算法 手写 最长无重复字符子串 链表的增、删、查、逆序 数组实现队列,要求可以动态扩展,保证较高的...

3347
来自专栏JackieZheng

把玩爬虫框架Gecco

如果你现在接到一个任务,获取某某行业下的分类。 作为一个非该领域专家,没有深厚的运营经验功底,要提供一套摆的上台面且让人信服的行业分类,恐怕不那么简单。 找不到...

5454
来自专栏数说工作室

【SAS Says】基础篇:SAS宏初步

特别说明:本节【SAS Says】基础篇:SAS宏初步,用的是数说君学习《The little SAS book》时的中文笔记,我们认为这是打基础的最好选择。 ...

2994

扫码关注云+社区