PMPP 第四章 计算架构和调度

01 Jun, 2026

计算架构和调度

4.1 Architecture of a modern GPU

第四章开始进入 GPU 计算架构与调度机制。前几章主要讲 CUDA C++ 如何写 kernel、如何组织 thread/block/grid；第四章开始解释这些程序在 GPU 硬件上到底是怎么执行的。4.1 节主要介绍现代 CUDA-capable GPU 的整体计算架构。

1. GPU 的基本组织：由多个 SM 组成

现代 GPU 可以看成是由大量 Streaming Multiprocessors，SM，流式多处理器 组成的阵列。

SM 是 CUDA 程序执行的核心硬件单元。一个 kernel 启动后，线程块最终会被分配到不同的 SM 上执行。每个 SM 内部又包含多个执行单元，书中称为 streaming processors，也就是通常容易被叫作 CUDA cores 的部分。

不过作者特意说明，为了避免和 CPU core 混淆，书中更倾向使用 streaming processor 这个说法。因为 GPU 的这些执行单元并不等价于 CPU 那种复杂、独立的核心。

2. SM 内部包含控制逻辑、计算单元和片上存储资源

从图 4.1 的高层视角看，每个 SM 主要包含三类资源：

一是 Control，也就是控制逻辑，负责指令调度、执行控制等。

二是很多小的计算单元，也就是 streaming processors，用于执行大量线程中的指令。

三是 Memory，这里指 SM 内部的片上存储资源。书中没有在 4.1 展开讲，因为这些片上存储结构会放到第 5 章详细讨论，例如 shared memory、cache 等。

所以可以先记成：

SM = 控制逻辑 + 多个计算单元 + 片上存储资源。

这也是后面理解 warp 调度、occupancy、shared memory 的基础。

3. GPU 还包含全局显存 Global Memory

除了 SM 内部的片上存储资源，GPU 还有容量更大的片外设备内存，也就是 Global Memory。

图 4.1 中，Global Memory 位于 GPU 架构图的底部，通过多个 Memory Controller / Memory Channel 与 SM 连接。这些 memory controller 共同提供高带宽的数据访问能力。

这里的 Global Memory 对应我们 CUDA 编程中经常使用的 cudaMalloc 分配出来的显存。它容量大，但访问延迟比 SM 内部的片上存储高很多。因此后面第 5、6 章会重点讲内存层次结构和访问优化。

4. GPU DRAM 类型：GDDR 与 HBM

书中提到，较老的 GPU 使用的是 GDDR SDRAM，也就是 Graphics Double Data Rate SDRAM。

从 NVIDIA Pascal 架构开始，一些 GPU 开始使用 HBM / HBM2 / HBM2E / HBM3，也就是 High-Bandwidth Memory。HBM 的特点是 DRAM 模块和 GPU 被更紧密地集成在同一个封装中，因此能够提供更高的带宽。

不过作者为了简化后续讨论，后面会统一把这些片外显存都称为 DRAM。

这里需要注意： Global Memory 是从 CUDA 编程模型角度说的地址空间；DRAM 是从硬件实现角度说的物理存储介质。

5. Hopper 架构引入 GPC 层级

从 Hopper 架构开始，GPU 的 SM 会进一步被组织成 GPU Processing Clusters，GPCs。

例如书中给出的 Hopper H100 GPU 有：

• 132 个 SM；
• 每个 SM 有 128 个 streaming processors；
• 总共 132 × 128 = 16896 个 streaming processors；
• 这些 SM 被组织成 8 个 GPC。

因此，现代 GPU 的层次可以粗略理解为：

GPU → GPC → SM → streaming processors

不过在一般 CUDA 编程中，我们最常直接接触的是 grid、block、thread 这些软件层次，以及 SM 这个硬件执行层次。GPC 这个层次在后面 thread block cluster、distributed shared memory 等机制中会更重要。

6. 一句话总结

第四章 4.1 节主要说明：现代 CUDA GPU 由多个 SM 组成，SM 是线程块执行的核心硬件单元；SM 内部包含计算单元、控制逻辑和片上存储资源，SM 通过内存控制器访问片外 Global Memory；从 Hopper 开始，SM 还会进一步组织成 GPC。

4.2 Thread block scheduling

1. 为什么需要 thread block scheduling？

当一个 kernel 被调用时，CUDA runtime 会启动一个由大量 threads 组成的 grid。虽然 GPU 有很多执行资源，但一个 grid 中的线程数量通常远远超过 GPU 能同时执行的线程数量。

因此，GPU 不可能一次性把所有线程都放到硬件上执行，而是需要一种调度机制，把这些线程逐批分配到可用的硬件资源上。

这就是 thread block scheduling，线程块调度 要解决的问题。

可以简单理解为：

kernel 启动的是整个 grid，但 GPU 实际执行时，是把 grid 中的 blocks 分批调度到 SM 上。

2. CUDA 以 block 为单位分配线程到 SM

CUDA 中，线程不是一个一个单独分配到 SM 的，而是以 thread block 为单位分配。

也就是说：

一个 block 中的所有 threads 会被同时分配到同一个 SM 上执行。

这是 4.2 节最核心的一句话。

它意味着，同一个 block 内的线程具有特殊关系：它们在同一个 SM 上执行，因此可以共享 SM 内部的一些资源，比如 shared memory，也可以进行 block-level synchronization，例如 __syncthreads()。

3. 一个 SM 可以同时运行多个 block

一个 SM 并不是一次只能执行一个 block。只要资源允许，一个 SM 可以同时驻留多个 block。

图 4.2 展示的就是 blocks 被分配到 SM 上的过程。图中每个 SM 同时被分配了多个 blocks。这里要注意，能同时分配多少个 block，不是随便决定的，而是受 SM 资源限制。

影响一个 SM 同时能驻留多少 block 的因素包括：

• 每个 block 有多少 threads；
• 每个 thread 使用多少 registers；
• 每个 block 使用多少 shared memory；
• SM 支持的最大 block 数；
• SM 支持的最大 thread 数。

这些内容后面会在 occupancy 和 resource partitioning 中详细展开。

4. grid 中的 blocks 通常分批执行

因为 SM 数量有限，并且每个 SM 同时能容纳的 block 数量也有限，所以整个 GPU 同一时刻能执行的 block 总数是有限的。

但是，一个 grid 往往包含比这个数量更多的 blocks。

因此 CUDA runtime 会维护一个待执行 block 列表：

当某些已经分配到 SM 的 blocks 执行完成后，runtime 再把新的 blocks 分配到空出来的 SM 上。

这就类似一个任务队列。SM 有空位时，从队列中取新的 block 执行。

所以 grid 的执行不是“所有 block 同时执行”，而是：

一部分 block 先执行，执行完后释放资源，然后其他 block 再被调度执行。

5. 同一个 block 内线程的特殊保证

CUDA 保证：

同一个 block 中的所有 threads 会被分配到同一个 SM，并且同时被分配资源。

这个保证非常重要。正因为如此，同一个 block 内的 threads 才可以进行一些跨线程协作。

主要包括两类：

第一，block-wide barrier synchronization。 也就是后面 4.3 节要讲的 __syncthreads()。它要求同一个 block 中的线程在某个位置等待，直到所有线程都到达该同步点。

第二，shared memory 数据交换。 Shared memory 位于 SM 内部，是一种低延迟 SRAM。因为同一个 block 的线程都在同一个 SM 上，所以它们可以通过 shared memory 交换数据。Shared memory 会在第 5 章重点讨论。

可以记成：

block 内线程能协作，是因为它们被放在同一个 SM 上；不同 block 之间一般不能直接这样协作，是因为它们可能被调度到不同 SM，甚至在不同时间执行。

6. 不同 block 之间没有默认同步关系

4.2 节暗含了一个非常重要的 CUDA 编程原则：

不同 blocks 之间不能假设执行顺序，也不能假设它们同时执行。

原因是 blocks 的调度顺序由 CUDA runtime 和硬件决定。一个 block 可能先执行，也可能后执行；两个 block 可能同时在不同 SM 上执行，也可能分批执行。

所以在普通 kernel 中，不能写出依赖 block 执行顺序的代码。例如不能假设 block 0 一定比 block 1 先完成，也不能在不同 block 之间直接用普通方式做同步。

这也是为什么 CUDA 早期的同步主要限制在 block 内部。

7. Hopper 架构中的 thread block cluster

从 Hopper 架构开始，CUDA 引入了一个可选的新层次：

thread block cluster，线程块簇

它是一组 thread blocks。

类似于同一个 block 中的 threads 会被调度到同一个 SM，同一个 thread block cluster 中的 blocks 会被共同调度到同一个 GPC（GPU Processing Cluster） 上。

这带来了两点能力：

第一，同一个 cluster 内的 threads 可以通过相应 API 进行同步。 第二，同一个 cluster 内的 blocks 可以访问 distributed shared memory。

Distributed shared memory 可以理解为：由参与同一个 cluster 的多个 blocks 的 shared memory 共同组成的一种分布式共享存储。具体内容后面第 5 章会展开。

暂时可以先记住：

普通 block 的协作范围是 SM 内部；thread block cluster 把协作范围扩展到了 GPC 内的一组 blocks。

9. 需要记住的术语

10. 一句话总结

CUDA 以 thread block 为单位把 threads 调度到 SM 上；同一个 block 的 threads 会同时驻留在同一个 SM，因此可以进行块内同步和 shared memory 数据交换，而不同 blocks 的执行顺序和并发关系由 runtime 决定，程序通常不能依赖它们之间的同步关系。

4.3 Synchronization and transparent scalability

1. CUDA 中的块内同步：__syncthreads()

CUDA 允许同一个 thread block 内的线程通过 block-wide barrier synchronization 来协调执行。最典型的同步函数是：

__syncthreads();

当一个线程执行到 __syncthreads() 时，它不会立刻继续往下执行，而是会停在这个位置等待。只有当同一个 block 中的所有线程都到达这个位置之后，这些线程才会一起越过这个同步点，继续执行后面的代码。

可以把 __syncthreads() 理解成一个“阶段分界线”：

// 阶段 1：所有线程各自完成一部分工作
...

__syncthreads();

// 阶段 2：只有当所有线程都完成阶段 1 后，才可以继续
...

它的作用是保证：

block 内所有线程都完成前一阶段之后，任何线程才能进入下一阶段。

这在使用 shared memory 时尤其重要。例如一个 block 内的线程先把数据写入 shared memory，然后再让其他线程读取这些数据，中间通常就需要 __syncthreads()。

__syncthreads() 是 CUDA 中的 intrinsic function。Intrinsic function 从代码形式上看像普通函数，但编译器会以特殊方式处理，通常会直接翻译成对应的硬件指令，而不是生成普通函数调用。CUDA 中双下划线 __ 常用于表示这类内建函数。__syncthreads() 对应的是 GPU 的 barrier synchronization 机制，用于让同一个 block 内所有线程在同一位置同步。

2. barrier 的含义：先到的线程必须等待后到的线程

书中用图 4.3 解释了 barrier synchronization 的执行过程。不同线程到达同步点的时间可能不同，有的线程先到，有的线程后到。先到的线程会在 barrier 处等待，直到最后一个线程也到达。最后一个线程到达后，所有线程才会被释放，继续执行后面的代码。

所以 barrier 的核心语义是：

no one is left behind 没有线程会被落下。

也就是说，同一个同步范围内的线程必须一起完成某个阶段，然后一起进入下一个阶段。

3. __syncthreads() 必须被 block 内所有线程执行到

__syncthreads() 有一个非常重要的使用规则：

如果一个 block 中出现了 __syncthreads()，那么这个 block 内的所有线程都必须能够执行到同一个同步点。

否则就可能出现错误结果，甚至死锁。

比如下面这种写法是错误的：

if (threadIdx.x % 2 == 0) {
    ...
    __syncthreads();
} else {
    ...
    __syncthreads();
}

表面看起来两个分支里都有 __syncthreads()，但实际上它们是两个不同的 barrier synchronization points。偶数线程到达的是 then 分支里的同步点，奇数线程到达的是 else 分支里的同步点。于是每个同步点都等不到 block 内所有线程，程序行为就是未定义的，可能直接死锁。

正确理解是：

不是“每个线程都调用过一次 __syncthreads() 就行”，而是“每个线程必须到达同一个 __syncthreads() 位置”。

4. 同步范围 scope 很重要

Barrier synchronization 有一个 scope，同步范围。这个范围决定哪些线程必须一起到达同步点。

在 CUDA 中，常见的同步范围包括：

• block-wide synchronization：同一个 block 内的线程同步；
• cluster-wide synchronization：同一个 thread block cluster 内的线程同步；
• grid-wide synchronization：整个 grid 内的线程同步。

本节重点讨论的是 block-wide synchronization，也就是 __syncthreads()。

CUDA 之所以能安全支持 block-wide barrier，是因为 CUDA runtime 会保证同一个 block 的所有线程被同时分配到同一个 SM 上。也就是说，一个 block 只有在 SM 已经分配了这个 block 所需的全部资源之后，才会开始执行。这样可以保证 block 内所有线程最终都有机会到达 barrier，从而避免因为资源没有分配齐导致无限等待。

5. 为什么普通 CUDA 程序不鼓励 block 之间同步？

4.3 节的另一个核心概念是 transparent scalability，透明可扩展性。

CUDA 程序通常只允许 block 内同步，而不允许普通情况下任意 block 之间同步。这样做的好处是：

不同 blocks 之间没有同步约束，因此 CUDA runtime 可以用任意顺序调度 blocks。

这带来了很强的硬件适配能力。假设同一个 kernel 有 8 个 blocks：

在低端 GPU 上，可能一次只能执行 2 个 blocks，于是需要分 4 批执行； 在高端 GPU 上，可能一次可以执行 4 个 blocks，于是只需要分 2 批执行。

但代码不需要修改。因为程序没有依赖 block 之间的执行顺序，也没有要求所有 blocks 必须同时存在。图 4.5 展示的就是这种透明可扩展性：同一个 grid 可以在不同资源规模的 GPU 上运行，只是并发执行的 block 数量不同。

这就是 CUDA 的一个重要设计原则：

block 之间尽量相互独立，硬件才能自由调度，程序才能适配不同规模的 GPU。

6. wave：同时执行的一批 blocks

书中把一组同时执行的 thread blocks 称为 wave。

如果一个 grid 中的 block 总数大于 GPU 同时能执行的 block 数量，那么这些 blocks 就会被分成多个 waves 依次执行。

例如：

如果 grid 有 8 个 blocks，某个 GPU 一次只能执行 2 个 blocks，那么一共需要 4 个 waves； 如果另一个 GPU 一次能执行 4 个 blocks，那么只需要 2 个 waves。

这里可以理解为：

wave 数量 = grid 中 block 总数 / GPU 同时能执行的 block 数量。

当然实际情况可能不是整数，如果最后一批 blocks 不满，就会出现后面说的 tail effect。

7. wave 数量多和少各有适用场景

书中提到，waves 的数量不是越少越好，也不是越多越好，要看 block 之间的负载是否均衡。

如果不同 blocks 的执行时间差异较大，那么拥有较多 waves 是有好处的。因为当某些 SM 提前完成自己的 block 后，还可以继续拿新的 block 来执行，从而帮助硬件做负载均衡。

但如果每个 block 的负载非常均衡，那么较少的 waves 也可以接受，有时甚至是有利的。后面第 6 章讲 thread coarsening 时会进一步讨论这种情况。

可以简单记成：

block 负载不均衡时，多一些 waves 更利于负载均衡； block 负载均衡时，少一些 waves 未必是问题。

8. tail effect：最后一个 wave 不满导致硬件利用率下降

如果 grid 的 block 数量不能整除 GPU 同时可执行的 block 数量，最后一个 wave 就可能是不完整的。这个不完整的 wave 无法填满整个 GPU，导致硬件资源利用率下降，这种现象叫 tail effect。

例如书中给出的例子：一个 grid 有 660 个 blocks，GPU 同时可以执行 264 个 blocks。那么执行过程是：

第 1 个 wave：264 blocks
第 2 个 wave：264 blocks
第 3 个 wave：132 blocks

最后一个 wave 只有 132 个 blocks，只占满了一半资源，因此会出现 underutilization。

所以在性能优化时，如果 block 负载比较均衡，最好让 grid 中的 block 数量尽量是“GPU 同时可执行 block 数”的整数倍，从而减少 tail effect。

不过这里也要注意：GPU 同时可执行多少 blocks 取决于具体 GPU 的 SM 数量、每个 SM 能驻留多少 blocks，以及 kernel 的资源使用情况，因此不是一个固定常数。

9. cluster-wide 和 grid-wide synchronization

除了 __syncthreads() 这种 block-wide barrier，CUDA 还提供了更大范围的同步机制。

在 Hopper 架构中，thread block cluster 支持 cluster-wide synchronization。同一个 cluster 中的线程可以通过 Cooperative Groups API 进行同步，并确保 distributed shared memory 操作完成后再继续执行。

CUDA 也提供 grid-wide synchronization，即整个 grid 范围内的同步。但这种同步比 block-wide 和 cluster-wide 更重，而且有严格限制。因为 grid-wide barrier 要求所有参与同步的线程都必须真的同时执行并且能够到达 barrier，否则就会有死锁风险。

因此，在普通 CUDA 编程中，最常用、最安全、最基础的还是 block-wide synchronization。

总结

4.3 节主要讲 CUDA 中的同步机制，以及为什么 CUDA 通常把同步范围限制在 thread block 内。

__syncthreads() 是 block-wide barrier synchronization。它要求同一个 block 中的所有线程都到达同一个同步点之后，才能继续执行后面的代码。这个机制常用于划分并行程序的执行阶段，尤其适合配合 shared memory 使用。

使用 __syncthreads() 时必须非常小心：它不能被放在导致 block 内线程分流到不同同步点的条件分支中。否则部分线程会等待永远不会到达的其他线程，造成未定义行为或死锁。

CUDA 把普通同步限制在 block 范围内，是为了保持不同 blocks 之间的独立性。这样 runtime 可以按照任意顺序调度 blocks，使同一份 CUDA 程序能够在低端 GPU 和高端 GPU 上不修改代码地运行，这就是 transparent scalability。grid 中同时执行的一批 blocks 称为 wave。如果最后一个 wave 不完整，就可能产生 tail effect，导致硬件资源利用率下降。

4.4 Warps and SIMD hardware

1. 从 block 调度进入 thread 调度

前面 4.2、4.3 主要讲的是 thread block 如何被调度到 SM 上，以及为什么 block 内可以同步、block 之间不能随便同步。

4.4 节开始进一步往下看：一个 block 被分配到 SM 之后，block 内部的 threads 是怎么被硬件执行的？

从 CUDA 编程模型上看，thread 是独立的执行单元。但从 GPU 硬件实现上看，SM 并不是一个线程一个线程地单独调度，而是会把 threads 进一步划分成更小的执行单位：warp。

2. Warp 是 SM 中的线程调度单位

在目前大多数 CUDA GPU 实现中，一个 thread block 被分配到 SM 后，会被进一步划分为若干个 warp。

一个 warp 通常包含 32 个 threads。

也就是说：

block 是 CUDA runtime 分配到 SM 的基本单位； warp 是 SM 内部进行 thread scheduling 的基本单位。

例如，一个 block 有 256 个 threads，而 warp size 是 32，那么这个 block 会被划分成：

256 / 32 = 8 warps

如果一个 SM 同时驻留 3 个这样的 blocks，那么这个 SM 上一共有：

3 × 8 = 24 warps

图 4.6 展示的就是 block 被划分成 warps 后驻留在 SM 中的情况。

3. Warp 的划分规则：按 threadIdx 线性连续划分

对于一维 block，warp 的划分很直接：

warp 0: thread 0  ~ thread 31
warp 1: thread 32 ~ thread 63
warp 2: thread 64 ~ thread 95
...

一般来说：

warp n: thread 32n ~ thread 32n + 31

如果 block 的 thread 数不是 32 的整数倍，最后一个 warp 会被补齐到 32 个位置，但多出来的是 inactive threads，也就是不参与实际执行的无效线程。

例如 block 只有 48 个 threads，那么会被划分成 2 个 warps：

warp 0: 32 个有效线程
warp 1: 16 个有效线程 + 16 个 inactive threads

这里的重点是：硬件仍然以 warp 为单位调度，哪怕最后一个 warp 不满 32 个有效线程。

4. 多维 block 会先线性化，再划分 warp

如果 block 是二维或三维的，CUDA 并不是直接按二维/三维形状划分 warp，而是会先把多维 thread layout 映射成一维线性顺序，然后再按每 32 个 threads 划分 warp。

对于二维 block，线性化方式类似 C/C++ 中二维数组的 row-major layout：

linear_tid = threadIdx.y * blockDim.x + threadIdx.x;

也就是说，先排列 threadIdx.y = 0 的所有线程，并且 threadIdx.x 从小到大；然后排列 threadIdx.y = 1 的所有线程；依次类推。

图 4.7 展示了一个二维 block 如何被线性化。上半部分是二维 thread 组织方式，下半部分是线性排列方式。比如 (x, y) 坐标的线程会按 y 行展开，每一行内部 x 递增。

对于三维 block，可以理解为：

linear_tid = threadIdx.z * blockDim.y * blockDim.x
           + threadIdx.y * blockDim.x
           + threadIdx.x;

即先展开 z = 0 的二维平面，再展开 z = 1 的二维平面，依次类推。

这里容易和前面第三章讲的 multidimensional grid/block 混淆。记住一个原则：

warp 划分看的是 thread 的线性编号，不是你视觉上画出来的二维/三维形状。

5. SIMD：同一个 warp 中的线程执行同一条指令

SM 执行 warp 时采用的是 SIMD，Single Instruction, Multiple Data 思想。

意思是：

同一时刻，硬件取出一条指令，然后让 warp 中的多个 threads 对各自的数据执行这条相同指令。

例如一个 warp 中的所有线程都执行：

C[i] = A[i] + B[i];

它们执行的是同一条加法相关指令，但每个线程的 i 不同，访问的数据也不同。

这就是 Single Instruction, Multiple Data：

Single Instruction: 同一条指令
Multiple Data: 不同线程处理不同数据

6. SM 内部有 processing blocks

图 4.8 展示了 SM 内部进一步被组织成多个 processing blocks。每个 processing block 有自己的 instruction fetch/dispatch 单元，并控制一组 streaming processors。

书中用一个简化例子说明：一个 SM 有两个 processing blocks，每个 processing block 有 8 个 streaming processors，共享一个 instruction fetch/dispatch 单元。

真实 GPU 中规模更大。例如 Hopper H100 的一个 SM 有 128 个 streaming processors，被组织成 4 个 processing blocks，每个 processing block 有 32 个 streaming processors。

这个设计的核心是：

多个执行单元共享同一套控制逻辑，从而减少控制硬件开销，把更多芯片面积用于计算吞吐。

这也是 GPU 和 CPU 的一个重要区别。CPU 更重视单线程低延迟，所以控制逻辑复杂；GPU 更重视吞吐量，所以让大量执行单元共享控制逻辑。

7. SIMT：CUDA 程序员看到的是线程，硬件执行的是 warp

书中提到，CUDA 的执行行为常被称为 SIMT，Single Instruction, Multiple Thread。

SIMT 和 SIMD 很接近，但侧重点不同：

• SIMD 是硬件执行模型：一条指令控制多个数据通路；
• SIMT 是 CUDA 的编程/执行抽象：程序员写的是独立 scalar threads，但硬件会把这些 threads 组成 warp 来执行。

也就是说，CUDA 程序员不需要像传统 CPU SIMD 编程那样手动写向量指令，例如显式使用 AVX/SSE intrinsic。你只需要写普通的 scalar thread 代码：

int i = blockIdx.x * blockDim.x + threadIdx.x;
C[i] = A[i] + B[i];

每个 thread 看起来都是独立执行的，但硬件会自动把相邻 threads 组成 warp，并用 SIMD 风格的硬件执行它们。

这就是 SIMT 的好处：

程序员写的是“每个线程做什么”； 硬件负责把多个线程打包成 warp，高效利用 SIMD 硬件。

8. Warp size 目前通常是 32，但本质上是硬件实现细节

书中强调，warp size 是硬件实现相关的。到目前为止，CUDA 设备基本都使用 32 threads per warp，但理论上未来架构可以改变这个值。

所以代码中如果需要查询 warp size，可以使用设备属性，而不是完全硬编码。不过在实际 CUDA 编程和性能分析中，32 仍然是非常常用的默认值。

9. Warp-level primitives：warp 级别的协作接口

因为同一个 warp 中的线程有特殊的调度关系，CUDA 提供了一些 warp-level primitives，用于 warp 内线程之间进行高效的数据交换和同步。

例如后面章节会讲到的：

__shfl_sync()
__ballot_sync()
__syncwarp()

这些接口允许程序员把 warp 当成一个比 block 更细粒度的协作单位。

传统层次是：

grid → block → thread

理解 warp 后，可以在性能优化时加入一个新的层次：

grid → block → warp → thread

这种 warp-centric programming 在高性能 CUDA kernel 中非常重要，比如 reduction、scan、matrix multiplication、softmax 等算子里都会经常见到。

Warp 和 SIMD 硬件

这部分主要解释：为什么 GPU 要用 warp 来执行线程。核心原因是为了让多个线程共享一套控制逻辑，提高硬件吞吐量。

冯·诺依曼模型的基本执行方式

普通计算机执行程序时，大致流程是：

程序和数据放在 Memory 中； Control Unit 根据 PC 取出下一条指令； 指令进入 IR； Processing Unit 根据指令执行计算。

其中：

• PC：记录下一条指令的地址；
• IR：保存当前正在执行的指令；
• Control Unit：负责取指令、解析指令、发控制信号；
• Processing Unit：负责真正执行计算。

SIMD 的思想

SIMD 的意思是：

Single Instruction, Multiple Data 单指令，多数据。

也就是：一个控制单元取出一条指令，让多个处理单元同时执行这条指令，但每个处理单元处理的数据不同。

例如大家都执行：

add r1, r2, r3

但是不同处理单元里的 r2、r3 值不同，所以算出来的结果也不同。

GPU 为什么适合这样做？

GPU 的目标是高吞吐，而不是像 CPU 那样尽量降低单个线程的延迟。

所以 GPU 会让多个 streaming processors 共享一个 instruction fetch/dispatch unit。这样可以减少控制逻辑占用的硬件面积，把更多芯片资源用于计算单元。

简单说：

CPU：控制逻辑复杂，擅长处理复杂流程； GPU：控制逻辑共享，擅长大量线程执行相同操作。

Warp 和 SIMD 的关系

一个 warp 中的多个 threads 会一起执行同一条指令，但每个 thread 使用自己的寄存器数据。

所以 warp 的执行方式本质上就是 SIMD 思想在 GPU 上的体现。

CUDA 又把这种模式包装成 SIMT：

程序员写的是一个个独立 thread 的代码； 硬件实际把这些 threads 组成 warp，用 SIMD 风格执行。

一句话记：

SIMD 是硬件执行方式，SIMT 是 CUDA 给程序员提供的线程抽象，warp 是二者连接起来的关键单位。

总结

4.4 节主要解释了 thread block 被分配到 SM 之后，SM 如何进一步组织和执行其中的 threads。

一个 block 会被划分成若干个 warp，目前常见的 warp size 是 32。对于一维 block，连续 32 个 threadIdx.x 组成一个 warp；对于二维或三维 block，线程会先按照 row-major 规则线性化，再每 32 个线程组成一个 warp。如果 block 的线程数不是 32 的整数倍，最后一个 warp 会包含 inactive threads。

Warp 是 SM 内部线程调度的基本单位。SM 执行 warp 时采用 SIMD 风格：同一个 warp 中的线程在同一时刻执行同一条指令，但作用在不同数据上。CUDA 把这种模式抽象为 SIMT，即程序员写的是独立的 scalar thread 代码，而硬件自动把 threads 组成 warp 来利用 SIMD 硬件。

这一节最关键的理解是：

CUDA 编程模型暴露的是 thread/block/grid； GPU 硬件真正调度执行时，还存在 warp 这一层； warp 是理解 GPU 执行效率、控制分支、内存访问合并和后续性能优化的核心概念。

4.5 Control divergence

1. 什么是 control divergence？

在 4.4 节中已经讲过，GPU 以 warp 为单位执行线程。同一个 warp 中的线程通常按照 SIMD/SIMT 方式执行：硬件取出一条指令，让 warp 内多个线程一起执行这条指令。

这种方式在所有线程执行相同代码路径时效率最高。比如一个 if-else 语句中，如果 warp 内所有线程都进入 if 分支，或者都进入 else 分支，那么执行效率就比较好。

但是，如果同一个 warp 内的一部分线程进入 if 分支，另一部分线程进入 else 分支，就会出现 control divergence，控制流分化。

简单说：

control divergence 指同一个 warp 内的线程走了不同的控制流路径。

2. 分支分化时，硬件会分多次执行不同路径

因为一个 warp 在硬件上共享同一条 instruction stream，所以它不能真正让一部分线程执行 A 指令、另一部分线程同时执行 B 指令。

当 warp 内线程走不同分支时，硬件通常会分多次执行这些路径。

比如：

if (threadIdx.x < 24) {
    A;
} else {
    B;
}
C;

假设一个 warp 包含 thread 0 到 thread 31：

• thread 0 到 23 执行 A；
• thread 24 到 31 执行 B。

那么硬件会先执行一遍 A 路径，此时 thread 0 到 23 是 active，thread 24 到 31 是 inactive；然后再执行一遍 B 路径，此时 thread 24 到 31 是 active，thread 0 到 23 是 inactive。最后大家再继续执行 C。

3. divergence 的代价

Control divergence 的主要代价有两个。

第一，需要多次执行不同路径。 原本一个 warp 只需要执行一条统一路径，现在要分别执行 then-path 和 else-path，执行时间可能变长。

第二，inactive threads 也占用执行资源。 当硬件执行 if 分支时，不走这个分支的线程会被屏蔽掉，不产生有效结果，但它们仍然占着 warp 的位置。因此 SIMD/SIMT 执行效率下降。

可以理解为：

warp 还是 32 个线程的位置，但某些路径上只有一部分线程真正干活，其余线程空转。

4. Volta 之后的 independent thread scheduling

书中提到，在 Pascal 及更早架构中，分化路径通常是顺序执行的：先完整执行一个路径，再执行另一个路径。

从 Volta 架构开始，NVIDIA 引入了 independent thread scheduling。这意味着不同路径的执行可能会交错进行，而不是严格地一个路径执行完再执行另一个路径。

这带来一个重要影响：

不能再默认认为 warp 内线程在分支结束后一定自动重新同步。

所以，如果程序要求 warp 内所有线程完成某个阶段之后再继续，就需要使用 warp-level barrier，例如：

__syncwarp();

这点和早期“warp 内线程天然同步”的想法不同。写现代 CUDA 代码时，不能随便依赖隐式 warp 同步。

5. 循环也可能产生 divergence

Control divergence 不只发生在 if-else，也可能发生在循环中。

比如：

int N = a[threadIdx.x];

for (int i = 0; i < N; i++) {
    A;
}

如果同一个 warp 中不同线程的 N 不同，那么每个线程需要执行的循环次数也不同。

图 4.10 中就是这种情况：有些线程循环 4 次，有些线程循环 8 次。前几轮可能所有线程都 active，但随着一些线程提前完成循环，后续迭代中就只有部分线程还在执行，其余线程 inactive。

所以循环条件如果依赖 threadIdx 或每个线程不同的数据，也可能造成 divergence。

6. 如何判断某个控制语句是否可能导致 divergence？

判断方法比较直接：

如果控制条件依赖于线程相关的值，那么它就可能导致 warp 内 divergence。

比如：

if (threadIdx.x < 3) {
    ...
}

这个条件依赖 threadIdx.x。在同一个 warp 中，thread 0、1、2 会进入分支，而其他线程不会进入，因此会产生 divergence。

类似地，如果循环次数依赖 threadIdx.x 或每个线程自己的数据，也可能产生 divergence。

不过要注意：条件依赖 threadIdx.x 并不一定总是造成 divergence，关键要看同一个 warp 内的线程是否会得到不同判断结果。例如：

if (blockIdx.x == 0) {
    ...
}

虽然也是条件判断，但同一个 block 内所有线程的 blockIdx.x 相同，所以通常不会造成 warp 内 divergence。

7. 边界检查是常见的 divergence 来源

CUDA 程序中最常见的一类 divergence 是边界检查。

例如向量加法中经常写：

int i = blockIdx.x * blockDim.x + threadIdx.x;

if (i < n) {
    C[i] = A[i] + B[i];
}

因为数据长度 n 不一定刚好是 block size 的整数倍，所以最后一个 block 中可能有些线程对应的 i 超出范围。为了避免越界访问，需要用 if (i < n) 保护。

这种写法会在最后一个 block 的最后几个 warp 中产生 divergence，但通常影响不大。因为只有边界处少量 warp 会分化，大部分 warp 仍然是完整工作的。

例如书中举例，向量长度为 1003，block size 为 64。为了处理全部元素，需要启动 16 个 blocks，共 1024 个线程。最后 21 个线程是不需要工作的，因此最后一个 block 的最后一个 warp 会有 divergence。但整个 grid 有 32 个 warps，只有最后一个 warp 受影响，所以总体影响很小。

8. 数据规模越大，边界 divergence 的影响越小

边界检查导致的 divergence 通常只发生在数据边界附近。因此数据规模越大，它对总体性能的影响越小。

例如：

• 向量长度 100 时，4 个 warps 中可能有 1 个 warp 受边界 divergence 影响，比例较高；
• 向量长度 1000 时，32 个 warps 中只有 1 个 warp 受影响，约 3%；
• 向量长度 10000 时，313 个 warps 中只有 1 个 warp 受影响，影响小于 1%。

所以边界判断虽然会造成 divergence，但在大规模数据并行任务中，通常不是最主要的性能瓶颈。

9. 二维数据中的边界 divergence

对于二维数据，比如第三章的彩色图像转灰度图，也常用 if 来处理图像边界。

如果图像尺寸不能被 block 的二维尺寸整除，那么位于图像右边界、下边界、右下角区域的 blocks 中会出现一些无效线程。这些线程需要通过边界判断屏蔽掉，因此也可能导致 warp divergence。

不过和一维向量类似，当图像尺寸较大时，边界区域只占整体很小一部分，因此 divergence 的总体影响也会下降。

总结

4.5 节主要讲 control divergence，控制流分化。

GPU 以 warp 为单位执行线程，同一个 warp 内的线程最好执行相同的控制流路径。如果 warp 内一部分线程进入 if 分支，另一部分线程进入 else 分支，硬件就需要分多次执行不同路径，并在每次执行时屏蔽掉不属于该路径的线程。这会降低 SIMD/SIMT 执行效率。

Control divergence 不只来自 if-else，也可能来自循环，尤其是循环次数依赖线程编号或线程私有数据时。判断一个控制语句是否可能产生 divergence，关键看它的条件是否会让同一个 warp 内的线程得到不同结果。

实际 CUDA 程序中，边界检查是最常见的 divergence 来源，例如 if (i < n)。但这类 divergence 往往只影响最后少量 warp，随着数据规模变大，其总体性能影响通常会变小。从 Volta 架构开始，由于 independent thread scheduling，不能再随便假设 warp 内线程在分支结束后自动同步；需要同步时应使用类似 __syncwarp() 的 warp-level barrier。

4.6 Warp scheduling and latency tolerance

1. 为什么 SM 上要驻留很多 warps？

一个 SM 内部虽然有很多 streaming processors，但在任意时刻，它只能执行一部分 resident warps 的指令。也就是说，分配到 SM 上的线程数量通常远大于 SM 同一时刻真正能执行的线程数量。

这看起来有点矛盾：既然不能同时执行所有 warp，为什么还要让这么多 warp 驻留在 SM 上？

原因是：

GPU 通过让更多 warps 驻留在 SM 上，来隐藏长延迟操作，比如 global memory 访问。

当一个 warp 因为等待内存访问结果而不能继续执行时，SM 可以立刻切换去执行另一个已经准备好的 warp。这样，硬件执行单元就不会一直空等。

2. Ready warp：可以被调度执行的 warp

如果一个 warp 当前不需要等待前面指令的结果，它就处于 ready for execution 状态。

当 SM 要选择下一个 warp 执行时，它会从 ready warps 中挑选一个。如果多个 warp 都 ready，硬件会用优先级逻辑决定执行哪个。

所以 warp 调度的基本逻辑是：

某个 warp 等待长延迟操作
        ↓
暂时不调度它
        ↓
选择另一个 ready warp 执行
        ↓
用其他 warp 的计算隐藏等待时间

这就是 GPU 的 latency hiding / latency tolerance。

3. Fine-grained multithreading：细粒度多线程

SM 能够在不同时间点选择不同 ready warp 来执行，这种机制叫 fine-grained multithreading。

它的核心作用是：

利用一个 warp 等待的时间，去执行其他 warp 的指令。

例如，一个 warp 发起 global memory load 后，可能需要等待很多周期。如果 SM 上只有这一个 warp，那执行单元就只能空等。但如果 SM 上还有其他 ready warps，SM 就可以继续执行它们，从而保持较高利用率。

这也是 GPU 能够容忍高延迟内存访问的重要原因。

4. Latency hiding 的直观理解

书中用邮局排队举了一个例子。

假设有一个人到柜台寄包裹，但还没填好表。如果服务员一直等他填完，后面所有人都会被阻塞。更合理的做法是：让他先到旁边填表，服务员继续服务后面已经准备好的顾客。

对应到 GPU：

服务员        → 硬件执行单元
顾客          → warp
没填好表的人  → 正在等待长延迟操作的 warp
准备好的顾客  → ready warp

所以 latency hiding 的本质就是：

不让执行单元等待某个 warp，而是让它继续服务其他 ready warps。

5. Zero-overhead scheduling：零开销调度

GPU 的 warp 切换与 CPU 的线程上下文切换不同。

CPU 在线程切换时，通常需要保存当前线程的状态，再恢复另一个线程的状态，这会产生额外开销。

而 GPU SM 会把已经驻留的 warps 的执行状态都保存在硬件寄存器中。因此，当 SM 从一个 warp 切换到另一个 warp 时，不需要像 CPU 那样保存和恢复上下文。

这称为 zero-overhead thread scheduling。

可以简单理解为：

GPU 切换 warp 很快，因为 resident warps 的状态已经在硬件里，不需要搬来搬去。

6. GPU 为什么不像 CPU 那样依赖大 cache 和复杂分支预测？

CPU 追求降低单个线程的执行延迟，所以会使用复杂的 cache 层次、分支预测、乱序执行等机制来减少等待。

GPU 的思路不同。GPU 不一定要让某个 warp 尽快完成，而是尽量让整个 SM 的执行单元一直有活干。

因此，当某个 warp 等待时，GPU 通过调度其他 warp 来隐藏延迟。这样 GPU 就可以把更多芯片面积用于：

• floating-point execution units；
• memory access channels；
• 大规模并行执行资源。

而不是像 CPU 那样投入大量面积到复杂控制逻辑和缓存机制上。

7. Oversubscription：给 SM 分配超过即时执行能力的线程

为了让 latency hiding 有效，SM 上必须有足够多的 resident warps。否则，当少数 warp 都在等待时，硬件就找不到 ready warp，执行单元仍然会空闲。

书中以 Hopper H100 为例：

• 一个 SM 有 128 个 streaming processors；
• 但一个 SM 最多可以同时驻留 2048 个 threads；
• 也就是驻留线程数最多可以是即时执行单元数量的 16 倍。

这种让 SM 驻留远多于其同一时刻能执行线程数的做法，叫 oversubscription of threads to SMs。它是 GPU 实现 latency tolerance 的关键。

总结

4.6 节主要讲 warp scheduling 如何帮助 GPU 隐藏长延迟操作。

SM 在任意时刻只能执行一部分 resident warps，但它会让很多 warps 同时驻留在 SM 上。当某个 warp 因为 global memory 访问、浮点流水线操作或分支等原因等待结果时，SM 会暂停调度它，转而执行其他 ready warps。这样，一个 warp 的等待时间就被其他 warp 的执行覆盖掉了，这就是 latency hiding。

GPU 的 warp 调度属于细粒度多线程，并且通常是 zero-overhead scheduling，因为 resident warps 的执行状态保存在硬件寄存器中，切换 warp 不需要像 CPU 线程上下文切换那样保存和恢复状态。

这一节的关键理解是：

GPU 不靠显著降低单个线程的延迟来提高性能，而是靠大量 resident warps 和快速 warp scheduling，把等待时间隐藏起来，从而维持整体吞吐量。

4.7 Resource partitioning and occupancy

1. Occupancy 的基本含义

前面 4.6 讲到，为了隐藏 global memory 访问等长延迟操作，一个 SM 上最好同时驻留很多 warps。这样当某些 warp 在等待时，硬件还能调度其他 ready warp 执行。

但是，SM 不一定总能达到最大可驻留 warp 数。书中把：

实际分配到一个 SM 上的 warps 数量 / 该 SM 支持的最大 warps 数量

称为 occupancy，占用率。

也可以近似理解为：

一个 SM 的线程槽位被利用了多少。

比如某个 SM 最多支持 2048 个线程，但某个 kernel 实际只能让每个 SM 驻留 1024 个线程，那么 occupancy 就是：

1024 / 2048 = 50%

Occupancy 高，通常意味着 SM 有更多 warps 可供调度，更容易隐藏延迟；但 occupancy 不是越高性能一定越好，它只是分析性能的重要指标之一。

2. SM 资源是动态分配给 blocks 和 threads 的

一个 SM 除了 streaming processors 之外，还有很多会限制并发执行的资源，例如：

• thread slots：线程槽位；
• block slots：线程块槽位；
• registers：寄存器；
• shared memory：共享内存。

这些资源会被动态分配给驻留在 SM 上的 blocks 和 threads。

以 Hopper H100 为例，书中给出的限制包括：

每个 SM 最多 32 个 blocks
每个 SM 最多 64 个 warps，也就是 2048 个 threads
每个 block 最多 1024 个 threads

如果 block size 是 1024，那么一个 SM 最多只能放 2 个 blocks，因为：

2048 / 1024 = 2

如果 block size 是 512、256、128、64，那么一个 SM 理论上可以分别放：

512  threads/block → 4  blocks/SM
256  threads/block → 8  blocks/SM
128  threads/block → 16 blocks/SM
64   threads/block → 32 blocks/SM

这些情况下都能刚好填满 2048 个 thread slots，因此 occupancy 可以达到 100%。

3. Block slots 也会限制 occupancy

并不是 block size 越小越好。

例如 H100 每个 SM 最多支持 32 个 blocks。如果每个 block 只有 32 个 threads，那么要填满 2048 个 thread slots，需要：

2048 / 32 = 64 blocks

但是一个 SM 只有 32 个 block slots，所以最多只能放：

32 blocks × 32 threads/block = 1024 threads

此时 occupancy 只有：

1024 / 2048 = 50%

这说明：block 太小可能会因为 block slots 不够而无法填满 thread slots。

所以选择 block size 时，不能只看单个 block 线程数少不少，还要看它能不能有效利用 SM 的资源。

4. Block size 不能整除 thread slots 也会降低 occupancy

另一种情况是 block size 和 SM 最大线程数不匹配。

例如 H100 每个 SM 最多支持 2048 个 threads。如果 block size 选择 768，那么一个 SM 最多只能放 2 个 blocks：

2 × 768 = 1536 threads

如果再放第 3 个 block，就需要：

3 × 768 = 2304 threads

超过 2048，因此放不下。

于是有：

2048 - 1536 = 512 thread slots

没有被利用，occupancy 只有：

1536 / 2048 = 75%

所以 block size 的选择会直接影响 occupancy。一个看似合理的 block size，如果不能很好地匹配 SM 的资源上限，也可能导致资源浪费。

5. Registers 也会限制 occupancy

除了 thread slots 和 block slots，registers 也是非常重要的限制因素。

CUDA kernel 中的 automatic variables 通常会被放入寄存器中。不同 kernel 每个线程使用的寄存器数量不同。如果每个线程使用的寄存器太多，那么一个 SM 能同时容纳的线程数就会下降。

以 H100 为例，书中给出每个 SM 有 65536 个 registers。若要满 occupancy，也就是驻留 2048 个 threads，那么每个 thread 最多大约只能使用：

65536 / 2048 = 32 registers/thread

如果一个 kernel 每个 thread 使用 64 个 registers，那么一个 SM 最多只能支持：

65536 / 64 = 1024 threads

此时 occupancy 最多只有：

1024 / 2048 = 50%

也就是说，即使 block size 配得很好，寄存器使用过多也会让 occupancy 上不去。

6. Register spilling 可能不是好事

当寄存器不够用时，编译器有时会进行 register spilling，把一部分本来应该放在寄存器里的值放到内存中。

这样可能降低每个线程对寄存器的需求，从而提高 occupancy。

但问题是，访问 spilled values 需要访问内存，比访问寄存器慢得多。因此 register spilling 虽然可能提高 occupancy，却可能增加每个线程的执行时间，最终让整体性能变差。

所以不能简单认为：

occupancy 提高 = 性能一定提高。

有时候较低 occupancy + 无 spilling，反而比高 occupancy + 大量 spilling 更快。

7. Performance cliff：资源使用稍微增加，性能可能突然下降

书中举了一个很典型的例子。

假设某个 kernel 使用：

31 registers/thread
512 threads/block

一个 SM 可以放 4 个 blocks：

4 × 512 = 2048 threads

总寄存器需求是：

2048 × 31 = 63488 registers

小于 65536，因此可以达到 100% occupancy。

但如果程序中多声明两个 automatic variables，使每个线程寄存器使用量从 31 增加到 33，那么：

2048 × 33 = 67584 registers

超过 65536。于是 runtime 可能只能让每个 SM 放 3 个 blocks：

3 × 512 = 1536 threads

此时 occupancy 从 100% 降到：

1536 / 2048 = 75%

这就是 performance cliff：资源使用只增加了一点点，但并发度和潜在性能却突然下降。

8. Occupancy 受多种资源共同限制

实际 occupancy 不是由单一因素决定的，而是多个资源约束共同作用的结果：

thread slots
block slots
registers
shared memory

只要其中某一种资源先达到上限，就会限制一个 SM 能同时驻留的 blocks/warps 数量。

所以分析 occupancy 时，需要同时考虑：

block size 多大？
每个 thread 用多少 registers？
每个 block 用多少 shared memory？
目标 GPU 每个 SM 的资源上限是多少？

这也是为什么准确手算 occupancy 有时候比较麻烦。实际开发中可以使用 NVIDIA Nsight Compute 的 Occupancy Calculator，或者 CUDA Occupancy API，例如 cudaOccupancyMaxActiveBlocksPerMultiprocessor 之类的接口来估算。

总结

4.7 节主要讲 SM 资源如何限制 occupancy。

Occupancy 表示一个 SM 上实际驻留的 warps 数量占硬件最大支持 warps 数量的比例。较高 occupancy 通常能提供更多 ready warps，帮助隐藏长延迟操作。

但是 occupancy 受到多种资源共同限制，包括 thread slots、block slots、registers 和 shared memory。Block size 太小可能受 block slots 限制，block size 不能整除 SM 最大线程数可能浪费 thread slots，每个线程使用太多 registers 也会降低可驻留线程数。

这一节最重要的理解是：

occupancy 不是单纯由 block size 决定的，而是由 SM 上多个有限资源共同决定的。 调 CUDA kernel 性能时，要关注每个 block 的线程数、每个 thread 的寄存器数、每个 block 的 shared memory 使用量，以及目标 GPU 的资源限制。

4.8 Querying device properties

1. 为什么要查询 device properties？

前面 4.7 讲到，GPU 的 occupancy 会受到很多硬件资源限制，比如：

• 每个 SM 最多能容纳多少 threads；
• 每个 SM 最多能容纳多少 blocks；
• 每个 SM 有多少 registers；
• warp size 是多少；
• grid/block 各维度的最大范围是多少。

这些限制并不是所有 GPU 都一样。不同 GPU 的资源规模、compute capability、SM 数量都可能不同。所以 CUDA 程序如果想适配多种 GPU，就不能完全依赖固定假设，而应该在运行时查询设备属性。

简单说：

4.7 讲“哪些资源会限制 occupancy”； 4.8 讲“怎么在代码里查这些资源”。

2. Compute capability 表示设备能力等级

CUDA 设备的很多资源规模都和 compute capability 有关。一般来说，compute capability 越高，GPU 架构越新，每个 SM 支持的资源也可能更多。

书中举例：

A100: compute capability 8.0
H100: compute capability 9.0

不过实际编程时，不能只看型号或架构名字，最好通过 CUDA Runtime API 查询设备属性。

3. 查询系统中有多少 CUDA 设备

CUDA Runtime 提供了：

cudaGetDeviceCount(&dev_count);

它可以返回当前系统中支持 CUDA 的设备数量。

这很重要，因为一台机器上可能有多个 CUDA-capable devices。有些设备可能是集成显卡，性能较弱；有些才是真正用于计算的独立 GPU。因此程序可以遍历所有设备，查看每个设备的资源和能力，再选择合适的 GPU 运行。

4. 查询某个设备的具体属性

CUDA Runtime 提供了：

cudaGetDeviceProperties(&dev_prop, device_id);

其中 dev_prop 的类型是：

cudaDeviceProp

它是一个 C struct，里面包含了大量描述 GPU 资源和能力的字段。

常见写法类似：

int dev_count;
cudaGetDeviceCount(&dev_count);

for (int dev = 0; dev < dev_count; dev++) {
    cudaDeviceProp dev_prop;
    cudaGetDeviceProperties(&dev_prop, dev);

    // 读取 dev_prop 中的字段
}

这类代码的作用是让程序在运行时知道当前 GPU 的资源限制，而不是把参数写死。

5. 常用 device properties 字段

这一节介绍了一些和执行资源分配密切相关的字段。

maxThreadsPerBlock

表示一个 block 中最多允许多少 threads。

有些设备最多支持 1024 threads/block，有些设备可能更少，未来设备也可能更多。因此如果程序需要较大的 block size，应该查询这个字段。

multiProcessorCount

表示设备中 SM 的数量。

SM 数量越多，通常代表 GPU 可并行执行的 blocks 更多。对于需要大量并行度的程序，这个字段很重要。

clockRate

表示 GPU 的时钟频率。

clockRate 和 multiProcessorCount 结合起来，可以粗略反映 GPU 的最大计算吞吐能力。当然实际性能还会受内存带宽、kernel 特征、occupancy 等因素影响。

maxThreadsDim[0/1/2]

表示一个 block 在 x、y、z 三个维度上分别允许的最大 thread 数。

例如：

dev_prop.maxThreadsDim[0]  // x 维最大线程数
dev_prop.maxThreadsDim[1]  // y 维最大线程数
dev_prop.maxThreadsDim[2]  // z 维最大线程数

这个字段在设置二维、三维 thread block 时很有用，可以避免 block 维度超过硬件限制。

maxGridSize[0/1/2]

表示 grid 在 x、y、z 三个维度上分别允许的最大 block 数。

例如：

dev_prop.maxGridSize[0]  // x 维最大 block 数
dev_prop.maxGridSize[1]  // y 维最大 block 数
dev_prop.maxGridSize[2]  // z 维最大 block 数

如果数据集很大，需要判断一次 kernel launch 的 grid 是否足够覆盖全部数据。如果不够，就可能需要循环处理或采用 grid-stride loop。

regsPerBlock

表示寄存器资源相关限制。

它可以帮助判断一个 kernel 是否可能因为 register 使用过多而降低 occupancy。不过书中也提醒，这个名字有点容易误导：在大多数 compute capability 下，一个 block 可使用的最大 registers 数量确实等于 SM 中总 registers 数量；但在某些 compute capability 下，block 可使用的最大 registers 数量可能小于 SM 总 registers 数量。

warpSize

表示当前设备的 warp size。

虽然目前 CUDA GPU 基本都是 32 threads/warp，但 warp size 是硬件实现相关的。理论上未来设备可能变化，所以代码中需要时可以查询 warpSize，而不是完全写死为 32。

6. Thread block cluster 的资源查询

对于使用 Hopper 架构中 thread block cluster 的 CUDA 程序，书中提到还应该查询最大 cluster size。

这个信息可以通过相关 API 获取，例如：

cudaOccupancyMaxPotentialClusterSize

这说明 CUDA 新架构中的层次不仅有 block 和 SM，还会涉及 cluster 和 GPC，因此资源查询也变得更重要。

总结

4.8 节主要讲 CUDA 程序如何在运行时查询 GPU 的资源和能力。

不同 CUDA 设备的 SM 数量、最大 block size、最大 grid size、register 数量、warp size 等都可能不同。这些属性会直接影响 kernel launch configuration、occupancy 和性能。因此，程序可以通过 cudaGetDeviceCount() 查询系统中 CUDA 设备数量，再用 cudaGetDeviceProperties() 获取每个设备的 cudaDeviceProp 信息。

这一节的核心理解是：

CUDA 程序不应该完全假设硬件资源固定不变。 如果程序需要适配不同 GPU，就应该查询 device properties，并根据设备资源选择合适的 block size、grid size 和优化策略。

4.9 Summary

4.9 总结

一个支持 CUDA 的 GPU 被组织成多个 Streaming Multiprocessors，SM，流式多处理器。每个 SM 又由多个包含 streaming processors 的 processing blocks 组成，这些 processing blocks 共享控制逻辑和存储资源。

当一个 grid 被启动时，它的 thread blocks 会以任意顺序被分配到 SM 上执行，从而使 CUDA 应用具备 transparent scalability，透明可扩展性。透明可扩展性也带来一个限制：不同 thread blocks 中的线程不应该彼此同步。

从 Hopper 架构开始，SM 会被组织成 GPCs，GPU Processing Clusters，blocks 也可以选择性地被组织成 thread block clusters。同一个 cluster 中的 blocks 会在同一个 GPC 上执行，它们可以彼此同步，并且可以访问彼此的 distributed shared memory，分布式共享内存。

线程会以 block-by-block 的方式被分配到 SM 上执行。一个 block 一旦被分配到某个 SM，它就会进一步被划分成 warps。一个 warp 中的线程按照 SIMD，Single Instruction Multiple Data，单指令多数据 模型执行。

如果同一个 warp 中的线程发生分化，走向不同的执行路径，那么 processing block 会分多次执行这些路径。在每一次执行某条路径时，只有选择这条路径的线程处于 active 状态，其他线程处于 inactive 状态。

一个 SM 上驻留的线程数量可能远多于它在同一时刻能够真正执行的线程数量。在任意时刻，SM 只会执行一部分 resident warps 的指令。这样做的好处是：当某些 warps 因为长延迟操作而等待时，SM 可以执行其他 warps，从而避免大量 processing units 空闲。

实际分配到 SM 上的线程数量与该 SM 最大支持线程数量之间的比例，称为 occupancy，占用率。SM 的 occupancy 越高，通常越有利于隐藏长操作延迟，并获得更高的执行吞吐量。

每个支持 CUDA 的设备都会对 SM 中可用资源数量施加不同限制。例如，每个 CUDA 设备都会限制每个 SM 能容纳的 blocks 数量、threads 数量、registers 数量，以及其他资源数量。对于每个 kernel 来说，这些资源限制中的一个或多个都可能成为限制 occupancy 的因素。

CUDA C++ 提供了 occupancy calculator 和相关 API 函数，帮助开发者为不同 GPU 选择更合适的 kernel 和 launch configuration。CUDA C++ 也允许程序员在运行时查询 GPU 中可用的资源。

第四章 4.10 练习题答案精简版

重点关注第一题第五问，第四题，第五题，第八题。

1. 代码分析题

已知：

N = 1024
block size = 128
grid size = 8 blocks
warp size = 32

所以：

每个 block = 128 / 32 = 4 warps
整个 grid = 8 × 4 = 32 warps

a. 4 个 warps/block。 b. 32 个 warps/grid。

c. 第 04 行：

条件：

threadIdx.x < 40 || threadIdx.x >= 104

每个 block 内：

warp 0: 0~31     全部 active，不 divergence
warp 1: 32~63    32~39 active，divergence
warp 2: 64~95    全部 inactive
warp 3: 96~127   104~127 active，divergence

所以每个 block 有 3 个 active warps，2 个 divergent warps。

答案：

i.   active warps = 8 × 3 = 24
ii.  divergent warps = 8 × 2 = 16
iii. warp 0 SIMD efficiency = 32/32 = 100%
iv.  warp 1 SIMD efficiency = 8/32 = 25%
v.   warp 3 SIMD efficiency = 24/32 = 75%

d. 第 07 行：

条件：

i % 2 == 0

每个 warp 中一半线程满足条件，一半不满足。

i.   active warps = 32
ii.  divergent warps = 32
iii. block 0 warp 0 SIMD efficiency = 16/32 = 50%

e. 第 09 行循环：

循环次数：

i % 3 == 0 → 5 次
i % 3 == 1 → 4 次
i % 3 == 2 → 3 次

所以：

j = 0,1,2：所有线程都执行，无 divergence
j = 3,4：部分线程执行，有 divergence

答案：

i.   无 divergence 的迭代 = 3 次
ii.  有 divergence 的迭代 = 2 次

2. 向量长度 2000，block size 512，grid 中有多少 threads？

blocks = ceil(2000 / 512) = 4
threads = 4 × 512 = 2048

答案：2048 个 threads。

3. 上一题中多少个 warps 因边界检查 divergence？

有效线程数是 2000，实际启动 2048 个线程。最后一个 block 中：

有效线程 = 2000 - 1536 = 464
464 = 14 × 32 + 16

所以最后一个 block 中有 1 个 warp 一半有效、一半无效。

答案：1 个 warp。

4. Barrier 等待时间百分比

执行时间：

2.0, 2.3, 3.0, 2.8, 2.4, 1.9, 2.6, 2.9 μs

最慢线程是 3.0 μs，所以总时间：

8 × 3.0 = 24 μs

总等待时间：

(3.0-2.0) + (3.0-2.3) + (3.0-3.0) + (3.0-2.8)
+ (3.0-2.4) + (3.0-1.9) + (3.0-2.6) + (3.0-2.9)
= 4.1 μs

等待比例：

4.1 / 24 ≈ 17.1%

答案：约 17.1%。

5. block 只有 32 threads，就可以省略 __syncthreads() 吗？

答案：不建议。

虽然 32 threads 通常正好是一个 warp，但不能依赖隐式 warp 同步。尤其 Volta 之后有 independent thread scheduling，warp 内线程不一定总是按旧假设同步执行。

如果需要 block 级同步，应使用：

__syncthreads();

如果只需要 warp 级同步，应使用：

__syncwarp();

6. SM 最多 1536 threads、4 blocks，哪种配置线程数最多？

逐项计算：

a. 128 threads/block  → 4 × 128 = 512
b. 256 threads/block  → 4 × 256 = 1024
c. 512 threads/block  → 3 × 512 = 1536
d. 1024 threads/block → 1 × 1024 = 1024

答案：c，512 threads/block。

7. 每个 SM 最多 64 blocks、2048 threads

a. 8 × 128  = 1024 threads → 可行，50%
b. 16 × 64  = 1024 threads → 可行，50%
c. 32 × 32  = 1024 threads → 可行，50%
d. 64 × 32  = 2048 threads → 可行，100%
e. 32 × 64  = 2048 threads → 可行，100%

8. 判断是否 full occupancy

硬件限制：

2048 threads/SM
32 blocks/SM
65536 registers/SM

a. 128 threads/block，30 registers/thread

需要 16 blocks
寄存器 = 2048 × 30 = 61440 < 65536

答案：可以 full occupancy。

b. 32 threads/block，29 registers/thread

达到 2048 threads 需要 64 blocks
但最多只有 32 blocks/SM

最多：

32 × 32 = 1024 threads

答案：不能 full occupancy，限制因素是 block slots，occupancy = 50%。

c. 256 threads/block，34 registers/thread

2048 × 34 = 69632 > 65536

寄存器不够。最多 7 个 blocks：

7 × 256 = 1792 threads
1792 / 2048 = 87.5%

答案：不能 full occupancy，限制因素是 registers，occupancy = 87.5%。

9. 32×32 thread block 做矩阵乘法是否合理？

32 × 32 = 1024 threads/block

但设备限制是：

最多 512 threads/block

所以这个配置非法，kernel 不应该能正常启动。

答案：学生说法有问题。32×32 block 超过该设备每个 block 512 threads 的限制。可以改成 16×16 或 32×16。