LithOS: An Operating System for Efficient Machine Learning on GPUs — Detailed Summary

Patrick H. Coppock, Brian Zhang, Eliot H. Solomon, Vasilis Kypriotis, Leon Yang†, Bikash Sharma†, Dan Schatzberg†, Todd C. Mowry, Dimitrios Skarlatos | Carnegie Mellon University / †Meta | SOSP '25 (31st ACM Symposium on Operating Systems Principles), Seoul, Oct 13–16 2025 | DOI: 10.1145/3731569.3764818

Per-section summary organized by the paper's heading structure. Each section includes paragraph-level bullet points and exact quantitative results where the paper provides them. All named components, equations, the Prelude algorithm, and the full experimental setup are preserved.

Abstract

• ML growth has made GPUs indispensable in datacenters and made GPU efficiency urgent; balancing diverse model demands against high utilization is a fundamental challenge.
• The authors argue that transparent, fine-grained GPU resource management that maximizes utilization, energy efficiency, and isolation requires an operating system approach. They introduce LithOS, a first step toward a GPU OS.
• LithOS contributes four new abstractions/mechanisms: (1) a TPC Scheduler for spatial scheduling at the granularity of individual TPCs, unlocking TPC stealing; (2) a transparent Kernel Atomizer to cut head-of-line blocking and allow mid-execution resource reallocation; (3) lightweight hardware right-sizing that determines the minimal TPC resources per atom; (4) transparent power management that lowers power based on in-flight work.
• Built in Rust, LithOS is evaluated against state-of-the-art NVIDIA and research solutions. Headline results: inference stacking cuts tail latency 13× vs MPS, and 4× vs the best SotA while raising aggregate goodput 1.3×; hybrid inference-training stacking cuts tail latency 4.7× vs MPS and 1.18× vs best SotA while improving aggregate throughput 1.35×; right-sizing yields 25% GPU capacity savings on average for under 4% performance loss; power management yields 25% energy savings on average for a 7% hit.

1. Introduction

• ML workloads drive massive GPU deployments, yet utilization stays low: 52% at Microsoft, 10% at Alibaba, below 30% for many Meta inference services, and ~40% for Llama 3 training. Figure 1 shows normalized utilization over a week. With per-GPU power now exceeding 1,000 W and high monetary cost, this is unsustainable.
• High utilization via GPU sharing is hard: dedicating a GPU to one workload still leaves cores idle on communication stalls, low batch sizes give insufficient parallelism, and dynamic request loads force overprovisioning. Rising SM counts and memory bandwidth make full utilization harder over time.
• Collocating latency-critical (LC) with best-effort (BE) tasks is one remedy, but existing systems cannot practically prioritize LC over BE under contention. Many lack transparency (incompatible with much of the ML stack); some are tied to unmaintained PyTorch/TVM versions. Temporal schedulers like TGS and Clockwork cannot run multiple models in parallel; spatial schemes (MPS, MIG, REEF, Orion) run in parallel but are coarse-grained (whole requests/batches/operators), causing low utilization and head-of-line (HoL) blocking.
• Beyond collocation, datacenter GPU management must move past static provisioning, which ignores fluctuating compute intensity and parallelism across models and execution phases, leaving GPUs underutilized while consuming significant power.
• This contrasts with CPUs, where time-sharing OSes context-switch cheaply and provide isolation, allocation, power management, and transparency. The extreme data-parallel nature of GPUs exposes the limits of abstractions built around compilers, frameworks, and drivers; the authors argue GPUs must evolve toward an OS model with first-class control, isolation, and resource management.

1.1 Our Approach: An Operating System for GPUs

• LithOS is fully transparent to the ML stack (no changes to models, runtimes, frameworks). It moves the bulk of GPU scheduling out of proprietary drivers/hardware into software, enabling fine-grained temporal and spatial scheduling. It works at the granularity of individual kernel thread blocks dynamically mapped onto Texture Processing Clusters (TPCs), decoupling work submission from thread-block execution.
• First, a fine-grained TPC Scheduler asynchronously decides compute-unit allocation and submission time per piece of work, at individual-TPC granularity for strong isolation, guided by an online kernel latency predictor and augmented with TPC Stealing for utilization.
• Second, to compensate for the absence of hardware preemption, kernel atomization transparently partitions kernels into schedulable atoms (subsets of thread blocks) with no compiler/runtime/PTX changes, reducing HoL blocking and enabling mid-execution TPC reconfiguration.
• Third, dynamic hardware right-sizing uses lightweight models to find the minimal TPC resources per kernel/atom, saving capacity.
• Finally, fine-grained power management adjusts GPU frequency to the characteristics of in-flight work, saving energy.
• Implemented in Rust. Restated headline numbers: inference stacking 13× tail vs MPS, 4× tail + 1.3× goodput vs best SotA; hybrid 4.7× vs MPS, 1.18× tail + 1.35× aggregate throughput vs best SotA; right-sizing <4% hit → 25% capacity savings; power management 7% hit → 25% energy savings.
• Contributions: a study of Meta inference services; a fine-grained spatial TPC Scheduler with TPC Stealing; a transparent Kernel Atomizer; dynamic hardware right-sizing; transparent power management; the LithOS design; and an evaluation across ML environments.

2. Background and Related Work

2.1 A Brief Background on GPUs

• GPU architecture (Figure 2): NVIDIA GPUs comprise several Graphics Processing Clusters (GPCs); each GPC has multiple TPCs; each TPC has a few Streaming Multiprocessors (SMs); each SM has tens of cores. The H100 has 8 GPCs, 9 TPCs/GPC, 2 SMs/TPC, 128 cores/SM.
• GPU programming: applications are composed of kernels running specific operators (e.g., convolution); a kernel fixes its resources (thread blocks, threads, registers, shared memory) at launch; each thread block runs on an SM.
• GPU streams: CUDA streams allow concurrent execution of independent tasks; work within a stream runs FIFO; some CUDA calls are asynchronous, others block.

2.2 Related Work

• Cooperative multitenancy: tenants coordinate to share resources, usually at the ML-framework level with all models in one process; limited because it cannot support arbitrary applications, often needs extensive offline profiling or kernel source modifications, and breaks when tenants don't cooperate.
• Transparent multitenancy: supports unmodified applications via runtime mechanisms (time slicing, MPS, MIG, or combinations). Most prior software solutions still need app/framework changes; TGS is an exception enabling transparent sharing across containers, but uncooperative tasks and limited application knowledge make it hard.
• Temporal multitenancy: dedicates the whole GPU to one task at a time. Some operate at whole-inference-request granularity (Clipper, Nexus, TensorFlow-Serving, Clockwork, INFaaS); others schedule kernels (PipeSwitch, AntMan, Gemini, KubeShare, TGS). NVIDIA's default time slicing gives each task exclusive access for several milliseconds round-robin. All run one job at a time → low utilization.
• Spatial multitenancy: usually built on MIG or MPS. MPS multiplexes GPU contexts onto one device for greater throughput but with performance interference. MIG partitions compute/memory along GPC boundaries for strong isolation, but is coarse-grained with steep reconfiguration overheads (>5 s), leaving resources idle. Hardware multitenancy works for Meta's production use, but fluctuations in Figure 1 occur at finer granularities, forcing overprovisioning.
• Existing spatial systems are coarse (inference-request or kernel level). They protect LC apps by restricting other jobs' kernels or capping BE resources (REEF, MuxFlow, PTask, others). Coarseness causes HoL blocking, low utilization, interference. Figure 3 illustrates the pitfalls: (a) a single workload issuing two requests finishes A and B fast but underutilizes the GPU; (b) MPS concurrency raises utilization but delays the original task's requests.
• Right-sizing: prior work often needs hardware modifications, lacks software transparency, depends on offline profiling, and operates at whole-job granularity, limiting fine-grained benefit.
• DVFS: recent work applied DVFS to minimize GPU power (esp. LLM inference clusters) but relies on extensive offline profiling across input lengths, trains dedicated output-length predictors, and operates at coarse (whole-request) granularity — missing finer opportunities and lacking transparency.

3. Motivation

• The section presents a detailed study of production GPU infrastructure challenges and opportunities.

3.1 Understanding GPU Utilization in Datacenters

• Meta inference services serve DL models partly on NVIDIA H100 nodes (8 GPUs per node), partitioned via software/hardware multitenancy into container shapes assigned by offline per-model analysis targeting tight tail-latency SLAs.
• Figure 1: device utilization ranges from under 25% to higher than 60%; SM utilization is lower, with lows of 15%; memory bandwidth has headroom (as little as 1/5 utilized on the low end); memory-capacity utilization is steady because models stay resident to meet SLAs. Tight SLAs enforce small batch sizes that prevent saturation even at high load; memory-intensive models drive SM utilization down. Multi-model stacking can lift utilization.
• Figure 4 (inference traffic): mean-normalized requests/sec over a week shows a diurnal pattern; RPS scales 2.2× (max/min = 2.23) between minimum and maximum, tracking the Figure 1 utilization trend.
• Figure 5 (model request frequencies): sampling thirteen common models on a log scale, the most popular model A gets several hundred times more requests than the least popular model M; overprovisioning for this wide distribution causes underutilization, especially for unpopular models.
• Figure 6 (model sizes): model sizes (weights/parameters/embeddings) vary by more than 10× between largest and smallest; half are large, half smaller, and both are heavily used (smallest model B has usage comparable to larger E and G), highlighting an opportunity to collocate models of different sizes while meeting each one's SLA.
• Takeaways: despite the need to raise utilization, datacenters rely on limited GPU sharing or MIG because of compatibility/transparency needs; non-transparent solutions are impractical at scale given the complexity of maintaining many frameworks/runtimes/compilers. Transparent solutions avoid locking infrastructure into outdated designs — motivating LithOS as a fully transparent OS for efficient ML multitenancy.

4. Abstractions, Interfaces, and Principles for a GPU Operating System

• LithOS is built on abstractions/interfaces/principles defining how a GPU OS should manage resources: the right granularity of control, flexible but predictable knobs, and a balance of efficiency, fairness, and robustness.

4.1 Resources and Isolation

• Scheduling granularity: GPU cores and memory bandwidth grew by orders of magnitude while GPC counts stayed nearly flat (six in P100 to eight in B100, Figure 7); with multi-die designs, coarse GPC-level partitioning (MIG) wastes more resources. Intra-SM control belongs to hardware/compilers (OS intervention would break transparency), so LithOS adopts the TPC as its scheduling abstraction. Current APIs do not expose TPC/SM control; LithOS uses reverse-engineering and argues native support is feasible. Principle: manage resources at the finest granularity where the OS is effective while preserving transparency.
• Resource allocation: one-app-per-GPU ignores that kernels scale differently; LithOS allocates TPCs by runtime scaling behavior and lets users specify tolerable performance loss for right-sizing. Principle: expose simple performance-tolerance knobs while hiding hardware complexity.
• Power management: device-wide DVFS assumes a single workload; in multitenant settings memory-bound kernels saturate bandwidth early while compute-bound kernels want high frequency. LithOS virtualizes frequency so each workload sees its preferred setting while the OS optimizes system-wide energy. Principle: virtualize frequency to preserve the illusion of dedicated control while optimizing system-wide power.
• Security and fault isolation: NVIDIA solutions are extremes — MIG (strong but coarse GPC isolation) vs MPS (flexible but unprotected). LithOS uses TPC-level isolation; each app runs in its own address space with hardware-enforced memory isolation; LithOS interposes on common GPU errors, killing only the faulty process and reinitializing the driver on unrecoverable errors. Principle: enforce isolation at the finest practical granularity, and ensure local faults degrade gracefully.

4.2 Closing the Gap

• These abstractions define LithOS's philosophy: virtualize resources at the right granularity, expose predictable interfaces, ensure robustness under multitenancy. The key challenge is bridging CPU and GPU OS design; LithOS applies proven CPU OS principles to GPU realities, transforming GPUs from single-model devices into virtualized multitenant platforms and laying a foundation for future GPU OS research.

5. LithOS Design

5.1 Architecture Overview

• Figure 8: LithOS runs on CPU cores and interposes at the driver level via a dynamically linked library, LibLithOS, that mimics the native CUDA library. It keeps a system-wide view of GPU state across applications of varying priority. The stack is: unmodified apps/frameworks (TensorRT High-Priority A, JAX High-Priority B, PyTorch Best-Effort) → LibLithOS → LithOS core with four components (TPC Scheduler, Kernel Atomizer, Hardware Right-sizing, Power Management) → GPU Device Driver → GPU hardware (TPCs + memory).
• Applications submit kernels following the CUDA model; LithOS decouples submission from execution, shifting scheduling from driver/hardware into software. The TPC Scheduler manages resources per TPC and unlocks TPC Stealing (idle TPCs lent to other tasks).
• The Kernel Atomizer transparently breaks kernels into thread-block chunks called atoms without source or PTX access, enabling finer scheduling and cutting HoL blocking.
• Hardware right-sizing uses lightweight models to shrink per-kernel/atom TPC allocations, yielding capacity savings.
• Transparent fine-grained DVFS adjusts frequency to in-flight work for energy savings. Together these enable intelligent scheduling policies, detailed with reference to Figure 9.

5.2 Interface with Userspace

• Kernel submission: apps interact with LithOS via launch queues that buffer work (Figure 9, Step ①), giving LithOS control even after dispatch — important because a kernel's priority/resources cannot change once submitted. A launch queue is created on cuStreamCreate; on async calls such as cuLaunchKernel, LithOS enqueues the kernel and returns control to the app.
• Compute quotas: users/admins enforce GPU limits via TPC quotas (Figure 9, Step ②), guaranteeing each app a number of TPCs when it has runnable work, analogous to CPU cores; a lightweight scheduler coordinates launch queues and quotas.

5.3 TPC Scheduler

• LithOS's scheduler operates per TPC, supporting dynamic on-the-fly allocation without MIG-style reconfiguration overhead, so kernels can run on different TPCs and high-priority apps get guaranteed resources. Because fixed allocations leave TPCs idle, LithOS adds dynamic scheduling and TPC Stealing; the authors view TPC scheduling as a foundation for evolving GPU scheduling policies, as CPU scheduling matured over time.
• Operation: dispatcher threads monitor launch queues (Step ①) and submit work, aiming to keep the GPU busy while preserving flexibility. Two challenges: varying kernel durations and balancing flexibility against GPU starvation. The first is handled by Kernel Atomization (Step ③) splitting long kernels into atoms; the second by tracking outstanding work via sync queues (Step ⑤), throttling submissions until backlog drops below a tunable threshold — a 100 µs limit that covers host-device latency. A dedicated Tracker thread monitors completion and updates scheduler state.
• TPC Stealing: the scheduler reassigns underutilized TPCs across apps. In Figure 10(a) static allocation leaves TPCs idle; in 10(b) stealing lets A₁ borrow TPCs from an idle workload. Stealing risks HoL blocking via priority inversion (a new request B delayed by C₂ on stolen TPCs). To mitigate, the scheduler keeps per-TPC timers informed by a latency prediction module estimating kernel/atom durations at submission, avoiding stealing from long-running TPCs; as tasks finish, sync queues clear and timers update. LithOS also caps outstanding atoms and uses lower hardware stream priorities for work on stolen TPCs.

5.4 Kernel Atomizer

• The Kernel Atomizer transforms kernels into atoms, each a subset of the grid's thread blocks (Step ③), without source/PTX access, enabling dispatch at thread-block rather than kernel granularity — critical because kernel times vary from microseconds to tens of milliseconds.

• Impact on latency (Figure 11): 11(a) shows P₉₉ kernel latency rising with training batch size (normalized by memory usage); most models quickly produce multi-millisecond kernels, and DLRM kernels exceed 30 ms — training is the main culprit. 11(b) shows LLM inference on a Microsoft Azure trace with small/medium/large prompts producing several-millisecond kernels for large prompts; with tight SLOs (low tens of ms), this motivates finer-grained scheduling that mitigates HoL blocking.

• Operation: before scheduling a long kernel, LithOS predicts its duration for the given TPC assignment (predictor, §5.7), then divides by a tunable atom_duration to get the number of atoms; setting it too low can make the atomized kernel slower, and limits of 250–500 µs are effective. Atoms are submitted and scheduled on TPCs dictated by the TPC Scheduler (Step ④). The atomizer works on any framework (including closed-source cuDNN) and any compiler.

• Granularity benefit (Figure 10(c)): dividing kernels into atoms packs work more tightly and lets TPC allocations adjust mid-kernel; B₁ is no longer blocked by C₂ because stealing is disabled for C₂'s subsequent atoms (Ĉ₂) once request B arrives.

• Prelude kernel demonstration: for a Conv kernel with grid {8,8,1} → 64 blocks (block_idx 0–63), LithOS launches a Prelude kernel with the same configuration that checks whether block_idx is in a range and either calls Conv or exits early. To make 2 atoms, the prelude launches twice over ranges [0,32) and [32,64); LithOS can split into up to 64 atoms. Non-overlapping ranges ensure each block runs exactly once.

  Algorithm 1 — Prelude Kernel:

  1  kernel fn prelude(*args):
  2      let atom : *const AtomMetadata = AtomMetadataAddr as _
  3      let block_idx = blockIdx.z * gridDim.y * gridDim.x
  4                    + blockIdx.y * gridDim.x
  5                    + blockIdx.x
  6      if atom->block_idx_lo <= block_idx < atom->block_idx_hi:
  7          atom->kernel_entrypoint(*args)

• Atomization considerations: the Prelude uses the same resources as the original kernel and needs the original entry point, passed via an AtomMetadata struct (Algorithm 1).

• Performance optimizations: LithOS monitors atomizer effectiveness — it may disable atomization for kernels with many short threads (to avoid prelude overhead) and dynamically adjusts atom_duration for kernels with many thread blocks to limit the penalty from early-exiting threads.

5.5 Right-Sizing Hardware Resources

• TPC-level scheduling unlocks fine-grained right-sizing. Figure 12 plots kernel speedup vs allocated TPCs for kernels covering 99% of execution time; for Llama inference, GEMM and multihead-attention kernels show diminishing returns while the token-frequency-penalty kernel does not scale. Whole-model right-sizing is suboptimal because no single configuration fits all kernels; kernels exhibit diverse scaling, and how execution time spreads across kernels varies by workload.

• Modeling kernel scaling: LithOS does on-the-fly per-kernel right-sizing (Step ⑥); atoms inherit their kernel's TPCs and scaling. From two measured points (latency on all TPCs and on one TPC), it fits

  l = m/t + b

  where l = predicted latency, t = number of TPCs, m and b constants. The form is consistent with Amdahl's law: b ≈ time for a single thread block on one SM; m ≈ how much the kernel exploits parallel processors.

• Filtering outliers: a few kernels (typically very short) deviate; a filtering heuristic based on thread-block occupancy estimates usable TPCs by dividing total thread blocks by occupancy per TPC (blocks a TPC runs concurrently). LithOS already tracks blocks per kernel from atomization and queries occupancy from the driver API, giving an upper bound that avoids overprovisioning hard-to-model kernels.

• Operation: on submission, the dispatch thread applies the filtering heuristic; if the estimate is below the job's allocated TPCs, it launches with the estimate. Otherwise it uses the learned model to find the minimum TPCs that raise latency by at most a multiplicative latency slip parameter k (e.g., configure for 10% acceptable degradation).

• Supporting hardware-aware optimizations: right-sizing is orthogonal to and complements framework/compiler intra-SM optimizations (Tensor Cores, warp-level techniques, memory-hierarchy tuning); LithOS works at the inter-SM kernel-to-TPC level, so the two coexist seamlessly.

5.6 Transparent Power Management

• LithOS enables transparent DVFS; like right-sizing scales resources horizontally, DVFS scales frequency vertically. Figure 13 shows kernels respond predictably to frequency scaling, enabling bounded-impact energy savings. Two challenges: (1) frequency switching is slow (~50 ms), making DVFS impractical for very short kernels, so LithOS considers cumulative impact across kernel sequences; (2) many kernels scale linearly, so gains must be balanced against latency.

• Modeling frequency scaling: LithOS uses a transparent sequence-based model (Step ⑦); atoms inherit the kernel's frequency target. Each kernel gets a weight w = its runtime fraction of the stream. Using a first-order Taylor approximation of relative slowdown vs fractional frequency drop:

  k = lat(f_th)/lat(f_max) − 1 = s · (f_max/f_th − 1)

  giving per-kernel sensitivity s = k / (f_max/f_th − 1), aggregate sensitivity S = Σ w·s, total slowdown S · (f_max/f_final − 1) ≤ k, and the assigned frequency

  f_final = f_max / (1 + k/S).

  Compute-bound kernels (slowdown linear in frequency reduction) skew f_final toward maximum; memory-bound kernels (frequency-insensitive) shift it lower per their weight.

• Operation: like right-sizing, DVFS uses the latency slip parameter k. Due to high switching latency, LithOS is conservative and extends its learning period: it first collects per-kernel metadata at maximum frequency, runs unseen kernels at max, assumes linear scaling and reduces frequency per k, then either lowers further or stops after confirming linearity, fitting collected data to the model over time.

5.7 Online Latency Prediction

• The latency prediction module learns kernel execution times for all LithOS components: it improves TPC Stealing (estimating outstanding durations), sets the atom count, and supplies latencies for right-sizing and DVFS speedup calculations — avoiding offline profiling impractical for a transparent OS.
• Prediction runs separately per launch queue to adapt to each application. It records kernel latencies during execution and refines predictions; because latency depends on TPCs, frequency, and atomization granularity, the module monitors these conditions and is conservative (assumes optimal linear scaling) when atom metadata is missing.
• A pitfall is assuming a kernel always has the same latency; duration depends on launch parameters and inputs (e.g., one Conv function across layers with varying tensor sizes). So the module tracks operators, not kernel functions: by recording explicit synchronization events it finds batch start/end and tags each launch with an ordinal index k (the kᵗʰ kernel after batch start), identifying operator nodes in the model's data flow graph (DFG) without explicit access to it.

6. Implementation

• The prototype targets NVIDIA GPUs in ~5000 lines of Rust (excluding macro-generated interposition code for the full CUDA Driver API), supports Ampere and Hopper, runs apps natively or in containers, and is built on top of MPS for concurrent execution across GPU contexts.
• Interposition architecture: LithOS is fully transparent, interposing at the CUDA Driver API (the lowest common denominator) so apps interact with LithOS while preserving CUDA semantics, supporting unmodified PyTorch, TensorRT, TensorFlow, JAX, and cuDNN. It implements only a small subset of Driver APIs (e.g., cuLaunchKernel) and auto-generates the rest; unlike prior CUDA interposition, it avoids cross-address-space marshaling, easing support for new CUDA versions.
• TPCs and atomization: LithOS extends the libsmctrl reverse-engineering work and adds Hopper support (new TPC masking layout); it reimplements TPC identification through the Queue MetaData (QMD) structure to enable dynamic TPC allocation at launch. On Hopper it reverse-engineers the new Thread Block Clusters abstraction, ensuring atoms are always multiples of the cluster size. Functionality is verified across Ampere (A30, A100), Hopper (H100), and Ada Lovelace (L4). Future drivers could expose these APIs.
• For atomization, LithOS injects Prelude logic by modifying the QMD struct: it first launches the original kernel so the CUDA driver configures the environment, then patches the QMD program address to the Prelude, which executes while retaining the original kernel's resources. QMD reverse-engineering is minimal (often days for a new architecture; details deferred to a separate technical report).
• Special kernels: for CUDA Graphs, LithOS interposes graph-creation APIs and atomizes graphs into subgraphs with correct ordering. Kernels with cross-block synchronization (e.g., grid_group::sync()) need a certain SM count; LithOS can return the allocated SM count for CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT. For cross-block-sync or persistent kernels, LithOS disables stealing and atomization.

7. Experimental Setup and Methodology

• Testbed: 1× A100 (SXM4) Lambda Labs instance, 30 CPU cores, 216 GB host memory; the A100 has 108 SMs and 40 GB. Software: Ubuntu 22.04, CUDA 12.6, Rust 1.83.0-nightly, Python 3.10, PyTorch 2.3, TensorRT 10.1, TensorRT-LLM 0.11.0, Triton 24.07.
• Baselines: all four NVIDIA sharing methods (Time slicing, MPS, stream Priority, MIG) plus SotA research: TGS, REEF (app modifications), and Orion (app modifications + offline profiling). TGS used directly; Orion and REEF reimplemented on LithOS's interposition infra and extended for multiple HP apps. For REEF, BE kernels are not launched if any HP app is running; for Orion, BE kernels are not launched if they contend with any HP kernel.
• Models/configs: HP inference runs on NVIDIA Triton with dynamic batching; RetinaNet on ONNX Runtime; others on TensorRT/TensorRT-LLM. Vision: RetinaNet, YOLOv5, ResNet-50 v1.5; language: Llama 3 8B, GPT-J 6B, BERT-Large. LLMs use a Microsoft Azure trace. BE training: ResNet-50, MobileNetV2, VGG-19, DLRM, BERT-Large, with batch size set to at most half the GPU DRAM so all models stay resident when stacking; BE training runs continuously.
• Latency constraints: from the MLPerf datacenter inference benchmark (Table 2), varying 2.3×–7.4× of baseline end-to-end request latency.

Table 1 — Training model parameters:

Table 2 — Inference services (inference-only multitenancy):

8. Evaluation

• Four questions: (1) does LithOS improve performance across multitenancy environments and SotA? (2) capacity savings from right-sizing? (3) energy savings from DVFS? (4) how do individual features contribute?

8.1 Performance in Multitenant Environments

• Right-sizing and power management are disabled here for an apples-to-apples comparison of scheduling efficiency alone.
• Inference-only multitenancy: two HP and one BE workload. HP A has a latency-oriented SLO; HP B a throughput-oriented SLO; HP/BE chosen from Table 2 combinations, with HP apps following Poisson load on Triton in closed loop, latencies measured end-to-end. For partitioning systems HP A/HP B get 75%/25%; MIG cannot do 25%-75%, so a 3/7–4/7 split is used; MIG and thread limits cannot run a provisioned BE; Priority/TGS/Orion cannot isolate multiple latency-sensitive apps, so both HP apps are high priority and BE low.
• Figure 14 (SLO attainment + throughput): MPS sets the throughput bar at 1.11 via intra-SM stacking but only 45% SLO attainment. MIG and thread limits meet SLOs but, without a BE app, aggregate throughput drops to 0.58 (thread limits) and 0.71 (MIG). Priority-only systems cannot attain SLOs (TGS leads at 84%). LithOS achieves 100% SLO attainment and throughput of 1.
• Figure 15 (goodput by app): LithOS leads in goodput while allowing significant BE throughput. Partitioning systems match HP A goodput but lag in HP B (MIG 0.37 vs LithOS 0.50) and support no BE. No SotA wins on all axes (Orion best in latency-sensitive throughput, TGS in HP throughput, REEF in best effort); only LithOS gives best HP throughput while sustaining high BE.
• Figure 16 (HP A tail latencies): averaged over combinations, only LithOS and partitioning systems hold latencies to constraints. MPS is worst; LithOS is 13× better. LithOS is 4× better than Orion (which cannot handle multiple HP apps) and 1.2× better than TGS.
• Hybrid inference/training multitenancy: stack an HP inference app (latency-oriented SLO) with a BE training job; idle resources donated to training without raising service latency. Inference from Llama 3 8B, GPT-J 6B, BERT-Large, RetinaNet, YOLOv4; training from Table 1; Poisson loads tuned to keep HP utilization around 80%.
• Figure 17 (P₉₉ service latency + aggregate throughput): MPS yields 5.83× ideal latency and lowest 60% service throughput. Time slicing fares better; MIG performs similarly by allocating 50% of the GPU, both failing peak HP throughput. Stream priority gives 2.89× latency and as low as 68% throughput. TGS averages 1.41× ideal latency (poor adaptive rate control assumes constant arrival rate) and REEF averages 2.89×, with REEF tails reaching 8.93×. LithOS holds tail latency within 20% of ideal, on average 2.34× and 1.18× better than REEF and TGS, and up to 13.54× / avg 4.7× better than MPS. LithOS keeps service throughput within 1% of load worst-case, improves training throughput 34× vs TGS, aggregate throughput 1.35× vs TGS, and overall 1.23×–1.57× (avg 1.38×).

8.2 Kernel-SM Right-Sizing

• Figure 18 (capacity savings): computed from time-weighted average TPC utilization before/after right-sizing — up to 51%, mean 26% across workloads; future higher-TPC GPUs should give more.
• Latency/throughput cost: with latency slip parameter 1.1, mean P₉₉ increase and throughput decrease are both 4%; the slip is conservative because not all end-to-end time is in-kernel.
• Accuracy: kernel-execution-time-weighted average R² of fitted curves ranges from 0.92 (Llama finetuning) to 0.99 (RetinaNet inference), showing linear models suffice.

8.3 Kernel-Dependent DVFS

• Energy compared to default settings (baseline mostly at maximum 1410 MHz), for a fixed number of requests/epochs; energy = average power × time, power sampled with nvidia-smi every 100 ms.
• Figure 19 (energy savings): difference between default frequency and LithOS DVFS — up to 46%, mean 26% across workloads, without offline profiling.
• Performance cost: slip parameter 1.1; mean P₉₉ increase 7%, showing a conservative policy that respects latency constraints; finer-grained control could unlock more savings.

8.4 Ablation and Case Studies

• Figure 20 (feature breakdown, inf-train): enabling the TPC scheduler improves HP tail latency to 1.38× ideal while keeping ideal HP throughput; adding Kernel Atomization reduces tail to 1.19× average (up to 1.55×) by splitting long BE kernels and improving stealing, at a 10% throughput overhead (LithOS trades BE throughput to protect HP).
• Figure 21 (atomization case study): HP BERT inference collocated with BE VGG training (vary batch size, (a)) or BE Llama 3 inference (vary sequence length, (b)), measuring HP P₉₅. LithOS beats REEF 6.5× in (a) and 3.9× in (b) (REEF reaches 47.9 ms and 26.9 ms). Disabling atomization shows it contributes 2× in (a) and 1.3× in (b). Full LithOS keeps HP tail within 14% (1 ms) for the largest batch size and 7% (0.45 ms) for the largest sequence length.
• Latency prediction module: comparing predicted atom latencies to CUDA events (absolute errors >50 µs counted as mispredictions): HP misprediction is only 0.9% (inf-inf) and 0.38% (inf-train), with error-tail P₉₉s of 49 µs and 31 µs; BE is higher at 14% and 11%, acceptable since BE is frequently preempted.
• Overheads: vs the vanilla NVIDIA driver without multitenancy, LithOS adds 4% (atomization <1%), vs ~2% for TGS and REEF and 6% for Orion.
• Memory contention: capacity contention is not a concern (models stay resident); bandwidth contention can matter — MIG/thread-limit estimates suggest bandwidth isolation would yield 4–13% gains, while compute isolation yields more than an order of magnitude improvement.

9. Discussion

• Other GPU resources: the same principles extend to memory, bandwidth, PCIe, SSDs, and networking targeted by prior systems; GPUfs is closest to an OS-like design (file-system extensions). LithOS complements these as a foundation to virtualize additional GPU resources.
• Driver and hardware support: further gains need kernel-to-SM assignment, preemption, cache/memory partitioning, NUMA-style placement, fine-grained (sub-ms) DVFS, per-SM power control, and richer context management — standard on CPUs and increasingly essential as GPUs scale, integrate multiple dies (e.g., Blackwell), and grow heterogeneous. Open-source drivers are seen as critical; intra-SM heterogeneity (tensor cores) opens further opportunities.
• Lessons learned: efficient GPU multitenancy needs both spatial and temporal partitioning (MPS lacks dedicated resources → interference; MIG lacks time-sharing → low utilization); fine-grained control is crucial (TPC scheduling yields many more virtual devices than MIG, atomization enables fast switching to HP tasks); power management is a key challenge demanding finer, sub-ms, spatially-applied mechanisms; the CUDA Driver API is a stable "narrow waist" for lightweight, portable interposition (easy retarget Ampere→Hopper). LithOS opens a new direction for GPU operating systems.

10. Conclusion

• LithOS is a first step toward an OS for efficient ML on GPUs, operating transparently across the ML stack; via TPC Scheduling, Kernel Atomization, hardware right-sizing, and power management it significantly improves GPU efficiency and lays a foundation for future GPU OS research.

System Architecture (block-diagram description)

+-----------------------------------------------------------------------+
| Unmodified Apps / Frameworks                                          |
|  TensorRT (HP A)   JAX (HP B)   PyTorch (BE)                          |
+-----------------------------------------------------------------------+
|                       LibLithOS  (mimics native CUDA)                 |
+-----------------------------------------------------------------------+
| LithOS core                                                          |
|  +-------------+  +-----------------+  +--------------+  +----------+ |
|  | TPC         |  | Kernel Atomizer |  | Hardware     |  | Power    | |
|  | Scheduler   |  | (atoms/Prelude) |  | Right-sizing |  | Mgmt/DVFS| |
|  | + Stealing  |  |                 |  | l=m/t+b      |  | f_final  | |
|  +-------------+  +-----------------+  +--------------+  +----------+ |
|        \_______________ Online Latency Prediction _______________/   |
+-----------------------------------------------------------------------+
|                       GPU Device Driver                              |
+-----------------------------------------------------------------------+
|        GPU Hardware:  [TPC0][TPC1][TPC2]...[TPCn] + Memory            |
+-----------------------------------------------------------------------+

Flow (Fig. 9):  app -> (1) launch queue buffer -> (2) TPC quota check ->
(3) kernel atomization -> (4) dispatch to TPC-mapped device queues ->
(5) sync-queue/Tracker throttle (100us) ; (6) right-sizing & (7) DVFS
consult the latency predictor.

Limitations and Future Work (author-stated)

• Built on top of MPS and on reverse-engineering (QMD, libsmctrl, TPC masking, Hopper Thread Block Clusters) because current APIs do not expose TPC/SM control; native driver support is argued feasible and desirable.
• Frequency switching is slow (~50 ms), so DVFS is impractical for very short kernels; sub-ms DVFS and per-SM power control are future hardware needs.
• No hardware preemption exists; atomization is a software workaround. Special kernels (cross-block sync, persistent) require disabling stealing/atomization.
• BE latency prediction is less accurate (14%/11% misprediction); more complex modeling is future work.
• Bandwidth/memory isolation is not implemented (estimated 4–13% gains); compute isolation is prioritized. Other resources (memory, PCIe, SSD, networking) are left as extensions.
• QMD reverse-engineering details are deferred to a separate technical report.

Named Components, Equations, and Algorithm (reference)

• Components/abstractions: LithOS, LibLithOS, TPC Scheduler, TPC Stealing, Kernel Atomizer, atoms, Prelude kernel, AtomMetadata struct, Hardware Right-sizing, Power Management (DVFS), Online Latency Prediction module, launch queues, sync queues, Tracker thread, dispatcher threads, Compute Quotas / TPC quotas, atom_duration, latency slip parameter (k), Queue MetaData (QMD), Thread Block Clusters, libsmctrl.
• Equations: kernel scaling l = m/t + b; frequency k = lat(f_th)/lat(f_max) − 1 = s·(f_max/f_th − 1); sensitivity s = k/(f_max/f_th − 1); aggregate S = Σ w·s; slowdown bound S·(f_max/f_final − 1) ≤ k; assigned f_final = f_max/(1 + k/S).
• Algorithm 1: Prelude Kernel pseudocode (reproduced in §5.4).

Note on Collective Communication

• The paper makes no mention of NCCL or collective communication. LithOS is single-GPU in scope (testbed is 1× A100; production nodes have 8 GPUs but LithOS operates per-GPU). Training workloads (DLRM, ResNet, VGG, MobileNet, BERT, Llama 3 finetuning) run as single-GPU best-effort jobs — there is no multi-GPU or distributed-training configuration, and no inter-GPU communication discussed. PCIe is mentioned only as a future extension; NVLink is not mentioned. The lone "communication stalls" reference is a general cause of intra-GPU core idling, not a collective-communication concern.