人工智能系统性能工程 (Chinese Edition)

第10章. 内核 流水线、线程组特化与协作线程块集群

本作品已使用人工智能进行翻译。欢迎您提供反馈和意见：[e-mail]

在前几章中，我们探讨了基础优化技术，包括调整内存访问、最大化并行度、计算与数据传输重叠、提升占用率以及减少线程组停滞。这些方法有助于隐藏延迟并消除瓶颈。然而，现代GPU提供了先进的硬件特性和执行模型，使我们能够将基础优化技术推向更深层次。

本章将介绍若干更高级的CUDA技术，包括：- 专属线程组流水线- 支持网格级与集群级同步的协作组- 循环处理动态工作队列的持久化内核- 基于分布式共享内存（DSMEM或DSM）与张量内存加速器（TMA）多播技术的线程块集群（即协作线程数组集群[CTA]）从高层次来看，线程块集群是一组被保证并行运行的线程块。它们可通过DSMEM相互读写共享内存并执行原子操作。

这些机制使我们能够在无需主机干预的情况下重叠内存访问与计算操作。同时支持跨线程块进行片上数据共享，确保每个流处理器（SM）始终处于满负荷运行状态。

掌握这些现代GPU执行模型后，您将准备好进入下一章：通过探索基于CUDA流的内核间管道，进一步扩展这些优化技术。下一章将在本章讨论的内核优化基础上，构建内核间管道体系。

内核内流水线技术

内核内流水线技术（ ）指在单次内核执行中实现内存操作与计算任务重叠的一系列技术。（下一章将探讨跨内核流水线技术，该技术可在不同流中运行的多个内核间实现任务重叠。）

其核心思想是将内核划分为并行阶段：当某段数据正在加载或存储时，先前加载的数据已进入处理阶段。这些阶段在不同瓦片或数据块上并行运行，从而提升吞吐量并有效隐藏延迟。

传统上，GPU依赖于波段级多线程来隐藏延迟。当一个波段因内存加载而停滞时，其他波段可继续执行计算。这是执行模型中单指令多线程（SIMT）延迟隐藏的基础。

内核流水线技术在此基础上更进一步，通过在同一波段或内核内重叠内存操作与计算任务实现深度优化。它采用精细化协调机制交错执行内存加载与计算操作——有时甚至在单个波段内同时进行。

基于CUDA管道API的内核流水线技术，无需调用__syncthreads()即可实现异步内存传输与计算的重叠执行。内核流水线主要有两种实现方式：双缓冲与波段专职化。

双缓冲（两阶段）流水线方案中，所有线程均采用统一协作模式。而线程组专职化流水线方案则将线程组划分为内存加载、计算、内存存储等不同角色。具体选择取决于工作负载与性能需求。表10-1总结了这两种<cuda/pipeline>变体 。

表 10-1. 基于 CUDA 流水线 API 在现代 GPU 上实现内核流水线化的两种方法

API变体	最佳适用场景	主要用途
双缓冲流水线	基于循环的分块与双缓冲	同一工作组或块内重叠加载与计算
专属工作组流水线（例如三阶段内存加载器、计算单元、内存存储器）	具有多个独立线程组角色的持久内核（本例中为3个）	将波束分配至独立角色/阶段（如内存加载、计算、内存存储）

使用CUDA管道API实现协作式分块与双缓冲

可通过C++管道API实现传统 双缓冲分块模式：创建两阶段管道以重叠内存加载与计算操作. 具体而言，可声明两阶段cuda::pipeline_shared_state<cuda::thread_scope_block, 2>对象，通过协作组（稍后详述）将其作用域限定于特定线程块。这本质上是生产者-消费者模式，如图10-1所示。

图10-1. 基于CUDA管道API的两阶段 生产者-消费者模式 ...