CUDA C编程权威指南-第三章：CUDA执行模型

系列导航：导读 | 上一篇：第2章 CUDA编程模型 | 下一篇：第4章 全局内存

系列第 3 篇。第 2 章里我们学会了组织网格与块、给核函数计时，也看到同一算法换一种执行配置性能差很多（比如矩阵加法里 32×16 的块比 16×16 快）。这一章从硬件视角解释原因：线程束怎么调度、分支分化怎么回事、占用率和延迟隐藏怎么指导配置，以及动态并行、循环展开等进阶内容。

---

前言：从「怎么写」到「怎么跑」

执行模型描述的是在特定计算架构上指令如何被执行。第 2 章的编程模型给出了「网格—块—线程」的逻辑视图与写法；本章则从硬件视角揭示：GPU 如何把块分配到 SM、如何把线程组织成线程束（warp）调度、为何有的配置更快、如何用占用率和延迟隐藏指导优化。从「怎么写」到「为什么这样写更好」，是迈向高性能 CUDA 的必经一步。

本章围绕指令吞吐量与计算资源展开；第 4、5 章将在此基础上深入内存访问与带宽。

---

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

1.1 为什么要学「执行模型」

第 2 章回答了「怎么写」：如何组织网格与块、如何写核函数、如何计时。但你一定会问：为什么 32×16 的块比 16×16 更快？有没有选择网格和块配置的准则？ 这一章会说明：尽管可以通过反复试验找到最佳配置，但若不了解硬件如何调度与执行线程，就很难系统性地优化。

执行模型描述的是在特定计算架构上指令如何被执行。它从硬件视角揭示：线程束（warp）是什么、SM 如何调度、为何会有分支分化、占用率与延迟隐藏如何影响性能。学完本章，你就能从「试参数」进阶到「看指标、懂原理、做权衡」。

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

学习目标	检验方式
理解 SM、线程束、SIMT 与「32」的含义	能解释为什么块最内层维度应为 32 的倍数
理解线程束分化的成因与代价	能写出避免分化的数据分区方式（如按 warp 分组）
理解占用率的定义及与寄存器/共享内存的关系	能用 -Xptxas -v 和占用率计算器分析内核
理解延迟隐藏（算术 vs 内存延迟）及所需并行量	能复述 Fermi 上隐藏内存延迟约需 36 个线程束/SM 的推导思路
掌握网格和块的启发式（块大小、最内层维度、块数）	能解释为何高占用率不一定带来最高性能
会用 nvprof 看 achieved_occupancy、gld_efficiency 等指标	能结合矩阵求和实验说明「指标平衡」
理解并行归约中分化来源及改进（相邻/交错、索引重组）	能写出 reduceNeighboredLess 式的条件避免同一 warp 内分化
了解循环展开与动态并行的作用与适用场景	能说出展开对延迟隐藏的意义、动态并行对递归的意义

1.3 博客阅读导图（本章架构）

阅读建议：本章较长，二、三、四、五节建议精读（执行模型与归约是后续内存与优化基础）；六节中「动态并行」为选读，首次学习可跳过，待掌握流与多 GPU 后再回看。

第 3 章  CUDA 执行模型
│
├── 二、CUDA 执行模型概述（书 3.1 节）【精读】
│   ├── GPU 架构概述（SM、线程束、SIMT、32、资源与同步）
│   ├── Fermi / Kepler 架构要点
│   └── 配置文件驱动优化（nvvp、nvprof、事件与指标）
│
├── 三、理解线程束执行的本质（书 3.2 节）【精读】
│   ├── 线程束与线程块、硬件视图下一维化
│   ├── 线程束分化、断定、避免分化的数据分区
│   ├── 资源分配、活跃线程束分类
│   ├── 延迟隐藏（算术/内存延迟、Fermi 公式）
│   ├── 占用率与网格/块准则
│   ├── 同步栅栏
│   └── 可扩展性
│
├── 四、并行性的表现（书 3.3 节）【精读】
│   └── 矩阵求和与 nvprof、网格与块启发式
│
├── 五、避免分支分化（书 3.4 节）【精读】
│   └── 并行归约、分化与改进（相邻/交错配对）
│
├── 六、展开循环与动态并行（书 3.5、3.6 节）【精读：展开 | 选读：动态并行】
│
└── 七、总结与重难点小结（书 3.7 节）

---

二、CUDA 执行模型概述（书 3.1 节）

2.1 GPU 架构概述（书 3.1.1）

一般来说，执行模型会提供一个操作视图，说明如何在特定的计算架构上执行指令。CUDA 执行模型揭示了 GPU 并行架构的抽象视图，使我们能够据此分析线程的并发。第2章已经介绍了 CUDA 编程模型中两个主要的抽象概念：内存层次结构和线程层次结构。本章将重点介绍指令吞吐量，第4章和第5章将介绍更多关于高效内存访问的内容。

GPU 架构是围绕**流式多处理器（Streaming Multiprocessor, SM）**的可扩展阵列搭建的。可以通过复制这种架构的构建块来实现 GPU 的硬件并行。

Fermi SM 的关键组件（原书图 3-1）包括：

• CUDA 核心（整数 ALU + 浮点 FPU）
• 共享内存 / 一级缓存
• 寄存器文件
• 加载/存储单元（LD/ST）
• 特殊功能单元（SFU，用于正弦、余弦、平方根、插值等）
• 线程束调度器

下面用框图概括 SM 内各组件的逻辑关系（建议配合原书图 3-1 理解）：

flowchart TB
    subgraph SM [流式多处理器 SM]
        WarpSched[线程束调度器]
        Cores[CUDA 核心]
        RegFile[寄存器文件]
        SharedL1[共享内存 / L1 缓存]
        LDST[加载存储单元]
        SFU[特殊功能单元 SFU]
    end
    WarpSched --> Cores
    WarpSched --> LDST
    WarpSched --> SFU
    Cores --> RegFile
    LDST --> SharedL1

GPU 中的每一个 SM 都能支持数百个线程并发执行，每个 GPU 通常有多个 SM，所以在一个 GPU 上并发执行数千个线程是有可能的。当启动一个内核网格时，线程块被分布在可用的 SM 上执行。线程块一旦被调度到一个 SM 上，其中的线程只会在那个指定的 SM 上并发执行。多个线程块可能会被分配到同一个 SM 上，而且是根据 SM 资源的可用性进行调度的。

同一线程中的指令利用指令级并行进行流水线化；此外，CUDA 还提供了线程级并行。CUDA 采用单指令多线程（SIMT）架构来管理和执行线程：每 32 个线程为一组，被称为线程束（warp）。线程束中的所有线程同时执行相同的指令。每个线程都有自己的指令地址计数器和寄存器状态，利用自身的数据执行当前的指令。每个 SM 都将分配给它的线程块划分到包含 32 个线程的线程束中，然后在可用的硬件资源上调度执行。

SIMT 与 SIMD 的区别

SIMT 架构与 SIMD（单指令多数据）架构相似——两者都是将相同的指令广播给多个执行单元来实现并行。一个关键的区别是：SIMD 要求同一个向量中的所有元素要在一个统一的同步组中一起执行，而 SIMT 允许属于同一线程束的多个线程独立执行。尽管一个线程束中的所有线程在相同的程序地址上同时开始执行，但是单独的线程仍有可能有不同的行为。

SIMT 模型包含三个 SIMD 所不具备的关键特征（需要掌握的三点）：

1. 每个线程可以有一个独立的执行路径
2. 每个线程都有自己的寄存器状态
3. 每个线程都有自己的指令地址计数器

因此 SIMT 确保可以编写独立的线程级并行代码、标量线程以及用于协调线程的数据并行代码。

神奇的数字：32

尤其要注意的是：在 CUDA 程序里，32 是一个神奇的数字。它来自于硬件系统，也对软件的性能有着重要的影响。从概念上讲，它是 SM 用 SIMD 方式所同时处理的工作粒度。优化工作负载以适应线程束（一组 32 个线程）的边界，一般这样会更有效地利用 GPU 计算资源。在后面的章节中将会多次看到「块的最内层维度应为线程束大小的倍数」这一准则。

资源与同步

一个线程块只能在一个 SM 上被调度。一旦线程块在一个 SM 上被调度，就会保存在该 SM 上直到执行完成。在同一时间，一个 SM 可以容纳多个线程块。

在 SM 中，共享内存和寄存器是非常重要的资源。共享内存被分配在 SM 上的常驻线程块中，寄存器在线程中被分配。线程块中的线程通过这些资源可以进行相互的合作和通信。

尽管线程块里的所有线程都可以逻辑地并行运行，但是并不是所有线程都可以同时在物理层面执行。因此，线程块里的不同线程可能会以不同的速度前进。在并行线程中共享数据可能会引起竞争：多个线程使用未定义的顺序访问同一个数据，从而导致不可预测的程序行为。CUDA 提供了一种用来同步线程块里的线程的方法（__syncthreads()），从而保证所有线程在进一步动作之前都达到执行过程中的一个特定点。然而，没有提供块间同步的原语。

尽管线程块里的线程束可以任意顺序调度，但活跃的线程束的数量还是会由 SM 的资源所限制。当线程束由于任何理由闲置的时候（如等待从设备内存中读取数值），SM 可以从同一 SM 上的常驻线程块中调度其他可用的线程束。在并发的线程束间切换并没有开销，因为硬件资源已经被分配到了 SM 上的所有线程和块中，所以最新被调度的线程束的状态已经存储在 SM 上。

理解与体会：可以把 SM 想象成「线程束的调度中心」——块一旦被分配到某个 SM，就一直在该 SM 上执行；块内的线程被拆成若干个 warp，调度器轮流执行「就绪」的 warp。正因为切换无开销，才需要足够多的活跃 warp 来隐藏内存延迟，这也是后面「占用率」和「延迟隐藏」的出发点。

---

2.2 Fermi 架构（书 3.1.2）

Fermi 架构是第一个完整的 GPU 计算架构，能够为大多数高性能计算应用提供所需要的功能，已被广泛应用于加速生产工作负载。

Fermi 架构的逻辑框图（原书图 3-3）（重点为 GPU 计算，忽略图形具体组成部分）。Fermi 的特征包括：

• 最多 512 个 CUDA 核心（加速器核心）。每个 CUDA 核心都有一个全流水线的整数算术逻辑单元（ALU）和一个浮点运算单元（FPU），每个时钟周期执行一个整数或浮点数指令。
• CUDA 核心被组织到 16 个 SM 中，每一个 SM 含有 32 个 CUDA 核心。
• 6 个 384 位的 GDDR5 DRAM 存储器接口，支持多达 6GB 的全局机载内存。
• GigaThread 引擎：全局调度器，用来分配线程块到 SM 的线程束调度器上。
• 768 KB 的二级缓存，被 16 个 SM 所共享。

每个多处理器有（原书图 3-1）：

• 16 个加载/存储单元：允许每个时钟周期内有 16 个线程（线程束的一半）计算源地址和目的地址。
• 特殊功能单元（SFU）：执行固有指令，如正弦、余弦、平方根和插值。每个 SFU 在每个时钟周期内的每个线程上执行一个固有指令。
• 2 个线程束调度器和 2 个指令调度单元。当一个线程块被指定给一个 SM 时，线程块中的所有线程被分成线程束。两个线程束调度器选择两个线程束，再把一个指令从线程束中发送到一个组上——组里有 16 个 CUDA 核心、16 个加载/存储单元或 4 个特殊功能单元（原书图 3-4）。

Fermi 架构（计算能力 2.x）可以在每个 SM 上同时处理 48 个线程束，即可在一个 SM 上同时常驻 1536 个线程。

Fermi 架构的一个关键特征是有一个 64 KB 的片内可配置存储器，它在共享内存与一级缓存之间进行分配。对于许多高性能的应用程序，共享内存是影响性能的一个关键因素。CUDA 提供了运行时 API 来调整共享内存和一级缓存的数量；根据给定的内核中共享内存或缓存的使用修改片内存储器的配置，可以提高性能（第 4、5 章详细介绍）。

Fermi 架构也支持并发内核执行：在相同的 GPU 上、相同应用程序的上下文中，同时启动多个内核。Fermi 允许多达 16 个内核同时在设备上运行。从程序员的角度看，并发内核执行使 GPU 表现得更像 MIMD 架构。

---

2.3 Kepler 架构（书 3.1.3）

发布于 2012 年秋季的 Kepler GPU 架构是一种快速、高效、高性能的计算架构。书中以 Kepler K20X 为例，其包含 15 个 SM 和 6 个 64 位的内存控制器。Kepler 的三大重要创新为：

1. 强化的 SM
2. 动态并行
3. Hyper-Q 技术

每个 Kepler SM 单元包含（原书图 3-7）：

• 192 个单精度 CUDA 核心、64 个双精度单元
• 32 个特殊功能单元（SFU）、32 个加载/存储单元（LD/ST）

每个 Kepler SM 包括 4 个线程束调度器和 8 个指令调度器，以确保在单一的 SM 上同时发送和执行 4 个线程束。Kepler K20X 架构（计算能力 3.5）可以同时在每个 SM 上调度 64 个线程束，即在一个 SM 上可同时常驻 2048 个线程。K20X 架构中寄存器文件容量达到 64 KB（Fermi 为 32 KB），同时 K20X 还允许片内存储器在共享内存和一级缓存间有更多的分区。K20X 能够提供超过 1 TFlop 的峰值双精度计算能力，相较于 Fermi 的设计，功率效率提高了 80%，每瓦性能也提升了三倍。

• 动态并行：允许 GPU 动态启动新的网格。任一内核都能启动其他的内核，并管理核间需要的依赖关系。这一特点也让你更容易创建和优化递归及与数据相关的执行模式。原书图 3-8 对比了「无动态并行时主机在 GPU 上启动每一个内核」与「有动态并行时 GPU 能够启动嵌套内核，消除与 CPU 通信的需求」。
• Hyper-Q 技术：增加了更多的 CPU 和 GPU 之间的同步硬件连接，确保 CPU 核心能够在 GPU 上同时运行更多的任务。Fermi 依赖一个单一的硬件工作队列，可能导致一个单独的任务阻塞队列中该任务之后的所有其他任务；Kepler Hyper-Q 消除了这个限制，在主机与 GPU 之间提供了 32 个硬件工作队列（原书图 3-9）。

书中表 3-1 简要总结了不同计算能力下的主要架构特点。下表根据书中内容整理，便于查阅与配置内核：

计算能力	架构类型	每 SM 最大线程数	每 SM 最大线程束数	每块最大线程数	每 SM 寄存器(32位)	每块共享内存
2.0	Fermi	1536	48	1024	32 K	48 KB
2.1	Fermi	1536	48	1024	32 K	48 KB
3.0	Kepler	2048	64	1024	64 K	48 KB
3.5	Kepler	2048	64	1024	64 K	48 KB

实际限制以 cudaGetDeviceProperties 为准；不同型号可能有细微差异。

---

2.4 配置文件驱动优化（书 3.1.4）

性能分析是通过检测来分析程序性能的行为，包括：

• 函数调用的频率和持续时间
• 特殊指令的使用
• 应用程序代码的空间（内存）或时间复杂度

性能分析是程序开发中的关键一步，特别是对于优化 HPC 应用程序代码。性能分析往往需要对平台的执行模型有一个基本的理解，才能制定应用程序的优化方法。这里提到：开发一个 HPC 应用程序通常包括两个主要步骤——提高代码的正确性和提高代码的性能。对于第二步，使用配置文件驱动的方法是很自然的。

配置文件驱动的发展对于 CUDA 编程尤为重要，原因主要有：

1. CUDA 提供了一个硬件架构的抽象，它能够让用户控制线程并发；性能分析工具可以检测和优化，并将优化可视化。
2. CUDA 将 SM 中的计算资源在该 SM 中的多个常驻线程块之间进行分配，这种分配形式导致一些资源成为了性能限制者；性能分析工具能帮助我们理解计算资源是如何被利用的。
3. 一个单纯的内核应用一般不会产生最佳的性能；性能分析工具能帮助你找到代码中影响性能的关键部分，也就是性能瓶颈。

CUDA 提供的两个主要性能分析工具：

• nvvp：独立的可视化分析器，可显示 CPU 与 GPU 上的程序活动时间表，并分析潜在的性能瓶颈并给出建议。既可独立使用，也可作为 Nsight Eclipse Edition 的一部分。
• nvprof：命令行分析器（随 CUDA 5 一起发布）。可以获得 CPU 与 GPU 上 CUDA 关联活动的时间表（内核执行、内存传输、CUDA API 调用），也可以获得硬件计数器和 CUDA 内核的性能指标。除了预定义的指标，还可以利用基于分析器获得的硬件计数器来自定义指标。

事件与指标（书中概念）：

• 事件（Event）：可计算的活动，对应一个在内核执行期间被收集的硬件计数器。
• 指标（Metric）：内核的特征，由一个或多个事件计算得到。

注意：大多数计数器是按流式多处理器报告的，而不是整个 GPU；一次运行只能获得几个计数器；有些计数器的获得是相互排斥的，因此多次性能分析运行往往需要获取所有相关的计数器。由于 GPU 执行中的变化（如线程块和线程束调度），经重复运行，计数器值可能不是完全相同的。

常见的有 3 种常见的限制内核性能的因素：

1. 指令和内存延迟
2. 计算资源（寄存器、共享内存等）
3. 存储带宽

本章主要介绍 1 和 2；后续章节将讨论其余因素。了解硬件资源的详细信息：作为 CUDA C 程序员，如果想改善内核的性能，必须对硬件资源有一定的了解；即使仅掌握最基本的 GPU 体系架构知识，也能够编写出更好的代码，并充分开发设备的性能。

---

三、理解线程束执行的本质（书 3.2 节）

3.1 线程束和线程块（书 3.2.1）

启动内核时，从软件的角度看，似乎所有的线程都是并行运行的——在逻辑上这是正确的。但从硬件的角度来看，不是所有线程在物理上都可以同时并行地执行。线程束是 SM 中基本的执行单元。

当一个线程块的网格被启动后，网格中的线程块分布在 SM 中。一旦线程块被调度到一个 SM 上，线程块中的线程会被进一步划分为线程束。一个线程束由 32 个连续的线程组成；在一个线程束中，所有的线程按照 SIMT 方式执行——即所有线程执行相同的指令，每个线程在私有数据上进行操作。原书图 3-10 展示了线程块的逻辑视图和硬件视图之间的关系。

重要：从硬件的角度来看，所有的线程都被组织成了一维的。线程块可以被配置为一维、二维或三维，但在一个块中，每个线程都有一个唯一的 ID；对于一维的线程块，唯一的线程 ID 存储在 threadIdx.x 中，并且 拥有连续 threadIdx.x 值的线程被分组到线程束中。

例如，一个有 128 个线程的一维线程块被组织成 4 个线程束：

• Warp 0：threadIdx.x = 0～31
• Warp 1：threadIdx.x = 32～63
• Warp 2：threadIdx.x = 64～95
• Warp 3：threadIdx.x = 96～127

因此，块的最内层维度（如一维块的 blockDim.x）若不是 32 的倍数，会导致部分线程束「未填满」，影响合并访问与占用率。

重点：硬件视角下线程在块内是一维排列的，连续 32 个 threadIdx（如 threadIdx.x=0～31）组成一个 warp。所以二维块 (16,16) 时，最内层维度是 16，每 16 个线程一组，半个 warp 一组，会导致加载效率下降（见 3.3 节矩阵求和实验）。

---

3.2 线程束分化（Warp Divergence）（书 3.2.2）

控制流是高级编程语言的基本构造。GPU 支持传统的、C 风格的、显式的控制流结构，如 if…then…else、for、while。CPU 拥有复杂的硬件以执行分支预测——在每个条件检查中预测应用程序的控制流会使用哪个分支。如果预测正确，CPU 中的分支只需付出很小的性能代价；如果预测不正确，CPU 可能会停止运行很多个周期，因为指令流水线被清空了。GPU 是相对简单的设备，它没有复杂的分支预测机制。

一个线程束中的所有线程在同一周期中必须执行相同的指令。如果一个线程执行一条指令，那么线程束中的所有线程都必须执行该指令。如果在同一线程束中的线程使用不同的路径通过同一个应用程序，就会产生问题。例如（书中示例）：

if (cond) {
    // 分支 A：16 个线程执行
} else {
    // 分支 B：另外 16 个线程执行
}

假设在一个线程束中有 16 个线程的 cond 为 true，另外 16 个为 false。一半的线程束需要执行 if 块中的指令，另一半需要执行 else 块中的指令。在同一线程束中的线程执行不同的指令，被称为线程束分化。

硬件如何解决这一「悖论」？如果一个线程束中的线程产生分化，线程束将连续执行每一个分支路径，而禁用不执行这一路径的线程。即：先执行 if 块（此时 else 对应的 16 个线程被禁用），再执行 else 块（此时 if 对应的 16 个线程被禁用）。线程束分化会导致性能明显地下降——在前面的例子中，线程束中并行线程的数量减少了一半。条件分支越多，并行性削弱越严重。

注意：线程束分化只发生在同一个线程束中。在不同的线程束中，不同的条件值不会引起线程束分化。

避免分化的数据分区示例

书中用简单的算术内核做了对比。可以用「偶数和奇数线程」的方法模拟导致分化的条件：

// 导致分化：同一 warp 内一半走 if、一半走 else
__global__ void mathKernel1(float *c) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid % 2 == 0) {
        c[tid] = 1.0f;  // 偶数线程
    } else {
        c[tid] = 2.0f;  // 奇数线程
    }
}

// 不导致分化：按 warp 分组，同一 warp 内走同一分支
__global__ void mathKernel2(float *c) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if ((tid / warpSize) % 2 == 0) {
        c[tid] = 1.0f;
    } else {
        c[tid] = 2.0f;
    }
}

mathKernel2 中，条件 (tid / warpSize) % 2 使同一线程束内的 32 个线程要么全为 0 要么全为 1，因此同一 warp 内不会分化。书中对应的示例可从 Wrox 网站下载 simpleDivergence.cu；测量时建议添加一次「预热」内核启动以去除首次运行的间接开销。

断定（Predication）与编译器优化

有一个常见现象：用 nvprof 测量时，没有报告显示出有分支分化（即分支效率是 100%）。这是因为 CUDA 编译器将短的、有条件的代码段用断定指令（predication）取代了分支指令。在断定中，根据条件把每个线程中的一个断定变量设置为 1 或 0；两种条件流路径被完全执行，但只有断定为 1 的指令的结果被保留，断定为 0 的指令不写回，但相应线程也不会停止。只有在条件语句的指令数小于某个阈值时，编译器才用断定指令替换分支指令；一段很长的代码路径肯定会导致线程束分化。

书中通过重写 mathKernel1 为 mathKernel3（将 if-else 分离为多个 if 语句），使分化分支的数量翻倍，从而让 nvprof 能直接显示出分支分化。重要提示（书中总结）：

• 当一个分化的线程采取不同的代码路径时，会产生线程束分化。
• 不同的 if-then-else 分支会连续执行。
• 尝试调整分支粒度以适应线程束大小的倍数，避免线程束分化。
• 不同线程束可以执行不同代码且无须以牺牲性能为代价。

难点与学习思考：分支分化是「同一 warp 内」的概念——不同 warp 走不同分支不会互相拖累。优化时优先考虑「让同一 warp 内的 32 个线程走同一条分支」：要么按数据分区（如 (tid/warpSize)%2），要么重组索引（如归约中让参与计算的线程 ID 连续）。另外，nvprof 可能显示 100% 分支效率，是因为短分支被编译器替换成了断定指令；只有分支体足够长时才会真正分化。

---

3.3 资源分配（书 3.2.3）

线程束的本地执行上下文主要由以下资源组成：

• 程序计数器
• 寄存器
• 共享内存

由 SM 处理的每个线程束的执行上下文，在整个线程束的生存期中是保存在芯片内的。因此，从一个执行上下文切换到另一个执行上下文没有损失。每个 SM 都有固定大小的寄存器组（存储在寄存器文件中，在线程中分配）和固定数量的共享内存（在线程块中分配）。对于一个给定的内核，同时存在于同一个 SM 中的线程块和线程束的数量取决于在 SM 中可用的且内核所需的寄存器和共享内存的数量。

书中图 3-13 显示：若每个线程消耗的寄存器越多，则可以放在一个 SM 中的线程束就越少。图 3-14 显示：若一个线程块消耗的共享内存越多，则在一个 SM 中可以被同时处理的线程块就会变少。资源可用性通常会限制 SM 中常驻线程块的数量；每个 SM 中寄存器和共享内存的数量因设备计算能力而不同。如果每个 SM 没有足够的寄存器或共享内存去处理至少一个块，那么内核将无法启动。

书中表 3-2 给出了不同计算能力下的一些关键限度。下表根据书中内容整理，便于与占用率计算器对照（实际以 cudaGetDeviceProperties 为准）：

计算能力	每线程最大寄存器数	每块最大共享内存	每块最大线程数	每 SM 最大线程数	每 SM 最大线程束数
2.0	63	48 KB	1024	1536	48
3.0	63	48 KB	1024	2048	64
3.5	255	48 KB	1024	2048	64

书中代码清单 3-2 提供了简单设备属性查询的完整示例（simpleDeviceQuery.cu）。下面给出书中风格的完整代码，可直接编译运行以查询当前 GPU 的上述限制：

#include <stdio.h>

int main(int argc, char **argv) {
    int deviceCount;
    cudaGetDeviceCount(&deviceCount);
    if (deviceCount == 0) {
        printf("No CUDA device found.\n");
        return 1;
    }
    int dev = 0;
    cudaSetDevice(dev);

    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, dev);

    printf("Device %d: \"%s\"\n", dev, deviceProp.name);
    printf("  Compute capability: %d.%d\n", deviceProp.major, deviceProp.minor);
    printf("  Multiprocessor count: %d\n", deviceProp.multiProcessorCount);
    printf("  Max threads per block: %d\n", deviceProp.maxThreadsPerBlock);
    printf("  Max threads per multiprocessor: %d\n", deviceProp.maxThreadsPerMultiProcessor);
    printf("  Max warps per multiprocessor: %d\n", deviceProp.maxThreadsPerMultiProcessor / 32);
    printf("  Registers per block: %d\n", deviceProp.regsPerBlock);
    printf("  Shared memory per block: %zu bytes\n", deviceProp.sharedMemPerBlock);

    return 0;
}

编译运行：nvcc simpleDeviceQuery.cu -o simpleDeviceQuery && ./simpleDeviceQuery。最大线程束数用于占用率计算：maxWarpsPerSM = maxThreadsPerMultiProcessor / 32。

活跃线程束的分类

当计算资源（如寄存器和共享内存）已分配给线程块时，线程块被称为活跃的块，它所包含的线程束被称为活跃的线程束。活跃的线程束可以进一步分为三种类型：

• 选定的线程束（Selected）：正在被调度到执行单元上执行的线程束。
• 符合条件的线程束（Eligible）：已就绪但尚未被调度。
• 阻塞的线程束（Blocked）：没有做好执行的准备（例如等待内存操作完成）。

线程束符合执行条件需要同时满足（书中）：(1) 32 个 CUDA 核心可用于执行当前指令；(2) 当前指令中所有的参数都已就绪。例如，Kepler SM 上活跃的线程束数量在任何时候都必须小于或等于 64（架构限度）；在任何周期中，选定的线程束数量都小于或等于 4。如果线程束阻塞，线程束调度器会令一个符合条件的线程束代替它去执行。由于计算资源是在线程束之间分配的，而且在线程束的整个生存期中都保持在芯片内，因此线程束上下文的切换是非常快的。为了隐藏由线程束阻塞造成的延迟，需要让大量的线程束保持活跃。

---

3.4 延迟隐藏（书 3.2.4）

SM 依赖线程级并行以最大化功能单元的利用率，因此利用率与常驻线程束的数量直接相关。指令延迟定义为：在指令发出和完成之间的时钟周期。当每个时钟周期中所有的线程调度器都有一个符合条件的线程束时，可以达到计算资源的完全利用——通过在其他常驻线程束中发布其他指令，可以隐藏每个指令的延迟。与在 CPU 上用 C 语言编程相比，延迟隐藏在 CUDA 编程中尤为重要：CPU 核心是为同时最小化一个或两个线程的延迟而设计的，而 GPU 则是为处理大量并发和轻量级线程以最大化吞吐量而设计的。GPU 的指令延迟被其他线程束的计算隐藏。

指令可分为两种基本类型（书中）：

• 算术指令：延迟约为 10～20 个周期。
• 内存指令：延迟约为全局内存访问 400～800 个周期。

书中图 3-15 表示线程束 0 阻塞时，线程束调度器选取其他线程束执行，当线程束 0 符合条件时再执行它。

带宽通常指理论峰值，吞吐量指已达到的值。对于算术运算，所需的并行可以表示成「隐藏算术延迟所需要的操作数量」。书中表 3-3 列出了 Fermi 和 Kepler 设备上示例算术运算（32 位浮点乘加）的每 SM 每周期操作数等；吞吐量因不同的算术指令而不同。

Fermi 内存延迟隐藏计算

书中以 Fermi 为例，给出了隐藏全局内存延迟所需并行量的具体计算：

• Fermi 内存频率（在 Tesla C2070 上测量）为 1.566 GHz。带宽 144 GB/s 可转换为每周期字节数：
  144 GB/s ÷ 1.566 GHz ≈ 92 字节/周期。
• 用内存延迟乘以每周期字节数：若取约 800 周期，则约 74 KB 的内存 I/O 运行才能实现充分利用（该值针对整个设备，因为内存带宽是针对整个设备的）。
• 假设每个线程将**一个 float（4 字节）**从全局内存移动到 SM 中用于计算，则：
  74 KB ÷ 4 字节/线程 ≈ 18 500 个线程
  18 500 ÷ 32 ≈ 579 个线程束
  Fermi 有 16 个 SM，因此 579 ÷ 16 ≈ 36 个线程束/SM，才能隐藏所有的内存延迟。

如果每个线程执行多个独立的 4 字节加载，隐藏内存延迟需要的线程就可以更少。延迟隐藏取决于每个 SM 中活跃线程束的数量，这一数量由执行配置和资源约束（内核中寄存器和共享内存的使用情况）隐式决定。选择一个最优执行配置的关键是在延迟隐藏和资源利用之间找到一种平衡。

显示充足的并行：计算所需并行的一个简单公式是，用每个 SM 核心的数量乘以在该 SM 上一条算术指令的延迟。例如，Fermi 有 32 个单精度浮点流水线，算术指令延迟约 20 个周期，所以每个 SM 至少需要有 32×20＝640 个线程使设备处于忙碌状态——然而这只是一个下边界。

延迟隐藏所需的并行量的显式公式与数值（以 Fermi 为例）如下。下面将关键公式与计算过程单独列出，便于对照书本与实验：

(1) 算术延迟隐藏（下界）

[
\text{每 SM 所需线程数} \geq \text{每 SM CUDA 核心数} \times \text{算术指令延迟（周期）}
]

例如 Fermi：(32 \times 20 = 640) 线程/SM。

(2) 全局内存延迟隐藏（书中 Fermi 示例）

• 内存带宽 144 GB/s，内存频率 1.566 GHz → 每周期字节数：
  [
  \frac{144\ \text{GB/s}}{1.566\ \text{GHz}} \approx 92\ \text{字节/周期}
  ]
• 取内存延迟约 800 周期，则充分隐藏延迟所需「在途」数据量约：
  [
  92 \times 800 \approx 74\ \text{KB（设备级）}
  ]
• 若每个线程只加载 1 个 float（4 字节）：
  [
  \text{所需线程数} \approx \frac{74\ \text{KB}}{4\ \text{字节}} \approx 18500,\quad
  \text{所需线程束数} \approx \frac{18500}{32} \approx 579
  ]
• Fermi 共 16 个 SM，故每 SM 约需 (579/16 \approx 36) 个活跃线程束才能较好隐藏全局内存延迟。

原书表 3-3 给出了 Fermi 与 Kepler 在示例算术运算（如 32 位浮点乘加）下的每 SM 每周期操作数等，用于与实测吞吐量对比。下表为典型含义（具体数值以原书表 3-3 为准）：

设备类型	算术指令延迟（周期）	每 SM 每周期操作数（示例）
Fermi	~20	见原书表 3-3
Kepler	~20	见原书表 3-3

理解与体会：延迟隐藏的本质是「用其他 warp 的计算把等待时间填满」。内存延迟（400～800 周期）远大于算术延迟（10～20 周期），所以内存型内核更需要大量活跃 warp。「每 SM 约 36 个线程束」是 Fermi 上的经验量级，实际以你设备的带宽和延迟为准；关键是建立「并行量 ↔ 延迟隐藏」的直觉，而不是死记数字。

---

3.5 占用率（Occupancy）（书 3.2.5）

在每个 CUDA 核心中指令是顺序执行的。当一个线程束阻塞时，SM 切换执行其他符合条件的线程束。理想情况下，我们想要有足够的线程束占用设备的核心。占用率定义为：每个 SM 中活跃的线程束占最大线程束数量的比值。公式为：

[
\text{占用率} = \frac{\text{每个 SM 上活跃的线程束数}}{\text{每个 SM 上最大线程束数}}
]

其中每个 SM 上最大线程束数为：

[
\text{最大线程束数} = \frac{\texttt{maxThreadsPerMultiProcessor}}{32}
]

例如，Fermi 的 maxThreadsPerMultiProcessor = 1536，则最大线程束数 = 1536/32 = 48。活跃的线程束数由执行配置（网格与块的维度）以及内核对寄存器和共享内存的占用共同决定：若每个线程用的寄存器过多，或每个块用的共享内存过多，则 SM 上能同时驻留的块和线程束会减少，从而占用率低于理论最大值。因此需要留意：为了提高占用率，需要调整线程块配置或重新调整资源的使用情况，以允许更多的线程束同时处于活跃状态。

CUDA 占用率计算器：CUDA 工具包包含一个电子表格（原书图 3-17），有助于选择网格和块的维数以使一个内核的占用率最大化。使用步骤大致为：

1. 提供 GPU 的计算能力；物理限制部分会自动填充。
2. 输入内核资源信息：每个块的线程数（执行配置，如 256、1024）、每个线程的寄存器数、每个块的共享内存（字节）。
3. 获取资源使用：编译时加 -Xptxas -v 可报告每个内核的寄存器和共享内存使用量，例如：
   nvcc -Xptxas -v -o app kernel.cu
   输出中会看到类似 Used 32 registers, 0 bytes smem 的信息。
4. 限制寄存器以调节占用率：使用 -maxrregcount=NUM 可限制每个线程使用的寄存器数上限，从而允许更多线程束常驻。例如：
   nvcc -maxrregcount=32 -o app kernel.cu
   占用率计算器会根据你填写的「每线程寄存器」和「每块共享内存」给出当前配置下的占用率及建议。

极端地操纵线程块会限制资源的利用（书中）：

• 小线程块：每个块中线程太少，会在所有资源被充分利用之前导致硬件达到每个 SM 的线程束数量限制。
• 大线程块：每个块中有太多的线程，会导致在每个 SM 中每个线程可用的硬件资源较少。

网格和线程块大小的准则（书中总结）：

• 保持每个块中线程数量是线程束大小（32）的倍数。
• 避免块太小：每个块至少要有 128 或 256 个线程。
• 根据内核资源的需求调整块大小。
• 块的数量要远远多于 SM 的数量，从而在设备中显示有足够的并行。
• 通过实验得到最佳执行配置和资源使用情况。

占用率只关注「每个 SM 里有多少并发线程/线程束」；但在每个 SM 中并发线程或线程束的数量；然而，充分的占用率不是性能优化的唯一目标。内核一旦达到一定级别的占用率，进一步增加占用率可能不会改进性能；还有许多其他因素需要在后续章节中考虑。

重难点：本章矩阵求和实验会直接说明——(16,16) 占用率最高但不是最快，因为最内层维度 16 导致加载效率（gld_efficiency）下降。所以优化时要在占用率、内存访问模式、分支效率之间找平衡，而不是盲目追求 100% 占用率。

---

3.6 同步栅栏（书 3.2.6）

同步是一个在许多并行编程语言中都很常见的原语。在 CUDA 中，同步可以在两个级别执行（书中）：

• 系统级：等待主机和设备完成所有的工作。
• 块级：在设备执行过程中等待一个线程块中所有线程到达同一点。

对于主机，由于许多 CUDA API 调用和所有的内核启动是异步的，可以使用：

cudaError_t cudaDeviceSynchronize(void);

阻塞主机应用程序，直到所有的 CUDA 操作（复制、核函数等）完成。

当 __syncthreads() 被调用时，在同一个线程块中每个线程都必须等待直至该线程块中所有其他线程都已经达到这个同步点。在栅栏之前所有线程产生的所有全局内存和共享内存访问，将会在栅栏后对线程块中所有其他的线程可见。该函数可以协调同一个块中线程之间的通信，但它强制线程束空闲，从而可能对性能产生负面影响。

竞争条件：线程块中的线程可以通过共享内存和寄存器来共享数据。当线程之间共享数据时，要避免竞争条件（race condition）。例如，写后读竞争：当一个位置的无序读发生在写操作之后时，读应该在写前还是写后加载值是未定义的。其他例子还有读后写、写后写。当线程块中的线程在逻辑上并行运行时，在物理上并不是所有的线程都可以在同一时间执行——如果线程 A 试图读取由线程 B（在不同线程束中）写的数据，必须使用适当的同步确定线程 B 已经写完，否则会出现竞争条件。在不同的块之间没有线程同步；块间同步唯一安全的方法是在每个内核执行结束时使用全局同步点：终止当前核函数，开始执行新的核函数。不同块中的线程不允许相互同步，因此 GPU 可以以任意顺序执行块，这使得 CUDA 程序在大规模并行 GPU 上是可扩展的。

---

3.7 可扩展性（书 3.2.7）

对于任何并行应用程序而言，可扩展性是一个理想的特性：为并行应用程序提供额外的硬件资源，相对于增加的资源，并行应用程序会产生加速。例如，若一个 CUDA 程序在两个 SM 中是可扩展的，则与在一个 SM 中运行相比，在两个 SM 中运行会使运行时间减半。可扩展性意味着增加的计算核心可以提高性能。串行代码本身是不可扩展的；并行代码有可扩展的潜能，但真正的可扩展性取决于算法设计和硬件特性。

透明可扩展性：能够在可变数量的计算核心上执行相同的应用程序代码的能力。CUDA 内核启动时，线程块分布在多个 SM 中；网格中的线程块以并行或连续或任意的顺序被执行，这种独立性使得 CUDA 程序在任意数量的计算核心间可以扩展。书中图 3-18 展示了可扩展性例子：左侧 GPU 有 2 个 SM，可同时执行 2 个块；右侧 GPU 有 4 个 SM，可同时执行 4 个块。不修改任何代码，应用程序可以在不同的 GPU 配置上运行，所需的执行时间根据可用的资源而改变。这里提到：可扩展性比效率更重要——一个可扩展但效率很低的系统可以通过简单添加硬件核心来处理更大的工作负载；一个效率很高但不可扩展的系统可能很快会达到可实现性能的上限。

---

四、并行性的表现（书 3.3 节）

书中通过矩阵求和核函数 sumMatrixOnGPU2D，使用不同的执行配置和 nvprof 指标，说明为什么有些网格/块的维数组合比其他的更好，并建立网格和块的启发式。

4.1 用 nvprof 检测活跃的线程束与性能基准（书 3.3.1）

对块大小为 (32,32)、(32,16)、(16,32)、(16,16)** 的配置进行测试（Tesla M2070 上）。从结果中可以观察到（书中）：

• 第二种情况（如 32×16）块数比第一种多，设备可以有更多活跃的线程束，可能有更高的可实现占用率和更好的性能。
• 第四种情况（16×16）有最高的可实现占用率，但它不是最快的——因此，更高的占用率并不一定意味着有更高的性能，肯定有其他因素限制 GPU 的性能。

4.2 用 nvprof 检测内存操作（书 3.3.2）

在 sumMatrix 内核（C[idx]=A[idx]+B[idx]）中有 3 个内存操作：两个加载、一个存储。可以用 gld_throughput（全局加载吞吐量）等指标检查内存读取效率。书中结果说明：最后两种情况（块最内层维度为 16）的加载效率是最前面两种情况的一半——这可以解释为什么最后两种配置下更高的加载吞吐量和可实现占用率没有产生较好的性能。尽管在最后两种情况下正在执行的加载数量（吞吐量）很多，但那些加载的有效性（效率）是较低的。对网格和块启发式算法来说，最内层的维数应该总是线程束大小的倍数（32）；第 4 章将讨论半个线程束大小的线程块是如何影响性能的。

4.3 增大并行性（书 3.3.3）

在建立性能基准后，书中通过测试更大范围的线程配置（如块 (64,2)、(64,4)、(128,2)、(256,2)、(256,8) 等）得到结论：

• 块大小为 (256,8) 时，一个块中线程总数超过 1024（GPU 硬件限制），配置无效。
• 最好的结果可能是块 (128,2) 等；块 (64,2) 虽然启动的线程块最多，但不是最快的。
• 线程块最内层维度的大小对性能起着关键的作用（重复了前一节的结论）；增大并行性仍然是性能优化的一个重要因素。

用 nvprof 检测 achieved_occupancy 可验证：最好的执行配置既不具有最高的可实现占用率，也不具有最高的加载吞吐量。

书中在 Tesla M2070 上对 sumMatrixOnGPU2D 做了多组测试。下表为书中风格的执行配置与性能对照表（数值为示例，实际以你本机 nvprof 为准），用于体会「块最内层维度」与「占用率、加载效率」的关系：

块大小 (block.x, block.y)	块数量（相对）	核函数时间（示例）	achieved_occupancy（示例）	gld_efficiency（示例）
(32, 32)	基准	较慢	较低	高
(32, 16)	2×	较快（约 2×）	较高	高
(16, 32)	2×	较快	较高	明显下降
(16, 16)	4×	介于中间	最高	明显下降

可以看到：(16, 16) 时块数最多、占用率最高，但加载效率因最内层维度为 16（线程束大小的一半）而变差，整体并非最快。(32, 16) 在书中实验中往往更优。这直接说明：最内层维度应为线程束大小（32）的倍数；单看占用率或单看吞吐量都不足以优化，需在多个指标间找平衡。

从这些实验中可以推断出（书中）：

• 没有一个单独的指标能直接优化性能。
• 需要在几个相关的指标间寻找一个恰当的平衡来达到最佳的总体性能。
• 指标与性能：在大部分情况下，一个单独的指标不能产生最佳的性能；与总体性能最直接相关的指标或事件取决于内核代码的本质；要在相关的指标与事件之间寻求一个好的平衡；从不同角度查看内核以寻找相关指标间的平衡。
• 网格/块启发式算法为性能调节提供了一个很好的起点。

学习思考：这一节的实验是全书第一次用多指标解释性能——占用率、加载吞吐量、加载效率要一起看。最内层维度 32 的倍数能保证「每个 warp 的线程在访问全局内存时更易合并」，第 4 章会从内存访问模式上严格说明。建议自己跑一遍 nvprof，对照 achieved_occupancy 和 gld_efficiency 体会「平衡」的含义。

---

五、避免分支分化（书 3.4 节）

书中以并行归约为例，介绍避免分支分化的基本技术。

5.1 并行归约问题（书 3.4.1）

假设要对一个有 N 个元素的整数数组求和。串行代码很容易实现；若有大量数据，则可以通过并行计算快速求和。鉴于加法的结合律和交换律，数组元素可以以任何顺序求和。常用方法（书中）：

1. 将输入向量划分到更小的数据块中。
2. 用一个线程计算一个数据块的部分和。
3. 对每个数据块的部分和再求和得出最终结果。

迭代成对实现：一个数据块只包含一对元素，一个线程对这两个元素求和产生一个局部结果；这些局部结果在最初的输入向量中就地保存，作为下一次迭代的输入。每次迭代后输出元素数量减半，当长度为 1 时得到最终和。根据每次迭代后输出元素就地存储的位置，可分为（书中）：

• 相邻配对：元素与它们直接相邻的元素配对（步长 1、2、4…）。
• 交错配对：根据给定的跨度配对元素（初始跨度块大小的一半，每次减半）。

任何满足交换律和结合律的运算都可以代替加法（如 max、min、乘积等）。在向量中执行这类运算称为归约问题；并行归约是这种运算的并行执行，是许多并行算法中的关键运算。书中实现了多个不同的并行归约核函数，并测试不同实现对内核性能的影响。

5.2 并行归约中的分化（书 3.4.2）

相邻配对方法的内核实现流程（原书图 3-21）如下：每个线程将相邻的两个元素相加产生部分和。内核中有两个全局内存数组：一个大数组存放整个数组用于归约，一个小数组存放每个线程块的部分和。每个线程块在数组的一部分上独立执行；循环中迭代一次执行一个归约步骤，归约是就地完成的。__syncthreads() 保证线程块中的任一线程在进入下一次迭代之前，当前迭代里每个线程的所有部分和都已写入全局内存。跨度初始为 1，每次归约循环结束后乘以 2。因为线程块间无法同步，每个线程块产生的部分和需要复制回主机并在主机上做串行求和（原书图 3-22）。完整源代码可从 Wrox 网站 reduceInteger.cu 获取。

分化来源：条件表达式 if ((tid % (2 * stride)) == 0) 只对偶数 ID 的线程为 true，会导致很高的线程束分化。在第一次迭代中只有 ID 为偶数的线程执行条件语句主体，但所有线程都必须被调度；第二次迭代中只有四分之一的线程活跃，但所有线程仍被调度。

难点：tid % (2*stride)==0 在同一 warp 内会交替为 true/false（例如 stride=1 时 tid=0,2,4,… 为 true），所以同一 warp 内一半线程走 if、一半闲置，然后反过来，造成明显的串行化。改进方向是让「参与计算的线程」在每轮中连续（如 tid=0～stride-1），这样同一 warp 内要么都参与要么都不参与。

相邻配对归约核函数（reduceNeighbored）的典型实现如下。注意：stride 从 1 开始，每轮翻倍；只有满足 (tid % (2*stride)) == 0 的线程参与加法，因此存在严重分化。

__global__ void reduceNeighbored(int *g_idata, int *g_odata, unsigned int n) {
    unsigned int tid = threadIdx.x;
    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;

    int *idata = g_idata + blockIdx.x * blockDim.x;

    if (idx >= n) return;

    // 就地归约，stride 1, 2, 4, ...
    for (int stride = 1; stride < blockDim.x; stride *= 2) {
        if ((tid % (2 * stride)) == 0) {
            idata[tid] += idata[tid + stride];
        }
        __syncthreads();
    }

    if (tid == 0)
        g_odata[blockIdx.x] = idata[0];
}

5.3 改善并行归约的分化（书 3.4.3）

通过重新组织每个线程的数组索引，让 ID 相邻的线程执行求和操作，线程束分化就能被减少。书中图 3-23 展示了这种实现：部分和的存储位置并没有改变，但工作线程的索引已更新，使同一线程束内线程走相同路径。修改后的内核（书中称为 reduceNeighboredLess）中，参与加法的线程在每轮中连续排列，从而同一 warp 内要么都参与要么都不参与，避免分化。核心改动：用 tid 直接作为写入下标，且只让「前一半」线程做加法：

__global__ void reduceNeighboredLess(int *g_idata, int *g_odata, unsigned int n) {
    unsigned int tid = threadIdx.x;
    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int *idata = g_idata + blockIdx.x * blockDim.x;

    if (idx >= n) return;

    for (int stride = 1; stride < blockDim.x; stride *= 2) {
        int index = 2 * stride * tid;
        if (index < blockDim.x) {
            idata[index] += idata[index + stride];
        }
        __syncthreads();
    }
    if (tid == 0)
        g_odata[blockIdx.x] = idata[0];
}

5.4 交错配对的归约（书 3.4.4）

与相邻配对相比，交错配对颠倒了元素的跨度：初始跨度是线程块大小的一半，然后在每次迭代中减半（原书图 3-24）。在每次循环中，每个线程对两个被当前跨度隔开的元素求和。交错配对归约核函数（reduceInterleaved）如下。注意：stride 从 blockDim.x / 2 开始，每轮减半；写入位置为 tid，同一 warp 内线程仍连续工作，分化与 reduceNeighboredLess 类似，但全局内存访问模式更利于合并（第 4 章）。

__global__ void reduceInterleaved(int *g_idata, int *g_odata, unsigned int n) {
    unsigned int tid = threadIdx.x;
    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int *idata = g_idata + blockIdx.x * blockDim.x;

    if (idx >= n) return;

    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride) {
            idata[tid] += idata[tid + stride];
        }
        __syncthreads();
    }
    if (tid == 0)
        g_odata[blockIdx.x] = idata[0];
}

书中实验表明：交错实现比第一个（reduceNeighbored）实现快了约 1.69 倍，比第二个（reduceNeighboredLess）实现快了约 1.34 倍。这种性能提升主要是由 reduceInterleaved 函数里的全局内存加载/存储模式导致的——第 4 章会介绍全局内存加载/存储模式对内核性能的影响。

理解与体会：归约这一节把「分支分化」和「内存访问模式」绑在一起：先通过索引重组消除分化（reduceNeighboredLess），再通过交错配对改善访存（reduceInterleaved）。性能优化往往是多步叠加的，先保证正确性和可读性，再逐步加「少分化 → 好访存 → 展开 → 共享内存」等优化。

---

六、展开循环与动态并行（书 3.5、3.6 节）

6.1 展开循环（书 3.5 节）

循环展开通过减少循环迭代次数和分支频率，增加可被调度器利用的独立指令，从而更好地隐藏指令或内存延迟。书中在归约等内核中应用了展开技术：例如每个线程在一次循环中处理多个元素（如 4 个），等价于展开；当剩余活动线程数等于一个 warp（32）时，可以完全展开最后一阶段，并去掉不必要的 __syncthreads()，减少同步开销。注意：展开会提高每线程寄存器使用量，需在占用率与指令级并行之间取得平衡；过度展开可能导致寄存器溢出或块内线程数下降，反而变慢。第 5 章还会结合共享内存进一步优化归约。

6.2 动态并行（书 3.6 节）

这里提到：到目前为止，所有核函数都是从主机线程中调用的，GPU 的工作负载完全在 CPU 的控制下。CUDA 的动态并行允许在 GPU 端直接创建和同步新的 GPU 内核——在一个核函数中的任意点动态增加 GPU 应用程序的并行性。这样就不需要把算法设计为单独的、大规模数据并行的内核启动；动态并行提供了一个更有层次结构的方法，使并发性可以在一个 GPU 内核的多个级别中表现出来。使用动态并行可以让递归算法更加清晰易懂；可以推迟到运行时决定需要在 GPU 上创建多少个块和网格；可以动态地利用 GPU 硬件调度器和负载平衡器，适应数据驱动或工作负载。在 GPU 端直接创建工作的能力可以减少主机和设备之间传输执行控制和数据的需求。书中通过使用动态并行实现递归归约核函数的例子，对如何利用动态并行做了基本介绍。实现有效的嵌套内核时，必须注意设备运行时的使用，包括子网格启动策略、父子同步和嵌套层的深度。计算能力 3.5 及以上的设备才支持动态并行。

---

七、总结与重难点小结（书 3.7 节）

书中总结：本章从硬件的角度分析了内核执行。在 GPU 设备上，CUDA 执行模型有两个最显著的特性：(1) 在线程块与线程中分配了硬件资源；(2) 使用 SIMT 方式在线程束中执行线程。这些执行模型的特征使得我们在提高并行性和性能时，能控制应用程序是如何让指令和内存带宽饱和的。不同计算能力的 GPU 设备有不同的硬件限制，因此网格和线程块的启发式算法在为不同平台优化内核性能方面发挥了非常重要的作用。动态并行使设备能够直接创建新的工作，确保我们可以用一种更自然和更易于理解的方式来表达递归或依赖数据并行的算法。本章也介绍了使用命令行分析工具 nvprof 详细分析内核性能的方法；因为一个单纯的内核实现可能不会产生很好的性能，所以配置文件驱动的方法在 CUDA 编程中尤其重要。性能分析对内核行为提供了详细的分析，并能找到产生最佳性能的主要因素。第 4 章和第 5 章将从 CUDA 内存模型的角度继续介绍内核执行的内容。

7.1 核心要点回顾

• 线程束是 SM 的基本执行单位，32 个连续线程一组；块的最内层维度应为 32 的倍数。
• 线程束分化会串行化同一 warp 内的分支执行；可通过数据分区（如 (tid/warpSize)%2）或重组索引减少分化；短分支可能被编译器用断定优化。
• 资源分配：寄存器和共享内存限制每 SM 可容纳的块和线程束；活跃线程束分为选定、符合条件的、阻塞的；线程束间切换几乎无开销。
• 延迟隐藏：算术延迟约 10～20 周期，全局内存约 400～800 周期；Fermi 上约需 36 个线程束/SM 才能较好隐藏内存延迟；需在延迟隐藏与资源利用之间平衡。
• 占用率 = 活跃线程束数 / 最大线程束数；高占用率不是唯一目标；准则包括块大小为 32 的倍数、至少 128～256 线程/块、块数远多于 SM 数。
• 同步：__syncthreads() 做块内同步，需避免分支内不同步导致死锁；块间无同步，只能通过结束内核再启动新内核。
• 可扩展性：块独立使同一代码可在不同 SM 数的 GPU 上扩展；可扩展性比单机效率更重要。
• 并行归约：相邻配对与交错配对、分化与访存模式共同决定性能；交错配对 + 少分化可带来明显加速。
• 动态并行：计算能力 3.5+，内核内启动子网格，适合递归与数据依赖型算法。

7.2 重难点自检与学习思考

主题	自检问题	要点
32 与 warp	为什么说「32 是神奇的数字」？块最内层维度为何要是 32 的倍数？	硬件按 32 线程一组调度；非 32 倍数会导致未填满的 warp、降低加载效率。
分支分化	分化发生在什么粒度？如何避免？为何 nvprof 有时显示 100% 分支效率？	同一 warp 内；数据分区或索引重组使同 warp 同路径；短分支可能被断定替代。
占用率 vs 性能	高占用率一定更快吗？矩阵求和 (16,16) 说明了什么？	不一定；要同时看 gld_efficiency 等；最内层 32 的倍数常比单纯高占用率更重要。
延迟隐藏	算术延迟与内存延迟量级？Fermi 上隐藏内存延迟约需多少活跃 warp/SM？	约 10～20 vs 400～800 周期；约 36 个线程束/SM（书中推导）。
网格与块启发式	块大小、最内层维度、块数量分别怎么选？	32 的倍数、至少 128～256 线程/块、块数远多于 SM 数；结合实验与 -Xptxas -v 调优。
归约优化	reduceNeighbored 的分化来自哪里？reduceNeighboredLess / reduceInterleaved 各解决什么？	tid % (2*stride)==0 导致同 warp 内交替；前者消除分化，后者进一步改善访存合并。

学习思考：第 2 章解决「怎么写」——网格、块、核函数；本章解决「为什么这样配置更好」——从 warp、SM、占用率、延迟隐藏到分支分化和归约改进。理解执行模型后，看 nvprof 的 achieved_occupancy、gld_efficiency 不再盲目，而是能对应到「活跃 warp 数量」「合并访问」等硬件行为。下一章从内存访问模式继续深入，解释「为何最内层 32 的倍数能提高加载效率」。

7.3 与前后章的联系

• 第 2 章：编程模型给出网格—块—线程与核函数写法；本章解释这些线程在硬件上如何被调度（warp、SM）、为何不同配置性能不同。
• 第 4 章：全局内存访问模式（对齐与合并）——矩阵求和中「最内层维度 32」对 gld_efficiency 的影响将在第 4 章从访存角度严格说明。
• 第 5 章：共享内存与归约的进一步优化（bank 冲突、层次化归约等），与本章的循环展开、减少分化形成完整优化链。

---

下一章预告

在下一篇博客中，我们将进入第 4 章：全局内存，从「执行与资源」转向「数据放在哪、怎么访问」：

• CUDA 内存模型与全局内存、缓存
• 内存管理：固定内存、零拷贝、统一寻址（Unified Memory）
• 内存访问模式：对齐与合并、对带宽的影响
• 矩阵转置等案例与带宽测量

理解全局内存的访问模式，是突破带宽瓶颈的第一步。

---

本章自测

1. 为什么块的最内层维度建议为 32 的倍数？与线程束有何关系？
2. 什么是线程束分化？减少分化的常见做法是什么？
3. 占用率与性能的关系如何？为何「高占用率」不一定是「最快配置」？

答案与解析

1. GPU 按线程束（32 线程）为单位调度；最内层维度为 32 的倍数时，同一 warp 内线程连续，利于合并访问与满载执行；否则会产生未填满的 warp，降低内存与计算利用率。
2. 同一 warp 内部分线程走 if、部分走 else，导致串行执行两条路径，称为分化。常见做法：数据分区使同 warp 内线程走相同分支、或用索引重组（如交错归约）避免同一 warp 内条件不一致。
3. 占用率是「活跃线程束数/最大线程束数」，高占用率有助于隐藏延迟，但性能还受内存带宽、合并访问等影响；若内存效率低，再高占用率也受限于带宽。矩阵求和中 (32,16) 比 (16,16) 更快说明需综合看 gld_efficiency 与占用率。

---

系列导航：导读 | 上一篇：第2章 CUDA编程模型 | 下一篇：第4章 全局内存

---

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