Hardware Compute Partitioning on NVIDIA GPUs\*

文章标题：NVIDIA GPU 上的硬件计算分区
作者/机构：Joshua Bakita and James H. Anderson, Department of Computer Science, University of North Carolina at Chapel Hill

A1 主要贡献

本文旨在解决当前 GPU 管理方法中存在的一个核心问题：GPU 通常被视为一个单一的、不可分割的设备，并强制执行互斥访问控制，这类似于在多核 CPU 上一次只调度一个任务，从而导致严重的计算容量损失，尤其是在单个任务无法饱和所有 GPU 计算核心时。研究目标是为 NVIDIA GPU 提供一种高粒度的计算单元空间分区机制，从而能够安全高效地共享 GPU，回收通常被闲置的计算容量。

为了实现这一目标，本文做出了以下贡献：
1. 揭示了一种自 2013 年以来所有 NVIDIA GPU 中都存在的用于计算单元空间分区的硬件机制。
2. 构建并演示了一个简单、有效且可移植的 GPU 空间分区 API，名为 libsmctrl。
3. 提供了迄今为止未公开的关于 NVIDIA GPU 硬件调度流水线的详细信息。
4. 详细描述了 NVIDIA GPU 中未曾公开的架构模式，包括计算单元和内存单元之间的布局及互连方式。
5. 评估了在 NVIDIA GPU 上有效使用空间分区的局限性，并为此制定了指导方针。
6. 通过一个案例研究，展示了如何应用空间分区来提高卷积神经网络 (CNN) 的 GPU 利用率并减少延迟。

下图直观地展示了该研究的价值。在传统的时间共享模式下（图 a），GPU 一次只能由一个任务使用，当任务无法占满所有计算核心（SMs）时会造成容量浪费。本文工作实现了空间分区（图 b），允许将 GPU 细分给多个任务，从而回收闲置容量。

A3 背景知识与相关工作

GPU 架构与 CUDA 概述

NVIDIA GPU架构。GPU 是高度并行的加速器，由多个离散的功能单元构成，每个单元内部都具备并行处理能力。以一个近期的 NVIDIA GPU 为例（如图 2 所示），其计算/图形引擎由八个 GPC（通用处理集群）组成，每个 GPC 包含十六个 SM（流式多处理器）。SM 以每两个为一组构成一个 TPC（线程处理集群），而每个 SM 包含 64 个 CUDA 核心。除了计算引擎，GPU 还包括五个异步复制引擎、三个视频编码引擎、一个视频解码引擎和一个 JPEG 解码引擎等。先前的研究【索引6，G. A. Elliott, “Real-time scheduling for GPUs with applications in advanced automotive systems,” 2015】表明，这些引擎可以在一定程度上独立于计算/图形引擎运行。

CUDA编程框架。为了简化对复杂加速器的编程，NVIDIA 开发了 CUDA 编程语言和 API。在 GPU 上执行的程序称为 CUDA 核函数（kernel）。一个典型的 CUDA 程序，如算法 1 所示的向量加法，同时利用了计算/图形引擎（第 13 行）和复制引擎（第 11 和 14 行）。启动核函数时，需要指定线程块（thread blocks）的数量和每个块内的线程数。例如，对于一个 2000 个元素的向量，可以设置 numBlocks = 2 和 threadsPerBlock = 1000。

// 算法1 CUDA中的向量加法
1: kernel VECADD(A: ptr to int, B: ptr to int, C: ptr to int, len: int)
2:   i := blockDim.x * blockIdx.x + threadIdx.x
3:   if i >= len then
4:     return
5:   end if
6:   C[i] := A[i] + B[i]
7: end kernel
8: procedure MAIN
9:   cudaMalloc(d_A, len)
10:  ...
11:  cudaMemcpy(d_A, h_A, len)
12:  ...
13:  vecAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C, len)
14:  cudaMemcpy(h_C, d_C, len)
15:  cudaFree(d_A)
16: end procedure

CUDA执行模型。所有 CUDA 应用程序都在各自的内存地址空间（称为上下文，context）中运行，默认情况下，系统使用时间分割复用（time-division multiplexing）来仲裁活跃的 CUDA 应用和其他 GPU 使用者（如显示任务）【索引7，N. Capodieci et al., “Deadline-based scheduling for GPU with preemption support,” 2018】。在单个 CUDA 应用内部，可以使用多个称为流（streams）的 FIFO 队列来实现工作负载的完全并发【索引10，T. Amert et al., “GPU scheduling on the NVIDIA TX2: Hidden details revealed,” 2017】。将所有使用 CUDA 的功能整合到单一上下文和多个流中，通常是为了避免多个小应用都无法充分利用 GPU 却被互斥地分时调度的情况【索引11，M. Yang et al., “Avoiding pitfalls when using NVIDIA GPUs for real-time tasks in autonomous systems,” 2018】。如图 1 所示，如果能实现硬件分区，将多个应用合并到单一上下文中，可以将 GPU 的繁忙时间减少超过 60%。

实时系统术语。任务（task）由作业（job）组成，其发布频率称为周期（period）。作业在发布后必须在截止时间（deadline）前完成。任务可以是周期性的（periodic）或零星的（sporadic）。任务还可能有关联的重要性（criticality），例如在自动驾驶汽车中，行人检测任务的重要性就高于屏幕显示更新任务。

相关工作

GPU管理的早期工作。自 GPU 用于通用计算以来，研究人员一直致力于开发可预测的 GPU 共享管理方法。早期的代表性工作是 TimeGraph【索引12，S. Kato et al., “TimeGraph: GPU scheduling for Real-Time Multi-Tasking environments,” 2011】，它通过一个拦截层和 GPU 中断驱动的调度器来串行仲裁 GPU 工作提交，但该方法需要完全开源的驱动程序且不支持并行执行。另一种更通用的方法是将 GPU 视为一种通过互斥锁保护的资源，如 GPUSync 框架【索引13，G. A. Elliott et al., “GPUSync: A framework for real-time GPU management,” 2013】及其扩展【索引6，G. A. Elliott, “Real-time scheduling for GPUs with applications in advanced automotive systems,” 2015】。这些工作允许对 GPU 上的辅助单元（如复制引擎）进行独立加锁，实现了引擎粒度的 GPU 内部共享。

GPU内部共享的进展。随后的工作集中于实现更可预测的 GPU 内部共享，特别是针对主要的计算/图形引擎。Otterness 和 Anderson 的近期工作【索引14，N. Otterness and J. H. Anderson, “Exploring AMD GPU scheduling details by experimenting with “worst practices”,” 2021】、【索引15，N. Otterness, “Developing real-time GPU-sharing platforms for artificial-intelligence applications,” 2022】在 AMD GPU 上通过透明的硬件空间分区实现了突破，允许多个应用同时共享计算核心。然而，这项技术尚未扩展到 NVIDIA GPU。针对 NVIDIA GPU 的近期工作仍局限于软件模拟分区【索引1，S. Jain et al., “Fractional GPUs: Softwarebased compute and memory bandwidth reservation for GPUs,” 2019】、【索引2，B. Wu et al., “Enabling and exploiting flexible task assignment on GPU through SM-centric program transformations,” 2015】、【索引3，T. Yandrofski et al., “Making powerful enemies on NVIDIA GPUs,” 2022】、【索引4，N. Feddal et al., “Toward precise real-time scheduling on NVIDIA GPUs,” 2022】或概念分析【索引16，H.-E. Zahaf et al., “Contention-aware GPU partitioning and task-to-partition allocation for real-time workloads,” 2021】。

GPU逆向工程。由于商业 GPU 设计的保密性，逆向工程工作为 GPU 管理研究提供了重要支持。重要的逆向工程工作包括：Otterness 等人【索引10，T. Amert et al., “GPU scheduling on the NVIDIA TX2: Hidden details revealed,” 2017】、【索引20，N. Otterness et al., “Inferring the scheduling policies of an embedded CUDA GPU,” 2017】、Amert 等人以及 Olmedo 等人【索引21，I. S. Olmedo et al., “Dissecting the CUDA scheduling hierarchy: a performance and predictability perspective,” 2020】揭示了 NVIDIA GPU 中用于计算工作的队列结构；Capodieci 等人【索引7，N. Capodieci et al., “Deadline-based scheduling for GPU with preemption support,” 2018】和 Spliet 等人【索引22，R. Spliet and R. Mullins, “The case for limited-preemptive scheduling in GPUs for real-time systems,” 2018】阐明了 NVIDIA GPU 的抢占能力；Jain 等人【索引1，S. Jain et al., “Fractional GPUs: Softwarebased compute and memory bandwidth reservation for GPUs,” 2019】阐明了 NVIDIA GPU 的内存层次结构。此外，Nouveau【索引23，Nouveau: Accelerated open source driver for nVidia cards, 2022】和 Mesa【索引24，The mesa 3d graphics library, 2022】等开源项目也提供了关于 GPU-CPU 接口的关键细节。尽管有这些努力，但目前尚无已发表的方法能够为 NVIDIA GPU 实现计算/图形引擎的透明空间分区。

A2 方法细节

实现计算分区

在 TPC 边界上进行分区

通过TMD结构中的SM掩码字段实现分区。我们通过激活一个名为 TMD（任务元数据）的 NVIDIA 数据结构中两个鲜为人知且未被使用的字段来实现分区。这个数据结构用于启动 GPU 计算工作【索引25，J. F. Duluk Jr et al., “Error checking in out-of-order task scheduling,” 2018】。我们发现的这两个字段名为 SM_DISABLE_MASK_LOWER 和 _UPPER。这些字段在从 TMD 结构 V1.6 (Kepler) 到 V3.0 (Ampere) 的公开文档【索引26，NVIDIA, “Open GPU documentation”】中均有列出，但在任何开源软件中都未使用。我们将这两个字段统称为 SM 掩码（SM mask）。

激活分区的机制与效果。修改由 CUDA 库创建的 TMD 相当困难，因为它在构建后几乎立即被上传到 GPU。为了解决这个问题，我们发现并利用了一个未文档化的 CUDA 调试回调 API，在 TMD 构建之后、上传到 GPU 之前拦截并修改它。通过这个机制，我们对 SM 掩码进行了实验，发现在较旧的 GPU 上，它是一个位掩码，其中置位的比特表示为该 TMD 描述的核函数禁用对应的 SM。这使得 SM 的空间分区成为可能。通过为每个核函数设置 SM 掩码，可以确保一个分区中的核函数不会在另一个分区的任何 SM 上运行。图 3 展示了这种分区效果。实验中，两个核函数 K1 和 K2 在 GTX 1060 3GB 上运行。在左侧的默认行为中，K1 的工作被分发到所有 SM 上，阻塞了 K2 的启动。在右侧，我们为 K1 和 K2 设置了互斥的 SM 掩码，使得 K1 和 K2 可以在各自的分区上并发执行，从而让 K2 能更早开始并完成。

SM掩码实为TPC掩码。在更新的 NVIDIA GPU 上进行实验时，我们发现了一个有趣的差异。如图 4 所示，我们在 GTX 1060 3GB（左）和 Tesla P100（右）上运行了相同的实验，使用了一个交替的奇偶位掩码。在 GTX 1060 上，掩码如预期那样禁用了奇数 SM。但在 P100 上，禁用模式的周期是原来的两倍。这是因为 P100 及更新的 GPU 每个 TPC（线程处理集群）包含两个 SM，而 GTX 1060 及更早的 GPU 每个 TPC 只有一个 SM。这一关键差异和大量支持性结果表明，所谓的“SM 掩码”实际上是作为 TPC 掩码来工作的。该字段的命名可能源于早期 TPC 和 SM 数量相同的时代。

在 GPC 边界上进行分区

由晶圆清扫（Floorsweeping）导致的不一致性。许多 GPU 资源是按 GPC（通用处理集群）实例化的，因此将 GPU 分区与 GPC 边界对齐可以避免资源争用。然而，与 AMD GPU 不同，NVIDIA GPU 的 TPC 到 GPC 的映射并不是固定的。我们发现，由于一种名为“晶圆清扫”（floorsweeping）的技术，即使是同一型号的两块 GPU，其内部 TPC 布局和编号也可能不同。该技术通过禁用有缺陷的芯片部分来利用不完美的芯片。由于制造缺陷是随机分布的，不同芯片上禁用的单元也不同，导致 TPC 在重新编号后其所属的 GPC 也可能发生变化，如图 5 所示。

TPC到GPC映射的提取。由于晶圆清扫的存在，要实现跨芯片的 GPC 边界分区，必须首先获取 TPC 到 GPC 的映射关系。我们确定了包含这些信息的 GPU 寄存器，并构建了一个名为 nvdebug 的 Linux 内核模块，通过 /proc 接口来提取和暴露这些信息。

我们的 API

libsmctrl库。结合 TPC 和 GPC 的分区能力，我们开发了一个灵活的用户空间 C 语言库和 API，名为 libsmctrl。API 详见表 I。我们支持在全局、每流（per-stream）和每核函数（per-kernel）级别设置 TPC 掩码，高粒度的设置会覆盖低粒度的设置。这允许默认情况下禁用大多数 TPC，仅对明确允许的流或核函数授予访问权限。我们的库支持 aarch64 和 x86_64 CPU，以及自 2016 年以来的 CUDA 版本，并且是完全用户空间的，二进制文件可移植。激活时，库对每个核函数调用的开销仅为几条指令。

表 I. 我们的 LIBSMCTRL 库的 API

基本分区示例。清单 1 展示了 API 的基本用法。我们建议默认禁用大多数 TPC，以避免 CUDA 内部核函数干扰分区。在示例中，我们在第 3 行设置了全局默认掩码，只允许工作在 TPC 0 上运行。然后创建了两个 CUDA 流，并在第 12 和 16 行为它们设置了不同的 TPC 分区，从而允许在 otherStream 中运行长任务，同时在 urgentStream 中能够立即响应零星任务。我们建议不要默认禁用所有 TPC（掩码为 ~0ull），因为这会导致使用默认掩码启动的核函数（包括 CUDA 内部的）无限期挂起。

// 清单1. 分区API使用示例（9 SM GPU）。
1:  unsigned long long default_mask;
2:  /* 除非另有指定，否则禁用所有TPC，除了TPC 0 */
3:  default_mask = ~(1ull);
4:  libsmctrl_set_global_mask(default_mask);
5:  ...
6:  /* 创建两个流以允许并发执行 */
7:  cudaStream_t otherStream;
8:  cudaStream_t urgentStream;
9:  cudaStreamCreate(&otherStream); cudaStreamCreate(&urgentStream);
10: /* 允许otherStream使用TPC 0-4 */
11: unsigned long long other_mask;
12: other_mask = ((1ull << 5) - 1);
13: libsmctrl_set_stream_mask(otherStream, other_mask);
14: /* 允许urgentStream使用TPC 5-8 */
15: unsigned long long urgent_mask;
16: urgent_mask = 0x1e0; /* 1 1110 0000 */
17: libsmctrl_set_stream_mask(urgentStream, urgent_mask);

GPC分区示例。为了支持在 GPC 边界上进行分区，我们提供了一个围绕 nvdebug 内核模块的易用包装函数。清单 2 展示了如何将两个 GPC 的 TPC 分配给下一个核函数启动。第 4 行获取一个位掩码数组，其中数组索引 i 对应 GPC i，掩码中的置位比特表示哪些 TPC 属于该 GPC。然后在第 10-12 行组合 GPC 0 和 1 的掩码，并在第 16 行应用。

// 清单2. GPC信息API使用示例。
1:  struct gpc_info gpcs;
2:  unsigned long long gpc_mask = 0;
3:  /* 获取每个GPC的TPC位掩码 */
4:  libsmctrl_get_gpc_info(&gpcs);
5:  ...
6:  /* 将GPC 0和1中的TPC分配给
7:     下一个核函数启动 */
8:  if (gpcs.num_gpcs >= 2) {
9:    /* 合并GPC 0和1的掩码 */
10:   gpc_mask |= gpcs.gpc_tpc_masks[0];
11:   gpc_mask |= gpcs.gpc_tpc_masks[1];
12: }
13: ...
14: /* 将组合掩码应用于下一个核函数 */
15: if (gpc_mask != 0)
16:   libsmctrl_set_next_mask(gpc_mask);

局限性

我们的系统通过在 CPU 上拦截 TMD 来工作，因此对于那些 TMD 在 GPU 上构建和执行（不涉及 CPU）的核函数，如使用 CUDA 动态并行（CDP）或 CUDA Graphs 的情况，该方法将不起作用。此外，由于 TMD 结构目前只为分区提供了 64 位，我们无法对超过 64 个 TPC 的 GPU 进行分区。最后，我们当前的库不支持基于每个 GPU 的全局分区配置，但可以轻松扩展以支持此功能。

NVIDIA GPU 如何调度计算任务

架构细节来源

多源信息交叉验证。我们通过交叉引用公共来源来收集架构信息。NVIDIA的专利涵盖了GPU计算优先级【索引27，T. J. Purcell et al., “Scheduling and management of compute tasks with different execution priority levels,” 2011】、块分发逻辑【索引28，J. F. Duluk Jr et al., “Dynamic partitioning of execution resources,” 2022】、【索引29，K. M. Abdalla et al., “Scheduling and execution of compute tasks,” 2015】、抢占设计【索引30，P. A. Cuadra et al., “Software-assisted instruction level execution preemption,” 2020】等。为了验证这些专利描述的是实际硬件而非假设设计，我们将其与NVIDIA的开源nvgpu驱动【索引35，NVIDIA, “nvgpu git repository”】、NVIDIA的公共头文件【索引26，NVIDIA, “Open GPU documentation”】、Nouveau【索引23，N. P. Authors, “Nouveau: Accelerated open source driver for nVidia cards,” 2022】和Mesa【索引24，M. P. Authors, “The mesa 3d graphics library,” 2022】项目的工作以及其他来源进行交叉引用和验证实验。以下讨论是我们对从Kepler到Ampere架构的NVIDIA GPU的综合理解。

计算调度流水线

我们遵循图 6 中编号的路径，追踪一个核函数在 GPU 调度流水线中的过程。

1. 核函数实例化。在 NVIDIA GPU 上，核函数通过前述的 TMD 结构进行内部描述。TMD 定义了核函数所需的线程块、每块线程数、共享内存资源，并包含核函数的入口地址和优先级等信息。用户空间库（如 CUDA）构建 TMD 后，会将一个包含 TMD 指针的启动命令放入命令段（command segment）中。

2. 主机接口 (Host Interface)。主机接口是 CPU 到 GPU 的桥梁。它包含一个或多个 PBDMA（Pushbuffer 直接内存访问）单元，这些单元通过一个间接缓冲区（indirect buffer）从用户空间加载命令，该缓冲区是一个指向命令段的指针的环形缓冲区。这种设计允许用户空间应用直接向 GPU 派发命令，无需系统调用或驱动开销。PBDMA 加载并解析命令后，主机接口会将命令转发到相应的引擎。对于核函数启动，它会将 TMD 传递给计算前端。

3. 计算前端 (Compute Front End)。计算前端将 TMD 指针从主机接口中继到任务管理单元。它还可以协调后续单元中的上下文切换。由于本文专注于单一上下文，我们不进一步探讨相关问题。

4. 任务管理单元 (TMU)。如图 7 所示，TMU 根据优先级和到达顺序对 TMD 进行排队，直到工作分发单元准备好接收它们。TMU 内部由一系列按优先级划分的单向链表构成，每个优先级一个链表，称为 TMD 组（TMD Group）。当 TMU 收到 TMD 时，它会读取 TMD 的 GROUP_ID 字段，并将其附加到相应 TMD 组的尾部。当下游单元发出就绪信号时，TMU 会从优先级最高的非空链表的头部取出一个 TMD 并传递出去。这使得高优先级的 TMD 能够插队到低优先级 TMD 之前。当用户空间派发核函数的速率超过其完成速率时，TMD 会在 TMU 中累积。

5. 工作分发单元 (WDU)。如图 8 所示，WDU 从任务槽（task slots）中将 TMD 分派到可用的 TPC。任务槽的数量是硬件限制的，这迫使 WDU 在就绪核函数数量超过任务槽数量时，基于不完全信息做出调度决策。当有任务槽变空时，WDU 会向 TMU 请求一个新的 TMD，并将其插入任务表（Task Table）和优先级排序任务表（Priority-Sorted Task Table）中。负载均衡器（Load Balancer）从优先级排序表的头部取出 TMD 并分派其线程块。当一个 TMD 的所有块都启动后，它会从优先级排序表中移除，但会保留在任务表中直到所有块完成。这个分派过程可能被两件事打断：TPC 分区和 TMU 中更高优先级的待处理工作。
* TPC分区的影响：如果排序表头部的 TMD 被禁止在某个空闲的 TPC 上执行，WDU 会在表中向前查找，直到找到一个允许在该 TPC 上执行的 TMD。
* 高优先级工作的影响：如果所有 WDU 任务槽都被占用，且其中任何一个槽中的 TMD 优先级低于 TMU 中任何待处理的 TMD，那么这个低优先级的 TMD 将被从 WDU 中驱逐，并替换为高优先级的 TMD。被驱逐的 TMD 会停止分派新块，并被重新插入到 TMU 中相应优先级链表的头部。

内存并行性

6. 执行引擎和内存子系统。图 9 展示了从 L1 缓存到 DRAM 芯片的数据和指令流水线。GPU 的内存系统具有大规模并行性。
* 内存分区单元：通常每个 DRAM 芯片配置一个，包含一个 DRAM 控制器和一部分 L2 缓存。每个内存分区单元通过一个交叉总线（crossbar bus）独立地连接到每个 GPC。
* GPC内部：交叉总线在每个 GPC 内连接到一个 MMU（内存管理单元），用于支持虚拟内存。MMU 再连接到每个 SM 中的 L1 数据缓存和一个 GPC 范围的 L1.5 指令缓存，后者又为每个 SM 的指令缓存提供数据。
* 资源隔离：理论上，每个 GPC 都可以配置为使用 GPU 缓存、总线和 DRAM 资源的独占子集。此外，每个 SM 可以在其 L1 缓存上操作而不会产生干扰。这表明 GPU 硬件完全有能力同时为多个分区提供作业和数据，并为这些分区提供无争用的缓存、总线和 DRAM 资源。

A4 实验环境与结果

实验环境

• 软件配置:
  • 实验框架: cuda_scheduling_examiner 的一个变体，集成了 libsmctrl 库。
  • 案例研究软件: DarkNet 框架的一个变体（用于 YOLOv2）、LITMUSRT（Linux 内核实时调度补丁）、libsmctrl。
• 硬件配置:
  • 测试的 GPU 型号见下表 II，涵盖了从桌面级到服务器级再到嵌入式平台的多种 GPU。
  • 案例研究硬件: 主机系统拥有充足的 CPU 核心，使用 NVIDIA Titan V GPU。

GPU 模型	架构	计算能力	TPC 数量	GPC 数量
GeForce GTX 970	Maxwell	5.2	13	4
GeForce GTX 1060 3GB	Pascal	6.1	9	2
Tesla P100	Pascal	6.0	56	6
Titan V	Volta	7.0	80	6
Jetson AGX Xavier	Volta	7.2	8	1
GeForce RTX 2060	Turing	7.5	30	3
GeForce RTX 3070	Ampere	8.6	46	6
NVIDIA A100	Ampere	8.0	68	8
表 II. 我们实验中测试的GPU

• 模型与数据集:
  • 模型: YOLOv2 卷积神经网络。
  • 数据集: PASCAL VOC 2012。

实验结果

计算分区的灵活性与可移植性

实验表明，TPC 分区在多核函数和多分区的复杂系统中具有良好的扩展性。图 10 展示了一个包含 38 个核函数、17 个独特 TPC 分区和两个流的实验。
* 结论:
1. 每核函数分区：分区的设置是编码在核函数的 TMD 中的，因此可以在运行时为每个核函数动态应用不同的分区。
2. 流顺序保留：即使后续核函数所需的分区与当前核函数互斥，GPU 也会严格遵守 CUDA 流的顺序语义，即同一个流中的后续核函数必须等待前面的核函数全部完成后才能启动。
3. 复杂与重叠分区：TPC 分区可以是不连续的（有“空洞”）、可以重叠，并且没有大小限制。
4. 可移植性：该分区方法在所有计算能力 3.5 (2013年) 及以上、CUDA 8.0 (2017年) 及以上的 NVIDIA GPU 上均有效，包括嵌入式 GPU，如 NVIDIA Xavier SoC。

硬件调度的失效情况

在某些特定情况下，硬件调度机制会妨碍分区的并行执行。
* 贪婪分配的后果: WDU 的调度策略是贪婪的，它按全局优先级和到达时间顺序分派块。当使用重叠分区时，这种策略可能导致非最优的块分配。如图 11 所示，两个核函数 K1 和 K2 在重叠的分区上运行，它们的完成时间取决于哪个核函数先到达 WDU，导致执行时间不可预测。

• 任务槽耗尽的后果: WDU 的任务槽数量有限（例如 GTX 1060 3GB 有 32 个）。当待处理的高优先级核函数数量超过任务槽数量时，低优先级的核函数会被从 WDU 中驱逐。如图 12 所示，一个在独立分区中运行的低优先级核函数 K33，因为 32 个高优先级核函数占满了所有任务槽，导致其被阻塞，无法分派新的块，直到有高优先级核函数完成并释放任务槽。
• 建议: 为防止此类跨分区阻塞，建议使用的 CUDA 流数量不超过 GPU 的 WDU 任务槽数量，或者禁止在所有分区中使用 CUDA 流优先级。

分区策略评估

先前在 AMD GPU 上的研究【索引14，N. Otterness and J. H. Anderson, “Exploring AMD GPU scheduling details by experimenting with “worst practices”,” 2021】发现，将计算单元分配给分区的方式至关重要。本文在 NVIDIA GPU 上复现了该实验，比较了 SE-packed（优先从同一 GPC 分配 TPC）和 SE-distributed（将 TPC 均匀分布到不同 GPC）两种策略。
* 结论: 如图 13 所示，与 AMD GPU 不同，在 NVIDIA GPU 上，向分区中增加 TPC 不会降低该分区的性能，无论这个 TPC 是来自同一个 GPC 还是不同的 GPC。尽管如此，为了最小化 GPC 内部的缓存和总线干扰，仍然推荐使用 SE-packed 分配策略。

案例研究：YOLOv2

我们通过一个案例研究来展示 TPC 分区的实际用途，其中两个 YOLOv2 实例共享一个 GPU。
* 对竞争者的保护: 表 III 的数据显示，当一个 YOLOv2 实例（主实例）单独运行时，平均耗时 24ms。与另一个正常实例共享 GPU（各分配 4 个 TPC）时，耗时约 40ms。但如果竞争实例发生故障并产生大量异常核函数，主实例的执行时间会飙升到 82ms。启用 TPC 分区后，即使竞争实例发生故障，主实例的执行时间仍能稳定在 44ms 左右，几乎恢复到无故障共享时的水平。
* 观察 1: TPC 分区可以保护共享 GPU 上下文的任务免受计算密集型任务的性能影响。

案例	竞争工作	分区启用	平均(ms)	标准差(ms)	最小(ms)	最大(ms)
1	无	否	24	0.99	22	27
2	YOLOv2 (正常)	否	40	1.09	38	44
3	YOLOv2 (故障)	否	82	15.5	62	119
4	YOLOv2 (故障)	是	44	8.36	39	102
表 III. YOLOv2 运行时

• 调整分区大小: 实验探究了在 8 个 TPC 上为两个 YOLOv2 实例分配不同数量 TPC 的影响。
  • 观察 2: 如图 14 所示，TPC 分区允许平滑地调整任务执行时间。增加分配给一个实例的 TPC 数量会减少其执行时间，反之亦然。这可以用作一种基于资源分配的优先级机制。
  • 观察 3: 为任务提供额外的 TPC 可能会带来微不足道的性能提升。图 14 显示，对于 YOLOv2，当 TPC 数量从 1 个增加到 2 个时性能提升显著（执行时间减半），但再增加 TPC 带来的性能提升则会递减。这表明任务的并行度存在瓶颈。

A5 结论

本文揭示了一种自 2013 年以来所有 NVIDIA GPU 都支持的计算核心空间分区新方法。我们开发的库 libsmctrl 使得分区变得简单，并暴露了如 TPC 到 GPC 映射等关键 GPU 细节。通过深入研究 NVIDIA GPU 硬件支持多个并行分区的能力，我们增强了分区的实用性。评估部分讨论了 GPU 调度硬件在某些情况下可能破坏分区边界或做出非最优调度决策的问题。与先前的工作相比，我们的方法更灵活、功能更强大，并通过案例研究证明了其在实际工作负载中的有效性。

未来的工作希望将我们的方法扩展到不需要共享 CUDA 上下文的共存分区。我们还希望进一步研究 GPU 如何执行 CUDA 流排序、如何调度复制等其他引擎，以及 GPU 主机接口在获取命令时究竟如何选择队列。