CUDA编程模型

CUDA编程模型

CUDA将CPU作为主机（Host），GPU作为设备（Device）。一个系统中可以有一个主机和多个设备。CPU负责逻辑性强的事务处理和串行计算，GPU专注于执行高度线程化的并行处理任务。它们拥有相互独立的存储器（主机端的内存和显卡端的显存）。

运行在GPU上的函数称为kernel（内核函数）。一个完整的CUDA程序是由一些列的kernel函数和主机端的串行处理步骤共同完成的。CPU串行代码的工作包括在kernel启动前进行的数据准备、设备初始化以及在kernel之间进行一些串行化计算。

kernel函数以函数类型限定符_global_定义，并且只能在主机端代码中调用。调用时需要制定kernel的Grid中的block数目以及每个block中thread的数目。每个线程有自己的blockID和threadID，可用来区分其它的线程，这两个内建变量是只读的，由专用寄存器提供，只能有kernel函数调用。

CUDA将计算任务映射为大量的可并行执行的线程，并由硬件动态调度和执行这些线程。kernel是以线程网格Grid的形式组织，每个Grid又若干个线程块block组成，每个block又由若干线程thread组成。本质上，kernel是以block为单位执行的，grid只是用来表示一些列并行block的集合。Block之间是不能彼此通信的。目前一个kernel只支持一个grid，在多指令多数据（MIMD）构架中会存在多个grid。

为方便编程，CUDA使用了dim3类型的内建变量threadIdx和threadIdx。

对于一维的block，线程的编号是threadIdx.x。

对于二维（Dx，Dy）的block，线程的编号是threadIdx.x+threadIdx.y * Dx。

对于三维（Dx，Dy，Dz）的block，线程的编号是threadIdx.x+threadIdx.y * Dx+hreadIdx.z* Dx * Dy。

GPU的计算核心是流多处理器（SM），每个SM包含8个标量流处理器（SP）以及其它的运算单元。Kernel是以block为单位执行的，同一个block的线程需要共享数据，因此它们共享同一个SM。一个block必须分配到一个SM，但是可以一个SM中同一个时刻有多个活动块（active block）执行等待，即同一个SM可以有多个block上下文。实际运行中，block会被分为更小的线程束（wrap），线程束的大小由硬件的计算能力决定，Tesla的构架中一个wrap由32个线程组成。

CUDA采用了单指令多线程执行模型。这个模型是对单指令多数据的改进。CUDA中执行宽度可以在1——512个线程之间变化，但是在单指令多数据中执行宽度必须是一个wrap（32）。

 

CUDA存储器模型：

 

        每一个线程拥有自己的私有存储器，每 一个线程块拥有一块共享存储器(Shared memory)；最后，grid中所有的线程都可以访问同一块全局存储器(global memory)。除此之外，还有两种可以被所有线程访问的只读存储器：常数存储器（constant memory）和纹理存储器(Texture memory),它们分别为不同的应用进行了优化。全局存储器、常数存储器和纹理存储器中的值在一个内核函数执行完成后将被继续保持，可以被同一程序中其 也内核函数调用。

      下表给出了这8种存储器的位置、缓存情况，访问权限及生存域

 

存储器	位置	拥有缓存	访问权限	变量生存周期
register	GPU片内	N/A	Device可读/写	与thread相同
Local memory	板载显存	无	Device可读/写	与thread相同
Shared memory	GPU片内	N/A	Device可读/写	与block相同
Constant memory	板载显存	有	Device可读，host要读写	可在程序中保持
Texture memory	板载显存	有	Device可读，host要读写	可在程序中保持
Global memory	板载显存	无	Device可读/写, host可读/写	可在程序中保持
Host memory	Host内存	无	host可读/写	可在程序中保持
Pinned memory	Host内存	无	host可读/写	可在程序中保持