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

在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<<>>(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__前缀，视需要编译出两个版本。

 

存储器类型前缀：

__local = __shared__

__global = __device__

__constant = __constant__

__private = 不带前缀的私有变量,register/local memory

 

值得注意的是__local（work group本地共享存储器）和CUDA中曾经短暂存在的__local__（显存中的线程私有存储器）意义不同。__global和__global__ ,__device__也很微妙。

 

此外,__local和__shared__的意义和使用方法基本相同，但是如果需要在调用内核函数时动态分配每个SM使用的共享存储器的大小，仍然有少许不同。假设我们需要在kernel中动态的定义几个__shared__数组的大小，那么

CUDA C语言中的方式是开辟一块不在kernel内指定大小extern __shared__ type a[]数组以及与需要的数组相应的几个__shared__指针，然后根据外部分配的大小使__shared__指针指向的a中的对应位置。

OpenCL的方式则是直接在调用函数时以参数的形式传入__local memory指针及大小，似乎比较灵活一些。

ref: http://blog.itpub.net/23172605/viewspace-627164/