Proposed Micro-Optimizations — wh-iron perf kernels

Generated: 2026-05-23
Branch: ek/t4-microopts
Base: feat/int-quant-perf

These proposals were researched during the T4 micro-optimization pass. Each entry explains why the pattern cannot be applied cleanly today (missing DSL primitives, codegen changes required, or benefit below threshold) and what would need to change before it could land.

---

Pattern 1 — Vectorized X loads (float4 / half4) in GEMV/GEMM hot loops

Status: Already implemented — no action required.

Targets: iron_qmv, iron_qmm{,_bm2,_bm4}, iron_qmv_int8_fast, iron_qmm_int8_fast{,_bm2,_bm4}, iron_qmm_mma, iron_qmm_mma_m16, iron_qmm_mma_int8, iron_qmm_mma_m16_int8.

Finding: The VectorizePass in crates/wh-iron-codegen/src/passes/vectorize.rs already detects consecutive Load ops at contiguous indices and promotes them to Op::VectorLoad with width up to 4 (MAX_VEC_RUN = 4). The MSL emitter (crates/wh-iron-codegen/src/msl/emit_block.rs) then lowers VectorLoad{len:4, dtype:F32} → float4, {len:4, dtype:F16} → half4, {len:4, dtype:BF16} → bfloat4.

All X-load sites in the GEMV/GEMM kernels are structured as consecutive scalar Load calls at base + 0 .. base + N — explicitly sequenced to be contiguous in the IR so the vectorize pass fuses them. The kernel comments confirm this (e.g. "16 X loads — consecutive in IR for vectorize fusion (4× float4)" in iron_qmv; "8 contiguous device loads → 2× vec4 after vectorize pass" in iron_qmm_mma).

Verification: Confirmed by inspecting the VectorizePass source, the MSL emitter, and the kernel DSL code — no gap.

---

Pattern 2 — simd_broadcast for scale/bias loads shared across lanes

Status: Proposal only — hardware already coalesces same-address loads; benefit unconfirmed; implementation adds complexity without clear profiling evidence.

Targets: iron_qmv, iron_qmm{,_bm2,_bm4}, iron_qmv_int8_fast, iron_qmm_int8_fast{,_bm2,_bm4}.

Background: The DSL exposes simd_broadcast(value, lane) → simd_broadcast(v, lane) in Metal Shading Language (see body_parser.rs:1315, ir.rs:986, emit_block.rs:1065). The primitive exists and is used in the AURA codebook kernels.

Where lanes share scale/bias addresses:

int4 kernels (iron_qmv, iron_qmm*): K-block = 512 elements. Each lane owns 16 X values; group_size = 64. Lanes 0–3 share group g = 0 for the first K-block, lanes 4–7 share group g = 1, etc. Every 4 consecutive lanes read the same scales[sb_base + g] and biases[sb_base + g] — 4 redundant device loads per 4-lane group, ×4 output rows = 32 redundant loads per K-block outer-iteration.

int8 kernels (iron_qmv_int8_fast, iron_qmm_int8_fast*): K-block = 128. Each lane owns 4 X values; group_size = 64. Lanes 0–15 share group g = _b / 64, lanes 16–31 share group g + 1. 16 redundant loads per K-block per output row.

What the implementation would look like:

// Determine the representative lane for this group (lowest lane in group).
// For int4 (16 X/lane, gs=64): base_lane = (lane / 4) * 4 = lane & !3
let base_lane = lane & !3u32;           // int4 groups: 4 lanes/group
// Only the base lane does the device load; others present 0.0.
let s0_raw = if lane == base_lane { load(scales[sb_base0 + g]).cast::<f32>() } else { 0.0f32 };
let s0 = simd_broadcast(s0_raw, base_lane);
let bi0_raw = if lane == base_lane { load(biases[sb_base0 + g]).cast::<f32>() } else { 0.0f32 };
let bi0 = simd_broadcast(bi0_raw, base_lane);

Why not applied now:

1. Apple GPU cache broadcast semantics: Apple Silicon GPUs (M1–M4) handle same-address loads from multiple threads in a simdgroup as a single cache-line fetch with broadcast. The 4-lane (int4) or 16-lane (int8) same-address loads already cost ≈1 effective device access. The simd_broadcast reformulation saves instruction count (fewer load instructions in the shader) but does not reduce L2 traffic — only a profiler can show whether the instruction pressure is the actual bottleneck on these kernels.

2. Implementation complexity: The if lane == base_lane … else 0.0f32 conditional introduces a divergent branch + the select pattern; the existing DSL if lowers to an if_conversion pass that can generate predicated instructions, but the branch itself adds IR nodes and may perturb the schedule pass for the load/compute interleaving.

3. base_lane is dynamic: simd_broadcast(val, lane) where lane is a non-constant variable requires Metal 2.1+ which all M-series hardware satisfies — but the pass must handle the dynamic-lane argument in the broadcast cost model (currently the VectorizePass and SchedulePass have no knowledge of simd_broadcast dependencies).

Prerequisite work: Profile iron_qmv_int8_fast with Metal GPU counter LOAD_CACHE_MISS_RATE to confirm same-address loads are not already coalesced by hardware. If they are not (i.e. older non-M-series hardware or future non-Apple targets), add a FastPath flag to the BenchSpec that enables the broadcast reformulation via a kernel-level constexpr toggle rather than restructuring all kernels.

---

Pattern 3 — fast:: math intrinsics in audio + numeric paths

Status: Proposal only — fast_exp, fast_log, fast_sin, fast_cos are not exposed by the DSL; adding them requires new UnaryOpKind variants + codegen + precision characterization.

Targets: mel_spectrogram (sin/cos inner DFT, log filterbank), iron_softmax (exp in both passes), iron_logsumexp (exp in loop, log at store), vocoder_istft (sin/cos inner iDFT).

Current emission: UnaryOpKind::{Exp,Log,Sin,Cos} in crates/wh-iron-core/src/ir.rs emit exp(arg), log(arg), sin(arg), cos(arg) — IEEE-754 precise Metal built-ins. No fast::exp, fast::log, fast::sin, fast::cos variants exist anywhere in the codegen stack.

What Metal’s fast math provides: Metal’s fast namespace (#include <metal_math>) provides fast::exp, fast::log, fast::sin, fast::cos, fast::sqrt — polynomial approximations that are:

• 1–3 ULP accurate (vs. ≤ 0.5 ULP for the IEEE versions)
• Approximately 1.5–2× faster in throughput-limited contexts
• Unsafe for edge cases: fast::log(-1.0) is undefined, fast::exp may not flush denormals

For mel_spectrogram: The direct DFT inner loop computes cos(angle) * xw and sin(angle) * xw for each of the n_fft ≈ 400 samples per frequency bin. The angle range is [-2π, 0]; inputs are well-defined reals — fast math edge cases don’t apply. Expected speedup: ~1.5–1.8× on the DFT inner loop (the dominant cost).

For iron_softmax: exp(v_i - max_i) — arguments are always ≤ 0 (subtracted maximum), so the output is always in (0, 1]. Fast exp is safe here. Expected speedup: ~1.4× on the reduction pass.

For iron_logsumexp: Same exp path as softmax (arguments ≤ 0), plus a final log(gs) where gs > 0 is always true. Safe.

For vocoder_istft: cos(angle) and sin(angle) inner loop — same argument range analysis as mel_spectrogram. Safe.

What needs to change:

1. New UnaryOpKind variants in wh-iron-core/src/ir.rs:

   FastExp,
   FastLog,
   FastSin,
   FastCos,
   FastSqrt,

   with msl_emit returning fast::exp(arg) etc.

2. New DSL keywords in wh-iron-macros/src/body_parser.rs:

   "fast_exp"  => quote! { UnaryOpKind::FastExp },
   "fast_log"  => quote! { UnaryOpKind::FastLog },
   "fast_sin"  => quote! { UnaryOpKind::FastSin },
   "fast_cos"  => quote! { UnaryOpKind::FastCos },
   "fast_sqrt" => quote! { UnaryOpKind::FastSqrt },

3. Precision validation: The mel_spectrogram_gpu_correctness test tolerates 1e-3; the audio end-to-end (waveform cosine similarity) is more forgiving. The softmax / logsumexp tests use 1e-4. A numerical sweep of fast::exp error against the IEEE path for the actual argument ranges would be needed before landing.

4. Kernel edits: Replace sin/cos/exp/log with fast_sin / fast_cos / fast_exp / fast_log at the relevant call sites in:

   • crates/wh-iron-std/src/iron/mel_spectrogram.rs
   • crates/wh-iron-std/src/mlx/softmax.rs
   • crates/wh-iron-std/src/mlx/logsumexp.rs
   • crates/wh-iron-std/src/iron/vocoder.rs

Risk: If the fast-math approximation error accumulates over many inner-loop iterations (e.g. 400 DFT taps), the cosine similarity of the Mel spectrogram output may drop below 0.999. The existing tests would catch this. A safe rollout would compare fast vs IEEE outputs on a representative audio clip before committing.

---

Pattern 4 — f16/bf16 accumulator for small-K shapes

Status: Not applicable — all production K shapes in the target kernels accumulate across ≥ 64 elements under quantization noise; switching to T-precision accumulators would degrade accuracy below the 0.999 cosine threshold.

Targets examined: iron_qmv, iron_qmm{,_bm2,_bm4}, iron_qmv_int8_fast, iron_qmm_int8_fast{,_bm2,_bm4}, mel_spectrogram, iron_softmax, iron_logsumexp, vocoder_istft.

Analysis:

Quantized GEMV/GEMM kernels: Every accumulator is f32 by design. The comment in iron_qmv is explicit: “accumulators stay in f32 regardless of T”. The reason is that int4 quantization introduces up to q_err = 15/(2^4) × scale = 0.93 × scale per element; with 64 elements per group and scale ≈ 1/16, accumulated error without f32 headroom exceeds the cosine similarity gate. For int8 the dynamic range is similar: 255/256 × scale per element, same argument applies.

mel_spectrogram: The inner DFT accumulates re += xw * cos(angle) over 400 samples. Running sum of floats at f16 precision would accumulate ≈ 0.001 × √400 ≈ 0.02 absolute error — audible in the output spectrogram. f32 accumulation is required.

iron_softmax, iron_logsumexp: Online softmax requires both a max accumulator and a sum accumulator. The numerical stability of the online algorithm depends on the precision of exp(v - m) where v and m may differ by many units. f16 would lose the mantissa bits needed to recover from the exp(large_neg_val) tail — catastrophic cancellation risk.

vocoder_istft: Same DFT accumulation argument as mel_spectrogram but the n_fft is smaller (Kokoro: 20) — at n_fft=20, f16 might be borderline, but the existing f32 path is not a bottleneck and the test tolerance is 1e-3.

Conclusion: fp32 accumulators are correctness-critical for all identified kernel shapes. The pattern is safe only if a kernel has both: (a) K ≤ 16 per lane per accumulation, and (b) no quantization noise. None of the production GEMV/GEMM targets satisfy (a); the audio paths violate (b).

---

Pattern 5 — K-loop software pipelining (prefetch-overlap)

Status: Proposal only — requires codegen support that does not exist; implementing it without codegen changes would require hand-unrolling the K-loop which defeats DSL portability.

Targets: iron_qmm_mma, iron_qmm_mma_m16, iron_qmm_mma_int8, iron_qmm_mma_m16_int8, iron_qmm_mma_mpp (via quantized_mpp.rs), iron_qmm_nax, iron_qmm_nax_int8, iron_moe_gather_qmm_mma_int4, iron_moe_gather_qmm_mma_int8 and the _bm{8,16,64}_mpp MoE variants.

Pattern: Apple’s hand-tuned mma_qmm Metal shaders (and typical CUTLASS-style GPU kernels) overlap loading the next K-block into shared/threadgroup memory while the MMA units process the previous K-block:

// Pipelined form (2-stage):
load_to_tg(K=0)
threadgroup_barrier()
for k in 1..K_blocks:
    async_issue: load_to_tg(K=k)   // start next load
    mma(K=k-1)                      // compute on current
    threadgroup_barrier()            // wait for next load
mma(K=K_blocks-1)

vs. our current sequential form:

for k in 0..K_blocks:
    load_to_tg(K=k)
    threadgroup_barrier()
    mma(K=k)
    threadgroup_barrier()

Why not done today: The DSL’s for k in range(...) block compiles into a monolithic loop in the IR; the codegen schedule pass (schedule.rs) is unaware of the “hoist-load-before-barrier” opportunity. Expressing software pipelining requires either:

1. A prefetch_async IR op — a new Op::PrefetchAsync { src, tg_dst, ... } that signals the schedule pass to emit the device load before the simdgroup barrier of the preceding block. The schedule pass would need a new cost model entry for memory-latency hiding.

2. A two-buffer (ping-pong) DSL primitive — expose a double_buffer wrapper in the body_parser that allocates two threadgroup arrays (Xs_a, Xs_b, Ws_a, Ws_b) and auto-alternates load / mma assignments, emitting a two-trip peeled loop. This is the ldmatrix / wmma pipelining pattern used by CUTLASS.

3. Explicit loop-peeling at the kernel DSL level — hand-peel the loop prologue and epilogue in the kernel source, using two named threadgroup buffers alternately. This gives correct pipelining without codegen changes but roughly doubles the kernel source length and is fragile across K shapes.

Expected benefit: On M3 Pro with iron_qmm_mma (K=4096, N=4096, M=32), an Apple hand-tuned pipelined MMA kernel typically runs 15–25% faster than the non-pipelined equivalent (Apple’s own steel_gemm docs reference this range for their pipelined GEMM over the non-pipelined baseline). The gain is larger for larger K (longer memory latency → more hiding opportunity).

Recommended approach: Implement Op::PrefetchAsync as a new codegen pass (prefetch.rs) that runs after schedule.rs and licm.rs, identifies the pattern [ThreadgroupStore* … ThreadgroupBarrier … MMA*] within a loop body, and hoists the store block one barrier ahead. This is a well-understood transformation; the risk is incorrect barrier placement (would produce wrong results, caught by existing MMA correctness tests). Estimated scope: 2–3 days of codegen work + 1 day of correctness testing.

---

Summary table

Pattern	Status	Blocker
1. float4/half4 X loads	Already done (VectorizePass)	—
2. simd_broadcast for scale/bias	Proposal only	Hardware coalesces same-address loads; no profiling evidence of bottleneck; implementation complexity
3. fast:: math in audio/softmax	Proposal only	Missing FastExp/FastLog/FastSin/FastCos UnaryOpKind variants + precision validation
4. f16/bf16 accumulator small-K	Not applicable	All targets have K ≥ 64 with quantization noise; f32 required for correctness
5. K-loop software pipelining	Proposal only	Needs Op::PrefetchAsync codegen pass or DSL double-buffer primitive

---

CUDA-specific extensions for exact decode

The Metal observations above transfer to CUDA only when the memory hierarchy, warp execution model, and numerical contract are treated explicitly. For checkpoint-native decode, a candidate may change load width, scheduling, ownership, or reuse, but it must not change stored weight values, conversion instructions, FMA order, accumulator precision, or the reduction tree unless it enters a separately qualified numerical policy.

The primary CUDA workloads are materially different:

• scalar autoregressive decode (M = 1) streams each projection weight once and is primarily a resident-weight bandwidth problem;
• block verification (M = 8 or M = 10) can reuse each weight tile across several activation rows and therefore has a stronger compute and software-pipeline opportunity;
• prefill is a conventional larger-M matrix workload and should retain its own tactics rather than inheriting a decode schedule.

CUDA Pattern 1 — direct-load register pipeline for scalar GEMV

Status: Proposal only — exact in principle; keep only if profiling shows memory-dependency stalls and a production-shape benchmark wins.

Target: checkpoint FP8-weight/F16-activation projection kernels at M = 1.

The scalar kernel should continue to load weights directly from global memory. Staging a tile in shared memory is normally a poor trade when that tile is consumed only once: it adds global-to-shared traffic, a shared-memory read, barriers, and an occupancy constraint without creating reuse.

The useful pipeline is instead a register-level ping-pong loop:

load weight and activation tile 0
for tile in 1..tile_count:
    issue direct loads for tile
    convert and accumulate the preceding tile
    rotate the two register slots
convert and accumulate the final tile

The generated CUDA should use aligned 16-byte transactions where alignment is proven. Two-stage and four-stage variants must be compared; deeper prefetch is not automatically better because additional live registers can reduce active warps or introduce spills. Inspect SASS before claiming overlap: source-level reordering is not evidence that the compiler retained the intended schedule.

Exactness contract:

• load the original checkpoint bytes;
• use the retained FP8 conversion instruction;
• preserve each output row’s FMA sequence;
• preserve the warp reduction tree;
• require raw-bit output equality, not a tolerance.

Kill gates: reject on any spill, reduced occupancy without a compensating kernel win, changed output bits, or less than a 5% complete-projection win on a hot production shape. A model-level candidate should additionally survive same-binary reversed-order timing.

CUDA Pattern 2 — reuse each weight stream across verifier rows

Status: High-value proposal for native MTP and block-draft target verification.

At verification width 8 or 10, every activation row uses the same projection weights. Loading and decoding a weight tile independently for every row wastes the most expensive stream. The CUDA kernel should load a weight vector once and apply it to a bounded group of activation rows before advancing K.

The search space should include two, four, and eight rows per weight stream. Ten-row verification should be expressed as explicit admitted row groups, not as a generic padded fallback. A full eight- or ten-row accumulator set can create a register cliff, so the widest row tile is not presumed fastest.

For these shapes, a shared-memory or asynchronous-copy pipeline can be useful because the staged weight tile is reused. Candidate stages are:

1. issue the next aligned weight tile into the inactive stage;
2. consume the current tile across the admitted row group;
3. wait only at the ownership boundary;
4. swap stages and continue.

The target verifier retains the checkpoint-native arithmetic contract. A draft-only kernel may use a different, explicitly named numerical contract, but it must be judged by acceptance-adjusted end-to-end throughput rather than isolated kernel time.

Correctness gates: raw-bit component equality, every proposal mismatch depth, full accept, rejected-tail poisoning, and equality of the next target logit plus committed KV and recurrent state after every possible commit depth.

CUDA Pattern 3 — lower vector IR operations explicitly

Status: Codegen gap — the backend-neutral vectorization pass creates VectorLoad, VectorStore, and VectorExtract operations, while the CUDA emitter does not yet provide the corresponding native lowering.

The CUDA emitter should lower proven-aligned vector operations to native raw transactions such as uint4, float4, or paired 64-bit loads, followed by explicit lane extraction. The scalar tail must remain available for unaligned or partial tiles. Alignment belongs in the tactic key or dispatch precondition; it must not be assumed from a pointer’s element type.

This work makes the existing vectorization pass useful to generated CUDA kernels. It is not, by itself, a new optimization for handwritten kernels that already use 16-byte loads.

Required tests:

• aligned and deliberately misaligned bases;
• all supported element widths and signedness;
• partial tails;
• generated-source assertions for the vector transaction;
• bitwise equivalence to scalar lowering.

CUDA Pattern 4 — warp-owned scale and bias metadata

Status: Profile first — use only where redundant load instructions, not merely repeated source expressions, are measurable.

The CUDA counterpart to simd_broadcast is a warp shuffle. One lane can load scale or bias metadata and distribute it with __shfl_sync. This is useful only when several active lanes consume the same value and the hardware/compiler does not already fold or broadcast the access effectively.

Avoid a divergent load branch. The owning lane should perform the load under a uniform warp protocol, followed by a shuffle executed by every active lane. Per-output scales used only by the lane that stores the final reduction do not benefit from an earlier broadcast.

Admission requires a profiler-confirmed reduction in executed load instructions or dependency stalls, unchanged register tier, and an end-to-end kernel win. A cache-line-traffic argument alone is insufficient.

CUDA Pattern 5 — fixed-tile lossless FP8 streaming

Status: Research proposal — potentially high bandwidth value, but a previous variable-length design is not a suitable CUDA decode format.

This is the exact alternative to converting checkpoint FP8 weights into a lower-precision target format. Store a fixed-size encoded tile, reconstruct the original FP8 bytes in registers, and feed those bytes into the retained conversion, FMA, and reduction sequence. Because reconstructed bytes must be identical, the model’s numerical contract does not change.

The format must provide:

• constant tile addresses;
• aligned fixed-width transactions;
• no escape stream or per-tile pointer chase;
• no data-dependent branch in the hot decode loop;
• exact reconstruction of every FP8 code and tail;
• enough physical-byte reduction to repay decode instructions.

Pipeline encoded tile N + 1 while reconstructing and accumulating tile N. Prefer register reconstruction; shared-memory expansion writes the full byte stream back out and can erase the bandwidth saving.

Reject before model integration if physical bytes saved are below 12–15%, PTXAS reports spills, or rotating real projection matrices do not beat the uncompressed path. Correctness requires exhaustive byte-code fixtures, production tails, teacher-forced logits and state, and exact greedy output fingerprints.

CUDA Pattern 6 — fast math and reduced accumulators are separate policies

Status: No-go for the exact target.

Approximate transcendental functions, F16/BF16 projection accumulation, reassociated reductions, and lower-precision recurrent state can change KV, hidden-state, and recurrent trajectories even when early output tokens agree. They cannot qualify an exact projection speedup.

These techniques may be evaluated inside a drafter because a verifier can reject incorrect proposals. Their receipt must still include acceptance length, proposal cost, target forwards per output token, and final target-state equality. A faster draft kernel that lowers acceptance can reduce overall throughput.

CUDA Pattern 7 — output-owned warp execution and stage elimination

Status: Design evidence for dense decode and block verification; direct candidate for future small-batch MoE decode.

Cursor’s Warp Decode reorganizes a small-batch MoE decode layer around output values rather than expert batches. Each warp owns one output scalar, streams the required weight rows directly, accumulates in private FP32 registers, and uses warp shuffles for the final reduction. Gate and up projections share one activation stream, while routing weights are folded into the down-projection accumulator. This removes expert padding, scatter, gather, combine, shared-memory handoffs, and per-expert output buffers. The reported 1.84x improvement was measured on an internal Qwen-3-style MoE workload on B200 hardware; it is evidence for the execution pattern, not a predicted speedup for a dense model or DGX Spark.

The parts that transfer to dense CUDA decode are:

• give each warp a stable output row or scalar for its lifetime;
• stream weights directly and keep partial sums private to the warp;
• use warp-shuffle reductions instead of shared-memory reductions when the reduction fits one warp;
• consume one activation load across related gate and up work;
• fold an exact epilogue into the owning projection rather than materializing a buffer that is immediately read once;
• audit every projection boundary for padding, layout conversion, staging, splitting, and combine passes that exist only to serve the next kernel.

This is primarily an audit checklist for scalar dense decode. A dense model has no expert packing, scatter, or expert-output combine path, so most of the published stage-elimination gain is unavailable. A kernel that already owns one output row per warp, uses direct vectorized loads, accumulates privately, and reduces with shuffles has already adopted the central mechanism.

Block verification at M = 8 or M = 10 has a stronger transferable form. Keep output ownership independent, but reuse a loaded weight tile across the verification rows and retain a private accumulator per admitted row group. Fuse adjacent gate/up consumption and exact epilogues only when doing so preserves the target reduction order. This combines the ownership rule here with the weight-reuse search in CUDA Pattern 2 without introducing expert-specific data structures.

For a future MoE target, admit a separate small-batch tactic that reads the routed expert identifiers directly and accumulates their weighted contributions without first forming expert-major batches. Retain an expert-centric tactic for prefill and larger batches, where packing overhead can be amortized. Admission must be based on batch size, expert count, active experts, row shape, and hardware rather than a model-wide switch.

The numerical result requires care. The published path removes an intermediate BF16-to-MXFP8 activation quantization and therefore measures closer to an FP32 reference, but its reported correctness gate is tolerance based rather than bitwise. Better similarity to FP32 does not establish identical logits, tokens, KV state, or recurrent state. The exact target lane must preserve checkpoint bytes, conversion behavior, accumulation order, and the reduction tree. Any implementation that changes those properties belongs to a separately named numerical policy and needs long-context drift testing.

Measurement gates:

• profile the removed stages and bytes before implementation;
• benchmark scalar decode and verifier widths separately;
• require raw-bit component equality for the exact lane;
• compare per-layer logits and recurrent state, not only final token text;
• run long teacher-forced and greedy sequences to expose accumulated drift;
• record registers, spills, active blocks, memory stalls, and achieved bandwidth;
• keep a dense candidate only when complete model decode improves, and keep a MoE candidate only when it beats the expert-centric path in its admitted small-batch range.

CUDA tactic selection and measurement

Do not choose one CUDA geometry globally. A tactic registry should key exact admission on at least:

(compute capability,
 weight format,
 activation format,
 batch rows,
 output features,
 input features,
 epilogue,
 alignment,
 numerical contract)

Each admitted tactic records its required CTA size, row tile, vector width, pipeline depth, register limit, and cache policy. Unsupported shapes retain the existing implementation.

Every CUDA experiment should report:

• complete projection time, not only the inner loop;
• achieved memory bandwidth and transaction efficiency;
• long-scoreboard and memory-dependency stalls;
• registers per thread, local-memory bytes, spills, and active blocks;
• raw-bit output comparison;
• rotating real-weight and activation corpora;
• warm medians with dispersion and reversed candidate/control ordering.

The immediate priority is scalar register pipelining and production-shape tactic selection for native FP8 projections. Weight-reuse pipelines become the priority for width-8 and width-10 verification. Fixed-tile lossless FP8 streaming should remain a fail-fast microbenchmark until it proves that byte savings exceed reconstruction cost.