Closing the Prediction Gap: Completing Machine Shape So That Predicted Time, Power, Energy, and Mapping Match What Real Hardware Does

arXiv:2610.03197v1 [cs.DC] 2 Oct 2026

Closing the Prediction Gap: Completing Machine Shape So That Predicted Time, Power, Energy, and Mapping Match What Real Hardware Does Lenore M. Mullin

Gaétan Hains

Professor Emerita, College of Nanotechnology, Science and Engineering University at Albany, SUNY, Albany, NY lmullin@albany.edu

LACL, Université Paris-Est Créteil Créteil, France gaetan.hains@u-pec.fr

Abstract—A companion paper derives a weighted, communication-free partition for heterogeneous, multiinstitution device ensembles from each device’s measured shape, ρmachine (d), and validates its predictions on real hardware. Its experiments, however, reveal that ρmachine is incomplete: fp32-to-fp16 speedup and power shifts are not fully indexed, and memory-capacity predictions depend on an assumed rather than measured reservation overhead. This paper extends the theory by indexing five previously unmeasured fields in the same ψ-selection vocabulary: memory capacity Micap , power draw Pi , execution-unit capability Xd , compiler-directive configuration Cd , and memory-reservation overhead Rd . We introduce reached(d), a Boolean indicating whether compiled code actually dispatches to specialized hardware. This distinction explains behavior such as NVIDIA’s fp16 power shift, which cannot be predicted from hardware availability alone. We restate the companion paper’s time and energy models in terms of reached(d), rather than merely the presence of specialized execution units. Because MoA fixes a computation independently of its target, we define a two-part admissibility criterion for compiler directives: a directive must select a real hardware quantity without discretionary reordering, and it must be deterministic rather than compiler-overridable. Across Intel, AMD, NVIDIA, OpenMP, OpenACC, and Open MPI, this leaves a small admissible class: cache-placement controls, deterministic vector or tile widths, and occupancy or placement caps. Heuristic mechanisms, including auto-vectorization and autotuning enable flags, are excluded and (k) represented by a residual ϵd . We validate the extended model on five production devices: NVIDIA A100 and H100, AMD MI100, NVIDIA V100, and Intel Max 1550. On each device, we directly measure or independently establish all five fields. Results show that specialized-hardware reachability can be real but partial, occupancy limits must be established per device and admissible realization, and memoryreservation overhead ranges from approximately 0.004 to 0.030. We also measure a negative residual of −9.8% for Triton’s excluded autotuner, demonstrating that an unvalidated heuristic can have an effect opposite to a presumed penalty. Together, these experiments close all five fields for every device in the companion paper’s ensemble and convert previously assumed or unindexed behavior into measured, target-specific quantities.

I. I NTRODUCTION A prediction that is only checked where it happens to already be right is not a validated theory; it is a coincidence reported as one. The companion paper to this one [15] does

not make that mistake: every one of its predicted numbers – a finish time, a memory ceiling, an energy figure – is checked against real, measured hardware, and two of those checks come back wrong. Both failures are diagnosed precisely, in closed form, rather than patched around, and both diagnoses point to the same conclusion stated only in that paper’s own future-work section: ρmachine (d), the tuple every prediction in that paper’s theory is built from, is missing fields that its own real experiments already show are load-bearing. This paper’s job is narrow and stated precisely: close those fields, giving the missing quantities the same formal status every already-indexed field of ρmachine has – a name, a ψselectable index position, and a stated role in a closed-form prediction – reasoning throughout from the companion paper’s own real numbers. A theory that names a gap in prose but leaves the corresponding equation silently incomplete has not closed anything; this paper’s contribution is closing it in the equations themselves, and, for one device, checking that closure against a first real measurement (Section VIII) rather than leaving the extended theory entirely unvalidated. Why any of this matters, stated once at the outset for a reader who has not yet been asked to care about a specific index position in a specific tuple: large language models can already translate source code from any language to any other, fluently, on request. What they cannot do is guarantee that the result is correct, memory-optimal, or predictable on the hardware it actually runs on – each translation is its own unverified event, earning nothing from the last one (Figure 1). A formal intermediate representation is what turns fluent translation into trustworthy translation: derive a computation’s meaning once, independent of target (MoA’s DNF), and mechanically produce a verified, memory-optimal realization per target (its ONF) – one proof reused, not re-earned per language, per vendor, per device. Prior work in this research program argues this case directly and tests it against production kernels on real hardware [14]. ρmachine (d) is the other half of that same argument: a formal derivation is only as trustworthy as the machine-shape facts it is derived against, and this paper’s job, narrow as it is, is making sure none of those facts are still missing.

Fig. 1. Why a formal intermediate representation, not just a capable translator: naive AI translation (top) produces a fresh, unverified guess per target language, with no relationship between one translation and the next. MoA-mediated translation (bottom) derives one formal specification once (the DNF) and mechanically produces each target’s realization (its ONF) from it, so correctness and the memory-optimal schedule both carry over rather than being re-earned per target – the argument developed and tested on real hardware in [14].

and instead of leaving a kernel’s performance to be discovered empirically after the fact. Every kernel measured in this paper and its companion is ONF code produced this way, not handtuned by trial and error. The companion paper’s own contribution, which this paper takes as given, is to apply this same DNF/ONF/γ machinery once more, one level up: instead of partitioning an array across the lanes of a single device, it partitions an array across the devices of a heterogeneous, multi-institution ensemble, with each device’s share sized by that device’s own measured shape, ρmachine (d) – exactly as Figure 4 shows for a toy instance and Figure 5 shows for a real, measured one. ρmachine (d) is, formally, nothing more exotic than another MoA array: a tuple, indexed by ψ, whose entries happen to be hardware facts (compute-unit count, memory ceiling, measured time) rather than tensor values. Table I below restates its five fields as the companion paper already indexes them; this paper’s entire contribution is a dimension lift on that same array – adding five more index positions to it, using that same ψselection notation, rather than introducing any new formalism – so that quantities the companion paper’s own real hardware

A. MoA in Brief: From Array Data to Machine Shape This subsection is for a reader meeting the Mathematics of Arrays (MoA) for the first time; a reader already fluent in it can skip directly to Section IV. MoA [13] is a formal array calculus: every array is a pair (shape, values), every array operation is a composition of a small, fixed set of algebraic operators – ψ (index/select), ι (generate an index set), ρ (shape), γ (index-substitution, i.e. partition or reshape), and rav (ravel to one dimension) – and a whole computation is first written once as a Denotational Normal Form (DNF): a single, target-independent composition of these operators that states what is computed, with no loop order, tiling, or device assumed. A mechanical rewrite step, justified by the algebra’s own identities rather than by handtuning, turns that DNF into an Operational Normal Form (ONF) – a loop nest, or a target-offload kernel – for a specific device. The goal this whole research program serves is stated plainly: derive correct, portable, and predictable low-level code from one formal specification, instead of hand-writing (and separately debugging) one implementation per vendor,

2

experiments already showed mattered – memory capacity, power draw, which execution units a compiled kernel actually reaches, which compiler directives are deterministic enough to trust, and how much memory is reserved rather than usable – have a formal place to live, instead of being silently folded into an existing field or left as an unindexed assumption. Formulating that lift precisely, deriving the predictions it feeds algebraically, and validating both against real hardware is this paper’s entire arc, restated once here: formulate the five fields as index positions (Section IV), derive what they complete in the companion paper’s own equations (Sections IV–V), state precisely what stays open to compiler or runtime optimization versus what is now shape-fixed (Section VI onward), and predict, then validate, every one of those five fields on real hardware for every device in the companion paper’s ensemble (Sections VIII–XII).

corrected evaluation on one device happened to need; it is not a measured property of memory reservation in general, and the companion paper does not claim it is. This paper names overhead as a per-device field for the same reason it names Micap itself: a quantity three different real bugs already turned out to depend on deserves an index position, not an assumed constant reused silently at every future call site. The four figures below make this concrete in order: Figure 2 states the single idea everything else in this paper builds on – that a shape is a shape, whether it names an array or a device – before any operator notation is assumed; Figure 3 then restates the operators themselves, Figure 4 shows a toy partition, and Figure 5 shows the same partition on a real, measured run. C. What this paper does and does not claim This paper’s own real content is the formulation: naming five fields precisely, stating exactly which existing equation each one completes, and showing algebraically that the resulting predictions reduce to the companion paper’s own equations in the special case where the new fields take the values that paper’s real measurements already established. Every number quoted from the companion paper is quoted, not rederived, and every claim about what a real compiler toolchain exposes is sourced to that toolchain’s own documentation [10], [7], [8], [2], [1], [19], not assumed. For one device, NCSA Delta’s A100, this paper also runs a new round of hardware experiments of its own (Section VIII), validating the extended theory directly rather than only formally, with all five fields measured. A second experiment, on Delta’s AMD MI100 (Section IX), validates Cd ’s occupancy-cap field and the Xd /Cd interaction specifically, under this paper’s own admissible vector_size realization rather than the companion paper’s excluded simdlen one, and is now followed, in the same section, by direct measurement of the remaining two fields, Rd and Xd , closing all five on this device. Two further real-hardware experiments, on PSC Bridges-2’s H100 (Section X) and SDSC Expanse’s V100 (Section XI), each close all five fields as well, using a different, vendor-neutral OpenACC realization instead of Triton – a test of the method itself, not only of one device or one kernel family. A fifth real-hardware experiment, on TACC Stampede3’s Intel Max 1550, the one remaining device in the companion paper’s ensemble (Section XII), closes this device in full as well: Micap and Cd ’s occupancy-cap are measured and pinned on real hardware, Pd reads consistently across repeated runs, reached(d) is confirmed directly at 99.88% busy via Level Zero Sysman engine activity, Xd = 0 follows from the kernel’s own construction, and Rd ≈ 0.004 is located by the same bracket-and-locate sweep used on the other three GPUs – all five devices in the companion paper’s ensemble now closed in full, the next step named precisely in Section XV.

B. The two real findings this paper closes The companion paper’s Section on lower-precision realization (its Section IX-G) measures fp32-to-fp16 speedup on all three vendors in its ensemble at the identical reference problem size: 14.7× on NVIDIA (via a Triton kernel whose fused tl.dot reaches the device’s tensor cores directly), 1.05× on Intel, and no speedup at all, within measurement noise, on AMD (both realized as portable OpenMP target-offload code with no path to either vendor’s specialized low-precision hardware). Its energy section (Section IX-H) measures the same three devices’ power draw at fp32 versus fp16 and finds the identical split restated on a second, independent axis: NVIDIA’s power draw rises 36.3% between precisions, while Intel’s and AMD’s shift by less than a percent each – not because one vendor’s silicon is more capable in kind, but because one realization actually reached specialized hardware built for that precision and the other two did not. Nothing in ρmachine (d) as defined – five fields, compute units, warp width, shared memory, measured time, memory ceiling – names this. The companion paper says so explicitly: it names the missing field Xd , execution-unit-type, and immediately qualifies it as conflating two different questions that need separating before either is trusted – whether the hardware exists at all, and whether a specific kernel’s compiled realization actually dispatches to it. That second question is a property of how the kernel was compiled, not of the device, and the companion paper names a second missing field for it, Cd , compiler directives, without populating it. This paper populates both, and states the interaction between them as the single quantity the energy and time predictions actually depend on. The companion paper’s memory-ceiling prediction (its Equation 39) has a second, quieter gap. Its overhead term – the fraction of a device’s memory capacity reserved for driver and runtime footprint, not available to hold the key/value cache – was initially assumed at a conservative 15%, found wrong (a unit error, decimal gigabytes read where the device reports binary gibibytes, was the larger effect), and re-evaluated at 0% once that unit error was corrected. Zero is what one

II. R ELATED W ORK : W HY A F ORMAL , D ETERMINISTIC ACCOUNT OF M ACHINE S HAPE Two mature literatures already predict or optimize GPU kernel performance; both differ from this paper (and its

3

⃗ = ⟨N, D⟩, and the shape of a Fig. 2. The idea this whole paper rests on, stated before any operator notation: MoA treats the shape of an array, ρ(Q) machine, ρmachine (d), as the same kind of object – a tuple with a real value at every position, indexed by the same operators (ψ, ι, ρ, γ). This is why deriving a kernel from an array’s shape and deriving a device’s share of the work from that device’s own shape is one algebra, not two. This paper’s contribution is ⃗ already is. exactly the right-hand side: giving ρmachine (d) the fields it needs to be as complete as ρ(Q)

Performance models bound; ρmachine (d) predicts, and is checked. The Roofline model [24] plots a kernel’s arithmetic intensity against a device’s peak bandwidth and peak compute to bound achievable performance – a real, widely used tool, but a bound, not a specific predicted time, and one that says nothing about whether a given compiled kernel actually dispatches to the execution units that peak figure assumes. That is precisely the gap reached(d) (Section V) closes: this paper’s five new fields, together with the companion paper’s original five, do not bound a family of possible kernels, they predict one specific kernel’s time, power, and memory ceiling on one specific device, and every validation section in this paper (Sections VIII–XI) checks that specific prediction against a live, direct hardware measurement – correctness, utilizationconfirmed dispatch, and timing/power/memory jointly – rather than checking only that a measured point falls under a bound. Hand-tuned kernel libraries achieve speed on one device at a time; this program states, and tests, whether the same derivation transfers. FlashAttention [4] is the clearest example in the attention-kernel space this paper’s own validation targets: an IO-aware, fused kernel, hand-designed around one specific memory hierarchy (on-chip SRAM versus HBM) to avoid materializing the full attention matrix, achieving real, substantial speedups. It is not, and does not claim to be, a portable derivation – porting its specific tiling and recompu-

companion) in the same structural way, which is worth stating precisely rather than only claiming. Autotuning compilers search; MoA derives. Halide [22] separates an image-processing pipeline’s algorithm from its schedule (loop order, tiling, fusion), then finds a good schedule by search – originally manual, later autotuned. TVM [3] and Ansor [25] generalize this to deep-learning operators: a large schedule space is defined, and a cost model plus ondevice measurement searches it, at real, nontrivial compute and wall-clock cost, and with no closed-form guarantee that the mechanism producing a fast schedule on one device transfers to another. This paper’s Cd -admissibility criterion (Section VI) makes the same distinction those systems draw operationally – some schedule choices are safe, deterministic layout decisions, others are heuristics the compiler or autotuner is free to override – but as a formal predicate checked once against a mechanism’s own documentation and hardware manual, not as a search that must be re-run per device, per kernel, and per compiler version. MoA’s DNF/ONF derivation (Section I-A) plays the role Halide’s schedule and TVM/Ansor’s search space play, except that the admissible transformations are fixed by the algebra’s own identities in advance, so a specific Cd configuration is a stated, checkable fact about a realization rather than the output of a search whose result is trusted because it measured fast on one occasion.

4

Fig. 3. The two MoA operators this paper leans on most, made concrete. Top: an array’s role (rl ∈ {0, 1}, partitioned vs. replicated across devices) and width (Dl ) are themselves just ψ-indexed MoA arrays – role and width are data, not metadata bolted on separately. Bottom: γ is the single substitution that turns a device-local index j ∈ ι ci into its corresponding position j + oi in the shared global range ι n – the same operator, applied once, that later reappears as the mechanism behind every device’s own (oi , ci ) slice below.

tation strategy to a new vendor’s memory hierarchy is itself new engineering work, repeated per device family (NVIDIA’s own later kernels for AMD and Intel targets confirm this: they are separate implementations, not one schedule retargeted). The companion paper’s own OpenMP/Triton/OpenACC realizations, and this paper’s extension of ρmachine (d) to explain why they behave differently per vendor, take the opposite bet: one DNF, mechanically retargeted, with the resulting crossdevice performance differences explained in closed form by measured machine-shape fields (Xd , Cd , reached(d)) rather than attributed to unexplained per-vendor engineering effort. Section X and Section XI are a direct test of that bet: the identical OpenACC-tiled DNF-derived kernel, unchanged in its algorithmic structure, run on two different NVIDIA devices (and, with the AMD-specific admissible realization of Section IX-E, a third vendor), with the resulting differences in Cd ’s occupancy-cap and in Rd traced to measured, named fields rather than left as an unexplained gap between one handtuned kernel and another. What is genuinely new here, relative to all three literatures, is not that machine shape matters – Roofline already says so, and every autotuner already searches over some proxy for it. It is that this paper’s own real-hardware findings (Section I-A onward) show a formally derived kernel’s performance can still be explained in closed form, after the fact, using only measured hardware facts and no per-kernel search: reached(d)’s boolean is read off a live profiler or a static disassembly, not learned from thousands of autotuning trials, and it explains a 36.3% power-draw shift and a 14.7× speedup exactly, on the

TABLE I ρMACHINE (d)’ S FIVE ALREADY- INDEXED FIELDS ( COMPANION PAPER , E QUATION 1), RESTATED FOR REFERENCE . T HIS PAPER ADDS TABLE II’ S FIVE FIELDS TO THIS SAME TUPLE . Index

Field

Meaning

Feeds

0 1 2 3 4

Ud Wd Md td nmax d

compute-unit count warp/wavefront width shared mem/block measured time at n0 measured memory ceiling

tile size, Focc tile size tile size wi , t̂i (n) capability check

one device where it is nonzero, and exactly zero shift on two devices where it is measured zero as well. III. BACKGROUND : T HE F IVE -F IELD T HEORY, AND W HERE I T WAS A LREADY S HOWN I NCOMPLETE The companion paper defines a device’s machine shape as an open-ended tuple,
ρmachine (d) = Ud , Wd , Md , td , nmax , ... , d with each field ψ-selectable at its own index position – Ud = ψ(⟨0⟩, ρmachine (d)) through nmax = ψ(⟨4⟩, ρmachine (d)) d – and derives, from this tuple alone, a weighted partition (its Equation 26), a predicted finish time (its Equation 11, restated as Equation 36 for a new problem size), a predicted memory ceiling (its Equation 39), and a predicted energy (its Equation 14). Table I restates the five fields already indexed, for reference; this paper adds to this table rather than replacing any of it.

5

Fig. 4. A concrete, minimal instance – a toy 64-row query array split across three invented devices, sized inversely to each device’s own measured time tdi ⃗ is cut into three disjoint, unequal ψ-selected slices, one per device, via the offset/count pair (oi , ci ). Bottom: (Equation 26 of the companion paper). Top: Q ⃗ slice, plus a full, identical, replicated copy of K ⃗ and V ⃗ (hatched), since attention’s key/value tensors what each device actually holds – its own unique Q are read whole by every device rather than partitioned. This is the same partition structure Section VIII onward measures on real hardware, at real problem sizes, rather than a toy count.

The companion paper’s own future-work section names four further quantities without indexing them: Micap (device i’s total memory capacity, distinct from Md ’s on-chip per-block figure), Pi (average power draw), Xd = ψ(⟨5⟩, ·) (execution-unittype), and Cd = ψ(⟨6⟩, ·) (compiler directives) – explicitly stating that assigning these their own index positions is real, substantial future work, deliberately out of that paper’s own scope. Reading its energy section closely surfaces a fifth: the memory-ceiling prediction’s overhead term, treated there as an assumed constant re-evaluated once by hand rather than as a field read from any device.

TABLE II T HE FIVE FIELDS THIS PAPER ADDS , AND THE EQUATION ( NUMBERED AS IN THE COMPANION PAPER ) EACH ONE COMPLETES . Index

Field

Meaning

Completes

5 6 7 8 9

cap Mi Pd Xd

total device memory (bytes) power draw, per realization execution-unit capability compiler-directive config. memory-reservation overhead

Eq. 39 Eq. 14 precision gap precision gap Eq. 39

Cd Rd

already uses this quantity, unindexed, in its Equation 39; indexing it here changes nothing about that equation’s arithmetic, only its provenance – Micap becomes a field γ reads, not a number substituted by hand.

IV. T HE F IVE F IELDS , F ORMULATED Table II states all five additions at once, each with the equation it completes; the rest of this section formulates each in turn, at the same level of rigor the companion paper already applies to Ud through nmax . d

B. Pd : power draw, indexed by realization, not device alone The companion paper’s own energy section states plainly that holding Pi constant as problem size grows is “a physically reasonable assumption... but it is exactly that: an assumption, not yet a validated one.” Its own real numbers show why that caution is warranted in a second, sharper sense: Pi is not even safely constant across precision for every device. NVIDIA’s measured power rises from 143.30 W (fp32) to 195.30 W (fp16) – a 36.3% increase, because its Triton realization’s tensor cores genuinely engage more silicon per cycle at fp16. Intel’s and AMD’s power draw is flat across precision, within

A. Micap : total device memory Micap = ψ(⟨5⟩, ρmachine (di )), device i’s own total memory capacity in bytes, distinct from Md = ψ(⟨2⟩, ·), which is on-chip per-block shared memory – a quantity two orders of magnitude smaller and governing a different limit (occupancy, not capacity). The companion paper

6

Fig. 5. Array data selected by array machine information, on a real, validated two-device run (Intel Max 1550 + AMD MI100, n = 65,536, companion paper). ⃗ (and Out) is produced from that same device’s own ψ-selected ρmachine (di ) entry – concretely, td = ψ(⟨3⟩, ρmachine (di )) Each device’s own row-slice of Q i ⃗ that device computes on. K ⃗ and V ⃗ , by contrast, are identical ⃗ i = ψ(oi +ι ci , Q) feeds the weight wdi , which in turn fixes (oi , ci ) and hence the exact slice Q and unsliced on both devices: no ρmachine (di ) entry participates in selecting them at all. This is the exact mechanism this paper’s five new fields plug into: each is a new ρmachine (di ) entry, selected by the identical ψ notation already driving the partition above, not a new kind of quantity requiring new notation.

measurement noise, because neither realization changes which execution units it dispatches to. Pi , indexed as a single perdevice field, would need to average over or arbitrarily pick one of these regimes; indexed instead per realization,

D. Cd : compiler-directive configuration, restricted to a shapeadmissibility criterion

Pd = ψ(⟨6⟩, ρmachine (d)),

a per-realization record of the compiler configuration that produced the measured td , Pd . Before stating what belongs in Cd , we state a criterion for what is allowed to: MoA’s own DNF (the companion paper’s Equation 4) fixes the computation once, independent of any target, and γ’s stated job (that paper’s Section II) is to derive an ONF from that DNF and a target’s own measured shape – never to alter what the DNF computes. A sub-quantity of Cd is therefore admissible only if it selects among values of a real, alreadyexisting hardware shape quantity – a cache level, a vector or tile width, a thread/block count – with no reordering of the DNF’s own operations beyond what that quantity’s native execution already forces. A sub-quantity is inadmissible if it gives the compiler discretionary license to reassociate, fuse, reorder, or restructure the computation in ways the chosen hardware does not itself require – because at that point the flag is no longer selecting a shape, it is silently editing the DNF. A single “no reordering” test is not sufficient on its own, however, and an earlier draft of this paper applied it alone. Some mechanisms neither reorder the DNF’s arithmetic nor edit it, yet still leave the actual outcome to the compiler’s own internal, undocumented cost model – a heuristic enabled rather than a value set. Recording such a mechanism in Cd would give a field with a name and an index position but

Cd = ψ(⟨8⟩, ρmachine (d)),

with ρmachine (d) itself understood – as Section V makes precise – as depending on which compiled kernel is asking, Pd is exactly as reliable as td already is, for the identical reason: both are measured once per realization, not assumed to transfer across a change the measurement never covered. C. Xd : execution-unit capability Xd = ψ(⟨7⟩, ρmachine (d)), a per-device record of which specialized low-precision execution paths physically exist – Intel’s AMX, AMD’s matrix cores, NVIDIA’s tensor cores – populated from a hardware fact (ISA documentation, or a profiler’s static capability query), never from whether any particular kernel happens to use them. The companion paper’s own Section X already states this distinction is necessary, not merely tidy: conflating “the device has this hardware” with “this realization reaches it” is precisely the error that would have made AMD’s flat fp16 result look like a hardware limitation, when the companion paper’s own analysis attributes it instead to the OpenMP realization never being written to reach AMD’s matrix cores at all.

7

pattern the second part excludes: setting n does not guarantee the compiler honors it. The deterministic mechanism for the identical quantity is different and real: GCC’s and Clang’s vector_size attribute specifies a vector’s width directly, in bytes, as part of its type – not a hint subject to a cost model, a value the compiler is obligated to honor [5]. NVIDIA’s Triton is not exposed to this ambiguity at all: BLOCK_M, BLOCK_N, BLOCK_K, and num_warps are kernel-language tiling and launch parameters the compiler lowers directly, not vectorizer hints subject to a separate legality/cost pass. Vector/tile width is therefore admissible on both counts, but only through these deterministic mechanisms – explicit vector-extension typing on Intel and AMD, kernel-language tile constants on NVIDIA – never through an auto-vectorization enable flag, on any vendor’s toolchain. Table III applies this criterion to every sub-quantity surveyed in Section VI. Three survive: memory-hierarchy placement (which cache level a prefetch targets), vector/tile/threadcount width (bounded exactly as argued above), and occupancy/launch-bound caps (a resource constraint on block count, not a reordering of operations within one). Loop-unroll factor and inline depth – included in this paper’s own earlier draft of Cd , before this criterion was applied – are excluded: unrolling produces the identical class of reassociation vector width does, but as a compiler-discretionary transformation available even on hardware that does not itself force it [23], and inline depth corresponds to no shape quantity ρmachine names at all. Cd is restricted accordingly:
Cd = cache-level, width, occupancy-cap ,

no known relationship to what the compiled kernel actually does, defeating the entire purpose of indexing it: a field of ρmachine exists so that γ can read a value and know, in closed form, what that value implies. We therefore add a second, independent requirement: a sub-quantity is admissible only if the value written to it is a deterministic directive – what the programmer specifies is exactly what the compiler does, with no compiler-internal cost-model step free to override, decline, or reinterpret it. A mechanism that merely enables an optimization and lets the compiler decide, case by case, whether and how to apply it fails this second test even when it does not reorder anything on its own, because Cd would then record a setting whose actual effect on td is not knowable from the setting alone. This second requirement resolves a case the first, reordering-only criterion left genuinely ambiguous in this paper’s own earlier draft: host-side prefetch insertion (Intel’s -qopt-prefetch[=n], AMD’s -mllvm -enable-x86-prefetching). Neither flag reassociates arithmetic, so the first criterion alone does not exclude it – but both are documented as heuristic enables: the compiler decides where, whether, and how far ahead to insert a prefetch, using its own internal analysis, not a value the programmer supplied. A field populated from such a flag would record “prefetching was turned on” without recording what the compiled kernel actually does, which is not a known parameterization in the sense the second criterion requires. It is excluded. Vector and tile width is the one sub-quantity that needs both parts of the criterion, not just one, so it is worth walking through in full. On the first part – no discretionary reordering – executing a reduction (the companion paper’s own S[i, j] = scale × +redk (. . .), Equation 4) on a w-wide vector unit necessarily reassociates that reduction’s summation order relative to scalar execution – concretely, ((0+4)+(2+6))+((1+5)+(3+7)) in place of ((((0+1)+2)+3)+4) for eight terms on a 4-wide unit [23], changing rounding error even though the sum is mathematically identical. That reassociation is forced by which real device is chosen, not discretionary within a fixed device – exactly the deviation the companion paper’s own correctness checks already tolerate, passing every device “to within what fp64 rounding alone accounts for” rather than requiring bitexact agreement across devices of different width. Width therefore passes the first part: its reordering effect is bounded by, and only by, the real hardware ρmachine (d) already names, not by anything the compiler is free to choose beyond that. The second part – a deterministic directive, not a heuristic enable – is where width looked settled in this paper’s own earlier draft but was not. LLVM’s own documentation is direct that its loop and SLP vectorizers apply an internal cost model and “may only vectorize... if forced,” declining even when a width hint or pragma requests it [12] – meaning AMD’s specific -mllvm -enable-epilog-vectorization with -min-width-epilog-vectorization=n, surveyed earlier as admissible, is exactly the heuristic-enable

a three-field structured record, smaller than this paper’s own first attempt at Cd , by design: a field excluded here for editing the DNF is not a smaller convenience, it is a correctness requirement this paper’s own admissibility criterion exists to enforce. E. Rd : memory-reservation overhead Rd = ψ(⟨9⟩, ρmachine (d)) ∈ [0, 1), the fraction of Micap unavailable to hold a served computation’s resident arrays, rewriting the companion paper’s memory ceiling (its Equation 39) as Micap × (1 − Rd ) , 4×D×b identical in form to that paper’s own equation, with Rd now a field rather than a number chosen once and reused. Every symbol on the right is itself a stated, deterministic quantity, not a free parameter: D is the attention head dimension (fixed by the served model’s own architecture, D=64 throughout this paper’s real experiments); b is bytes per element at the realization’s chosen precision (b=4 for fp32, b=8 for fp64); and the factor 4 counts the number of identically (n, D)shaped resident device arrays this kernel family holds live per ⃗ K, ⃗ V ⃗ , and Out – so 4 × D × b is the query row at once – Q, exact per-row byte footprint the memory boundary is measured nmax ≈ d

8

against, not an empirically fit constant. The companion paper’s own real debugging history is the argument for indexing this: its first evaluation used Rd = 0.15, a conservative guess; its corrected evaluation used Rd = 0, a value that happened to fit once a unit error was separately fixed. Neither value was measured, and nothing says 0% generalizes to a different device, driver version, or runtime. Indexing Rd does not resolve what its real value is on any given device – that is real-hardware work, named in Section XV – it only stops the theory from silently treating an unmeasured guess as a closedform constant.

shape-associated and a deterministic directive, admissible) or X (exclude – either a discretionary reordering that edits the DNF, or a heuristic enable whose actual effect the compiler, not the programmer, decides). An earlier draft of this survey also carried a B (borderline) verdict for mechanisms that do not themselves reorder arithmetic but leave the compiler free to decide whether and how to act; the second, determinism requirement resolves every such case in this table to a firm K or X, so no row below is left borderline. Each row is cited to the issuing vendor’s own documentation, never inferred from observed behavior. Reading the table requires nothing beyond the two-part test just restated. Each row names one vendor-exposed subquantity in the left-most column, states in one sentence what the compiler or runtime actually does with it, and closes with a single verdict letter; no row is scored from this paper’s own inference, since every citation names the issuing vendor’s documentation directly, so a reader unfamiliar with any one of the three toolchains can check a verdict without taking it on this paper’s word. A reader new to ρmachine and Cd should read the K rows first: cache-level placement, vector/tile width, and occupancy cap are exactly the three fields Section IV-D already named as Cd ’s structured record, restated here across three real vendors rather than assumed to generalize from the one this paper’s companion measured. The X rows are not failures of the criterion – they are the compiler’s own discretionary territory, named precisely so that Section VII-E can measure its effect directly rather than leave it folded silently into a device’s constant. The three toolchains surveyed here are not an arbitrary sample: Intel, AMD, and NVIDIA are the only vendor toolchains this paper’s companion measured directly, on real hardware, elsewhere in this ensemble (Sections VIII–XII), so every verdict below is checked against a compiler this paper’s own experiments actually ran, never against a toolchain assumed by name alone. A fourth toolchain – a new vendor’s compiler, or a future revision of one already surveyed – extends this table by the same two-question test, never by re-deriving Cd itself: the criterion stated in Section IV-D does not name Intel, AMD, or NVIDIA, only the shape-associated/deterministic distinction those three vendors’ own documentation happens to make checkable today. This survey is not a detour from the paper’s own arc. Sections IV-D–V formulated and derived the two-part criterion in the abstract, stating exactly what a shape-associated, deterministic sub-quantity must look like without yet checking whether any real toolchain actually exposes one. Table III is where that derivation is checked against real vendor documentation for the first time – the state-admissible step of this paper’s own dimension-lift arc, applied here to Cd specifically, before Table IV repeats the same check one abstraction layer up, against OpenMP, OpenACC, and MPI rather than a single compiler’s command-line flags. Neither table changes the criterion itself; both discharge the same formula against a wider vocabulary each time, which is the sense in which this section – like every other survey in the paper – validates a

V. reached(d): T HE I NTERACTION Xd AND Cd A LONE C ANNOT S TATE Naming Xd and Cd separately is necessary but not sufficient: neither field alone predicts the fp16 speedup or powerdraw asymmetry Section IV opened with. What predicts it is whether Cd ’s own configuration actually dispatches to whichever path Xd says exists – a joint fact about the pair, not a property of either field read in isolation. We name this
reached(d) = dispatches Cd , Xd ∈ {0, 1}, a boolean (or, where a realization partially engages specialized hardware, a degree in [0, 1]) read directly from a profiler – Intel VTune and the Level-Zero Sysman interface report whether AMX was reached; NVIDIA Nsight Compute reports tensorcore utilization directly; AMD’s rocprof reports MFMA (matrix-core) utilization – never inferred from precision alone, since precision change without reached(d) = 1 is exactly the flat-power, flat-speedup case the companion paper measured on Intel and AMD. With reached(d) named, the companion paper’s time and energy predictions (its Equations 11 and 14) are restated, not replaced:
W (ci , n) , ti = ti reached(di ) , (1) t̂i (n) = ti × W (n0 , n0 )
Êi (n) = Pi × t̂i (n), Pi = Pi reached(di ) . (2) Both equations are unchanged in form; what changes is that ti and Pi are now understood as functions of reached(di ) – of which realization measured them – rather than as fixed perdevice constants a new kernel is free to assume transfer unchanged. This is exactly the discipline the companion paper’s own Section IX-G recommends in prose (“the fix... is adding this missing shape dimension to γ’s own input... re-deriving Intel’s and AMD’s ONF from a complete shape vector”) stated here as the closed form that recommendation was missing. VI. A S URVEY OF W HAT T HREE R EAL T OOLCHAINS E XPOSE , S CORED AGAINST THE A DMISSIBILITY C RITERION Every candidate sub-quantity found across the three toolchains this paper’s companion already measured – nine in total, wider than Cd ’s own three admitted fields – is listed in Table III with an explicit verdict, scored against both requirements of Section IV-D’s criterion: K (keep –

9

derivation already stated rather than introducing a new one. One row in Table III earns neither a flat K nor a flat X, and is worth walking through on its own because the reasoning generalizes to any future row that looks like it. ISA targeting – Intel’s -x/-march= family – selects an instruction-set width, which by itself is a real, shape-associated hardware fact: it fixes how many lanes a vector instruction actually spans, exactly like the vector/tile-width row directly above it. What keeps it from being a flat K is that several of these flags silently enable fused multiply-add contraction as a side effect of raising the target ISA, and FMA contraction reassociates a reduction’s rounding exactly as the excluded fast-math row does – a discretionary reorder riding along with an otherwise admissible directive, not named anywhere in the flag the programmer actually wrote. The fix is not to exclude the row; it is to state the second requirement explicitly rather than leave it implicit, which is what “conditional” means throughout this table: ISA targeting is admissible if and only if it is paired with -fp-model=strict -ffp-contract=off (or the equivalent anti-contraction flag on the compiler in question), so that the width selection and the rounding behavior are pinned independently. Omit the pairing and the row silently reverts to X – the compiler is once again free to reassociate – with nothing in the command line saying so. That conditional case is also the cleanest preview of what Section VII-E does with every X row in this table: each one (k) is a candidate for a paired on/off run under ϵd , holding the admissible fields fixed and toggling only the excluded flag, so that whatever effect it has on td or Ed is measured rather than silently absorbed into a device constant the way it was before this table existed. Table IV carries the same three-column shape – sub-quantity, real action, verdict – one abstraction layer up, and Section VI-A immediately below applies both tables together to a fourth, automated source of Cd candidates: an autotuner’s own search space, which mixes admissible and inadmissible axes in exactly the pattern this section’s K/X split was built to separate. A caution belongs here before the table itself, so it is read once rather than re-derived per row: a K verdict states that a sub-quantity is admissible, not that any particular device has measured it. Cache-level placement, vector/tile width, and occupancy cap are the three fields this paper’s companion actually closed on real hardware; the remaining K rows surfaced by this survey – for instance a second vendor’s own cache-placement enum, or a third’s tile-width attribute – are admissible by the same criterion but not yet folded into Cd ’s structured record for any device this ensemble measures, precisely the gap Section XV’s roadmap exists to close in order, one field at a time, rather than all at once on the strength of the criterion alone. Table III is therefore a map of what is allowed, not a report of what has already been measured – the two are kept separate throughout this paper on purpose, since conflating them is exactly the failure mode the companion paper’s own Rd = 0.15 guess (Section IV) illustrates: a plausible, admissible-looking number substituted for one nobody had actually gone and read off a real device.

10

TABLE III E VERY COMPILER - DIRECTIVE SUB - QUANTITY SURVEYED ACROSS I NTEL , AMD, AND NVIDIA, SCORED AGAINST S ECTION IV-D’ S TWO - PART ADMISSIBILITY CRITERION . K = KEEP ( SHAPE - ASSOCIATED and A DETERMINISTIC DIRECTIVE ), X = EXCLUDE ( DISCRETIONARY REORDERING , OR A HEURISTIC ENABLE THE COMPILER IS FREE TO OVERRIDE ). Sub-quantity

Vendor mechanism

Real action

Cache-level placement

Intel LSC_LDCC enum (L1UC_L3C default, L1C_L3C recommended), chosen in source, not by a compile flag [10], [7] Intel: fixed at 512-bit, 16 fp32 lanes [8], or explicit vector_size attribute typing. AMD: vector_size attribute typing (GCC/Clang front end) [5] – not -mllvm -enable-epilogvectorization with -min-width-epilogvectorization=n, a cost-modelgated auto-vectorizer enable that the compiler may decline even when set [12], excluded below under Host-side prefetch’s own reasoning. NVIDIA: 32-thread warp, BLOCK_*/num_warps, kernellanguage constants the compiler lowers directly [19] AMD --gpu-max-threadsper-block=n [2]

Selects which level of the existing memory hierarchy a prefetch targets – a placement choice, no arithmetic reordering

K

Selects a real hardware SIMD/warp/tile width; reassociates a reduction’s summation order, but bounded exactly by which real device is chosen, not by compiler discretion within one device – and, via vector_size/kernel-language constants, is a value the compiler is obligated to honor, not a hint it may decline (Section IV-D)

K

Caps concurrent thread blocks the compiled kernel can support – a resource constraint, not an operation reorder; the compiler-side counterpart to the companion paper’s own data-side occupancy-collapse finding (Section IX-D) Inserts early, non-blocking loads ahead of dependent accesses – does not itself reassociate arithmetic, but is a heuristic enable: the compiler’s own internal analysis decides where, whether, and how far ahead to insert a prefetch, so the flag does not fix a known value the way the second admissibility requirement demands Duplicates the loop body, reordering instruction issue and, for a reduction loop, reassociating summation order exactly as vector width does – but discretionarily, not forced by any chosen hardware [23] Restructures control flow across call sites; corresponds to no shape quantity ρmachine names

K

Explicitly permits fusing multiply-add and reordering FP reductions; Intel’s own diagnostics document -fp-model=precise disabling vectorization the same compiler performs under fast [6] – the clearest direct violation of the criterion

X

Selects instruction-set width, but AVX-512 targeting silently enables FMA contraction by default [11] – width selection alone is K; admissible only when explicitly paired with -fp-model=strict/-ffp-contract=off to strip the coupled reassociation Cross-function inlining and specialization beyond a single translation unit – a restructuring, by its own documentation

K, conditional

Reorders instruction issue across loop iterations – the same class of reassociation as loop unrolling, discretionary rather than hardware-forced

X

Vector/tile width

Occupancy / launch-bound cap Host-side prefetch insertion

Loop-unroll factor Inline depth

FP contraction / numeric strictness ISA targeting

IPO / wholeprogram optimization Softwarepipelining depth

Intel -qopt-prefetch[=n], -qopt-prefetchdistance [7]; AMD -mllvm -enable-x86prefetching [2] No dedicated flag documented on Intel or NVIDIA; AMD’s related loopsplitting flags [2] AMD -mllvm -inline-recursion=[1–4], -finlineaggressive [2] Intel -fp-model= {fast,precise,strict}; AMD -ffast-math, -ffp-contract= {fast,on,off} Intel -x/-march=

Intel -ipo [9]; AMD -flto [2] NVIDIA Triton num_stages (default 3) [19]

Verdict

Restricting Cd to the three K rows – and treating ISA targeting as admissible only paired with an explicit anticontraction flag – gives the structured record already stated in Section IV-D. Table III’s own excluded rows are still useful, but in a different role than tuning: run once with a subquantity enabled and once without, holding every admissible field fixed, each excluded flag isolates whether a measured timing difference is attributable to that specific discretionary transformation – a diagnostic use, never a production configuration this paper’s own predictions should depend on.

X

X X

X

A. Autotuning: admissible only over the restricted subset Triton’s @triton.autotune does not set one parameter; it searches a declared grid – tile shape (BLOCK_M, BLOCK_N, BLOCK_K), num_warps, num_stages – compiling and measuring each candidate and caching whichever configuration measured fastest. Under Table III’s verdicts, this search mixes an admissible axis (tile shape and num_warps, both real width quantities) with an inadmissible one (num_stages, software-pipelining depth) in a single tool. Using autotuning to populate Cd therefore requires

11

restricting its search grid to hold num_stages fixed at its own minimum (the least pipelining the toolchain permits) and varying only tile shape and num_warps – recovering an empirical search over exactly Cd ’s own admissible width field, nothing wider. Run unrestricted, autotuning remains useful only as the same kind of diagnostic named above: its full search log, including the num_stages dimension, measures how sensitive td is to pipelining depth specifically, informative for deciding whether Section XV’s real-hardware validation needs to treat pipelining as a fourth, separatelyjustified admissible field in a future revision of this criterion – not evidence to fold into Cd as defined here.

removes the correctness objection to vectorizing that plain auto-vectorization has to infer on its own, but it does not remove the compiler’s discretion over whether or how wide to vectorize – so it is excluded for the identical reason AMD’s auto-vectorization enable flag already is, even though it is written as an explicit pragma rather than a command-line flag. B. OpenACC: parallel is a directive, kernels is not OpenACC’s parallel construct takes num_gangs, num_workers, and vector_length clauses that hold one gang/worker/vector configuration fixed for the entire region [18] – the same real-hardware-width argument already used for SIMD width, extended from an ISA-level width to an accelerator-level one. The kernels construct accepts the identical three clauses syntactically, but the specification states that the compiler “will split the code in the kernels region into a sequence of accelerator kernels,” whose “number and configuration of gangs of workers and vector length may be different for each kernel” [18]. A single clause value on kernels therefore does not fix one configuration for the region the way the identical clause does under parallel: the compiler’s own kernel-splitting decision, not the clause, determines what a given loop nest actually runs with. kernels is excluded for the same reason schedule(auto) is – the actual configuration is a compiler decision the clause does not pin down – while parallel with explicit clauses is admissible.

VII. T HE S AME C RITERION A PPLIED B EYOND THE C OMPILER : O PEN MP, O PENACC, AND MPI Sections IV-D–VI-A restricted Cd using compiler command-line flags and one kernel-language extension (Triton) as the surveyed vocabulary. Nothing in Section IV-D’s two-part criterion is specific to a compiler’s own flags, however: a test of “does this value select among a real hardware quantity, and is it a deterministic directive or a heuristic enable” applies to any directive written into source, including the pragma-based APIs that annotate the companion paper’s own loop nests (OpenMP, OpenACC) and the communication library a realization spanning more than one device must eventually call (MPI). This section applies the same two questions to those three real, cited APIs and finds the identical split.

C. MPI: default collective selection is not a directive; forcing an algorithm or a placement is

A. OpenMP: thread count and static scheduling are directives; dynamic scheduling and simd are not

Only once a realization spans more than one device does a communication library enter Cd at all – the companion paper’s own partition is explicitly communication-free, so this subsection documents the criterion’s next application rather than a mechanism this paper’s ensemble currently uses. Open MPI’s default collective-selection mode is, by its own documentation, “a decision tree, essentially a large set of nested if-then-else-if blocks with baked in comm and message size thresholds derived by measuring performance on existing clusters,” that “selects one of a number of available algorithms” internally [17] – a heuristic exactly like LLVM’s vectorizer cost model, excluded for the same reason. Setting coll_tuned_<op>_algorithm to an explicit number is different: the documentation states that doing so means “the fixed decision rules associated with that collective are short circuited” [17], so the algorithm actually run is exactly the one named, known before the call executes. Rank-tocore placement splits the same way: --bind-to core and --map-by core bind processes to cores “in a roundrobin manner,” a fixed algorithm computable in advance from process and core counts, and rankfile assigns each rank to an explicitly named core [16] – both admissible, selecting among the real hardware topology ρmachine already names for occupancy, on the same grounds as Table III’s occupancy-cap row.

num_threads(n) fixes the exact thread count a parallel region uses – a real hardware core/thread count, set exactly, admissible under both criteria. The schedule clause splits cleanly on the second criterion alone. schedule(static[,chunk_size]) assigns chunks to threads “in a round-robin fashion in the order of the thread number” – a fixed mapping computable from thread count and chunk size alone, before the loop runs [21]. schedule(dynamic) and schedule(guided) instead distribute chunks to threads as they become available, and schedule(auto) “delegate[s] the decision regarding scheduling... to the compiler and/or runtime system” [21] – three heuristic-enable patterns, differing only in mechanism (a runtime work queue rather than a compiler cost model) from the auto-vectorization enable flags Table III already excludes: a field populated from dynamic or guided would record that adaptive scheduling was requested, not which thread actually executed which iteration. The simd construct fails the same test for a subtler reason. It asserts a correctness property – no loop-carried dependency prevents concurrent execution – but the specification is explicit that “the number of iterations that are executed concurrently at any given time is implementation defined,” and that the companion simdlen(n) clause sets only “the preferred number of iterations,” not a guaranteed one [20]. omp simd

12

TABLE IV T HE SAME TWO - PART ADMISSIBILITY CRITERION (S ECTION IV-D) APPLIED TO O PEN MP, O PENACC, AND O PEN MPI. K = KEEP ( SHAPE - ASSOCIATED and A DETERMINISTIC DIRECTIVE ), X = EXCLUDE . Sub-quantity

Vendor mechanism

Real action

Thread count

OpenMP num_threads(n)

K

Static loop scheduling

OpenMP schedule(static [,chunk_size]) [21]

Dynamic / guided / auto loop scheduling

OpenMP schedule(dynamic), schedule(guided), schedule(auto) [21] OpenMP simd construct, simdlen(n) [20]

Fixes the exact thread count a parallel region uses – a real hardware core/thread count, set exactly Assigns chunks to threads in round-robin order by thread number – a fixed mapping computable from thread count and chunk size alone, before the loop runs Distributes chunks to threads as they become available, or delegates the mapping decision to the compiler/runtime entirely – which thread executes which iteration is not knowable from the clause alone Asserts loop-dependence safety, but the actual concurrent-iteration count is implementation defined; simdlen sets only a preferred width, not a guaranteed one Holds one accelerator-level gang/worker/vector configuration fixed for the entire region – the same real-hardware-width argument already used for SIMD width

X

SIMD (OpenMP)

width

Gang / worker / vector width

Accelerator kernel decomposition

OpenACC parallel, num_gangs/ num_workers/ vector_length [18] OpenACC kernels (same clauses) [18]

Collective algorithm (default)

Open MPI coll/tuned fixed-decision tree [17]

Collective algorithm (forced)

Open MPI coll_tuned_<op>_ algorithm=n [17] Open MPI --bind-to core, --map-by core, rankfile [16]

Process / placement

rank

Verdict

Compiler splits the region into a sequence of accelerator kernels whose own gang/worker/vector configuration may differ per kernel – the clause value does not fix the region’s actual configuration Selects a collective’s algorithm from an internal, message-size-keyed decision tree built from prior benchmarking – a heuristic, not a value the caller supplied Short-circuits the fixed-decision tree; the algorithm actually run is exactly the one named, known before the call executes Selects among the real hardware topology ρmachine already names for occupancy – a placement choice computable in advance, not a reorder

K X

K

X X K K

change what Cd can honestly record: with the algorithm shortcircuited, Cd ’s value is the algorithm that runs, not a request the runtime is free to reinterpret per call. That is the same distinction Section IV-D draws for every other row in both tables, stated here once more because it is easy to mistake for a performance argument (“the forced algorithm is faster”) when it is actually a determinism argument: Cd is admissible only when its written value and its realized effect are the same fact seen twice, not because the written value happens to measure well on one run.

Table IV restates this same two-part test – does the value select among a real, already-existing hardware or topology quantity with no discretionary reordering, and is the value written a fixed directive rather than a preferred, compileroverridable approximation – across all three subsections above, row by row. The pattern each subsection just argued in prose is visible there at a glance: every admissible (K) row is a shape component this paper’s ρmachine (d) can name and γ can substitute a real value into before the kernel runs – thread count, static chunk assignment, gang/worker/vector width, a forced collective algorithm, a fixed rank-to-core mapping – while every excluded (X) row hands that same decision to the compiler or runtime at call time, leaving nothing for γ to read in advance. The table applies no new criterion; it is the one already stated in Section IV-D, run once each against OpenMP, OpenACC, and Open MPI in turn. Take the MPI collective row as one concrete instance of the general pattern. Open MPI’s own default – the decision tree, keyed off message size and cluster topology it measured once, at build time, on hardware that need not be the hardware this paper runs on – is excluded for exactly the reason schedule(auto) and LLVM’s auto-vectorizer cost model already were: the value γ would need to substitute is not fixed until the call executes, so no ρmachine (d) entry can name it in advance. Forcing coll_tuned_<op>_algorithm to an explicit number changes nothing about the arithmetic MPI performs – a reduction still sums the same terms – but it does

D. What this confirms Table III’s split was not an artifact of surveying compiler command-line flags specifically. Table IV’s same two questions – does the value select among a real, already-existing hardware or topology quantity with no discretionary reordering, and does the compiler or runtime honor exactly the value written, not a preferred or fixed-decision approximation of it – separate deterministic from heuristic directives identically across a compiler’s own flags, a loop-annotation API, an accelerator-offload API, and a communication library. Cd is not redefined by this section; it is shown to generalize to every directive class a real MoA-derived realization actually passes through, not only the one this paper happened to survey first. E. ϵd : the residual once determinism is fully isolated Every X row in Table III and Table IV is a real mechanism a real toolchain exposes; excluding it from Cd does not make

13

2000-repeat window (fp32: 5.114 ms → 3.969 ms; fp16: 0.346 ms → 0.265 ms) – warm-up residue not cleared by three untimed calls alone. Pd needed the same fix for a different reason: at 20 repeats the cumulative energy counter (nvmlDeviceGetTotalEnergyConsumption) produced a physically impossible 1114 W average, above the device’s own 400 W rated limit; at ≥ 2000 repeats (a several-second window) the derived average agrees with an independent instantaneous reading (nvmlDeviceGetPowerUsage) to within 1–2 W at both precisions, which is what makes Table V’s Pd values trustworthy rather than merely present. Both corrections are reported here, rather than silently discarded, because they are themselves evidence for this paper’s own governing principle (Section IV-D): a measurement protocol can look deterministic and still carry a hidden heuristic-like sensitivity – here, to sample count – that only shows up once checked against a second, independent method. Rd is itself a real finding, not only a replaced placeholder: it is higher at fp16 (0.500) than fp32 (0.250) – the reservationoverhead fraction did not shrink with the smaller dtype, it doubled. A single device cannot say whether this is a general property of reduced-precision execution or specific to this kernel and this GPU generation; it is named here as a concrete target for the remaining four devices. reached(d), measured directly (NVIDIA Nsight Compute, metric sm__pipe_tensor_op_hmma_cycles_active.avg. pct_of_peak_sustained_elapsed), is 20.73% and 20.74% tensor-pipe active-cycle fraction across two profiled fp16 calls, against a 0.55–0.56% noise floor measured on unrelated background kernels in the same trace. This confirms the tensor-core path is dispatched – nonzero, well above the noise floor – but not saturated, real evidence that reached(d) ∈ {0, 1} as formulated in Section V does not fit the measured behavior cleanly on this device. This paper does not resolve that tension: whether to keep reached(d) boolean with a principled dispatch-fraction threshold, or restate it as a continuous quantity in [0, 1] and re-derive how it composes with Equations (1)–(2), is left as a concrete, real-data-motivated question for the remaining ensemble. (k) Finally, ϵd is measured, not merely named, for one excluded lever this device’s kernel language exposes directly – Triton’s own autotuner (Section VI-A), run unrestricted (its full search grid, including the inadmissible num_stages dimension) against the pinned Cd baseline:

TABLE V NCSA D ELTA A100: MEASURED ρMACHINE (d) FIELDS , Cd PINNED Field

fp32

fp16

td Pd Rd correctness

3.969 ms 179.99 W 0.250 PASS (3.9 × 10−7 )

0.265 ms 256.9 W 0.500 PASS (6.1 × 10−5 )

its effect on td vanish, it only stops the theory from pretending that effect is known when it is not. Once the admissible (K) mechanisms are identified in full – the task this section and Section VI complete – the excluded mechanisms name a second, independent family of quantities worth indexing formally: for each heuristic-controlled lever k (host-side prefetch, an auto-vectorization enable, OpenMP’s dynamic/guided scheduling or simd width, OpenACC’s kernels decomposition, Open MPI’s default collective selection), (k)

ϵd

(k)-on

= t̂d

(k)-off

− t̂d

,

the measured difference in realized time between running with lever k left to its default, heuristic behavior and running with every admissible Cd field pinned at a stated value and lever k pinned off (or at its own most conservative setting). Unlike Cd , (k) ϵd is not written by the programmer – it cannot be, precisely because k is not a deterministic directive – it is measured, and it is only meaningful once a fully deterministic baseline exists to measure it against. Establishing that baseline is the (k) work of Sections IV-D–VII; ϵd itself is stated here only as notation, with its first real measurement – one excluded lever, Triton’s autotuner, on NCSA Delta’s A100 – reported in Section VIII. That single measurement is a proof of the method, not its completion: every other excluded row of Table III and Table IV, on every device this paper measures, remains real-hardware work, named precisely in Section XV. VIII. R EAL -H ARDWARE VALIDATION : NCSA D ELTA A100 This section reports the first real-hardware experiment this (k) paper runs: all five fields, reached(d), and ϵd for one excluded lever, measured directly on NCSA Delta (NSF ACCESS allocation CIS261342), NVIDIA A100-SXM4-40GB, using the companion paper’s own reference kernel (fused, single-pass online-softmax attention, Triton, n = 8192, B = 1, D = 64). Cd is pinned to this paper’s own admissibleonly configuration throughout: fixed tile shape (BLOCK_M=64, BLOCK_N=64), Triton’s autotuner off. Correctness against PyTorch’s fused scaled_dot_product_attention is checked before any timing or power number is trusted, at both precisions (Table V). Micap = 40960 MiB, Xd = 3rd-generation Tensor Cores (confirmed via compute_cap 8.0; a second, independent cross-check via deviceQuery was unavailable on this system). Both td figures required a methodology correction mid-measurement: an initial 20-repeat timed window over-measured td by 20–25% relative to a

(autotune)

ϵd

off = t̂on d − t̂d

= 0.239 ms − 0.265 ms = −0.026 ms (≈ −9.8%)

(3)

reproduced identically across three independent paired runs on a dedicated allocation with a fixed random seed. The heuristic lever is real, repeatable, and faster here – the opposite sign from an earlier, differently-conditioned measurement recorded in this same kernel’s own source comments. This is direct

14

(k)

evidence for Section VII-E’s premise: ϵd ’s sign is an empirical question, not an assumption to carry forward from one measurement to the next, even on the same device and the same kernel.

TABLE VI NCSA D ELTA AMD MI100: MEASURED ρMACHINE (d) FIELDS , V E C T O R _ S I Z E - REALIZED Cd , OCCUPANCY- CAP RE - DERIVED

IX. R EAL -H ARDWARE VALIDATION : NCSA D ELTA AMD MI100 This section reports the second real-hardware experiment this paper runs, and the first for a device where the companion paper’s own width mechanism, OpenMP’s simdlen(n), is exactly the inadmissible case Table III excludes: a preferred, not guaranteed, vector width. This paper’s own criterion (Section IV-D) instead admits vector_size-typed C attribute vectors as the deterministic alternative. All numbers below are measured from that admissible realization, not the companion paper’s excluded one, on NCSA Delta (NSF ACCESS allocation CIS261342), AMD Instinct MI100 (gfx908), using the companion paper’s own reference kernel and problem size (n = 8192, B = 1, D = 64, single-pass online-softmax attention, OpenMP target-offload). Correctness against a CPU double-precision reference is checked at both precisions before any timing is trusted (Table VI).

Field

fp32

fp16

td (occupancy-cap= 10) td (occupancy-cap= 32, tuned) Pd correctness

1432.250 ms 453.149 ms 142.8 W PASS (5.5 × 10−6 )

1417.692 ms 447.357 ms 136.2 W PASS (4.9 × 10−4 )

493.75 ms and 1405.15 ms (1.3% and 0.1% apart). The admissible vector_size realization, compiled with the same occupancy-cap, width, and tile shape carried over unchanged, measured 2.9× slower: fp32 1432.250 ms. A controlled sweep of occupancy-cap × tile width × vector width (24 combinations, every cell correctness-PASS and GPUdispatch-confirmed by the same live-utilization method as Section IX-A) found the realization’s true saturation point at occupancy-cap = 32 – not 10 – with tile width and vector width unchanged: fp32 446.377 ms, reproduced at 453.149 ms on an independent re-run. Occupancy-cap, once re-derived for the realization actually in use, does not merely close the gap to the simdlen baseline – it surpasses it. This is itself evidence for this paper’s own construction: occupancy-cap is admitted into Cd as a field co-determined with width and tile shape (Section IV-D), not a per-device constant reusable across whichever mechanism populates the width field. The identical occupancy-cap value that was correct for one admissible-width realization was wrong by a factor of three for another, at the identical width value, on the identical device.

A. A methodological finding before the numeric one: confirming reached(d) requires live dispatch evidence, not compile success Delta’s gpuMI100x8 partition is a mixed node: eight real MI100 GPUs and one MI210 GPU share a single host (sinfo: gpu:mi100:8,gpu:mi210:1). A generic resource request (--gpus-per-node=1) does not pin GPU type, and three consecutive allocations landed on the lone MI210 slot. Because the kernel was compiled with -march=gfx908 (MI100’s ISA, distinct from MI210’s gfx90a), the architecture-mismatched code object did not dispatch – and the OpenMP runtime fell back to host execution silently: no compile error, no runtime error, a correctnessPASS result, and a plausible-looking timing (23.5 s for the width-64 realization at n = 8192). Only direct GPU utilization evidence (rocm-smi, polled once per second across the run) exposed the fallback: GPU use (%) read 0 for the run’s entire duration. This is a direct, real instance of Section V’s own argument – reached(d) is a property of whether a specific compiled realization actually dispatches to the intended hardware, not a property inferable from correctness or compile success alone – and it cost a full order of magnitude in the reported time before being caught. Every timing in Table VI below is confirmed by the same live GPU use (%) > 0 evidence, not by correctness or successful compilation alone.

C. Pd : resolved via an independent-method substitution As on the A100, Pd here required a correction before it could be trusted, though for a different reason than a samplecount artifact. ROCm-SMI’s cumulative energy-counter API, rsmi_dev_energy_count_get, returned an identical before/after reading – a zero delta – across a confirmed ∼ 9 s, 20call timed window during which rocm-smi’s own concurrent polling recorded sustained 93–177 W draw at 100% GPU utilization throughout. A zero-delta counter under nearly nine seconds of confirmed load is not a plausible physical reading; it indicates the accumulator itself does not increment on this device/driver combination, not that no energy was drawn. Pd was instead derived from rocm-smi’s instantaneous power reporting (--showpower), polled at 0.1 s resolution throughout each timed run and averaged over samples independently confirmed at 100% GPU utilization: Pd = 142.8 W (fp32, n=58 samples) and 136.2 W (fp16, n=57 samples), consistent across two independently polled runs and well within the device’s ∼ 300 W TDP. This is the same governing principle the A100 section states (Section IV-D) reappearing in a different form: a measurement API can present as reliable while silently returning nothing usable, and the fix is the same in kind – a second, independent method, not a longer window.

B. Occupancy-cap is realization-specific, not device-specific The companion paper’s simdlen-realized AMD baseline was measured at occupancy-cap (OpenMP thread_limit) = 10. Re-measured here on confirmed real MI100 hardware at the identical occupancy-cap, tile shape, and width, that baseline reproduces almost exactly: fp32 500.000 ms and fp64 1406.546 ms against the companion paper’s cited

15

D. Micap and Rd : memory boundary

TABLE VII PSC B RIDGES -2 H100: MEASURED ρMACHINE (d) FIELDS , O PENACC- TILED Cd , OCCUPANCY- CAP CHECKED

Micap

= 34,342,961,152 B (≈ 32 GiB), read directly from rocm-smi rather than assumed from the device’s nominal specification. Bracketing the companion paper’s pre-rewrite kernel boundary (≈ 28.5M elements, fp32) against this admissible realization’s own per-team local buffers (k_tile/v_tile/s_tile, now vector_sizetyped), N = 20,000,000 is confirmed correctness-PASS (max_err= 1.3905 × 10−2 , within the relaxed large-N tolerance this test uses). The remaining bracket, needed to locate nmax and solve Rd from nmax ≈ Micap (1 − Rd )/(4Db) d d (D = 64, b = 4 bytes, fp32), could not be completed on first attempt: the gpuMI100x8 partition’s sole node suffered a hardware NODE_FAIL mid-sweep and remained down for the balance of that measurement window, with no fallback node available in the partition. Once the partition recovered, a six-point bracket located the boundary directly: N ∈ {32,200,000, 32,500,000, 32,800,000, 33,100,000} correctness-PASS; N ∈ {33,300,000, 33,450,000} fails (HSA_STATUS_ERROR_OUT_OF_ RESOURCES), giving Rd ≈ 0.013 (1.3%) – notably smaller than H100’s independently measured Rd ≈ 0.030 (Section X), a first data point on how much Rd varies across vendors rather than a value that could have been guessed once and reused.

Field

fp64

td (num_workers=16, vector_length=32) Pd cap Mi correctness

25.445 ms 135.05 W 81,559 MiB PASS (7.77 × 10−15 )

static form of evidence: the compiled device code object is extracted from the binary’s .llvm.offloading section (via llvm-objcopy, falling back to a raw ELF-magic-byte scan when llvm-offload-binary’s triple-matching silently failed) and disassembled at gfx908 with llvm-objdump. The resulting 8,973-instruction disassembly contains 169 ordinary floating-point FMA opcodes (v_fma_f32) – confirming this is genuine device code and not an empty or host-code artifact – and zero matrix-core v_mfma_* opcodes of any kind. Xd = 0 for this device is therefore closed by direct static evidence at the same evidentiary standard as a live profiler counter, not merely reported as consistent with the flat-timing result. X. R EAL -H ARDWARE VALIDATION : PSC B RIDGES -2 H100

E. The flat AMD fp16 result holds under the admissible realization

This section reports a third real-hardware experiment, on a device the companion paper’s own ensemble includes but this paper had not yet independently remeasured: PSC Bridges-2’s H100 (NSF ACCESS allocation CIS261342, account cis260272p). Unlike the A100 and MI100 sections above, the reference realization here is not the Triton kernel but the portable, vendor-neutral OpenACC tiled kernel from prior work [14] – deliberately, as a test of whether this paper’s five-field method itself is realization-agnostic, not only device-agnostic. That kernel’s occupancy-cap (num_workers/vector_length) was hardcoded rather than overridable; it was made a compiletime macro here, with its original values (16/32) preserved as the default, specifically so Cd ’s occupancy-cap field could be checked rather than assumed transferable, the same discipline Section IX-B applies to MI100. Correctness against a CPU double-precision reference (fp64 throughout – this realization has no lower-precision variant) and reached(d) via live nvidia-smi utilization polling are checked before any timing or power number is trusted, the same order of operations as every device above. reached(d) is confirmed directly, not inferred: live nvidia-smi polling at 0.1 s resolution shows GPU utilization ramping 0% → 68% → 100% → 96% across the timed window, with power climbing from a ∼ 69 W idle baseline to a 186.85 W peak – the same live-dispatch discipline Section IX-A used to catch MI100’s silent host fallback, applied here from the start rather than after a problem surfaced. Pd = 135.05 W is derived as the average power over exactly the utilization-confirmed-active samples (n=6),

At the tuned, GPU-dispatch-confirmed occupancy-cap, fp16 measures 447.357 ms against fp32’s 453.149 ms – a 1.3% difference, within the same measurement noise the companion paper reports for its own (inadmissible-realization) AMD measurement. This closes a specific open question this paper’s own theory raised but did not itself answer: whether the companion paper’s flat AMD fp16 result (its Section IX-G, “no speedup at all, within measurement noise, on AMD”) was a real hardware property, or an artifact of measuring it through simdlen, the mechanism Table III itself excludes as non-deterministic. It is not an artifact: switching to the admissible vector_size realization, at its own correctly re-derived occupancy-cap, reproduces the identical flat result. This closes the timing side of the question, but Xd itself required direct evidence, not inference from timing alone, per Section IX-A’s own argument that reached(d)adjacent claims should rest on live dispatch evidence. rocprofv3 was the available profiler on this system (not legacy rocprof/rocprofv2); both its --stats and bare --kernel-trace instrumentation reliably triggered a GPU memory-access fault (write to a read-only page) during the kernel’s single dispatch, a fault absent from every unprofiled run of the identical binary, including every correctness- and utilization-confirmed run used elsewhere in this section. This indicates a latent synchronization issue in the realization’s target-offload region, sensitized by rocprofv3’s HSA-level dispatch-tracing overhead but not present at normal execution cadence. Rather than leave Xd resting on inference while that fault is outstanding, this paper substitutes an equally direct but

16

well within the H100 SXM variant’s 700 W TDP – the same independent-method substitution used for MI100’s Pd (Section IX-C), applied here as the primary method from the outset rather than as a fallback after a cumulative counter proved unreliable. Micap is read directly from nvidia-smi, not assumed from nominal specification.

TABLE VIII SDSC E XPANSE V100: MEASURED ρMACHINE (d) FIELDS , O PENACC- TILED Cd , OCCUPANCY- CAP CHECKED

A. Occupancy-cap, checked rather than assumed: flat on H100, unlike MI100 A controlled sweep of num_workers× vector_length (9 combinations, the tenth – num_workers=32, vector_length=64, 2048 threads/block – exceeds the hardware’s 1024-thread/block limit and was skipped; every remaining cell correctnessPASS and reached(d)-confirmed by the same live-utilization method) finds this realization’s timing essentially flat across the entire swept range: 25.445–25.635 ms, under 1% spread, with the minimum tied at (16, 16) and the paper-cited baseline (16, 32). This is a genuinely different outcome from MI100’s occupancy-cap finding, not merely a confirmation of the same one: on MI100, the carried-over occupancy-cap value was wrong by a factor of three once actually swept instead of assumed; on H100, the carried-over value turns out to already be at (or statistically indistinguishable from) this realization’s saturation point. Both results support the same underlying claim – occupancy-cap is a field that must be checked per realization and device, not assumed transferable by citing a shared hardware constant – but only because both were actually measured. Had only the MI100 result been available, “occupancy-cap is realization-specific” could have been overstated as “occupancy-cap is never transferable”; the H100 result shows that conclusion would have been wrong in general, even though it was right for MI100. Rd required the same QUERY_CHUNK/QUERY_OFFSET rewrite MI100’s kernel needed before a memoryboundary sweep was possible; once rewritten, a six-point bracket located the boundary directly: N = 40,500,000 correctness-PASS, N = 41,500,000 fails, giving Rd ≈ 1 − 40,500,000 × 4 × 64 × 8/(Micap in bytes) ≈ 0.030 (3.0%). Xd was confirmed directly via NVIDIA Nsight Compute (sm__pipe_tensor_op_hmma_cycles_active.avg. pct_of_peak_sustained_elapsed): 0% across all profiled passes of this fp64-only realization – direct profiler evidence, not an inference from H100’s tensor cores being fp64-incapable, that no dispatch to them occurs. XI. R EAL -H ARDWARE VALIDATION : SDSC E XPANSE V100

Field

fp64

td (num_workers=8, vector_length=64, min) td (num_workers=16, vector_length=32, baseline) Pd cap Mi Rd Xd correctness

60.405 ms 59.769–60.490 ms 131.69 W 32,768 MiB 0.028 (2.8%) 0% PASS (7.88 × 10−15 )

and reached(d) via live nvidia-smi utilization polling are checked before any timing, power, or memory-boundary number is trusted, the same discipline applied to every device above. reached(d) is confirmed directly: live nvidia-smi polling at 0.1 s resolution shows GPU utilization ramping 0% → 43% → 98% → 100% across the timed window, power climbing from a ∼ 42 W idle baseline to a 158 W peak. Pd = 131.69 W is the average power over the utilization-confirmedactive samples (n=15), well within the V100 SXM2 variant’s 300 W TDP. Micap = 34,359,738,368 B (32,768 MiB) is read directly from nvidia-smi. A. Occupancy-cap: flat, as on H100 A controlled sweep of num_workers× vector_length (8 combinations; the ninth, num_workers=32/vector_length=64, exceeds the 1024-thread/block limit and was skipped; every cell correctness-PASS and reached(d)-confirmed) finds this realization’s timing flat on V100 as well: 60.405–63.059 ms, a 4.4% spread, minimum at (num_workers=8, vector_length=64), with the paper-cited baseline (16, 32) statistically indistinguishable from it. This is the same qualitative result as H100 (Section X-A) and the opposite of MI100’s: a third data point supporting the same underlying claim – occupancy-cap must be checked, not assumed transferable – while also showing that result is not uniform across NVIDIA hardware by fiat; it is a property this realization happens to share on both NVIDIA devices measured so far, confirmed independently on each rather than assumed from one to the other. B. Rd : memory boundary, an independent NVIDIA measurement A six-point bracket, chosen from the accumulated experience of the H100 (Rd ≈ 0.030) and MI100 (Rd ≈ 0.013) sweeps rather than a wide first pass, located the boundary in a single round: N ∈ {15,900,000, 16,100,000, 16,300,000} correctnessPASS; N ∈ {16,500,000, 16,650,000, 16,750,000} fail, each with an identical, unambiguous signature (cuMemAlloc

This section reports a fourth real-hardware experiment, on the last NVIDIA device in the companion paper’s ensemble this paper had not yet independently re-measured: SDSC Expanse’s Tesla V100-SXM2-32GB (NSF ACCESS allocation CIS261342). As on H100, the reference realization is the portable OpenACC tiled kernel, unchanged from Section X apart from the memory-boundary rewrite already required there. Correctness against a CPU double-precision reference

17

error 2, out of memory, on the kernel’s third device-side array = 16,300,000: allocation). Taking nmax d

TABLE IX TACC S TAMPEDE 3 I NTEL M AX 1550: MEASURED ρMACHINE (d) FIELDS , V E C T O R _ S I Z E - REALIZED Cd , OCCUPANCY- CAP SWEPT, ALL FIVE FIELDS

16,300,000 × 4 × 64 × 8 ≈ 0.028 (2.8%), 34,359,738,368 close to H100’s independently measured 0.030 and noticeably larger than MI100’s 0.013 – a second NVIDIA data point at a similar magnitude, and a third confirmation that Rd is a real, device-dependent quantity worth indexing rather than a correction term whose value can be guessed once and reused.

CLOSED BY DIRECT MEASUREMENT

Rd ≈ 1 −

C. Xd : confirmed directly, without the workaround MI100 needed Unlike MI100, where rocprofv3 faulted under every instrumentation mode and Xd required static disassembly of the compiled device code object (Section IX-E), NVIDIA Nsight Compute ran cleanly on V100 with no workaround needed: the same metric used on H100, sm__pipe_tensor_op_hmma_cycles_active.avg. pct_of_peak_sustained_elapsed, read 0% across all eight profiled passes. V100’s tensor cores are fp16-input-only and this kernel is fp64 throughout, so Xd = 0 is the expected result; confirming it directly, by the same live-profiler standard applied everywhere else in this paper, closes the field rather than leaving it as an assumption consistent with the device’s known architecture.

Field

fp64

td (nw=8, vl=16) td (nw=16, vl=16, pinned) Pd (pinned, 4-run avg) cap Mi (per tile) Xd reached(d) Rd correctness

338.27 ms 180.22 ms 306.40 W 63.98 GiB 0 (by construction, see text) 99.88% (Sysman engine activity) 0.004 (0.4%) PASS (1.02 × 10−14 )

A. Occupancy-cap: a plateau at 16, and a real admissiblewidth ceiling at vector_length=64 A controlled sweep of num_workers× vector_length over {8, 16, 32, 64} × {16, 32, 64} (12 cells, block_m=block_n=16 held fixed) finds two real results, not one. First, an occupancy plateau: num_workers ∈ {16, 32, 64} are statistically indistinguishable at fixed vector_length (e.g. 180.221, 180.210, 180.246 ms at width 16 – a 0.02% spread), while num_workers=8 undersaturates the device by nearly a factor of two (338.266 ms at the same width). The pinned value is the smallest that reaches the plateau, num_workers=16, the same reasoning already applied to every device above. Second, and unlike every NVIDIA and AMD device measured so far, width does not help here: vector_length=16 is 3.1× faster than vector_length=32 at every tied num_workers value (180 vs. 556 ms), the opposite direction from MI100’s own width result. Third, vector_length=64 fails outright – a reproducible segmentation fault at every num_workers value tried, not an intermittent one. At D = 64, vector_length=64 means a single vec_t spans the entire feature dimension: sizeof(vec_t) = 64 × 8 = 512 B, a 4096-bit vector type (comp_t is double throughout). The 2048-bit realization (vector_length=32) compiles and runs correctly; the 4096-bit one does not, on any occupancy-cap tried. This is a real, empirically located boundary on icx’s vector_size codegen for this target, not a harness defect – every width that does execute is correctness-PASS at 1.02 × 10−14 – and it belongs in Table III as a device-specific limit on an otherwise portable mechanism: the same C attribute vector type that is admissible in principle on both AMD and Intel is not admissible at every width on both, and that ceiling can only be located by running, not read off a spec sheet. The pinned configuration is num_workers=16, vector_length=16, block_m=block_n=16.

XII. R EAL -H ARDWARE VALIDATION : TACC S TAMPEDE 3 I NTEL M AX 1550 This section reports a fifth real-hardware experiment, on the one device in the companion paper’s ensemble this paper had not yet attempted: TACC Stampede3’s pvc partition (NSF ACCESS allocation CIS261342), Intel Data Center GPU Max 1550 (Ponte Vecchio). The kernel is ported, not reused: same DNF, same tiling, same online-softmax structure as the MI100 realization (Section IX), recompiled for OpenMP target-offload to Level Zero via Intel’s icx (oneAPI DPC++/C++ Compiler, LLVM/Clang-based, -fopenmp-targets=spir64). The vector_size attribute mechanism itself required no change at all: icx and AMD’s amdclang are both Clang-derived front ends, so the identical admissible Cd realization compiles unmodified on a second, independent vendor toolchain – a genuine finding for Table III, distinct from NVIDIA’s nvc (PGI-derived), which required the separate num_workers/vector_length pragma-clause mechanism instead. As on every device above, correctness against a CPU double-precision reference is checked before any timing or power number is trusted. Level Zero enumerates eight level_zero:gpu devices on the allocated node (confirmed via sycl-ls); every run below is pinned to device 0 via ZE_AFFINITY_MASK=0, the same one-GPU-per-run discipline used on every other device in this paper. Micap = 63.98 GiB is read directly by the kernel itself, via Level Zero Sysman’s zesDeviceEnumMemoryModules – physically consistent with the Max 1550’s 128 GB HBM2e capacity split across two tiles, one tile per Level-Zero device.

B. Pd : consistent across four confirmatory runs, plausible against the device’s per-tile power budget Unlike A100 and MI100, where Pd required an independent-method substitution before the reading could be trusted (Sections VIII, IX-C), Level Zero Sysman’s

18

zesPowerGetEnergyCounter returned a real, physically plausible, and internally consistent reading on the first attempt: 306.182, 306.268, 306.181, and 306.981 W across four runs of the pinned configuration – the last a separately dispatched confirmatory run (num_workers=8 and vector_length=32 read 303.1 and 302.3–302.9 W respectively, all within the same tight band). Averaged, Pd = 306.40 W – consistent with one tile of the Max 1550’s dual-tile OAM module drawing roughly half the card’s rated board power.

V100’s does from its fp16-only tensor cores against an fp64 kernel, but without V100’s live profiler confirmation available to double-check it. This is stated plainly as the one field on this device closed by argument rather than by direct measurement, together with reached(d)’s own confirmation (Section XII-C) that the GPU is genuinely doing the work – just not through a matrix engine this kernel never asks for. E. Rd : memory boundary, the smallest overhead measured on any device in this paper The same bracket-and-locate technique already used for MI100 (Section IX-D) and, in the same pattern, for H100 and V100 (Sections X and XI) applies here at the pinned Cd configuration (num_workers=16, vector_length=16, block_m=block_n=16, B = 1): sweep N upward until the device-side allocation of the kernel’s four mapped arrays – Q, K, V (map(to)) and Out (map(from)), each N × D × 8 B at D = 64, fp64 – fails, and read Rd off how much of Micap the largest passing N leaves on the table. The bracket located a clean, reproducible boundary: N = 33,400,000 completes the device kernel call and returns normally (a real, plausible 130.9 s timed call at 295.5 W); N = 33,545,166 – the theoretical exact-capacity point, ⌈Micap /(4 × D × 8)⌉ for this device’s confirmed 63.98 GiB – fails outright with a real Level Zero/OpenMP allocation error ("Call to getTargetPointer returned null pointer (device failure or illegal mapping)’’, aborting before the kernel launches), the same evidentiary standard (a real runtime error, not a timeout or an assumption) as every other device’s own Rd closing. From nmax = 33,400,000, d

C. reached(d): confirmed via Level Zero Sysman engine activity, not VTune This paper’s Section V names two tools for this device, VTune and Level Zero Sysman. VTune’s gpu-hotspots collection, the one used for genuine hardware-counter profiling on other vendors’ equivalents, failed outright on this shared TACC allocation: "Cannot collect GPU hardware metrics due to a lack of permissions’’, requiring root or render-group membership plus /proc/sys/dev/i915/perf_stream_paranoid=0, neither available here and not worth a support ticket to chase. Level Zero Sysman’s own engine-activity API, zesEngineGetActivity, needs no elevated permissions – the same is already true of the power-domain read above – and was added to the kernel’s timed-loop bracket exactly as the energy counter already is. Over a confirmed 45.31 ms wall-clock window, the compute engine group reports 45.26 ms of active time: 99.88% busy, unambiguous, the same evidentiary standard as nvidia-smi/rocm-smi utilization polling on every other device in this paper – a real GPU busy fraction, not an assumption, confirming reached(d) = 1.

Rd ≈ 1 −

D. Xd : zero by construction, not by static disassembly

33,400,000 × 4 × 64 × 8 ≈ 0.004 (0.4%), 68,702,699,437

the smallest overhead measured on any device in this paper’s ensemble (against A100’s 0.250, MI100’s 0.013, H100’s 0.030, and V100’s 0.028) – plausible given the Max 1550’s Sysman-reported Micap is already a flat, per-tile HBM2e figure with, evidently, very little driver/context reservation subtracted from it before a kernel can use it, unlike a discrete GPU’s own CUDA/ROCm context overhead. One incidental artifact surfaced at N = 33,400,000, just below the boundary: the same run’s own reached(d) engine-activity read (Section XII-C’s API, read again in this sweep) reported an unphysical value from an unsigned 64-bit counter wraparound (a negative activeTime delta read as a huge positive one) – a Sysman engine-counter quirk under near-capacity memory pressure, not a device-dispatch failure: the kernel call itself completed, returned, and reported a real, physically plausible energy reading in the same run. This does not affect the Rd measurement above, which depends only on the kernel call’s own success or failure, not on that counter; it is noted here as a boundarycondition caveat on Sysman’s engine-activity API specifically, distinct from the clean 99.88% reading Section XII-C reports at the pinned, small-N configuration.

Unlike MI100, where a broken profiler was worked around by static disassembly of the compiled device code object (Section IX-E), that same workaround is not available on this system: ocloc, the offline AOT compiler needed to regenerate a native device binary from the kernel’s SPIR-V intermediate and disassemble it, is present (bundled inside VTune’s own profiler tooling) but depends on Intel’s Graphics Compiler backend, libigc.so.1, which is not installed anywhere on this allocation – confirmed by exhausting the library cache (ldconfig -p), the RPM database, and a filesystem search, all turning up no IGC or GPU compute-runtime package at all, only the CPU-side OpenCL runtime. Xd is instead reported the way V100’s was motivated (Section XI-C) but stated here more precisely, by construction rather than by profiling: the kernel is fp64 throughout and every arithmetic op uses the vec_t (vector_size) mechanism directly – plain multiply-accumulate over C-level vector types, with no XMX/DPAS intrinsic, subgroup-matrix call, or any other matrix-engine-targeting construct anywhere in the source. No compiler backend can dispatch to hardware a kernel gives it no path to reach; Xd = 0 follows from the source the same way

19

TABLE X A LL FIVE NEW FIELDS , reached(d), AND CORRECTNESS , ACROSS ALL FIVE DEVICES MEASURED IN THIS PAPER . td /Pd AT EACH DEVICE ’ S OWN MEASURED PRECISION ( FP 32 FOR A100/MI100, FP 64- ONLY FOR H100/V100/M AX 1550); Rd /MiCAP ARE PRECISION - INDEPENDENT AND DIRECTLY COMPARABLE . Field

A100 (NCSA Delta)

MI100 (NCSA Delta)

H100 (Bridges-2)

V100 (Expanse)

Max 1550 (Stampede3)

td Pd Rd Micap Xd Cd reached(d) correctness

3.969 ms (fp32) 179.99 W 0.250 (25.0%) 40,960 MiB 3rd-gen Tensor Cores BLOCK_M/N=64 20.7% (partial) PASS (3.9×10−7 )

453.149 ms (fp32) 142.8 W 0.013 (1.3%) ≈32 GiB 0 (static disasm.) vector_size, occ.-cap=32 confirmed (live util.) PASS (5.5×10−6 )

25.445 ms (fp64) 135.05 W 0.030 (3.0%) 81,559 MiB 0% num_workers=16/vl=32 confirmed (live util.) PASS (7.77×10−15 )

60.405 ms (fp64) 131.69 W 0.028 (2.8%) 32,768 MiB 0% num_workers=8/vl=64 confirmed (live util.) PASS (7.88×10−15 )

180.22 ms (fp64) 306.40 W 0.004 (0.4%) 63.98 GiB (per tile) 0 (by construction) vector_size, occ.-cap=16/vl=16 confirmed (99.88%, Sysman) PASS (1.02×10−14 )

All five new fields, plus reached(d), are now closed on this device: Micap , Cd , and Rd by direct measurement, Pd by a real and consistent Sysman reading, reached(d) by real engineactivity evidence, and Xd by a stated construction argument rather than direct measurement – TACC Stampede3’s Intel Max 1550 is the fifth and last device in the companion paper’s ensemble to close in full, joining NCSA Delta’s A100, AMD’s MI100, PSC Bridges-2’s H100, and SDSC Expanse’s V100 (Table X).

(Eq. 39), energy (Eq. 14) – as a function of a ρmachine (d) that is now, by construction, complete with respect to every gap that paper’s own real hardware experiments surfaced. This is real theoretical progress: before this paper, a reader implementing the companion paper’s derivation for a new device would have no formal place to record “does this device’s fastest kernel actually reach its tensor cores,” and would be forced either to omit that fact or smuggle it into td silently, exactly the conflation the companion paper’s own Section X already warns against in prose. For every device in the companion paper’s ensemble, this paper now reports full empirical progress: Micap , Pd , Xd , Cd , and Rd are formulated here with the same rigor the companion paper applies to Ud through nmax , and all five d fields are closed by direct measurement on all five devices, rather than left open or quoted from elsewhere. NCSA Delta’s A100 (Section VIII) was full from the start: real values for every field, measured by this paper’s own new experiment. AMD’s MI100 (Section IX) began partial and closed in two stages during the same measurement window: Cd ’s occupancycap, Micap , and Pd (the latter via an independent-method substitution after the device’s own cumulative energy-counter API proved non-functional) were measured first; Rd ≈ 0.013 followed once the allocation’s failed node recovered and a completed memory-boundary bracket located the boundary directly; and Xd = 0 was closed not by rocprofv3, which continued to fault under every instrumentation mode tried, but by static disassembly of the compiled device code object – zero matrix-core opcodes against a nonzero ordinaryFMA sanity check, the same evidentiary standard as a live profiler counter reached by a different route. PSC Bridges2’s H100 (Section X) shows the same closing pattern: td , Pd , Micap , and a checked (not assumed) occupancy-cap came first, finding it already at saturation rather than 3× off as on MI100; Rd ≈ 0.030 and Xd = 0% followed once the kernel’s QUERY_CHUNK/QUERY_OFFSET rewrite made a memory-boundary sweep possible and Nsight Compute was run for the first time. SDSC Expanse’s V100 (Section XI) is the fourth device closed in full, and the cleanest of the three GPUs measured this session: the same occupancy-capflat result as H100, Rd ≈ 0.028 at a similar magnitude to both GPUs measured directly, and Xd = 0% confirmed via Nsight Compute with no workaround needed at all – a

XIII. A LL F IVE D EVICES , E VERY F IELD , IN O NE V IEW Sections VIII–XII report each device on its own; Table X restates all five new fields, reached(d), and correctness for all five side by side, so the pattern across devices is visible without paging back through five separate tables. td and Pd are read at the precision and realization each device was actually measured under – Triton fp32 for A100, OpenMP fp32 for MI100, OpenACC fp64 for H100 and V100, OpenMP fp64 for Intel Max 1550 – and are not a cross-vendor race for that reason: A100 and MI100 also have fp16 figures (Tables V and VI), while H100, V100, and Max 1550 have no lower-precision variant to compare against. Rd and Micap , by contrast, are precision-independent device properties, so all five entries in those two rows are directly comparable as they stand. Every entry of Table X is a shape component in the literal sense Section I-A sets up: ψ(⟨k⟩, ρmachine (d)) for the field’s own index k and the device’s own d. Read column by column, the table makes visible the same correspondence Figure 5 draws for td alone: a device’s own sub-array of data – its ⃗ i = ψ(oi +ιci , Q) ⃗ – is produced by that same device’s slice Q own sub-array of machine shape, the column ρmachine (di ) this ⃗ into per-device slices and table restricts to. Partitioning Q partitioning ρmachine ’s own array of devices into per-device columns are the same operation, applied to two different arrays; this table exists so a reader can see that second partition directly, the way Figure 4 already makes the first one visible rather than leaving it as a formula to evaluate mentally. XIV. D ISCUSSION : W HAT T HIS C LOSES , AND W HAT I T D OES N OT Sections IV and V together restate every prediction the companion paper makes – time (Eq. 11/36), memory ceiling

20

useful contrast with MI100, where the equivalent field required substituting an entirely different evidentiary method. TACC Stampede3’s Intel Max 1550 (Section XII) is the fifth device, and closes in full as well: Micap and Cd measured and pinned directly, Pd reading consistently across four confirmatory runs, reached(d) confirmed at 99.88% busy via Level Zero Sysman engine activity after VTune’s own hardware-metric collection proved blocked by permissions on this allocation, Xd = 0 closed by a stated construction argument rather than the static disassembly used for MI100, whose own backend dependency is not installed on this system, and Rd ≈ 0.004 located by the same bracket-and-locate sweep used on the other three GPUs – the smallest overhead of any device measured, and a clean Level Zero/OpenMP allocation-failure signature at the theoretical exact-capacity point, not an assumption. Closing a gap in the theory and closing the gap between predicted and measured numbers are two different achievements; this paper claims both, in full, for all five devices.

simd, dynamic/guided scheduling, OpenACC kernels, default MPI collective selection – on every remaining device, turns those tables’ own excluded rows from a diagnostic list into a measured error bound across the whole ensemble: not whether a heuristic lever matters, but by how much, on real hardware, once every field this paper can state in closed form is already accounted for. XVI. C ONCLUSION A prediction that quietly assumes away the one variable a real experiment already showed matters is not wrong by accident; it is incomplete by construction. This paper closes that construction gap, not the empirical one: Micap , Pd , Xd , Cd , and Rd now have index positions in ρmachine (d), a stated equation each one completes, and, for Xd and Cd together, a named interaction, reached(d), that the companion paper’s own real numbers already show is necessary. For one device, NCSA Delta’s A100, Section VIII shows directly that measuring these five fields narrows the gap: real td , Pd , and Rd values in hand, (k) reached(d) confirmed dispatched but not saturated, and ϵd measured, reproducibly, with a sign an unvalidated assumption would have gotten backwards. For a second device, Delta’s AMD MI100, Section IX shows the same narrowing in part: Cd ’s occupancy-cap re-derived directly rather than assumed transferable across width realizations, reached(d) confirmed, concretely, as a question live GPU-utilization evidence must answer – a realization that compiles cleanly, passes correctness, and reports a plausible time can still be running entirely on host – and Micap and Pd measured directly, the latter via a second, independent method after the device’s own cumulative energy-counter API proved silently non-functional under real, confirmed load. Rd ≈ 0.013 and Xd = 0 close out MI100 in full: the former once the allocation’s failed node recovered and a completed memory-boundary bracket located the boundary directly, the latter not from the only profiler available on this system, which faulted under every instrumentation mode tried, but from static disassembly of the compiled device code object – zero matrix-core opcodes against a nonzero ordinary-FMA sanity check, the same evidentiary standard as a live profiler counter reached by a different route when the intended one proved unusable. For a third device, PSC Bridges-2’s H100, Section X shows the same narrowing again, and adds a genuinely new result rather than a repeated one: occupancy-cap, checked by the identical sweep discipline used on MI100, turns out already near-optimal on H100 rather than 3× off. Measuring the same field on two devices and finding opposite outcomes – wrong by a factor of three on one, already correct on the other – is stronger evidence for treating occupancy-cap as a field that must be checked, not assumed, than either result would have been alone; a single device’s finding could not have distinguished “this is a general property of occupancy-cap” from “this is a fact about MI100.” Rd ≈ 0.030 and Xd = 0% close out H100 in full as well, the former once the kernel was rewritten for a memory-boundary sweep, the latter confirmed directly via Nsight Compute. For a fourth device, SDSC Expanse’s V100, Section XI closes all

XV. F UTURE W ORK Sections VIII–XII close all five fields, by direct measurement, on all five devices in the companion paper’s ensemble; Table X (Section XIII) and the Discussion (Section XIV) already restate what each device’s own closing looked like, so this section states only what is actually still to do, rather than repeating that record a third time. The direct next step is the companion paper’s own predicted-versus-measured checks (its Tables VI, VII, VIII, IX): re-run them under the extended equations (1)–(2) to test directly whether reached(d) closes the prediction gap those tables already report, or whether a further, currently-unnamed field is still missing – and whether the A100’s own open question, whether reached(d) is better stated as boolean or continuous, holds across vendors or is itself device-specific. A second, controlled experiment the companion paper’s own Section X already proposes – same device, same precision, only the kernel implementation changed (hand-written intrinsics or a vendor-tuned library in place of the generic OpenMP target-offload realization used throughout that paper) – is the direct test of whether Table III’s uneven toolchain coverage is a real hardware ceiling or a realization gap software alone can close, separating Xd ’s hardware claim from Cd ’s implementation claim empirically rather than only formally. Both experiments require the same real ACCESSallocated hardware the companion paper already used, and neither is attempted in this paper. A third experiment follows directly from Section VII-E, and Section VIII reports its first instance: once the admissible fields of Cd are pinned to stated values on a given device, measure (k) ϵd for each excluded lever in Tables III and IV in turn – one controlled run with k pinned off, one with k left at its default heuristic behavior, holding every admissible field fixed – using the same profiling tools already named above. The remaining work is breadth, not method: the A100 measurement covers one lever (Triton’s autotuner) on one device; every other row of Tables III and IV – prefetch insertion, omp

21

five fields independently in a single experiment: occupancycap flat as on H100, Rd ≈ 0.028 at a magnitude close to both other GPUs measured directly, and Xd = 0% confirmed via Nsight Compute with no workaround needed at all – underscoring, by contrast, that MI100’s disassembly route was a genuine substitution for a broken profiler, not this paper’s default method. For a fifth and final device, TACC Stampede3’s Intel Max 1550, Section XII closes the same narrowing in full: Micap and Cd ’s occupancy-cap measured and pinned directly on real hardware, locating a genuine admissible-width ceiling – a reproducible segmentation fault at vector_length=64, avoided by a narrower, faster, working width – that no prior device in this paper’s ensemble surfaced, and Pd reading consistently across four confirmatory runs. reached(d) is confirmed directly, at 99.88% busy, via Level Zero Sysman engine activity, once VTune’s own hardwaremetric collection proved blocked by permissions on this shared allocation – the same paper-named tool pairing (Section V) reaching the same conclusion by whichever of its two named routes an allocation actually allows. Xd = 0 closes by a stated construction argument, the static-disassembly route MI100 used being unavailable here for want of an installed compiler backend, not for want of trying. Rd ≈ 0.004 closes the fifth field, located by the same bracket-and-locate sweep used on the other three GPUs and, at 0.4%, the smallest memoryreservation overhead measured on any device in this paper’s ensemble – full five-field measurement now closes on all five devices, the real work named precisely above, not assumed.

[3] Tianqi Chen, Thierry Moreau, Ziheng Jiang, Lianmin Zheng, Eddie Yan, Haichen Shen, Meghan Cowan, Leyuan Wang, Yuwei Hu, Luis Ceze, Carlos Guestrin, and Arvind Krishnamurthy. TVM: An automated end-to-end optimizing compiler for deep learning. In 13th USENIX Symposium on Operating Systems Design and Implementation (OSDI), 2018. arXiv preprint arXiv:1802.04799. [4] Tri Dao, Daniel Y. Fu, Stefano Ermon, Atri Rudra, and Christopher Ré. FlashAttention: Fast and memory-efficient exact attention with IO-awareness. In Advances in Neural Information Processing Systems (NeurIPS), 2022. arXiv preprint arXiv:2205.14135. [5] GNU Project. Vector extensions. Using the GNU Compiler Collection (GCC), online documentation, 2024. Documents the vector size attribute: the resulting vector width is fully specified by the programmer in bytes, not subject to the compiler’s own vectorization cost-model heuristics – a deterministic directive rather than a heuristic hint. [6] Intel Corporation. Diagnostic 15331: Using FP model: Precise prevents vectorization. Intel Software Developer Zone, 2024. Documents that -fp-model=precise (strict IEEE floating-point conformance) can itself disable vectorization the compiler would otherwise perform under -fpmodel=fast, coupling a numerics flag to a vectorization outcome. [7] Intel Corporation. qopt-prefetch, qopt-prefetch. Intel oneAPI DPC++/C++ Compiler Developer Guide and Reference, version 20242, 2024. Host-code prefetch-insertion flag, distinct from the GPU-side LSC LDCC mechanism; related option -qopt-prefetch-distance controls how far ahead compiler-generated prefetches are issued. [8] Intel Corporation. Intel Xe GPU architecture. Intel oneAPI GPU Optimization Guide, version 2025-0, 2025. States the Xe-core vector engine is 512 bits wide, 16 fp32 lanes per engine with fused FMA, 128 fp64 operations/cycle across 8 vector engines per Xe-core. [9] Intel Corporation. ipo, Qipo. Intel Fortran Compiler Developer Guide and Reference, version 2025-0, 2025. Interprocedural (whole-program) optimization flag; enables cross-function inlining and specialization decisions beyond a single translation unit. [10] Intel Corporation. Prefetch. Intel oneAPI GPU Optimization Guide, version 2025-0, 2025. Documents the LSC LDCC cache-control enum (L1UC L3C, L1C L3C, L1S L3C) for Intel Data Center GPU Max Series prefetch, chosen in source code rather than via a compile flag. [11] LLVM Project. D10617: Enable fma and f16c features when avx512f is enabled. LLVM Phabricator code review, 2015. Confirms that enabling AVX-512 (-march/-x ISA targeting) silently enables fusedmultiply-add contraction by default, coupling an instruction-set-width flag to a floating-point reassociation side effect. [12] LLVM Project. Auto-vectorization in LLVM. LLVM Documentation, 2024. States that LLVM’s loop and SLP vectorizers use an internal cost model to decide whether and how to vectorize, and may decline even when a width hint or pragma requests it – confirming that hint-based vectorization flags are heuristic, not deterministic, controls. [13] Lenore M. Mullin. Mathematics of Arrays. PhD thesis, Syracuse University, 1988. [14] Lenore M. Mullin. Scalable, Portable, Verifiable, Predictable: What AI Systems Are Missing, and What a Formal Array Calculus Already Delivers. Submitted to IPDPS 2027; source of the portable OpenACCtiled attention kernel used as the vendor-neutral reference realization in the H100 and V100 real-hardware validation sections of this paper, 2027. [15] Lenore M. Mullin and Gaétan Hains. One derivation, any machine: Deriving a weighted, communication-free partition for attention across a heterogeneous, multi-institution device ensemble. This project’s own companion paper; status (preprint venue, posting date) as of this writing to be confirmed before this paper is finalized. Names the compilerdirectives field Cd this paper formulates, as future work (Section X). [16] Open MPI Project. mpirun / mpiexec. Open MPI Documentation, Version 5.0.x, man-openmpi(1), 2024. Documents --bind-to core and --map-by core, which bind/map processes to cores “in a roundrobin manner” – a fixed algorithm computable in advance from process and core counts – and rankfile, which assigns each rank to an explicitly named core. [17] Open MPI Project. Tuning collectives. Open MPI Documentation, Version 5.0.4, 2024. Documents Open MPI’s default “fixed decision” mode – “a decision tree, essentially a large set of nested if-then-elseif blocks with baked in comm and message size thresholds derived by measuring performance on existing clusters” – and states that explicitly setting coll_tuned_<collective>_algorithm to a specific number means “the fixed decision rules associated with that collective

ACKNOWLEDGMENTS This work reasons primarily from real hardware results already reported in [15], and additionally reports five new real-hardware experiments of its own, on NCSA Delta’s A100 (Section VIII), AMD MI100 (Section IX), PSC Bridges2’s H100 (Section X), SDSC Expanse’s V100 (Section XI), and TACC Stampede3’s Intel Max 1550 (Section XII). This work used NCSA Delta GPU, NCSA Delta CPU, TACC Stampede3, SDSC Expanse GPU, and PSC Bridges-2 GPU through allocation CIS261342 from the Advanced Cyberinfrastructure Coordination Ecosystem: Services & Support (ACCESS) program, which is supported by U.S. National Science Foundation grants #2138259, #2138286, #2138307, #2137603, and #2138296, and Chameleon Cloud project CHI261724. This manuscript’s text was developed with the assistance of Claude (Anthropic), an AI language model, used throughout for drafting, editing, and technical fact-checking against source material. R EFERENCES [1] Advanced Micro Devices, Inc. AMD matrix cores. AMD Lab Notes, GPUOpen, 2024. Describes CDNA’s 64-lane wavefront and MFMA matrix-core instruction tile shapes on gfx908 (MI100). [2] Advanced Micro Devices, Inc. ROCmCC compiler reference guide. ROCm Documentation, version 6.1.1, 2024. Documents -mllvm enable-x86-prefetching (software prefetch insertion), -mllvm -inlinerecursion=[1-4] and -finline-aggressive (inlining depth), and -mllvm enable-epilog-vectorization with -min-width-epilog-vectorization (vector width) for amdclang/amdclang++.

22

are short circuited,” making the algorithm choice deterministic and known in advance. [18] OpenACC-Standard.org. The OpenACC application programming interface, version 2.7. OpenACC Specification, 2018. Documents num_gangs/num_workers/vector_length clauses on both the parallel and kernels constructs; for kernels states that the compiler “will split the code... into a sequence of accelerator kernels” whose “number and configuration of gangs of workers and vector length may be different for each kernel” – so a single clause value on kernels does not fix one configuration for the region the way it does under parallel. [19] OpenAI. triton.config. Triton Documentation, 2026. Documents num warps (threads per kernel instance, 32 threads/warp), num stages (software-pipelining depth, the mechanism standing in for prefetch), num ctas (SM90+ block clustering), and maxnreg (per-thread register cap). [20] OpenMP Architecture Review Board. SIMD directives (simd construct). OpenMP Application Programming Interface Specification, Version 5.0, 2018. States that under the simd construct “the number of iterations that are executed concurrently at any given time is implementation defined,” and that simdlen(n) sets only “the preferred number of iterations” – the construct asserts loop-dependence safety but does not fix the actual vector width or whether vectorization occurs, the same implementationdefined outcome already excluded for auto-vectorization enable flags. [21] OpenMP Architecture Review Board. Worksharing-loop construct (schedule clause). OpenMP Application Programming Interface Specification, Version 5.0, 2018. States that schedule(static[,chunk_size]) assigns chunks to threads “in a round-robin fashion in the order of the thread number,” a fixed assignment computable in advance, while dynamic and guided distribute chunks to threads “as they become available” and auto delegates the mapping decision to the compiler/runtime entirely – the identical heuristic-versus-directive split already found in compiler auto-vectorization. [22] Jonathan Ragan-Kelley, Connelly Barnes, Andrew Adams, Sylvain Paris, Frédo Durand, and Saman Amarasinghe. Halide: A language and compiler for optimizing parallelism, locality, and recomputation in image processing pipelines. In Proceedings of the 34th ACM SIGPLAN Conference on Programming Language Design and Implementation (PLDI), 2013. [23] solidpixel. Creating invariant floating-point accumulators. Technical note, 2023. Shows explicitly that vector-width execution reassociates a reduction’s summation order relative to scalar execution (e.g. (((0+4)+(2+6))+((1+5)+(3+7))) versus ((((0+1)+2)+3)+4)), and that loop unrolling produces the identical class of reassociation, changing rounding error even though the computation is mathematically identical. [24] Samuel Williams, Andrew Waterman, and David Patterson. Roofline: An insightful visual performance model for multicore architectures. Communications of the ACM, 52(4):65–76, 2009. [25] Lianmin Zheng, Chengfan Jia, Minmin Sun, Zhao Wu, Cody Hao Yu, Ameer Haj-Ali, Yida Wang, Jun Yang, Danyang Zhuo, Koushik Sen, Joseph E. Gonzalez, and Ion Stoica. Ansor: Generating highperformance tensor programs for deep learning. In 14th USENIX Symposium on Operating Systems Design and Implementation (OSDI), 2020. arXiv preprint arXiv:2006.06762.

23

Related documents

• Closing the Prediction Gap: Completing Machine Shape So That Predicted Time, Power, Energy, and Mapping Match What Real Hardware Does#1187875
• Adaptation Fidelity of SPEC CPU2026#781007
• Adaptation Fidelity of SPEC CPU2026#779632
• DeepFedNAS: Efficient Hardware-Aware Architecture Adaptation for Heterogeneous IoT Federations via Pareto-Guided Supernet Training#1104943
• Ensemble Complexity in Photovoltaic Forecasting#1084941