diff --git a/docs/architecture.md b/docs/architecture.md
index ca078ed..00ce2e0 100644
--- a/docs/architecture.md
+++ b/docs/architecture.md
@@ -112,6 +112,8 @@ OptiScaler 0.9.4 checks a relative `NvngxPath` against the process working direc
 
 ## Logs
 
+For startup stalls, GPU resets and model-build failures, see [Linux startup diagnostics](linux-startup-diagnostics.md).
+
 The launcher writes each run's logs to `~/.cache/d4r-dlss-captures/game-<exe>-<time>/`. `d4r_nvngx.log` holds the shim's log. With `[Debug] Profile = true` it measures stage timings every frame and records sampled `D4R_PROFILE` lines. `ProfileLogEvery = 60` (the default) logs the first five frames and then every 60th frame; set it to `1` to log every frame, or another positive integer for a different interval. Restart the game after changing it. `D4R_PROFILE_LOG_EVERY` overrides the INI setting. FSR RR stage summaries average that many completed GPU samples; native Vulkan profiling uses the same log interval. Sampling reduces log writes while retaining profiling measurements and per-frame throughput intervals. `D4R_PROFILE_LOG=0` suppresses timing rows and FSR RR summaries.
 
 Output completion uses event queries with 200 us NtDelayExecution sleeps by default. `D4R_SHIM_BLOCKING_SYNC = 0`
diff --git a/docs/linux-startup-diagnostics.md b/docs/linux-startup-diagnostics.md
new file mode 100644
index 0000000..aabe989
--- /dev/null
+++ b/docs/linux-startup-diagnostics.md
@@ -0,0 +1,65 @@
+# Linux startup and GPU compatibility diagnostics
+
+Keep one complete runtime together when comparing GPUs: the shim, nvcuda bridge, ZLUDA, patched D3D12 pair and
+native kernels. New shim/bridge logs include a source build identifier. Record file SHA-256 hashes as well when
+reporting failures; an old log format is not a controlled comparison with a new build.
+
+## Capture a failing run
+
+Add these entries to the existing `[Env]` section of `d4r/d4r.ini` (or add that section once):
+
+```ini
+[Env]
+D4R_CUDA_VERBOSE = 1
+D4R_CUDA_TRACE = Z:\tmp\d4r-cuda-trace.log
+```
+
+Use no shell quotes in INI values. This saves the bridge trace to `/tmp/d4r-cuda-trace.log` on the machine running
+the game. The bridge also writes to stderr. ZLUDA's own compiler/native-override messages require Proton's stderr
+log separately. Verbose tracing can affect timing; remove it after collecting the run.
+
+Save the CUDA trace and `d4r_nvngx.log` separately for each successful/failed launch. The trace appends, while the
+default shim log is replaced on a fresh process. For a driver recovery or corrupted display, also save the kernel
+journal covering that run:
+
+```sh
+journalctl -b -k -o short-monotonic --no-pager > d4r-kernel.log
+```
+
+`CUDA_CREATE_BEGIN`, `CUDA_CREATE_WORKER_BEGIN`, `CUDA_CREATE_NGX_BEGIN/END` and `CUDA_CREATE_END` bracket feature
+creation. A worker still waiting logs progress every 15 seconds without abandoning resources still in use.
+`CUDA_BEGIN api=...` is recorded before module loading and, with verbose tracing, uploads, allocations, launches
+and synchronization. The following ordinary API line records its result. A launch return proves submission, not
+correct pixels or completed GPU work. Bridge records include a monotonic timestamp and thread ID.
+
+## GPU selection
+
+The game adapter's PCI address is matched to the compute device. Other compute devices are not given its LUID,
+and native kernels use the matching KFD topology entry rather than the largest GPU. When no PCI mapping is
+available, one visible compute GPU is unambiguous; multiple devices are refused instead of guessed. Selection
+failures report the rendering PCI address and visible compute devices. `D4R_GPU_ARCH` cannot override a known
+target with an incompatible architecture.
+
+RDNA2 keeps ZLUDA's translation fallback and its existing software-matrix K kernels. It is not assigned a native
+FP8/WMMA Ray Reconstruction model. This does not establish that every translated model is stable or fast there.
+
+## Model installer
+
+The installer builds portable FP8 RR-D/E/F models plus separate `rr-d-gfx11` and `rr-e-gfx11` variants for compatible
+RDNA3 devices. The runtime checks model arithmetic and physical matrix capabilities before constructing RR
+pipelines. Unsupported models fall back to the existing CUDA path; advertised NGX availability is not a runtime
+validation. Engine `auto` has the same model-directory detection in the portable shim and Python launcher.
+
+Model build logs are printed as they start and progress every 15 seconds. Each child has a 900-second deadline;
+`D4R_MODEL_TIMEOUT=SECONDS` changes it. The capture harness also limits individual GPU fence waits to 60 seconds
+(`D4R_HARNESS_GPU_TIMEOUT_MS`, 1–600000, changes that bound). A failed process or incomplete/nonfinite capture
+cannot be accepted solely because a replay filename exists.
+
+Diagnostics are retained under `${XDG_CACHE_HOME:-$HOME/.cache}/d4r/model-builds/build.*`; `D4R_MODEL_LOG_DIR`
+changes that root. Successful builds retain logs and remove temporary weights/captures/prefixes. Failed builds
+retain their workspace, including proprietary model data; keep those artifacts local and remove them when no
+longer needed. Model folders carry `files.sha256.json`; cached builds are checked before reuse, and compiler
+arithmetic options participate in the cache key. Old caches without these manifests rebuild once.
+
+These checks prevent known selection/setup failures and improve fault localization. They do not by themselves
+prove that an intermittent GPU reset or corrupted frame is fixed; that requires matching traces and GPU testing.
diff --git a/docs/native-kernels.md b/docs/native-kernels.md
index 1b640a6..7148754 100644
--- a/docs/native-kernels.md
+++ b/docs/native-kernels.md
@@ -44,7 +44,7 @@ The release is built this way, with hashes from DLSS 310.7.0 and 310.9.1. The PT
   - dec5 runs as four dependent phase kernels (`pwin_phase.h`, post kernels, needs patch 0008).
   - enc1 keeps one LDS buffer and reloads its input for the residual (`pwin_pos_recompute.h`); dec1 computes its Q, K and V chains two at a time.
   - Tried and not kept (slower or not exact): interleaved WMMA chains, a GELU lookup table, streaming attention, f16-accumulating WMMA, persistent or megakernel layers, phase kernels for the other deep layers, fast math.
-- **Texture kernels as native code:** `kernels/native/ptx2hip.py` translates a DLSS PTX kernel into HIP, each instruction as ZLUDA's lowering of it (`ptx_rt.h`, `native_rt.h`). `kernels/build.sh` uses it, at build time and from your own DLL, for two kernels of the fast set: `hiluma_engine_output_depthinv_mvhi_hdr_max_v2_rel` (`--ctile`, the path-B LDS tile as a function in `output_tile.h`, no unused depth prefetch, global-memory redirect stores) and `hiluma_engine_input_depthinv_mvhi_hdr_v2_rel` (`--ifconv --sink`, `in_fuse.py` moves its scratch arrays into registers, `in_regs.hip`, wave32). Both match ZLUDA's compile byte for byte in the harness. The generated sources are never stored. Other flag combinations of these kernels still use the PTX route above, and the accuracy set keeps ZLUDA's compile because the native code flushes f32 denormals.
+- **Texture kernels as native code:** `kernels/native/ptx2hip.py` translates a DLSS PTX kernel into HIP, each instruction as ZLUDA's lowering of it (`ptx_rt.h`, `native_rt.h`). `kernels/build.sh` uses it, at build time and from your own DLL, for two kernels of the fast set: `hiluma_engine_output_depthinv_mvhi_hdr_max_v2_rel` (`--ctile`, the path-B LDS tile as a function in `output_tile.h`, no unused depth prefetch, global-memory redirect stores) and `hiluma_engine_input_depthinv_mvhi_hdr_v2_rel` (`--ifconv --sink`, `in_fuse.py` moves its scratch arrays into registers, `in_regs.hip`, wave32). Both match ZLUDA's compile byte for byte in the harness. The generated sources are never stored. `cvt.rn.f16.f32` goes through `f2h_rn` (`ptx_rt.h`), whose empty asm stops LLVM from folding the conversion into the f32 operation that produced the value (`v_fma_mixlo/mixhi_f16`, which rounds the exact result to f16 once where PTX rounds to f32 and then to f16; that differed in 1 of ~29,000 random f16 operand triples on an RX 7700 XT). Before that, the output kernels left single 1-ulp values (3 history and 1 output value of a 960×540 → 1920×1080 frame) that the temporal history spread to 72 dB after 30 frames; since, every buffer the mvlo output kernel writes equals ZLUDA's compile over 12 frames, and the presented frames of all four HDR max variants are byte-identical over 30 frames (performance, quality 1129×636, ultra performance, DLAA, exposure changes, history resets, feature recreation). The input kernels compile to the same code as before. Other flag combinations of these kernels still use the PTX route above, and the accuracy set keeps ZLUDA's compile because the native code flushes f32 denormals.
 
 **DLSS 4.5, presets L/M (`kernels/m`, `rrlite_*`).**
 - **Network:** the Swin blocks of DLSS 4.5, whose weights are FP8. Their prep kernels expand the weights to f16 WMMA operands. The enc3 tube layer prepares on every launch: NGX can recycle a weights address for a different tube block after feature recreation, so persistent address-only reuse can return stale weights and corrupt the image after quality changes. Other layers retain their existing preparation policy.
@@ -263,9 +263,9 @@ kernels/build.sh all DIR    # everything into DIR (default kernels/out/native)
 D4R_PREFER_ACCURACY=1 kernels/build.sh all DIR/accuracy  # conservative versions of every family
 ```
 
-`D4R_GPU_ARCH` selects one target for this standalone build (default gfx1101). The release script builds the network layers for gfx1100–gfx1103 and gfx1200–gfx1201 by default and puts each set in its own directory, plus K's layers for the RDNA2 targets gfx1030–gfx1036. The network kernels use gfx11 or gfx12 WMMA through `kernels/common/wmma_layout.h`.
+`D4R_GPU_ARCH` selects one target for this standalone build (default gfx1101). The release script builds the network layers for gfx1100–gfx1103 and gfx1200–gfx1201 by default and puts each set in its own directory, plus K's layers and K's natively translated hiluma kernels (below) for the RDNA2 targets gfx1030–gfx1036. The network kernels use gfx11 or gfx12 WMMA through `kernels/common/wmma_layout.h`.
 
-RDNA2 has no WMMA instruction. On gfx103x (or with `D4R_WMMA_SOFT=1` on an RDNA3 target, for validation) `wm_mma` keeps the gfx11 register layout and does each 16×16×16 product with 64 `v_dot2c_f32_f16` instructions whose `row_share` DPP broadcast supplies the A row, after one `row_ror` per A dword that offsets half 1 by a row. Every accumulator sums its K pairs in the WMMA order: on the RX 7700 XT the result is bit-identical to the WMMA instruction for random 64-step chains (denormals included) and for K's whole network over a moving 24-frame harness sequence. There K takes 3.6 ms per 1080p frame against 1.48 ms with WMMA and 46.8 ms fully translated without WMMA (ZLUDA's dot-product fallback, what RDNA2 ran before). M's layers and the texture tails still call WMMA directly, so RDNA2 builds only K; LLVM does not fold a `v_mov_b32_dpp` into `v_dot2c`, so the fused instruction is inline assembly. Texture-kernel builds use `D4R_ZLUDA_EMIT` to generate code objects for the selected target without that GPU. Set `D4R_NATIVE_FP8=1` on gfx12 for the matching FP8 M and texture variants.
+RDNA2 has no WMMA instruction. On gfx103x (or with `D4R_WMMA_SOFT=1` on an RDNA3 target, for validation) `wm_mma` keeps the gfx11 register layout and does each 16×16×16 product with 64 `v_dot2c_f32_f16` instructions whose `row_share` DPP broadcast supplies the A row, after one `row_ror` per A dword that offsets half 1 by a row. Every accumulator sums its K pairs in the WMMA order: on the RX 7700 XT the result is bit-identical to the WMMA instruction for random 64-step chains (denormals included) and for K's whole network over a moving 24-frame harness sequence. There K takes 3.6 ms per 1080p frame against 1.48 ms with WMMA and 46.8 ms fully translated without WMMA (ZLUDA's dot-product fallback, what RDNA2 ran before). M's layers and the texture tails still call WMMA directly, so RDNA2 builds only K's layers and the hiluma kernels that `ptx2hip.py` translates (`kernels/build.sh tex` on a gfx103x target: the four HDR max output variants and the eight input variants; `m`, `l`, `rr` and `rrswin` are refused). Those contain no matrix products. The output kernels are wave32 there (RDNA2 issues wave64 in two passes; wave64 was RDNA3's FP32 dual-issue gain) and WGP mode: at 96 VGPRs a gfx103x SIMD holds 10 waves, and a 256-thread workgroup takes 2 of them per SIMD in WGP mode instead of 4, so 10 fit instead of 8. With RDNA2's occupancy imposed on the RX 7700 XT (a register clobber to 144 VGPRs) that measured 0.596 → 0.568 ms, same bytes; capping registers to 80 for 12 waves spills and was slower (0.567 against 0.553 ms with the occupancy imposed through the LDS size). On the RX 7700 XT at 960×540 → 1920×1080 the output kernel takes 0.79 ms translated and 0.44 ms native. ZLUDA's own gfx1032 compile of it uses 131 VGPRs (7 waves); the native one 96. RDNA2 timings are not measured yet; LLVM does not fold a `v_mov_b32_dpp` into `v_dot2c`, so the fused instruction is inline assembly. Texture-kernel builds use `D4R_ZLUDA_EMIT` to generate code objects for the selected target without that GPU. Set `D4R_NATIVE_FP8=1` on gfx12 for the matching FP8 M and texture variants.
 
 The release includes both fast and accuracy network sets. Full releases also include accuracy texture sets: supply them in `D4R_BUNDLE_TEX/accuracy/<target>` or set `D4R_ZLUDA_EMIT` so packaging builds them. Build accuracy sets into a separate directory; the builder refuses to mark a directory containing unmarked older binaries as accurate.
 
diff --git a/engine/README.md b/engine/README.md
index fbe9a43..33939b7 100644
--- a/engine/README.md
+++ b/engine/README.md
@@ -1093,6 +1093,32 @@ on the Vulkan shaders unless `[Engine] Network = hip`:
 A model built with no options was compared shader by shader with the explicitly optioned build of each row, and the
 exact switches still reproduce the captures (`build/rr/f/nv/defaults.sh`).
 
+### Byte-identical speed-ups (2026-10-19)
+
+RX 9070 XT, RADV 26.2.2, `d4r-m` 1280x720 -> 2560x1440 with `D4R_ENGINE_M_LAYER_STAMPS=1`. Every output frame, and
+everything `D4R_M_DUMP` writes, is byte-identical to the previous build (M and L, HDR and LDR, the three
+reconstruction tile tiers and the stress inputs).
+
+| change | stage | before -> after (us) |
+|---|---|---|
+| `layer_m.comp`: hardware FP8 saturation where a value cannot be infinite, and the V tiles in a buffer that is dead during the attention (more waves per SIMD) | frame | M 1583 -> 1566, L 2349 -> 2293 |
+| `post_m.comp`: the second history candidate reuses the first one's sample when both have the same vector | reconstruction, uniform motion | 325 -> 288 |
+| `post_m.comp`: both rows of a pixel pair take their Gaussian taps in one pass | reconstruction | 287 -> 265 (L 355 -> 320) |
+| `down_m.comp`: the horizontal pass reads its texels into registers, no shared source tile | downsample | L 102 -> 91, M 117 -> 113 |
+| `layer_m.comp`, `compile_m.py`: L's expansion stage at two windows per wave (`M_DEFS_DEC0L=-DWPW=1` for one) | expansion | 383 -> 366 |
+
+The network layers are separated by an execution-only barrier (`emptyStage`). A memory barrier there gives the same
+output and costs about 60 us (M 1520 -> 1592), so it stays.
+
+### `M_RESIDUAL_F32` (opt-in, 2026-10-19)
+
+`M_RESIDUAL_F32=1 python3 engine/compile_m.py ...` keeps the residual stream in f32 across the MLP's 32-channel chunks
+and rounds it to binary16 once after them, instead of after every chunk. It needs the relaxed build (not
+`M_EXACT=1`) and is off by default. 1280x720 -> 1440p: M 1521 -> 1410 us, L 2195 -> 2077. Against the default build,
+final RGB, peak 1: 43 - 46 dB (L 45 - 49) on smooth frames, 17 - 18 dB on noise. The exact build is as far from the
+default, so this moves no further from NVIDIA's arithmetic than `M_RELAXED` already does. Not checked against an
+NVIDIA card.
+
 ## Preset L (runtime and shim path exist; experimental)
 
 Preset L (render preset 12, NVIDIA's model for Ultra Performance) is M's pipeline with its own input, expansion and
diff --git a/engine/compile_m.py b/engine/compile_m.py
index b35e38c..a5293bf 100644
--- a/engine/compile_m.py
+++ b/engine/compile_m.py
@@ -32,6 +32,13 @@ from test_m import pack_layer, defines
 EXACT = os.environ.get('M_EXACT') == '1'
 FAST = [] if EXACT or '-DM_RELAXED' in os.environ.get('M_DEFS', '').split() else ['-DM_RELAXED']
 FAST_WPW = [] if EXACT or os.environ.get('M_WPW') or 'WPW=' in os.environ.get('M_DEFS', '') else ['-DWPW=2']
+# preset L's expansion stage (dec0l) two windows per wave as well; its input stage (enc0l) is slower that way
+L_WPW = [] if EXACT or 'WPW=' in os.environ.get('M_DEFS', '') + os.environ.get('M_DEFS_DEC0L', '') else ['-DWPW=2']
+# M_RESIDUAL_F32=1 (opt-in, not with M_EXACT=1): every layer_m.comp stage keeps the residual stream f32 across the MLP's
+# 32-channel chunks, one binary16 rounding after them instead of one per chunk (engine/README.md)
+RESF32 = ['-DM_RESIDUAL_F32'] if os.environ.get('M_RESIDUAL_F32') == '1' else []
+if RESF32 and EXACT:
+    raise SystemExit('M_RESIDUAL_F32=1 is part of the relaxed arithmetic: not with M_EXACT=1')
 SHADERS = {'enc1': 'm_enc1', 'enc2': 'm_enc2', 'enc3_tube': 'm_tube', 'dec2': 'm_dec2', 'dec1': 'm_dec1'}
 
 
@@ -105,7 +112,7 @@ def main():
     # the network in the arithmetic of ZLUDA's lowering of NVIDIA's PTX (TEXL): byte-identical to the CUDA backend
     # with NVIDIA's translated kernels in the harness (2026-10-11), where the native accuracy kernels' K32_MIX order
     # drifted to 36-41 dB through the FP8 tube
-    for name, out in SHADERS.items(): glsl('layer_m.comp', out + '.spv', defines(name) + ['-DTEXL=1'] + FAST + FAST_WPW)
+    for name, out in SHADERS.items(): glsl('layer_m.comp', out + '.spv', defines(name) + ['-DTEXL=1'] + FAST + FAST_WPW + RESF32)
     for s in ('enc0_m', 'dec0_m', 'post_m', 'down_m'): glsl(s + '.comp', s + '.spv', ['-DLDR=1'] if a.ldr and s != 'dec0_m' else [])
     glsl('down_m.comp', 'down_m_rgb10.spv', ['-DOUTPUT_RGB10=1'] + (['-DLDR=1'] if a.ldr else []))
     glsl('post_m.comp', 'post_m_q.spv', ['-DTILE_W=15', '-DTILE_H=15'] + (['-DLDR=1'] if a.ldr else []))
@@ -156,10 +163,10 @@ def compile_l(a, one, R):
         subprocess.run(['glslangValidator', '-V', '--target-env', 'vulkan1.3', *flags, str(HERE / src), '-o', str(a.output / out)], check=True, stdout=subprocess.DEVNULL)
     ldr = ['-DLDR=1'] if a.ldr else []
     # the network in the arithmetic of ZLUDA's lowering of NVIDIA's PTX (TEXL), as for preset M (L_NOTES.md)
-    for name, out in SHADERS.items(): glsl('layer_m.comp', out + '.spv', defines(name) + ['-DTEXL=1'] + FAST + FAST_WPW)
+    for name, out in SHADERS.items(): glsl('layer_m.comp', out + '.spv', defines(name) + ['-DTEXL=1'] + FAST + FAST_WPW + RESF32)
     for planes, suffix in ((0, ''), (1, '_planes')):
-        glsl('layer_m.comp', f'enc0_l{suffix}.spv', defines('enc0l') + ldr + [f'-DPLANES={planes}'] + FAST)
-        glsl('layer_m.comp', f'dec0_l{suffix}.spv', defines('dec0l') + [f'-DPLANES={planes}'] + FAST)
+        glsl('layer_m.comp', f'enc0_l{suffix}.spv', defines('enc0l') + ldr + [f'-DPLANES={planes}'] + FAST + RESF32)
+        glsl('layer_m.comp', f'dec0_l{suffix}.spv', defines('dec0l') + [f'-DPLANES={planes}'] + FAST + L_WPW + RESF32)
     # the reconstruction one render texel wider per tier than M's (4x4 taps): 11 (reconstruction/render >= 3),
     # 16 (the whole post_3_2 range), 12 (the compact tier: render/reconstruction <= 0.45, jitter within half a texel)
     for out, tile in (('post_l', 11), ('post_l_q', 16), ('post_l_c', 12)):
diff --git a/engine/compile_rr2.py b/engine/compile_rr2.py
index 90e07a5..0d8fed8 100644
--- a/engine/compile_rr2.py
+++ b/engine/compile_rr2.py
@@ -12,7 +12,7 @@ D4R_CUDA_REPLAY_DUMP_FILTER=rr2 D4R_CUDA_LAUNCH_STATS_FILTER=rr2 and preset hint
                        table of its PTX module; enc0-aux.bin is encoder 0's second table (argument 104)
   <stage>.args         the argument block of the capture's frame, pointers and handles cleared: the runtime keeps its
                        constants and fills what depends on the frame or the geometry (engine/runtime_rr2.inc)
-  exposure_f0.spv, exposure_f1.spv   engine/rr/exposure_f.comp
+  exposure_f0.spv, exposure_f1.spv   engine/rr/exposure_f.comp (exposure_f2.spv: both in one, 64-bit atomics)
   rr.txt               the manifest RModel reads ("D4RR0002")
 With --native, the stages engine/rr/rr2_m.comp implements (enc1 ... enc5, dec4 ... dec1, dec0's first kernel) are that
 shader with weights read from the DLL itself (engine/rr/rr2_pack.py) instead of the translation: bit-identical to
@@ -80,10 +80,18 @@ GROUP = int(os.environ.get('D4R_RR2_GROUP', '1'))
 # D4R_RR2_SPLIT=1 / 0 forces the gather pass / the single shader.
 SPLIT = os.environ.get('D4R_RR2_SPLIT', '0' if RELAXED >= 3 and 'enc0' in ONE_WAVE else '1') != '0'
 # The pooled histories as RGBA32F holding q / 1023, as ZLUDA backs them (16 bytes a pixel). D4R_RR2_POOLED=rgb10_a2
-# makes them real UNORM 10:10:10:2 images (0.1 ms faster at 4K), but the hardware's bilinear filter then rounds
-# differently: identical at 3x and 2x scaling, 44 - 51 dB on temporal frames at 1.5x (capFQ, 2026-10-13)
+# makes them real UNORM 10:10:10:2 images (0.1 ms faster at 4K, 70 MB less), but the hardware's bilinear filter then
+# rounds differently: identical at 3x and 2x scaling, 44 - 51 dB on temporal frames at 1.5x (capFQ, 2026-10-13); on
+# random-motion noise also at 2x (29 dB on the history frames, engine/rr/README.md)
 POOLED = os.environ.get('D4R_RR2_POOLED', 'rgba32f')
 SURFACES = {'dec0_1': {176: 'rgba8_snorm', 184: 'rgba8', 192: 'rgba8'}, 'post': {120: POOLED}, 'pooled': {8: POOLED}}
+# D4R_RR2_UNGUARDED=1 (opt-in, D4R_RR_RELAXED=3 only): the translated code (encoder 0's and decoder 0's per-pixel parts,
+# the output passes) without its exactness guards: no `precise` registers, no opaque zero OR-ed into binary16 stores and
+# conversions (ptx2rr.py --unguarded) and no NoContraction pass (spv_exact_convert.py) on those modules, so the driver may
+# contract, fold and fuse there as it does in the relaxed native stages (engine/rr/README.md)
+UNGUARDED = os.environ.get('D4R_RR2_UNGUARDED') == '1'
+if UNGUARDED and RELAXED < 3:
+    raise SystemExit('D4R_RR2_UNGUARDED=1 is for D4R_RR_RELAXED=3')
 
 
 def find_module(ptx_dir, kernel):
@@ -170,8 +178,8 @@ def main():
         cmd += [f'--surface={off}={fmt}' for off, fmt in SURFACES.get(stage, {}).items()]
         # the hybrid stages' and the output passes' modules go through spv_exact_convert.py, so their conversions need
         # only the lighter guard (ptx2rr.py --exact-convert)
-        exact = (bool(a.native) and stage in HYBRID) or not network
-        cmd += ['--exact-convert'] if exact else []
+        exact = ((bool(a.native) and stage in HYBRID) or not network) and not UNGUARDED
+        cmd += ['--exact-convert'] if exact else ['--unguarded'] if UNGUARDED else []
         subprocess.run(cmd, check=True, stderr=subprocess.DEVNULL)
         info = json.loads(generated.with_suffix('.json').read_text())
         weights = '-'
@@ -195,8 +203,9 @@ def main():
                 subprocess.run(capped + ['glslangValidator', '-V', '--target-env', 'vulkan1.3', *defs, f'-I{work}',
                                          f'-I{HERE / "rr"}', str(work / 'enc0_g.comp'), '-o', str(out / 'enc0g.spv')],
                                check=True, stdout=subprocess.DEVNULL)
-                subprocess.run([sys.executable, str(HERE / 'rr' / 'spv_exact_convert.py'), str(out / 'enc0g.spv')],
-                               check=True, stdout=subprocess.DEVNULL)
+                if not UNGUARDED:
+                    subprocess.run([sys.executable, str(HERE / 'rr' / 'spv_exact_convert.py'), str(out / 'enc0g.spv')],
+                                   check=True, stdout=subprocess.DEVNULL)
                 record = int((work / 'enc0_g.record').read_text())
             if RELAXED >= 3 and stage in ONE_WAVE and (split or GROUP == 1):
                 how = how + ['--one-wave']
@@ -206,8 +215,9 @@ def main():
             subprocess.run(capped + ['glslangValidator', '-V', '--target-env', 'vulkan1.3', *defs, f'-I{work}',
                                      f'-I{HERE / "rr"}', str(work / f'{stage}_m.comp'), '-o', str(out / f'{stage}.spv')],
                            check=True, stdout=subprocess.DEVNULL)
-            subprocess.run([sys.executable, str(HERE / 'rr' / 'spv_exact_convert.py'), str(out / f'{stage}.spv')],
-                           check=True, stdout=subprocess.DEVNULL)
+            if not UNGUARDED:
+                subprocess.run([sys.executable, str(HERE / 'rr' / 'spv_exact_convert.py'), str(out / f'{stage}.spv')],
+                               check=True, stdout=subprocess.DEVNULL)
         else:
             subprocess.run(capped + ['glslangValidator', '-V', '--target-env', 'vulkan1.3', str(generated), '-o',
                                      str(out / f'{stage}.spv')], check=True, stdout=subprocess.DEVNULL)
@@ -249,7 +259,7 @@ def main():
         if hybrid and split:   # the same block, bindings and regions; "rN": N words per window record, "s3": in three ranges
             lines.append(f'stage enc0g enc0g.spv 32 - {stage}.args {tex} {surf} {buf} {",".join(ptrs) or "-"} r{record} s3')
         print(f'{stage}: {kernel}, {len(args)} argument bytes' + (' (native block)' if hybrid else ''))
-    for p in (0, 1):
+    for p in (0, 1, 2):
         subprocess.run(['glslangValidator', '-V', '--target-env', 'vulkan1.3', f'-DPASS={p}',
                         str(HERE / 'rr' / 'exposure_f.comp'), '-o', str(out / f'exposure_f{p}.spv')],
                        check=True, stdout=subprocess.DEVNULL)
diff --git a/engine/down_m.comp b/engine/down_m.comp
index d445b0b..2a675af 100644
--- a/engine/down_m.comp
+++ b/engine/down_m.comp
@@ -4,6 +4,7 @@
 // five-tap filter whose weights depend on the output pixel's parity, limits each pass to the range of its taps
 // (no ringing), and undoes the range compression and the exposure.
 #extension GL_EXT_shader_explicit_arithmetic_types : require
+#extension GL_EXT_control_flow_attributes : require
 #extension GL_GOOGLE_include_directive : require
 #include "output_store.glsl"
 // LDR (NGX's *_ldr_* kernel): the packed image is the display-referred colour itself.
@@ -13,8 +14,7 @@
 // D4RM0003: the runtime dispatches 8x64 output pixels per workgroup.
 #define BLOCK_X 8
 #define BLOCK_Y 64
-#define TILE_W (3 * BLOCK_X / 2 + 4)
-#define TILE_H (3 * BLOCK_Y / 2 + 4)
+#define TILE_H (3 * BLOCK_Y / 2 + 4)     // source rows of the block
 layout(local_size_x = BLOCK_X, local_size_y = BLOCK_Y) in;
 layout(set = 0, binding = 0, std430) readonly buffer Params
 {
@@ -27,8 +27,19 @@ layout(set = 0, binding = 0, std430) readonly buffer Params
 layout(set = 0, binding = 1) uniform sampler2D colorTex;        // render resolution; its alpha is passed through
 layout(set = 0, binding = 2, FINAL_OUTPUT_FORMAT) uniform writeonly image2D outImg;
 layout(set = 0, binding = 3, std430) readonly buffer Src { uint w[]; } src;
-shared f16vec4 tile[TILE_W * TILE_H];       // source texels around the block
 shared f16vec4 rows[TILE_H * BLOCK_X];      // after the horizontal pass
+// The horizontal pass reads the source texels straight into registers, without a shared tile: an invocation filters
+// SEG output columns of one source row from the 3 SEG / 2 + 4 texels around them.
+#define SEG 4
+#define SPAN (3 * SEG / 2 + 4)
+f16vec3 unpack(uint w)
+{
+#if LDR
+    return f16vec3(vec3(float(w & 2047u) * (1.0 / 2047.0), float((w >> 11) & 2047u) * (1.0 / 2047.0), float(w >> 22) * (1.0 / 1023.0)));
+#else
+    return f16vec3(vec3(float(w & 2047u) * 0.0011968734906986356, float((w >> 11) & 2047u) * 0.0011968734906986356, float(w >> 22) * 0.002394916955381632));
+#endif
+}
 // weights of taps -2 .. +2 for an even output coordinate; an odd one uses them mirrored
 // as binary16 bit patterns: NGX rounds its f32 constants to the nearest binary16 (cvt.rn), while glslang folds
 // float16_t(<literal>) toward zero (W0 became -0.041321 instead of -0.041351: 6% of the captured frame's outputs
@@ -49,25 +60,27 @@ void main()
 {
     const ivec2 lid = ivec2(gl_LocalInvocationID.xy), pix = ivec2(gl_GlobalInvocationID.xy);
     const ivec2 base = ivec2(3 * BLOCK_X / 2, 3 * BLOCK_Y / 2) * ivec2(gl_WorkGroupID.xy) - 2;
-    for (int i = lid.y * BLOCK_X + lid.x; i < TILE_W * TILE_H; i += BLOCK_X * BLOCK_Y)
+    // the horizontal pass, SEG output columns x0 .. x0 + SEG - 1 of source row y per item
+    const int zero = int(gl_NumWorkGroups.z) - 1;
+    for (int item = lid.y * BLOCK_X + lid.x; item < TILE_H * (BLOCK_X / SEG); item += BLOCK_X * BLOCK_Y)
     {
-        const ivec2 at = clamp(base + ivec2(i % TILE_W, i / TILE_W), ivec2(0), p.srcSize - 1);
-        const uint w = src.w[at.y * p.srcSize.x + at.x];
-#if LDR
-        tile[i] = f16vec4(vec4(float(w & 2047u) * (1.0 / 2047.0), float((w >> 11) & 2047u) * (1.0 / 2047.0), float(w >> 22) * (1.0 / 1023.0), 0.0));
-#else
-        tile[i] = f16vec4(vec4(float(w & 2047u) * 0.0011968734906986356, float((w >> 11) & 2047u) * 0.0011968734906986356, float(w >> 22) * 0.002394916955381632, 0.0));
-#endif
+        const int y = item / (BLOCK_X / SEG), x0 = item % (BLOCK_X / SEG) * SEG;
+        const int row = clamp(base.y + y, 0, p.srcSize.y - 1) * p.srcSize.x;
+        f16vec3 t[SPAN];
+        [[unroll]] for (int k = 0; k < SPAN; ++k)
+            t[k] = unpack(src.w[row + clamp(base.x + 3 * x0 / 2 + k, 0, p.srcSize.x - 1)]);
+        [[unroll]] for (int q = 0; q < SEG; ++q)
+        {
+            // column x0 + q (as even as q) is centred on texel 3 (q / 2) + 2 (q & 1) + 2 of the item's. Its parity stays a
+            // run-time value: with constant weights the compiler shares products between neighbouring filters and
+            // splits their fused multiply-adds (30% of the outputs change)
+            const int o = 3 * (q >> 1) + 2 * (q & 1) + 2;
+            rows[y * BLOCK_X + x0 + q].rgb = filter5(t[o - 2], t[o - 1], t[o], t[o + 1], t[o + 2], ((q + zero) & 1) != 0);
+        }
     }
     memoryBarrierShared(); barrier();
     const ivec2 c = ivec2(floor((vec2(pix) + 0.5) * 1.5)) - base;
-    for (int y = lid.y; y < TILE_H; y += BLOCK_Y)
-    {
-        const int o = y * TILE_W + c.x;
-        rows[y * BLOCK_X + lid.x].rgb = filter5(tile[o - 2].rgb, tile[o - 1].rgb, tile[o].rgb, tile[o + 1].rgb, tile[o + 2].rgb, (pix.x & 1) != 0);
-    }
-    memoryBarrierShared(); barrier();
-    // After both shared barriers: borrowed images may extend beyond the feature's true output rectangle.
+    // After the shared barrier: borrowed images may extend beyond the feature's true output rectangle.
     // M accepts even output dimensions and srcSize = outputSize/2*3, so this bounds check is exact.
     if (3 * pix.x >= 2 * p.srcSize.x || 3 * pix.y >= 2 * p.srcSize.y) return;
     const int o = c.y * BLOCK_X + lid.x;
diff --git a/engine/game.cpp b/engine/game.cpp
index c983b35..cc22d14 100644
--- a/engine/game.cpp
+++ b/engine/game.cpp
@@ -22,7 +22,7 @@ bool gameMotionFormat(VkFormat format) {
 }
 bool gameDepthFormat(VkFormat format, VkImageAspectFlags aspect) {
     if (aspect & VK_IMAGE_ASPECT_DEPTH_BIT) return depthAspect(format);
-    return format == VK_FORMAT_R32_SFLOAT || format == VK_FORMAT_R16_SFLOAT || format == VK_FORMAT_R16_UNORM;
+    return format == VK_FORMAT_R32_SFLOAT || format == VK_FORMAT_R32_UINT || format == VK_FORMAT_R16_SFLOAT || format == VK_FORMAT_R16_UNORM;
 }
 bool gameOutputFormat(VkFormat format) { return gameColorFormat(format); }
 
diff --git a/engine/game_rr.cpp b/engine/game_rr.cpp
index 79aa379..f3a9bac 100644
--- a/engine/game_rr.cpp
+++ b/engine/game_rr.cpp
@@ -1,5 +1,6 @@
 #include "game_rr.h"
 #include "game_internal.h"
+#include <algorithm>
 #include <cstdlib>
 #include <cstring>
 #include <fstream>
@@ -15,6 +16,12 @@ struct GameDenoiser::Impl {
     game_detail::CanonicalInputs fallback;                 // colour, motion, depth
     std::array<game_detail::Plane, 8> guides;              // the guides and optional inputs, when one cannot be sampled directly
     bool started = false;
+    // Per-input ring-buffer detection: tracks distinct views for each of the 11 inputs
+    // (color, motion, depth, 8 guides). When an input exceeds the threshold, it is
+    // copied to a canonical plane so the engine sees a stable set of views.
+    static constexpr int NUM_INPUTS = 11;
+    std::array<std::vector<VkImageView>, NUM_INPUTS> seenViews;
+    std::array<bool, NUM_INPUTS> needsCopy = {};
     // debugging: D4R_ENGINE_RR_DUMP=DIR copies every engine image after each frame into a host-visible buffer and
     // writes it at the next frame's recording (by then the game has waited for the frame), as DIR/F-NAME.bin
     VkPhysicalDevice physical;
@@ -95,10 +102,28 @@ void GameDenoiser::record(VkCommandBuffer cb, const GameFrameR& frame) {
         throw std::runtime_error("unsupported game texture format");
     // The stages sample every input at its texel coordinates from the origin, so only views at origin zero are
     // borrowed; anything else is copied first. Every input must be in GENERAL when borrowed (one layout for all).
-    auto direct = [](const GameImage& image) {
-        return image.view != VK_NULL_HANDLE && !image.x && !image.y && image.layout == VK_IMAGE_LAYOUT_GENERAL;
+    // Per-input ring-buffer detection: track distinct views and flag inputs that rotate.
+    // The engine caches descriptor sets per unique combination of input views (limit 16).
+    // If K inputs each rotate through N views, there are up to N^K combinations.
+    // Copying any rotating input to a canonical plane removes it from the combination.
+    const VkImageView inputViews[11] = {
+        frame.color.view, frame.motion.view, frame.depth.view,
+        frame.normals.view, frame.roughness.view, frame.diffuseAlbedo.view,
+        frame.specularAlbedo.view, frame.specularHitDistance.view,
+        frame.disocclusionMask.view, frame.subsurfaceGuide.view, frame.colorBeforeParticles.view
     };
-    const bool directColor = direct(frame.color), directMotion = direct(frame.motion), directDepth = direct(frame.depth);
+    for (int i = 0; i < 11; ++i) {
+        if (inputViews[i] == VK_NULL_HANDLE) continue;
+        auto& seen = g.seenViews[i];
+        if (std::find(seen.begin(), seen.end(), inputViews[i]) == seen.end()) {
+            seen.push_back(inputViews[i]);
+            if (seen.size() > 4) g.needsCopy[i] = true;
+        }
+    }
+    auto direct = [&](const GameImage& image, int idx) {
+        return !g.needsCopy[idx] && image.view != VK_NULL_HANDLE && !image.x && !image.y && image.layout == VK_IMAGE_LAYOUT_GENERAL;
+    };
+    const bool directColor = direct(frame.color, 0), directMotion = direct(frame.motion, 1), directDepth = direct(frame.depth, 2);
     // the four guides, then the two optional inputs (kept in their own one-channel formats when copied)
     const GameImage* guides[8] = {&frame.normals, &frame.roughness, &frame.diffuseAlbedo, &frame.specularAlbedo,
                                   &frame.specularHitDistance, &frame.disocclusionMask, &frame.subsurfaceGuide,
@@ -114,7 +139,7 @@ void GameDenoiser::record(VkCommandBuffer cb, const GameFrameR& frame) {
     }
     g.fallback.record(cb, frame.color, frame.motion, frame.depth, w, h, w, h, directColor, directMotion, directDepth);
     for (int i = 0; i < 8; ++i) {
-        if (!present(i) || direct(*guides[i])) continue;
+        if (!present(i) || direct(*guides[i], 3 + i)) continue;
         if (!g.guides[i].image) {
             g.fallback.ensure(cb, g.guides[i], guideFormats[i], w, h);
         }
@@ -132,7 +157,7 @@ void GameDenoiser::record(VkCommandBuffer cb, const GameFrameR& frame) {
                                   &images.specularHitDistance, &images.disocclusionMask, &images.subsurfaceGuide,
                                   &images.colorBeforeParticles};
     for (int i = 0; i < 8; ++i)
-        if (present(i)) *guideViews[i] = direct(*guides[i]) ? guides[i]->view : g.guides[i].view;
+        if (present(i)) *guideViews[i] = direct(*guides[i], 3 + i) ? guides[i]->view : g.guides[i].view;
     const bool directOutput = frame.output.view != VK_NULL_HANDLE && frame.output.format == VK_FORMAT_R16G16B16A16_SFLOAT &&
         frame.output.layout == VK_IMAGE_LAYOUT_GENERAL && !frame.output.x && !frame.output.y &&
         frame.output.image != frame.color.image && frame.output.image != frame.motion.image && frame.output.image != frame.depth.image;
diff --git a/engine/layer_m.comp b/engine/layer_m.comp
index 6617327..2d85faa 100644
--- a/engine/layer_m.comp
+++ b/engine/layer_m.comp
@@ -64,8 +64,8 @@
 #ifndef WPW
 #define WPW 1
 #endif
-#if WPW > 1 && KIND >= 3
-#error layer_m.comp: L's input and expansion stages stage their pixels per thread (WPW = 1)
+#if WPW > 2 && KIND >= 3
+#error layer_m.comp: L's input and expansion stages handle their 64 pixels one per invocation (WPW <= 2)
 #endif
 layout(local_size_x = 128 / WPW) in;
 // o: 0 g1, 1 bo, 2 g2, 3 b2 (f16); per head h at 4 + 4h: q, v (FP8), bias (f16), wo (FP8);
@@ -213,6 +213,15 @@ HAC bias16(uint off, uint stride) { HAC b; coopMatLoad(b, w16.w, off / 4u, strid
 // straight from f32 into FP8, two elements per saturating conversion; the activation in f32 as well
 #define SAT2_(dst, i, v) { fe4m3vec2 t_; saturatedConvertEXT(t_, (v)); (dst)[i] = t_.x; (dst)[(i) + 1] = t_.y; }
 #define MK8_F8(dst, a) { SAT2_(dst, 0, vec2((a)[0], (a)[1])) SAT2_(dst, 2, vec2((a)[2], (a)[3])) SAT2_(dst, 4, vec2((a)[4], (a)[5])) SAT2_(dst, 6, vec2((a)[6], (a)[7])) }
+// SATF2_: the same for a value that is never infinite, written as "an infinity becomes NaN, then saturate", which RADV
+// turns into the hardware conversion with FP16_OVFL instead of a NaN-preserving clamp per element and the conversion
+// (f2e4m3fn_satfn; the same result for every finite value and NaN, engine/rr/rr_exact.glsl FINITE). Used for the
+// attention's f32 sums of FP8 products (MK8_F8F) and the normalised tokens: the norm sums the squares of the very
+// values it scales, so an infinite channel makes the reciprocal 0 and the product NaN, never infinite. The binary16
+// accumulators (q, the MLP, the expansion, the residual) can overflow to infinity and keep SAT2_.
+vec2 finite2_(vec2 v) { return vec2(isinf(v.x) ? uintBitsToFloat(0x7fc00000u) : v.x, isinf(v.y) ? uintBitsToFloat(0x7fc00000u) : v.y); }
+#define SATF2_(dst, i, v) { fe4m3vec2 t_; saturatedConvertEXT(t_, finite2_(v)); (dst)[i] = t_.x; (dst)[(i) + 1] = t_.y; }
+#define MK8_F8F(dst, a) { SATF2_(dst, 0, vec2((a)[0], (a)[1])) SATF2_(dst, 2, vec2((a)[2], (a)[3])) SATF2_(dst, 4, vec2((a)[4], (a)[5])) SATF2_(dst, 6, vec2((a)[6], (a)[7])) }
 #define GELU2_(dst, i, h) { const vec2 x_ = vec2((h)[i], (h)[(i) + 1]); const vec2 y_ = clamp(x_, vec2(-2.0), vec2(2.0)); \
     const vec2 g_ = fma(y_, fma(-abs(y_), vec2(0.08108133), vec2(0.41216981)), vec2(0.5)); SAT2_(dst, i, x_ * g_) }
 #define GELU8(dst, h) { GELU2_(dst, 0, h) GELU2_(dst, 2, h) GELU2_(dst, 4, h) GELU2_(dst, 6, h) }
@@ -222,12 +231,13 @@ HAC bias16(uint off, uint stride) { HAC b; coopMatLoad(b, w16.w, off / 4u, strid
     const float16_t q4_ = CL8(Q16((a)[4])), q5_ = CL8(Q16((a)[5])), q6_ = CL8(Q16((a)[6])), q7_ = CL8(Q16((a)[7])); \
     (dst)[0] = floate4m3_t(q0_); (dst)[1] = floate4m3_t(q1_); (dst)[2] = floate4m3_t(q2_); (dst)[3] = floate4m3_t(q3_); \
     (dst)[4] = floate4m3_t(q4_); (dst)[5] = floate4m3_t(q5_); (dst)[6] = floate4m3_t(q6_); (dst)[7] = floate4m3_t(q7_); }
+#define MK8_F8F MK8_F8
 #endif
 // the normalised tokens into FP8 (inside NM1 / N2: x0..x7 the token, g0 / g1 the gains)
 #ifdef M_RELAXED
 #define H1N(N, r_) { const float f_ = float(r_); \
-        SAT2_(h1[N], 0, vec2(x0, x1) * (f_ * vec2(g0.xy))) SAT2_(h1[N], 2, vec2(x2, x3) * (f_ * vec2(g0.zw))) \
-        SAT2_(h1[N], 4, vec2(x4, x5) * (f_ * vec2(g1.xy))) SAT2_(h1[N], 6, vec2(x6, x7) * (f_ * vec2(g1.zw))) }
+        SATF2_(h1[N], 0, vec2(x0, x1) * (f_ * vec2(g0.xy))) SATF2_(h1[N], 2, vec2(x2, x3) * (f_ * vec2(g0.zw))) \
+        SATF2_(h1[N], 4, vec2(x4, x5) * (f_ * vec2(g1.xy))) SATF2_(h1[N], 6, vec2(x6, x7) * (f_ * vec2(g1.zw))) }
 #else
 #define H1N(N, r_) { \
         h1[N][0] = floate4m3_t(CL8(x0 * float16_t(r_ * g0.x))); h1[N][1] = floate4m3_t(CL8(x1 * float16_t(r_ * g0.y))); \
@@ -356,7 +366,6 @@ float16_t gelu1(float16_t x)
     precise f16vec2 C_ = f16vec2(EV(T, 4), EV(T, 5)), D_ = f16vec2(EV(T, 6), EV(T, 7)); \
     P0_[T] = (A_ * A_) + (B_ * B_); P1_[T] = (C_ * C_) + (D_ * D_); }
 
-shared f16vec4 VSH[4 * 2 * 16 * VWD / 4];   // per wave and head: the V tiles, transposed through shared memory
 #define MERGE (KIND == 1 || KIND == 3)
 #define EXPAND (KIND == 2 || KIND == 4)
 #if MERGE
@@ -375,6 +384,34 @@ shared float16_t X0S[64 * XS];              // the embedding's output x0, the bl
 const uint HS = 40u;
 shared float16_t HO[64 * HS];               // the block's f16 output, window-major, for the per-pixel head
 #endif
+// VSH: per wave and head the V tiles, transposed through shared memory (4 KiB, 512 bytes a tile). It is only accessed
+// as matrices, and only during the attention, so it takes the place of a buffer that is dead then: the merge rows
+// (written after the attention), the expansion (read before it) or, for KIND 4, the head's input (written after it).
+// Shared memory caps how many waves this GPU runs, and ACO gives a shader the registers that number of waves leaves.
+#if MERGE
+#define VSH NB
+#define VSH_UNIT 1u
+#if 32 * (4 * C + 8) < 4096
+#error layer_m.comp: the V tiles no longer fit in NB
+#endif
+#elif KIND == 2
+#define VSH NEX
+#define VSH_UNIT 1u
+#if 16 * (4 * C + 8) < 4096
+#error layer_m.comp: the V tiles no longer fit in NEX
+#endif
+#elif KIND == 4
+#define VSH HO
+#define VSH_UNIT 2u
+#if 64 * 40 * 2 < 4096
+#error layer_m.comp: the V tiles no longer fit in HO
+#endif
+#else
+shared floate4m3_t VSH[4 * 2 * 16 * VWD * 2];
+#define VSH_UNIT 1u
+#endif
+#define VSH_AT(t) ((t) * (512u / VSH_UNIT))      // tile t, in the buffer's elements
+#define VSH_STRIDE (VWD * 2u / VSH_UNIT)          // a row of VWD binary16
 #define SYNC memoryBarrierShared(); barrier();
 #if KIND == 4
 // NVIDIA's f32 activations of the head (rrlite_dec0_4x4's output code): a sigmoid rounded to binary16, and a
@@ -459,8 +496,7 @@ void main()
 #endif
 #if KIND == 3
     // ------------------------------------------------------------ L's input: features, embedding -> x0
-    FORW {
-        WINV
+    {
         const float e1 = ep.expo == 0.0 ? 1.0 : ep.expo, e2 = ep.expoPrev == 0.0 ? 1.0 : ep.expoPrev;
         const uint li = gl_LocalInvocationID.x;
         if (li < 64u)
@@ -473,7 +509,10 @@ void main()
             const uint at = (((py >> 2) * 2u + (px >> 2)) * 16u + (py & 3u) * 4u + (px & 3u)) * FS;
             for (uint k = 0u; k < 32u; ++k) FEW[at + k] = fe[k];
         }
-        SYNC
+    }
+    SYNC
+    FORW {
+        WINV
         // the embedding has the shape of an MLP chunk: 32 features -> 32 hidden (ReLU, FP8) -> C
         B8 f0, f1;
         coopMatLoad(f0, FEW, mt * 16u * FS, FS, CM);
@@ -495,8 +534,8 @@ void main()
         for (int i = 0; i < 8; ++i) X0S[xb + 16u * (N) + uint(i)] = float16_t(x_[i]); }
         EACH_T(EMB_N)
 #undef EMB_N
-        SYNC
     }
+    SYNC
 #endif
     // ------------------------------------------------------------ norm 1: sum of squares, residual and h1
     FORW {
@@ -542,6 +581,9 @@ void main()
     H1N(N, rn) }
     EACH_T(NM1)
     }
+#if KIND == 2
+    SYNC   // every wave has read the block's input from NEX: the V tiles take its place
+#endif
     // ------------------------------------------------------------ attention, one head at a time
     for (uint hd = 0u; hd < uint(NH); ++hd)
     {
@@ -593,8 +635,8 @@ void main()
         // V of every window goes to its shared buffer now, so it does not occupy registers during the scores
         FORW {
             const uint mt = mt0 + w;
-            coopMatStore(v[w][0], VSH, (mt * 2u + 0u) * 64u, VWD / 4u, RM);
-            coopMatStore(v[w][1], VSH, (mt * 2u + 1u) * 64u, VWD / 4u, RM);
+            coopMatStore(v[w][0], VSH, VSH_AT(mt * 2u + 0u), VSH_STRIDE, RM);
+            coopMatStore(v[w][1], VSH, VSH_AT(mt * 2u + 1u), VSH_STRIDE, RM);
         }
         }
         // scores: q8(Q) against itself, rows are the queries and columns the keys, plus the relative position bias
@@ -624,8 +666,8 @@ void main()
         // V through this wave's shared buffer, then O^T = V^T x P^T: rows are the head channels, columns the tokens
         ACC o[2];
         A16 vt0, vt1;
-        coopMatLoad(vt0, VSH, (mt * 2u + 0u) * 64u, VWD / 4u, RM);
-        coopMatLoad(vt1, VSH, (mt * 2u + 1u) * 64u, VWD / 4u, RM);
+        coopMatLoad(vt0, VSH, VSH_AT(mt * 2u + 0u), VSH_STRIDE, RM);
+        coopMatLoad(vt1, VSH, VSH_AT(mt * 2u + 1u), VSH_STRIDE, RM);
 #if TEXL
         // P V is the block's one product of two f16 operands, and its f32 sum over the 16 keys is not exact: ZLUDA's
         // lowering of m16n8k16 puts key pair (j & 1) * 4 + 2 h + (j >> 1) in K dword j of half h (swin_tok.h
@@ -635,7 +677,7 @@ void main()
         o[0] = coopMatMulAdd(vt0, pb, ACC(0.0));
         o[1] = coopMatMulAdd(vt1, pb, ACC(0.0));
         R16(o[0]); R16(o[1]);
-        MK8_F8(OB_[w][0], o[0]); MK8_F8(OB_[w][1], o[1]);
+        MK8_F8F(OB_[w][0], o[0]); MK8_F8F(OB_[w][1], o[1]);
         }
         // the output projection of this head's 32 attention channels, straight into the residual stream
 #define WO_N(N) { \
@@ -683,6 +725,20 @@ void main()
     EACH_T(N2)
     }
     // ------------------------------------------------------------ the MLP, 32 hidden channels at a time
+    // M_RESIDUAL_F32 (opt-in, part of M_RELAXED): the residual stays f32 across the MLP's chunks and is rounded to
+    // binary16 once after them, instead of after every chunk
+#ifdef M_RESIDUAL_F32
+#ifndef M_RELAXED
+#error layer_m.comp: M_RESIDUAL_F32 needs M_RELAXED
+#endif
+    ACC RF_[WPW][NT];
+#define RF_IN(N) RF_[w][N] = ACC(res[N]);
+    FORW { EACH_T(RF_IN) }
+#undef RF_IN
+#define K32R(N, a0, b0, a1, b1) { RF_[w][N] = coopMatMulAdd((a1), (b1), coopMatMulAdd((a0), (b0), RF_[w][N])); }
+#else
+#define K32R(N, a0, b0, a1, b1) K32H(res[N], a0, b0, a1, b1)
+#endif
     for (uint c = 0u; c < UC / 8u; ++c)
     {
         const uint ow1 = p.o[20u + 3u * c], ob1 = p.o[21u + 3u * c], ow2 = p.o[22u + 3u * c];
@@ -717,17 +773,26 @@ void main()
         A8 a0, a1; \
         LDA8(a0, ow2, 0u, UC, 16u * (N)); \
         LDA8(a1, ow2, 1u, UC, 16u * (N)); \
-        FORW { K32H(res[N], a0, G0_[w], a1, G1_[w]) } }
+        FORW { K32R(N, a0, G0_[w], a1, G1_[w]) } }
         EACH_T(MLP2_N)
     }
+#undef K32R
+#ifdef M_RESIDUAL_F32
+#define RF_OUT(N) res[N] = HAC(RF_[w][N]);
+    FORW { EACH_T(RF_OUT) }
+#undef RF_OUT
+#endif
 #if KIND == 4
     // ------------------------------------------------------------ L's expansion head: per-pixel reconstruction inputs
+    SYNC   // every wave is done with the V tiles in HO
     FORW {
         WINV
         const uint hb = (mt * 16u + row) * HS + ch8;
         for (uint nt = 0u; nt < NT; ++nt)
             for (int i = 0; i < 8; ++i) HO[hb + 16u * nt + uint(i)] = float16_t(res[nt][i]);
-        SYNC
+    }
+    SYNC
+    {
         const uint li = gl_LocalInvocationID.x;
         const uint px = li & 7u, py = li >> 3;
         const ivec2 at = ivec2(X0 + int(px), Y0 + int(py));
@@ -763,6 +828,7 @@ void main()
 #endif
     // ------------------------------------------------------------ patch merge: 2x2 tokens (4C) -> XC, half resolution
 #if MERGE
+    SYNC   // every wave is done with the V tiles in NB
     FORW {
         WINV
         // the 16 merged rows are gathered from the four waves: the row is the merged token of the lane's pixel
diff --git a/engine/post_m.comp b/engine/post_m.comp
index ab02dc6..a347ae2 100644
--- a/engine/post_m.comp
+++ b/engine/post_m.comp
@@ -154,6 +154,42 @@ void main()
     // three render texel rows and its own-cell residual rounding exactly as in the 16x16 shader
     const int y0 = 2 * gid.y;
     const bool inside0 = gid.x < p.hiSize.x && y0 < p.hiSize.y;
+    const vec2 pos0 = fma(vec2(gid.x, y0) + 0.5, p.hi2rnd, p.rndOffs), pos1 = fma(vec2(gid.x, y0 + 1) + 0.5, p.hi2rnd, p.rndOffs);
+    // The Gaussians of both rows at once: they have the same column (dx) and texels and differ only in dy, so one set
+    // of tile reads serves both, and each row's operations are the ones it had on its own, lane by lane (dy, the
+    // exponents and the third accumulator are binary16 pairs, element 0 the even row).
+    f16vec4 a0[2] = {f16vec4(0.0), f16vec4(0.0)}, a1[2] = {f16vec4(0.0), f16vec4(0.0)};
+    f16vec2 a2 = f16vec2(0.0);
+    float W[2] = {0.0, 0.0};
+    if (inside0)
+    {
+        // NVIDIA's kernel handles rows in pairs and gives the odd row the even row's three texel rows
+        const ivec2 base = ivec2(floor(pos0 - 0.5));
+        const f16vec2 d0 = f16vec2((vec2(base) + 0.5) - pos0);
+        // Native paired rows round the odd row's own-cell residual first, then adjust it in f16.
+        // Casting the even-cell residual once changes Gaussian weights at half-rounding boundaries.
+        const int ownBaseY = int(floor(pos1.y - 0.5));
+        const float16_t d1y = ownBaseY != base.y ? float16_t((float(ownBaseY) + 0.5) - pos1.y) - float16_t(1.0)
+                                                 : float16_t((float(base.y) + 0.5) - pos1.y);
+        const f16vec2 dy0 = f16vec2(d0.y, d1y);
+        for (int i = TAP_LO; i <= TAP_HI; ++i)
+        {
+            const float16_t dx = d0.x + float16_t(i), sq = dx * dx;
+            for (int j = TAP_LO; j <= TAP_HI; ++j)
+            {
+                const int k = (base.y - origin.y + j) * TILE_W + (base.x - origin.x + i);
+                const f16vec4 t0 = tile0[k], t1 = tile1[k], t2 = tile2[k];
+                const f16vec2 dy = dy0 + float16_t(j);
+                precise f16vec2 e = (f16vec2(sq * t2.y) + (dy * dy) * t2.z) + (dx * dy) * t2.w;
+                const float w0 = exp2(float(e.x)), w1 = exp2(float(e.y));
+                const float16_t wh0 = float16_t(w0), wh1 = float16_t(w1);
+                W[0] += w0; W[1] += w1;
+                a0[0] = fma(f16vec4(wh0), t0, a0[0]); a1[0] = fma(f16vec4(wh0), t1, a1[0]);
+                a0[1] = fma(f16vec4(wh1), t0, a0[1]); a1[1] = fma(f16vec4(wh1), t1, a1[1]);
+                a2 = fma(f16vec2(wh0, wh1), f16vec2(t2.x), a2);
+            }
+        }
+    }
     vec3 outA0 = vec3(0.0), outA1 = vec3(0.0);
     for (int r = 0; r < 2; ++r)
     {
@@ -163,39 +199,10 @@ void main()
         if (inside)
         {
             const vec2 centre = vec2(pix) + 0.5;
-            const vec2 pos = fma(centre, p.hi2rnd, p.rndOffs);
-            // NVIDIA's kernel handles rows in pairs and gives the odd row the even row's three texel rows
-            const float evenY = fma(float(pix.y & ~1) + 0.5, p.hi2rnd.y, p.rndOffs.y);
-            const ivec2 near = ivec2(floor(pos)), base = ivec2(floor(pos.x - 0.5), floor(evenY - 0.5));
-            f16vec2 d0 = f16vec2((vec2(base) + 0.5) - pos);
-            if ((pix.y & 1) != 0)
-            {
-                // Native paired rows round the odd row's own-cell residual first, then adjust it in f16.
-                // Casting the even-cell residual once changes Gaussian weights at half-rounding boundaries.
-                const int ownBaseY = int(floor(pos.y - 0.5));
-                if (ownBaseY != base.y)
-                    d0.y = float16_t((float(ownBaseY) + 0.5) - pos.y) - float16_t(1.0);
-            }
-            f16vec4 a0 = f16vec4(0.0), a1 = f16vec4(0.0);
-            float16_t a2 = float16_t(0.0);
-            float W = 0.0;
-            for (int i = TAP_LO; i <= TAP_HI; ++i)
-            {
-                const float16_t dx = d0.x + float16_t(i), sq = dx * dx;
-                for (int j = TAP_LO; j <= TAP_HI; ++j)
-                {
-                    const int k = (base.y - origin.y + j) * TILE_W + (base.x - origin.x + i);
-                    const f16vec4 t2 = tile2[k];
-                    const float16_t dy = d0.y + float16_t(j);
-                    precise float16_t e = (sq * t2.y + (dy * dy) * t2.z) + (dx * dy) * t2.w;
-                    const float w = exp2(float(e));
-                    const float16_t wh = float16_t(w);
-                    W += w;
-                    a0 = fma(f16vec4(wh), tile0[k], a0); a1 = fma(f16vec4(wh), tile1[k], a1); a2 = fma(wh, t2.x, a2);
-                }
-            }
-            const float inv = 1.0 / W;
-            const vec3 cur = inv * vec3(a0.rgb), m = inv * vec3(a0.w, a1.x, a1.y), g = inv * vec3(a1.z, a1.w, a2);
+            const vec2 pos = r == 0 ? pos0 : pos1;
+            const ivec2 near = ivec2(floor(pos));
+            const float inv = 1.0 / W[r];
+            const vec3 cur = inv * vec3(a0[r].rgb), m = inv * vec3(a0[r].w, a1[r].x, a1[r].y), g = inv * vec3(a1[r].z, a1[r].w, a2[r]);
             // the two most likely motion candidates
             const vec2 mvPos = any(notEqual(textureSize(mvCTex, 0), p.rndSize))
                 ? clamp(pos, vec2(0.5), vec2(p.rndSize) - 0.5) : pos;
@@ -214,6 +221,9 @@ void main()
             // the fallback colour is only fetched where the history sample leaves the image or the history was reset:
             // the value is the same either way, so computing it in the branch saves the fetch on ordinary pixels
             const bool out1 = q1.x < 0.0 || q1.y < 0.0 || q1.x > size.x || q1.y > size.y || p.reset != 0u;
+            // the second candidate often has the first one's vector (where the motion is uniform): the same position,
+            // so the first one's sample or fallback (a NaN is never equal, and +0 and -0 give the same position)
+            const bool same = all(equal(second.xy, big.xy));
 #if PRESET_L
             // linear light: a sample's expansion (when the previous exposure is positive) or the colour texel
             const vec3 h1 = out1 ? texelFetch(colorTex, colourTexel(near), 0).rgb : historyLinear(q1, p.expoPrev == 0.0 ? 1.0 : p.expoPrev);
@@ -221,7 +231,9 @@ void main()
             if (b2 > 0.0)
             {
                 const vec2 q2 = centre + second.xy;
-                if (q2.x < 0.0 || q2.y < 0.0 || q2.x > size.x || q2.y > size.y || p.reset != 0u)
+                if (same)
+                    h2 = h1;
+                else if (q2.x < 0.0 || q2.y < 0.0 || q2.x > size.x || q2.y > size.y || p.reset != 0u)
                     h2 = out1 ? h1 : texelFetch(colorTex, colourTexel(near), 0).rgb;
                 else h2 = historyLinear(q2, p.expoPrev == 0.0 ? 1.0 : p.expoPrev);
             }
@@ -232,14 +244,16 @@ void main()
             if (b2 > 0.0)
             {
                 const vec2 q2 = centre + second.xy;
-                if (q2.x < 0.0 || q2.y < 0.0 || q2.x > size.x || q2.y > size.y || p.reset != 0u)
+                if (same)
+                    h2 = h1;
+                else if (q2.x < 0.0 || q2.y < 0.0 || q2.x > size.x || q2.y > size.y || p.reset != 0u)
                     h2 = out1 ? h1 : compress(e2 * texelFetch(colorTex, colourTexel(near), 0).rgb);
                 else h2 = history(q2);
             }
             const vec3 hist = e1 * (expand(fma(vec3(b1), h1, b2 * h2)) * (1.0 / e2));
 #endif
             const vec3 lin = expand(cur);
-            const float g0 = clamp(p.gain * W * g.x, 0.0, 1.0);
+            const float g0 = clamp(p.gain * W[r] * g.x, 0.0, 1.0);
             const float ga = fma(g0, 1.0 - g.y, g.y), gb = fma(g0, 1.0 - g.z, g.z);
             outA = fma(lin - hist, vec3(ga), hist);
             if (isinf(outA.x + outA.y + outA.z) || isnan(outA.x + outA.y + outA.z)) outA = lin;
diff --git a/engine/replay_m.cpp b/engine/replay_m.cpp
index 817e22b..d9ff6f7 100644
--- a/engine/replay_m.cpp
+++ b/engine/replay_m.cpp
@@ -107,7 +107,7 @@ int main(int argc, char** argv) try {
     CK(vkCreateCommandPool(c.device, &pi, nullptr, &c.pool));
     VkCommandBufferAllocateInfo cai{VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO}; cai.commandPool = c.pool; cai.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY; cai.commandBufferCount = 1;
     VkCommandBuffer cb; CK(vkAllocateCommandBuffers(c.device, &cai, &cb));
-    VkQueryPoolCreateInfo qpi{VK_STRUCTURE_TYPE_QUERY_POOL_CREATE_INFO}; qpi.queryType = VK_QUERY_TYPE_TIMESTAMP; qpi.queryCount = 7;
+    VkQueryPoolCreateInfo qpi{VK_STRUCTURE_TYPE_QUERY_POOL_CREATE_INFO}; qpi.queryType = VK_QUERY_TYPE_TIMESTAMP; qpi.queryCount = 17;
     VkQueryPool qp; CK(vkCreateQueryPool(c.device, &qpi, nullptr, &qp));
     VkPhysicalDeviceProperties props; vkGetPhysicalDeviceProperties(c.physical, &props);
     auto submit = [&] {
@@ -151,7 +151,7 @@ int main(int argc, char** argv) try {
                 if (data.size() != sizes[k]) throw std::runtime_error("wrong input size");
                 frameInputs.push_back(std::move(data));
             }
-        std::vector<double> times, walls; std::vector<std::vector<double>> stage(6);
+        std::vector<double> times, walls; std::vector<std::vector<double>> stage(6), layer(10);
         const auto now = [] { return std::chrono::duration<double, std::micro>(std::chrono::steady_clock::now().time_since_epoch()).count(); };
         // D4R_M_RESET_ALL=1: every frame is a reset frame, as the D3D12 harness passes them without its temporal option
         const bool resetAll = std::getenv("D4R_M_RESET_ALL") != nullptr;
@@ -169,7 +169,7 @@ int main(int argc, char** argv) try {
                 s.renderWidth = rw; s.renderHeight = rh;
                 const double t0 = now();
                 engine.setFrame({color.view, motion.view, depth.view}, s);
-                vkCmdResetQueryPool(cb, qp, 0, 7);
+                vkCmdResetQueryPool(cb, qp, 0, 17);
                 double networkUs = 0;
                 if (hip) {
 #ifdef D4R_ENGINE_HIP
@@ -186,12 +186,16 @@ int main(int argc, char** argv) try {
                 if (pass == 0) vkCmdCopyImageToBuffer(cb, engine.outputImage(), VK_IMAGE_LAYOUT_GENERAL, out.buffer, 1, &region);
                 submit();
                 if (pass) walls.push_back(now() - t0);
-                uint64_t t[7]; CK(vkGetQueryPoolResults(c.device, qp, 0, 7, sizeof(t), t, 8, VK_QUERY_RESULT_64_BIT | VK_QUERY_RESULT_WAIT_BIT));
+                uint64_t t[17] = {};
+                // D4R_ENGINE_M_LAYER_STAMPS=1: the ten network layers' stamps as well (queries 7 - 16)
+                static const bool layerStamps = std::getenv("D4R_ENGINE_M_LAYER_STAMPS") != nullptr && !hip;
+                CK(vkGetQueryPoolResults(c.device, qp, 0, layerStamps ? 17 : 7, sizeof(t), t, 8, VK_QUERY_RESULT_64_BIT | VK_QUERY_RESULT_WAIT_BIT));
                 const double tick = props.limits.timestampPeriod / 1000.0;
                 double d[6]; for (int k = 0; k < 6; ++k) d[k] = double(t[k + 1] - t[k]) * tick;
                 if (hip) d[2] = networkUs; // measured HIP network time; the complete span includes host handoffs
                 const double us = double(t[6] - t[0]) * tick;
                 if (pass) for (int k = 0; k < 6; ++k) stage[k].push_back(d[k]);
+                if (pass && layerStamps) for (int l = 0; l < 10; ++l) layer[l].push_back(double(t[7 + l] - (l ? t[6 + l] : t[2])) * tick);
                 if (pass == 0 && std::getenv("D4R_M_DUMP")) {
                     // development: every intermediate of this frame, to PREFIX-NAME-FRAME.bin
                     for (const auto& d : engine.debugResources()) {
@@ -216,6 +220,11 @@ int main(int argc, char** argv) try {
             std::sort(walls.begin(), walls.end()); std::printf("  wall clock, record to completion: median %.1f us, min %.1f us\n", walls[walls.size() / 2], walls[0]);
             const char* names[6] = {"exposure", "input", "network", "expansion", "reconstruction", "downsample"};
             for (int k = 0; k < 6; ++k) { std::sort(stage[k].begin(), stage[k].end()); std::printf("  %-14s median %.1f us\n", names[k], stage[k][stage[k].size() / 2]); }
+            if (!layer[0].empty()) {
+                std::printf("  layers (us):");
+                for (auto& v : layer) { std::sort(v.begin(), v.end()); std::printf(" %.1f", v[v.size() / 2]); }
+                std::printf("\n");
+            }
         }
         CK(vkDeviceWaitIdle(c.device));
     }
diff --git a/engine/replay_rr.cpp b/engine/replay_rr.cpp
index 4e3d605..c591d79 100644
--- a/engine/replay_rr.cpp
+++ b/engine/replay_rr.cpp
@@ -158,6 +158,8 @@ int main(int argc, char** argv) try {
     if (std::getenv("D4R_RR_TEST_NO_ROBUST")) { f2.features.robustBufferAccess = VK_FALSE; v13.robustImageAccess = VK_FALSE; }
     // D4R_RR_TEST_NO_COOPMAT=1: a device like vkd3d-proton's, without the cooperative-matrix and float8 features
     if (std::getenv("D4R_RR_TEST_NO_COOPMAT")) { di.enabledExtensionCount = 0; v13.pNext = nullptr; }
+    // D4R_RR_TEST_NO_INT64_ATOMICS=1: no 64-bit buffer atomics (preset F's exposure then takes its two passes)
+    if (std::getenv("D4R_RR_TEST_NO_INT64_ATOMICS")) v12.shaderBufferInt64Atomics = v12.shaderSharedInt64Atomics = VK_FALSE;
     CK(vkCreateDevice(c.physical, &di, nullptr, &c.device));
     vkGetDeviceQueue(c.device, family, 0, &c.queue);
     VkCommandPoolCreateInfo pi{VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO}; pi.queueFamilyIndex = family;
@@ -169,7 +171,7 @@ int main(int argc, char** argv) try {
     std::string key;
     uint32_t rw, rh, ow, oh;
     frames >> key >> rw >> rh >> ow >> oh;
-    d4r::RModel model(c.physical, c.device, modelDir);
+    d4r::RModel model(c.physical, c.device, modelDir, VK_NULL_HANDLE, v12.shaderBufferInt64Atomics == VK_TRUE);
     d4r::REngine engine(model, rw, rh, ow, oh);
     const char* names[] = {"color", "motion", "depth", "normals", "roughness", "diffuse", "specular"};
     const VkFormat formats[] = {VK_FORMAT_R16G16B16A16_SFLOAT, VK_FORMAT_R16G16_SFLOAT, VK_FORMAT_R32_SFLOAT,
diff --git a/engine/rr/README.md b/engine/rr/README.md
index f7efb44..3e1f6a9 100644
--- a/engine/rr/README.md
+++ b/engine/rr/README.md
@@ -615,3 +615,30 @@ was measured and not built for D / E (`build/rr/f/nv/dstage.sh`, `dnl.sh`, level
   loads between a window's waves) decoder 0 takes the same time (803-818 us), encoder 1 the same, decoder 1 10% less.
   These windows also have 96 key tokens (3 KiB each for the tokens, the keys and the values at C = 32), which does
   not fit the 4 KiB a single wave can hold at full rate.
+
+### Byte-identical speed-ups (2026-10-19)
+
+RX 9070 XT, RADV 26.2.2, 1280x720 -> 2560x1440 with `D4R_ENGINE_RR_STAGE_STAMPS=1`. Every image of every frame is
+byte-identical to the previous model on the replay and stress inputs, at levels 0 to 3.
+
+- `ptx2rr.py`: OCML's tanh (decoder 0's second kernel) in f32 where f64 was not needed. Exp's error term becomes
+  Dekker's exact product error; the polynomial below 0.625 stays in f64, because the driver folds away the f32
+  rounding between its f64 steps. Identical to the previous function on all 2^32 inputs. dec0_1 475 -> 446 us, frame
+  3.568 -> 3.537 ms.
+- `exposure_f.comp` PASS 2: the exposure in one pass as NGX does it, each block adding its fixed-point sum to a 64-bit
+  total with an atomic add and the last block finishing. It needs `shaderBufferInt64Atomics`, which vkd3d-proton
+  enables; the shim reads the device's features through `GetDeviceFeatures`. Without it, or with a model that has no
+  `exposure_f2.spv`, the two passes run (`D4R_RR_TEST_NO_INT64_ATOMICS=1` forces them). Exposure 25.2 -> 14.9 us.
+
+### Opt-in builds for preset F (2026-10-19)
+
+Against the default build, presented image, RGB, peak 1, 1280x720 -> 1440p. For scale, the exact build
+(`D4R_RR_RELAXED=0`) is 43.0 - 44.4 dB from the default.
+
+- `D4R_RR2_UNGUARDED=1` (level 3 only): the translated code (encoder 0's and decoder 0's per-pixel parts, the output
+  passes) without its exactness guards: no `precise`, no opaque zero on binary16 stores, no `spv_exact_convert.py`.
+  56.0 dB on the reset frame, then 43.0 - 45.6 dB, as far from the exact build as the default is. Frame 3.526 -> 3.501
+  ms.
+- `D4R_RR2_POOLED=rgb10_a2` (since 2026-10-13): the pooled histories as UNORM 10:10:10:2 images, 31 MB less at 1440p.
+  Identical at 1280x720 -> 1440p and on reset frames; 43.7 - 47.2 dB on the history frames at 1114x626 -> 1080p and
+  28.8 dB with random motion. Frame 3.526 -> 3.484 ms, 3.456 with `D4R_RR2_UNGUARDED=1` as well.
diff --git a/engine/rr/exposure_f.comp b/engine/rr/exposure_f.comp
index fb4e308..fa456c6 100644
--- a/engine/rr/exposure_f.comp
+++ b/engine/rr/exposure_f.comp
@@ -1,33 +1,39 @@
 #version 460
-// Ray Reconstruction preset F (rr2_*): the auto exposure of rr2_histogram_auto_exposure_basic_kernel, in two passes.
+// Ray Reconstruction preset F (rr2_*): the auto exposure of rr2_histogram_auto_exposure_basic_kernel.
 //
 // NVIDIA's kernel sums a clamped luma per 16x16 block (a pairwise tree over the block's threads), converts each block
 // sum to fixed point (x 2^32 / 10000, round to nearest) and adds them with a 64-bit atomic; the last block to finish
-// turns the total into the exposure. The integer total does not depend on the order of the blocks, so PASS 0 writes
-// every block's fixed-point sum and PASS 1 (one workgroup) adds them and does what the last block does:
+// turns the total into the exposure:
 //   exposure = 0.23547088 / max(mean luma, 0.005)               -> the 1x1 exposure image of this frame
 //   values   = { e, p, 1 / e, clamp(log2(e / p), -6, 6) / 6 }   -> the 16-byte buffer encoder 0 and the output passes read
 // with e the exposure truncated to binary16 and p the previous frame's (1 where either is zero).
+// PASS 2 is that kernel (64-bit atomics on a storage buffer: shaderBufferInt64Atomics). Without them the runtime runs
+// two passes: PASS 0 writes every block's fixed-point sum and PASS 1 (one workgroup) adds them and does what the last
+// block does. The total is an integer: it does not depend on the order of the blocks, so both give the same values.
 #extension GL_EXT_shader_explicit_arithmetic_types : require
 #extension GL_EXT_control_flow_attributes : require
 #ifndef PASS
 #define PASS 0
 #endif
+#if PASS == 2
+#extension GL_EXT_shader_atomic_int64 : require
+#extension GL_KHR_memory_scope_semantics : require
+#endif
 layout(set = 0, binding = 0, std430) readonly buffer Params {
     ivec2 last;        // render size - 1
     uint blocksX, blocks;
     float invCount;    // 10000 / (2^32 * render pixels)
     uint reset;        // the kernel's argument 32: no log ratio when set
 } p;
-#if PASS == 0
+#if PASS != 1
 // One thread per column of the 16x16 block. NVIDIA's tree over the block's 256 threads (index 16 y + x: i += i + 128,
 // then 64, 32, ... 1) is, per column, ((r0 + r8) + (r4 + r12)) + ((r2 + r10) + (r6 + r14)) of the even rows plus the
 // same of the odd rows, and then the columns (c + c8, + 4, + 2, + 1).
 layout(local_size_x = 16) in;
 layout(set = 0, binding = 1) uniform sampler2D color;
-layout(set = 0, binding = 2, std430) writeonly buffer Sums { uint64_t s[]; } sums;
 shared float col[16];
-void main()
+// the block's fixed-point sum, in thread 0
+uint64_t blockSum()
 {
     const uint tx = gl_LocalInvocationID.x;
     const int x = int(gl_GlobalInvocationID.x), y0 = int(gl_WorkGroupID.y) * 16;
@@ -49,33 +55,37 @@ void main()
     precise float c = w1 + w0;
     col[tx] = c;
     barrier();
-    if (tx == 0u) {
-        precise float d[8];
-        [[unroll]] for (int k = 0; k < 8; ++k) d[k] = col[k + 8] + col[k];
-        precise float e0 = d[4] + d[0], e1 = d[5] + d[1], e2 = d[6] + d[2], e3 = d[7] + d[3];
-        precise float f0 = e2 + e0, f1 = e3 + e1;
-        precise float g = f1 + f0;
-        sums.s[gl_WorkGroupID.y * p.blocksX + gl_WorkGroupID.x] = uint64_t(roundEven(g * 429496.71875));
-    }
+    if (tx != 0u) return 0ul;
+    precise float d[8];
+    [[unroll]] for (int k = 0; k < 8; ++k) d[k] = col[k + 8] + col[k];
+    precise float e0 = d[4] + d[0], e1 = d[5] + d[1], e2 = d[6] + d[2], e3 = d[7] + d[3];
+    precise float f0 = e2 + e0, f1 = e3 + e1;
+    precise float g = f1 + f0;
+    return uint64_t(roundEven(g * 429496.71875));
+}
+#endif
+#if PASS == 0
+layout(set = 0, binding = 2, std430) writeonly buffer Sums { uint64_t s[]; } sums;
+void main()
+{
+    const uint64_t s = blockSum();
+    if (gl_LocalInvocationID.x == 0u) sums.s[gl_WorkGroupID.y * p.blocksX + gl_WorkGroupID.x] = s;
 }
 #else
+#if PASS == 1
 layout(local_size_x = 64) in;
 shared uint64_t part[64];
 layout(set = 0, binding = 1, std430) readonly buffer Sums { uint64_t s[]; } sums;
-layout(set = 0, binding = 2, std430) writeonly buffer Values { vec4 v; } values;
-layout(set = 0, binding = 3) uniform sampler2D previous;
-layout(set = 0, binding = 4, r16f) uniform writeonly image2D current;
-void main()
+#else
+// s[0]: the total, s[1]: the blocks done; the last block leaves both 0 for the next frame
+layout(set = 0, binding = 2, std430) buffer Sums { uint64_t s[]; } sums;
+#endif
+layout(set = 0, binding = PASS + 1, std430) writeonly buffer Values { vec4 v; } values;
+layout(set = 0, binding = PASS + 2) uniform sampler2D previous;
+layout(set = 0, binding = PASS + 3, r16f) uniform writeonly image2D current;
+// the last block's work, from the frame's total
+void finish(uint64_t total)
 {
-    // the block sums are integers: any order. 64 strided partial sums, then thread 0 finishes
-    const uint tid = gl_LocalInvocationID.x;
-    uint64_t mine = 0ul;
-    for (uint b = tid; b < p.blocks; b += 64u) mine += sums.s[b];
-    part[tid] = mine;
-    barrier();
-    if (tid != 0u) return;
-    uint64_t total = 0ul;
-    for (uint b = 0u; b < 64u; ++b) total += part[b];
     precise float mean = p.invCount * float(total);
     precise float e = 0.23547088 / max(mean, 0.005);
     imageStore(current, ivec2(0), vec4(e));
@@ -90,4 +100,33 @@ void main()
     if (p.reset == 0u) ratio = clamp(log2(max(cur, 1e-8) / max(prev, 1e-8)), -6.0, 6.0) * 0.16666667;
     values.v = vec4(cur, prev, 1.0 / cur, ratio);
 }
+#if PASS == 1
+void main()
+{
+    // the block sums are integers: any order. 64 strided partial sums, then thread 0 finishes
+    const uint tid = gl_LocalInvocationID.x;
+    uint64_t mine = 0ul;
+    for (uint b = tid; b < p.blocks; b += 64u) mine += sums.s[b];
+    part[tid] = mine;
+    barrier();
+    if (tid != 0u) return;
+    uint64_t total = 0ul;
+    for (uint b = 0u; b < 64u; ++b) total += part[b];
+    finish(total);
+}
+#else
+void main()
+{
+    const uint64_t s = blockSum();
+    if (gl_LocalInvocationID.x != 0u) return;
+    atomicAdd(sums.s[0], s, gl_ScopeDevice, 0, gl_SemanticsRelaxed);
+    // the release orders this block's add before its count, and the last block's acquire every block's add before
+    // its read of the total
+    if (atomicAdd(sums.s[1], 1ul, gl_ScopeDevice, gl_StorageSemanticsBuffer, gl_SemanticsAcquireRelease) != uint64_t(p.blocks - 1u))
+        return;
+    const uint64_t total = atomicExchange(sums.s[0], 0ul, gl_ScopeDevice, 0, gl_SemanticsRelaxed);
+    atomicExchange(sums.s[1], 0ul, gl_ScopeDevice, 0, gl_SemanticsRelaxed);
+    finish(total);
+}
+#endif
 #endif
diff --git a/engine/rr/ptx2rr.py b/engine/rr/ptx2rr.py
index 6f93ce5..d22a00f 100755
--- a/engine/rr/ptx2rr.py
+++ b/engine/rr/ptx2rr.py
@@ -139,31 +139,49 @@ void SA32(uint a, uint v) { CHECKS(a + 2u);
 // UNORM 10:10:10:2 surfaces (the pooled histories) are RGBA32F images holding the quantised values, like ZLUDA's.
 // tanh.approx.f32 as the reference computes it: ROCm's OCML tanh (device-libs ocml/src/tanhF.cl) with LLVM's AMDGPU
 // lowering of exp (lowerFEXP, fast-FMA path): the same fused operations, hardware exp2, ldexp and reciprocal.
-// (a fused f32 a * b + c through f64: the product is exact there and the sum is rounded to 53 bits before the f32
-// rounding. The driver's own f32 fma is a product and a sum whenever its product also exists on its own, as in
-// exp's ph / fma(x, c, -ph) pair.)
-float fmad(float a, float b, float c) { precise double s = double(a) * double(b) + double(c); return float(s); }
+// Up to 2026-10-19 every fused operation went through f64 (the product exact there, the sum rounded to 53 bits, then
+// to f32), which this GPU runs at a small fraction of the f32 rate. They are the driver's f32 fma now, with two
+// exceptions that keep every result bit for bit what it was, on all 2^32 arguments (RX 9070 XT, RADV; NaN payloads
+// included):
+// - the polynomial stays in f64: the driver drops the f32 rounding between nested f64 operations (it folds
+//   f2f64(f2f32(x)) to x; the conversions are not NoContraction), so its four nested steps always ran in f64 with one
+//   rounding at the end. In f32 (true fma), 24,836 arguments round one f32 step differently, 2 of them in binary16.
+// - exp's error term fma(w, c, -ph) next to ph = w * c: the driver makes that the product it already has and a sum,
+//   exactly 0, even with every operand precise or opaque. It is Dekker's exact product error: w split into two halves
+//   of 12 bits (w * 4097 = w * (2^12 + 1)), log2(e) into 1.4423828125 + 0.0003122091293334961 (11 and 12 bits), so
+//   every partial product is exact and so is each sum.
+// A NaN argument keeps its payload, quieted and positive, as the f64 conversion left it.
 float tanh_ocml(float x)
 {
     const float y = abs(x);
     precise float z;
     if (y < 0.625) {
         precise float y2 = y * y;
-        precise float p = fmad(y2, fmad(y2, fmad(y2, fmad(y2, -0.005700020585209131, 0.02063407190144062), -0.053737930953502655), 0.13331416249275208), -0.3333328068256378);
+        precise double y2d = double(y2);
+        precise double pd = (((y2d * -0.005700020585209131 + 0.02063407190144062) * y2d - 0.053737930953502655) * y2d +
+                             0.13331416249275208) * y2d - 0.3333328068256378;
+        precise float p = float(pd);
         precise float yp = y * p;
-        z = fmad(y2, yp, y);
+        z = fma(y2, yp, y);
     } else {
         const float w = 2.0 * y;
         precise float ph = w * 1.4426950216293335;
-        precise float f0 = fmad(w, 1.4426950216293335, -ph);
-        precise float pl = fmad(w, 1.925962855864327e-08, f0);
+        precise float wt = w * 4097.0;
+        precise float wh = wt - (wt - w);
+        precise float wl = w - wh;
+        precise float f1 = wh * 1.4423828125 - ph;
+        precise float f2 = f1 + wh * 0.0003122091293334961;
+        precise float f3 = f2 + wl * 1.4423828125;
+        precise float f0 = f3 + wl * 0.0003122091293334961;
+        precise float pl = fma(w, 1.925962855864327e-08, f0);
         const float e = roundEven(ph);
         precise float d = ph - e;
         precise float a = d + pl;
         float t = ldexp(exp2(a), int(e));
         if (w > 88.72283172607422) t = uintBitsToFloat(0x7f800000u);
         precise float s1 = t + 1.0;
-        z = fmad(-2.0, rcp_a(s1), 1.0);
+        z = fma(-2.0, rcp_a(s1), 1.0);
+        if (isnan(z)) z = uintBitsToFloat(floatBitsToUint(y) | 0x00400000u);
     }
     return x < 0.0 || floatBitsToUint(x) == 0x80000000u ? -z : z;
 }
@@ -1032,7 +1050,11 @@ class RRT(base.T):
         w(f'shared uint mms[{384 * (self.bx * self.by * self.bz // 32)}];   // mma_e4m3_f16 scratch, per subgroup')
 
     def runtime(self):
-        return (RR_RUNTIME + (ROOT / 'engine' / 'round_half.glsl').read_text() +
+        runtime = RR_RUNTIME
+        if getattr(self, 'unguarded', False):   # --unguarded: no opaque zero
+            assert '#define SUST_Z (gl_NumWorkGroups.z - 1u)' in RR_RUNTIME
+            runtime = RR_RUNTIME.replace('#define SUST_Z (gl_NumWorkGroups.z - 1u)', '#define SUST_Z 0u')
+        return (runtime + (ROOT / 'engine' / 'round_half.glsl').read_text() +
                 (ROOT / 'engine' / 'cnn_sr' / 'mma_f16.glsl').read_text())
 
 
@@ -1061,6 +1083,9 @@ def main():
     ap.add_argument('--oob-report', action='store_true', help='debugging: count out-of-range buffer accesses')
     ap.add_argument('--exact-convert', action='store_true',
                     help='the module goes through spv_exact_convert.py: lighter guards on binary16 conversions (CVT_GUARD)')
+    ap.add_argument('--unguarded', action='store_true',
+                    help="none of the exactness guards: no 'precise' registers and no opaque zero OR-ed into binary16 "
+                         "stores and conversions, so the driver may contract, fold and fuse (compile_rr2.py D4R_RR2_UNGUARDED)")
     ns = ap.parse_args()
     block = tuple(int(x) for x in ns.block.split(',')) if ns.block else default_block(ns.kernel)
     formats = {int(k): v for k, v in (x.split('=') for x in ns.surface)}
@@ -1077,8 +1102,9 @@ def main():
         t.probes = probes
         t.probe_block = tuple(int(x) for x in ns.probe_block.split(','))
         t.oob = ns.oob_report
+        t.unguarded = ns.unguarded
         return t
-    t = base.build(make, False, False, False)
+    t = base.build(make, False, ns.unguarded, False)
     for off, fmt in formats.items():
         if off in t.surfaces:
             t.surfaces[off] = fmt
diff --git a/engine/runtime.h b/engine/runtime.h
index e40763e..b511eef 100644
--- a/engine/runtime.h
+++ b/engine/runtime.h
@@ -273,7 +273,10 @@ struct RFrameSettings {
 };
 class RModel {
 public:
-    RModel(VkPhysicalDevice physical, VkDevice device, const std::string& directory, VkPipelineCache cache = VK_NULL_HANDLE);
+    // int64Atomics: the device has shaderBufferInt64Atomics enabled; preset F then measures the exposure in one pass
+    // (engine/rr/exposure_f.comp PASS 2) instead of two
+    RModel(VkPhysicalDevice physical, VkDevice device, const std::string& directory, VkPipelineCache cache = VK_NULL_HANDLE,
+           bool int64Atomics = false);
     ~RModel();
     RModel(const RModel&) = delete;
     RModel& operator=(const RModel&) = delete;
diff --git a/engine/runtime_m.inc b/engine/runtime_m.inc
index 803f18d..ca63544 100644
--- a/engine/runtime_m.inc
+++ b/engine/runtime_m.inc
@@ -512,9 +512,12 @@ void MEngine::recordFrame(VkCommandBuffer cb, VkQueryPool stamps, uint32_t first
     auto& e = *impl; auto& m = e.model;
     if (e.external) throw std::runtime_error("this engine's network is external: record the front and the back");
     recordFront(cb, stamps, first);
+    // timing only: D4R_ENGINE_M_LAYER_STAMPS=1 also stamps after every network layer, from query first + 7 on
+    static const bool layerStamps = std::getenv("D4R_ENGINE_M_LAYER_STAMPS") != nullptr;
     for (unsigned l = 0; l < MLayers; ++l) {
         const uint32_t w = e.gw / 2 / MScale[l], h = e.gh / 2 / MScale[l];
         dispatch(cb, m.network[l], e.layerSets[l], (w + (MShift[l] ^ ((e.frame & 1u) * 2u)) + 7) / 8, (h + MShift[l] + 7) / 8); emptyStage(cb);
+        if (stamps && layerStamps) vkCmdWriteTimestamp(cb, VK_PIPELINE_STAGE_ALL_COMMANDS_BIT, stamps, first + 7 + l);
     }
     recordBack(cb, stamps, stamps ? first + 3 : 0);
 }
@@ -525,11 +528,24 @@ std::vector<MEngine::DebugResource> MEngine::debugResources() const {
     std::vector<DebugResource> r{{"enc0-params", e.enc0P.handle, VK_NULL_HANDLE, e.enc0P.size, 0, 0},
                                  {"tokens", e.tokens.handle, VK_NULL_HANDLE, e.tokens.size, 0, 0},
                                  {"head", e.dec1Out.handle, VK_NULL_HANDLE, e.dec1Out.size, 0, 0},
+                                 // the network's own activations, layer by layer
+                                 {"skip1", e.skip1.handle, VK_NULL_HANDLE, e.skip1.size, 0, 0},
+                                 {"merged1", e.merged1.handle, VK_NULL_HANDLE, e.merged1.size, 0, 0},
+                                 {"skip2", e.skip2.handle, VK_NULL_HANDLE, e.skip2.size, 0, 0},
+                                 {"merged2", e.merged2.handle, VK_NULL_HANDLE, e.merged2.size, 0, 0},
+                                 {"tubeA", e.tubeA.handle, VK_NULL_HANDLE, e.tubeA.size, 0, 0},
+                                 {"tubeB", e.tubeB.handle, VK_NULL_HANDLE, e.tubeB.size, 0, 0},
+                                 {"dec2", e.dec2Out.handle, VK_NULL_HANDLE, e.dec2Out.size, 0, 0},
                                  {"gates", VK_NULL_HANDLE, e.gates.handle, VkDeviceSize(w) * h * 4, w, h},
                                  {"mix", VK_NULL_HANDLE, e.mix.handle, VkDeviceSize(w) * h * 4, w, h},
                                  {"cov", VK_NULL_HANDLE, e.cov.handle, VkDeviceSize(w) * h * 8, w, h},
                                  {"mvA", VK_NULL_HANDLE, e.mvA[1 - e.current].handle, VkDeviceSize(w) * h * 4, w, h},
-                                 {"feat", VK_NULL_HANDLE, e.feat[1 - e.current].handle, VkDeviceSize(w) * h * 4, w, h}};
+                                 {"feat", VK_NULL_HANDLE, e.feat[1 - e.current].handle, VkDeviceSize(w) * h * 4, w, h},
+                                 // the reconstruction's outputs: the history, its half-resolution mean and the packed
+                                 // image the downsample reads
+                                 {"hist", VK_NULL_HANDLE, e.hist[1 - e.current].handle, VkDeviceSize(e.hw) * e.hh * 8, e.hw, e.hh},
+                                 {"half", VK_NULL_HANDLE, e.half[1 - e.current].handle, VkDeviceSize(e.hw / 2) * (e.hh / 2) * 16, e.hw / 2, e.hh / 2},
+                                 {"packed", e.packed.handle, VK_NULL_HANDLE, e.packed.size, 0, 0}};
     if (e.model.presetL) r.push_back({"skip", e.skipL.handle, VK_NULL_HANDLE, e.skipL.size, 0, 0});
     return r;
 }
diff --git a/engine/runtime_rr.inc b/engine/runtime_rr.inc
index 825035e..04908c6 100644
--- a/engine/runtime_rr.inc
+++ b/engine/runtime_rr.inc
@@ -90,7 +90,7 @@ struct RModel::Impl {
     bool uploaded = false;
     Impl(VkPhysicalDevice physical, VkDevice dev) : device(physical, dev) {}
 };
-RModel::RModel(VkPhysicalDevice physical, VkDevice device, const std::string& directory, VkPipelineCache cache)
+RModel::RModel(VkPhysicalDevice physical, VkDevice device, const std::string& directory, VkPipelineCache cache, bool int64Atomics)
     : impl(std::make_unique<Impl>(physical, device)) {
     auto& m = *impl;
     // rr.txt: "D4RR0001", "preset N", one "stage NAME SHADER.spv SUBGROUP WEIGHTS TEXTURES SURFACES" line per stage in
@@ -103,7 +103,7 @@ RModel::RModel(VkPhysicalDevice physical, VkDevice device, const std::string& di
     if (line.rfind("D4RR0002", 0) == 0) {
         m.preset = 6;
         m.v2 = std::make_unique<R2Model>();
-        m.v2->load(m.device, directory, lines, cache);
+        m.v2->load(m.device, directory, lines, cache, int64Atomics);
         return;
     }
     if (line.rfind("D4RR0001", 0) != 0) throw std::runtime_error("not a Ray Reconstruction model (rr.txt D4RR0001)");
diff --git a/engine/runtime_rr2.inc b/engine/runtime_rr2.inc
index 85e8fcc..dfdf963 100644
--- a/engine/runtime_rr2.inc
+++ b/engine/runtime_rr2.inc
@@ -44,9 +44,10 @@ struct R2Model {
     Buffer aux;                   // encoder 0's second table (argument 104)
     std::vector<char> auxPending;
     std::array<Pipeline, 2> exposure;
+    Pipeline exposureOne;         // both exposure passes in one (exposure_f2.spv, 64-bit atomics); else null
     bool particles = false;       // the model has R2DownP
     bool gather = false;          // the model has R2Gather
-    void load(const Device& d, const std::string& directory, std::istringstream& lines, VkPipelineCache cache) {
+    void load(const Device& d, const std::string& directory, std::istringstream& lines, VkPipelineCache cache, bool int64Atomics) {
         const auto b = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, t = VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER, s = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE;
         auto split = [](const std::string& list) {
             std::vector<std::string> items;
@@ -108,6 +109,8 @@ struct R2Model {
         blob(aux, auxPending, read(directory + "/enc0-aux.bin"));
         exposure[0].create(d.handle, directory + "/exposure_f0.spv", {b, t, b}, 64, cache);
         exposure[1].create(d.handle, directory + "/exposure_f1.spv", {b, b, b, t, s}, 64, cache);
+        if (int64Atomics && exists(directory + "/exposure_f2.spv"))
+            exposureOne.create(d.handle, directory + "/exposure_f2.spv", {b, t, b, b, t, s}, 64, cache);
     }
     void recordUpload(VkCommandBuffer cb) {   // the weights into the command buffer (recordContents)
         for (auto& st : stages) recordContents(cb, st.weights, st.pending);
@@ -172,7 +175,7 @@ struct R2Engine {
         std::vector<VkDescriptorSet> down;
         std::vector<std::array<VkDescriptorSet, 2>> downP;
         VkDescriptorSet pooled = VK_NULL_HANDLE, exposure0 = VK_NULL_HANDLE;
-        std::array<VkDescriptorSet, 2> exposure1{};
+        std::array<VkDescriptorSet, 2> exposure1{}, exposureOne{};
     };
     std::vector<Bindings> bindings;
     const Bindings* bound = nullptr;
@@ -337,7 +340,7 @@ struct R2Engine {
         exposure.blocksX = (rw + 15) / 16; exposure.blocks = exposure.blocksX * ((rh + 15) / 16);
         exposure.last[0] = static_cast<int32_t>(rw - 1); exposure.last[1] = static_cast<int32_t>(rh - 1);
         exposure.invCount = static_cast<float>(10000.0 / (4294967296.0 * rw * rh));
-        sums.create(d, exposure.blocks * 8, storage);
+        sums.create(d, std::max(exposure.blocks, 2u) * 8, storage);   // the one-pass exposure: the total and the count
         const size_t tiles = src.size();
         auto place = [&](size_t bytes) { const VkDeviceSize at = paramBytes; paramBytes += (bytes + 255) & ~VkDeviceSize(255); return at; };
         for (unsigned i = 0; i < R2Network; ++i) { blocks[i] = model.stages[i].args; paramAt[i] = place(blocks[i].size()); }
@@ -593,17 +596,29 @@ struct R2Engine {
                     b.downP.push_back({stageSet(b, R2DownP, downPAt[i], 0, output), stageSet(b, R2DownP, downPAt[i], 1, output)});
             }
             b.pooled = stageSet(b, R2Pooled, pooledAt, 0, output);
-            b.exposure0 = descriptorSet(d.handle, b.pool, model.exposure[0].descriptors);
-            bindRange(d.handle, b.exposure0, 0, paramAll, expAt);
-            bindImage(d.handle, b.exposure0, 1, images.color, nearest, images.layout);
-            bindBuffer(d.handle, b.exposure0, 2, sums);
-            for (uint32_t par = 0; par < 2; ++par) {
-                VkDescriptorSet set = b.exposure1[par] = descriptorSet(d.handle, b.pool, model.exposure[1].descriptors);
-                bindRange(d.handle, set, 0, paramAll, expAt);
-                bindBuffer(d.handle, set, 1, sums);
-                bindBuffer(d.handle, set, 2, values);
-                bindImage(d.handle, set, 3, e[par ^ 1].view, nearest);
-                bindImage(d.handle, set, 4, e[par].view);
+            if (model.exposureOne.handle)
+                for (uint32_t par = 0; par < 2; ++par) {
+                    VkDescriptorSet set = b.exposureOne[par] = descriptorSet(d.handle, b.pool, model.exposureOne.descriptors);
+                    bindRange(d.handle, set, 0, paramAll, expAt);
+                    bindImage(d.handle, set, 1, images.color, nearest, images.layout);
+                    bindBuffer(d.handle, set, 2, sums);
+                    bindBuffer(d.handle, set, 3, values);
+                    bindImage(d.handle, set, 4, e[par ^ 1].view, nearest);
+                    bindImage(d.handle, set, 5, e[par].view);
+                }
+            else {
+                b.exposure0 = descriptorSet(d.handle, b.pool, model.exposure[0].descriptors);
+                bindRange(d.handle, b.exposure0, 0, paramAll, expAt);
+                bindImage(d.handle, b.exposure0, 1, images.color, nearest, images.layout);
+                bindBuffer(d.handle, b.exposure0, 2, sums);
+                for (uint32_t par = 0; par < 2; ++par) {
+                    VkDescriptorSet set = b.exposure1[par] = descriptorSet(d.handle, b.pool, model.exposure[1].descriptors);
+                    bindRange(d.handle, set, 0, paramAll, expAt);
+                    bindBuffer(d.handle, set, 1, sums);
+                    bindBuffer(d.handle, set, 2, values);
+                    bindImage(d.handle, set, 3, e[par ^ 1].view, nearest);
+                    bindImage(d.handle, set, 4, e[par].view);
+                }
             }
             bindings.push_back(std::move(b));
             found = bindings.end() - 1;
@@ -675,9 +690,15 @@ struct R2Engine {
             vkCmdUpdateBuffer(cb, paramAll.handle, at, std::min<VkDeviceSize>(65536, paramBytes - at), all.data() + at);
         memoryBarrier(cb, VK_PIPELINE_STAGE_TRANSFER_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, VK_ACCESS_TRANSFER_WRITE_BIT,
                       VK_ACCESS_SHADER_READ_BIT | VK_ACCESS_SHADER_WRITE_BIT);
-        dispatch(cb, model.exposure[0], b.exposure0, exposure.blocksX, exposure.blocks / exposure.blocksX);
-        step();
-        dispatch(cb, model.exposure[1], b.exposure1[cur], 1, 1);
+        // the exposure: one pass with 64-bit atomics (the last block finishes, as NGX's kernel), or the block sums and
+        // one workgroup that adds them
+        if (model.exposureOne.handle)
+            dispatch(cb, model.exposureOne, b.exposureOne[cur], exposure.blocksX, exposure.blocks / exposure.blocksX);
+        else {
+            dispatch(cb, model.exposure[0], b.exposure0, exposure.blocksX, exposure.blocks / exposure.blocksX);
+            step();
+            dispatch(cb, model.exposure[1], b.exposure1[cur], 1, 1);
+        }
         step();
         stamp(1);
         for (unsigned stage = 0; stage < R2Network; ++stage) {
diff --git a/engine/test_m.py b/engine/test_m.py
index 257e4c0..ffa273e 100644
--- a/engine/test_m.py
+++ b/engine/test_m.py
@@ -90,7 +90,7 @@ def defines(name):
     # compile_m.py's arithmetic (ZLUDA's lowering); M_TEXL=0 gives the native accuracy kernels' for the network layers.
     # M_DEFS: extra defines for every layer_m.comp stage; M_DEFS_ENC1, M_DEFS_ENC3_TUBE, ...: for one layer
     texl = ['-DTEXL=1'] if kind >= 3 or os.environ.get('M_TEXL', '1') != '0' else []
-    # M_WPW=2: two 4x4-token windows per wave in the network layers (layer_m.comp WPW; not L's input / expansion)
+    # M_WPW=2: two 4x4-token windows per wave in the network layers (layer_m.comp WPW; L's input / expansion: M_DEFS_*)
     wpw = [f"-DWPW={os.environ['M_WPW']}"] if kind < 3 and os.environ.get('M_WPW') else []
     return [f'-DC={C}', f'-DNH={heads}', f'-DKIND={kind}', f'-DXC={NPM or CIN or 16}'] + texl + wpw + os.environ.get('M_DEFS', '').split() + os.environ.get('M_DEFS_' + name.upper(), '').split()
 
diff --git a/install.sh b/install.sh
index c16cf92..5a77ba7 100755
--- a/install.sh
+++ b/install.sh
@@ -168,15 +168,18 @@ if [[ ! -f "$BUILDER" ]]; then
   done_msg; exit 0
 fi
 # the Super Resolution models depend on the GPU (RDNA2 builds no L and M), so the builder lists them
-models=(rr-d rr-e rr-f $(bash "$BUILDER" --list-sr-models))
+models=(rr-d rr-e rr-f rr-d-gfx11 rr-e-gfx11 $(bash "$BUILDER" --list-sr-models))
 # One build serves every game with the same DLLs, builder and model set: the key covers them and the accuracy choice.
 key() {   # key EXACT(0|1)
   { echo "exact=$1"
     echo "models=${models[*]}"
+    # Model compilers accept per-layer K_DEFS_*, M_* arithmetic/tile options, and D4R_RR_* choices.
+    # A cache hit must not silently replace a model built with different compiler options.
+    env | LC_ALL=C sort | sed -n -e '/^K_/p' -e '/^M_/p' -e '/^D4R_RR_/p'
     sha256sum -- "$RUNTIME/nvngx_dlss.dll" "$RUNTIME/nvngx_dlssd.dll" 2>/dev/null | cut -d' ' -f1
     (cd "$RUNTIME/tools" && find . -type f ! -name '*.pyc' -print0 | LC_ALL=C sort -z | xargs -0 sha256sum); } | sha256sum | cut -c1-16
 }
-have() { local m; for m in "${models[@]}"; do [[ -f "$1/$m/model.json" || -f "$1/$m/rr.txt" || -f "$1/$m/e.txt" ]] || return 1; done; }
+have() { python3 "$RUNTIME/tools/models/model_tools.py" validate "$1" "${models[@]}" 2>/dev/null; }
 STAMP="$RUNTIME/engine/.d4r-models"
 if [[ "$EXACT" == ask ]]; then
   EXACT=0
diff --git a/kernels/build.sh b/kernels/build.sh
index f170b30..7340150 100755
--- a/kernels/build.sh
+++ b/kernels/build.sh
@@ -5,8 +5,10 @@
 # usage: kernels/build.sh [all|k|l|m|tex|rr|rrswin] [OUT_DIR] (default: all, kernels/out/native)
 #
 #   D4R_ROCM_DIR    ROCm installation with clang and the HIP device libraries (default /opt/rocm)
-#   D4R_GPU_ARCH    target GPU (default gfx1101): an RDNA3 (gfx110x) or RDNA4 (gfx120x) target; the kernels
-#                   need wave32 WMMA, and kernels/common/wmma_layout.h picks the register layout per target
+#   D4R_GPU_ARCH    target GPU (default gfx1101): an RDNA3 (gfx110x) or RDNA4 (gfx120x) target, where
+#                   kernels/common/wmma_layout.h picks the WMMA register layout per target, or an RDNA2 (gfx103x)
+#                   target: there only K's layers (software matrix products) and the texture kernels that
+#                   need no WMMA (tex: K's natively translated hiluma output and input kernels) are built
 #   D4R_NATIVE_FP8  1 (RDNA4 only): the variant for ZLUDA's native FP8 WMMA (d4r.ini [Kernels] NativeFp8):
 #                   M's layers on FP8 WMMA with NVIDIA's e4m3 activations, texture tails matching ZLUDA's
 #                   D4R_ZLUDA_WMMA_FP8_NATIVE lowering. K is the same in both variants.
@@ -45,6 +47,12 @@ CLANG="$ROCM/lib/llvm/bin/clang++"
 [[ -x "$CLANG" ]] || { echo "clang++ not found in $ROCM/lib/llvm/bin (set D4R_ROCM_DIR)" >&2; exit 2; }
 mkdir -p "$OUT"
 case "$WHAT" in all|k|l|m|tex|rr|rrswin) ;; *) echo "unknown kernel family: $WHAT" >&2; exit 2 ;; esac
+# RDNA2 has no WMMA instruction: M/L's Swin layers, their texture tails and Ray Reconstruction call it directly
+RDNA2=0
+if [[ "$ARCH" == gfx103* ]]; then
+    RDNA2=1
+    case "$WHAT" in m|l|rr|rrswin) echo "$WHAT kernels need WMMA, which RDNA2 ($ARCH) does not have" >&2; exit 2 ;; esac
+fi
 # Never certify a directory containing older fast binaries as an accuracy set.
 if [[ "$ACCURACY" == 1 ]] && compgen -G "$OUT/*.hsaco" >/dev/null &&
     { [[ ! -f "$OUT/d4r-accuracy.txt" ]] || [[ "$(cat "$OUT/d4r-accuracy.txt")" != 1 ]]; }; then
@@ -82,14 +90,15 @@ if [[ "$WHAT" == all || "$WHAT" == k ]]; then
     echo "== DLSS 4 (preset K) transformer layers"
     for src in "$HERE"/k/dltss_pwin_*.hip; do build_hip "$src"; done
 fi
-if [[ "$WHAT" == all || "$WHAT" == m || "$WHAT" == l ]]; then
+if [[ "$RDNA2" == 0 && ( "$WHAT" == all || "$WHAT" == m || "$WHAT" == l ) ]]; then
     echo "== DLSS 4.5 (presets L/M) shared Swin layers"
     for src in "$HERE"/m/rrlite_*.hip; do build_hip "$src" "$([[ "$FP8" == 1 ]] && echo -DD4R_FP8_WMMA)"; done
 fi
 if [[ "$WHAT" == all || "$WHAT" == tex || "$WHAT" == l ]]; then
     echo "== texture kernels (NVIDIA PTX with native parts)"
     : "${D4R_DLSS_DLL:?set D4R_DLSS_DLL to nvngx_dlss.dll (310.7)}"
-    : "${D4R_ZLUDA_EMIT:?set D4R_ZLUDA_EMIT to d4r_emit from a ZLUDA build with patches/zluda}"
+    # RDNA2 builds only the PTX-to-HIP translations below, which ZLUDA does not compile
+    [[ "$RDNA2" == 1 ]] || : "${D4R_ZLUDA_EMIT:?set D4R_ZLUDA_EMIT to d4r_emit from a ZLUDA build with patches/zluda}"
     PTX_DIR="${D4R_DLSS_PTX_DIR:-$HERE/extracted/ptx}"
     if [[ -z "${D4R_DLSS_PTX_DIR:-}" ]]; then
         python3 "$HERE/tools/extract_dlss_ptx.py" "$D4R_DLSS_DLL" "$PTX_DIR"
@@ -152,7 +161,10 @@ if [[ "$WHAT" == all || "$WHAT" == tex || "$WHAT" == l ]]; then
             for range in hdr ldr; do native_ins+=("hiluma_engine_input_${depth}_${mv}_${range}_v2_rel"); done
         done
     done
-    for spec in "${specs[@]}"; do
+    # RDNA2: none of the ZLUDA-compiled texture builds above (their tails call WMMA, and the sust_only builds
+    # are not validated there); only the natively translated hiluma kernels below
+    [[ "$RDNA2" == 1 ]] && specs=()
+    for spec in "${specs[@]+"${specs[@]}"}"; do
         IFS=: read -r kernel src mode <<< "$spec"
         [[ "$ACCURACY" == 0 && " ${native_outs[*]} " == *" $kernel "* ]] && continue
         # accuracy sets are wave32, except M's post and downsample kernels on RDNA4: their wave64 builds keep the
@@ -180,8 +192,13 @@ if [[ "$WHAT" == all || "$WHAT" == tex || "$WHAT" == l ]]; then
             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
+            # RDNA3/4: wave64 (FP32 dual issue; 0.49 -> 0.44 ms on an RX 7700 XT, same bytes). RDNA2 issues wave64
+            # in two passes, so it gets wave32 (no stack frame, 96 VGPRs), and WGP mode: a 256-thread workgroup
+            # then takes 2 of each SIMD's 10 wave slots instead of 4, so 10 waves fit per SIMD instead of 8
+            # (measured with RDNA2's occupancy on the RX 7700 XT: 0.596 -> 0.568 ms, same bytes).
+            wave=(); [[ "$RDNA2" == 1 ]] && wave=(WAVE=32 CUMODE=0)
+            env "${wave[@]}" 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; }
@@ -192,7 +209,7 @@ if [[ "$WHAT" == all || "$WHAT" == tex || "$WHAT" == l ]]; then
         done
     fi
 fi
-if [[ "$WHAT" == rr || ( "$WHAT" == all && -n "${D4R_DLSSD_DLL:-}" ) ]]; then
+if [[ "$WHAT" == rr || ( "$WHAT" == all && "$RDNA2" == 0 && -n "${D4R_DLSSD_DLL:-}" ) ]]; then
     echo "== Ray Reconstruction texture kernels"
     : "${D4R_DLSSD_DLL:?set D4R_DLSSD_DLL to the CUDA-capable nvngx_dlssd.dll}"
     : "${D4R_ZLUDA_EMIT:?set D4R_ZLUDA_EMIT to d4r_emit from a patched ZLUDA build}"
diff --git a/kernels/k/pwin_common.h b/kernels/k/pwin_common.h
index e112170..0cd2d73 100644
--- a/kernels/k/pwin_common.h
+++ b/kernels/k/pwin_common.h
@@ -61,6 +61,17 @@ __device__ __forceinline__ uint32_t lane_id()
 #undef PWIN_NATIVE_F16ACC
 #undef PWIN_F16ACC_NATIVE
 #endif
+// RDNA2: the weight images use plain global loads. The raw-buffer path below builds its resource with gfx11's
+// dword3 (0x31004000: unified format 4, 8_SSCALED), which gfx11 ignores for untyped loads; gfx10 descriptors want
+// a 32-bit format there (AMD's own libraries use 0x31014000 on gfx103x), and no gfx10 encoding has been validated
+// on RDNA2 hardware, so RDNA2 takes no descriptor at all. Keyed on D4R_WMMA_SOFT (always set on gfx10) so that the
+// RDNA3 validation build of the RDNA2 kernels runs the same loads.
+#ifdef D4R_WMMA_SOFT
+#ifndef PWIN_NO_BUFFER_IMG
+#define PWIN_NO_BUFFER_IMG
+#endif
+#undef PWIN_GELU_LUT_BUFFER
+#endif
 typedef wm_op op_t; // WMMA operand: u8v (gfx11) or u4v (gfx12)
 
 // one k16 step with the f16 accumulator of NVIDIA's f16 wmma (rounded after the step)
diff --git a/kernels/native/build_native.sh b/kernels/native/build_native.sh
index 73d2ff3..4a8beaf 100755
--- a/kernels/native/build_native.sh
+++ b/kernels/native/build_native.sh
@@ -1,5 +1,6 @@
 #!/usr/bin/env bash
 # build_native.sh SRC.hip OUT.hsaco [extra clang flags]: natively compiled DLSS texture kernel (gfx1101, wave64 default)
+# WAVE=32|64 (default 64); CUMODE=1|0 (default 1): CU mode, or WGP mode (a workgroup's waves on all four SIMDs)
 set -euo pipefail
 ND=$(cd "$(dirname "$0")" && pwd)
 SRC=$(realpath "$1") OUT=$(realpath -m "$2"); shift 2
@@ -8,9 +9,10 @@ R=${D4R_ROCM_DIR:-/opt/rocm}
 ARCH=${D4R_GPU_ARCH:-gfx1101}
 W=${WAVE:-64}
 WF=$([[ $W == 64 ]] && echo -mwavefrontsize64 || echo -mno-wavefrontsize64)
+CU=$([[ ${CUMODE:-1} == 1 ]] && echo -mcumode || echo -mno-cumode)
 D=$(cd "$(dirname "$SRC")" && pwd)
 T=$(mktemp -d); trap 'rm -rf "$T"' EXIT
-(cd "$T" && "$R/lib/llvm/bin/clang++" -x hip -std=c++20 --offload-arch=$ARCH --offload-device-only -O3 $WF -mcumode \
+(cd "$T" && "$R/lib/llvm/bin/clang++" -x hip -std=c++20 --offload-arch=$ARCH --offload-device-only -O3 $WF $CU \
   -ffp-contract=off -fgpu-flush-denormals-to-zero -fno-fast-math -I"$D" -I"$ND" -I"$R/include" --rocm-path="$R" \
   --rocm-device-lib-path="$R/amdgcn/bitcode" "$@" -o "$OUT" "$SRC" -save-temps=cwd --no-gpu-bundle-output)
 S=$(ls "$T"/*.s | head -1)
diff --git a/kernels/native/output_tile.h b/kernels/native/output_tile.h
index dd9c786..0130d96 100644
--- a/kernels/native/output_tile.h
+++ b/kernels/native/output_tile.h
@@ -24,8 +24,9 @@ RT void output_color_pack(f4 color, float exposure, uint32_t* a, uint32_t* b)
         const float q = div_a(ln, 9.0f) + asf(0xBC4FAE29u);
         weight = rcp_a(y) * q;
     }
-    *a = pack16(ashu((h1)(y * weight)), ashu((h1)((exposure * co) * weight)));
-    *b = pack16(ashu((h1)((exposure * cg) * weight)), ashu((h1)color[3]));
+    // f2h_rn: rounded to f32 and then to f16, as ZLUDA's compile does (no v_fma_mixlo of the product)
+    *a = pack16(ashu(f2h_rn(y * weight)), ashu(f2h_rn((exposure * co) * weight)));
+    *b = pack16(ashu(f2h_rn((exposure * cg) * weight)), ashu(f2h_rn(color[3])));
 }
 
 RT int output_tile_coord(int relative, int lo, int hi)
diff --git a/kernels/native/ptx2hip.py b/kernels/native/ptx2hip.py
index af5ec51..acac466 100755
--- a/kernels/native/ptx2hip.py
+++ b/kernels/native/ptx2hip.py
@@ -593,10 +593,10 @@ class Translator:
         if dt == 'f16' and stp == 'f32':
             if 'rn' not in mods:
                 raise Err(op)
-            A(a[0], f'(h1){V(a[1], "f32")}', 'f16'); return
+            A(a[0], f'f2h_rn({V(a[1], "f32")})', 'f16'); return
         if dt == 'f16x2' and stp == 'f32':
             # d = {cvt(a) in the high half, cvt(b) in the low half}
-            A(a[0], f'pack16(ashu((h1){V(a[2], "f32")}), ashu((h1){V(a[1], "f32")}))', 'u32'); return
+            A(a[0], f'pack16(ashu(f2h_rn({V(a[2], "f32")})), ashu(f2h_rn({V(a[1], "f32")})))', 'u32'); return
         if dt == 'f16' and stp == 'f16':
             x = V(a[1], 'f16')
             if 'rni' in mods:
diff --git a/kernels/native/ptx_rt.h b/kernels/native/ptx_rt.h
index 3761833..690dd03 100644
--- a/kernels/native/ptx_rt.h
+++ b/kernels/native/ptx_rt.h
@@ -77,6 +77,16 @@ RT uint32_t f2u_rz(float x)
 }
 
 // ---------------------------------------------------------------- f16 (denormals kept)
+// cvt.rn.f16.f32: the f32 value rounded to f16. The empty asm keeps LLVM from folding the conversion into the
+// f32 operation that produced the value (fptrunc(fma(fpext ...)) -> v_fma_mixlo/mixhi_f16), which rounds the exact
+// result to f16 once instead of to f32 and then to f16: on an RX 7700 XT that differed from the two-step result for
+// 1 in ~29,000 random f16 operand triples, and in the hiluma output kernel it left single 1-ulp seeds that the
+// temporal history then spread (ZLUDA keeps both roundings).
+RT h1 f2h_rn(float x)
+{
+    __asm__ volatile("" : "+v"(x));
+    return (h1)x;
+}
 RT h1 hfma(h1 a, h1 b, h1 c) { return __builtin_fmaf16(a, b, c); }
 RT h2 h2fma(h2 a, h2 b, h2 c) { return __builtin_elementwise_fma(a, b, c); }
 RT h1 hmaxn(h1 a, h1 b) { return __builtin_fmaximum_numf16(a, b); }
diff --git a/packaging/D4R_README.txt b/packaging/D4R_README.txt
index 2fe6142..2d67c9d 100644
--- a/packaging/D4R_README.txt
+++ b/packaging/D4R_README.txt
@@ -70,6 +70,10 @@ that archive's d4r/README.txt: setup runs automatically when the ASI loads.
 @installer    compatibility tools (or set D4R_PROTON=/path/to/GE-Proton11-3). It takes a minute or two once;
 @installer    the models are kept in ~/.cache/d4r/engine-models for the next game. bash install.sh --help lists
 @installer    options for scripted installs.
+@installer    RR D/E also have separate gfx11 models; devices without compatible matrix support retain CUDA
+@installer    translation. Builds report progress and time out after 900 seconds per command (override with
+@installer    D4R_MODEL_TIMEOUT). Failed captures and logs are retained under ~/.cache/d4r/model-builds;
+@installer    D4R_MODEL_LOG_DIR changes that location. Model checksums are verified before cache reuse.
 @installer 3. In Steam, open the game's Properties:
 @clean 3. In Steam, open the game's Properties:
 @full 2. In Steam, open the game's Properties:
diff --git a/packaging/build_engine_models.sh b/packaging/build_engine_models.sh
index 8629a89..6bbc359 100755
--- a/packaging/build_engine_models.sh
+++ b/packaging/build_engine_models.sh
@@ -84,8 +84,28 @@ export PATH="$KIT/bin:$PATH"   # the bundled glslangValidator
 glslangValidator --version > /dev/null 2>&1 || fail "glslangValidator ($KIT/bin) does not run on this system"
 sha() { sha256sum -- "$1" | cut -d' ' -f1; }
 
-WORK="$(mktemp -d "${TMPDIR:-/tmp}/d4r-models.XXXXXX")"
-trap 'kill $(jobs -p) 2> /dev/null || true; rm -rf -- "$WORK"' EXIT
+LOG_ROOT="${D4R_MODEL_LOG_DIR:-${XDG_CACHE_HOME:-$HOME/.cache}/d4r/model-builds}"
+mkdir -p -- "$LOG_ROOT"
+BUILD_LOGS="$(mktemp -d "$LOG_ROOT/build.XXXXXX")"
+WORK="$BUILD_LOGS/work"; mkdir -p "$WORK"
+MODEL_TOOLS="$KIT/model_tools.py"
+[[ -f "$MODEL_TOOLS" ]] || fail "missing $MODEL_TOOLS; update the model builder tools"
+cleanup() {
+  local status=$?
+  kill $(jobs -p) 2> /dev/null || true
+  if (( status == 0 )); then
+    while IFS= read -r -d '' log; do
+      relative="${log#"$WORK/"}"
+      mkdir -p -- "$BUILD_LOGS/$(dirname -- "$relative")"
+      cp -- "$log" "$BUILD_LOGS/$relative"
+    done < <(find "$WORK" -name '*.log' -type f -print0)
+    rm -rf -- "$WORK"
+  else
+    printf 'build-engine-models: failed build and captures preserved at %s\n' "$WORK" >&2
+  fi
+  printf 'Model build diagnostics: %s\n' "$BUILD_LOGS"
+}
+trap cleanup EXIT
 trap 'exit 130' INT
 trap 'exit 143' TERM
 
@@ -95,6 +115,7 @@ place() {   # place NAME DIR FILE...
   shift 2
   mkdir -p "$RUNTIME/engine" "$WORK/place/$name"
   for f in "$@"; do cp -- "$src/$f" "$WORK/place/$name/"; done
+  python3 "$MODEL_TOOLS" manifest "$WORK/place/$name"
   rm -rf -- "$dst.new"; mv -- "$WORK/place/$name" "$dst.new"
   rm -rf -- "$dst.old"; [[ ! -e "$dst" ]] || mv -- "$dst" "$dst.old"
   mv -- "$dst.new" "$dst"; rm -rf -- "$dst.old"
@@ -106,22 +127,28 @@ model_files() {   # rr.txt and the files it names, for model directory $1
 }
 run() {   # run LOG COMMAND...: the command's output in WORK/LOG, shown when it fails
   local log="$WORK/$1"; shift
-  "$@" > "$log" 2>&1 || { tail -n 20 "$log" >&2; fail "failed: $(basename -- "$2") (log above)"; }
+  python3 "$MODEL_TOOLS" run --log "$log" --timeout "${D4R_MODEL_TIMEOUT:-900}" -- "$@" ||
+    fail "failed: $1 (see $log)"
 }
 
 if (( DO_RR )); then
   DLL="$RUNTIME/nvngx_dlssd.dll"
   [[ -f "$DLL" && "$(sha "$DLL")" == "$DLSSD_SHA256" ]] || fail "Ray Reconstruction models need nvngx_dlssd.dll 310.9.1 in $RUNTIME (run install.sh)"
   echo 'Building the Ray Reconstruction models (presets D, E and F) from nvngx_dlssd.dll...'
+  # Build both portable arithmetic variants; runtime selects by device capabilities. Do not let the build
+  # host (or an inherited developer switch) silently determine what another game/GPU receives from the cache.
   for p in D E; do
-    l=$(tr DE de <<< "$p")
-    bash "$KIT/engine/rr/build_native_model.sh" --dll "$DLL" "$p" "$WORK/rr-$l" > "$WORK/rr-$l.log" 2>&1 ||
-      { tail -n 20 "$WORK/rr-$l.log" >&2; fail "preset $p model failed"; }
-    place "rr-$l" "$WORK/rr-$l" $(model_files "$WORK/rr-$l") exposure_rr0.spv exposure_rr1.spv
+    for arithmetic in fp8 gfx11; do
+      l=$(tr DE de <<< "$p"); name="rr-$l"; gfx11=0
+      if [[ "$arithmetic" == gfx11 ]]; then name+="-gfx11"; gfx11=1; fi
+      run "$name.log" env D4R_RR_GFX11="$gfx11" bash "$KIT/engine/rr/build_native_model.sh" \
+        --dll "$DLL" "$p" "$WORK/$name"
+      place "$name" "$WORK/$name" $(model_files "$WORK/$name") exposure_rr0.spv exposure_rr1.spv
+    done
   done
   run ptx-f.log python3 "$KIT/kernels/tools/extract_dlss_ptx.py" "$DLL" "$WORK/ptx-f"
   run rr-f.log python3 "$KIT/engine/compile_rr2.py" --dll "$DLL" "$WORK/rr-f" --ptx-dir "$WORK/ptx-f"
-  place rr-f "$WORK/rr-f" $(model_files "$WORK/rr-f") enc0-aux.bin exposure_f0.spv exposure_f1.spv
+  place rr-f "$WORK/rr-f" $(model_files "$WORK/rr-f") enc0-aux.bin exposure_f0.spv exposure_f1.spv exposure_f2.spv
 fi
 
 if (( DO_SR )); then
@@ -149,16 +176,17 @@ if (( DO_SR )); then
     local w="$WORK/$1"
     mkdir -p "$w/replay"
     printf '  rendering one frame with model %s through d4r (GE-Proton %s)...\n' "$1" "$(basename "$PROTON")"
-    (cd "$APP" && env STEAM_COMPAT_CLIENT_INSTALL_PATH="$STEAM_ROOT" STEAM_COMPAT_DATA_PATH="$WORK/prefix" \
+    (cd "$APP" && python3 "$MODEL_TOOLS" run --log "$w/run.log" --timeout "${D4R_MODEL_TIMEOUT:-900}" -- \
+      env STEAM_COMPAT_CLIENT_INSTALL_PATH="$STEAM_ROOT" STEAM_COMPAT_DATA_PATH="$WORK/prefix" \
       PROTON_FORCE_NVAPI=1 DXVK_NVAPI_GPU_ARCH=AD100 WINEDLLOVERRIDES='d3d12=n,b;d3d12core=n,b;d4r12core=n,b' \
       D4R_ENGINE=0 D4R_DLSS_PRESET="$2" D4R_HARNESS_CREATE_FLAGS="${4:-0x4B}" D4R_SHIM_LOG="$(winpath "$w/shim.log")" \
+      D4R_CUDA_TRACE="$(winpath "$w/cuda.log")" D4R_CUDA_VERBOSE=1 \
       D4R_CUDA_REPLAY_DUMP_DIR="$w/replay" D4R_CUDA_REPLAY_DUMP_LIMIT=60 \
       "$PROTON/proton" run "$APP/d3d12_dlss_harness.exe" "$(winpath "$RUNTIME/nvngx.dll")" "$(winpath "$w/out.raw")" \
-      1 640 360 1280 720) > "$w/run.log" 2>&1 || true
-    if ! ls "$w/replay" | grep -q "$3"; then
-      tail -n 15 "$w/run.log" >&2
+      1 640 360 1280 720) || fail "model $1: capture process failed; diagnostics in $w"
+    if ! python3 "$MODEL_TOOLS" capture "$w" "$1"; then
       [[ ! -f "$w/shim.log" ]] || tail -n 15 "$w/shim.log" >&2
-      fail "model $1: the test frame did not run through DLSS (logs above)"
+      fail "model $1: incomplete or invalid test frame (diagnostics in $w)"
     fi
   }
   mkdir -p "$WORK/prefix"
diff --git a/packaging/model_tools.py b/packaging/model_tools.py
new file mode 100644
index 0000000..17922b5
--- /dev/null
+++ b/packaging/model_tools.py
@@ -0,0 +1,167 @@
+#!/usr/bin/env python3
+"""Bounded model-build commands and integrity checks (also shipped in d4r/tools)."""
+import argparse
+import hashlib
+import json
+import os
+from pathlib import Path
+import signal
+import subprocess
+import sys
+import time
+
+MANIFEST = 'files.sha256.json'
+MODELS = ('rr-d', 'rr-e', 'rr-f', 'rr-d-gfx11', 'rr-e-gfx11', 'k', 'k-ldr', 'm', 'm-ldr', 'l', 'l-ldr')
+
+
+def digest(path):
+    with path.open('rb') as stream:
+        return hashlib.file_digest(stream, 'sha256').hexdigest()
+
+
+def manifest(directory, write=False):
+    directory = Path(directory)
+    files = {str(p.relative_to(directory)): digest(p) for p in sorted(directory.rglob('*'))
+             if p.is_file() and p.name != MANIFEST}
+    if not files or not any(n in files for n in ('model.json', 'rr.txt')):
+        raise ValueError(f'{directory}: missing model metadata')
+    if write:
+        (directory / MANIFEST).write_text(json.dumps(files, indent=2, sort_keys=True) + '\n')
+    elif files != json.loads((directory / MANIFEST).read_text()):
+        raise ValueError(f'{directory}: missing, changed or unexpected model files; rebuild the models')
+
+
+def capture(directory, name):
+    """A completed, finite frame and all files referenced by each network-layer capture."""
+    import numpy as np
+    directory = Path(directory)
+    output = directory / 'out.raw'
+    if output.stat().st_size != 1280 * 720 * 8:
+        raise ValueError('capture did not produce a complete RGBA16F frame')
+    pixels = np.fromfile(output, dtype=np.float16).reshape(-1, 4)
+    if not np.isfinite(pixels).all() or not np.any(pixels[:, :3]):
+        raise ValueError('capture produced a nonfinite or empty frame')
+    if name == 'k':
+        kernels = ['dltss_pwin_' + s + '_layer' for s in
+                   ('enc0', 'enc1', 'enc2', 'enc3', 'enc4', 'dec5', 'dec4', 'dec3', 'dec2', 'dec1', 'dec0')]
+    else:
+        kernels = ['rrlite_' + s + '_4x4' for s in ('enc1', 'enc2', 'enc3', 'enc4', 'enc5', 'dec4', 'dec3', 'dec2', 'dec1')]
+        kernels += ['rrlite_enc0_4x4_mvlo_hdr' + ('_folded' if name == 'm' else ''),
+                    'rrlite_dec0_4x4' + ('_folded' if name == 'm' else '')]
+    for kernel in kernels:
+        matches = sorted((directory / 'replay').glob('replay-*-' + kernel))
+        if not matches:
+            raise ValueError(f'missing captured layer {kernel}')
+        launch = matches[0]
+        rows = [line.split() for line in (launch / 'manifest.txt').read_text().splitlines()]
+        if ['kernel', kernel] not in rows:
+            raise ValueError(f'{kernel}: incorrect capture metadata')
+        args = next(int(r[1]) for r in rows if r[0] == 'args')
+        if args <= 0 or (launch / 'args.bin').stat().st_size != args:
+            raise ValueError(f'{kernel}: truncated arguments')
+        allocations = {r[1]: int(r[3]) for r in rows if r[0] == 'alloc'}
+        if not allocations:
+            raise ValueError(f'{kernel}: no weight allocations captured')
+        for index, size in allocations.items():
+            if size <= 0 or (launch / f'alloc-{index}.bin').stat().st_size != size:
+                raise ValueError(f'{kernel}: truncated allocation {index}')
+        for r in rows:
+            if r[0] == 'pointer' and (r[2] not in allocations or not 0 <= int(r[3]) < allocations[r[2]]):
+                raise ValueError(f'{kernel}: invalid captured pointer')
+
+
+def run(log, timeout, command):
+    if command and command[0] == '--':
+        command = command[1:]
+    if not command or not 0 < timeout <= 86400:
+        raise ValueError('command required; timeout must be between 0 and 86400 seconds')
+    log = Path(log)
+    log.parent.mkdir(parents=True, exist_ok=True)
+    started = time.monotonic()
+    print(f'  starting {log.stem} (deadline {timeout:g}s; log {log})', flush=True)
+    with log.open('wb') as output:
+        child = subprocess.Popen(command, stdout=output, stderr=subprocess.STDOUT, start_new_session=True)
+        def stop():
+            for sig in (signal.SIGTERM, signal.SIGKILL):
+                try:
+                    os.killpg(child.pid, sig)
+                except ProcessLookupError:
+                    break
+                try:
+                    child.wait(timeout=3)
+                except subprocess.TimeoutExpired:
+                    continue
+                # Kill remaining members too, even when the group leader has exited.
+                if sig == signal.SIGTERM:
+                    try:
+                        os.killpg(child.pid, signal.SIGKILL)
+                    except ProcessLookupError:
+                        pass
+                break
+        previous = {}
+        def interrupted(sig, frame):
+            raise KeyboardInterrupt
+        for sig in (signal.SIGTERM, signal.SIGINT):
+            previous[sig] = signal.signal(sig, interrupted)
+        result = None
+        try:
+            heartbeat = 0
+            while child.poll() is None:
+                elapsed = time.monotonic() - started
+                if elapsed >= timeout:
+                    stop()
+                    print(f'  {log.stem}: timed out after {elapsed:.0f}s', file=sys.stderr)
+                    result = 124
+                    break
+                if elapsed >= heartbeat + 15:
+                    print(f'  {log.stem}: still running ({elapsed:.0f}s, pid {child.pid}); log {log}', flush=True)
+                    heartbeat = elapsed
+                time.sleep(min(.25, max(.01, timeout - elapsed)))
+            if result is None:
+                result = child.returncode
+        except KeyboardInterrupt:
+            stop()
+            result = 130
+        finally:
+            for sig, handler in previous.items():
+                signal.signal(sig, handler)
+    if result:
+        print('\n'.join(log.read_text(errors='replace').splitlines()[-30:]), file=sys.stderr)
+    print(f'  {log.stem}: exit={result}, elapsed={time.monotonic() - started:.1f}s; log {log}', flush=True)
+    return result if result >= 0 else 128 - result
+
+
+def main():
+    parser = argparse.ArgumentParser(description=__doc__)
+    sub = parser.add_subparsers(dest='action', required=True)
+    for action in ('manifest', 'validate'):
+        action_parser = sub.add_parser(action)
+        action_parser.add_argument('directory', type=Path)
+        if action == 'validate':
+            action_parser.add_argument('models', nargs='*', choices=MODELS)
+    p = sub.add_parser('capture')
+    p.add_argument('directory', type=Path)
+    p.add_argument('name', choices=('k', 'm', 'l'))
+    p = sub.add_parser('run')
+    p.add_argument('--log', required=True)
+    p.add_argument('--timeout', type=float, default=900)
+    p.add_argument('command', nargs=argparse.REMAINDER)
+    a = parser.parse_args()
+    try:
+        if a.action == 'run':
+            return run(a.log, a.timeout, a.command)
+        if a.action == 'capture':
+            capture(a.directory, a.name)
+        elif a.action == 'manifest':
+            manifest(a.directory, write=True)
+        else:
+            for name in (a.models or MODELS):
+                manifest(a.directory / name)
+    except (OSError, ValueError, KeyError, StopIteration) as error:
+        print(f'model-tools: {error}', file=sys.stderr)
+        return 1
+    return 0
+
+
+if __name__ == '__main__':
+    sys.exit(main())
diff --git a/scripts/build_d4r_nvngx_shim.sh b/scripts/build_d4r_nvngx_shim.sh
index e8bbf6c..924774a 100755
--- a/scripts/build_d4r_nvngx_shim.sh
+++ b/scripts/build_d4r_nvngx_shim.sh
@@ -6,6 +6,7 @@ set -euo pipefail
 ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)"
 MINGW_CXX="${MINGW_CXX:-x86_64-w64-mingw32-g++}"
 CLANG_CL="${CLANG_CL:-clang-cl}"
+BUILD_ID="$(bash "$ROOT/scripts/runtime_build_id.sh")"
 
 VULKAN_INCLUDE="${VULKAN_INCLUDE:-/usr/include}"
 
@@ -22,7 +23,7 @@ PARAM_OBJECT="$ROOT/build/d4r_ngx_param_msvc.obj"
 "$CLANG_CL" --target=x86_64-pc-windows-msvc /nologo /std:c++20 /O2 /c /GS- /GR- /EHs-c- /Zl \
   "/Fo$PARAM_OBJECT" -- "$ROOT/tools/d4r_ngx_param_msvc.cpp"
 # The engine (D4R_ENGINE) reaches Vulkan through a function table: MinGW has no Vulkan import library.
-"$MINGW_CXX" -std=c++20 -O2 -Wall -Wextra -Wno-missing-field-initializers -DD4R_ENGINE_VK_TABLE -I"$ROOT/build/vulkan-include" -I"$ROOT/build/vulkan-staging" -shared -static -static-libgcc -static-libstdc++ \
+"$MINGW_CXX" -std=c++20 -O2 -Wall -Wextra -Wno-missing-field-initializers -DD4R_ENGINE_VK_TABLE "-DD4R_RUNTIME_BUILD_ID=\"$BUILD_ID\"" -I"$ROOT/build/vulkan-include" -I"$ROOT/build/vulkan-staging" -shared -static -static-libgcc -static-libstdc++ \
   "$ROOT/tools/d4r_nvngx_shim.cpp" "$ROOT/tools/d4r_ngx_param_host.cpp" "$PARAM_OBJECT" \
   "$ROOT/engine/runtime.cpp" "$ROOT/engine/game.cpp" "$ROOT/engine/game_m.cpp" "$ROOT/engine/game_rr.cpp" "$ROOT/engine/game_e.cpp" "$ROOT/engine/vk_dispatch.cpp" \
   -ldxgi -lbcrypt -Wl,--entry,d4r_dll_entry -o "$ROOT/build/d4r_nvngx.dll"
diff --git a/scripts/build_release_glibc241.sh b/scripts/build_release_glibc241.sh
index f2b7aad..277d100 100755
--- a/scripts/build_release_glibc241.sh
+++ b/scripts/build_release_glibc241.sh
@@ -76,8 +76,14 @@ export D4R_DLSS_PTX_DIR="$BUILD/ptx"
 python3 "$ROOT/kernels/tools/extract_dlss_ptx.py" "$D4R_BUNDLE_DLSS" "$D4R_DLSS_PTX_DIR" > "$BUILD/logs/extract.log"
 build_textures() {
     local arch="$1" folder fp8 accuracy out
-    # RDNA2 has no texture tails (they need WMMA); package_release.sh builds its K layers
-    [[ "$arch" == gfx103* ]] && return 0
+    # RDNA2: only K's natively translated hiluma output/input kernels (no WMMA), fast set only (the accuracy
+    # set keeps ZLUDA's compile of them); package_release.sh builds the K layers and copies these
+    if [[ "$arch" == gfx103* ]]; then
+        D4R_PREFER_ACCURACY=0 D4R_GPU_ARCH="$arch" D4R_DLSS_DLL="$D4R_BUNDLE_DLSS" \
+            "$ROOT/kernels/build.sh" tex "$D4R_BUNDLE_TEX/$arch" > "$BUILD/logs/tex-$arch-0.log" 2>&1
+        echo "Texture kernels rebuilt: $arch (native hiluma output/input only)"
+        return 0
+    fi
     folders=("$arch")
     [[ "$arch" == gfx12* ]] && folders+=("$arch-fp8")
     for folder in "${folders[@]}"; do
diff --git a/scripts/build_wine_nvcuda_bridge.sh b/scripts/build_wine_nvcuda_bridge.sh
index ebc9bdc..6e88a36 100755
--- a/scripts/build_wine_nvcuda_bridge.sh
+++ b/scripts/build_wine_nvcuda_bridge.sh
@@ -5,6 +5,7 @@ ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)"
 WINEGCC="${WINEGCC:-winegcc}"
 WINEBUILD="${WINEBUILD:-winebuild}"
 BRIDGE_ROOT="$ROOT/build/wine-nvcuda"
+BUILD_ID="$(bash "$ROOT/scripts/runtime_build_id.sh")"
 
 mkdir -p "$BRIDGE_ROOT/x86_64-unix" "$BRIDGE_ROOT/x86_64-windows"
 MOTION_FLAGS=()
@@ -13,6 +14,7 @@ if [[ -x "${D4R_ROCM_DIR:-/opt/rocm}/lib/llvm/bin/clang++" ]]; then
   MOTION_FLAGS=(-DD4R_MOTION_KERNELS -I"$BRIDGE_ROOT/generated")
 fi
 "$WINEGCC" -m64 -shared -o "$BRIDGE_ROOT/x86_64-unix/nvcuda.dll" \
+  "-DD4R_RUNTIME_BUILD_ID=\"$BUILD_ID\"" \
   "${MOTION_FLAGS[@]}" \
   "$ROOT/tools/wine_nvcuda_bridge.c" \
   -Wb,--export="$ROOT/tools/wine_nvcuda_bridge.spec" -ldl -lpthread
diff --git a/scripts/d4r_config.py b/scripts/d4r_config.py
index 08b927e..95cb2de 100755
--- a/scripts/d4r_config.py
+++ b/scripts/d4r_config.py
@@ -206,6 +206,12 @@ def main(argv):
         env["D4R_SHIM_MARKER_POLL_US"] = value
 
     value = flag("Engine", "Enabled")
+    if value is None:
+        models = Path(os.environ.get("D4R_RUNTIME_DIR", path.parent)) / "engine"
+        value = any((models / model).is_dir() for model in
+                    ("k", "k-ldr", "m", "m-ldr", "l", "l-ldr", "rr-d", "rr-e", "rr-f",
+                     "rr-d-gfx11", "rr-e-gfx11"))
+        env["D4R_ENGINE_K"] = "0"
     env["D4R_ENGINE"] = "1" if value else "0"
     if path_value("Engine", "ModelDir"):
         env["D4R_ENGINE_MODEL_DIR"] = path_value("Engine", "ModelDir")
diff --git a/scripts/package_release.sh b/scripts/package_release.sh
index f006618..e70c702 100755
--- a/scripts/package_release.sh
+++ b/scripts/package_release.sh
@@ -16,7 +16,8 @@
 #   D4R_ROCM_RUNTIME ROCm runtime bundled as d4r/rocm (scripts/fetch_rocm_runtime.sh)
 #                                                                           default ~/.cache/d4r-rocm-runtime
 #   D4R_MAX_GLIBC     highest GLIBC ABI the Linux libraries may require     default 2.41
-#   D4R_GPU_ARCHS    GPU targets to build kernels for          default gfx1030-gfx1036 (K only), gfx1100-gfx1103, gfx1200, gfx1201
+#   D4R_GPU_ARCHS    GPU targets to build kernels for          default gfx1030-gfx1036 (K and its hiluma kernels only),
+#                    gfx1100-gfx1103, gfx1200, gfx1201
 #                    (RDNA4 targets also get a <target>-fp8 folder: the kernels for d4r.ini NativeFp8 = on)
 #   D4R_OPTISCALER_LICENSE  OptiScaler's LICENSE (GPL-3.0) text; default: system SPDX or common-licenses copy
 #   D4R_VKD3D_SRC    vkd3d-proton source checkout, for its license files     default ~/.cache/d4r-vkd3d-proton
@@ -216,15 +217,30 @@ ensure_rr_textures() {
   fi
 }
 build_target() {
-  local folder="$1" arch fp8 tex_dir accurate accurate_tex f
+  local folder="$1" arch fp8 tex_dir tmp_tex accurate accurate_tex f
   arch="${folder%-fp8}"
   fp8=0
   [[ "$folder" == *-fp8 ]] && fp8=1
-  # RDNA2 (gfx103x): only DLSS 4's layers (preset K) have native kernels there, built with software matrix
-  # products (kernels/common/wmma_layout.h); M, L, the texture tails and Ray Reconstruction stay translated by ZLUDA
+  # RDNA2 (gfx103x): DLSS 4's layers (preset K), built with software matrix products
+  # (kernels/common/wmma_layout.h), and K's hiluma output/input kernels translated from the PTX
+  # (kernels/native, no WMMA; wave32, byte-identical to ZLUDA's compile). M, L, the texture tails and Ray
+  # Reconstruction call WMMA and stay translated by ZLUDA. The accuracy set keeps ZLUDA's compile of hiluma.
   if [[ "$arch" == gfx103* ]]; then
     D4R_PREFER_ACCURACY=0 D4R_ROCM_DIR="$ROCM" D4R_GPU_ARCH="$arch" "$ROOT/kernels/build.sh" k "$STAGE/d4r/kernels/$folder" >/dev/null
     D4R_PREFER_ACCURACY=1 D4R_ROCM_DIR="$ROCM" D4R_GPU_ARCH="$arch" "$ROOT/kernels/build.sh" k "$STAGE/d4r/kernels/accuracy/$folder" >/dev/null
+    if [[ "$VARIANT" != clean ]]; then
+      tex_dir="$D4R_BUNDLE_TEX/$folder" tmp_tex=""
+      if [[ ! -d "$tex_dir" ]]; then
+        tex_dir="$(mktemp -d)" tmp_tex="$tex_dir"
+        D4R_PREFER_ACCURACY=0 D4R_ROCM_DIR="$ROCM" D4R_GPU_ARCH="$arch" D4R_DLSS_DLL="$D4R_BUNDLE_DLSS" \
+          "$ROOT/kernels/build.sh" tex "$tex_dir" >/dev/null
+      fi
+      # only the kernels validated for RDNA2, whatever else an older bundle folder holds
+      for f in "$tex_dir"/hiluma_engine_output_*_hdr_max_v2_rel.hsaco "$tex_dir"/hiluma_engine_input_*_v2_rel.hsaco; do
+        [[ -f "$f" ]] && cp "$f" "$STAGE/d4r/kernels/$folder/"
+      done
+      [[ -z "$tmp_tex" ]] || rm -rf "$tmp_tex"
+    fi
     rm -f "$STAGE/d4r/kernels/$folder"/*.resolution.txt "$STAGE/d4r/kernels/accuracy/$folder"/*.resolution.txt
     python3 "$ROOT/kernels/tools/kernel_manifest.py" "$STAGE/d4r/kernels/$folder" "${DLLS[@]}" "${RR_DLLS[@]}"
     python3 "$ROOT/kernels/tools/kernel_manifest.py" "$STAGE/d4r/kernels/accuracy/$folder" "${DLLS[@]}" "${RR_DLLS[@]}"
diff --git a/scripts/runtime_build_id.sh b/scripts/runtime_build_id.sh
new file mode 100644
index 0000000..751f9b3
--- /dev/null
+++ b/scripts/runtime_build_id.sh
@@ -0,0 +1,7 @@
+#!/usr/bin/env bash
+set -euo pipefail
+cd -- "$(dirname -- "${BASH_SOURCE[0]}")/.."
+revision=$(git rev-parse --short HEAD 2>/dev/null || echo unknown)
+sources=$(sha256sum tools/*.h tools/d4r_nvngx_shim.cpp tools/d4r_vulkan_frontend.inl tools/wine_nvcuda_bridge.c \
+  engine/*.h engine/*.cpp engine/*.inc | sha256sum | cut -c1-16)
+printf '%s-%s\n' "$revision" "$sources"
diff --git a/scripts/stage_engine_kit.sh b/scripts/stage_engine_kit.sh
index 8f52579..fc50973 100755
--- a/scripts/stage_engine_kit.sh
+++ b/scripts/stage_engine_kit.sh
@@ -49,5 +49,6 @@ cp "$GL/bin/glslangValidator" "$KIT/bin/glslangValidator"
 chmod +x "$KIT/bin/glslangValidator"
 cp "$D4R_GLSLANG_LICENSE" "$D4R/licenses/glslang-LICENSE.txt"
 cp "$ROOT/packaging/build_engine_models.sh" "$D4R/tools/build-engine-models.sh"
+cp "$ROOT/packaging/model_tools.py" "$KIT/model_tools.py"
 chmod +x "$D4R/tools/build-engine-models.sh"
 printf 'Staged the engine-model builder in %s (%s)\n' "$D4R/tools" "$(du -sh "$D4R/tools" | cut -f1)"
diff --git a/tests/test_gpu_compatibility.py b/tests/test_gpu_compatibility.py
new file mode 100644
index 0000000..910c7d8
--- /dev/null
+++ b/tests/test_gpu_compatibility.py
@@ -0,0 +1,149 @@
+"""GPU identity and model capability decisions without a GPU or Wine process."""
+from pathlib import Path
+import subprocess
+import tempfile
+import unittest
+
+ROOT = Path(__file__).resolve().parents[1]
+
+
+class GpuCompatibilityTests(unittest.TestCase):
+    def compile_run(self, source, cxx=False):
+        with tempfile.TemporaryDirectory() as t:
+            executable = Path(t) / 'check'
+            result = subprocess.run(['c++' if cxx else 'cc', '-x', 'c++' if cxx else 'c',
+                            '-std=c++17' if cxx else '-std=gnu11', '-Wall', '-Wextra', '-Werror',
+                            '-pthread',
+                            '-I', str(ROOT / 'tools'), '-o', str(executable), '-'], input=source,
+                           capture_output=True, text=True)
+            self.assertEqual(result.returncode, 0, result.stderr)
+            subprocess.run([str(executable)], check=True, timeout=5)
+
+    def test_pci_identity_is_not_a_vendor_id_or_gpu_size_guess(self):
+        self.compile_run(r'''
+#include "d4r_gpu_identity.h"
+#include <assert.h>
+int main(void) {
+    assert(d4r_same_pci("0000:03:00.0", "00000000:03:00.0"));
+    assert(!d4r_same_pci("0000:03:00.0", "0000:04:00.0"));
+    assert(!d4r_same_pci("0000:03:00.0", "0001:03:00.0"));
+    assert(!d4r_same_pci("0000:03:00.0garbage", "0000:03:00.0"));
+    assert(d4r_pci_matches_kfd("0000:03:00.0", 0, 0x300));
+    assert(!d4r_pci_matches_kfd("0000:03:00.0", 0, 0x400));
+    assert(!d4r_pci_matches_kfd("0000:03:00.0", 1, 0x300));
+}
+''')
+
+    def test_bridge_maps_nonzero_compute_ordinal_and_only_its_luid(self):
+        source = (ROOT / 'tools/wine_nvcuda_bridge.c').read_text()
+        selection = source[source.index('static CUresult selected_cuda_device('):
+                           source.index('static void create_default_zluda_context(')]
+        luid = source[source.index('CUresult WINAPI cuDeviceGetLuid('):
+                      source.index('CUresult WINAPI cuDeviceGetAttribute(')]
+        self.compile_run(r'''
+#include "d4r_gpu_identity.h"
+#include <assert.h>
+#include <stdlib.h>
+#include <string.h>
+#define WINAPI
+typedef int CUresult;
+typedef int CUdevice;
+typedef int (*CUDEVICEGETCOUNT_FN)(int*);
+typedef int (*CUDEVICEGET_FN)(int*, int);
+typedef int (*CUDEVICEGETLUID_FN)(char*, unsigned*, int);
+enum { CUDA_SUCCESS=0, CUDA_ERROR_NOT_INITIALIZED=3, CUDA_ERROR_NOT_SUPPORTED=801 };
+static const char* render = "0000:03:00.0";
+static int devices = 2;
+static unsigned GetEnvironmentVariableA(const char* key, char* out, unsigned size) {
+    (void)key; snprintf(out, size, "%s", render); return (unsigned)strlen(render);
+}
+static void tracef(const char* fmt, ...) { (void)fmt; }
+static int count(int* n) { *n=devices; return 0; }
+static int get(int* d,int n) { *d=100+n; return 0; }
+static int pci(char* out,int size,int device) {
+    snprintf(out, (size_t)size, "%s", device==101 ? "00000000:03:00.0" : "0000:04:00.0"); return 0;
+}
+static int original_luid(char* out,unsigned* mask,int device) { (void)device; memset(out,0,8);*mask=0;return 0; }
+static void* find_zluda_symbol(const char* name) {
+    if (!strcmp(name,"cuDeviceGetCount")) return (void*)count;
+    if (!strcmp(name,"cuDeviceGet")) return (void*)get;
+    if (!strcmp(name,"cuDeviceGetPCIBusId")) return (void*)pci;
+    return (void*)original_luid;
+}
+static int missing(const char* name) { (void)name; return 999; }
+static unsigned get_process_u32(const char* name,unsigned fallback) {
+    (void)fallback;
+    return strstr(name,"LOW") ? 0x3f2 : strstr(name,"HIGH") ? 0 : 1;
+}
+''' + selection + luid + r'''
+int main(void) {
+    assert(d4r_pci_matches_kfd("0000:03:00.0",0,0x300));
+    int selected=-1; assert(selected_cuda_device(&selected)==0 && selected==101);
+    char id[8]; unsigned mask=9;
+    assert(cuDeviceGetLuid(id,&mask,100)==0 && mask==0);
+    assert(cuDeviceGetLuid(id,&mask,101)==0 && mask==1);
+    unsigned low=0;memcpy(&low,id,4);assert(low==0x3f2);
+    render="0000:05:00.0";assert(selected_cuda_device(&selected)==801);
+    render="";assert(selected_cuda_device(&selected)==801);
+    devices=1;assert(selected_cuda_device(&selected)==0 && selected==100);
+}
+''')
+
+    def test_rr_model_cannot_reach_unsupported_fp8_pipelines(self):
+        self.compile_run(r'''
+#include "d4r_engine_compat.h"
+#include <cassert>
+int main() {
+    const char* fp8 = "D4RR0001\npreset 4\nstage enc0 x.spv\n";
+    const char* wide = "D4RR0001\npreset 4\narithmetic gfx11\nstage enc0 x.spv\n";
+    assert(!d4r_rr_model_compatible(fp8, true, false));
+    assert(!d4r_rr_model_compatible(wide, false, false));
+    assert(d4r_rr_model_compatible(wide, true, false));
+    assert(d4r_rr_model_compatible(fp8, true, true));
+    assert(!d4r_rr_model_compatible("D4RR0002\npreset 6\n", true, false));
+    assert(!d4r_rr_model_compatible("D4RR0001\narithmetic unknown\n", true, true));
+    VkCooperativeMatrixPropertiesKHR p{};
+    p.MSize=p.NSize=p.KSize=16; p.scope=VK_SCOPE_SUBGROUP_KHR;
+    p.AType=p.BType=VK_COMPONENT_TYPE_FLOAT16_KHR;
+    p.CType=p.ResultType=VK_COMPONENT_TYPE_FLOAT32_KHR;
+    assert(d4r_matrix_tuple(p, VK_COMPONENT_TYPE_FLOAT16_KHR));
+    p.KSize=32; assert(!d4r_matrix_tuple(p, VK_COMPONENT_TYPE_FLOAT16_KHR));
+}
+''', cxx=True)
+
+    def test_existing_context_is_preserved_and_bootstrap_binds_each_thread(self):
+        text = (ROOT / 'tools/wine_nvcuda_bridge.c').read_text()
+        ensure = text[text.index('static void ensure_context(void)'):text.index('/* D4R_NGX_STREAM_PRIORITY')]
+        self.compile_run(r'''
+#include <assert.h>
+#include <pthread.h>
+#include <stdint.h>
+#include <string.h>
+typedef int CUresult;
+typedef void* CUcontext;
+typedef CUresult (*CUCTX_CURRENT_FN)(CUcontext);
+enum { CUDA_SUCCESS=0 };
+static pthread_once_t context_once=PTHREAD_ONCE_INIT;
+static int context_bootstrap_result=3, bootstrap_calls=0;
+static _Thread_local int context_setup_result=3;
+static _Thread_local void* active;
+static void* cuda_context;
+static int current(void** out) { *out=active;return 0; }
+static int set(void* context) { active=context;return 0; }
+static void* find_zluda_symbol(const char* key) {
+    return !strcmp(key,"cuCtxGetCurrent") ? (void*)current : (void*)set;
+}
+static void create_default_zluda_context(void) {
+    ++bootstrap_calls;cuda_context=(void*)(uintptr_t)0x456;context_bootstrap_result=0;
+}
+''' + ensure + r'''
+static void* worker(void* ignored) {
+    (void)ignored;ensure_context();assert(active==cuda_context && context_setup_result==0);return NULL;
+}
+int main(void) {
+    active=(void*)(uintptr_t)0x123;ensure_context();assert(active==(void*)(uintptr_t)0x123 && bootstrap_calls==0);
+    pthread_t a;assert(!pthread_create(&a,NULL,worker,NULL));pthread_join(a,NULL);
+    assert(bootstrap_calls==1);ensure_context();assert(active==(void*)(uintptr_t)0x123);
+    assert(!pthread_create(&a,NULL,worker,NULL));pthread_join(a,NULL);assert(bootstrap_calls==1);
+}
+''')
diff --git a/tests/test_model_builder.py b/tests/test_model_builder.py
new file mode 100644
index 0000000..a4cd054
--- /dev/null
+++ b/tests/test_model_builder.py
@@ -0,0 +1,97 @@
+"""Model builds must time out, retain failures, and reject incomplete/corrupt artifacts."""
+import contextlib
+import importlib.util
+import io
+import json
+import os
+from pathlib import Path
+import subprocess
+import sys
+import tempfile
+import unittest
+
+ROOT = Path(__file__).resolve().parents[1]
+spec = importlib.util.spec_from_file_location('model_tools', ROOT / 'packaging/model_tools.py')
+model_tools = importlib.util.module_from_spec(spec)
+spec.loader.exec_module(model_tools)
+
+
+class ModelBuildTests(unittest.TestCase):
+    def test_installer_key_tracks_model_arithmetic_options(self):
+        text = (ROOT / 'install.sh').read_text()
+        key = text[text.index('key() {'):text.index('have() {')]
+        with tempfile.TemporaryDirectory() as t:
+            root = Path(t)
+            (root / 'tools').mkdir()
+            (root / 'tools/compiler.py').write_text('compiler fixture')
+            for name in ('nvngx_dlss.dll', 'nvngx_dlssd.dll'): (root / name).write_bytes(b'DLL fixture')
+            env = {k: v for k, v in os.environ.items() if not k.startswith(('D4R_RR_', 'K_', 'M_'))}
+            command = ['bash', '-c', 'RUNTIME=$1\n' + key + '\nkey 0', '_', t]
+            baseline = subprocess.check_output(command, env=env)
+            for name in ('K_DEFS_ENC0', 'M_EXACT', 'D4R_RR_MIX'):
+                self.assertNotEqual(baseline, subprocess.check_output(command, env=dict(env, **{name: '1'})))
+
+    def test_child_exit_status_and_timeout_are_preserved(self):
+        with tempfile.TemporaryDirectory() as t, contextlib.redirect_stdout(io.StringIO()), contextlib.redirect_stderr(io.StringIO()):
+            log = Path(t) / 'capture.log'
+            self.assertEqual(model_tools.run(log, 5, [sys.executable, '-c', 'print("failure detail");raise SystemExit(37)']), 37)
+            self.assertIn('failure detail', log.read_text())
+            self.assertEqual(model_tools.run(log, .15, [sys.executable, '-c', 'import time;time.sleep(60)']), 124)
+            self.assertTrue(log.exists())
+
+    def test_manifest_rejects_missing_modified_and_extra_files(self):
+        with tempfile.TemporaryDirectory() as t:
+            directory = Path(t)
+            (directory / 'model.json').write_text('{}')
+            weights = directory / 'weights.bin'
+            weights.write_bytes(b'weights')
+            model_tools.manifest(directory, write=True)
+            model_tools.manifest(directory)
+            weights.write_bytes(b'corrupt')
+            with self.assertRaises(ValueError): model_tools.manifest(directory)
+            weights.unlink()
+            with self.assertRaises(ValueError): model_tools.manifest(directory)
+            weights.write_bytes(b'weights')
+            (directory / 'unexpected.spv').write_bytes(b'extra')
+            with self.assertRaises(ValueError): model_tools.manifest(directory)
+
+    def test_capture_requires_completed_frame_and_all_layers(self):
+        import numpy as np
+        with tempfile.TemporaryDirectory() as t:
+            directory = Path(t)
+            pixels = np.ones((720, 1280, 4), dtype=np.float16)
+            pixels.tofile(directory / 'out.raw')
+            replay = directory / 'replay'
+            replay.mkdir()
+            for layer in ('enc0', 'enc1', 'enc2', 'enc3', 'enc4', 'dec5', 'dec4', 'dec3', 'dec2', 'dec1', 'dec0'):
+                kernel = 'dltss_pwin_' + layer + '_layer'
+                d = replay / ('replay-000001-' + kernel)
+                d.mkdir()
+                (d / 'manifest.txt').write_text(f'kernel {kernel}\nargs 8\nalloc 0 0x1000 4\npointer 0 0 0\n')
+                (d / 'args.bin').write_bytes(bytes(8))
+                (d / 'alloc-0.bin').write_bytes(bytes(4))
+            model_tools.capture(directory, 'k')
+            (d / 'alloc-0.bin').write_bytes(b'')
+            with self.assertRaises(ValueError): model_tools.capture(directory, 'k')
+            (d / 'alloc-0.bin').write_bytes(bytes(4))
+            pixels[0, 0, 0] = np.nan
+            pixels.tofile(directory / 'out.raw')
+            with self.assertRaises(ValueError): model_tools.capture(directory, 'k')
+
+    def test_dump_does_not_accept_a_failed_child_with_a_replay_marker(self):
+        text = (ROOT / 'packaging/build_engine_models.sh').read_text()
+        start = text.index('  dump() {')
+        function = text[start:text.index('\n  }\n', start) + len('\n  }\n')]
+        with tempfile.TemporaryDirectory() as t:
+            root = Path(t)
+            (root / 'proton').mkdir(); (root / 'app').mkdir()
+            proton = root / 'proton/proton'
+            proton.write_text('#!/bin/sh\nmkdir -p "$D4R_CUDA_REPLAY_DUMP_DIR/replay-000001-dltss_pwin_dec0_layer"\nexit 37\n')
+            proton.chmod(0o755)
+            prefix = 'set -euo pipefail\nWORK=$1; APP=$WORK/app; PROTON=$WORK/proton; STEAM_ROOT=$WORK; RUNTIME=$WORK\n'
+            prefix += 'MODEL_TOOLS=$2; winpath() { printf "%s" "$1"; }; fail() { echo "$*" >&2; exit 1; }\n'
+            result = subprocess.run(['bash', '-c', prefix + function + '\ndump k 11 ignored\n', '_', t,
+                                     str(ROOT / 'packaging/model_tools.py')], capture_output=True, text=True, timeout=10)
+            self.assertNotEqual(result.returncode, 0)
+            self.assertIn('capture process failed', result.stderr)
+            self.assertTrue((root / 'k/run.log').exists())
diff --git a/tests/test_rdna2_native_output.py b/tests/test_rdna2_native_output.py
new file mode 100644
index 0000000..e8633ed
--- /dev/null
+++ b/tests/test_rdna2_native_output.py
@@ -0,0 +1,193 @@
+"""RDNA2 (gfx103x) native hiluma kernels: build gating, packaging allowlist and the f16 rounding contract.
+
+Run: python3 -m unittest tests.test_rdna2_native_output -v
+The compile checks need ROCm's clang (D4R_ROCM_DIR, default /opt/rocm) and are skipped without it. The full
+texture build additionally needs the extracted DLSS 310.7 PTX (kernels/extracted/ptx) and
+D4R_RUN_KERNEL_BUILD_TESTS=1, since it compiles all twelve kernels.
+Arithmetic equality with ZLUDA's compile is checked in the harness (docs/native-kernels.md), not here.
+"""
+import os
+from pathlib import Path
+import re
+import shutil
+import subprocess
+import tempfile
+import unittest
+
+ROOT = Path(__file__).resolve().parents[1]
+ROCM = Path(os.environ.get("D4R_ROCM_DIR", "/opt/rocm"))
+CLANG = ROCM / "lib/llvm/bin/clang++"
+PTX_DIR = ROOT / "kernels/extracted/ptx"
+TOOLS = [shutil.which(t) for t in ("llvm-readelf", "llvm-objdump")]
+OUTPUTS = {f"hiluma_engine_output_{d}_{mv}_hdr_max_v2_rel" for d in ("depthinv", "depthreg") for mv in ("mvhi", "mvlo")}
+INPUTS = {f"hiluma_engine_input_{d}_{mv}_{r}_v2_rel"
+          for d in ("depthinv", "depthreg") for mv in ("mvhi", "mvlo") for r in ("hdr", "ldr")}
+# cvt.rn.f16.f32 of an fma (and the packed form of a product) of widened f16 values: LLVM would fold the
+# conversion into v_fma_mixlo/mixhi_f16, which rounds once instead of twice
+TINY_PTX = """.version 8.0
+.target sm_80
+.address_size 64
+.visible .entry tiny(
+	.param .align 8 .b8 tiny_param_0[16]
+)
+{
+	.reg .b16 	%rs<6>;
+	.reg .f32 	%f<8>;
+	.reg .b32 	%r<3>;
+	.reg .b64 	%rd<3>;
+	ld.param.u64 	%rd1, [tiny_param_0];
+	cvta.to.global.u64 	%rd2, %rd1;
+	ld.global.u16 	%rs2, [%rd2+8];
+	ld.global.u16 	%rs3, [%rd2+10];
+	ld.global.u16 	%rs4, [%rd2+12];
+	cvt.f32.f16 	%f3, %rs2;
+	cvt.f32.f16 	%f4, %rs3;
+	cvt.f32.f16 	%f5, %rs4;
+	fma.rn.f32 	%f2, %f3, %f4, %f5;
+	cvt.rn.f16.f32 	%rs1, %f2;
+	mul.rn.f32 	%f6, %f3, %f4;
+	cvt.rn.f16x2.f32 	%r1, %f6, %f2;
+	st.global.u16 	[%rd2], %rs1;
+	st.global.u32 	[%rd2+4], %r1;
+	ret;
+}
+"""
+
+
+def run(*args, **kwargs):
+    return subprocess.run([str(a) for a in args], text=True, capture_output=True, **kwargs)
+
+
+class Rdna2BuildGating(unittest.TestCase):
+    def setUp(self):
+        self.temp = tempfile.TemporaryDirectory()
+        self.addCleanup(self.temp.cleanup)
+        self.root = Path(self.temp.name)
+        # a compiler that fails if anything reaches it: the refusals must come first
+        compiler = self.root / "rocm/lib/llvm/bin/clang++"
+        compiler.parent.mkdir(parents=True)
+        compiler.write_text("#!/bin/sh\nexit 99\n")
+        compiler.chmod(0o755)
+
+    def test_wmma_families_are_refused_on_every_rdna2_target(self):
+        for arch in ("gfx1030", "gfx1032", "gfx1036"):
+            for family in ("m", "l", "rr", "rrswin"):
+                env = dict(os.environ, D4R_ROCM_DIR=str(self.root / "rocm"), D4R_GPU_ARCH=arch)
+                result = run("bash", ROOT / "kernels/build.sh", family, self.root / "out", env=env)
+                self.assertEqual(result.returncode, 2, (arch, family, result.stderr))
+                self.assertIn("need WMMA", result.stderr)
+
+    def test_rdna2_texture_build_does_not_need_zluda(self):
+        env = dict(os.environ, D4R_ROCM_DIR=str(self.root / "rocm"), D4R_GPU_ARCH="gfx1032",
+                   D4R_DLSS_DLL=str(self.root / "missing.dll"), D4R_DLSS_PTX_DIR=str(self.root / "no-ptx"))
+        env.pop("D4R_ZLUDA_EMIT", None)
+        result = run("bash", ROOT / "kernels/build.sh", "tex", self.root / "out", env=env)
+        self.assertNotIn("D4R_ZLUDA_EMIT", result.stderr)
+        self.assertIn("missing extracted PTX directory", result.stderr)
+
+    def test_rdna3_texture_build_still_needs_zluda(self):
+        env = dict(os.environ, D4R_ROCM_DIR=str(self.root / "rocm"), D4R_GPU_ARCH="gfx1101",
+                   D4R_DLSS_DLL=str(self.root / "missing.dll"))
+        env.pop("D4R_ZLUDA_EMIT", None)
+        result = run("bash", ROOT / "kernels/build.sh", "tex", self.root / "out", env=env)
+        self.assertNotEqual(result.returncode, 0)
+        self.assertIn("D4R_ZLUDA_EMIT", result.stderr)
+
+
+class Rdna2Packaging(unittest.TestCase):
+    def rdna2_branch(self, path):
+        text = (ROOT / path).read_text()
+        start = text.index('if [[ "$arch" == gfx103* ]]; then')
+        return text[start:text.index("\n  fi\n", start)]
+
+    def test_package_copies_only_the_validated_hiluma_kernels(self):
+        branch = self.rdna2_branch("scripts/package_release.sh")
+        copies = re.findall(r'"\$tex_dir"/([^ ;]+\.hsaco)', branch)
+        self.assertEqual(sorted(copies), ["hiluma_engine_input_*_v2_rel.hsaco",
+                                          "hiluma_engine_output_*_hdr_max_v2_rel.hsaco"])
+        self.assertIn('"$VARIANT" != clean', branch)  # the no-NVIDIA zip ships no PTX-derived code
+        self.assertNotIn("build.sh\" m", branch)
+        # the accuracy set gets K's layers only (ZLUDA keeps compiling hiluma there)
+        self.assertNotRegex(branch, r'PREFER_ACCURACY=1[^\n]*\n?[^\n]*build.sh" tex')
+        self.assertIn("kernel_manifest.py", branch)
+
+    def test_release_builder_builds_rdna2_fast_textures_only(self):
+        text = (ROOT / "scripts/build_release_glibc241.sh").read_text()
+        start = text.index('if [[ "$arch" == gfx103* ]]; then')
+        branch = text[start:text.index("fi\n", start)]
+        self.assertIn("D4R_PREFER_ACCURACY=0", branch)
+        self.assertNotIn("D4R_PREFER_ACCURACY=1", branch)
+        self.assertIn('kernels/build.sh" tex', branch)
+        self.assertIn("return 0", branch)
+
+
+class F16RoundingContract(unittest.TestCase):
+    def translate(self, directory):
+        ptx = directory / "tiny.ptx"
+        ptx.write_text(TINY_PTX)
+        result = run("python3", ROOT / "kernels/native/ptx2hip.py", ptx, "tiny", check=True)
+        return result.stdout
+
+    def test_translator_routes_f16_conversions_through_f2h_rn(self):
+        with tempfile.TemporaryDirectory() as temp:
+            source = self.translate(Path(temp))
+        self.assertEqual(source.count("f2h_rn("), 3)
+        self.assertNotRegex(source, r"\(h1\)f\d")
+
+    @unittest.skipUnless(CLANG.is_file(), "ROCm clang not found (set D4R_ROCM_DIR)")
+    def test_no_single_rounding_mix_instructions(self):
+        with tempfile.TemporaryDirectory() as temp:
+            temp = Path(temp)
+            (temp / "tiny.hip").write_text(self.translate(temp))
+            for arch, wave in (("gfx1032", "32"), ("gfx1101", "64")):
+                env = dict(os.environ, D4R_ROCM_DIR=str(ROCM), D4R_GPU_ARCH=arch, WAVE=wave)
+                run("bash", ROOT / "kernels/native/build_native.sh", temp / "tiny.hip", temp / f"{arch}.hsaco",
+                    env=env, check=True)
+                asm = (temp / f"{arch}.s").read_text()
+                self.assertNotRegex(asm, r"v_fma_mix(lo|hi)_f16", arch)
+            # canary: a plain cast in the same kernel is fused, so the check above can fail
+            (temp / "cast.hip").write_text((temp / "tiny.hip").read_text().replace("f2h_rn(", "(h1)("))
+            env = dict(os.environ, D4R_ROCM_DIR=str(ROCM), D4R_GPU_ARCH="gfx1032", WAVE="32")
+            run("bash", ROOT / "kernels/native/build_native.sh", temp / "cast.hip", temp / "cast.hsaco",
+                env=env, check=True)
+            self.assertRegex((temp / "cast.s").read_text(), r"v_fma_mix(lo|hi)_f16")
+
+
+@unittest.skipUnless(os.environ.get("D4R_RUN_KERNEL_BUILD_TESTS") == "1", "set D4R_RUN_KERNEL_BUILD_TESTS=1")
+@unittest.skipUnless(CLANG.is_file() and PTX_DIR.is_dir() and all(TOOLS), "needs ROCm clang, LLVM tools and PTX")
+class Rdna2TextureBuild(unittest.TestCase):
+    @classmethod
+    def setUpClass(cls):
+        cls.temp = tempfile.TemporaryDirectory()
+        cls.out = Path(cls.temp.name) / "gfx1032"
+        env = dict(os.environ, D4R_ROCM_DIR=str(ROCM), D4R_GPU_ARCH="gfx1032", D4R_DLSS_DLL="unused",
+                   D4R_DLSS_PTX_DIR=str(PTX_DIR))
+        run("bash", ROOT / "kernels/build.sh", "tex", cls.out, env=env, check=True)
+
+    @classmethod
+    def tearDownClass(cls):
+        cls.temp.cleanup()
+
+    def notes(self, name):
+        return run(TOOLS[0], "--notes", "-h", self.out / f"{name}.hsaco", check=True).stdout
+
+    def test_exactly_the_validated_kernels(self):
+        self.assertEqual({p.stem for p in self.out.glob("*.hsaco")}, OUTPUTS | INPUTS)
+
+    def test_code_objects_are_wave32_gfx1032_without_scratch_or_wmma(self):
+        for name in OUTPUTS | INPUTS:
+            notes = self.notes(name)
+            self.assertIn("gfx1032", notes, name)
+            self.assertRegex(notes, r"\.wavefront_size:\s+32", name)
+            self.assertRegex(notes, r"\.private_segment_fixed_size:\s+0\b", name)
+            code = run(TOOLS[1], "-d", "--mcpu=gfx1032", self.out / f"{name}.hsaco", check=True).stdout
+            self.assertNotRegex(code, r"v_wmma|v_fma_mix(lo|hi)_f16|v_dot\d", name)
+
+    def test_output_kernels_use_wgp_mode(self):
+        for name in OUTPUTS:
+            asm = (self.out / f"{name}.s").read_text()
+            self.assertIn(".amdhsa_workgroup_processor_mode 1", asm, name)
+
+
+if __name__ == "__main__":
+    unittest.main()
diff --git a/tests/test_shim_config.py b/tests/test_shim_config.py
index be1644d..126efb2 100644
--- a/tests/test_shim_config.py
+++ b/tests/test_shim_config.py
@@ -60,6 +60,20 @@ class ShimConfigTests(unittest.TestCase):
             result = subprocess.run(cmd, env=env, text=True, capture_output=True)
             self.assertEqual(result.returncode, 2)
 
+    def test_auto_engine_detects_models_without_overriding_environment(self):
+        with tempfile.TemporaryDirectory() as t:
+            root = Path(t)
+            ini = root / 'd4r.ini'
+            ini.write_text('[Engine]\nEnabled = auto\n')
+            (root / 'engine/rr-d-gfx11').mkdir(parents=True)
+            env = {k: v for k, v in os.environ.items() if not k.startswith('D4R_')}
+            cmd = ['python3', str(ROOT / 'scripts/d4r_config.py'), '--config', str(ini)]
+            output = subprocess.check_output(cmd, env=env, text=True)
+            self.assertIn('export D4R_ENGINE=1\n', output)
+            self.assertIn('export D4R_ENGINE_K=0\n', output)
+            output = subprocess.check_output(cmd, env=dict(env, D4R_ENGINE='0'), text=True)
+            self.assertNotIn('export D4R_ENGINE=', output)
+
     def test_launcher_marker_poll_defaults_overrides_and_validation(self):
         with tempfile.TemporaryDirectory() as temporary:
             ini = Path(temporary) / 'd4r.ini'
diff --git a/tools/d3d12_dlss_harness.cpp b/tools/d3d12_dlss_harness.cpp
index f2465ea..43c0ac7 100644
--- a/tools/d3d12_dlss_harness.cpp
+++ b/tools/d3d12_dlss_harness.cpp
@@ -127,14 +127,29 @@ static bool check(HRESULT hr, const char* what)
 
 static void submit_and_wait()
 {
-    g_list->Close();
+    if (!check(g_list->Close(), "Close command list")) std::_Exit(1);
     ID3D12CommandList* lists[] = {g_list};
     g_queue->ExecuteCommandLists(1, lists);
-    g_queue->Signal(g_fence, ++g_fenceValue);
+    if (!check(g_queue->Signal(g_fence, ++g_fenceValue), "Signal fence")) std::_Exit(1);
     if (g_fence->GetCompletedValue() < g_fenceValue)
     {
-        g_fence->SetEventOnCompletion(g_fenceValue, g_event);
-        WaitForSingleObject(g_event, INFINITE);
+        if (!check(g_fence->SetEventOnCompletion(g_fenceValue, g_event), "SetEventOnCompletion")) std::_Exit(1);
+        DWORD timeout = 60000;
+        if (const char* value = std::getenv("D4R_HARNESS_GPU_TIMEOUT_MS"))
+        {
+            char* end = nullptr;
+            const unsigned long parsed = std::strtoul(value, &end, 10);
+            if (end != value && *end == '\0' && parsed >= 1 && parsed <= 600000) timeout = parsed;
+        }
+        const DWORD waited = WaitForSingleObject(g_event, timeout);
+        if (waited != WAIT_OBJECT_0)
+        {
+            std::fprintf(stderr, "GPU fence wait failed: wait=0x%lx timeout=%lu ms fence=%llu completed=%llu device=0x%08lx\n",
+                         waited, timeout, g_fenceValue, g_fence->GetCompletedValue(), g_device->GetDeviceRemovedReason());
+            std::fflush(stderr);
+            // The capture is an isolated diagnostic process. Do not enter teardown with outstanding GPU work.
+            std::_Exit(waited == WAIT_TIMEOUT ? 124 : 1);
+        }
     }
     g_allocator->Reset();
     g_list->Reset(g_allocator, nullptr);
@@ -1043,6 +1058,9 @@ static float halton(uint32_t index, uint32_t base)
 
 int main(int argc, char** argv)
 {
+    // The installer redirects output to a file. Keep the last stage visible even if HIP/Wine stalls or crashes.
+    std::setvbuf(stdout, nullptr, _IONBF, 0);
+    std::setvbuf(stderr, nullptr, _IONBF, 0);
     trace_stage("entered main");
     if (argc < 3)
     {
diff --git a/tools/d4r_build.h b/tools/d4r_build.h
new file mode 100644
index 0000000..e78d56d
--- /dev/null
+++ b/tools/d4r_build.h
@@ -0,0 +1,4 @@
+#pragma once
+#ifndef D4R_RUNTIME_BUILD_ID
+#define D4R_RUNTIME_BUILD_ID "unidentified (built without runtime_build_id.sh)"
+#endif
diff --git a/tools/d4r_engine_compat.h b/tools/d4r_engine_compat.h
new file mode 100644
index 0000000..d80e8b9
--- /dev/null
+++ b/tools/d4r_engine_compat.h
@@ -0,0 +1,35 @@
+#pragma once
+#include <vulkan/vulkan_core.h>
+#include <sstream>
+#include <string>
+
+// Inspect metadata before constructing any pipeline. Old default D/E and F models require FP8;
+// only the explicitly marked gfx11 D/E models use widened f16 cooperative matrices.
+inline bool d4r_rr_model_compatible(const std::string& metadata, bool f16Matrices, bool fp8Matrices)
+{
+    std::istringstream lines(metadata);
+    std::string line;
+    if (!std::getline(lines, line)) return false;
+    if (!line.empty() && line.back() == '\r') line.pop_back();
+    if (line == "D4RR0002") return fp8Matrices;
+    if (line != "D4RR0001") return false;
+    bool widened = false;
+    while (std::getline(lines, line))
+    {
+        std::istringstream fields(line);
+        std::string key, value; fields >> key >> value;
+        if (key == "arithmetic")
+        {
+            if (value != "gfx11") return false;
+            widened = true;
+        }
+    }
+    return widened ? f16Matrices : fp8Matrices;
+}
+
+inline bool d4r_matrix_tuple(const VkCooperativeMatrixPropertiesKHR& p, VkComponentTypeKHR operand)
+{
+    return p.MSize == 16 && p.NSize == 16 && p.KSize == 16 && p.scope == VK_SCOPE_SUBGROUP_KHR &&
+        p.AType == operand && p.BType == operand && p.CType == VK_COMPONENT_TYPE_FLOAT32_KHR &&
+        p.ResultType == VK_COMPONENT_TYPE_FLOAT32_KHR && !p.saturatingAccumulation;
+}
diff --git a/tools/d4r_fsr_rr_identity.h b/tools/d4r_fsr_rr_identity.h
index 6e65e73..5052c23 100644
--- a/tools/d4r_fsr_rr_identity.h
+++ b/tools/d4r_fsr_rr_identity.h
@@ -73,7 +73,7 @@ inline bool d4r_read_hex_file(const char* path, uint32_t& value)
 
 // The real PCI identity of the GPU behind a vkd3d-proton device: Vulkan's PCI bus address, then that
 // device's sysfs vendor/device files.
-inline D4rPciIdentity d4r_true_identity(ID3D12Device* device, std::string& detail)
+inline D4rPciIdentity d4r_true_identity(ID3D12Device* device, std::string& detail, std::string* address = nullptr)
 {
     D4rPciIdentity identity;
     ID3D12DXVKInteropDevice1* interop = nullptr;
@@ -126,6 +126,12 @@ inline D4rPciIdentity d4r_true_identity(ID3D12Device* device, std::string& detai
         return identity;
     }
     identity.known = true;
+    if (address != nullptr)
+    {
+        char pci[32];
+        std::snprintf(pci, sizeof(pci), "%04x:%02x:%02x.%x", bus.pciDomain, bus.pciBus, bus.pciDevice, bus.pciFunction);
+        *address = pci;
+    }
     char text[256];
     std::snprintf(text, sizeof(text), "PCI %04x:%02x:%02x.%x is %04x:%04x (Vulkan reports %04x:%04x, driver %s)",
                   bus.pciDomain, bus.pciBus, bus.pciDevice, bus.pciFunction, identity.vendor, identity.device,
diff --git a/tools/d4r_gpu_identity.h b/tools/d4r_gpu_identity.h
new file mode 100644
index 0000000..61c8200
--- /dev/null
+++ b/tools/d4r_gpu_identity.h
@@ -0,0 +1,26 @@
+/* PCI identity comparisons shared by the Wine bridge and CPU tests. */
+#pragma once
+#include <stdio.h>
+
+typedef struct D4rPciAddress { unsigned domain, bus, device, function; } D4rPciAddress;
+
+static int d4r_parse_pci(const char* text, D4rPciAddress* out)
+{
+    char extra;
+    return text != NULL && sscanf(text, "%x:%x:%x.%x%c", &out->domain, &out->bus,
+        &out->device, &out->function, &extra) == 4 && out->bus <= 255 && out->device <= 31 && out->function <= 7;
+}
+
+static int d4r_same_pci(const char* a, const char* b)
+{
+    D4rPciAddress x, y;
+    return d4r_parse_pci(a, &x) && d4r_parse_pci(b, &y) && x.domain == y.domain &&
+        x.bus == y.bus && x.device == y.device && x.function == y.function;
+}
+
+static int d4r_pci_matches_kfd(const char* pci, unsigned long domain, unsigned long location)
+{
+    D4rPciAddress address;
+    return d4r_parse_pci(pci, &address) && address.domain == domain &&
+        ((address.bus << 8) | (address.device << 3) | address.function) == location;
+}
diff --git a/tools/d4r_nvngx_shim.cpp b/tools/d4r_nvngx_shim.cpp
index 3b34a2d..28eb11a 100644
--- a/tools/d4r_nvngx_shim.cpp
+++ b/tools/d4r_nvngx_shim.cpp
@@ -101,6 +101,8 @@
 #include "d4r_fsr_rr.h"
 #include "d4r_fsr_rr_identity.h"
 #include "d4r_fsr4_identity.h"
+#include "d4r_engine_compat.h"
+#include "d4r_build.h"
 #include "d4r_fsr_rr_passes.h"
 #include "../engine/game.h"
 #include "../engine/game_m.h"
@@ -753,7 +755,7 @@ static void load_portable_config()
         bool on = false;
         if (enabled.empty() || enabled == "auto")
         {
-            for (const wchar_t* model : {L"k", L"k-ldr", L"m", L"m-ldr", L"l", L"l-ldr", L"rr-d", L"rr-e", L"rr-f"})
+            for (const wchar_t* model : {L"k", L"k-ldr", L"m", L"m-ldr", L"l", L"l-ldr", L"rr-d", L"rr-e", L"rr-f", L"rr-d-gfx11", L"rr-e-gfx11"})
             {
                 const DWORD attributes = GetFileAttributesW((dir + L"\\engine\\" + model).c_str());
                 if (attributes != INVALID_FILE_ATTRIBUTES && (attributes & FILE_ATTRIBUTE_DIRECTORY))
@@ -957,12 +959,16 @@ public:
         return jobs_.size();
     }
 
-    template <typename Function> auto call(Function function) -> decltype(function())
+    template <typename Function> auto call(Function function, const char* label = nullptr) -> decltype(function())
     {
         using Result = decltype(function());
         auto task = std::make_shared<std::packaged_task<Result()>>(std::move(function));
         auto future = task->get_future();
         post([task] { (*task)(); });
+        const auto started = std::chrono::steady_clock::now();
+        while (label && future.wait_for(std::chrono::seconds(15)) == std::future_status::timeout)
+            logf("worker wait: %s still pending after %.0f s (resources retained)", label,
+                 std::chrono::duration<double>(std::chrono::steady_clock::now() - started).count());
         return future.get();
     }
 
@@ -1975,6 +1981,7 @@ static NgxResult initialize(unsigned long long applicationId, const wchar_t* dat
     ensure_portable_config();
     if (g_portable.active && !g.core)
     {
+        logf("d4r runtime build %s", D4R_RUNTIME_BUILD_ID);
         logf("portable install in %ls (settings from d4r.ini, Linux path %s)", g_portable.dir.c_str(),
              g_portable.unixDir.c_str());
         for (const std::string& note : g_portable.notes)
@@ -1994,6 +2001,11 @@ static NgxResult initialize(unsigned long long applicationId, const wchar_t* dat
         LUID luid;
         device->GetAdapterLuid(&luid);
         set_cuda_adapter_luid(luid);
+        std::string adapterDetail, adapterPci;
+        d4r_true_identity(device, adapterDetail, &adapterPci);
+        SetEnvironmentVariableA("D4R_RENDER_PCI", adapterPci.empty() ? nullptr : adapterPci.c_str());
+        logf("CUDA adapter mapping: %s; PCI match %s", adapterDetail.c_str(),
+             adapterPci.empty() ? "unavailable (multiple compute devices will be refused)" : adapterPci.c_str());
         // OptiScaler's FSR 4 option needs the real AMD identity for OptiScaler and AMD's modules (d4r_fsr4_identity.h)
         if (env_string("D4R_FSR4_IDENTITY") != "0")
         {
@@ -2004,6 +2016,8 @@ static NgxResult initialize(unsigned long long applicationId, const wchar_t* dat
     }
 
     g.paths.clear();
+    logf("engine selection: D4R_ENGINE=%s D4R_ENGINE_K=%s network=%s (environment overrides d4r.ini)",
+         env_string("D4R_ENGINE").c_str(), env_string("D4R_ENGINE_K").c_str(), env_string("D4R_ENGINE_NETWORK").c_str());
     wchar_t featureDirectory[MAX_PATH] = {};
     DWORD length = GetEnvironmentVariableW(L"D4R_NGX_FEATURE_DIR", featureDirectory, MAX_PATH);
     if (length > 0 && length < MAX_PATH)
@@ -2047,6 +2061,22 @@ static NgxResult initialize(unsigned long long applicationId, const wchar_t* dat
         result = g.worker.call([&] { return g.ngx.init(id, g.dataPath.c_str(), &g.featureInfo, apiVersion); });
         logf("NVSDK_NGX_CUDA_Init(app=%llu, sdk=0x%x) -> 0x%08x", id, version, result);
     }
+    const bool hasProject = project != nullptr && !project->id.empty();
+    if (result != NGX_SUCCESS && !hasProject && g.ngx.initProjectId != nullptr)
+    {
+        // NGX 617.14 rejects application-ID initialization on this path (Cyberpunk 2077:
+        // 0xBAD0000C, or SR unavailable with 0xBAD00004). Streamline's RR plugin reads
+        // capabilities right after that first Init, so a failure there hides RR for the
+        // session even though the game's later project-ID Init works. Retry with d4r's
+        // own project identity (the diagnostics' default) so the first Init succeeds.
+        logf("D4R_NGX_IDENTITY_FALLBACK application=%llu: 0x%08x", id, result);
+        result = g.worker.call([&] {
+            return g.ngx.initProjectId("24480451-f00d-face-1304-0308dabad187", 0, "1.0", g.dataPath.c_str(),
+                                       apiVersion, &g.featureInfo);
+        });
+        logf("NVSDK_NGX_CUDA_Init_ProjectID(d4r fallback project=24480451-f00d-face-1304-0308dabad187, engine=0 1.0, sdk=0x%x) -> 0x%08x",
+             version, result);
+    }
     using LoadErrorFn = const char*(WINAPI*)();
     const auto loadError = reinterpret_cast<LoadErrorFn>(reinterpret_cast<void*>(GetProcAddress(g.cuda, "d4rLoadError")));
     if (loadError != nullptr && loadError()[0] != '\0')
@@ -3388,6 +3418,10 @@ struct VulkanInterop
     // VK_EXT_shader_float8 (vkd3d-proton enables it when the GPU has it): M's and L's engine layers are FP8
     // cooperative-matrix shaders, which RADV's compiler crashes on without it (RDNA3)
     bool float8 = false;
+    // shaderBufferInt64Atomics is enabled on vkd3d-proton's device (it enables it when the GPU has it): Ray
+    // Reconstruction preset F then measures the exposure in one pass, as NGX does, instead of two
+    bool int64Atomics = false;
+    bool matrixF16 = false;
     // The engine runs on the game's own Vulkan device (native Vulkan NGX frontend): no vkd3d interop objects.
     bool native = false;
     PFN_vkGetPhysicalDeviceFormatProperties formatProperties = nullptr;
@@ -8936,6 +8970,9 @@ static int start_engine_model_rr(const Feature& feature);
 // Vulkan-only K features retain just their creation settings until a frame actually needs CUDA.
 static NgxResult create_cuda_feature(Feature& feature)
 {
+    const auto createStarted = std::chrono::steady_clock::now();
+    logf("CUDA_CREATE_BEGIN preset=%u render=%ux%u output=%ux%u flags=0x%x", feature.preset,
+         feature.width, feature.height, feature.outWidth, feature.outHeight, feature.flags);
     static const char* const kSuperSamplingPresets[] = {
         "DLSS.Hint.Render.Preset.DLAA", "DLSS.Hint.Render.Preset.Quality", "DLSS.Hint.Render.Preset.Balanced",
         "DLSS.Hint.Render.Preset.Performance", "DLSS.Hint.Render.Preset.UltraPerformance",
@@ -8947,6 +8984,7 @@ static NgxResult create_cuda_feature(Feature& feature)
     const bool rayReconstruction = feature.rayReconstruction;
     const char* const* presetNames = rayReconstruction ? kRayReconstructionPresets : kSuperSamplingPresets;
     const NgxResult result = g.worker.call([&]() -> NgxResult {
+        logf("CUDA_CREATE_WORKER_BEGIN preset=%u", feature.preset);
         NgxResult allocation = g.ngx.allocateParameters(&feature.cudaParams);
         if (allocation != NGX_SUCCESS)
             return allocation;
@@ -8986,8 +9024,10 @@ static NgxResult create_cuda_feature(Feature& feature)
             feature.scratch = 0;
         // The core owns the denoiser: it loads nvngx_dlssd.dll, knows its kernel map, and
         // dispatches feature 13 to it. Only the feature id differs from Super Sampling.
+        logf("CUDA_CREATE_NGX_BEGIN preset=%u scratch=0x%llx", feature.preset, feature.scratch);
         const NgxResult created = g.ngx.createFeature(
             rayReconstruction ? NGX_FEATURE_RAY_RECONSTRUCTION : NGX_FEATURE_SUPER_SAMPLING, p, &feature.cudaHandle);
+        logf("CUDA_CREATE_NGX_END preset=%u result=0x%08x", feature.preset, created);
         if (created != NGX_SUCCESS)
         {
             if (feature.cudaHandle != nullptr) g.ngx.releaseFeature(feature.cudaHandle);
@@ -9046,9 +9086,11 @@ static NgxResult create_cuda_feature(Feature& feature)
             }
         }
         return created;
-    });
+    }, "CUDA feature creation");
     logf("NVSDK_NGX_CUDA_CreateFeature(feature=%u) -> 0x%08x",
          rayReconstruction ? NGX_FEATURE_RAY_RECONSTRUCTION : NGX_FEATURE_SUPER_SAMPLING, result);
+    logf("CUDA_CREATE_END preset=%u result=0x%08x elapsed_ms=%.1f", feature.preset, result,
+         std::chrono::duration<double, std::milli>(std::chrono::steady_clock::now() - createStarted).count());
     if (result == NGX_SUCCESS) feature.cudaCreationDeferred = false;
     return result;
 }
@@ -9323,8 +9365,46 @@ static void init_engine_vulkan()
             for (uint32_t index = 0; index < count; ++index)
                 g_vk.float8 |= std::strcmp(extensions[index].extensionName, "VK_EXT_shader_float8") == 0;
     }
+    auto features2 = reinterpret_cast<PFN_vkGetPhysicalDeviceFeatures2>(instanceProc(instance, "vkGetPhysicalDeviceFeatures2"));
+    auto matrices = reinterpret_cast<PFN_vkGetPhysicalDeviceCooperativeMatrixPropertiesKHR>(
+        instanceProc(instance, "vkGetPhysicalDeviceCooperativeMatrixPropertiesKHR"));
+    bool fp8 = false;
+    if (features2 && matrices)
+    {
+        VkPhysicalDeviceShaderFloat8FeaturesEXT f8 = {VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SHADER_FLOAT8_FEATURES_EXT};
+        VkPhysicalDeviceCooperativeMatrixFeaturesKHR cm = {VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_COOPERATIVE_MATRIX_FEATURES_KHR};
+        cm.pNext = g_vk.float8 ? &f8 : nullptr;
+        VkPhysicalDeviceFeatures2 features = {VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_FEATURES_2, &cm};
+        features2(g_vk.physical, &features);
+        uint32_t n = 0;
+        if (cm.cooperativeMatrix && matrices(g_vk.physical, &n, nullptr) == VK_SUCCESS)
+        {
+            std::vector<VkCooperativeMatrixPropertiesKHR> properties(n, {VK_STRUCTURE_TYPE_COOPERATIVE_MATRIX_PROPERTIES_KHR});
+            if (matrices(g_vk.physical, &n, properties.data()) == VK_SUCCESS)
+                for (uint32_t i = 0; i < n; ++i)
+                {
+                    g_vk.matrixF16 |= d4r_matrix_tuple(properties[i], VK_COMPONENT_TYPE_FLOAT16_KHR);
+                    fp8 |= d4r_matrix_tuple(properties[i], VK_COMPONENT_TYPE_FLOAT8_E4M3_EXT);
+                }
+        }
+        fp8 = fp8 && f8.shaderFloat8 && f8.shaderFloat8CooperativeMatrix;
+    }
+    g_vk.float8 = g_vk.float8 && fp8;
+    logf("engine backend: f16 matrix tuple %s; FP8 matrix tuple %s", g_vk.matrixF16 ? "supported" : "unsupported",
+         g_vk.float8 ? "supported" : "unsupported");
     logf("engine backend: VK_EXT_shader_float8 %s", g_vk.float8 ? "available (M and L engine layers)" : "unavailable: "
          "M and L run on the CUDA backend");
+    // the features vkd3d-proton created its device with
+    const VkPhysicalDeviceFeatures2* features = nullptr;
+    if (SUCCEEDED(g_vk.interop->GetDeviceFeatures(&features)) && features != nullptr)
+        for (auto* next = static_cast<const VkBaseInStructure*>(features->pNext); next != nullptr; next = next->pNext)
+        {
+            if (next->sType == VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_VULKAN_1_2_FEATURES)
+                g_vk.int64Atomics = reinterpret_cast<const VkPhysicalDeviceVulkan12Features*>(next)->shaderBufferInt64Atomics == VK_TRUE;
+            if (next->sType == VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SHADER_ATOMIC_INT64_FEATURES)
+                g_vk.int64Atomics = reinterpret_cast<const VkPhysicalDeviceShaderAtomicInt64Features*>(next)->shaderBufferInt64Atomics == VK_TRUE;
+        }
+    logf("engine backend: 64-bit buffer atomics %s", g_vk.int64Atomics ? "enabled" : "not enabled: preset F's exposure in two passes");
 }
 
 static void init_engine_backend()
@@ -10487,6 +10567,8 @@ static std::string rr_model_directory(int slot)
         char folder[MAX_PATH * 3] = {};
         WideCharToMultiByte(CP_ACP, 0, (g_portable.dir + folders[slot]).c_str(), -1, folder, sizeof(folder), nullptr, nullptr);
         directory = folder;
+        if (slot < 2 && !g_vk.float8)
+            directory += "-gfx11";
     }
     return directory[0] == '/' ? "Z:" + directory : directory;
 }
@@ -10505,7 +10587,11 @@ static void init_engine_backend_rr(int slot)
     const auto start = std::chrono::steady_clock::now();
     try
     {
-        auto model = std::make_unique<d4r::RModel>(g_vk.physical, g_vk.device, path);
+        std::string metadata;
+        if (!read_whole_file(widen(path + "\\rr.txt"), metadata) ||
+            !d4r_rr_model_compatible(metadata, g_vk.matrixF16, g_vk.float8))
+            throw std::runtime_error("model arithmetic is unsupported by this GPU (no pipelines created)");
+        auto model = std::make_unique<d4r::RModel>(g_vk.physical, g_vk.device, path, VK_NULL_HANDLE, g_vk.int64Atomics);
         if (model->preset() != g_rrPreset[slot])
             throw std::runtime_error("the directory holds a model of another preset");
         g_engineR.models[slot] = std::move(model);
diff --git a/tools/d4r_vulkan_frontend.inl b/tools/d4r_vulkan_frontend.inl
index 0ded9d5..b14c4b7 100644
--- a/tools/d4r_vulkan_frontend.inl
+++ b/tools/d4r_vulkan_frontend.inl
@@ -620,6 +620,13 @@ static NgxResult native_vk_init(unsigned long long id, const wchar_t* path, VkIn
         VkPhysicalDeviceIDProperties identity = {VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_ID_PROPERTIES};
         VkPhysicalDeviceProperties2 properties = {VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PROPERTIES_2, &identity};
         properties2(physical, &properties);
+        VkPhysicalDevicePCIBusInfoPropertiesEXT bus = {VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PCI_BUS_INFO_PROPERTIES_EXT};
+        properties.pNext = &bus;
+        properties2(physical, &properties);
+        char pci[32] = {};
+        if (bus.pciDomain || bus.pciBus || bus.pciDevice || bus.pciFunction)
+            std::snprintf(pci, sizeof(pci), "%04x:%02x:%02x.%x", bus.pciDomain, bus.pciBus, bus.pciDevice, bus.pciFunction);
+        SetEnvironmentVariableA("D4R_RENDER_PCI", pci[0] ? pci : nullptr);
         if (identity.deviceLUIDValid)
         {
             uint32_t low, high;
diff --git a/tools/wine_nvcuda_bridge.c b/tools/wine_nvcuda_bridge.c
index 7af8670..cc1c6ee 100644
--- a/tools/wine_nvcuda_bridge.c
+++ b/tools/wine_nvcuda_bridge.c
@@ -13,6 +13,8 @@
 #include <unistd.h>
 #include <time.h>
 #include "d4r_native_selection.h"
+#include "d4r_gpu_identity.h"
+#include "d4r_build.h"
 #include "d4r_motion_dilation.h"
 #include "d4r_native_motion.h"
 #include "d4r_wire_history.h"
@@ -430,46 +432,54 @@ static void native_cleanup(void)
         remove_directory(native_served);
 }
 
-/* The gfx target ("gfx1101") of the GPU DLSS runs on, "" when unknown: D4R_GPU_ARCH if set, else the KFD
-   topology's GPU with the most SIMDs. With an integrated and a discrete GPU (Ryzen 7000/9000 desktops) the
-   first KFD GPU node can be the integrated one, whose kernels would not match the discrete GPU. */
+/* Use the rendering adapter's PCI address, never the first/largest KFD GPU. Without an authenticated
+   address a single GPU is unambiguous; multiple GPUs use translated kernels instead of a guessed target. */
 static void gpu_architecture(char* out, size_t size)
 {
     out[0] = '\0';
-    const char* forced = getenv("D4R_GPU_ARCH");
-    if (forced != NULL && strncmp(forced, "gfx", 3) == 0)
+    char pci[64] = {}, only[32] = {};
+    GetEnvironmentVariableA("D4R_RENDER_PCI", pci, sizeof(pci));
+    int gpus = 0, matches = 0;
+    DIR* nodes = opendir("/sys/class/kfd/kfd/topology/nodes");
+    if (nodes == NULL) return;
+    for (struct dirent* node; (node = readdir(nodes)) != NULL;)
     {
-        snprintf(out, size, "%s", forced);
-        return;
-    }
-    unsigned long best_simds = 0;
-    int gpus = 0;
-    for (int node = 0; node < 16; ++node)
-    {
-        char path[128];
-        snprintf(path, sizeof(path), "/sys/class/kfd/kfd/topology/nodes/%d/properties", node);
+        if (node->d_name[0] == '.') continue;
+        char path[512];
+        snprintf(path, sizeof(path), "/sys/class/kfd/kfd/topology/nodes/%s/properties", node->d_name);
         FILE* properties = fopen(path, "r");
         if (properties == NULL)
             continue;
         char line[256];
-        unsigned long simds = 0, target = 0;
+        unsigned long simds = 0, target = 0, domain = ~0ul, location = ~0ul;
         while (fgets(line, sizeof(line), properties) != NULL)
         {
             sscanf(line, "simd_count %lu", &simds);
             sscanf(line, "gfx_target_version %lu", &target);
+            sscanf(line, "domain %lu", &domain);
+            sscanf(line, "location_id %lu", &location);
         }
         fclose(properties);
         if (simds == 0 || target == 0)
             continue;
         ++gpus;
-        if (simds > best_simds)
+        snprintf(only, sizeof(only), "gfx%lu%lu%lx", target / 10000, (target / 100) % 100, target % 100);
+        if (d4r_pci_matches_kfd(pci, domain, location))
         {
-            best_simds = simds;
-            snprintf(out, size, "gfx%lu%lu%lx", target / 10000, (target / 100) % 100, target % 100);
+            ++matches;
+            snprintf(out, size, "%s", only);
         }
     }
-    if (gpus > 1)
-        tracef("native kernels: %d GPUs; using the one with the most SIMDs (%s); D4R_GPU_ARCH overrides", gpus, out);
+    closedir(nodes);
+    if (!pci[0] && gpus == 1) snprintf(out, size, "%s", only);
+    else if (matches != 1) out[0] = '\0';
+    const char* forced = getenv("D4R_GPU_ARCH");
+    if (forced != NULL && strcmp(forced, out) != 0)
+    {
+        tracef("D4R_GPU_ARCH=%s does not match the rendering GPU (%s); native kernels disabled", forced, out);
+        out[0] = '\0';
+    }
+    tracef("GPU target: render PCI=%s KFD GPUs=%d matches=%d architecture=%s", pci, gpus, matches, out);
 }
 
 static void prepare_native_kernels(const char* cache_home)
@@ -480,6 +490,11 @@ static void prepare_native_kernels(const char* cache_home)
     /* a release directory holds one manifest and set of code objects per GPU target */
     char architecture[32], source[1024], path[1200];
     gpu_architecture(architecture, sizeof(architecture));
+    if (!architecture[0])
+    {
+        unsetenv("D4R_ZLUDA_NATIVE_DIR");
+        return;
+    }
     FILE* manifest = NULL;
     /* RDNA4 with native FP8 WMMA (D4R_ZLUDA_WMMA_FP8_NATIVE=1, d4r.ini NativeFp8): the <target>-fp8 variant,
        whose kernels match ZLUDA's FP8 lowering; no other GPU has such a directory */
@@ -1385,8 +1400,12 @@ static void tracef(const char* format, ...)
     vsnprintf(line, sizeof(line), format, args);
     va_end(args);
 
+    struct timespec now;
+    clock_gettime(CLOCK_MONOTONIC, &now);
+    const double seconds = (double)now.tv_sec + (double)now.tv_nsec / 1000000000.0;
+    const unsigned long thread = GetCurrentThreadId();
     pthread_mutex_lock(&trace_lock);
-    fprintf(stderr, "[d4r nvcuda] %s\n", line);
+    fprintf(stderr, "[d4r nvcuda] [%.6f tid=%lu] %s\n", seconds, thread, line);
     fflush(stderr);
 
     char path[MAX_PATH] = {0};
@@ -1396,9 +1415,14 @@ static void tracef(const char* format, ...)
         FILE* output = fopen(path, "a");
         if (output != NULL)
         {
-            fprintf(output, "%s\n", line);
+            fprintf(output, "[%.6f tid=%lu] %s\n", seconds, thread, line);
             fclose(output);
         }
+        else
+        {
+            static int warned;
+            if (!warned++) fprintf(stderr, "[d4r nvcuda] cannot open D4R_CUDA_TRACE=%s: %s\n", path, strerror(errno));
+        }
     }
     pthread_mutex_unlock(&trace_lock);
 }
@@ -1473,7 +1497,8 @@ typedef CUresult(__attribute__((sysv_abi)) *CUGETERRORSTRING_FN)(CUresult, const
 typedef CUresult(__attribute__((sysv_abi)) *CUGETPROCADDRESS_FN)(const char*, void**, int, uint64_t, int*);
 
 static pthread_once_t context_once = PTHREAD_ONCE_INIT;
-static CUresult context_setup_result = CUDA_ERROR_NOT_INITIALIZED;
+static CUresult context_bootstrap_result = CUDA_ERROR_NOT_INITIALIZED;
+static _Thread_local CUresult context_setup_result = CUDA_ERROR_NOT_INITIALIZED;
 static CUcontext cuda_context;
 
 static unsigned int get_process_u32(const char* name, unsigned int fallback)
@@ -1487,6 +1512,36 @@ static unsigned int get_process_u32(const char* name, unsigned int fallback)
     return end != value && *end == '\0' ? (unsigned int)parsed : fallback;
 }
 
+/* This calls ZLUDA directly: bridge wrappers would recurse through context bootstrap. */
+static CUresult selected_cuda_device(CUdevice* selected)
+{
+    CUDEVICEGETCOUNT_FN count_fn = (CUDEVICEGETCOUNT_FN)find_zluda_symbol("cuDeviceGetCount");
+    CUDEVICEGET_FN get_fn = (CUDEVICEGET_FN)find_zluda_symbol("cuDeviceGet");
+    typedef CUresult(__attribute__((sysv_abi)) * PCI_FN)(char*, int, CUdevice);
+    PCI_FN pci_fn = (PCI_FN)find_zluda_symbol("cuDeviceGetPCIBusId");
+    int count = 0;
+    if (!count_fn || !get_fn || count_fn(&count) != CUDA_SUCCESS || count < 1)
+        return CUDA_ERROR_NOT_INITIALIZED;
+    char expected[64] = {};
+    GetEnvironmentVariableA("D4R_RENDER_PCI", expected, sizeof(expected));
+    if (!expected[0] && count == 1) return get_fn(selected, 0);
+    if (!expected[0] || !pci_fn) return CUDA_ERROR_NOT_SUPPORTED;
+    int matches = 0;
+    for (int ordinal = 0; ordinal < count; ++ordinal)
+    {
+        CUdevice device;
+        char actual[64] = {};
+        if (get_fn(&device, ordinal) == CUDA_SUCCESS && pci_fn(actual, sizeof(actual), device) == CUDA_SUCCESS)
+        {
+            tracef("compute adapter ordinal=%d device=%d PCI=%s render=%s", ordinal, device, actual, expected);
+            if (d4r_same_pci(expected, actual)) { *selected = device; ++matches; }
+        }
+    }
+    if (matches != 1)
+        tracef("CUDA adapter mapping refused: render PCI=%s matches=%d; no GPU will be guessed", expected, matches);
+    return matches == 1 ? CUDA_SUCCESS : CUDA_ERROR_NOT_SUPPORTED;
+}
+
 static void create_default_zluda_context(void)
 {
     CUINIT_FN init = (CUINIT_FN)find_zluda_symbol("cuInit");
@@ -1499,30 +1554,30 @@ static void create_default_zluda_context(void)
         return;
     }
 
-    context_setup_result = init(0);
-    if (context_setup_result != CUDA_SUCCESS)
+    context_bootstrap_result = init(0);
+    if (context_bootstrap_result != CUDA_SUCCESS)
     {
-        tracef("cuInit returned %d", context_setup_result);
+        tracef("cuInit returned %d", context_bootstrap_result);
         return;
     }
 
     int device_count = 0;
-    context_setup_result = get_device_count(&device_count);
-    if (context_setup_result != CUDA_SUCCESS || device_count < 1)
+    context_bootstrap_result = get_device_count(&device_count);
+    if (context_bootstrap_result != CUDA_SUCCESS || device_count < 1)
     {
-        tracef("cuDeviceGetCount returned %d, count %d", context_setup_result, device_count);
-        if (context_setup_result == CUDA_SUCCESS)
-            context_setup_result = CUDA_ERROR_NOT_INITIALIZED;
+        tracef("cuDeviceGetCount returned %d, count %d", context_bootstrap_result, device_count);
+        if (context_bootstrap_result == CUDA_SUCCESS)
+            context_bootstrap_result = CUDA_ERROR_NOT_INITIALIZED;
         return;
     }
 
     CUdevice device = 0;
-    context_setup_result = get_device(&device, 0);
-    if (context_setup_result == CUDA_SUCCESS)
-        context_setup_result = create_context(&cuda_context, 0, device);
-    if (context_setup_result != CUDA_SUCCESS)
+    context_bootstrap_result = selected_cuda_device(&device);
+    if (context_bootstrap_result == CUDA_SUCCESS)
+        context_bootstrap_result = create_context(&cuda_context, 0, device);
+    if (context_bootstrap_result != CUDA_SUCCESS)
     {
-        tracef("context creation returned %d", context_setup_result);
+        tracef("context creation returned %d", context_bootstrap_result);
         return;
     }
 
@@ -1531,7 +1586,21 @@ static void create_default_zluda_context(void)
 
 static void ensure_context(void)
 {
+    typedef CUresult(__attribute__((sysv_abi)) * CURRENT_FN)(CUcontext*);
+    CURRENT_FN current = (CURRENT_FN)find_zluda_symbol("cuCtxGetCurrent");
+    CUcontext active = NULL;
+    if (current && current(&active) == CUDA_SUCCESS && active != NULL)
+    {
+        context_setup_result = CUDA_SUCCESS;
+        return; /* NGX already selected a context; do not replace it with a second context. */
+    }
     pthread_once(&context_once, create_default_zluda_context);
+    context_setup_result = context_bootstrap_result;
+    if (context_setup_result == CUDA_SUCCESS && cuda_context != NULL)
+    {
+        CUCTX_CURRENT_FN set = (CUCTX_CURRENT_FN)find_zluda_symbol("cuCtxSetCurrent");
+        if (set) context_setup_result = set(cuda_context);
+    }
 }
 
 /* D4R_NGX_STREAM_PRIORITY=high|normal (default off): the null-stream work of NGX, the shim and this bridge
@@ -2567,6 +2636,7 @@ static int before_teardown(void)
 
 CUresult WINAPI cuInit(unsigned int flags)
 {
+    tracef("bridge build %s; CUDA_BEGIN api=cuInit flags=%u", D4R_RUNTIME_BUILD_ID, flags);
     CUINIT_FN function = (CUINIT_FN)find_zluda_symbol("cuInit");
     CUresult result = function != NULL ? function(flags) : CUDA_ERROR_NOT_INITIALIZED;
     tracef("cuInit flags=%u result=%d", flags, result);
@@ -2618,6 +2688,14 @@ CUresult WINAPI cuCtxGetDevice(CUdevice* device)
 
 CUresult WINAPI cuCtxCreate_v2(CUcontext* context, unsigned int flags, CUdevice device)
 {
+    CUdevice selected = -1;
+    CUresult mapping = selected_cuda_device(&selected);
+    if (mapping != CUDA_SUCCESS || selected != device)
+    {
+        if (context) *context = NULL;
+        tracef("CUDA context refused: device=%d mapped render device=%d status=%d", device, selected, mapping);
+        return mapping == CUDA_SUCCESS ? CUDA_ERROR_NOT_SUPPORTED : mapping;
+    }
     CUCTXCREATE_FN function = (CUCTXCREATE_FN)find_zluda_symbol("cuCtxCreate_v2");
     CUresult result = function != NULL ? function(context, flags, device) : CUDA_ERROR_NOT_INITIALIZED;
     tracef("cuCtxCreate_v2 flags=0x%x device=%d result=%d context=%p", flags, device,
@@ -2659,6 +2737,7 @@ static int elide_ngx_sync(void)
 
 CUresult WINAPI d4rCtxSynchronize(void)
 {
+    if (trace_verbose()) tracef("CUDA_BEGIN api=d4rCtxSynchronize");
     if (before_teardown())
         return CUDA_SUCCESS;
     CUCTX_SYNCHRONIZE_FN function = (CUCTX_SYNCHRONIZE_FN)find_zluda_symbol("cuCtxSynchronize");
@@ -2671,6 +2750,7 @@ CUresult WINAPI d4rCtxSynchronize(void)
 
 CUresult WINAPI cuCtxSynchronize(void)
 {
+    if (trace_verbose()) tracef("CUDA_BEGIN api=cuCtxSynchronize");
     if (before_teardown())
         return CUDA_SUCCESS;
     if (elide_ngx_sync())
@@ -2695,6 +2775,15 @@ CUresult WINAPI cuDeviceGetLuid(char* luid, unsigned int* device_node_mask, CUde
     if ((result == CUDA_SUCCESS || result == CUDA_ERROR_NOT_SUPPORTED) &&
         luid != NULL && device_node_mask != NULL)
     {
+        CUdevice selected = -1;
+        CUresult mapping = selected_cuda_device(&selected);
+        if (mapping != CUDA_SUCCESS || selected != device)
+        {
+            if (luid) memset(luid, 0, 8);
+            if (device_node_mask) *device_node_mask = 0;
+            tracef("cuDeviceGetLuid device=%d is not the mapped render adapter (selected=%d status=%d)", device, selected, mapping);
+            return mapping == CUDA_SUCCESS ? CUDA_SUCCESS : mapping;
+        }
         unsigned int low = get_process_u32("D4R_CUDA_LUID_LOW", 0xffffffffu);
         unsigned int high = get_process_u32("D4R_CUDA_LUID_HIGH", 0xffffffffu);
         unsigned int node_mask = get_process_u32("D4R_CUDA_NODE_MASK", 0xffffffffu);
@@ -2791,6 +2880,7 @@ static const void* replacement_image(const void* image, const char* api)
 
 CUresult WINAPI cuModuleLoadData(CUmodule* module, const void* image)
 {
+    tracef("CUDA_BEGIN api=cuModuleLoadData image=%p", image);
     CUMODULELOADDATA_FN function = (CUMODULELOADDATA_FN)find_zluda_symbol("cuModuleLoadData");
     verify_native_kernels(image);
     capture_ptx_if_present(image, "cuModuleLoadData");
@@ -2804,6 +2894,7 @@ CUresult WINAPI cuModuleLoadData(CUmodule* module, const void* image)
 CUresult WINAPI cuModuleLoadDataEx(CUmodule* module, const void* image, unsigned int option_count,
                                    int* options, void** option_values)
 {
+    tracef("CUDA_BEGIN api=cuModuleLoadDataEx image=%p options=%u", image, option_count);
     CUMODULELOADDATAEX_FN function = (CUMODULELOADDATAEX_FN)find_zluda_symbol("cuModuleLoadDataEx");
     verify_native_kernels(image);
     capture_ptx_if_present(image, "cuModuleLoadDataEx");
@@ -2817,6 +2908,7 @@ CUresult WINAPI cuModuleLoadDataEx(CUmodule* module, const void* image, unsigned
 
 CUresult WINAPI cuModuleLoad(CUmodule* module, const char* filename)
 {
+    tracef("CUDA_BEGIN api=cuModuleLoad file=%s", filename != NULL ? filename : "<null>");
     CUMODULELOAD_FN function = (CUMODULELOAD_FN)find_zluda_symbol("cuModuleLoad");
     CUresult result = function != NULL ? function(module, filename) : missing("cuModuleLoad");
     tracef("cuModuleLoad file=%s result=%d module=%p", filename != NULL ? filename : "<null>",
@@ -2837,6 +2929,7 @@ CUresult WINAPI cuModuleUnload(CUmodule module)
 
 CUresult WINAPI cuModuleGetFunction(CUfunction* function_out, CUmodule module, const char* name)
 {
+    tracef("CUDA_BEGIN api=cuModuleGetFunction module=%p name=%s", module, name ? name : "<null>");
     CUMODULEGETFUNCTION_FN function = (CUMODULEGETFUNCTION_FN)find_zluda_symbol("cuModuleGetFunction");
     CUresult result = function != NULL ? function(function_out, module, name) : missing("cuModuleGetFunction");
     if (result == CUDA_SUCCESS && function_out != NULL)
@@ -2882,6 +2975,7 @@ CUresult WINAPI cuLaunchKernel(CUfunction function_handle, unsigned int grid_x,
                                unsigned int block_z, unsigned int shared_bytes, CUstream stream,
                                void** kernel_params, void** extra)
 {
+    if (trace_verbose()) tracef("CUDA_BEGIN api=cuLaunchKernel function=%p", function_handle);
     static CULAUNCHKERNEL_FN function;
     if (function == NULL)
         function = (CULAUNCHKERNEL_FN)find_zluda_symbol("cuLaunchKernel");
@@ -2964,6 +3058,7 @@ CUresult WINAPI cuLaunchKernel(CUfunction function_handle, unsigned int grid_x,
 
 CUresult WINAPI cuMemAlloc(CUdeviceptr* pointer, size_t bytes)
 {
+    if (trace_verbose()) tracef("CUDA_BEGIN api=cuMemAlloc bytes=%zu", bytes);
     CUMEMALLOC_FN function = (CUMEMALLOC_FN)find_zluda_symbol("cuMemAlloc_v2");
     CUresult result = function != NULL ? function(pointer, bytes) : missing("cuMemAlloc_v2");
     if (result == CUDA_SUCCESS && pointer != NULL)
@@ -3022,6 +3117,8 @@ CUresult WINAPI cuMemcpy2D(const void* copy)
 
 CUresult WINAPI cuMemcpyHtoDAsync(CUdeviceptr destination, const void* source, size_t bytes, CUstream stream)
 {
+    if (trace_verbose()) tracef("CUDA_BEGIN api=cuMemcpyHtoDAsync bytes=%zu dst=0x%llx stream=%p", bytes,
+                              (unsigned long long)destination, stream);
     CUMEMCPYHTODASYNC_FN function = (CUMEMCPYHTODASYNC_FN)find_zluda_symbol("cuMemcpyHtoDAsync_v2");
     if (microcuda_k_write) microcuda_k_write(destination, bytes);
     stream = d4r_ngx_stream(stream);
@@ -3033,6 +3130,7 @@ CUresult WINAPI cuMemcpyHtoDAsync(CUdeviceptr destination, const void* source, s
 
 CUresult WINAPI cuMemcpyDtoH(void* destination, CUdeviceptr source, size_t bytes)
 {
+    if (trace_verbose()) tracef("CUDA_BEGIN api=cuMemcpyDtoH bytes=%zu src=0x%llx", bytes, (unsigned long long)source);
     CUMEMCPYDTOH_FN function = (CUMEMCPYDTOH_FN)find_zluda_symbol("cuMemcpyDtoH_v2");
     CUresult result = function != NULL ? function(destination, source, bytes) : missing("cuMemcpyDtoH_v2");
     TRACE_CALL(result, "cuMemcpyDtoH->v2 dst=%p src=0x%llx bytes=%zu result=%d",
@@ -3378,6 +3476,7 @@ static int d4r_forward_is_teardown(const char* target)
 #define D4R_FORWARD(name, target, count, params, args) \
     CUresult WINAPI name params \
     { \
+        if (trace_verbose()) tracef("CUDA_BEGIN api=" #name); \
         if (d4r_forward_is_teardown(#target) && before_teardown()) \
             return CUDA_SUCCESS; \
         typedef CUresult(__attribute__((sysv_abi)) * function_type) params; \
@@ -3394,6 +3493,7 @@ D4R_FORWARD(cuEventRecord, cuEventRecord, 2, (D4rArg a0, D4rArg a1), (a0, (D4rAr
 D4R_FORWARD(cuEventRecordWithFlags, cuEventRecordWithFlags, 3, (D4rArg a0, D4rArg a1, D4rArg a2), (a0, (D4rArg)d4r_ngx_stream((CUstream)a1), a2))
 CUresult WINAPI cuEventSynchronize(D4rArg a0)
 {
+    if (trace_verbose()) tracef("CUDA_BEGIN api=cuEventSynchronize event=%p", (void*)a0);
     if (before_teardown())
         return CUDA_SUCCESS;
     typedef CUresult(__attribute__((sysv_abi)) * function_type)(D4rArg);
@@ -3411,6 +3511,7 @@ CUresult WINAPI cuEventSynchronize(D4rArg a0)
    game's split command list only after this event has actually completed. */
 CUresult WINAPI d4rEventSynchronize(D4rArg a0)
 {
+    if (trace_verbose()) tracef("CUDA_BEGIN api=d4rEventSynchronize event=%p", (void*)a0);
     if (before_teardown())
         return CUDA_SUCCESS;
     typedef CUresult(__attribute__((sysv_abi)) * function_type)(D4rArg);