AI Infra 硬核拆解GPU 的块、线程和线程束——从软件抽象到硬件本质在我们的 SM 架构系列文章中我们已经深入解析了 SM 的整体架构、CUDA Core、Tensor Core 和 Warp 调度的工作原理。今天我们终于要把这些零散的硬件概念串联起来讲解 CUDA 编程模型的核心——线程Thread、线程块Block和线程束Warp。毫不夸张地说这三个概念是所有 GPU 编程和 AI 性能优化的基石。如果你不能清晰地理解它们之间的关系以及它们如何映射到硬件上那么你写的 CUDA 代码永远只能是能跑而不可能跑得快。很多人学 CUDA 时对这三个概念的理解停留在线程是最小执行单位块是线程的集合网格是块的集合这样的表面定义上。但这远远不够。今天我们就从软件抽象到硬件本质彻底搞清楚这三个概念到底是什么以及它们是如何决定 GPU 真实性能的。一、为什么需要这样的层级结构在开始之前我们先思考一个最根本的问题为什么 NVIDIA 要设计这样一个复杂的三层编程模型直接让程序员管理几千个核心不好吗答案很简单为了在硬件演进时保持软件的兼容性。如果 CUDA 编程模型直接暴露硬件核心的数量那么每当 NVIDIA 发布新一代 GPU核心数量变了所有的旧代码都需要重新编写才能发挥新硬件的性能。这显然是不可接受的。而通过引入 Grid、Block、Thread 这样的软件抽象NVIDIA 成功地将软件逻辑和硬件实现分离开来程序员只需要根据问题的并行性将计算任务分解为 Grid、Block 和 Thread硬件驱动负责将这些软件抽象映射到具体的硬件核心上当硬件升级时只需要更新驱动旧代码就能自动在新硬件上获得更好的性能这就是为什么 15 年前写的 CUDA 代码今天仍然能在最新的 Blackwell GPU 上运行并且性能比原来好得多的原因。二、CUDA 编程模型的三层结构CUDA 编程模型将并行计算任务组织成严格的三层结构Grid网格→ Block线程块→ Thread线程。2.1 整体结构概览一个 Kernel 启动 一个 Grid ├── Block 0 │ ├── Thread 0 │ ├── Thread 1 │ ├── ... │ └── Thread N-1 ├── Block 1 ├── ... └── Block M-1核心定义Kernel在 GPU 上执行的函数Grid一个 Kernel 启动的所有线程的集合BlockGrid 被划分为多个线程块每个线程块是一组可以相互协作的线程Thread最小的执行单位每个线程执行 Kernel 函数的一个实例关键规则同一个 Grid 中的所有线程执行相同的 Kernel 函数同一个 Block 中的线程可以通过共享内存和同步原语相互协作不同 Block 中的线程不能直接通信也不能相互同步2.2 线程Thread最小的执行单位线程是 CUDA 编程模型中最小的执行单位。每个线程有自己的程序计数器PC寄存器集合栈空间线程索引threadIdx每个线程执行的代码完全相同但处理的数据不同。这就是所谓的单程序多数据SPMD编程模型。线程索引每个线程在 Block 内有一个唯一的索引threadIdx它是一个三维向量可以表示为threadIdx.x、threadIdx.y、threadIdx.z。三维索引主要是为了方便处理二维和三维数据如图像、体素。利用内置变量执行内核的每个线程都可以确定自身在其所属块内的位置以及该块在其所属网格内的位置。线程还可以使用这些内置变量来确定线程块和内核启动所在网格的尺寸。这使得每个线程在运行内核的所有线程中都拥有唯一的身份标识。该身份标识通常用于确定线程负责的数据或操作。2.3 线程块Block协作的基本单位线程块是一组可以相互协作的线程的集合。这是 CUDA 编程模型中最重要的一层因为它提供了线程间通信和同步的能力。线程块的核心特性共享内存同一个 Block 中的所有线程共享一块大小有限的片上内存共享内存同步原语同一个 Block 中的线程可以使用__syncthreads()函数进行同步原子操作同一个 Block 中的线程可以对共享内存和全局内存执行原子操作线程块内的所有线程都在同一个子线程SM中执行。这使得线程块内的线程能够高效地相互通信和同步。线程块内的所有线程都可以访问片上共享内存用于在线程块内的线程之间交换信息。线程块索引每个 Block 在 Grid 内有一个唯一的索引blockIdx同样是一个三维向量。全局索引计算要计算一个线程在整个 Grid 中的全局索引需要结合blockIdx和threadIdx// 一维 Grid 和一维 Blockintglobal_idxblockIdx.x*blockDim.xthreadIdx.x;// 二维 Grid 和二维 Blockintglobal_idx_xblockIdx.x*blockDim.xthreadIdx.x;intglobal_idx_yblockIdx.y*blockDim.ythreadIdx.y;2.4 网格Grid一个 Kernel 的所有线程当应用程序启动内核时它会使用许多线程通常是数百万个线程。这些线程被组织成块。顾名思义一个线程块被称为线程块。线程块被组织成网格。网格中的所有线程块都具有相同的大小和尺寸。下图展示了一个线程块网格的示意图。网格是一个 Kernel 启动的所有线程的集合。一个 Grid 由多个线程块组成这些线程块可以并行执行在 GPU 的不同 SM 上。网格的大小Grid 的大小由启动 Kernel 时指定的网格维度gridDim决定。gridDim 也是一个三维向量所以 Grid 可以是一维、二维或三维的。关键限制不同 Block 中的线程不能直接通信也不能相互同步。这是因为不同的 Block 可能执行在不同的 SM 上而 SM 之间没有直接的同步机制。网格可能包含数百万个线程块而执行该网格的GPU可能只有几十或几百个SM。线程块中的所有线程都由单个SM执行并且在大多数情况下[ 1 ]这些线程会在该SM上运行至完成。线程块之间的调度无法保证因此线程块不能依赖其他线程块的结果因为其他线程块可能要等到当前线程块完成后才能被调度。图4展示了网格中的线程块如何分配给SM的示例。CUDA 编程模型使得任意规模的网格能够在任意大小的 GPU 上运行无论其只有一个 SM 还是数千个 SM。为了实现这一点CUDA 编程模型除少数例外情况外要求不同线程块中的线程之间不存在数据依赖关系。也就是说一个线程不应该依赖于同一网格中不同线程块中线程的结果也不应该与该线程块中的线程同步。同一线程块中的所有线程同时在同一个 SM 上运行。网格中的不同线程块会在可用的 SM 之间进行调度并且可以按任意顺序执行。简而言之CUDA 编程模型要求线程块可以按任意顺序并行或串行执行。2.5 线程块簇Thread Block Clusters除了线程块之外计算能力为 9.0 及以上的 GPU 还提供了一种称为“集群”的可选分组级别。集群是一组线程块与线程块和网格一样可以布局在一维、二维或三维空间中。图 5展示了一个由线程块组成的网格该网格也由集群构成。指定集群不会改变网格的维度或网格内线程块的索引。指定集群会将相邻的线程块分组到集群中并为集群级别的同步和通信提供一些额外的机会。具体来说集群中的所有线程块都在单个 GPC 中执行。图 6显示了指定集群时线程块如何调度到 GPC 中的 SM。由于线程块在单个 GPC 内同时调度因此同一集群中不同块的线程可以使用协作组提供的软件接口相互通信和同步。集群中的线程可以访问集群中所有块的共享内存这被称为分布式共享内存。集群的最大大小取决于硬件并且因设备而异。图 6展示了集群内的线程块如何在 GPC 内的 SM 上同时调度。集群内的线程块在网格中始终彼此相邻。三、线程束Warp硬件层面的执行单位上面讲的 Thread、Block、Grid 都是软件层面的抽象。而在硬件层面GPU 并不以线程为单位调度执行而是以Warp线程束为单位。在一个线程块内线程被组织成32个线程一组的线程组称为线程束warp。一个线程束以单指令多线程SIMT模式执行内核代码。在SIMT中线程束内的所有线程都执行相同的内核代码但每个线程可以执行不同的代码分支。也就是说尽管程序的所有线程执行相同的代码但它们不必遵循相同的执行路径。3.1 Warp 的定义Warp 是 GPU 硬件调度和执行的基本单位。一个 Warp 包含 32 个线程这些线程会被同时调度锁步执行同一条指令。关键区别线程是软件层面的最小执行单位Warp 是硬件层面的最小执行单位Warp 的划分方式线程块中的线程会被按照线程索引连续的方式划分为多个 Warp。也就是说线程 0-31 组成 Warp 0线程 32-63 组成 Warp 1线程 64-95 组成 Warp 2以此类推重要结论Warp 的划分是完全确定的只和线程索引有关和线程块的维度无关。无论你的线程块是一维、二维还是三维的Warp 都是按照线程索引的线性顺序划分的。3.2 Warp 与 Block 的关系一个线程块会被划分为若干个 Warp。如果线程块的大小不是 32 的整数倍那么最后一个 Warp 会有部分线程是无效的但它们仍然会被调度执行只是执行结果会被丢弃。线程束执行的一个影响是线程块的最佳线程总数应为 32 的倍数。虽然可以使用任意数量的线程但如果线程总数不是 32 的倍数则线程块的最后一个线程束在整个执行过程中都会有一些通道未被使用。这很可能导致该线程束的功能单元利用率和内存访问效率低下。示例线程块大小为 256256 ÷ 32 8 个完整的 Warp线程块大小为 100100 ÷ 32 3 个完整的 Warp 1 个不完整的 Warp只有 4 个有效线程性能影响不完整的 Warp 会导致硬件利用率下降。例如一个大小为 100 的线程块最后一个 Warp 只有 4 个有效线程CUDA Core 的利用率只有 4/32 12.5%。最佳实践永远让线程块的大小是 32 的整数倍。这是 CUDA 编程的第一条铁律。3.3 Warp 与 SM 的关系Warp 是 SM 调度的基本单位。一个 SM 可以同时驻留多个 Warp这些 Warp 会被 Warp 调度器轮流调度执行。各代 GPU 的 SM 最大驻留 Warp 数AmpereA10064 Warps/SMHopperH10064 Warps/SMBlackwellB10064 Warps/SM关键机制零开销上下文切换。当一个 Warp 因为等待数据而停顿时Warp 调度器会立即切换到另一个就绪的 Warp 执行切换过程没有任何开销。这是 GPU 实现高吞吐量的核心。3.4 线程束发散线程束中的所有线程同时执行同一条指令。如果线程束中的某些线程在执行过程中遵循某个控制流分支而其他线程则不遵循那么不遵循分支的线程将被屏蔽而遵循分支的线程则会被执行。例如如果某个条件仅对线程束中一半的线程成立那么另一半线程束将被屏蔽而活动的线程则会执行这些指令。这种情况如图7所示。当线程束中的不同线程遵循不同的代码路径时这种情况有时被称为线程束发散。因此当线程束中的线程遵循相同的控制流路径时GPU 的利用率最高。四、从软件到硬件完整的映射过程现在我们来完整地跟踪一个 Kernel 从启动到执行完成的整个过程看看软件层面的 Grid、Block、Thread 是如何映射到硬件层面的 SM 和 Warp 的。4.1 阶段 1Kernel 启动与 Grid 创建你在 CPU 代码中调用kernelgridDim, blockDim(args)启动一个 KernelCUDA 驱动在 GPU 上创建一个 Grid包含gridDim.x * gridDim.y * gridDim.z个线程块每个线程块包含blockDim.x * blockDim.y * blockDim.z个线程4.2 阶段 2线程块分配到 SMGPU 的全局调度器将空闲的线程块分配给空闲的 SM一个线程块一旦被分配到某个 SM就会一直在该 SM 上执行直到完成一个 SM 可以同时驻留多个线程块具体数量取决于线程块使用的资源寄存器和共享内存各代 GPU 的 SM 最大驻留线程块数AmpereA10032 Blocks/SMHopperH10032 Blocks/SMBlackwellB10032 Blocks/SM4.3 阶段 3线程块划分为 WarpSM 将分配到的线程块按照线程索引连续的方式划分为多个 WarpSM 为每个 Warp 分配所需的寄存器资源SM 为每个线程块分配所需的共享内存资源4.4 阶段 4Warp 调度与执行Warp 被加入 SM 的就绪队列Warp 调度器从就绪队列中选择一个 WarpWarp 调度器将指令发射给对应的执行单元CUDA Core、Tensor Core 等当 Warp 遇到长延迟操作时被切换出执行另一个 Warp 被调度执行当 Warp 执行完所有指令后它占用的资源被释放4.5 阶段 5Kernel 完成当 Grid 中的所有线程块都执行完成后Kernel 执行结束控制权返回给 CPU五、关键硬件限制与 SM 占用率理解了映射过程我们就可以理解 GPU 性能的关键限制因素SM 占用率。5.1 什么是 SM 占用率SM 占用率是指 SM 上实际活跃的 Warp 数与 SM 最大可支持的活跃 Warp 数的比值。占用率 实际活跃 Warp 数 / SM 最大可支持活跃 Warp 数为什么占用率重要因为 GPU 通过大量的并行 Warp 来隐藏内存延迟。如果占用率太低就没有足够的 Warp 来隐藏延迟计算单元就会空闲性能就会下降。5.2 影响占用率的三大因素占用率受到三个因素的限制形成木桶效应最终的占用率由最严格的限制因素决定寄存器限制每个线程使用的寄存器数量越多SM 能容纳的 Warp 数就越少共享内存限制每个线程块使用的共享内存越多SM 能容纳的线程块数就越少线程块大小限制每个线程块的线程数越多SM 能容纳的线程块数就越少各代 GPU 的 SM 资源配置架构每个 SM 寄存器文件每个 SM 共享内存每个 SM 最大 Warps每个 SM 最大 BlocksA100256 KB164 KB6432H100256 KB228 KB6432B100256 KB256 KB64325.3 占用率计算示例让我们以 H100 为例计算一个 Kernel 的理论占用率已知条件线程块大小256 线程/块每个线程使用寄存器40 个每个线程块使用共享内存16 KB计算过程寄存器限制每个 Warp 使用寄存器40 寄存器/线程 × 32 线程/Warp 1280 寄存器/Warp每个 SM 最多 Warp 数256 KB ÷ 1280 寄存器/Warp 204800 ÷ 1280 160 Warps受最大 Warp 数限制64 Warps共享内存限制每个 SM 最多线程块数228 KB ÷ 16 KB/块 14.25 → 14 块每个块有 8 Warps256 ÷ 32共享内存限制 Warp 数14 × 8 112 Warps → 64 Warps线程块大小限制每个 SM 最多 32 个线程块32 × 8 256 Warps → 64 Warps最终占用率64 ÷ 64 100%5.4 占用率与性能的关系这是最容易被误解的一点占用率不是越高越好而是足够高就好。25% 占用率通常无法有效隐藏内存延迟性能很差25%-50% 占用率可以隐藏大部分延迟性能较好50% 占用率继续提高占用率对性能的提升非常有限有时候降低占用率反而能提高性能例如减少寄存器使用量可以提高占用率但如果每个线程的指令数增加整体性能可能会下降不要盲目追求 100% 的占用率。通过实验找到最佳的占用率平衡点通常在 30%-70% 之间。六、性能优化的核心原则基于对线程、块和线程束的理解我们可以总结出以下核心性能优化原则6.1 选择合适的线程块大小线程块大小是影响性能的最重要参数之一。选择合适的线程块大小需要考虑多个因素必须是 32 的整数倍这是第一条铁律通常在 128-512 之间这是经过实践验证的最佳范围256 是一个很好的默认值在大多数情况下都能获得不错的性能避免使用小于 64 或大于 1024 的线程块大小小于 64 会导致 Warp 数量太少大于 1024 会导致寄存器压力过大不同工作负载的推荐线程块大小逐元素操作256 或 512 线程/块规约操作128 或 256 线程/块矩阵乘法128 或 256 线程/块卷积操作256 线程/块6.2 提供足够的并行性GPU 是一个吞吐量导向的处理器它需要足够多的并行工作来保持所有计算单元忙碌。线程块数量应该至少是 SM 数量的 4-8 倍这样才能让所有 SM 都忙碌起来避免过小的 KernelKernel 启动有一定的开销约 1-10 微秒如果 Kernel 执行时间太短开销占比会很高尽量将多个小 Kernel 融合成一个大 Kernel这样可以减少 Kernel 启动开销提高数据重用率6.3 优化内存访问模式内存访问是大多数 AI 工作负载的瓶颈。优化内存访问模式可以带来数倍的性能提升。合并内存访问让同一个 Warp 内的线程访问连续的内存地址。这样可以将 32 个独立的内存访问合并成一个内存事务带宽利用率提高 32 倍利用共享内存将频繁访问的数据加载到共享内存中减少对全局显存的访问避免非对齐的内存访问非对齐的内存访问会导致多个内存事务带宽利用率下降使用异步内存拷贝重叠计算和数据传输隐藏内存延迟6.4 避免分支发散分支发散是 SIMT 执行模型的最大弱点会严重降低 CUDA Core 的利用率。尽量让同一个 Warp 内的线程执行相同的代码路径使用掩码操作代替分支result condition ? a : b如果必须使用分支尽量让分支条件基于线程块索引而不是线程索引对于边界检查可以只对边界 Warp 进行检查6.5 合理使用共享内存和寄存器共享内存和寄存器是 GPU 上最快的存储资源合理使用它们可以显著提高性能。优先使用寄存器寄存器比共享内存更快并且没有 Bank 冲突问题合理使用共享内存共享内存用于线程间通信和数据重用避免共享内存 Bank 冲突通过添加填充或改变数据布局来避免 Bank 冲突不要过度使用寄存器过多的寄存器使用会降低 SM 占用率七、常见误区与最佳实践误区 1线程越多性能越好真相过多的线程会导致寄存器和共享内存不足反而降低 SM 占用率和性能。最佳实践通过实验找到最佳的线程块大小和网格大小通常在 128-512 个线程/块之间。误区 2线程块越大性能越好真相过大的线程块会导致寄存器压力过大降低 SM 占用率。同时过大的线程块也会导致负载不均衡。最佳实践256 线程/块是一个很好的默认值在大多数情况下都能获得不错的性能。误区 3占用率越高性能越好真相当占用率超过 50% 时继续提高占用率对性能的提升非常有限。有时候为了提高每个线程的性能降低占用率反而会带来更好的整体性能。最佳实践关注整体吞吐量而不是单一的占用率指标。误区 4不同 Block 中的线程可以相互通信真相不同 Block 中的线程不能直接通信也不能相互同步。这是 CUDA 编程模型的基本规则。最佳实践如果需要线程间通信将它们放在同一个 Block 中。如果必须在不同 Block 之间通信使用全局内存和原子操作。八、总结与学习建议线程、块和线程束是 CUDA 编程模型的核心也是连接软件和硬件的桥梁。理解它们之间的关系以及它们如何映射到硬件上是写出高性能 CUDA 代码的前提。核心要点回顾CUDA 编程模型采用三层结构Grid → Block → Thread线程是软件层面的最小执行单位Warp 是硬件层面的最小执行单位一个 Warp 包含 32 个线程按照线程索引连续划分一个线程块只能在一个 SM 上执行一个 SM 可以同时执行多个线程块SM 占用率是衡量 GPU 资源利用率的重要指标但不是越高越好性能优化的核心是选择合适的线程块大小、提供足够的并行性、优化内存访问模式、避免分支发散学习建议动手写几个简单的 CUDA Kernel观察不同线程块大小对性能的影响使用 NVIDIA Nsight Compute 工具分析 Kernel 的执行情况查看 SM 占用率、Warp 调度效率、内存带宽利用率等指标阅读 CUTLASS 和 FlashAttention 等高性能库的源码学习它们是如何组织线程和块的尝试自己实现一个简单的矩阵乘法 Kernel并与 cuBLAS 的实现进行性能对比理解了线程、块和线程束你就真正掌握了 CUDA 编程的精髓。它能让你透过现象看本质快速定位和解决各种性能问题写出真正高效的 AI 系统。
AI Infra 硬件体系与编程模型：7. GPU的块、线程和线程束详解
发布时间：2026/6/8 20:22:14
AI Infra 硬核拆解GPU 的块、线程和线程束——从软件抽象到硬件本质在我们的 SM 架构系列文章中我们已经深入解析了 SM 的整体架构、CUDA Core、Tensor Core 和 Warp 调度的工作原理。今天我们终于要把这些零散的硬件概念串联起来讲解 CUDA 编程模型的核心——线程Thread、线程块Block和线程束Warp。毫不夸张地说这三个概念是所有 GPU 编程和 AI 性能优化的基石。如果你不能清晰地理解它们之间的关系以及它们如何映射到硬件上那么你写的 CUDA 代码永远只能是能跑而不可能跑得快。很多人学 CUDA 时对这三个概念的理解停留在线程是最小执行单位块是线程的集合网格是块的集合这样的表面定义上。但这远远不够。今天我们就从软件抽象到硬件本质彻底搞清楚这三个概念到底是什么以及它们是如何决定 GPU 真实性能的。一、为什么需要这样的层级结构在开始之前我们先思考一个最根本的问题为什么 NVIDIA 要设计这样一个复杂的三层编程模型直接让程序员管理几千个核心不好吗答案很简单为了在硬件演进时保持软件的兼容性。如果 CUDA 编程模型直接暴露硬件核心的数量那么每当 NVIDIA 发布新一代 GPU核心数量变了所有的旧代码都需要重新编写才能发挥新硬件的性能。这显然是不可接受的。而通过引入 Grid、Block、Thread 这样的软件抽象NVIDIA 成功地将软件逻辑和硬件实现分离开来程序员只需要根据问题的并行性将计算任务分解为 Grid、Block 和 Thread硬件驱动负责将这些软件抽象映射到具体的硬件核心上当硬件升级时只需要更新驱动旧代码就能自动在新硬件上获得更好的性能这就是为什么 15 年前写的 CUDA 代码今天仍然能在最新的 Blackwell GPU 上运行并且性能比原来好得多的原因。二、CUDA 编程模型的三层结构CUDA 编程模型将并行计算任务组织成严格的三层结构Grid网格→ Block线程块→ Thread线程。2.1 整体结构概览一个 Kernel 启动 一个 Grid ├── Block 0 │ ├── Thread 0 │ ├── Thread 1 │ ├── ... │ └── Thread N-1 ├── Block 1 ├── ... └── Block M-1核心定义Kernel在 GPU 上执行的函数Grid一个 Kernel 启动的所有线程的集合BlockGrid 被划分为多个线程块每个线程块是一组可以相互协作的线程Thread最小的执行单位每个线程执行 Kernel 函数的一个实例关键规则同一个 Grid 中的所有线程执行相同的 Kernel 函数同一个 Block 中的线程可以通过共享内存和同步原语相互协作不同 Block 中的线程不能直接通信也不能相互同步2.2 线程Thread最小的执行单位线程是 CUDA 编程模型中最小的执行单位。每个线程有自己的程序计数器PC寄存器集合栈空间线程索引threadIdx每个线程执行的代码完全相同但处理的数据不同。这就是所谓的单程序多数据SPMD编程模型。线程索引每个线程在 Block 内有一个唯一的索引threadIdx它是一个三维向量可以表示为threadIdx.x、threadIdx.y、threadIdx.z。三维索引主要是为了方便处理二维和三维数据如图像、体素。利用内置变量执行内核的每个线程都可以确定自身在其所属块内的位置以及该块在其所属网格内的位置。线程还可以使用这些内置变量来确定线程块和内核启动所在网格的尺寸。这使得每个线程在运行内核的所有线程中都拥有唯一的身份标识。该身份标识通常用于确定线程负责的数据或操作。2.3 线程块Block协作的基本单位线程块是一组可以相互协作的线程的集合。这是 CUDA 编程模型中最重要的一层因为它提供了线程间通信和同步的能力。线程块的核心特性共享内存同一个 Block 中的所有线程共享一块大小有限的片上内存共享内存同步原语同一个 Block 中的线程可以使用__syncthreads()函数进行同步原子操作同一个 Block 中的线程可以对共享内存和全局内存执行原子操作线程块内的所有线程都在同一个子线程SM中执行。这使得线程块内的线程能够高效地相互通信和同步。线程块内的所有线程都可以访问片上共享内存用于在线程块内的线程之间交换信息。线程块索引每个 Block 在 Grid 内有一个唯一的索引blockIdx同样是一个三维向量。全局索引计算要计算一个线程在整个 Grid 中的全局索引需要结合blockIdx和threadIdx// 一维 Grid 和一维 Blockintglobal_idxblockIdx.x*blockDim.xthreadIdx.x;// 二维 Grid 和二维 Blockintglobal_idx_xblockIdx.x*blockDim.xthreadIdx.x;intglobal_idx_yblockIdx.y*blockDim.ythreadIdx.y;2.4 网格Grid一个 Kernel 的所有线程当应用程序启动内核时它会使用许多线程通常是数百万个线程。这些线程被组织成块。顾名思义一个线程块被称为线程块。线程块被组织成网格。网格中的所有线程块都具有相同的大小和尺寸。下图展示了一个线程块网格的示意图。网格是一个 Kernel 启动的所有线程的集合。一个 Grid 由多个线程块组成这些线程块可以并行执行在 GPU 的不同 SM 上。网格的大小Grid 的大小由启动 Kernel 时指定的网格维度gridDim决定。gridDim 也是一个三维向量所以 Grid 可以是一维、二维或三维的。关键限制不同 Block 中的线程不能直接通信也不能相互同步。这是因为不同的 Block 可能执行在不同的 SM 上而 SM 之间没有直接的同步机制。网格可能包含数百万个线程块而执行该网格的GPU可能只有几十或几百个SM。线程块中的所有线程都由单个SM执行并且在大多数情况下[ 1 ]这些线程会在该SM上运行至完成。线程块之间的调度无法保证因此线程块不能依赖其他线程块的结果因为其他线程块可能要等到当前线程块完成后才能被调度。图4展示了网格中的线程块如何分配给SM的示例。CUDA 编程模型使得任意规模的网格能够在任意大小的 GPU 上运行无论其只有一个 SM 还是数千个 SM。为了实现这一点CUDA 编程模型除少数例外情况外要求不同线程块中的线程之间不存在数据依赖关系。也就是说一个线程不应该依赖于同一网格中不同线程块中线程的结果也不应该与该线程块中的线程同步。同一线程块中的所有线程同时在同一个 SM 上运行。网格中的不同线程块会在可用的 SM 之间进行调度并且可以按任意顺序执行。简而言之CUDA 编程模型要求线程块可以按任意顺序并行或串行执行。2.5 线程块簇Thread Block Clusters除了线程块之外计算能力为 9.0 及以上的 GPU 还提供了一种称为“集群”的可选分组级别。集群是一组线程块与线程块和网格一样可以布局在一维、二维或三维空间中。图 5展示了一个由线程块组成的网格该网格也由集群构成。指定集群不会改变网格的维度或网格内线程块的索引。指定集群会将相邻的线程块分组到集群中并为集群级别的同步和通信提供一些额外的机会。具体来说集群中的所有线程块都在单个 GPC 中执行。图 6显示了指定集群时线程块如何调度到 GPC 中的 SM。由于线程块在单个 GPC 内同时调度因此同一集群中不同块的线程可以使用协作组提供的软件接口相互通信和同步。集群中的线程可以访问集群中所有块的共享内存这被称为分布式共享内存。集群的最大大小取决于硬件并且因设备而异。图 6展示了集群内的线程块如何在 GPC 内的 SM 上同时调度。集群内的线程块在网格中始终彼此相邻。三、线程束Warp硬件层面的执行单位上面讲的 Thread、Block、Grid 都是软件层面的抽象。而在硬件层面GPU 并不以线程为单位调度执行而是以Warp线程束为单位。在一个线程块内线程被组织成32个线程一组的线程组称为线程束warp。一个线程束以单指令多线程SIMT模式执行内核代码。在SIMT中线程束内的所有线程都执行相同的内核代码但每个线程可以执行不同的代码分支。也就是说尽管程序的所有线程执行相同的代码但它们不必遵循相同的执行路径。3.1 Warp 的定义Warp 是 GPU 硬件调度和执行的基本单位。一个 Warp 包含 32 个线程这些线程会被同时调度锁步执行同一条指令。关键区别线程是软件层面的最小执行单位Warp 是硬件层面的最小执行单位Warp 的划分方式线程块中的线程会被按照线程索引连续的方式划分为多个 Warp。也就是说线程 0-31 组成 Warp 0线程 32-63 组成 Warp 1线程 64-95 组成 Warp 2以此类推重要结论Warp 的划分是完全确定的只和线程索引有关和线程块的维度无关。无论你的线程块是一维、二维还是三维的Warp 都是按照线程索引的线性顺序划分的。3.2 Warp 与 Block 的关系一个线程块会被划分为若干个 Warp。如果线程块的大小不是 32 的整数倍那么最后一个 Warp 会有部分线程是无效的但它们仍然会被调度执行只是执行结果会被丢弃。线程束执行的一个影响是线程块的最佳线程总数应为 32 的倍数。虽然可以使用任意数量的线程但如果线程总数不是 32 的倍数则线程块的最后一个线程束在整个执行过程中都会有一些通道未被使用。这很可能导致该线程束的功能单元利用率和内存访问效率低下。示例线程块大小为 256256 ÷ 32 8 个完整的 Warp线程块大小为 100100 ÷ 32 3 个完整的 Warp 1 个不完整的 Warp只有 4 个有效线程性能影响不完整的 Warp 会导致硬件利用率下降。例如一个大小为 100 的线程块最后一个 Warp 只有 4 个有效线程CUDA Core 的利用率只有 4/32 12.5%。最佳实践永远让线程块的大小是 32 的整数倍。这是 CUDA 编程的第一条铁律。3.3 Warp 与 SM 的关系Warp 是 SM 调度的基本单位。一个 SM 可以同时驻留多个 Warp这些 Warp 会被 Warp 调度器轮流调度执行。各代 GPU 的 SM 最大驻留 Warp 数AmpereA10064 Warps/SMHopperH10064 Warps/SMBlackwellB10064 Warps/SM关键机制零开销上下文切换。当一个 Warp 因为等待数据而停顿时Warp 调度器会立即切换到另一个就绪的 Warp 执行切换过程没有任何开销。这是 GPU 实现高吞吐量的核心。3.4 线程束发散线程束中的所有线程同时执行同一条指令。如果线程束中的某些线程在执行过程中遵循某个控制流分支而其他线程则不遵循那么不遵循分支的线程将被屏蔽而遵循分支的线程则会被执行。例如如果某个条件仅对线程束中一半的线程成立那么另一半线程束将被屏蔽而活动的线程则会执行这些指令。这种情况如图7所示。当线程束中的不同线程遵循不同的代码路径时这种情况有时被称为线程束发散。因此当线程束中的线程遵循相同的控制流路径时GPU 的利用率最高。四、从软件到硬件完整的映射过程现在我们来完整地跟踪一个 Kernel 从启动到执行完成的整个过程看看软件层面的 Grid、Block、Thread 是如何映射到硬件层面的 SM 和 Warp 的。4.1 阶段 1Kernel 启动与 Grid 创建你在 CPU 代码中调用kernelgridDim, blockDim(args)启动一个 KernelCUDA 驱动在 GPU 上创建一个 Grid包含gridDim.x * gridDim.y * gridDim.z个线程块每个线程块包含blockDim.x * blockDim.y * blockDim.z个线程4.2 阶段 2线程块分配到 SMGPU 的全局调度器将空闲的线程块分配给空闲的 SM一个线程块一旦被分配到某个 SM就会一直在该 SM 上执行直到完成一个 SM 可以同时驻留多个线程块具体数量取决于线程块使用的资源寄存器和共享内存各代 GPU 的 SM 最大驻留线程块数AmpereA10032 Blocks/SMHopperH10032 Blocks/SMBlackwellB10032 Blocks/SM4.3 阶段 3线程块划分为 WarpSM 将分配到的线程块按照线程索引连续的方式划分为多个 WarpSM 为每个 Warp 分配所需的寄存器资源SM 为每个线程块分配所需的共享内存资源4.4 阶段 4Warp 调度与执行Warp 被加入 SM 的就绪队列Warp 调度器从就绪队列中选择一个 WarpWarp 调度器将指令发射给对应的执行单元CUDA Core、Tensor Core 等当 Warp 遇到长延迟操作时被切换出执行另一个 Warp 被调度执行当 Warp 执行完所有指令后它占用的资源被释放4.5 阶段 5Kernel 完成当 Grid 中的所有线程块都执行完成后Kernel 执行结束控制权返回给 CPU五、关键硬件限制与 SM 占用率理解了映射过程我们就可以理解 GPU 性能的关键限制因素SM 占用率。5.1 什么是 SM 占用率SM 占用率是指 SM 上实际活跃的 Warp 数与 SM 最大可支持的活跃 Warp 数的比值。占用率 实际活跃 Warp 数 / SM 最大可支持活跃 Warp 数为什么占用率重要因为 GPU 通过大量的并行 Warp 来隐藏内存延迟。如果占用率太低就没有足够的 Warp 来隐藏延迟计算单元就会空闲性能就会下降。5.2 影响占用率的三大因素占用率受到三个因素的限制形成木桶效应最终的占用率由最严格的限制因素决定寄存器限制每个线程使用的寄存器数量越多SM 能容纳的 Warp 数就越少共享内存限制每个线程块使用的共享内存越多SM 能容纳的线程块数就越少线程块大小限制每个线程块的线程数越多SM 能容纳的线程块数就越少各代 GPU 的 SM 资源配置架构每个 SM 寄存器文件每个 SM 共享内存每个 SM 最大 Warps每个 SM 最大 BlocksA100256 KB164 KB6432H100256 KB228 KB6432B100256 KB256 KB64325.3 占用率计算示例让我们以 H100 为例计算一个 Kernel 的理论占用率已知条件线程块大小256 线程/块每个线程使用寄存器40 个每个线程块使用共享内存16 KB计算过程寄存器限制每个 Warp 使用寄存器40 寄存器/线程 × 32 线程/Warp 1280 寄存器/Warp每个 SM 最多 Warp 数256 KB ÷ 1280 寄存器/Warp 204800 ÷ 1280 160 Warps受最大 Warp 数限制64 Warps共享内存限制每个 SM 最多线程块数228 KB ÷ 16 KB/块 14.25 → 14 块每个块有 8 Warps256 ÷ 32共享内存限制 Warp 数14 × 8 112 Warps → 64 Warps线程块大小限制每个 SM 最多 32 个线程块32 × 8 256 Warps → 64 Warps最终占用率64 ÷ 64 100%5.4 占用率与性能的关系这是最容易被误解的一点占用率不是越高越好而是足够高就好。25% 占用率通常无法有效隐藏内存延迟性能很差25%-50% 占用率可以隐藏大部分延迟性能较好50% 占用率继续提高占用率对性能的提升非常有限有时候降低占用率反而能提高性能例如减少寄存器使用量可以提高占用率但如果每个线程的指令数增加整体性能可能会下降不要盲目追求 100% 的占用率。通过实验找到最佳的占用率平衡点通常在 30%-70% 之间。六、性能优化的核心原则基于对线程、块和线程束的理解我们可以总结出以下核心性能优化原则6.1 选择合适的线程块大小线程块大小是影响性能的最重要参数之一。选择合适的线程块大小需要考虑多个因素必须是 32 的整数倍这是第一条铁律通常在 128-512 之间这是经过实践验证的最佳范围256 是一个很好的默认值在大多数情况下都能获得不错的性能避免使用小于 64 或大于 1024 的线程块大小小于 64 会导致 Warp 数量太少大于 1024 会导致寄存器压力过大不同工作负载的推荐线程块大小逐元素操作256 或 512 线程/块规约操作128 或 256 线程/块矩阵乘法128 或 256 线程/块卷积操作256 线程/块6.2 提供足够的并行性GPU 是一个吞吐量导向的处理器它需要足够多的并行工作来保持所有计算单元忙碌。线程块数量应该至少是 SM 数量的 4-8 倍这样才能让所有 SM 都忙碌起来避免过小的 KernelKernel 启动有一定的开销约 1-10 微秒如果 Kernel 执行时间太短开销占比会很高尽量将多个小 Kernel 融合成一个大 Kernel这样可以减少 Kernel 启动开销提高数据重用率6.3 优化内存访问模式内存访问是大多数 AI 工作负载的瓶颈。优化内存访问模式可以带来数倍的性能提升。合并内存访问让同一个 Warp 内的线程访问连续的内存地址。这样可以将 32 个独立的内存访问合并成一个内存事务带宽利用率提高 32 倍利用共享内存将频繁访问的数据加载到共享内存中减少对全局显存的访问避免非对齐的内存访问非对齐的内存访问会导致多个内存事务带宽利用率下降使用异步内存拷贝重叠计算和数据传输隐藏内存延迟6.4 避免分支发散分支发散是 SIMT 执行模型的最大弱点会严重降低 CUDA Core 的利用率。尽量让同一个 Warp 内的线程执行相同的代码路径使用掩码操作代替分支result condition ? a : b如果必须使用分支尽量让分支条件基于线程块索引而不是线程索引对于边界检查可以只对边界 Warp 进行检查6.5 合理使用共享内存和寄存器共享内存和寄存器是 GPU 上最快的存储资源合理使用它们可以显著提高性能。优先使用寄存器寄存器比共享内存更快并且没有 Bank 冲突问题合理使用共享内存共享内存用于线程间通信和数据重用避免共享内存 Bank 冲突通过添加填充或改变数据布局来避免 Bank 冲突不要过度使用寄存器过多的寄存器使用会降低 SM 占用率七、常见误区与最佳实践误区 1线程越多性能越好真相过多的线程会导致寄存器和共享内存不足反而降低 SM 占用率和性能。最佳实践通过实验找到最佳的线程块大小和网格大小通常在 128-512 个线程/块之间。误区 2线程块越大性能越好真相过大的线程块会导致寄存器压力过大降低 SM 占用率。同时过大的线程块也会导致负载不均衡。最佳实践256 线程/块是一个很好的默认值在大多数情况下都能获得不错的性能。误区 3占用率越高性能越好真相当占用率超过 50% 时继续提高占用率对性能的提升非常有限。有时候为了提高每个线程的性能降低占用率反而会带来更好的整体性能。最佳实践关注整体吞吐量而不是单一的占用率指标。误区 4不同 Block 中的线程可以相互通信真相不同 Block 中的线程不能直接通信也不能相互同步。这是 CUDA 编程模型的基本规则。最佳实践如果需要线程间通信将它们放在同一个 Block 中。如果必须在不同 Block 之间通信使用全局内存和原子操作。八、总结与学习建议线程、块和线程束是 CUDA 编程模型的核心也是连接软件和硬件的桥梁。理解它们之间的关系以及它们如何映射到硬件上是写出高性能 CUDA 代码的前提。核心要点回顾CUDA 编程模型采用三层结构Grid → Block → Thread线程是软件层面的最小执行单位Warp 是硬件层面的最小执行单位一个 Warp 包含 32 个线程按照线程索引连续划分一个线程块只能在一个 SM 上执行一个 SM 可以同时执行多个线程块SM 占用率是衡量 GPU 资源利用率的重要指标但不是越高越好性能优化的核心是选择合适的线程块大小、提供足够的并行性、优化内存访问模式、避免分支发散学习建议动手写几个简单的 CUDA Kernel观察不同线程块大小对性能的影响使用 NVIDIA Nsight Compute 工具分析 Kernel 的执行情况查看 SM 占用率、Warp 调度效率、内存带宽利用率等指标阅读 CUTLASS 和 FlashAttention 等高性能库的源码学习它们是如何组织线程和块的尝试自己实现一个简单的矩阵乘法 Kernel并与 cuBLAS 的实现进行性能对比理解了线程、块和线程束你就真正掌握了 CUDA 编程的精髓。它能让你透过现象看本质快速定位和解决各种性能问题写出真正高效的 AI 系统。
 ；YGGFLRRIRPKLK)
