feat(fp8): §19 MEASURED on gb10 sm_121 — fp8 GEMM RUNS; TE per-tensor e4m3 = real 1.6-1.8x tensor-core win; in-house cuda twin correct but unoptimized

Closes the §18 "fp8 UNRUNNABLE on sm_121" gap with MEASURED numbers at the four
prod bs4 transformer-block GEMM shapes (M=16384). Two distinct fp8 paths now run:

R1 TE per-tensor E4M3 (DelayedScaling -> cuBLASLt fp8 MMA): the REAL fp8
tensor-core win — 1.57-1.84x over bf16 MEASURED (52.1-61.8 fp8 TFLOPs vs
33.3-34.3 bf16), rel-err 0.0375 honest over ALL elements (<=0.10 ceiling). Only
MXFP8 block-scaling stays upstream-blocked ("not supported on 12.0+"); per-tensor
works once libnvrtc-builtins.so.13.3 is on the loader path.

R2 in-house dequant->half->T.gemm CUDA fragment-C twin
(fp8_scaled_matmul_path_c_cuda_prim): COMPILES + RUNS + parity-passes (0.0375,
native fp8_e4_t->half_t decode verified in emitted CUDA) — but 0.18-0.42x as-is:
its MMA runs on the decoded fp16 values (mma_sync<kFloat16,...>, not an fp8-input
MMA) on an untuned 32x64x32 single-128-thread tile. It is the proven MEMORY lever
(2.0x operand-byte halving MEASURED) + the on-CUDA optimization target, NOT a
speed lever as-is. Attribution explicit: throughput belongs to TE-cuBLASLt fp8.

Two port fixes were needed for R2 to run on CUDA (RULE #1, no fallback):
- fp8_matmul_path_c.py: the Metal T.alloc_var("float32") decode scratch lowered
  to a non-assignable float[1] on codegen_cuda ("expression must be a modifiable
  lvalue"). Replaced with a register-fragment stage (T.copy e4m3 global->fragment
  with disable_tma=True, then per-element T.cast->fp16) — keeps the native
  per-element decode, emits a valid CUDA lvalue (F2-forward CUDA-twin pattern).
- fp8_gemm_microbench.py: the prim's operand param is typed float8_e4m3; feeding
  a .view(torch.uint8) tensor segfaults the tvm_ffi adapter on the dtype
  mismatch. Fixed _quantize_pertensor_fp8 to pass native torch.float8_e4m3fn
  (the 1-byte storage IS the 2.0x halving regardless of the dtype tag).

cppmega_mlx/_gb10_nvrtc_env.py (NEW, stdlib-only): self-healing NVRTC loader-path
guard — prepends /usr/local/cuda-13.3/.../lib and re-execs once so TE's NVRTC can
dlopen libnvrtc-builtins.so.13.3 (the R1 tensorwise route's only loader blocker);
no-op on non-gb10. Validated end-to-end on gb10 with env -u LD_LIBRARY_PATH.

docs: RELAX-GRAPH-VS-MEGATRON.md §19 (full table + GO verdict + e2e EXTRAPOLATION
+ corrected memory synergy bs4-fp8 ~= 2x bs1-fp16); FP8-ACTIVATIONS-PATHC.md
TL;DR/status updated with the MEASURED fp8 GEMM numbers.

GO: a REAL fp8 e4m3 tensor-core GEMM runs on sm_121 with a measured ~1.8x
speedup and honest 0.0375 parity. Measured 2026-06-04 gb10 tvm.cuda(0).

Author: apstenku123

cppmega_mlx/_gb10_nvrtc_env.py

@@ -0,0 +1,110 @@
+"""gb10 NVRTC loader-path guard — persist the CUDA 13.3 builtins dir for EVERY run.
+
+WHY (the measured gb10 sm_121 bug, 2026-06-04). TransformerEngine 2.16 (source
+build at /home/dave/TransformerEngine) links ``libnvrtc.so.13`` built against
+CUDA 13.3. At JIT-compile time that NVRTC dlopen's a companion
+``libnvrtc-builtins.so.<MAJOR>.<MINOR>`` whose MINOR must match EXACTLY — i.e.
+``libnvrtc-builtins.so.13.3``. That file exists
+(``/usr/local/cuda-13.3/targets/sbsa-linux/lib/libnvrtc-builtins.so.13.3``) but
+its directory is not on the login-shell ``LD_LIBRARY_PATH`` (empty) and was
+historically absent from the ldconfig cache, so the TE tensorwise (Float8
+delayed-scaling) route failed with::
+
+    failed to open libnvrtc-builtins.so.13.3
+
+This is a TOOLCHAIN/loader-path problem, NOT a code bug — fixed purely by making
+that directory findable. The system-level real fix is a one-time
+``sudo -n ldconfig`` (the cache now resolves 13.3). This module is the COMMITTED,
+self-healing, per-run belt-and-suspenders that makes the fix survive a future
+ldconfig regression WITHOUT relying on a transient one-shell ``export`` (RULE #1:
+no silent dependence on a shell var that a future operator might forget).
+
+HOW. ``LD_LIBRARY_PATH`` is parsed by glibc ONCE at process startup; a late
+``os.environ`` mutation is NOT honored by later ``dlopen`` calls for the dynamic
+search path. So this helper, when the 13.3 dir is missing from the *startup*
+``LD_LIBRARY_PATH``, prepends it and RE-EXECS the interpreter so the loader picks
+it up at startup. The re-exec is idempotent (guarded by an env sentinel) and runs
+at most once. RULE #1: if the dir does not exist or the re-exec cannot be
+performed, this RAISES with the precise where+what — it never silently proceeds
+with a half-applied path that would later fail deep inside an NVRTC dlopen.
+
+NON-gb10 hosts (no such directory) are a no-op: the helper only acts when the
+13.3 builtins dir is actually present, so importing it on a Mac/M4 is harmless.
+
+This module is deliberately a TOP-LEVEL ``cppmega_mlx`` module (NOT under
+``cppmega_mlx.runtime``) and imports ONLY the stdlib, so importing it does not
+pull torch / mlx / TransformerEngine — the re-exec MUST happen before any of
+those build their loader link map.
+"""
+
+from __future__ import annotations
+
+import os
+import sys
+
+# The directory that holds libnvrtc-builtins.so.13.3 on gb10 (CUDA 13.3 toolkit,
+# sbsa-linux / aarch64). This is the loader path TE's NVRTC needs.
+GB10_CUDA133_LIB = "/usr/local/cuda-13.3/targets/sbsa-linux/lib"
+
+# Sentinel so the re-exec happens at most once (guards against an exec loop).
+_REEXEC_SENTINEL = "CPPMEGA_GB10_NVRTC_PATH_APPLIED"
+
+
+def _dir_has_builtins(path: str) -> bool:
+    """True iff *path* exists and contains the libnvrtc-builtins.so.13.3 SONAME."""
+    soname = os.path.join(path, "libnvrtc-builtins.so.13.3")
+    return os.path.isdir(path) and os.path.exists(soname)
+
+
+def ensure_nvrtc_builtins_path(*, lib_dir: str = GB10_CUDA133_LIB) -> str:
+    """Guarantee *lib_dir* is on the STARTUP LD_LIBRARY_PATH for this process.
+
+    Returns one of: ``"noop-not-gb10"`` (dir absent — not this host, harmless),
+    ``"already-present"`` (startup env already had it), or it RE-EXECS the
+    interpreter (does not return in that case). RULE #1: raises RuntimeError if
+    the 13.3 builtins dir exists but the re-exec cannot be performed — never a
+    silent half-fix.
+
+    MUST be called BEFORE importing torch / transformer_engine (i.e. at the very
+    top of any gb10 run entrypoint) so the corrected env is in place when NVRTC
+    is later dlopen'd.
+    """
+    if not _dir_has_builtins(lib_dir):
+        # Not gb10 (or the 13.3 toolkit is not installed here). Nothing to do —
+        # this is the correct no-op on a Mac/M4 dev box. We do NOT fabricate the
+        # path on a host that lacks it.
+        return "noop-not-gb10"
+
+    current = os.environ.get("LD_LIBRARY_PATH", "")
+    entries = [p for p in current.split(os.pathsep) if p]
+    if lib_dir in entries:
+        # Startup env already carries it (e.g. the run command's LD_LIBRARY_PATH
+        # prefix, or a prior re-exec). The loader honors it — done.
+        return "already-present"
+
+    if os.environ.get(_REEXEC_SENTINEL) == "1":
+        # We already re-exec'd once yet the dir is still not on LD_LIBRARY_PATH.
+        # That should be impossible; fail loud rather than loop or limp on.
+        raise RuntimeError(
+            "gb10_nvrtc_env: re-exec sentinel set but "
+            f"{lib_dir!r} still absent from LD_LIBRARY_PATH={current!r}; refusing "
+            "to proceed — TE's NVRTC would later fail to open "
+            "libnvrtc-builtins.so.13.3 (RULE #1: no silent half-applied path)."
+        )
+
+    # Prepend the 13.3 dir and re-exec so the loader parses it at startup.
+    new_ldlp = os.pathsep.join([lib_dir, *entries]) if entries else lib_dir
+    new_env = dict(os.environ)
+    new_env["LD_LIBRARY_PATH"] = new_ldlp
+    new_env[_REEXEC_SENTINEL] = "1"
+    try:
+        os.execve(sys.executable, [sys.executable, *sys.argv], new_env)
+    except OSError as exc:  # pragma: no cover - exec almost never fails
+        raise RuntimeError(
+            "gb10_nvrtc_env: failed to re-exec the interpreter to apply "
+            f"LD_LIBRARY_PATH={new_ldlp!r} (needed so TE's NVRTC can dlopen "
+            f"libnvrtc-builtins.so.13.3): {exc!r}. RULE #1: not falling through "
+            "with a path that would crash inside NVRTC later — fix the run env."
+        ) from exc
+    # os.execve does not return on success.
+    raise AssertionError("unreachable: os.execve returned without raising")  # pragma: no cover

cppmega_mlx/nn/_tilelang/fp8_matmul_path_c.py

@@ -259,15 +259,35 @@ def _fp8_scaled_matmul2d_cuda_kernel_template(
         C_frag = T.alloc_fragment((_FP8_MM_BM, _FP8_MM_BN), _FP8_MM_C_DTYPE)  # type: ignore[name-defined]  # noqa: F821
         C_shared = T.alloc_shared((_FP8_MM_BM, _FP8_MM_BN), _FP8_MM_C_DTYPE, scope="shared")  # type: ignore[name-defined]  # noqa: F821
         T.clear(C_frag)
+        A_fp8_local = T.alloc_fragment((_FP8_MM_BM, _FP8_MM_BK), "float8_e4m3")  # type: ignore[name-defined]  # noqa: F821
+        B_fp8_local = T.alloc_fragment((_FP8_MM_BN, _FP8_MM_BK), "float8_e4m3")  # type: ignore[name-defined]  # noqa: F821
         for ko in T.serial(T.ceildiv(_FP8_MM_K, _FP8_MM_BK)):  # type: ignore[name-defined]  # noqa: F821
+            # Stage the e4m3 bytes into register fragments first (disable_tma on
+            # the global load), then cast element-wise. On CUDA codegen the prior
+            # ``T.alloc_var("float32")`` scratch lowered to a non-assignable
+            # ``float[1]`` array (lvalue error); the per-element fragment->fragment
+            # ``T.cast(e4m3 -> float16)`` here lowers to the NATIVE
+            # ``__nv_fp8_e4m3 -> half`` decode (codegen_cuda ``GetTileLangFP8Type``
+            # -> ``fp8_e4*_t``), NOT a vectorized ``(half4)(uchar4)`` reinterpret —
+            # the per-element scalar cast keeps the decode honest (RULE #1) while
+            # producing a valid CUDA lvalue. (F2 forward CUDA twin pattern: copy
+            # global->local, then operate on the fragment under T.Parallel.)
+            T.copy(  # type: ignore[name-defined]  # noqa: F821
+                A_fp8[by * _FP8_MM_BM : by * _FP8_MM_BM + _FP8_MM_BM,
+                      ko * _FP8_MM_BK : ko * _FP8_MM_BK + _FP8_MM_BK],
+                A_fp8_local,
+                disable_tma=True,
+            )
+            T.copy(  # type: ignore[name-defined]  # noqa: F821
+                B_fp8[bx * _FP8_MM_BN : bx * _FP8_MM_BN + _FP8_MM_BN,
+                      ko * _FP8_MM_BK : ko * _FP8_MM_BK + _FP8_MM_BK],
+                B_fp8_local,
+                disable_tma=True,
+            )
             for i, kk in T.Parallel(_FP8_MM_BM, _FP8_MM_BK):  # type: ignore[name-defined]  # noqa: F821
-                a_val = T.alloc_var("float32")  # type: ignore[name-defined]  # noqa: F821
-                a_val = T.cast(A_fp8[by * _FP8_MM_BM + i, ko * _FP8_MM_BK + kk], "float32")  # type: ignore[name-defined]  # noqa: F821
-                A_shared[i, kk] = T.cast(a_val, "float16")  # type: ignore[name-defined]  # noqa: F821
+                A_shared[i, kk] = T.cast(A_fp8_local[i, kk], "float16")  # type: ignore[name-defined]  # noqa: F821
             for j, kk in T.Parallel(_FP8_MM_BN, _FP8_MM_BK):  # type: ignore[name-defined]  # noqa: F821
-                b_val = T.alloc_var("float32")  # type: ignore[name-defined]  # noqa: F821
-                b_val = T.cast(B_fp8[bx * _FP8_MM_BN + j, ko * _FP8_MM_BK + kk], "float32")  # type: ignore[name-defined]  # noqa: F821
-                B_shared[j, kk] = T.cast(b_val, "float16")  # type: ignore[name-defined]  # noqa: F821
+                B_shared[j, kk] = T.cast(B_fp8_local[j, kk], "float16")  # type: ignore[name-defined]  # noqa: F821
             T.gemm(A_shared, B_shared, C_frag, transpose_B=True)  # type: ignore[name-defined]  # noqa: F821
         sa = A_scale[0]
         sb = B_scale[0]

docs/FP8-ACTIVATIONS-PATHC.md

@@ -1,6 +1,26 @@
 # FP8 ACTIVATIONS END-TO-END for the path_c pipeline (design + microbench spec)
 
-**Status: DESIGN (read-only round), 2026-06-04. Target: gb10 sm_121 (Grace-Blackwell, `tvm.cuda(0)`).**
+**§19 MEASURED UPDATE (gb10 sm_121, 2026-06-04) — fp8 GEMM now RUNS; the SPEED is MEASURED.**
+The §18 "fp8 UNRUNNABLE on sm_121" gap is CLOSED on two fronts. (1) **TE per-tensor E4M3
+(`DelayedScaling`, cuBLASLt fp8 MMA) is the REAL fp8 tensor-core win: 1.57–1.84× over bf16
+MEASURED** at the four prod bs4 GEMM shapes (52.1–61.8 fp8 TFLOPs vs 33.3–34.3 bf16), rel-err
+**0.0375 honest over ALL elements** (≤0.10 ceiling; only the MXFP8 *block-scaling* recipe stays
+upstream-blocked "not supported on 12.0+"; the per-tensor NVRTC-builtins loader gap is fixed by the
+committed self-healing `cppmega_mlx/_gb10_nvrtc_env.py`). (2) **Our in-house Metal dequant→half→
+`T.gemm` route is ported to a `target="cuda"` fragment-C twin** (`fp8_scaled_matmul_path_c_cuda_prim`)
+that COMPILES + RUNS + passes parity (0.0375, native `fp8_e4_t`→`half_t` decode verified in the
+emitted CUDA) — **but it is NOT yet fast (0.18–0.42×): its MMA runs on the decoded fp16 values
+(`mma_sync<kFloat16,…>`, not an fp8-input MMA) on an untuned 32×64×32 single-128-thread tile**, so it
+is the proven *memory* lever (2.0× operand-byte halving MEASURED) and the on-CUDA optimization target,
+NOT a speed lever as-is. Attribution is explicit: the throughput belongs to TE-cuBLASLt fp8; our
+kernel gives the memory halving + a correct on-CUDA fp8 path. Full numbers + reproduce: RELAX §19.
+(Two port edits were required: the Metal `T.alloc_var` decode scratch lowered to a non-assignable
+`float[1]` on codegen_cuda → replaced with a register-fragment stage; and the microbench must pass the
+native `torch.float8_e4m3fn` not a `.view(uint8)` reinterpret, which segfaults the tvm_ffi adapter.)
+
+---
+
+**Status: DESIGN (read-only round) + §19 MEASURED microbench, 2026-06-04. Target: gb10 sm_121 (Grace-Blackwell, `tvm.cuda(0)`).**
 This is Track 2 of the 3-lever Megatron-gap workflow (lever 1 = bs1→bs4; lever 2 = THIS doc, fp8
 activations end-to-end; lever 3 = B2 grid-restructure). It is the user's thesis verbatim: *"why hold
 bf16 activations at all if every kernel is fp8? that is BOTH speed (fp8 tensor cores) AND memory
@@ -22,11 +42,15 @@ any failure it RAISES with where+what — never an fp8→bf16 or bs4→bs1 silen
   (4 tokens × 1 byte = 2× of 1 token × 2 bytes). Concretely ~12.8 GB @8L / ~26 GB @28L for bs4-fp8
   vs the naive ~25.6/~52 GB bs4-fp16 — back to Megatron-class (~26 GB) at 8L. PROJECTION. (Correction
   per §18 adversarial verify: the earlier "bs4-fp8 ≈ bs1-fp16" headline was arithmetically wrong.)
-* **Speed (the de-risk target).** The dominant GEMMs at bs4 are the transformer-block MLP/attention
+* **Speed (NOW MEASURED, §19).** The dominant GEMMs at bs4 are the transformer-block MLP/attention
   GEMMs (M = bs·seq = 16384), NOT the SSD scan. On Blackwell these run on the **regular FP8 Tensor
-  Cores** (MXFP8 / per-tensor cuBLASLt e4m3), which — unlike the FP4 SR/RHT path — DO exist on
-  sm_121 (no tcgen05/TMEM dependency). The microbench measures the realized fp8-vs-bf16 TFLOPs at
-  those exact shapes before any rewrite.
+  Cores** (per-tensor cuBLASLt e4m3), which — unlike the FP4 SR/RHT path — DO exist on sm_121 (no
+  tcgen05/TMEM dependency): **MEASURED 1.57–1.84× over bf16** (TE `DelayedScaling` E4M3, 52.1–61.8
+  fp8 TFLOPs vs 33.3–34.3 bf16, rel-err 0.0375 ≤ 0.10). MXFP8 block-scaling stays upstream-blocked on
+  12.0+. Our in-house `target="cuda"` dequant→`T.gemm` twin RUNS + parity-passes (0.0375) but is
+  **0.18–0.42× as-is** (decoded-fp16 MMA on an untuned tile) — a memory lever (2.0× MEASURED) + the
+  optimization target, not yet a speed lever. (Was: "the microbench measures … before any rewrite" —
+  now measured; full table in RELAX §19.)
 * **Scaling recommendation.** Use **MXFP8 (E4M3, 32-element block / E8M0 scale) for ALL operands
   incl. activation gradients** as the primary recipe, with **per-tensor delayed-amax E4M3** (reusing
   `fp8_amax.py`) as the simpler fallback-FREE alternative for the in-house TileLang kernels that do
@@ -79,6 +103,13 @@ target for the SSD GEMMs that live inside our own grid kernels.
 
 ### 2b. Scaling strategy — RECOMMENDATION: MXFP8 (E4M3, block-32) primary; per-tensor delayed-amax secondary
 
+> **§2b-status (MEASURED 2026-06-04, gb10 sm_121 / CC 12.1) — MXFP8 is UPSTREAM-BLOCKED here today; per-tensor delayed-amax (R1 tensorwise) is the route that RUNS.**
+> The MXFP8 "primary" recommendation below is the *target* recipe for Blackwell, but it does **not** run on gb10 (GB10, sm_121) this round, and we do **not** fake it (RULE #1). NVIDIA TransformerEngine's own Python gate `transformer_engine/pytorch/quantization.py::_compute_mxfp8_support()` (line 69) hard-raises `"MXFP8 (for all gemm layouts) is not supported on 12.0+ architectures yet."` for compute capability ≥ (12,0). Root cause (TE tech lead *ptrendx*, issue [#2668](https://github.com/NVIDIA/TransformerEngine/issues/2668), 2026-02-11): cuBLAS lacks the **non-TN GEMM layouts** MXFP8 backward needs on SM120/SM121 — a cuBLAS gap, **not** hardware/CUTLASS and **not** our code. The fix, PR [#3050](https://github.com/NVIDIA/TransformerEngine/pull/3050) (still an **OPEN DRAFT** as of 2026-05-29, complemented by PR [#2833](https://github.com/NVIDIA/TransformerEngine/pull/2833)), relaxes the gate **only** when `tex.get_cublasLt_version() >= 130600` (**cuBLASLt 13.6.0.2**) — and that cuBLASLt does **not exist anywhere yet** (PyPI max + box max are both 13.5.1.27 = `130501 < 130600`). So even a perfect cherry-pick of #3050 today would still `fail-closed` (`is_mxfp8_available()` → False), which is correct: forcing the gate lower would silently miscompute the MXFP8 backward (a RULE #1 violation). **VERDICT: MXFP8 on sm_121 is genuinely upstream-blocked on (a) an unmerged draft PR and (b) an unreleased cuBLASLt 13.6.0.2 dependency. Recorded as a measured FAIL, never silently degraded to bf16 or tensorwise.**
+>
+> **Forward plan (when NVIDIA ships cuBLASLt 13.6.0.2):** `pip install 'nvidia-cublas>=13.6.0.2'` into `cppmega-venv`, then update `/home/dave/TransformerEngine` to a SHA carrying BOTH PR #2833 and PR #3050 once merged (or cherry-pick `git fetch origin pull/2833/head:pr2833 && pull/3050/head:pr3050` onto a throwaway branch off local `main`, rebuild with `NVTE_CUDA_ARCHS=121a pip install -e . --no-build-isolation`). Expect conflicts: the installed TE is a fork-tainted build (HEAD `8d1d79bf`, v2.16.0.dev0, jewelmusicee remote + local cppmega MXFP8 transpose commits, 31 behind origin) — merge origin/main first. Re-verify forward **and** backward numerically before labeling the route fp8-pass; the drafts are unfinished ("TODO: code/git history clean up", tests box unchecked).
+>
+> **What runs on gb10 today:** **R1 tensorwise** (per-tensor Float8 delayed-amax, the "secondary" below) — TE supports it on sm_121 (NVIDIA forum: "FP8 with delayed scaling works correctly on these architectures as a workaround"); its only blocker was the NVRTC `libnvrtc-builtins.so.13.3` loader path, now fixed durably (system `ldconfig` cache + the committed `cppmega_mlx/_gb10_nvrtc_env.py` re-exec guard + the run-command `LD_LIBRARY_PATH` prefix). And **R2** (our CUDA `T.gemm` e4m3 dequant twin). §19 MEASURED: R1 tensorwise = real fp8 cuBLASLt MMA, 1.58–1.83× over bf16, rel_err 0.0375; R2 compiles+runs+parity 0.0375 but 0.18–0.42× (untuned tile + fp16-MMA).
+
 * **Primary: MXFP8 block-32 E4M3 (E8M0 power-of-2 scale).** On Blackwell this is the native recipe.
   Decisive grounding from the 2025 literature + NVIDIA docs: **E4M3 for ALL tensors including
   activation gradients** (E5M2-for-gradients measurably regresses perplexity on Blackwell; the