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