Nova: An End-to-End MLIR Compiler for Deep Learning

Adwaid Suresh; Aparna A; Harshini V M; Jona Delcy C A; Killi Uma Maheswara Rao; Ram Charan Golla; Surendra Vendra

REVIEW 4 major objections 4 minor 39 references

Nova: An End-to-End MLIR Compiler for Deep Learning

T0 review · 4 major / 4 minor · reviewed 2026-08-04 · deepseek-v4-flash

Challenge this review Re-run · record.json Download PDF Read on arXiv ↗

Pith's one-line read A fused whole-training-step compiler claims to match hand-tuned GPU libraries while cutting memory up to 29%.

desk verdict A plausible training-step compiler with a real memory-fusion idea, but the evaluation doesn't back the headline claims — numbers are internally inconsistent and 'trains' is never verified. read the letter →

arxiv 2608.00029 v1 pith:QS4PD7Z7 submitted 2026-07-15 cs.AI cs.ARcs.LGcs.PL

Adwaid Suresh , Aparna A , Harshini V M , Jona Delcy C A , Killi Uma Maheswara Rao , Ram Charan Golla , Surendra Vendra This is my paper · ORCID

classification cs.AIcs.ARcs.LGcs.PL

keywords deeplearningcompilerMLIRJITcompilationoperatorfusionanalyticschedulingtensorcoresmemoryfootprintarithmeticintensity

verification ladder T0 review T1 audit T2 compute T3 formal

The pith

A machine-rendered reading of the paper's core claim, the machinery that carries it, and where it could break.

The reading

Nova is an end-to-end JIT compiler that treats a full training step—forward and backward passes together—as one fused computation graph. It claims that by deriving execution schedules deterministically from arithmetic intensity and hardware parameters, rather than searching over configurations, it matches or modestly exceeds cuBLAS and XLA on TF32 matrix multiplications for most shapes. On a 42-million-parameter model, Nova reports up to 10.6% higher throughput than PyTorch and 4.4% higher than XLA; and with up to 29% lower memory use than PyTorch, it trains a 144-million-parameter model on a 12 GB GPU where PyTorch runs out of memory. The paper is a systems implementation report; the load-bearing idea is that whole-graph visibility plus analytic scheduling removes the need for autotuning and hand-written kernels.

What carries the argument

The central machinery is the interplay of three components. The nova dialect represents the full training step, forward and backward, as a single MLIR function with strict value semantics, dissolving the eager-execution boundary and enabling whole-graph fusion. The Analytic Configurator then classifies each contraction by arithmetic intensity (AI = 2MNK / (MK + KN + MN)) and, using queried device parameters, deterministically embeds an execution schedule—tile sizes, warp grids, shared-memory promotion, and MMA intrinsics—into the IR as a #nova.lowering_config attribute. A structural-hashing runtime binds compiled executables to live memory via address-independent hashes, ensuring compilation

What would settle it

A concrete test: run the Analytic Configurator on a grid of matmul shapes (e.g., all combinations of M, N, K ∈ {128, 512, 2048, 8192}) and compare the generated TF32 kernels against cuBLAS or a short autotuning search (e.g., 100 configurations) on the same GPU. If Nova's deterministically chosen schedule loses by more than 5% on a substantial fraction of shapes, or if a different GPU generation (e.g., A100 or H100) yields a different optimal schedule that the analytic model cannot predict from its parameters, the central premise fails.

Watch

Extended reading notes

Core claim

Nova's central discovery is that compiling the entire training step—forward activations and backward gradients expressed natively in a single value-semantic MLIR dialect—enables aggressive cross-operator fusion and memory reuse that eager frameworks cannot achieve. A deterministic 'Analytic Configurator' reads device limits (SM count, shared memory, tensor-core shapes) and uses arithmetic intensity to select tile sizes, warp distributions, and MMA intrinsics, eliminating search. With a structural-hashing runtime that caches compiled executables by graph topology rather than memory addresses, Nova synthesizes kernels directly from computation structure. Measured on an RTX 3060, the resulting

Load-bearing premise

The paper's central claim rests on the assumption that a static analytic model—arithmetic intensity plus queried device parameters—is sufficient to deterministically select near-optimal execution schedules (tile sizes, warp distributions, MMA intrinsics) without any search; if that fails on other shapes or GPU architectures, the 'zero-search optimal scheduling' contribution reduces to a hand-tuned configuration for the one benchmarked GPU (RTX 3060).

Editorial extensions

If this is right

If correct, eager deep learning frameworks can gain the benefits of whole-graph fusion without requiring users to write custom kernels or DSL code.
The analytic configurator implies that per-shape autotuning (as in TVM/Ansor or Triton's heuristic search) can be replaced by a deterministic derivation from arithmetic intensity and device limits, reducing compile time from hours to milliseconds.
The memory reductions (14–29% vs PyTorch) suggest that fused training steps can train larger models on the same GPU, directly addressing OOM constraints on consumer hardware.
The unified forward-backward IR and static DDP callback injection imply that distributed data-parallel overlap can be achieved without runtime hooks, preserving fused execution in multi-GPU training.
The measured parity with cuBLAS/XLA on TF32 matmuls suggests a path to hardware portability: since schedules derive from queried device parameters, retargeting requires updating the hardware database rather than rewriting libraries.

Reading between the lines

Editorial extensions of the paper, not claims the author makes directly.

The 3–5% gap versus empirical search conceded for microarchitectural effects implies that purely analytic scheduling will likely plateau below search-based methods on complex fusions, until the model incorporates L2 queuing or similar effects.
The single worked example for the configurator (4096³ → 128×128 tiles, K-step 32, 2×2 warps) and the absence of the mapping equations leaves the generality of 'deterministic optimal' scheduling untested; a natural extension would be to check whether the same derivation holds across a shape grid and across GPU generations.
The structural-hashing cache could extend to shape-generic kernels: if compiled executables were parameterized by symbolic shapes rather than specialized per shape, the runtime would avoid recompilation on dynamic shapes, addressing the paper's own limitation on autoregressive workloads.
Since the compiler targets PTX only, one testable prediction is whether the same analytic core transfers to CDNA or Hopper by updating only the hardware parameter database and the low-level vector/swizzle passes; the paper states this as future work but does not demonstrate it.

Share X Bluesky LinkedIn Reddit HN

Editorial analysis

A structured set of objections, weighed in public.

Desk editor's note, referee report, and a circularity audit.

Desk Editor's Note

read the letter

You should know this paper has a real architectural idea and a hole in its evidence big enough to walk through. The idea: compile the entire training step (forward and backward) into one fused MLIR graph, schedule it analytically with no autotuning, and cache by structural hash. That's a sensible direction, and the memory mechanism — fusing the LM head and softmax so the logits tensor never materializes — is concrete and plausible.

What the paper does well: the lowering pipeline is described clearly, the standalone matmul benchmark compares against cuBLAS, CUTLASS/XLA, and AutoTVM, which is the right comparison. The paper also honestly admits a 3–5% gap versus empirical search and lists limitations. The architectural choices are well motivated, and the writing is straightforward.

The soft spots are substantial. The load-bearing claim is that Nova 'successfully trains a 144M model.' Nowhere do they show a loss curve, a gradient comparison against autograd, or any evidence that the fused backward ops (nova.sce_backward, nova.gelu_backward) compute what their names say. Standalone matmul accuracy does not validate a whole training graph. In a compiler paper that omission is a precondition, not a luxury — an optimally scheduled wrong graph is still wrong. I agree with the stress-test note.

The numbers also clash internally. The abstract says 4.4% greater than XLA, but Table 4 gives roughly 19% for 42M. The abstract says 17,900 tokens/s for 144M, but Table 4 says 33–32 thousand. The model size jumps between 42M and 46M. The '10⁻⁴ relative error band' in the text conflicts with Table 2 showing 4.6e-4. These are the kind of inconsistencies that make me doubt every other metric until I see the artifact.

The Analytic Configurator is also under-specified: one worked example, a vague arithmetic-intensity formula, and no mapping equations. The claim of 'optimal' schedules without search is plausible but not defensible from what's written.

For whom is this paper? Compiler and ML-systems researchers who care about training-time compilation and the evidence bar for performance claims. It deserves a serious referee — the architecture is substantive and the idea is worth engaging — but the paper needs major revision: show the training actually converges, fix the numbers, and specify or release the configurator. I would not cite it in its current form.

Referee Report

4 major / 4 minor

Summary. The paper introduces Nova, an end-to-end JIT compiler for deep learning training built on MLIR, which traces a full eager training step, unifies forward and backward passes into a single custom dialect (nova), and lowers it to GPU kernels through a pipeline that fuses operations, performs deterministic analytic scheduling, and generates NVPTX code via structured codegen. The key claims are: (1) an Analytic Configurator can deterministically choose optimal tile sizes, warp distributions, and MMA intrinsics from arithmetic intensity with zero search; (2) a structural-hashing runtime makes JIT recompilation effectively zero-overhead; (3) Nova matches or exceeds cuBLAS/XLA/AutoTVM on standalone TF32 matmuls; (4) end-to-end, Nova trains MLP-only language models up to 10.6% faster than PyTorch and 4.4% faster than XLA, and uses 14–29% less memory, fitting a 144M-parameter model on a 12 GB GPU where PyTorch OOMs. The evaluation also includes preliminary distributed (DDP) and mixed-precision results. The central claim is that whole-graph fusion plus analytic, search-free scheduling yields both higher throughput and lower memory than eager frameworks.

Significance. If the central claims were fully substantiated, the paper would be a meaningful contribution to the deep-learning compiler literature: it would demonstrate that a fully fused forward+backward training step can be scheduled analytically to match or beat hand-tuned kernel libraries, while lowering memory footprint and eliminating autotuning. The paper has real strengths: it is built on a concrete MLIR-based implementation, it compares against strong baselines (cuBLAS via PyTorch, CUTLASS via XLA, AutoTVM), it reports numerical accuracy of generated TF32 matmuls against an FP64 reference, and it measures memory and throughput on a modern consumer GPU. However, the load-bearing elements of the central claim—training correctness and the analytic optimality of the Configurator—are not established. The absence of any loss curves, gradient checks, or comparison against eager autograd for the fused training step means the reported throughput and memory numbers could in principle measure a silently incorrect computation. The Analytic Configurator's decision procedure is described only through one example and a 3–5% performance gap is admitted, so the 'deterministically derived optimal schedu

major comments (4)

[§6.2, Table 4] The abstract and §6.2 claim that Nova 'successfully trains' a 144M-parameter model, and the introduction states that whole-graph fusion preserves numerical fidelity. However, no evidence is provided that the compiled fused forward+backward graph computes correct gradients or reduces loss: there is no loss curve, no gradient norm check, and no comparison against PyTorch's autograd. The microbenchmarks in §6.1 validate only standalone matmuls; the end-to-end section reports throughput and memory only. Since Nova introduces custom fused backward ops (nova.sce_backward, nova.gelu_backward) and a fused LM-head/softmax, a silent lowering bug would invalidate every training throughput and memory claim. Please add an end-to-end numerical correctness check: for example, compare loss/gradients from the compiled step to a reference eager implementation across a few steps, and show that the training
[§4.3, §7.1] The Analytic Configurator is the key differentiator, but the paper never states the mapping from arithmetic intensity (AI) and device parameters to tile sizes, warp distributions, and MMA intrinsics. Only one worked example is given (4096³ → 128×128 blocks, K-step 32, 2×2 warp grid), with no equations, thresholds, or algorithmic description. §7.1 admits a 3–5% performance gap versus empirical search due to unmodeled microarchitectural effects, which is reasonable as a limitation but directly contradicts the text's claims that schedules are 'optimal' and 'derived, not guessed.' As written, the Configurator could be a small set of hand-picked rules tuned to the RTX 3060, and the 'zero-search optimal scheduling' contribution would not be substantiated. Please provide the full analytic model (or a verifiable reference to it), including the decision criteria, and validate it across a wider se
[§6.2, Table 3] The table header reads '46M-Parameter MLP' but the abstract, §6.2, and Table 4 all refer to a '42-million parameter model.' More importantly, the given configuration (vocab 50304, nembd 384, nlayers 6, MLP-only) does not produce 42M parameters: the embedding alone is 50304×384 ≈ 19.3M, the six MLP blocks each with 384→1536 and 1536→384 projections add roughly 1.2M per block, and the LM head (if not tied) adds another 19.3M, yielding roughly 26–46M depending on tying and additional norms. This inconsistency undermines the reproducibility of the model-level results and the specific throughput and memory numbers. Please correct the parameter counts and provide a clear model architecture table that matches the experimental results.
[§6.3, §7.4] The distributed (DDP) and mixed-precision sections are explicitly described as 'under active development' in the first paragraph of §6.3, yet the paper reports concrete memory and throughput numbers for DDP and claims that Nova 'sustains the highest throughput' and achieves 'perfect network-compute overlap.' This is contradictory: preliminary, in-development results should not be presented with the same confidence as the single-device results. Moreover, §7.4 later concedes that static communication scheduling causes thread blocks to stall entirely during transient network lags, which undermines the 'perfect overlap' claim. Either remove or explicitly relabel these results as preliminary feasibility demonstrations, and temper the language accordingly.

minor comments (4)

[§6.1, text after Table 2] The text says 'All tested implementations remained within the 10−4 relative error band,' but Table 2 reports values up to 4.59×10−4, which is not within 10−4. The abstract's '<5e−4 relative error' is the correct framing; please make the in-text statement consistent.
[§5.1] The claim of a 'guaranteed 100% cache hit rate' is definitional given that the hash is defined to ignore all inputs that change between iterations. It would be clearer to state that for static graphs the hash is invariant across steps and therefore recompilation is avoided, rather than presenting a 'guarantee' that could be read as an empirical result.
[§6.2, Table 3] The table title '46M-Parameter MLP' appears to be a typo for '42M-Parameter MLP'; the same table lists 'Global batch 65,536' but it is unclear how this relates to batch size B=8 and context length T=1024—please clarify the global batching configuration.
[Throughout] Several claims are phrased with absolute certainty ('guarantees optimal network-compute overlap', 'exactly once', 'absolute control') in places where the paper itself later qualifies the claims (e.g., §7.1, §7.2, §7.4). Please harmonize the language so the contributions are presented accurately without overclaiming.

Circularity Check

1 steps flagged · score 2.0 of 10

No load-bearing circularity: headline performance claims are externally benchmarked; only minor definitional cache-hit statement.

self definitional [§5.1 (Structural Hashing for JIT Caching)]
"the hasher explicitly ignores all concrete tensor pointers. Instead, it computes a hash based entirely on graph topology (opcodes, normalized IDs, sorted attributes) and static metadata (tensor shapes, strides, dtypes, and target devices). ... Because a model’s architecture remains static while its memory shifts, this structural hash guarantees a 100% cache hit rate after the first iteration."

The 100% cache-hit guarantee is entailed by the hash's construction: because concrete tensor addresses are deliberately excluded from the hash, address changes cannot change the hash for a static graph. Thus the 'guarantee' is a restatement of the design rather than an independently measured property. It is not load-bearing for the headline performance results, which rest on measured comparisons against cuBLAS/XLA/PyTorch/eager baselines.

full rationale

The paper's central claims are empirical and externally anchored: standalone TF32 matmuls are compared against cuBLAS (via PyTorch), CUTLASS (via XLA), and AutoTVM; end-to-end throughput and memory are compared against PyTorch, XLA, and an eager baseline on the same RTX 3060. These results are measured, not derived from a fitted model or from a self-citation chain, so they do not reduce to the paper's inputs. The Analytic Configurator's 'optimal' scheduling is an internal modeling claim and is explicitly qualified in §7.1 by a conceded 3–5% gap versus empirical search; that is an overclaim/limitation, not circularity. The 'trains' claim for the 144M model lacks a loss curve or convergence check, but this is an omitted verification, not a circular derivation. There are no load-bearing self-citations and no imported uniqueness theorem. The only definitional element is the 100% cache-hit guarantee in §5.1, which follows by construction from the address-independent hash; it is minor and does not affect the external performance comparisons. Overall circularity is low.

Assumptions & free parameters 4 free parameters · 5 assumptions · 0 invented entities

The paper contributes no new physical entities — the nova dialect, analytic configurator, and structural hashing runtime are software constructs evidenced only by internal IR snippets. The load-bearing postulates are about the sufficiency of a static scheduling model, the soundness of inherited MLIR/IREE lowering, and the correctness of the paper's own autodiff ops. Four hand-picked quantities (configurator thresholds, pipeline-depth fallback, DDP bucket size, swizzle constants) function as free parameters. The central new claim's burden is carried by the domain assumption that arithmetic intensity plus device parameters determines the optimal schedule — an assumption the paper partially retracts in §7.1.

free parameters (4)

Analytic Configurator schedule thresholds = 128x128 block, K-step 32, 2x2 warp grid, 3-stage pipeline (AI≈2730 example)
The mapping from arithmetic intensity + device parameters to tile sizes/warp layout/pipeline depth is asserted as deterministic (§4.3) but only one worked example is given; the thresholds are effectively hand-picked and unstated, so the schedule choice is not independently reproducible.
Adaptive pipeline depth fallback = 3 stages, falling back to 2 on shared-memory overflow
Pipeline depth is chosen by a capacity check (§4.5); the fallback rule is a heuristic, not derived from a stated analytic model.
DDP communication bucket size = 25 MB
Bucket size in §3.2 is stated without justification; it trades overlap granularity against all-reduce efficiency.
Shared-memory XOR swizzle constants = bank = column XOR ((row & 7) << 2)
The mask/shift constants are given for the 128x32-tile example (§4.5); the rule for choosing them across tile shapes is not specified.

assumptions (5)

domain assumption Arithmetic intensity (Roofline) plus device parameters is sufficient to determine the near-optimal execution schedule for a contraction.
Central premise of the Analytic Configurator (§4.3). The paper concedes a 3–5% gap versus empirical search from microarchitectural effects the model ignores (§7.1), so this is an unverified sufficiency assumption.
domain assumption MLIR/IREE structured codegen (tiling, vectorization, nvgpu lowering) is a sound base on which to build the backend.
The pipeline inherits correctness from MLIR/IREE passes [11] and PTX semantics (§2, §4); the paper adds passes on top without machine-checked verification.
domain assumption RTX 3060 device parameters queried at compile time fully capture scheduling-relevant hardware behavior.
All evaluations are on one Ampere GPU (§6); generality to Hopper/Blackwell/AMD is a promise (§7.1, §9), so the parameter model is validated only on a single device generation.
domain assumption Structural hashing over opcodes/shapes/strides/dtypes/device captures all execution-relevant state, and generated kernels never depend on tensor address identity.
The "100% cache hit rate" claim (§5.1) requires address-independent code and that no in-place alias or side-effect hidden from the hash changes behavior across iterations; this is asserted, not verified.
domain assumption Reverse-mode autodiff semantics of the nova dialect ops (nova.sce_backward, nova.gelu_backward, nova.linear_backward) are correct.
Numerical fidelity is checked only against FP64 matmul references (§6.1); training convergence is never shown, so gradient correctness is assumed (§3.1).

how reviews work

0 comments

Cite this review

Pith. "Pith review of Nova: An End-to-End MLIR Compiler for Deep Learning." pith.science (2026). https://pith.science/paper/QS4PD7Z7

@misc{pith2026260800029,
  author       = {Pith},
  title        = {Pith review of: Nova: An End-to-End MLIR Compiler for Deep Learning},
  year         = {2026},
  howpublished = {\url{https://pith.science/paper/QS4PD7Z7}},
  note         = {Machine review of arXiv:2608.00029}
}

read the original abstract

The performance of deep learning models at scale relies heavily on how effectively high-level mathematical operations are mapped to underlying physical hardware. While high-level tensor frameworks provide flexible abstractions for model design, their eager execution models inherently lack the whole-graph visibility and granular control over hardware and memory required to maximize physical hardware utilization natively. To bridge this gap, we designed Nova, an automated end-to-end JIT compiler whose defining purpose is to achieve absolute control over this hardware mapping: fusing operations across operation boundaries, optimizing complex memory hierarchies, and tuning execution down to the register level. By capturing eager executions and unifying forward and backward passes into a single value-semantic dialect, Nova unlocks aggressive whole-graph optimizations. It then utilizes an Analytic Configurator to deterministically derive optimal execution schedules based on arithmetic intensity, dropping search time to zero. Backed by a structural hashing runtime, Nova synthesizes fine-grained kernels directly from the computation's structure. In our evaluations on an RTX 3060, Nova matches or modestly exceeds cuBLAS and XLA on TF32 matmuls on most shapes, maintaining a stringent < 5e-4 relative error. At the model level, Nova achieves up to 10.6% greater throughput than PyTorch and 4.4% greater than XLA on a 42-million parameter model, without compromising on numerical fidelity. Crucially, by reducing the memory footprint by up to 29% relative to PyTorch, Nova successfully trains a 144-million parameter model at 17,900 tokens/s where PyTorch encounters Out-Of-Memory (OOM) failures on the same 12 GB consumer GPU.

Figures

Figures reproduced from arXiv: 2608.00029 by the authors.

**Figure 2.** Compiled DDP architecture. By statically injecting an llvm.call @nova_ddp_bucket_ready(k) the exact moment a communication bucket’s final gradient is computed, Nova triggers asynchronous NCCL all-reduces to achieve perfect network-compute overlap without eager framework hooks. By enforcing synchronization via static IR-level callback injection, Nova guarantees optimal network-compute overlap and distributed scala… view at source ↗

**Figure 4.** Software-pipelined K-loop execution. By maintaining depth − 1 asynchronous copy groups in flight, the tensor cores compute on resident slab i while the memory controller simultaneously prefetches future slabs, effectively hiding HBM latency. 2. Shared Memory Swizzling: A tensor core’s throughput relies heavily on feeding it data from shared memory as quickly as possible via ldmatrix. However, naive column-major r… view at source ↗

Figures from the paper (5 more)

**Figure 6.** Figure 6: Standalone matrix-multiplication throughput (TFLOP/s, TF32) comparing Nova against XLA, PyTorch, and AutoTVM 1. Numerical Accuracy We verified the generated TF32 kernel against a doubleprecision (FP64) reference across the six shapes. All tested implementations remain…

**Figure 7.** Figure 7: Device memory consumption (MB) across models from 42M to 144M parameters 2. Sustained Throughput Simultaneously, Nova matches or exceeds the sustained training throughput (tokens/sec) of PyTorch, the eager baseline, and XLA across all viable model sizes. By operating …

**Figure 8.** Figure 8: Sustained training throughput (tokens/sec) 6.3 Distributed Training and Mixed Precision Note that both Distributed Data Parallel (DDP) and hardwarenative mixed precision are currently under active development. However, we report our preliminary metrics below. We exte…

**Figure 9.** Figure 9: Device memory consumption (MB) in DDP training 2. Sustained Throughput (DDP) In addition to memory savings, Nova sustains the highest throughput. It efficiently overlaps gradient synchronization with computation. Params Nova PyTorch Eager XLA 42M 137–132 127–124 135–13…

**Figure 10.** Figure 10: Device memory consumption (MB) in Mixed precision training 4. Sustained Throughput (Mixed Precision) Nova’s throughput remains competitive but falls slightly behind PyTorch in mixed precision execution. 7 Discussions and limitations In this section, we analyze Nova’…

Discussion (0). Continue with ORCID to comment.

Reference graph

Works this paper leans on

39 extracted references · 5 linked inside Pith

[1]

By capturing both forward activations and backward gradients into a unified graph, it unlocks the global visibility required for whole-step optimizations

Front-End Tracing:Built to overcome the opaque dis- patch of eager frameworks, the front-end dynamically traces BluTrain API calls. By capturing both forward activations and backward gradients into a unified graph, it unlocks the global visibility required for whole-step optimizations
[2]

This represents the full training step natively within 2 MLIR, serving as the definitive foundation for both cross- operator fusion and compiler-native gradient synchronization

Intermediate Representation (IR) (§3):Designed to provide a mathematically formal structure for optimization, the IR Emitter lowers the captured trace into the custom nova dialect. This represents the full training step natively within 2 MLIR, serving as the definitive foundation for both cross- operator fusion and compiler-native gradient synchronization
[3]

Lowering and Code Generation (§4):Built to determin- istically map algorithmic logic to physical hardware with no autotuning search, this five-stage pipeline enforces the com- piler’s absolute control. The Analytic Configurator injects exact tile sizes and warp mappings, which are subsequently preserved through memory hierarchy passes and final vector- iz...
[4]

parallel

JIT Runtime (§5):Designed to ensure aggressive compila- tion does not bottleneck the dynamic training loop, the LLVM ORC JIT caches the generated artifacts. An ABI adapter binds the executable directly to live device memory, guaranteeing that the heavy cost of whole-graph optimization is paid exactly once. 3 Graph Tracing and Intermediate Representation T...
[5]

Nova time-shifts nvgpu.device_async_copy operations by depth − 1 iterations, guaranteeing that itera- tion i+2 prefetches from HBM while iteration i computes on resident memory

Adaptive Multi-Buffering & Software Pipelining:The central K-loop is software-pipelined to maintain continuous data flow. Nova time-shifts nvgpu.device_async_copy operations by depth − 1 iterations, guaranteeing that itera- tion i+2 prefetches from HBM while iteration i computes on resident memory. Crucially, this depth adaptively scales; if the footprint...
[6]

Shared Memory Swizzling:A tensor core’s through- put relies heavily on feeding it data from shared mem- ory as quickly as possible via ldmatrix. However, naive column-major reads by a warp cause severe 32-way mem- ory bank conflicts because multiple threads attempt to ac- cess the same physical memory bank simultaneously, forc- ing the hardware to seriali...
[7]

The hardware ldmatrix instruction is highly rigid, and its ldmatrix.trans variant is strictly limited to 16-bit granularity

Precision-Aware ldmatrix Routing:When loading frag- ments from shared memory to the tensor cores, Nova dynam- ically alters its instruction emission based on operand preci- sion and tensor layout. The hardware ldmatrix instruction is highly rigid, and its ldmatrix.trans variant is strictly limited to 16-bit granularity. Because full-precision float32 (TF3...
[8]

To optimize them, Nova pins the accumulators directly in registers and expands clustered subgroup reductions into gpu.shuffle 6 butterfly trees

Butterfly Warp Reductions:Reductions bypass the MMA pipeline and lack a native tensor-core intrinsic. To optimize them, Nova pins the accumulators directly in registers and expands clustered subgroup reductions into gpu.shuffle 6 butterfly trees. This maps the reduction directly onto native PTX warp-level instructions, collapsing the row entirely in regis...

Show all 39 references

[9]

All tested implementations remained within the 10−4 relative error band

Numerical Accuracy We verified the generated TF32 kernel against a double- precision (FP64) reference across the six shapes. All tested implementations remained within the 10−4 relative error band. 7 Nova’s maximum relative error (4.6×10 −4) falls well inside the expected TF32...
[10]

At 144M parameters, both PyTorch and the eager baseline run out of memory (OOM) on the 12 GB card, whereas Nova and XLA successfully fit and continue training

Memory Consumption Nova dictates the lowest memory floor at nearly every model size, consuming 14–29% less steady-state device memory than PyTorch and consistently less than the eager baseline. At 144M parameters, both PyTorch and the eager baseline run out of memory (OOM) on ...
[11]

By operating en- tirely on a single fused graph, Nova proves that on a memory- constrained GPU, it is possible to lower the memory floor and raise throughput simultaneously

Sustained Throughput Simultaneously, Nova matches or exceeds the sustained training throughput (tokens/sec) of PyTorch, the eager base- line, and XLA across all viable model sizes. By operating en- tirely on a single fused graph, Nova proves that on a memory- constrained GPU, ...
[12]

At 144M parame- ters, PyTorch and XLA both encounter OOM errors, whereas Nova successfully fits the DDP step

Memory Consumption (DDP) Nova maintains the lowest memory footprint across all tested model scales in a distributed setting. At 144M parame- ters, PyTorch and XLA both encounter OOM errors, whereas Nova successfully fits the DDP step. 8 Figure 9:Device memory consumption (MB) ...
[13]

It efficiently overlaps gradient synchronization with computation

Sustained Throughput (DDP) In addition to memory savings, Nova sustains the highest throughput. It efficiently overlaps gradient synchronization with computation. Params Nova PyTorch Eager XLA 42M137–132127–124 135–132 115–112 92M66–6464–63 64–63 OOM 106M53–5152–51 51–50 OOM 1...
[14]

Figure 10:Device memory consumption (MB) in Mixed preci- sion training

Mixed Precision In mixed-precision scenarios, PyTorch compile mode main- tains a slightly lower memory footprint across all model scales. Figure 10:Device memory consumption (MB) in Mixed preci- sion training
[15]

7 Discussions and limitations In this section, we analyze Nova’s architectural trade-offs, evaluating the limits of its performance characteristics and open challenges for scaling

Sustained Throughput (Mixed Precision)Nova’s through- put remains competitive but falls slightly behind PyTorch in mixed precision execution. 7 Discussions and limitations In this section, we analyze Nova’s architectural trade-offs, evaluating the limits of its performance cha...
[16]

IREE: Intermediate Representation Execution En- vironment

The IREE Authors. “IREE: Intermediate Representation Execution En- vironment.” GitHub repository, https://github.com/iree-org/ iree
[17]

OpenXLA: A compiler for machine learning

The OpenXLA Authors. “OpenXLA: A compiler for machine learning.” GitHub repository,https://github.com/openxla/xla
[18]

Operator Fusion in XLA: Analysis and Eval- uation

D. Snider and R. Liang. “Operator Fusion in XLA: Analysis and Eval- uation.” arXiv:2301.13062, 2023

2023 arXiv
[19]

MLIR: Scaling Compiler Infrastructure for Domain- Specific Computation

C. Lattner et al. “MLIR: Scaling Compiler Infrastructure for Domain- Specific Computation.” CGO, 2021

2021
[20]

Halide: A Language and Compiler for Opti- mizing Parallelism, Locality, and Recomputation in Image Processing Pipelines

J. Ragan-Kelley et al. “Halide: A Language and Compiler for Opti- mizing Parallelism, Locality, and Recomputation in Image Processing Pipelines.” PLDI, 2013

2013
[21]

Tensor Comprehensions: Framework- Agnostic High-Performance Machine Learning Abstractions

N. Vasilache et al. “Tensor Comprehensions: Framework- Agnostic High-Performance Machine Learning Abstractions.” arXiv:1802.04730, 2018

2018 arXiv
[22]

cuDNN: Efficient Primitives for Deep Learning

S. Chetlur et al. “cuDNN: Efficient Primitives for Deep Learning.” arXiv:1410.0759, 2014

2014 arXiv
[23]

TVM: An Automated End-to-End Optimizing Compiler for Deep Learning

T. Chen et al. “TVM: An Automated End-to-End Optimizing Compiler for Deep Learning.” OSDI, 2018. arXiv:1802.04799

2018 arXiv
[24]

PyTorch: An Imperative Style, High-Performance Deep Learning Library

A. Paszke et al. “PyTorch: An Imperative Style, High-Performance Deep Learning Library.” NeurIPS, 2019. 10

2019
[25]

TensorFlow: A System for Large-Scale Machine Learning

M. Abadi et al. “TensorFlow: A System for Large-Scale Machine Learning.” OSDI, 2016

2016
[26]

Composable and Modular Code Generation in MLIR: A Structured and Retargetable Approach to Tensor Compiler Construction

N. Vasilache et al. “Composable and Modular Code Generation in MLIR: A Structured and Retargetable Approach to Tensor Compiler Construction.” arXiv:2202.03293, 2022

2022 arXiv
[27]

Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations

P. Tillet, H. T. Kung, and D. Cox. “Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations.” MAPL, 2019

2019
[28]

Data Movement Is All You Need: A Case Study on Optimizing Transformers

A. Ivanov et al. “Data Movement Is All You Need: A Case Study on Optimizing Transformers.” MLSys, 2021

2021
[29]

The MLIR Transform Dialect: Your Compiler Is More Powerful Than You Think

M. Lücke, O. Zinenko, W. S. Moses, M. Steuwer, and A. Cohen. “The MLIR Transform Dialect: Your Compiler Is More Powerful Than You Think.” CGO, 2024

2024
[30]

Ansor: Generating High-Performance Tensor Programs for Deep Learning

L. Zheng et al. “Ansor: Generating High-Performance Tensor Programs for Deep Learning.” OSDI, 2020

2020
[31]

cuBLAS Library User Guide

NVIDIA. “cuBLAS Library User Guide.” NVIDIA Corporation. https://docs.nvidia.com/cuda/cublas/
[32]

CUTLASS: CUDA Templates for Linear Algebra Subroutines

V . Thakkar et al. “CUTLASS: CUDA Templates for Linear Algebra Subroutines.” NVIDIA, GitHub repository, https://github.com/ NVIDIA/cutlass
[33]

FlashAttention: Fast and Memory-Efficient Exact Atten- tion with IO-Awareness

T. Dao et al. “FlashAttention: Fast and Memory-Efficient Exact Atten- tion with IO-Awareness.” NeurIPS, 2022

2022
[34]

PyTorch Distributed: Experiences on Accelerating Data Parallel Training

S. Li et al. “PyTorch Distributed: Experiences on Accelerating Data Parallel Training.” VLDB, 2020

2020
[35]

Automatic Differentiation in Machine Learning: a Survey

A. G. Baydin, B. A. Pearlmutter, A. A. Radul, and J. M. Siskind. “Automatic Differentiation in Machine Learning: a Survey.” JMLR, 2018

2018
[36]

JAX: Composable Transformations of Python+NumPy Programs

J. Bradbury et al. “JAX: Composable Transformations of Python+NumPy Programs.” 2018. http://github.com/google/ jax

2018
[37]

Roofline: An Insightful Visual Performance Model for Multicore Architectures

S. Williams, A. Waterman, and D. Patterson. “Roofline: An Insightful Visual Performance Model for Multicore Architectures.” Communica- tions of the ACM, 2009

2009
[38]

CUDA C++ Programming Guide

NVIDIA. “CUDA C++ Programming Guide” (Warp Matrix Functions: mma.sync, ldmatrix). NVIDIA Corporation. https://docs.nvidia. com/cuda/cuda-c-programming-guide/
[39]

LLVM/MLIR 21.1.6; CUDA Toolkit 13.0

The LLVM Project. LLVM/MLIR 21.1.6; CUDA Toolkit 13.0. (Soft- ware artifact / tooling versions.) 11

Pith tools

Reviewed August 4, 2026 · model on record in the stance chip above.

[1] [1]

By capturing both forward activations and backward gradients into a unified graph, it unlocks the global visibility required for whole-step optimizations

Front-End Tracing:Built to overcome the opaque dis- patch of eager frameworks, the front-end dynamically traces BluTrain API calls. By capturing both forward activations and backward gradients into a unified graph, it unlocks the global visibility required for whole-step optimizations

[2] [2]

This represents the full training step natively within 2 MLIR, serving as the definitive foundation for both cross- operator fusion and compiler-native gradient synchronization

Intermediate Representation (IR) (§3):Designed to provide a mathematically formal structure for optimization, the IR Emitter lowers the captured trace into the custom nova dialect. This represents the full training step natively within 2 MLIR, serving as the definitive foundation for both cross- operator fusion and compiler-native gradient synchronization

[3] [3]

Lowering and Code Generation (§4):Built to determin- istically map algorithmic logic to physical hardware with no autotuning search, this five-stage pipeline enforces the com- piler’s absolute control. The Analytic Configurator injects exact tile sizes and warp mappings, which are subsequently preserved through memory hierarchy passes and final vector- iz...

[4] [4]

parallel

JIT Runtime (§5):Designed to ensure aggressive compila- tion does not bottleneck the dynamic training loop, the LLVM ORC JIT caches the generated artifacts. An ABI adapter binds the executable directly to live device memory, guaranteeing that the heavy cost of whole-graph optimization is paid exactly once. 3 Graph Tracing and Intermediate Representation T...

[5] [5]

Nova time-shifts nvgpu.device_async_copy operations by depth − 1 iterations, guaranteeing that itera- tion i+2 prefetches from HBM while iteration i computes on resident memory

Adaptive Multi-Buffering & Software Pipelining:The central K-loop is software-pipelined to maintain continuous data flow. Nova time-shifts nvgpu.device_async_copy operations by depth − 1 iterations, guaranteeing that itera- tion i+2 prefetches from HBM while iteration i computes on resident memory. Crucially, this depth adaptively scales; if the footprint...

[6] [6]

Shared Memory Swizzling:A tensor core’s through- put relies heavily on feeding it data from shared mem- ory as quickly as possible via ldmatrix. However, naive column-major reads by a warp cause severe 32-way mem- ory bank conflicts because multiple threads attempt to ac- cess the same physical memory bank simultaneously, forc- ing the hardware to seriali...

[7] [7]

The hardware ldmatrix instruction is highly rigid, and its ldmatrix.trans variant is strictly limited to 16-bit granularity

Precision-Aware ldmatrix Routing:When loading frag- ments from shared memory to the tensor cores, Nova dynam- ically alters its instruction emission based on operand preci- sion and tensor layout. The hardware ldmatrix instruction is highly rigid, and its ldmatrix.trans variant is strictly limited to 16-bit granularity. Because full-precision float32 (TF3...

[8] [8]

To optimize them, Nova pins the accumulators directly in registers and expands clustered subgroup reductions into gpu.shuffle 6 butterfly trees

Butterfly Warp Reductions:Reductions bypass the MMA pipeline and lack a native tensor-core intrinsic. To optimize them, Nova pins the accumulators directly in registers and expands clustered subgroup reductions into gpu.shuffle 6 butterfly trees. This maps the reduction directly onto native PTX warp-level instructions, collapsing the row entirely in regis...