线程块

3秒看懂

线程块（Thread Block / CTA） 是 GPU 并行编程模型中最核心的组织单元。它将数百个线程编为一组，这些线程共享同一块高速片上存储器（共享内存），能够在块内快速同步与协作。在深度学习工作负载中，矩阵乘法、卷积等关键算子正是通过精心划分线程块，把海量计算映射到上千个核心上执行，从而实现远超 CPU 的训练/推理吞吐。

3分钟产业解释

GPU 之所以能颠覆深度学习，关键在于它把「成千上万并行线程」变成了现实。但直接管理几万个线程既困难又低效。于是 NVIDIA 在 CUDA 中引入 线程块（Cooperative Thread Array, CTA），将线程划分成一个二维或三维的逻辑块。一个 kernel 启动时，会生成一个线程块网格（Grid），每个线程块执行相同的指令，但处理不同的数据块。块内线程可以通过 __syncthreads() 轻量级同步，并通过共享内存交换数据。块与块之间则无依赖关系，可独立调度于流多处理器（SM）上。

产业意义：

• 所有主流深度学习框架（PyTorch、TensorFlow）的底层算子均由线程块实现，例如 cuBLAS 的 GEMM、cuDNN 的卷积。
• GPU 架构演进（Volta、Ampere、Hopper）中，线程块的资源上限（最大线程数、共享内存、寄存器）是决定算子性能的关键约束。
• 大规模模型训练（千卡集群）的通信与计算重叠，依赖于对线程块执行顺序与资源占用的精细控制。

换句话说，不了解线程块，就很难理解算子优化为何如此复杂，以及为何一张 GPU 的“实际算力”与理论峰值存在差距。

15分钟专家深入

从专家视角看，线程块远不止“线程的篮子”，它直接制约着 GPU 的 占用率 (occupancy)、延迟隐藏 和 数据局部性。

关键机制：

1. 调度与 Warp：硬件调度器以 32 个线程为一组（一个 warp）作为最小执行单元。一个线程块内包含多个 warp（例如 256 个线程即 8 个 warp）。SM 会同时驻留多个线程块，当某个 warp 因访存延迟或同步停滞时，SM 可瞬时切换到另一个 ready 的 warp，将算术操作的吞吐填补延迟空洞。这是 GPU 隐藏内存延迟的核心手段。
2. 资源约束：每个线程块会消耗一定量的寄存器和共享内存。SM 的资源总量固定，因此单一线程块申请过多资源，会降低 SM 上可同时驻留的线程块数量，从而降低占用率，减少延迟隐藏的机会。“用共享内存换全局内存访问”与“保持高占用率”之间往往存在折衷。
3. 层级并行映射：在深度学习算子中，常采用三层次分解：任务 → 线程块网格（Grid）→ 单个线程块。例如矩阵乘法，Grid 维度对应输出矩阵的分块，每个线程块负责计算一个小输出块；块内线程再协同从全局内存加载数据到共享内存，然后执行内积循环。卷积的 Implicit GEMM 或 Winograd 算法同样遵循此分层。
4. 同步代价：__syncthreads() 是块内 barrier，会强迫 warp 停滞，等待块内所有线程到达该点。若块内线程因分支（如不同 if-else 路径）到达 barrier 的时间参差，产生显著的同步开销。因此高性能 kernel 设计会极力避免线程块内的控制流分歧。

与深度学习负载的深度绑定：

• 矩阵乘法 (GEMM)：线程块处理的输出子块常设计为 128×128 或 256×128 等大小，而线程块本身的维度通常为 16×16 或 32×8（总线程数 256）等，每个线程计算输出矩阵的一个或几个元素，通过共享内存分块（tiling）复用数据，减少 HBM 流量。
• FlashAttention：通过将注意力计算拆分为多个线程块，每个块独自在片上 SRAM 完成部分 softmax 归一化，然后合并，避免写入昂贵的 HBM 中间张量。线程块的精确尺寸与共享内存使用量决定了该算子能否在 40KB~100KB 的有限空间里完成高效计算。
• Megatron-LM 张量并行：通信原语（AllReduce/ReduceScatter）通常预先调用精心优化的 kernel，其线程块参数针对不同矩阵大小和数据布局调校，以在计算与通信之间重叠最优。

因此，一个深度学习系统工程师或编译器开发者，必须能从算子计算图的形状，推演出合理的 Grid 和 Block 维度，并在寄存器、共享内存、占用率三者间寻找 Pareto 最优点。

技术原理（最深）

硬件执行模型

GPU 设备包含多个 SM，每个 SM 内部有：

• 整数/浮点执行单元、加载/存储单元、张量核心（Tensor Core）
• 寄存器文件（如 65536 个 32 位寄存器）
• 共享内存/L1 缓存（可配置大小，如 128KB 组合可选）

一个 kernel 启动时，Host 指定 Grid 维度（gridDim）和 Block 维度（blockDim），生成 gridDim.x × gridDim.y × gridDim.z 个线程块，每个块内含 blockDim.x × blockDim.y × blockDim.z 个线程。

  Grid (二维，处理整个输出张量)
+---------------------+
|  B0  |  B1  |  B2  |
+---------------------+
|  B3  |  B4  |  B5  |   <-- 每个框是一个线程块
+---------------------+

   Block (二维，例如 16x16 线程)
+------------------+
| t0  t1 ... t15 |
| t16 ...     t31|
| ...            |
| t240 ...  t255 |
+------------------+

资源映射与占用率

SM 的资源空间被划分为“槽位”(slot)，每个线程块占据一个槽位，可容纳一个或多个 warp。

• 最大驻留线程块数 = min( max_blocks_per_SM, total_registers / (regs_per_thread × block_size), shared_mem_per_SM / shmem_per_block )
• 占用率 = (活跃 warp 数) / (SM 最大支持 warp 数)

例如，某 SM 可容纳 64 warp（2048 线程）。若每个线程块含 256 线程（8 warp），且 SM 最多可驻留 8 个块，则最大活跃 warp 数为 64，占用率 100%。若某个块使用过多共享内存，导致 SM 只能驻留 4 个块，活跃 warp 数降至 32，占用率 50%，于是 SM 隐藏延迟的能力腰斩，尽管单块计算强度很高。

深度学习算子的线程块设计范式

矩阵乘法 C = A × B（A: M×K, B: K×N）

• Grid 维度：(N / TILE_N, M / TILE_M)
• Block 维度：(TILE_M, TILE_N) 中的线程数，通常二维布局，如 16×16。
• 每个线程块分步将 A 的子块和 B 的子块加载到共享内存，块内所有线程协作完成该输出子块的计算。
• 使用 float4 向量加载、寄存器预取、双缓冲等技术进一步提升吞吐。

卷积（Im2col + GEMM 或 Winograd）：

• 线程块的折叠与循环展开极为复杂，常由自动调谐器（如 Triton、TVM）搜索最佳配置。
• 一个趋势是使用 Hopper 架构的 Thread Block Cluster，将多个线程块声明为一个集群，允许跨块直接访问对方的共享内存（通过 SM-to-SM 网络），进一步减少全局内存往返，这用于加速注意力机制等算子。

与 Tensor Core 的耦合

从 Volta 起，张量核心 (Tensor Core) 提供 warp 级矩阵乘累加指令（mma.sync）。线程块需要按 warp 为单位组织数据以喂入张量核心。例如，一个线程块包含 8 个 warp，每个 warp 可能需要多次调用 mma 指令（单次 mma 仅完成较小形状，如 Ampere 的 fp16 规格为 m16n8k16），整个 warp 通过多次协作计算才能覆盖 16×16×16 等更大形状，然后协作存储结果。这里的同步与数据排布精确到每个 warp 的 lane（线程），设计难度极高，但也是达到接近峰值 TFLOPS 的必经之路。

技术演进史

• G80 (2006)：首次引入线程块模型，块内共享内存与 __syncthreads()。可启动的 Grid 为二维，Block 为三维，最大线程块 512 线程，共享内存 16KB。基于此实现早期卷积网络加速。
• Fermi (2010)：增加 L1/共享内存可配置（16KB/48KB），最大线程数提升至 1024，更灵活的并行调度使得大规模的矩阵库（如 cuBLAS）开始成熟。
• Kepler (2012)：引入动态并行（Dynamic Parallelism），允许 kernel 内部启动新的 kernel 和 Grid，但线程块模型本身保持不变。多 warp 调度器改进，增强了对不规则计算的适应性。
• Pascal (2016)：统一内存、NVLink 助力多 GPU 通信，但线程块模型无本质变化，XLNet、早期 GPT 的分布式训练仍基于传统的 Block-Grid 映射。
• Volta (2017)：推出 Tensor Core，引入了 warp 级协作概念，线程块内需要小心翼翼地将数据布局适配到 mma.sync。线程调度引入独立线程调度，使每线程可拥有自己的程序计数器，但 __syncthreads() 语义不变。共享内存增大至最高 96KB。
• Ampere (2020)：第三代 Tensor Core，支持稀疏矩阵加速，共享内存带宽进一步提升。同时引入了 异步复制（cp.async），允许直接从全局内存加载数据到共享内存而无需中间寄存器，大幅缓解了线程块加载流水线的停滞。
• Hopper (2022)：革命性变化——线程块集群（Thread Block Cluster）。多个线程块可以在硬件层面进行同步（Cluster Group Sync），并相互访问对方的共享内存。这使得跨线程块协同成为可能，显著降低了如 FlashAttention 这类跨块归约算子的 HBM 流量。此外还有分布式共享内存 (DSMEM) 概念，为分布式推理和 MoE 中的 all-to-all 通信提供底层支持。
• 未来：随着 Chiplet 和 3D 封装演进，线程块可能跨越物理 Die，统一的地址空间与更细粒度的同步将打开新的并行范式，进一步影响大模型的训练效率。

技术路线对比

另类对比：CPU 的并行任务划分（如 OpenMP 并行 for）不提供等价于线程块内共享内存的显式抽象，因而数据复用靠缓存自动管理，无法精细控制，难以达到 GPU 级别的能效比。TPU 等专用架构则采用更粗粒度的脉动阵列，其控制单元对标线程块网格，但内部可编程性弱于 CUDA 线程块。

上下游

• 上游：
  • 编译器与工具链：NVCC 编译器负责将 CUDA 代码中的 <<<grid, block 启动语法翻译为驱动可以理解的 ABI 调用，并分配寄存器和共享内存。PTX（Parallel Thread Execution）作为中间表示，向上承接高层语言的线程块抽象，向下映射到 SM 微架构。
  • 算子库与框架：cuBLAS、cuDNN、CUTLASS 等库内部高度模板化，通过 C++ 模板参数在编译期确定 Block 维度、Tile 大小等，供 PyTorch、JAX 调用。新一代 DSL（如 Triton）允许用户在类似 NumPy 的语义上定义寻址与分块，编译器自动生成线程块配置，极大降低开发门槛。
• 下游：
  • 大模型训练（GPT‑4、Llama 3 等）：所有混合精度训练、TP/PP/DP 拆分后的矩阵运算，最终都分解为无数个线程块的执行。线程块的资源利用率直接影响单卡训练成本。
  • 推理服务：batch 推理中的 GEMM 与 Attention 算子，需要低延迟优化，往往使用持久化线程块（Persistent Thread Block）技术避免 kernel 重复启动开销。
  • 边缘与自动驾驶：低功耗 GPU（如 Orin）上的线程块配置需极度节省寄存器，以在有限的 SM 带宽上跑实时模型。

关键指标

• 线程块尺寸：典型值为 128、256、512 个线程，二维布局（如 16×16 或 32×8）常用在矩阵运算中。过小则 warp 浪费，过大则资源不足。
• 共享内存使用量 (per block)：直接影响占用率。深度学习 kernel 通常在 16KB~48KB 范围内取平衡。若需超过硬件上限，必须退化为多次小分块或使用全局内存。
• 寄存器压占 (register pressure)：编译器分配给每个线程的寄存器数量。例如，每个线程使用 128 个寄存器，块大小 256，则消耗 256×128 = 32768 寄存器，可能成为 SM 容纳块数的瓶颈。
• 占用率：工程师常通过 NVIDIA Nsight Compute 工具查看 “Occupancy” 百分比，理想值 >50%，但低占用率且高算数强度时也可能仍达到高效（如某些 GEMM kernel 仅需 2-3 块即可填满计算单元）。
• warp 调度效率：由于 warp 是调度单元，块内 warp 数量最好是 2 的幂且≥ 4（128 线程），保证调度器有足够的“待命 warp”。
• 同步开销：块内 barrier 使用频繁或分歧严重，将导致 SM 空闲周期。无硬数字，需通过 profiler 测量 stall 原因。

供需与市场数据

由于检索失败，无法提供具体的 GPU 出货量或利用率数据。但可以定性描述：

• GPU 供应直接影响线程块并行规模：A100/H100 等数据中心 GPU 持续供不应求，每张卡的 SM 数量、共享内存大小决定了可以同时驻留的线程块总量。全球 AI 算力需求剧增的情况下，云厂商与科技巨头对 Hopper 架构 GPU 的争夺，实质上是在争夺可并发线程块的数量与访问带宽。
• 架构代际供给：NVIDIA 的安培→霍普→布莱克威尔的迭代，每次带来 ~1.5‑2.7 倍的线程块硬件资源提升（如 H100 的集群块、更大共享内存带宽）。这直接缩短大模型单步训练时间，降低总拥有成本（TCO）。
• 软件生态的“线程块效率红利”：架构代际更新往往需要算子重新调优（例如 FlashAttention-2、CUTLASS 3.x 专门为 Hopper 的集群块设计），因此能快速适配新架构的框架和团队将获得显著的算力利用优势，推动头部 AI 公司的资本回报。

代表公司与资本映射

• NVIDIA：CUDA 线程块模型的创建者和持续推动者。从 G80 到 H200，线程块实现细节的逐代进化构筑了其强大的软件护城河。资本市场上，英伟达数据中心收入的爆发式增长，直接受益于其 GPU 上高效线程块执行带来的深度学习训练成本优势。
• AMD：通过 ROCm 生态追赶，其 work‑group 模型与 CUDA 线程块高度兼容。Instinct MI300X 等产品若能在软件栈层面弥合 CUTLASS/hipBLAS 的性能差距，有望从 AI 算力市场中分得份额。
• 新兴 GPU/加速器公司（摩尔线程、壁仞等）：均需实现一套等价于线程块的并行机制。能否提供与之匹配的编译器及标准库，决定了市场对其“国产替代”性能的认可。
• 云厂商：AWS Trainium、Google TPU 等专用芯片虽然不直接暴露线程块概念，但其底层多维并行单元的设计参考了相似的分层映射思想。资本关注点在于它们能否以更低成本承载原先由 GPU 线程块承担的注意力/矩阵运算负载。

产业验证逻辑

• 线程块效率验证：GPU 算力不会自动转化为训练吞吐，取决于算子库对线程块资源的使用效率。评估深耕 CUDA 生态、能快速推出适配新架构库的团队时，应重点看其 kernel 适配速度、吞吐提升幅度和客户验证结果。
• 关注软硬协同：每次 GPU 架构引入新的线程块特性（如 Hopper 的 Tensor Memory Accelerator、Clusters），都要求软件重写关键 kernel。能够快速完成这一过程的公司，将在同期硬件采购中获得先发成本优势。
• 风险点：若未来算法趋势（如 MoE、动态路由）导致计算模式朝着更细粒度、不规则的方向发展，传统密排线程块的优势可能减弱，部分专用稀疏加速器也可能侵蚀 GPU 在特定推理负载中的份额，需持续跟踪。
• 估值锚点：算力即吞吐，而线程块是吞吐的基本计费单元。评估一个 AI 数据中心的 CAPEX 效率，最终会落到每个 SM 每时钟执行的线程块数量上。任何导致该指标恶化的架构/软件缺陷都会反映在运维成本中。

常见误读纠偏

误读1：「线程块数量越多，并行度越高，性能越好」 正解：性能取决于有效隐藏延迟的足够多 warp，而非线程块数量本身。若每块分配的寄存器或共享内存过多，SM 只驻留 1-2 个块，虽然块数少，但每块极高效地利用 Tensor Core，仍可能跑满峰值；相反，大量小块可能引起寄存器溢出或调度开销，导致性能下降。优化内核本质上是在资源限制内寻找“块数×每块效率”的最大积。

误读2：「线程块之间可以通过全局内存随意通信」 正解：线程块之间没有执行顺序保证，不能使用 __syncthreads() 同步；它们虽然可以通过全局内存通信，但缺乏高效同步原语（Hopper 之前）。如果在需要跨块一致性的场景下（如全局归约）随意操作，可能导致数据竞争和不确定结果。正确做法是单个 kernel 内不进行跨块同步，要么拆成多个 kernel，要么使用 Hopper 的集群块或协作组（Cooperative Groups）的网格同步。

误读3：「所有线程块使用相同的共享内存大小」 正解：在同一个 kernel 内，所有线程块确实使用相同尺寸的静态共享内存（编译期确定）或最大动态共享内存（运行时指定）。但是没有任何机制强制它们实际使用量一致，开发者有责任不越界。如果某个块内使用超过声明大小，会导致未定义行为，极难调试。

学习路径

1. 基础：阅读《CUDA C++ Programming Guide》的“Execution Model”和“Memory Hierarchy”章节，理解 Grid、Block、Thread 的索引计算及共享内存声明。
2. 动手：手写一个朴素的矩阵乘法（naive GEMM），然后逐步添加共享内存分块、寄存器预取、向量化加载，使用 Nsight Compute 观察线程块占用率和warp stall 原因。
3. 进阶：研读 CUTLASS 的 gemm_device_level 和 gemm_block_level 抽象，分析如何通过模板参数静态配置线程块形状和计算过程。推荐《CUTLASS: Fast Linear Algebra in CUDA C++》论文。
4. 现代实践：学习 Triton 语言，通过 triton.autotune 探索线程块配置的自动搜索，理解为何编译器能产生接近手写 CUDA 性能的代码。
5. 前沿：深入 Hopper TMA 与 Thread Block Cluster 示例（如 NVIDIA 官方 cuda-samples 中的 cudaTensorCoreGemm 和 FlashAttention-2 源码），掌握跨块直接共享内存的编程模型。
6. 系统思考：结合 Megatron-LM 或 vLLM 源代码，理解大模型系统中的 kernel 调用图谱，建立“线程块 → 算子 → 分布式训练”的全局观。

一句话总结

线程块是 GPU 上最小但最关键的执行协作单元，它的尺寸、资源分配与同步方式直接划定了深度学习计算的效率边界；掌握线程块，就掌握了挖掘 AI 算力的钥匙。

延伸阅读与来源

• NVIDIA CUDA C Programming Guide (Programming Model) – 官方定义与资源约束
• NVIDIA Hopper Architecture Whitepaper – Thread Block Cluster、DSMEM 等新特性
• CUTLASS Documentation & Repository (github.com/nvidia/cutlass) – 矩阵乘法的线程块分解范式
• FlashAttention-2: Faster Attention with Better Parallelism and Work Partitioning – 跨块/跨 warp 负载均衡的实现
• 各代 GPU 架构微手册（NVIDIA Fermi/Kepler/Pascal/Volta/Ampere Tuning Guides） – 占用率与资源调优数值
• 《Programming Massively Parallel Processors》 (David Kirk & Wen-mei Hwu) – 线程块设计案例与习题
• Nsight Compute Kernel Profiling Guide – 如何测量线程块的实际性能瓶颈

（注：因联网检索失败，具体产品数据（如某架构的共享内存精确大小）未引出处；文中所述均为行业共识性概念，不构成投资建议。）