最近在搞TensorRT 自定义层 把一些后处理过程放在CUDA中计算加速,CUDA代码***.cu中
遇到了__global__ 函数前缀
百度了一下 整理
CUDA C语言对C语言的扩展之一就是加入了一些函数前缀,在CUDA中有三种函数前缀:device,global,host。如下图即为Faster R-CNN代码global前缀添加后的函数定义部分:
三种前缀分别用于在定义函数时限定该函数的调用和执行方式,例如:
- __host__ int foo(int a){}与C或者C++中的foo(int a){}相同,是由CPU调用,由CPU执行的函数
- __global__ int foo(int a){}表示一个内核函数,是一组由GPU执行的并行计算任务,以foo<<>>(a)的形式或者driver API的形式调用。目前__global__函数必须由CPU调用,并将并行计算任务发射到GPU的任务调用单元。随着GPU可编程能力的进一步提高,未来可能可以由GPU调用。由于
- __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