5
图 1-3.fork/join 模式
第四节、CUDA C 语言
对一个语言来说，最简单的，也是用得最多的当然是它的语法了，下面简单的介绍一下
CUDA 的语法。
CUDA C 不是 C 语言，而是对 C 语言进行扩展。CUDA 对 C 的扩展主要包括以下四个方
面：
函数类型限定符，用来确定函数是在 CPU 还是在 GPU 上执行，以及这个函数是从
CPU 调用还是从 GPU 调用。
__device__ ， 表示从 GPU 上调用，在 GPU 上执行；也就是说其可以被
__global__ 或者__device__修饰的函数调用。此限定符修饰的函数使用有限
制，比如在 G80/GT200 架构上不能使用递归，不能使用函数指针等，具体可参
见我翻译的编程指南。
__global__，表示在 CPU 上调用，在 GPU 上执行，也就是所谓的内核(kernel)
函数；内核只能够被主机调用，内核并不是一个完整的程序，它只是一个数据
并行步骤，其指令流由多个线程执行。
__host__，表示在 CPU 上调用，在 CPU 上执行，这是默认时的情况，也就是传
统的 C 函数。CUDA 支持__host__和__device__的联用，表示同时为主机和设备
编译。
变量类型限定符，用来规定变量存储什么位置上。在传统的 CPU 程序上，这个任务
由编译器承担。在 CUDA 中，不仅要使用主机端的内存，还要使用设备端的显存和
GPU 片上的寄存器、共享存储器和缓存。在 CUDA 存储器模型中，一共抽象出来了 8
种不同的存储器。复杂的存储器模型使得必须要使用限定符要说明变量的存储位
置。
__device__，__device__表明声明的数据存放在显存中，所有的线程都可以访
问，而且主机也可以通过运行时库访问；
__shared__，__shared__表示数据存放在共享存储器在，只有在所在的块内的
线程可以访问，其它块内的线程不能访问；
__constant__，__constant__表明数据存放在常量存储器中，可以被所有的线
程访问，也可以被主机通过运行时库访问；
Texture，texture 表明其绑定的数据可以被纹理缓存加速存取，其实数据本身
的存放位置并没有改变，纹理是来源于图形学的一介概念，CUDA 使用它的原因
一部分在于支持图形处理，另一方面也可以利用它的一些特殊功能。
如果在 GPU 上执行的函数内部的变量没有限定符，那表示它存放在寄存器或者
本地存储器中，在寄存器中的数据只归线程所有，其它线程不可见。
如果 SM 的寄存器用完，那么编译器就会将本应放到寄存器中的变量放到本地存
储器中。
执行配置运算符<<< >>>，用来传递内核函数的执行参数。执行配置有四个参数，第
一个参数声明网格的大小，第二个参数声明块的大小，第三个参数声明动态分配的
共享存储器大小，默认为 0，最后一个参数声明执行的流，默认为 0。
五个内建变量，用于在运行时获得网格和块的尺寸及线程索引等信息
gridDim, gridDim 是一个包含三个元素 x,y,z 的结构体，分别表示网格在
x,y,z 三个方向上的尺寸，虽然其有三维，但是目前只能使用二维；
blockDim, blockDim 也是一个包含三个元素 x,y,z 的结构体，分别表示块在
