Add CUDA NGX Ray Reconstruction support

Run the signed CUDA-capable denoiser through feature 13 with independent D/E presets, pre-Init discovery, per-frame guide and camera snapshots, normalized integer guide staging, paired alpha output, and recorded-resource retirement. Add the required ZLUDA/LLVM patch, behavioral configuration tests, GPU harness coverage, optional denoiser installation and packaging, and tested support limitations. Preserve upstream attribution and distinguish LLVM licensing from the ZLUDA portions of the patch.
2026-10-12 12:47:53 +02:00
1 parent 503061d568
commit 7117c38674
@@ -1,5 +1,6 @@
 d4r
 Copyright 2026 countervolts and the d4r contributors
 Modified in this fork for CUDA Ray Reconstruction support and validation (2026).
 This product is licensed under the GNU General Public License, version 3 only
 (GPL-3.0-only; see LICENSE), except for the files
@@ -7,7 +8,9 @@ listed below.
 patches/zluda/*.patch
    Changes to ZLUDA (https://github.com/vosen/ZLUDA). ZLUDA is dual-licensed under the Apache
-    License, Version 2.0 and the MIT license; these patches are offered under the same terms.
+    License, Version 2.0 and the MIT license; the ZLUDA portions of these patches are offered
    under the same terms. Changes to LLVM within these patches retain the Apache License,
    Version 2.0 with LLVM exceptions.
 patches/vkd3d-proton/*.patch
    Changes to vkd3d-proton (https://github.com/HansKristian-Work/vkd3d-proton), which is licensed
@@ -19,12 +22,12 @@ LLVM under Apache-2.0 with LLVM exceptions), the patched vkd3d-proton (LGPL-2.1)
 (https://github.com/optiscaler/OptiScaler, GPL-3.0, unmodified). Each zip carries their license
 texts in d4r/licenses and says where every file comes from in d4r/source/SOURCES.txt.
 Release zips include install.sh to download NVIDIA's nvngx_dlss.dll and _nvngx.dll separately.
-Optional builds may bundle those DLLs. Release zips also include DLSS kernels compiled from NVIDIA's
+Optional builds may bundle those DLLs and nvngx_dlssd.dll. Release zips also include DLSS kernels compiled from NVIDIA's
 PTX. Those are NVIDIA's property, are not covered by this repository's license, and are included
 by whoever publishes the zip. Downloaded DLLs remain governed by NVIDIA's license terms.
-This repository contains no NVIDIA software. The NVIDIA DLSS feature library (nvngx_dlss.dll) and
+This repository contains no NVIDIA software. The NVIDIA DLSS feature libraries (nvngx_dlss.dll and
-the NGX core (_nvngx.dll) are NVIDIA's; you supply your own copies, and their use is governed by
+nvngx_dlssd.dll) and the NGX core (_nvngx.dll) are NVIDIA's; you supply your own copies, and their use is governed by
 NVIDIA's license terms. At build time, kernels/tools/extract_dlss_ptx.py reads PTX from your copy
 of nvngx_dlss.dll into kernels/extracted/, which is not part of the repository and must not be
 redistributed. The same holds for the model files the native engine's tools write (engine/README.md):
@@ -1,14 +1,17 @@
 <!-- Modified in this fork for CUDA Ray Reconstruction support and validation (2026). -->
 # d4r (dlss 4 radeon)
 want to run xess (xmx) on radeon? check out [x4r](https://github.com/countervolts/x4r)!
-d4r runs NVIDIA's official DLSS Super Resolution library (`nvngx_dlss.dll`) in Windows games on AMD Radeon GPUs under Linux and Proton. The game asks for DLSS as usual; the DLSS network runs on the AMD GPU through [ZLUDA](https://github.com/vosen/ZLUDA) (CUDA on ROCm/HIP), with the heaviest DLSS kernels replaced by hand-written RDNA3 and RDNA4 code.
+d4r runs NVIDIA's official DLSS Super Resolution (`nvngx_dlss.dll`) and experimental Ray Reconstruction (`nvngx_dlssd.dll`) in Windows games on AMD Radeon GPUs under Linux and Proton. The game asks for DLSS as usual; the DLSS network runs on the AMD GPU through [ZLUDA](https://github.com/vosen/ZLUDA) (CUDA on ROCm/HIP), with the heaviest Super Resolution kernels replaced by hand-written RDNA3 and RDNA4 code.
 **Supported DLSS models:** DLSS 3 CNN (E), DLSS 4 transformer (K, default), and DLSS 4.5 transformer (M, plus experimental L). DLSS 5 is PURPOSELY not supported.
 Preset L is NVIDIA's DLSS 4.5 model for Ultra Performance, especially at 4K. Set `[DLSS] Model = L` in `d4r.ini` and select Ultra Performance in game or OptiScaler; selecting L alone does not change the input resolution. L shares M's native Swin layers and has its own texture-kernel build. Its unfolded input/output arithmetic remains translated, so performance and image quality need game testing; the M benchmarks below do not apply to L.
-**Proof of concept:** d4r shows that DLSS can run on an AMD GPU, but it is not really that practical for everyday use yet. It has been tested on one GPU in a handful of games and depends on unreleased patches to ZLUDA and vkd3d-proton.
+Ray Reconstruction presets D and E use the signed CUDA-capable 3.10.7 denoiser and the translated NVIDIA network, independently of the Super Resolution model. Temporal reconstruction and separate alpha output have been verified on Radeon RX 9070 XT (gfx1201) with native FP8 WMMA enabled. Newer DLSS 4.5 RR preset F is not implemented or validated. The scalar FP8 path produced nonfinite temporal RR output and is not validated; other GPU targets and image-quality parity with RTX hardware remain unverified. See [building](docs/building.md) for the required runtime, denoiser placement, and FP8 settings.
 **Proof of concept:** d4r shows that DLSS can run on an AMD GPU, but it is not really that practical for everyday use yet. Super Resolution game testing primarily uses Radeon RX 7700 XT; Ray Reconstruction validation additionally uses RX 9070 XT. It depends on unreleased patches to ZLUDA and vkd3d-proton.
 See [supported games](SUPPORTED_GAMES.md) for the tested games and DLSS models.
@@ -1,6 +1,7 @@
 <!-- Modified in this fork for CUDA Ray Reconstruction support and validation (2026). -->
 # Supported games
-The following DirectX 12 games run DLSS Super Resolution through d4r on the tested Radeon RX 7700 XT Linux setup with GE-Proton 11-3. All three documented DLSS models run in each game; image-quality issues are noted in the table.
+The following DirectX 12 games run DLSS Super Resolution through d4r on the tested Radeon RX 7700 XT Linux setup with GE-Proton 11-3. All three Super Resolution models in the table run in each game; image-quality issues are noted below. Experimental Ray Reconstruction has a separate RX 9070 XT validation described under Cyberpunk 2077.
 | Game | E: DLSS 3 CNN | K: DLSS 4 | M: DLSS 4.5 |
 |---|---|---|---|
@@ -22,16 +23,18 @@ The following DirectX 12 games run DLSS Super Resolution through d4r on the test
 General notes:
 - Select DLSS in the game's graphics settings where it is offered. OptiScaler can also feed DLSS from a game's FSR or XeSS inputs, but those are a fallback for games that hide DLSS and can ghost or show a black screen.
- d4r supports DLSS Super Resolution only. Keep DLSS Ray Reconstruction off in games that offer it (such as Cyberpunk 2077 and Alan Wake 2); Control requests it whenever ray tracing is on and crashes.
+- Experimental Ray Reconstruction requires a CUDA-capable 3.10.7 denoiser, matching lifetime-aware d3d12.dll/d3d12core.dll, and native FP8 WMMA for the validated gfx1201 path. Cyberpunk 2077 is verified below; RR in Alan Wake 2 and Control remains unverified, so keep it off for their documented setups.
 - Saving settings from the OptiScaler overlay ("Save INI") rewrites `OptiScaler.ini` and resets d4r's values to `auto`, including frame generation (`[FrameGen] Enabled=false`, `FGInput=nofg`, `FGOutput=nofg`) and `[Inputs] EnableDlssInputs=true`. Restore them afterwards, or keep a copy of the original file.
 Game-specific setup:
 - **The Last of Us Part I:** the game's Streamline integration crashes on boot with OptiScaler. Rename `sl.common.dll` in the game folder to `sl.common.dll.bak` and select FSR 3.1 in the game's graphics settings; OptiScaler passes it to DLSS through d4r.
- **Control Ultimate Edition:** turn ray tracing off. With ray tracing on, the game requests DLSS Ray Reconstruction, which d4r does not support, and crashes when loading into gameplay.
+- **Control Ultimate Edition:** keep ray tracing off for the documented setup. Older builds crashed when loading gameplay with ray tracing enabled; this game's integration has not been revalidated against the new RR route.
 - **Cyberpunk 2077:** GE-Proton 11-3 loads Wine's own ICU libraries instead of the game's, and the game aborts at launch (`unimplemented function icuuc.dll.u_setMemoryFunctions_65`). Add `WINEDLLOVERRIDES="icuuc,icuin=n,b"` to the launch options (fixed in GE-Proton 11-4).
  - **Ray Reconstruction:** Cyberpunk 2077 2.31 has been verified in loaded local saves on RX 9070 XT (gfx1201), GE-Proton 11-3, at 1920×1080 and 3840×2160 Ultra Performance with path tracing and RR enabled. The actual backend is `DLSSD 310.7.0`, CUDA feature 13, preset D; observed gameplay snapshots showed about 53 FPS at 1080p and 20 FPS at 4K, not comparative benchmarks.
  - Enable `[Kernels] NativeFp8 = true` and supply the CUDA-capable 3.10.7 `nvngx_dlssd.dll` beside the shim. Update game-local d3d12.dll **and** d3d12core.dll too: local overrides shadow a patched Proton prefix. RR preset E has separate temporal/alpha harness validation, not a Cyberpunk performance result. Newer DLSS 4.5 RR preset F is not implemented or validated. Scalar FP8 temporal RR is not validated after producing nonfinite output; RTX image-quality parity remains unverified.
 - **Alan Wake 2:** with the game's own DLSS inputs, every model shows heavy ghosting. Set `AutoExposure=true` under `[InitFlags]` in `OptiScaler.ini` (or enable auto exposure in the OptiScaler overlay), or select FSR 3.1 in the game instead so OptiScaler passes FSR inputs to DLSS.
@@ -1,3 +1,4 @@
 ; Modified in this fork for CUDA Ray Reconstruction support and validation (2026).
 ; d4r configuration: official NVIDIA DLSS on AMD (ZLUDA + Wine CUDA bridge + NGX shim)
 ;
 ; Read by scripts/run_game_d4r_optiscaler_proton.sh at every game launch. The first launch copies
@@ -32,6 +33,18 @@ GameExe =
 ; Aliases: CNN = E, DLSS4 = K, DLSS4.5 = M. A/B/C/D/G/J use K in this DLL version.
 ; auto = use the game/OptiScaler choice. Benchmarks: docs/performance.md.
 Model = K
 [RayReconstruction]
 ; Runs NVIDIA's nvngx_dlssd.dll denoiser, separate from Super Resolution.
 ; auto enables the CUDA route when the NGX core reports the denoiser initialized and available.
 ; false disables it. Select Ray Reconstruction and ray tracing in the game.
 ; RR D/E temporal validation on gfx1201 requires native FP8 WMMA ([Kernels] NativeFp8 = true).
 ; The scalar FP8 RR path is not validated.
 Enable = auto
 ; auto preserves the game's choice; default = NVIDIA's default, D/E = transformer models.
 ; These presets are independent of [DLSS] Model. E requires a DLL that contains that model.
 Model = auto
 [Latency]
 ; How many frames old the upscaled image the game presents may be.
 ;   0 = same frame (no added latency, no trailing artefacts; needs VkD3DDir below)
@@ -126,7 +139,7 @@ RuntimeDir = auto
 ; ZLUDA build (directory with libcuda.so / libnvcuda.so)
 ZludaDir = auto
 RocmDir = auto
-; d4r-patched vkd3d-proton (d3d12.dll + d3d12core.dll); required for FrameAge = 0
+; Matching lifetime-aware d4r vkd3d-proton (both d3d12.dll + d3d12core.dll); required for all shim paths.
 VkD3DDir = auto
 ; JIT cache directory (XDG_CACHE_HOME): auto = the default ~/.cache; per-zluda = one per ZLUDA build,
 ; ~/.cache/d4r-perf-game-jit-<ZludaDir name> (a new ZLUDA build compiles every kernel once on first use)
@@ -1,3 +1,4 @@
 <!-- Modified in this fork for CUDA Ray Reconstruction support and validation (2026). -->
 # Architecture
 d4r keeps NVIDIA's DLSS library unmodified and gives it what it expects: an NGX D3D12 or Vulkan entry point in the game process, a CUDA driver underneath, and CUDA kernels that run. Everything between those points is translation.
@@ -11,7 +12,7 @@ host staging and command-buffer event synchronization; it does not use vkd3d-pro
 |---|---|---|
 | OptiScaler (in GE-Proton) | external | Hooks the game's upscaler call and forwards DLSS requests to an NGX DLL (`NvngxPath`) |
 | NGX shim `d4r_nvngx.dll` | `tools/d4r_nvngx_shim.cpp` | Implements the NGX D3D12 API. Moves inputs and output between the game's D3D12 resources and CUDA, and drives the official NGX core through its CUDA API |
-| NGX core + `nvngx_dlss.dll` | NVIDIA, supplied by you | The real DLSS: parameter handling, network selection, the CUDA kernels |
+| NGX core + `nvngx_dlss.dll` and optional `nvngx_dlssd.dll` | NVIDIA, supplied by you | The real DLSS: parameter handling, Super Resolution or Ray Reconstruction network selection, and CUDA kernels |
 | CUDA bridge `nvcuda.dll` | `tools/wine_nvcuda_bridge.c` | A Wine builtin whose Windows exports forward to ZLUDA's Linux `libcuda.so`, plus d4r helpers (Vulkan memory import, asynchronous 2D copies, GPU-side waits, output redirect, linear textures) |
 | ZLUDA | `patches/zluda` | CUDA driver API on HIP; compiles NVIDIA's PTX to AMDGPU code; serves native kernels in place of PTX kernels |
 | Native kernels | `kernels/` | RDNA3 implementations of the DLSS network layers and parts of texture kernels |
@@ -37,6 +38,16 @@ The watchdog reports delayed output but never signals an unfinished producer as
 The shim requires the matching lifetime-aware d3d12.dll and d3d12core.dll, even for host staging. It rejects older runtimes rather than recording references whose lifetime it cannot protect. FrameAge > 0 shows the newest finished result; without VRAM interop, inputs and output are staged through host memory.
 ### Ray Reconstruction path
 Feature 13 uses the NGX core's CUDA lifecycle, not the denoiser's native D3D12 backend. A CUDA-capable `nvngx_dlssd.dll` must sit beside the calling shim; the installer stages it there as well as in the feature directory. Availability follows the core's denoiser availability and initialization result. `[RayReconstruction] Enable = auto` does not manufacture support; `Model` selects its own default/D/E preset independently of the Super Resolution model.
 Feature requirements may be queried before NGX initialization. The shim loads the providers without initializing NGX with a provisional application identity; when the core declines feature-13 discovery, it asks the CUDA-capable denoiser's D3D12 requirements export for the adapter's eligibility. This discovery-only call does not select the native D3D12 rendering backend. Initialized capability checks and feature creation still require the core's successful denoiser initialization.
 The shim snapshots denoiser guide bindings, scalar settings, camera matrices, and subrect origins per submitted frame. Normals and roughness use the SDK's `GBuffer.Normals` and `GBuffer.Roughness` resource keys; diffuse and specular albedo bindings are translated to the CUDA denoiser's `DLSS.Input.*Albedo` keys. Guide readback and CUDA upload follow the main input marker, and the backing resources remain held until that frame retires. Uploads use per-slot host staging because direct ROCm image copies from Vulkan-mapped memory truncated guide rows on gfx1201. CUDA parameters point at stable texture-handle storage, not D3D12 resource pointers. Normals preserve the caller's component values rather than being renormalized during format conversion.
 Separate alpha output follows the same produced-frame slot as RGB, including its output resource and subrect. Requesting separate alpha cuts over through the existing drained host-staging path instead of publishing mismatched VRAM RGB and alpha results.
 ## Why each piece exists
 - **CUDA NGX instead of D3D12 NGX.** NVIDIA's D3D12 path needs NVIDIA's driver. The CUDA path only needs a CUDA driver API, which ZLUDA provides.
@@ -72,7 +83,7 @@ The release zip (`scripts/package_release.sh`) is unpacked into the folder that
 | `d4r\zluda\libcuda.so` | ZLUDA |
 | `d4r\kernels\<gfx target>\` | native kernels and their manifest `d4r-kernels.txt` |
 | `d4r\d4r.ini` | this game's settings |
-| `d4r\nvngx_dlss.dll`, `d4r\ngx\_nvngx.dll` | NVIDIA's files, downloaded by `install.sh` (`D4R_BUNDLE_NVIDIA=installer`, the default), bundled (`=1`) or added by the user (`=0`) |
+| `d4r\nvngx_dlss.dll`, `d4r\ngx\_nvngx.dll`, optional `d4r\nvngx_dlssd.dll` | NVIDIA's files, downloaded by `install.sh` (`D4R_BUNDLE_NVIDIA=installer`, the default), bundled (`=1`) or added by the user (`=0`); the optional denoiser is bundled only through `D4R_BUNDLE_DLSSD` or added by the user |
 **NGX routing.** OptiScaler tries `_nvngx.dll` before `nvngx.dll`, with system fallback for each name. `OptiDllPath=d4r` alone can therefore select Proton's system `_nvngx.dll` before reaching the shim. A file-valued `NvngxPath` takes priority on the first probe, regardless of the probe's name. It must select `d4r\nvngx.dll`, never `d4r\ngx\_nvngx.dll`: the latter is NVIDIA's core, which the shim loads internally. See OptiScaler 0.9.4's [NGX loader](https://github.com/optiscaler/OptiScaler/blob/v0.9.4/OptiScaler/proxies/NVNGX_Proxy.h) and [override handling](https://github.com/optiscaler/OptiScaler/blob/v0.9.4/OptiScaler/Util.cpp).
@@ -1,3 +1,4 @@
 <!-- Modified in this fork for CUDA Ray Reconstruction support and validation (2026). -->
 # Building d4r
 These steps produce the pieces the launcher needs: a patched ZLUDA, a patched vkd3d-proton, the NGX shim and CUDA bridge, a staged runtime directory with your NVIDIA files, and the native kernels. The versions below are the ones that were tested; others may need changes.
@@ -45,6 +46,7 @@ What the patches add:
 - **0006**: `m16n8k8` f16 MMAs (DLSS 3 CNN, presets E/F) on RDNA3 WMMA like the k16 ones (`D4R_ZLUDA_WMMA_K8=0` disables it); weight-image slots for native prep kernels (`d4r_prep_key_at`, `d4r_prep_key_slots`); and wave64 compilation for offline texture-kernel builds (`D4R_ZLUDA_WAVE64=1`, with wave64 builds of the helper bitcode). `scripts/build_zluda_ptx_helpers.sh` now writes the `_w64` helper variants too.
 - **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.
 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.
@@ -54,7 +56,9 @@ When linking on a system with ROCm libraries outside the default search path, in
 scripts/build_vkd3d_proton_d4r.sh ~/.cache/d4r-vkd3d-d4r
 ```
-Patches `0001` and `0002` provide same-frame command-list splitting and resource-lifetime tracking; both are required by the shim. Patch `0003` adds experimental exportable linear textures, enabled only by `VKD3D_D4R_LINEAR_TEXTURES=1`. Point `D4R_VKD3D_DIR` (or `VkD3DDir` in d4r.ini) at the output and update `d3d12.dll` and `d3d12core.dll` together.
+Patches `0001` and `0002` provide same-frame command-list splitting and resource-lifetime tracking; both are required by the shim, which rejects feature creation without the lifetime interface, even when CPU staging or previous-frame output is selected. Patch `0003` adds experimental exportable linear textures, enabled only by `VKD3D_D4R_LINEAR_TEXTURES=1`. Point `D4R_VKD3D_DIR` (or `VkD3DDir` in d4r.ini) at the output and update `d3d12.dll` and `d3d12core.dll` together.
 Game-local `d3d12.dll` and `d3d12core.dll` override the prefix's copies. Update both if the game already contains an older d4r runtime; changing only `D4R_VKD3D_DIR` does not replace those local overrides.
 ## 3. Shim, bridge and NVAPI identity
@@ -69,9 +73,13 @@ The launcher builds the small NVAPI identity bridge (`scripts/build_d4r_nvapi_id
 ```sh
 scripts/install_d4r_runtime.sh /path/to/_nvngx.dll /path/to/nvngx_dlss.dll
 # With the CUDA-capable Ray Reconstruction feature library:
 scripts/install_d4r_runtime.sh /path/to/_nvngx.dll /path/to/nvngx_dlss.dll /path/to/nvngx_dlssd.dll
 ```
-This stages the shim, the bridge, the NGX core and the DLSS feature library in `D4R_RUNTIME_DIR` (default `~/.local/share/d4r-dlss`, or `RuntimeDir` in d4r.ini). The NVIDIA files come from your own sources, for example an NVIDIA driver package (NGX core) and a game or the DLSS SDK (feature library).
+This stages the shim, the bridge, the NGX core and the DLSS feature libraries in `D4R_RUNTIME_DIR` (default `~/.local/share/d4r-dlss`, or `RuntimeDir` in d4r.ini). The NVIDIA files come from your own sources, for example an NVIDIA driver package (NGX core) and a game or the DLSS SDK (feature libraries). Ray Reconstruction requires a CUDA-capable `nvngx_dlssd.dll`, such as 310.7; the older D3D12-only denoiser is not a replacement. The installer also places the denoiser beside the calling shim, where the NGX core discovers it.
 RR D/E temporal evaluation and separate alpha output have been verified on gfx1201 with native FP8 WMMA (`[Kernels] NativeFp8 = true`). For developer launches without an INI, set `D4R_ZLUDA_WMMA=1`, `D4R_ZLUDA_WMMA_FP8=1` and `D4R_ZLUDA_WMMA_FP8_NATIVE=1`. The scalar FP8 fallback produced nonfinite temporal RR output and is not a validated RR path. Other GPU targets and image-quality parity with RTX hardware remain unverified.
 ## 5. Native kernels
@@ -110,7 +118,7 @@ The script:
 - builds the shim, the bridge and both fast and accuracy network-layer kernels for gfx1100–gfx1103 and gfx1200–gfx1201 by default (`D4R_GPU_ARCHS` can select fewer targets); each gfx12 target also gets an `<arch>-fp8` variant;
 - writes the kernel manifest from the DLLs you list (it records hashes of their PTX, nothing else);
 - stages OptiScaler as `dxgi.dll` with the settings in `packaging/optiscaler.settings`, applied by `scripts/configure_optiscaler.py`; both full and clean ZIPs set `NvngxPath=d4r\nvngx.dll` explicitly, retaining `OptiDllPath=d4r` for other libraries;
- includes `install.sh` instead of NVIDIA's two DLLs by default, retaining any supplied texture kernels (built from NVIDIA's PTX by `kernels/build.sh tex`); builds accuracy texture sets with `D4R_ZLUDA_EMIT` unless marked prebuilt sets are supplied in `D4R_BUNDLE_TEX/accuracy/<target>`;
+- includes `install.sh` instead of NVIDIA's two DLLs by default, retaining any supplied texture kernels (built from NVIDIA's PTX by `kernels/build.sh tex`); builds accuracy texture sets with `D4R_ZLUDA_EMIT` unless marked prebuilt sets are supplied in `D4R_BUNDLE_TEX/accuracy/<target>`; the full variant adds the optional `nvngx_dlssd.dll` supplied through `D4R_BUNDLE_DLSSD`;
 - with `D4R_BUNDLE_ENGINE`, adds the native engine's shaders as `d4r/engine/k`; the model's weights and tables are never packaged;
 - zips the result together with the ZLUDA and vkd3d-proton builds, the ROCm runtime (as `d4r/rocm`), `packaging/d4r.ini`, the licenses and the patches.
@@ -141,7 +149,7 @@ Mount the repository at `/work`, ROCm's compiler/headers/device libraries at `/o
 Run `bash scripts/build_release_glibc241.sh` inside the container. The script checks that the build system actually reports `glibc 2.41`, rejects external source-cache checkouts, builds with the locked Cargo dependencies offline, and runs the configuration tests before packaging. The ZIP and checksum go in `$D4R_RELEASE_BUILD_ROOT/dist/`; logs and `toolchain.txt` remain in the build directory. AMD's bundled ROCm runtime, OptiScaler and NVIDIA's libraries remain supplied binaries; the Linux ABI check covers the bundled ROCm libraries too.
-`D4R_BUNDLE_TEX` accepts the older flat directory for gfx1101, or a directory with per-target subdirectories (`gfx1100/`–`gfx1103/`, `gfx1200/`, `gfx1201/`, plus `gfx1200-fp8/` and `gfx1201-fp8/`). Build each texture set with `D4R_GPU_ARCH` and `D4R_ZLUDA_EMIT`; the `-fp8` sets also need `D4R_NATIVE_FP8=1`. `d4r_emit` targets those GPUs offline. Missing texture kernels fall back to ZLUDA on that target; the network-layer kernels are still included. Only the RX 7700 XT has been tested on real hardware.
+`D4R_BUNDLE_TEX` accepts the older flat directory for gfx1101, or a directory with per-target subdirectories (`gfx1100/`–`gfx1103/`, `gfx1200/`, `gfx1201/`, plus `gfx1200-fp8/` and `gfx1201-fp8/`). Build each texture set with `D4R_GPU_ARCH` and `D4R_ZLUDA_EMIT`; the `-fp8` sets also need `D4R_NATIVE_FP8=1`. `d4r_emit` targets those GPUs offline. Missing texture kernels fall back to ZLUDA on that target; the network-layer kernels are still included. The published Super Resolution benchmarks use RX 7700 XT; separate native-FP8 Ray Reconstruction D/E temporal validation and Cyberpunk gameplay use RX 9070 XT (gfx1201). This does not validate every packaged native-kernel target.
 ## Checks
@@ -1,14 +1,17 @@
 Modified in this fork for CUDA Ray Reconstruction support and validation (2026).
 d4r @VERSION@: NVIDIA DLSS on AMD Radeon RDNA3 and RDNA4 under Linux and Proton
 =====================================================================
-d4r runs NVIDIA's own DLSS Super Resolution (DLSS 4 by default) on AMD Radeon RDNA3 and RDNA4 GPUs in
+d4r runs NVIDIA's own DLSS Super Resolution (DLSS 4 by default), plus experimental Ray Reconstruction,
-DirectX 12 games under Proton. OptiScaler (included) catches the game's DLSS calls and hands them to
+on AMD Radeon GPUs in DirectX 12 games under Proton. OptiScaler (included) catches the game's DLSS calls
-d4r. d4r then runs NVIDIA's DLSS on the Radeon through ZLUDA, with the heaviest parts replaced by
+and hands them to d4r. d4r runs NVIDIA's DLSS through ZLUDA; the heaviest Super Resolution kernels are
-kernels written for RDNA3 and RDNA4.
+replaced by kernels written for RDNA3 and RDNA4.
-This is an early release, tested on one GPU (Radeon RX 7700 XT) in about a dozen DirectX 12 games.
+This is an early release. Super Resolution game testing primarily uses Radeon RX 7700 XT.
-SUPPORTED_GAMES.md in the source repository lists them, with per-game setup notes. Expect problems
+Ray Reconstruction presets D/E, temporal history, and separate alpha were validated on RX 9070 XT
-in other games.
+(gfx1201) with native FP8 WMMA enabled. Other RR GPU targets and RTX image-quality parity are unverified.
 SUPPORTED_GAMES.md in the source repository lists per-game setup notes. Expect problems in other games.
 What you need
@@ -144,7 +147,9 @@ If something goes wrong
  can be expected with the shim under d4r/; do not move it into the game root to silence that check.
 - A wrong image: try Model = E in d4r/d4r.ini, and report the problem with both logs.
 - Do not use d4r in games with anti-cheat. OptiScaler's DLL injection can get an account banned.
- Not supported: DLSS Frame Generation, DLSS Ray Reconstruction and DirectX 11.
+- Ray Reconstruction requires the CUDA-capable 3.10.7 nvngx_dlssd.dll beside the shim and the matching
  patched d3d12.dll/d3d12core.dll. Enable [Kernels] NativeFp8; the scalar FP8 RR path is not validated.
 - Not supported: DLSS Frame Generation and DirectX 11.
 - Development builds expose an experimental native Vulkan NGX path using CPU staging.
  Youngblood is the first intended game test; gameplay remains unverified. See the
  project documentation (docs/vulkan.md) for its separate Vulkan setup and limits.
@@ -1,3 +1,4 @@
 <!-- Modified in this fork for CUDA Ray Reconstruction support and validation (2026). -->
 # Unreleased
 - Native engine, preset L (DLSS 4.5, render preset 12): the engine now runs L ([Engine] Enabled, models in `engine/l` and `engine/l-ldr`, built by `engine/compile_m.py` from a capture of preset L; `scripts/install_d4r_runtime.sh` stages `D4R_ENGINE_MODEL_L` / `D4R_ENGINE_MODEL_L_LDR`). Every stage, `d4r-m` on captured frames and the D3D12 harness are byte-identical to the CUDA backend running NVIDIA's translated kernels, apart from isolated single-frame knife edges (96-122 dB). 1280x720 -> 3840x2160 takes 3.8 ms on an RX 9070 XT. Display-resolution motion vectors keep the CUDA backend, as for M. Details: `engine/L_NOTES.md`.
@@ -15,6 +16,28 @@
 - Prevent stale in-place output mappings when games destroy and recreate output textures. The harness now supports `D4R_HARNESS_RECREATE_OUTPUT=1` for this regression.
 - Integration validation on RX 9070 XT: moving 12-frame RGBA16F/R10G10B10A2 copy, packed and in-place sequences matched byte-for-byte; fast and accuracy paths, token-disabled reference layers, MicroCUDA tube-cache reuse, discarded input recordings and four feature-recreation cycles were exercised. gfx1101 and gfx1200 compatibility builds also passed; those targets were not runtime-tested.
 ## Ray Reconstruction
 - Run the signed CUDA-capable 3.10.7 DLSS-Denoiser through the NGX core's feature-13 lifecycle.
  Presets D and E select distinct network weights independently of the Super Resolution preset.
 - Preserve per-frame guide resources, camera matrices, subrects, and stable CUDA texture handles.
  Support normalized 8-bit albedo guides with the CUDA array-format ABI and stage full guide rows.
 - Support pre-Init feature discovery through the denoiser's requirements export without initializing
  NGX with a provisional application identity. Initialized capability checks remain authoritative.
 - Keep separate alpha and RGB outputs paired, including padded output subrects; consume alpha
  readback synchronously before CPU conversion.
 - Bound R8_UINT alpha stores to one byte and balance guide/alpha references when evaluation
  cannot place its frame marker. Publish cached denoiser capabilities as a coherent snapshot.
 - Extend the ZLUDA patch stack for real denoiser PTX, LLVM lowering, half texture/store operations,
  and release/acquire ordering around workgroup barriers.
 - Verified D/E temporal evaluation and separate alpha on RX 9070 XT (gfx1201) with native FP8 WMMA.
  Scalar FP8 temporal RR produced nonfinite output and is not validated. Cold compilation can stall
  initial rendering; other GPU targets and image-quality parity with RTX hardware remain unverified.
 - Cyberpunk 2077 2.31 rendered loaded local saves with path tracing and RR preset D enabled at
  1920×1080 and 3840×2160 Ultra Performance on gfx1201. The actual overlay reported DLSSD 310.7.0,
  and CUDA feature-13 evaluations completed. The observed ~53 FPS at 1080p and ~20 FPS at 4K are
  snapshots, not comparative benchmarks. Newer DLSS 4.5 RR preset F is not implemented or validated.
 # d4r 0.1.3
 This release adds an optional accuracy mode and fixes the GLIBC compatibility failure, DLSS artifacts, and stale output-buffer redirects reported after 0.1.2. Accuracy mode is **off by default**.
@@ -1,3 +1,5 @@
 Modified in this fork for CUDA Ray Reconstruction support and validation (2026).
 d4r @VERSION@: where every file in this release comes from
 ============================================================
@@ -35,7 +37,7 @@ Third-party software, unmodified or patched by d4r:
  OptiScaler.ini         OptiScaler 0.9.4's OptiScaler.ini with the values d4r needs set
                         (packaging/optiscaler.settings in the d4r repository)
-@clean Not included: NVIDIA's DLSS library (nvngx_dlss.dll) and NGX runtime (_nvngx.dll). You supply them
+@clean Not included: NVIDIA's DLSS libraries (nvngx_dlss.dll, nvngx_dlssd.dll) and NGX runtime (_nvngx.dll). You supply them
@clean and are responsible for complying with NVIDIA's license terms.
@full NVIDIA's software, not covered by any license in d4r/licenses and included by whoever built this
@full zip, not by NVIDIA. NVIDIA's license does not permit redistributing it; do not share this zip:
@@ -48,6 +50,7 @@ Third-party software, unmodified or patched by d4r:
@installer   d4r/kernels/*/hiluma_engine_*, rrlite_enc0_*, rrlite_dec0_*, rrlite_post_*, rrlite_downsample_*
@installer                        NVIDIA's PTX with parts replaced by d4r code (kernels/tex) or translated to
@installer                        HIP (kernels/native), compiled by ZLUDA or ROCm's clang.
@full   Optional d4r/nvngx_dlssd.dll: NVIDIA DLSS Ray Reconstruction library, when supplied through D4R_BUNDLE_DLSSD
@full   d4r/kernels/*/hiluma_engine_*, rrlite_enc0_*, rrlite_dec0_*, rrlite_post_*, rrlite_downsample_*
@full                        DLSS kernels: NVIDIA's PTX from nvngx_dlss.dll, with parts replaced by d4r code
@full                        (kernels/tex in the d4r repository) or translated to HIP (kernels/native),
@@ -1,3 +1,4 @@
 ; Modified in this fork for CUDA Ray Reconstruction support and validation (2026).
 ; d4r settings for this game. This file lives in the game's d4r folder and applies to this game only.
 ; "auto" means the built-in default. Any setting can also be overridden with the environment variable
 ; that config/d4r.ini.default in the source repository lists for it, e.g. in the game's launch options.
@@ -21,6 +22,15 @@ Model = K
 ; true shows it.
 ShowWatermark = false
 [RayReconstruction]
 ; Requires a CUDA-capable nvngx_dlssd.dll (tested: 310.7) beside nvngx.dll and RR enabled in game.
 ; auto uses the denoiser when available; false disables the route.
 ; RR D/E temporal validation on gfx1201 uses [Kernels] NativeFp8 = on; scalar FP8 RR is not validated.
 Enable = auto
 ; Independent denoiser presets: auto (game choice), default, D or E (transformer).
 ; E requires a denoiser DLL containing that model.
 Model = auto
 [Latency]
 ; 0: every frame shows its own DLSS result, as with a native upscaler (needs the d3d12.dll that came
 ;    with d4r next to the game).
@@ -0,0 +1,716 @@
 From: d4r
 Subject: [PATCH] Real DLSS-Denoiser PTX and AMDGPU compiler fixes
 Apply after 0002-0007, with the LLVM submodule initialized. Adds the denoiser
 surface stores, floating-coordinate half texture sampling, finite-saturating
 half conversion, typed half exponentiation, wide-register byte stores,
 half unordered comparison and atomic reduction. CUDA block barriers release
 writes before the hardware barrier and acquire visibility after it.
 Fixes strict division operand selection in AMDGPU and quadratic instruction-order
 rebuilding in CodeGenPrepare operand sinking. Scalar sinking checks for a viable
 target before scanning memory dependencies, retaining write tracking for reads.
 Device helper bitcode must be regenerated before building ZLUDA.
 ---
 --- a/llvm_zluda/build.rs
 +++ b/llvm_zluda/build.rs
@@ -116,15 +116,24 @@
 // We resort to this garbage because Cargo is ideologically opposed to
 // being useful, see here: https://github.com/rust-lang/cargo/issues/11054
 fn try_get_build_profile_name() -> String {
 +    // Walk up from OUT_DIR until a directory naming a build profile turns up, rather than popping
 +    // a fixed number of times. Cargo has added levels to the target/<profile>/build/<crate>/<hash>
 +    // /out layout more than once (current Cargo emits build/<crate>/<hash>/out), and each time it
 +    // does, the fixed count lands on "build" rather than the profile and panics with
 +    // "Unknown build profile: build". Stopping at the first ancestor whose name is a profile this
 +    // script knows survives that.
 +    const PROFILES: [&str; 4] = ["debug", "dev-llvm", "release", "release-lto"];
     let mut path = PathBuf::from(env::var("OUT_DIR").unwrap());
 -    path.pop();
 -    path.pop();
 -    path.pop();
 -    let name = path.file_name().unwrap().to_str().unwrap();
 -    if !["debug", "dev-llvm", "release", "release-lto"].contains(&name) {
 -        panic!("Unknown build profile: {}", name);
 -    }
 -    name.to_string()
 +    loop {
 +        if let Some(name) = path.file_name().and_then(|name| name.to_str()) {
 +            if PROFILES.contains(&name) {
 +                return name.to_string();
 +            }
 +        }
 +        if !path.pop() {
 +            panic!("no build profile directory above OUT_DIR");
 +        }
 +    }
 }
 fn get_llvm_build_path(llvm_build_dir: &PathBuf, cmake_profile: &str) -> PathBuf {
 --- a/ptx/lib/zluda_ptx_impl.cpp
 +++ b/ptx/lib/zluda_ptx_impl.cpp
@@ -96,6 +96,8 @@
 typedef int s32;
 typedef int v1s32 __attribute__((ext_vector_type(1)));
 typedef short v2s16 __attribute__((ext_vector_type(2)));
 +typedef short v4s16 __attribute__((ext_vector_type(4)));
 +typedef unsigned char v4u8 __attribute__((ext_vector_type(4)));
 typedef int v2s32 __attribute__((ext_vector_type(2)));
 typedef int v4s32 __attribute__((ext_vector_type(4)));
 typedef int v8s32 __attribute__((ext_vector_type(8)));
@@ -125,6 +127,8 @@
 extern "C"
 {
 +    __device__ int __ockl_image_width_2D(unsigned int ADDRESS_SPACE_CONSTANT* image);
 +    __device__ int __ockl_image_height_2D(unsigned int ADDRESS_SPACE_CONSTANT* image);
     extern "C" __attribute__((constant)) uint32_t __oclc_ISA_version __device__;
     // d4r: kernels may be compiled as wave64 (D4R_ZLUDA_WAVE64, offline texture-kernel builds, linked with
@@ -312,8 +316,9 @@
     void FUNC(bar_sync)(uint32_t)
     {
 -        __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");
 +        __builtin_amdgcn_fence(__ATOMIC_RELEASE, "workgroup");
         __builtin_amdgcn_s_barrier();
 +        __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup");
     }
     int32_t __ockl_wgred_and_i32(int32_t) __device__;
@@ -729,6 +734,22 @@
     Fp8x2 FUNC(cvt_rn_satfinite_e5m2x2_f32)(float a, float b)
     {
         return {cvt_float_to_fp8(b, __HIP_E5M2), cvt_float_to_fp8(a, __HIP_E5M2)};
 +    }
 +
 +    // cvt.rn.satfinite.f16.f32. Round to nearest even first, then clamp: a plain
 +    // f32->f16 conversion turns an out-of-range magnitude into infinity, while
 +    // .satfinite wants the largest finite half (65504) with the input's sign. A
 +    // NaN input stays NaN rather than being clamped, which is the point of the
 +    // .satfinite qualifier. Rounding before clamping is what the ISA specifies,
 +    // so 65520 (which rounds to 65536) still ends at 65504.
 +    f16 FUNC(cvt_rn_satfinite_f16_f32)(float a)
 +    {
 +        if (__builtin_isnan(a))
 +            return f16(__builtin_nanf(""));
 +        f16 rounded = f16(a);
 +        if (__builtin_isinf(static_cast<float>(rounded)))
 +            return f16(__builtin_copysignf(65504.0f, a));
 +        return rounded;
     }
     // .relu clamps negative results (including -0) to +0 and keeps NaN.
@@ -2231,6 +2252,17 @@
         return FUNC_CALL(texobj_2d_v4_f16_s32)(uint64_t(texref->textureObject), coord);
     }
 +    v4f16 FUNC(texobj_2d_v4_f16_f32)(uint64_t texobj, v2f32 coord)
 +    {
 +        auto [i, s] = get_image_and_sampler(texobj);
 +        auto result = sample_or_fetch_2D(i, s, coord);
 +        return {f16(result.x), f16(result.y), f16(result.z), f16(result.w)};
 +    }
 +    v4f16 FUNC(texref_2d_v4_f16_f32)(struct textureReference GLOBAL_SPACE *texref, v2f32 coord)
 +    {
 +        return FUNC_CALL(texobj_2d_v4_f16_f32)(uint64_t(texref->textureObject), coord);
 +    }
 +
 #define tex_2d_level(RETURN_TYPE, COORD_TYPE)                                                                                                                                           \
     v4##RETURN_TYPE FUNC(texobj_level_2d_v4_##RETURN_TYPE##_##COORD_TYPE)(uint64_t texobj, v2##COORD_TYPE coord, float lod)                                                           \
     {                                                                                                                                                                                   \
@@ -2612,7 +2644,7 @@
     }
     // Replaces the bytes of one channel that overlap the stored bytes.
 -    static uint32_t surface_merge_channel(uint32_t bits, int channel, int channel_bytes, uint32_t data,
 +    static uint32_t surface_merge_channel(uint32_t bits, int channel, int channel_bytes, uint64_t data,
                                           int byte_offset, int byte_count) __device__
     {
         for (int b = 0; b < 4; b++)
@@ -2636,10 +2668,11 @@
         return surface_decode_channel(surface_select_channel(c0, c1, c2, c3, channel), data_type);
     }
 -    // Stores byte_count (at most 4) little-endian bytes of data at the byte
 -    // offset coord.x of row coord.y.
 +    // Stores byte_count (at most 8) little-endian bytes of data at the byte
 +    // offset coord.x of row coord.y. 64 bits wide, because a v4 packed-half
 +    // store writes four halves in one go.
     static void surface_store_raw_2d(unsigned int ADDRESS_SPACE_CONSTANT *image, v2s32 coord,
 -                                     uint32_t data, int byte_count) __device__
 +                                     uint64_t data, int byte_count) __device__
     {
         int data_type = __ockl_image_channel_data_type_2D(image);
         int order = __ockl_image_channel_order_2D(image);
@@ -2659,7 +2692,7 @@
             int chunk_bytes = pixel_bytes - byte_offset;
             if (chunk_bytes > byte_count - consumed)
                 chunk_bytes = byte_count - consumed;
 -            uint32_t fragment = data >> (8 * consumed);
 +            uint64_t fragment = data >> (8 * consumed);
             int2 coordinates{pixel_x, coord.y};
             auto native_coordinates = get_native_vector(coordinates);
             uint32_t c0 = 0, c1 = 0, c2 = 0, c3 = 0;
@@ -2694,6 +2727,41 @@
         surface_store_raw_2d(image, coord, data, 4);
     }
 +    // Four packed halves are one 64-bit texel write. The lanes go out in
 +    // increasing address order, matching the v2 form.
 +    void FUNC(surfobj_b_2d_v4_b16_zero)(uint64_t surface_object, v2s32 coord, v4s16 value)
 +    {
 +        auto wrapper = load_surface_handle(surface_object);
 +        auto image = reinterpret_cast<unsigned int ADDRESS_SPACE_CONSTANT *>(
 +            static_cast<uintptr_t>(wrapper.native));
 +        uint64_t data = uint64_t(uint16_t(value.x))
 +            | (uint64_t(uint16_t(value.y)) << 16)
 +            | (uint64_t(uint16_t(value.z)) << 32)
 +            | (uint64_t(uint16_t(value.w)) << 48);
 +        surface_store_raw_2d(image, coord, data, 8);
 +    }
 +
 +    // A .trap store must fault instead of dropping an out-of-range write.
 +    void FUNC(surfobj_b_2d_v4_b8_trap)(uint64_t surface_object, v2s32 coord, v4u8 value)
 +    {
 +        auto wrapper = load_surface_handle(surface_object);
 +        auto image = reinterpret_cast<unsigned int ADDRESS_SPACE_CONSTANT *>(
 +            static_cast<uintptr_t>(wrapper.native));
 +        const int channel_bytes = surface_channel_bytes(__ockl_image_channel_data_type_2D(image));
 +        const int channels = surface_channel_count(__ockl_image_channel_order_2D(image));
 +        const uint64_t row_bytes = static_cast<uint64_t>(__ockl_image_width_2D(image)) *
 +                                   static_cast<uint64_t>(channel_bytes * channels);
 +        if (channel_bytes == 0 || coord.x < 0 || coord.y < 0 ||
 +            coord.y >= __ockl_image_height_2D(image) ||
 +            static_cast<uint64_t>(coord.x) + sizeof(uint32_t) > row_bytes)
 +            __builtin_trap();
 +        uint32_t data = uint32_t(value.x)
 +            | (uint32_t(value.y) << 8)
 +            | (uint32_t(value.z) << 16)
 +            | (uint32_t(value.w) << 24);
 +        surface_store_raw_2d(image, coord, data, 4);
 +    }
 +
     void FUNC(surfobj_b_2d_b32_zero)(uint64_t surface_object, v2s32 coord, uint32_t value)
     {
         auto wrapper = load_surface_handle(surface_object);
 --- a/ptx/src/pass/insert_post_saturation.rs
 +++ b/ptx/src/pass/insert_post_saturation.rs
@@ -76,6 +76,7 @@
         | ast::Instruction::MadExtended { .. }
         | ast::Instruction::And { .. }
         | ast::Instruction::Atom { .. }
 +        | ast::Instruction::Red { .. }
         | ast::Instruction::AtomCas { .. }
         | ast::Instruction::BarWarp { .. }
         | ast::Instruction::Bar { .. }
 --- a/ptx/src/pass/instruction_mode_to_global_mode/mod.rs
 +++ b/ptx/src/pass/instruction_mode_to_global_mode/mod.rs
@@ -2210,6 +2210,7 @@
         | ast::Instruction::CvtPack { .. }
         | ast::Instruction::Cvta { .. }
         | ast::Instruction::Atom { .. }
 +        | ast::Instruction::Red { .. }
         | ast::Instruction::Mul24 { .. }
         | ast::Instruction::Nanosleep { .. }
         | ast::Instruction::AtomCas { .. }
 --- a/ptx/src/pass/llvm/emit.rs
 +++ b/ptx/src/pass/llvm/emit.rs
@@ -683,6 +683,9 @@
             ast::Instruction::Rsqrt { data, arguments } => self.emit_rsqrt(data, arguments),
             ast::Instruction::Selp { data, arguments } => self.emit_selp(data, arguments),
             ast::Instruction::Atom { data, arguments } => self.emit_atom(data, arguments),
 +            ast::Instruction::Red { data, arguments } => {
 +                self.emit_atomic_rmw(data, arguments.src1, arguments.src2).map(|_| ())
 +            }
             ast::Instruction::AtomCas { data, arguments } => self.emit_atom_cas(data, arguments),
             ast::Instruction::Div { data, arguments } => self.emit_div(data, arguments),
             ast::Instruction::Neg { data, arguments } => self.emit_neg(data, arguments),
@@ -1321,9 +1324,20 @@
         data: ast::AtomDetails,
         arguments: ast::AtomArgs<SpirvWord>,
     ) -> Result<(), TranslateError> {
 +        let value = self.emit_atomic_rmw(data, arguments.src1, arguments.src2)?;
 +        self.resolver.register(arguments.dst, value);
 +        Ok(())
 +    }
 +
 +    fn emit_atomic_rmw(
 +        &mut self,
 +        data: ast::AtomDetails,
 +        src1: SpirvWord,
 +        src2: SpirvWord,
 +    ) -> Result<LLVMValueRef, TranslateError> {
         let builder = self.builder;
 -        let src1 = self.resolver.value(arguments.src1)?;
 -        let src2 = self.resolver.value(arguments.src2)?;
 +        let src1 = self.resolver.value(src1)?;
 +        let src2 = self.resolver.value(src2)?;
         let op = match data.op {
             ast::AtomicOp::And => LLVMZludaAtomicRMWBinOp::LLVMZludaAtomicRMWBinOpAnd,
             ast::AtomicOp::Or => LLVMZludaAtomicRMWBinOp::LLVMZludaAtomicRMWBinOpOr,
@@ -1344,7 +1358,7 @@
             ast::AtomicOp::FloatMin => LLVMZludaAtomicRMWBinOp::LLVMZludaAtomicRMWBinOpFMin,
             ast::AtomicOp::FloatMax => LLVMZludaAtomicRMWBinOp::LLVMZludaAtomicRMWBinOpFMax,
         };
 -        self.resolver.register(arguments.dst, unsafe {
 +        Ok(unsafe {
             LLVMZludaBuildAtomicRMW(
                 builder,
                 op,
@@ -1353,8 +1367,7 @@
                 get_scope(data.scope)?,
                 get_ordering(data.semantics),
             )
 -        });
 -        Ok(())
 +        })
     }
     fn emit_atom_cas(
 --- a/ptx/src/pass/replace_instructions_with_functions.rs
 +++ b/ptx/src/pass/replace_instructions_with_functions.rs
@@ -191,7 +191,8 @@
                                     data: ast::CvtDetails {
                                         from: ast::ScalarType::U32,
                                         to: ast::ScalarType::Pred,
 -                                        mode: ast::CvtMode::Truncate
 +                                        mode: ast::CvtMode::Truncate,
 +                                        satfinite: false
                                     },
                                     arguments: ast::CvtArgs {
                                         dst: dst_pred,
@@ -241,7 +242,8 @@
                                     data: ast::CvtDetails {
                                         from: ast::ScalarType::U32,
                                         to: ast::ScalarType::Pred,
 -                                        mode: ast::CvtMode::Truncate
 +                                        mode: ast::CvtMode::Truncate,
 +                                        satfinite: false
                                     },
                                     arguments: ast::CvtArgs {
                                         dst: dst_pred,
@@ -259,6 +261,7 @@
                             from: from @ (ast::ScalarType::E4m3x2 | ast::ScalarType::E5m2x2),
                             to: ast::ScalarType::F16x2,
                             mode: _,
 +                            satfinite: _,
                         },
                     arguments:
                         ast::CvtArgs {
@@ -310,7 +313,8 @@
                             data: ast::CvtDetails {
                                 from: ast::ScalarType::B32,
                                 to: ast::ScalarType::F16x2,
 -                                mode: ast::CvtMode::Bitcast
 +                                mode: ast::CvtMode::Bitcast,
 +                                satfinite: false
                             },
                             arguments: ast::CvtArgs {
                                 dst,
@@ -449,6 +453,14 @@
         } => to_call(resolver, fn_declarations, "surfobj_p_2d_v4_b32_zero".into(), i)?,
         i @ ptx_parser::Instruction::Sust {
             data: ast::SustData {
 +                vector_width: Some(4),
 +                scalar_type: ast::ScalarType::B16,
 +                formatted: false,
 +            },
 +            ..
 +        } => to_call(resolver, fn_declarations, "surfobj_b_2d_v4_b16_zero".into(), i)?,
 +        i @ ptx_parser::Instruction::Sust {
 +            data: ast::SustData {
                 vector_width: Some(2),
                 scalar_type: ast::ScalarType::B16,
                 formatted: false,
@@ -696,6 +708,7 @@
                     from: ast::ScalarType::F16x2,
                     to: to @ (ast::ScalarType::E4m3x2 | ast::ScalarType::E5m2x2),
                     mode,
 +                    satfinite: _,
                 },
             arguments: _,
         } => {
@@ -718,6 +731,7 @@
                     from: ast::ScalarType::F32,
                     to: to @ (ast::ScalarType::E4m3x2 | ast::ScalarType::E5m2x2),
                     mode,
 +                    satfinite: _,
                 },
             arguments: _,
         } => {
@@ -733,6 +747,57 @@
                 resolver,
                 fn_declarations,
                 format!("cvt_rn_satfinite{}_{}_f32", relu, to).into(),
 +                i,
 +            )?
 +        }
 +        // sust.b.2d.v4.b8.trap. In-bounds this is exactly the v4 packed store; only the
 +        // out-of-range behaviour differs from .zero, and AMD image stores cannot trap on
 +        // an out-of-range coordinate. The helper drops an out-of-range write instead of
 +        // trapping, which is what the .zero path already does - so this is correct for
 +        // every in-bounds store and differs only where the program would have faulted.
 +        i @ ptx_parser::Instruction::Sust {
 +            data: ast::SustData {
 +                vector_width: Some(4),
 +                scalar_type: ast::ScalarType::B8,
 +                formatted: false,
 +            },
 +            ..
 +        } => to_call(
 +            resolver,
 +            fn_declarations,
 +            "surfobj_b_2d_v4_b8_trap".into(),
 +            i,
 +        )?,
 +        // cvt.satfinite.f16.f32: finite saturation to half. A plain f32->f16
 +        // conversion turns an out-of-range magnitude into infinity; .satfinite
 +        // clamps to +/-65504 instead and leaves a NaN input as NaN. Rounding is
 +        // applied before clamping, so 65520 - which would round up past the
 +        // largest finite half - still ends at 65504.
 +        //
 +        // Only .rn is implemented: the helper rounds with the hardware's
 +        // round-to-nearest-even and then clamps. Accepting .rz/.rm/.rp and
 +        // quietly rounding the other way would miscompute, so they are rejected
 +        // rather than silently mistranslated.
 +        i @ ptx_parser::Instruction::Cvt {
 +            data:
 +                ptx_parser::CvtDetails {
 +                    from: ast::ScalarType::F32,
 +                    to: ast::ScalarType::F16,
 +                    satfinite: true,
 +                    mode: ast::CvtMode::FPTruncate { rounding, .. },
 +                    ..
 +                },
 +            arguments: _,
 +        } => {
 +            if rounding != ast::RoundingMode::NearestEven {
 +                return Err(error_todo_msg(
 +                    "cvt.satfinite.f16.f32 is only implemented with .rn rounding",
 +                ));
 +            }
 +            to_call(
 +                resolver,
 +                fn_declarations,
 +                "cvt_rn_satfinite_f16_f32".into(),
                 i,
             )?
         }
 --- a/ptx_parser/src/ast.rs
 +++ b/ptx_parser/src/ast.rs
@@ -108,6 +108,17 @@
                 src2: T,
             }
         },
 +        Red {
 +            type: &data.type_,
 +            data: AtomDetails,
 +            arguments<T>: {
 +                src1: {
 +                    repr: T,
 +                    space: { data.space },
 +                },
 +                src2: T,
 +            }
 +        },
         AtomCas {
             type: Type::Scalar(data.type_),
             data: AtomCasDetails,
@@ -350,8 +361,12 @@
                 },
             }
         },
 +        // ex2.approx is typed by the instruction suffix (.f32, .f16, .f16x2,
 +        // .bf16, .bf16x2). Hardcoding .f32 made the operand check compare a
 +        // 16-bit register against a 32-bit instruction type, which rejects
 +        // `ex2.approx.f16` outright.
         Ex2 {
 -            type: Type::Scalar(ScalarType::F32),
 +            type: Type::Scalar(data.type_.clone()),
             data: TypeFtz,
             arguments<T>: {
                 dst: T,
@@ -1035,6 +1050,13 @@
                             None => Type::Scalar(data.scalar_type),
                         }
                     },
 +                    // PTX has no 8-bit registers: a `.b8` operand lives in the
 +                    // low bits of a wider register, and the upper bits are
 +                    // undefined. DLSS emits `ld.local.v4.u8` into `.b16`
 +                    // registers and feeds the same registers to
 +                    // `sust.b.2d.v4.b8`, so the source must be allowed to be
 +                    // wider than the instruction type.
 +                    relaxed_type_check: true,
                 },
             }
         },
@@ -2475,6 +2497,11 @@
     pub from: ScalarType,
     pub to: ScalarType,
     pub mode: CvtMode,
 +    /// `.satfinite`: out-of-range results saturate to the largest finite value
 +    /// of the destination rather than to infinity, and a NaN input stays NaN.
 +    /// Distinct from `.sat`, the integer saturation modifier, and meaningful
 +    /// only where the destination is a float type.
 +    pub satfinite: bool,
 }
 #[derive(Clone, Copy)]
@@ -2530,6 +2557,7 @@
         rnd: Option<RawRoundingMode>,
         ftz: bool,
         saturate: bool,
 +        satfinite: bool,
         relu: bool,
         dst: ScalarType,
         src: ScalarType,
@@ -2660,6 +2688,7 @@
             mode,
             to: dst,
             from: src,
 +            satfinite,
         }
     }
 }
 --- a/ptx_parser/src/lib.rs
 +++ b/ptx_parser/src/lib.rs
@@ -2607,6 +2607,7 @@
         }
     }
 +
     set.CmpOp.BoolOp{.ftz}.dtype.stype  d, a, b, {!}c => {
         let (negate_src3, c) = c;
         let base = ast::SetpData::try_parse(state, cmpop, ftz, stype);
@@ -2628,7 +2629,10 @@
                                   .lo, .ls, .hi, .hs, // signed
                                   .equ, .neu, .ltu, .leu, .gtu, .geu, .num, .nan }; // float-only
     .BoolOp: SetpBoolPostOp = { .and, .or, .xor };
 -    .dtype: ScalarType = { .u32, .s32, .f32 };
 +    // .dtype is the type of the compared operands. The denoiser compares halves with
 +    // set.nan.f16.f16, which needs .f16 here; .bf16 is the other narrow float the ISA
 +    // allows and comes along for free because the compare itself is type-generic.
 +    .dtype: ScalarType = { .u32, .s32, .f32, .f16, .bf16 };
     .dtype_f16x2: ScalarType = { .u32, .f16x2 };
     .stype: ScalarType = { .b16, .b32, .b64, .u16, .u32, .u64, .s16, .s32, .s64, .f16, .f32, .f64 };
@@ -2703,10 +2707,15 @@
         if satfinite && !matches!(
             (dtype, atype),
             (ScalarType::E4m3x2 | ScalarType::E5m2x2, ScalarType::F32 | ScalarType::F16x2)
 +                // Finite saturation to a narrow float: out-of-range results become the
 +                // largest finite value of the destination (f16: +/-65504) instead of
 +                // infinity, and a NaN input stays NaN. The denoiser's encoder uses this
 +                // on the way down to f16.
 +                | (ScalarType::F16, ScalarType::F32)
         ) {
 -            state.errors.push(PtxError::Todo("cvt.satfinite is only implemented for E4M3/E5M2 destinations".to_string()));
 -        }
 -        let data = ast::CvtDetails::new(&mut state.errors, ifrnd, ftz, sat, false, dtype, atype);
 +            state.errors.push(PtxError::Todo("cvt.satfinite is only implemented for E4M3/E5M2/F16 destinations".to_string()));
 +        }
 +        let data = ast::CvtDetails::new(&mut state.errors, ifrnd, ftz, sat, satfinite, false, dtype, atype);
         let arguments = ast::CvtArgs { dst: d, src: a, src2: b };
         ast::Instruction::Cvt {
             data, arguments
@@ -2719,7 +2728,7 @@
         if satfinite {
             state.errors.push(PtxError::Todo("cvt.frnd2 instruction with satfinite modifier".to_string()));
         }
 -        let data = ast::CvtDetails::new(&mut state.errors, Some(frnd2), false, false, relu, x2_to_type, x2_from_type);
 +        let data = ast::CvtDetails::new(&mut state.errors, Some(frnd2), false, false, satfinite, relu, x2_to_type, x2_from_type);
         ast::Instruction::Cvt {
             data,
             arguments:  ast::CvtArgs { dst: d, src: a, src2: b }
@@ -2729,7 +2738,7 @@
     // cvt.rn.satfinite{.relu}.f8x2type.f16x2     d, a;
     // The form without .relu is handled by the generic rule above.
     cvt.f8rnd.satfinite.relu.f8x2type.f8_from_type    d, a {, b} => {
 -        let data = ast::CvtDetails::new(&mut state.errors, Some(f8rnd), false, false, true, f8x2type, f8_from_type);
 +        let data = ast::CvtDetails::new(&mut state.errors, Some(f8rnd), false, false, true, true, f8x2type, f8_from_type);
         ast::Instruction::Cvt {
             data,
             arguments:  ast::CvtArgs { dst: d, src: a, src2: b }
@@ -3654,6 +3663,25 @@
     ScalarType =                    { .b16, .b128, .f32 };
     StateSpace =                    { .global };
     RawAtomicOp =                   { .exch };
 +    // A reduction preserves atomic ordering and scope, but does not return the old value.
 +    red{.sem}{.scope}{.space}.op.type [a], b => {
 +        ast::Instruction::Red {
 +            data: AtomDetails {
 +                semantics: sem.unwrap_or(AtomSemantics::Relaxed),
 +                scope: scope.unwrap_or(MemScope::Gpu),
 +                space: space.unwrap_or(StateSpace::Generic),
 +                op: ast::AtomicOp::new(op, type_.kind()),
 +                type_: type_.into()
 +            },
 +            arguments: RedArgs { src1: a, src2: b }
 +        }
 +    }
 +    .sem: AtomSemantics = { .relaxed, .release };
 +    .scope: MemScope = { .cta, .cluster, .gpu, .sys };
 +    .space: StateSpace = { .global, .shared };
 +    .op: RawAtomicOp = { .and, .or, .xor, .add, .inc, .dec, .min, .max };
 +    .type: ScalarType = { .b32, .b64, .u32, .u64, .s32, .s64 };
 +
     // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-div
     // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-div
@@ -4571,6 +4599,42 @@
     sust.b.2d.v2.b16.zero [a, b], c => {
         Instruction::Sust {
             data: SustData { vector_width: Some(2), scalar_type: ScalarType::B16, formatted: false },
 +            arguments: SustArgs { src_surf: a, src_coord: b, src_value: c }
 +        }
 +    }
 +
 +    // Unformatted packed-half store of a whole 64-bit word. The denoiser's swin kernels write
 +    // four halves per texel, so this is not a spelling of the v2 form.
 +    sust.b.2d.v4.b16.zero [a, b], c => {
 +        Instruction::Sust {
 +            data: SustData { vector_width: Some(4), scalar_type: ScalarType::B16, formatted: false },
 +            arguments: SustArgs { src_surf: a, src_coord: b, src_value: c }
 +        }
 +    }
 +
 +    // `.cs` is the streaming cache policy. It is a hint about how the write is cached, not about
 +    // what is written: the stored bytes and the zero-extension are identical to the default
 +    // (`.wb`) spelling, so these lower to the same helper rather than carrying a cache-policy
 +    // field that nothing would consume.
 +    sust.b.2d.cs.v2.b16.zero [a, b], c => {
 +        Instruction::Sust {
 +            data: SustData { vector_width: Some(2), scalar_type: ScalarType::B16, formatted: false },
 +            arguments: SustArgs { src_surf: a, src_coord: b, src_value: c }
 +        }
 +    }
 +
 +    sust.b.2d.cs.v4.b16.zero [a, b], c => {
 +        Instruction::Sust {
 +            data: SustData { vector_width: Some(4), scalar_type: ScalarType::B16, formatted: false },
 +            arguments: SustArgs { src_surf: a, src_coord: b, src_value: c }
 +        }
 +    }
 +    // sust.b.2d.v4.b8.trap: four bytes written as one 32-bit texel store. Same store as
 +    // the .zero forms; only the out-of-range behaviour differs, and that is handled by the
 +    // lowering rather than by the parser.
 +    sust.b.2d.v4.b8.trap [a, b], c => {
 +        Instruction::Sust {
 +            data: SustData { vector_width: Some(4), scalar_type: ScalarType::B8, formatted: false },
             arguments: SustArgs { src_surf: a, src_coord: b, src_value: c }
         }
     }
 --- a/ext/llvm-project/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
 +++ b/ext/llvm-project/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -12501,8 +12501,13 @@
   SDLoc SL(Op);
   EVT VT = Op.getValueType();
 -  SDValue LHS = Op.getOperand(0);
 -  SDValue RHS = Op.getOperand(1);
 +  // ZLUDA changes start
 +  // A strict node carries the chain in operand 0, so the value operands are
 +  // shifted by one (see FPExpansionBuilder).
 +  FPExpansionBuilder B(Op, DAG);
 +  SDValue LHS = B.getOperand(0);
 +  SDValue RHS = B.getOperand(1);
 +  // ZLUDA changes end
   SDValue LHSExt = DAG.getNode(ISD::FP_EXTEND, SL, MVT::f32, LHS);
   SDValue RHSExt = DAG.getNode(ISD::FP_EXTEND, SL, MVT::f32, RHS);
@@ -12611,8 +12616,13 @@
   Flags.setNoFPExcept(true);
   SDLoc SL(Op);
 -  SDValue LHS = Op.getOperand(0);
 -  SDValue RHS = Op.getOperand(1);
 +  // ZLUDA changes start
 +  // A strict node carries the chain in operand 0, so the value operands are
 +  // shifted by one (see FPExpansionBuilder).
 +  FPExpansionBuilder B(Op, DAG);
 +  SDValue LHS = B.getOperand(0);
 +  SDValue RHS = B.getOperand(1);
 +  // ZLUDA changes end
   const SDValue One = DAG.getConstantFP(1.0, SL, MVT::f32);
@@ -12741,8 +12751,13 @@
     return FastLowered;
   SDLoc SL(Op);
 -  SDValue X = Op.getOperand(0);
 -  SDValue Y = Op.getOperand(1);
 +  // ZLUDA changes start
 +  // A strict node carries the chain in operand 0, so the value operands are
 +  // shifted by one (see FPExpansionBuilder).
 +  FPExpansionBuilder B(Op, DAG);
 +  SDValue X = B.getOperand(0);
 +  SDValue Y = B.getOperand(1);
 +  // ZLUDA changes end
   const SDValue One = DAG.getConstantFP(1.0, SL, MVT::f64);
 --- a/ext/llvm-project/llvm/lib/CodeGen/CodeGenPrepare.cpp
 +++ b/ext/llvm-project/llvm/lib/CodeGen/CodeGenPrepare.cpp
@@ -7815,17 +7815,13 @@
   bool Changed = false;
   SmallVector<Use *, 4> ToReplace;
   Instruction *InsertPoint = I;
 -  DenseMap<const Instruction *, unsigned long> InstOrdering;
 -  unsigned long InstNumber = 0;
 -  for (const auto &I : *TargetBB)
 -    InstOrdering[&I] = InstNumber++;
   for (Use *U : reverse(OpsToSink)) {
     auto *UI = cast<Instruction>(U->get());
     if (isa<PHINode>(UI))
       continue;
     if (UI->getParent() == TargetBB) {
 -      if (InstOrdering[UI] < InstOrdering[InsertPoint])
 +      if (UI->comesBefore(InsertPoint))
         InsertPoint = UI;
       continue;
     }
 --- a/ext/llvm-project/llvm/lib/Transforms/Scalar/Sink.cpp
 +++ b/ext/llvm-project/llvm/lib/Transforms/Scalar/Sink.cpp
@@ -30,11 +30,6 @@
 static bool isSafeToMove(Instruction *Inst, AliasAnalysis &AA,
                          SmallPtrSetImpl<Instruction *> &Stores) {
 -  if (Inst->mayWriteToMemory()) {
 -    Stores.insert(Inst);
 -    return false;
 -  }
 -
   // Don't sink static alloca instructions.  CodeGen assumes allocas outside the
   // entry block are dynamically sized stack objects.
   if (AllocaInst *AI = dyn_cast<AllocaInst>(Inst))
@@ -104,9 +99,12 @@
                             SmallPtrSetImpl<Instruction *> &Stores,
                             DominatorTree &DT, LoopInfo &LI, AAResults &AA) {
 -  // Check if it's safe to move the instruction.
 -  if (!isSafeToMove(Inst, AA, Stores))
 -    return false;
 +  // Record writes even when this instruction has no sink target: preceding
 +  // reads must not be moved across them.
 +  if (Inst->mayWriteToMemory()) {
 +    Stores.insert(Inst);
 +    return false;
 +  }
   // FIXME: This should include support for sinking instructions within the
   // block they are currently in to shorten the live ranges.  We often get
@@ -155,6 +153,11 @@
   if (!SuccToSinkTo)
     return false;
 +  // Alias checks can scan every following write. Avoid doing that work for
 +  // instructions whose uses already rule out sinking.
 +  if (!isSafeToMove(Inst, AA, Stores))
 +    return false;
 +
   LLVM_DEBUG(dbgs() << "Sink" << *Inst << " (";
              Inst->getParent()->printAsOperand(dbgs(), false); dbgs() << " -> ";
              SuccToSinkTo->printAsOperand(dbgs(), false); dbgs() << ")\n");
@@ -1,4 +1,5 @@
 #!/usr/bin/env bash
 # Modified in this fork for CUDA Ray Reconstruction support and validation (2026).
 # Run inside packaging/build/Dockerfile.glibc241 with a recent official Rust toolchain and ROCm mounted.
 # Supply isolated, patched ZLUDA/vkd3d checkouts and the usual package_release.sh inputs.
 set -euo pipefail
@@ -1,4 +1,5 @@
 #!/usr/bin/env python3
 # Modified in this fork for CUDA Ray Reconstruction support and validation (2026).
 """Translate d4r.ini (see config/d4r.ini.default) into shell exports for the game launcher.
 usage: d4r_config.py [--config PATH] [--print | --launch]
@@ -86,6 +87,20 @@ def main(argv):
        else:
            fail(f"[DLSS] Model {model!r}: use one of {', '.join(list(PRESETS) + list(MODEL_ALIASES))} or a preset number")
    rr_enabled = flag("RayReconstruction", "Enable")
    if rr_enabled is not None:
        env["D4R_RR_ENABLE"] = "1" if rr_enabled else "0"
    rr_model = get("RayReconstruction", "Model")
    if rr_model is not None:
        rr_presets = {"DEFAULT": 0, "D": 4, "E": 5}
        rr_model = rr_model.upper()
        if rr_model in rr_presets:
            env["D4R_RR_PRESET"] = str(rr_presets[rr_model])
        elif rr_model in ("0", "4", "5"):
            env["D4R_RR_PRESET"] = rr_model
        else:
            fail(f"[RayReconstruction] Model {rr_model!r}: use auto, default, D or E")
    age = get("Latency", "FrameAge")
    if age is not None:
        if age not in ("0", "1", "2", "3"):
@@ -1,4 +1,5 @@
 #!/usr/bin/env bash
 # Modified in this fork for CUDA Ray Reconstruction support and validation (2026).
 set -euo pipefail
 # Stages the pieces a Proton game needs under D4R_RUNTIME_DIR
@@ -7,6 +8,7 @@ set -euo pipefail
 #   bin/d4r_nvngx.dll   D3D12-to-CUDA NGX core shim (OptiScaler NvngxPath target)
 #   ngx/_nvngx.dll      official NGX core
 #   dlss/nvngx_dlss.dll official DLSS SR feature DLL (also copied into bin/)
 #   dlss/nvngx_dlssd.dll optional DLSS Ray Reconstruction feature DLL (also in bin/)
 #   engine/k/, k-ldr/   models of preset K for the native engine ([Engine] in d4r.ini), when D4R_ENGINE_MODEL /
 #                       D4R_ENGINE_MODEL_LDR name the folders engine/compile_k.py wrote (--ldr for the second)
 #   engine/m/, m-ldr/   models of preset M, when D4R_ENGINE_MODEL_M / D4R_ENGINE_MODEL_M_LDR name the folders
@@ -14,9 +16,9 @@ set -euo pipefail
 #   engine/l/, l-ldr/   models of preset L, when D4R_ENGINE_MODEL_L / D4R_ENGINE_MODEL_L_LDR name the folders
 #                       engine/compile_m.py wrote from a capture of preset L
 # Each model folder is copied whole, with the native HIP network's hipnet.bin and hip/ when built with --hip.
-# usage: install_d4r_runtime.sh PATH_TO_NGX_CORE_DLL PATH_TO_NVNGX_DLSS_DLL
+# usage: install_d4r_runtime.sh PATH_TO_NGX_CORE_DLL PATH_TO_NVNGX_DLSS_DLL [PATH_TO_NVNGX_DLSSD_DLL]
-if [[ $# -ne 2 ]]; then
+if [[ $# -lt 2 || $# -gt 3 ]]; then
-  printf 'usage: %s PATH_TO_NGX_CORE_DLL PATH_TO_NVNGX_DLSS_DLL\n' "$0" >&2
+  printf 'usage: %s PATH_TO_NGX_CORE_DLL PATH_TO_NVNGX_DLSS_DLL [PATH_TO_NVNGX_DLSSD_DLL]\n' "$0" >&2
  exit 2
 fi
 ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)"
@@ -51,5 +53,10 @@ for variant in k k-ldr m m-ldr l l-ldr; do
  cp -r "$source_dir" "$RUNTIME/engine/$variant"
  printf 'Installed the native engine model in %s\n' "$RUNTIME/engine/$variant"
 done
 if [[ $# == 3 ]]; then
  copy_unless_same "$3" "$RUNTIME/dlss/nvngx_dlssd.dll"
  copy_unless_same "$3" "$RUNTIME/bin/nvngx_dlssd.dll"
  sha256sum "$RUNTIME/dlss/nvngx_dlssd.dll"
 fi
 sha256sum "$RUNTIME"/ngx/_nvngx.dll "$RUNTIME"/dlss/nvngx_dlss.dll
 printf 'Installed d4r runtime in %s\n' "$RUNTIME"
@@ -1,4 +1,5 @@
 #!/usr/bin/env bash
 # Modified in this fork for CUDA Ray Reconstruction support and validation (2026).
 # Builds the drag-in release: a zip whose contents are extracted into the folder that holds a game's
 # main .exe, like an OptiScaler release. packaging/D4R_README.txt describes the result.
 #
@@ -26,6 +27,7 @@
 # those is up to whoever publishes it (they are not covered by d4r's license):
 #   D4R_BUNDLE_DLSS  local nvngx_dlss.dll for texture builds (pinned by install.sh in installer mode)
 #   D4R_BUNDLE_NGX   local _nvngx.dll (only required for D4R_BUNDLE_NVIDIA=1)
 #   D4R_BUNDLE_DLSSD optional nvngx_dlssd.dll to include for Ray Reconstruction (full variant only)
 #   D4R_BUNDLE_TEX   directory with texture-kernel code objects (kernels/build.sh tex), one subdirectory
 #                    per target folder (gfx1101, gfx1201, gfx1201-fp8, ...), or a flat gfx1101 directory for
 #                    older builds
@@ -111,6 +113,7 @@ if [[ "$VARIANT" != clean ]]; then
    cp "$D4R_BUNDLE_DLSS" "$STAGE/d4r/nvngx_dlss.dll"
    : "${D4R_BUNDLE_NGX:?set D4R_BUNDLE_NGX}"
    cp "$D4R_BUNDLE_NGX" "$STAGE/d4r/ngx/_nvngx.dll"
    [[ -z "${D4R_BUNDLE_DLSSD:-}" ]] || cp "$D4R_BUNDLE_DLSSD" "$STAGE/d4r/nvngx_dlssd.dll"
  else
    cp "$ROOT/install.sh" "$STAGE/install.sh"
    chmod +x "$STAGE/install.sh"
@@ -0,0 +1,67 @@
 """Preset isolation, invalid-setting rejection, and environment precedence."""
 import os
 from pathlib import Path
 import subprocess
 import tempfile
 import unittest
 ROOT = Path(__file__).resolve().parents[1]
 CONFIG = ROOT / "scripts" / "d4r_config.py"
 RR_VARIABLES = ("D4R_RR_ENABLE", "D4R_RR_PRESET")
 def run_config(text, **overrides):
    """Runs d4r_config.py over `text`; returns (exit code, exports of interest).
    d4r_config.py also emits its own defaults for unrelated variables, so the result is
    filtered to the Ray Reconstruction and Super Sampling preset variables - the ones whose
    independence from each other is under test here.
    """
    with tempfile.TemporaryDirectory() as directory:
        ini = Path(directory) / "d4r.ini"
        ini.write_text(text)
        env = dict(os.environ)
        for name in RR_VARIABLES + ("D4R_DLSS_PRESET",):
            env.pop(name, None)
        env.update(overrides)
        result = subprocess.run(["python3", str(CONFIG), "--config", str(ini)], env=env,
                                capture_output=True, text=True)
    interesting = set(RR_VARIABLES) | {"D4R_DLSS_PRESET"}
    decided = {}
    for line in result.stdout.splitlines():
        if line.startswith("export "):
            key, _, value = line[len("export "):].partition("=")
            if key in interesting:
                decided[key] = value
    return result.returncode, decided
 class RayReconstructionConfigTests(unittest.TestCase):
    def test_rr_and_super_sampling_presets_do_not_collide(self):
        code, decided = run_config("[RayReconstruction]\nEnable = true\nModel = E\n[DLSS]\nModel = K\n")
        self.assertEqual(code, 0)
        self.assertEqual(decided["D4R_RR_ENABLE"], "1")
        self.assertEqual(decided["D4R_RR_PRESET"], "5")     # Ray Reconstruction preset E
        self.assertEqual(decided["D4R_DLSS_PRESET"], "11")  # Super Sampling preset K
    def test_invalid_rr_settings_are_rejected_rather_than_guessed(self):
        for text in ("[RayReconstruction]\nModel = F\n",
                     "[RayReconstruction]\nModel = K\n",
                     "[RayReconstruction]\nEnable = maybe\n"):
            with self.subTest(ini=text):
                code, decided = run_config(text)
                self.assertNotEqual(code, 0, f"accepted {text!r} as {decided}")
                for name in RR_VARIABLES:
                    self.assertNotIn(name, decided)
    def test_environment_wins_over_the_ini(self):
        # An explicitly exported variable is the user's decision and must survive.
        code, decided = run_config("[RayReconstruction]\nEnable = true\nModel = E\n",
                                   D4R_RR_ENABLE="0", D4R_RR_PRESET="4")
        self.assertEqual(code, 0)
        self.assertNotIn("D4R_RR_ENABLE", decided)
        self.assertNotIn("D4R_RR_PRESET", decided)
 if __name__ == "__main__":
    unittest.main()
@@ -1,3 +1,4 @@
 // Modified in this fork for CUDA Ray Reconstruction support and validation (2026).
 // Drives d4r_nvngx.dll through the NGX D3D12 API the way a game (or
 // OptiScaler's DLSS backend) does: a D3D12 device and command queue, input
 // textures in NON_PIXEL_SHADER_RESOURCE, an output UAV, and one
@@ -9,6 +10,7 @@
 #define WIDL_EXPLICIT_AGGREGATE_RETURNS
 #include <windows.h>
 #include <d3d12.h>
 #include <dxgi1_4.h>
 #include <algorithm>
 #include <array>
@@ -37,6 +39,7 @@ extern "C"
    void d4r_ngx_set_uint(void* parameters, const char* name, unsigned int value);
    void d4r_ngx_set_int(void* parameters, const char* name, int value);
    void d4r_ngx_set_d3d12_resource(void* parameters, const char* name, ID3D12Resource* value);
    void d4r_ngx_set_void(void* parameters, const char* name, void* value);
    NgxResult d4r_ngx_get_int(void* parameters, const char* name, int* value);
 }
@@ -47,6 +50,36 @@ using PFN_Evaluate = NgxResult (*)(ID3D12GraphicsCommandList*, const NgxHandle*,
 using PFN_Release = NgxResult (*)(NgxHandle*);
 using PFN_Shutdown = NgxResult (*)();
 struct NgxApplicationIdentifier
 {
    unsigned int type;
    union
    {
        unsigned long long applicationId;
        struct
        {
            const char* id;
            int engineType;
            const char* engineVersion;
        } project;
    } value;
 };
 struct NgxFeatureDiscovery
 {
    unsigned int sdkVersion;
    unsigned int feature;
    NgxApplicationIdentifier identifier;
    const wchar_t* dataPath;
    const void* featureInfo;
 };
 struct NgxFeatureRequirement
 {
    unsigned int supported;
    unsigned int minArchitecture;
    char minOsVersion[255];
 };
 using PFN_Requirements = NgxResult (*)(void*, const NgxFeatureDiscovery*, NgxFeatureRequirement*);
 static ID3D12Device* g_device;
 static ID3D12CommandQueue* g_queue;
 static ID3D12CommandAllocator* g_allocator;
@@ -618,6 +651,59 @@ int main(int argc, char** argv)
    }
    const auto srv = D3D12_RESOURCE_STATE_NON_PIXEL_SHADER_RESOURCE;
    const char* featureSetting = std::getenv("D4R_HARNESS_FEATURE_ID");
    const unsigned int featureId = featureSetting ? static_cast<unsigned int>(std::strtoul(featureSetting, nullptr, 0)) : 1;
    const bool reconstruction = featureId == 13;
    std::array<ID3D12Resource*, 4> guides{};
    if (reconstruction)
    {
        // Planar synthetic scene: world-space +Z normals, rough diffuse material.
        // Distinct guides exercise denoising inputs rather than aliasing the noisy colour.
        std::vector<uint16_t> guide(static_cast<size_t>(resourceWidth) * inHeight * 4);
        const std::array<std::array<float, 4>, 4> values{{
            {0.0f, 0.0f, 1.0f, 0.0f}, {0.5f, 0.0f, 0.0f, 0.0f},
            {0.7f, 0.4f, 0.2f, 1.0f}, {0.04f, 0.04f, 0.04f, 1.0f}}};
        for (size_t input = 0; input < guides.size(); ++input)
        {
            // Cyberpunk supplies UNORM8 albedos, while normals remain floating point.
            const bool unorm = input >= 2;
            guides[input] = create_texture(resourceWidth, inHeight,
                                           unorm ? DXGI_FORMAT_R8G8B8A8_UNORM : DXGI_FORMAT_R16G16B16A16_FLOAT,
                                           D3D12_RESOURCE_FLAG_NONE);
            if (!guides[input]) return 1;
            for (size_t pixel = 0; pixel < guide.size() / 4; ++pixel)
                for (size_t component = 0; component < 4; ++component)
                    if (unorm)
                        reinterpret_cast<uint8_t*>(guide.data())[pixel * 4 + component] =
                            static_cast<uint8_t>(std::lround(values[input][component] * 255.0f));
                    else
                        guide[pixel * 4 + component] =
                            std::bit_cast<uint16_t>(static_cast<_Float16>(values[input][component]));
            upload(guides[input], guide.data(), resourceWidth * (unorm ? 4 : 8), srv);
        }
    }
    const bool alphaScenario = reconstruction && std::getenv("D4R_HARNESS_RR_ALPHA") != nullptr;
    const UINT alphaBaseX = 4, alphaBaseY = 2;
    const UINT alphaWidth = outWidth + 8, alphaHeight = outHeight + 4;
    ID3D12Resource* alphaInput = nullptr;
    ID3D12Resource* alphaOutput = nullptr;
    if (alphaScenario)
    {
        // A varying coverage guide and a padded destination expose dropped alpha writes and
        // copies that overwrite pixels outside the caller's nonzero output subrect.
        std::vector<float> coverage(static_cast<size_t>(resourceWidth) * inHeight);
        for (UINT y = 0; y < inHeight; ++y)
            for (UINT x = 0; x < resourceWidth; ++x)
                coverage[static_cast<size_t>(y) * resourceWidth + x] =
                    0.25f + 0.5f * static_cast<float>(x) / std::max(1u, resourceWidth - 1);
        alphaInput = create_texture(resourceWidth, inHeight, DXGI_FORMAT_R32_FLOAT, D3D12_RESOURCE_FLAG_NONE);
        alphaOutput = create_texture(alphaWidth, alphaHeight, DXGI_FORMAT_R32_FLOAT,
                                     D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS);
        if (!alphaInput || !alphaOutput) return 1;
        upload(alphaInput, coverage.data(), resourceWidth * sizeof(float), srv);
        std::vector<float> sentinel(static_cast<size_t>(alphaWidth) * alphaHeight, -2.0f);
        upload(alphaOutput, sentinel.data(), alphaWidth * sizeof(float), D3D12_RESOURCE_STATE_UNORDERED_ACCESS);
    }
    ID3D12Resource* colorTexture = create_texture(resourceWidth, inHeight,
        rgba8 ? DXGI_FORMAT_R8G8B8A8_UNORM : DXGI_FORMAT_R16G16B16A16_FLOAT, D3D12_RESOURCE_FLAG_NONE);
    ID3D12Resource* depthTexture = create_texture(resourceWidth, inHeight, DXGI_FORMAT_R32_FLOAT, D3D12_RESOURCE_FLAG_NONE);
@@ -686,6 +772,32 @@ int main(int argc, char** argv)
    GetTempPathW(MAX_PATH, dataPath);
    const char* appIdSetting = std::getenv("D4R_HARNESS_APP_ID");
    const unsigned long long appId = appIdSetting != nullptr ? std::strtoull(appIdSetting, nullptr, 10) : 241534723ULL;
    // Games discover RR before Init; an initialized-only capability answer hides the mode.
    auto requirements = reinterpret_cast<PFN_Requirements>(
        reinterpret_cast<void*>(GetProcAddress(shim, "NVSDK_NGX_D3D12_GetFeatureRequirements")));
    NgxFeatureDiscovery discovery{};
    discovery.sdkVersion = 0x15;
    discovery.feature = reconstruction ? 13 : 1;
    discovery.identifier.value.applicationId = appId;
    discovery.dataPath = dataPath;
    NgxFeatureRequirement requirement{};
    using CreateFactoryFn = HRESULT (*)(REFIID, void**);
    const auto createFactory = reinterpret_cast<CreateFactoryFn>(
        reinterpret_cast<void*>(GetProcAddress(LoadLibraryA("dxgi.dll"), "CreateDXGIFactory1")));
    IDXGIFactory4* factory = nullptr;
    IDXGIAdapter* adapter = nullptr;
    LUID adapterLuid{};
    g_device->GetAdapterLuid(&adapterLuid);
    if (FAILED(createFactory(__uuidof(IDXGIFactory4), reinterpret_cast<void**>(&factory))) ||
        FAILED(factory->EnumAdapterByLuid(adapterLuid, __uuidof(IDXGIAdapter), reinterpret_cast<void**>(&adapter))))
        return 1;
    const NgxResult discoveryResult = requirements(adapter, &discovery, &requirement);
    adapter->Release();
    factory->Release();
    std::printf("GetFeatureRequirements before Init -> 0x%08x, feature=%u supported=%u\n",
                discoveryResult, discovery.feature, requirement.supported);
    if (discoveryResult != NGX_SUCCESS || requirement.supported != 0)
        return 1;
    NgxResult result = init(appId, dataPath, g_device, 0x15, nullptr);
    trace_stage(result == NGX_SUCCESS ? "NGX init succeeded" : "NGX init failed");
    std::printf("NVSDK_NGX_D3D12_Init_Ext -> 0x%08x\n", result);
@@ -694,9 +806,10 @@ int main(int argc, char** argv)
    void* capabilities = nullptr;
    result = capability(&capabilities);
    int available = -1;
    const char* availabilityKey = reconstruction ? "SuperSamplingDenoising.Available" : "SuperSampling.Available";
    if (result == NGX_SUCCESS)
-        d4r_ngx_get_int(capabilities, "SuperSampling.Available", &available);
+        d4r_ngx_get_int(capabilities, availabilityKey, &available);
-    std::printf("GetCapabilityParameters -> 0x%08x, SuperSampling.Available=%d\n", result, available);
+    std::printf("GetCapabilityParameters -> 0x%08x, %s=%d\n", result, availabilityKey, available);
    void* parameters = nullptr;
    result = allocate(&parameters);
@@ -716,13 +829,32 @@ int main(int argc, char** argv)
    d4r_ngx_set_int(parameters, "PerfQualityValue", quality != nullptr && *quality != '\0' ? std::atoi(quality) : 2);
    // D4R_HARNESS_CREATE_FLAGS mirrors a game's NVSDK_NGX_DLSS_Feature_Flags
    // (e.g. 0x49 = HDR | DepthInverted | AutoExposure).
    // Transformer RR requires HDR color and low-resolution motion vectors.
    const char* createFlags = std::getenv("D4R_HARNESS_CREATE_FLAGS");
    d4r_ngx_set_int(parameters, "DLSS.Feature.Create.Flags",
-                    createFlags != nullptr ? static_cast<int>(std::strtol(createFlags, nullptr, 0)) : 0);
+                    (createFlags != nullptr ? static_cast<int>(std::strtol(createFlags, nullptr, 0)) :
                                             (reconstruction ? (motionHighRes ? 1 : 3) : 0)) |
                    (alphaScenario ? 0x80 : 0));
    if (alphaScenario)
        d4r_ngx_set_int(parameters, "DLSS.Enable.Output.Subrects", 1);
    std::array<float, 16> identity{1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1};
    if (reconstruction)
    {
        d4r_ngx_set_int(parameters, "DLSS.Denoise.Mode", 1);
        d4r_ngx_set_int(parameters, "DLSS.Roughness.Mode", 0);
        d4r_ngx_set_int(parameters, "DLSS.Use.HW.Depth", 0);
        d4r_ngx_set_void(parameters, "WorldToViewMatrix", identity.data());
        d4r_ngx_set_void(parameters, "ViewToClipMatrix", identity.data());
        const char* preset = std::getenv("D4R_HARNESS_RR_PRESET");
        if (preset)
            for (const char* mode : {"DLAA", "Quality", "Balanced", "Performance", "UltraPerformance", "UltraQuality"})
                d4r_ngx_set_uint(parameters, (std::string("RayReconstruction.Hint.Render.Preset.") + mode).c_str(),
                                 static_cast<unsigned int>(std::strtoul(preset, nullptr, 0)));
    }
    NgxHandle* feature = nullptr;
    if (std::getenv("D4R_HARNESS_LIFETIME_TEST") != nullptr && !test_external_lifetime())
        return 1;
-    result = createFeature(g_list, 1, parameters, &feature);
+    result = createFeature(g_list, featureId, parameters, &feature);
    std::printf("CreateFeature -> 0x%08x handle=%p\n", result, static_cast<void*>(feature));
    if (result != NGX_SUCCESS)
        return 1;
@@ -732,6 +864,23 @@ int main(int argc, char** argv)
    d4r_ngx_set_d3d12_resource(parameters, "MotionVectors", motionTexture);
    d4r_ngx_set_d3d12_resource(parameters, "Output", outputTexture);
    d4r_ngx_set_d3d12_resource(parameters, "ExposureTexture", exposureTexture);
    if (reconstruction)
    {
        // The SDK's caller-facing names (nvsdk_ngx_defs.h: NVSDK_NGX_Parameter_GBuffer_Normals,
        // _Roughness, _DiffuseAlbedo, _SpecularAlbedo), which is what a game registers. The
        // shim republishes the two albedos under the names the denoiser reads.
        const char* names[] = {"GBuffer.Normals", "GBuffer.Roughness",
                               "GBuffer.DiffuseAlbedo", "GBuffer.SpecularAlbedo"};
        for (size_t input = 0; input < guides.size(); ++input)
            d4r_ngx_set_d3d12_resource(parameters, names[input], guides[input]);
    }
    if (alphaScenario)
    {
        d4r_ngx_set_d3d12_resource(parameters, "DLSSD.Alpha", alphaInput);
        d4r_ngx_set_d3d12_resource(parameters, "DLSSD.OutputAlpha", alphaOutput);
        d4r_ngx_set_uint(parameters, "DLSSD.OutputAlpha.Subrect.Base.X", alphaBaseX);
        d4r_ngx_set_uint(parameters, "DLSSD.OutputAlpha.Subrect.Base.Y", alphaBaseY);
    }
    d4r_ngx_set_float(parameters, "Jitter.Offset.X", 0.0f);
    d4r_ngx_set_float(parameters, "Jitter.Offset.Y", 0.0f);
    d4r_ngx_set_float(parameters, "MV.Scale.X", 1.0f);
@@ -989,6 +1138,13 @@ int main(int argc, char** argv)
                             1e6 * static_cast<double>(counterStart.QuadPart) / counterFrequency.QuadPart);
                std::fclose(timingFile);
            }
        // NGX rejecting an evaluation outright is reported here; only a failure raised later,
        // on the shim's worker, is invisible in this return value and shows up in the image.
        if (result != NGX_SUCCESS)
        {
            std::fprintf(stderr, "EvaluateFeature rejected frame %d: 0x%08x\n", frame, result);
            return 1;
        }
        Sleep(frameWaitMs); // let the CUDA worker finish before the next presentation
        if (qualityScene && !poleScene && frame > frames - 17)
        {
@@ -1086,6 +1242,122 @@ int main(int argc, char** argv)
    for (uint8_t value : output)
        nonzero += value != 0;
    std::printf("output read back: %zu of %zu bytes nonzero, written to %s\n", nonzero, output.size(), argv[2]);
    // The shim queues each evaluation on its own worker, so a failure inside the denoiser is
    // invisible in this thread's EvaluateFeature return value: the frame is accepted and the
    // output texture is simply never written. The image is the only observable. On the static
    // scene the denoiser reconstructs, so the output must carry that scene's structure:
    // compare block-averaged luma against the same average of the input. Averaging first makes
    // this about scene structure rather than per-pixel detail a denoiser is free to change. An
    // untouched output buffer (the initial fill) has no variance and cannot correlate.
    // Temporal scenes move content between the input and the frame presented, and replay
    // frames come from a capture, so neither is checked here.
    if (reconstruction && !rgba8 && replayDir == nullptr && !motionScene && !jitterScene && !qualityScene)
    {
        const uint16_t* const halves = reinterpret_cast<const uint16_t*>(output.data());
        constexpr UINT block = 8;
        const UINT blocksX = (outWidth + block - 1) / block, blocksY = (outHeight + block - 1) / block;
        std::vector<double> outBlocks, inBlocks;
        outBlocks.reserve(static_cast<size_t>(blocksX) * blocksY);
        inBlocks.reserve(static_cast<size_t>(blocksX) * blocksY);
        for (UINT by = 0; by < blocksY; ++by)
            for (UINT bx = 0; bx < blocksX; ++bx)
            {
                double outSum = 0.0, inSum = 0.0;
                size_t count = 0;
                for (UINT y = by * block; y < std::min(by * block + block, outHeight); ++y)
                    for (UINT x = bx * block; x < std::min(bx * block + block, outWidth); ++x)
                    {
                        const size_t outPixel = (static_cast<size_t>(y) * outWidth + x) * 4;
                        // The same output block mapped back onto the render grid: the denoiser
                        // upsamples exactly this scene, so no filter assumption is needed.
                        const UINT sourceY = static_cast<UINT>(static_cast<uint64_t>(y) * inHeight / outHeight);
                        const UINT sourceX = static_cast<UINT>(static_cast<uint64_t>(x) * inWidth / outWidth);
                        const size_t inPixel = (static_cast<size_t>(sourceY) * resourceWidth + sourceX) * 4;
                        for (int channel = 0; channel < 3; ++channel)
                        {
                            outSum += static_cast<double>(std::bit_cast<_Float16>(halves[outPixel + channel]));
                            inSum += static_cast<double>(std::bit_cast<_Float16>(color[inPixel + channel]));
                        }
                        ++count;
                    }
                outBlocks.push_back(outSum / (3.0 * static_cast<double>(count)));
                inBlocks.push_back(inSum / (3.0 * static_cast<double>(count)));
            }
        const double n = static_cast<double>(outBlocks.size());
        double sumOut = 0.0, sumIn = 0.0, sumOutOut = 0.0, sumInIn = 0.0, sumProduct = 0.0;
        for (size_t index = 0; index < outBlocks.size(); ++index)
        {
            sumOut += outBlocks[index];
            sumIn += inBlocks[index];
            sumOutOut += outBlocks[index] * outBlocks[index];
            sumInIn += inBlocks[index] * inBlocks[index];
            sumProduct += outBlocks[index] * inBlocks[index];
        }
        const double outVariance = sumOutOut - sumOut * sumOut / n;
        const double inVariance = sumInIn - sumIn * sumIn / n;
        const double correlation = (sumProduct - sumOut * sumIn / n) /
                                   (std::sqrt(outVariance) * std::sqrt(inVariance));
        std::printf("RR output block correlation with the input scene: %.6f over %zu blocks "
                    "(mean luma %.6f vs %.6f)\n",
                    correlation, outBlocks.size(), sumOut / n, sumIn / n);
        // A scene-structure floor, not a quality comparison with NVIDIA output. Unwritten
        // output has zero variance and produces NaN here; nonfinite or unrelated output also
        // fails. Denoising may change individual pixels without losing the scene's structure.
        if (!(correlation >= 0.5))
        {
            std::fprintf(stderr,
                         "Ray Reconstruction produced no usable image: block correlation with the "
                         "input scene is %.6f (unwritten, unrelated, or nonfinite output)\n",
                         correlation);
            return 1;
        }
    }
    if (alphaScenario)
    {
        const std::vector<uint8_t> raw =
            read_back(alphaOutput, D3D12_RESOURCE_STATE_UNORDERED_ACCESS, alphaWidth * sizeof(float));
        const float* values = reinterpret_cast<const float*>(raw.data());
        size_t invalid = 0, overwritten = 0, leftCount = 0, rightCount = 0;
        double left = 0, right = 0;
        for (UINT y = 0; y < alphaHeight; ++y)
            for (UINT x = 0; x < alphaWidth; ++x)
            {
                const float value = values[static_cast<size_t>(y) * alphaWidth + x];
                const bool inside = x >= alphaBaseX && x < alphaBaseX + outWidth &&
                                    y >= alphaBaseY && y < alphaBaseY + outHeight;
                if (!inside)
                {
                    overwritten += value != -2.0f;
                    continue;
                }
                invalid += !std::isfinite(value) || value < -0.001f || value > 1.001f;
                const UINT localX = x - alphaBaseX;
                if (localX < outWidth / 4)
                {
                    left += value;
                    ++leftCount;
                }
                if (localX >= outWidth * 3 / 4)
                {
                    right += value;
                    ++rightCount;
                }
            }
        const double leftMean = leftCount ? left / leftCount : 0;
        const double rightMean = rightCount ? right / rightCount : 0;
        std::printf("RR alpha: invalid=%zu outside-overwritten=%zu left=%.6f right=%.6f\n",
                    invalid, overwritten, leftMean, rightMean);
        if (FILE* alphaFile = std::fopen((std::string(argv[2]) + ".alpha.raw").c_str(), "wb"))
        {
            std::fwrite(raw.data(), 1, raw.size(), alphaFile);
            std::fclose(alphaFile);
        }
        if (invalid || overwritten || !leftCount || !rightCount || !(rightMean - leftMean > 0.1))
        {
            std::fprintf(stderr, "RR alpha coverage or subrect preservation failed\n");
            return 1;
        }
    }
    if (std::getenv("D4R_HARNESS_DISCARD_EVALUATION") != nullptr)
    {
@@ -1101,7 +1373,7 @@ int main(int argc, char** argv)
        g_list->Close();
        g_allocator->Reset();
        g_list->Reset(g_allocator, nullptr);
-        result = createFeature(g_list, 1, parameters, &feature);
+        result = createFeature(g_list, featureId, parameters, &feature);
        if (result != NGX_SUCCESS) return 1;
        lifecycle_report("DISCARD evaluation: release completed without submitting inputs");
    }
@@ -1121,7 +1393,7 @@ int main(int argc, char** argv)
        d4r_ngx_set_uint(parameters, "Height", height);
        d4r_ngx_set_uint(parameters, "DLSS.Render.Subrect.Dimensions.Width", width);
        d4r_ngx_set_uint(parameters, "DLSS.Render.Subrect.Dimensions.Height", height);
-        result = createFeature(g_list, 1, parameters, &feature);
+        result = createFeature(g_list, featureId, parameters, &feature);
        if (result != NGX_SUCCESS)
        {
            std::printf("RECREATE cycle %d: CreateFeature -> 0x%08x\n", cycle, result);
@@ -1167,5 +1439,9 @@ int main(int argc, char** argv)
    release(feature);
    shutdown();
    for (ID3D12Resource* guide : guides)
        if (guide) guide->Release();
    if (alphaInput) alphaInput->Release();
    if (alphaOutput) alphaOutput->Release();
    return 0;
 }
@@ -1,3 +1,4 @@
 // Modified in this fork for CUDA Ray Reconstruction support and validation (2026).
 #include <dlfcn.h>
 #include <cstddef>
@@ -85,10 +86,8 @@ static CUresult report(const char* stage, CUresult result,
    return result;
 }
-// One unformatted 4-byte store (sust.b.2d.v2.b16.zero, and the same bytes via
+// Unformatted packed stores address X in bytes. Check every byte, including the
-// sust.b.2d.b32.zero) into a surface of the given CUDA array format. X is a byte offset (unformatted surface addressing). The expected
+// untouched pixels around a write and row boundaries.
 // result is the fill pattern with the four stored bytes replaced at X, or the
 // untouched fill pattern when X is out of bounds.
 struct StoreCase
 {
    const char* name;
@@ -100,6 +99,10 @@ struct StoreCase
    uint16_t value0;
    uint16_t value1;
    uint8_t fill[16];     // one pixel of initial contents
    uint32_t bytes = 4;
    uint16_t value2 = 0;
    uint16_t value3 = 0;
    bool byteVector = false;
 };
 int main()
@@ -146,6 +149,44 @@ int main()
    sust.b.2d.b32.zero [%rd1, {%r0,%r1}], %r2;
    ret;
 }
 .visible .entry d4r_surface_v4_half_store_probe(
 .param .b64 surface,
 .param .b32 x,
 .param .b32 y,
 .param .b64 packed
 )
 {
    .reg .b64 %rd<3>;
    .reg .b16 %rs<4>;
    .reg .b32 %r<4>;
    ld.param.b64 %rd1, [surface];
    ld.param.b32 %r0, [x];
    ld.param.b32 %r1, [y];
    ld.param.b64 %rd2, [packed];
    mov.b64 {%rs0, %rs1, %rs2, %rs3}, %rd2;
    sust.b.2d.v4.b16.zero [%rd1, {%r0,%r1}], {%rs0,%rs1,%rs2,%rs3};
    ret;
 }
 .visible .entry d4r_surface_wide_byte_store_probe(
 .param .b64 surface,
 .param .b32 x,
 .param .b32 y,
 .param .b64 packed
 )
 {
    .reg .b64 %rd<3>;
    .reg .b16 %rs<4>;
    .reg .b32 %r<2>;
    ld.param.b64 %rd1, [surface];
    ld.param.b32 %r0, [x];
    ld.param.b32 %r1, [y];
    ld.param.b64 %rd2, [packed];
    mov.b64 {%rs0, %rs1, %rs2, %rs3}, %rd2;
    sust.b.2d.v4.b8.trap [%rd1, {%r0,%r1}], {%rs0,%rs1,%rs2,%rs3};
    ret;
 }
 )PTX";
    // Fill values: 0x3400 is half 0.25, 0x3e800000 is float 0.25.
@@ -180,6 +221,18 @@ int main()
         {0x00, 0x00, 0x80, 0x3e}},
        {"RG16F infinity/quiet NaN", 0x10, 2, 2, 13 * 4, 13, 0xfc00, 0x7e00,
         {0x00, 0x34, 0x00, 0x34}},
        {"RGBA16F four distinct halves", 0x10, 4, 2, 3 * 8, 2, 0x3555, 0xc248,
         {0x00, 0x34, 0x00, 0x34, 0x00, 0x34, 0x00, 0x34}, 8, 0x4900, 0xbc00},
        {"R16F four-pixel row end", 0x10, 1, 2, width * 2 - 8, height - 1, 0x0001, 0x7bff,
         {0x00, 0x34}, 8, 0x8000, 0x0400},
        {"RGBA16 uint full packed word", 0x02, 4, 2, 0, 0, 0xfffe, 0x1234,
         {0x11, 0x11, 0x22, 0x22, 0x33, 0x33, 0x44, 0x44}, 8, 0xabcd, 0x8001},
        {"RGBA16F eight-byte out-of-bounds", 0x10, 4, 2, width * 8, 7, 0x3555, 0xc248,
         {0x00, 0x34, 0x00, 0x34, 0x00, 0x34, 0x00, 0x34}, 8, 0x4900, 0xbc00},
        {"R8 low bytes from wide registers at row end", 0x01, 1, 1, width - 4, height - 1, 0xaaff, 0xbb80,
         {0x5a}, 4, 0xcc01, 0xdd7f, true},
        {"RGBA8 low bytes from wide registers", 0x01, 4, 1, 3 * 4, 2, 0x1200, 0x347f,
         {0x11, 0x22, 0x33, 0x44}, 4, 0x5680, 0x78ff, true},
    };
    void* library = dlopen("libcuda.so", RTLD_NOW | RTLD_LOCAL);
@@ -234,8 +287,9 @@ int main()
    CUcontext context = nullptr;
    CUmodule module = nullptr;
-    CUfunction functions[2] = {};
+    CUfunction functions[4] = {};
-    const char* functionNames[2] = {"sust.b.2d.v2.b16", "sust.b.2d.b32"};
+    const char* functionNames[4] = {"sust.b.2d.v2.b16", "sust.b.2d.b32", "sust.b.2d.v4.b16", "sust.b.2d.v4.b8.trap"};
    const uint32_t storeBytes[4] = {4, 4, 8, 4};
    CUresult result = report("cuInit", init(0), getErrorString);
    CUdevice device = 0;
    if (result == 0)
@@ -250,12 +304,21 @@ int main()
    if (result == 0)
        result = report("cuModuleGetFunction(d4r_surface_b32_store_probe)",
                        moduleGetFunction(&functions[1], module, "d4r_surface_b32_store_probe"), getErrorString);
    if (result == 0)
        result = report("cuModuleGetFunction(d4r_surface_v4_half_store_probe)",
                        moduleGetFunction(&functions[2], module, "d4r_surface_v4_half_store_probe"), getErrorString);
    if (result == 0)
        result = report("cuModuleGetFunction(d4r_surface_wide_byte_store_probe)",
                        moduleGetFunction(&functions[3], module, "d4r_surface_wide_byte_store_probe"), getErrorString);
    size_t passed = 0;
-    size_t total = 2 * sizeof(cases) / sizeof(cases[0]);
+    size_t total = 0;
-    for (size_t functionIndex = 0; functionIndex < 2; ++functionIndex)
+    for (size_t functionIndex = 0; functionIndex < 4; ++functionIndex)
    for (const StoreCase& storeCase : cases)
    {
        if (storeCase.bytes != storeBytes[functionIndex] || storeCase.byteVector != (functionIndex == 3))
            continue;
        ++total;
        CUfunction function = functions[functionIndex];
        if (result != 0)
            break;
@@ -265,11 +328,20 @@ int main()
        for (size_t offset = 0; offset < expected.size(); ++offset)
            expected[offset] = storeCase.fill[offset % pixelBytes];
        std::vector<uint8_t> surfaceBytes = expected;
-        const uint8_t stored[4] = {
+        uint8_t stored[8] = {
            static_cast<uint8_t>(storeCase.value0), static_cast<uint8_t>(storeCase.value0 >> 8),
-            static_cast<uint8_t>(storeCase.value1), static_cast<uint8_t>(storeCase.value1 >> 8)};
+            static_cast<uint8_t>(storeCase.value1), static_cast<uint8_t>(storeCase.value1 >> 8),
-        if (storeCase.x + 4 <= rowBytes && storeCase.y < height)
+            static_cast<uint8_t>(storeCase.value2), static_cast<uint8_t>(storeCase.value2 >> 8),
-            std::memcpy(expected.data() + storeCase.y * rowBytes + storeCase.x, stored, 4);
+            static_cast<uint8_t>(storeCase.value3), static_cast<uint8_t>(storeCase.value3 >> 8)};
        if (storeCase.byteVector)
        {
            stored[0] = static_cast<uint8_t>(storeCase.value0);
            stored[1] = static_cast<uint8_t>(storeCase.value1);
            stored[2] = static_cast<uint8_t>(storeCase.value2);
            stored[3] = static_cast<uint8_t>(storeCase.value3);
        }
        if (static_cast<uint64_t>(storeCase.x) + storeCase.bytes <= rowBytes && storeCase.y < height)
            std::memcpy(expected.data() + storeCase.y * rowBytes + storeCase.x, stored, storeCase.bytes);
        CUarray array = nullptr;
        CUsurfObject surface = 0;
@@ -300,7 +372,10 @@ int main()
            uint64_t surfaceArgument = surface;
            uint32_t x = storeCase.x;
            uint32_t y = storeCase.y;
-            uint32_t packed = storeCase.value0 | (static_cast<uint32_t>(storeCase.value1) << 16);
+            uint64_t packed = static_cast<uint64_t>(storeCase.value0) |
                              (static_cast<uint64_t>(storeCase.value1) << 16) |
                              (static_cast<uint64_t>(storeCase.value2) << 32) |
                              (static_cast<uint64_t>(storeCase.value3) << 48);
            void* arguments[] = {&surfaceArgument, &x, &y, &packed};
            caseResult = launchKernel(function, 1, 1, 1, 1, 1, 1, 0, nullptr, arguments, nullptr);
            if (caseResult == 0)
@@ -344,7 +419,7 @@ int main()
    }
    if (result == 0)
-        std::printf("sust.b.2d unformatted 4-byte surface store tests: %zu/%zu passed\n", passed, total);
+        std::printf("sust.b.2d unformatted surface store tests: %zu/%zu passed\n", passed, total);
    const int exitCode = result == 0 && passed == total ? 0 : 1;
    if (module != nullptr)
        moduleUnload(module);
@@ -1,3 +1,4 @@
 // Modified in this fork for CUDA Ray Reconstruction support and validation (2026).
 #include <dlfcn.h>
 #include <algorithm>
@@ -123,7 +124,8 @@ int main()
 {
    .reg .b64 %rd<5>;
    .reg .b32 %r<5>;
-    .reg .f32 %f<18>;
+    .reg .f32 %f<22>;
    .reg .b16 %h<4>;
    ld.param.b64 %rd1, [texture_object];
    ld.param.b64 %rd2, [output_pointer];
    mov.u32 %r1, %tid.x;
@@ -138,13 +140,19 @@ int main()
    mul.f32 %f9, %f9, 0f3E000000;
    tex.base.2d.v4.f32.f32 {%f10,%f11,%f12,%f13}, [%rd1, {%f8,%f9}];
    tex.level.2d.v4.f32.f32 {%f14,%f15,%f16,%f17}, [%rd1, {%f8,%f9}], 0f00000000;
    tex.base.2d.v4.f16.f32 {%h0,%h1,%h2,%h3}, [%rd1, {%f8,%f9}];
    cvt.f32.f16 %f18, %h0;
    cvt.f32.f16 %f19, %h1;
    cvt.f32.f16 %f20, %h2;
    cvt.f32.f16 %f21, %h3;
    mad.lo.u32 %r3, %r2, 8, %r1;
-    mul.wide.u32 %rd3, %r3, 64;
+    mul.wide.u32 %rd3, %r3, 80;
    add.s64 %rd4, %rd2, %rd3;
    st.global.v4.f32 [%rd4], {%f0,%f1,%f2,%f3};
    st.global.v4.f32 [%rd4+16], {%f4,%f5,%f6,%f7};
    st.global.v4.f32 [%rd4+32], {%f10,%f11,%f12,%f13};
    st.global.v4.f32 [%rd4+48], {%f14,%f15,%f16,%f17};
    st.global.v4.f32 [%rd4+64], {%f18,%f19,%f20,%f21};
    ret;
 }
@@ -250,7 +258,7 @@ int main()
        const uint16_t mantissa = static_cast<uint16_t>((fraction * 2.0f - 1.0f) * 1024.0f);
        return static_cast<uint16_t>(((exponent - 1 + 15) << 10) | mantissa);
    };
-    const size_t outputFloats = static_cast<size_t>(width) * height * 16;
+    const size_t outputFloats = static_cast<size_t>(width) * height * 20;
    if (result == 0)
        result = report("cuMemAlloc_v2(output)", memAlloc(&deviceOutput, outputFloats * sizeof(float)), getErrorString);
    if (result == 0)
@@ -364,10 +372,10 @@ int main()
            size_t mismatches = 0;
            for (unsigned int y = 0; status == 0 && y < height; ++y)
                for (unsigned int x = 0; x < width; ++x)
-                    for (unsigned int form = 0; form < (samplerCase.flags == 2 ? 4u : 2u); ++form)
+                    for (unsigned int form = 0; form < (samplerCase.flags == 2 ? 5u : 2u); ++form)
                        for (unsigned int channel = 0; channel < formatCase.channels; ++channel)
                        {
-                            const float got = sampled[(y * width + x) * 16 + form * 4 + channel];
+                            const float got = sampled[(y * width + x) * 20 + form * 4 + channel];
                            float want = expected_texel(x, y)[channel];
                            if (form >= 2 && samplerCase.filterMode == 1)
                            {
@@ -378,10 +386,10 @@ int main()
                                want = 0.75f * (0.25f * topLeft + 0.75f * topRight)
                                     + 0.25f * (0.25f * bottomLeft + 0.75f * bottomRight);
                            }
-                            if (std::fabs(got - want) > 1e-6f && mismatches++ < 4)
+                            if ((!std::isfinite(got) || std::fabs(got - want) > 1e-6f) && mismatches++ < 4)
                                std::printf("  %s %s %s (%u,%u) channel %u: got %g expected %g\n", formatCase.name,
                                            samplerCase.name,
-                                            form == 0 ? "tex.base" : form == 1 ? "tex" : form == 2 ? "tex.base normalized" : "tex.level LOD0",
+                                            form == 0 ? "tex.base" : form == 1 ? "tex" : form == 2 ? "tex.base normalized" : form == 3 ? "tex.level LOD0" : "tex.base f16 normalized",
                                            x, y, channel, got, want);
                        }
            const bool casePass = status == 0 && mismatches == 0;
@@ -436,12 +444,12 @@ int main()
        size_t mismatches = 0;
        for (unsigned int y = 0; status == 0 && y < height; ++y)
            for (unsigned int x = 0; x < width; ++x)
-                for (unsigned int form = 0; form < 4; ++form)
+                for (unsigned int form = 0; form < 5; ++form)
                    for (unsigned int channel = 0; channel < 4; ++channel)
                    {
                        const float expected = x == 3 && y == 4 ? expectedStored[channel]
                                                                  : expected_texel(x, y)[channel];
-                        const float got = sampled[(y * width + x) * 16 + form * 4 + channel];
+                        const float got = sampled[(y * width + x) * 20 + form * 4 + channel];
                        if (got != expected && mismatches++ < 5)
                            std::printf("  surface->texture (%u,%u) form %u channel %u: got %g expected %g\n",
                                        x, y, form, channel, got, expected);