CUDA C编程权威指南-第一章：基于CUDA的异构并行计算

系列导航：导读 | 下一篇：第2章 CUDA编程模型

本篇是《CUDA C编程权威指南》读书笔记系列的第 1 篇，从并行计算与异构架构讲起：为什么需要 GPU、什么是异构计算、CUDA 如何定位。先把「并行思维」和「CPU+GPU 分工」搞清楚，后面学编程模型、执行模型和内存优化会顺很多。

一、本章在全书中的位置与学习目标

第 1 章不写复杂代码，主要解决观念与语境问题——先把「并行计算」和「异构计算」讲清楚，是因为：

  • CUDA 编程的本质是数据并行 + 在 GPU 上执行，若不清楚并行性的类型和 CPU/GPU 的差异，后面学线程网格、共享内存时会缺少动机。
  • 主机（Host）与设备（Device）、主机代码与设备代码的区分将贯穿全书，第1章就必须建立清晰术语。

学完本章，你应该能回答：

目标 检验方式
理解串行与并行的区别，以及数据相关性对并行性的约束 能解释「为什么某段代码难以并行」
区分任务并行与数据并行，并知道 CUDA 主攻哪一类 能说出块划分与周期划分的大致思路
理解弗林分类中的 SISD/SIMD/MIMD，以及 GPU 的 SIMT 能简述 SIMT 与 SIMD 的联系与区别
理解同构与异构、主机与设备、GPU 作为协处理器与硬件加速器 能画出一个简单异构节点（CPU + GPU + PCIe）
理解 CPU 与 GPU 的设计哲学差异（延迟 vs 吞吐量、线程轻重） 能说明「为何 GPU 适合大规模数据并行」
会编写、编译并运行一个最简单的核函数（Hello World） 能解释 __global__ 和 <<<1,10>>> 的含义

二、并行计算（书 1.1 节）

2.1 串行编程与并行编程

这一节里把概念厘清一下：用程序解决问题时，我们会自然地把问题划分成许多运算块，每个块执行一个指定任务。若这些块按先后顺序依次执行，就是串行程序。

两个计算块之间要么存在执行次序约束（必须串行），要么没有次序约束（可以并发）。所有包含并发执行任务的程序都是并行程序。尤其要注意的是：一个并行程序中可能既有并行部分也有串行部分（例如先串行初始化，再并行计算，再串行汇总）。

理解与体会：并行化不是「全盘同时执行」，而是找出可以同时执行的子集，其余部分仍保持顺序。这对后面理解「主机上串行、设备上并行」的分工很重要。

数据相关性（这一章会反复用到）：

  • 一个任务若依赖另一个任务的输出，则两者相关；否则独立。
  • 在并行算法实现中，分析数据相关性是最基本的内容，因为相关性是限制并行性的主要因素。

难点：很多性能问题源于「以为能并行、实则存在隐式依赖」。写 CUDA 时，要习惯先问：这段计算依赖哪些数据、是否被其他线程写入？

并行计算涉及两个紧密相关的层面（书中 1.1 节）：

层面 关注点
计算机架构（硬件） 在结构级别上支持并行执行
并行程序设计（软件） 充分利用硬件能力并发地解决问题

这里提到：编写串行程序时，未必需要深入理解计算机架构；但在多核/众核上编写并行程序时，对底层架构的基本认识就显得非常重要。现代处理器多采用哈佛体系结构：内存（指令与数据）、中央处理单元（控制单元与算术逻辑单元）、输入/输出接口。高性能计算的关键是 CPU；早期单芯片单 CPU 为单核，现在多核处理器将多个核心集成在同一芯片上，并行程序设计可以看作把问题的计算分配给可用核心以实现并行的过程。

2.2 并行性的两种基本类型（书 1.1.2）

并行性有两种基本类型：

类型 含义 重点
任务并行 许多任务或函数可以独立地、大规模地并行执行 利用多核系统对任务进行分配
数据并行 可以同时处理许多数据 利用多核系统对数据进行分配

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

数据并行程序设计的第一步是把数据依据线程进行划分，使每个线程处理一部分数据。两种常见划分方式（原书图 1-4、1-5）：

划分方式 做法 特点
块划分（block partitioning） 一组连续的数据被分到一个块内；每个数据块以任意次序分配给一个线程，线程通常同一时间只处理一个数据块 每个线程只处理数据的一部分
周期划分（cyclic partitioning） 更少的数据被分到一个块内；相邻线程处理相邻数据块，每个线程可以处理多个数据块；选新块时相当于「跳过与现有线程数相同数量的数据块」 每个线程处理数据的多个部分

对二维数据，可以沿 x 轴、y 轴或两者做块划分或周期划分（原书图 1-5 给出了沿 y 轴块划分、沿 x 和 y 块划分、沿 x 轴周期划分等例子）。

注意：数据通常在物理上是一维存储的，即便逻辑上是多维，也要映射到一维地址空间。如何在线程中分配数据，不仅与数据的物理存储方式密切相关，而且与每个线程的执行次序也有很大关系；组织线程的方式对程序性能有很大影响。 这一点在第 2 章（网格与块）、第 4 章（全局内存访问模式）会反复出现。

2.3 计算机架构分类：弗林分类法（书 1.1.3）

弗林分类法（Flynn’s Taxonomy）根据指令和数据进入 CPU 的方式，将计算机架构分为四类（原书图 1-6）：

类型 全称 说明
SISD 单指令单数据 传统串行架构：一个核心，任一时刻只有一个指令流处理一个数据流
SIMD 单指令多数据 多个核心，同一时刻所有核心只有一个指令流，处理不同数据流；向量机是典型代表；现代很多机器都采用 SIMD
MISD 多指令单数据 多个指令流处理同一数据流，实际中较少见
MIMD 多指令多数据 多个核心用多个指令流异步处理多个数据流，实现空间上的并行；许多 MIMD 架构还包含 SIMD 子组件

这里提到：SIMD 的一大优势是，程序员可以继续按串行逻辑思考，而对并行数据的操作由编译器等实现并行加速。

延迟、带宽与吞吐量（书中 1.1.3）：

指标 含义 常用单位
延迟（latency） 一个操作从开始到完成所需的时间 微秒（μs）
带宽（bandwidth） 单位时间内可处理的数据量 MB/s、GB/s
吞吐量（throughput） 单位时间内成功处理的运算数量 GFlops（每秒十亿次浮点运算）等

延迟衡量完成单次操作的时间，吞吐量衡量单位时间内处理的操作量。架构设计会在降低延迟、提高带宽、提高吞吐量之间做权衡。

按内存组织的分类（书中 1.1.3）：

  • 分布式内存的多节点系统（集群）：多个处理器通过网络连接，各有本地内存。
  • 共享内存的多处理器系统：多个处理器共享同一物理内存，或通过低延迟链路（如 PCIe）相连；共享内存指共享地址空间，不一定是单一物理内存。这类系统包括多核单片系统，以及多芯片、每芯片多核的系统。

从多核到众核：「众核」（many-core） 通常指具有很多核心（几十到几百个）的多核架构。近年来架构正从多核转向众核。GPU 代表众核架构，几乎包含前述所有并行结构：多线程、MIMD、SIMD 以及指令级并行。NVIDIA 将这种架构称为 SIMT（单指令多线程）。

理解与体会：SIMT 可以理解为「在 GPU 上一条指令对应很多线程，每条线程处理自己的数据」，既保留「单指令」的编程简洁性，又实现大规模多线程与 MIMD 的灵活性。后续学到 warp、线程块时会再具体化。

CPU 核心与 GPU 核心（书中原意）：尽管可用多核与众核来区分 CPU 与 GPU，但两种「核心」本质不同。CPU 核心较重，用于处理非常复杂的控制逻辑，优化串行程序执行；GPU 核心较轻，用于优化具有简单控制逻辑的数据并行任务，注重吞吐量。

三、异构计算（书 1.2 节）

3.1 从同构到异构（书 1.2.1）

最初计算机只有 CPU；近年来，高性能计算主流机器不断加入其他处理单元，其中最主要的就是 GPU。GPU 最初为图形并行计算设计，现已演化为强大的、多用途的、完全可编程的处理器，适合任务与数据并行的大规模计算。

CPU 和 GPU 是两个独立的处理器，通过单个计算节点中的 PCI-Express 总线相连。 在这种典型架构中，GPU 作为离散设备存在。值得留意的是：从同构系统到异构系统的转变，是高性能计算史上的一个里程碑。

概念 含义
同构计算 使用同一架构下的一个或多个处理器执行应用
异构计算 使用不同架构的处理器（如 CPU + GPU）执行应用，为任务选择适合的架构，以改进性能

书中也指出：异构系统虽能带来优势，但有效利用受限于应用程序设计的复杂性；若你已有并行编程基础，适应异构并行会相对容易。

典型异构计算节点：包含两个多核 CPU 插槽和两个或更多众核 GPU。GPU 不是独立运行平台，而是 CPU 的协处理器，必须通过 PCIe 总线与基于 CPU 的主机相连才能工作（原书图 1-9）。因此：

  • 主机端（Host）：CPU 所在位置
  • 设备端（Device）：GPU 所在位置

异构应用的组成：

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

应用通常由 CPU 初始化；在把计算密集型任务加载到设备之前，CPU 负责管理设备端的环境、代码和数据。当使用与 CPU 物理上分离的硬件来加速应用中计算密集型部分时，该组件就称为硬件加速器；GPU 是最常见的硬件加速器之一。

3.2 GPU 产品线与性能指标（书 1.2.1）

书中简要列举了应用 NVIDIA GPU 计算平台的产品线：

产品线 面向
Tegra 移动与嵌入式（如平板、手机）
GeForce 消费级图形
Quadro 专业绘图与设计
Tesla 大规模并行计算（如 HPC 加速器）

Fermi 是 Tesla 系列中的一种 GPU 加速器，在高性能计算中应用广泛。NVIDIA 于 2010 年发布的 Fermi 架构是世界上第一款完整的 GPU 计算架构；Kepler 于 2012 年发布，处理能力相比以往有较大提升。本书多数示例可在 Fermi 与 Kepler 上运行；部分示例需要 Kepler 中特殊架构支持。

描述 GPU 能力的两个重要特征（书中 1.2.1）：

  • CUDA 核心数量：决定并行规模
  • 内存大小：决定可处理的数据规模

对应的两个性能指标：

指标 含义 单位
峰值计算性能 每秒能处理的单精度/双精度浮点运算数量 GFlops、TFlops
内存带宽 从内存读取或写入数据的速率 GB/s

计算能力（Compute Capability）：NVIDIA 用「计算能力」描述 Tesla 系列 GPU 的硬件版本。主版本号相同的设备具有相同的核心架构。例如：主版本号 3 对应 Kepler 类架构，主版本号 2 对应 Fermi 类架构，主版本号 1 对应早期 Tesla 类架构。本书所有示例都需要计算能力 2.0 及以上。

3.3 异构计算范例：CPU 与 GPU 的分工（书 1.2.2）

这里区分清楚：GPU 计算并不是要取代 CPU 计算。 对特定程序而言，各有适用场景；CPU 与 GPU 结合能有效提高大规模计算问题的处理速度与性能。

  • CPU：针对动态工作负载优化——短序列计算、不可预测的控制流。
  • GPU：针对由计算主导、控制流简单的工作负载。

可从并行级别和数据规模两方面区分（原书图 1-10）：

特征 更适合 CPU 更适合 GPU
数据规模 较小 较大
控制逻辑 复杂 简单
并行级别 低（如指令级并行） 高（大规模数据并行）

最佳实践（原书图 1-11）：为获得最佳性能，应同时使用 CPU 和 GPU——在 CPU 上执行串行部分或任务并行部分，在 GPU 上执行数据密集型并行部分。这种分工使 CPU + GPU 系统得以充分利用。

CPU 线程与 GPU 线程（书中 1.2.2 重要对比）：

  • CPU 上的线程：通常是重量级实体；操作系统需要切换线程以提供多线程能力，上下文切换缓慢且开销大。四核 CPU 上可同时运行约 16 个线程（若支持超线程则可达约 32 个）。
  • GPU 上的线程：高度轻量级；典型系统中有成千上万个线程在排队。若 GPU 需等待一组线程结束，可以转而执行另一组线程。
  • 设计目标：CPU 核心旨在尽可能减少一个或少量线程的延迟；GPU 核心旨在处理大量并发、轻量级的线程，以最大化吞吐量。
  • 数量对比：现代 NVIDIA GPU 每个多处理器最多可并发支持 1536 个活跃线程；一块有 16 个多处理器的 GPU 可同时支持超过 24000 个活跃线程。

3.4 CUDA：一种异构计算平台（书 1.2.3）

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

CUDA 平台可通过多种方式使用（原书图 1-12）：

  • CUDA 加速库（如 cuBLAS、cuFFT 等）
  • 编译器指令（如 OpenACC）
  • 行业标准语言的扩展（C、C++、Fortran、Python 等）

本书重点介绍 CUDA C。 CUDA C 是标准 ANSI C 的扩展，通过少量语言扩展使异构编程成为可能，并通过 API 管理设备、内存等。CUDA 还是可扩展的编程模型：程序能在具有不同数量核心的 GPU 上透明地扩展并行性，对熟悉 C 的程序员也较易上手。

一个 CUDA 程序包含主机代码（在 CPU 上运行）与设备代码（在 GPU 上运行）的混合。NVIDIA 的 nvcc 编译器在编译过程中将设备代码从主机代码中分离出来（原书图 1-14）：

  • 主机代码：标准 C，由 C 编译器编译。
  • 设备代码（核函数）：用带有关键字标记的 CUDA C 编写，由 nvcc 编译。
  • 链接阶段：在核函数调用与显式 GPU 操作中链接 CUDA 运行时库。
1
2
3
4
5
6
7
8
9
10
11
12
13
14
源代码(.cu)
    │
    ▼
  nvcc ──→ 分离主机代码与设备代码
    │              │
    ▼              ▼
 C 编译器      CUDA 编译器
 (主机代码)    (设备代码/核函数)
    │              │
    ▼              ▼
    链接 + CUDA 运行时库
              │
              ▼
          可执行程序

四、用 GPU 输出 Hello World（书 1.3 节）

下面是书中风格的第一个 CUDA 程序：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
#include <stdio.h>

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

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

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

    cudaDeviceReset();
    return 0;
}

要点说明：

  • __global__：修饰符，表示这是一个核函数（kernel）——在设备上执行，由主机调用。
  • <<<1, 10>>>：执行配置。第一个数 1 表示使用 1 个线程块，第二个数 10 表示每个块 10 个线程，因此共 10 个 GPU 线程执行该核函数。每个线程都会执行一次 helloFromGPU，因此会打印 10 次 “Hello World from GPU!”。
  • cudaDeviceReset()：释放当前设备上的资源，并促使设备端输出缓冲区被刷新（便于在主机端看到 GPU 的 printf 输出）。

编译与运行：

1
2
nvcc hello.cu -o hello
./hello

典型输出（顺序可能因线程调度略有差异）：

1
2
3
4
Hello World from CPU!
Hello World from GPU!
Hello World from GPU!
…（共 10 次）

理解与体会：同一段核函数代码被 10 个线程「同时」执行，正是数据并行的直观体现——所有线程执行相同代码，区别仅在于线程索引（后续会用 threadIdx 等区分每个线程负责的数据）。

五、使用 CUDA C 编程难吗？（书 1.4 节）

不难，但需要思维转换。 主要区别在于对 GPU 架构的熟悉程度；一旦学会用并行思维思考并对 GPU 架构有基本了解，编写面向众多核心的并行程序会变得自然。

数据局部性（书中 1.4）：高性能编程中的核心概念——通过数据重用降低内存访问延迟。

  • 时间局部性：在较短时间内重复访问同一数据。
  • 空间局部性：访问存储空间中相邻的数据。

CPU 通过大容量缓存自动利用局部性；在 CUDA 中，程序员可通过共享内存等机制显式控制数据局部性，这是性能优化的关键之一。

CUDA 的三个核心抽象（书中 1.4）：

  1. 线程组的层次结构：线程 → 线程块 → 网格
  2. 内存的层次结构：寄存器、共享内存、全局内存等
  3. 同步机制：如屏障同步

编写 CUDA 程序时，你主要编写被单个线程执行的一小段串行代码；GPU 负责启动大量线程并行执行这段代码。概念上，把串行代码中的循环「展开」成由不同线程执行，就得到核函数。

开发工具链（书中 1.4）：NVIDIA 提供 Nsight 集成开发环境、CUDA-GDB 调试器、Visual Profiler 性能分析、CUDA-MEMCHECK 内存检查、nvidia-smi 设备管理等，后续章节会逐步用到。

六、重难点与学习体会小结

下表集中梳理本章容易混淆或需要反复强化的点，便于复习与后续章节衔接。

主题 书中要点 个人理解/易错点
数据相关性 限制并行性的主要因素；分析数据相关性是并行算法实现的基本工作 写核函数前先想清楚：每个线程读哪些数据、写哪些数据，是否有写冲突或未定义依赖
块划分 vs 周期划分 块划分：连续数据成块，每线程一块；周期划分：块更小，相邻线程处理相邻块，每线程多块 与第 2 章 grid/block 对应：块划分更易对应「连续内存访问」，周期划分在某些负载均衡场景有用；数据物理上一维存储，分配方式影响访问模式与性能
SIMT GPU 结合多线程、MIMD、SIMD 等，NVIDIA 称为单指令多线程 与 SIMD 类似处是「一条指令、多份数据」；不同处在于 GPU 线程可有自己的分支、更灵活，后续会学到 warp 内的 SIMD 执行
主机 vs 设备 CPU=主机，GPU=设备；主机代码在 CPU 运行，设备代码在 GPU 运行；GPU 是协处理器，通过 PCIe 连接 全书术语统一：分配设备内存、拷贝到设备、在设备上启动核函数；不要混淆主机指针与设备指针
CPU vs GPU 设计哲学 CPU：少线程、低延迟、复杂控制；GPU：海量轻量级线程、高吞吐、简单控制 不要用「跑得快」笼统比较；要问「延迟敏感还是吞吐敏感」「数据规模与并行度如何」
计算能力 主版本号相同则核心架构相同；本书示例需 2.0+ 选 GPU 或写条件编译时要注意 Compute Capability

七、本章小结

第1章建立了三个层次的认知（对应书中全章结构）：

层次 内容
计算思维 串行 vs 并行；数据相关性限制并行性；任务并行 vs 数据并行；块划分与周期划分；弗林分类与 SIMT；延迟、带宽、吞吐量；多核到众核
异构架构 同构→异构；主机/设备；主机代码/设备代码；GPU 为协处理器与硬件加速器；NVIDIA 产品线与计算能力；CPU 与 GPU 分工；CPU 线程与 GPU 线程的差异
CUDA 平台 CUDA 为异构计算平台与编程模型；CUDA C 为 C 的扩展；nvcc 分离主机/设备代码；第一个核函数与执行配置 <<<1,10>>>

本章对我而言最重要的三点：

  1. 数据相关性是并行性的天花板——能并行多少，首先取决于依赖关系，其次才是硬件核心数。
  2. 主机与设备、主机代码与设备代码——从第1章就要建立清晰概念，否则后面内存管理、拷贝、同步都会乱。
  3. CPU 优化延迟、GPU 优化吞吐量；GPU 用海量轻量级线程换高吞吐——这决定了什么问题该放 CPU、什么问题该放 GPU，以及为什么 CUDA 适合数据并行。

下一章预告

第 2 章：CUDA 编程模型 将把第1章的概念落地为可编写、可度量的程序。你将系统学习：

  • CUDA 程序的生命周期：主机与设备内存、数据拷贝、核函数启动
  • 内存管理 API（cudaMalloc、cudaMemcpy 等）
  • 线程层次：网格、线程块、线程索引（gridDim、blockIdx、blockDim、threadIdx）
  • 如何给核函数计时、如何选择网格与块的大小
  • 设备查询与选择（为多 GPU 与正确性验证打基础）

从「为什么用 GPU」到「怎么写、怎么测」——我们下一章见。

本章自测

  1. 数据并行与任务并行的区别是什么？CUDA 主要针对哪一种？
  2. 弗林分类中，SIMT 与 SIMD 各属哪一类？GPU 属于哪种？
  3. 写出一个最简单的 __global__ 核函数并在主机上用 <<<1, 10>>> 启动，说明两个数字的含义。

答案与解析

  1. 数据并行是「同时处理多份数据」，任务并行是「同时执行多个不同任务」。CUDA 主要针对数据并行，通过大量线程对数据划分并行处理。
  2. SIMD 属「单指令多数据」，SIMT 是 NVIDIA 对「单指令多线程」的命名，本质也是单指令多数据的一种实现方式；GPU 采用 SIMT，每 32 线程为一 warp 执行相同指令。
  3. 例如 __global__ void hello() { }，主机调用 hello<<<1, 10>>>()。<<<1, 10>>> 表示 1 个块、每块 10 个线程，即共 10 个线程执行该核函数。

系列导航：导读 | 下一篇：第2章 CUDA编程模型

本文为「CUDA C编程权威指南」系列博客第 1 篇，共 10 章。基于《Professional CUDA C Programming》by John Cheng, Max Grossman, Ty McKercher。