__device__ __global__ __host__ __shared__

CUDA编码

对于函数部分前缀

1.  __device__

使用 _device_ 限定符声明的函数具有以下特征：

n         在设备上执行；

n         仅可通过设备调用。

2. __global__

使用 _global_ 限定符可将函数声明为内核。此类函数：

n         在设备上执行；

n         仅可通过主机调用。

3. __host__

使用 _host_ 限定符声明的函数具有以下特征：

n         在主机上执行；

n         仅可通过主机调用。

仅使用 _host_ 限定符声明函数等同于不使用 _host_、_device_ 或 _global_ 限定符声明函数，这两种情况下，函数都将仅为主机进行编译。

函数前缀的一些限制：

_device_ 和 _global_ 函数不支持递归。

_device_ 和 _global_ 函数的函数体内无法声明静态变量。

_device_ 和 _global_ 函数不得有数量可变的参数。

_device_ 函数的地址无法获取，但支持 _global_ 函数的函数指针。

_global_ 和 _host_ 限定符无法一起使用。

_global_ 函数的返回类型必须为空。

对 _global_ 函数的任何调用都必须按规定指定其执行配置。

_global_ 函数的调用是异步的，也就是说它会在设备执行完成之前返回。

_global_ 函数参数将同时通过共享存储器传递给设备，且限制为 256 字节。

对于变量前缀：

1.__device__

_device_ 限定符声明位于设备上的变量。

在接下来的三节中介绍的其他类型限定符中，最多只能有一种可与 _device_ 限定符一起使用，以更具体地指定变量属于哪个存储器空间。如果未出现其他任何限定符，则变量具有以下特征：

n         位于全局存储器空间中；

n         与应用程序具有相同的生命周期；

可通过网格内的所有线程访问，也可通过运行时库从主机访问。

2.__constant__

_constant_ 限定符可选择与 _device_ 限定符一起使用，所声明的变量具有以下特征：

n         位于固定存储器空间中；

n         与应用程序具有相同的生命周期；

可通过网格内的所有线程访问，也可通过运行时库从主机访问。

3.__shared__

_shared_ 限定符可选择与 _device_ 限定符一起使用，所声明的变量具有以下特征：

n         位于线程块的共享存储器空间中；

n         与块具有相同的生命周期；

n         尽可通过块内的所有线程访问。

只有在 _syncthreads()_（参见第 4.4.2 节）的执行写入之后，才能保证共享变量对其他线程可见。除非变量被声明为瞬时变量，否则只要之前的语句完成，编译器即可随意优化共享存储器的读写操作。

限制：

不允许为在主机上执行的函数内的 struct 和 union 成员、形参和局部变量使用这些限定符。

_shared_ 和 _constant_ 变量具有隐含的静态存储。

_device_、_shared_ 和 _constant_ 变量无法使用 extern 关键字定义为外部变量。

_device_ 和 _constant_ 变量仅允许在文件作用域内使用。

不可为设备或从设备指派 _constant_ 变量，仅可通过主机运行时函数从主机指派（参见第 4.5.2.3 节和第 4.5.3.6 节）。

_shared_ 变量的声明中不可包含初始化。

 

下面是具体的一个应用：

将共享存储器中的变量声明为外部数组时，例如：

extern __shared__ float shared[];

数组的大小将在启动时确定（参见第 4.2.3 节）。所有变量均以这种形式声明，在存储器中的同一地址开始，因此数组中的变量布局必须通过偏移显式管理。例如，如果一名用户希望在动态分配的共享存储器内获得与以下代码对应的内容：

short array0[128];

float array1[64];

int array2[256];

则应通过以下方法声明和初始化数组：

extern __shared__ char array[];

__device__ void func() // __device__ or __global__ function

{

    short* array0 = (short*)array;

     float* array1 = (float*)&array0[128];

     int* array2 = (int*)&array1[64];

}

 

在设备代码中声明、不带任何限定符的自动变量通常位于寄存器中。但在某些情况下，编译器可能选择将其置于本地存储器中。

只要编译器能够确定在设备上执行的代码中的指针指向的是共享存储器空间还是全局存储器空间，此类指针即受支持，否则将仅限于指向在全局存储器空间中分配或声明的存储器。

通过获取 _device_、_shared_ 或 _constant_ 变量的地址而获得的地址仅可在设备代码中使用。通过 cudaGetSymbolAddress()获取的 _device_ 或 _constant_ 变量的地址仅可在主机代码中使用。

 

对global函数进行配置

对 _global_ 函数的任何调用都必须指定该调用的执行配置。

执行配置定义将用于在该设备上执行函数的网格和块的维度，以及相关的流。可通过在函数名称和括号参数列表之间插入 <<<Dg, Db, Ns, s>>>形式的表达式来指定，其中：

Dg 的类型为 dim3，指定网格的维度和大小，Dg.x * Dg.y 等于所启动的块数量，Dg.z 无用；

Db 的类型为 dim3，指定各块的维度和大小，Db.x * Db.y * Db.z 等于各块的线程数量；

Ns 的类型为 size_t，指定各块为此调用动态分配的共享存储器（除静态分配的存储器之外），这些动态分配的存储器可供声明为外部数组的其他任何变量使用，Ns 是一个可选参数，默认值为 0；

S 的类型为 cudaStream_t，指定相关流；S 是一个可选参数，默认值为 0。

举例来说，一个函数的声明如下：

__global__ void Func(float* parameter);

必须通过如下方法来调用此函数：

Func<<< Dg, Db, Ns >>>(parameter);

执行配置的参数将在实际函数参数之前被评估，与函数参数相同，通过共享存储器同时传递给设备。

如果 Dg 或 Db 大于设备允许的最大大小，或 Ns 大于设备上可用的共享存储器最大值，或者小于静态分配、函数参数和执行配置所需的共享存储器数量，则函数将失败。