diff --git a/NOTICE b/NOTICE index 1837ac3..2a3cebb 100644 --- a/NOTICE +++ b/NOTICE @@ -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 -the NGX core (_nvngx.dll) are NVIDIA's; you supply your own copies, and their use is governed by +This repository contains no NVIDIA software. The NVIDIA DLSS feature libraries (nvngx_dlss.dll and +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): diff --git a/README.md b/README.md index 90401f8..029e4fe 100644 --- a/README.md +++ b/README.md @@ -1,14 +1,17 @@ + # 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. diff --git a/SUPPORTED_GAMES.md b/SUPPORTED_GAMES.md index e7aa511..f58eda8 100644 --- a/SUPPORTED_GAMES.md +++ b/SUPPORTED_GAMES.md @@ -1,6 +1,7 @@ + # 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. diff --git a/config/d4r.ini.default b/config/d4r.ini.default index 734fbcb..d5ad753 100644 --- a/config/d4r.ini.default +++ b/config/d4r.ini.default @@ -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- (a new ZLUDA build compiles every kernel once on first use) diff --git a/docs/architecture.md b/docs/architecture.md index 3cc797f..bc9897e 100644 --- a/docs/architecture.md +++ b/docs/architecture.md @@ -1,3 +1,4 @@ + # 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\\` | 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). diff --git a/docs/building.md b/docs/building.md index a48dad9..d554de4 100644 --- a/docs/building.md +++ b/docs/building.md @@ -1,3 +1,4 @@ + # 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 `-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/`; +- 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/`; 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 diff --git a/packaging/D4R_README.txt b/packaging/D4R_README.txt index 6a019ff..ba3c06d 100644 --- a/packaging/D4R_README.txt +++ b/packaging/D4R_README.txt @@ -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 -DirectX 12 games under Proton. OptiScaler (included) catches the game's DLSS calls and hands them to -d4r. d4r then runs NVIDIA's DLSS on the Radeon through ZLUDA, with the heaviest parts replaced by -kernels written for RDNA3 and RDNA4. +d4r runs NVIDIA's own DLSS Super Resolution (DLSS 4 by default), plus experimental Ray Reconstruction, +on AMD Radeon GPUs in DirectX 12 games under Proton. OptiScaler (included) catches the game's DLSS calls +and hands them to d4r. d4r runs NVIDIA's DLSS through ZLUDA; the heaviest Super Resolution kernels are +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. -SUPPORTED_GAMES.md in the source repository lists them, with per-game setup notes. Expect problems -in other games. +This is an early release. Super Resolution game testing primarily uses Radeon RX 7700 XT. +Ray Reconstruction presets D/E, temporal history, and separate alpha were validated on RX 9070 XT +(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. diff --git a/packaging/RELEASE_NOTES.md b/packaging/RELEASE_NOTES.md index 49a5d4c..adb698e 100644 --- a/packaging/RELEASE_NOTES.md +++ b/packaging/RELEASE_NOTES.md @@ -1,3 +1,4 @@ + # 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**. diff --git a/packaging/SOURCES.txt b/packaging/SOURCES.txt index 3de2044..f05c135 100644 --- a/packaging/SOURCES.txt +++ b/packaging/SOURCES.txt @@ -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), diff --git a/packaging/d4r.ini b/packaging/d4r.ini index bb17512..b715539 100644 --- a/packaging/d4r.ini +++ b/packaging/d4r.ini @@ -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). diff --git a/patches/zluda/0008-ray-reconstruction-ptx.patch b/patches/zluda/0008-ray-reconstruction-ptx.patch new file mode 100644 index 0000000..25e0fdf --- /dev/null +++ b/patches/zluda/0008-ray-reconstruction-ptx.patch @@ -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//build// ++ // /out layout more than once (current Cargo emits build///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(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( ++ static_cast(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( ++ static_cast(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(__ockl_image_width_2D(image)) * ++ static_cast(channel_bytes * channels); ++ if (channel_bytes == 0 || coord.x < 0 || coord.y < 0 || ++ coord.y >= __ockl_image_height_2D(image) || ++ static_cast(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, + ) -> 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 { + 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: { ++ 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: { + 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, + 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 ToReplace; + Instruction *InsertPoint = I; +- DenseMap InstOrdering; +- unsigned long InstNumber = 0; +- for (const auto &I : *TargetBB) +- InstOrdering[&I] = InstNumber++; + + for (Use *U : reverse(OpsToSink)) { + auto *UI = cast(U->get()); + if (isa(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 &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(Inst)) +@@ -104,9 +99,12 @@ + SmallPtrSetImpl &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"); diff --git a/scripts/build_release_glibc241.sh b/scripts/build_release_glibc241.sh index a9781ae..d6c8ec0 100755 --- a/scripts/build_release_glibc241.sh +++ b/scripts/build_release_glibc241.sh @@ -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 diff --git a/scripts/d4r_config.py b/scripts/d4r_config.py index 5bc166b..d850488 100755 --- a/scripts/d4r_config.py +++ b/scripts/d4r_config.py @@ -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"): diff --git a/scripts/install_d4r_runtime.sh b/scripts/install_d4r_runtime.sh index 1484bcf..dc88d1b 100755 --- a/scripts/install_d4r_runtime.sh +++ b/scripts/install_d4r_runtime.sh @@ -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 -if [[ $# -ne 2 ]]; then - printf 'usage: %s PATH_TO_NGX_CORE_DLL PATH_TO_NVNGX_DLSS_DLL\n' "$0" >&2 +# usage: install_d4r_runtime.sh PATH_TO_NGX_CORE_DLL PATH_TO_NVNGX_DLSS_DLL [PATH_TO_NVNGX_DLSSD_DLL] +if [[ $# -lt 2 || $# -gt 3 ]]; then + 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" diff --git a/scripts/package_release.sh b/scripts/package_release.sh index 154ebb5..9134165 100755 --- a/scripts/package_release.sh +++ b/scripts/package_release.sh @@ -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" diff --git a/tests/test_ray_reconstruction.py b/tests/test_ray_reconstruction.py new file mode 100644 index 0000000..cfe4a39 --- /dev/null +++ b/tests/test_ray_reconstruction.py @@ -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() \ No newline at end of file diff --git a/tools/d3d12_dlss_harness.cpp b/tools/d3d12_dlss_harness.cpp index cafb0cc..a4a8fe3 100644 --- a/tools/d3d12_dlss_harness.cpp +++ b/tools/d3d12_dlss_harness.cpp @@ -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 #include +#include #include #include @@ -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(std::strtoul(featureSetting, nullptr, 0)) : 1; + const bool reconstruction = featureId == 13; + std::array 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 guide(static_cast(resourceWidth) * inHeight * 4); + const std::array, 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(guide.data())[pixel * 4 + component] = + static_cast(std::lround(values[input][component] * 255.0f)); + else + guide[pixel * 4 + component] = + std::bit_cast(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 coverage(static_cast(resourceWidth) * inHeight); + for (UINT y = 0; y < inHeight; ++y) + for (UINT x = 0; x < resourceWidth; ++x) + coverage[static_cast(y) * resourceWidth + x] = + 0.25f + 0.5f * static_cast(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 sentinel(static_cast(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( + reinterpret_cast(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( + reinterpret_cast(GetProcAddress(LoadLibraryA("dxgi.dll"), "CreateDXGIFactory1"))); + IDXGIFactory4* factory = nullptr; + IDXGIAdapter* adapter = nullptr; + LUID adapterLuid{}; + g_device->GetAdapterLuid(&adapterLuid); + if (FAILED(createFactory(__uuidof(IDXGIFactory4), reinterpret_cast(&factory))) || + FAILED(factory->EnumAdapterByLuid(adapterLuid, __uuidof(IDXGIAdapter), reinterpret_cast(&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); - std::printf("GetCapabilityParameters -> 0x%08x, SuperSampling.Available=%d\n", result, available); + d4r_ngx_get_int(capabilities, availabilityKey, &available); + std::printf("GetCapabilityParameters -> 0x%08x, %s=%d\n", result, availabilityKey, available); void* parameters = nullptr; result = allocate(¶meters); @@ -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(std::strtol(createFlags, nullptr, 0)) : 0); + (createFlags != nullptr ? static_cast(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 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(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(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(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(output.data()); + constexpr UINT block = 8; + const UINT blocksX = (outWidth + block - 1) / block, blocksY = (outHeight + block - 1) / block; + std::vector outBlocks, inBlocks; + outBlocks.reserve(static_cast(blocksX) * blocksY); + inBlocks.reserve(static_cast(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(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(static_cast(y) * inHeight / outHeight); + const UINT sourceX = static_cast(static_cast(x) * inWidth / outWidth); + const size_t inPixel = (static_cast(sourceY) * resourceWidth + sourceX) * 4; + for (int channel = 0; channel < 3; ++channel) + { + outSum += static_cast(std::bit_cast<_Float16>(halves[outPixel + channel])); + inSum += static_cast(std::bit_cast<_Float16>(color[inPixel + channel])); + } + ++count; + } + outBlocks.push_back(outSum / (3.0 * static_cast(count))); + inBlocks.push_back(inSum / (3.0 * static_cast(count))); + } + const double n = static_cast(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 raw = + read_back(alphaOutput, D3D12_RESOURCE_STATE_UNORDERED_ACCESS, alphaWidth * sizeof(float)); + const float* values = reinterpret_cast(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(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; } diff --git a/tools/d4r_nvngx_shim.cpp b/tools/d4r_nvngx_shim.cpp index 6be7d51..cf6e92c 100644 --- a/tools/d4r_nvngx_shim.cpp +++ b/tools/d4r_nvngx_shim.cpp @@ -1,4 +1,5 @@ // Experimental native Vulkan NGX frontend: d4r_vulkan_frontend.inl. +// Modified in this fork for CUDA Ray Reconstruction support and validation (2026). // d4r_nvngx.dll: an NGX core replacement that implements the D3D12 DLSS // entry points on top of the official NGX core's CUDA path (which runs on AMD // through the Wine nvcuda bridge and ZLUDA). OptiScaler loads it through its @@ -38,6 +39,16 @@ // D4R_SHIM_INPLACE=1 experimental, with VKD3D_D4R_LINEAR_TEXTURES=1 in the patched vkd3d-proton: the output // kernel writes the game's output texture itself (no game-side output copy). // D4R_SHIM_OUTPUT_PACKED=1 lets M's output kernel store R10G10B10A2 texels itself. +// (the DLSS-Denoiser library nvngx_dlssd.dll, which serves Ray Reconstruction, must sit +// next to this DLL: that is where the NGX core looks for it. It needs +// a build exporting NVSDK_NGX_CUDA_* (3.10.7); without it Ray +// Reconstruction reports itself unavailable and Super Sampling is +// unaffected.) +// D4R_RR_ENABLE=1/0 Ask for Ray Reconstruction / forbid it. Unset (auto) offers it +// exactly when the denoiser loaded and initialised. +// D4R_RR_PRESET Force the Ray Reconstruction render preset +// (NVSDK_NGX_RayReconstruction_Hint_Render_Preset: 4 = D, 5 = E). +// Independent of D4R_DLSS_PRESET, which only affects Super Sampling. #define WIDL_EXPLICIT_AGGREGATE_RETURNS #include #include @@ -57,6 +68,7 @@ #include #include #include +#include #include #include #include @@ -82,6 +94,9 @@ constexpr NgxResult NGX_FAIL_NOT_INITIALIZED = NGX_FAIL | 7; constexpr NgxResult NGX_FAIL_UNSUPPORTED_FORMAT = NGX_FAIL | 14; constexpr unsigned int NGX_FEATURE_SUPER_SAMPLING = 1; +// nvngx_dlssd.dll (DLSS-Denoiser / Ray Reconstruction), NVSDK_NGX_Feature. +constexpr unsigned int NGX_FEATURE_RAY_RECONSTRUCTION = 13; + struct NgxHandle { unsigned int Id; @@ -231,6 +246,7 @@ struct CudaApi int(WINAPI* surfObjectCreate)(CudaObject*, const CudaResourceDesc*) = nullptr; int(WINAPI* surfObjectDestroy)(CudaObject) = nullptr; int(WINAPI* ctxSynchronize)() = nullptr; + int(WINAPI* ctxGetDevice)(int*) = nullptr; int(WINAPI* eventCreate)(CudaEvent*, unsigned int) = nullptr; int(WINAPI* eventRecord)(CudaEvent, void*) = nullptr; int(WINAPI* eventQuery)(CudaEvent) = nullptr; @@ -272,6 +288,33 @@ struct CoreApi NgxResult (*createFeature)(unsigned int, void*, NgxHandle**) = nullptr; NgxResult (*releaseFeature)(NgxHandle*) = nullptr; NgxResult (*evaluateFeature)(const NgxHandle*, void*, void*) = nullptr; + // The core's CUDA entry point takes a CUdevice ordinal, not an IDXGIAdapter* (its + // disassembly rejects a negative value in the first argument, which is the ordinal + // check). The D3D12 export this file implements takes an adapter and must translate. + NgxResult (*getFeatureRequirements)(int, const void*, NgxFeatureRequirement*) = nullptr; +}; + +// --- DLSS-Denoiser (Ray Reconstruction) ---------------------------------------- +// +// nvngx_dlssd.dll carries the Ray Reconstruction models, and the NGX core loads and +// dispatches it itself: the core is the only thing that knows the denoiser's module name, +// its kernel map and its lifetime. So both features are created, evaluated and released +// through g.ngx, with feature id 13 for Ray Reconstruction, and this file's job is only to +// answer honestly whether that route exists. +// +// Existence is decided by loading the denoiser and looking for its CUDA entry points, the +// only reliable external evidence - builds genuinely differ. The 3.10.7 build exports the +// whole NVSDK_NGX_CUDA_* set; the DLAA-only build exports D3D11/D3D12/Vulkan/DirectSR and +// no CUDA entry points, and cannot serve Ray Reconstruction however it is loaded. +// +// The pointers here are resolved only as that evidence and are never called. In particular +// NVSDK_NGX_CUDA_Init in the 310.7 build is `mov eax, 0xBAD00001; ret` (RVA 0x22b80), a stub +// that only ever reports "not supported"; the core's own init is what brings the denoiser up. +// A single yes/no, not a table of entry points: nothing here is ever called. The denoiser is +// loaded by the NGX core, which is the only thing that may drive it. +struct DenoiserProbe +{ + bool cudaCapable = false; }; // --- logging ----------------------------------------------------------------- @@ -527,6 +570,45 @@ static void load_portable_config() if (!model.empty()) portable_set("D4R_DLSS_PRESET", model); + // Ray Reconstruction is configured on its own section and drives its own feature, so + // nothing here reaches [DLSS] Model and nothing there reaches this. + // + // Enable: true -> D4R_RR_ENABLE=1, false -> 0, auto/unset -> the variable is left alone, + // which is the same thing as auto here: Ray Reconstruction is offered exactly when a + // CUDA-capable denoiser really loaded and initialised. + const std::string rrEnable = ascii_lower(ini_value(ini, "rayreconstruction", "Enable")); + if (rrEnable == "true" || rrEnable == "on" || rrEnable == "1") + portable_set("D4R_RR_ENABLE", "1"); + else if (rrEnable == "false" || rrEnable == "off" || rrEnable == "0") + portable_set("D4R_RR_ENABLE", "0"); + else if (!rrEnable.empty() && rrEnable != "auto") + g_portable.notes.push_back("d4r.ini: [RayReconstruction] Enable must be true, false or auto, not '" + + rrEnable + "'; treating it as auto"); + + // Model: auto/unset leaves the choice to the game; "default" is the SDK's Default preset + // (0); D (4) and E (5) are the two the SDK gives a meaning to. F and above are + // documented as reverting to default behaviour, so they are not offered. E is the model + // to pick when the depth-of-field guide is in use. These are the same values + // scripts/d4r_config.py accepts, so both ways of reading d4r.ini agree. + const std::string rrModel = ascii_lower(ini_value(ini, "rayreconstruction", "Model")); + if (!rrModel.empty() && rrModel != "auto" && rrModel != "game") + { + std::string upper = rrModel; + for (char& c : upper) + c = static_cast(toupper(static_cast(c))); + if (upper == "DEFAULT") + portable_set("D4R_RR_PRESET", "0"); + else if (upper == "D") + portable_set("D4R_RR_PRESET", "4"); + else if (upper == "E") + portable_set("D4R_RR_PRESET", "5"); + else if (upper == "0" || upper == "4" || upper == "5") + portable_set("D4R_RR_PRESET", upper); + else + g_portable.notes.push_back("d4r.ini: [RayReconstruction] Model must be auto, default, D or E, not '" + + rrModel + "'; leaving the model to the game"); + } + const std::string age = ini_value(ini, "latency", "FrameAge"); if (age.empty() || age == "0") { @@ -609,6 +691,12 @@ static void load_portable_config() if (!file_exists(dir + L"\\nvngx_dlss.dll")) g_portable.notes.push_back("nvngx_dlss.dll is missing: copy NVIDIA's DLSS library (310.7 or 310.9 recommended) " "into the d4r folder"); + // Only a note: without the denoiser Super Sampling is unaffected and Ray Reconstruction + // simply reports itself unavailable, which is the same thing the capability answer says. + if (ini_flag(ini, "rayreconstruction", "Enable", 1) && !file_exists(dir + L"\\nvngx_dlssd.dll")) + g_portable.notes.push_back("nvngx_dlssd.dll is missing: copy NVIDIA's DLSS-Denoiser library (3.10.7, the " + "build that exports NVSDK_NGX_CUDA_*) into the d4r folder to get Ray " + "Reconstruction"); wchar_t corePath[MAX_PATH] = {}; const DWORD coreLength = GetEnvironmentVariableW(L"D4R_NGX_CORE", corePath, MAX_PATH); if (coreLength > 0 && coreLength < MAX_PATH && !file_exists(corePath)) @@ -1241,6 +1329,11 @@ struct Global NgxFeatureCommonInfo featureInfo = {}; std::wstring dataPath; std::atomic nextHandleId{0x7200}; + // Ray Reconstruction: nvngx_dlssd.dll, loaded and initialised only when asked for. + HMODULE denoiserModule = nullptr; + DenoiserProbe denoiser; + bool denoiserAttempted = false; + std::wstring denoiserPath; }; static Global g; @@ -1381,6 +1474,9 @@ static bool load_libraries() ok &= load_export(g.core, "NVSDK_NGX_CUDA_CreateFeature", g.ngx.createFeature); ok &= load_export(g.core, "NVSDK_NGX_CUDA_ReleaseFeature", g.ngx.releaseFeature); ok &= load_export(g.core, "NVSDK_NGX_CUDA_EvaluateFeature", g.ngx.evaluateFeature); + // Optional: the core answers feature 13 by consulting the denoiser it loaded. + if (!load_export(g.core, "NVSDK_NGX_CUDA_GetFeatureRequirements", g.ngx.getFeatureRequirements)) + g.ngx.getFeatureRequirements = nullptr; g.cuda = LoadLibraryA("nvcuda.dll"); if (g.cuda == nullptr) @@ -1399,6 +1495,8 @@ static bool load_libraries() ok &= load_export(g.cuda, "cuSurfObjectCreate", g.cu.surfObjectCreate); ok &= load_export(g.cuda, "cuSurfObjectDestroy", g.cu.surfObjectDestroy); ok &= load_export(g.cuda, "cuCtxSynchronize", g.cu.ctxSynchronize); + if (!load_export(g.cuda, "cuCtxGetDevice", g.cu.ctxGetDevice)) + g.cu.ctxGetDevice = nullptr; // The d4r bridge's own export keeps waiting when the application's syncs are elided // (D4R_ELIDE_NGX_SYNC); cuCtxSynchronize above stays the fallback. { @@ -1447,6 +1545,131 @@ static bool load_libraries() return ok; } +// Loads nvngx_dlssd.dll and looks for the CUDA entry points that mark a build able to serve +// Ray Reconstruction. Nothing here calls into the denoiser: the NGX core loads and +// dispatches it, so this is a presence-and-capability check, not a second lifecycle. +// Without it Super Sampling is unaffected and Ray Reconstruction reports itself unavailable. +// +// The 3.10.7 build of this library carries the whole snippet stack and exports +// NVSDK_NGX_CUDA_*. Other builds - notably the DLAA-only "rel_tot" one, which ships D3D11, +// D3D12, Vulkan and DirectSR and nothing else - do not, and are rejected here rather than +// half-supported: a missing entry point means there is no CUDA path to bridge to. +static bool load_denoiser_library() +{ + if (g.denoiserAttempted) + return g.denoiser.cudaCapable; + g.denoiserAttempted = true; + ensure_portable_config(); + + // The core finds nvngx_dlssd.dll beside the DLL that called into it and ignores the + // feature search paths, so the only place worth looking is here: offering a path that + // could not redirect the core's own load would be claiming a selection it does not have. + wchar_t self[MAX_PATH] = {}; + const DWORD selfLength = GetModuleFileNameW(g_selfModule, self, MAX_PATH); + if (selfLength == 0 || selfLength >= MAX_PATH) + { + logf("Ray Reconstruction: cannot locate this DLL, so the denoiser cannot be found"); + return false; + } + std::wstring directory(self, selfLength); + directory.resize(directory.find_last_of(L"\\/")); + g.denoiserPath = directory + L"\\nvngx_dlssd.dll"; + if (!file_exists(g.denoiserPath)) + { + logf("Ray Reconstruction: no denoiser at %ls; the NGX core loads it from beside this DLL, so copy " + "NVIDIA's nvngx_dlssd.dll (3.10.7) there", + g.denoiserPath.c_str()); + return false; + } + g.denoiserModule = load_library_below_hooks(g.denoiserPath.c_str()); + if (g.denoiserModule == nullptr) + { + logf("Ray Reconstruction: loading %ls failed: %lu", g.denoiserPath.c_str(), GetLastError()); + return false; + } + // Which entry points exist is only ever evidence - none of them is called from here. + // Create/Evaluate/Release is what "can serve Ray Reconstruction" means. NVSDK_NGX_CUDA_Init + // is deliberately not required and is never called: in the 310.7 build it is a stub that + // always reports 0xBAD00001. + void* create = nullptr, *evaluate = nullptr, *release = nullptr; + load_export(g.denoiserModule, "NVSDK_NGX_CUDA_CreateFeature", create); + load_export(g.denoiserModule, "NVSDK_NGX_CUDA_EvaluateFeature", evaluate); + load_export(g.denoiserModule, "NVSDK_NGX_CUDA_ReleaseFeature", release); + g.denoiser.cudaCapable = create != nullptr && evaluate != nullptr && release != nullptr; + wchar_t mapped[MAX_PATH] = {}; + GetModuleFileNameW(g.denoiserModule, mapped, MAX_PATH); + logf("found the DLSS-Denoiser library %ls (mapped as %ls): %s", g.denoiserPath.c_str(), mapped, + g.denoiser.cudaCapable ? "CUDA entry points present" + : "no CUDA entry points - not a CUDA-capable denoiser build"); + return g.denoiser.cudaCapable; +} + +// D4R_RR_ENABLE: 1 asks for Ray Reconstruction, 0 forbids it, and unset (auto) offers it +// exactly when the denoiser really loaded. Auto is the default because the only honest answer +// to a game asking whether Ray Reconstruction is available is what this process can actually +// run - never a fixed yes. +static bool rr_enabled() +{ + return env_uint("D4R_RR_ENABLE", 1) != 0; +} + +// What the NGX core last said about the denoiser it loaded. The core brings the denoiser up +// during its own init and reports the outcome here; a file merely existing on disk proves +// nothing, because the core can come up perfectly well for Super Sampling with a denoiser +// that failed to initialise. +struct DenoiserCapability +{ + bool known = false; + int available = 0; + int initResult = NGX_FAIL_NOT_INITIALIZED; +}; +static DenoiserCapability g_denoiserCapability; +static std::mutex g_denoiserCapabilityMutex; + +static DenoiserCapability denoiser_capability() +{ + std::lock_guard lock(g_denoiserCapabilityMutex); + return g_denoiserCapability; +} + +// Whether a Ray Reconstruction feature can be created right now: the switch is on, a +// CUDA-capable denoiser build is present, and the core itself reported that denoiser +// initialised. Anything less is reported as unavailable rather than discovered at +// CreateFeature, where a game has already committed to the mode. +// Asks the core what it made of the denoiser. Must run on the worker, like every other core +// call, and must not be called from inside one. +static void refresh_denoiser_capability(); + +static bool ray_reconstruction_available() +{ + const DenoiserCapability capability = denoiser_capability(); + return rr_enabled() && g.ngxInitialized && g.denoiser.cudaCapable && + capability.known && capability.available != 0; +} + +// Why the route is not on offer, for the log and for the capability answer. +static const char* rr_denial_reason() +{ + if (!rr_enabled()) + return "disabled by D4R_RR_ENABLE=0"; + if (!g.ngxInitialized) + return "the NGX core is not initialised"; + if (!g.denoiser.cudaCapable) + return g.denoiserAttempted ? "nvngx_dlssd.dll has no NVSDK_NGX_CUDA_* entry points (needs a 3.10.7 build)" + : "no nvngx_dlssd.dll was found next to this DLL"; + const DenoiserCapability capability = denoiser_capability(); + if (!capability.known) + return "the NGX core has not reported the denoiser's state yet"; + if (capability.available == 0) + { + thread_local char text[160]; + snprintf(text, sizeof(text), "the NGX core reports the denoiser unavailable (init 0x%08x)", + static_cast(capability.initResult)); + return text; + } + return "available"; +} + // A game that initialises NGX with its project ID (engine integrations such as Unreal's) is identified // to the NGX core the same way. struct ProjectIdentity @@ -1456,18 +1679,13 @@ struct ProjectIdentity std::string engineVersion; }; -static NgxResult initialize(unsigned long long applicationId, const wchar_t* dataPath, ID3D12Device* device, - const NgxFeatureCommonInfo* featureInfo, unsigned int sdkVersion, - const ProjectIdentity* project = nullptr) +// Discovery is legal before NGX Init; use the same pinned worker for both entry points. +// The caller holds g.mutex. +static NgxResult ensure_workers_started() { - std::lock_guard lock(g.mutex); - g.shutdownPending = false; - if (g.ngxInitialized) - return NGX_SUCCESS; if (!g.started) { - // Allocators can retain COM callbacks and detached workers past NGX - // shutdown. Keep their code loaded until process exit. + // Allocators can retain COM callbacks and detached workers past NGX shutdown. HMODULE pinned = nullptr; if (!GetModuleHandleExW(GET_MODULE_HANDLE_EX_FLAG_FROM_ADDRESS | GET_MODULE_HANDLE_EX_FLAG_PIN, reinterpret_cast(g_selfModule), &pinned)) @@ -1479,6 +1697,19 @@ static NgxResult initialize(unsigned long long applicationId, const wchar_t* dat g.cleanup.start(); g.started = true; } + return NGX_SUCCESS; +} + +static NgxResult initialize(unsigned long long applicationId, const wchar_t* dataPath, ID3D12Device* device, + const NgxFeatureCommonInfo* featureInfo, unsigned int sdkVersion, + const ProjectIdentity* project = nullptr) +{ + std::lock_guard lock(g.mutex); + g.shutdownPending = false; + if (g.ngxInitialized) + return NGX_SUCCESS; + if (ensure_workers_started() != NGX_SUCCESS) + return NGX_FAIL_PLATFORM_ERROR; ensure_portable_config(); if (g_portable.active && !g.core) { @@ -1558,6 +1789,19 @@ static NgxResult initialize(unsigned long long applicationId, const wchar_t* dat if (loadError != nullptr && loadError()[0] != '\0') logf("nvcuda bridge: %s", loadError()); g.ngxInitialized = result == NGX_SUCCESS; + + // The denoiser is not initialised here. The NGX core loads and dispatches + // nvngx_dlssd.dll itself, so a second init would duplicate that state rather than add + // anything; this only looks for the library and checks it is a build that can serve + // Ray Reconstruction. Its own NVSDK_NGX_CUDA_Init is a stub that always fails, so + // calling it would report "unavailable" for a denoiser that is in fact working. + if (rr_enabled()) + { + load_denoiser_library(); + logf("Ray Reconstruction: %s", rr_denial_reason()); + } + else + logf("Ray Reconstruction: disabled by D4R_RR_ENABLE=0"); return result; } @@ -1629,10 +1873,365 @@ struct CudaImage { CudaArray array = nullptr; CudaObject object = 0; // texture or surface object + uint32_t format = 0, channels = 0; // the layout the array carries, so a format change + // recreates it rather than only reacting to a resize UINT width = 0; UINT height = 0; }; +// --- Ray Reconstruction auxiliary inputs -------------------------------------- +// +// Super Sampling needs four surfaces (colour, depth, motion vectors, exposure) and the +// file above already moves those. Ray Reconstruction additionally runs a denoiser network +// over the ray-traced G-buffer, so a game that enables it registers a further set of +// surfaces: normals, roughness, diffuse and specular albedo, emissive, the ray hit +// distances and directions, and the effect guides for particles, transparency, fog and +// depth of field. +// +// Those surfaces are read in whatever format the game renders them in, and the denoiser +// samples them through texture objects minted by this file. Rather than force them +// through the canonical planes above - which would re-encode a UNORM normal as half and +// silently change the values the network sees - each is copied with its own format and +// channel count, so the texture object describes exactly what the game drew. + +// CUarray_format values beyond the two above (nvsdk: CU_AD_FORMAT_*). +constexpr uint32_t CUDA_FORMAT_UNSIGNED_INT8 = 0x01; +constexpr uint32_t CUDA_FORMAT_SIGNED_INT8 = 0x08; +constexpr uint32_t CUDA_FORMAT_UNSIGNED_INT16 = 0x2; +constexpr uint32_t CUDA_FORMAT_SIGNED_INT16 = 0x09; +constexpr uint32_t CUDA_FORMAT_UNSIGNED_INT32 = 0x03; +constexpr uint32_t CUDA_FORMAT_SIGNED_INT32 = 0x0a; + +// How one denoiser input reaches the GPU. +enum class AuxLayout +{ + Verbatim, // the bytes are already exactly what the CUDA array must hold + Converted // the bytes must be unpacked into floats first (see convert_aux_row) +}; + +// The CUDA array a denoiser input uses. +// +// Only formats that are byte-for-byte correct go down the verbatim path, and that is a +// shorter list than it looks. A CUDA array stores element_size * channels bytes per texel +// and only ever has 1, 2 or 4 channels, so a packed format has no verbatim layout at all: +// R10G10B10A2_UNORM is 4 bytes per texel, but a 4-channel uint array is 16, and copying +// "4 channels' worth" would read four times past the end of every row. R11G11B10_FLOAT has +// three channels, which no CUDA array has. +// +// Nor is memory layout the only question. Sampling an integer-typed CUDA array normalizes it, +// which is right for a UNORM or SNORM surface and wrong for a UINT one, where a D3D12 shader +// reads the raw integer. And B8G8R8A8_UNORM has RGBA bytes in BGRA order, so a verbatim copy +// would hand the denoiser a texture whose red and blue channels are swapped. +// +// Everything that is not in this table is converted to a float array instead, which is +// correct for all of those cases at the cost of a per-row unpack. +static bool aux_array_layout(DXGI_FORMAT format, uint32_t& arrayFormat, uint32_t& channels, AuxLayout& how) +{ + struct Entry + { + DXGI_FORMAT format; + uint32_t arrayFormat; + uint32_t channels; + }; + // UNORM and SNORM map onto signed/unsigned integer arrays of the same width: sampling + // normalizes them, which is exactly how a D3D12 shader reads the resource. + static const Entry entries[] = { + {DXGI_FORMAT_R8_UNORM, CUDA_FORMAT_UNSIGNED_INT8, 1}, + {DXGI_FORMAT_R8_SNORM, CUDA_FORMAT_SIGNED_INT8, 1}, + {DXGI_FORMAT_R8G8_UNORM, CUDA_FORMAT_UNSIGNED_INT8, 2}, + {DXGI_FORMAT_R8G8_SNORM, CUDA_FORMAT_SIGNED_INT8, 2}, + {DXGI_FORMAT_R8G8B8A8_UNORM, CUDA_FORMAT_UNSIGNED_INT8, 4}, + {DXGI_FORMAT_R8G8B8A8_SNORM, CUDA_FORMAT_SIGNED_INT8, 4}, + {DXGI_FORMAT_R16_UNORM, CUDA_FORMAT_UNSIGNED_INT16, 1}, + {DXGI_FORMAT_R16_SNORM, CUDA_FORMAT_SIGNED_INT16, 1}, + {DXGI_FORMAT_R16_FLOAT, CUDA_FORMAT_HALF, 1}, + {DXGI_FORMAT_R16G16_UNORM, CUDA_FORMAT_UNSIGNED_INT16, 2}, + {DXGI_FORMAT_R16G16_SNORM, CUDA_FORMAT_SIGNED_INT16, 2}, + {DXGI_FORMAT_R16G16_FLOAT, CUDA_FORMAT_HALF, 2}, + {DXGI_FORMAT_R16G16B16A16_UNORM, CUDA_FORMAT_UNSIGNED_INT16, 4}, + {DXGI_FORMAT_R16G16B16A16_SNORM, CUDA_FORMAT_SIGNED_INT16, 4}, + {DXGI_FORMAT_R16G16B16A16_FLOAT, CUDA_FORMAT_HALF, 4}, + {DXGI_FORMAT_R32_FLOAT, CUDA_FORMAT_FLOAT, 1}, + {DXGI_FORMAT_R32G32_FLOAT, CUDA_FORMAT_FLOAT, 2}, + {DXGI_FORMAT_R32G32B32A32_FLOAT, CUDA_FORMAT_FLOAT, 4}, + }; + for (const Entry& entry : entries) + if (entry.format == format) + { + arrayFormat = entry.arrayFormat; + channels = entry.channels; + how = AuxLayout::Verbatim; + return true; + } + // Converted formats all end up as float arrays, so a shader read of the denoiser's + // texture returns the same numbers a D3D12 read of the caller's resource would. + struct Converted + { + DXGI_FORMAT format; + uint32_t channels; + }; + static const Converted converted[] = { + {DXGI_FORMAT_R8_UINT, 1}, + {DXGI_FORMAT_R8G8_UINT, 2}, + {DXGI_FORMAT_R8G8B8A8_UINT, 4}, + {DXGI_FORMAT_R16_UINT, 1}, + {DXGI_FORMAT_R32_UINT, 1}, + {DXGI_FORMAT_R32_SINT, 1}, + {DXGI_FORMAT_R32G32_UINT, 2}, + {DXGI_FORMAT_R32G32_SINT, 2}, + {DXGI_FORMAT_R32G32B32A32_UINT, 4}, + {DXGI_FORMAT_B8G8R8A8_UNORM, 4}, + {DXGI_FORMAT_R10G10B10A2_UNORM, 4}, + {DXGI_FORMAT_R11G11B10_FLOAT, 3}, + }; + for (const Converted& entry : converted) + if (entry.format == format) + { + arrayFormat = CUDA_FORMAT_FLOAT; + channels = entry.channels; + how = AuxLayout::Converted; + return true; + } + return false; +} + +// Unpacks one row of a converted denoiser input into `channels` floats per texel. Every +// branch reads exactly the bytes of one texel and no more. +static void convert_aux_row(DXGI_FORMAT format, const uint8_t* source, float* destination, UINT width, + uint32_t channels) +{ + switch (format) + { + case DXGI_FORMAT_R8_UINT: + for (UINT x = 0; x < width; ++x) + destination[x] = static_cast(source[x]); + return; + case DXGI_FORMAT_R8G8_UINT: + for (UINT x = 0; x < width; ++x) + { + destination[x * 2] = static_cast(source[x * 2]); + destination[x * 2 + 1] = static_cast(source[x * 2 + 1]); + } + return; + case DXGI_FORMAT_R8G8B8A8_UINT: + for (UINT x = 0; x < width; ++x) + for (uint32_t c = 0; c < 4; ++c) + destination[x * 4 + c] = static_cast(source[x * 4 + c]); + return; + case DXGI_FORMAT_R16_UINT: + for (UINT x = 0; x < width; ++x) + { + uint16_t value; + std::memcpy(&value, source + x * 2, sizeof(value)); + destination[x] = static_cast(value); + } + return; + case DXGI_FORMAT_R32_UINT: + for (UINT x = 0; x < width; ++x) + { + uint32_t value; + std::memcpy(&value, source + x * 4, sizeof(value)); + destination[x] = static_cast(value); + } + return; + case DXGI_FORMAT_R32_SINT: + for (UINT x = 0; x < width; ++x) + { + int32_t value; + std::memcpy(&value, source + x * 4, sizeof(value)); + destination[x] = static_cast(value); + } + return; + case DXGI_FORMAT_R32G32_UINT: + for (UINT x = 0; x < width; ++x) + { + uint32_t value[2]; + std::memcpy(value, source + x * 8, sizeof(value)); + destination[x * 2] = static_cast(value[0]); + destination[x * 2 + 1] = static_cast(value[1]); + } + return; + case DXGI_FORMAT_R32G32_SINT: + for (UINT x = 0; x < width; ++x) + { + int32_t value[2]; + std::memcpy(value, source + x * 8, sizeof(value)); + destination[x * 2] = static_cast(value[0]); + destination[x * 2 + 1] = static_cast(value[1]); + } + return; + case DXGI_FORMAT_R32G32B32A32_UINT: + for (UINT x = 0; x < width; ++x) + for (uint32_t c = 0; c < 4; ++c) + { + uint32_t value; + std::memcpy(&value, source + x * 16 + c * 4, sizeof(value)); + destination[x * 4 + c] = static_cast(value); + } + return; + case DXGI_FORMAT_B8G8R8A8_UNORM: + // Stored B, G, R, A; the denoiser's texture is read as R, G, B, A. + for (UINT x = 0; x < width; ++x) + { + const uint8_t* texel = source + x * 4; + destination[x * 4] = unorm(texel[2], 8); + destination[x * 4 + 1] = unorm(texel[1], 8); + destination[x * 4 + 2] = unorm(texel[0], 8); + destination[x * 4 + 3] = unorm(texel[3], 8); + } + return; + case DXGI_FORMAT_R10G10B10A2_UNORM: + for (UINT x = 0; x < width; ++x) + { + uint32_t packed; + std::memcpy(&packed, source + x * 4, sizeof(packed)); + destination[x * 4] = unorm(packed & 0x3FF, 10); + destination[x * 4 + 1] = unorm((packed >> 10) & 0x3FF, 10); + destination[x * 4 + 2] = unorm((packed >> 20) & 0x3FF, 10); + destination[x * 4 + 3] = unorm((packed >> 30) & 0x3, 2); + } + return; + case DXGI_FORMAT_R11G11B10_FLOAT: + for (UINT x = 0; x < width; ++x) + { + uint32_t packed; + std::memcpy(&packed, source + x * 4, sizeof(packed)); + destination[x * 3] = small_float_to_float(packed & 0x7FF, 6); + destination[x * 3 + 1] = small_float_to_float((packed >> 11) & 0x7FF, 6); + destination[x * 3 + 2] = small_float_to_float((packed >> 22) & 0x3FF, 5); + } + return; + default: + break; + } + // Unreachable: aux_array_layout only marks a format Converted when this switch handles it. + for (UINT x = 0; x < width * channels; ++x) + destination[x] = 0.0f; +} + +// The texel size `format` occupies, for the pitch of an interop buffer holding it verbatim. +static size_t aux_texel_bytes(uint32_t arrayFormat, uint32_t channels) +{ + const size_t element = arrayFormat == CUDA_FORMAT_FLOAT || arrayFormat == CUDA_FORMAT_UNSIGNED_INT32 || + arrayFormat == CUDA_FORMAT_SIGNED_INT32 + ? 4 + : (arrayFormat == CUDA_FORMAT_HALF || arrayFormat == CUDA_FORMAT_UNSIGNED_INT16 || + arrayFormat == CUDA_FORMAT_SIGNED_INT16) + ? 2 + : 1; + return element * channels; +} + +// One auxiliary surface: the name to ask the caller for, and the name the denoiser reads +// it under. They differ for the two albedos, which the SDK and the shipped library spell +// differently, in which case `source` names the caller's parameter and `publish` the +// denoiser's. +struct AuxInput +{ + const char* source; + const char* publish; +}; + +// Every auxiliary read name nvngx_dlssd.dll looks up, verified against the names in the +// shipped library and the SDK's nvsdk_ngx_defs_dlssd.h. A caller may bind several of these +// to one texture, so surfaces are keyed on the resource and carry every name bound to it. +// +// The publish column is the name the denoiser reads, taken from the shipped library's own +// lookup table (its .rdata strings, each referenced exactly once from +// NGXDLSSD::EvaluateFeature). It is not always the SDK's spelling: nvngx_dlssd.dll declares +// DLSS.Input.Normals.Subrect.Base.X/Y and DLSS.Input.Roughness.Subrect.Base.X/Y, but it reads +// the normals and roughness textures themselves as GBuffer.Normals and GBuffer.Roughness. It +// does read the two albedos as DLSS.Input.DiffuseAlbedo and DLSS.Input.SpecularAlbedo, which a +// caller may instead supply as GBuffer.DiffuseAlbedo / GBuffer.SpecularAlbedo - the SDK defines +// both spellings, so both are accepted and both publish the name the denoiser asks for. The +// subrect origins in kAuxSubrectNames are a separate matter and keep the SDK's own names. +static const AuxInput kAuxInputs[] = { + // Ray-traced G-buffer + {"GBuffer.Normals", "GBuffer.Normals"}, + {"GBuffer.Roughness", "GBuffer.Roughness"}, + {"DLSS.Input.DiffuseAlbedo", "DLSS.Input.DiffuseAlbedo"}, + {"GBuffer.DiffuseAlbedo", "DLSS.Input.DiffuseAlbedo"}, + {"GBuffer.IndirectAlbedo", "GBuffer.IndirectAlbedo"}, + {"DLSS.Input.SpecularAlbedo", "DLSS.Input.SpecularAlbedo"}, + {"GBuffer.SpecularAlbedo", "DLSS.Input.SpecularAlbedo"}, + {"GBuffer.Emissive", "GBuffer.Emissive"}, + {"GBuffer.SpecularMvec", "GBuffer.SpecularMvec"}, + // Ray-tracing noise hints + {"RayTracingHitDistance", "RayTracingHitDistance"}, + {"DLSSD.DiffuseHitDistance", "DLSSD.DiffuseHitDistance"}, + {"DLSSD.SpecularHitDistance", "DLSSD.SpecularHitDistance"}, + {"DLSSD.DiffuseRayDirection", "DLSSD.DiffuseRayDirection"}, + {"DLSSD.SpecularRayDirection", "DLSSD.SpecularRayDirection"}, + {"DLSSD.DiffuseRayDirectionHitDistance", "DLSSD.DiffuseRayDirectionHitDistance"}, + {"DLSSD.SpecularRayDirectionHitDistance", "DLSSD.SpecularRayDirectionHitDistance"}, + {"DLSSD.ReflectedAlbedo", "DLSSD.ReflectedAlbedo"}, + // Per-effect guides, read only when the game preserves that effect + {"DLSSD.ColorBeforeParticles", "DLSSD.ColorBeforeParticles"}, + {"DLSSD.ColorAfterParticles", "DLSSD.ColorAfterParticles"}, + {"DLSSD.ColorBeforeTransparency", "DLSSD.ColorBeforeTransparency"}, + {"DLSSD.ColorAfterTransparency", "DLSSD.ColorAfterTransparency"}, + {"DLSSD.ColorBeforeFog", "DLSSD.ColorBeforeFog"}, + {"DLSSD.ColorAfterFog", "DLSSD.ColorAfterFog"}, + {"DLSSD.DepthOfFieldGuide", "DLSSD.DepthOfFieldGuide"}, + {"DLSSD.ColorBeforeDepthOfField", "DLSSD.ColorBeforeDepthOfField"}, + {"DLSSD.ColorAfterDepthOfField", "DLSSD.ColorAfterDepthOfField"}, + {"DLSSD.ScreenSpaceRefractionGuide", "DLSSD.ScreenSpaceRefractionGuide"}, + {"DLSSD.ColorBeforeScreenSpaceRefraction", "DLSSD.ColorBeforeScreenSpaceRefraction"}, + {"DLSSD.ColorAfterScreenSpaceRefraction", "DLSSD.ColorAfterScreenSpaceRefraction"}, + {"DLSSD.ScreenSpaceSubsurfaceScatteringGuide", "DLSSD.ScreenSpaceSubsurfaceScatteringGuide"}, + {"DLSSD.ColorBeforeScreenSpaceSubsurfaceScattering", "DLSSD.ColorBeforeScreenSpaceSubsurfaceScattering"}, + {"DLSSD.ColorAfterScreenSpaceSubsurfaceScattering", "DLSSD.ColorAfterScreenSpaceSubsurfaceScattering"}, + {"DLSSD.Alpha", "DLSSD.Alpha"}, + // Masks and secondary motion + {"DLSS.DisocclusionMask", "DLSS.DisocclusionMask"}, + {"DLSS.Input.Reduce.Ghost.Mask", "DLSS.Input.Reduce.Ghost.Mask"}, + {"DLSSD.ResponsivityMask", "DLSSD.ResponsivityMask"}, + {"MotionVectorsReflection", "MotionVectorsReflection"}, + {"TransparencyMask", "TransparencyMask"}, + {"DepthHighRes", "DepthHighRes"}, +}; + +// Scalars the denoiser reads per frame, forwarded from the caller's own parameters. +struct AuxScalar +{ + const char* name; +}; + +static const AuxScalar kAuxScalars[] = { + {"DLSS.Denoise.Mode"}, + {"DLSS.Roughness.Mode"}, + {"DLSS.Use.HW.Depth"}, + {"DLSS.Use.Folded.Network"}, + {"DLSSD.IndicatorLevel"}, + {"DLSS.Translucency.Type"}, + {"DLSS.Subsurface.Type"}, + {"DLSS.Specular.Type"}, +}; + +// The camera matrices the denoiser reprojects with, forwarded as the float[16] the caller +// registered (the SDK passes matrices by pointer, not as sixteen scalars). +static const char* const kAuxMatrices[] = {"WorldToViewMatrix", "ViewToClipMatrix"}; + +static unsigned int get_uint_or(void* parameters, const char* name, unsigned int fallback); + +// The Subrect origins nvngx_dlssd.dll declares for its auxiliary inputs +// (nvsdk_ngx_defs_dlssd.h:60-67 and 99-146). A name here has both a .Subrect.Base.X and a +// .Subrect.Base.Y parameter; the parallel FrameParams array holds one origin per entry, and +// the strings are static so no per-frame allocation is needed to forward them. +static const char* const kAuxSubrectNames[] = { + "DLSS.Input.DiffuseAlbedo", "DLSS.Input.SpecularAlbedo", "DLSS.Input.Normals", + "DLSS.Input.Roughness", "DLSSD.Alpha", "DLSSD.ReflectedAlbedo", + "DLSSD.ColorBeforeParticles", "DLSSD.ColorAfterParticles", "DLSSD.ColorBeforeTransparency", + "DLSSD.ColorAfterTransparency", "DLSSD.ColorBeforeFog", "DLSSD.ColorAfterFog", + "DLSSD.ScreenSpaceRefractionGuide", "DLSSD.ColorBeforeScreenSpaceRefraction", + "DLSSD.ColorAfterScreenSpaceRefraction", "DLSSD.ScreenSpaceSubsurfaceScatteringGuide", + "DLSSD.ColorBeforeScreenSpaceSubsurfaceScattering", + "DLSSD.ColorAfterScreenSpaceSubsurfaceScattering", "DLSSD.DepthOfFieldGuide", + "DLSSD.ColorBeforeDepthOfField", "DLSSD.ColorAfterDepthOfField", "DLSSD.DiffuseHitDistance", + "DLSSD.SpecularHitDistance", "DLSSD.DiffuseRayDirection", "DLSSD.SpecularRayDirection", + "DLSSD.DiffuseRayDirectionHitDistance", "DLSSD.SpecularRayDirectionHitDistance", +}; + struct FrameParams { float jitterX = 0, jitterY = 0, mvScaleX = 1, mvScaleY = 1, sharpness = 0; @@ -1651,6 +2250,19 @@ struct FrameParams // D4R_SHIM_OUTPUT_PACKED: the output kernel stores the game's R10G10B10A2 texels itself, so the game-side // copy (or the in-place texture) takes them without the RGBA16F conversion image and blit bool packedOutput = false; + // Ray Reconstruction: the denoiser's own per-frame settings and camera matrices, copied + // off the caller's parameter object on the game thread so the worker never reads an + // object the game may have destroyed by then. A scalar the caller did not set this frame + // carries the creation-time value, not an invented one and not last frame's override; a + // matrix the caller did not set is published as absent, because the pointer names this + // frame's storage and nothing outlives it. + int auxScalar[sizeof(kAuxScalars) / sizeof(kAuxScalars[0])] = {}; + bool auxScalarSet[sizeof(kAuxScalars) / sizeof(kAuxScalars[0])] = {}; + float auxMatrix[sizeof(kAuxMatrices) / sizeof(kAuxMatrices[0])][16] = {}; + bool auxMatrixSet[sizeof(kAuxMatrices) / sizeof(kAuxMatrices[0])] = {}; + // One Subrect origin per entry of kAuxSubrectNames, snapshotted per frame. Zero means the + // caller did not register one, and zero is what is forwarded - never the previous frame's. + unsigned int auxSubrect[sizeof(kAuxSubrectNames) / sizeof(kAuxSubrectNames[0])][2] = {}; }; struct FrameTiming @@ -1687,6 +2299,41 @@ struct HostPlane size_t size() const { return rowBytes * height; } }; +// Ray Reconstruction auxiliary surface: one per distinct D3D12Resource the caller registers, +// keyed on the resource because several names routinely resolve to one texture. The surface +// carries the caller's own DXGI format - it is never converted to a canonical plane - so the +// texture object the denoiser samples describes exactly what was rendered. +// +// The surface outlives the frames that name it: `holds` counts queued frames still referring +// to it, and the resource is AddRef'd so a caller that releases it while a frame is in flight +// cannot pull it out from under the copy recorded on its command list. +struct AuxSurface +{ + ID3D12Resource* resource = nullptr; + std::vector publish; + DXGI_FORMAT format = DXGI_FORMAT_UNKNOWN; + UINT width = 0, height = 0; + uint32_t arrayFormat = 0, channels = 0; + AuxLayout how = AuxLayout::Verbatim; // whether the bytes reach the array as they are + uint32_t lastSeen = 0; // frame the caller last registered this resource + unsigned int holds = 0; // queued frames still naming this surface + bool dropping = false; // no longer registered; freed once holds reaches zero + CudaImage image; // worker only + CudaObject handle = 0; // what the parameters point at; its address must stay put + Staging staging[kSlots]; + HostPlane host[kSlots]; // unpacked floats, only for a converted surface (worker only) +}; + +// What one frame's slot bound: which surface answered which denoiser name, for that frame +// only. Publish reads this, never the live surface list, so a surface the caller registered +// last frame is never handed to the denoiser for a frame that did not register it. +struct AuxBinding +{ + AuxSurface* surface = nullptr; + std::vector names; + uint32_t frame = 0; +}; + struct InputSlot { Staging color, depth, motion, exposure; @@ -1701,11 +2348,36 @@ struct InputSlot CudaObject dilatedTexture = 0; UINT dilatedWidth = 0, dilatedHeight = 0; size_t dilatedPitch = 0; + std::vector auxBindings; // this slot's denoiser inputs, for one frame + // The DLSSD.OutputAlpha texture the caller asked for on this frame, held until the frame + // retires. Present only on a frame that actually registered it. + ID3D12Resource* alphaResource = nullptr; + UINT alphaWidth = 0, alphaHeight = 0, alphaBaseX = 0, alphaBaseY = 0; + UINT alphaExtentW = 0, alphaExtentH = 0; + DXGI_FORMAT alphaFormat = DXGI_FORMAT_UNKNOWN; + uint32_t alphaWriteback = 0; + bool alphaRequested = false; std::atomic busy{false}; // readback staging owned by the pipeline HostPlane host[4]; std::atomic hostBusy{false}; // host planes not yet uploaded by the worker }; +// DLSSD.OutputAlpha is an OUTPUT the denoiser writes (nvsdk_ngx_defs_dlssd.h:98, with its own +// Subrect base at 101/102). It is delivered back through its own ring of slots, modelled on +// OutputSlot above, because it has exactly the same lifetime problem as the colour output: a +// result staged for one frame is read by a later frame's command list, so it must not be +// rewritten until every reader has passed. +struct AlphaSlot +{ + Staging staging; // upload buffer holding one result + std::atomic lastReadFrame{0}; // reserved/last-read frame, as OutputSlot's + std::atomic producedFrame{0}; // frame that filled it + DXGI_FORMAT format = DXGI_FORMAT_UNKNOWN; + UINT baseX = 0, baseY = 0; // where this result sits in the staging buffer + UINT width = 0, height = 0; // the region it covers + uint32_t writeback = 0; // destination texel kind (see alpha_output_layout) +}; + struct OutputSlot { Staging staging; @@ -1838,6 +2510,34 @@ struct Feature // Serializes picking the slot to present (game thread) against picking // the slot to overwrite (worker). std::mutex outputMutex; + // Ray Reconstruction: feature 13, which the NGX core dispatches to the denoiser it + // loaded. Super Sampling features leave this false and differ only in their feature id + // and their preset names. + bool rayReconstruction = false; + // The denoiser's settings as the caller stated them at creation time. A setting the + // caller varies per frame must fall back to this, not to the previous frame's value: + // the parameter object outlives the frame, so an untouched one would silently carry + // over an override that is no longer being asked for. + int auxScalarBaseline[sizeof(kAuxScalars) / sizeof(kAuxScalars[0])] = {}; + + // The denoiser's alpha output. Its CUDA surface is float32 single channel: the vendor's + // transformer output kernel stores it as a nullable scalar float32 pixel + // (sust.p.2d.b32.zero at param+248), so that is what the surface is created as. + CudaImage alphaImage; + HostPlane alphaHost; // worker: the read-back float32 result + AlphaSlot alpha[kOutputSlots]; + std::atomic latestAlpha{-1}; // the slot holding the newest unpresented result + + // A list, not a vector or a deque: the parameter object holds &AuxSurface.handle for the + // duration of the denoiser's evaluate, so every element must keep its address. A vector + // moves elements on growth and a deque invalidates *all* references when an element is + // erased from the middle - which is exactly what dropping a surface does while other + // surfaces are still queued. A list's erase touches only the element removed. + std::list aux; // the caller's registered surfaces, under auxMutex + // The game thread rewrites this list from the caller's parameter object while the worker + // reads it, and the parameter object holds &aux.handle for the duration of the + // denoiser's evaluate, so every aux access is under this lock. + std::mutex auxMutex; }; static std::mutex g_featuresMutex; @@ -1952,7 +2652,8 @@ static CudaTextureDesc input_sampler(UINT width, uint32_t channels, bool linear) static bool ensure_cuda_image(CudaImage& image, UINT width, UINT height, uint32_t format, uint32_t channels, bool surface, bool linear) { - if (image.array != nullptr && image.width == width && image.height == height) + if (image.array != nullptr && image.width == width && image.height == height && + image.format == format && image.channels == channels) return true; if (image.object != 0) surface ? g.cu.surfObjectDestroy(image.object) : g.cu.texObjectDestroy(image.object); @@ -2006,6 +2707,8 @@ static bool ensure_cuda_image(CudaImage& image, UINT width, UINT height, uint32_ } image.width = width; image.height = height; + image.format = format; + image.channels = channels; return true; } @@ -3483,6 +4186,635 @@ static void write_output(Staging& staging, const HostPlane& host, FrameTiming* t timing->outputConvert = profile_ms(conversionStart, ProfileClock::now()); } +static ID3D12Resource* get_resource(void* parameters, const char* name) +{ + ID3D12Resource* resource = nullptr; + if (d4r_ngx_get_d3d12_resource(parameters, name, &resource) == NGX_SUCCESS && resource != nullptr) + return resource; + void* raw = nullptr; + if (d4r_ngx_get_void(parameters, name, &raw) == NGX_SUCCESS) + return static_cast(raw); + return nullptr; +} + +// --- DLSSD.OutputAlpha: the denoiser's own output --------------------------------- +// +// nvsdk_ngx_defs_dlssd.h:98 makes this an OUTPUT the denoiser writes, not another G-buffer +// input, and it carries its own Subrect base (101/102), so it has to come back to the caller +// rather than be read from it. +// +// The surface is float32, single channel: the vendor's transformer output kernel stores this +// value as a nullable scalar float32 pixel (sust.p.2d.b32.zero against param+248), so the +// surface object is created over a 1-channel float array and converted once on the way out. +// +// Everything else is borrowed from the colour output's proven lifecycle - a ring of slots, a +// claim so a staged result is not overwritten while a recorded command list still reads it, +// and the result presented a frame later - rather than being a second, asymmetric mechanism. + +// The single-channel formats a caller may register, and how to write each back. +static bool alpha_output_layout(DXGI_FORMAT format, uint32_t& writeback) +{ + switch (format) + { + case DXGI_FORMAT_R8_UNORM: writeback = 1; return true; + case DXGI_FORMAT_R8_UINT: writeback = 2; return true; + case DXGI_FORMAT_R16_FLOAT: writeback = 3; return true; + case DXGI_FORMAT_R16_UNORM: writeback = 4; return true; + case DXGI_FORMAT_R16_UINT: writeback = 5; return true; + case DXGI_FORMAT_R32_FLOAT: writeback = 6; return true; + default: return false; + } +} + +static size_t alpha_texel_bytes(uint32_t writeback) +{ + return writeback <= 2 ? 1 : (writeback <= 5 ? 2 : 4); +} + +static void write_alpha_texel(uint32_t writeback, uint8_t* destination, float value) +{ + switch (writeback) + { + case 1: + { + const uint32_t v = to_unorm(value, 8); + std::memcpy(destination, &v, 1); + return; + } + case 2: + { + const uint8_t v = value <= 0.0f ? 0u : (value >= 255.0f ? 255u : static_cast(value)); + std::memcpy(destination, &v, sizeof(v)); + return; + } + case 3: + { + const uint16_t v = float_to_half(value); + std::memcpy(destination, &v, sizeof(v)); + return; + } + case 4: + { + const uint16_t v = static_cast(to_unorm(value, 16)); + std::memcpy(destination, &v, sizeof(v)); + return; + } + case 5: + { + const uint32_t v = value <= 0.0f ? 0u : (value >= 65535.0f ? 65535u : static_cast(value)); + const uint16_t narrow = static_cast(v); + std::memcpy(destination, &narrow, sizeof(narrow)); + return; + } + default: + { + std::memcpy(destination, &value, sizeof(value)); + return; + } + } +} + +// Game thread. Records which DLSSD.OutputAlpha texture this frame wants written, and holds a +// reference on it until the frame retires. Refuses loudly a texture this file cannot fill: +// accepting the frame and leaving the caller's surface untouched is exactly the silent wrong +// answer this whole route exists to avoid. Sets nothing on the CUDA parameter object - that is +// the worker's, and this runs on the game thread. +static bool capture_alpha_output(Feature& feature, void* parameters, int slotIndex, uint32_t frame) +{ + InputSlot& slot = feature.inputs[slotIndex]; + ID3D12Resource* const resource = get_resource(parameters, "DLSSD.OutputAlpha"); + if (resource == nullptr) + { + slot.alphaRequested = false; + return true; + } + D3D12_RESOURCE_DESC desc; + resource->GetDesc(&desc); + uint32_t writeback = 0; + if (desc.Dimension != D3D12_RESOURCE_DIMENSION_TEXTURE2D || desc.DepthOrArraySize != 1 || + desc.Width == 0 || desc.Height == 0 || desc.Width > 0xFFFFFFFFull || desc.SampleDesc.Count != 1 || + !alpha_output_layout(desc.Format, writeback)) + { + logf("Ray Reconstruction: DLSSD.OutputAlpha is %s %llux%u array=%u samples=%u fmt=0x%x, which this " + "file cannot write back; this frame's evaluation is refused rather than accepted with the " + "texture left untouched", + desc.Dimension == D3D12_RESOURCE_DIMENSION_TEXTURE2D ? "a texture" : "not a texture", desc.Width, + desc.Height, desc.DepthOrArraySize, desc.SampleDesc.Count, desc.Format); + return false; + } + // The caller's Subrect. The surface is allocated at the texture's FULL extent and the base + // is forwarded, so the denoiser stores where the caller expects; only the evaluated region + // is copied back, leaving the caller's pixels outside it untouched. + const UINT baseX = get_uint_or(parameters, "DLSSD.OutputAlpha.Subrect.Base.X", 0); + const UINT baseY = get_uint_or(parameters, "DLSSD.OutputAlpha.Subrect.Base.Y", 0); + // Guard the subtractions themselves: a base outside the texture underflows below. + if (baseX > desc.Width || baseY > desc.Height) + { + logf("Ray Reconstruction: DLSSD.OutputAlpha subrect origin (%u,%u) is outside its %llux%u texture; " + "this frame's evaluation is refused", + baseX, baseY, desc.Width, desc.Height); + return false; + } + // The evaluated region is the output size placed at that origin, and it must fit; it is not + // clamped silently into whatever happens to fit. + const UINT remainingW = static_cast(desc.Width) - baseX; + const UINT remainingH = desc.Height - baseY; + UINT extentW = get_uint_or(parameters, "OutWidth", feature.outWidth); + UINT extentH = get_uint_or(parameters, "OutHeight", feature.outHeight); + if (extentW == 0 || extentW > remainingW || extentH == 0 || extentH > remainingH) + { + logf("Ray Reconstruction: DLSSD.OutputAlpha needs a %ux%u region at (%u,%u) inside a %llux%u texture, " + "which does not fit; this frame's evaluation is refused", + extentW, extentH, baseX, baseY, desc.Width, desc.Height); + return false; + } + if (slot.alphaResource != nullptr && slot.alphaResource != resource) + { + slot.alphaResource->Release(); + slot.alphaResource = nullptr; // so the new resource is always AddRef'd below + } + if (slot.alphaResource == nullptr) + resource->AddRef(); + slot.alphaResource = resource; + slot.alphaFormat = desc.Format; + slot.alphaWidth = static_cast(desc.Width); // full extent: where the vendor stores + slot.alphaHeight = desc.Height; + slot.alphaBaseX = baseX; + slot.alphaBaseY = baseY; + slot.alphaExtentW = extentW; // what is actually copied back + slot.alphaExtentH = extentH; + slot.alphaWriteback = writeback; + slot.alphaRequested = true; + if (frame <= 2) + logf("Ray Reconstruction: DLSSD.OutputAlpha %ux%u fmt=0x%x, subrect (%u,%u) region %ux%u", desc.Width, + desc.Height, desc.Format, baseX, baseY, extentW, extentH); + return true; +} + +// Worker. Gives the denoiser a float32 single-channel CUDA surface object for this frame's +// alpha and forwards the subrect origin, so its stores land where the caller expects. Clearing +// the parameter when no texture was registered keeps the denoiser from writing into a surface +// belonging to a frame that no longer wants one; the image itself is kept, since dropping it +// would mean a fresh GPU allocation on the next frame that does ask. +static bool publish_alpha_output(Feature& feature, const InputSlot& slot) +{ + if (!slot.alphaRequested) + { + d4r_ngx_set_void(feature.cudaParams, "DLSSD.OutputAlpha", nullptr); + return true; + } + if (!ensure_cuda_image(feature.alphaImage, slot.alphaWidth, slot.alphaHeight, CUDA_FORMAT_FLOAT, 1, true, false)) + return false; + d4r_ngx_set_void(feature.cudaParams, "DLSSD.OutputAlpha", &feature.alphaImage.object); + d4r_ngx_set_uint(feature.cudaParams, "DLSSD.OutputAlpha.Subrect.Base.X", slot.alphaBaseX); + d4r_ngx_set_uint(feature.cudaParams, "DLSSD.OutputAlpha.Subrect.Base.Y", slot.alphaBaseY); + return true; +} + +// Worker. The colour output's exact claim: the GPU marker is the only thing that proves a +// recorded copy has run, so a slot is reusable only once the GPU is past the frame that last +// read it, and the newest result is never taken from under a reader. +static int claim_alpha_slot(Feature& feature) +{ + const auto start = std::chrono::steady_clock::now(); + for (;;) + { + { + std::lock_guard lock(feature.outputMutex); + const uint32_t gpuFrame = *feature.markerValue; + for (int index = 0; index < kOutputSlots; ++index) + { + if (index == feature.latestAlpha.load()) + continue; + if (static_cast(gpuFrame - feature.alpha[index].lastReadFrame.load()) > 0) + { + feature.alpha[index].lastReadFrame = 0x7fffffffu + gpuFrame; + return index; + } + } + } + if (std::chrono::steady_clock::now() - start > std::chrono::milliseconds(50)) + return -1; + d4r_sleep_us(200); + } +} + +// Worker. Reads back the surface the denoiser just wrote and converts the evaluated region into +// the claimed slot's upload buffer. A failure fails the frame: an accepted frame with a missing +// alpha would leave the caller believing it was written. +static bool download_alpha_output(Feature& feature, uint32_t frame) +{ + const InputSlot& slot = feature.inputs[frame % kSlots]; + if (!slot.alphaRequested) + return true; + if (feature.alphaImage.array == nullptr) + return false; + const int index = claim_alpha_slot(feature); + if (index < 0) + { + logf("frame %u: no free DLSSD.OutputAlpha result slot; the evaluation is dropped", frame); + return false; + } + AlphaSlot& target = feature.alpha[index]; + if (!ensure_staging(target.staging, slot.alphaResource, D3D12_HEAP_TYPE_UPLOAD)) + { + logf("frame %u: no upload buffer for the DLSSD.OutputAlpha result; the evaluation is dropped", frame); + return false; + } + reserve_host(feature.alphaHost, sizeof(float) * slot.alphaWidth * slot.alphaHeight); + if (feature.alphaHost.bytes == nullptr || + feature.alphaHost.capacity < sizeof(float) * slot.alphaWidth * slot.alphaHeight) + { + logf("frame %u: no room to read the DLSSD.OutputAlpha surface back; the evaluation is dropped", frame); + return false; + } + CudaMemcpy2D copy = {}; + copy.srcMemoryType = CUDA_MEMORY_ARRAY; + copy.srcArray = feature.alphaImage.array; + copy.dstMemoryType = CUDA_MEMORY_HOST; + copy.dstHost = feature.alphaHost.bytes; + copy.dstPitch = sizeof(float) * slot.alphaWidth; + copy.WidthInBytes = sizeof(float) * slot.alphaWidth; + copy.Height = slot.alphaHeight; + // Conversion reads the host bytes immediately, so this transfer must complete on return. + if (g.cu.memcpy2D(©) != 0) + { + logf("frame %u: reading the DLSSD.OutputAlpha surface back failed; the evaluation is dropped", frame); + return false; + } + // Only the evaluated rectangle is written, at its position in the buffer. The footprint + // offset from GetCopyableFootprints is left alone - it must stay 512-byte aligned - and the + // region is instead selected by the source box of the copy below. + const float* const source = reinterpret_cast(feature.alphaHost.bytes); + uint8_t* const destination = target.staging.mapped + target.staging.layout.Offset; + const size_t rowPitch = target.staging.layout.Footprint.RowPitch; + const size_t texel = alpha_texel_bytes(slot.alphaWriteback); + for (UINT row = 0; row < slot.alphaExtentH; ++row) + for (UINT x = 0; x < slot.alphaExtentW; ++x) + write_alpha_texel(slot.alphaWriteback, + destination + static_cast(slot.alphaBaseY + row) * rowPitch + + static_cast(slot.alphaBaseX) * texel + x * texel, + source[static_cast(slot.alphaBaseY + row) * slot.alphaWidth + slot.alphaBaseX + x]); + { + // Published under the same lock the claim uses, exactly as the colour output publishes. + std::lock_guard lock(feature.outputMutex); + target.format = slot.alphaFormat; + target.baseX = slot.alphaBaseX; + target.baseY = slot.alphaBaseY; + target.width = slot.alphaExtentW; + target.height = slot.alphaExtentH; + target.writeback = slot.alphaWriteback; + target.lastReadFrame = 0; + target.producedFrame = frame; + feature.latestAlpha = index; + } + return true; +} + +// Game thread. Presents the newest finished alpha result into the caller's texture. The slot's +// lastReadFrame is stamped before the copy is recorded and producedFrame is NOT cleared: the +// colour output's claim relies on both, and clearing here would let the worker overwrite a +// result whose command list has not run yet. +static bool record_alpha_copyback(Feature& feature, ID3D12GraphicsCommandList* list, InputSlot& slot, + D3D12_RESOURCE_STATES outputState, uint32_t colorFrame) +{ + if (!slot.alphaRequested || slot.alphaResource == nullptr) + return true; + if (colorFrame == 0) + return true; // no colour result to pair with this frame + int index = -1; + { + std::lock_guard lock(feature.outputMutex); + // The frame whose colour is about to be presented is `colorFrame`; only an alpha result + // produced by that very same frame may go with it. Picking "the newest alpha" instead + // would pair alpha N with colour N-1, since the alpha read-back completes before the + // colour's, and the caller would composite two different frames. No match, no copy: the + // colour is presented on its own and the alpha waits for its own frame to come round. + for (int candidate = 0; candidate < kOutputSlots; ++candidate) + if (feature.alpha[candidate].producedFrame.load() == colorFrame) + { + index = candidate; + break; + } + if (index < 0) + return true; // the matching frame has not been read back yet; colour goes alone + AlphaSlot& chosen = feature.alpha[index]; + if (chosen.staging.buffer == nullptr || chosen.producedFrame.load() == 0) + return true; + if (chosen.width != slot.alphaExtentW || chosen.height != slot.alphaExtentH || + chosen.format != slot.alphaFormat || chosen.baseX != slot.alphaBaseX || + chosen.baseY != slot.alphaBaseY) + { + // The caller's rectangle changed under a result already in flight; presenting it + // would write where the caller no longer expects. Leave it for the next result. + chosen.producedFrame = 0; + return true; + } + chosen.lastReadFrame = feature.frame; + } + AlphaSlot& chosen = feature.alpha[index]; + transition(list, slot.alphaResource, outputState, D3D12_RESOURCE_STATE_COPY_DEST); + D3D12_TEXTURE_COPY_LOCATION destination = {}; + destination.pResource = slot.alphaResource; + destination.Type = D3D12_TEXTURE_COPY_TYPE_SUBRESOURCE_INDEX; + destination.SubresourceIndex = 0; + D3D12_TEXTURE_COPY_LOCATION source = {}; + source.pResource = chosen.staging.buffer; + source.Type = D3D12_TEXTURE_COPY_TYPE_PLACED_FOOTPRINT; + source.PlacedFootprint = chosen.staging.layout; + // The source box picks out the evaluated rectangle inside the full-extent staging buffer, + // and the destination origin places it at the caller's subrect. Together they leave every + // pixel of the caller's texture outside the rectangle untouched. + D3D12_BOX box = {}; + box.left = chosen.baseX; + box.top = chosen.baseY; + box.right = chosen.baseX + chosen.width; + box.bottom = chosen.baseY + chosen.height; + box.back = 1; + list->CopyTextureRegion(&destination, chosen.baseX, chosen.baseY, 0, &source, &box); + transition(list, slot.alphaResource, D3D12_RESOURCE_STATE_COPY_DEST, outputState); + return true; +} + +// --- Ray Reconstruction: moving the auxiliary surfaces -------------------------- + +// How long a caller may stop registering a surface before it is treated as dropped. Long +// enough that a game which registers its G-buffer lazily, or only on the frames where it +// changes, keeps it; short enough that a resolution change - which reallocates every resource +// and can reuse a freed address - does not leave a surface bound to somebody else's texture. +constexpr uint32_t kAuxStaleFrames = 300; + +static AuxSurface* find_aux_surface(Feature& feature, ID3D12Resource* resource) +{ + for (AuxSurface& surface : feature.aux) + if (surface.resource == resource) + return &surface; + return nullptr; +} + +static std::string join_aux_names(const std::vector& names) +{ + std::string joined; + for (const char* name : names) + joined += std::string(joined.empty() ? "" : ", ") + name; + return joined; +} + +// Worker. Frees the CUDA objects of a surface and the caller's reference to its resource. +// Only ever called once no queued frame names it any more. +static void destroy_aux_surface(AuxSurface& surface) +{ + if (surface.image.object != 0) + g.cu.texObjectDestroy(surface.image.object); + if (surface.image.array != nullptr) + g.cu.arrayDestroy(surface.image.array); + for (Staging& staging : surface.staging) + staging.release(); + if (surface.resource != nullptr) + surface.resource->Release(); + logf("Ray Reconstruction: released denoiser input %s (%ux%u fmt=0x%x)", join_aux_names(surface.publish).c_str(), + surface.width, surface.height, surface.format); +} + +// Game thread. Finds the auxiliary surfaces the caller registered this frame and records a +// readback copy of each into this frame's slot, ahead of the frame marker so the worker can +// read them as soon as that marker is in. The result is this slot's binding snapshot: the +// worker publishes exactly these, so a surface the caller did not register this frame is +// neither sampled nor named, whatever it did last frame. +// +// A texture the caller did supply but this file cannot hand to the denoiser (a shape or a +// format with no CUDA layout) fails the evaluation. Silently omitting a guide the denoiser +// expects is exactly the failure that produces a plausible-looking wrong image, so the frame +// is refused and the reason named instead. +static bool gather_aux_inputs(Feature& feature, ID3D12GraphicsCommandList* list, void* parameters, int slotIndex, + D3D12_RESOURCE_STATES inputState, uint32_t frame) +{ + InputSlot& slot = feature.inputs[slotIndex]; + std::lock_guard lock(feature.auxMutex); + std::vector bindings; + for (const AuxInput& input : kAuxInputs) + { + ID3D12Resource* const resource = get_resource(parameters, input.source); + if (resource == nullptr) + continue; + AuxSurface* surface = find_aux_surface(feature, resource); + if (surface == nullptr) + { + feature.aux.emplace_back(); + surface = &feature.aux.back(); + surface->resource = resource; + resource->AddRef(); // held until every frame naming it has retired + logf("Ray Reconstruction: denoiser input %s is resource %p", input.source, static_cast(resource)); + } + surface->dropping = false; + surface->lastSeen = frame; + // One binding per name: two names for the same texture are two denoiser parameters. + AuxBinding* binding = nullptr; + for (AuxBinding& candidate : bindings) + if (candidate.surface == surface) + { + binding = &candidate; + break; + } + if (binding == nullptr) + { + bindings.emplace_back(); + binding = &bindings.back(); + binding->surface = surface; + binding->frame = frame; + } + bool listed = false; + for (const char* name : surface->publish) + listed = listed || strcmp(name, input.publish) == 0; + if (!listed) + surface->publish.push_back(input.publish); + bool bound = false; + for (const char* name : binding->names) + bound = bound || strcmp(name, input.publish) == 0; + if (!bound) + binding->names.push_back(input.publish); + + D3D12_RESOURCE_DESC desc; + resource->GetDesc(&desc); + uint32_t arrayFormat = 0, channels = 0; + AuxLayout how = AuxLayout::Verbatim; + if (desc.DepthOrArraySize != 1 || desc.Width > 0xFFFFFFFFull || + !aux_array_layout(desc.Format, arrayFormat, channels, how)) + { + logf("Ray Reconstruction: denoiser input %s is %llux%u array=%u fmt=0x%x, which has no single-plane " + "CUDA layout; this frame's evaluation is refused rather than run without it", + join_aux_names(binding->names).c_str(), desc.Width, desc.Height, desc.DepthOrArraySize, desc.Format); + return false; + } + surface->format = desc.Format; + surface->arrayFormat = arrayFormat; + surface->channels = channels; + surface->how = how; + surface->width = static_cast(desc.Width); + surface->height = static_cast(desc.Height); + if (!ensure_staging(surface->staging[slotIndex], resource, D3D12_HEAP_TYPE_READBACK)) + return false; + copy_to_staging(list, resource, surface->staging[slotIndex], inputState); + } + // Drop surfaces the caller stopped registering, but only once no queued frame names one: + // erasing a surface whose handle is still in a parameter object would leave the denoiser + // sampling a freed texture object. + for (AuxSurface& surface : feature.aux) + { + if (surface.holds != 0 || frame - surface.lastSeen < kAuxStaleFrames) + continue; + surface.dropping = true; + logf("Ray Reconstruction: denoiser input %s was not registered for %u frames", + join_aux_names(surface.publish).c_str(), kAuxStaleFrames); + } + // Committed only now. An earlier failure path returns from this function with `bindings` + // half-built, and a hold taken there would never be given back - the surface could then + // never be dropped, and its reference to the caller's resource would never be released. + for (const AuxBinding& binding : bindings) + ++binding.surface->holds; // given back by retire_aux_bindings when this frame passes + slot.auxBindings = std::move(bindings); + return true; +} + +// Worker. Copies this frame's bound surfaces into their CUDA arrays verbatim - same bytes, +// same row order, only the destination is the GPU - and mints the texture objects the +// denoiser samples. A failure fails the evaluation: a denoiser handed fewer guides than the +// caller registered produces a wrong image, not a visible error. +static bool upload_aux_inputs(Feature& feature, int slotIndex, uint32_t frame) +{ + for (const AuxBinding& binding : feature.inputs[slotIndex].auxBindings) + { + AuxSurface& surface = *binding.surface; + const Staging& staging = surface.staging[slotIndex]; + if (!ensure_cuda_image(surface.image, surface.width, surface.height, surface.arrayFormat, + surface.channels, false, false)) + { + logf("frame %u: Ray Reconstruction input %s: a CUDA texture object over %ux%u fmt=0x%x could not be " + "created", + frame, join_aux_names(binding.names).c_str(), surface.width, surface.height, surface.format); + return false; + } + CudaMemcpy2D copy = {}; + copy.srcMemoryType = CUDA_MEMORY_HOST; + if (surface.how == AuxLayout::Verbatim) + { + // Use HostPlane staging: direct ROCm image copies from Vulkan-mapped memory + // truncate guide rows on gfx1201, even when the source footprint is tightly packed. + HostPlane& host = surface.host[slotIndex]; + const size_t rowBytes = aux_texel_bytes(surface.arrayFormat, surface.channels) * surface.width; + reserve_host(host, rowBytes * surface.height); + if (host.bytes == nullptr || host.capacity < rowBytes * surface.height) + return false; + host.width = surface.width; + host.height = surface.height; + host.rowBytes = rowBytes; + for (UINT row = 0; row < surface.height; ++row) + memcpy(host.bytes + static_cast(row) * rowBytes, + staging.mapped + staging.layout.Offset + + static_cast(row) * staging.layout.Footprint.RowPitch, rowBytes); + copy.srcHost = host.bytes; + copy.srcPitch = rowBytes; + } + else + { + // Packed, integer and BGRA inputs are unpacked into floats first: a CUDA array + // has no three-channel layout, a uint array is sampled normalized, and BGRA bytes + // are in the wrong order for a read as RGBA. + HostPlane& host = surface.host[slotIndex]; + const size_t rowBytes = sizeof(float) * surface.channels * surface.width; + reserve_host(host, rowBytes * surface.height); + if (host.bytes == nullptr || host.capacity < rowBytes * surface.height) + { + logf("frame %u: Ray Reconstruction input %s: no room for %zu unpacked bytes", + frame, join_aux_names(binding.names).c_str(), rowBytes * surface.height); + return false; + } + host.width = surface.width; + host.height = surface.height; + host.rowBytes = rowBytes; + const uint8_t* const base = staging.mapped + staging.layout.Offset; + for (UINT row = 0; row < surface.height; ++row) + convert_aux_row(surface.format, base + static_cast(row) * staging.layout.Footprint.RowPitch, + reinterpret_cast(host.bytes + static_cast(row) * rowBytes), + surface.width, surface.channels); + copy.srcHost = host.bytes; + copy.srcPitch = rowBytes; + } + copy.dstMemoryType = CUDA_MEMORY_ARRAY; + copy.dstArray = surface.image.array; + copy.WidthInBytes = aux_texel_bytes(surface.arrayFormat, surface.channels) * surface.width; + copy.Height = surface.height; + if (copy_2d(copy) != 0) + { + logf("frame %u: Ray Reconstruction input %s: the upload of %ux%u failed", frame, + join_aux_names(binding.names).c_str(), surface.width, surface.height); + return false; + } + surface.handle = surface.image.object; + } + return true; +} + +// Worker. Points the denoiser's parameters at this frame's texture objects and forwards the +// settings and camera matrices the caller registered for this frame. Every name this file +// ever published is cleared first: the parameter object outlives the frame, so a name left +// holding last frame's handle would be read as a live texture. +static void publish_aux_parameters(Feature& feature, FrameParams& p, const InputSlot& slot) +{ + void* params = feature.cudaParams; + for (size_t index = 0; index < sizeof(kAuxScalars) / sizeof(kAuxScalars[0]); ++index) + if (p.auxScalarSet[index]) + d4r_ngx_set_int(params, kAuxScalars[index].name, p.auxScalar[index]); + // Always written, never conditionally: the parameter object outlives this frame, so a + // matrix pointer left over from the previous one would dangle the moment this frame's + // FrameParams goes out of scope. Absent, it is cleared - the denoiser then knows the + // caller supplied none, rather than reading a stale one. + for (size_t index = 0; index < sizeof(kAuxMatrices) / sizeof(kAuxMatrices[0]); ++index) + d4r_ngx_set_void(params, kAuxMatrices[index], + p.auxMatrixSet[index] ? static_cast(p.auxMatrix[index]) : nullptr); + for (size_t index = 0; index < sizeof(kAuxSubrectNames) / sizeof(kAuxSubrectNames[0]); ++index) + { + static std::string nameX, nameY; + nameX.assign(kAuxSubrectNames[index]).append(".Subrect.Base.X"); + nameY.assign(kAuxSubrectNames[index]).append(".Subrect.Base.Y"); + d4r_ngx_set_uint(params, nameX.c_str(), p.auxSubrect[index][0]); + d4r_ngx_set_uint(params, nameY.c_str(), p.auxSubrect[index][1]); + } + for (const AuxSurface& surface : feature.aux) + for (const char* name : surface.publish) + d4r_ngx_set_void(params, name, nullptr); + // Pointers to the handles, which is what a CUDA resource parameter is: a variable holding + // the object, not the object itself. Two surfaces naming one denoiser input (a caller that + // binds GBuffer.DiffuseAlbedo and DLSS.Input.DiffuseAlbedo to different textures) leave the + // later one winning, which is the caller's ambiguity to resolve, not this file's. + for (const AuxBinding& binding : slot.auxBindings) + if (binding.surface->handle != 0) + for (const char* name : binding.names) + d4r_ngx_set_void(params, name, &binding.surface->handle); +} + +// Worker, once the frame this slot was gathering for has passed. Each bound surface keeps the +// caller's resource alive until here, and a dropped surface is freed only once its last frame +// is gone. +static void retire_aux_bindings(Feature& feature, int slotIndex) +{ + std::lock_guard lock(feature.auxMutex); + for (AuxBinding& binding : feature.inputs[slotIndex].auxBindings) + if (binding.surface->holds > 0) + --binding.surface->holds; + feature.inputs[slotIndex].auxBindings.clear(); + for (auto it = feature.aux.begin(); it != feature.aux.end();) + { + if (!it->dropping || it->holds != 0) + { + ++it; + continue; + } + destroy_aux_surface(*it); + it = feature.aux.erase(it); + } +} + static void set_evaluation_parameters(Feature& feature, const FrameParams& p) { void* params = feature.cudaParams; @@ -3565,6 +4897,12 @@ static void frame_retired(Feature* feature, uint32_t frame) { if (feature->split) release_split_frame(feature, frame); + // The denoiser inputs this frame bound belong to the slot it used, and the caller stops + // owing them a reference only now that its evaluate has returned. + if (feature->rayReconstruction) + { + retire_aux_bindings(*feature, static_cast(frame % kSlots)); + } feature->inFlight.fetch_sub(1); } @@ -4191,7 +5529,9 @@ static bool recreate_for_render_size(Feature& feature, FrameParams& params, uint void* p = feature.cudaParams; d4r_ngx_set_uint(p, "Width", params.renderWidth); d4r_ngx_set_uint(p, "Height", params.renderHeight); - NgxResult created = g.ngx.createFeature(NGX_FEATURE_SUPER_SAMPLING, p, &feature.cudaHandle); + const unsigned int featureId = + feature.rayReconstruction ? NGX_FEATURE_RAY_RECONSTRUCTION : NGX_FEATURE_SUPER_SAMPLING; + NgxResult created = g.ngx.createFeature(featureId, p, &feature.cudaHandle); unsigned int width = params.renderWidth, height = params.renderHeight; if (created != NGX_SUCCESS) { @@ -4202,7 +5542,7 @@ static bool recreate_for_render_size(Feature& feature, FrameParams& params, uint d4r_ngx_set_uint(p, "Width", width); d4r_ngx_set_uint(p, "Height", height); if (width != 0 && height != 0) - created = g.ngx.createFeature(NGX_FEATURE_SUPER_SAMPLING, p, &feature.cudaHandle); + created = g.ngx.createFeature(featureId, p, &feature.cudaHandle); feature.refusedWidth = created == NGX_SUCCESS ? params.renderWidth : 0; feature.refusedHeight = created == NGX_SUCCESS ? params.renderHeight : 0; if (created != NGX_SUCCESS) @@ -4254,6 +5594,24 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram true, false); const auto dilationStart = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{}; const CudaObject dilatedMotion = ok ? prepare_dilated_motion(*feature, slot, params) : 0; + // Ray Reconstruction: the auxiliary surfaces are copied in the same window as the four + // main planes, so the one synchronize below covers them as well. The lock is taken only + // around the copies and, later, around the evaluate - never across the synchronize, which + // is the one long wait here and would otherwise stall the game thread's next gather. + bool auxUploaded = false; + if (ok && feature->rayReconstruction) + { + { + std::lock_guard lock(feature->auxMutex); + auxUploaded = upload_aux_inputs(*feature, slotIndex, frame); + } + if (!auxUploaded) + { + slot.busy = false; + frame_retired(feature, frame); + return; + } + } // Device-to-array copies (VRAM interop, staged uploads) return before they // finish, and NGX does not evaluate on the stream they were issued on, so // without this DLSS can sample a mix of this frame's and the previous @@ -4261,7 +5619,7 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram // thin geometry in motion. D4R_SHIM_INPUT_SYNC=0 restores the old behaviour. static const bool inputSync = env_uint("D4R_SHIM_INPUT_SYNC", 1) != 0; // Linear inputs are already ready; other copies need completion before NGX samples them. - if (ok && (dilatedMotion != 0 || (inputSync && !linearInputs))) + if (ok && (dilatedMotion != 0 || auxUploaded || (inputSync && !linearInputs))) { const int syncResult = synchronize_default_stream(*feature); if (syncResult != 0) @@ -4291,6 +5649,21 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram if (feature->evaluatedFrames++ == 0) params.reset = 1; set_evaluation_parameters(*feature, params); + // Held across the denoiser's evaluate too: the parameter object it is reading names + // &auxSurface.handle, which the game thread must not be able to move or erase meanwhile. + std::unique_lock auxLock(feature->auxMutex, std::defer_lock); + if (feature->rayReconstruction) + { + auxLock.lock(); + publish_aux_parameters(*feature, params, slot); + if (!publish_alpha_output(*feature, slot)) + { + auxLock.unlock(); + slot.busy = false; + frame_retired(feature, frame); + return; + } + } if (linearInputs) { // the parameters point at these variables @@ -4339,9 +5712,13 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram const int eventStartResult = timing.enabled && feature->profileEventsReady ? g.cu.eventRecord(feature->profileStart, nullptr) : -1; const auto evalStart = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{}; + // The core dispatches feature 13 to the denoiser it loaded, so this call is the same + // either way; only the feature id that created it differs. NgxResult result = g.ngx.evaluateFeature(feature->cudaHandle, feature->cudaParams, nullptr); if (result == NGX_FAIL_INVALID_PARAMETER && recreate_for_render_size(*feature, params, frame)) result = g.ngx.evaluateFeature(feature->cudaHandle, feature->cudaParams, nullptr); + if (auxLock.owns_lock()) + auxLock.unlock(); if (feature->outputRedirected && g.cu.outputKernelNative != nullptr && g.cu.outputKernelNative() != 1) { // this frame's output kernel ignored the redirect: publish from the array as usual @@ -4370,6 +5747,15 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram frame_retired(feature, frame); return; } + // The alpha surface is complete once the synchronize above has passed. A failure here drops + // the whole frame, exactly as a failed synchronize does: an accepted frame with an unwritten + // alpha would leave the caller believing it was written. + if (feature->rayReconstruction && !download_alpha_output(*feature, frame)) + { + slot.busy = false; + frame_retired(feature, frame); + return; + } if (timing.enabled) { timing.ngxHost = profile_ms(evalStart, evalReturned); @@ -4716,6 +6102,28 @@ static void* new_parameters() // Copies the official capability values (availability, driver requirements, // and the optimal-settings/stats callbacks, which work on any parameter // implementation) into one of our parameter objects. +// Reads the denoising family out of the core's own capability parameters. Worker only. +static void refresh_denoiser_capability() +{ + void* official = nullptr; + if (g.ngx.getCapabilityParameters == nullptr || g.ngx.getCapabilityParameters(&official) != NGX_SUCCESS || + official == nullptr) + return; + DenoiserCapability capability; + int available = 0; + capability.available = + d4r_ngx_get_int(official, "SuperSamplingDenoising.Available", &available) == NGX_SUCCESS ? available : 0; + int initResult = 0; + capability.initResult = + d4r_ngx_get_int(official, "SuperSamplingDenoising.FeatureInitResult", &initResult) == NGX_SUCCESS + ? initResult + : static_cast(NGX_FAIL_NOT_INITIALIZED); + capability.known = true; + g.ngx.destroyParameters(official); + std::lock_guard lock(g_denoiserCapabilityMutex); + g_denoiserCapability = capability; +} + static NgxResult fill_capabilities(void* parameters) { return g.worker.call([&]() -> NgxResult { @@ -4738,18 +6146,60 @@ static NgxResult fill_capabilities(void* parameters) if (d4r_ngx_get_uint(official, name, &value) == NGX_SUCCESS) d4r_ngx_set_uint(parameters, name, value); } - const char* pointers[] = {"DLSSOptimalSettingsCallback", "DLSSGetStatsCallback"}; + // DLSSDOptimalSettingsCallback is not optional decoration: the SDK's helpers call it to + // work out the render dimensions, so a Ray Reconstruction route that dropped it would + // misbehave in a way that looks like a game bug. Forwarded exactly as the Super + // Sampling pair already is - no new ABI is invented here. + const char* pointers[] = {"DLSSOptimalSettingsCallback", "DLSSGetStatsCallback", + "DLSSDOptimalSettingsCallback", "DLSSDGetStatsCallback"}; for (const char* name : pointers) { void* value = nullptr; if (d4r_ngx_get_void(official, name, &value) == NGX_SUCCESS && value != nullptr) d4r_ngx_set_void(parameters, name, value); } + // The denoising family is the core's own account of the denoiser it loaded, so it is + // passed through rather than re-derived here - and cached, because it is also what + // decides whether Ray Reconstruction may be created at all. + refresh_denoiser_capability(); + const char* denoisingInts[] = {"SuperSamplingDenoising.Available", + "SuperSamplingDenoising.NeedsUpdatedDriver", + "SuperSamplingDenoising.FeatureInitResult"}; + for (const char* name : denoisingInts) + { + int value = 0; + if (d4r_ngx_get_int(official, name, &value) == NGX_SUCCESS) + d4r_ngx_set_int(parameters, name, value); + } + const char* denoisingUnsigned[] = {"SuperSamplingDenoising.MinDriverVersionMajor", + "SuperSamplingDenoising.MinDriverVersionMinor"}; + for (const char* name : denoisingUnsigned) + { + unsigned int value = 0; + if (d4r_ngx_get_uint(official, name, &value) == NGX_SUCCESS) + d4r_ngx_set_uint(parameters, name, value); + } + // D4R_RR_ENABLE=0 is this file's decision and outranks what the core reports: with the + // route switched off, availability reads 0 and the callbacks are cleared, so a game + // cannot reach a denoiser route that is not being offered. + if (!rr_enabled()) + { + d4r_ngx_set_int(parameters, "SuperSamplingDenoising.Available", 0); + d4r_ngx_set_int(parameters, "SuperSamplingDenoising.FeatureInitResult", + static_cast(NGX_FAIL_FEATURE_NOT_SUPPORTED)); + d4r_ngx_set_void(parameters, "DLSSDOptimalSettingsCallback", nullptr); + d4r_ngx_set_void(parameters, "DLSSDGetStatsCallback", nullptr); + } + const DenoiserCapability capability = denoiser_capability(); + logf("denoiser capabilities from the core: available=%d init=0x%08x (Ray Reconstruction %s)", + capability.available, static_cast(capability.initResult), + rr_denial_reason()); g.ngx.destroyParameters(official); return NGX_SUCCESS; }); } + D4R_EXPORT NgxResult NVSDK_NGX_D3D12_GetParameters(void** parameters) { if (!g.ngxInitialized) @@ -4772,13 +6222,16 @@ D4R_EXPORT NgxResult NVSDK_NGX_D3D12_GetCapabilityParameters(void** parameters) return NGX_FAIL_NOT_INITIALIZED; *parameters = new_parameters(); const NgxResult result = fill_capabilities(*parameters); - int available = -1, initResult = -1, needsDriver = -1; + int available = -1, initResult = -1, needsDriver = -1, denoising = -1, denoisingInit = -1; d4r_ngx_get_int(*parameters, "SuperSampling.Available", &available); d4r_ngx_get_int(*parameters, "SuperSampling.FeatureInitResult", &initResult); d4r_ngx_get_int(*parameters, "SuperSampling.NeedsUpdatedDriver", &needsDriver); + d4r_ngx_get_int(*parameters, "SuperSamplingDenoising.Available", &denoising); + d4r_ngx_get_int(*parameters, "SuperSamplingDenoising.FeatureInitResult", &denoisingInit); logf("NVSDK_NGX_D3D12_GetCapabilityParameters -> 0x%08x, SuperSampling.Available=%d FeatureInitResult=0x%08x " - "NeedsUpdatedDriver=%d", - result, available, static_cast(initResult), needsDriver); + "NeedsUpdatedDriver=%d, SuperSamplingDenoising.Available=%d FeatureInitResult=0x%08x", + result, available, static_cast(initResult), needsDriver, denoising, + static_cast(denoisingInit)); return result; } @@ -4805,14 +6258,99 @@ D4R_EXPORT NgxResult NVSDK_NGX_D3D12_GetScratchBufferSize(unsigned int, const vo return NGX_SUCCESS; } -D4R_EXPORT NgxResult NVSDK_NGX_D3D12_GetFeatureRequirements(IUnknown*, const void*, NgxFeatureRequirement* requirement) +// nvsdk_ngx.h: NVSDK_NGX_D3D12_GetFeatureRequirements(IDXGIAdapter *Adapter, +// const NVSDK_NGX_FeatureDiscoveryInfo *FeatureDiscoveryInfo, +// NVSDK_NGX_FeatureRequirement *OutSupported). The feature is NOT a separate argument: it is +// FeatureDiscoveryInfo->FeatureID (defs.h, the field after SDKVersion), and the first argument +// is the display adapter, which must never be read as a feature id. +D4R_EXPORT NgxResult NVSDK_NGX_D3D12_GetFeatureRequirements(void* adapter, const void* discovery, + NgxFeatureRequirement* requirement) { - if (requirement != nullptr) + if (requirement == nullptr) + return NGX_SUCCESS; + // NVSDK_NGX_FeatureDiscoveryInfo: { NVSDK_NGX_Version SDKVersion; NVSDK_NGX_Feature FeatureID; ... }, + // both enums, so FeatureID is the unsigned int at offset 4. + unsigned int feature = 0; + if (discovery != nullptr) + std::memcpy(&feature, static_cast(discovery) + sizeof(unsigned int), sizeof(feature)); + std::memset(requirement, 0, sizeof(*requirement)); + ensure_portable_config(); + // The local switch is checked before the core is asked: with D4R_RR_ENABLE=0 this process + // is not offering Ray Reconstruction, and the core's authoritative "supported" must not + // be passed on to a game that would then create a feature this shim would refuse. + if (feature == NGX_FEATURE_RAY_RECONSTRUCTION && !rr_enabled()) { - std::memset(requirement, 0, sizeof(*requirement)); - requirement->FeatureSupported = 0; // supported - requirement->MinHWArchitecture = 0x160; + requirement->FeatureSupported = 1; // not supported + logf("Ray Reconstruction is not supported: %s", rr_denial_reason()); + return NGX_SUCCESS; } + { + std::lock_guard lock(g.mutex); + // Requirements must not depend on a prior Init. Load the provider without + // initializing NGX here: the later Init retains the game's identity and device. + if (ensure_workers_started() != NGX_SUCCESS || (g.core == nullptr && !load_libraries())) + return NGX_FAIL_PLATFORM_ERROR; + if (feature == NGX_FEATURE_RAY_RECONSTRUCTION) + load_denoiser_library(); + } + // Otherwise the core's answer is authoritative - it consulted the denoiser it loaded - and + // this file only steps in when the core declines. + if (g.ngx.getFeatureRequirements != nullptr) + { + // The adapter this export receives is not what the core's CUDA entry point wants: it + // takes a CUdevice ordinal. Passing the adapter pointer through would read its low 32 + // bits as an ordinal, so the current context's device is asked for instead. + int device = 0; + const NgxResult result = g.worker.call([&] { + if (g.cu.ctxGetDevice != nullptr && g.cu.ctxGetDevice(&device) != 0) + device = 0; // no current context: device 0 is this process's only CUDA device + return g.ngx.getFeatureRequirements(device, discovery, requirement); + }); + if (result == NGX_SUCCESS) + { + logf("GetFeatureRequirements(feature=%u, cuda device %d) from the NGX core: supported=%u minHW=0x%x " + "'%s'", + feature, device, requirement->FeatureSupported, requirement->MinHWArchitecture, + requirement->MinOSVersion); + return result; + } + logf("GetFeatureRequirements(feature=%u) from the NGX core -> 0x%08x; answering here instead", feature, + result); + } + if (feature == NGX_FEATURE_RAY_RECONSTRUCTION) + { + // Requirements are legal before Init. Ask the denoiser's discovery export, + // not its Init stub, without binding NGX to an application prematurely. + if (!g.ngxInitialized && g.denoiser.cudaCapable) + { + using RequirementsFn = NgxResult (*)(void*, const void*, NgxFeatureRequirement*); + RequirementsFn denoiserRequirements = nullptr; + if (load_export(g.denoiserModule, "NVSDK_NGX_D3D12_GetFeatureRequirements", denoiserRequirements)) + { + const NgxResult result = g.worker.call([&] { + return denoiserRequirements(adapter, discovery, requirement); + }); + logf("GetFeatureRequirements before Init from DLSS-Denoiser -> 0x%08x, supported=%u", + result, requirement->FeatureSupported); + return result; + } + } + const bool supported = ray_reconstruction_available(); + if (!supported) + { + requirement->FeatureSupported = 1; + logf("Ray Reconstruction is not supported: %s", + g.ngxInitialized ? rr_denial_reason() : "no CUDA-capable DLSS-Denoiser for pre-Init discovery"); + return NGX_SUCCESS; + } + } + if (feature != NGX_FEATURE_SUPER_SAMPLING && feature != NGX_FEATURE_RAY_RECONSTRUCTION) + { + requirement->FeatureSupported = 1; // not supported + return NGX_SUCCESS; + } + requirement->FeatureSupported = 0; // supported + requirement->MinHWArchitecture = 0x160; return NGX_SUCCESS; } @@ -4854,18 +6392,37 @@ static bool engine_requested(); // Vulkan-only K features retain just their creation settings until a frame actually needs CUDA. static NgxResult create_cuda_feature(Feature& feature) { - const char* presetNames[] = {"DLSS.Hint.Render.Preset.DLAA", "DLSS.Hint.Render.Preset.Quality", - "DLSS.Hint.Render.Preset.Balanced", "DLSS.Hint.Render.Preset.Performance", - "DLSS.Hint.Render.Preset.UltraPerformance", "DLSS.Hint.Render.Preset.UltraQuality"}; + static const char* const kSuperSamplingPresets[] = { + "DLSS.Hint.Render.Preset.DLAA", "DLSS.Hint.Render.Preset.Quality", "DLSS.Hint.Render.Preset.Balanced", + "DLSS.Hint.Render.Preset.Performance", "DLSS.Hint.Render.Preset.UltraPerformance", + "DLSS.Hint.Render.Preset.UltraQuality"}; + static const char* const kRayReconstructionPresets[] = { + "RayReconstruction.Hint.Render.Preset.DLAA", "RayReconstruction.Hint.Render.Preset.Quality", + "RayReconstruction.Hint.Render.Preset.Balanced", "RayReconstruction.Hint.Render.Preset.Performance", + "RayReconstruction.Hint.Render.Preset.UltraPerformance", "RayReconstruction.Hint.Render.Preset.UltraQuality"}; + const bool rayReconstruction = feature.rayReconstruction; + const char* const* presetNames = rayReconstruction ? kRayReconstructionPresets : kSuperSamplingPresets; const NgxResult result = g.worker.call([&]() -> NgxResult { NgxResult allocation = g.ngx.allocateParameters(&feature.cudaParams); if (allocation != NGX_SUCCESS) return allocation; void* p = feature.cudaParams; - d4r_ngx_set_uint(p, "Width", feature.width); - d4r_ngx_set_uint(p, "Height", feature.height); - d4r_ngx_set_uint(p, "OutWidth", feature.outWidth); - d4r_ngx_set_uint(p, "OutHeight", feature.outHeight); + if (rayReconstruction) + { + // The CUDA denoiser reads creation dimensions through Get(int*), + // unlike Super Sampling's Get(unsigned int*) contract. + d4r_ngx_set_int(p, "Width", static_cast(feature.width)); + d4r_ngx_set_int(p, "Height", static_cast(feature.height)); + d4r_ngx_set_int(p, "OutWidth", static_cast(feature.outWidth)); + d4r_ngx_set_int(p, "OutHeight", static_cast(feature.outHeight)); + } + else + { + d4r_ngx_set_uint(p, "Width", feature.width); + d4r_ngx_set_uint(p, "Height", feature.height); + d4r_ngx_set_uint(p, "OutWidth", feature.outWidth); + d4r_ngx_set_uint(p, "OutHeight", feature.outHeight); + } d4r_ngx_set_int(p, "PerfQualityValue", feature.quality); d4r_ngx_set_int(p, "DLSS.Feature.Create.Flags", feature.flags); d4r_ngx_set_int(p, "DLSS.Enable.Output.Subrects", feature.subrects); @@ -4873,6 +6430,9 @@ static NgxResult create_cuda_feature(Feature& feature) d4r_ngx_set_uint(p, "VisibilityNodeMask", 1); for (int index = 0; index < 6; ++index) d4r_ngx_set_uint(p, presetNames[index], feature.creationPresets[index]); + if (rayReconstruction) + for (size_t index = 0; index < sizeof(kAuxScalars) / sizeof(kAuxScalars[0]); ++index) + d4r_ngx_set_int(p, kAuxScalars[index].name, feature.auxScalarBaseline[index]); if (g.cu.memAlloc(&feature.scratch, 64ull * 1024 * 1024) == 0) { d4r_ngx_set_void(p, "Scratch", reinterpret_cast(static_cast(feature.scratch))); @@ -4880,7 +6440,10 @@ static NgxResult create_cuda_feature(Feature& feature) } else feature.scratch = 0; - const NgxResult created = g.ngx.createFeature(NGX_FEATURE_SUPER_SAMPLING, p, &feature.cudaHandle); + // The core owns the denoiser: it loads nvngx_dlssd.dll, knows its kernel map, and + // dispatches feature 13 to it. Only the feature id differs from Super Sampling. + const NgxResult created = g.ngx.createFeature( + rayReconstruction ? NGX_FEATURE_RAY_RECONSTRUCTION : NGX_FEATURE_SUPER_SAMPLING, p, &feature.cudaHandle); if (created != NGX_SUCCESS) { if (feature.cudaHandle != nullptr) g.ngx.releaseFeature(feature.cudaHandle); @@ -4926,7 +6489,8 @@ static NgxResult create_cuda_feature(Feature& feature) } return created; }); - logf("NVSDK_NGX_CUDA_CreateFeature -> 0x%08x", result); + logf("NVSDK_NGX_CUDA_CreateFeature(feature=%u) -> 0x%08x", + rayReconstruction ? NGX_FEATURE_RAY_RECONSTRUCTION : NGX_FEATURE_SUPER_SAMPLING, result); if (result == NGX_SUCCESS) feature.cudaCreationDeferred = false; return result; } @@ -4946,11 +6510,31 @@ static NgxResult create_feature(unsigned int featureId, void* parameters, NgxHan logf("DLSS requires lifetime-aware d4r d3d12.dll/d3d12core.dll; update both files with the shim"); return NGX_FAIL_PLATFORM_ERROR; } - if (featureId != NGX_FEATURE_SUPER_SAMPLING || parameters == nullptr || handle == nullptr) + // Two features are served here. Super Sampling is feature 1 and Ray Reconstruction + // feature 13; the NGX core dispatches the second to the denoiser it loaded, and is + // offered feature 13 only once it has reported that denoiser working. + const bool rayReconstruction = featureId == NGX_FEATURE_RAY_RECONSTRUCTION; + if (featureId != NGX_FEATURE_SUPER_SAMPLING && !rayReconstruction) { logf("NVSDK_NGX_D3D12_CreateFeature(feature=%u) unsupported", featureId); return NGX_FAIL_FEATURE_NOT_SUPPORTED; } + if (rayReconstruction) + { + // A game may create feature 13 without ever asking for capabilities, so the core's + // account of the denoiser is fetched here if it has not been already. + g.worker.call([] { refresh_denoiser_capability(); return 0; }); + } + if (rayReconstruction && !ray_reconstruction_available()) + { + logf("NVSDK_NGX_D3D12_CreateFeature(feature=%u) unavailable: %s", featureId, rr_denial_reason()); + return NGX_FAIL_FEATURE_NOT_SUPPORTED; + } + if (parameters == nullptr || handle == nullptr) + { + logf("NVSDK_NGX_D3D12_CreateFeature(feature=%u): no parameters or handle", featureId); + return NGX_FAIL_INVALID_PARAMETER; + } auto* feature = new Feature(); feature->nativeVulkan = nativeVulkan; feature->width = get_uint_or(parameters, "Width", 0); @@ -4962,15 +6546,24 @@ static NgxResult create_feature(unsigned int featureId, void* parameters, NgxHan feature->quality = quality; feature->flags = flags; const int subrects = get_int_or(parameters, "DLSS.Enable.Output.Subrects", 0); - const char* presetNames[] = {"DLSS.Hint.Render.Preset.DLAA", "DLSS.Hint.Render.Preset.Quality", - "DLSS.Hint.Render.Preset.Balanced", "DLSS.Hint.Render.Preset.Performance", - "DLSS.Hint.Render.Preset.UltraPerformance", "DLSS.Hint.Render.Preset.UltraQuality"}; + // The render preset is chosen by two independent knobs. D4R_DLSS_PRESET forces the + // Super Sampling model (NVSDK_NGX_DLSS_Hint_Render_Preset, e.g. 5 = E) and touches only + // Super Sampling features; D4R_RR_PRESET forces the Ray Reconstruction model + // (NVSDK_NGX_RayReconstruction_Hint_Render_Preset, D = 4, E = 5) and touches only + // Ray Reconstruction features. Neither one silently becomes the other. + static const char* const kSuperSamplingPresets[] = { + "DLSS.Hint.Render.Preset.DLAA", "DLSS.Hint.Render.Preset.Quality", "DLSS.Hint.Render.Preset.Balanced", + "DLSS.Hint.Render.Preset.Performance", "DLSS.Hint.Render.Preset.UltraPerformance", + "DLSS.Hint.Render.Preset.UltraQuality"}; + static const char* const kRayReconstructionPresets[] = { + "RayReconstruction.Hint.Render.Preset.DLAA", "RayReconstruction.Hint.Render.Preset.Quality", + "RayReconstruction.Hint.Render.Preset.Balanced", "RayReconstruction.Hint.Render.Preset.Performance", + "RayReconstruction.Hint.Render.Preset.UltraPerformance", "RayReconstruction.Hint.Render.Preset.UltraQuality"}; + const char* const* presetNames = rayReconstruction ? kRayReconstructionPresets : kSuperSamplingPresets; unsigned int presets[6]; for (int index = 0; index < 6; ++index) presets[index] = get_uint_or(parameters, presetNames[index], 0); - // D4R_DLSS_PRESET forces one render preset (NVSDK_NGX_DLSS_Hint_Render_Preset - // value, e.g. 5 = E) for every quality mode. - const std::string forced = env_string("D4R_DLSS_PRESET"); + const std::string forced = env_string(rayReconstruction ? "D4R_RR_PRESET" : "D4R_DLSS_PRESET"); if (!forced.empty()) { const unsigned int value = static_cast(strtoul(forced.c_str(), nullptr, 0)); @@ -4985,8 +6578,8 @@ static NgxResult create_feature(unsigned int featureId, void* parameters, NgxHan feature->preset = presets[1]; // Direct output needs the native output kernel of the preset in use (verified: K = 11 through // hiluma_engine_output, M = 13 through rrlite_downsample_kernel, its last kernel). Only a forced - // preset qualifies. - if (!forced.empty()) + // Super Sampling preset qualifies; a Ray Reconstruction preset never does here. + if (!forced.empty() && !rayReconstruction) { char allowed[128] = "11,13"; if (const std::string list = env_string("D4R_SHIM_OUTPUT_DIRECT_PRESETS"); !list.empty()) @@ -4994,8 +6587,13 @@ static NgxResult create_feature(unsigned int featureId, void* parameters, NgxHan for (char* token = strtok(allowed, ","); token != nullptr; token = strtok(nullptr, ",")) feature->outputDirectAllowed |= strtoul(token, nullptr, 0) == presets[1]; } - logf("NVSDK_NGX_D3D12_CreateFeature: %ux%u -> %ux%u quality=%d flags=0x%x subrects=%d preset=%u", feature->width, - feature->height, feature->outWidth, feature->outHeight, quality, flags, subrects, presets[1]); + feature->rayReconstruction = rayReconstruction; + if (rayReconstruction) + for (size_t index = 0; index < sizeof(kAuxScalars) / sizeof(kAuxScalars[0]); ++index) + feature->auxScalarBaseline[index] = get_int_or(parameters, kAuxScalars[index].name, 0); + logf("NVSDK_NGX_D3D12_CreateFeature(feature=%u%s): %ux%u -> %ux%u quality=%d flags=0x%x subrects=%d preset=%u", + featureId, rayReconstruction ? ", Ray Reconstruction" : "", feature->width, feature->height, + feature->outWidth, feature->outHeight, quality, flags, subrects, presets[1]); feature->subrects = subrects; std::copy(std::begin(presets), std::end(presets), feature->creationPresets.begin()); @@ -5048,16 +6646,6 @@ static Feature* find_feature(const NgxHandle* handle) return nullptr; } -static ID3D12Resource* get_resource(void* parameters, const char* name) -{ - ID3D12Resource* resource = nullptr; - if (d4r_ngx_get_d3d12_resource(parameters, name, &resource) == NGX_SUCCESS && resource != nullptr) - return resource; - void* raw = nullptr; - if (d4r_ngx_get_void(parameters, name, &raw) == NGX_SUCCESS) - return static_cast(raw); - return nullptr; -} // --- Native engine backend (engine/): no NGX evaluation, CUDA, worker or frame split ------------- // [Engine] Enabled in d4r.ini (D4R_ENGINE=1) runs presets K and M with d4r's own shaders inside the game's command @@ -6295,6 +7883,42 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand p.outputBaseY = get_uint_or(parameters, "DLSS.Output.Subrect.Base.Y", 0); p.hasExposure = exposure != nullptr; + // Ray Reconstruction: the denoiser's own per-frame settings and camera matrices are read + // here, on the game thread, while the caller's parameter object is certainly alive, and + // carried to the worker by value. + if (feature->rayReconstruction) + { + for (size_t index = 0; index < sizeof(kAuxScalars) / sizeof(kAuxScalars[0]); ++index) + { + int value = 0; + const bool given = d4r_ngx_get_int(parameters, kAuxScalars[index].name, &value) == NGX_SUCCESS; + // Not given this frame: back to what the caller asked for at creation, so a + // setting it has stopped varying does not keep an old per-frame override. + p.auxScalar[index] = given ? value : feature->auxScalarBaseline[index]; + p.auxScalarSet[index] = true; + } + // Each guide's Subrect origin, read here while the caller's object is alive. An origin + // the caller stops registering resets to zero rather than lingering. + for (size_t index = 0; index < sizeof(kAuxSubrectNames) / sizeof(kAuxSubrectNames[0]); ++index) + { + static std::string nameX, nameY; + nameX.assign(kAuxSubrectNames[index]).append(".Subrect.Base.X"); + nameY.assign(kAuxSubrectNames[index]).append(".Subrect.Base.Y"); + p.auxSubrect[index][0] = get_uint_or(parameters, nameX.c_str(), 0); + p.auxSubrect[index][1] = get_uint_or(parameters, nameY.c_str(), 0); + } + for (size_t index = 0; index < sizeof(kAuxMatrices) / sizeof(kAuxMatrices[0]); ++index) + { + // A matrix is registered as a pointer to float[16] (nvsdk: the SDK's helpers + // set these with SetVoidPointer), never as sixteen separate scalars. + void* value = nullptr; + p.auxMatrixSet[index] = + d4r_ngx_get_void(parameters, kAuxMatrices[index], &value) == NGX_SUCCESS && value != nullptr; + if (p.auxMatrixSet[index]) + std::memcpy(p.auxMatrix[index], value, sizeof(p.auxMatrix[index])); + } + } + if (engine_requested() && engine_evaluate(*feature, list, color, depth, motion, output, exposure, p, frame)) return NGX_SUCCESS; @@ -6340,7 +7964,11 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand const Plane planes[4] = {Plane::Color, Plane::Depth, Plane::Motion, Plane::Exposure}; const int inputCount = exposure != nullptr ? 4 : 3; VramCopy vramInputs[4], vramOutput; - if ((!feature->vramDecided || feature->vram) && vram_interop_available()) + // Alpha readback uses the host result ring; color must use the same pipeline so both + // outputs can present one produced frame. The existing cutover below drains VRAM work. + const bool alphaOutputRequested = + feature->rayReconstruction && get_resource(parameters, "DLSSD.OutputAlpha") != nullptr; + if (!alphaOutputRequested && (!feature->vramDecided || feature->vram) && vram_interop_available()) { p.vram = describe_vram_copy(output, Plane::Color, vramOutput, true); for (int index = 0; index < inputCount && p.vram; ++index) @@ -6503,8 +8131,6 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand if (exposure != nullptr) copy_to_staging(list, exposure, slot.exposure, inputState); } - if (timing.enabled) - timing.inputRecord = profile_ms(inputRecordStart, ProfileClock::now()); // Frame marker, written once the copies above have completed. ID3D12GraphicsCommandList2* list2 = nullptr; @@ -6513,6 +8139,29 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand logf("ID3D12GraphicsCommandList2 unavailable; cannot place frame marker"); return NGX_FAIL_PLATFORM_ERROR; } + // Ray Reconstruction's auxiliary surfaces go into their own readback buffers here, still + // ahead of the frame marker so the worker can read them the moment that marker lands. + if (feature->rayReconstruction) + { + if (!capture_alpha_output(*feature, parameters, slotIndex, frame) || + !gather_aux_inputs(*feature, list, parameters, slotIndex, inputState, frame)) + { + // capture_alpha_output may already have taken a reference on the caller's texture, + // and this frame is never queued, so nothing else would release it. + InputSlot& failed = feature->inputs[slotIndex]; + if (failed.alphaResource != nullptr) + { + failed.alphaResource->Release(); + failed.alphaResource = nullptr; + failed.alphaRequested = false; + } + list2->Release(); + return NGX_FAIL_PLATFORM_ERROR; + } + } + if (timing.enabled) + timing.inputRecord = profile_ms(inputRecordStart, ProfileClock::now()); + const D3D12_GPU_VIRTUAL_ADDRESS markerAddress = feature->marker->GetGPUVirtualAddress(); D3D12_WRITEBUFFERIMMEDIATE_PARAMETER markers[] = { {markerAddress + sizeof(uint32_t) * (1 + slotIndex), frame}, @@ -6526,14 +8175,16 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand // Present the most recent finished DLSS result. int latest; + uint32_t colorFrame = 0; // the frame whose colour is presented below, 0 if none { std::lock_guard lock(feature->outputMutex); latest = feature->latestOutput.load(); if (latest >= 0) { feature->outputs[latest].lastReadFrame = frame; + colorFrame = feature->outputs[latest].producedFrame.load(); if (timing.enabled) - timing.presentedFrame = feature->outputs[latest].producedFrame.load(); + timing.presentedFrame = colorFrame; } } const auto outputRecordStart = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{}; @@ -6547,8 +8198,14 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand } if (p.split) { - // Runs in the second half of the list, once frame's result is in place. + // Runs in the second half of the list, once frame's result is in place. Split frames + // present their own result rather than the newest one, so the alpha must pair with + // that one instead. const int target = static_cast(frame % kOutputSlots); + { + std::lock_guard lock(feature->outputMutex); + colorFrame = feature->outputs[target].producedFrame.load(); + } if (p.outputInPlace == 0) { if (p.packedOutput) @@ -6583,6 +8240,8 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand list->CopyTextureRegion(&destination, 0, 0, 0, &source, nullptr); transition(list, output, D3D12_RESOURCE_STATE_COPY_DEST, outputState); } + if (feature->rayReconstruction) + record_alpha_copyback(*feature, list, slot, outputState, colorFrame); if (timing.enabled) { timing.outputRecord = profile_ms(outputRecordStart, ProfileClock::now()); @@ -6636,6 +8295,20 @@ static void release_feature(Feature* feature) g.ngx.releaseFeature(feature->cudaHandle); if (feature->cudaParams != nullptr) g.ngx.destroyParameters(feature->cudaParams); + // Every denoiser input dies with its feature, drains to nothing first. + for (AuxSurface& surface : feature->aux) + destroy_aux_surface(surface); + feature->aux.clear(); + if (feature->alphaImage.object != 0) + g.cu.surfObjectDestroy(feature->alphaImage.object); + if (feature->alphaImage.array != nullptr) + g.cu.arrayDestroy(feature->alphaImage.array); + release_host(feature->alphaHost); + for (AlphaSlot& slot : feature->alpha) + slot.staging.release(); + for (InputSlot& slot : feature->inputs) + if (slot.alphaResource != nullptr) + slot.alphaResource->Release(); for (InputSlot& slot : feature->inputs) { for (CudaObject& texture : slot.linearTexture) @@ -6744,6 +8417,7 @@ static void finish_shutdown() drain_pipeline(); g.worker.call([] { return g.ngx.shutdown(); }); } + // No separate denoiser shutdown: the core loaded it and the core's shutdown releases it. g.ngxInitialized = false; g.shutdownPending = false; if (g.device != nullptr) diff --git a/tools/replay_dlss_surface_half_store.cpp b/tools/replay_dlss_surface_half_store.cpp index 458a093..473bf3f 100644 --- a/tools/replay_dlss_surface_half_store.cpp +++ b/tools/replay_dlss_surface_half_store.cpp @@ -1,3 +1,4 @@ +// Modified in this fork for CUDA Ray Reconstruction support and validation (2026). #include #include @@ -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 -// sust.b.2d.b32.zero) into a surface of the given CUDA array format. X is a byte offset (unformatted surface addressing). The expected -// result is the fill pattern with the four stored bytes replaced at X, or the -// untouched fill pattern when X is out of bounds. +// Unformatted packed stores address X in bytes. Check every byte, including the +// untouched pixels around a write and row boundaries. 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] = {}; - const char* functionNames[2] = {"sust.b.2d.v2.b16", "sust.b.2d.b32"}; + CUfunction functions[4] = {}; + 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]); - for (size_t functionIndex = 0; functionIndex < 2; ++functionIndex) + size_t total = 0; + 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 surfaceBytes = expected; - const uint8_t stored[4] = { + uint8_t stored[8] = { static_cast(storeCase.value0), static_cast(storeCase.value0 >> 8), - static_cast(storeCase.value1), static_cast(storeCase.value1 >> 8)}; - if (storeCase.x + 4 <= rowBytes && storeCase.y < height) - std::memcpy(expected.data() + storeCase.y * rowBytes + storeCase.x, stored, 4); + static_cast(storeCase.value1), static_cast(storeCase.value1 >> 8), + static_cast(storeCase.value2), static_cast(storeCase.value2 >> 8), + static_cast(storeCase.value3), static_cast(storeCase.value3 >> 8)}; + if (storeCase.byteVector) + { + stored[0] = static_cast(storeCase.value0); + stored[1] = static_cast(storeCase.value1); + stored[2] = static_cast(storeCase.value2); + stored[3] = static_cast(storeCase.value3); + } + if (static_cast(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(storeCase.value1) << 16); + uint64_t packed = static_cast(storeCase.value0) | + (static_cast(storeCase.value1) << 16) | + (static_cast(storeCase.value2) << 32) | + (static_cast(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); diff --git a/tools/replay_dlss_texture_sample.cpp b/tools/replay_dlss_texture_sample.cpp index 9f1208f..a1209da 100644 --- a/tools/replay_dlss_texture_sample.cpp +++ b/tools/replay_dlss_texture_sample.cpp @@ -1,3 +1,4 @@ +// Modified in this fork for CUDA Ray Reconstruction support and validation (2026). #include #include @@ -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((fraction * 2.0f - 1.0f) * 1024.0f); return static_cast(((exponent - 1 + 15) << 10) | mantissa); }; - const size_t outputFloats = static_cast(width) * height * 16; + const size_t outputFloats = static_cast(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);