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<<<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__前缀，视需要编译出两个版本。
存储器类型前缀：
__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指针及大小，似乎比较灵活一些。
OpenCL 的方式则是直接在调用函数时以参数的形式传入__local memory指针及大小，似乎比较灵活一些。
from : http://space.itpub.net/23172605/viewspace-627164