ROCm kernel调度执行过程详解

1.ROCm软件平台

参考：通过 "最差实践 "实验探索 AMD GPU 调度细节(翻译) (qq.com)

本文使用的ROCm（Radeon Open Compute）软件堆栈4.2版本。

在NVIDIA GPU上，术语“CUDA”通常是指GPU编程编译器、API和运行时库，但ROCm不那么单一，通常根据其组件进行描述。图1显示了ROCm中涉及的主要组件堆栈。ROCm的顶部面向用户的组件通常是HIPAPI可移植性接口，它与CUDA几乎相同，主要的实际区别只是API函数的名称。HIP程序中的GPU内核使用LLVM编译器的AMD GPU后端进行编译，并使用ROCclr（ROCm公共语言运行时）运行时库运行。ROCclr下面是一个实现HSA（异构系统架构）API的低级用户空间库，它创建和管理与驱动程序和硬件接口的内存映射队列和命令。

图1 ：ROCm 软件栈的组成部分。

2.GPU编程模型

将任务下发到GPU的程序通常使用以下模式（省略分配内存等设置）：

1. 将输入数据从CPU内存复制到GPU内存.

2. GPU执行一段被称为kernel的GPU代码.

3. 等待GPU代码（kernel）执行完毕.

4. 将结果数据从GPU内存复制到CPU内存. 

从用户空间来看，所有这些步骤都是使用更高级别的API来控制GPU进行的。例如，著名的CUDA API为NVIDIA GPU提供了这种功能。CUDA不支持AMD GPU，因此在本文中我们使用了与CUDA非常相似但适用于AMD的HIP API。还要注意，我们使用术语"kernel"来指代在GPU上运行的代码段（这是GPU文献中常用的术语）。当需要指代操作系统的"kernel code"时，我们使用替代术语，如"driver code"。

3.AMD gpu的计算内核调度

本文将详细说明了kernel如何到达硬件以开始执行。

3.1 调度流程概述

熟悉NVIDIA GPU的读者，应对于了解kernel-launch请求在到达AMD GPU硬件之前经过一系列的过程应该比较熟悉。

图2描述了此请求可能经过的路径。为了帮助简化后面解释的复杂性，首先总结了一个概括性的描述：

1. 用户程序调用hipLaunchKernelGGL API函数来启动一个内核。

2. HIP运行时将一个内核启动命令插入由ROCclr运行时库管理的软件队列中。

3. ROCclr将内核启动命令转换为AQL（体系结构队列语言）数据包。

4. ROCclr将AQL数据包插入HSA（异构系统架构）队列中。

5. 在硬件中，异步计算引擎（ACE）处理HSA队列，将内核分配给计算硬件。

kernel到达GPU计算硬件的旅程始于hipLaunchKernelGGL API调用，如图2顶部所示，负责将kernel启动请求排入队列。

程序员与队列结构的典型接触点是通过HIP的“流”接口。简而言之，HIP流是程序员在调用hipLaunchKernelGGL时可以指定的几个参数之一。每个HIP流由ROCclr管理的软件队列支持，ROCclr是HIP使用的后端运行时库。ROCclr将hipLaunchKernelGGL的参数存储在一个C++对象中，然后将此对象插入软件队列中。

图 2：通过 ROCm 队列结构的路径

• HSA队列

一旦kernel-launch的C++对象到达其软件队列的开头，ROCclr（ROCm公共语言运行时）将其转换为一个kernel调度的AQL（Architected Queueing Language）数据包。

AQL数据包用于请求单个GPU操作，例如kernel启动或内存传输。

为了将AQL数据包发送到GPU，ROCclr将AQL数据包复制到一个HSA（Heterogeneous System Architecture 异构系统架构）队列。HSA队列是AQL数据包的环形缓冲区，并且直接在GPU和用户空间内存之间共享。这种直接的内存共享允许用户程序在不进行系统调用的情况下发出GPU命令。

• HSA队列池

每个ROCclr队列均使用一个专用的HSA队列来支持可能是一个更好的选择，但是图2可能已经揭示了ROCclr更加复杂的行为。

ROCclr的软件队列在内部共享一个HSA队列池：一个ROCclr软件队列可能会向多个不同的HSA队列提交任务，而每个HSA队列可能包含来自多个ROCclr软件队列的任务。即使共享HSA队列有时会阻止并发的kernel-launch，但它却不会破坏顶层"流"抽象背后的所有排序保证。ROCm采用一种硬件（即"屏障" AQL 数据包）和软件机制（即ROCclr的软件队列）的组合来确保来自单个流的命令有序完成。

图2所示的ROCclr软件队列比HSA队列少，这只是为了在图中节省空间。实际上，使用共享的HSA队列池旨在减少应用程序创建的HSA队列的总数；即使有几十个HIP流，ROCclr仍然会使用相同的小型HSA队列池。在ROCm 4.2的默认配置中，该池限制为四个HSA队列。然而，要理解这个限制的原因，需要深入了解调度层次结构。

3.2 为GPU硬件分配队列

如前所述，HSA队列（即kernel-launch packets）的内容可以直接在用户应用程序和GPU硬件之间共享，因此在启动任何一个内核函数时不需要驱动程序代码（这就是为什么在图2中没有显示驱动程序的原因）。

即便如此，在初始化HSA队列和通知GPU其存在时，仍然需要Linux的amdgpu驱动程序。因此，驱动程序代码仍然控制着AMD GPU性能的关键方面，并揭示了关于GPU调度内部的有用细节。

不过，目前我们主要关注启动kernel所需的功能。在驱动程序中，这始于HSA队列的初始化。即使在这里，大部分队列创建逻辑都在ROCm的用户空间代码中处理：HSA API层实际上负责通过mmap和其他标准Linux系统调用为HSA队列预留环形缓冲区、设置操作系统信号等。然而，驱动程序必须负责将这些信息传达给硬件。在内部，它通过填充名为内存队列描述符（MQD）的数据结构来实现，该数据结构包括HSA队列缓冲区的虚拟地址以及其他元数据。MQD之所以被命名为MQD，是因为它们是从CPU内存的GPU可访问区域分配的。然而，为了让GPU真正开始从队列中运行工作，MQD必须分配给GPU上的硬件队列描述符（HQD）。

在我们测试系统的默认配置中，amdgpu驱动程序通过向GPU发送一个runlist来通知GPU新队列——一个包含系统上所有MQD列表的缓冲区。有趣的是，“发送runlist”本身的行为需要将runlist写入一个特殊的GPU命令队列，在驱动程序代码中称为HIQ（HSA接口队列）。驱动程序为系统中的每个GPU创建一个HIQ，与在用户空间中创建的HSA队列不同，这个命令队列被映射到内核空间内存中，并手动分配给GPU硬件，允许它被初始化，而不需要成为runlist本身的一部分。

3.3 队列到达硬件

图3给出了计算工作负载中涉及的GPU硬件的大致表示。如图2所示，异步计算引擎（ACE）是负责处理AQL数据包队列的硬件单元。由于图2关注的是内核函数而不是队列，因此图2不包括MQD分配给HQD的过程。

图 3：Radeon VII 的计算相关组件

3.4 线程块调度

现在我们将描述如何将一个位于HSA队列的头部的内核函数分配给计算硬件运行。

回想一下，线程块是GPU计算的基本可调度实体，因此当内核调度的AQL数据包到达HSA队列的头部时，问题就变成了GPU如何决定运行哪些块，以及在哪里运行它们。

从一个异步计算引擎（ACE）处理单个HSA队列的视角来看，图4实际上延续了图2结束之后的内核启动过程。

为了简化图4，我们只包括了一个HSA队列和一个ACE。如果多个HSA队列被分配给同一个ACE，ACE会以时间片轮询的方式在每个队列的头部调度数据包。

图 4：将数据块分派到 CU 所涉及的硬件。(本图将 "工作负载管理器 "缩写为 WLM）。

CU：计算单元，包含大量的矢量处理逻辑，负责在AMD gpu中执行并行计算。类似于nvidia的流媒体多处理器（SMs）

• 分配块Blocks到SEs

图3和图4中更显著的特征之一是将GPU的计算资源划分为四个着色器引擎（SEs）。

如图4所示，异步计算引擎（ACE）的主要作用是将位于HSA队列头部的内核AQL数据包派发给SE。

然而，在事先没有解释某些设计决策背后原因的情况下，ACE在将块Blocks派发给SEs时的行为可能会显得奇怪。为了避免这种困惑，我们首先描述了HSA规范中提供的一个线程排序保证，AMD在其GPU计算架构实现了该线程排序。

HSA规范指出， GPU线程等待具有较低块索引(即blockIdx提供的值)的任何GPU线程完成必须是安全的，如blockIdx.x提供的值。 从技术上讲， 块索引是一个三维元组， 因此HSA规范实际上根据块的扁平ID（考虑到特殊的blockIdx变量是三维的）来说明其保证。 例如， block 1中的线程(具体来说， ID为1的block)等待block 0中的线程完成一定是安全的， 但是block 0中的线程等待block 1的完成可能是不安全的——block 0可能会占用block 1所需的资源， 从而阻止block 1永远开始执行。 块排序保证有一个实际应用：先前的研究正式证明它可以实现单个内核中的块之间的生产者-消费者关系。 即使本文中的实验没有使用如此复杂的内核逻辑，但块排序保证在调度中起着重要作用， 具有显著的性能影响.

• AMD硬件在将块分配给SE时强制实施块排序

方法很简单：ACE必须按顺序将块Blocks分配给SE。

例如，直到将块0分配给SE 0后，ACE才能将块1分配给SE 1。图5说明了这个概念，ACE将四个连续的块分派给SE。图5显示循环继续，直到将块4分配给SE 0，并且只在内核中的所有块都被分派后结束。

用一个比喻来说：块派发行为可以类比为一个发牌员给四个玩家发牌的扑克牌游戏。就像大多数现实生活中的扑克牌游戏一样，即使有一个慢玩家也可能迫使发牌员等待，从而减慢整个游戏的速度！然而，跳过慢玩家会破坏游戏。将这一概念应用到AMD GPU上， "发牌员"就是ACE，四个 "玩家" 对应四个SE，而 "牌" 就是一个内核的块。但是是什么原因导致SE，也就是比喻中的 "玩家"，变得异常缓慢呢？答案既与CUmask有关，也与工作负载管理器的行为有关

图 5：ACE 向 SE 调度连续块时的简化图。

• 工作负载管理器（WLM）

如图3和图4所示， 在每个SE硬件中将块Blocks分配给cu的模块，称为工作负载管理器（WLM）。

功能：为了让异步计算引擎（ACE）将块分配给SE， 并且将该块分配给SE上的特定CU。 

每个工作负载管理器（WLM）都有四个专用Slot "插槽"，用于暂存传入的块Blocks：GPU 的四个 ACE 各有一个专用插槽。

这种设计意味着一个ACE的活动不会阻止另一个ACE访问工作负载管理器。

理想情况下，工作负载管理器将以循环方式将这四个插槽中的块分配给CU。但是如果没有CU有足够的可用资源（如寄存器或线程）来处理来自特定插槽的快，则可能会出现工作负载管理器将资源需求较小的块提前切入。

附录 A：详细的内核启动流程

第 1.1 节试图用人类可读的语言描述 ROCm 的内核启动行为，但试图在现有 ROCm 代码基础上进行修改的研究人员可能会受益于更详细的描述，包括更具体的源代码参考。本附录假定读者已经熟悉第 1 节讨论的内容。虽然这不是理解本文主体内容的必要条件，但我们还是将其作为希望修改或使用 AMD ROCm 软件栈的研究人员的参考资料。

图 7：HIP 内核启动所涉及的 ROCm 源代码概览。

图 7 更详细地概括了内核启动过程。图 7 流程图中的每个组件都包含四项信息：ROCm 源代码中的函数名称、描述函数用途的（简短）注释、定义函数的 ROCm 组件，以及组件中包含函数定义的特定源代码文件。

如图 1 所示，ROCm 有几个组件，都是开源的。每个相关组件的源代码都可在线获取：本文基于 ROCm 4.2 版，但仍适用于 ROCm 4.3 版（提交本文时的版本），并且自 ROCm 3.7 版以来一直保持相对稳定。

图 7 沿用了图 1 中的流程，但更加详细。如第 1.1 节所述，内核启动请求首先在用户空间 C++ HostQueue 对象中排队，最终转换为 AQL 数据包并插入 HSA 队列。图 7 显示的一个新细节是，从 HostQueue 条目到 AQL 数据包的转换是由异步线程完成的，即调用 hipLaunchKernelGGL 的线程之外的线程。考虑到启动内核时的预期异步行为，这一点不足为奇，但在设计实时系统时，这可能是一个重要的细节，因为这样的线程可能会阻塞等待内核完成的任何其他线程。

好奇的读者可能会注意到，图 7 并没有详细介绍队列创建。相反，我们在附录 B 中对队列创建进行了更深入的解释，包括负责将队列分配给 GPU 硬件的代码。请注意，Stream::Create 代码块同时出现在两个流程图中，这表明如果有必要，给出了队列创建将在整个内核启动过程中的哪个位置进行。

与附录 A 一样，本附录面向已经熟悉第 1.1 和 1.2 节材料的读者。图 8 提供了创建队列并将其分配给 GPU 硬件所需的 ROCm 源代码的更详细视图。再次强调，这些信息是为希望使用 AMD 代码的研究人员准备的，尤其是那些希望在用户层面或驱动程序内部修改 AMD GPU 队列管理的研究人员。如本文正文所述，AMD GPU 的内核启动不需要驱动程序的干预，因此附录 A 中的图 8 不包括任何驱动程序代码。遗憾的是，创建队列背后更复杂的逻辑并非如此，这意味着图 8 也必须包含 AMD 驱动程序代码的一部分，这些代码位于 Linux 内核中。除了用各自的 ROCm 组件标注图 8 中的每个单独组件外，流程图中的驱动程序部分也用虚线矩形括起来。遗憾的是，流程图中驱动程序组件的完整 "文件 "路径太长，无法在流程图中清晰显示，因此我们在此指出，"amdgpu 驱动程序 "框中给出的所有路径都位于 Linux 5.14 源代码树中的 drivers/gpu/drm/amd 目录下。

在随意观察 ROCm 的代码时，最初可能很难分辨出 HSA 队列是在哪里创建的。关键点在于 createVirtualDevice 函数，当 ROCm 启动一个新线程来处理 HIP 流的内核启动时会调用该函数。实际上，"虚拟设备 "是一个允许访问 GPU 的 C++ 接口；许多虚拟设备可能与单个底层 GPU 相关联。图 8 中一个有趣的特征是驱动程序代码中出现了带有 cpsch 后缀的函数名称。cpsch "后缀代表命令处理器调度（Command Processor Scheduling），指的是 HWS 的使用。这些函数的替代版本（后缀为 nocpsch）也可以在内核代码中找到，并在禁用 HWS 时使用。在内部，驱动程序可以通过使用函数指针列表间接调用这些函数，从而交替使用这些函数的不同版本。与图 8 中某些函数取决于 HWS 是否启用的情况类似，某些函数也会根据 GPU 架构的不同而发生变化。这些函数同样是通过函数指针列表间接调用的。我们在图 8 中只包含了一个这样的特定架构函数（init mqd 函数），但还有几个函数用于更低级的细节，例如填充运行列表数据包的内容或 unmap-queues 请求。

图 8：创建 GPU 队列所涉及的 ROCm 源代码概览。