{"id":"75044f7c-4609-4e02-a202-3b7c52d95ecd","arxiv_id":"2504.15465","paper_version":1,"verdict":"CONDITIONAL","confidence":"MODERATE","novelty_score":7.0,"correctness_risk":"medium","formal_verification":"none","parameter_count":4,"one_line_summary":"LithOS splits kernels into atom chunks, schedules them at TPC granularity, right-sizes TPC counts, and applies DVFS, reporting 13x lower tail latency than MPS and roughly 26% GPU capacity and energy savings in A100 experiments.","lead":"LithOS is a software layer for NVIDIA GPUs that splits machine-learning kernels into small chunks and schedules each chunk on individual GPU compute clusters, so multiple workloads can share one GPU with less interference. In experiments on an A100 it reports much lower tail latencies than NVIDIA's MPS sharing and about a quarter of GPU capacity and energy savings for a few percent performance cost.","discovery_kind":"new_method","skeptic_critique":{"model":"deepseek-v4-flash","headline":"Kernel atomization is the central enabler, but as described it cannot deliver fine-grained scheduling: each atom is a full-grid launch, semantic correctness for cross-block kernels is unverified, and the prelude's device-side call to a __global__ entrypoint is not expressible in CUDA.","rationale":"The reader's weakest_assumption correctly identifies the block-independence premise behind atomization. My review sharpens this into a broader set of deficiencies: (i) the atoms as launched are full-grid prelude executions, so the claimed TPC-granular spatial control is not delivered by the published mechanism; (ii) semantic correctness for grid-wide or cross-block kernels is unverified; and (iii) the device-side call to the original __global__ entrypoint appears infeasible under CUDA without an unpublished binary-rewriting trick. The first and third issues are internal-consistency problems that the withheld technical report might resolve, but as the paper stands they leave the central novelty under-specified. This does not change the reader's CONDITIONAL verdict, because the condition is now clearer: the authors must release the technical report or artifact and demonstrate atomization on at least one grid-synchronizing kernel, while also explaining how a full-grid prelude launch can be mapped onto specific TPCs. If the CUDA-callability issue cannot be resolved, the verdict should move to REJECT, since the central mechanism would be unimplementable as described. I do not see a separate load-bearing concern in the right-sizing or DVFS evaluations; they are secondary and depend on the same atomization foundation.","tokens_in":21577,"tokens_out":15056,"duration_ms":152409,"concrete_test":"Release a minimal reproducer (or the LithOS prototype) and run a two-part check: (1) attempt to compile Algorithm 1 as written with nvcc, taking the address of a __global__ function and calling it inside a prelude kernel; (2) if it compiles, atomize a small kernel that uses cooperative_groups::this_grid().sync() and a second kernel with an atomic-counter grid barrier, comparing completion and output against unatomized launches. If compilation fails, or if either kernel hangs or produces different output, the published atomization mechanism is not semantically or mechanically valid as described.","verdict_should_be":"UNCHANGED","load_bearing_attack":"The paper's headline results (§1, §7.1) all depend on §4.4's Kernel Atomizer: transparently splitting kernels into atoms so the TPC Scheduler can place each atom on chosen TPCs. As written, Algorithm 1 launches a prelude 'using the same launch configuration' as the original kernel and gates execution on block_idx ranges. This has three unaddressed consequences. First, the atom is not a physically smaller launch: every atom submits the full grid, and unselected blocks early-return only after occupying hardware block-scheduling slots. The TPC Scheduler therefore cannot confine an atom to a subset of TPCs; the hardware still distributes all blocks, undermining TPC Stealing and per-TPC isolation. Second, semantics are preserved only if thread blocks are independent and have no grid-wide interactions. Kernels using cooperative-groups grid.sync(), atomic-counter grid barriers, persistent-CTA work queues, or cross-block producer-consumer dependencies can deadlock, double-execute work, or skip required finalization when split by block-index ranges. The paper reports no audit of the §7 cuDNN/TensorRT kernels for such patterns and no output-equivalence check against unatomized execution. Third, Algorithm 1 invokes `atom->kernel_entrypoint(*args)` from inside device code; CUDA does not permit calling a __global__ function as a regular device function or taking its address for device-side invocation, and §4.4 explicitly excludes source/PTX access, so it is unclear what `kernel_entrypoint` denotes. Section 5 defers these low-level details to a separate technical report that is not part of this submission. Until that mechanism is described and validated, the atomization-based efficiency gains are unsupported.","agreement_with_reader":"partial"},"referee_report":{"model":"deepseek-v4-flash","summary":"The paper presents LithOS, a userspace GPU resource-management layer that interposes on the CUDA driver API and claims to provide transparent, fine-grained spatial scheduling of ML workloads. Its main components are a TPC-level scheduler with TPC stealing, a kernel atomizer that splits kernels into thread-block ranges, a right-sizing mechanism that reduces TPC allocations, and a DVFS power-management mechanism. The evaluation, on an A100 with inference and training workloads, claims that LithOS reduces tail latencies by 13x versus MPS and by up to 3x versus the best prior software baseline for inference stacking, while also providing capacity and energy savings of roughly a quarter. The central design claim is that atomization is transparent and requires no source, PTX, compiler, or runtime modifications.","tokens_in":21917,"tokens_out":7171,"duration_ms":69026,"significance":"If the central mechanism worked, LithOS would be a notable step toward OS-like GPU management: it targets a real production utilization problem, evaluates across many models and frameworks, and compares against four NVIDIA mechanisms plus three research baselines. The evaluation breadth is a genuine strength, as is the attempt to make all ML-stack components unmodified. However, the contribution stands or falls on the kernel atomizer and TPC-level placement, and the manuscript as written does not provide a credible mechanism for either. The right-sizing and DVFS results are also of limited independent value until the scheduling foundation is established. The paper is therefore potentially significant but currently lacks support for its load-bearing premise.","major_comments":[{"comment":"The load-bearing atomization mechanism is not expressible in CUDA as described. In Algorithm 1, the prelude kernel calls `atom->kernel_entrypoint(*args)` from device code; a __global__ function cannot be called as an ordinary device function and its address cannot be taken for device-side invocation in the CUDA programming model, and the paper explicitly states (§4.4, §5) that LithOS has no source or PTX access. This makes it impossible to \"call into the original kernel\" for arbitrary cuDNN or TensorRT kernels. Furthermore, every atom is launched with the original full grid configuration, so the GPU's hardware block scheduler still distributes blocks across all TPCs; therefore the TPC Scheduler cannot confine an atom to a chosen subset of TPCs, which is exactly what TPC Stealing and per-TPC isolation require. The headline claims in §1 and §7.1 (13x and 3x tail-latency reductions, 1.35x aggregate throughput) all depend on this mechanism.","section":"§4.4, Algorithm 1; also §1 and §7.1"},{"comment":"Atomization by block-index range preserves semantics only if thread blocks are independent and have no grid-wide interactions. Kernels using cooperative-groups grid.sync(), persistent-CTA work queues, atomic-counter barriers, or cross-block producer-consumer patterns can deadlock, double-execute work, or skip required finalization when split by block-index ranges. The paper reports no audit of the §7 kernels for such patterns and no output-equivalence check against unatomized execution; without this, the evaluated results may reflect incorrect executions. Please provide evidence of semantic preservation or restrict the claim accordingly.","section":"§4.4"},{"comment":"The reported right-sizing accuracy is computed in-sample: the text says the model fits each kernel curve and then computes R² for the \"curves we fit\" on the same data. With the two-parameter form l = m/t + b and the two-point interpolation described in §4.5, R² would be 1 by construction, which is inconsistent with the reported 0.92–0.99 range. Please clarify whether the R² values come from held-out data and report prediction error on unseen TPC allocations; otherwise the 26% capacity-savings claim is not supported by the accuracy evidence.","section":"§7.2"},{"comment":"The comparison to Orion and REEF uses the authors' reimplementations instead of the original systems, with no validation of those reimplementations (e.g., against the papers' reported numbers), and all multitenancy results are point estimates without confidence intervals. Moreover, §5 states that \"some low-level details\" of TPC-level control are deferred to a technical report, so neither the TPC placement mechanism nor the reimplemented baselines can be independently checked from the manuscript. Given that the 3x and 1.18x \"best SotA\" claims depend on exact baseline numbers, this is a reproducibility gap for the central evaluation.","section":"§6 and §7.1"}],"minor_comments":[{"comment":"The symbol k is used both for the configured latency slip parameter and for the per-kernel relative slowdown in the DVFS formula, which is confusing; please use a distinct symbol for the slowdown.","section":"§4.6"},{"comment":"The discussion of prelude resources should state explicitly how shared memory, registers, and dynamic shared memory are replicated from the original kernel, since the atomization overhead claims depend on this.","section":"§4.4"},{"comment":"The legend \"MPS + TPC Scheduling + Kernel Atomization\" is not labeled as cumulative; please clarify that each line adds one feature.","section":"Figure 19"},{"comment":"Reference [23] contains a typo: \"unjie Qian\" should be \"Junjie Qian.\"","section":"References"}],"recommendation":"reject","confidential_remarks":"I am not recommending rejection lightly: the evaluation breadth and the production motivation are strong. However, the mechanism on which all headline results depend is, as written, not implementable in the CUDA programming model, and the manuscript itself defers the crucial details to an unavailable technical report. If the authors can release an artifact that demonstrates TPC-confined atomized execution and provides the actual driver-level mechanism, the paper would be worth reconsidering."},"author_rebuttal":null,"desk_editor":{"model":"deepseek-v4-flash","letter":"The central mechanism here is kernel atomization, and as described in Algorithm 1 it cannot do what the paper claims. The prelude is launched with the full grid and gates on block index, so every atom is still a full-grid launch; unselected blocks just exit. That does not give the TPC Scheduler the ability to place atoms on particular TPCs, and it does not provide per-TPC isolation. On top of that, the pseudocode calls a __global__ function as a regular device function, which CUDA does not allow, and the paper explicitly says it has no source or PTX access. The low-level details are deferred to a technical report, so the mechanism as written is unsound. This is a load-bearing flaw because the 13x tail-latency improvements and the per-TPC scheduling claims all depend on it.\n\nThe paper deserves credit for some things. The production utilization study of the Meta ads service is concrete and useful. TPC stealing and kernel-level right-sizing are sensible extensions of prior work, and the DVFS model is simple but reasonable. The evaluation is broad: multiple baselines, multiple models, an ablation, and an honest attempt to compare against prior systems.\n\nBut the evaluation does not repair the central flaw. There is no artifact, no output-equivalence check between atomized and unatomized kernels, no audit of the evaluated kernels for grid-wide synchronization or cross-block dependencies, and confidence intervals are absent. The right-sizing 'accuracy' is computed on the same data used to fit the curves, so the reported R2 is circular. Orion and REEF were reimplemented by the authors, which adds uncertainty. These are all addressable, but together they mean the paper's main claims are unverified.\n\nThis paper is for systems researchers working on GPU resource management and datacenter ML efficiency. The problem is real and the design is novel, so I would send it to peer review—but reviewers should ask for the technical report, an artifact, and a correctness check on atomization before accepting. As it stands, I would treat it as a conditional reject with an invitation to resubmit once the implementation is substantiated.","headline":"The system design is ambitious and the evaluation broad, but the kernel atomization mechanism as written cannot deliver the fine-grained spatial scheduling it claims, leaving the headline results unsupported.","tokens_in":22530,"tokens_out":5187,"would_cite":false,"duration_ms":46789,"reading_group":"yes","serious_thinker":"yes","would_accept_peer_review":true},"rs_alignment":null,"lean_confirmation":null,"pith_extraction":{"msc":[],"pacs":[],"model":"deepseek-v4-flash","headline":"Atomizing GPU kernels into schedulable chunks cuts tail latency 13x","keywords":["GPU operating system","kernel atomization","thread processing cluster scheduling","GPU multitenancy","CUDA driver interposition","GPU right-sizing","GPU power management","tail latency"],"falsifier":"Run LithOS on kernels that use grid-wide synchronization, cooperative groups, or cross-block communication and compare outputs bit-for-bit against monolithic launches; any divergence shows atomization is not universally correct. Separately, with MPS active, use hardware counters to check whether a high-priority workload's TPC allocation is truly respected while a best-effort kernel is running.","tokens_in":1693,"feed_emoji":"⚙️","tokens_out":4385,"duration_ms":83937,"temperature":0.7,"pith_summary":"LithOS argues that GPU underutilization in datacenters can be attacked at the operating-system layer rather than the framework or hardware layer. It interposes at the CUDA driver boundary and, without modifying applications, splits long-running kernels into atoms, contiguous ranges of thread blocks, and schedules those atoms onto individual thread processing clusters (TPCs), lending idle TPCs to other workloads. The paper reports that this transparent scheduling reduces P99 tail latency by 13x versus MPS and 3x versus the best prior software baseline for inference stacking, and by 4.7x versus MPS and 1.18x versus TGS for hybrid inference/training, while aggregate throughput rises by 1.6x and 1.35x respectively. It also claims that right-sizing TPC counts saves a quarter of GPU capacity on average for under a 4% latency cost, and that transparent frequency scaling saves a quarter of GPU energy for a 7% cost. A sympathetic reader would care because the result suggests that an OS-style resource manager, not new hardware or rewritten models, could provide MIG-like isolation together with high utilization.","feed_headline":"Atomizing GPU kernels into schedulable chunks cuts tail latency 13x","feed_subtitle":"A CUDA-level OS layer splits kernels into atoms and lends idle compute clusters to other workloads.","key_machinery":"The load-bearing mechanism is the Kernel Atomizer, implemented by a prelude kernel (Algorithm 1): a wrapper launched with the same grid configuration and resources as the original kernel that checks each block index against an atom's metadata and invokes the original kernel entry point only for blocks in the atom's range. Because the grid is partitioned into non-overlapping block-index ranges, a monolithic kernel becomes many independently schedulable atoms, which the TPC Scheduler places onto individual thread processing clusters with TPC stealing. Two supporting models carry the resource-efficiency claims: a scaling curve of the form $l = m/t + b$ fitted from measured latencies at one TPC and at all TPCs, used to right-size each kernel under a user-specified latency slip $k$, and a sequence-based DVFS model that weights each kernel's frequency sensitivity by its share of a stream's total runtime to pick a safe final frequency.","core_discovery":"The central claim is that a transparent OS layer for GPUs can decouple kernel submission from execution by atomizing kernels at runtime and scheduling atoms at TPC granularity, eliminating head-of-line blocking without relying on hardware preemption. Atomization works by launching a prelude kernel that reads each thread block's index and calls the original kernel's entry point only for blocks in the atom's index range, which is correct provided thread blocks are independent. The TPC scheduler uses predicted atom durations and per-TPC timers to steal idle TPCs from underutilized workloads while protecting latency-critical work through layered priorities and limits on outstanding atoms. On top of this, a two-point Amdahl-style scaling model right-sizes TPC allocation per kernel, and a sequence-weighted frequency model guides DVFS, together yielding the paper's reported latency, capacity, and energy numbers.","pith_inferences":["If atomization is safe in production, GPU scheduling gains a software-only analogue of CPU preemption, which could unlock standard OS policies such as priority inheritance, fairness, and work conservation on GPUs; the paper only hints at these policies.","Atomization correctness depends on thread-block independence, so the natural test is to run grids that use grid-wide synchronization or cooperative groups; the paper does not identify or test such kernels.","The reported right-sizing savings should grow on future GPUs with more TPCs, since finer scheduling granularity makes it easier to match allocation to a kernel's intrinsic parallelism.","The same interposition layer could extend the atom/TPC abstraction to other contended resources such as memory bandwidth or L2 capacity, but the paper does not evaluate that direction."],"forward_implications":["Datacenter GPUs can be shared among latency-critical and best-effort ML workloads without application, framework, or compiler changes; a driver-level library is sufficient.","A long-running training kernel no longer blocks latency-critical inference for its full duration, because atomization bounds head-of-line blocking to the duration of a single atom.","Kernel-level right-sizing captures capacity savings that whole-model partitioning misses, since individual kernels within one model scale very differently with TPC count.","Transparent DVFS can save about a quarter of GPU energy at a 7% P99 latency cost without offline profiling, by weighting kernels according to their contribution to stream runtime.","TPC-level spatial isolation plus stealing gives SLO attainment comparable to hardware partitioning while still allowing best-effort throughput that static partitions cannot support."],"supporting_citations":[{"why":"Supplies the MPS baseline that LithOS must beat and the concurrency substrate on which LithOS builds.","marker":"[7]"},{"why":"Supplies the MIG hardware-partitioning baseline whose coarse granularity and slow reconfiguration motivate dynamic TPC allocation.","marker":"[10]"},{"why":"Supplies TGS, the best prior transparent-sharing baseline for the hybrid inference/training comparison.","marker":"[48]"},{"why":"Supplies REEF, a spatial-sharing baseline that throttles best-effort work and is compared against atomization.","marker":"[20]"},{"why":"Supplies Orion, an interference-aware fine-grained GPU sharing baseline requiring application modification and offline profiling.","marker":"[44]"},{"why":"Supplies the Azure LLM inference trace used to generate the LLM prompt-length workloads in the motivation and evaluation.","marker":"[43]"},{"why":"Supplies the MLPerf inference latency constraints used as SLO targets in the multitenancy experiments.","marker":"[38]"},{"why":"Supplies the CUDA occupancy query used by the right-sizing filtering heuristic to bound useful TPC allocations.","marker":"[33]"},{"why":"Documents the over-5-second MIG reconfiguration overhead that motivates dynamic, on-the-fly TPC allocation.","marker":"[47]"}],"fun_headline_variants":["GPU OS atomizes kernels, slashes tail latency 13x","LithOS: TPC-level kernel scheduling cuts GPU latency 13x","Kernel atomization on GPUs cuts tail latency by 13x","GPU OS steals idle compute clusters to cut latency 13x"],"cache_read_input_tokens":24448,"weakest_assumption_plain":"Two premises carry LithOS: that thread blocks in a kernel are independent enough that splitting a grid by block-index ranges and letting out-of-range blocks exit early never changes the result, and that a CUDA driver-interposition library on top of MPS can actually enforce per-TPC allocation; if either fails, the reported latency and isolation numbers do not generalize.","fun_headline_variants_meta":{"raw":{"variants":["GPU OS atomizes kernels, slashes tail latency 13x","LithOS: TPC-level kernel scheduling cuts GPU latency 13x","Kernel atomization on GPUs cuts tail latency by 13x","GPU OS steals idle compute clusters to cut latency 13x"]},"model":"deepseek-v4-flash","effort":"low","cost_usd":0.000551,"raw_usage":{"total_tokens":2693,"prompt_tokens":1073,"completion_tokens":1620,"prompt_tokens_details":{"cached_tokens":384},"prompt_cache_hit_tokens":384,"prompt_cache_miss_tokens":689,"completion_tokens_details":{"reasoning_tokens":1544}},"tokens_in":689,"tokens_out":1620,"duration_ms":9559,"temperature":1.0,"reasoning_tokens":1544,"cache_read_input_tokens":384,"cache_creation_input_tokens":0},"cache_creation_input_tokens":0},"created_at":"2026-08-16T11:26:08.919169+00:00","model_set":{"reader":"deepseek-v4-flash"},"falsifier":"Run LithOS on kernels that use grid-wide synchronization, cooperative groups, or cross-block communication and compare outputs bit-for-bit against monolithic launches; any divergence shows atomization is not universally correct. Separately, with MPS active, use hardware counters to check whether a high-priority workload's TPC allocation is truly respected while a best-effort kernel is running.","supporting_citations":[{"cited_title":null,"cited_arxiv_id":null,"evidence_quote":"Supplies the MPS baseline that LithOS must beat and the concurrency substrate on which LithOS builds."},{"cited_title":null,"cited_arxiv_id":null,"evidence_quote":"Supplies the MIG hardware-partitioning baseline whose coarse granularity and slow reconfiguration motivate dynamic TPC allocation."},{"cited_title":null,"cited_arxiv_id":null,"evidence_quote":"Supplies TGS, the best prior transparent-sharing baseline for the hybrid inference/training comparison."},{"cited_title":null,"cited_arxiv_id":null,"evidence_quote":"Supplies REEF, a spatial-sharing baseline that throttles best-effort work and is compared against atomization."},{"cited_title":null,"cited_arxiv_id":null,"evidence_quote":"Supplies the CUDA occupancy query used by the right-sizing filtering heuristic to bound useful TPC allocations."}],"review_version":1}