CUDA编程之函数前缀：device/global/host

最近在搞TensorRT 自定义层 把一些后处理过程放在CUDA中计算加速，CUDA代码***.cu中

遇到了__global__ 函数前缀

百度了一下 整理

CUDA C语言对C语言的扩展之一就是加入了一些函数前缀，在CUDA中有三种函数前缀：device，global，host。如下图即为Faster R-CNN代码global前缀添加后的函数定义部分：

三种前缀分别用于在定义函数时限定该函数的调用和执行方式，例如：

1. __host__ int foo(int a){}与C或者C++中的foo(int a){}相同，是由CPU调用，由CPU执行的函数 
2. __global__ int foo(int a){}表示一个内核函数，是一组由GPU执行的并行计算任务，以foo<<>>(a)的形式或者driver API的形式调用。目前__global__函数必须由CPU调用，并将并行计算任务发射到GPU的任务调用单元。随着GPU可编程能力的进一步提高，未来可能可以由GPU调用。由于
3. __device__ int foo(int a){}则表示一个由GPU中一个线程调用的函数。由于Tesla架构的GPU允许线程调用函数，因此实际上是将__device__ 函数以__inline形式展开后直接编译到二进制代码中实现的，并不是真正的函数。

具体来说，device前缀定义的函数只能在GPU上执行，所以device修饰的函数里面不能调用一般常见的函数；global前缀，CUDA允许能够在CPU，GPU两个设备上运行，但是也不能运行CPU里常见的函数；host前缀修饰的事普通函数，默认缺省，可以调用普通函数。

__global__函数执行配置

任何对__global__函数的调用都必须指定该调用的执行配置。执行配置定义将用于在该设备上执行函数的网格和块的维度，以及相关的流。

运行时API通过在函数名称和参数列表之间插入<<<Dg, Db, Ns, S>>>的形式来指定。其中：

• Dg 的类型为dim3，指定网格的维度和大小，Dg.x * Dg.y 等于所发射的块数量；
• Db 的类型为dim3，指定各块的维度和大小，Db.x * Db.y *Db.z 等于各块的线程数量；
• Ns 的类型为size_t，指定各块为此调用动态分配的共享存储器（除静态分配的存储器之外），这些动态分配的存储器可供声明为动态数组的其他任何变量使用，Ns 是一个可选参数，默认值为0；
• S 的类型为cudaStream t，指定相关流；S 是一个可选参数，默认值为0。

例如一个函数的声明如下:

__global__ void CalDetection(const float *input, float *output, int num_elem, int step, int anchor, int output_elem)

则必须通过类似如下的方式调用：

CalDetection<<< (num_elem + thread_count - 1) / thread_count, thread_count>>>

参考资料：

https://blog.csdn.net/u012905422/article/details/52904853

https://stackoverflow.com/questions/12373940/difference-between-global-and-device-functions

https://blog.csdn.net/heiheiya/article/details/82019309