Add MicroCUDA backend, adaptive output wait, wider native K coverage (improve perf)

Shim
- Adaptive output-event wait: sleep to ~120 us before the predicted
  completion (smoothed per wait site), spin near it, back off when late.
  D4R_SHIM_OUTPUT_POLL_US restores a fixed interval.

- Defer the input-marker wait into the worker thread with VRAM interop
  (D4R_SHIM_DEFER_MARKER, default on), taking the host setup off the
  critical path.

MicroCUDA
- Experimental ZLUDA-free backend in wine_nvcuda_bridge, plus
  tools/microcuda and scripts/build_microcuda.sh. Serves accuracy, fast
  and native kernels; unregistered modules fall back to the reference
  path. Unclassified writes invalidate prepared weights.

M kernels
- SWIN_PREP_IDENTITY_AT: key weight prep by allocation identity
  (MicroCUDA); enc3 tube uses it with 8 slots. ZLUDA ignores it.

- dec2 (gfx1101): SWIN_PACKED_RESIDUAL, byte-identical, ~1.6% faster.

- Opt-in gfx12 hardware fp8 conversion (SWIN_FP8_HWCVT), unverified
  on hardware.

Native K translators
- ctile: handle variants that re-materialize the parameter base (mvlo).

- rewrite_output: check live-in registers, not just outputs.

- build.sh: native input covers all 8 flag variants, output the 4 HDR
  "max" variants; depth-tile skip stays mvhi-only.

Tools
- dump_runner and bench_dlss_kernel harness additions.

kernels/build.sh

@@ -125,13 +125,23 @@ if [[ "$WHAT" == all || "$WHAT" == tex || "$WHAT" == l ]]; then
     for mode in static dynamic; do
         for range in hdr ldr; do specs+=("rrlite_downsample_kernel_${mode}_${range}:sust_only${downsample_mode}"); done
     done
-    # The fast set serves two K kernels from native code translated from the PTX (kernels/native): these builds
-    # replace the sust_only build of the same name. The accuracy set keeps ZLUDA's compile (denormals preserved).
-    native_out=hiluma_engine_output_depthinv_mvhi_hdr_max_v2_rel
-    native_in=hiluma_engine_input_depthinv_mvhi_hdr_v2_rel
+    # The fast set serves K's hiluma input/output kernels from native code translated from the PTX
+    # (kernels/native): these builds replace the sust_only build of the same name. Every flag variant whose
+    # code the translators prove compatible is covered (they refuse anything else): input = all 8
+    # depth x motion x range variants; output = the 4 HDR "max" variants (the LDR max kernels are a different
+    # kernel structure and keep the sust_only build). D4R_K_NO_DEPTH_TILE (skip the depth tile fill) is proven
+    # for mvhi only: mvlo kernels read the depth tile, so they keep the fill (measured: byte-exact only then).
+    # The accuracy set keeps ZLUDA's compile (denormals preserved).
+    native_outs=() native_ins=()
+    for depth in depthinv depthreg; do
+        for mv in mvhi mvlo; do
+            native_outs+=("hiluma_engine_output_${depth}_${mv}_hdr_max_v2_rel")
+            for range in hdr ldr; do native_ins+=("hiluma_engine_input_${depth}_${mv}_${range}_v2_rel"); done
+        done
+    done
     for spec in "${specs[@]}"; do
         IFS=: read -r kernel src mode <<< "$spec"
-        [[ "$ACCURACY" == 0 && "$kernel" == "$native_out" ]] && continue
+        [[ "$ACCURACY" == 0 && " ${native_outs[*]} " == *" $kernel "* ]] && continue
         # accuracy sets are wave32, except M's post and downsample kernels on RDNA4: their wave64 builds keep the
         # denormal handling and gave the same image on an RX 9070 XT (post 0.79 -> 0.67 ms, downsample 0.43 -> 0.39 ms)
         w64=0
@@ -145,24 +155,28 @@ if [[ "$WHAT" == all || "$WHAT" == tex || "$WHAT" == l ]]; then
     done
     if [[ "$WHAT" != l && "$ACCURACY" == 0 ]]; then
         # NVIDIA's PTX translated to HIP (native/ptx2hip.py) and rebuilt natively; the generated sources stay in a
-        # temporary directory. Both kernels run through the same texture/surface helpers as the other texture
-        # kernels and are byte-identical to ZLUDA's compile of the PTX.
+        # temporary directory. They run through the same texture/surface helpers as the other texture kernels
+        # and are byte-identical to ZLUDA's compile of the PTX (validated per variant in the harness).
         ptx_of() { grep -l -E "\.entry[[:space:]]+$1[[:space:]]*\(" "$PTX_DIR"/*.ptx | head -1; }
         NAT="$HERE/native"
         NT="$(mktemp -d)"
         trap 'rm -rf "$NT"' EXIT
-        ptx="$(ptx_of "$native_out")"; [[ -n "$ptx" ]] || { echo "no PTX defines $native_out" >&2; exit 2; }
-        # hiluma output (path B for Quality..UltraPerf), LDS tile as a function, global-memory redirect stores
-        python3 "$NAT/ptx2hip.py" --ctile --as1 "$ptx" "$native_out" > "$NT/output_translated.hip"
-        python3 "$NAT/rewrite_output.py" "$NT/output_translated.hip" > "$NT/$native_out.hip"
-        D4R_ROCM_DIR="$ROCM" D4R_GPU_ARCH="$ARCH" "$NAT/build_native.sh" "$NT/$native_out.hip" "$OUT/$native_out.hsaco" \
-            -DD4R_NATIVE_F2I -DD4R_K_NO_DEPTH_TILE -DD4R_AS1_STORES
-        # hiluma input: scratch arrays in registers, wave32
-        ptx="$(ptx_of "$native_in")"; [[ -n "$ptx" ]] || { echo "no PTX defines $native_in" >&2; exit 2; }
-        python3 "$NAT/ptx2hip.py" --ifconv --sink "$ptx" "$native_in" > "$NT/in_ifs.hip"
-        python3 "$NAT/in_fuse.py" "$NT/in_ifs.hip" > "$NT/in_unit.h"
-        cp "$NAT/in_regs.hip" "$NT/in_regs.hip"
-        D4R_ROCM_DIR="$ROCM" D4R_GPU_ARCH="$ARCH" WAVE=32 "$NAT/build_native.sh" "$NT/in_regs.hip" "$OUT/$native_in.hsaco"
+        for native_out in "${native_outs[@]}"; do
+            ptx="$(ptx_of "$native_out")"; [[ -n "$ptx" ]] || { echo "no PTX defines $native_out" >&2; exit 2; }
+            # hiluma output (path B for Quality..UltraPerf), LDS tile as a function, global-memory redirect stores
+            python3 "$NAT/ptx2hip.py" --ctile --as1 "$ptx" "$native_out" > "$NT/output_translated.hip"
+            python3 "$NAT/rewrite_output.py" "$NT/output_translated.hip" > "$NT/$native_out.hip"
+            depth_tile=(); [[ "$native_out" == *_mvhi_* ]] && depth_tile=(-DD4R_K_NO_DEPTH_TILE)
+            D4R_ROCM_DIR="$ROCM" D4R_GPU_ARCH="$ARCH" "$NAT/build_native.sh" "$NT/$native_out.hip" "$OUT/$native_out.hsaco" \
+                -DD4R_NATIVE_F2I "${depth_tile[@]}" -DD4R_AS1_STORES
+        done
+        for native_in in "${native_ins[@]}"; do
+            ptx="$(ptx_of "$native_in")"; [[ -n "$ptx" ]] || { echo "no PTX defines $native_in" >&2; exit 2; }
+            # hiluma input: per-thread LDS scratch and local array in registers (in_lds2reg.py), wave32
+            python3 "$NAT/ptx2hip.py" --ifconv --sink "$ptx" "$native_in" > "$NT/in_ifs.hip"
+            python3 "$NAT/in_lds2reg.py" "$NT/in_ifs.hip" > "$NT/$native_in.hip"
+            D4R_ROCM_DIR="$ROCM" D4R_GPU_ARCH="$ARCH" WAVE=32 "$NAT/build_native.sh" "$NT/$native_in.hip" "$OUT/$native_in.hsaco"
+        done
     fi
 fi
 [[ "$ACCURACY" == 1 ]] && printf '1\n' > "$OUT/d4r-accuracy.txt"

kernels/m/rrlite_dec2_4x4.hip

@@ -1,6 +1,8 @@
 // fast numerics: intermediate activations f16 (no FP8 re-quantisation), f32 accumulation through GEMM chains
-// RX 7700 XT: memory-clause scheduling at the 192 VGPR cap; byte-identical to the default build, gfx1101 only
-// d4r-build-flags-gfx1101: -mllvm -misched=gcn-max-memory-clause
+// RX 7700 XT: memory-clause scheduling at the 192 VGPR cap; byte-identical to the default build, gfx1101 only.
+// SWIN_PACKED_RESIDUAL (residual kept as already-rounded halves until the output projection): byte-identical,
+// -1.6% on the captured dec2 launch (0.768 -> 0.756 ms, 3 alternating reps of 100); enc1/enc2/dec1/tube: no gain
+// d4r-build-flags-gfx1101: -mllvm -misched=gcn-max-memory-clause -DSWIN_PACKED_RESIDUAL
 #define SWIN_NO_Q8
 #define SWIN_F32ACC
 #ifndef SWIN_MAX_VGPR

kernels/m/rrlite_enc3_tube_4x4.hip

@@ -8,9 +8,13 @@
 #define SWIN_MLP_R(NW) 2
 #endif
 // NGX can recycle a weights address for another tube block after feature
-// recreation. Address-only caching then serves stale packed weights. Prepare
-// each launch; the default single image is sufficient on the launch stream.
+// recreation. Address-only caching then serves stale packed weights, so
+// d4r_prep_key_at stays 0 (ZLUDA prepares every launch). MicroCUDA instead keys
+// preparation by allocation identity (d4r_prep_identity_at): one image per tube
+// weight set (6 per evaluation), re-prepared on any free/reallocation or write.
 #define SWIN_PREP_KEY_AT 0
+#define SWIN_PREP_IDENTITY_AT 1
+#define SWIN_PREP_SLOTS 8
 #include "swin_block.h"
 #ifndef SWIN_NW
 #define SWIN_NW 8

kernels/m/swin_block.h

@@ -1050,6 +1050,12 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
 #ifndef SWIN_PREP_KEY_AT
 #define SWIN_PREP_KEY_AT 1 // weights pointer at offset 0; 0 forces preparation on every launch
 #endif
+#ifndef SWIN_PREP_IDENTITY_AT
+// weights pointer at offset (value - 1) for runtimes that key preparation by allocation identity (pointer +
+// allocation generation + write epoch: MicroCUDA), never by address alone; 0 = no such key. Runtimes that only
+// know d4r_prep_key_at (ZLUDA) ignore it.
+#define SWIN_PREP_IDENTITY_AT 0
+#endif
 
 // Optional weight-image slots for several weight sets with stable address identity. The ZLUDA
 // hook can run prep only for a weights pointer it has not seen (d4r_prep_key_at / d4r_prep_key_slots);
@@ -1090,13 +1096,14 @@ template <int S> __device__ __forceinline__ int prep_slot_claim(uint64_t* keys,
     using NAME##_L = SwinLayout<C, NH, NPM, CIN>;                                                                       \
     constexpr int NAME##_S = SWIN_PREP_SLOTS;                                                                           \
     __device__ wslot g_w16[(NAME##_S + 1) * NAME##_L::TOTAL];                                                           \
-    __device__ uint64_t g_prep_keys[NAME##_S + 1];                                                                      \
+    extern "C" __device__ __attribute__((used)) uint64_t g_prep_keys[NAME##_S + 1] = {};                                \
     __constant__ SwinDescs<C, NH, NPM, CIN> g_descs = make_descs<C, NH, NPM, CIN>();                                    \
     extern "C" __device__ __attribute__((used)) uint32_t d4r_block_z = NWAVES;                                         \
     extern "C" __device__ __attribute__((used)) uint32_t d4r_prep_blocks = (NAME##_L::TOTAL + 127) / 128;              \
     /* enc3 disables address-only reuse: a recreated feature may recycle a weights address */                        \
     extern "C" __device__ __attribute__((used)) uint32_t d4r_prep_key_at = SWIN_PREP_KEY_AT;                             \
     extern "C" __device__ __attribute__((used)) uint32_t d4r_prep_key_slots = NAME##_S;                                \
+    extern "C" __device__ __attribute__((used)) uint32_t d4r_prep_identity_at = SWIN_PREP_IDENTITY_AT;                 \
     extern "C" __global__ void __launch_bounds__(128) NAME##_prep(PARAMS p)                                             \
     {                                                                                                                   \
         const int idx = blockIdx.x * 128 + threadIdx.x;                                                                 \

kernels/m/swin_common.h

@@ -337,6 +337,20 @@ __device__ __forceinline__ sop lds16(const half_t* p)
 // e4m3 bytes of an f16 operand whose values are e4m3 values (K order kept: byte j = element j)
 __device__ __forceinline__ gop to_fp8(sop v)
 {
+#if defined(SWIN_FP8_HWCVT) && defined(__GFX12__)
+    // Opt-in, UNVERIFIED ON HARDWARE: the values already are e4m3 values (finite, |x| <= 448), so the hardware
+    // conversion has nothing to round or saturate; byte j = element j as below. Open question before enabling
+    // it by default: whether e4m3 subnormals (k * 2^-9) survive the gfx12 denormal mode in effect.
+    uint32_t w[2];
+#pragma unroll
+    for (int j = 0; j < 2; ++j)
+    {
+        int packed = __builtin_amdgcn_cvt_pk_fp8_f32((float)v[4 * j], (float)v[4 * j + 1], 0, false);
+        packed = __builtin_amdgcn_cvt_pk_fp8_f32((float)v[4 * j + 2], (float)v[4 * j + 3], packed, true);
+        w[j] = (uint32_t)packed;
+    }
+    return (u2v){w[0], w[1]};
+#endif
     uint32_t c[4];
 #pragma unroll
     for (int j = 0; j < 4; ++j)

kernels/native/ctile.py

@@ -224,8 +224,15 @@ def apply(tr, tokenize):
     body = tr.body
     items = parse(body, tokenize)
     # color texture handle: ld.param.u64 [%rd1+-56] (param+200, %rd10 / %rd69); taps use it or its mov aliases
+    # (some variants, e.g. mvlo, re-materialize the parameter base: %rdN = param_0 + 256, the same address as %rd1)
+    param0 = {it['dst'][0] for it in items if it['kind'] == 'stmt' and it['op'] == 'mov.b64'
+              and re.search(r',\s*\w+_param_0\s*$', it['rest'])}
+    bases = {it['dst'][0] for it in items if it['kind'] == 'stmt' and it['op'] == 'add.s64'
+             and len(it['src']) == 1 and it['src'][0] in param0 and re.search(r',\s*256\s*$', it['rest'])}
+    if '%rd1' not in bases:
+        raise CtileErr(f'parameter base %rd1 not found ({sorted(bases)})')
     color = {it['dst'][0] for it in items if it['kind'] == 'stmt' and it['op'] == 'ld.param.u64'
-             and re.search(r'\[%rd1\+-56\]', it['rest'])}
+             and any(re.search(r'\[' + re.escape(b) + r'\+-56\]', it['rest']) for b in bases)}
     aliases = set(color)
     for it in items:
         if it['kind'] == 'stmt' and it['op'] == 'mov.u64' and len(it['src']) == 1 and it['src'][0] in color:

kernels/native/in_lds2reg.py

@@ -0,0 +1,192 @@
+#!/usr/bin/env python3
+"""in_lds2reg.py IN.hip > OUT.hip: a hiluma input kernel (any variant, translated by ptx2hip.py --ifconv --sink)
+with its per-thread scratch kept in registers. Generalizes in_fuse.py + in_regs.hip, which matched one variant's
+register numbers and its output address formula (mvlo variants store to a different address).
+
+The kernel stays the translated kernel: same thread ids, same arithmetic, same final stores. Only two
+per-thread scratch areas change, both value-preserving:
+- LDS: 4 planes x 3 floats at S + 3072 p + 4 j, where S = T * 12 + base is the thread's own slot, written in
+  straight-line code before any branch, read back at fixed offsets and once as plane k (address
+  (k << 8) * 12 + S) -> 12 registers + a 4-way select. The shared array is removed (12 KB LDS per block).
+- the 32-byte local depot: 4 float pairs at depot + 0..28 written before any branch, read once as the pair at
+  depot + 8 k -> 8 registers + a 4-way select (it compiled to scratch memory).
+The plane index k must be the same register for both reads, and is proven to lie in 0..3 by evaluating its
+defining expression (which may depend only on workgroup/workitem ids and constants) for every thread of
+grids up to 65536 x 16384 threads. Anything else is refused.
+"""
+import re
+import sys
+
+import numpy as np
+
+src = open(sys.argv[1]).read()
+lines = src.split('\n')
+k0 = next(i for i, l in enumerate(lines) if l.startswith('extern "C" __attribute__((global))'))
+assert lines[k0 + 1] == '{'
+body_start = k0 + 2
+first_label = next(i for i in range(body_start, len(lines)) if re.match(r'^L__BB0_\d+:;$', lines[i]))
+
+defs = {}
+for i, l in enumerate(lines):
+    m = re.match(r'^    ((?:r|rd)\d+) = (.*);$', l)
+    if m:
+        defs.setdefault(m.group(1), []).append((i, m.group(2)))
+
+
+def single(reg):
+    d = defs.get(reg, [])
+    assert len(d) == 1, f'{reg}: {len(d)} definitions'
+    return d[0]
+
+
+# ---- LDS slot S, base, plane address P = (k << 8) * 12 + S
+ST = re.compile(r'^    \*\(AS3 float\*\)\(uintptr_t\)\((r\d+) \+ (\d+)u\) = \(float\)\((f\d+)\);$')
+LD = re.compile(r'^    (f\d+) = \*\(AS3 float\*\)\(uintptr_t\)\((r\d+) \+ (\d+)u\);$')
+stores = [(i, ST.match(l)) for i, l in enumerate(lines) if ST.match(l)]
+slot = {m.group(1) for _, m in stores}
+assert len(slot) == 1, slot
+S = slot.pop()
+_, sdef = single(S)
+m = re.fullmatch(r'\(uint32_t\)\((r\d+) \* 12u \+ (r\d+)\)', sdef)
+assert m, sdef
+T, BASE = m.groups()
+_, bdef = single(BASE)
+shared = re.fullmatch(r'\(uint32_t\)\(\(uint32_t\)\(uintptr_t\)\(AS3 uint8_t\*\)(\w+)\)', bdef)
+assert shared, bdef
+shared_name = shared.group(1)
+offsets = {}
+for i, m in stores:
+    assert i < first_label, f'LDS store inside control flow: {lines[i]}'
+    off = int(m.group(2))
+    assert off % 4 == 0 and off // 3072 < 4 and (off % 3072) // 4 < 3 and off not in offsets, lines[i]
+    offsets[off] = m.group(3)
+assert len(offsets) == 12, sorted(offsets)
+loads = [(i, LD.match(l)) for i, l in enumerate(lines) if LD.match(l)]
+plane_regs = {m.group(2) for _, m in loads if m.group(2) != S}
+assert len(plane_regs) == 1, plane_regs
+P = plane_regs.pop()
+_, pdef = single(P)
+m = re.fullmatch(r'\(uint32_t\)\((r\d+) \* 12u \+ ' + S + r'\)', pdef)
+assert m, pdef
+_, shdef = single(m.group(1))
+m = re.fullmatch(r'\(uint32_t\)\((r\d+) << \(8u & 31\)\)', shdef)
+assert m, shdef
+K = m.group(1)
+for i, m in loads:
+    if m.group(2) == P:
+        assert m.group(3) == '0', lines[i]
+
+# ---- local depot
+DBASE = {r for r, d in defs.items() if any(e == '(uint64_t)(l2_SPL + (uint64_t)0ull)' for _, e in d)}
+DST = re.compile(r'^    \*\(AS5 float\*\)\(uintptr_t\)\(uint32_t\)\((rd\d+) \+ (\d+)\) = \(float\)\((f\d+)\);$')
+DLD = re.compile(r'^    (f\d+) = \*\(AS5 float\*\)\(uintptr_t\)\(uint32_t\)\((rd\d+) \+ ([04])\);$')
+dstores = {}
+for i, l in enumerate(lines):
+    m = DST.match(l)
+    if m:
+        assert m.group(1) in DBASE and i < first_label, l
+        off = int(m.group(2))
+        assert off % 4 == 0 and off < 32 and off not in dstores, l
+        dstores[off] = m.group(3)
+assert len(dstores) == 8, sorted(dstores)
+dloads = [(i, DLD.match(l)) for i, l in enumerate(lines) if DLD.match(l)]
+assert len(dloads) == 2, len(dloads)
+dl = {m.group(2) for _, m in dloads}
+assert len(dl) == 1, dl
+DL = dl.pop()
+_, ddef = single(DL)
+m = re.fullmatch(r'\(uint64_t\)\((rd\d+) \+ (rd\d+)\)', ddef)
+assert m and m.group(1) in DBASE, ddef
+_, kdef = single(m.group(2))
+assert kdef == f'(uint64_t)((int64_t)(int32_t){K} * (int64_t)(int32_t)8)', kdef
+assert all(re.search(r'\bAS5\b', lines[i]) is None or i in [j for j, _ in dloads] or DST.match(lines[i])
+           or 'l2_SPL' in lines[i] or '__local_depot0' in lines[i] for i in range(len(lines))), 'other depot use'
+assert sum('AS3' in l for l in lines) == len(stores) + len(loads) + 1, 'other LDS use'  # + the base address
+
+# ---- prove k in 0..3 for every thread: evaluate its expression over the thread-id range
+IDS = {'workgroup_id_x': 'wgx', 'workitem_id_x': 'wix', 'workgroup_id_y': 'wgy', 'workitem_id_y': 'wiy'}
+
+
+def to_np(reg, env, seen=()):
+    assert reg not in seen
+    if reg in env:
+        return env[reg]
+    _, e = single(reg)
+    m = re.fullmatch(r'\(uint32_t\)\(\(uint32_t\)__builtin_amdgcn_(\w+)\(\)\)', e)
+    if m:
+        env[reg] = env[IDS[m.group(1)]]
+        return env[reg]
+    expr = e
+    for r in sorted(set(re.findall(r'\br\d+\b', e)), key=len, reverse=True):
+        to_np(r, env, seen + (reg,))
+    expr = re.sub(r'\(uint32_t\)', '', expr)
+    expr = re.sub(r'\(int32_t\)(r\d+)', r's32(\1)', expr)
+    expr = re.sub(r'(\d+)u\b', r'np.uint32(\1)', expr)
+    expr = re.sub(r'\(np\.uint32\((\d+)\) & np\.uint32\(31\)\)', r'np.uint32(\1)', expr)
+    assert re.fullmatch(r'[\w\s().,+\-*&|<>]*', expr), e
+    env[reg] = u32(eval(expr, {'np': np, 's32': s32, 'bfi32': bfi32}, env))
+    return env[reg]
+
+
+def u32(x):
+    return np.asarray(x).astype(np.int64).astype(np.uint32) if np.asarray(x).dtype.kind == 'i' else np.asarray(x, np.uint32)
+
+
+def s32(x):
+    return x.view(np.int32) if isinstance(x, np.ndarray) else np.int32(x)
+
+
+def bfi32(a, b, offset, width):
+    mask = np.uint32(((1 << int(width)) - 1) << int(offset))
+    return (u32(a) << np.uint32(offset)) & mask | (u32(b) & ~mask)
+
+
+gx = np.arange(65536, dtype=np.uint32)
+for gy0 in range(0, 16384, 256):
+    gy = np.arange(gy0, gy0 + 256, dtype=np.uint32)[:, None]
+    env = {'wgx': np.broadcast_to(gx >> 4, (256, 65536)).copy(), 'wix': np.broadcast_to(gx & 15, (256, 65536)).copy(),
+           'wgy': np.broadcast_to(gy >> 4, (256, 65536)).copy(), 'wiy': np.broadcast_to(gy & 15, (256, 65536)).copy()}
+    k = to_np(K, env)
+    assert k.max() <= 3, f'plane index {K} reaches {k.max()} at rows {gy0}..'
+
+# ---- emit
+out = []
+sel = lambda name: (f'{K} == 0u ? {name(0)} : {K} == 1u ? {name(1)} : {K} == 2u ? {name(2)} : {name(3)}')
+for i, l in enumerate(lines):
+    if re.match(rf'^__attribute__\(\(shared\)\).*\b{shared_name}\[', l):
+        continue  # the LDS scratch array
+    if i == k0 + 1:
+        out.append(l)
+        out.append('    float ' + ', '.join(f'shv{p}_{j}' for p in range(4) for j in range(3)) + ';')
+        out.append('    float ' + ', '.join(f'dpv{p}_{j}' for p in range(4) for j in range(2)) + ';')
+        continue
+    if i == single(BASE)[0]:
+        out.append(f'    {BASE} = 0u; // (LDS scratch kept in registers)')
+        continue
+    m = ST.match(l)
+    if m:
+        off = int(m.group(2))
+        out.append(f'    shv{off // 3072}_{(off % 3072) // 4} = {m.group(3)};')
+        continue
+    m = LD.match(l)
+    if m:
+        if m.group(2) == P:
+            out.append(f'    {m.group(1)} = {sel(lambda p: f"shv{p}_0")}; // {K} in 0..3 (proven)')
+        else:
+            off = int(m.group(3))
+            assert off in offsets, f'load before store: {l}'
+            out.append(f'    {m.group(1)} = shv{off // 3072}_{(off % 3072) // 4};')
+        continue
+    m = DST.match(l)
+    if m:
+        off = int(m.group(2))
+        out.append(f'    dpv{off // 8}_{(off % 8) // 4} = {m.group(3)};')
+        continue
+    m = DLD.match(l)
+    if m:
+        j = int(m.group(3)) // 4
+        out.append(f'    {m.group(1)} = {sel(lambda p: f"dpv{p}_{j}")};')
+        continue
+    out.append(l)
+print(f'// Generated by kernels/native/in_lds2reg.py from {sys.argv[1].split("/")[-1]} (do not edit).')
+print('\n'.join(out), end='')

kernels/native/rewrite_output.py

@@ -43,6 +43,12 @@ def rewrite(source):
     live_out = definitions & uses_outside
     if live_out != {'r215', 'r217'}:
         raise ValueError(f'tile register contract changed: {sorted(live_out)}')
+    # The replacement passes r3, r4, r7, r8 and f1 by name: the region's inputs must be exactly the ones it was
+    # written for (other flag variants of the kernel reuse this region; the checks above alone only cover outputs).
+    region_text = ''.join(lines[i] for i in sorted(removed))
+    live_in = set(REGISTER.findall(region_text)) - definitions
+    if live_in != {'r3', 'r4', 'r5', 'r6', 'r7', 'r8', 'rd1', 'f1', 'f4'}:
+        raise ValueError(f'tile input contract changed: {sorted(live_in)}')
     replacement = '''L__BB0_124:;
     output_tile_fill(KARGP, (int32_t)r3, (int32_t)r4, (int32_t)r7, (int32_t)(r8 + 1u),
                      f1, (AS3 uint8_t*)sharedColor, (AS3 uint8_t*)sharedDepth);

kernels/tools/dump_runner.cpp

@@ -133,17 +133,87 @@ int main(int argc, char** argv)
     size_t argSize = args.size();
     void* cfg[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, args.data(), HIP_LAUNCH_PARAM_BUFFER_SIZE, &argSize,
                    HIP_LAUNCH_PARAM_END};
+    // phase kernels NAME_post1..NAME_post8, launched after the main kernel with the same parameters, exactly as
+    // the ZLUDA hook and MicroCUDA do: u32 globals d4r_postK_grid_x/_grid_y (0 = the launch's grid; if either is
+    // set, grid z is 1) and d4r_postK_block_x/_y/_z (0 = 32/1/1)
+    struct Post { hipFunction_t fn; unsigned gx, gy, gz, bx, by, bz; };
+    std::vector<Post> posts;
+    auto global_u32 = [&](const std::string& name) -> uint32_t {
+        hipDeviceptr_t p;
+        size_t n = 0;
+        uint32_t v = 0;
+        if (hipModuleGetGlobal(&p, &n, mod, name.c_str()) == hipSuccess && n == 4 && hipMemcpyDtoH(&v, p, 4) != hipSuccess)
+            v = 0;
+        return v;
+    };
+    for (int k = 1; k <= 8; ++k)
+    {
+        Post q{};
+        const std::string name = std::string(argv[2]) + "_post" + std::to_string(k);
+        if (hipModuleGetFunction(&q.fn, mod, name.c_str()) != hipSuccess)
+            break;
+        const std::string g = "d4r_post" + std::to_string(k) + "_";
+        const uint32_t px = global_u32(g + "grid_x"), py = global_u32(g + "grid_y");
+        q.gx = px ? px : grid[0];
+        q.gy = py ? py : grid[1];
+        q.gz = (px || py) ? 1 : grid[2];
+        q.bx = global_u32(g + "block_x");
+        q.by = global_u32(g + "block_y");
+        q.bz = global_u32(g + "block_z");
+        if (!q.bx) q.bx = 32;
+        if (!q.by) q.by = 1;
+        if (!q.bz) q.bz = 1;
+        posts.push_back(q);
+    }
+    if (!posts.empty())
+        printf("%s: %zu post phase(s)\n", argv[2], posts.size());
     const bool prepOnce = getenv("PREP_ONCE") != nullptr;
     int launches = 0;
-    auto launch = [&]() {
+    // PHASE_TIMES=1: an event before and after every phase (prep, main, post1..) for per-phase GPU times
+    const bool phaseTimes = getenv("PHASE_TIMES") != nullptr;
+    std::vector<hipEvent_t> marks(posts.size() + 3);
+    for (auto& m : marks)
+        CHECK(hipEventCreate(&m));
+    std::vector<std::vector<float>> phaseMs(posts.size() + 2);
+    bool recordPhases = false;
+    auto launch = [&]() -> hipError_t {
+        bool prepped = false;
+        if (recordPhases) (void)hipEventRecord(marks[0]);
         if (prep != nullptr && !(prepOnce && launches++ > 0))
         {
             hipError_t e = hipModuleLaunchKernel(prep, prepGrid, 1, 1, 128, 1, 1, 0, nullptr, nullptr, cfg);
             if (e != hipSuccess)
                 return e;
+            prepped = true;
         }
-        return hipModuleLaunchKernel(fn, grid[0], grid[1], grid[2], block[0], block[1], block[2], shared, nullptr,
-                                     nullptr, cfg);
+        if (recordPhases) (void)hipEventRecord(marks[1]);
+        hipError_t e = hipModuleLaunchKernel(fn, grid[0], grid[1], grid[2], block[0], block[1], block[2], shared,
+                                             nullptr, nullptr, cfg);
+        if (e != hipSuccess)
+            return e;
+        if (recordPhases) (void)hipEventRecord(marks[2]);
+        for (size_t k = 0; k < posts.size(); ++k)
+        {
+            const Post& q = posts[k];
+            e = hipModuleLaunchKernel(q.fn, q.gx, q.gy, q.gz, q.bx, q.by, q.bz, 0, nullptr, nullptr, cfg);
+            if (e != hipSuccess)
+                return e;
+            if (recordPhases) (void)hipEventRecord(marks[3 + k]);
+        }
+        if (recordPhases)
+        {
+            if (hipError_t w = hipEventSynchronize(marks.back()); w != hipSuccess)
+                return w;
+            for (size_t k = 0; k < phaseMs.size(); ++k)
+            {
+                if (k == 0 && !prepped)
+                    continue;
+                float t;
+                if (hipEventElapsedTime(&t, marks[k], marks[k + 1]) == hipSuccess)
+                    phaseMs[k].push_back(t);
+            }
+        }
+        return hipSuccess;
     };
     CHECK(launch());
     CHECK(hipDeviceSynchronize());
@@ -178,6 +248,22 @@ int main(int argc, char** argv)
         CHECK(hipEventElapsedTime(&t, a, b));
         ms.push_back(t);
     }
+    if (phaseTimes)
+    {
+        recordPhases = true;
+        for (int i = 0; i < iters; ++i)
+            CHECK(launch());
+        CHECK(hipDeviceSynchronize());
+        for (size_t k = 0; k < phaseMs.size(); ++k)
+        {
+            auto& v = phaseMs[k];
+            if (v.empty())
+                continue;
+            std::sort(v.begin(), v.end());
+            const std::string phase = k == 0 ? "prep" : k == 1 ? "main" : "post" + std::to_string(k - 1);
+            printf("%s %s: median %.4f ms, min %.4f ms\n", argv[2], phase.c_str(), v[v.size() / 2], v[0]);
+        }
+    }
     std::sort(ms.begin(), ms.end());
     if (!ms.empty())
         printf("%s: median %.4f ms, min %.4f ms over %d launches\n", argv[2], ms[ms.size() / 2], ms[0], iters);