zluda: coalesce pair shadow stores, signed native FP8 conversion

Patch 0011 merges a complete accumulator-register pair overwrite into one shadow-layout conversion (reads, rewritten shadow accesses and block boundaries prevent it; partial writes keep read/modify/write), and requests eight waves per SIMD for gfx12 native-FP8 functions. Patch 0012 keeps signs in the native conversion: finite overflow clamped, NaNs mapped to signed infinity, post-conversion sign masks gone; plain and ReLU remain exhaustive-table identical to the software codec. Eighty-frame harness at 720p->4K, excluding the first twenty: control 24.764 ms, coalescing 23.340 ms, plus eight waves 20.135 ms, plus signed conversion 16.234 ms (15.144 ms with the harness idle gap removed); all eighty saved frames byte-identical to the matching control. Twelve- and ten-wave budgets regress (25.56/24.61 ms); fast math changes the image (18.92 ms) and is not selected.
2026-10-12 12:47:53 +02:00
1 parent 406c0286c0
commit 798e25611c
@@ -138,7 +138,7 @@ ROCm does not need to be installed: the zip includes its runtime (from AMD's Ubu
 The full sequence is in [docs/building.md](docs/building.md). In short:
-1. Build ZLUDA `ee2f25a` with `patches/zluda/0002`–`0008` applied, including its `d4r_emit` example for offline texture builds.
+1. Build ZLUDA `ee2f25a` with `patches/zluda/0002`–`0012` applied, including its `d4r_emit` example for offline texture builds.
 2. Build the patched vkd3d-proton: `scripts/build_vkd3d_proton_d4r.sh OUT_DIR`.
 3. Build the shim and bridge: `scripts/build_d4r_nvngx_shim.sh`, `scripts/build_wine_nvcuda_bridge.sh`.
 4. Stage the runtime with your NVIDIA files: `scripts/install_d4r_runtime.sh _nvngx.dll nvngx_dlss.dll`.
@@ -5,7 +5,7 @@ These steps produce the pieces the launcher needs: a patched ZLUDA, a patched vk
 | Component | Tested version |
 |---|---|
-| ZLUDA | `ee2f25a` (upstream), plus `patches/zluda/0002` through `0008` in order |
+| ZLUDA | `ee2f25a` (upstream), plus `patches/zluda/0002` through `0012` in order |
 | vkd3d-proton | `3dfc6f07` (the base GE-Proton11-3 ships), plus `patches/vkd3d-proton/0001` through `0003` in order |
 | ROCm | 7.2 (HIP runtime, clang, device libraries) |
 | Proton | GE-Proton11-3 (with its OptiScaler integration) |
@@ -18,7 +18,7 @@ git clone https://github.com/vosen/ZLUDA zluda && cd zluda
 git checkout ee2f25a
 git submodule update --init --recursive
 git lfs pull
-for p in 0002 0003 0004 0005 0006 0007 0008; do git apply /path/to/d4r/patches/zluda/$p-*.patch; done
+for p in 0002 0003 0004 0005 0006 0007 0008 0009 0010 0011 0012; do git apply /path/to/d4r/patches/zluda/$p-*.patch; done
 # rebuild the device helpers the patches changed (ptx/lib/zluda_ptx_impl*.bc)
 ZLUDA_SOURCE_ROOT=$PWD ROCM_ROOT=/opt/rocm /path/to/d4r/scripts/build_zluda_ptx_helpers.sh
 LIBRARY_PATH=/opt/rocm/lib cargo build --release -p zluda
@@ -47,6 +47,10 @@ What the patches add:
 - **0007**: gfx12 WMMA layout lowering, optional native e4m3 FP8 WMMA (`D4R_ZLUDA_WMMA_FP8_NATIVE=1`), and an architecture argument for `d4r_emit`. `D4R_ZLUDA_WMMA_LAYOUT=12` on gfx11 is a validation shim, not a release setting.
 - **0008**: post-phase kernels for native overrides (`NAME_post1` … `NAME_post8`, see [native-kernels.md](native-kernels.md)). Preset K's dec5 layer runs as four phase kernels, so the K set needs this patch: an older `libnvcuda.so` would skip the phases.
 - **0008**: CUDA denoiser PTX support: four-half surface stores, floating-coordinate half texture sampling, streaming-cache store spellings, byte-vector trapping stores (including low-byte truncation from wider registers), finite-saturating half conversion, correctly typed half exponentiation, half unordered comparisons, and atomic reductions. CUDA block barriers use a workgroup release fence before the hardware barrier and an acquire fence afterward. It also fixes strict floating-point division operand selection in the LLVM AMDGPU backend and Cargo build-profile detection. Two LLVM sinking fixes remove repeated instruction-order rebuilding and avoid alias scans when no sink target exists; write tracking and memory-dependency checks remain intact. Initialize the LLVM submodule before applying this patch; rebuild device helpers and ZLUDA afterward.
 - **0009**: gfx12 native-FP8 half-to-e4m3 requantization, including the ReLU variant. Hardware packed FP8 conversion replaces the software rounding path only when `D4R_ZLUDA_WMMA_FP8_NATIVE=1` selects gfx12 native FP8. Explicit saturation and sign restoration preserve finite overflow, infinities, signed zero and NaNs. Rebuild both the device helpers and ZLUDA.
 - **0010**: gfx12 native FP8 operands stay packed through the NVIDIA-to-WMMA lane gathers, removing fake-half splitting and repacking. Paired and unpaired operations preserve the previous K permutation, accumulation order and FP32-shadow scaling. Unpaired gathers execute in every source lane before zero-padding; masking EXEC around them loses valid source values. Rebuild both the device helpers and ZLUDA.
 - **0011**: coalesce complete accumulator-register overwrites into one shadow-layout conversion. Intervening reads, rewritten shadow accesses and basic-block boundaries prevent coalescing; partial writes retain read/modify/write behavior. Gfx12 native-FP8 functions request eight waves per SIMD to limit register pressure. Rebuild ZLUDA and `d4r_emit`.
 - **0012**: preserve signs through the native FP8 conversion instead of reconstructing them afterward. Explicitly clamp finite overflow and infinities, and map NaNs to signed infinity before conversion: native infinity produces the required `0x7f`/`0xff`, whereas native NaN always produces `0xff`. Plain and ReLU paths remain exhaustive-table identical to software. Rebuild device helpers and ZLUDA.
 When linking on a system with ROCm libraries outside the default search path, include their library directory in `LIBRARY_PATH`. The build also needs the appropriate ROCm link libraries. `CARGO_BUILD_JOBS=8` caps parallel Rust compilation if memory is limited.
@@ -160,7 +164,7 @@ Mount the repository at `/work`, ROCm's compiler/headers/device libraries at `/o
 | Variable | Input |
 |---|---|
 | `D4R_RELEASE_BUILD_ROOT` | writable directory under `/work/build/` |
-| `D4R_ZLUDA_SRC` | isolated checkout inside that directory, with both pinned LLVM and HiGHS submodules, the real OCKL LFS payload, and patches `0002`–`0008` applied |
+| `D4R_ZLUDA_SRC` | isolated checkout inside that directory, with both pinned LLVM and HiGHS submodules, the real OCKL LFS payload, and patches `0002`–`0012` applied |
 | `D4R_VKD3D_SRC` | isolated vkd3d-proton checkout inside that directory, with its submodules and patches `0001`–`0003` applied |
 | `CARGO_HOME` | writable build-local Cargo cache; prefetch the locked dependencies for an offline build |
 | `D4R_ROCM_DIR` | `/opt/rocm` |
@@ -74,33 +74,116 @@ Initial validation: a gfx1101 fast L build completed two Ultra Performance harne
 **Ray Reconstruction, the denoiser (`kernels/tex` + `cuda_dldn_engine_*`).**
 - **Network:** `nvngx_dlssd.dll` defines the DLSS-D denoiser: six Swin encoders (`cuda_dldn_engine_swin_enc0_kernel` … `enc5`), five decoders (`dec0` … `dec4`), a high-fidelity kernel-prediction network whose output kernel is `cuda_dldn_engine_hkpn_output_kernel_transformer`, and the luma/auto-exposure/`reduce_sum` helpers. It shares no kernel with the Super Resolution libraries.
 - **Native today:** the hkpn output kernel (and its `_diamond_wallaby` twin), built by `kernels/build.sh rr` as a wave64 surface-store replacement. `make_ptx.py` replaces `sust.b.2d.v4.b16.zero` — a whole RGBA16F texel — with `d4r_sust_v4b16`; the neural body and the texture reads are unchanged.
- **Still translated:** all eleven Swin layers. ZLUDA compiles NVIDIA's PTX at load time, and that compile is what costs the time (see the table below). A native replacement has to recover each layer's attention, norm and merge stages, the way `kernels/m/rrlite_*` does for the upscaler.
+- **Still translated:** all eleven Swin layers. ZLUDA compiles NVIDIA's PTX at load time; the table below measures the resulting kernels' GPU execution, not JIT latency. A native replacement must preserve each layer's attention, norm and merge stages, as `kernels/m/rrlite_*` does for the upscaler.
 - **Manifest:** the denoiser DLL is passed to `kernel_manifest.py` as `--rr-dll`, because its hashes may only authorize `cuda_dldn_engine_*` replacements. Without that split, a name both libraries define — `dl4rt_input_kernel` — would let a super-resolution binary be served for the denoiser's different implementation of it.
-Measured on the RX 9070 XT (gfx1201, RDNA4), Ray Reconstruction preset E, 1280×720 → 3840×2160, `D4R_CUDA_KERNEL_PROFILE` in the D3D12 harness, 16 frames after 6 warm-up, with the shipped `[Kernels]` settings (`NativeFp8`, `IgnoreDenormals`; `IgnoreDenormals` alone is worth 11% here, and `FastMath`, the f32 accumulator shadows, a 1024- or 128-thread bound, both scheduler strategies and the wave64 switch are all neutral or worse):
+Measured on the RX 9070 XT (gfx1201, RDNA4), Ray Reconstruction preset E,
 1280×720 → 3840×2160, with `D4R_CUDA_KERNEL_PROFILE` in the D3D12 harness:
 The earlier 20-frame run excluded the first 6 frames, with native hkpn output
 and gfx12 native FP8 enabled.
 The first implementation of patch `0009` replaced software half-to-e4m3
 requantization with hardware packed FP8 conversion: **33.90 ms** for the
 same-runtime software control, **26.28 ms** with native conversion, and
 **26.65 ms** without the diagnostic A/B gate. The compact integer saturation
 in `0009` plus `0010`'s packed FP8 operand gathers measured **25.33 ms**.
 All four final RGBA16F frames are byte-identical.
 An 80-frame run, excluding the first 20, measures **24.764 ms** with `0009` +
 `0010`. Patch `0011`'s complete-pair shadow-store coalescing reduces that to
 **23.340 ms**. Eight-wave register budgeting alone measures **20.135 ms**;
 both changes together measure **19.146 ms**. All 80 saved
 frames from each variant are byte-identical to the matching 80-frame control.
 Do not compare final images from different frame counts: the 20- and 80-frame
 captures differ even with the unchanged control.
 With `0012`'s signed native conversion, the same 80-frame configuration measures
 **16.234 ms**, with a byte-identical final frame. Signs stay in the half inputs;
 clamped finite values convert directly, and NaNs become signed infinity before
 conversion to obtain PTX's signed NaN bytes without post-conversion sign masks.
 The breakdown below uses that configuration. With the harness's 200 ms
 inter-frame idle gap removed, the same path measures **15.144 ms** and retains
 the identical final image; this steady-state run is not the idle-gap control.
 | kernel | ms/frame | kernel | ms/frame |
 |---|---|---|---|
-| dec0 | 7.71 | dec4 | 1.16 |
+| enc0 | 3.79 | dec4 | 0.52 |
-| enc0 | 6.49 | dec3 | 1.17 |
+| dec0 | 3.63 | dec3 | 0.60 |
-| dec1 | 3.66 | enc4 | 1.02 |
+| dec1 | 1.76 | enc4 | 0.45 |
-| enc1 | 3.48 | enc3 | 0.92 |
+| enc1 | 1.75 | enc3 | 0.49 |
-| dec2 | 2.18 | enc5 | 0.65 |
+| dec2 | 1.01 | enc5 | 0.41 |
-| enc2 | 1.74 | luma, exposure, reduce_sum | 0.05 |
+| enc2 | 0.82 | luma, exposure, reduce_sum | 0.05 |
-| hkpn output (translated → native) | 1.44 → 1.00 | **total** | **31.6 → 31.1** |
+| hkpn output | 0.96 | **total** | **16.23** |
 The actual PTX conversion probe checks all 65,536 half encodings in both packed
 positions, including overflow, infinities, signed zero and NaNs. Native and
 software conversions agree for both the plain and ReLU variants; the plain
 variant also matches an independent nearest-even E4M3 oracle. The gfx11 fallback
 compiles offline without reaching the gfx12-only intrinsic; gfx11 execution
 has not been checked. The older installed-runtime FTZ baseline (~31.5 ms)
 produces a different image and is not the arithmetic control for this gain.
 The packed-MMA GPU probe agrees with the previous implementation for all
 81,920 half values across paired, chained and unpaired operations, including
 subnormals, overflow, infinities and NaNs. The standalone
 `rrswin_fp8_unpaired_test.cpp` checks zero-padding with every source lane active,
 nonuniform bias and partial accumulator writes separated by a dependent read,
 against 512 independently calculated, exactly representable half values.
 A gfx12 register-permutation replacement for the LDS gathers was correct but
 slower (27.44 ms), and was reverted.
 A twelve-wave register budget preserved the final image but increased time to
 25.560 ms; it is not selected by `0011`.
 An isolated-cache fast-math run measured 18.920 ms versus 19.146 ms without
 relaxed math, but changed the image; relaxed math is not selected.
 Cyberpunk 2077 with `0009` + `0010` rendered a loaded save at 1280×720→3840×2160 with CUDA feature 13,
 preset E, DLSSD 310.7.0 and split-frame VRAM interop. Moving views showed coherent
 geometry and reflections; the overlay's average upscaler time was about 26.0 ms.
 This is a functional visual check, not RTX image-quality parity. Native-Wayland
 window shutdown reported a swapchain-creation error; the test prefix was stopped
 and the original game settings restored byte-for-byte.
 The eleven layers' cut ranges, from `rr_layer_spec.py <layer>` (file lines of the layer's corpus module, the region a native body replaces; everything outside stays in NVIDIA's PTX):
 | layer | corpus | cut lines | mma | phases | weight image |
 |---|---|---|---|---|---|
 | enc0 | 0013 | 3964-44524 | 536 | 22 | 54560 |
 | enc1 | 0014 | 13343-66941 | 1072 | 32 | 96468 |
 | enc2 | 0015 | 7669-33946 | 630 | 38 | 212244 |
 | enc3 | 0016 | 9812-41980 | 936 | 48 | 315732 |
 | enc4 | 0017 | 11883-52668 | 1386 | 62 | 533844 |
 | enc5 | 0018 | 11883-50752 | 1306 | 61 | 431124 |
 | dec0 | 0023 | 1393-46263 | 544 | 19 | 47872 |
 | dec1 | 0022 | 1486-38469 | 632 | 26 | 96788 |
 | dec2 | 0021 | 1535-36889 | 726 | 38 | 212756 |
 | dec3 | 0020 | 1658-46562 | 1112 | 48 | 316436 |
 | dec4 | 0019 | 1707-58076 | 1606 | 62 | 534804 |
 The denoiser's Swin block is the same block as presets L/M's: `cuda_dldn_engine_swin_enc0_kernel` computes position-only attention `S = bias + Q*Q^T` with no K projection, a 16x16 f16 bias table at 512 bytes per head, and one fused Q/V projection - the score `mma` at line 8406 takes both of its operands from the same accumulator registers, `%r1900`/`%r1920`, through their `cvt.rn.satfinite.e4m3x2.f16x2` at 8126/8129 and 8294/8297. That is `kernels/m/swin_block.h`'s block (`T16 S = bias; k32(S, a0, a0, a1, a1);`), not `kernels/k`'s learned-attention one. So the native denoiser layers can be built by instantiating that template - a prep that permutes the denoiser's own weight image into its slot layout, a parameter struct for the 144-byte block, and the surface stores through `tex_common.h` - rather than by writing a Swin block from scratch.
-Measured launch shape, from the replay manifest of a real harness frame: `cuda_dldn_engine_swin_enc0_kernel` and `..._dec0_kernel` are both launched with grid (161, 93, 1) and block (32, 1, 1) - one warp per block, 14973 blocks, one 8x8 token patch each (161 = (1280+8)/8, 93 = 736/8) - so there is no wave index in either; tid.y and tid.z are zero. The parameter block is 1032 bytes for enc0 and 152 for dec0, so it is not one shared struct: enc0 takes the guide textures at +160 through +512, dec0 the three at +88/+96/+104. dec0's dumped allocations are the 65 945 600-byte plane arena and a 47872-byte weight image (the spec tool's 47376 again plus an unpadded tail). The native override can set block.z and flatten the grid per kernel (`d4r_block_z`, `d4r_grid_x`), so a native replacement is free to use more waves per block than one.
+`rr_layer_spec.py`'s M and N come from the distinct A and B fragment counts of a phase. That is unambiguous only when the phase makes one k32 step: dec0's first phase has 96 mma over 6 distinct A and 32 distinct B fragments, which is 48x128 with K=64 (two k32 steps of 3 A x 16 B), while the tool prints it as 96x256 with K=64 because it attributes both steps' fragments to M and N. Check the arithmetic `M/16 * N/8 * K/32 == mma` before trusting a row; enc0's rows all satisfy it and its first phase is verified bit-exactly against the GPU.
 Measured launch shape, from the replay manifest of a real harness frame. Block z is the number of waves per block and varies by layer, which is why the decoders' weight addressing carries a per-`tid.z` stride; the grid shrinks as the U-Net deepens. The parameter block is 144 bytes for the interior layers, 152 for dec0 and 1032 for enc0, which takes the guide textures:
 | layer | grid | block | waves | kernarg |
 |---|---|---|---|---|
 | enc0 | 161x93 | 32x1x1 | 1 | 1032 |
 | enc1 | 81x47 | 32x1x1 | 1 | 144 |
 | enc3 | 21x13 | 32x1x4 | 4 | 144 |
 | dec0 | 161x93 | 32x1x1 | 1 | 152 |
 | dec1 | 81x47 | 32x1x2 | 2 | 144 |
 | dec2 | 41x24 | 32x1x4 | 4 | 144 |
 | dec4 | 11x6 | 32x1x8 | 8 | 144 |
 enc0 and dec0 read their input from CUDA textures rather than the plane arena, and write to CUDA surfaces; the other nine read the arena. The native override can set block z and flatten the grid per kernel (`d4r_block_z`, `d4r_grid_x`), so a native replacement is free to choose its own shape.
 `kernels/rr/rr_layer_spec.py` recovers a layer's structure from the corpus PTX alone: `rr_layer_spec.py enc0` prints the phase table (each maximal run of adjacent `mma`, with M/N/K derived from the fragment counts and the accumulator dataflow, the weight and bias byte offsets it reads, and the size of the epilogue between phases) and the stage classification; `--geometry` does the same for all eleven layers and prints each one's weight image extent. It is how enc0's layout above was established, and it is the reference the native layers are written against.
-The eleven Swin layers are 30.2 of those 31.6 ms. They are instruction-bound, not MMA-bound: the compiled `enc0` kernel is 83 373 instructions per thread for 536 FP8 WMMAs (0.6% of the stream), and the frame reaches about 6% of the card's FP8 WMMA peak. The stream is 13 k operand-gather shuffles (`ds_bpermute`/`v_perm`/`permlane`), 15 k packed 16-bit ops from the software `cvt.rn.satfinite.e4m3x2.f16x2` requantization, 9 k selects and 13 k ALU-dependency waits. `D4R_ZLUDA_IMPLICIT_MAX_BLOCK` (256 in the shipped config), `D4R_ZLUDA_SCHED_STRATEGY` and a 128-thread bound move the total by under 3%; a 128-thread bound additionally fails the launches of kernels that use larger blocks, and the whole-module wave64 switch produces no launches at all. Removing the gather and the codec — what a native layer does — is the remaining work.
+Before native FP8 conversion, enc0's translated shader contained 83,373 instructions for 536 FP8 WMMAs. Operand gathering and software requantization dominated that static stream; those counts are not measurements of the patch `0009` shader. Hardware conversion removes the software codec but leaves operand-layout conversion and the surrounding neural computation. A global 128-thread bound is invalid for the 256-thread layers; the captured E launches use 32 threads for enc0, enc1 and dec0, 64 for dec1, 128 for enc2/enc3/dec2/dec3, and 256 for enc4/enc5/dec4. The remaining network must be optimized and measured rather than treating the codec gain as meeting the target.
 ## RDNA4
 `kernels/common/wmma_layout.h` selects the WMMA operand and accumulator layout. gfx12 uses eight f16 values per lane in each K half; the accumulator rows are `i + 8·half`. The gfx11 path retains its original layout. `D4R_WMMA_LAYOUT=12` on gfx11 is a test shim: it exercises gfx12 indexing but executes gfx11 WMMA, so its replay output can be compared byte for byte with the existing gfx11 build.
 A native kernel can call that unit directly: declaring `__device__ float8 f(int2, int2, float8) __asm("llvm.amdgcn.wmma.f32.16x16x16.fp8.fp8.v8f32.v2i32");` compiles for gfx1201 and emits the instruction, which is what lets a denoiser layer read its weights straight out of NVIDIA's fragment order instead of gathering them. Verified by compiling a kernel that uses it.
 The gfx12 M variant can use native e4m3 FP8 WMMA with `D4R_NATIVE_FP8=1`. Its activations are requantized to e4m3 at operand load, and its prepared weights remain bytes. The matching ZLUDA and texture-tail variant uses `D4R_ZLUDA_WMMA_FP8_NATIVE=1` / `D4R_TEX_FP8=1`; the release bridge chooses the `-fp8` kernel folder. `NativeFp8` in d4r.ini controls that choice and defaults to on. The f16 widening path remains available.
 In that build the GEMM activations live in LDS as e4m3 bytes (`SWIN_A8` in `swin_common.h`). Every value is encoded once where it is produced, four at a time with `v_cvt_pk_fp8_f32`, and an operand load is one 8-byte LDS read; the earlier build kept f16 values and encoded them again at every operand load (up to 16 times per value in the MLP). The instruction rounds like NVIDIA's conversion but encodes overflow as `0x7f`, so values are clamped to ±448 first with the NaN-propagating packed minimum and maximum (`codes4`). It equals `half_to_e4m3` for every non-NaN f16 input on an RX 9070 XT; a NaN encodes as the NaN code `0xff`, where the reference keeps its sign (`0x7f` or `0xff`). The accuracy build of these layers rounds each k32 step with `v_fma_mix` (`K32_MIX`; `-DK32_NO_MIX` restores the convert-and-round form) and its stored codes still equal `swin_model.py`. The patch merge reads the stored output codes from a byte image. `-DSWIN_NO_A8` restores the previous operand path. Builds without `D4R_FP8_WMMA` (RDNA3, and RDNA4 with `NativeFp8` off) compile to the same code as before.
@@ -184,7 +267,7 @@ The release includes both fast and accuracy network sets. Full releases also inc
   ```
   It runs the prep and main kernels, saves every buffer after the first launch, and times 30 launches. The replay tool refuses captures with texture or surface objects; use the full harness for those kernels.
-3. **A per-layer oracle for kernels with texture and surface objects.** The replay capture above cannot describe those kernels (the replayers refuse them), but the bridge's other instrumentation can: with `D4R_CUDA_LAUNCH_STATS=1` and `D4R_CUDA_LAUNCH_DUMP_DIR=<dir>` (with `D4R_CUDA_LAUNCH_STATS_FILTER=<substring>` to keep only the kernels being worked on, checked before the synchronize) every `cuLaunchKernel` is followed by a synchronize and a raw dump of every allocation, surface object and texture object named in its packed argument buffer, one file per launch and argument (`launch-NNN-<kernel>-argNNN-{buffer,surface,texture}-<shape>-fmtN.bin`). Run the harness once with the translated kernel and once with the native one and compare that kernel's files byte for byte: the surface files are the layer's own output, and a mismatch localises to that layer rather than to the frame. `cuda_dldn_engine_swin_enc0_kernel` dumps its weight image at arg040, its activation buffer at arg048 and four surfaces (1280x736x4 and 640x368x4 RGBA16F) as its outputs. The dump is large - 3 frames cost 3.7 GB - so keep the frame count small.
+3. **A per-layer oracle for kernels with texture and surface objects.** The replay capture above cannot describe those kernels (the replayers refuse them), but the bridge's other instrumentation can: with `D4R_CUDA_LAUNCH_STATS=1` and `D4R_CUDA_LAUNCH_DUMP_DIR=<dir>` (with `D4R_CUDA_LAUNCH_STATS_FILTER=<substring>` to keep only the kernels being worked on, checked before the synchronize) every `cuLaunchKernel` is followed by a synchronize and a raw dump of every allocation, surface object and texture object named in its packed argument buffer, one file per launch and argument (`launch-NNN-<kernel>-argNNN-{buffer,surface,texture}-<shape>-fmtN.bin`). Run the harness once with the translated kernel and once with the native one and compare that kernel's files byte for byte: the surface files are the layer's own output, and a mismatch localises to that layer rather than to the frame. `cuda_dldn_engine_swin_enc0_kernel` dumps its weight image at arg040, its activation buffer at arg048 and four surfaces (1280x736x4 and 640x368x4 RGBA16F) as its outputs. The dump is large - 3 frames cost 3.7 GB - so keep the frame count small. A native denoiser layer is checked at a finer grain still: `D4R_RRSWIN_DEBUG=<stage>` builds NVIDIA's own body with a store of that stage's registers spliced in, and the native body built with `D4R_TEX_CFLAGS=-DD4R_RRSWIN_DEBUG_STAGE=<n>` stores its own; both land in the plane arena at 46,000,000, so the two are compared byte for byte with the frame out of the way. Verified on phase 4: both sides agree on all 524288 bytes of the region (md5 788f4274c7ff), two runs of one build reproduce it exactly, and the only arguments that differ between runs are the downstream textures the pipeline writes, not the layer's inputs.
 4. **Compare against a reference.** `kernels/tools/pwin_model.py` (K) and `swin_model.py` (M) are numpy models of the layers, written stage by stage against `ptxsim.py`, a small vectorised PTX interpreter that runs NVIDIA's kernels on the CPU. A bit-exact variant should match the model within f16 rounding; after any change, compare the new build's buffers with the previous build's.
 5. **Check the whole pipeline.** Replay captured frames through the harness with both kernel sets and compare the outputs with `kernels/tools/psnr.py`. `D4R_CUDA_KERNEL_PROFILE=1` makes the bridge time every kernel; `kernels/tools/kprof.py` summarises the log per frame.
@@ -13,7 +13,7 @@ Built from source by d4r (GPL-3.0-only, d4r/licenses/d4r-LICENSE.txt):
 Third-party software, unmodified or patched by d4r:
  d4r/zluda/libcuda.so   ZLUDA (Apache License 2.0 or MIT), https://github.com/vosen/ZLUDA
-                         commit @ZLUDA_COMMIT@ with d4r/source/patches/zluda/0002 to 0008 applied.
+                         commit @ZLUDA_COMMIT@ with d4r/source/patches/zluda/0002 to 0012 applied.
                         It contains LLVM (Apache License 2.0 with LLVM exceptions).
                         Licenses: d4r/licenses/ZLUDA-LICENSE-*.txt, d4r/licenses/LLVM-LICENSE.txt
  d3d12.dll, d3d12core.dll
@@ -0,0 +1,92 @@
 diff --git a/llvm_zluda/src/lib.cpp b/llvm_zluda/src/lib.cpp
 --- a/llvm_zluda/src/lib.cpp
 +++ b/llvm_zluda/src/lib.cpp
@@ -485,6 +485,8 @@
     bool AnyFp8 = false;
     for (const Site &S : Sites)
         AnyFp8 |= S.Fp8;
 +    if (PackedFp8 && AnyFp8)
 +        F.addFnAttr("amdgpu-waves-per-eu", "8");
     Type *F32Ty = Type::getFloatTy(F.getContext());
     Value *Scale = ConstantFP::get(F32Ty, AnyFp8 ? 0x1p-16 : 1.0);
     Value *Inverse = ConstantFP::get(F32Ty, AnyFp8 ? 0x1p16 : 1.0);
@@ -758,6 +760,70 @@
         Call->eraseFromParent();
     }
 +    // A complete pair overwrite needs one conversion. Reads of either
 +    // NVIDIA register, accesses to its rewritten shadow, and block boundaries
 +    // prevent coalescing. Scan after MMA rewriting to include hoisted reads.
 +    DenseMap<AllocaInst *, int> ShadowOwners;
 +    for (size_t Index = 0; Index < Shadows.size(); ++Index)
 +        if (Shadows[Index] != nullptr)
 +            ShadowOwners[Shadows[Index]] = static_cast<int>(Index);
 +    SmallVector<std::pair<StoreInst *, StoreInst *>, 64> StorePairs;
 +    DenseMap<int, StoreInst *> PendingStores;
 +    for (BasicBlock &Block : F)
 +    {
 +        PendingStores.clear();
 +        for (Instruction &I : Block)
 +        {
 +            auto *Load = dyn_cast<LoadInst>(&I);
 +            auto *Store = dyn_cast<StoreInst>(&I);
 +            if (Load == nullptr && Store == nullptr)
 +                continue;
 +            Value *Pointer = Load != nullptr ? Load->getPointerOperand() : Store->getPointerOperand();
 +            auto *Register = dyn_cast<AllocaInst>(Pointer);
 +            auto Shadow = ShadowOwners.find(Register);
 +            if (Shadow != ShadowOwners.end())
 +            {
 +                PendingStores.erase(Shadow->second);
 +                continue;
 +            }
 +            auto Found = Owner.find(Register);
 +            if (Found == Owner.end() || Found->second < 0 || Shadows[Found->second] == nullptr)
 +                continue;
 +            int Index = Found->second;
 +            if (Load != nullptr)
 +            {
 +                PendingStores.erase(Index);
 +                continue;
 +            }
 +            auto Previous = PendingStores.find(Index);
 +            if (Previous != PendingStores.end() && Previous->second->getPointerOperand() != Pointer)
 +            {
 +                StorePairs.push_back({Previous->second, Store});
 +                PendingStores.erase(Previous);
 +            }
 +            else
 +                PendingStores[Index] = Store;
 +        }
 +    }
 +    DenseSet<StoreInst *> CoalescedStores;
 +    for (auto [First, Second] : StorePairs)
 +    {
 +        if (ResultStores.contains(First) || ResultStores.contains(Second))
 +            continue;
 +        auto *Register = cast<AllocaInst>(First->getPointerOperand());
 +        const RegisterPair &Pair = Shadowed[Owner[Register]];
 +        const bool FirstHigh = Register == Pair.High;
 +        IRBuilder<> Builder(Second);
 +        Value *Nv = UndefValue::get(Pair32);
 +        Nv = Builder.CreateInsertElement(Nv, First->getValueOperand(), Builder.getInt8(FirstHigh));
 +        Nv = Builder.CreateInsertElement(Nv, Second->getValueOperand(), Builder.getInt8(!FirstHigh));
 +        Builder.CreateStore(fromNv(Builder, Nv), shadowFor(Pair));
 +        CoalescedStores.insert(First);
 +        CoalescedStores.insert(Second);
 +        First->eraseFromParent();
 +        Second->eraseFromParent();
 +    }
 +
     // Every other access now goes through the shadow.
     for (LoadInst *Load : Loads)
     {
@@ -770,6 +836,8 @@
     }
     for (StoreInst *Store : Stores)
     {
 +        if (CoalescedStores.contains(Store))
 +            continue;
         if (!ResultStores.contains(Store))
         {
             auto *Register = cast<AllocaInst>(Store->getPointerOperand());
@@ -0,0 +1,34 @@
 diff --git a/ptx/lib/zluda_ptx_impl.cpp b/ptx/lib/zluda_ptx_impl.cpp
 --- a/ptx/lib/zluda_ptx_impl.cpp
 +++ b/ptx/lib/zluda_ptx_impl.cpp
@@ -877,21 +877,18 @@
         typedef _Float16 half_vector __attribute__((ext_vector_type(2)));
         typedef uint16_t magnitude_vector __attribute__((ext_vector_type(2)));
         const uint32_t bits = std::bit_cast<uint32_t>(input);
 -        // Positive half encodings are ordered by magnitude. Clamp both lanes
 -        // with one packed integer min, then retain NaNs as positive quiet NaNs.
 -        // Restore original signs after conversion, including negative zero.
         const magnitude_vector magnitude = std::bit_cast<magnitude_vector>(bits & 0x7fff7fffu);
 -        const magnitude_vector saturated =
 +        // Clamp finite overflow and infinities explicitly. Map NaNs to signed
 +        // infinity: native infinity encodes as 0x7f/0xff, whereas native NaN is
 +        // always 0xff. This preserves the sign without post-conversion masks.
 +        const magnitude_vector bounded =
             __builtin_elementwise_min(magnitude, magnitude_vector{0x5f00u, 0x5f00u});
         const magnitude_vector kept{
 -            magnitude[0] > 0x7c00u ? uint16_t(0x7e00u) : saturated[0],
 -            magnitude[1] > 0x7c00u ? uint16_t(0x7e00u) : saturated[1]};
 -        const half_vector halves = std::bit_cast<half_vector>(kept);
 -        const float x = float(halves[0]), y = float(halves[1]);
 -        const uint32_t encoded = d4r_cvt_pk_fp8_f32(x, y, 0, false);
 -        // AMD canonicalizes NaNs to 0xff. PTX retains each input's sign.
 -        const uint16_t packed = uint16_t((encoded & 0x7f7fu) |
 -                                        ((bits >> 8) & 0x80u) | ((bits >> 16) & 0x8000u));
 +            magnitude[0] > 0x7c00u ? uint16_t(0x7c00u) : bounded[0],
 +            magnitude[1] > 0x7c00u ? uint16_t(0x7c00u) : bounded[1]};
 +        const uint32_t signed_bits = std::bit_cast<uint32_t>(kept) | (bits & 0x80008000u);
 +        const half_vector halves = std::bit_cast<half_vector>(signed_bits);
 +        const uint16_t packed = uint16_t(d4r_cvt_pk_fp8_f32(float(halves[0]), float(halves[1]), 0, false));
         return std::bit_cast<Fp8x2>(packed);
     }
@@ -5,7 +5,7 @@
 set -euo pipefail
 ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)"
 : "${D4R_RELEASE_BUILD_ROOT:?set the build output directory}"
-: "${D4R_ZLUDA_SRC:?set a build-local ZLUDA checkout with patches 0002-0008 applied}"
+: "${D4R_ZLUDA_SRC:?set a build-local ZLUDA checkout with patches 0002-0012 applied}"
 : "${D4R_VKD3D_SRC:?set a build-local vkd3d-proton checkout with patches 0001-0003 applied}"
 : "${D4R_ROCM_DIR:?set ROCm with clang and device libraries}"
 BUILD="$(realpath -m "$D4R_RELEASE_BUILD_ROOT")"