CUDA是一种通用的并行计算平台和编程模型，是在C语言上扩展的。借助于CUDA，你可以像编写C语言程序一样实现并行算法。你可以在NIVDIA的GPU平台上用CUDA为多种系统编写应用程序，范围从嵌入式设备、平板电脑、笔记本电脑、台式机工作站到HPC集群。在CUDA编程平台中，GPU并不是一个独立运行的计算平台，而需要与CPU协同工作，可以看成是CPU的协处理器，因此当我们在说GPU并行计算时，其实是指的基于CPU+GPU的异构计算架构。在异构计算架构中，GPU与CPU通过PCIe总线连接在一起来协同工作，CPU所在位置称为为主机端（host），而GPU所在位置称为设备端（device）。

第一章 基于CUDA的异构并行计算

并行计算

• 一个大的问题可以被分解成多个小问题，然后在不同的计算资源上并行处 理这些小问题。
• 并行计算的软件和硬件层面是紧密联系的。
• 事实上，并行计算通常涉及两个不同的计算技术领域。
• 串行编程和并行编程
• 大多数现代处理器的哈佛体系结构（Harvard architecture）
• 计算机架构（硬件方面）
• 并行程序设计（软件方面）

并行性

• 多层次的并行性设计是架构设计的驱动力。在应用程序中有两种基本的并行类型。

  CUDA编程非常适合解决数据并行计算的问题。本书的重点便是如何使用CUDA编程 解决数据并行问题。

数据并行程序设计

• • 块划分（block partitioning）
  • 周期划分（cyclic partitioning）
  • 一组连续的数据被分到一个块内
  • 每个数据块以任 意次序被安排给一个线程
  • 线程通常在同一时间只处理一个数据块
  • 更少的数据被分到一个块内
  • 相邻的线程处理相邻的数据块
  • 每个线程可以处理多个数据块
  • 为一个待处理的线程选择一个新的块，就意味着要跳过和现有线程一样多的数据块
  • 第一步是把数据依据线程进行划分，使每个线程处理一部分数据
  • 利用多核系统对数据进行分配
  • 利用多核系统对任务进行分配
  • 任务并行
  • 数据并行

计算机架构

• 弗林分类法（Flynn’s Taxonomy）

• • 降低延迟
  • 提高带宽
  • 提高吞吐量
  • 延迟是一个操作从开始到完成所需要的时间，常用微秒来表示
  • 延迟用来衡量完成一次操作的时间
  • 带宽是单位时间内可 处理的数据量，通常表示为MB/s或GB/s。
  • 吞吐量是单位时间内成功处理的运算数量，通常表示为gflops（即每秒十亿次的浮点运算数量），特别是在重点使用浮点计算的科学计算领域经常用到
  • 而吞吐量用来衡量在给定的单位时 间内处理的操作量
  • 单指令单数据（SISD）(传统计算机，一种串行架构)
  • 单指令多数据（SIMD） (按串行逻辑思考但对并行数据操作实现并行加速，而其他细节则由编译器来负责)
  • 多指令单数据（MISD） (每个核心通过使用多个指令流处理同一个数据流)
  • 多指令多数据（MIMD）(多个核心使用多个指令流来异步处理多个数据流，从而实现空间上的并行性)
  • 根据指令和数据进入CPU的方式
  • 实现以下目的，架构取得了许多进展

• 计算机架构也能根据内存组织方式进行进一步划分

• • 要么是与同一个物理内存相关联

  • 要么共用一个低延迟的链路（如PCIExpress或PCIe）

  • 多核架构已经永久地取代了单核架构

  • 多线程、 MIMD（多指令多数据）、SIMD（单指令多数据），以及指令级并行。

    NVIDIA公司称这 种架构为SIMT（单指令多线程）

  • “众核”（many-core）通常是指有很多核心（几十或几百个）的多核架构

  • GPU代表了一种众核架构，几乎包括了前文描述的所有并行结构：

  • 这种系统常被称作集群

  • 分布式内存的多节点系统

  • 共享内存的多处理器系统

• GPU核心和CPU核心

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

异构计算

GPU指的是离散的设备从同构系统到异构系统的转变是高性能计算 史上的一个里程碑。

• 同构计算使用的是同一架构下的一个或多个处理器来执行一个应用。

• 而异构计算则使用一个处理器架构来执行一个应用，为任务选择适合它的架构，使其最终对性能有所改进。

• • 系统的有效利用 受限于增加应用程序设计的复杂性

异构架构

• 一个典型的异构计算节点

• • GPU不 是一个独立运行的平台而是CPU的协处理器
  • 通过PCIe总线与基于CPU的 主机相连来进行操作
  • 两个多核CPU插槽（主机端）
  • 两个或更多个的众核GPU （设备端）

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

• • 在计算密集型应用中，往往有很多并行数据的程序段。

    GPU就是用来提高这些并行数据的执行速度的。

  • 在设备端加载计算密集型任务之前，CPU代码负责管理设备端的环境、代码和数据。

  • 主机代码（CPU上运行）

  • 设备代码（GPU上运行）

• NVIDIA公司的GPU计算平台

• • Fermi是Tesla系列产品中的一种，用作GPU加速器，近来在高性能计算中获得了广泛应用
  • Fermi之后的新一代GPU 计算架构Kepler，于2012年秋季发布，其处理能力相比以往的GPU有很大提升，并且提供 了新的方法来优化和提高GPU并行工作的执行，有望将高性能计算提升到新的高度
  • Tegra系列产品是专为移动和嵌入式设备而设计的，如平板电脑和手机
  • GeForce面向图形用户
  • Quadro用于专业绘图设计
  • Tesla用于大规模的并行计算

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

• • CUDA核心数量
  • 内存大小

• 两种不同的指标来评估GPU的性能

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

• 计算能力

• • NVIDIA使用一个术语“计算能力”（compute capability）来描述整个Tesla系列的GPU加 速器的硬件版本。

异构计算范例

GPU计算并不是要取代CPU计算

Tip：这一段我就不写了，本章节的主要内容用一张图即可概括

CUDA：一种异构计算平台

API

CUDA提供了两层API来管理GPU设备和组织线程。

这两种API是相互排斥的，你必须使用两者之一，从两者中混合函数调用是不可能的。

• CUDA驱动API

• • 驱动API是一种低级API，它相对来说较难编程
  • 对于在GPU设备使用上提供了更多的控制

• CUDA运行时API

• • 运行时API是一个高级API
  • 它在驱动API的上层实现
  • 每个运行时API函数都被分解为更多传给驱动API的基本运算

也就是说Runtime api 可以看作是由Driver api 封装而成的

CUDA程序

• 一个CUDA程序包含了以下两个部分的混合。

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

• 一个典型的CUDA编程结构包括5个主要步骤

• • 分配GPU内存
  • 从CPU内存中拷贝数据到GPU内存
  • 调用CUDA内核函数来完成程序指定的运算
  • 将数据从GPU拷回CPU内存
  • 释放GPU内存空间

进阶

• 数据局部性在并行编程中是一个非常重要的概念。

• • 时间局部性是指在相对较短的时间段内数据和/或资源的重用。
  • 空间局部性是指在相对较接近的存储空间内数据元素的重用。
  • 数据局部性指的是数据重用，以降低内存访问的延迟。
  • 数据局部性有两种基本类型

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

• • 内存层次结构

  • 线程层次结构

  • 例如：

    在CUDA编程模型中使用的共享内存（一个特殊的内存）。

    共享内存可以视为 一个被软件管理的高速缓存，通过为主内存节省带宽来大幅度提高运行速度。

    有了共享内存，你可以直接控制代码的数据局部性。

编译环境：本代码将使用nvcc编译器来编译，你可以使用以下命令来检查CUDA是否正确安装:

1
2

$ which nvcc
/usr/local/cuda-8.0/bin/nvcc  # cuda-8.0 版本

用GPU输出 Hello World

不妨先写一个cuda C程序，命名为helloFromGPU，用它来输出字符串 “Hello World from GPU！”

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18

========================== 代码清单 1-1 Hello World from GPU (hello.cu) ==========================
// hello.cu
#include <stdio.h>

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

int main(void)
{
    // hello from cpu
    printf("Hello World from CPU!\n");

    helloFromGPU <<<1, 10>>>();
    cudaDeviceReset();
    return 0;
}

在linux终端下使用以下命令进行编译hello.cu，然后执行程序得到

1
2
3
4
5
6
7
8
9
10
11
12
13

$ nvcc -arch sm_20 hello.cu -o hello
$ ./hello
Hello World from CPU!
Hello World from GPU!
Hello World from GPU!
Hello World from GPU!
Hello World from GPU!
Hello World from GPU!
Hello World from GPU!
Hello World from GPU!
Hello World from GPU!
Hello World from GPU!
Hello World from GPU!

在上面的代码中，cudaDeviceReset表示重置当前线程所关联过的当前设备的所有资源；修饰符__global__告诉编译器这是一个内核函数，它将从CPU中调用，然后在GPU上执行，在CPU上通过下面的代码启动内核函数

1

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

三重尖号意味着从主线程到端代码的调用。1和10分别表示有1个块区域和10个线程，后续会作相关介绍。