在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指针及大小,似乎比较灵活一些。