Float8E4M3: fn (float8_e4m3fn) is now the DEFAULT overflow convention - 4.14.0-local.2

Correctness over compatibility (lib <4mo, no external users): the cast operator
and the IR-level convert now use the fn convention - finite overflow AND +-Inf
-> NaN - making Float8E4M3 bit-exact to PyTorch/JAX/ml_dtypes float8_e4m3fn, the
dtype it is named after. Was: saturate to +-448. Since the change is at the IR
convert, the cast, PrecisionConvert, and the generic INumber<T> path are all
consistently fn.

Changed the E4M3 overflow branch in the managed conversion AND all 6 backend
emitters (PTX/OpenCL/WGSL/GLSL/Wasm bit-manip): e>8 -> NaN, and the post-round
carry reaching the 0x7F slot -> NaN (check the FULL result, not the masked low 7
bits, so the 0x80 carry is caught). 449..464 round DOWN to 448 under both
conventions; only |x|>464 differs (fn: NaN, saturating: +-448).

Saturating is now opt-in via FromSingleSaturating / FromSingle(x, saturate:true):
finite overflow -> +-448, +-Inf -> NaN. Built from a BIT-level finite check
(exponent != all-ones, NOT a float compare vs MaxValue - that is unreliable for
+-Inf on WebGL) + the fn cast + a >464 redirect, so it transpiles with no
per-backend codegen.

Gates: fp8-oracle cast/FromSingleFn/FromSingleSaturating all 1099/1099 vs the
reference (decode 0/256, E5M2 723/723); fp8-verify desktop 24/24; PMT
Float8E4M3_FromSingleFn_OverflowToNaN 9/0 all lanes (cast fn + saturating, incl
the WebGL Inf fix); no regression: PrecisionConvert 37/0, Fp8Radix 44/0. Forks 2.0.29.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

Author: LostBeard

CHANGELOG.md

@@ -2,13 +2,15 @@
 
 This file tracks notable changes per release. The README's "Recent Highlights" section links here for the full version history.
 
-## 4.14.0-local.1 (2026-06-17) - Float8E4M3 selectable overflow convention (float8_e4m3fn parity)
+## 4.14.0-local.2 (2026-06-17) - Float8E4M3 is now bit-exact to float8_e4m3fn (overflow → NaN), saturating opt-in
 
-Additive new API on `Float8E4M3` (forks bump to `2.0.28`). No change to any existing behavior - the cast operator is unchanged.
+`Float8E4M3` float→fp8 conversion changed from saturating to the `fn` (`float8_e4m3fn`) convention as the DEFAULT, matching the dtype it is named after. Forks bump to `2.0.29`.
 
-- **Validated `Float8E4M3` / `Float8E5M2` conversions against the `ml_dtypes` reference** (the impl PyTorch / JAX `float8_e4m3fn` / `float8_e5m2` share) with a new evidence harness `DemoConsole -- fp8-oracle` (generators in `_research/fp8_oracle/`). Result: decode is bit-exact (0/256) and encode rounding/subnormal is bit-exact (0 divergences across 1099 / 723 probes) for both types. The **only** divergence was E4M3 finite overflow: ILGPU saturated to ±448 (the NVIDIA Transformer Engine / OCP saturating cast), whereas the dtype literally named `e4m3fn` overflows to **NaN**. Both are real-world conventions; they agree everywhere except `|x| > 464` (the region that rounds up past the 448 slot).
-- **Made the overflow convention selectable.** The bare cast `(Float8E4M3)x`, `Float8E4M3.FromSingleSaturating(x)`, and `FromSingle(x, saturate: true)` keep the **saturating** behavior (finite overflow → ±448, ±Inf → NaN). New **`Float8E4M3.FromSingleFn(x)`** / `FromSingle(x, saturate: false)` use the **fn** convention (finite overflow AND ±Inf → NaN), **bit-exact to PyTorch / JAX / ml_dtypes `float8_e4m3fn`** - use it for reference-matching ML (loading/comparing FP8 checkpoints). `FromSingleFn` is composed only of existing intrinsics (compare + the saturating cast + Neg + cast-of-NaN), so it transpiles with **no per-backend conversion codegen** and is bit-exact on all 6 backends.
-- Gates: `DemoConsole -- fp8-oracle` (managed `FromSingleFn` 1099/1099 vs `float8_e4m3fn`; saturating cast's 62 overflow points reported as the documented convention) + `fp8-verify` desktop kernel (`FromSingleFn` 24/24 bit-exact on CPU/OpenCL/CUDA) + **PMT `Float8E4M3_FromSingleFn_OverflowToNaN` 9/0 across all backend lanes** (CPU/CUDA/OpenCL/WebGPU/WebGPU-NoSubgroups/WebGL/Wasm). No regression to existing FP8/bf16/Half gates. `Float8E5M2` already matched its reference (overflow → ±Inf); its canonical NaN byte is `0x7F` (ml_dtypes uses `0x7E` - both valid NaN patterns).
+- **Validated `Float8E4M3` / `Float8E5M2` conversions against the `ml_dtypes` reference** (the impl PyTorch / JAX `float8_e4m3fn` / `float8_e5m2` share) with a new evidence harness `DemoConsole -- fp8-oracle` (generators in `_research/fp8_oracle/`). Decode is bit-exact (0/256) and encode rounding/subnormal bit-exact (0 divergences across 1099 / 723 probes) for both. The one prior divergence was E4M3 finite overflow: ILGPU **saturated to ±448** while `e4m3fn` **overflows to NaN**. The two agree everywhere except `|x| > 464` (the region that rounds up past the 448 slot; `449..464` round **down** to 448 under both).
+- **Default flipped to `fn` (overflow → NaN), made selectable.** The cast operator `(Float8E4M3)x`, `Float8E4M3.FromSingleFn(x)`, and `FromSingle(x, saturate: false)` now map finite overflow AND ±Inf → NaN - **bit-exact to PyTorch / JAX / ml_dtypes `float8_e4m3fn`**. The change is at the **IR-level convert**, so every float→fp8 path (the cast, `PrecisionConvert`, the generic `INumber<T>` path) is consistently `fn`. This required updating the E4M3 overflow branch in the managed conversion AND all 6 backend emitters (PTX / OpenCL / WGSL / GLSL / Wasm bit-manip): `e>8` → NaN, and the post-round carry that reaches the `0x7F` slot → NaN (checking the full result, not the masked low 7 bits, so the `0x80` carry is caught).
+- **Saturating is now opt-in.** `Float8E4M3.FromSingleSaturating(x)` / `FromSingle(x, saturate: true)` clamp finite overflow to ±448 (±Inf → NaN) - the NVIDIA Transformer Engine saturating cast / OCP mode. It is composed of existing intrinsics (a **bit-level** finite check - exponent ≠ all-ones, NOT a float compare against `MaxValue`, which is unreliable for ±Inf on WebGL - plus the fn cast and a `>464` redirect), so it transpiles with no per-backend codegen and is bit-exact on all 6 backends.
+- Gates: `DemoConsole -- fp8-oracle` (cast operator / `FromSingleFn` / `FromSingleSaturating` all 1099/1099 vs the reference; decode 0/256; E5M2 723/723) + `fp8-verify` desktop kernels (`FromSingleFn` 24/24 bit-exact CPU/OpenCL/CUDA) + **PMT `Float8E4M3_FromSingleFn_OverflowToNaN` 9/0 all backend lanes** (cast operator fn + `FromSingleSaturating`, incl WebGPU/WebGL/Wasm) + no regression: `PMT_FILTER=PrecisionConvert` 37/0, `PMT_FILTER=Fp8Radix` 44/0. `Float8E5M2` already matched its reference (overflow → ±Inf); its canonical NaN byte is `0x7F` (ml_dtypes uses `0x7E` - both valid NaN patterns).
+- _Supersedes the unreleased local.1 (which had kept the cast operator saturating); local.1 was never promoted._
 
 ## 4.13.2 (2026-06-16) - Packaging fix (no code changes)
 

CLAUDE.md

@@ -147,7 +147,7 @@ If total > 10: `InvalidOperationException` at dispatch time (v4.9.1+). Before v4
 
 `ArrayView<byte>`, `ArrayView<sbyte>`, `ArrayView<short>`, `ArrayView<ushort>`, `ArrayView<Half>` (ILGPU.Half), `ArrayView<BFloat16>` (ILGPU.BFloat16) supported on all 6 backends.
 
-**Use `ILGPU.Half`, NOT `System.Half`** in kernel signatures. Implicit conversion operators exist for interop. Same for **`ILGPU.BFloat16`** (the "brain float": top 16 bits of an fp32, so fp32's full dynamic range - the ML-weights trade vs `Half`) and the two FP8 types **`ILGPU.Float8E4M3`** (forward/inference, no Inf; **selectable overflow**: the cast/`FromSingleSaturating` clamps to ±448 = NVIDIA TE/OCP, `FromSingleFn` → NaN = bit-exact PyTorch/JAX `float8_e4m3fn`) + **`ILGPU.Float8E5M2`** (backward/gradient, IEEE Inf/NaN). bf16/FP8 detail: [Docs/data-type-support.md](Docs/data-type-support.md). On CUDA bf16 + FP8 use an f32-register-compute model (no native PTX bf16/fp8 arithmetic); the load/store conversion is **PORTABLE bit-manipulation (basic integer ops on every CUDA arch incl. pre-Ampere)** - 4.13.0+ replaced the sm_80-only `cvt.*.bf16` shortcut that broke on older cards. The browser/OpenCL/Wasm backends emulate the same exact conversion, byte-identical to CUDA.
+**Use `ILGPU.Half`, NOT `System.Half`** in kernel signatures. Implicit conversion operators exist for interop. Same for **`ILGPU.BFloat16`** (the "brain float": top 16 bits of an fp32, so fp32's full dynamic range - the ML-weights trade vs `Half`) and the two FP8 types **`ILGPU.Float8E4M3`** (forward/inference, no Inf; **selectable overflow, default = fn**: the cast operator / `FromSingleFn` map overflow → NaN = bit-exact PyTorch/JAX `float8_e4m3fn`; `FromSingleSaturating` clamps to ±448 = NVIDIA TE/OCP) + **`ILGPU.Float8E5M2`** (backward/gradient, IEEE Inf/NaN). bf16/FP8 detail: [Docs/data-type-support.md](Docs/data-type-support.md). On CUDA bf16 + FP8 use an f32-register-compute model (no native PTX bf16/fp8 arithmetic); the load/store conversion is **PORTABLE bit-manipulation (basic integer ops on every CUDA arch incl. pre-Ampere)** - 4.13.0+ replaced the sm_80-only `cvt.*.bf16` shortcut that broke on older cards. The browser/OpenCL/Wasm backends emulate the same exact conversion, byte-identical to CUDA.
 
 **Per-backend implementation:**
 - **WebGPU:** Packed into `array<atomic<u32>>`. Load via atomicLoad + shift + mask. Store via atomicAnd + atomicOr (thread-safe sub-word writes). Float16 load/store calls `_f16_to_f32` / `_f32_to_f16` helpers from `WGSLEmulationLibrary.F16Functions` when `!shader-f16`; native WGSL `f16` type otherwise. `WebGPUBackend.ForceEmulatedF16` test flag forces the emulation path for verification.

Docs/data-type-support.md

@@ -187,22 +187,24 @@ backend** (CPU-verified idempotence 0/256 for all representable values).
 | **CUDA** | f32-register model. The FP8<->f32 conversion is **inline PTX bit-manipulation** (branchless `setp`/`selp`, unrolled normalize) using only basic integer ops - FP8 has no portable native PTX cvt (`cvt.*.e4m3` is sm_89/Hopper only), so this works on every CUDA arch. Load = `ld.global.u8` + convert; store = convert + `st.global.u8`. |
 | **CPU** | Native - the managed `Float8E4M3`/`Float8E5M2` structs run directly. |
 
-> **Convention note (E4M3 overflow) - SELECTABLE.** E4M3 has two real-world overflow conventions and
-> both are exposed; the conversion is otherwise **bit-exact** to the `ml_dtypes` reference (the impl
-> PyTorch / JAX `float8_e4m3fn` share) - verified by `DemoConsole -- fp8-oracle`: decode 0/256, encode
-> rounding/subnormal 0 divergences across 1099 probes.
+> **Convention note (E4M3 overflow).** The conversion is **bit-exact** to the `ml_dtypes` reference (the
+> impl PyTorch / JAX `float8_e4m3fn` share) - verified by `DemoConsole -- fp8-oracle`: decode 0/256,
+> encode rounding/subnormal/overflow 0 divergences across 1099 probes, on all 6 backends. The overflow
+> behavior is **selectable**, with the reference-matching `fn` convention as the default:
 >
-> | Entry point | Finite overflow | ±Inf | Matches |
+> | Entry point | Finite overflow (`\|x\|>464`) | ±Inf | Matches |
 > |---|---|---|---|
-> | `(Float8E4M3)x` cast / `FromSingleSaturating(x)` / `FromSingle(x, saturate: true)` | clamps to ±448 | → NaN | NVIDIA Transformer Engine default cast / OCP saturating-forward |
-> | `FromSingleFn(x)` / `FromSingle(x, saturate: false)` | → NaN | → NaN | **PyTorch / JAX / ml_dtypes `float8_e4m3fn`** (bit-exact) |
+> | `(Float8E4M3)x` cast / `FromSingleFn(x)` / `FromSingle(x, saturate: false)` — **DEFAULT** | → NaN | → NaN | **PyTorch / JAX / ml_dtypes `float8_e4m3fn`** (bit-exact) |
+> | `FromSingleSaturating(x)` / `FromSingle(x, saturate: true)` | clamps to ±448 | → NaN | NVIDIA Transformer Engine saturating cast / OCP saturating-forward |
 >
-> The two agree everywhere except `|x| > 464` (the region that rounds up past the 448 slot): saturating
-> gives ±448, fn gives NaN. Every *representable* value round-trips exactly under both. `FromSingleFn` is
-> composed only of existing intrinsics (compare + the saturating cast + Neg + cast-of-NaN), so it
-> transpiles and is bit-exact on **all 6 backends** (PMT `Float8E4M3_FromSingleFn_OverflowToNaN`). Use
-> `FromSingleFn` for reference-matching ML (e.g. loading/comparing PyTorch FP8 checkpoints); use the
-> saturating cast when you want overflow clamped rather than NaN-poisoning a downstream reduction.
+> The cast operator and the IR-level convert (so `PrecisionConvert` and the generic `INumber<T>` path too)
+> are all `fn`. `449..464` round **down** to 448 under both conventions; the two differ only for `|x|>464`,
+> which rounds up past the 448 slot (`fn` → NaN, saturating → ±448). Every *representable* value round-trips
+> exactly. `FromSingleSaturating` is composed only of existing intrinsics (a bit-level finite check + the fn
+> cast + a `>464` redirect), so it transpiles and is bit-exact on **all 6 backends** (PMT
+> `Float8E4M3_FromSingleFn_OverflowToNaN`). Use the default for reference-matching ML (loading/comparing
+> PyTorch FP8 checkpoints); use `FromSingleSaturating` when you want overflow clamped rather than
+> NaN-poisoning a downstream reduction.
 >
 > `Float8E5M2` is IEEE-754-style (has ±Inf): overflow → ±Inf, bit-exact to `float8_e5m2` (decode 0/256,
 > encode 723/723); its canonical NaN byte is `0x7F` (ml_dtypes uses `0x7E` - both are valid NaN patterns).

ILGPU.Algorithms/ILGPU.Algorithms.csproj

@@ -12,7 +12,7 @@
          SpawnDev.ILGPU.Fork* PackageReference Versions inside SpawnDev.ILGPU.csproj.
          Run `_check-fork-version-sync.bat` at repo root. See the banner comment in
          SpawnDev.ILGPU.csproj for the full procedure. -->
-    <Version>2.0.28</Version>
+    <Version>2.0.29</Version>
     <IsPackable>true</IsPackable>
     <GeneratePackageOnBuild>true</GeneratePackageOnBuild>
   </PropertyGroup>

ILGPU/Backends/OpenCL/CLBackend.cs

@@ -220,7 +220,7 @@ public CLBackend(
             extensionBuilder.AppendLine("    int f32Exp = (int)((rest >> 23) & 0xFFu);");
             extensionBuilder.AppendLine("    uint f32Mant = rest & 0x7FFFFFu;");
             extensionBuilder.AppendLine("    int e = f32Exp - 127;");
-            extensionBuilder.AppendLine("    if (e > 8 || (e == 8 && f32Mant > 0x600000u)) return (uchar)(sign | 0x7Eu);");
+            extensionBuilder.AppendLine("    if (e > 8) return (uchar)(sign | 0x7Fu);"); // fn: e>8 unconditional overflow -> NaN; e==8 rounds below
             extensionBuilder.AppendLine("    if (e < -6) {");
             extensionBuilder.AppendLine("        if (f32Exp == 0) return (uchar)sign;");
             extensionBuilder.AppendLine("        uint signif = f32Mant | 0x800000u;");
@@ -238,7 +238,7 @@ public CLBackend(
             extensionBuilder.AppendLine("    uint eField = (uint)(e + 7);");
             extensionBuilder.AppendLine("    uint outBits = (eField << 3) | mant3;");
             extensionBuilder.AppendLine("    if (round == 1u && (stick == 1u || (mant3 & 1u) == 1u)) outBits += 1u;");
-            extensionBuilder.AppendLine("    if ((outBits & 0x7Fu) >= 0x7Fu) outBits = 0x7Eu;");
+            extensionBuilder.AppendLine("    if (outBits >= 0x7Fu) outBits = 0x7Fu;"); // fn: full outBits (incl 0x80 carry) reaching the 0x7F slot -> NaN
             extensionBuilder.AppendLine("    return (uchar)(sign | (outBits & 0x7Fu));");
             extensionBuilder.AppendLine("}");
             extensionBuilder.AppendLine();

ILGPU/Backends/PTX/PTXCodeGenerator.Values.cs

@@ -773,8 +773,8 @@ void EmitI(string op, HardwareRegister d, HardwareRegister a, long imm)
 
         /// <summary>
         /// Emits a PORTABLE f32 -&gt; FP8 raw-byte (low 8 bits in dst .b16) conversion using only basic
-        /// integer ops (every CUDA arch). Branchless (setp/selp), RNE rounding; E4M3 saturates finite
-        /// overflow to +-448 + maps Inf-&gt;NaN, E5M2 overflows to Inf. Byte-identical to the managed/Wasm
+        /// integer ops (every CUDA arch). Branchless (setp/selp), RNE rounding; E4M3 = fn (finite
+        /// overflow AND Inf -&gt; NaN, float8_e4m3fn), E5M2 overflows to Inf. Byte-identical to the managed/Wasm
         /// ConvertFloatToFloat8E*M* (CPU-verified). The subnormal shift is clamped (PTX shr is UB for
         /// shift&gt;=32) and edge-guarded to match the managed return-0 cases.
         /// </summary>
@@ -844,10 +844,10 @@ private void EmitF32ToFP8Bits(HardwareRegister srcF32, HardwareRegister dstByte,
             Selp(nrm, t, nrm, p);
             if (isE4M3)
             {
-                // if (nrm & 0x7F) >= 0x7F -> 0x7E (avoid NaN slot)
-                EmitI("and.b32", t, nrm, 0x7F);
-                using (var c = BeginCommand("setp.ge.u32")) { c.AppendArgument(p); c.AppendArgument(t); c.AppendConstant(0x7F); }
-                MovI(t2, 0x7E); Selp(nrm, t2, nrm, p);
+                // fn: if nrm (FULL, incl a 0x80 carry) reaches the 0x7F slot -> 0x7F (NaN).
+                // Compare nrm directly (not masked) so the round-up-past-448 carry is caught.
+                using (var c = BeginCommand("setp.ge.u32")) { c.AppendArgument(p); c.AppendArgument(nrm); c.AppendConstant(0x7F); }
+                MovI(t2, 0x7F); Selp(nrm, t2, nrm, p);
             }
             EmitI("and.b32", nrm, nrm, 0x7F);
             Emit("or.b32", nrm, sign, nrm);
@@ -899,20 +899,10 @@ private void EmitF32ToFP8Bits(HardwareRegister srcF32, HardwareRegister dstByte,
             // overflow
             if (isE4M3)
             {
-                // (ev>8) || (ev==8 && f32Mant>0x600000) -> sign|0x7E
+                // fn: only ev>8 is unconditional overflow -> sign|0x7F (NaN). ev==8 is handled by
+                // the normal RNE path + its full-outBits>=0x7F clamp (449->448, >464->NaN).
                 SetpI("setp.gt.s32", p, ev, 8);
-                SetpI("setp.eq.s32", p2, ev, 8);
-                using (var c = BeginCommand("setp.gt.u32")) { c.AppendArgument(p2); c.AppendArgument(f32Mant); c.AppendConstant(0x600000); }
-                // need ev==8 AND f32Mant>0x600000: recompute (p2 got overwritten); use a temp predicate via and.pred
-                SetpI("setp.eq.s32", p2, ev, 8);
-                {
-                    var p3 = AllocateRegister(BasicValueType.Int1, PTXRegisterKind.Predicate);
-                    using (var c = BeginCommand("setp.gt.u32")) { c.AppendArgument(p3); c.AppendArgument(f32Mant); c.AppendConstant(0x600000); }
-                    using (var c = BeginCommand("and.pred")) { c.AppendArgument(p2); c.AppendArgument(p2); c.AppendArgument(p3); }
-                    using (var c = BeginCommand("or.pred")) { c.AppendArgument(p); c.AppendArgument(p); c.AppendArgument(p2); }
-                    FreeRegister(p3);
-                }
-                EmitI("or.b32", t, sign, 0x7E);
+                EmitI("or.b32", t, sign, 0x7F);
                 Selp(result, t, result, p);
             }
             else