CUDA_C_NOTES [2]

CH02 CUDA编程模型

2.1 CUDA编程模型概述

CUDA编程模型提供了一个计算机架构抽象作为应用程序和其可用硬件之间的桥梁。

CUDA编程模型还利用GPU架构的计算能力提供了以下几个特有功能:

• 一种通过层次结构在GPU中组织线程的方法(2.3)
• 一种通过层次结构在GPU中访问内存的方法(4.5)

程序员可以通过以下几个不同层面来看待并行计算:

• 领域层：如何解析数据和函数，以便在并行环境中正确高效的解决问题（在并行编程中高效的使用pthreads或者OpemMP技术显式地管理线程）
• 逻辑层：如何组织并发线程
• 硬件层：理解线程如何映射到核心以帮助提高其性能

2.1.1 CUDA编程

在一个异构环境中包含多个CPU和GPU， 每个GPU和CPU的内存都由一条PCI-Express总线分隔开。

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

“统一寻址”（Unified Memory） 的编程模型的改进， 它连接了主机内存和设备内存空间， 可使用单个指针访问CPU和GPU内存， 无须彼此之间手动拷贝数据。

什么是“统一寻址”（Unified Memory)? CUDA 6.0提出了统一寻址， 使用一个指针来访问CPU和GPU的内存。(详见第4章)

内核（kernel） 是CUDA编程模型的一个重要组成部分， 其代码在GPU上运行。

CUDA编程模型主要是异步的， 因此在GPU上进行的运算可以与主机-设备通信重叠。 一个典型的CUDA程序包 括由并行代码互补的串行代码。

串行代码在cpu上执行，并行代码在GPU上执行。

一个典型的CUDA程序实现流程遵循以下模式：

1. 把数据从CPU内存拷贝到GPU内存；
2. 调用核函数对存储在GPU内存中的数据进行操作；
3. 将数据从GPU内存传送回到CPU内存。

2.1.2 内存管理

CUDA运行时负责分配与释放设备内存， 并且在主机内存和设备内存之间传输数据。

表2-1 主机和设备内存函数

cudaMalloc函数负责在GPU的内存里分配内存； cudaMemcpy函数负责主机和设备之间的数据传输；

• 1	cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind)

• 从src指向的源存储区复制一定数量的字节到dst指向的目标存储区
• kind有以下几种:
  • cudaMemcpyHostToHost
  • cudaMemcpyHostToDevice
  • cudaMemcpyDeviceToHost
  • cudaMemcpyDeviceToDevice

CUDA编程模型从GPU架构中抽象出一个内存层次结构：全局内存和共享内存。

内存层次结构

• 全局内存
• 共享内存

为什么CPU和GPU是异步的？ 当数据被转移到GPU的全局内存后， 主机端调用核函数在GPU上进行数组求和。 一旦内核被调用， 控制权立刻被传回主机， 这样的话， 当核函数在GPU上运行时， 主机可以执行其他函数。 因此， 内核与主机是异步的。

不同的存储空间

2.1.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

网格和线程块的维度

一个线程格会被组织成线程块的二维数组形式， 一个线程块会被组织成线程的三维数组形式

在CUDA程序中有两组不同的网格和块变量： 手动定义的dim3数据类型和预定义的uint3数据类型。

手动定义的dim3类型的网格和块变量仅在主机端可见， 而unit3类型的内置预初始化的网格和块变量仅在设备端可见。

从主机端和设备端访问网格/块变量

区分主机端和设备端的网格和块变量的访问是很重要的。

例如， 声明一个主机端的块变量， 你按如下定义它的坐标并对其进行访问：

block.x, block.y, block.z

在设备端， 你已经预定义了内置块变量的大小：

blockDim.x, blockDim.y, and blockDim.z

在启动内核之前就定义了主机端的网格和块变量， 并从主机端通过由x、 y、 z三个字段决定的矢量结构来访问它们。 当内核启动时， 可以使用内核中预初始化的内置变量。

总之， 在启动内核之前就定义了主机端的网格和块变量， 并从主机端通过由x、 y、 z三个字段决定的矢量结构来访问它们。 当内核启动时， 可以使用内核中预初始化的内置变量.

对于一个给定的数据大小， 确定网格和块尺寸的一般步骤为：

• 确定块的大小
• 在已知数据大小和块大小的基础上计算网格维度

要确定块尺寸， 通常需要考虑：

• 内核的性能特性
• GPU资源的限制

线程层次结构

CUDA的特点之一就是通过编程模型揭示了一个两层的线程层次结构（grid->block->thread）。 由于一个内核 启动的网格和块的维数会影响性能， 这一结构为程序员优化程序提供了一个额外的途径。

2.1.4 启动一个CUDA核函数

CUDA内核调用是对C语言函数调用语句的延伸， «<»>运算符内是核函数的执行配置。

kernel_name <<<grid, block>>>(argument list)

利用执行配置可以指定线程在GPU上调度运行的方式。 执行配置的第一个值是网格维度， 也就是启动块的数目。 第二个值是块维度， 也就是每个块中线程的数目。 通过指定网格和块的维度， 你可以进行以下 配置：

• 内核中线程的数目
• 内核中使用的线程布局

同一个块(block)中的线程之间可以相互协作， 不同块内的线程不能协作。

假设你有32个数据元素用于计算， 每8个元素一个块， 需要启动4个块：

kernel_name<<<4, 8>>>(argument list)

由于数据在全局内存中是线性存储的， 因此可以用变量blockIdx.x和threadId.x来进行以下操作。

• 在网格中标识一个唯一的线程
• 建立线程和数据元素之间的映射关系

如果把所有32个元素放到一个块里， 那么只会得到一个块:

kernel_name<<<1, 32>>>(argument list)

如果每个块只含有一个元素， 那么会有32个块：

kernel_name<<<32, 1>>>(argument list)

核函数的调用与主机线程是异步的。 核函数调用结束后， 控制权立刻返回给主机端.

你可以调用以下函数来强制主机端程序等待所有的核函数执行结束：

cudaError_t cudaDeivceSynchronize(void);

一些CUDA运行时API在主机和设备之间是隐式同步的。 当使用cudaMemcpy函数在主 机和设备之间拷贝数据时， 主机端隐式同步， 即主机端程序必须等待数据拷贝完成后才能 继续执行程序。

异步行为

不同于C语言的函数调用， 所有的CUDA核函数的启动都是异步的。 CUDA内核调用完成后， 控制权立刻返回给CPU。

2.1.5 编写核函数

核函数是在设备端执行的代码。

用__global__声明定义核函数:

__global__ void kernel_name(argument list);

核函数必须有一个void返回类型。

表2-2总结了CUDA C程序中的函数类型限定符

CUDA核函数的限制 以下限制适用于所有核函数:

• 只能访问设备内存
• 必须具有void返回类型
• 不支持可变数量的参数
• 不支持静态变量
• 显示异步行为

2.3 组织并行线程 (以阅读为主)

从前面的例子可以看出， 如果使用了合适的网格和块大小来正确地组织线程， 那么可以对内核性能产生很大的影响。

2.3.1 使用块和线程建立矩阵索引

在一个矩阵加法核函数中，一个线程通常被分配一个数据元素来处理。首先要完成的任务是使用块和线程索引从全局内存中访问指定的数据。