1. 论文基本信息

1.1. 标题

Miriam: Exploiting Elastic Kernels for Real-time Multi-DNN Inference on Edge GPU

1.2. 作者

Zhihe Zhao†, Neiwen Ling†, Nan Guan§, and Guoliang Xing†

• † The Chinese University of Hong Kong (香港中文大学)
• § City University of Hong Kong (香港城市大学)

1.3. 发表期刊/会议

该论文发布于 arXiv，是一个预印本 (preprint) 平台。arXiv 是一个开放获取的论文预印本库，在计算机科学、物理学、数学等多个领域具有重要的影响力，通常用于快速分享最新的研究成果。

1.4. 发表年份

2023年7月10日

1.5. 摘要

许多应用，如自动驾驶和增强现实，需要并发运行多个具有不同实时性能需求的深度神经网络（DNN）。然而，在边缘 GPU 上协调具有不同关键性级别的多个 DNN 任务的研究相对有限。与服务器级 GPU 不同，边缘 GPU 资源受限，并且缺乏硬件级资源管理机制以避免资源争用。因此，本文提出了 Miriam，一个针对边缘 GPU 上多 DNN 推理的争用感知任务协调框架。Miriam 整合了两个主要组件：一个弹性内核生成器（elastic-kernel generator）和一个运行时动态内核协调器（runtime dynamic kernel coordinator），以支持混合关键 DNN 推理。为了评估 Miriam，作者基于 CUDA 构建了一个新的 DNN 推理基准测试，其中包含多样化的代表性 DNN 工作负载。在两个边缘 GPU 平台上的实验表明，与现有基线相比，Miriam 可以将系统吞吐量提高 92%，同时对关键任务仅引入不到 10% 的延迟开销。

1.6. 原文链接

• 原文链接: https://arxiv.org/abs/2307.04339
• PDF 链接: https://arxiv.org/pdf/2307.04339v1.pdf

2. 整体概括

2.1. 研究背景与动机

2.1.1. 论文试图解决的核心问题

论文旨在解决在边缘 GPU 上同时运行多个深度神经网络（DNN）推理任务时，如何有效地协调这些任务，以满足不同任务的实时性能需求，并在资源受限的环境下最大化系统吞吐量，同时保证关键任务的低延迟。

2.1.2. 为什么这个问题在当前领域是重要的？现有研究存在哪些具体的挑战或空白（Gap）？

• 应用需求驱动： 自动驾驶、增强现实等边缘应用普遍要求并发执行多个 DNN 任务，且这些任务具有不同的实时性要求（例如，障碍物检测是 critical task，用户行为分析可以是 normal task）。
• 边缘设备限制： 边缘 GPU 资源有限，与服务器级 GPU 不同，它们缺乏硬件层面的资源管理机制（如 NVIDIA Multi-Process Service (MPS) 和 Multi-Instance GPU (MIG)），这导致多个 DNN 并发运行时容易出现资源争用，降低性能。
• 现有方案不足：
  • 串行执行： 如果为了保证 critical task 的低延迟而串行执行任务，会导致 GPU 利用率低，整体吞吐量大幅下降。
  • 直接并发： 如果直接将多个 DNN 任务并发执行而不进行有效管理，会因为严重的资源争用导致 critical task 的延迟显著增加，无法满足实时性要求。
  • 其他现有方法： 许多现有方法要么牺牲模型精度，要么需要耗时的离线性能分析，或者需要修改底层 GPU 驱动（在闭源设备上不切实际），或者无法很好地处理内核间的资源争用，或不能适应动态变化的运行时环境。
• 冲突的目标： multi-DNN inference 面临两个潜在冲突的目标：
  1. 关键任务优先级： critical DNN task 必须优先执行，以最小化端到端延迟，并避免其他任务的干扰。
  2. 高整体吞吐量： 所有 co-running DNN task 应以最佳努力方式并发执行，以实现高整体吞吐量。

2.1.3. 这篇论文的切入点或创新思路是什么？

论文的创新点在于提出了 elastic kernel（弹性内核）这一抽象概念。通过将传统内核分解为更小、更灵活的单元，这些单元可以根据优先级和关键性动态地调度和重新映射到不同的 GPU 资源。这种 elasticization 方法允许在不引起显著资源争用的情况下，将其他 GPU 内核“填充”到空闲资源中，从而在优先处理关键任务的同时最大化 GPU 利用率和整体吞吐量。其核心在于通过细粒度的内核级 GPU 资源映射来解决 intra-multi-processor (SM) 和 inter-multi-processor 两种资源争用问题。

2.2. 核心贡献/主要发现

2.2.1. 论文最主要的贡献

1. 提出 Miriam 框架： 设计并实现了 Miriam，一个争用感知（contention-aware）的任务协调框架，专门用于边缘 GPU 上的实时 multi-DNN inference。
2. 引入弹性内核（Elastic Kernel）： 提出了 elastic kernel 的概念，这是一种具有可调节 grid size 和 block size 的 GPU 内核，能够实现更细粒度的 GPU 资源映射。
3. 双组件系统设计： Miriam 包含两个主要组件：
   • 弹性内核生成器（Elastic-kernel Generator）： 负责离线生成资源可控的弹性 GPU 内核，包括 elastic grid/block generator 和 source-to-source kernel transformer，确保计算一致性。
   • 运行时动态内核协调器（Runtime Dynamic Kernel Coordinator）： 负责在线调度弹性内核，动态控制 co-running kernel 的执行，以解决 intra-SM 和 inter-SM contention。
4. 构建新的基准测试： 基于 CUDA 构建了一个新的 DNN inference benchmark (MDTB)，包含多样化的代表性 DNN 工作负载，用于评估系统在混合关键性任务场景下的性能。

2.2.2. 论文得出了哪些关键的结论或发现？这些发现解决了什么具体问题？

1. 显著提高系统吞吐量： 实验结果表明，与现有最先进的基线方法相比，Miriam 可以将系统吞吐量提高高达 92%。这解决了边缘设备上多 DNN 并发执行时 GPU 利用率低和整体性能不足的问题。
2. 保证关键任务低延迟： 对于 critical task，Miriam 仅引入了不到 10% 的延迟开销。这解决了在追求高吞吐量的同时，如何不牺牲关键任务实时性要求的难题。
3. 有效管理资源争用： 通过 elastic kernel 设计和动态协调机制，Miriam 能够有效缓解 intra-SM 和 inter-SM 资源争用，提高 GPU 的利用率（Achieved Occupancy）。例如，在深度分析中，Miriam 的平均 layer-wise achieved occupancy 达到 65.25%，远高于 Multi-stream 的 32.9%。
4. 实际应用潜力： 在自动驾驶 LGSVL 模拟器场景的案例研究中，Miriam 在保持 critical task 低延迟的同时，将整体吞吐量提高了 89%，证明了其在真实世界边缘应用中的有效性。

3. 预备知识与相关工作

3.1. 基础概念

为了理解 Miriam 框架，需要对 NVIDIA GPU 的 CUDA 编程模型和 GPU 资源管理有一些基本认识。

3.1.1. CUDA 编程模型 (CUDA Programming Model)

• GPU (Graphics Processing Unit)： 图形处理器，最初设计用于加速图形渲染，但因其并行计算能力而广泛应用于通用计算（GPGPU）。

• CUDA (Compute Unified Device Architecture)： NVIDIA 推出的一种并行计算平台和编程模型，允许开发者使用 C、C++ 和 Fortran 等语言在 NVIDIA GPU 上执行高性能计算任务。

• SM (Streaming Multiprocessor，流多处理器)： GPU 中的核心处理单元，每个 SM 包含多个 core（处理核心）。所有 core 在同一个 SM 内共享寄存器集，并通过共享内存（shared memory）进行通信。

• Core (处理核心)： SM 内的最小处理单元，执行线程指令。

• Threads (线程)： GPU 上可并行执行的最小工作单元。

• Blocks (线程块)： 一组线程的集合。每个 block 内的线程可以在同一个 SM 上并发执行，并可以通过共享内存和同步机制进行协作。一个 block 通常会被调度到一个 SM 上。

• Grid (网格)： 一组线程块的集合。grid 定义了数据处理的整体结构，以及如何将其划分为线程块。一个 grid 可以是三维数组。

• GPU Kernels (GPU 内核)： 由 GPU 执行的代码，通常是并行执行在大量数据上的函数。当在 CUDA 中启动一个 kernel 时，需要指定 grid 和 block 的维度。例如，Conv (卷积) 和 MemCopy (内存拷贝) 都是常见的 kernel。

• GPU Streams (GPU 流)： 一种在 GPU 上组织和执行异步任务的方式。每个 stream 都是一系列 kernel（如 Conv、MemCopy）的序列，这些 kernel 可以独立于其他 stream 执行。同一个 stream 中的 kernel 以 FIFO（先进先出）方式执行。

  下图（原文 Figure 1）展示了 NVIDIA Jetson TX2 GPU 的布局，由两个 SM 组成，每个 SM 都能运行多个 GPU 线程，并共享全局内存。

  该图像是一个示意图，展示了NVIDIA Pascal GPU的CUDA编程范式及其计算硬件的结构。图中包括了CUDA核心、寄存器、L2缓存及全局内存的布局，说明了内核启动和线程块的关系。

Figure 1: An overview of CUDA programming paradigm and the computation hardware in NVIDIA TX2.

3.1.2. GPU 内核执行时的资源争用

当 kernel 在 SM 上执行时，它会与其他已调度到同一 SM 上并正在执行的 kernel 争用 on-SM resources（如线程槽和共享内存）。这种竞争会极大地影响 kernel 在 SM 上的执行时间。此外，一个 block 在队列中等待调度的时间，以及其在 SM 上执行工作负载的时间，共同决定了 kernel 经历的整体延迟。

• Intra-SM Contention (SM 内部争用)： 发生在单个 SM 内部的资源争用。当来自不同 kernel 的多个线程块被调度到同一个 SM 并竞争共享资源时（例如寄存器、共享内存和执行单元），就会出现 intra-SM contention。

• Inter-SM Contention (SM 之间争用)： 发生在不同 SM 之间的资源争用。当来自不同 kernel 的多个线程块被调度到不同的 SM 并竞争共享资源时（例如全局内存和内存控制器），就会出现 inter-SM contention。

  这些争用都会导致性能下降和延迟波动。

下图（原文 Figure 2）展示了 ResNet50 在与其他 DNN 模型共运行时的延迟分布，以及 intra-SM 和 inter-SM contention 的示意图。

该图像是示意图，展示了如何通过弹性内核生成器实现GPU资源的灵活重映射。左侧为原始内核结构，右侧为弹性内核，其中包含了网格和块的重新配置，以适应动态资源分配。图中展示了数据获取及其在DRAM/L2缓存中的干扰。

Figure 2: (left) The latency distribution of ResNet50 when co-running with other DNN models. (right) Illustration of intra-SM and inter-SM contention.

3.1.3. 混合关键性任务 (Mixed-criticality Tasks)

• Critical Task (关键任务)： 具有严格实时性要求的任务。例如，自动驾驶中的障碍物检测，必须在特定截止日期前完成，以留出足够时间进行车辆操作。
• Normal Task (普通任务)： 不具有严格实时性截止日期的任务。例如，监测驾驶员情绪和疲劳，可以以 best-effort（尽力而为）的方式执行，以改善驾驶体验。

3.2. 前人工作

论文回顾了当前边缘设备上 multi-DNN inference 相关的几种主要方法，并指出它们的局限性。

3.2.1. 模型压缩和新编译技术

• 联合 DNN 模型压缩 (Joint DNN model compression)： NestDNN [8] 等方法通过牺牲每个模型少量精度来减少混合 DNN 工作负载的计算成本。
  • 局限性： 牺牲精度。
  • Miriam 的不同： Miriam 不妥协精度，可作为正交方法与模型压缩结合。
• 新编译技术： Veltair [24] 提出了生成不同资源争用强度的编译 DNN 模型版本，以在运行时进行调度加速 multi-DNN inference。
  • 局限性： 存储开销高，需要离线性能分析，难以扩展到更多用例。
  • Miriam 的不同： Miriam 避免了高昂的离线分析和存储开销。

3.2.2. 算子交织与内核级抢占

• 算子交织 (Interleaving of operators)： DeepEye [25]、Abacus [6] 和 Dart [35] 等系统通过交织不同“争用通道”（memory-bound 或 compute-bound）的算子来解决问题。
  • 局限性： 需要耗时的离线性能分析，且难以推广到新的 DNN 任务。
• 内核级抢占 (Kernel-level preemption)： REEF [11] 针对混合关键性 multi-DNN inference 协调问题，实现了 kernel-level preemption（内核级抢占）。
  • 局限性： 需要修改 GPU 驱动库，这在许多流行的闭源设备（如 NVIDIA Jetson）上不切实际。
• 其他资源争用解决： Heimdall [38] 和 Band [19] 也旨在解决 multi-DNN inference 的资源争用问题，但它们的应用场景与 Miriam 不同。

3.2.3. 性能建模与资源管理

• 性能与计算单元占用率曲线： Warped-Slicer [36] 使用性能与计算单元占用率曲线来选择优化的并发内核模式。
  • 局限性： 未能解决内核间的资源争用。
• 并发 GPU 内核执行延迟建模： HSM [40] 和 [31] 基于硬件信息对并发 GPU kernel 执行的延迟下降进行建模。
  • 局限性： 这些工作中构建的预测器难以适应实际场景中非确定性 kernel overlapping 的 multi-DNN inference 场景。
• 资源管理： Smcentric [33] 和 Effisha [4] 从资源管理角度通过 space-multiplexing（空间复用）方式 [18, 34] 解决 GPU 多任务问题。
  • Miriam 的不同： Miriam 采用的 elastic kernel 方法与 space-multiplexing 是正交的。

3.3. 技术演进

从传统的 GPU 计算模型来看，通常假设 GPU 有足够的资源来独立或简单地调度任务。然而，随着深度学习模型在边缘设备上的普及，以及这些设备资源受限的特点，简单的调度或依赖服务器级硬件特性的方案不再适用。领域内的技术演进逐渐从模型层面的优化（压缩）转向更底层的运行时调度和资源管理，以适应边缘设备对实时性和效率的双重需求。Miriam 正是处于这种演进的脉络中，通过编译器-运行时协同设计，尝试在不修改底层硬件驱动的情况下，实现细粒度的资源控制和动态调度。

3.4. 差异化分析

Miriam 的核心创新和差异化在于：

• 无需牺牲精度： 与模型压缩方法不同，Miriam 不影响 DNN 模型的精度。
• 避免离线分析和驱动修改： 与需要大量离线分析或修改底层 GPU 驱动的方法不同，Miriam 的 elastic kernel 生成器和运行时协调器在不依赖这些复杂机制的情况下工作。
• 细粒度资源管理： 提出了 elastic kernel 概念，能够同时解决 intra-SM 和 inter-SM 两种类型的资源争用，这是许多现有方法未能全面处理的。通过调节 grid size 和 block size，Miriam 实现了对 GPU 资源的更精细控制。
• 编译器-运行时协同设计： 通过离线的 source-to-source kernel transformer 和在线的 dynamic kernel coordinator 的协同工作，实现了对 kernel 的灵活控制和动态调度。
• 针对边缘 GPU 优化： 明确地针对边缘 GPU 的资源受限特性进行设计，而许多现有方案（如 NVIDIA MPS 和 MIG）在边缘平台上不可用。

4. 方法论

4.1. 方法原理

Miriam 的核心思想是引入“弹性内核”（elastic kernel），这是一种具有可调节 grid size 和 block size 的 GPU 内核。通过将传统内核分解成更小、更灵活的单元，并根据优先级和关键性动态地调度和重新映射这些单元，Miriam 能够在不引起显著资源争用的情况下，将其他 GPU 内核“填充”到空闲资源中。这使得关键任务能够优先执行，同时最大化 GPU 利用率和整体吞吐量。

其设计基于一个关键观察：并发 DNN 内核的延迟下降主要由两种主导因素引起：intra-multi-processor (SM) 资源争用和 inter-multi-processor 资源争用。Miriam 利用弹性内核来解决这两种资源争用。

4.2. 核心方法详解

Miriam 框架整合了两个主要组件：elastic-kernel generator（弹性内核生成器）和 runtime dynamic kernel coordinator（运行时动态内核协调器），以支持混合关键 DNN 推理。

下图（原文 Figure 3）展示了 Miriam 的系统架构。绿色框是贡献部分。虚线框内的模块是处理 DNN 推理请求的关键路径上的。

该图像是一个示意图，展示了物理线程和逻辑线程之间的关系，以及如何通过共享内存和寄存器进行矩阵乘法优化。图中强调了 Tile ext{ }size riangleq TB ext{ }size 的概念，并显示了不同内存层次结构的布局。

Figure 3: Architecture of Miriam. Green boxes are our contributions. Modules in boxes with dashed border on the right are on the crucial route for handling DNN inference requests. The source-to-source kernel transformer module enables the elastic kernel design without incurring computation inconsistency.

4.2.1. 离线弹性内核生成 (Offline Elastic Kernel Generation)

此阶段负责将原始的 GPU 内核转化为弹性内核，使其具备细粒度资源控制的能力。它包含两个子组件：elastic grid/block generator 和 source-to-source kernel transformer。

4.2.1.1. 通过弹性块实现可控的 SM 内部资源 (Controllable Intra-SM Resource by Elastic Block)

intra-SM contention 是指多个线程块在同一个 SM 内部争用寄存器、共享内存和执行单元等资源。为了解决这个问题，Miriam 调整目标线程块内的线程数量，为每个线程块生成 elastic blocks（弹性块）。

• 传统内核与问题： 传统内核中，线程在完成执行后终止。当多个线程块从不同内核竞争同一 SM 内的资源时，会导致部分块执行失败或延迟，降低整体吞吐量和增加延迟。
• 优化方法： 传统的优化方法包括代码级优化（如内存访问模式优化、减少不必要计算），以降低 intra-SM 资源使用。然而，这耗时、难以泛化，且需要平衡执行性能、内存使用和代码复杂性。
• Miriam 的解决方案： 采用 persistent thread（持久线程）技术 [10]，该技术能够调整 kernel 在 SM 上的常驻块大小。与传统内核不同，persistent thread 在 kernel 函数执行期间保持活跃。Miriam 将每个 elastic block 的大小限制在 1 到最大常驻块大小之间。同时，将默认的 1:1 逻辑到物理线程映射方案转换为 N:1 映射方案，同时保持初始程序语义。
• 优点： 相比静态块融合（static block fusion）[32]，persistent thread 设计不需要预编译所有可能的内核组合，这使得 SM 级别的资源映射在运行时更加灵活。
• 内存占用： elastic kernel 设计在共享内存限制内，通过修改 intra-SM 资源（包括共享内存）的控制方式，其内存占用与原始内核相同或更少。
• 挑战与设计空间缩小： persistent thread 机制虽提供了细粒度的 intra-SM parallelism 控制，但有不可忽略的开销。最优启动的 persistent thread 数量不总是等于单个 SM 能负担的最大并发线程数。因此，需要缩小 elastic block 的设计空间（详见 4.2.1.3）。

4.2.1.2. 用于 SM 之间争用的弹性网格 (Elastic Grid for Inter-SM Contention)

inter-SM contention 是指 DRAM、L2 Cache 等全局资源在多个 SM 之间被争用。DNN kernel 通常使用大量 block 来隐藏数据访问造成的停顿周期。当多个 DNN inference request 快速连续到达时，多个 SM 争用总线资源，导致性能下降。

• Miriam 的解决方案： elastic grid generator（弹性网格生成器）将初始 grid 切片成多个更小的 grid。

• 核心思想： 这种方法通过允许跨多个 SM 进行更高效的内存访问，从而提高资源利用率并减少 inter-SM contention。

• 内核切片计划 (Kernel Slicing Plan)： 对于给定内核 $K$，一个切片计划 P(K) 将 $K$ 切片成 $n$ 个切片序列 $[s_0, s_1, s_2, ..., s_{n-1}]$，基于线程块粒度划分。

• 切片方案公式： 对于具有 $M$ 个线程块的 DNN kernel $K$，可以使用基于二分算法的切片计划 S(K)：

  $$S ( K ) = ( { \frac { M } { 2 ^ { n } } } , { \frac { M } { 2 ^ { n - 1 } } } . . . , M ) , n = \operatorname* { m a x } _ { i } \{ M m o d 2 ^ { i } = 0 \}$$

  • 符号解释：
    • S(K): 内核 $K$ 的切片计划。
    • $M$: 内核 $K$ 中初始线程块的总数。
    • $n$: 一个整数，表示 2 的幂次，使得 $M$ 可以被 $2^n$ 整除。$\operatorname*{max}_i \{ M \pmod{2^i} = 0 \}$ 表示找到满足 $M$ 能被 $2^i$ 整除的最大 $i$ 值，即 $n$ 为 $M$ 所含 2 的最大因子幂次。
    • 这个公式表示切片计划将 $M$ 分割成一系列块数，例如 $M/2^n, M/2^{n-1}, ..., M/2^1, M$。这使得可以在 SM 上以灵活的线程块数量发出 normal kernel，并与 critical kernel 共存。

• 优点： 通过将单个 kernel 划分为多个，这些切片的 grid 可以由 GPU 独立调度运行，允许 GPU 将它们的执行与其他 critical kernel 的执行交织，从而减少 co-locating kernel 的 inter-SM memory contention。

4.2.1.3. 基于工作负载均衡的设计空间缩小 (Workload-balanced-guided Design Space Shrinking)

elastic kernel 的执行参数包括 grid 数量 ($N_{blk\_be}$) 和 block size ($S_{blk\_be}$)。每个参数对被称为一个调度。由于可行的调度数量巨大，难以在运行时枚举或启发式地找到最优调度。

• 挑战： 针对 AlexNet 模型，对于所有 Conv kernel，可行的调度数量可达 $2.2 \times 10^{25}$ [21]。

• Miriam 的解决方案： 通过消除因严重资源争用可能导致调度失败的 elastic grid 和 block size 组合来缩小设计空间。换句话说，Miriam 移除那些预期性能较低的配置。

• 两个约束条件 (Eq. 2)： 当多个 kernel 共运行时，来自不同 kernel 的线程块可能导致 SM 级别的争用或效率低下。Miriam 提出以下两个约束来解决这些问题：

  $$\left\{ \begin{array} { l l } & { N _ { b l k \_ b e } \leqslant N _ { S M } - \ N _ { b l k \_ r t } \ m o d \ N _ { S M } } \\ & { S _ { b l k \_ b e } \leqslant L _ { t h r e a d s } - \ b l k \_ s i z e _ { r t } } \end{array} \right.$$

  • 符号解释（参考 Table 1）：
    • $N_{SM}$: GPU 上的流多处理器 (Streaming Multiprocessor) 数量。
    • $N_{blk\_rt}$: 已调度 critical kernel 中的线程块数量。
    • $N_{blk\_be}$: 已调度 elastic normal kernel 中的线程块数量。
    • $S_{blk\_rt}$: 已调度 critical kernel 中每个线程块启动的工作线程数量。
    • $S_{blk\_be}$: 已调度 elastic normal kernel 中每个线程块的工作线程数量。
    • $L_{threads}$: 工作线程数量的限制。
  • 第一个约束： 解决 SM 之间工作负载不平衡问题。当线程块数量不是 GPU 内部 SM 数量的倍数时，会出现这种不平衡。这个约束剪除了 elastic kernel 的线程块数量超过在调度所有 critical kernel 的线程块后剩余可用 SM 数量的情况。
  • 第二个约束： 解决 SM 内部工作负载平衡问题。旨在减少来自不同 kernel 的线程块在 SM 内部竞争资源时的争用。它剪除了 elastic kernel 的工作线程数量在被 critical kernel 的块占用后，过多超出剩余 intra-SM 资源的情况。

• 工作负载不平衡指标 WIScore (Eq. 3)： 为了量化这两个效率低下情况，定义 WIScore 如下： $$W I S c o r e = \frac { N _ { b l k _ { - } r t } \ m o d \ N _ { S M } + N _ { b l k _ { - } b e } } { N _ { S M } } * \frac { S _ { b l k _ { - } b e } + S _ { b l k _ { - } b e } } { L _ { t h r e a d s } }$$

  • 符号解释：
    • WIScore: 工作负载不平衡得分，值域为 [0, 1]。
    • $N_{blk\_rt} \pmod{N_{SM}}$: 已调度 critical kernel 的线程块数量除以 SM 数量的余数，表示 critical kernel 对 SM 资源的占用不平衡部分。
    • $N_{SM}$: GPU 上的流多处理器数量。
    • $N_{blk\_be}$: elastic normal kernel 中的线程块数量。
    • $S_{blk\_be}$: elastic normal kernel 中每个线程块的工作线程数量。
    • $L_{threads}$: 工作线程数量的限制。
    • 公式左侧部分表示 SM 间负载分布，右侧部分表示 SM 内负载分布。
    • 注意：公式中的 $S_{blk\_be} + S_{blk\_be}$ 疑似排版错误，通常应表示与 $S_{blk\_rt}$ 的关系，或仅为 $S_{blk\_be}$。但根据指令，此处严格忠实原文，不作修改，并解释为对 $S_{blk\_be}$ 的某种聚合或重复计算。

• 调度开销指标 OScore (Eq. 4)： 另一个考虑因素是 elastic kernel 的调度开销。为确保生成的潜在调度方案可行且不违反关键决策要求，Miriam 使用 OScore 进行剪枝： $$O S c o r e = \left\{ \begin{array} { r l } & { 1 \sum L O _ { b l k } ( k _ { b e \_ i } ) < M A X _ { b l k } , \forall i \in [ 1 , N _ { s h a r d } ] } \\ & { a n d \sum L O _ { p t } ( k _ { b e \_ i } ) < M A X _ { p t } , \forall i \in [ 1 , N _ { s h a r d } ] } \\ & { 0 \ O t h e r w i s e } \end{array} \right.$$

  • 符号解释：
    • OScore: 开销得分，取值为 0 或 1。
    • LO(): 表示启动开销的函数，其值为每个 elastic kernel fragment 的启动时间之和减去初始 normal kernel 的启动时间。
    • $k_{be\_i}$: 第 $i$ 个 elastic kernel fragment。
    • $MAX_{blk}$: 块的最大可接受启动开销。
    • $MAX_{pt}$: persistent thread 的最大可接受启动开销。
    • $N_{shard}$: elastic kernel 的分片数量。
    • 当开销超过预设的最大可接受阈值时，OScore 设置为 0。

• 设计空间导航： WIScore 和 OScore 的乘积 (WIScore * OScore) 作为设计空间缩小的导航器。Miriam 为每个 elastic kernel implementation setting 的可能组合计算这个乘积。通过测试代表性的张量操作（如 CifarNet 中的卷积和 GRU 中的矩阵乘法），Miriam 选取所有候选组合中性能排名前 20% 的作为下一阶段运行时内核协调的候选。

4.2.1.4. 源到源弹性内核转换器 (Source-to-Source Elastic Kernel Transformer)

直接修改 DNN kernel 的 grid 或 block size 可能会导致计算错误，因为这会影响线程的组织和执行方式。实验表明，Tango 基准中只有 7.4% 的内核可以直接调整 grid/block size 而无需修改内部计算调度。

下图（原文 Figure 6）展示了 grid/block size 不能直接修改可能导致重新计算/计算错误的情况。

该图像是图表，展示了多流和Miriam方法下两种AlexNet模型的活跃内核级时效性以及各层的平均占用率。上部分比较了多流和Miriam的时效性表现，底部则展示了不同层中取得的占用率，表明Miriam方法有效减少了延迟。

Figure 6: Grid/Block size cannot be directly modified in case of recomputation/computation error.

• Miriam 的解决方案： 设计了一个 source-to-source kernel transformer（源到源内核转换器），通过在每个 kernel 开始处注入一段代码来等效地转换 DNN kernel。
• 工作原理： 这段代码检查计算和内存偏移量，以确定 kernel 在被逐出后从何处开始和结束。具体来说，计算一个 global thread identifier（全局线程标识符），并将其作为 SM 级别工作负载分布的基础。这个标识符以 thread ID 为输入，生成线程访问的数据元素的相应索引。Miriam 用逻辑等效项替换原始 kernel 代码中关于物理线程（例如 GridDim）和标识符变量（例如 threadIdx.x）的引用。
• 索引函数实现：
  • 计算型 (Computation-based)： 在 kernel 内部计算索引，当线程访问相应数据元素时。
  • 内存型 (Memory-based)： 在 kernel 启动前在主机端（CPU）预计算索引，并存储在共享内存中供 kernel 执行时使用。

4.2.2. 运行时动态内核协调 (Runtime Dynamic Kernel Coordination)

此阶段负责在线调度 elastic kernel，以最大化整体实时性能并缓解资源争用。每个 elastic kernel（即 elastic grid 和 elastic block）被称为一个 elastic kernel shard（弹性内核分片）。

• 指导原则： 最大化整体实时性能和缓解资源争用。
• 工作机制： 运行时协调器持续监控可用的 GPU 资源（包括 critical kernel 和 elastic kernel 的资源）。然后，它决定哪些 elastic kernel shard 可以与 critical kernel 有效地共运行。
• 共运行内核的执行时间线：

  1. 接收到多个 normal task request ($b_1...b_n$) 后，Miriam 将所有 kernel 推入 normal tasks queue，并通过多个 stream 分派到 GPU。

  2. 当 critical task 到达时，Miriam 会立即以“bin-packing”（装箱）的方式，根据当前的 intra-SM 和 inter-SM 级别资源分布，选择后续 normal kernel 的适当 elastic kernel fragment。

  3. critical kernel 执行完成后，来自 normal task 的所有 kernel 将重新占用 GPU。 下图（原文 Figure 4）展示了 Miriam 中的时间线示例。

     该图像是一个示意图，展示了用于内核碎片形成的阴影二叉树构建。图中标注了关键内核、正常内核及其在不同相位的调度过程，同时展示了网格大小和分片程度。弹性块大小（Elastic Block Size）和可能的碎片（Possible Shard）与实际碎片（Actual Shard）也被说明。

Figure 4: An example of timeline in Miriam. The yellow star represents the critical kernel, while the green one represents the normal kernel.

• 弹性内核的网格/块大小确定：
  • 挑战： 运行时 elastic grid 和 block setting 的固定大小会随着与不同 critical kernel 共调度的 elastic kernel shard 的变化而变得低效。例如，如果一个 critical kernel 完成了，但 co-locating elastic kernel 还有一半计算未完成，剩余的线程块在与后续 critical kernel 共存时可能导致严重的资源争用或利用率不足。
  • 传统方法的局限性： 通过运行时性能事件（如缓存未命中和全局内存带宽）构建算子组持续时间预测模型来控制内核重叠 [12, 37]，但这些运行时事件在 Nvidia Jetson 等边缘 GPU 上不支持，且工具报告的硬件事件开销高。
• Miriam 的解决方案： 采用 greedy scheduling policy（贪婪调度策略）。
  • 当 elastic kernel 与 critical kernel 部分重叠时，协调器必须仔细平衡分配给每个 kernel 的资源。确保填充的 elastic kernel 不干扰 critical kernel 的执行，同时尽可能多地利用可用资源。
  • 当填充的 kernel 独立运行时，协调器可以将其所有可用资源分配给它，因为没有其他任务在 GPU 上运行。
• 动态大小的阴影二叉树方法 (Dynamic-sized Shade Binary Tree Approach)： 为了高效管理这些 elastic kernel 并实现目标，Miriam 提出了 dynamic-sized shade binary tree（动态大小的阴影二叉树）方法用于 elastic kernel shard 的形成，以实现高运行时效率和低资源争用。

  • 结构抽象： 类似于 complete binary tree（完全二叉树）的 shard 结构。树的根代表来自 normal task 的 kernel，其初始 grid size 为 $M$。

  • 节点与边： 每个节点对应一部分计算，或 kernel 内部待调度的潜在线程块。每个节点的“阴影属性”是线程块的 elastic block size。有向边表示前任遗留的未完成计算的潜在切片对等体。

  • 组成： 由 actual shard（实际分片，最终被调度的 elastic kernel shard）和 virtual shard（虚拟分片，不会被调度的 elastic kernel 的潜在片段）组成。

  • 策略： 当 critical kernel 和 normal kernel 都存在时，协调器根据来自 critical kernel 和 normal kernel 的线程块数量，利用动态阴影内核二叉树结构来选择 elastic kernel shard。策略是从阴影内核二叉树的头部选择一组 elastic block，与常驻 critical kernel 的 co-locating thread block 共享 SM 级别资源，且只使用 critical kernel 剩余的资源。

    下图（原文 Figure 7）展示了用于内核碎片形成的阴影二叉树构建。

    该图像是示意图，展示了不同 DNN 模型的弹性候选设计空间收缩情况。图中标出了帕累托前沿（Pareto Frontier）、弹性内核候选（Elastic Kernel Candidates）和不合格候选（Unqualified Candidates），并且提供了在修剪后的空间中各个模型的对应效果。该图为理解 Miriam 整体框架的性能优化提供了可视化参考。

Figure 7: Shaded Binary Tree Construction for Kernel Shards Formations. ES refers to the elastic kernel, and EBS refers to the elastic block size. The sharding degree represents the degree of elastic kernel splitting depth.

5. 实验设置

5.1. 数据集

论文使用了名为 MDTB (Mixed-critical DNN Task Benchmarks) 的新基准测试，该基准测试基于 CUDA 实现的内核，以充分展示 Miriam 框架的性能和泛化能力。

• DNN 模型： 选择了六个流行的 DNN 模型，涵盖计算机视觉和语言处理领域，包括：AlexNet [22]、SqueezeNet [15]、GRU [7]、LSTM [14]、ResNet [13] 和 CifarNet [30]。所有模型均使用 CUDA 实现。
• 输入数据： 神经推理使用 $224 \times 224 \times 3$ 的单批图像作为输入，以模拟真实应用中的推理场景。
• 推理任务模式： MDTB 模拟了三种用户请求的推理任务模式：

  1. 均匀分布到达 (Uniform distribution arrival)： 客户端以固定频率（例如 10 reqs/s）发送推理请求，模拟如姿态估计等关键应用。

  2. 泊松分布到达 (Poisson distribution arrival)： 模拟事件驱动型应用，例如障碍物检测。

  3. 闭环工作负载 (Closed-loop workloads)： 客户端持续发送推理请求。

     以下是原文 Table 2 中 MDTB 工作负载的详细描述：

     MDTB	A	B	C	D
     Critical Task	AlexNet	SqueezeNet	GRU	LSTM
     Frequency (req/s)	Closed-loop Uniform (10 reqs/s) Poisson (10 reqs/s) Uniform (10 reqs/s)
     Normal Tasks	CifarNet	AlexNet	ResNet	SqueezeNet
     Frequency (req/s)	Closed-loop	Closed-loop	Closed-loop	Closed-loop

5.2. 评估指标

论文使用以下三个指标来评估 Miriam 的性能：

5.2.1. 关键任务的端到端延迟 (End-to-end Latency of Critical Tasks)

• 概念定义： 此指标衡量具有实时性要求的关键任务从开始到完成的整个推理过程所需的时间。它关注的是关键任务的及时完成能力，直接反映了系统对实时性要求的满足程度。
• 数学公式： 无标准化公式提供，通常指任务从提交到结果返回的总时间。
• 符号解释： 无特定符号。

5.2.2. 整体吞吐量 (Overall Throughput)

• 概念定义： 此指标表示 Miriam 在目标边缘 GPU 上每秒能够服务的用户请求数量。它衡量了系统的整体效率和处理能力，反映了在给定时间内系统能够完成的工作量。
• 数学公式： 无标准化公式提供，通常计算为： $$\text{Throughput} = \frac{\text{Total completed requests}}{\text{Total time}}$$
• 符号解释：
  • $\text{Total completed requests}$: 在测量时间内完成的请求总数。
  • $\text{Total time}$: 测量持续的总时间。

5.2.3. 实现占用率 (Achieved Occupancy)

• 概念定义： 根据定义，achieved occupancy 是 SM 上活跃 warp 的平均比例与 SM 支持的最大活跃 warp 数量之比。它是一个细粒度的 GPU 利用率指标，反映了 SM 在执行 kernel 时的忙碌程度。高占用率通常意味着资源利用更充分，但并非 100% 占用率总是代表最佳性能，因为过高的占用率可能导致资源争用和调度开销。
• 数学公式： $$AchievedOccupancy = \frac { Active\_warps/Active\_cycles } { MAX\_warps\_per\_SM }$$
• 符号解释：
  • $Active\_warps$: 在给定周期内处于活跃状态的 warp 数量。
  • $Active\_cycles$: SM 处于活跃状态的周期数。
  • $MAX\_warps\_per\_SM$: 单个 SM 支持的最大活跃 warp 数量。

5.3. 对比基线

论文将 Miriam 与以下几种 DNN 调度方法在边缘 GPU 上进行了比较：

• Sequential (顺序执行)：
  • 描述： 从 critical task 和 normal task 队列中以 round-robin（轮询）方式选择模型，并逐个执行推理。
  • 特点： 关键任务独立运行，占用 GPU 资源，因此能为关键任务提供最优的端到端延迟。
• GPU Multi-stream with Priority (带优先级的 GPU 多流)：
  • 描述： 同时将来自 critical task 和 normal task 的 kernel 入队，模型并行执行。这是 NVIDIA Triton [3] 采用的一种常见方法。
  • 特点： 试图通过并发来提高吞吐量，但缺乏精细的争用管理，可能导致关键任务延迟增加。
• Inter-stream Barrier (IB，流间屏障)：
  • 描述： 最先进的基于 multi-stream 的 multi-DNN operator scheduling 方法 [39]。它使用 inter-stream barrier 手动同步不同 kernel 之间的 kernel dispatch。
  • 特点： 通过流和同步机制控制 kernel 之间的并发性。

5.4. 实现和测试平台

• 实现： Miriam 基于 NVIDIA CUDA 11.2 实现，用于 elastic kernel generation 和 online kernel scheduling。source-to-source kernel transformer 则使用 Python3.6 实现。
• 测试平台：
  • NVIDIA GeForce RTX 2060：配备 1920 个 CUDA 核心。
  • NVIDIA Jetson AGX Xavier：配备 Pascal GPU architecture 和 256 个 NVIDIA CUDA 核心。
• 可扩展性： 论文指出 Miriam 具有可扩展性，可以在官方支持 OpenCL、HIP 或其他类似 CUDA 编程范式的 GPU 平台（如 AMD Embedded RadeonTM E9170）上良好工作。

6. 实验结果与分析

6.1. 核心结果分析

为了反映系统在保证关键任务实时性能的同时提高整体吞吐量的效果，论文将 Miriam 与其他 GPU 调度方法在 MDTB A-D 工作负载下进行了比较。

6.1.1. 整体性能 (Overall Performance)

论文在 RTX 2060 和 Jetson AGX Xavier 两个平台进行了实验。

6.1.1.1. 闭环关键任务 (MDTB A)

• 工作负载： AlexNet 作为 critical task，CifarNet 作为 normal task，两者均为闭环工作负载。

• 性能对比：

  • 关键任务延迟： 相较于 Sequential 方法，Multi-stream 和 IB 在 2060 平台上分别使关键任务延迟增加了 $1.95 \times$ 和 $1.52 \times$，在 Xavier 平台上则分别增加了 $2.02 \times$ 和 $1.77 \times$。而 Miriam 仅对关键任务造成 21% (2060) 和 28% (Xavier) 的开销。
  • 整体吞吐量： Miriam 在两个平台上分别将整体吞吐量提高了 64% (2060) 和 83% (Xavier)，显著优于其他方法。
  • IB 的吞吐量问题： IB 的吞吐量甚至低于 Sequential，这是因为频繁的关键任务请求需要插入更多同步屏障来管理内核组，导致显著开销。

• 占用率 (Achieved Occupancy)： Miriam 相比其他基线方法，实现了更高的 SM 级别 GPU 资源利用率。论文指出，对于 DNN 推理任务，由于其大型线程块容易导致资源闲置或 SM 无法覆盖内存访问延迟，因此很难达到接近 100% 的理论占用率 [20]。

  以下是原文 Figure 8 展示了在 2060 和 Xavier 平台上，不同 GPU 调度方法在 MDTB A 工作负载下的关键任务延迟、整体吞吐量和平均实现占用率的比较。

  该图像是实验结果的图表，展示了在2060平台上，Miriam方法在端到端关键任务延迟、整体吞吐量以及平均实现占用率上的表现。图(a)显示关键任务的标准化延迟，图(b)展示整体吞吐量，而图(c)则表现了实现的占用率。Miriam方法在延迟和吞吐量方面均优于其他基准。

Figure 8: Performance comparison of critical task latency, overall throughput, and average achieved occupancy among different GPU scheduling approaches under MDTB A workload on 2060 and Xavier platforms.

6.1.1.2. 均匀/泊松分布关键任务 (MDTB B, C, D)

• 工作负载： 关键工作负载的启动频率降低（MDTB B, C, D），为 normal task 提供了更多共享 GPU 资源的机会。
• 性能对比（以 Xavier 平台为例）：
  • 整体吞吐量： Miriam 在 MDTB B、C 和 D 上，相比 Sequential 分别将整体吞吐量提高了 $1.85 \times$、$1.79 \times$ 和 $1.91 \times$，远超其他基线。Multi-stream 和 IB 也提高了吞吐量（例如 $1.34 \times$ 到 $1.73 \times$），但对关键任务造成了严重的延迟下降（32% 到 88%）。
  • 关键任务延迟： Miriam 对这些基准的关键任务仅造成不到 21% 的延迟开销。
• GPU 利用率： Miriam 增加了每个周期的平均活跃 warp 数量，从而提高了 SM 利用率。

6.1.1.3. Xavier 平台上的性能观察

论文观察到 Miriam 在 Jetson Xavier 上的性能提升可能不总是导致更高的 SM 占用率。

• 原因： Xavier 相比 2060 具有更少的板载资源和更少的 SM 数量。其相对较低的内存带宽也可能限制内存和 SM 之间的数据传输量，导致复杂模型出现性能瓶颈。
• 热设计功耗 (TDP)： Xavier 的 TDP 也较低，限制了 GPU 功耗和产生的热量。这可能负面影响处理器核心的时钟速度和可实现的并行度，进而影响 SM 占用率与性能之间的关系。

6.1.2. Miriam 的深入分析 (In-depth Analysis of Miriam)

为了更好地理解 Miriam 在严重争用情况下表现优于其他 GPU 调度方法的原因，论文进行了深入分析，选取两个 AlexNet 模型在单个 2060 GPU 上共同运行，其中一个 AlexNet-C 作为 critical task，另一个 AlexNet-N 作为 normal task。两者都以闭环方式启动。

• 时间线分析 (Figure 9 上部)：
  • 图示基于 NVIDIA Nsight Sys [2] 的实际分析结果。蓝色代表 critical task，绿色代表 Multi-stream 启动的 normal task，粉色代表 Miriam 的 normal task 弹性内核。
  • 结果显示，粉色块明显多于绿色块，且这些粉色块与蓝色块紧密填充，这展示了 elastic kernel shard 与 critical kernel 的填充效果。
  • Miriam 中 AlexNet-C 的端到端延迟远低于 Multi-stream。
• 占用率分析 (Figure 9 下部)：

  • Miriam 的平均 layer-wise achieved occupancy 为 65.25%，而 Multi-stream 为 32.9%。

  • 这证实了每个周期更多的平均活跃 warp 和更少的争用开销是提高并行度并保持关键任务速度的关键。

    以下是原文 Figure 9 展示了两个共运行的 AlexNet 模型的活跃内核级及时性（上部）和关键 AlexNet 每层的平均实现占用率（下部）。

    该图像是图表，展示了从LGSVL模拟器收集的真实轨迹，其中(a)为基于图像数据的目标检测结果，(b)为基于激光雷达点云数据的结果，(c)展示了我们收集轨迹的设置。

Figure 9: (Upper) The active kernel-level timeliness of two co-running AlexNet models with mixed-criticality, which is profiled from the Nsight System. (Lower) The average achieved occupancy for each layer of the critical AlexNet.

6.1.3. 设计空间缩小的评估 (Evaluations on Design Space Shrinking)

Miriam 通过应用硬件限制器，过滤掉了 80% 的确定性慢速情况（如 6.3 节所述）。论文进一步探讨了 elasticized scale（动态阴影二叉树的深度）与调度粒度之间的权衡，以指导进一步的缩小过程。

• 缩小效果： 图示（原文 Figure 10）总结了 MDTB 中模型候选弹性内核的剪枝空间，范围从 84% 到 95.2%。

• 原因： 剪枝空间可能因模型复杂度（即使用的算子类型）和输入大小等多种因素而异。

• 有效性： 通过这些测试，论文没有发现任何模型剪枝掉最佳性能参数集的情况。

  以下是原文 Figure 10 展示了不同 DNN 模型弹性候选的设计空间缩小情况。

  该图像是图表，左侧展示了不同深度神经网络（DNN）模型在共运行时的延迟分布，包含模型如ResNet50、VGG11、VGG16、MobileViT、ResNet18和Inception_v3的CDF曲线。右侧则以示意图形式展示了单线程块和两个内核之间的块，在内存调度中可能产生的争用因素，包括“内部”和“外部”争用因素，突出了资源的低利用率和任务调度失败的问题。

Figure 10: Shrinking the design space of elastic candidates for different DNN Models. Miriam picks up elastic kernels lying in the Pareto Frontier (for visualization) of tradeoff space between the elasticized scale and the scheduling granularity.

6.1.4. 案例研究：LGSVL 自动驾驶 (Case Study: Autonomous Driving with LGSVL)

论文使用来自开放自动驾驶平台 LGSVL [16] 的真实世界轨迹作为工作负载进行案例研究，模拟了自动驾驶中 critical task（障碍物检测）和 normal task（姿态估计）的实际到达分布。

• 工作负载：

  • Normal task：SqueezeNet，模拟姿态估计（激光雷达数据），频率 $12.5 \mathrm{Hz}$。
  • Critical task：ResNet，模拟障碍物检测（摄像头数据），频率 $10 \mathrm{Hz}$。
  • 客户端以均匀分布发送推理请求。
  • 实验平台：GTX 2060。

• 实验结果 (Figure 11)：

  • 整体吞吐量： Multi-stream 和 IB 相比 Sequential 分别提高了 $1.41 \times$ 和 $1.25 \times$ 的整体吞吐量。
  • 关键任务延迟： 但同时将 critical task 延迟分别增加了 82% 和 56%。
  • Miriam 性能： Miriam 相比 Sequential 实现了 89% 的整体吞吐量提升，而对 critical task 仅引入 11% 的延迟开销。
  • 高 SM 占用率： Figure 12 (c) 中展示的高 SM 占用率也证实了这一点。

• 结论： 这证明了 Miriam 能够通过 elastic kernel 设计在对 critical task 延迟影响很小的情况下，大幅提高吞吐量。

  以下是原文 Figure 11 展示了使用 LGSVL 模拟工作负载时，Miriam 与不同调度方案的性能比较。

  该图像是Miriam架构的示意图。左侧展示离线阶段的模块，包括弹性块生成器与源到源内核变换器，右侧为在线阶段的关键任务调度。各模块通过优先级流调度机制连接，以便高效处理正常与关键任务的DNN推理请求。

F1o--elieskate uvehp different scheduling schemes with our LGSVL simulated workloads. 以下是原文 Figure 12 展示了从 LGSVL 模拟器收集的真实轨迹。

该图像是Miriam中的时间线示例，展示了关键内核和普通内核的调度。黄色星星表示关键内核，绿色星星代表普通内核。图中还显示了线程块的分配和GPU的占用情况，关键模式下的协调过程。该框架通过弹性内核生成和在线内核协调提高了边缘GPU的资源管理效率。

Figure 12: Real-world trace collected from LGSVL simulator, where (a) is the object detection result based on image data, (b) is the result with lidar point cloud data, and (c) depicts the setting of our collected trace.

6.1.5. 系统开销 (System Overhead)

Miriam 的调度开销主要由两部分组成：

1. 运行时弹性内核分片选择 (Runtime elastic kernel shards selection)：
   • 开销： 扫描分片候选的复杂度为 $O(N)$。由于 Miriam 中调度机制的低复杂度，每个 DNN 模型服务的平均总开销小于 $0.35 \mathrm{ms}$。
2. 关键内核的启动时间开销 (Launch time overhead for critical kernels)：
   • 开销： 这是由于 elastic kernel 的填充造成的。在大多数（超过 80%）情况下，开销小于 $15 \mathrm{\mu s}$。
   • 主要原因： 这种延迟开销主要源于 texture cache 和 L2 memory 上的争用，这被列为未来的工作。

6.2. 数据呈现 (表格)

以下是原文 Table 2 的结果：

6.3. 消融实验/参数分析

论文中没有明确的“消融实验”章节，但设计空间缩小 (workload-balanced-guided design space shrinking) 的评估（Section 8.4）以及 in-depth analysis 中对 elastic kernel shards 填充效果的展示（Section 8.3）可以视为对 Miriam 核心组件有效性的验证。

• 设计空间缩小： 证明了通过 WIScore 和 OScore 剪枝大部分低性能配置是有效的，在不牺牲最优解的情况下，大幅减少了运行时搜索空间（剪枝 84% 到 95.2% 的候选）。
• 弹性内核填充： Figure 9 清晰展示了 Miriam 如何通过紧密填充 elastic kernel shards 来提高 SM 占用率并降低关键任务延迟，这证明了 elastic kernel 和 dynamic kernel coordinator 的有效性。

7. 总结与思考

7.1. 结论总结

论文提出了一种名为 Miriam 的新型系统，旨在解决边缘 GPU 上并发运行多个深度神经网络（DNN）推理任务所面临的延迟和吞吐量问题。该系统通过引入“弹性内核”（elastic kernel）实现了细粒度的 GPU 资源重映射，并通过一个“运行时动态内核协调器”（runtime dynamic kernel coordinator）支持动态的 multi-DNN inference task。在作者构建的基准测试和两个边缘 GPU 平台上的实验结果表明，与仅将 GPU 用于关键任务的基线方法相比，Miriam 能够显著提高整体系统吞吐量（最高达 92%），同时仅对关键任务引入极小的延迟开销（不到 10%）。这证明了 Miriam 在平衡实时性要求和资源利用效率方面的卓越性能。

7.2. 局限性与未来工作

论文作者指出了 Miriam 的一些局限性以及未来的研究方向：

• 可扩展性 (Scalability)：
  • Miriam 目前主要针对成对的 DNN 任务共运行场景进行了深入研究。虽然有潜力扩展到更通用的多任务场景，但需要额外的考虑。
  • 挑战包括为具有相同优先级的 normal task 建立调度策略，以及寻找高效的离线内核分析方法，因为设计空间将呈指数级增长。
• 特定资源争用 (Specific Resource Contention)：
  • 在系统开销分析中提到，elastic kernel 填充导致的关键内核启动时间开销（尽管很小，<15 us）主要源于 texture cache 和 L2 memory 上的争用。这部分争用机制的解决留待未来的工作。
• 与 DNN 编译器的集成 (Integrated with DNN Compiler)：
  • TVM [5] 等代表性 DNN 编译器可以通过 auto-tuning [42] 生成高性能、低延迟的 DNN 内核。然而，它们是离线编译且耗时，生成的内核难以在运行时修改。
  • Miriam 可以作为 post-compiling runtime，填补静态编译和动态运行时场景之间的空白，确保边缘设备资源在运行时以自适应方式被充分利用。
• 与其他方法的正交性 (Orthogonal to Other Approaches)：
  • Miriam 可以与模型压缩 [23]、边缘-云卸载 [41] 等其他优化 DNN 执行方法协同工作，以更有效地执行 multi-DNN workload。这种协作方法有望在资源受限的边缘计算环境中实现更高的运行时性能和更好的资源利用率。

7.3. 个人启发与批判

7.3.1. 个人启发

1. 弹性内核的创新性： elastic kernel 的概念非常巧妙，它从底层资源调度入手，通过改造内核的执行方式来适应动态的边缘环境。这种编译器-运行时协同设计，在不修改硬件或驱动的前提下实现细粒度的资源管理，是解决边缘 AI 挑战的有效途径。
2. 细粒度争用管理： 论文清晰地识别并解决了 intra-SM 和 inter-SM 两种核心资源争用，这是许多现有工作未能全面覆盖的。这为理解 GPU 上的并发性能瓶颈提供了深刻的洞察。
3. 实践与理论结合： WIScore 和 OScore 等指标的设计，以及基于二叉树的 elastic kernel shard 管理，体现了理论分析与实践工程的良好结合。它们有效地将复杂的资源争用问题量化并转化为可操作的调度策略。
4. 对混合关键性任务的洞察： 论文强调了 critical task 和 normal task 之间不同的性能需求，并提供了一个优雅的解决方案，即在保证前者实时性的同时最大化后者的吞吐量。这对于自动驾驶等真实世界应用具有重要意义。

7.3.2. 批判性思考

1. WIScore 公式中的潜在笔误： 在公式 (3) WIScore 中，$S_{blk\_be} + S_{blk\_be}$ 部分可能是一个笔误，直观上更合理的形式可能是 $S_{blk\_rt} + S_{blk\_be}$，或者是其他含义。虽然论文严格按照原文呈现，但这可能会对初学者理解公式的实际计算意图造成困扰。如果确实是笔误，澄清这一点将有助于提高论文的严谨性。
2. 设计空间缩小策略的鲁棒性： 论文提到通过测试代表性的张量操作来选择前 20% 的候选作为设计空间缩小后的集合。这种启发式方法在不同 DNN 架构和任务组合下，是否始终能保证不剪枝掉最佳性能的参数集？其泛化能力值得进一步探讨和更严格的理论分析。
3. texture cache 和 L2 memory 争用： 论文明确指出 critical kernel 的启动时间开销主要源于 texture cache 和 L2 memory 上的争用，并将其列为未来工作。这意味着 Miriam 在 SM 级别虽然实现了高效填充，但在更高级别的 GPU 内存层次结构中仍可能存在瓶颈，这限制了其在某些内存密集型任务上的最终性能。
4. persistent thread 技术的额外开销： 论文提到 persistent thread 机制虽提供细粒度控制，但伴随“不可忽略的开销”。虽然 Miriam 通过缩小设计空间来缓解，但这种固有开销在不同场景下对性能的影响程度，以及如何进一步优化，是值得深入研究的方向。
5. 实际部署的复杂性： 尽管 Miriam 避免了修改 GPU 驱动，但 source-to-source kernel transformer 的实现可能仍需要深入理解 CUDA 编译流程和各种 DNN kernel 的内部结构，这对于广泛部署而言可能具有一定的工程复杂性。