fxd0h · Nov 7, 2026
diff --git a/.venv b/.venv
@@ -0,0 +1 @@
+../nxnative/.venv
diff --git a/dm/gfx/nvn/dmsh_chain.h b/dm/gfx/nvn/dmsh_chain.h
@@ -1,3 +1,3 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header): the B shader library translation chain version this source tree expects (default options), one for every game;
 // the host loads only <library dir>/b-<V>/. Run dmsh_chain.py --header after changing the chain's tools and commit this file with the change
-#define DMSH_CHAIN "d403eb0683df"
+#define DMSH_CHAIN "1b36081a67bf"
diff --git a/dm/gfx/tests/hfma_orin_fixture.u32 b/dm/gfx/tests/hfma_orin_fixture.u32
diff --git a/dm/gfx/tests/test_dmsh_hfma.py b/dm/gfx/tests/test_dmsh_hfma.py
@@ -0,0 +1,196 @@
+"""HFMA2 on Metal (dmsh_build.py hfma_msl): the rewrite of the `fma` on unpacked halves that ends in a pack to half, and on this Mac's GPU the rewritten kernel against hfma_orin_fixture.u32, the output
+of PTX fma.rn.f16x2 on a Jetson Orin (sm_87) for 4570 (a, b, c) half pairs (the 70 of 1,048,576 random triples where the float fma + pack of the old translation differs from the Orin, 3000 other random ones and
+1500 of the 24,389 edge triples: denormals, Inf, NaN, powers of two). The old kernel rounds twice (float, then half) and is one half ULP off on the 70 the fixture
+keeps; the rewritten one is the Orin's bits on all of them, with the chain's flags and with fast, safe and relaxed math spelled out"""
+import array
+import ast
+from pathlib import Path
+import re
+import shutil
+import struct
+import subprocess
+import tempfile
+import unittest
+
+TOOLS = Path(__file__).resolve().parents[1] / 'tools'
+FIXTURE = Path(__file__).resolve().parent / 'hfma_orin_fixture.u32'
+NAMES = {'hfma_msl', 'HFMA_HELPERS', 'HFMA_DEF', 'HFMA_WRAP'}
+tree = ast.parse((TOOLS / 'dmsh_build.py').read_text())
+selected = [n for n in tree.body if (isinstance(n, ast.FunctionDef) and n.name in NAMES)
+            or (isinstance(n, ast.Assign) and any(getattr(t, 'id', None) in NAMES for t in n.targets))]
+scope = {'re': re, 'struct': struct}
+exec(compile(ast.Module(body=selected, type_ignores=[]), str(TOOLS / 'dmsh_build.py'), 'exec'), scope)
+hfma_msl = scope['hfma_msl']
+
+HEAD = '#include <metal_stdlib>\nusing namespace metal;\n\n'
+# HFMA2 as spirv-cross writes it (ids made up): the three sources unpacked, one fma a lane, the pack. The second kernel adds .SAT (precise::clamp) and a negated source, the third an FFMA whose result is packed
+KERNEL = '''kernel void check(device const uint *ia [[buffer(0)]], device const uint *ib [[buffer(1)]], device const uint *ic [[buffer(2)]], device uint *out [[buffer(3)]], uint gid [[thread_position_in_grid]])
+{
+    uint _1 = ia[gid];
+    uint _2 = ib[gid];
+    uint _3 = ic[gid];
+    float _10 = float2(as_type<half2>(_1)).x;
+    float _11 = float2(as_type<half2>(_1)).y;
+    float _12 = float2(as_type<half2>(_2)).x;
+    float _13 = float2(as_type<half2>(_2)).y;
+    float _14 = float2(as_type<half2>(_3)).x;
+    float _15 = float2(as_type<half2>(_3)).y;
+    float _16 = fma(_10, _12, _14);
+    float _17 = fma(_11, _13, _15);
+    uint _18 = as_type<uint>(half2(float2(_16, _17)));
+    out[gid] = _18;
+}
+'''
+OTHER = '''kernel void check(device const uint *ia [[buffer(0)]], device const float *fb [[buffer(1)]], device uint *out [[buffer(3)]], uint gid [[thread_position_in_grid]])
+{
+    uint _1 = ia[gid];
+    float _10 = float2(as_type<half2>(_1)).x;
+    float _11 = float2(as_type<half2>(_1)).y;
+    float _12 = fb[gid];
+    float _13 = -_10;
+    float _14 = fma(_10, _12, _11);
+    float _15 = fma(_13, 0.5, _11);
+    float _16 = fma(_10, _11, _10);
+    float _17 = precise::clamp(_15, 0.0, 1.0);
+    float _18 = fma(_11, _10, _10);
+    float _19 = max(_18, 0.0);
+    float _20 = fma(_11, 1.0009765625, _10);
+    float _21 = fma(_11, 0.1, _10);
+    float _22 = fma(_10, _10, _11);
+    uint _30 = as_type<uint>(half2(float2(_14, _17)));
+    uint _31 = as_type<uint>(half2(float2(_16, _19)));
+    uint _32 = as_type<uint>(half2(float2(_20, _21)));
+    out[gid] = _30 + _31 + _32 + as_type<uint>(_22);
+}
+'''
+RUNNER = r'''#import <Foundation/Foundation.h>
+#import <Metal/Metal.h>
+int main(int argc, char **argv) {
+    @autoreleasepool {
+        id<MTLDevice> dev = MTLCreateSystemDefaultDevice();
+        if (!dev) { printf("no Metal device\n"); return 2; }
+        NSError *e = nil;
+        id<MTLLibrary> lib = [dev newLibraryWithURL:[NSURL fileURLWithPath:@(argv[1])] error:&e];
+        id<MTLComputePipelineState> ps = lib ? [dev newComputePipelineStateWithFunction:[lib newFunctionWithName:@"check"] error:&e] : nil;
+        if (!ps) { printf("pipeline: %s\n", e.localizedDescription.UTF8String); return 1; }
+        NSData *d = [NSData dataWithContentsOfFile:@(argv[2])]; NSUInteger n = ((const uint32_t *)d.bytes)[0];
+        const uint32_t *w = (const uint32_t *)d.bytes + 1;
+        id<MTLBuffer> in[3]; for (int k = 0; k < 3; k++) in[k] = [dev newBufferWithBytes:w + k * n length:n * 4 options:MTLResourceStorageModeShared];
+        id<MTLBuffer> out = [dev newBufferWithLength:n * 4 options:MTLResourceStorageModeShared];
+        id<MTLCommandQueue> q = [dev newCommandQueue]; id<MTLCommandBuffer> cb = [q commandBuffer]; id<MTLComputeCommandEncoder> c = [cb computeCommandEncoder];
+        [c setComputePipelineState:ps]; for (int k = 0; k < 3; k++) [c setBuffer:in[k] offset:0 atIndex:k]; [c setBuffer:out offset:0 atIndex:3];
+        [c dispatchThreads:MTLSizeMake(n, 1, 1) threadsPerThreadgroup:MTLSizeMake(64, 1, 1)]; [c endEncoding]; [cb commit]; [cb waitUntilCompleted];
+        [[NSData dataWithBytes:out.contents length:n * 4] writeToFile:@(argv[3]) atomically:YES];
+    }
+    return 0;
+}
+'''
+MODES = ['', '-fmetal-math-mode=fast', '-fmetal-math-mode=safe', '-fmetal-math-mode=relaxed']
+
+
+def fixture():
+    w = array.array('I'); w.frombytes(FIXTURE.read_bytes())
+    n = w[0]
+    return n, list(w[1:1 + n]), list(w[1 + n:1 + 2 * n]), list(w[1 + 2 * n:1 + 3 * n]), list(w[1 + 3 * n:1 + 4 * n])
+
+
+def gpu(test, msl, mode):   # the kernel's output word for each fixture triple
+    with tempfile.TemporaryDirectory() as t:
+        d = Path(t)
+        (d / 'h.metal').write_text(msl)
+        (d / 'run.m').write_text(RUNNER)
+        (d / 'in.bin').write_bytes(FIXTURE.read_bytes())
+        flags = ['-O1', '-fpreserve-invariance'] + ([mode] if mode else [])
+        subprocess.run(['xcrun', '-sdk', 'macosx', 'metal', '-c', *flags, str(d / 'h.metal'), '-o', str(d / 'h.air')], check=True, capture_output=True)
+        subprocess.run(['xcrun', '-sdk', 'macosx', 'metallib', str(d / 'h.air'), '-o', str(d / 'h.metallib')], check=True)
+        subprocess.run(['clang', '-fobjc-arc', '-O1', str(d / 'run.m'), '-framework', 'Metal', '-framework', 'Foundation', '-o', str(d / 'run')], check=True)
+        r = subprocess.run([str(d / 'run'), str(d / 'h.metallib'), str(d / 'in.bin'), str(d / 'out.bin')], capture_output=True, text=True)
+        if r.returncode == 2: test.skipTest(r.stdout.strip())
+        test.assertEqual(r.returncode, 0, r.stdout + r.stderr)
+        o = array.array('I'); o.frombytes((d / 'out.bin').read_bytes())
+        return list(o)
+
+
+def is_nan16(h): return (h >> 10) & 31 == 31 and h & 1023 != 0
+def same_lane(x, y): return x == y or (is_nan16(x) and is_nan16(y))
+def lanes_equal(got, orin):   # lanes that differ from the Orin (any NaN equals any NaN)
+    return sum(not same_lane(g >> s & 0xFFFF, o >> s & 0xFFFF) for g, o in zip(got, orin) for s in (0, 16))
+
+
+class HfmaRewrite(unittest.TestCase):
+    def test_fma_of_unpacked_halves_that_is_packed(self):
+        out = hfma_msl(HEAD + KERNEL)
+        self.assertIn('float _16 = dmsh_hfma(_10, _12, _14);', out)
+        self.assertIn('float _17 = dmsh_hfma(_11, _13, _15);', out)
+
+    def test_helper_once_after_the_using(self):
+        out = hfma_msl(HEAD + KERNEL)
+        self.assertEqual(out.count('static inline float dmsh_hfma('), 1)
+        self.assertLess(out.index('using namespace metal;'), out.index('static inline float dmsh_hfma('))
+        self.assertLess(out.index('dmsh_hfma(float'), out.index('kernel void check'))
+        self.assertIn('math_mode(safe)', out)
+
+    def test_saturate_negation_and_literals_count(self):
+        out = hfma_msl(HEAD + OTHER)
+        self.assertIn('float _15 = dmsh_hfma(_13, 0.5, _11);', out)   # negated half, literal, then precise::clamp, then the pack
+        self.assertIn('float _18 = dmsh_hfma(_11, _10, _10);', out)   # max(x, 0) (.RELU) then the pack
+        self.assertIn('float _20 = dmsh_hfma(_11, 1.0009765625, _10);', out)   # 1 + 2^-10 is a half
+
+    def test_relu_that_keeps_a_nan(self):
+        # .RELU as it keeps a NaN result (the Orin gives the NaN): isnan(v) ? v : max(v, 0.0), then the pack
+        relu = HEAD + '''kernel void check(device const uint *ia [[buffer(0)]], device uint *out [[buffer(3)]], uint gid [[thread_position_in_grid]])
+{
+    uint _1 = ia[gid];
+    float _10 = float2(as_type<half2>(_1)).x;
+    float _11 = float2(as_type<half2>(_1)).y;
+    float _63 = fma(_10, _11, _10);
+    bool _69 = isnan(_63);
+    float _71 = fast::max(_63, 0.0);
+    float _73 = _69 ? _63 : _71;
+    uint _81 = as_type<uint>(half2(float2(_73, _10)));
+    out[gid] = _81;
+}
+'''
+        self.assertIn('float _63 = dmsh_hfma(_10, _11, _10);', hfma_msl(relu))
+
+    def test_pragmas_survive_a_line_at_the_start_of_the_body(self):
+        # a rewrite that puts a line at the start of every function (loop_guard) must leave the helper's pragmas at the start of a compound statement
+        out = hfma_msl(HEAD + KERNEL)
+        body = out[out.index('static inline float dmsh_hfma'):]
+        body = body.replace('{\n', '{\n    uint dm_lg = 0u; (void)dm_lg;\n', 1)
+        self.assertRegex(body, r'\{[^\n]*\n#pragma METAL fp math_mode\(safe\)')
+
+    def test_what_is_not_a_half_fma_stays(self):
+        out = hfma_msl(HEAD + OTHER)
+        self.assertIn('float _14 = fma(_10, _12, _11);', out)   # a float read from memory is not a half
+        self.assertIn('float _16 = dmsh_hfma(_10, _11, _10);', out)
+        self.assertIn('float _21 = fma(_11, 0.1, _10);', out)   # 0.1 is not a half
+        self.assertIn('float _22 = fma(_10, _10, _11);', out)   # kept as a float (HFMA2.F32): not packed
+
+    def test_other_shaders_unchanged(self):
+        plain = HEAD + 'kernel void k() { float _1 = 0.0; float _2 = fma(_1, _1, _1); }\n'
+        self.assertEqual(hfma_msl(plain), plain)
+
+    def test_no_using_namespace_refused(self):
+        with self.assertRaises(ValueError):
+            hfma_msl(KERNEL)
+
+
+@unittest.skipUnless(shutil.which('xcrun') and shutil.which('clang') and FIXTURE.exists(), 'needs Xcode and the Orin fixture')
+class HfmaOnGpu(unittest.TestCase):
+    def test_matches_the_orin_in_every_math_mode(self):
+        n, a, b, c, orin = fixture()
+        msl = hfma_msl(HEAD + KERNEL)
+        for mode in MODES:
+            with self.subTest(mode=mode or 'default'):
+                self.assertEqual(lanes_equal(gpu(self, msl, mode), orin), 0, 'lanes that differ from the Orin [%s]' % mode)
+
+    def test_without_the_rewrite_the_same_lanes_fail(self):
+        n, a, b, c, orin = fixture()
+        wrong = lanes_equal(gpu(self, HEAD + KERNEL, ''), orin)
+        self.assertGreaterEqual(wrong, 60, 'the old fma + pack is a half ULP off on the 70 lanes the fixture keeps (%d here)' % wrong)
+        self.assertLessEqual(wrong, 80)
+
+
+if __name__ == '__main__':
+    unittest.main()
diff --git a/dm/gfx/tools/dmsh_build.py b/dm/gfx/tools/dmsh_build.py
@@ -163,8 +163,8 @@
    for c, args in fcalls: writes[args[HMMA_ARGS.index('d1')]] += 1
    for name in writers:
        for w in re.finditer(r'\b%s\(' % name, body):
            depth, j = 1, w.end()
            while depth and j < len(body): depth += {'(': 1, ')': -1}.get(body[j], 0); j += 1
            writes.update(re.findall(r'\b[A-Za-z_]\w*\b', body[w.end():j]))
    braces = [(b.start() + open_at, 1 if b.group() == '{' else -1) for b in BRACE.finditer(body)]
    at, depth_after = [p for p, _ in braces], list(itertools.accumulate(d for _, d in braces))
@@ -317,6 +317,69 @@
         if u < 0: raise ValueError('dmsh_shared_store: no using namespace metal')
         u += len('using namespace metal;\n'); src = src[:u] + SHARED_ALIAS + src[u:]
     return src
+# HFMA2 on Metal. The translator writes HFMA2 as `fma(a, b, c)` on the two halves unpacked to float and packs the result with half2(float2(..)): one rounding to float, then another to half. The Orin rounds the exact
+# a * b + c once, so when the float result lands on the middle of two halves and the exact one does not (the product is exact in float, the small addend is lost in the first rounding), Metal is one half ULP off.
+# Measured on a Jetson Orin (sm_87, PTX fma.rn.f16x2, 1,048,576 random triples with 15 percent denormals and 6 percent Inf / NaN): 70 of 2,097,152 lanes
+# differ (99.9967 percent identical, e.g. -3136 * -8.125 + 0.000124693: Orin 25488, Metal 25472). dmsh_hfma returns the float rounded to odd (the neighbour of the exact value that has an odd last bit, which is the
+# float itself when the sum was exact or already odd), and the conversion to half then rounds the exact value once: round to odd costs no double rounding when the float has two more bits than the half. The error of
+# the sum comes from a TwoSum of the exact product (11 bit halves: 22 bits) and the addend, which fast math folds away (a - (b - a) is 0 for it) unless the function is compiled in safe math mode and without contraction.
+# Only an fma whose three operands are halves (unpacked, negated, or a literal a half holds) and whose every use ends in a pack to half (through the saturate and the ReLU clamp) is rewritten: an FFMA that is
+# then converted with F2FP rounds twice on the hardware as well, and HFMA2.F32 keeps its float
+HFMA_HELPERS = """// HFMA2 rounded once to half (see hfma_msl)
+static inline float dmsh_hfma(float a, float b, float c)
+{
+  {   // the pragmas open a compound statement of their own: a rewrite that puts a line at the start of every function body (loop_guard) leaves them valid
+#pragma METAL fp math_mode(safe)
+#pragma METAL fp contract(off)
+    float p = a * b;
+    float s = fma(a, b, c);
+    float t = s - p;
+    float err = (p - (s - t)) + (c - t);
+    uint sb = as_type<uint>(s);
+    uint step = ((as_type<uint>(err) ^ sb) >> 31) != 0u ? 0xFFFFFFFFu : 1u;
+    bool odd = err != 0.0 && (sb & 1u) == 0u && (sb & 0x7F800000u) != 0x7F800000u;
+    return select(s, as_type<float>(sb + step), odd);
+  }
+}
+"""
+HFMA_DEF = re.compile(r'^([ \t]*float (_\d+) = )fma\(([^,()]+), ([^,()]+), ([^,()]+)\);$', re.M)
+HFMA_WRAP = re.compile(r'^[ \t]*float (_\d+) = (?:precise::clamp|clamp|fast::clamp|max|fmax|precise::max|fast::max)\((_\d+), [-\d.]+f?(?:, [-\d.]+f?)?\);$', re.M)
+def hfma_msl(msl):
+    defs = {}
+    for m in re.finditer(r'^[ \t]*float (_\d+) = ([^;]*);$', msl, re.M): defs.setdefault(m.group(1), []).append(m.group(2))
+    def is_half(x, depth=0):   # an operand that holds a half: an unpacked lane, a literal a half holds, or the negation or absolute value of one
+        x = x.strip()
+        if depth > 6: return False
+        if re.fullmatch(r'-?\d+\.\d*(?:e[-+]?\d+)?f?|-?\d+\.?', x) is not None:
+            try: v = float(x.rstrip('f')); return struct.unpack('<e', struct.pack('<e', v))[0] == v
+            except (OverflowError, struct.error): return False
+        if x.startswith('-'): return is_half(x[1:], depth + 1)
+        if x.startswith('abs(') and x.endswith(')'): return is_half(x[4:-1], depth + 1)
+        d = defs.get(x)
+        if not d or len(d) != 1: return False
+        return re.fullmatch(r'float2\(as_type<half2>\([^;]*\)\)\.[xy]', d[0]) is not None or is_half(d[0], depth + 1)
+    def packs(v, depth=0):   # every use of v is an argument of a pack to half, or of a clamp / max whose result is
+        uses = [l for l in msl.split('\n') if re.search(r'\b%s\b' % v, l) and not re.match(r'\s*float %s = ' % v, l)]
+        if not uses or depth > 4: return False
+        for l in uses:
+            if re.search(r'half2\(float2\((?:%s, [^;]*?|[^;]*?, %s)\)\)' % (v, v), l): continue
+            w = HFMA_WRAP.match(l)
+            if w and re.search(r'\(%s,' % v, l) and packs(w.group(1), depth + 1): continue
+            if re.fullmatch(r'\s*bool _\d+ = isnan\(%s\);' % v, l): continue   # .RELU keeping a NaN: isnan(v) ? v : max(v, 0)
+            sel = re.fullmatch(r'\s*float (_\d+) = _\d+ \? (_\d+) : (_\d+);', l)
+            if sel and v in sel.group(2, 3) and packs(sel.group(1), depth + 1): continue
+            return False
+        return True
+    out, last = [], 0
+    for m in HFMA_DEF.finditer(msl):
+        if len(defs.get(m.group(2), [])) == 1 and all(is_half(m.group(i)) for i in (3, 4, 5)) and packs(m.group(2)):
+            out.append(msl[last:m.start()] + m.group(1) + 'dmsh_hfma(%s, %s, %s);' % m.group(3, 4, 5)); last = m.end()
+    if not out: return msl
+    out = ''.join(out) + msl[last:]
+    u = out.find('using namespace metal;\n')
+    if u < 0: raise ValueError('hfma: no using namespace metal')
+    u += len('using namespace metal;\n')
+    return out[:u] + HFMA_HELPERS + out[u:]
 # Compute shared memory past Metal's limit: Apple's GPUs give a kernel 32 KB of threadgroup memory (maxThreadgroupMemoryLength), and Metal refuses a pipeline that declares more, while Ampere
 # gives up to 100 KB. Mario Kart World's DrivePassShaderDefault has 168 compute shaders of 48 KB (spvUnsafeArray<uint, 12288>), seven of them run on the way to the first race; refused, their
 # dispatches were skipped. Such a kernel takes its threadgroup arrays from a device buffer the host binds (abi.TG_SLOT): one slice of the arrays' total per workgroup, at the workgroup's index in
@@ -1756,6 +1819,8 @@
     except ValueError as e: return 'gather offsets rewrite failed', str(e)[:160], None
     try: src = hmma_msl(src)   # Ampere HMMA through simdgroup_matrix (see hmma_msl)
     except ValueError as e: return 'HMMA rewrite failed', str(e)[:160], None
+    try: src = hfma_msl(src)   # HFMA2 rounded once to half (see hfma_msl)
+    except ValueError as e: return 'HFMA2 rewrite failed', str(e)[:160], None
     try: src = shared_store_msl(src)   # shared memory byte and half-word stores as Metal byte stores (see shared_store_msl)
     except ValueError as e: return 'shared store rewrite failed', str(e)[:160], None
     try: src = pool_arrays(src)   # handles traced to no slot read the pools the replay keeps (see pool_arrays)

diff --git a/docs/agents/symptoms.md b/docs/agents/symptoms.md
@@ -954,6 +954,7 @@ One line per symptom in the [subsystem references](../reference/): what you see
 - `[RED] failed to create compute pipeline: Threadgroup memory size (49152) exceeds the maximum threadgroup memory allowed (32768)`, thousands a run; the dispatch is skipped. Once moved to device memory: with more than 512 workgroups in a dispatch, or an indirect one, workgroups found other workgroups' data in their shared memory (no game seen doing so). → [Compute shared memory past 32 KB](../reference/shaders.md#compute-shared-memory-past-32-kb)
 - `[RED] shader <key>_b reflection conversion failed: fetch call 0: frag_pos uniform reference out of range`; the shader does not load and its draws are skipped. → [Uniform references past the first kilobyte](../reference/shaders.md#uniform-references-past-the-first-kilobyte)
 - the GPU hangs (`kIOGPUCommandBufferCallbackErrorHang`) from the first 3D scene on, or frames of about 88 ms, in a shader with a half-float loop. → [HADD2 and HMUL2 with an immediate](../reference/shaders.md#hadd2-and-hmul2-with-an-immediate)
+- none seen in a game. Against a Jetson Orin (sm_87, PTX `fma.rn.f16x2`, 1,048,576 random triples with denormals, Inf and NaN, plus 24,389 edge triples), 70 of 2,097,152 lanes differed by one half ULP (99.9967 percent identical), e.g. -3136 × -8.125 + 0.000124693: Orin 25488, Metal 25472. → [HFMA2 rounds twice to half](../reference/shaders.md#hfma2-rounds-twice-to-half)
 - wrong shading where a shader needs a cross product or a 2D rotation of half pairs; on Eden-NS2, which first read the mode as a plain negation of the high lane, sun and cloud shadows were lost. → [HFMA2 .H0_NH1](../reference/shaders.md#hfma2-h0_nh1)
 - a global load, store or atomic at the wrong address in code that uses `desc[URb][Ra.64]` (the descriptor's value added to the address); a shader with `ATOMG.CAS` refused (`Ampere instruction 3a9 ... is not implemented`); an `ATOMG.INC` / `ATOMG.DEC` / `RED.INC` / `RED.DEC` with no effect. → [Global memory: descriptors, ATOMG.CAS, INC / DEC](../reference/shaders.md#global-memory-descriptors-atomgcas-inc--dec)
 - a shared-memory counter that never moves (`ATOMS.INC`, `ATOMS.DEC`: CUDA's atomicInc / atomicDec, the old value read as 0); a lock-free update on shared memory that loops or takes a wrong branch after `ATOMS.CAS` (its result 0 or 1 instead of the old value). → [ATOMS.INC / ATOMS.DEC, and ATOMS.CAS's result](../reference/shaders.md#atomsinc--atomsdec-and-atomscass-result)