CUDA_C_NOTES [1]

Ch01 基于CUDA的异构并行计算

1.1 并行计算

并行计算通常设计两个不同的计算机领域

• 计算机架构(硬件)：在结构级别上支持并行性
• 并行程序设计(软件)：充分使用计算机架构的计算能力来并发地解决问题

1.1.1 串行编程和并行编程

1.1.2 并行性

并行性方式

• 任务并行： 当许多任务或函数可以独立地、大规模地并行执行时，这就是任务并行。任务并行的核心是在于利用多核系统对任务进行分配。

• 数据并行： 当可以处理许多数据的时候，就是数据并行。数据并行的重点是利用多核系统对数据进行分配。

CUDA编程非常适合解决数据并行问题。

数据划分方式：

• 块划分：
  • 每个线程作用于一部分数据， 通常这些数据具有相同大小。
  • 一组连续数据被分到一个块内，每个数据块以任意次序被安排给一个线程，线程通常在同一时间只处理一个数据块。
• 周期划分：
  • 每个线程作用于数据的多部分。
  • 在周期划分中，更少的数据被分到一个块内。相邻的线程处理相邻的数据块，每个线程可以处理多个数据块。为一个待处理的线程选择一个新的块，就意味着要跳过和现有线程一样多的数据块。

1.1.3 计算机架构

计算机机构分类(弗林分类(Flynn’s Taxonomy))：根据指令和数据进入CPU的方式进行分类

• 单指令单数据（SISD）
  • 一种串行架构。 在这种计算机上只有一个核心。在任何时间点上只有一个指令流在处理一个数据流。
• 单指令多数据（SIMD）
  • 一种并行架构类型。在这种计算机上有多个核心。 在任何时间点上所有的核心只有一个指令流处理不同的数据流，例如向量机。
  • 优势: 在CPU上编写代码时， 程序员可以继续按串行逻辑思考但对并行数据操作实现并行加速，而其他细节则由编译器来负责。
• 多指令单数据（MISD）
  • 比较少见, 每个核心通过使用多个指令流处理同一个数据流
• 多指令多数据（MIMD）
  • 一种并行架构， 在这种架构中，多个核心使用多个指令流来异步处理多个数据流，从而实现空间上的并行性。 许多MIMD架构还包括SIMD执行的子组件。

计算机架构优劣的评价指标：

• 降低延迟
  • 延迟是一个操作从开始到完成所需要的时间， 常用微秒来表示
• 提高带宽
  • 带宽是单位时间内可处理的数据量， 通常表示为MB/s或GB/s。
• 提高吞吐量
  • 吞吐量是单位时间内成功处理的运算数量， 通常表示为gflops（即每秒十亿次的浮点运算数量） ， 特别是在重点使用浮点计算的科学计算领域经常用到
  • 延迟用来衡量完成一次操作的时间， 而吞吐量用来衡量在给定的单位时间内处理的操作量

根据内存组织方式进一步划分计算机架构:

• 分布式内存的多节点系统
  • 大型计算引擎是由许多网络连接的处理器构成的。 每个处理器有自己的本地内存， 而且处理器之间可以通过网络进行通信(类似于多机多卡)
• 共享内存的多处理器系统

GPU代表了一种众核架构，几乎包括了前文描述的所有并行结构： 多线程、MIMD（多指令多数据）、 SIMD（单指令多数据）， 以及指令级并行。 NVIDIA公司称这 种架构为SIMT（单指令多线程）。

GPU核心和CPU核心

尽管可以使用多核和众核来区分CPU和GPU的架构， 但这两种核心是完全不同的。

• CPU核心比较重， 用来处理非常复杂的控制逻辑， 以优化串行程序执行。
• GPU核心较轻， 用于优化具有简单控制逻辑的数据并行任务， 注重并行程序的吞吐量。

1.2 异构计算

CPU和GPU是两个独立的处理器， 它们通过单个计算节点中的PCI-Express总线相连。 在这种典型的架构中， GPU指的是离散的设备，从同构系统到异构系统的转变是高性能计算 史上的一个里程碑。 同构计算使用的是同一架构下的一个或多个处理器来执行一个应用。 而异构计算则使用一个处理器架构来执行一个应用，为任务选择适合它的架构，使其最终 对性能有所改进.

1.2.1 异构架构

一个典型的异构计算节点包括两个多核CPU插槽和两个或更多个的众核GPU。 GPU不 是一个独立运行的平台而是CPU的协处理器。 因此， GPU必须通过PCIe总线与基于CPU的 主机相连来进行操作， 如图1-9所示。 这就是为什么CPU所在的位置被称作主机端(host)而GPU 所在的位置被称作设备端(device)。

一个异构应用包括两部分：

• 主机代码：在CPU上运行
• 设备代码：在GPU上运行.

描述GPU容量的两个重要特征

• CUDA核心数量
• 内存大小

相应的， 有两种不同的指标来评估GPU的性能:

• 峰值计算性能：用来评估计算容量的一个指标， 通常定义为每秒能处理的单精度或双精度浮点运算的数量，通常用GFlops（每秒十亿次浮点运算） 或TFlops（每秒万 亿次浮点运算） 来表示
• 内存带宽：从内存中读取或写入数据的比率。 内存带宽通常用GB/s表示

计算能力

1.2.2 异构计算范例

GPU与CPU结合后， 能有效提高大规模计算问题的处理速度与性能。 CPU针对动态工作负载进行了优化， 这些动态工作负载是由短序列的计算操作和不可预测的控制流程标 记的； 而GPU在其他领域内的目的是： 处理由计算任务主导的且带有简单控制流的工作负载。

CPU线程与GPU线程

CPU上的线程通常是重量级的实体。 操作系统必须交替线程使用启用或关闭CPU执行通道以提供多线程处理功能。 上下文的切换缓慢且开销大。

GPU上的线程是高度轻量级的。 在一个典型的系统中会有成千上万的线程排队等待工作。 如果GPU必须等待一组线程执行结束， 那么它只要调用另一组线程执行其他任务即可

1.2.3 CUDA： 一种异构计算平台

CUDA是一种通用的并行计算平台和编程模型，它利用NVIDIA GPU中的并行计算引擎能更有效地解决复杂的计算问题。通过使用CUDA，你可以像在CPU上那样，通过GPU来进行计算。

CUDA提供了两层API来管理GPU设备和组织线程， 如图1-13所示。

• CUDA驱动API：驱动API是一种低级API， 它相对来说较难编程， 但是它对于在GPU设备使用上提供了更多的控制。
• CUDA运行时API：运行时API是一个高级API， 它在驱动API的上层实现。 每个运行时API函数都被分解为更多传给驱动API的基本运算。

一个CUDA程序包含了以下两个部分:

• 在CPU上运行的主机代码
• 在GPU上运行的设备代码

NVIDIA的CUDA nvcc编译器在编译过程中将设备代码从主机代码中分离出来. 主机代码是标准的C代码，使用C编译器进行编译。 设备代码，也就是核函数， 是用扩展的带有标记数据并行函数关键字的CUDA C语言编写的. 设备代码通过nvcc进行编译。 在链接阶段，在内核程序调用和显示GPU设备操作中添加CUDA运行时库。

1.3 用GPU输出Hello World

1. 用专用扩展名.cu来创建一个源文件

2. 使用CUDA nvcc编译器来编译程序

3. 从命令行运行可执行文件， 这个文件有可在GPU上运行的内核代码。 首先， 我们编写一个C语言程序来输出“Hello World”， 如下所示

1
2
3
4
5

#include<stdio.h>

int main(void) {
  printf("Hello World from CPU!\n")；
}

把代码保存到hello.cu中， 然后使用nvcc编译器来编译。 CUDA nvcc编译器和gcc编译器及其他编译器有相似的语义

1

nvcc hello.cu -o hello

如果你运行可执行文件hello， 将会输出： Hello World from CPU!

接下来， 编写一个内核函数， 命名为helloFromGPU， 用它来输出字符串“Hello World from GPU！ ”。

1
2
3

__global__ void helloFromGPU(void) {
  printf("Hello World from GPU!\n");
}

修饰符__global__告诉编译器这个函数将会从CPU中调用， 然后在GPU上执行。用下面的代码启动内核函数.

1

helloFromGPU <<<1, 10>>>()

三重尖括号意味着从主线程到设备端代码的调用。 一个内核函数通过一组线程来执行， 所有线程执行相同的代码。

三重尖括号里面的参数是执行配置， 用来说明使用多少线程来执行内核函数。 在这个例子中，有10个GPU线程被调用。

cudaDeviceReset()用来显式地释放和清空当前进程中与当前设备有关的所有资源。

一个典型的CUDA编程结构包括5个主要步骤: 1. 分配GPU内存 2. 从CPU内存中拷贝数据到GPU内存 3. 调用CUDA内核函数来完成程序指定的运算 4. 将数据从GPU拷回CPU内存 5. 释放GPU内存空间

1.4 使用CUDA C编程难吗

数据局部性: 指的是数据重用， 以降低内存访问的延迟

• 时间局部性：指在相对较短的时间段内数据或资源的重用
• 空间局部性：指在相对较接近的存储空间内数据元素的重用。

CUDA中有内存层次和线程层次的概念

• 内存层次结构
• 线程层次结构

CUDA核中有3个关键抽象

• 线程组的层次结构
• 内存的层次结构
• 障碍同步

1.5 总结

CPU + GPU的异构系统成为高性能计算的主流架构: 在GPU上执行数据并行工作， 在CPU上执行串行和任务并行的工作。