CUDA函数前缀与存储器前缀讨论

偏执的太偏执、 2022-08-08 04:29 149阅读 0赞

在CUDA C语言对C语言的扩展之一就是加入了一些函数前缀和存储器前缀，它们是：

函数前缀：

\_\_device\_\_ , \_\_global\_\_, \_\_host\_\_

存储器类型前缀：

\_\_device\_\_, \_\_shared\_\_, \_\_constant\_\_(constant, 不是const)

其中\_\_是两条下划线。

值得注意的是函数前缀和存储器前缀中都有\_\_device\_\_，但表达的意思不同。

在早期版本的CUDA中曾经还有\_\_local\_\_前缀用于限定将某些变量放在local memory中，但是后来的版本决定由编译器决定将变量放在local memory还是register中。

下面详细介绍一下函数前缀和存储器类型前缀的意义和使用方法。

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

\_\_host\_\_ int foo(int a)\{\}与C或者C++中的foo(int a)\{\}相同，是由CPU调用，由CPU执行的函数

\_\_global\_\_ int foo(int a)\{\}表示一个内核函数，是一组由GPU执行的并行计算任务，以foo<<<grid, dim, sharedsize, streamid>>>(a)的形式或者driver API的形式调用。目前\_\_global\_\_函数必须由CPU调用，并将并行计算任务发射到GPU的任务调用单元。随着GPU可 编程能力的进一步提高，未来可能可以由GPU调用。

\_\_device\_\_ int foo(int a)\{\}则表示一个由GPU中一个线程调用的函数。由于Tesla架构的GPU允许线程调用函数，因此实际上是将\_\_device\_\_ 函数以\_\_inline形式展开后直接编译到二进制代码中实现的，并不是真正的函数。而fermi则允许GPU线程调用函数，\_\_device\_\_也就成了名副其实的函数了。

\_\_host\_\_和\_\_device\_\_关键字可以连用，例如

\_\_host\_\_ \_\_device\_\_ int foo(int a)\{\}会被编译成两个版本，分别可以由CPU和GPU线程调用。

数据类型前缀\_\_device\_\_， \_\_shared\_\_和 \_\_constant\_\_使用场合有：

定义显存中的一个数组或者常数存储器中的一个数组时，例如\_\_device\_\_ a\[100\], \_\_constant\_\_ b\[100\] = \{0\}, \_\_shared\_\_ a\[\]。调用以这种方式声明的数组时需要注意数组的作用范围，必须对调用它的核函数可见。

\_\_shared\_\_有一些特别的地方，\_\_shared\_\_数组在定义时不能赋值，并且只能在\_\_global\_\_函数内部定义。

定义一个指针的类型时，例如\_\_device\_\_ a\*, \_\_device\_\_ b\*等。在定义\_\_global\_\_和\_\_device\_\_函数时有时也要在参数前加上数据类型前缀。

随着Fermi引入统一编址，未来不需要在指针类型前加数据类型前缀也可以工作。

CUDA前缀与OpenCL前缀的对应关系：

函数前缀：

OpenCL中只有\_\_kernel函数前缀与CUDA的\_\_global\_\_函数前缀对应

而所有的由CPU和GPU线程调用的函数都等同于CUDA中的\_\_host\_\_ \_\_device\_\_前缀，视需要编译出两个版本。

存储器类型前缀：