Tally: Non-Intrusive Performance Isolation for Concurrent Deep Learning Workloads

• ASPLOS 25

• LC推理+BE训练

• https://github.com/tally-project/tally.git

引入

• DL的优势

  • DL在广泛的应用中表现卓越，为促进DL应用的研究与部署，各机构正在构建大规模GPU基础设施

• DL的挑战

  • 计算资源需求增加，高效利用GPU至关重要，来降低成本并保障AI应用的可扩展性

  • 训练负载中，数据获取和处理这种CPU执行成为瓶颈

  • 推理任务受请求率大幅波动影响，资源超额配置以满足峰值QoS，导致不可预测的GPU闲置

• 挑战的解决

  • 大量研究和实践证明GPU共享能答大幅减少资源浪费，提高GPU利用率，提升集群效率

• GPU共享目前的局限

  • GPU共享还未大规模投入生产，因为缺乏健壮的性能隔离，LC、BE任务打包共置将产生长尾延迟

  • 许多GPU共享方法是侵入式的，要么修改用户代码，要么修改DL框架，限制开发灵活性，提高适配难度

  • 这类方法要求负载符合某些特性，限制了适用性，因为生产集群通常有大量的异构特征负载

• 由局限得到优秀方法应满足的条件

  • 非侵入，便于生产集群适配

  • 性能隔离保障，满足LC任务的SLA

  • 普遍适用各种DL应用

• 提出解决方案Tally

  • 非侵入的健壮性能隔离保障的生产集群GPU共享系统，能最大化硬件效率和吞吐量

  • 通过拦截GPU内核（指GPU上执行的任务，包含thread、wrap、block、grid）启动API调用，采用任务无关调度策略，优先执行LC内核，同时利用空闲调度BE内核

  • 创新提出块级内核调度原语：分片与抢占，分片将大内核拆成小片段，抢占允许中断正在执行的内核

  • 可通过GPU编程模型将内核组织为独立的线程块的集合，通过非侵入的任务无关的内核转换实现上述原语

  • 上述原语各有缺陷，分片产生更多内核启动开销，抢占涉及复杂转换和同步开销，需要动态透明地分析负载的内核实时性能特征选择原语

• 实验与结论

  • 通过模拟真实生产推理负载的流量模式评估了性能，达成了略低于sota的吞吐量但保障了极低的99%延迟开销

背景

• GPU共享技术分为应用级和运行时级

  • 应用级引入专为并行深度学习执行设计的新编程接口

  • 运行时级将共享方案集成至深度学习框架、容器编排系统、GPU驱动程序等实现透明化运作

实践中的GPU共享

• GPU架构与编程模型，略

• 当前GPU共享的局限

  • 高集成与维护成本

    • 应用级修改DL框架或用户代码，运行时级需要集群管理员实现并维护

    • 运行时级可能隐式要求用户修改代码以满足预剖析或加入检查点与重启以实现实时迁移

  • 缺少性能隔离保障

    • LC任务SLA对响应时间要求严苛，负载共置时性能退化严重

    • GPU上下文切换速度受调度粒度影响，粒度越细（迭代 —> 核 —> 块 —> 线程）速度越快，粗粒度时切换时间远大于推理任务完成时间

    • 依赖于适用狭窄的负载特征，如需要预测内核执行模式，或要求所有内核幂等（输出与执行次数无关）

系统设计

• 层次：应用层和GPU之间

• 调度算法：任务无关优先级调度，优先调度LC任务、空闲时机会性地调度BE任务

• 性能隔离：分片与抢占原语，块级调度在内核间通用且开销可接受，两个原语可通过拦截PTX码转换后再应用于内核实现

• 组成

  • 内核转换器

    • 分块转换：将大内核拆成多个小内核，增加偏移量参数计算原blockIdx

    • 抢占转换：

      • 启动固定个数的工作块以迭代方式动态处理原来的所有block，通过全局计数器获取当前处理块的索引

      • 执行过程被计数器记录，可以中断与恢复，每个工作块执行直到完成或收到抢占信号

      • 将原始内核包装外层循环，返回变为跳转到下一循环来执行下一个block，被抢占时退出循环

    • 统一同步转换

      • 上述抢占将导致块内线程同步异常，因为有的线程会被抢占提前退出而未到达同步点

      • 将__syncthreads() 和 return 指令修改为跳转到一个同步点，仅在块内线程全部return时返回，否则，将return线程挂起，将__syncthreads() 送回循环继续迭代

  • 透明分析器

    • 通过周转延迟估计算法，预估分块内核和抢占内核的执行时间，决定选择哪种原语以及如何配置BE任务内核

  • 非侵入式虚拟化层

    • 在客户端共享库中实现替身函数，动态链接器优先调用这些函数，同时捕获设备代码，将调用信息与设备代码序列化发送给服务端，服务端进行内核转换后重编译并调用GPU执行，将结果发送给客户端，信息交互过程通过共享内存等优化

实验

准备

• 数据集

  • 负载自建，6训练，6推理

  • 请求到达模式：MAF2，选取频率最高的函数，50%load，https://github.com/Azure/AzurePublicDataset/blob/master/AzureFunctionsInvocationTrace2021.md

• 指标

  • P99延迟：推理任务第99%位的响应时间

  • 吞吐量：训练和推理任务单位时间处理样本数

  • 系统吞吐量：所有任务归一化吞吐量之和

• baseline：

  • Time-Slicing、MP、MPS-Priority、TGS

实验

• 双任务场景

  • 将推理负载与某个不同的训练负载并置，探究不同组合下，50%load时推理负载的延迟以及系统吞吐量

  • 将推理负载与某个不同的训练负载并置，探究不同组合下推理负载的延迟以及系统吞吐量随GPU空闲时间的变化

  • Bert推理与Bert训练共置时推理请求数、推理延迟、BE任务吞吐量随时间的变化

• 多任务场景

  • 10%load，将一组相同但只有一个标为高优先级的resnet50推理负载共置，探究推理负载延迟和系统吞吐量随共置数量变化关系

• 消融实验

  • 对优先级调度、内核转换进行消融实验

  • 对周转延迟界限进行调优