arXiv CSopen access

Instant GPU Efficiency Visibility at Fleet Scale

Instant GPU Efficiency Visibility at Fleet Scale Connor Pedersen, Dong H. Ahn, Michel Migdal, Collin Neale, Nik Konyuchenko

arXiv:2605.20799v1 [cs.DC] 20 May 2026

NVIDIA {connorp, donga, mmigdal, cneale, nkonyuchenko}@nvidia.com

Abstract—We present Overall FLOP Utilization (OFU), a hardware-level, precision-agnostic GPU efficiency metric for AI workloads on HPC systems, derived from two on-chip performance counters: Tensor Pipe Activity and SM clock frequency. OFU requires no application instrumentation and works across GPU generations and numeric precisions. We characterize five properties of OFU approximation—tile quantization, floatingpoint precision scaling, clock sampling noise, Tensor Core clock domains, and non-tensor undercounting—through controlled GEMM experiments on H100 and GB200 across FP16, TF32, FP8, and NVFP4. After tile-quantization correction, OFU predicts application-level MFU to within ≤2 percentage points. Against 608 production training jobs, OFU achieves r = 0.78 correlation with application-level MFU and surfaces two framework-level FLOPs miscalculations. Deployed across large-scale GPU fleets, OFU has detected a 2.5× efficiency regression and tracked precisiondependent utilization changes in mixed-precision pretraining. Our evaluation and operational experience suggest OFU is a practical, deployment-ready complement to application-level MFU for continuous fleet-wide efficiency monitoring. Index Terms—MFU, GPU utilization, tensor cores, hardware performance counters, large-scale training

I. I NTRODUCTION The current push toward one of the largest capitalexpenditure cycles in technology history has made GPU compute efficiency a critical factor in economic viability. Microsoft spent $64.6B in capex in FY2025, Alphabet invested $91.4B in 2025, and Meta spent $69.7B for the same period— with cumulative Big Tech Artificial Intelligence (AI) High Performance Computing (HPC) infrastructure investment on track to exceed $2.8T through 2029 [1]–[4]. At this scale, each percentage point of GPU utilization recovered across a large HPC fleet1 represents millions of dollars in value; conversely, undetected inefficiencies can silently double the effective cost of compute. Fleet-wide measurement of GPU utilization which is accurate, continuous, and actionable is therefore essential to realizing the full potential of AI HPC infrastructure investment. Achieving this requires a utilization metric that can be deployed across every workload on a massive, heterogeneous fleet—without instrumenting applications, without modifying the software stack, with well-characterized accuracy, and scalable across a wide range of GPU generations—while remaining simple enough to integrate into fleet-wide resilience and goodput services at multiple levels that detect inefficiencies and drive optimization. Existing measurement techniques, however, fall short of these demands. 1We use fleet to refer to the aggregate of all GPUs deployed across one or more HPC or data centers under a single administrative domain.

In measuring GPU utilization, three broad classes of techniques exist. First, user-level performance profiling tools such as NVIDIA Nsight Compute, Nsight Systems, and open-source alternatives like PyTorch Profiler [5], DeepSpeed Flops Profiler [6], and HPCToolkit [7], measure floating-point utilization with high fidelity. However, these are per-workload tools: each user must instrument and profile their own job, they impose non-trivial overhead on the profiled application, and they are designed for targeted profiling runs, not continuous fleet-wide monitoring. Second, framework-level algorithmic throughput estimates represent an improvement: once implemented within a framework such as Megatron-LM [8], PyTorch Lightning [9], or NeMo [10], any workload built on that framework can report Model FLOPs Utilization (MFU) derived from model-architecture FLOPs2 counts [11]. However, coverage remains fragmented across frameworks, and the FLOPs formulas must be updated for each new training modality— dense, mixture-of-experts, latent-space routing, multimodal pipelines—making them brittle and error-prone as architectures evolve. Third, hardware performance counters, exposed through DCGM [12], offer a non-intrusive alternative requiring no application instrumentation and no runtime overhead, but lack a rigorous study of their accuracy and nuance, validation against controlled benchmarks, and demonstrated scalability across GPU generations. In this paper, we present GPU utilization techniques with well-characterized accuracy, designed for instant fleet-wide visibility and downstream optimization, collectively called Overall FLOP Utilization (OFU). OFU comprises four complementary elements: a hardware-counter-based metric that combines Tensor Pipe Activity with SM clock frequency to instantly expose floating-point utilization across any NVIDIA GPU generation, without application instrumentation or softwarestack modifications; error characterization against well-defined computational kernels (GEMMs) across precisions and GPU architectures, establishing bounded accuracy for the metric; a deep analysis of how these kernels are mapped to hardware, including the nuances of over- and undercounting of floatingpoint operations due to tile quantization and peak-TFLOP/s specification discrepancies; and operational experience deploying these techniques at fleet scale, demonstrating that they indeed unearth actionable areas of inefficiency. Specifically, this paper makes the following contributions: •

A first-principles derivation of the OFU metric as a simple

2 Throughout this paper, we use FLOPs to denote a count of floating-point operations and FLOP/s to denote a rate (floating-point operations per second).

product of Tensor Pipe Activity and normalized SM clock to be one of the fleet’s best-performing workloads. When frequency, yielding a precision-agnostic, architecture- measured with hardware counters, however, the job showed agnostic proxy for MFU that requires no model-specific only 25.58% MFU, a 112.2% relative discrepancy. Investigainformation. tion revealed that the framework’s FLOPs counter did not • Correlation with application-level MFU across 608 pro- account for the model’s latent-space projections, inflating duction training jobs on H100 GPUs, achieving r = 0.78. the reported FLOPs by ∼3×. Such errors are increasingly • Error characterization through controlled GEMM exper- common as emerging and diverse model architectures—mixtureiments on H100 and GB200 GPUs across FP16, TF32, of-experts, latent-space routing, multimodal pipelines—outpace FP8, and NVFP4, quantifying all major sources of error the assumptions baked into framework-level FLOPs formulas. in our approximation. Researchers understandably prioritize model quality over • Operational experiences of deploying OFU to a large throughput-calculation accuracy, and may lack the domain GPU fleet including three case studies that led to spot expertise to maintain correct FLOPs accounting as architectures and address inefficient GPU usage. evolve. Given these limitations, the review team demanded a solution Our evaluation shows that controlled GEMM experiments on that collects utilization without modifying training code or H100 and GB200 across FP16, TF32, FP8, and NVFP4 predict injecting profiling hooks; deploys with minimal changes to application-level MFU to within 1–3 percentage points, tightthe AI software stack, preserving backward compatibility; ening to ≤2 percentage points after tile-quantization correction. covers all workloads—training and inference—regardless of Against 608 production training jobs on H100 GPUs, OFU framework, model architecture, or GPU generation; integrates achieves r = 0.78 correlation with application-level MFU and readily into fleet-wide resilience and goodput services that surfaces two distinct framework-level FLOPs miscalculations. detect inefficiencies and drive optimization; provides bounded, Operationally, OFU has been deployed across large-scale GPU fleets at multiple integration levels—per-job dashboards, cluster- well-characterized accuracy against established benchmarks; wide resilience services, and automated goodput monitoring— and identifies where its approximations fall short, charting a where it detected a 2.5× efficiency regression in embodied- path toward higher-fidelity measurement in future hardware agent training and tracked precision-dependent utilization and software stacks. changes across mixed-precision foundation-model pretraining. III. OVERALL FLOP U TILIZATION Overall, our techniques demonstrate that hardware performance counters, when carefully validated and operationalized, The requirements from Section II—no application instruprovide a practical path to instant, fleet-wide GPU utilization mentation, no software-stack modifications, coverage across all visibility—meeting the demands that existing measurement workloads and GPU generations, bounded accuracy, and inteapproaches cannot. gration into fleet-wide automation—demand a metric grounded entirely in signals the hardware already exposes. Meeting II. M ANY T ECHNIQUES , N O F LEET-W IDE S OLUTION From the perspective of fleet infrastructure management— these requirements simultaneously demands three things: (a) a where the goal is to maximize the return on GPU investment universally available measurement substrate that is independent across all users and workloads—we prepared a fleet-wide of any training framework or model architecture, (b) a mapping efficiency review for a collection of large internal GPU clusters. from that substrate to floating-point throughput rooted in the The measured training MFU averaged approximately 20% over physical execution pipelines of the GPU, and (c) a quantitative a two-week window, well below the 35–50% range considered understanding of where the resulting approximation is tight and healthy. Understanding where the waste originated required where it breaks down. We satisfy (a) with on-chip performance a utilization metric available across all workloads, but from counters available through DCGM [12], satisfy (b) by deriving this infrastructure vantage point, we met with the following an MFU estimator from the first-principles structure of GPU floating-point pipelines, and satisfy (c) through controlled limitations to answering this fundamental question. Application-reported MFU requires explicit per-framework kernel-level experiments and production-scale validation. The integration: each training codebase must compute and emit remainder of this section develops the core derivation: we FLOPs counts via a reporting tool such as OneLogger [13]. first describe the floating-point execution model of modern At the time of our review, only approximately 20% of fleet NVIDIA GPUs from first principles, then show how a single workloads had been onboarded. The remaining 80% of GPU- hardware counter—Tensor Pipe Activity—combined with SM hours had no MFU measurement at all, leaving the majority clock frequency can yield a precision-agnostic utilization metric of the fleet invisible to efficiency monitoring. Onboarding we call Overall FLOP Utilization (OFU). additional workloads requires application-level code changes— a process that does not scale across a diverse fleet with hundreds of users and configurations. Even where application MFU was available, we found it could be significantly incorrect. A 288-GPU DeepSeekstyle [14] MoE training job reported 54.27% MFU—appearing

A. Floating-Point Execution in Modern GPUs A modern NVIDIA GPU organizes its compute resources into an array of Streaming Multiprocessors (SMs), each capable of issuing instructions to several independent execution pipelines in parallel [15]. Two pipeline families are relevant to

floating-point throughput: the CUDA-core (FP/INT) pipelines and the Tensor Core pipelines. Each SM contains scalar floating-point/integer (FP32, FP64, INT32) functional units organized into processing blocks, each executing one operation per thread per cycle across a warp of 32 threads. The number of FP32 CUDA cores per SM has grown with each generation—64 on Volta and Ampere (A100), and 128 on both Hopper and Blackwell [15]–[18]. These units handle general-purpose arithmetic: element-wise activations, reductions, address computation, and any floating-point work that is not matrix multiplication. Beginning with Volta, NVIDIA introduced dedicated Tensor Cores: fixed-function matrix-multiply-accumulate (MMA) units that operate at the warp level on small tiles to compute D = A×B+C in a single instruction [15], [18]. Because each MMA instruction retires an entire tile of output elements rather than one element per thread, the per-cycle throughput is substantially higher than the scalar pipeline for the same silicon area. On the H100, for example, each SM delivers 4,096 FP16 FLOPs/cycle on the Tensor Core pipeline versus 256 FP32 FLOPs/cycle on the CUDA-core pipeline—a 16× ratio even before accounting for the 2× element width difference [17]. This gap has widened with every generation as NVIDIA has invested transistor budget predominantly in Tensor Core throughput. In any workload dominated by matrix multiplications— which describes virtually all modern deep-learning training and inference—the vast majority of achieved FLOPs flow through the Tensor Core pipeline. The CUDA-core contribution to total FLOPs is negligible by comparison. This asymmetry is the physical basis for our approach: monitoring Tensor Core pipeline activity alone captures the dominant term in GPU floating-point utilization, and any FLOPs missed by ignoring the scalar pipeline are well within the noise floor of practical measurement. B. Tensor Core Evolution Across Architectures

V100 [15] to 989 TFLOP/s on H100 [17], while FP32 CUDAcore throughput has increased far more modestly. This gap makes Tensor Core activity the dominant signal for GPU compute utilization in modern AI applications. C. Deriving the MFU Estimator We use Tensor Pipe Activity (TPA) as the base signal, which reports the fraction of cycles during which the GPU is executing Tensor Core floating-point instructions across all GPU generations: Cycles GPU is executing Tensor instructions TPA = Total Cycles TPA is a hardware-averaged counter: the GPU accumulates active and total cycle counts over the collection window and reports their ratio, so a single readout already reflects the true mean activity over that interval. Thus, the per-cycle approximation is given by: Actual FLOPs/cycle TPA ≈ . Peak FLOPs/cycle TPA is measured in cycles, but MFU is a throughput ratio in FLOPs per second. The link between the two is the SM clock frequency. The theoretical peak FLOP/s of a GPU is max defined at the maximum boost clock fSM . When the GPU max runs at a lower clock fSM < fSM , every cycle still delivers the same FLOPs per cycle, but fewer cycles elapse per second, so the realized FLOP/s is reduced proportionally. Multiplying the per-cycle ratio by the clock ratio therefore converts from a cycle-domain metric to a time-domain one: fSM Actual FLOP/s Actual FLOPs/cycle × max = ≈ MFU. Peak FLOPs/cycle fSM Peak FLOP/s In practice, fSM is not constant. The GPU’s power and thermal management continuously adjusts the SM clock in response to workload intensity, power draw, and GPU temperature— even during a sustained single-kernel workload, the clock can fluctuate by hundreds of megahertz multiple times per second. Unlike TPA, which is hardware-averaged over the collection window, the SM clock reported by the hardware (DCGM_FI_DEV_SM_CLOCK) is an instantaneous point sample. A single readout can therefore over- or underestimate the true mean clock for that interval, introducing noise into the OFU estimate. Averaging OFU over many collection points mitigates this sampling bias, but characterizing the residual error is critical to establishing bounded accuracy for any hardwarecounter-based metric.

Tensor Cores first appeared in Volta (V100). At the warp level, a single HMMA instruction computes a 16 × 16 × 16 mixed-precision MMA—multiplying two FP16 input tiles and accumulating into FP32—delivering 8,192 FLOPs per instruction [15]. Subsequent architectures broadened support. Ampere introduced BF16 and TF32 on the Tensor Cores [16], [19] and replaced HMMA with the mma.sync instruction, whose native warp-level tile is 16 × 8 × 16 (FP16/BF16 inputs, FP32 accumulation). Ampere also added cp.async for asynchronous global-to-shared-memory loads, enabling multistage pipelining crucial to saturating Tensor Cores. D. Overall FLOP Utilization (OFU) Ada and Hopper added native FP8 (E4M3/E5M2). Hopper We refer to this hardware-derived estimate as Overall FLOP introduced asynchronous WGMMA instructions that increase Utilization (OFU): parallelism and efficiency by accepting inputs in shared memory fSM and supporting even larger tile sizes [20]–[22]. Most recently, OFU = TPA × max (1) fSM Blackwell added native FP4 via NVFP4 and tensor memory with tcgen05 instructions [23]–[25]. We use the term Overall FLOP Utilization to distinguish it A key trend across all generations is the widening gap from application-level MFU: OFU captures all FLOPs executed between Tensor Core and CUDA-core throughput: dense FP16 at the Tensor Core hardware level, not only the forward and Tensor Core throughput has grown from 125 TFLOP/s on backward pass FLOPs that application-level MFU typically

20 10000

Matrix Size (N)

0

15000

(a) H100 FP16

60 40 20 0

0

5000

10000

Matrix Size (N)

(e) GB200 FP16

5000

10000

Matrix Size (N)

100

NVJet

80

0

60 40 20 0

15000

0

(b) H100 FP16 (random)

Overhead (%)

Overhead (%)

100

20

CUTLASS XMMA

80

15000

60 40 20 0

0

5000

10000

Matrix Size (N)

15000

(f) GB200 TF32

10000

Matrix Size (N)

60 40 20 0

15000

NVJet

80

0

(c) H100 TF32 100

CUTLASS

80

5000

100

60 40 20 0

0

5000

10000

Matrix Size (N)

(g) GB200 FP8

10000

Matrix Size (N)

100

NVJet

80

5000

15000

(d) H100 FP8

Overhead (%)

5000

40

Overhead (%)

0

60

100

Overhead (%)

40

CUTLASS NVJet

80

Overhead (%)

60

0

100

NVJet

80

Overhead (%)

Overhead (%)

100

15000

NVJet

80 60 40 20 0

0

5000

10000

Matrix Size (N)

15000

(h) GB200 NVFP4

Fig. 1: FLOP overhead for square matrices due to tiling and cuBLAS kernel selection.

counts; and OFU is precision-agnostic, since the hardware counter measures Tensor Core activity regardless of numeric format. IV. C HARACTERIZING OFU P ROPERTIES

FP16 and TF32, we used PyTorch torch.matmul; for FP8, we used torch._scaled_mm; for NVFP4 on GB200, we used an internal matrix multiplication benchmark, as PyTorch does not yet have full native NVFP4 GEMM support at the time of writing. For each (M, K, N ) triple we executed a single matrix multiplication, collected the precision-specific tensor-op counter, and computed the FLOP overhead, the fraction of extra FLOPs the hardware executes beyond the theoretical 2M N K due to tile padding:

Before deploying OFU in production, we must understand its behavior across hardware configurations, numeric formats, and operating conditions. We examine five properties: tile quantization and cuBLAS kernel selection, where fixed-size tiles and zero-padding cause the hardware to execute more FLOPsprofiled − 2M N K FLOPs than the theoretical 2M N K cost; floating-point preciOverhead = × 100% (2) 2M N K sion, where we validate that OFU correctly tracks utilization across numeric formats (FP16, TF32, FP8, NVFP4) and GPU We tested FP16, TF32, FP8, and NVFP4 using square matrices architectures (H100, GB200); SM clock sampling noise, where from N = 128 to N = 16384 in increments of 128 (tensorthe instantaneous clock sample can diverge from the true mean, core-friendly alignment [27]), plus ∼1000 square matrices widening the confidence interval of any single OFU reading; with randomly chosen dimensions (not necessarily multiples theoretical peak FLOPs, where Tensor Core pipelines may of 128). Fig. 1 shows the measured overhead for H100 and GB200. operate at a different maximum clock frequency than the SM boost clock; and non-tensor undercounting, where OFU Two patterns are clear: overhead decreases with matrix size as padding waste becomes a smaller fraction of total work, and monitors only the Tensor Core pipeline. overhead varies by precision and GPU due to differing tile A. Tile Quantization and cuBLAS Kernel Selection sizes selected by cuBLAS [28]. For well-aligned matrices (multiples of 128) with N ≥ Tile quantization is a software constraint imposed by GEMM kernels to maximize Tensor Core throughput. Because virtually 4096, the maximum overhead observed was ∼9% across both all AI training computation reduces to matrix multiplications GPUs and all precisions, with means of 2–3%. For non-aligned (attention, linear layers, convolutions) [26], GEMM is the matrices (Fig. 1b; GB200 random omitted as it exhibits a similar natural workload for isolating this error. For C = A × B pattern), overhead at N ≥ 4096 reached up to ∼12%, though where A is M × K and B is K × N , the theoretical cost the mean remained around 5%. At small sizes (N < 512, rarely is exactly 2M N K FLOPs. Because the workload is fully used in large-model training), overhead can exceed 50% due specified, any discrepancy between OFU and the true utilization to severe tile quantization. must originate in the hardware execution—not in uncertainty As shown in Fig. 1, FP16, FP8, and NVFP4 exhibit nearly about what the application is computing. identical overhead curves on GB200: all are routed exclusively We profiled matrix multiplications on H100 and GB200 to nvJet kernels and converge to approximately 2–4% overhead GPUs using NVIDIA Nsight Compute (NCU) inside the for matrices above N = 4096. TF32 is a notable outlier— nvcr.io/nvidia/pytorch:25.11-py3 container. For cuBLAS selects XMMA and CUTLASS kernels instead of

(a) Tiled matrix multiplication.

(b) Tile quantization impact.

Fig. 2: Tile-quantization overhead in GEMM execution [27]. The cuBLAS heuristics (which among other heuristics leverages nvMatmulHeuristics [30]) select from several kernel families depending on matrix shape and precision. The primary families observed in our experiments are: • nvJet: NVIDIA’s proprietary high-performance GEMM kernels, which seem to efficiently expose a very large number of tile count, precisions and fused epilogues. Selected for most well-aligned shapes that can leverage TMA. • XMMA: CUDA C++ template-based kernels focused on Cooperative Thread Array (CTA)-level decomposition, used by cuBLAS and cuDNN.

8 7

Speedup over TF32

nvJet, producing systematically higher overhead (up to 33% at small sizes) that converges more slowly. This suggests that cuBLAS kernel selection, a software optimization, is another significant factor in the tile-quantization overhead. cuBLAS employs a comprehensive set of heuristics that optimize kernel dispatch based on matrix shape, kernel implementation, precision GPU architecture, clocks, and available pipelines [29]. These heuristics can spread computation across multiple execution pipelines in ways that do not correspond to the user’s chosen datatype—for example, TF32 operations may be dispatched to kernels that heavily utilize the FP16 (HMMA) pipeline. This observation further motivates a hardware-level metric like OFU: optimization opportunities arise not only from application-level choices and hardware capabilities, but also from intermediate library layers whose behavior is opaque to both the user and the training framework. Application-level MFU, which derives FLOPs from model architecture, cannot capture these library-level effects. OFU, by measuring what the GPU actually executes, reflects the true utilization regardless of how cuBLAS maps the workload to hardware pipelines. GPU GEMM kernels partition the output matrix C into rectangular tiles assigned to thread blocks [27] (Fig. 2a). When M , N , or K does not divide evenly into the tile dimensions (TM , TN , TK ), the last tiles are zero-padded in shared memory and computed in full (Fig. 2b), yielding effective dimensions:       M N K Meff = TM , Neff = TN , Keff = TK TM TN TK (3) The actual FLOPs executed are 2 Meff Neff Keff ≥ 2M N K. This is known as tile quantization [27].

BF16 FP8 NVFP4

6 5 4 3 2 109

1010

1011

Total operand bytes

1012

Fig. 3: Throughput speedup over TF32 on GB200. CUTLASS 2/3: Open-source GEMM templates; CUTLASS 2 lacks Cooperative Grid Array (CGA) support and is typically selected for small or poorly aligned matrices. Modern kernels (nvJet, XMMA, CUTLASS 3) use CGAs [20], which group (CM , CN ) thread blocks into clusters that share distributed shared memory across SMs. This introduces a two-level tiling hierarchy: at the first level, each thread block computes one TM × TN output tile, rounding the matrix dimensions up to tile boundaries as described above; at the second level, thread blocks are grouped into CM × CN clusters, and the number of tiles must itself be rounded up to a whole number of clusters. The effective dimension therefore undergoes two successive ceiling operations:   ⌈M/TM ⌉ Meff = · C M · TM (4) CM •

(and analogously for N ). When CM > 1, a matrix that fits exactly into an integer number of tiles can still incur padding at the cluster level, adding an extra CM − 1 tiles of waste in the worst case. For nvJet kernels, the tile dimensions and CGA configuration are encoded in the kernel name (e.g. nvjet_sm90_hsh_256x160_64x4_2x1), enabling a closed-form FLOP prediction that matched NCU measurements to within < 1000 FLOPs for all tested cases. For XMMA and CUTLASS kernels the K-dimension tiling and CGA configuration are not statically visible, so exact prediction requires runtime introspection. B. Floating-Point Precision Precision format affects realized throughput relative to the theoretical peak. To characterize this, we profiled square GEMMs on a GB200 GPU across increasing matrix sizes in BF16, FP8, and NVFP4, with N ranging up to 18,432 for TF32/BF16, 24,576 for FP8, and 32,768 for NVFP4, measuring the ratio of achieved throughput (FLOP/s) relative to TF32 (Fig. 3). At large matrix sizes the curves approach their theoretical speedups: 2× for BF16, 4× for FP8, and 8× for NVFP4, with BF16 converging most cleanly because 2× is the smallest multiplier. Lower precisions show progressively

TABLE I: Errors from clock frequency sampling rates. Int. (s)

N =4096 σ 95%

5 10 20 30

0.01 0.02 0.04 0.02

±0.01 ±0.04 ±0.07 ±0.03

N =8192 σ 95% 0.03 0.07 0.08 0.09

±0.05 ±0.13 ±0.15 ±0.18

N =16384 σ 95%

σ

Alt. 95%

±0.07 ±0.18 ±0.06 ±0.22

0.01 0.03 0.08 0.09

±0.03 ±0.06 ±0.15 ±0.17

0.03 0.09 0.03 0.11

more deviation, partly due to scaling-factor (SF) overhead in block-scaled formats: FP8 requires one 512-byte SF block per 128 × 128 input tile, while NVFP4 requires one per 128 × 64 input tile. Considering typical FP8 tiles are 128×256×128 and NVFP4 are 128 × 256 × 256 we quadruple the SF overhead, going from 3 to 12 SF blocks per tile. NVFP4’s speedup falls below 8× at small matrix sizes where SF bookkeeping dominates, recovering as matrix dimensions grow. OFU-derived speedup, computed as (OFUp × Peakp )/(OFUTF32 × PeakTF32 ), closely tracks the measured curves: at N ≥ 4096, OFU-derived speedups of 1.85×, 3.51×, and 6.75× for BF16, FP8, and NVFP4 agree with the measured 1.78×, 3.27×, and 6.10×, confirming that OFU correctly captures precision-dependent throughput scaling.

estimation error. Thus, a collection interval should be at most 30 seconds. D. Theoretical Peak FLOPs and Tensor Core Clock Domains Computing OFU requires a correct peak TFLOP/s denominator, which depends on the clock frequency of the Tensor Core pipeline. The theoretical peak throughput of a GPU is determined by three architectural parameters: the number of Streaming Multiprocessors (SMs), the number of floating-point operations each SM can perform per clock cycle on the relevant execution pipeline, and the maximum clock frequency of that pipeline: Peak TFLOP/s =

SMs × FLOPs/cycle/SM × f max 1012

(5)

A subtlety arises because Tensor Core pipelines do not necessarily run at the same maximum clock frequency as the rest of the SM [17]. On the H100 SXM, Tensor Core operations in lower-precision formats (FP8, FP16, BF16, TF32) boost to a maximum of 1,830 MHz, whereas the SM boost clock is 1,980 MHz. FP32 and FP64 operations, including FP32 and C. SM Clock Sampling Noise FP64 Tensor Core/Non Tensor Core instructions, run at the Clock sampling introduces noise into OFU estimates. OFU full 1,980 MHz SM clock. This distinction is critical for computing correct peak is computed from two hardware counters, tensor core activity and SM clock frequency, that are polled at discrete throughput values. Using the primary clock for Tensor Core intervals. Because the SM clock reported by the hardware precisions on H100: is an instantaneous sample rather than a hardware-averaged 132 × 4,096 × 1,830 × 106 value (unlike tensor pipe activity, which is averaged over PeakH100, FP16 = (6) 1012 the collection window), coarser scrape intervals introduce = 989.4 TFLOP/s sampling noise into the OFU estimate. For example, during a sustained 16384 × 16384 BF16 GEMM on an H100, the SM which agrees with the published specification of clock sampled at 1 kHz via Nsight Systems fluctuates between 989 TFLOP/s [31]. The remaining Tensor Core precisions ∼1,201 MHz and ∼1,558 MHz (mean 1,352 MHz, std 32 MHz), scale proportionally from this base rate: driven by power and thermal management. • FP8: 2 × 989 = 1,978 TFLOP/s. To quantify the impact on OFU, we collected both counters at • TF32: 989 / 2 = 494.5 TFLOP/s. 1-second intervals, the minimum supported by nvidia-smi For the GB200, no public documentation currently specifies dmon, over 3,000 seconds of sustained FP16 matrix multiplia separate Tensor Core clock frequency. Using the published cation on a GB200 GPU, then subsampled at coarser intervals SM boost clock of 2,062 MHz as the Tensor Core frequency: (5–30 s) and measured the deviation from the 1-second baseline. (DCGM supports collection intervals as low as 100 ms, which 148 × 8,192 × 2,062 × 106 would further reduce sampling error.) Three steady-state matrix PeakGB200, FP16 = (7) 1012 sizes (N = 4096, 8192, 16384) and an alternating workload = 2,499.9 TFLOP/s (16384 ↔ 4096, switching every 10 s) were tested. Table I reports the results. which matches the published specification of Even at 30-second intervals (∼100 samples over a 50- 2,500 TFLOP/s [32]. minute window), the 95% confidence interval remains below ±0.22 percentage points—negligible relative to the OFU values E. Non-Tensor Undercounting themselves (∼55%). At 5-second intervals the bound drops below ±0.07 percentage points. Sampling noise is therefore not OFU monitors only the Tensor Core pipeline, excluding a material source of error for production OFU measurements, CUDA-core work (activations, normalization, softmax). This provided the collection window spans at least several minutes. omission is negligible: matrix multiplications account for A practical constraint is that the DCGM hardware counter for 99.8% of total FLOPs in a transformer encoder layer [26]. tensor pipe activity averages over at most 30-second windows. Standard MFU definitions—PaLM [11], Megatron-LM [33], Collecting at intervals longer than 30 seconds yields an average and the OpenAI scaling laws [34]—follow the same convention, of averages rather than a true window average, compounding deriving FLOPs exclusively from matrix-multiplication terms.

TABLE II: Prediction accuracy.

V. P RACTICAL ACCURACY In this section, we evaluate OFU’s practical accuracy. We first apply tile-quantization corrections to controlled GEMM workloads, establishing bounded accuracy on fully specified workloads. We then compare OFU against application-reported MFU on 608 production training jobs, measuring correlation and surfacing cases where the two metrics diverge. A. Predicting MFU from Hardware Counters

GPU

Prec

Estimator

MAE

≤2 pp

≤5 pp

H100

FP16

OFU Adj OFU

1.90 0.06

64% 100%

96% 100%

H100

TF32

OFU Adj OFU

3.46 0.50

44% 99%

86% 99%

H100

FP8

OFU Adj OFU

1.58 0.07

73% 100%

99% 100%

OFU

1.08

86%

99%

GB200 FP16 Using the tile-quantization corrections from Section IV-A Adj OFU 1.04 100% 100% and the Tensor Core clock frequencies from Section IV-D, we OFU 2.10 65% 88% evaluate how accurately OFU tracks application-level MFU on GB200 TF32 Adj OFU 1.03 100% 100% sustained matrix multiplications. OFU 0.64 96% 100% We compare three quantities for each matmul: GB200 FP8 Adj OFU 0.70 100% 100% • OFU (unadjusted), as defined in (1), computed from OFU 1.21 87% 98% Nsight Systems GPU metrics sampled at 10 kHz. GB200 NVFP4 Adj OFU 1.15 95% 100% • Adjusted OFU, which corrects OFU for tile-quantization overhead (Section IV-A): Adjusted OFU achieves ≤2 percentage-point error for 95–100% FLOPstheoretical 2M N K of matmuls and ≤5 percentage points for 100% across all OFUadj = OFU × = OFU × FLOPsprofiled FLOPsNCU four precisions (FP16, TF32, FP8, NVFP4), though a small (8) systematic underestimate of ∼1 percentage point remains, likely • App MFU (ground truth), computed as measured TFLOP/s from sampling overhead in the 10 kHz hardware counters. divided by the architecturally derived peak TFLOP/s for For pure GEMM workloads, OFU predicts application-level the given precision (Section IV-D). MFU to within 1–3 percentage points without any modelFor each GPU and precision we profiled 500 random specific information. The NCU-adjusted estimator further tight(M, K, N ) matrix multiplications, where each dimension was ens this to ≤2 percentage points for all tested configurations. a random multiple of 16. Each matmul ran for 5 minutes under Nsight Systems profiling, collecting both application throughput B. Production Workloads Validation and hardware counters (tensor pipe activity and SM clock We evaluated OFU as a proxy for MFU across 608 producfrequency) at 10 kHz. A single-iteration NCU pass was then run tion training jobs on H100 GPUs at a commercial GPU cluster to measure the actual FLOP count for the tile-quantization cor(August 27–September 10, 2025), run by an internal research rection. For FP16 and TF32 we used PyTorch torch.matmul group using Megatron-LM. Jobs ranged from 8 to 5,888 GPUs and for FP8 we used torch._scaled_mm, inside the across 80 distinct configurations from 26 users. nvcr.io/nvidia/pytorch:25.11-py3 container; for MFU was sourced from OneLogger, computed from NVFP4 on GB200 NVL we used an internal matrix multiMegatron-LM’s reported total FLOPs and training loop wallplication benchmark; NVFP4 dimensions were restricted to clock time: multiples of 128. We profiled 500 random (M, K, N ) GEMMs train_tflop × gpu_count per configuration: FP16, TF32, and FP8 on H100 SXM, and MFU = × 100% (10) FP16, TF32, FP8, and NVFP4 on GB200 NVL, using the 989 architecturally derived peak TFLOP/s from Section IV-D. where 989 TFLOP/s is the H100 BF16 Tensor Core peak [31] Fig. 4a and Fig. 4b show the distribution of prediction error (see Section IV-D). (estimate – App MFU, in percentage points) across all 500 DCGM telemetry was scraped via Prometheus at 30-second matmuls per configuration. Table II provides summary statistics, intervals, aligned to each job’s training window. OFU was where the mean absolute error (MAE) is defined as: computed as: n   X SM_CLOCK 1 × 100 MAE = |OFUi − App MFUi | (9) OFU = mean PIPE_TENSOR_ACTIVE × 1830 n i=1 (11) Raw OFU consistently overestimates App MFU by 1– averaged across all GPUs and time samples, where 1830 MHz 2 percentage points across both GPUs and all precisions, is the H100 Tensor Core maximum clock frequency (Secas expected from the tile-quantization overhead identified in tion IV-D). Section IV-A. After NCU correction, Adjusted OFU centres Across all 608 jobs, MFU and OFU show a moderate positive near zero with substantially reduced variance. correlation (Pearson r = 0.53). Mean MFU was 25.1%±10.9% On H100, Adjusted OFU achieves ≤2 percentage-point error versus mean OFU of 25.0% ± 8.3%. Mean absolute error was for 99–100% of matmuls across all three precisions, with 6.2%. Of all jobs, 79.4% fell within 10% absolute error, while mean absolute error under 0.6 percentage points. On GB200, 6.7% exceeded 20% error. Fig. 5 shows the per-job relationship;

OFU Adj. OFU

7

OFU Adj. OFU

6

5

Prediction Error (pp)

Prediction Error (pp)

6

4 3 2

4 2 0

1 2

0

FP16

TF32

FP8

FP16

TF32

(a) H100

FP8

NVFP4

(b) GB200

Fig. 4: OFU prediction error.

70

TABLE III: Absolute error between MFU and OFU.

Other (526) Case study jobs (82) y=x

60

OFU (%)

50 40 30 20 10 0

r = 0.53 (all) r = 0.78 (excl. case studies)

0

10

20

30

40

App MFU (%)

50

60

GPUs

Jobs

MFU (%)

Abs Err (%)

8 16 64 128 256 288 512 736 768 1024 1536 2944 5888

6 48 52 48 76 65 144 11 57 49 10 33 9

28.7 ± 6.9 23.8 ± 3.3 23.6 ± 2.5 24.3 ± 8.7 20.1 ± 12.6 40.1 ± 16.3 23.9 ± 5.6 24.2 ± 0.4 16.9 ± 4.1 35.0 ± 9.1 12.4 ± 2.3 24.0 ± 3.7 13.6 ± 0.1

7.5 ± 3.9 12.2 ± 2.0 2.2 ± 2.4 4.5 ± 2.9 9.1 ± 4.9 18.0 ± 14.4 3.6 ± 2.2 3.6 ± 0.1 1.2 ± 0.7 4.1 ± 0.7 0.3 ± 0.2 2.6 ± 0.3 1.5 ± 0.2

70

inflating the reported FLOPs by a factor of ∼3×. This produced Fig. 5: App MFU vs. OFU for 608 production training jobs. an application-reported MFU of 54.27% against OFU of 25.58% (relative error: 112.2%). Correcting the FLOPs count to account most jobs cluster near the y = x line, with the 288-GPU MoE for the down- and up-projections reduced the reported MFU to 18.45%, cutting the relative error to 27.9%. group as a clear outlier. A second miscalculation affected hybrid MoE jobs trainTable III summarizes results by GPU count. Agreement improves substantially at large scale: jobs with ≥768 GPUs ing an 8B-parameter model on 300B tokens that interleave consistently achieve sub-5% absolute error. The most significant attention, Mamba [35], dense MLP, and sparse MoE layers. outlier is the 288-GPU group, which exhibits a mean absolute The Megatron-LM branch used for these experiments did not error of 18.0%, driven by a specific architectural issue described support hybrid architectures in its FLOPs counter. As a result, every layer’s FLOPs were counted as if it were a self-attention below. and dense MLP layer cost. This inflated the reported FLOPs, C. Detecting Production FLOPs Miscalculations producing MFU of 24.51% against OFU of 15.56% (relative To illustrate the kinds of production problems our techniques error: 57.5%). After the FLOPs function was updated with perhelp detect and remediate, we examined jobs where MFU and layer-type accounting, subsequent runs of a similar architecture OFU diverged most significantly. This analysis surfaced two (1,536 GPUs) reported MFU of 17.8–18.0% versus OFU of distinct FLOPs miscalculations in Megatron-LM. Excluding 18.5–18.7% (relative error 3–4%). the affected 82 jobs improves overall correlation from r = 0.53 Together, these case studies indicate that significant diverto r = 0.78, and reduces the fraction of jobs exceeding 10% gence between OFU and application-reported MFU consistently absolute error from 21.8% to 16.7%. traced back to incorrect FLOPs calculations in the training A representative 288-GPU job training a 16B-parameter framework rather than OFU measurement error. This is DeepSeek-style MoE where activations were down-projected expected: application-level MFU depends on manually derived from hidden dimension 2048 to latent dimension 512 before FLOPs counts that are brittle for novel architectures, whereas expert routing. Megatron-LM’s FLOPs counter incorrectly OFU is computed directly from hardware counters with no assumed experts operated at the full hidden dimension (2048), model-specific assumptions.

OFU has been integrated at multiple levels of a large-scale GPU fleet infrastructure—from per-job dashboards visible to individual researchers, to cluster-wide resilience and goodput services that flag inefficiencies and drive optimization. In each deployment, OFU discovered problems that were hidden due to the lack of instant visibility into floating-point behavior with well-defined accuracy properties. This section describes how OFU is operationalized in practice and presents the problems it uncovered.

% of Pre-Fix Mean OFU

VI. O PERATIONAL E XPERIENCES

200 100 0 0.0

2.5

5.0

7.5

10.0 Time (hours)

12.5

15.0

17.5

Fig. 6: OFU before and after removing debug overhead, normalized to the pre-fix mean. 250

% of Mean MFU

A. Embodied Agent Training To evaluate whether OFU is useful in practice, we operationalized the metric for an internal research lab that develops foundation models for robotic embodied agents. We integrated OFU into Osmo [36], the lab’s Kubernetes-native orchestration platform for Physical AI workloads, so that researchers could monitor GPU utilization at the job level without deriving modelspecific FLOPs counts. For each training job managed by Osmo, we compute OFU from DCGM metrics and display it as a time-series dashboard, both per GPU individually and as a job-level aggregate. Separately, the lab’s training infrastructure optionally runs a few iterations with PyTorch Profiler at job start and uploads the trace to S3 for later analysis. Because OFU is model-architecture independent, it works for every training experiment without requiring researchers to manually derive FLOPs per token. This is especially valuable for labs working with novel architectures where applicationlevel MFU is either not calculated or calculated incorrectly, as demonstrated in the internal case studies (Section V-C). Even for smaller experiments where teams do not typically invest in performance tuning, OFU provides zero-effort visibility that can surface low-hanging-fruit misconfigurations. A robot foundation model experimental training run on 256 H100 GPUs (32 nodes) observed lower OFU (Fig. 6) than expected. Because OFU flagged the issue immediately, the team collected a PyTorch Profiler trace to investigate. The trace revealed that the environment variable TORCH_DISTRIBUTED_DEBUG=DETAIL had been set and merged to the main repository, which causes PyTorch to inject gloo:all_gather validation calls during every NCCL collective operation. The before-fix trace contained 15,360 gloo:send events, 7,560 gloo:recv events, and 240 gloo:all_gather events, all running over CPU-based Gloo transport. These validation collectives serialized with the NCCL all-reduce operations, dominating wall-clock time and leaving the Tensor Cores idle for the vast majority of each training step. After removing the debug flag, the Gloo operations disappeared entirely from the trace and OFU improved by 2.5× (Fig. 6). This case study illustrates two points. First, it is often impractical for researchers to derive FLOPs per token for every experiment, making issues like this difficult to surface without a hardware-level metric like OFU. Second, OFU is

Before fix After fix

300

App MFU OFU

200 150 100 50 0

0

100

200

300

400

Wall-clock time (hours)

500

600

700

Fig. 7: OFU and MFU relative to mean MFU on a pretraining at 6,144 GB200 GPUs. best understood as a coarse utilization signal. It identifies that a problem exists and quantifies its severity, but diagnosing why utilization is low requires profiling the job to identify the specific bottleneck. B. Large-Scale Mixed-Precision Pretraining Extreme-scale training provided another opportunity to evaluate both the usefulness and practical performance of OFU on real-world workloads. We incorporated OFU into a NVIDIA Mission Control Autonomous Job Recovery Service that monitors large-scale training jobs and takes corrective action to improve productivity and reliability. Using this service, we monitored training runs for an internal research lab focused on scaling large language models and GPU-intensive deep learning workloads, pretraining a mixed-precision Mixtureof-Experts hybrid Mamba-Transformer on 6,144-GPU Slurm jobs. We evaluated 711 hours (∼30 days) of wall-clock training time on 6,144 GPUs (Fig. 7). These represent all jobs from the training run that used 6,144 GPUs and had both OFU data and application-level throughput data emitted at least every 90 seconds. Because this workload mixes multiple precisions (BF16, FP8, NVFP4), which have different hardware peak throughputs, the single-precision denominator in (10) does not apply. Instead, we define an effective peak as the FLOPsweighted harmonic mean of per-precision peaks: P Fi Peff = P i Fi , (12) i Pi

where Fi is the FLOPs executed at precision i and Pi is the corresponding hardware peak from [32]. Application MFU is then computed as in (10) with Peff replacing the singleprecision peak. Across the 711-hour span, the point-by-point correlation between OFU and application MFU was r = 0.718, with shortterm noise from clock sampling and transient workload phases

reducing the agreement at individual time steps. When we average both metrics per job, the noise cancels out and the correlation across the 174 jobs rises to r = 0.977, indicating that OFU reliably distinguishes efficient jobs from inefficient ones. OFU is precision-agnostic by construction. The tensor pipe activity counter measures cycles executing tensor instructions regardless of numeric format. This run provided a natural test of this property, and provides empirical real-world backing for the results in IV-A. Multiple debugging periods required switching from mixed precision (NVFP4, FP8, and BF16) to BF16-only. Observed compute throughput (TFLOP/s/GPU) remained roughly constant between these modes. However, because BF16 has a lower theoretical peak than FP8 or NVFP4, the effective peak for BF16-only periods was lower. With roughly constant throughput and a lower denominator, application-reported MFU increased accordingly. Figure 7 shows that OFU exhibited a corresponding increase (both metrics are shown as a percentage of mean MFU in the figure). In both mixed-precision and BF16-only modes, the two metrics agreed within 1 absolute percentage points on average, confirming that OFU correctly reflects precisiondependent utilization changes despite having no knowledge of the numeric format in use. C. World Foundation Model Training

variants of this approach. Google has published TPU v4 MFU benchmarks using a similar methodology [38]. MegaScale [39] reports MFU at 10,000+ GPU scale, and MLPerf [40] provides standardized training throughput benchmarks across hardware platforms. While these estimates require no hardware-level access, they depend on manually derived FLOPs formulas that must be updated for each new architecture—mixture-of-experts, latent-space routing, multimodal pipelines—and can silently become incorrect as models evolve, as demonstrated in our production case studies. Hardware performance counters exposed through NVIDIA DCGM [12] provide non-intrusive utilization signals at negligible overhead. Prior work has used these counters primarily for coarse utilization metrics such as SM Activity [41] rather than as a first-principles MFU estimator. To our knowledge, OFU is the first systematic study that derives, characterizes, and validates a hardware-counter-based MFU proxy across multiple GPU generations and precisions, with bounded accuracy established through controlled experiments. The roofline model [42] bounds throughput by arithmetic intensity but requires per-kernel analysis and does not yield a continuous fleet-wide utilization metric. VIII. C ONCLUSION

Fleet-wide GPU utilization measurement demands a metric An 8B-parameter world foundation model for generating that is non-intrusive, precision-agnostic, and accurate across physics-aware images and videos was trained on 256 GB200 GPU generations. We derived Overall FLOP Utilization GPUs. The run reported application-level MFU of 26%, while (OFU) from first principles of GPU floating-point execution OFU measured 34%—a larger discrepancy than the 1–2 per- pipelines, grounding it in Tensor Pipe Activity and SM centage points expected from tile-quantization overhead alone. clock frequency—two signals universally available through Investigation revealed that the framework’s FLOPs formula hardware performance counters—and characterized five sources did not account for the additional forward-pass recomputation of divergence from application-level MFU: tile quantization and introduced by activation checkpointing [37]. With full activation cuBLAS kernel selection, floating-point precision portability, checkpointing enabled, each training step performs roughly 4F SM clock sampling noise, Tensor Core clock domains, and nonFLOPs (F forward + F recomputed forward + 2F backward) tensor undercounting. Evaluation on 3,500 controlled GEMM rather than the standard 3F , a 33% increase. After correcting experiments across FP16, TF32, FP8, and NVFP4 on both H100 the FLOPs formula, application-reported MFU rose from 26% and GB200 demonstrated ≤2 percentage-point accuracy after to 33%, aligning with the 34% OFU to within 1 percentage tile-quantization correction, and against 608 production training point. jobs OFU achieved r = 0.78 correlation with applicationlevel MFU while surfacing two distinct framework-level VII. R ELATED W ORK FLOPs miscalculations. Deployed across large-scale GPU GPU utilization measurement spans three broad categories, fleets, OFU has detected a 2.5× efficiency regression in each with distinct trade-offs between fidelity, coverage, and embodied-agent training, and when integrated into fleet-wide deployment cost. resilience and goodput services, has surfaced performance drops Performance profiling tools such as NVIDIA Nsight Com- invisible to application-reported throughput. Our operational pute, Nsight Systems, PyTorch Profiler [5], DeepSpeed Flops experience yields three lessons: (1) OFU is less error-prone Profiler [6], and HPCToolkit [7] provide detailed per-kernel than application-level MFU and can pinpoint workloads with metrics including instruction counts, memory bandwidth, and significant optimization opportunities; (2) OFU works well at occupancy. These tools offer the highest measurement fidelity very large scale, faithfully tracking mixed-precision and lowerbut require per-job instrumentation, impose runtime overhead, precision training—a typical optimization evolution—without and are designed for targeted profiling rather than continuous any code instrumentation; (3) instant, fleet-wide visibility fleet-wide monitoring. enables rapid diagnosis of performance bugs, leading to Framework-level MFU estimation, popularized by improvements as large as 2.5× in our case studies. Overall, our PaLM [11], derives throughput from model-architecture evaluation and operational experience indicate that OFU can FLOPs counts divided by wall-clock time. Megatron-LM [8], significantly lower the barrier to unearthing GPU inefficiency NeMo [10], and PyTorch Lightning [9] each implement for AI workloads on HPC systems.

ACKNOWLEDGMENTS AI-Generated Content Disclosure: AI writing tools were used for editorial assistance in preparing this manuscript. R EFERENCES [1] Microsoft, “Microsoft fy2025 q4 earnings,” Microsoft Investor Relations, 2025. [Online]. Available: https://www.microsoft.com/en-us/investor/eve nts/fy-2025/earnings-fy-2025-q4 [2] Alphabet Inc., “Alphabet announces fourth quarter and fiscal year 2025 results,” February 2026. [Online]. Available: https: //www.sec.gov/Archives/edgar/data/1652044/000165204426000012/goo gexhibit991q42025.htm [3] Meta Platforms, “Meta reports fourth quarter and full year 2025 results,” January 2026. [Online]. Available: https://investor.atmeta.com/investor-n ews/press-release-details/2026/Meta-Reports-Fourth-Quarter-and-Ful l-Year-2025-Results/default.aspx [4] Reuters, “Citigroup forecasts big tech’s ai spending to cross $2.8 trillion by 2029,” September 2025. [Online]. Available: https://www.reuters.com/world/china/citigroup-forecasts-big-techs-ai-s pending-cross-28-trillion-by-2029-2025-09-30/ [5] PyTorch Contributors, “Pytorch profiler,” PyTorch Documentation, 2024. [Online]. Available: https://pytorch.org/tutorials/recipes/recipes/profiler r ecipe.html [6] Microsoft DeepSpeed Team, “Deepspeed flops profiler,” DeepSpeed Documentation, 2024. [Online]. Available: https://www.deepspeed.ai/tut orials/flops-profiler/ [7] L. Adhianto, S. Banerjee, M. Fagan, M. Krentel, G. Marin, J. MellorCrummey, and N. R. Tallent, “Hpctoolkit: Tools for performance analysis of optimized parallel programs,” in Concurrency and Computation: Practice and Experience, vol. 22, no. 6, 2010, pp. 685–701. [Online]. Available: https://doi.org/10.1002/cpe.1553 [8] NVIDIA, “Nvidia/megatron-lm: Ongoing research training transformer models at scale,” GitHub README, 2025. [Online]. Available: https://github.com/NVIDIA/Megatron-LM [9] Lightning AI, “Pytorch lightning,” GitHub repository, 2024. [Online]. Available: https://github.com/Lightning-AI/pytorch-lightning [10] NVIDIA, “Nvidia nemo: A toolkit for building ai models,” GitHub repository, 2024. [Online]. Available: https://github.com/NVIDIA/NeMo [11] A. Chowdhery et al., “Palm: Scaling language modeling with pathways,” arXiv:2204.02311, 2022. [Online]. Available: https: //doi.org/10.48550/arXiv.2204.02311 [12] NVIDIA, “Nvidia dcgm exporter,” GitHub repository, 2025. [Online]. Available: https://github.com/NVIDIA/dcgm-exporter [13] ——, “nv-one-logger: Nvidia’s distributed metrics logging system,” GitHub repository, 2024. [Online]. Available: https://github.com/NVIDI A/nv-one-logger [14] DeepSeek-AI, “Deepseek-v2: A strong, economical, and efficient mixture-of-experts language model,” arXiv preprint arXiv:2405.04434, 2024. [Online]. Available: https://doi.org/10.48550/arXiv.2405.04434 [15] NVIDIA, “Nvidia tesla v100 gpu architecture,” NVIDIA, Tech. Rep. WP-08608-001 v1.1, 2017. [Online]. Available: https://images.nvidia.co m/content/volta-architecture/pdf/volta-architecture-whitepaper.pdf [16] ——, “Nvidia a100 tensor core gpu architecture,” NVIDIA, Tech. Rep., 2020. [Online]. Available: https://images.nvidia.com/aem-dam/en-zz/So lutions/data-center/nvidia-ampere-architecture-whitepaper.pdf [17] ——, “Nvidia h100 tensor core gpu architecture,” NVIDIA, Tech. Rep., 2022, whitepaper. [Online]. Available: https://www.advancedclustering.c om/wp-content/uploads/2022/03/gtc22-whitepaper-hopper.pdf [18] SemiAnalysis, “Nvidia tensor core evolution: From volta to blackwell,” SemiAnalysis Newsletter, June 2025. [Online]. Available: https: //newsletter.semianalysis.com/p/nvidia-tensor-core-evolution-from-vol ta-to-blackwell [19] D. Stosic and P. Micikevicius, “Accelerating ai training with nvidia tf32 tensor cores,” NVIDIA Technical Blog, January 2021. [Online]. Available: https://developer.nvidia.com/blog/accelerating-ai-training-wit h-tf32-tensor-cores/ [20] NVIDIA, “Nvidia h100 tensor core gpu architecture,” NVIDIA, Tech. Rep., 2022, whitepaper. [Online]. Available: https://www.advancedcluste ring.com/wp-content/uploads/2022/03/gtc22-whitepaper-hopper.pdf

[21] K. Sevegnani, G. Fiameni, U. Uppal, S. Perez, and A. Pilzer, “Floating-point 8: An introduction to efficient, lower-precision ai training,” NVIDIA Technical Blog, June 2025. [Online]. Available: https://developer.nvidia.com/blog/floating-point-8-an-introduction-to-e fficient-lower-precision-ai-training/ [22] Colfax Research, “Cutlass tutorial: Fast matrix-multiplication with wgmma on nvidia hopper gpus,” Colfax Research Blog, August 2024. [Online]. Available: https://research.colfax-intl.com/cutlass-tutorial-wgm ma-hopper/ [23] E. Alvarez, O. Almog, E. Chung, S. Layton, D. Stosic, R. Krashinsky, and K. Aubrey, “Introducing nvfp4 for efficient and accurate lowprecision inference,” NVIDIA Technical Blog, June 2025. [Online]. Available: https://developer.nvidia.com/blog/introducing-nvfp4-for-effic ient-and-accurate-low-precision-inference/ [24] NVIDIA, “Nvidia rtx blackwell gpu architecture,” NVIDIA, Tech. Rep., 2025, architecture Whitepaper. [Online]. Available: https: //images.nvidia.com/aem-dam/Solutions/geforce/blackwell/nvidia-rtx-b lackwell-gpu-architecture.pdf [25] P. Micikevicius, D. Stosic, N. Burgess, M. Cornea, P. Dubey, R. Grisenthwaite, S. Ha, A. Heinecke, P. Judd, J. Kamalu, N. Mellempudi, S. Oberman, M. Shoeybi, M. Renber, K. Siu, and H. Wu, “Pretraining large language models with NVFP4,” arXiv preprint arXiv:2509.25149, 2025. [Online]. Available: https://doi.org/10.48550/arXiv.2509.25149 [26] A. Ivanov, N. Dryden, T. Ben-Nun, S. Li, and T. Hoefler, “Data movement is all you need: A case study on optimizing transformers,” in Proceedings of Machine Learning and Systems (MLSys), vol. 3, 2021, pp. 711–732. [Online]. Available: https://doi.org/10.48550/arXiv.2007.00072 [27] NVIDIA, “Matrix multiplication background user’s guide,” NVIDIA Deep Learning Performance Documentation, 2024. [Online]. Available: https://docs.nvidia.com/deeplearning/performance/dl-performance-matri x-multiplication/index.html [28] ——, “cublas library user’s guide,” CUDA Toolkit Documentation, 2025. [Online]. Available: https://docs.nvidia.com/cuda/cublas/index.html [29] ——, “nvmatmulheuristics,” 2025. [Online]. Available: https://developer. nvidia.com/blog/improving-gemm-kernel-auto-tuning-efficiency-on-nvi dia-gpus-with-heuristics-and-cutlass-4-2/ [30] H. Zhao, D. Yan, A. Wang, A. Kerr, and M. Yan, “Improving gemm kernel auto-tuning efficiency on nvidia gpus with heuristics and cutlass 4.2,” NVIDIA Technical Blog, January 2025. [Online]. Available: https://developer.nvidia.com/blog/improving-gemm-kernel-auto-tunin g-efficiency-on-nvidia-gpus-with-heuristics-and-cutlass-4-2/ [31] NVIDIA, “Nvidia h100 tensor core gpu,” NVIDIA Data Center GPU Product Page, 2024. [Online]. Available: https://www.nvidia.com/en-us/ data-center/h100/ [32] ——, “Nvidia gb200 tensor core gpu,” NVIDIA Data Center GPU Product Page, 2024. [Online]. Available: https://www.nvidia.com/en-us/ data-center/gb200-nvl72/ [33] V. Korthikanti, J. Casper, S. Lym, L. McAfee, M. Andersch, M. Shoeybi, and B. Catanzaro, “Reducing activation recomputation in large transformer models,” arXiv preprint arXiv:2205.05198, 2022. [Online]. Available: https://doi.org/10.48550/arXiv.2205.05198 [34] J. Kaplan, S. McCandlish, T. Henighan, T. B. Brown, B. Chess, R. Child, S. Gray, A. Radford, J. Wu, and D. Amodei, “Scaling laws for neural language models,” arXiv preprint arXiv:2001.08361, 2020. [Online]. Available: https://doi.org/10.48550/arXiv.2001.08361 [35] A. Gu and T. Dao, “Mamba: Linear-time sequence modeling with selective state spaces,” arXiv preprint arXiv:2312.00752, 2024. [Online]. Available: https://doi.org/10.48550/arXiv.2312.00752 [36] NVIDIA, “Nvidia osmo: Cloud-native orchestration platform for physical ai,” Product page, 2025. [Online]. Available: https: //us-west-2-aws.osmo.nvidia.com/ [37] T. Chen, B. Xu, C. Zhang, and C. Guestrin, “Training deep nets with sublinear memory cost,” arXiv preprint arXiv:1604.06174, 2016. [Online]. Available: https://doi.org/10.48550/arXiv.1604.06174 [38] J. Bradbury, Q. Zhang, and A. Selvan, “Benchmarking flops utilization on tpu v4,” Google Cloud (whitepaper), 2023. [Online]. Available: https://services.google.com/fh/files/blogs/tpu v4 benchmarking.pdf [39] Z. Jiang, H. Lin et al., “Megascale: Scaling large language model training to more than 10,000 gpus,” in NSDI 2024, 2024. [Online]. Available: https://doi.org/10.48550/arXiv.2402.15627 [40] P. Mattson, C. Cheng, G. Diamos, C. Coleman, P. Micikevicius, D. Patterson, H. Tang, G.-Y. Wei, P. Bailis, V. Bittorf, D. Brooks, D. Chen, D. Dutta, U. Gupta, K. Hazelwood, A. Hock, X. Huang, A. Ike, B. Jia, D. Kang, D. Kanter, N. Kumar, J. Liao, G. Ma, D. Narayanan,

T. Oguntebi, G. Pekhimenko, L. Pentecost, V. J. Reddi, T. Robie, T. St. John, T. Tabber, C.-J. Wu, L. Xu, M. Yamazaki, C. Young, and M. Zaharia, “MLPerf training benchmark,” in Proceedings of Machine Learning and Systems (MLSys), 2020. [Online]. Available: https://doi.org/10.48550/arXiv.1910.01500 [41] C.-J. Wu, D. Brooks, K. Chen, D. Chen, S. Choudhury, M. Dukhan, K. Hazelwood, E. Isaac, Y. Jia, B. Jia, T. Leesatapornwongsa, H. Li, Y. Liang et al., “Datacenter-scale analysis and optimization of GPU machine learning workloads,” IEEE Transactions on Parallel and Distributed Systems, vol. 32, no. 11, pp. 2766–2780, 2021. [Online]. Available: https://doi.org/10.1109/TPDS.2021.3084540 [42] S. Williams, A. Waterman, and D. Patterson, “Roofline: An insightful visual performance model for multicore architectures,” Communications of the ACM, vol. 52, no. 4, pp. 65–76, 2009. [Online]. Available: https://doi.org/10.1145/1498765.1498785