From 69b284b8b711474cbe474d23d06578ee5ab75b3f Mon Sep 17 00:00:00 2001 From: woutermans Date: Wed, 7 Oct 2026 22:39:37 +0200 Subject: [PATCH] Fix PR integration: token prep slots, input ordering and output lifetime --- docs/building.md | 9 ++-- docs/native-kernels.md | 1 + docs/performance.md | 2 + kernels/m/swin_tok.h | 19 +++++--- packaging/RELEASE_NOTES.md | 8 ++++ packaging/SOURCES.txt | 2 +- scripts/build_release_glibc241.sh | 2 +- tests/test_input_readiness.py | 73 ------------------------------- tools/d3d12_dlss_harness.cpp | 14 ++++++ tools/d4r_nvngx_shim.cpp | 41 ++++++++++------- 10 files changed, 69 insertions(+), 102 deletions(-) delete mode 100644 tests/test_input_readiness.py diff --git a/docs/building.md b/docs/building.md index 3747463..3a818db 100644 --- a/docs/building.md +++ b/docs/building.md @@ -5,7 +5,7 @@ These steps produce the pieces the launcher needs: a patched ZLUDA, a patched vk | Component | Tested version | |---|---| | ZLUDA | `ee2f25a` (upstream), plus `patches/zluda/0002` through `0008` in order | -| vkd3d-proton | `3dfc6f07` (the base GE-Proton11-3 ships), plus `patches/vkd3d-proton/0001` and `0002` | +| vkd3d-proton | `3dfc6f07` (the base GE-Proton11-3 ships), plus `patches/vkd3d-proton/0001` through `0003` in order | | ROCm | 7.2 (HIP runtime, clang, device libraries) | | Proton | GE-Proton11-3 (with its OptiScaler integration) | | DLSS | `nvngx_dlss.dll` 310.7.0 | @@ -54,7 +54,7 @@ When linking on a system with ROCm libraries outside the default search path, in scripts/build_vkd3d_proton_d4r.sh ~/.cache/d4r-vkd3d-d4r ``` -The patch lets the shim split the game's command list at the DLSS call, so the rest of the frame waits (at queue level) for DLSS instead of showing the previous frame's result. Point `D4R_VKD3D_DIR` (or `VkD3DDir` in d4r.ini) at the output. Without it the shim falls back to showing a one-frame-old result. +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. ## 3. Shim, bridge and NVAPI identity @@ -133,7 +133,7 @@ Mount the repository at `/work`, ROCm's compiler/headers/device libraries at `/o |---|---| | `D4R_RELEASE_BUILD_ROOT` | writable directory under `/work/build/` | | `D4R_ZLUDA_SRC` | isolated checkout inside that directory, with both pinned LLVM and HiGHS submodules, the real OCKL LFS payload, and patches `0002`–`0008` applied | -| `D4R_VKD3D_SRC` | isolated vkd3d-proton checkout inside that directory, with its submodules and patches `0001` and `0002` applied | +| `D4R_VKD3D_SRC` | isolated vkd3d-proton checkout inside that directory, with its submodules and patches `0001`–`0003` applied | | `CARGO_HOME` | writable build-local Cargo cache; prefetch the locked dependencies for an offline build | | `D4R_ROCM_DIR` | `/opt/rocm` | | `D4R_ROCM_LINK_STUBS` | optional directory for additional ROCm link libraries | @@ -147,4 +147,7 @@ Run `bash scripts/build_release_glibc241.sh` inside the container. The script ch - `python3 -m unittest discover -s tests -v` includes OptiScaler INI generation, full/clean package staging and NGX-routing checker regressions. Package tests use fixture binaries and GPU-build stand-ins; real DLL loading still needs the OptiScaler/Proton harness. - `scripts/check_environment.sh` lists the tools, GPU and Proton builds it finds. - The D3D12 harness (`scripts/run_d3d12_dlss_harness_proton.sh`) drives DLSS outside a game. + - `D4R_HARNESS_MOTION_SCENE=1 D4R_HARNESS_SAVE_FRAMES=1` saves a temporal sequence for byte comparisons. Compare the regular output path against `D4R_SHIM_OUTPUT_PACKED=1` and/or `D4R_SHIM_INPLACE=1 VKD3D_D4R_LINEAR_TEXTURES=1`, with VRAM interop, split frames and direct output enabled; `D4R_HARNESS_OUTPUT_RGB10A2=1` selects packed output. + - `D4R_HARNESS_RECREATE_OUTPUT=1` releases and recreates the game's output texture between completed frames; compare its copy and in-place sequences to catch imported-memory handle reuse. + - `D4R_HARNESS_RECREATE=4 D4R_HARNESS_VERIFY_RECREATION=1` checks repeated feature/quality recreation. `D4R_HARNESS_SKIP_RECORDED_INPUT=1` checks discarded command-list input readiness without injecting stale history. - Native kernels have their own validation path; see [native-kernels.md](native-kernels.md). diff --git a/docs/native-kernels.md b/docs/native-kernels.md index acc0388..6ca72e5 100644 --- a/docs/native-kernels.md +++ b/docs/native-kernels.md @@ -74,6 +74,7 @@ In that build the GEMM activations live in LDS as e4m3 bytes (`SWIN_A8` in `swin - The MLP clamp uses the NaN-propagating minimum / maximum; it differs from minnum / maxnum only for a NaN input, whose product is NaN either way. - `mix2` (all exact FP8 layers, enc2 and dec2 included) rounds in place: the accumulator is both addend and destination, which removes a register copy per pair. - `-DSWIN_NO_TOKEN_LANES` builds these layers from `swin_block.h` again. enc2 and dec2 (96 channels, four heads) use the same template: byte-identical on the same recorded launches and harness sequences, 0.22 → 0.14 and 0.20 → 0.14 ms at 1280×720 render, 0.75 → 0.47 and 0.69 → 0.47 ms at 2560×1440. +- The tube retains upstream's eight weight-image slots plus an overflow slot. Weights and per-channel constants use the same slot; `d4r_prep_identity_at` lets MicroCUDA reuse each image only while its allocation identity and write epoch remain valid. ZLUDA still prepares each tube launch. A 12-frame RX 9070 XT harness sequence matched the token-disabled layers byte-for-byte; MicroCUDA prepared the six tube images once and skipped the following 66 preparations, and repeated feature recreation remained byte-identical. - Tried, no measurable gain: interleaving the WMMAs of independent k32 steps. Skipping windows that lie entirely outside the image was also exact but below the measurement noise (at most 2.4% of the windows of every other tube launch). Not tried: true16 code generation (the half-register conversions would save about 3% of the instructions). **Texture tails on the native-FP8 path.** The enc0 tail and the dec0 head encode their results with `v_cvt_pk_fp8_f32` on the f32 copies of the f16 values (`e4m3x4_f32` in `tex_common.h`) instead of ZLUDA's integer `e4m3x2`. With the tails' clamps (±2π for enc0, ±448 for dec0, as a median of three) the codes are equal for every non-NaN f16, checked over all 65536 patterns on an RX 9070 XT; a NaN in dec0 now encodes as −448 instead of the NaN code. Harness frames are byte-identical in both sets (LDR and HDR, motion vectors at render and display resolution). dec0: 0.21 → 0.15 ms at 1280×720 render and 0.74 → 0.53 ms at 2560×1440; enc0 changes by 1–2% (its time is ZLUDA's compile of the feature code). `-DD4R_TEX_SOFT_E4M3` (`D4R_TEX_CFLAGS`) keeps `e4m3x2`. diff --git a/docs/performance.md b/docs/performance.md index 2a43e21..aae56ff 100644 --- a/docs/performance.md +++ b/docs/performance.md @@ -65,6 +65,8 @@ After DLSS, the result normally goes from the shared buffer into the game's outp - `D4R_SHIM_OUTPUT_PACKED=1` (preset M): the native output kernel stores R10G10B10A2 texels itself, rounding as the RGBA16F store followed by the blit does, so the conversion image and blit are not needed. - `D4R_SHIM_INPLACE=1` with `VKD3D_D4R_LINEAR_TEXTURES=1` (vkd3d-proton patch 0003): committed single-level RGBA16F and R10G10B10A2 textures are created linear on exportable memory, the shim imports the output texture into HIP, and the output kernel writes it directly. +Both require VRAM interop, split frames and `D4R_SHIM_OUTPUT_DIRECT=1`; packed/in-place output remains off by default. In-place imports retain their source textures for the process lifetime to prevent recycled Vulkan memory handles from selecting stale CUDA mappings. Recreating output textures therefore retains their VRAM until exit; this is an experimental path, not a recommendation for general game use. + Measured on an RX 9070 XT (gfx1201) in the D3D12 harness, preset M at 1280×720 → 3840×2160, median frame interval of 4 alternating runs of 400 frames. The harness runs DLSS back to back with no game rendering. | Output texture | Passes after DLSS | Frame interval | diff --git a/kernels/m/swin_tok.h b/kernels/m/swin_tok.h index f6dc808..4335be6 100644 --- a/kernels/m/swin_tok.h +++ b/kernels/m/swin_tok.h @@ -656,28 +656,33 @@ __device__ __forceinline__ void swin_tok_block(const CommonParams& p, const Tube // Module boilerplate: the weight and constant images, their prep kernel, and the 4-wave main kernel #define SWIN_TOK_MODULE(NAME, C, NH, NPM, TUBE, PARAMS, CIN) \ using NAME##_K = TokLayout; \ - static_assert(SWIN_PREP_SLOTS == 0, "one prepared image"); \ - __device__ wslot g_w16[NAME##_K::TOTAL]; \ - __device__ h8 g_c16[NAME##_K::CTOTAL]; \ + constexpr int NAME##_S = SWIN_PREP_SLOTS; \ + __device__ wslot g_w16[(NAME##_S + 1) * NAME##_K::TOTAL]; \ + __device__ h8 g_c16[(NAME##_S + 1) * NAME##_K::CTOTAL]; \ + extern "C" __device__ __attribute__((used)) uint64_t g_prep_keys[NAME##_S + 1] = {}; \ __constant__ TokTables g_tok = make_tok_tables(); \ extern "C" __device__ __attribute__((used)) uint32_t d4r_block_z = 4; \ extern "C" __device__ __attribute__((used)) uint32_t d4r_prep_blocks = \ ((NAME##_K::TOTAL > NAME##_K::CTOTAL ? NAME##_K::TOTAL : NAME##_K::CTOTAL) + 127) / 128; \ extern "C" __device__ __attribute__((used)) uint32_t d4r_prep_key_at = SWIN_PREP_KEY_AT; \ - extern "C" __device__ __attribute__((used)) uint32_t d4r_prep_key_slots = 0; \ + extern "C" __device__ __attribute__((used)) uint32_t d4r_prep_key_slots = NAME##_S; \ + extern "C" __device__ __attribute__((used)) uint32_t d4r_prep_identity_at = SWIN_PREP_IDENTITY_AT; \ extern "C" __global__ void __launch_bounds__(128) NAME##_prep(PARAMS p) \ { \ const int idx = blockIdx.x * 128 + threadIdx.x; \ const uint8_t* w = ((const CommonParams*)&p)->w; \ + const int slot = NAME##_S ? prep_slot_claim(g_prep_keys, (uint64_t)w, idx == 0) : 0; \ if (idx < NAME##_K::TOTAL) \ - expand_tok_weights(w, g_w16, g_tok.d, NAME##_K::NDESC, idx); \ + expand_tok_weights(w, g_w16 + slot * NAME##_K::TOTAL, g_tok.d, NAME##_K::NDESC, idx); \ if (idx < NAME##_K::CTOTAL) \ - expand_tok_consts(w, g_c16, g_tok.c, NAME##_K::NCONST, idx); \ + expand_tok_consts(w, g_c16 + slot * NAME##_K::CTOTAL, g_tok.c, NAME##_K::NCONST, idx); \ } \ extern "C" __device__ __attribute__((used)) uint32_t d4r_grid_x = 0; \ extern "C" __global__ void __launch_bounds__(128) SWIN_TOK_VGPR_ATTR NAME(PARAMS p) \ { \ const CommonParams& cp = *(const CommonParams*)&p; \ + const int slot = NAME##_S ? prep_slot_find(g_prep_keys, (uint64_t)cp.w) : 0; \ const int gx = (cp.tw + cp.sx + 7) / 8; \ - swin_tok_block(cp, (const TubeParams*)&p, g_w16, g_c16, blockIdx.x, blockIdx.y, gx); \ + swin_tok_block(cp, (const TubeParams*)&p, g_w16 + slot * NAME##_K::TOTAL, \ + g_c16 + slot * NAME##_K::CTOTAL, blockIdx.x, blockIdx.y, gx); \ } diff --git a/packaging/RELEASE_NOTES.md b/packaging/RELEASE_NOTES.md index 58aad49..9b87dbf 100644 --- a/packaging/RELEASE_NOTES.md +++ b/packaging/RELEASE_NOTES.md @@ -1,3 +1,11 @@ +# Unreleased + +- Merge the RDNA4 token-lane Swin and hardware-FP8 texture-tail optimizations, preserving MicroCUDA's allocation-identity tube weight cache. +- Add opt-in packed R10G10B10A2 and in-place DLSS output, with vkd3d-proton patches `0001`–`0003`. In-place mappings retain their source textures until process exit; output recreation can increase retained VRAM. +- Fix deferred input-marker ordering: wait for the game's GPU copies before CUDA uploads or diagnostic readbacks, not just before NGX evaluation. +- 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. + # 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 aea1ba1..7a932c0 100644 --- a/packaging/SOURCES.txt +++ b/packaging/SOURCES.txt @@ -15,7 +15,7 @@ Third-party software, unmodified or patched by d4r: Licenses: d4r/licenses/ZLUDA-LICENSE-*.txt, d4r/licenses/LLVM-LICENSE.txt d3d12.dll, d3d12core.dll vkd3d-proton (GNU LGPL 2.1), https://github.com/HansKristian-Work/vkd3d-proton - commit @VKD3D_COMMIT@ with d4r/source/patches/vkd3d-proton/0001 and 0002 applied, + commit @VKD3D_COMMIT@ with d4r/source/patches/vkd3d-proton/0001 to 0003 applied, built with its build-win64.txt cross file (meson --buildtype release). Licenses: d4r/licenses/vkd3d-proton-*.txt d4r/rocm/lib/ ROCm 7.2.4 runtime, unmodified, from AMD's Ubuntu 22.04 packages diff --git a/scripts/build_release_glibc241.sh b/scripts/build_release_glibc241.sh index 28e5b31..95a1d7b 100755 --- a/scripts/build_release_glibc241.sh +++ b/scripts/build_release_glibc241.sh @@ -5,7 +5,7 @@ set -euo pipefail ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)" : "${D4R_RELEASE_BUILD_ROOT:?set the build output directory}" : "${D4R_ZLUDA_SRC:?set a build-local ZLUDA checkout with patches 0002-0008 applied}" -: "${D4R_VKD3D_SRC:?set a build-local vkd3d-proton checkout with patches 0001 and 0002 applied}" +: "${D4R_VKD3D_SRC:?set a build-local vkd3d-proton checkout with patches 0001-0003 applied}" : "${D4R_ROCM_DIR:?set ROCm with clang and device libraries}" BUILD="$(realpath -m "$D4R_RELEASE_BUILD_ROOT")" [[ "$(getconf GNU_LIBC_VERSION)" == 'glibc 2.41' ]] || { diff --git a/tests/test_input_readiness.py b/tests/test_input_readiness.py deleted file mode 100644 index 4719bd3..0000000 --- a/tests/test_input_readiness.py +++ /dev/null @@ -1,73 +0,0 @@ -"""Exercise the shim's marker loop with discarded and delayed recordings.""" -from pathlib import Path -import subprocess -import tempfile -import unittest - -ROOT = Path(__file__).resolve().parents[1] - - -class InputReadinessTests(unittest.TestCase): - def test_exact_input_slot_before_evaluation(self): - source = (ROOT / 'tools/d4r_nvngx_shim.cpp').read_text() - start = source.index(' // A later input marker') - end = source.index(' const auto ready = ProfileClock::now();', start) - runner_source = r''' -#include -#include -#include -#include -#include -using ProfileClock = std::chrono::steady_clock; -struct Feature { uint32_t* markerValue; std::atomic retiring{false}; }; -struct InputSlot { bool busy = true; }; -static unsigned retired = 0, evaluated = 0, sleeps = 0; -static uint32_t* arrival = nullptr; -static uint32_t arrivalFrame = 0; -static void frame_retired(Feature*, uint32_t) { ++retired; } -static void logf(const char*, ...) {} -static unsigned env_uint(const char*, unsigned fallback) { return fallback; } -static void d4r_sleep_us(unsigned) { - assert(++sleeps <= 2); - assert(arrival != nullptr); - *arrival = arrivalFrame; -} -static void prepare(Feature* feature, int slotIndex, uint32_t frame, InputSlot& slot) { - const auto start = ProfileClock::now(); -''' + source[start:end] + r''' - ++evaluated; -} -int main() { - // A later submission advances the aggregate marker, but not frame 1's - // slot. Old aggregate-only logic would evaluate unsubmitted inputs here. - uint32_t markers[] = {2, 0, 0, 2}; - Feature feature{markers}; - InputSlot skipped; - prepare(&feature, 1, 1, skipped); - assert(retired == 1 && evaluated == 0 && sleeps == 0 && !skipped.busy); - InputSlot ready; - prepare(&feature, 2, 2, ready); - assert(retired == 1 && evaluated == 1 && sleeps == 0); - // Delayed input stays pending until that slot's exact tag arrives. - markers[0] = 3; - arrival = &markers[1]; - arrivalFrame = 3; - InputSlot delayed; - prepare(&feature, 0, 3, delayed); - assert(evaluated == 2 && sleeps == 1); - feature.retiring = true; - InputSlot cancelled; - prepare(&feature, 1, 4, cancelled); - assert(retired == 2 && evaluated == 2 && !cancelled.busy); -} -''' - with tempfile.TemporaryDirectory() as temporary: - runner = Path(temporary) / 'input-readiness' - subprocess.run(['c++', '-std=c++17', '-Wall', '-Wextra', '-Werror', - '-x', 'c++', '-', '-o', str(runner)], input=runner_source, - text=True, capture_output=True, check=True) - subprocess.run([str(runner)], check=True, timeout=5) - - -if __name__ == '__main__': - unittest.main() diff --git a/tools/d3d12_dlss_harness.cpp b/tools/d3d12_dlss_harness.cpp index 8fa0702..2a925d8 100644 --- a/tools/d3d12_dlss_harness.cpp +++ b/tools/d3d12_dlss_harness.cpp @@ -784,6 +784,20 @@ int main(int argc, char** argv) for (int frame = 1; frame <= frames; ++frame) { + // Recreate the game's output between completed frames to exercise imported-memory handle reuse. + if (frame > 1 && std::getenv("D4R_HARNESS_RECREATE_OUTPUT") != nullptr) + { + outputTexture->Release(); + outputTexture = create_texture(outWidth, outHeight, + rgba8 ? DXGI_FORMAT_R8G8B8A8_UNORM + : rgb10 ? DXGI_FORMAT_R10G10B10A2_UNORM : DXGI_FORMAT_R16G16B16A16_FLOAT, + D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS); + upload(outputTexture, rgba8 ? static_cast(outputInit8.data()) + : rgb10 ? static_cast(outputInit10.data()) + : static_cast(outputInit.data()), + outWidth * (rgba8 || rgb10 ? 4 : 8), D3D12_RESOURCE_STATE_UNORDERED_ACCESS); + d4r_ngx_set_d3d12_resource(parameters, "Output", outputTexture); + } if (qualityScene) { // Halton(2,3) jitter in render pixels, [-0.5, 0.5); a render pixel diff --git a/tools/d4r_nvngx_shim.cpp b/tools/d4r_nvngx_shim.cpp index 979b811..cdd9617 100644 --- a/tools/d4r_nvngx_shim.cpp +++ b/tools/d4r_nvngx_shim.cpp @@ -2487,15 +2487,21 @@ static size_t vk_texel_bytes(VkFormat format) } // D4R_SHIM_INPLACE (experimental): a game texture that vkd3d-proton created linear and exportable is imported -// into CUDA once, so DLSS reads or writes its texels without a game-side copy. Imports are kept for the -// process lifetime (ROCm does not free mapped external memory, see g_vramPool). +// into CUDA once, so DLSS writes its texels without a game-side copy. Imports and their source resources +// stay alive for the process lifetime (ROCm does not free mapped external memory, see g_vramPool). struct InPlaceTexture { CudaDevicePtr texels = 0; size_t rowPitch = 0; }; +struct InPlaceImport +{ + uint64_t memory; + CudaDevicePtr base; + ID3D12Resource* owner; // retained: prevents reuse of the cached VkDeviceMemory handle +}; static std::mutex g_inPlaceMutex; -static std::vector> g_inPlaceImports; // VkDeviceMemory -> CUDA mapping +static std::vector g_inPlaceImports; // Game thread. static bool describe_in_place(ID3D12Resource* resource, InPlaceTexture& texture) @@ -2507,8 +2513,8 @@ static bool describe_in_place(ID3D12Resource* resource, InPlaceTexture& texture) { std::lock_guard lock(g_inPlaceMutex); for (const auto& entry : g_inPlaceImports) - if (entry.first == memory) - base = entry.second; + if (entry.memory == memory) + base = entry.base; } if (base == 0) { @@ -2525,7 +2531,8 @@ static bool describe_in_place(ID3D12Resource* resource, InPlaceTexture& texture) static_cast(size), static_cast(offset), static_cast(pitch)); std::lock_guard lock(g_inPlaceMutex); - g_inPlaceImports.emplace_back(memory, base); + resource->AddRef(); + g_inPlaceImports.push_back({memory, base, resource}); } texture.texels = base + offset; texture.rowPitch = static_cast(pitch); @@ -4141,6 +4148,14 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram timing.workerWait = profile_ms(timing.workerQueued, start); FrameTiming* stages = timing.enabled ? &timing : nullptr; const bool overlapped = async_copies() && copy_streams() != nullptr; + if (timing.markerDeferred) + { + const auto markerStart = ProfileClock::now(); + if (!wait_for_input_marker(feature, slotIndex, frame)) + return; + if (timing.enabled) + timing.markerWait = profile_ms(markerStart, ProfileClock::now()); + } if (params.vram && env_uint("D4R_SHIM_VRAM_VERIFY", 0) == 0) dump_vram_inputs(slot, frame, params, params.hasExposure ? 4 : 3); const bool linearInputs = params.vram && feature->linearInputs; @@ -4239,14 +4254,6 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram } feature->outputRedirected = redirect; } - if (timing.markerDeferred) - { - const auto markerStart = ProfileClock::now(); - if (!wait_for_input_marker(feature, slotIndex, frame)) - return; - if (timing.enabled) - timing.markerWait = profile_ms(markerStart, ProfileClock::now()); - } const int eventStartResult = timing.enabled && feature->profileEventsReady ? g.cu.eventRecord(feature->profileStart, nullptr) : -1; const auto evalStart = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{}; @@ -4370,9 +4377,9 @@ static void prepare_inputs(Feature* feature, int slotIndex, uint32_t frame, Fram { InputSlot& slot = feature->inputs[slotIndex]; const auto start = ProfileClock::now(); - // D4R_SHIM_DEFER_MARKER (default on): with VRAM interop nothing else happens here before the marker lands, so - // the worker thread waits for it itself, after the host-side setup of the evaluation (which no longer queues - // behind the marker: ~0.1 ms off the critical path) and without a thread wake-up between marker and NGX. + // D4R_SHIM_DEFER_MARKER (default on): the worker waits for the game's copies itself, avoiding a thread + // wake-up between readiness and evaluation. The wait must precede every input read, including CUDA + // uploads and diagnostic readbacks; waiting only before NGX would copy stale data into its arrays. static const bool deferMarker = env_uint("D4R_SHIM_DEFER_MARKER", 1) != 0 && env_uint("D4R_MOTION_DILATION", 0) == 0; if (params.vram && deferMarker && env_uint("D4R_SHIM_VRAM_VERIFY", 0) == 0)