CUDA编程模型

本篇博客主要对CUDA的编程模型、内存层次结构、线程管理进行简单的介绍。

1、编程结构

CUDA是一种通用的并行计算平台和编程模型，是在C语言基础上扩展。
一些术语区分：

设备	GPU及其内存（设备内存）
主机	CPU及其内存（主机内存 ）

CUDA编程模型主要是异步的，因此在GPU上进行的运算可以与主机-设备通信重叠。一个典型的CUDA程序包括由并行代码互补的串行代码。如图所示，串行代码（及任务并行代码）在主机CPU上执行，而并行代码在GPU上执行。主机代码按照ANSI C标准进行编写，而设备代码使用CUDA 进行编写。你可以将所有的代码统一放在一个源文件中，也可以使用多个源文件来构建应用程序和库。
一个典型的CUDA程序实现流程遵循以下模式。

1. 把数据从CPU内存拷贝到GPU内存。
2. 调用核函数对存储在GPU内存中的数据进行操作。
3. 将数据从GPU内存传送到CPU内存。
   
   图表 1 一个典型的CUDA程序

2 内存层次结构成

CUDA编程模型从GPU架构中抽象出一个内存层次结构。如下图所示的是一个简化的GPU内存结构，它主要包含两部分：全局内存和共享内存。CUDA编程模型最显著的一个特点就是揭示了内存层次结构。每一个GPU设备都有用于不同用途的存储类型。
在GPU内存层次结构中，最主要的两种内存是全局内存和共享内存。全局类似于CPU的系统内存，而共享内存类似于CPU的缓存。然而GPU的共享内存可以由CUDA C的内核直接控制。

图表2 简化的GPU内存结构

3 线程管理

核函数在主机端启动时，它的执行会移动到设备上，此时设备中会产生大量的线程并且每个线程都执行由核函数指定的语句。CUDA明确了线程层次抽象的概念以便于组织线程，这是一个两层的线程层次结构，由线程块和线程块网格构成，如下图所示。

图表 3 线程层次结构示例
由一个内核启动所产生的所有线程统称为一个网格。同一网格中的所有线程共享相同的全局内存空间。一个网格由多个线程块构成，一个线程块包含一组线程，同一线程块内的线程协作可以通过以下方式来实现。不同块内的线程不能协作。
 同步
 共享内存
线程依靠以下两个坐标变量来区分彼此。
 blockIdx（线程块在线程格内的索引）
 threadIdx（块内的线程索引）
这些变量是核函数中需要预初始化的内置变量。当执行一个核函数时，CUDA运行时为每个线程分配坐标变量blockIdx和threadIdx。基于这些坐标，你可以将部分数据分配给不同的线程。
该坐标变量是基于uint3定义的CUDA内置的向量类型，是一个包含3个无符号整数的结构，可以通过x、y、z三个字段来指定。

blockIdx.x
blockIdx.y
blockIdx.z
threadIdx.x
threadIdx.y
threadIdx.z

CUDA可以组织三维的网格和块。线程层次结构(见图线程层次结构示例)是一个包含二维块的二维网格。网格和块的维度由下列两个内置变量指定。
.blockDim(线程块的维度，用每个线程块中的线程数来表示)
.gridDim (线程格的维度，用每个线程格中的线程数来表示)
它们是dim3类型的变量，是基于uint3定义的整数型向量，用来表示维度。当定义一个dim3类型的变量时，所有未指定的元素都被初始化为1。dim3类型变量中的每个组件可以通过它的x、y、z字段获得。如下所示。

blockDim.x
blockDim.y
blockDim.z