diff --git a/README.md b/README.md
index 2c2d131..44d266f 100644
--- a/README.md
+++ b/README.md
@@ -1,8 +1,15 @@
 # d4r (dlss 4 radeon)
 
+**This fork includes a native Windows 11 / RX 9070 XT (gfx1201) port.**
+Start with [Windows build and validation](docs/windows-rdna4-port.md) and
+[OptiScaler installation](docs/windows-game.md). K/M have hardware validation;
+K performance is still being improved. The Linux upstream description follows.
+
 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.
 
-**Supported DLSS models:** DLSS 3 CNN (E), DLSS 4 transformer (K, default), and DLSS 4.5 transformer (M). DLSS 5 is PURPOSELY not supported.
+**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.
 
@@ -14,15 +21,19 @@ submission, and returns output in the same frame without CPU image copies.
 RGBA16/32F, R11G11B10, RGBA/BGRA8, depth, exposure and motion formats have
 independent D3D12 typed SRV/UAV conversion tests. Four-frame K/M runs with packed
 output match the FP16 baseline exactly after resource-format quantisation.
-The build incorporates upstream updates through `f0d1a65`. Patched OptiScaler now
+The build incorporates upstream updates through `aa447b2`. Patched OptiScaler now
 passes standalone K/M native-network tests with exact RGB agreement against
 the direct-d4r baseline; see [the source patch](patches/optiscaler/README.md).
 Silent Hill 2 now renders K through the native Windows backend; M also renders
 a saved gameplay level with GPU output NaN/Inf checks. Startup depth/stencil
 and swap-chain lifetime issues found in the game have regression checks.
 Both K and M now pass 3840x2160 game output scans. K reaches a user-reported
-49-51 FPS in the measured scene after native output stores and command-hook
-lookup caching. Native M FP8 is separately replay/network validated; the
+62 FPS / 63% GPU with optional async interop, compared with the earlier
+49-51 FPS / 53% synchronous run. The fixed async build completes 16941 finite
+4K frames including feature recreation; queued-frame/recreation outputs match
+synchronous controls exactly. These user readings are not a controlled FSR4
+benchmark; the K performance gate remains open. Windows preset L is not yet
+validated. Native M FP8 is separately replay/network validated; the
 default keeps the FP16-equivalent baseline based on measured performance.
 All sixteen CTest gates pass with the final native Windows ZLUDA runtime. See
 [the reversible game package](docs/windows-game.md) and
@@ -67,7 +78,7 @@ d4r builds for RDNA3 and RDNA4. A newly built release compiles native DLSS 4 and
 | RX 7900 GRE / XT / XTX, Radeon PRO W7800 / W7900 | gfx1100 | widened to f16 | RX 7900 XTX: [external video](https://www.youtube.com/watch?v=_GLjJ2Dn5pU) (DLSS 4.5); other cards untested |
 | RX 7600 / 7600 XT, RX 7650 GRE | gfx1102 | widened to f16 | untested |
 | RDNA3 integrated GPUs | gfx1103 | widened to f16 | untested |
-| RX 9070 XT / 9070 / 9070 GRE, Radeon AI PRO R9700 | gfx1201 | native FP8 (`NativeFp8`, default on) | emulator only |
+| RX 9070 XT / 9070 / 9070 GRE, Radeon AI PRO R9700 | gfx1201 | native FP8 (`NativeFp8`, default on) | RX 9070 XT: [external video](https://www.youtube.com/watch?v=lZ0BLYqAtCs), other cards untested |
 | RX 9060 XT / 9060 | gfx1200 | native FP8 (`NativeFp8`, default on) | compile only |
 | RDNA2 and older | | | unsupported |
 
@@ -106,6 +117,8 @@ The release zip works like an OptiScaler release: its contents go into the folde
    - an `d4r` folder with the shim, the CUDA bridge, ZLUDA, the ROCm 7.2.4 runtime, the native kernels, this game's `d4r.ini`, and NVIDIA's `nvngx_dlss.dll` (310.7) and `_nvngx.dll`.
 2. In Steam, select GE-Proton 11 for the game and set these launch options: `PROTON_FORCE_NVAPI=1 DXVK_NVAPI_GPU_ARCH=AD100 %command%`.
 
+When updating, replace `d3d12.dll` and `d3d12core.dll` together with `d4r/nvngx.dll`. The shim requires their matching resource-lifetime extension.
+
 ROCm does not need to be installed: the zip includes its runtime (from AMD's Ubuntu 22.04 packages, which run under Steam's container runtime on any distribution). The zip's NVIDIA files and the kernels built from NVIDIA's code are not covered by this repository's license (see [NOTICE](NOTICE)). The Proton prefix and the system are not changed. [packaging/D4R_README.txt](packaging/D4R_README.txt) is the full guide that ships in the zip; `scripts/package_release.sh` builds the zip (see [docs/building.md](docs/building.md#7-package-a-release)).
 
 ## Requirements (building from source)
diff --git a/config/d4r.ini.default b/config/d4r.ini.default
index bdd17db..e37e406 100644
--- a/config/d4r.ini.default
+++ b/config/d4r.ini.default
@@ -17,14 +17,20 @@ CompatDataDir =
 GameExe =
 
 [DLSS]
-; Which DLSS network ("preset") to run. fps measured on an RX 7700 XT, Townfall 2560x1440, same street
-; walk, for Quality / Balanced / Performance / Ultra Performance (FSR 4: 76 / 85 / 94 / 107):
-;   E (or CNN)     = DLSS 3 CNN, the fastest at Quality..Performance  72 / 81 / 88 / 89
-;   K (or DLSS4)   = DLSS 4 transformer, sharper; best at Ultra Perf  69 / 77 / 84 / 95
-;   M (or DLSS4.5) = DLSS 4.5 transformer, best image, slowest        52 / 60 / 70 / 90
-;   F = the other CNN (same speed as E); A, B, C, D, G, J run the K network in this DLSS version;
-;   L (DLSS 4.5 Ultra Performance model) is not accelerated here and runs very slowly.
-;   auto = whatever the game/OptiScaler asks for
+; DLSS model recommendations:
+;
+; Model | DLSS  | Recommended modes               | Notes
+; ------+-------+---------------------------------+----------------------------------
+; K     | 4     | DLAA / Quality / Balanced       | Default; faster alternative to M/L
+; M     | 4.5   | Performance                     | More demanding on RDNA3
+; L     | 4.5   | Ultra Performance               | Especially 4K; experimental
+; E     | 3 CNN | Quality / Balanced / Performance | Legacy faster fallback
+; F     | 3 CNN | DLAA / Ultra Performance        | Deprecated; prefer K or L
+;
+; Select the upscaling mode separately in game/OptiScaler. These are recommendations;
+; other modes are allowed. L requires the L texture build.
+; 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
 [Latency]
 ; How many frames old the upscaled image the game presents may be.
@@ -42,6 +48,11 @@ FrameAge = 0
 ; directory directly). Missing accuracy kernels use translated PTX, never the fast set. Slower; this
 ; aims for fidelity, but RTX image parity is not proven. (D4R_PREFER_ACCURACY)
 PreferAccuracy = false
+; false: use original translated enc1/enc2 layers for L/M, retaining other native kernels.
+; Reduces the native encoders' motion trails seen in Townfall; best paired with PreferAccuracy.
+; Costs some GPU time. Manifest required for selective fallback; flat sets fall back entirely.
+; (D4R_NATIVE_SWIN_ENCODERS)
+NativeSwinEncoders = true
 ; Native RDNA3 / RDNA4 kernels for the DLSS 4 and 4.5 network layers (replace ZLUDA's translated code).
 ;   fast  = native kernels with relaxed numerics (unless PreferAccuracy is true)
 ;   exact = a separately built directory with conservative FP8/f16 rounding; not an RTX parity guarantee
@@ -60,6 +71,12 @@ IgnoreDenormals = true
 ImplicitMaxBlock = 256
 
 [Interop]
+; Experimental foreground motion coverage for thin geometry. 0 = original vectors;
+; 1/2 = opaque-surface dilation radius in render pixels; both automatically fill sky gaps
+; from the nearest foreground within 8 render pixels. Needs render- or display-resolution, unjittered motion vectors
+; and same-frame linear VRAM inputs; unsupported features retain original vectors.
+; Default off. May affect disocclusion edges. (D4R_MOTION_DILATION)
+MotionVectorDilation = 0
 ; Keep DLSS inputs and output in VRAM (vkd3d-proton Vulkan interop) instead of copying through
 ; system memory. Falls back automatically when the game's resources do not qualify.
 VramInterop = true
@@ -67,19 +84,18 @@ VramInterop = true
 InputSync = true
 ; Faster hand-offs between the game and DLSS (same frame, same result). Need FrameAge = 0 and VRAM
 ; interop, and a runtime built from this checkout (scripts/install_d4r_runtime.sh).
-;   GpuWait: DLSS is queued ahead and starts on the GPU the moment the game's inputs are ready
-;   LinearInputs: DLSS reads the game's inputs in place (no input copies; needs GpuWait)
+; Input readiness is checked with CPU sleeps; recorded-but-discarded lists never queue GPU waits.
+;   LinearInputs: DLSS reads the game's inputs in place once the input marker is ready
 ;   DirectOutput: the result is written straight into the game-side buffer (preset K, and only while
 ;     NGX runs the native output kernel; drops NVIDIA's "DLSS SDK" watermark)
 ;   ElideNgxSync: removes NGX's internal CPU waits between its GPU stages
 ;   EvalSync = false: no CPU wait between DLSS and the output copy
-;   MarkerPollUs: how often (microseconds) the CPU checks for the game's inputs without GpuWait
-GpuWait = true
+;   MarkerPollUs: how often (microseconds) the CPU checks for the game's inputs
 LinearInputs = true
 DirectOutput = true
 ElideNgxSync = true
 EvalSync = false
-MarkerPollUs = 20
+MarkerPollUs = 200
 
 [Paths]
 ; auto = defaults of scripts/d4r_proton_env.sh
@@ -133,4 +149,4 @@ LogDir = auto
 [Env]
 ; Anything else, passed through as environment variables (one per line):
 ;   D4R_SHIM_SPLIT_TIMEOUT_MS = 200
-;   D4R_SHIM_BLOCKING_SYNC = 1  ; yield CPU while waiting for DLSS output
+;   D4R_SHIM_BLOCKING_SYNC = 0  ; A/B only: restore context waits instead of default event queries/sleeps
diff --git a/docs/architecture.md b/docs/architecture.md
index 3a1533d..d1b40ba 100644
--- a/docs/architecture.md
+++ b/docs/architecture.md
@@ -12,7 +12,7 @@ d4r keeps NVIDIA's DLSS library unmodified and gives it what it expects: an NGX
 | 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 |
-| vkd3d-proton patch | `patches/vkd3d-proton` | Lets the shim split a D3D12 command list around the DLSS call |
+| vkd3d-proton patch | `patches/vkd3d-proton` | Splits a D3D12 command list around DLSS and retains external resources through command allocator completion |
 | NVAPI identity | `tools/d4r_nvapi_identity.c` | Reports an Ada (sm_89) GPU so NGX selects the network weights matching the PTX ZLUDA compiles |
 
 ## One frame
@@ -24,23 +24,22 @@ With VRAM interop and split frames (the defaults when the patched vkd3d-proton i
    - A frame marker.
    - A split. vkd3d-proton submits everything the game records after the DLSS call separately, gated on a timeline semaphore reaching this frame's number.
    - At the start of that second part, a copy of the DLSS result into the game's output texture. Linear R11G11B10, RGB10A2, and RGBA8/BGRA8 outputs are converted from RGBA16F on the GPU. sRGB formats continue through host staging to preserve their existing byte interpretation.
-2. **The CUDA worker queues DLSS.** A worker thread owns the CUDA context. For each frame it queues on the GPU, without waiting on the CPU:
-   - a wait for the marker (`GpuWait`);
+2. **The CUDA worker queues DLSS.** The prep thread checks the GPU-written input marker with short CPU sleeps. It drops recorded-but-unsubmitted frames when they time out or their feature retires, before any DLSS work is queued. The CUDA worker then queues:
    - the input copies into CUDA arrays, or none when NGX samples the buffers in place (`LinearInputs`);
    - NGX's evaluation (NGX's own CPU synchronisations are elided with `ElideNgxSync`);
    - the result, written straight into the output buffer by the native output kernel (`DirectOutput`, preset K) or copied there.
 3. **Release.** When the GPU finishes, the worker signals the timeline semaphore and the rest of the game's frame runs, now with this frame's result.
 
-A watchdog releases any frame that never completes (for example a command list that is recorded but never executed), so the game cannot stall on DLSS.
+The watchdog reports delayed output but never signals an unfinished producer as ready. Retiring features remain registered until their CPU/CUDA jobs drain. The lifetime-aware vkd3d extension retains their staging buffers, imported buffers, conversion images and split semaphore through command-allocator reset or destruction after GPU completion. Discarded recordings are covered too; cleanup no longer relies on a 50 ms delay.
 
-Without the vkd3d-proton patch the shim shows the newest finished result instead (one or more frames old). Without VRAM interop, inputs and output are staged through host memory.
+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.
 
 ## 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.
 - **The Ada identity.** NGX picks network weights by GPU architecture. The weights have to match the PTX that ZLUDA compiles, which is sm_89 for DLSS 310.
 - **Native kernels.** Translating NVIDIA's warp-level matrix code is slow (register pressure, lane shuffles, per-instruction mode switches). Earlier translated K layers also produced non-finite outputs that the final kernel masked. Native layers avoid that path; the compiler now preserves FP16 denormal requirements too, which fixes the translated failure in recorded captures. Numerical and image-quality parity with RTX hardware remains unverified.
-- **The GPU-side hand-offs.** With DLSS itself fast, most of the remaining gap to FSR 4 was the GPU idling while CPU threads waited on each other. Moving the waits onto the GPU raised GPU utilisation from about 95% to 99%.
+- **Sleeping readiness checks.** Waiting for input submission before queueing DLSS avoids an uncancellable GPU wait. Explicit output event queries with short sleeps avoid HIP event/context waits that can consume a CPU core. Submillisecond waits use NtDelayExecution because MinGW std::sleep_for can truncate them to zero. Linear inputs and direct output remain available.
 
 ## Configuration
 
@@ -53,7 +52,6 @@ Without the vkd3d-proton patch the shim shows the newest finished result instead
 | `[Kernels] NativeKernels` | `D4R_ZLUDA_NATIVE_DIR` | native kernel directory (`fast` / `exact` / `off`) |
 | `[Kernels] PreferAccuracy` | `D4R_PREFER_ACCURACY` | default off; accuracy variants for every native family and conservative translation/synchronization settings |
 | `[Interop] VramInterop` | `D4R_SHIM_VRAM_INTEROP` | keep inputs and output in VRAM |
-| `[Interop] GpuWait` | `D4R_SHIM_GPU_WAIT` | GPU-side wait for the game's inputs |
 | `[Interop] LinearInputs` | `D4R_SHIM_LINEAR_INPUTS` | NGX samples the input buffers in place |
 | `[Interop] DirectOutput` | `D4R_SHIM_OUTPUT_DIRECT` | the output kernel writes the game-side buffer directly (preset K) |
 | `[Interop] ElideNgxSync` | `D4R_ELIDE_NGX_SYNC` | skip NGX's internal CPU waits |
@@ -90,11 +88,11 @@ The release zip (`scripts/package_release.sh`) is unpacked into the folder that
 
 The launcher writes each run's logs to `~/.cache/d4r-dlss-captures/game-<exe>-<time>/`. `d4r_nvngx.log` holds the shim's log. With `[Debug] Profile = true` it also records one `D4R_PROFILE` line per frame, with the stage timings.
 
-For a CPU-use A/B check, set `D4R_SHIM_BLOCKING_SYNC = 1` under `[Env]` in the game's `d4r.ini`.
-The VRAM path then requests a blocking event wait for its null-stream output. If ZLUDA and HIP
-honour the event flag, the waiting thread yields the CPU; an unavailable or failed event falls back
-to the original context sync. This is opt-in because it changes the wait scope and may add wakeup
-latency. The profile reports `output_sync_wall`, `output_sync_cpu`, and `output_sync_blocking`;
-compare CPU use and frame rate in the same scene.
+Output completion uses event queries with 200 us NtDelayExecution sleeps by default. `D4R_SHIM_BLOCKING_SYNC = 0`
+under `[Env]` restores context synchronization for an A/B check. The profile records
+`output_sync_wall`, `output_sync_cpu`, and `output_sync_method=event_query_sleep` for the new path.
+An unavailable or failed event falls back to context synchronization. The old `GpuWait` setting
+is no longer used: input readiness is checked on the CPU before CUDA work is queued.
+
 The `d2h` field on this path includes the output synchronization wait, even when direct output
 skips the array-to-buffer copy.
diff --git a/docs/building.md b/docs/building.md
index ac43fdb..ed84278 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 `0007` in order |
-| vkd3d-proton | `3dfc6f07` (the base GE-Proton11-3 ships), plus `patches/vkd3d-proton/0001` |
+| vkd3d-proton | `3dfc6f07` (the base GE-Proton11-3 ships), plus `patches/vkd3d-proton/0001` and `0002` |
 | ROCm | 7.2 (HIP runtime, clang, device libraries) |
 | Proton | GE-Proton11-3 (with its OptiScaler integration) |
 | DLSS | `nvngx_dlss.dll` 310.7.0 |
@@ -81,7 +81,7 @@ D4R_ZLUDA_EMIT=/path/to/zluda/target/release/examples/d4r_emit \
 kernels/build.sh all kernels/out/native
 ```
 
-`kernels/out/native` then holds one code object per replaced DLSS kernel; set `NativeKernelDirFast` in d4r.ini (or `D4R_ZLUDA_NATIVE_DIR`) to it. `kernels/build.sh k` or `m` builds only the network layers and needs neither the DLL nor ZLUDA. The texture kernels (`tex`) extract PTX from your DLL into `kernels/extracted/` and compile it offline for `D4R_GPU_ARCH` through `D4R_ZLUDA_EMIT`; that directory and the build output are git-ignored and must not be redistributed. On gfx12, build a second variant with `D4R_NATIVE_FP8=1` and place it in an `<arch>-fp8` folder.
+`kernels/out/native` then holds one code object per replaced DLSS kernel; set `NativeKernelDirFast` in d4r.ini (or `D4R_ZLUDA_NATIVE_DIR`) to it. `kernels/build.sh k` or `m` builds only the network layers and needs neither the DLL nor ZLUDA. `kernels/build.sh l` builds L's shared network layers and unfolded texture variants; it needs the DLL and emitter. The texture kernels (`tex`, including L) extract PTX from your DLL into `kernels/extracted/` and compile it offline for `D4R_GPU_ARCH` through `D4R_ZLUDA_EMIT`; that directory and the build output are git-ignored and must not be redistributed. On gfx12, build a second variant with `D4R_NATIVE_FP8=1` and place it in an `<arch>-fp8` folder.
 
 To add the accuracy option to a developer set, repeat the build with `D4R_PREFER_ACCURACY=1` and output `kernels/out/native/accuracy` (multi-target builds: `kernels/accuracy/<arch>` and `<arch>-fp8`). Use a separate empty directory; successfully built accuracy sets receive `d4r-accuracy.txt`. Enable `[Kernels] PreferAccuracy = true` and restart. The option defaults to false and never substitutes the fast set when accuracy binaries are absent.
 
@@ -132,7 +132,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`–`0007` applied |
-| `D4R_VKD3D_SRC` | isolated vkd3d-proton checkout inside that directory, with its submodules and patch `0001` applied |
+| `D4R_VKD3D_SRC` | isolated vkd3d-proton checkout inside that directory, with its submodules and patches `0001` and `0002` 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 |
diff --git a/docs/native-kernels.md b/docs/native-kernels.md
index 2f2ebb1..b53dc6f 100644
--- a/docs/native-kernels.md
+++ b/docs/native-kernels.md
@@ -35,12 +35,18 @@ The release is built this way, with hashes from DLSS 310.7.0 and 310.9.1. The PT
   - `pwin_wide.h`: the deep layers, which have only 6–77 windows per frame. It uses 4·NG waves per window, split by token tile and channel group, so the GPU stays busy.
 - **Weights:** prep kernels expand the weights into WMMA operand images once.
 
-**DLSS 4.5, preset M (`kernels/m`, `rrlite_*`).**
+**DLSS 4.5, presets L/M (`kernels/m`, `rrlite_*`).**
 - **Network:** the Swin blocks of DLSS 4.5, whose weights are FP8. Their prep kernels expand the weights to f16 WMMA operands, once per weight set.
 - **Template:** `swin_block.h` covers encoders, the tube-shaped enc3 and decoders.
 - **Output encoding:** the FP8 output is encoded two values at a time with packed 16-bit operations (`enc8x2`, exhaustively equal to the scalar encoder), and the 2×2 patch merge reads the rounded f16 values the codes decode to, from a row layout without LDS bank conflicts, instead of decoding the bytes again in every wave.
 - **Weight loads:** enc1 forms each weight tile's address in scalar registers (`SWIN_SCALAR_BLOAD`); the 8-wave layers are faster without it.
 
+**Preset L texture variants.** L uses unfolded `rrlite_enc0_4x4_*`, `rrlite_dec0_4x4`, and `rrlite_post_3_*` kernels. The L build keeps their full neural arithmetic and replaces surface stores with the existing native format-conversion functions. M's folded enc0 tail and dec0 head cannot be used for these variants. All four input flag combinations and all eight post variants are built, along with the shared downsample kernels. `all` and `tex` include these variants; `l` builds the shared Swin layers plus only the L texture set. L remains experimental; these replacements do not establish RTX parity or game performance.
+
+**Motion mitigation.** `[Kernels] NativeSwinEncoders = false` (`D4R_NATIVE_SWIN_ENCODERS=0`) bypasses only `rrlite_enc1_4x4` and `rrlite_enc2_4x4` in the native manifest. L/M then use the original translated PTX for those encoders, retaining native tube, decoder and texture kernels. Keep `PreferAccuracy = true` for the tested conservative path. It reduced the extra building-edge trails in a 64-frame moving Townfall capture at 853×480 → 2560×1440 on gfx1101; it does not recover detail absent from the low-resolution input or establish RTX parity. The default remains `true` for speed. Flat developer sets without manifests fall back entirely when this switch is off; add a manifest to keep selective acceleration.
+
+Initial validation: a gfx1101 fast L build completed two Ultra Performance harness frames (192×108 → 576×324, DLSS 310.7). Its finite, nonzero RGBA16F output matched the original translated L texture kernels byte for byte with the same native Swin layers. A gfx1201 FP8 accuracy L set compiled successfully; RDNA4 hardware and 4K game output have not been tested. Full packaging rebuilds missing L variants when given an older texture bundle.
+
 ## RDNA4
 
 `kernels/common/wmma_layout.h` selects the WMMA operand and accumulator layout. gfx12 uses eight f16 values per lane in each K half; the accumulator rows are `i + 8·half`. The gfx11 path retains its original layout. `D4R_WMMA_LAYOUT=12` on gfx11 is a test shim: it exercises gfx12 indexing but executes gfx11 WMMA, so its replay output can be compared byte for byte with the existing gfx11 build.
@@ -79,6 +85,7 @@ This mode aims to preserve NVIDIA's arithmetic, not to promise bit-identical RTX
 ```sh
 kernels/build.sh k          # DLSS 4 layers   (ROCm clang only)
 kernels/build.sh m          # DLSS 4.5 layers
+kernels/build.sh l          # L shared layers + unfolded textures (needs DLL and D4R_ZLUDA_EMIT)
 kernels/build.sh tex        # texture kernels (needs D4R_DLSS_DLL and D4R_ZLUDA_BUILD)
 kernels/build.sh all DIR    # everything into DIR (default kernels/out/native)
 D4R_PREFER_ACCURACY=1 kernels/build.sh all DIR/accuracy  # conservative versions of every family
@@ -119,3 +126,11 @@ Median GPU time per frame at Quality (1706×960 → 2560×1440, RX 7700 XT), fro
 | **total** | **3.10** | **2.97** | **total** | **9.04** | **8.23** |
 
 Runs of the same build vary by about ±3% per kernel. M additionally skips its output copy now (direct output through the downsample kernel), which is not a kernel and not in the table. Games whose flags differ from Townfall's (motion vectors at render resolution, LDR input, regular depth) gain more: those kernel variants used to run as plain ZLUDA compiles (K 3.6–4.6 ms → 2.8–3.0 ms, M −0.5 to −0.7 ms per frame).
+
+## Experimental thin-feature motion coverage
+
+`[Interop] MotionVectorDilation = 1` or `2` (`D4R_MOTION_DILATION`) extends nearby foreground motion into pixels whose depth otherwise belongs to the background. The value selects the opaque-surface radius in render pixels. Both enabled values also fill sky gaps using the nearest foreground pixel within an eight-render-pixel circle (far depth 0 for reversed Z, 1 for normal Z). This is the wider sky fill from the comparison; it can strengthen trails in some frames and remains opt-in through `MotionVectorDilation`. This is a depth-guided approximation to better thin-feature motion coverage, not engine-side conservative rasterization. It improved the captured Townfall rooftop where the earlier native-encoder fallback did not resolve the artifact.
+
+The d4r-owned GPU kernel preserves the original depth/motion buffers and writes a separate RG16F velocity texture. The worker waits for its completion before NGX evaluates. No per-frame CPU readback is needed. It requires same-frame linear VRAM inputs, full-resolution unjittered motion vectors and zero subrect bases. Unsupported features or unavailable kernels keep the original vectors. The default is `0` (off); larger radii can affect disocclusion and unrelated foreground edges. The method does not recreate geometry missing from every input frame or establish RTX parity.
+
+`scripts/build_wine_nvcuda_bridge.sh` embeds the support kernels when `D4R_ROCM_DIR` points to ROCm clang. All six gfx110x/gfx120x targets compile; physical validation is currently gfx1101 only. A bridge built without the support kernels reports the feature unavailable rather than changing inputs. Optional GPU regressions cover forward/reversed depth, two radii, jitter, padded rows and preservation of input buffers; see `tests/test_motion_kernel_gpu.py`.
diff --git a/docs/windows-rdna4-port.md b/docs/windows-rdna4-port.md
index 364ddc5..138c7f3 100644
--- a/docs/windows-rdna4-port.md
+++ b/docs/windows-rdna4-port.md
@@ -99,7 +99,11 @@ All sixteen CTest gates also pass with `D4R_ASYNC_INTEROP=1` after the release
 fix. The previously archived synchronous package described below remains a
 separate artifact; it does not contain this async scheduling change.
 
-Upstream refreshed and merged on 2026-10-01 through
+Upstream refreshed and merged on 2026-10-02 through `aa447b2` (experimental L,
+Linux resource lifetimes and event waits). Its Linux-only runtime helpers do
+not enter the native Windows build; K/M native sources and the Windows
+backend are unchanged by that update. Windows L remains outside the validated
+manifest and game runner. The prior refresh on 2026-10-01 incorporated
 [`f0d1a65`](https://github.com/countervolts/d4r/commit/f0d1a65e27aff6cbe7eeaa13227c80facb9056ab)
 and artifact fix `65dfc9f`. Windows retains strict f16 rounding and the exact
 M attention baseline. Upstream's optional denormal override now drops only
diff --git a/kernels/build.sh b/kernels/build.sh
index aca13b3..3e279e4 100755
--- a/kernels/build.sh
+++ b/kernels/build.sh
@@ -2,7 +2,7 @@
 # Builds the native replacements for DLSS kernels into one directory that ZLUDA serves them from
 # (D4R_ZLUDA_NATIVE_DIR, or NativeKernelDirFast in d4r.ini). See docs/native-kernels.md.
 #
-# usage: kernels/build.sh [all|k|m|tex] [OUT_DIR]      (default: all, kernels/out/native)
+# usage: kernels/build.sh [all|k|l|m|tex] [OUT_DIR]      (default: all, kernels/out/native)
 #
 #   D4R_ROCM_DIR    ROCm installation with clang and the HIP device libraries (default /opt/rocm)
 #   D4R_GPU_ARCH    target GPU (default gfx1101): an RDNA3 (gfx110x) or RDNA4 (gfx120x) target; the kernels
@@ -37,7 +37,7 @@ esac
 CLANG="$ROCM/lib/llvm/bin/clang++"
 [[ -x "$CLANG" ]] || { echo "clang++ not found in $ROCM/lib/llvm/bin (set D4R_ROCM_DIR)" >&2; exit 2; }
 mkdir -p "$OUT"
-case "$WHAT" in all|k|m|tex) ;; *) echo "unknown kernel family: $WHAT" >&2; exit 2 ;; esac
+case "$WHAT" in all|k|l|m|tex) ;; *) echo "unknown kernel family: $WHAT" >&2; exit 2 ;; esac
 # Never certify a directory containing older fast binaries as an accuracy set.
 if [[ "$ACCURACY" == 1 ]] && compgen -G "$OUT/*.hsaco" >/dev/null &&
     { [[ ! -f "$OUT/d4r-accuracy.txt" ]] || [[ "$(cat "$OUT/d4r-accuracy.txt")" != 1 ]]; }; then
@@ -67,11 +67,11 @@ if [[ "$WHAT" == all || "$WHAT" == k ]]; then
     echo "== DLSS 4 (preset K) transformer layers"
     for src in "$HERE"/k/dltss_pwin_*.hip; do build_hip "$src"; done
 fi
-if [[ "$WHAT" == all || "$WHAT" == m ]]; then
-    echo "== DLSS 4.5 (preset M) Swin layers"
+if [[ "$WHAT" == all || "$WHAT" == m || "$WHAT" == l ]]; then
+    echo "== DLSS 4.5 (presets L/M) shared Swin layers"
     for src in "$HERE"/m/rrlite_*.hip; do build_hip "$src" "$([[ "$FP8" == 1 ]] && echo -DD4R_FP8_WMMA)"; done
 fi
-if [[ "$WHAT" == all || "$WHAT" == tex ]]; then
+if [[ "$WHAT" == all || "$WHAT" == tex || "$WHAT" == l ]]; then
     echo "== texture kernels (NVIDIA PTX with native parts)"
     : "${D4R_DLSS_DLL:?set D4R_DLSS_DLL to nvngx_dlss.dll (310.7)}"
     : "${D4R_ZLUDA_EMIT:?set D4R_ZLUDA_EMIT to d4r_emit from a ZLUDA build with patches/zluda}"
@@ -86,18 +86,30 @@ if [[ "$WHAT" == all || "$WHAT" == tex ]]; then
     # post 0.89 -> 0.77 ms, hiluma output 0.905 -> 0.887 ms (downsample was slower as wave64).]
     # Every flag combination DLSS selects (motion vectors hi/lo, HDR/LDR, depth inverted/regular, ...) gets
     # the same treatment, so games other than the one measured run the native versions too.
-    specs=(rrlite_dec0_4x4_folded:dec0_head)
-    for mv in mvhi mvlo; do
-        for range in hdr ldr; do
-            specs+=("rrlite_enc0_4x4_${mv}_${range}_folded:enc0_tail")
-            for v in 3_1 3_2; do specs+=("rrlite_post_${v}_${mv}_${range}_folded:sust_only:w64"); done
-            for depth in depthinv depthreg; do
-                for kind in "" _max; do
-                    [[ -z "$kind" ]] && specs+=("hiluma_engine_output_${depth}_${mv}_${range}_v1_rel:sust_only:w64")
-                    specs+=("hiluma_engine_output_${depth}_${mv}_${range}${kind}_v2_rel:sust_only:w64")
+    specs=()
+    if [[ "$WHAT" != l ]]; then
+        specs+=(rrlite_dec0_4x4_folded:dec0_head)
+        for mv in mvhi mvlo; do
+            for range in hdr ldr; do
+                specs+=("rrlite_enc0_4x4_${mv}_${range}_folded:enc0_tail")
+                for v in 3_1 3_2; do specs+=("rrlite_post_${v}_${mv}_${range}_folded:sust_only:w64"); done
+                for depth in depthinv depthreg; do
+                    for kind in "" _max; do
+                        [[ -z "$kind" ]] && specs+=("hiluma_engine_output_${depth}_${mv}_${range}_v1_rel:sust_only:w64")
+                        specs+=("hiluma_engine_output_${depth}_${mv}_${range}${kind}_v2_rel:sust_only:w64")
+                    done
                 done
             done
         done
+    fi
+    # L's unfolded variants keep their complete neural arithmetic. The folded M
+    # GEMM heads/tails are incompatible; substitute only the surface stores.
+    specs+=(rrlite_dec0_4x4:sust_only)
+    for mv in mvhi mvlo; do
+        for range in hdr ldr; do
+            specs+=("rrlite_enc0_4x4_${mv}_${range}:sust_only")
+            for v in 3_1 3_2; do specs+=("rrlite_post_${v}_${mv}_${range}:sust_only:w64"); done
+        done
     done
     for mode in static dynamic; do
         for range in hdr ldr; do specs+=("rrlite_downsample_kernel_${mode}_${range}:sust_only"); done
diff --git a/kernels/support/motion_dilate.hip b/kernels/support/motion_dilate.hip
new file mode 100644
index 0000000..5d09d60
--- /dev/null
+++ b/kernels/support/motion_dilate.hip
@@ -0,0 +1,64 @@
+#include "../../tools/d4r_motion_dilation.h"
+#pragma clang fp contract(off)
+
+__attribute__((device)) static int clamp_index(int x, uint32_t limit)
+{
+    return x < 0 ? 0 : x >= (int)limit ? (int)limit - 1 : x;
+}
+
+extern "C" __attribute__((global, amdgpu_flat_work_group_size(128, 128)))
+void d4r_motion_dilate(D4rMotionDilationParams p)
+{
+    const uint32_t id = __builtin_amdgcn_workgroup_id_x() * 128 + __builtin_amdgcn_workitem_id_x();
+    if ((uint64_t)id >= (uint64_t)p.motion_width * p.motion_height)
+        return;
+    const uint32_t x = id % p.motion_width, y = id / p.motion_width;
+    const int cx = clamp_index((int)__builtin_floorf((x + 0.5f) * p.render_width / p.motion_width + p.jitter_x), p.render_width);
+    const int cy = clamp_index((int)__builtin_floorf((y + 0.5f) * p.render_height / p.motion_height + p.jitter_y), p.render_height);
+    const uint8_t* depth = (const uint8_t*)p.depth;
+    const uint8_t* motion = (const uint8_t*)p.motion;
+    uint8_t* output = (uint8_t*)p.output;
+    const float current = *(const float*)(depth + (uint64_t)cy * p.depth_pitch + 4 * cx);
+    float best = current;
+    int bx = cx, by = cy;
+    for (int dy = -(int)p.radius; dy <= (int)p.radius; ++dy)
+        for (int dx = -(int)p.radius; dx <= (int)p.radius; ++dx)
+        {
+            const int nx = clamp_index(cx + dx, p.render_width), ny = clamp_index(cy + dy, p.render_height);
+            const float candidate = *(const float*)(depth + (uint64_t)ny * p.depth_pitch + 4 * nx);
+            if (p.reversed_depth ? candidate > best : candidate < best)
+                best = candidate, bx = nx, by = ny;
+        }
+    // Sky gaps need wider coverage than opaque edges. Select the closest
+    // foreground pixel within eight render pixels, rather than the closest
+    // depth (which can borrow a distant roof/chimney's velocity).
+    const bool sky = p.reversed_depth ? current <= 1.e-7f : current >= 1.0f - 1.e-7f;
+    bool sky_filled = false;
+    if (sky)
+    {
+        int best_distance = 65;
+        for (int dy = -8; dy <= 8; ++dy)
+            for (int dx = -8; dx <= 8; ++dx)
+            {
+                const int nx = cx + dx, ny = cy + dy;
+                const int distance = dx * dx + dy * dy;
+                if (nx < 0 || ny < 0 || nx >= (int)p.render_width || ny >= (int)p.render_height ||
+                    distance > 64 || distance > best_distance)
+                    continue;
+                const float candidate = *(const float*)(depth + (uint64_t)ny * p.depth_pitch + 4 * nx);
+                const bool foreground = p.reversed_depth ? candidate > 1.e-7f : candidate < 1.0f - 1.e-7f;
+                // Stable coordinate tie-break matches nearest-foreground fill.
+                if (foreground && (distance < best_distance || nx < bx || (nx == bx && ny < by)))
+                    best_distance = distance, best = candidate, bx = nx, by = ny, sky_filled = true;
+            }
+    }
+    uint32_t value = *(const uint32_t*)(motion + (uint64_t)y * p.motion_pitch + 4 * x);
+    const bool nearer = p.reversed_depth ? best > current * 1.02f + 1.e-7f : best < current * 0.98f - 1.e-7f;
+    if (sky_filled || nearer)
+    {
+        const int mx = clamp_index((int)__builtin_floorf((bx + 0.5f - p.jitter_x) * p.motion_width / p.render_width), p.motion_width);
+        const int my = clamp_index((int)__builtin_floorf((by + 0.5f - p.jitter_y) * p.motion_height / p.render_height), p.motion_height);
+        value = *(const uint32_t*)(motion + (uint64_t)my * p.motion_pitch + 4 * mx);
+    }
+    *(uint32_t*)(output + (uint64_t)y * p.motion_pitch + 4 * x) = value;
+}
diff --git a/kernels/tex/make_ptx.py b/kernels/tex/make_ptx.py
index abe0f42..02e7704 100755
--- a/kernels/tex/make_ptx.py
+++ b/kernels/tex/make_ptx.py
@@ -211,7 +211,9 @@ def replace_rz_round(lines):
 ENC0_REF = "rrlite_enc0_4x4_mvhi_hdr_folded"
 # the other flag combinations of the kernels whose only native part is the surface-store rewrite
 SUST_ONLY_RE = re.compile(r"^(hiluma_engine_output_depth(inv|reg)_mv(hi|lo)_(hdr|ldr)(_max)?_v[12]_rel|"
-                          r"rrlite_post_3_[12]_mv(hi|lo)_(hdr|ldr)_folded|rrlite_downsample_kernel_(static|dynamic)_(hdr|ldr))$")
+                          r"rrlite_post_3_[12]_mv(hi|lo)_(hdr|ldr)(_folded)?|"
+                          r"rrlite_enc0_4x4_mv(hi|lo)_(hdr|ldr)|rrlite_dec0_4x4|"
+                          r"rrlite_downsample_kernel_(static|dynamic)_(hdr|ldr))$")
 ENC0_RE = re.compile(r"^rrlite_enc0_4x4_mv(hi|lo)_(hdr|ldr)_folded$")
 
 
diff --git a/packaging/D4R_README.txt b/packaging/D4R_README.txt
index da65cad..26c52cb 100644
--- a/packaging/D4R_README.txt
+++ b/packaging/D4R_README.txt
@@ -37,6 +37,7 @@ Install
 1. Extract everything in this zip into the folder that holds the game's main .exe. For Unreal
    Engine games that is <game>/<Project>/Binaries/Win64/, next to <Project>-Win64-Shipping.exe.
    You get dxgi.dll, OptiScaler.ini, d3d12.dll, d3d12core.dll, this file and an d4r folder.
+   When updating, replace both d3d12 DLLs with the shim: it requires their resource-lifetime extension.
    If the folder already has a dxgi.dll or OptiScaler.ini (another OptiScaler install), move those
    out of the way first. Also remove PROTON_USE_OPTISCALER from the game's launch options if you
    used GE-Proton's built-in OptiScaler: d4r brings its own.
@@ -65,10 +66,16 @@ Settings
   start may compile a separate cache. Missing accuracy variants use translated kernels, never the
   fast native set. The mode aims to match NVIDIA's output, but 1:1 RTX image quality is not proven.
 - d4r/d4r.ini: d4r's settings for this game. The main one is the DLSS model:
-  K (DLSS 4, default), E (DLSS 3 CNN, a little faster) or M (DLSS 4.5, slowest).
+  K (DLSS 4, default), E (DLSS 3 CNN, a little faster), M (DLSS 4.5), or L (DLSS 4.5
+  for Ultra Performance, especially 4K; experimental). Set Model = L and select Ultra Performance
+  in game or OptiScaler. Choosing the model does not change the render resolution.
 - On RDNA4, [Kernels] NativeFp8 is on by default. It selects native FP8 WMMA and the matching
   gfx12-fp8 kernel folder; set it to false to use f16 widening. The FP8 path has not been tested
   on a real RDNA4 GPU.
+- If L/M shows extra trails in motion, try [Kernels] NativeSwinEncoders = false with
+  PreferAccuracy = true. This uses original translated enc1/enc2 layers while retaining other
+  native acceleration. It reduced trails in a captured Townfall sequence on RX 7700 XT and costs
+  some GPU time. Restart after changing it; other scenes and RDNA4 hardware still need testing.
 - OptiScaler.ini: OptiScaler's settings, such as its fps overlay and the render resolution of each
   quality mode. Press Insert in game for OptiScaler's menu.
 
diff --git a/packaging/SOURCES.txt b/packaging/SOURCES.txt
index 70a5e0d..8caff54 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 applied,
+                         commit @VKD3D_COMMIT@ with d4r/source/patches/vkd3d-proton/0001 and 0002 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/packaging/d4r.ini b/packaging/d4r.ini
index 1f27567..e6d5b7c 100644
--- a/packaging/d4r.ini
+++ b/packaging/d4r.ini
@@ -3,12 +3,19 @@
 ; that config/d4r.ini.default in the source repository lists for it, e.g. in the game's launch options.
 
 [DLSS]
-; Which DLSS network runs:
-;   K     DLSS 4 transformer (default)
-;   E     DLSS 3 CNN: a little faster, softer image
-;   M     DLSS 4.5 transformer: best image, much slower on RDNA3
-;   game  whatever the game (or OptiScaler's preset override) asks for
-; Frame rates measured on an RX 7700 XT at 2560x1440, Quality mode: see D4R_README.txt.
+; DLSS model recommendations:
+;
+; Model | DLSS  | Recommended modes               | Notes
+; ------+-------+---------------------------------+----------------------------------
+; K     | 4     | DLAA / Quality / Balanced       | Default; faster alternative to M/L
+; M     | 4.5   | Performance                     | More demanding on RDNA3
+; L     | 4.5   | Ultra Performance               | Especially 4K; experimental
+; E     | 3 CNN | Quality / Balanced / Performance | Legacy faster fallback
+; F     | 3 CNN | DLAA / Ultra Performance        | Deprecated; prefer K or L
+;
+; Select the upscaling mode separately in game/OptiScaler. These are recommendations;
+; other modes are allowed. L requires the L texture build.
+; game = use the game/OptiScaler choice. Benchmarks: D4R_README.txt.
 Model = K
 ; NVIDIA's library marks the image "DLSS SDK - DO NOT DISTRIBUTE" in this setup; d4r turns that off.
 ; true shows it.
@@ -25,6 +32,9 @@ FrameAge = 0
 ; intermediate FP8/f16 rounding and denormal handling, and keeps NGX's synchronization calls.
 ; Slower but aims to match RTX DLSS.
 PreferAccuracy = false
+; false reduces L/M motion trails by using original enc1/enc2 layers; other native kernels stay on.
+; Best paired with PreferAccuracy = true; costs some GPU time. Restart after changing this.
+NativeSwinEncoders = true
 ; Hand-written RDNA3 and RDNA4 kernels for the DLSS 4 and 4.5 networks (d4r\kernels). Each one is used
 ; only when your DLSS library's code for that kernel matches the code it was written for; otherwise
 ; that kernel is compiled from NVIDIA's code instead, which is slower but still correct.
@@ -35,10 +45,13 @@ NativeKernels = on
 NativeFp8 = on
 
 [Interop]
-; These keep DLSS's inputs and output in video memory and move the waits between the game and DLSS
-; onto the GPU. Leave them on unless you are chasing a problem.
+; Experimental thin-feature tracking: 0 = off; 1/2 = opaque-surface dilation radius.
+; Both enabled values fill sky gaps from the nearest foreground within 8 render pixels.
+; Requires render- or display-resolution, unjittered motion vectors and same-frame linear VRAM inputs. Restart to change.
+; Unsupported features keep original vectors; larger radii may affect disocclusion edges.
+MotionVectorDilation = 0
+; Keep inputs and output in video memory. Input and output readiness waits use CPU sleeps.
 VramInterop = true
-GpuWait = true
 LinearInputs = true
 ElideNgxSync = true
 ; Lets DLSS's last kernel write straight into the game's output buffer (DLSS 4 / K only). It takes
@@ -59,5 +72,5 @@ Profile = false
 
 [Env]
 ; Extra environment variables for d4r, NAME = value.
-; Use a blocking CUDA event while waiting for DLSS output.
-; D4R_SHIM_BLOCKING_SYNC = 1
+; Sleeping event queries are the default; 0 restores context waits for an A/B comparison.
+; D4R_SHIM_BLOCKING_SYNC = 0
diff --git a/patches/vkd3d-proton/0002-retain-external-resources.patch b/patches/vkd3d-proton/0002-retain-external-resources.patch
new file mode 100644
index 0000000..7c11cb2
--- /dev/null
+++ b/patches/vkd3d-proton/0002-retain-external-resources.patch
@@ -0,0 +1,148 @@
+diff --git a/include/vkd3d_device_vkd3d_ext.idl b/include/vkd3d_device_vkd3d_ext.idl
+--- a/include/vkd3d_device_vkd3d_ext.idl
++++ b/include/vkd3d_device_vkd3d_ext.idl
+@@ -120,6 +120,19 @@
+ }
+ 
+ [
++    uuid(4f8c9462-8197-4dc1-a584-cf44d9706db2),
++    object,
++    local,
++    pointer_default(unique)
++]
++interface ID3D12DXVKInteropDeviceD4R2 : ID3D12DXVKInteropDeviceD4R
++{
++    /* Retain external resources until the allocator can safely reset or is
++     * destroyed. Pending submissions already retain the allocator. */
++    HRESULT RetainExternalResources(ID3D12CommandList *list, IUnknown *resources);
++}
++
++[
+     uuid(f3112584-41f9-348d-a59b-00b7e1d285d6),
+     object,
+     local,
+diff --git a/libs/vkd3d/command.c b/libs/vkd3d/command.c
+--- a/libs/vkd3d/command.c
++++ b/libs/vkd3d/command.c
+@@ -2791,6 +2791,24 @@
+     }
+ }
+ 
++HRESULT d3d12_command_list_retain_external_resources(struct d3d12_command_list *list, IUnknown *resources)
++{
++    struct d3d12_command_allocator *allocator = list->allocator;
++    size_t i;
++
++    if (!list->is_recording || !allocator)
++        return E_INVALIDARG;
++    for (i = 0; i < allocator->external_resources_count; i++)
++        if (allocator->external_resources[i] == resources)
++            return S_OK;
++    if (!vkd3d_array_reserve((void **)&allocator->external_resources, &allocator->external_resources_size,
++            allocator->external_resources_count + 1, sizeof(*allocator->external_resources)))
++        return E_OUTOFMEMORY;
++    IUnknown_AddRef(resources);
++    allocator->external_resources[allocator->external_resources_count++] = resources;
++    return S_OK;
++}
++
+ static bool d3d12_command_allocator_add_view(struct d3d12_command_allocator *allocator,
+         struct vkd3d_view *view)
+ {
+@@ -2858,6 +2876,10 @@
+     struct d3d12_device *device = allocator->device;
+     const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
+     unsigned int i;
++
++    for (i = 0; i < allocator->external_resources_count; ++i)
++        IUnknown_Release(allocator->external_resources[i]);
++    allocator->external_resources_count = 0;
+ 
+     for (i = 0; i < allocator->buffer_view_count; ++i)
+     {
+@@ -2993,6 +3015,7 @@
+             d3d12_command_list_allocator_destroyed(allocator->current_command_list);
+ 
+         d3d12_command_allocator_free_resources(allocator);
++        vkd3d_free(allocator->external_resources);
+         vkd3d_free(allocator->buffer_views);
+         vkd3d_free(allocator->views);
+         vkd3d_free(allocator->pipelines);
+diff --git a/libs/vkd3d/device.c b/libs/vkd3d/device.c
+--- a/libs/vkd3d/device.c
++++ b/libs/vkd3d/device.c
+@@ -4658,7 +4658,8 @@
+             || IsEqualGUID(riid, &IID_ID3D12DXVKInteropDevice1)
+             || IsEqualGUID(riid, &IID_ID3D12DXVKInteropDevice2)
+             || IsEqualGUID(riid, &IID_ID3D12DXVKInteropDevice3)
+-            || IsEqualGUID(riid, &IID_ID3D12DXVKInteropDeviceD4R))
++            || IsEqualGUID(riid, &IID_ID3D12DXVKInteropDeviceD4R)
++            || IsEqualGUID(riid, &IID_ID3D12DXVKInteropDeviceD4R2))
+     {
+         d3d12_dxvk_interop_device_AddRef(&device->ID3D12DXVKInteropDevice_iface);
+         *object = &device->ID3D12DXVKInteropDevice_iface;
+@@ -11380,7 +11381,7 @@
+ }
+ 
+ extern CONST_VTBL struct ID3D12DeviceExt5Vtbl d3d12_device_vkd3d_ext_vtbl;
+-extern CONST_VTBL struct ID3D12DXVKInteropDeviceD4RVtbl d3d12_dxvk_interop_device_vtbl;
++extern CONST_VTBL struct ID3D12DXVKInteropDeviceD4R2Vtbl d3d12_dxvk_interop_device_vtbl;
+ extern CONST_VTBL struct ID3DLowLatencyDeviceVtbl d3d_low_latency_device_vtbl;
+ extern CONST_VTBL struct IAmdExtAntiLagApiVtbl d3d_amd_ext_anti_lag_vtbl;
+ 
+diff --git a/libs/vkd3d/device_vkd3d_ext.c b/libs/vkd3d/device_vkd3d_ext.c
+--- a/libs/vkd3d/device_vkd3d_ext.c
++++ b/libs/vkd3d/device_vkd3d_ext.c
+@@ -1157,7 +1157,16 @@
+     return d3d12_command_list_split_for_external_wait(cmd_list, (VkSemaphore)vk_semaphore, value);
+ }
+ 
+-CONST_VTBL struct ID3D12DXVKInteropDeviceD4RVtbl d3d12_dxvk_interop_device_vtbl =
++static HRESULT STDMETHODCALLTYPE d3d12_dxvk_interop_device_RetainExternalResources(
++        d3d12_dxvk_interop_device_iface *iface, ID3D12CommandList *list, IUnknown *resources)
++{
++    struct d3d12_command_list *cmd_list = d3d12_command_list_from_iface(list);
++    if (!cmd_list || !resources)
++        return E_INVALIDARG;
++    return d3d12_command_list_retain_external_resources(cmd_list, resources);
++}
++
++CONST_VTBL struct ID3D12DXVKInteropDeviceD4R2Vtbl d3d12_dxvk_interop_device_vtbl =
+ {
+     /* IUnknown methods */
+     d3d12_dxvk_interop_device_QueryInterface,
+@@ -1192,6 +1201,9 @@
+ 
+     /* ID3D12DXVKInteropDeviceD4R methods */
+     d3d12_dxvk_interop_device_SplitCommandListForExternalWait,
++
++    /* ID3D12DXVKInteropDeviceD4R2 methods */
++    d3d12_dxvk_interop_device_RetainExternalResources,
+ };
+ 
+ static inline struct d3d12_device *d3d12_device_from_ID3DLowLatencyDevice(d3d_low_latency_device_iface *iface)
+diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h
+--- a/libs/vkd3d/vkd3d_private.h
++++ b/libs/vkd3d/vkd3d_private.h
+@@ -2758,6 +2758,10 @@
+     LONG refcount;
+     LONG internal_refcount;
+ 
++    IUnknown **external_resources;
++    size_t external_resources_size;
++    size_t external_resources_count;
++
+     D3D12_COMMAND_LIST_TYPE type;
+ 
+     struct vkd3d_view **views;
+@@ -5474,7 +5478,9 @@
+ HRESULT d3d12_command_list_split_for_external_wait(struct d3d12_command_list *list,
+         VkSemaphore semaphore, uint64_t value);
+ 
+-typedef ID3D12DXVKInteropDeviceD4R d3d12_dxvk_interop_device_iface;
++HRESULT d3d12_command_list_retain_external_resources(struct d3d12_command_list *list, IUnknown *resources);
++
++typedef ID3D12DXVKInteropDeviceD4R2 d3d12_dxvk_interop_device_iface;
+ 
+ /* ID3DLowLatencyDevice */
+ typedef ID3DLowLatencyDevice d3d_low_latency_device_iface;
diff --git a/scripts/build_d3d12_dlss_harness.sh b/scripts/build_d3d12_dlss_harness.sh
index 513a579..a6ef3d7 100755
--- a/scripts/build_d3d12_dlss_harness.sh
+++ b/scripts/build_d3d12_dlss_harness.sh
@@ -5,7 +5,7 @@ ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)"
 MINGW_CXX="${MINGW_CXX:-x86_64-w64-mingw32-g++}"
 
 "$ROOT/scripts/build_d4r_nvngx_shim.sh"
-"$MINGW_CXX" -std=c++20 -O2 -Wall -Wextra -static -static-libgcc -static-libstdc++ \
+"$MINGW_CXX" -std=c++20 -O2 -Wall -Wextra -I"$ROOT/build/vulkan-include" -static -static-libgcc -static-libstdc++ \
   "$ROOT/tools/d3d12_dlss_harness.cpp" "$ROOT/tools/d4r_ngx_param_host.cpp" "$ROOT/build/d4r_ngx_param_msvc.obj" \
   -o "$ROOT/build/d3d12_dlss_harness.exe"
 printf 'Built %s\n' "$ROOT/build/d3d12_dlss_harness.exe"
diff --git a/scripts/build_d4r_motion_kernels.sh b/scripts/build_d4r_motion_kernels.sh
new file mode 100755
index 0000000..4d71821
--- /dev/null
+++ b/scripts/build_d4r_motion_kernels.sh
@@ -0,0 +1,30 @@
+#!/usr/bin/env bash
+# Build d4r-owned motion processing kernels and embed them in the Wine bridge.
+set -euo pipefail
+ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)"
+OUT="${1:-$ROOT/build/wine-nvcuda/generated}"
+ROCM="${D4R_ROCM_DIR:-/opt/rocm}"
+CLANG="$ROCM/lib/llvm/bin/clang++"
+[[ -x "$CLANG" ]] || { echo "motion kernels need ROCm clang (D4R_ROCM_DIR)" >&2; exit 2; }
+mkdir -p "$OUT"
+for arch in gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201; do
+    "$CLANG" -x hip -nogpuinc -nogpulib --offload-device-only --offload-arch="$arch" \
+        --no-gpu-bundle-output -O3 -mno-wavefrontsize64 -o "$OUT/$arch.hsaco" \
+        "$ROOT/kernels/support/motion_dilate.hip"
+done
+python3 - "$OUT" <<'PY'
+from pathlib import Path
+import sys
+root = Path(sys.argv[1])
+lines = ['/* Generated from d4r-owned HIP source; do not edit. */', '#include <string.h>']
+for p in sorted(root.glob('gfx*.hsaco')):
+    data = p.read_bytes()
+    lines.append(f'static const unsigned char motion_{p.stem}[] __attribute__((aligned(64))) = {{')
+    lines.extend(','.join(f'0x{x:02x}' for x in data[i:i+24]) + ',' for i in range(0,len(data),24))
+    lines.append('};')
+lines.append('static const void* d4r_motion_image(const char* arch) {')
+for p in sorted(root.glob('gfx*.hsaco')):
+    lines.append(f'  if (strcmp(arch, "{p.stem}") == 0) return motion_{p.stem};')
+lines.append('  return NULL;\n}')
+(root/'d4r_motion_kernels.h').write_text('\n'.join(lines)+'\n')
+PY
diff --git a/scripts/build_release_glibc241.sh b/scripts/build_release_glibc241.sh
index 46ed9f4..69be734 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-0007 applied}"
-: "${D4R_VKD3D_SRC:?set a build-local vkd3d-proton checkout with patch 0001 applied}"
+: "${D4R_VKD3D_SRC:?set a build-local vkd3d-proton checkout with patches 0001 and 0002 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/scripts/build_wine_nvcuda_bridge.sh b/scripts/build_wine_nvcuda_bridge.sh
index 1bcf498..ebc9bdc 100755
--- a/scripts/build_wine_nvcuda_bridge.sh
+++ b/scripts/build_wine_nvcuda_bridge.sh
@@ -7,7 +7,13 @@ WINEBUILD="${WINEBUILD:-winebuild}"
 BRIDGE_ROOT="$ROOT/build/wine-nvcuda"
 
 mkdir -p "$BRIDGE_ROOT/x86_64-unix" "$BRIDGE_ROOT/x86_64-windows"
+MOTION_FLAGS=()
+if [[ -x "${D4R_ROCM_DIR:-/opt/rocm}/lib/llvm/bin/clang++" ]]; then
+  "$ROOT/scripts/build_d4r_motion_kernels.sh" "$BRIDGE_ROOT/generated"
+  MOTION_FLAGS=(-DD4R_MOTION_KERNELS -I"$BRIDGE_ROOT/generated")
+fi
 "$WINEGCC" -m64 -shared -o "$BRIDGE_ROOT/x86_64-unix/nvcuda.dll" \
+  "${MOTION_FLAGS[@]}" \
   "$ROOT/tools/wine_nvcuda_bridge.c" \
   -Wb,--export="$ROOT/tools/wine_nvcuda_bridge.spec" -ldl -lpthread
 "$WINEBUILD" --dll --fake-module -E "$ROOT/tools/wine_nvcuda_bridge.spec" \
diff --git a/scripts/d4r_config.py b/scripts/d4r_config.py
index 8115eee..2ac1710 100755
--- a/scripts/d4r_config.py
+++ b/scripts/d4r_config.py
@@ -133,18 +133,24 @@ def main(argv):
     # RDNA4's native FP8 WMMA (on unless set off; ZLUDA and the bridge ignore it on other GPUs)
     value = flag("Kernels", "NativeFp8")
     env["D4R_ZLUDA_WMMA_FP8_NATIVE"] = "0" if value is False else "1"
+    value = flag("Kernels", "NativeSwinEncoders")
+    env["D4R_NATIVE_SWIN_ENCODERS"] = "0" if value is False else "1"
     value = get("Kernels", "ImplicitMaxBlock")
     if value:
         env["D4R_ZLUDA_IMPLICIT_MAX_BLOCK"] = value
 
     for key, var in (("VramInterop", "D4R_SHIM_VRAM_INTEROP"), ("InputSync", "D4R_SHIM_INPUT_SYNC"),
-                     ("GpuWait", "D4R_SHIM_GPU_WAIT"), ("LinearInputs", "D4R_SHIM_LINEAR_INPUTS"),
+                     ("LinearInputs", "D4R_SHIM_LINEAR_INPUTS"),
                      ("DirectOutput", "D4R_SHIM_OUTPUT_DIRECT"), ("ElideNgxSync", "D4R_ELIDE_NGX_SYNC"),
                      ("EvalSync", "D4R_SHIM_EVAL_SYNC")):
         value = flag("Interop", key)
         if value is not None:
             env[var] = "1" if value else "0"
     value = get("Interop", "MarkerPollUs")
+    dilation = get("Interop", "MotionVectorDilation")
+    if dilation is not None and dilation not in ("0", "1", "2"):
+        fail("[Interop] MotionVectorDilation must be 0, 1 or 2")
+    env["D4R_MOTION_DILATION"] = dilation or "0"
     if value:
         if not value.isdigit():
             fail(f"[Interop] MarkerPollUs must be a number of microseconds, not {value!r}")
diff --git a/scripts/package_release.sh b/scripts/package_release.sh
index d04621d..c467f69 100755
--- a/scripts/package_release.sh
+++ b/scripts/package_release.sh
@@ -30,7 +30,7 @@
 #                    older builds
 #                    Accuracy texture sets go in accuracy/<target>/, built with D4R_PREFER_ACCURACY=1.
 #   D4R_ZLUDA_EMIT  d4r_emit from the patched ZLUDA build; required in full builds when an accuracy
-#                    texture set is not supplied in D4R_BUNDLE_TEX/accuracy/<target>.
+#                    texture set or L's unfolded texture variants are not supplied in D4R_BUNDLE_TEX.
 # D4R_BUNDLE_NVIDIA=0 leaves them out (d4r-VERSION-nonvidia.zip; users then add the two DLLs themselves).
 set -euo pipefail
 
@@ -126,6 +126,24 @@ for arch in $ARCHS; do
   folders+=("$arch")
   [[ "$arch" == gfx12* ]] && folders+=("$arch-fp8")
 done
+# Older texture bundles predate L. Complete them from the user's DLSS DLL rather
+# than silently shipping only M's folded variants under the new L support claim.
+ensure_l_textures() {
+  local dir="$1" accuracy="$2" missing=0 mv range v name
+  local names=(rrlite_dec0_4x4)
+  for mv in mvhi mvlo; do
+    for range in hdr ldr; do
+      names+=("rrlite_enc0_4x4_${mv}_${range}")
+      for v in 3_1 3_2; do names+=("rrlite_post_${v}_${mv}_${range}"); done
+    done
+  done
+  for name in "${names[@]}"; do [[ -f "$dir/$name.hsaco" ]] || missing=1; done
+  if [[ "$missing" == 1 ]]; then
+    : "${D4R_ZLUDA_EMIT:?set D4R_ZLUDA_EMIT or supply the unfolded L texture variants in D4R_BUNDLE_TEX}"
+    D4R_PREFER_ACCURACY="$accuracy" D4R_ROCM_DIR="$ROCM" D4R_GPU_ARCH="$arch" D4R_NATIVE_FP8="$fp8" \
+      D4R_DLSS_DLL="$D4R_BUNDLE_DLSS" "$ROOT/kernels/build.sh" l "$dir" >/dev/null
+  fi
+}
 build_target() {
   local folder="$1" arch fp8 tex_dir accurate accurate_tex f
   arch="${folder%-fp8}"
@@ -148,6 +166,7 @@ build_target() {
       done
     fi
   fi
+  [[ "$VARIANT" != full ]] || ensure_l_textures "$STAGE/d4r/kernels/$folder" 0
   python3 "$ROOT/kernels/tools/kernel_manifest.py" "$STAGE/d4r/kernels/$folder" "${DLLS[@]}"
 
   # Every release target also gets conservative network and texture variants. Do not copy fast
@@ -169,6 +188,7 @@ build_target() {
     fi
   fi
   if [[ "$VARIANT" == full ]]; then
+    ensure_l_textures "$accurate" 1
     for f in "$STAGE/d4r/kernels/$folder"/*.hsaco; do
       [[ -f "$accurate/$(basename "$f")" ]] || {
         echo "missing accuracy variant of $(basename "$f") for $folder" >&2; exit 2;
diff --git a/tests/test_accuracy.py b/tests/test_accuracy.py
index d15f3f0..e96e8b1 100644
--- a/tests/test_accuracy.py
+++ b/tests/test_accuracy.py
@@ -31,6 +31,9 @@ int main(int argc, char** argv) {
         puts("");
         return 0;
     }
+    if (argc == 3 && strcmp(argv[1], "kernel") == 0) {
+        printf("%d\n", d4r_native_kernel_allowed(argv[2])); return 0;
+    }
     if (argc != 5) return 2;
     char path[1024] = {0};
     int kind = d4r_select_native_source(argv[1], argv[2], atoi(argv[3]), atoi(argv[4]), path, sizeof(path));
@@ -140,6 +143,37 @@ int main(int argc, char** argv) {
         result = self.config("[Kernels]\nPreferAccuracy=perhaps\n", expected=2)
         self.assertIn("PreferAccuracy must be true or false", result.stderr)
 
+    def test_swin_encoder_fallback_keeps_other_layers(self):
+        encoders = ("rrlite_enc1_4x4", "rrlite_enc2_4x4")
+        others = ("rrlite_enc3_tube_4x4", "rrlite_dec1_4x4", "rrlite_dec2_4x4",
+                  "rrlite_enc0_4x4_mvhi_hdr", "rrlite_post_3_1_mvhi_hdr", "dltss_pwin_enc1_layer")
+        for value in ("0", "1", None):
+            env = {k: v for k, v in os.environ.items() if k != "D4R_NATIVE_SWIN_ENCODERS"}
+            if value is not None:
+                env["D4R_NATIVE_SWIN_ENCODERS"] = value
+            for name in encoders + others:
+                got = subprocess.check_output([str(self.runner), "kernel", name], env=env, text=True).strip()
+                self.assertEqual(got, "0" if value == "0" and name in encoders else "1")
+
+    def test_swin_encoder_config_and_environment_override(self):
+        for value, expected in (("false", "0"), ("true", "1"), ("auto", "1")):
+            self.assertIn(f"export D4R_NATIVE_SWIN_ENCODERS={expected}",
+                          self.config(f"[Kernels]\nNativeSwinEncoders = {value}\n").stdout)
+        result = self.config("[Kernels]\nNativeSwinEncoders = false\n",
+                             {"D4R_NATIVE_SWIN_ENCODERS": "1"})
+        self.assertNotIn("export D4R_NATIVE_SWIN_ENCODERS", result.stdout)
+        self.config("[Kernels]\nNativeSwinEncoders = maybe\n", expected=2)
+
+    def test_motion_dilation_config_bounds_and_environment_precedence(self):
+        for value, expected in (("0", "0"), ("1", "1"), ("2", "2"), ("auto", "0")):
+            self.assertIn(f"export D4R_MOTION_DILATION={expected}",
+                          self.config(f"[Interop]\nMotionVectorDilation = {value}\n").stdout)
+        for value in ("-1", "3", "true", "1.5"):
+            self.config(f"[Interop]\nMotionVectorDilation = {value}\n", expected=2)
+        self.assertNotIn("export D4R_MOTION_DILATION",
+                         self.config("[Interop]\nMotionVectorDilation = 2\n",
+                                     {"D4R_MOTION_DILATION": "0"}).stdout)
+
     def test_shipped_configs_default_off(self):
         for file in ("config/d4r.ini.default", "packaging/d4r.ini"):
             self.assertRegex((ROOT / file).read_text(), r"(?m)^PreferAccuracy\s*=\s*false$")
diff --git a/tests/test_event_wait.py b/tests/test_event_wait.py
new file mode 100644
index 0000000..00e93ee
--- /dev/null
+++ b/tests/test_event_wait.py
@@ -0,0 +1,36 @@
+"""Completion/error contracts for the sleeping CUDA event wait (no GPU needed)."""
+from pathlib import Path
+import subprocess
+import tempfile
+import unittest
+
+ROOT = Path(__file__).resolve().parents[1]
+
+
+class EventWaitTests(unittest.TestCase):
+    def test_completion_pending_and_error(self):
+        source = r'''
+#include "d4r_event_wait.h"
+#include <cassert>
+int main() {
+    int sleeps = 0, queries = 0;
+    assert(d4r_wait_event([] { return 0; }, [&] { ++sleeps; }) == 0);
+    assert(sleeps == 0);
+    assert(d4r_wait_event([&] { return ++queries <= 7 ? 600 : 0; },
+                         [&] { ++sleeps; }) == 0);
+    assert(queries == 8 && sleeps == 7);
+    sleeps = 0;
+    assert(d4r_wait_event([] { return 719; }, [&] { ++sleeps; }) == 719);
+    assert(sleeps == 0);
+    queries = 0;
+    assert(d4r_wait_event([&] { return ++queries == 1 ? 600 : 700; },
+                         [&] { ++sleeps; }) == 700);
+    assert(sleeps == 1);
+}
+'''
+        with tempfile.TemporaryDirectory() as tmp:
+            runner = Path(tmp) / "wait"
+            subprocess.run(["c++", "-x", "c++", "-std=c++17", "-Wall", "-Wextra", "-Werror",
+                            "-I", str(ROOT / "tools"), "-o", str(runner), "-"],
+                           input=source, text=True, check=True, capture_output=True)
+            subprocess.run([str(runner)], check=True, timeout=5)
diff --git a/tests/test_model_l.py b/tests/test_model_l.py
new file mode 100644
index 0000000..4733579
--- /dev/null
+++ b/tests/test_model_l.py
@@ -0,0 +1,48 @@
+"""L must retain its unfolded neural math when native surface stores are substituted."""
+import os
+from pathlib import Path
+import subprocess
+import tempfile
+import unittest
+
+ROOT = Path(__file__).resolve().parents[1]
+
+
+class ModelLTests(unittest.TestCase):
+    def test_model_l_config_and_environment_precedence(self):
+        with tempfile.TemporaryDirectory() as directory:
+            ini = Path(directory) / 'd4r.ini'
+            ini.write_text('[DLSS]\nModel = L\n')
+            env = dict(os.environ)
+            env.pop('D4R_DLSS_PRESET', None)
+            cmd = ['python3', str(ROOT / 'scripts/d4r_config.py'), '--config', str(ini)]
+            self.assertIn('export D4R_DLSS_PRESET=12', subprocess.check_output(cmd, env=env, text=True))
+            env['D4R_DLSS_PRESET'] = '11'
+            self.assertNotIn('export D4R_DLSS_PRESET=', subprocess.check_output(cmd, env=env, text=True))
+
+    def test_unfolded_variants_preserve_neural_arithmetic(self):
+        names = ['rrlite_dec0_4x4']
+        for mv in ('mvhi', 'mvlo'):
+            for dynamic_range in ('hdr', 'ldr'):
+                names.append(f'rrlite_enc0_4x4_{mv}_{dynamic_range}')
+                names.extend(f'rrlite_post_{v}_{mv}_{dynamic_range}' for v in ('3_1', '3_2'))
+        math = 'mma.sync.aligned.m16n8k32.row.col.f16.e4m3.e4m3.f16 {%r1, %r2}, {%r3, %r4, %r5, %r6}, {%r7, %r8}, {%r9, %r10};'
+        store = 'sust.b.2d.v2.b16.zero [%rd1, {%r1,%r2}], {%rs1,%rs2};'
+        with tempfile.TemporaryDirectory() as directory:
+            ptx = Path(directory)
+            env = dict(os.environ, D4R_DLSS_PTX_DIR=directory, D4R_SUST_KINDS='b16,p,b32')
+            for name in names:
+                with self.subTest(kernel=name):
+                    (ptx / 'input.ptx').write_text(f'.visible .entry {name}(\n)\n{{\n{math}\n{store}\nret;\n}}\n')
+                    subprocess.run(['python3', str(ROOT / 'kernels/tex/make_ptx.py'), name, str(ptx / 'out.ptx')],
+                                   env=env, check=True, capture_output=True)
+                    output = (ptx / 'out.ptx').read_text()
+                    self.assertIn(math, output)
+                    self.assertNotIn(store, output)
+                    self.assertIn('call d4r_sust_v2b16', output)
+                    self.assertNotIn('d4r_enc0_tail', output)
+                    self.assertNotIn('d4r_dec0_head', output)
+
+
+if __name__ == '__main__':
+    unittest.main()
diff --git a/tests/test_motion_geometry.py b/tests/test_motion_geometry.py
new file mode 100644
index 0000000..972469c
--- /dev/null
+++ b/tests/test_motion_geometry.py
@@ -0,0 +1,45 @@
+"""Motion-grid selection and padded low-resolution allocation regression."""
+from pathlib import Path
+import subprocess
+import tempfile
+import unittest
+
+ROOT = Path(__file__).resolve().parents[1]
+
+
+class MotionGeometryTests(unittest.TestCase):
+    def test_low_res_padding_flags_and_display_res_mapping(self):
+        source = r'''
+#include "d4r_motion_dilation.h"
+#include <cassert>
+int main() {
+    D4rMotionDilationParams p = {};
+    p.render_width = 854; p.render_height = 480;
+    p.motion_width = 856; p.motion_height = 480;
+    p.motion_pitch = 4096; p.jitter_x = 0.375f; p.jitter_y = -0.25f;
+    auto low = p;
+    assert(d4r_motion_dilation_geometry(&low, 0xb, 2560, 1440));
+    assert(low.motion_width == 854 && low.motion_height == 480);
+    assert(low.motion_pitch == 4096 && low.jitter_x == 0 && low.jitter_y == 0);
+    low = p;
+    assert(d4r_motion_dilation_geometry(&low, 0x4b, 2560, 1440));
+    low = p; low.motion_width = 853;
+    assert(!d4r_motion_dilation_geometry(&low, 0xb, 2560, 1440));
+    low = p; low.motion_height = 479;
+    assert(!d4r_motion_dilation_geometry(&low, 0xb, 2560, 1440));
+    low = p;
+    assert(!d4r_motion_dilation_geometry(&low, 0xf, 2560, 1440));
+    assert(!d4r_motion_dilation_geometry(&p, 0x49, 2560, 1440));
+    p.motion_width = 2560; p.motion_height = 1440;
+    assert(d4r_motion_dilation_geometry(&p, 0x49, 2560, 1440));
+    assert(p.jitter_x == 0.375f && p.jitter_y == -0.25f);
+    p.render_width = 0;
+    assert(!d4r_motion_dilation_geometry(&p, 0x49, 2560, 1440));
+}
+'''
+        with tempfile.TemporaryDirectory() as tmp:
+            runner = Path(tmp) / "geometry"
+            subprocess.run(["c++", "-x", "c++", "-std=c++17", "-Wall", "-Wextra", "-Werror",
+                            "-I", str(ROOT / "tools"), "-o", str(runner), "-"],
+                           input=source, text=True, check=True, capture_output=True)
+            subprocess.run([str(runner)], check=True, timeout=5)
diff --git a/tests/test_motion_kernel_gpu.py b/tests/test_motion_kernel_gpu.py
new file mode 100644
index 0000000..b1f5758
--- /dev/null
+++ b/tests/test_motion_kernel_gpu.py
@@ -0,0 +1,123 @@
+"""Optional GPU regression: independently specified thin-foreground velocity coverage.
+
+D4R_RUN_GPU_TESTS=1 D4R_DUMP_RUNNER=/path/to/dump_runner python3 -m unittest discover -s tests -v
+Build the embedded kernels first with scripts/build_d4r_motion_kernels.sh.
+"""
+import os
+from pathlib import Path
+import struct
+import subprocess
+import tempfile
+import unittest
+
+ROOT = Path(__file__).resolve().parents[1]
+
+
+@unittest.skipUnless(os.environ.get('D4R_RUN_GPU_TESTS') == '1', 'opt-in GPU test')
+class MotionKernelGpuTests(unittest.TestCase):
+    def test_low_resolution_opaque_edges_and_padded_rows(self):
+        runner = os.environ['D4R_DUMP_RUNNER']
+        kernel = ROOT / 'build/wine-nvcuda/generated' / (os.environ.get('D4R_GPU_ARCH', 'gfx1101') + '.hsaco')
+        width, height, stride = 5, 3, 8
+        sentinel = 0xA5A5A5A5
+        for reversed_depth in (False, True):
+            for radius in (1, 2):
+                with self.subTest(reversed_depth=reversed_depth, radius=radius), tempfile.TemporaryDirectory() as t:
+                    directory = Path(t); output = directory / 'out'; output.mkdir()
+                    depth = [0.1 if reversed_depth else 0.9] * (stride*height)
+                    depth[stride+2] = 0.9 if reversed_depth else 0.1
+                    motion = list(range(stride*height))
+                    arrays = [struct.pack('<24f', *depth), struct.pack('<24I', *motion), struct.pack('<24I', *([sentinel]*24))]
+                    addresses = (0x100000000, 0x200000000, 0x300000000)
+                    args = struct.pack('<5Q4I2f2I', *addresses, stride*4, stride*4,
+                                       width, height, width, height, 0.0, 0.0, radius, reversed_depth)
+                    (directory/'args.bin').write_bytes(args)
+                    manifest = ['kernel d4r_motion_dilate', 'launch 1 1 1 128 1 1 0', 'args 72']
+                    for i, data in enumerate(arrays):
+                        (directory/f'alloc-{i}.bin').write_bytes(data)
+                        manifest += [f'alloc {i} {hex(addresses[i])} {len(data)}', f'pointer {i*8} {i} 0']
+                    (directory/'manifest.txt').write_text('\n'.join(manifest)+'\n')
+                    subprocess.run([runner, str(kernel), 'd4r_motion_dilate', str(directory), str(output), '1'], check=True, capture_output=True)
+                    self.assertEqual((output/'alloc-0.bin').read_bytes(), arrays[0])
+                    self.assertEqual((output/'alloc-1.bin').read_bytes(), arrays[1])
+                    expected = [sentinel] * 24
+                    for y in range(height):
+                        for x in range(width):
+                            expected[y*stride+x] = motion[stride+2] if abs(x-2) <= radius and abs(y-1) <= radius else motion[y*stride+x]
+                    self.assertEqual((output/'alloc-2.bin').read_bytes(), struct.pack('<24I', *expected))
+
+    def test_foreground_coverage_padding_and_input_immutability(self):
+        runner = os.environ['D4R_DUMP_RUNNER']
+        kernel = ROOT / 'build/wine-nvcuda/generated' / (os.environ.get('D4R_GPU_ARCH', 'gfx1101') + '.hsaco')
+        sentinel = 0xA5A5A5A5
+        for reversed_depth in (False, True):
+            for radius in (1, 2):
+                for jitter in ((0.0, 0.0), (0.25, -0.25)):
+                    with self.subTest(reversed_depth=reversed_depth, radius=radius, jitter=jitter), tempfile.TemporaryDirectory() as t:
+                        directory = Path(t); output = directory / 'out'; output.mkdir()
+                        depth = [1000.0] * 16
+                        for y in range(2):
+                            for x in range(3):
+                                depth[y*8+x] = 0.1 if reversed_depth else 0.9
+                        depth[1] = 0.9 if reversed_depth else 0.1
+                        motion = list(range(96))
+                        arrays = [struct.pack('<16f', *depth), struct.pack('<96I', *motion), struct.pack('<96I', *([sentinel]*96))]
+                        addresses = (0x100000000, 0x200000000, 0x300000000)
+                        args = struct.pack('<5Q4I2f2I', *addresses, 32, 64, 3, 2, 9, 6, *jitter, radius, reversed_depth)
+                        (directory/'args.bin').write_bytes(args)
+                        manifest = ['kernel d4r_motion_dilate', 'launch 1 1 1 128 1 1 0', 'args 72']
+                        for i, data in enumerate(arrays):
+                            (directory/f'alloc-{i}.bin').write_bytes(data)
+                            manifest += [f'alloc {i} {hex(addresses[i])} {len(data)}', f'pointer {i*8} {i} 0']
+                        (directory/'manifest.txt').write_text('\n'.join(manifest)+'\n')
+                        subprocess.run([runner, str(kernel), 'd4r_motion_dilate', str(directory), str(output), '1'], check=True, capture_output=True)
+                        self.assertEqual((output/'alloc-0.bin').read_bytes(), arrays[0])
+                        self.assertEqual((output/'alloc-1.bin').read_bytes(), arrays[1])
+                        expected = [sentinel] * 96
+                        # Known single foreground cell: all neighbouring background follows its velocity.
+                        # Foreground itself preserves its original per-pixel velocity; row padding is untouched.
+                        for y in range(6):
+                            for x in range(9):
+                                inside = (3 <= x <= 5 and y <= 2) if jitter == (0.0, 0.0) else (2 <= x <= 4 and y <= 3)
+                                expected[y*16+x] = motion[y*16+x] if inside else motion[1*16+4] if jitter == (0.0, 0.0) else motion[2*16+3]
+                        self.assertEqual((output/'alloc-2.bin').read_bytes(), struct.pack('<96I', *expected))
+
+    def test_automatic_sky_fill_circle_nearest_ties_and_empty_sky(self):
+        runner = os.environ['D4R_DUMP_RUNNER']
+        kernel = ROOT / 'build/wine-nvcuda/generated' / (os.environ.get('D4R_GPU_ARCH', 'gfx1101') + '.hsaco')
+        width, height, stride = 23, 23, 32
+        sentinel = 0xA5A5A5A5
+        for reversed_depth in (False, True):
+            for radius in (1, 2):
+                for empty, near_far in ((False, False), (True, False), (False, True)):
+                    with self.subTest(reversed_depth=reversed_depth, radius=radius, empty=empty, near_far=near_far), tempfile.TemporaryDirectory() as t:
+                        directory = Path(t); output = directory / 'out'; output.mkdir()
+                        depth = [1000.0] * (stride*height)
+                        for y in range(height):
+                            for x in range(width):
+                                depth[y*stride+x] = (1.e-7 if near_far else 0.0) if reversed_depth else 1.0
+                        sources = [] if empty else [(8, 11), (14, 11)]
+                        for x, y in sources:
+                            depth[y*stride+x] = (1.5e-7 if near_far else 0.9) if reversed_depth else (0.99 if near_far else 0.1)
+                        motion = list(range(stride*height))
+                        arrays = [struct.pack(f'<{len(depth)}f', *depth), struct.pack(f'<{len(motion)}I', *motion), struct.pack(f'<{len(motion)}I', *([sentinel]*len(motion)))]
+                        addresses = (0x100000000, 0x200000000, 0x300000000)
+                        args = struct.pack('<5Q4I2f2I', *addresses, stride*4, stride*4, width, height, width, height, 0.0, 0.0, radius, reversed_depth)
+                        (directory/'args.bin').write_bytes(args)
+                        manifest = ['kernel d4r_motion_dilate', 'launch 5 1 1 128 1 1 0', 'args 72']
+                        for i, data in enumerate(arrays):
+                            (directory/f'alloc-{i}.bin').write_bytes(data)
+                            manifest += [f'alloc {i} {hex(addresses[i])} {len(data)}', f'pointer {i*8} {i} 0']
+                        (directory/'manifest.txt').write_text('\n'.join(manifest)+'\n')
+                        subprocess.run([runner, str(kernel), 'd4r_motion_dilate', str(directory), str(output), '1'], check=True, capture_output=True)
+                        self.assertEqual((output/'alloc-0.bin').read_bytes(), arrays[0])
+                        self.assertEqual((output/'alloc-1.bin').read_bytes(), arrays[1])
+                        expected = [sentinel] * len(motion)
+                        for y in range(height):
+                            for x in range(width):
+                                candidates = sorted(((x-sx)**2+(y-sy)**2, sx, sy) for sx, sy in sources)
+                                expected[y*stride+x] = motion[y*stride+x]
+                                if candidates and candidates[0][0] <= 64:
+                                    _, sx, sy = candidates[0]
+                                    expected[y*stride+x] = motion[sy*stride+sx]
+                        self.assertEqual((output/'alloc-2.bin').read_bytes(), struct.pack(f'<{len(expected)}I', *expected))
diff --git a/tools/d3d12_dlss_harness.cpp b/tools/d3d12_dlss_harness.cpp
index 3dba5f5..2217e71 100644
--- a/tools/d3d12_dlss_harness.cpp
+++ b/tools/d3d12_dlss_harness.cpp
@@ -20,6 +20,9 @@
 #include <cstring>
 #include <string>
 #include <vector>
+#include <atomic>
+#include <chrono>
+#include "d4r_vkd3d_interop.h"
 
 using NgxResult = unsigned int;
 constexpr NgxResult NGX_SUCCESS = 1;
@@ -103,6 +106,90 @@ static void submit_and_wait()
     g_list->Reset(g_allocator, nullptr);
 }
 
+class LifetimeProbe final : public IUnknown
+{
+public:
+    explicit LifetimeProbe(std::atomic<bool>& destroyed) : destroyed_(destroyed) {}
+    HRESULT STDMETHODCALLTYPE QueryInterface(REFIID iid, void** out) override
+    {
+        if (!out) return E_POINTER;
+        *out = nullptr;
+        if (iid != __uuidof(IUnknown)) return E_NOINTERFACE;
+        *out = this; AddRef(); return S_OK;
+    }
+    ULONG STDMETHODCALLTYPE AddRef() override { return ++refs_; }
+    ULONG STDMETHODCALLTYPE Release() override
+    {
+        ULONG refs = --refs_;
+        if (!refs) { destroyed_ = true; delete this; }
+        return refs;
+    }
+private:
+    std::atomic<ULONG> refs_{1};
+    std::atomic<bool>& destroyed_;
+};
+
+static void lifecycle_report(const char* message)
+{
+    std::printf("%s\n", message);
+    std::fflush(stdout);
+    if (const char* path = std::getenv("D4R_HARNESS_TEST_REPORT"))
+        if (FILE* file = std::fopen(path, "a"))
+        {
+            std::fprintf(file, "%s\n", message);
+            std::fclose(file);
+        }
+}
+
+static bool test_external_lifetime()
+{
+    ID3D12DXVKInteropDeviceD4R2* interop = nullptr;
+    if (!check(g_device->QueryInterface(__uuidof(ID3D12DXVKInteropDeviceD4R2),
+                                       reinterpret_cast<void**>(&interop)), "lifetime interface"))
+        return false;
+    for (bool submitted : {false, true})
+    {
+        std::atomic<bool> destroyed{false};
+        auto* owner = new LifetimeProbe(destroyed);
+        const HRESULT retained = interop->RetainExternalResources(g_list, owner);
+        const HRESULT repeated = interop->RetainExternalResources(g_list, owner);
+        owner->Release();
+        if (!check(retained, "retain resources") || !check(repeated, "deduplicate resources") || destroyed)
+            return false;
+        g_list->Close();
+        ID3D12Fence* gate = nullptr;
+        if (submitted)
+        {
+            if (!check(g_device->CreateFence(0, D3D12_FENCE_FLAG_NONE, __uuidof(ID3D12Fence),
+                                            reinterpret_cast<void**>(&gate)), "lifetime gate")) return false;
+            g_queue->Wait(gate, 1);
+            ID3D12CommandList* lists[] = {g_list};
+            g_queue->ExecuteCommandLists(1, lists);
+            Sleep(20);
+            g_allocator->Reset(); // Pending submissions must keep their external owners.
+            if (destroyed) { gate->Signal(1); return false; }
+            gate->Signal(1);
+            g_queue->Signal(g_fence, ++g_fenceValue);
+            if (g_fence->GetCompletedValue() < g_fenceValue)
+            {
+                g_fence->SetEventOnCompletion(g_fenceValue, g_event);
+                if (WaitForSingleObject(g_event, 5000) != WAIT_OBJECT_0) return false;
+            }
+            gate->Release();
+        }
+        // The fence callback may take a moment to drop the submission's allocator ref.
+        const ULONGLONG start = GetTickCount64();
+        do { g_allocator->Reset(); if (!destroyed) Sleep(1); }
+        while (!destroyed && GetTickCount64() - start < 5000);
+        if (!destroyed) return false;
+        if (!check(g_list->Reset(g_allocator, nullptr), "lifetime list reset")) return false;
+        lifecycle_report(submitted ? "LIFETIME submitted: retained until safe allocator reset"
+                                   : "LIFETIME discarded: retained until safe allocator reset");
+    }
+    interop->Release();
+    return true;
+}
+
 static ID3D12Resource* create_buffer(D3D12_HEAP_TYPE type, UINT64 size, D3D12_RESOURCE_STATES state)
 {
     D3D12_HEAP_PROPERTIES heap = {};
@@ -606,6 +693,8 @@ int main(int argc, char** argv)
     d4r_ngx_set_int(parameters, "DLSS.Feature.Create.Flags",
                     createFlags != nullptr ? static_cast<int>(std::strtol(createFlags, nullptr, 0)) : 0);
     NgxHandle* feature = nullptr;
+    if (std::getenv("D4R_HARNESS_LIFETIME_TEST") != nullptr && !test_external_lifetime())
+        return 1;
     result = createFeature(g_list, 1, parameters, &feature);
     std::printf("CreateFeature -> 0x%08x handle=%p\n", result, static_cast<void*>(feature));
     if (result != NGX_SUCCESS)
@@ -785,6 +874,8 @@ int main(int argc, char** argv)
             d4r_ngx_set_int(parameters, "Reset", forceReset || frame == 1 ? 1 : 0);
         const DWORD start = GetTickCount();
         result = evaluate(g_list, feature, parameters, nullptr);
+        if (const char* delay = std::getenv("D4R_HARNESS_DELAY_SUBMIT_MS"))
+            Sleep(static_cast<DWORD>(std::clamp(std::atoi(delay), 0, 4000)));
         submit_and_wait();
         std::printf("frame %d: EvaluateFeature -> 0x%08x (%lu ms incl. submit)\n", frame, result, GetTickCount() - start);
         Sleep(frameWaitMs); // let the CUDA worker finish before the next presentation
@@ -885,6 +976,25 @@ int main(int argc, char** argv)
         nonzero += value != 0;
     std::printf("output read back: %zu of %zu bytes nonzero, written to %s\n", nonzero, output.size(), argv[2]);
 
+    if (std::getenv("D4R_HARNESS_DISCARD_EVALUATION") != nullptr)
+    {
+        const NgxResult queued = evaluate(g_list, feature, parameters, nullptr);
+        const ULONGLONG start = GetTickCount64();
+        const NgxResult retired = release(feature);
+        feature = nullptr;
+        if (queued != NGX_SUCCESS || retired != NGX_SUCCESS || GetTickCount64() - start > 2000)
+        {
+            std::fprintf(stderr, "DISCARD evaluation release failed or blocked\n");
+            return 1;
+        }
+        g_list->Close();
+        g_allocator->Reset();
+        g_list->Reset(g_allocator, nullptr);
+        result = createFeature(g_list, 1, parameters, &feature);
+        if (result != NGX_SUCCESS) return 1;
+        lifecycle_report("DISCARD evaluation: release completed without submitting inputs");
+    }
+
     // D4R_HARNESS_RECREATE=N: N more release/create cycles at alternating render sizes (as when a game's DLSS
     // quality setting changes), a few evaluations each, logging this process's VRAM to find leaks per cycle.
     const int recreateCycles = std::getenv("D4R_HARNESS_RECREATE") != nullptr ? std::atoi(std::getenv("D4R_HARNESS_RECREATE")) : 0;
diff --git a/tools/d4r_event_wait.h b/tools/d4r_event_wait.h
new file mode 100644
index 0000000..b048ce2
--- /dev/null
+++ b/tools/d4r_event_wait.h
@@ -0,0 +1,16 @@
+#pragma once
+
+// Query completion rather than trusting HIP event flags to put the CPU to
+// sleep. Never treat NOT_READY as completion, including during retirement.
+template <typename Query, typename Sleep>
+int d4r_wait_event(Query query, Sleep sleep)
+{
+    constexpr int cudaNotReady = 600;
+    for (;;)
+    {
+        const int result = query();
+        if (result != cudaNotReady)
+            return result;
+        sleep();
+    }
+}
diff --git a/tools/d4r_motion_dilation.h b/tools/d4r_motion_dilation.h
new file mode 100644
index 0000000..ece2180
--- /dev/null
+++ b/tools/d4r_motion_dilation.h
@@ -0,0 +1,37 @@
+#pragma once
+#include <stdint.h>
+
+/* Shared Wine/HIP ABI. Pitches are bytes; motion pixels are packed RG16F. */
+typedef struct D4rMotionDilationParams
+{
+    uint64_t depth, motion, output;
+    uint64_t depth_pitch, motion_pitch;
+    uint32_t render_width, render_height, motion_width, motion_height;
+    float jitter_x, jitter_y;
+    uint32_t radius, reversed_depth;
+} D4rMotionDilationParams;
+
+/* Low-resolution velocity and depth share the render raster grid. Allocation
+   padding is not part of that grid; retain the original NGX flags and MV scale.
+   Display-resolution velocity still needs the jittered depth-grid mapping. */
+static inline int d4r_motion_dilation_geometry(D4rMotionDilationParams* p, uint32_t flags,
+    uint32_t output_width, uint32_t output_height)
+{
+    if ((flags & 4u) != 0 || p->render_width == 0 || p->render_height == 0)
+        return 0; /* Jittered motion values remain unsupported. */
+    if ((flags & 2u) != 0)
+    {
+        if (p->motion_width < p->render_width || p->motion_height < p->render_height)
+            return 0;
+        p->motion_width = p->render_width;
+        p->motion_height = p->render_height;
+        p->jitter_x = p->jitter_y = 0;
+        return 1;
+    }
+    return p->motion_width == output_width && p->motion_height == output_height;
+}
+#ifdef __cplusplus
+static_assert(sizeof(D4rMotionDilationParams) == 72, "motion dilation ABI");
+#else
+_Static_assert(sizeof(D4rMotionDilationParams) == 72, "motion dilation ABI");
+#endif
diff --git a/tools/d4r_native_selection.h b/tools/d4r_native_selection.h
index 4b91c4b..3556b21 100644
--- a/tools/d4r_native_selection.h
+++ b/tools/d4r_native_selection.h
@@ -5,6 +5,18 @@
 #include <stdlib.h>
 #include <string.h>
 
+static int d4r_native_swin_encoders_enabled(void)
+{
+    const char* value = getenv("D4R_NATIVE_SWIN_ENCODERS");
+    return value == NULL || strcmp(value, "0") != 0;
+}
+
+static int d4r_native_kernel_allowed(const char* name)
+{
+    return d4r_native_swin_encoders_enabled() ||
+        (strcmp(name, "rrlite_enc1_4x4") != 0 && strcmp(name, "rrlite_enc2_4x4") != 0);
+}
+
 static int d4r_apply_accuracy_policy(void)
 {
     const char* prefer = getenv("D4R_PREFER_ACCURACY");
diff --git a/tools/d4r_nvngx_shim.cpp b/tools/d4r_nvngx_shim.cpp
index 213c97a..31255f6 100644
--- a/tools/d4r_nvngx_shim.cpp
+++ b/tools/d4r_nvngx_shim.cpp
@@ -38,6 +38,7 @@
 #include <windows.h>
 #include <d3d12.h>
 #include <vulkan/vulkan_core.h>
+#include "d4r_motion_dilation.h"
 
 #include <algorithm>
 #include <bit>
@@ -57,6 +58,9 @@
 #include <string>
 #include <thread>
 #include <vector>
+#include "d4r_event_wait.h"
+#include "d4r_win32_wait.h"
+#include "d4r_vkd3d_interop.h"
 
 // --- NGX types (declared locally: the core's exports differ from the SDK) --
 
@@ -221,6 +225,7 @@ struct CudaApi
     int(WINAPI* ctxSynchronize)() = nullptr;
     int(WINAPI* eventCreate)(CudaEvent*, unsigned int) = nullptr;
     int(WINAPI* eventRecord)(CudaEvent, void*) = nullptr;
+    int(WINAPI* eventQuery)(CudaEvent) = nullptr;
     int(WINAPI* eventSynchronize)(CudaEvent) = nullptr;
     int(WINAPI* eventElapsedTime)(float*, CudaEvent, CudaEvent) = nullptr;
     int(WINAPI* eventDestroy)(CudaEvent) = nullptr;
@@ -241,6 +246,7 @@ struct CudaApi
     int(WINAPI* outputKernelNative)() = nullptr;
     // Optional (d4r nvcuda bridge): report a pitch-linear texture to NGX as an array of that size/format.
     int(WINAPI* registerLinearTexture)(CudaObject, size_t, size_t, uint32_t, uint32_t) = nullptr;
+    int(WINAPI* dilateMotion)(const D4rMotionDilationParams*) = nullptr;
 };
 
 // --- official NGX core CUDA API --------------------------------------------
@@ -461,7 +467,7 @@ static std::string preset_number(const std::string& model)
     for (const auto& [letter, number] : letters)
         if (name == letter)
             return std::to_string(number);
-    g_portable.notes.push_back("d4r.ini: [DLSS] Model '" + model + "' is not a preset (use K, E, M or a number); "
+    g_portable.notes.push_back("d4r.ini: [DLSS] Model '" + model + "' is not a preset (use K, E, L, M or a number); "
                                "leaving the preset to the game");
     return std::string();
 }
@@ -531,6 +537,7 @@ static void load_portable_config()
     const std::string native = ascii_lower(ini_value(ini, "kernels", "NativeKernels"));
     // The bridge selects accuracy variants and enforces their compiler/sync policy before cuInit.
     portable_set_unix("D4R_PREFER_ACCURACY", ini_flag(ini, "kernels", "PreferAccuracy", 0) ? "1" : "0");
+    portable_set_unix("D4R_NATIVE_SWIN_ENCODERS", ini_flag(ini, "kernels", "NativeSwinEncoders", 1) ? "1" : "0");
     const bool nativeOn = native.empty() || native == "on" || native == "true" || native == "1" || native == "fast";
     if (!nativeOn && native != "off" && native != "false" && native != "0")
         g_portable.notes.push_back("d4r.ini: [Kernels] NativeKernels must be on or off, not '" + native + "'");
@@ -550,7 +557,7 @@ static void load_portable_config()
         const char* variable;
         int fallback;
     } interop[] = {{"VramInterop", "D4R_SHIM_VRAM_INTEROP", 1}, {"InputSync", "D4R_SHIM_INPUT_SYNC", 1},
-                   {"GpuWait", "D4R_SHIM_GPU_WAIT", 1},         {"LinearInputs", "D4R_SHIM_LINEAR_INPUTS", 1},
+                   {"LinearInputs", "D4R_SHIM_LINEAR_INPUTS", 1},
                    {"EvalSync", "D4R_SHIM_EVAL_SYNC", 0}};
     for (const auto& setting : interop)
         portable_set(setting.variable, ini_flag(ini, "interop", setting.key, setting.fallback) ? "1" : "0");
@@ -562,6 +569,11 @@ static void load_portable_config()
     portable_set("D4R_SHIM_OUTPUT_DIRECT", ini_flag(ini, "interop", "DirectOutput", nativeOn) ? "1" : "0");
 
     portable_set("D4R_SHIM_WATERMARK", ini_flag(ini, "dlss", "ShowWatermark", 0) ? "1" : "0");
+    const std::string dilation = ini_value(ini, "interop", "MotionVectorDilation");
+    if (dilation.empty() || dilation == "0" || dilation == "1" || dilation == "2")
+        portable_set("D4R_MOTION_DILATION", dilation.empty() ? "0" : dilation);
+    else
+        g_portable.notes.push_back("MotionVectorDilation must be 0, 1 or 2; using original motion vectors");
     if (ini_flag(ini, "debug", "Profile", 0))
         portable_set("D4R_PROFILE", "1");
     if (const std::string level = ini_value(ini, "debug", "NgxLogLevel"); !level.empty())
@@ -1121,6 +1133,9 @@ struct Global
     Worker worker;        // every NGX-core and CUDA call
     std::thread::id workerThread;
     Worker prep, finish;  // CPU-only pipeline stages around it
+    Worker cleanup;       // deferred GPU resource destruction, never under allocator locks
+    std::atomic<unsigned int> resourceOwners{0};
+    bool shutdownPending = false; // under mutex
     HMODULE core = nullptr;
     CoreApi ngx;
     HMODULE cuda = nullptr;
@@ -1135,6 +1150,7 @@ struct Global
 };
 
 static Global g;
+static std::mutex g_featureApiMutex;
 
 template <typename Function> static bool load_export(HMODULE module, const char* name, Function& function)
 {
@@ -1305,17 +1321,20 @@ static bool load_libraries()
         g.cu.outputKernelNative = nullptr;
     if (!load_export(g.cuda, "d4rRegisterLinearTexture", g.cu.registerLinearTexture))
         g.cu.registerLinearTexture = nullptr;
-    if (profile_enabled() || env_uint("D4R_SHIM_BLOCKING_SYNC", 0) != 0)
+    g.cu.dilateMotion = reinterpret_cast<decltype(g.cu.dilateMotion)>(
+        reinterpret_cast<void*>(GetProcAddress(g.cuda, "d4rDilateMotionVectors")));
+    if (profile_enabled() || env_uint("D4R_SHIM_BLOCKING_SYNC", 1) != 0)
     {
         const bool create = load_export(g.cuda, "cuEventCreate", g.cu.eventCreate);
         const bool record = load_export(g.cuda, "cuEventRecord", g.cu.eventRecord);
-        const bool wait = env_uint("D4R_SHIM_BLOCKING_SYNC", 0) != 0
+        const bool wait = env_uint("D4R_SHIM_BLOCKING_SYNC", 1) != 0
                               ? load_export(g.cuda, "d4rEventSynchronize", g.cu.eventSynchronize) : true;
+        const bool query = load_export(g.cuda, "cuEventQuery", g.cu.eventQuery);
         const bool destroy = load_export(g.cuda, "cuEventDestroy", g.cu.eventDestroy);
         const bool elapsed = profile_enabled() ? load_export(g.cuda, "cuEventElapsedTime", g.cu.eventElapsedTime) : true;
         if (profile_enabled() && !(create && record && destroy && elapsed))
             logf("D4R_PROFILE: CUDA event timing unavailable; CPU stage timings remain enabled");
-        if (env_uint("D4R_SHIM_BLOCKING_SYNC", 0) != 0 && !(create && record && wait && destroy))
+        if (env_uint("D4R_SHIM_BLOCKING_SYNC", 1) != 0 && !(create && record && wait && query && destroy))
             logf("blocking output wait unavailable; using context sync");
     }
     logf("loaded NGX core %ls and nvcuda.dll: %s", corePath, ok ? "all exports present" : "exports missing");
@@ -1336,14 +1355,22 @@ static NgxResult initialize(unsigned long long applicationId, const wchar_t* dat
                             const ProjectIdentity* project = nullptr)
 {
     std::lock_guard<std::mutex> 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.
+        HMODULE pinned = nullptr;
+        if (!GetModuleHandleExW(GET_MODULE_HANDLE_EX_FLAG_FROM_ADDRESS | GET_MODULE_HANDLE_EX_FLAG_PIN,
+                               reinterpret_cast<LPCWSTR>(g_selfModule), &pinned))
+            return NGX_FAIL_PLATFORM_ERROR;
         g.worker.start();
         g.workerThread = g.worker.call([] { return std::this_thread::get_id(); });
         g.prep.start();
         g.finish.start();
+        g.cleanup.start();
         g.started = true;
     }
     ensure_portable_config();
@@ -1463,10 +1490,24 @@ struct Staging
     UINT width = 0;
     UINT height = 0;
 
+    std::vector<ID3D12Resource*> retired;
+
+    void retire()
+    {
+        if (buffer != nullptr)
+            retired.push_back(buffer);
+        buffer = nullptr;
+        mapped = nullptr;
+        width = height = 0;
+        format = DXGI_FORMAT_UNKNOWN;
+    }
+
     void release()
     {
         if (buffer != nullptr)
             buffer->Release();
+        for (ID3D12Resource* old : retired)
+            old->Release();
         *this = Staging{};
     }
 };
@@ -1490,6 +1531,7 @@ struct FrameParams
     bool hasExposure = false;
     bool vram = false; // inputs and output stay in VRAM (VRAM interop)
     bool split = false; // this frame presents its own result (split frames)
+    unsigned motionDilation = 0; // actually applied, worker only
 };
 
 struct FrameTiming
@@ -1501,6 +1543,7 @@ struct FrameTiming
     double download = 0, outputConvert = 0;
     double prep = 0, worker = 0, workerWait = 0, finishWait = 0, markerWait = 0;
     double ngxHost = 0, ctxSync = 0, gpuEval = -1;
+    double motionDilationWall = 0;
     bool gpuEventsRecorded = false;
     double outputSyncWall = 0, outputSyncCpu = -1;
     bool outputSyncBlocking = false;
@@ -1533,6 +1576,11 @@ struct InputSlot
     CudaDevicePtr linearPointer[4] = {};
     UINT linearWidth[4] = {}, linearHeight[4] = {};
     size_t linearPitch[4] = {};
+    CudaDevicePtr dilatedMotion = 0;
+    size_t dilatedBytes = 0;
+    CudaObject dilatedTexture = 0;
+    UINT dilatedWidth = 0, dilatedHeight = 0;
+    size_t dilatedPitch = 0;
     std::atomic<bool> busy{false}; // readback staging owned by the pipeline
     HostPlane host[4];
     std::atomic<bool> hostBusy{false}; // host planes not yet uploaded by the worker
@@ -1549,6 +1597,8 @@ struct OutputSlot
 struct Feature
 {
     NgxHandle handle = {}; // returned to the caller
+    std::atomic<bool> retiring{false};
+    IUnknown* resources = nullptr; // main reference plus command-allocator references
     NgxHandle* cudaHandle = nullptr;
     void* cudaParams = nullptr;
     CudaDevicePtr scratch = 0;
@@ -1586,26 +1636,20 @@ struct Feature
     VramImage outputConversion;
     // Split frames: the rest of frame N's command list is submitted only once
     // splitSemaphore reaches N (signalled when N retires, or by the watchdog).
-    bool split = false;
+    std::atomic<bool> split{false};
     VkSemaphore splitSemaphore = VK_NULL_HANDLE;
     std::mutex splitMutex;
     uint64_t splitSignalled = 0; // under splitMutex
-    // GPU-side start (D4R_SHIM_GPU_WAIT, split frames): the input copies end with a vkCmdFillBuffer of the
-    // frame number into gpuMarker, and the null stream waits for it on the GPU instead of the CPU polling
-    // the frame marker before any DLSS work is queued.
-    bool gpuWait = false;
-    VramBuffer gpuMarker;
-    // Frame the null stream currently waits for (0: none) and since when; the watchdog releases a wait for a
-    // frame whose command list never reaches the GPU (e.g. recorded but not executed), so DLSS cannot stall.
-    std::atomic<uint32_t> gpuWaitFrame{0};
-    std::atomic<int64_t> gpuWaitSince{0};
+    // Inputs become ready on the CPU before CUDA work is queued. GPU-side waits
+    // for recorded-but-discarded command lists can otherwise wedge the device.
     // D4R_SHIM_OUTPUT_DIRECT: this frame's result is stored by the native output kernel straight into its
     // destination buffer (worker only), so the array -> buffer copy is skipped.
     bool outputRedirected = false;
     bool outputDirectAllowed = false; // the forced preset's output kernel honours the redirect
-    // D4R_SHIM_LINEAR_INPUTS (with gpuWait): NGX samples the interop buffers directly through pitch-linear
+    // D4R_SHIM_LINEAR_INPUTS: NGX samples the interop buffers directly through pitch-linear
     // texture objects (rows padded to 256 bytes), so the buffer -> array copies disappear.
     bool linearInputs = false;
+    bool dilationReported = false;
     std::vector<VramBuffer> retiredBuffers; // may still be read by queued command lists
     std::vector<VramImage> retiredImages;
     ID3D12Resource* marker = nullptr;
@@ -1666,7 +1710,7 @@ static bool ensure_staging(Staging& staging, ID3D12Resource* resource, D3D12_HEA
     if (staging.buffer != nullptr && staging.format == desc.Format && staging.width == desc.Width &&
         staging.height == desc.Height)
         return true;
-    staging.release();
+    staging.retire();
     UINT64 rowSize = 0;
     g.device->GetCopyableFootprints(&desc, 0, 1, 0, &staging.layout, &staging.rows, &rowSize, &staging.total);
     if (!create_buffer(heapType, staging.total, &staging.buffer, &staging.mapped))
@@ -2002,59 +2046,12 @@ static uint32_t plane_channels(Plane plane)
 // worker then only moves data between them and the CUDA arrays on the GPU.
 // Frames whose formats do not qualify use the readback/upload path.
 
-MIDL_INTERFACE("39da4e09-bd1c-4198-9fae-86bbe3be41fd")
-ID3D12DXVKInteropDevice : public IUnknown
-{
-    virtual HRESULT STDMETHODCALLTYPE GetDXGIAdapter(REFIID iid, void** object) = 0;
-    virtual HRESULT STDMETHODCALLTYPE GetInstanceExtensions(UINT* count, const char** extensions) = 0;
-    virtual HRESULT STDMETHODCALLTYPE GetDeviceExtensions(UINT* count, const char** extensions) = 0;
-    virtual HRESULT STDMETHODCALLTYPE GetDeviceFeatures(const VkPhysicalDeviceFeatures2** features) = 0;
-    virtual HRESULT STDMETHODCALLTYPE GetVulkanHandles(VkInstance* instance, VkPhysicalDevice* physical,
-                                                       VkDevice* device) = 0;
-    virtual HRESULT STDMETHODCALLTYPE GetVulkanQueueInfo(ID3D12CommandQueue* queue, VkQueue* vkQueue,
-                                                         UINT32* family) = 0;
-    virtual void STDMETHODCALLTYPE GetVulkanImageLayout(ID3D12Resource* resource, D3D12_RESOURCE_STATES state,
-                                                        VkImageLayout* layout) = 0;
-    virtual HRESULT STDMETHODCALLTYPE GetVulkanResourceInfo(ID3D12Resource* resource, UINT64* handle,
-                                                            UINT64* offset) = 0;
-    virtual HRESULT STDMETHODCALLTYPE LockCommandQueue(ID3D12CommandQueue* queue) = 0;
-    virtual HRESULT STDMETHODCALLTYPE UnlockCommandQueue(ID3D12CommandQueue* queue) = 0;
-};
-
-MIDL_INTERFACE("902d8115-59eb-4406-9518-fe00f991ee65")
-ID3D12DXVKInteropDevice1 : public ID3D12DXVKInteropDevice
-{
-    virtual HRESULT STDMETHODCALLTYPE GetVulkanResourceInfo1(ID3D12Resource* resource, UINT64* handle, UINT64* offset,
-                                                             VkFormat* format) = 0;
-    virtual HRESULT STDMETHODCALLTYPE CreateInteropCommandQueue(const D3D12_COMMAND_QUEUE_DESC* desc, UINT32 family,
-                                                                ID3D12CommandQueue** queue) = 0;
-    virtual HRESULT STDMETHODCALLTYPE CreateInteropCommandAllocator(D3D12_COMMAND_LIST_TYPE type, UINT32 family,
-                                                                    ID3D12CommandAllocator** allocator) = 0;
-    virtual HRESULT STDMETHODCALLTYPE BeginVkCommandBufferInterop(ID3D12CommandList* list, VkCommandBuffer* buffer) = 0;
-    virtual HRESULT STDMETHODCALLTYPE EndVkCommandBufferInterop(ID3D12CommandList* list) = 0;
-};
-__CRT_UUID_DECL(ID3D12DXVKInteropDevice1, 0x902d8115, 0x59eb, 0x4406, 0x95, 0x18, 0xfe, 0x00, 0xf9, 0x91, 0xee, 0x65)
-
-// Added by the d4r vkd3d-proton patch (patches/vkd3d-proton/): ends the list's
-// current Vulkan command buffer; what is recorded afterwards is submitted
-// separately, with a queue-level wait for the timeline semaphore value.
-MIDL_INTERFACE("5a7c8b3e-2f61-4d0e-9c1a-7e3b52d4a901")
-ID3D12DXVKInteropDeviceD4R : public ID3D12DXVKInteropDevice1
-{
-    virtual HRESULT STDMETHODCALLTYPE LockVulkanQueue(ID3D12CommandQueue* queue) = 0;
-    virtual HRESULT STDMETHODCALLTYPE UnlockVulkanQueue(ID3D12CommandQueue* queue) = 0;
-    virtual HRESULT STDMETHODCALLTYPE GetVulkanHeapInfo(ID3D12Heap* heap, UINT64* memory, UINT64* offset,
-                                                        UINT32* type) = 0;
-    virtual HRESULT STDMETHODCALLTYPE SplitCommandListForExternalWait(ID3D12CommandList* list, UINT64 semaphore,
-                                                                      UINT64 value) = 0;
-};
-__CRT_UUID_DECL(ID3D12DXVKInteropDeviceD4R, 0x5a7c8b3e, 0x2f61, 0x4d0e, 0x9c, 0x1a, 0x7e, 0x3b, 0x52, 0xd4, 0xa9, 0x01)
-
 struct VulkanInterop
 {
     std::once_flag once;
     bool ready = false;
     ID3D12DXVKInteropDevice1* interop = nullptr;
+    ID3D12DXVKInteropDeviceD4R2* lifetime = nullptr;
     ID3D12DXVKInteropDeviceD4R* split = nullptr; // only with the d4r vkd3d-proton patch
     VkDevice device = VK_NULL_HANDLE;
     VkPhysicalDeviceMemoryProperties memory = {};
@@ -2159,6 +2156,8 @@ static void init_vram_interop()
     load(g_vk.signalSemaphore, "vkSignalSemaphore");
     if (FAILED(g.device->QueryInterface(__uuidof(ID3D12DXVKInteropDeviceD4R), reinterpret_cast<void**>(&g_vk.split))))
         g_vk.split = nullptr;
+    if (FAILED(g.device->QueryInterface(__uuidof(ID3D12DXVKInteropDeviceD4R2), reinterpret_cast<void**>(&g_vk.lifetime))))
+        g_vk.lifetime = nullptr;
     g_vk.ready = ok;
     logf("VRAM interop: %s (VkDevice %p), split frames %s", ok ? "ready" : "missing Vulkan entry points",
          static_cast<void*>(g_vk.device), g_vk.split != nullptr ? "available" : "unavailable (stock vkd3d-proton)");
@@ -2492,7 +2491,7 @@ static VkBufferImageCopy full_region(const VramCopy& copy)
 // Records the input copies into the slot's buffers. The resources are in
 // COPY_SOURCE state already.
 static bool record_vram_inputs(Feature& feature, ID3D12GraphicsCommandList* list, InputSlot& slot,
-                               const VramCopy* copies, int count, uint32_t frame)
+                               const VramCopy* copies, int count)
 {
     VkCommandBuffer cmd = VK_NULL_HANDLE;
     if (FAILED(g_vk.interop->BeginVkCommandBufferInterop(list, &cmd)))
@@ -2548,9 +2547,6 @@ static bool record_vram_inputs(Feature& feature, ID3D12GraphicsCommandList* list
     }
     g_vk.barrier(cmd, VK_PIPELINE_STAGE_TRANSFER_BIT, VK_PIPELINE_STAGE_ALL_COMMANDS_BIT, 0, 1, &kAfterTransfer, 0,
                  nullptr, 0, nullptr);
-    // After the barrier: the copies are complete and their writes available when the value lands.
-    if (feature.gpuWait)
-        g_vk.fill(cmd, feature.gpuMarker.buffer, 0, 4, frame);
     return SUCCEEDED(g_vk.interop->EndVkCommandBufferInterop(list));
 }
 
@@ -2623,7 +2619,8 @@ static void signal_split(Feature* feature, uint64_t value, const char* reason)
     info.semaphore = feature->splitSemaphore;
     info.value = value;
     const VkResult result = g_vk.signalSemaphore(g_vk.device, &info);
-    feature->splitSignalled = value;
+    if (result == VK_SUCCESS)
+        feature->splitSignalled = value;
     if (result != VK_SUCCESS || reason != nullptr)
         logf("split frame %llu released%s%s (VkResult %d)", static_cast<unsigned long long>(value),
              reason != nullptr ? " by " : "", reason != nullptr ? reason : "", result);
@@ -2661,22 +2658,17 @@ static void split_watchdog()
                 std::lock_guard<std::mutex> splitLock(feature->splitMutex);
                 signalled = feature->splitSignalled;
             }
-            // GPU-side input wait for a frame the game's queue has not reached for over a second
-            const uint32_t waiting = feature->gpuWaitFrame.load();
-            if (feature->gpuWait && waiting != 0 && static_cast<int32_t>(marker - waiting) < 0 &&
-                now.time_since_epoch().count() - feature->gpuWaitSince.load() >
-                    std::chrono::duration_cast<std::chrono::steady_clock::duration>(std::chrono::seconds(1)).count() &&
-                [&] { uint32_t expected = waiting; return feature->gpuWaitFrame.compare_exchange_strong(expected, 0u); }())
-            {
-                logf("frame %u: its inputs never reached the GPU (marker %u); releasing the GPU-side wait", waiting, marker);
-                g.cu.writeValue32(feature->gpuMarker.device, waiting);
-            }
             if (marker <= signalled)
                 continue; // nothing on the GPU is waiting on us
             if (it->second.marker != marker)
                 it->second = Seen{marker, now};
             else if (now - it->second.at > timeout)
-                signal_split(feature, marker, "the watchdog");
+            {
+                // A slow producer is not a completed producer. Releasing this
+                // wait would race the game's output copy with a late CUDA write.
+                logf("split frame %u still waiting for DLSS; preserving output synchronization", marker);
+                it->second.at = now;
+            }
         }
         // Forget released features.
         seen.erase(std::remove_if(seen.begin(), seen.end(),
@@ -2735,7 +2727,6 @@ static void release_vram(Feature& feature)
             recycle_vram_buffer(buffer);
     for (OutputSlot& slot : feature.outputs)
         recycle_vram_buffer(slot.vram);
-    recycle_vram_buffer(feature.gpuMarker);
     for (VramBuffer& buffer : feature.retiredBuffers)
         recycle_vram_buffer(buffer);
     feature.retiredBuffers.clear();
@@ -3306,7 +3297,7 @@ static bool wait_for_release(const std::atomic<bool>& flag)
     {
         if (std::chrono::steady_clock::now() - start > std::chrono::seconds(2))
             return false;
-        std::this_thread::sleep_for(std::chrono::microseconds(100));
+        d4r_sleep_us(100);
     }
     return true;
 }
@@ -3516,6 +3507,91 @@ static bool ensure_linear_inputs(InputSlot& slot, bool hasExposure)
     return true;
 }
 
+// VRAM interop, worker thread: moves the planes from the imported buffers
+// Opt-in thin-feature motion coverage. Keep the original imported velocity buffer
+// immutable; the new texture is owned by this input slot until its GPU work retires.
+static CudaObject prepare_dilated_motion(Feature& feature, InputSlot& slot, const FrameParams& p)
+{
+    const unsigned radius = env_uint("D4R_MOTION_DILATION", 0);
+    if (radius == 0)
+        return 0;
+    const auto& depth = slot.host[1];
+    const auto& motion = slot.host[2];
+    D4rMotionDilationParams args = {slot.vram[1].device, slot.vram[2].device, slot.dilatedMotion,
+        depth.rowBytes, motion.rowBytes, p.renderWidth, p.renderHeight, motion.width, motion.height,
+        p.jitterX, p.jitterY, radius, (feature.flags & 8u) != 0};
+    const bool motionGeometry = d4r_motion_dilation_geometry(&args, feature.flags,
+                                                            feature.outWidth, feature.outHeight);
+    const bool supported = radius <= 2 && p.vram && feature.linearInputs && g.cu.dilateMotion != nullptr &&
+        motionGeometry && p.colorBaseX == 0 && p.colorBaseY == 0 && p.depthBaseX == 0 &&
+        p.depthBaseY == 0 && p.mvBaseX == 0 && p.mvBaseY == 0 && p.outputBaseX == 0 && p.outputBaseY == 0 &&
+        p.renderWidth > 0 && p.renderHeight > 0 && std::isfinite(p.jitterX) && std::isfinite(p.jitterY) &&
+        p.renderWidth <= depth.width && p.renderHeight <= depth.height;
+    if (!supported)
+    {
+        if (!feature.dilationReported)
+            logf("motion-vector dilation unavailable for this feature; original motion vectors retained");
+        feature.dilationReported = true;
+        return 0;
+    }
+    const size_t bytes = motion.size();
+    const bool recreate = slot.dilatedBytes < bytes || slot.dilatedWidth != motion.width ||
+        slot.dilatedHeight != motion.height || slot.dilatedPitch != motion.rowBytes;
+    if (recreate)
+    {
+        if (slot.dilatedTexture != 0)
+            g.cu.texObjectDestroy(slot.dilatedTexture), slot.dilatedTexture = 0;
+        if (slot.dilatedMotion != 0)
+            g.cu.memFree(slot.dilatedMotion), slot.dilatedMotion = 0;
+        slot.dilatedBytes = 0;
+        if (g.cu.memAlloc(&slot.dilatedMotion, bytes) != 0)
+            return 0;
+        slot.dilatedBytes = bytes;
+        struct Pitch2D { CudaDevicePtr pointer; uint32_t format, channels; size_t width, height, pitch; };
+        const Pitch2D pitch = {slot.dilatedMotion, plane_format(Plane::Motion), 2, motion.width, motion.height, motion.rowBytes};
+        CudaResourceDesc resource = {};
+        resource.resType = 3;
+        static_assert(sizeof(pitch) == 40);
+        std::memcpy(&resource.res, &pitch, sizeof(pitch));
+        const CudaTextureDesc sampler = input_sampler(motion.width, 2, false);
+        int result = g.cu.texObjectCreate(&slot.dilatedTexture, &resource, &sampler, nullptr);
+        if (result == 0)
+            result = g.cu.registerLinearTexture(slot.dilatedTexture, motion.width, motion.height, pitch.format, 2);
+        if (result != 0)
+        {
+            if (slot.dilatedTexture != 0)
+                g.cu.texObjectDestroy(slot.dilatedTexture), slot.dilatedTexture = 0;
+            return 0;
+        }
+        slot.dilatedWidth = motion.width, slot.dilatedHeight = motion.height, slot.dilatedPitch = motion.rowBytes;
+    }
+    if (slot.dilatedTexture == 0)
+        return 0;
+    args.output = slot.dilatedMotion;
+    // Preserve allocation padding, which DLSS may sample at the active edge.
+    // The dilation kernel only writes the active render grid for low-res MVs.
+    if (args.motion_width != motion.width || args.motion_height != motion.height)
+    {
+        CudaMemcpy2D copy = {};
+        copy.srcMemoryType = copy.dstMemoryType = CUDA_MEMORY_DEVICE;
+        copy.srcDevice = args.motion;
+        copy.dstDevice = args.output;
+        copy.srcPitch = copy.dstPitch = motion.rowBytes;
+        copy.WidthInBytes = motion.rowBytes;
+        copy.Height = motion.height;
+        if (copy_2d(copy) != 0)
+            return 0;
+    }
+    const int result = g.cu.dilateMotion(&args);
+    if (!feature.dilationReported)
+        logf("motion-vector dilation %s: radius=%u render pixels, GPU status=%d, %s active=%ux%u allocation=%ux%u",
+             result == 0 ? "enabled" : "unavailable", radius, result,
+             (feature.flags & 2u) != 0 ? "low-resolution" : "display-resolution",
+             args.motion_width, args.motion_height, motion.width, motion.height);
+    feature.dilationReported = true;
+    return result == 0 ? slot.dilatedTexture : 0;
+}
+
 // VRAM interop, worker thread: moves the planes from the imported buffers
 // into the CUDA arrays on the GPU. With D4R_SHIM_VRAM_VERIFY also compares
 // them with the host-staged planes of the same frame.
@@ -3608,6 +3684,32 @@ static bool upload_inputs_vram(Feature& feature, InputSlot& slot, bool hasExposu
     return result == 0;
 }
 
+// Only called on the CUDA worker. A completion event covers queued work on
+// the default stream; do not rely on HIP's blocking flag to release the CPU.
+static int synchronize_default_stream(Feature& feature, bool* querySleep = nullptr)
+{
+    if (feature.outputReadyEvent != nullptr)
+    {
+        const int recorded = g.cu.eventRecord(feature.outputReadyEvent, nullptr);
+        if (recorded == 0)
+        {
+            const int result = d4r_wait_event(
+                [&] { return g.cu.eventQuery(feature.outputReadyEvent); },
+                [] { d4r_sleep_us(200); });
+            if (result == 0)
+            {
+                if (querySleep != nullptr) *querySleep = true;
+                // The event is already complete; this also flushes bridge kernel profiles.
+                return g.cu.eventSynchronize(feature.outputReadyEvent);
+            }
+            logf("output event query failed (%d); using context synchronization", result);
+        }
+        else
+            logf("output event record failed (%d); using context synchronization", recorded);
+    }
+    return g.cu.ctxSynchronize();
+}
+
 // VRAM interop, worker thread: copies the finished result into a free output
 // slot's imported buffer and publishes it; evaluate() copies it on to the game.
 static void publish_vram(Feature* feature, uint32_t frame, const FrameParams& params, FrameTiming& timing,
@@ -3638,23 +3740,7 @@ static void publish_vram(Feature* feature, uint32_t frame, const FrameParams& pa
     {
         const auto syncStart = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{};
         const double cpuStart = timing.enabled ? profile_thread_cpu_ms() : -1.0;
-        if (feature->outputReadyEvent != nullptr)
-        {
-            const int recordResult = g.cu.eventRecord(feature->outputReadyEvent, nullptr);
-            const int waitResult = recordResult == 0 ? g.cu.eventSynchronize(feature->outputReadyEvent) : recordResult;
-            if (waitResult == 0)
-                timing.outputSyncBlocking = true;
-            else
-            {
-                logf("frame %u: blocking output wait failed (record %d, wait %d); using context sync", frame,
-                     recordResult, waitResult);
-                result = g.cu.ctxSynchronize();
-                g.cu.eventDestroy(feature->outputReadyEvent);
-                feature->outputReadyEvent = nullptr;
-            }
-        }
-        else
-            result = g.cu.ctxSynchronize();
+        result = synchronize_default_stream(*feature, &timing.outputSyncBlocking);
         if (timing.enabled)
         {
             timing.outputSyncWall = profile_ms(syncStart, ProfileClock::now());
@@ -3746,22 +3832,9 @@ 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;
-    const bool gpuWait = params.vram && feature->gpuWait;
-    if (params.vram && !gpuWait && env_uint("D4R_SHIM_VRAM_VERIFY", 0) == 0)
+    if (params.vram && env_uint("D4R_SHIM_VRAM_VERIFY", 0) == 0)
         dump_vram_inputs(slot, frame, params, params.hasExposure ? 4 : 3);
-    if (gpuWait)
-    {
-        feature->gpuWaitSince = std::chrono::steady_clock::now().time_since_epoch().count();
-        feature->gpuWaitFrame = frame;
-    }
-    if (gpuWait && g.cu.streamWaitValue32(feature->gpuMarker.device, frame) != 0)
-    {
-        logf("frame %u: queueing the GPU-side input wait failed", frame);
-        slot.busy = false;
-        frame_retired(feature, frame);
-        return;
-    }
-    const bool linearInputs = gpuWait && feature->linearInputs;
+    const bool linearInputs = params.vram && feature->linearInputs;
     bool ok = linearInputs ? ensure_linear_inputs(slot, params.hasExposure)
               : params.vram ? upload_inputs_vram(*feature, slot, params.hasExposure, frame, stages)
               : overlapped
@@ -3773,22 +3846,30 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram
                          upload_plane(Plane::Exposure, slot.host[3], feature->exposure, frame, stages, 3));
     ok = ok && ensure_cuda_image(feature->output, feature->outputWidth, feature->outputHeight, CUDA_FORMAT_HALF, 4,
                                  true, false);
+    const auto dilationStart = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{};
+    const CudaObject dilatedMotion = ok ? prepare_dilated_motion(*feature, slot, params) : 0;
     // 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
     // frame's colour/depth/motion: invisible when the camera is still, beaded
     // 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;
-    // With gpuWait everything up to the output copy is queued in order on the null stream.
-    if (ok && inputSync && !gpuWait)
+    // Linear inputs are already ready; other copies need completion before NGX samples them.
+    if (ok && (dilatedMotion != 0 || (inputSync && !linearInputs)))
     {
-        const int syncResult = g.cu.ctxSynchronize();
+        const int syncResult = synchronize_default_stream(*feature);
         if (syncResult != 0)
         {
             logf("frame %u: input copy synchronize failed: %d", frame, syncResult);
             ok = false;
         }
     }
+    if (dilatedMotion != 0 && ok)
+    {
+        params.motionDilation = env_uint("D4R_MOTION_DILATION", 0);
+        if (timing.enabled)
+            timing.motionDilationWall = profile_ms(dilationStart, ProfileClock::now());
+    }
     slot.hostBusy = false; // the host planes are no longer needed
     // The previous frame's download ran alongside these uploads.
     complete_pending_download();
@@ -3809,7 +3890,7 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram
         // the parameters point at these variables
         feature->colorHandle = slot.linearTexture[0];
         feature->depthHandle = slot.linearTexture[1];
-        feature->motionHandle = slot.linearTexture[2];
+        feature->motionHandle = dilatedMotion != 0 ? dilatedMotion : slot.linearTexture[2];
         feature->exposureHandle = params.hasExposure ? slot.linearTexture[3] : feature->exposureHandle;
     }
     static const bool outputDirect = env_uint("D4R_SHIM_OUTPUT_DIRECT", 0) != 0;
@@ -3855,24 +3936,18 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram
     // D4R_SHIM_EVAL_SYNC=0 (VRAM interop): no wait between NGX's kernels and the output copy queued
     // behind them on the same stream; publish_vram synchronises before the result is released.
     static const bool evalSync = env_uint("D4R_SHIM_EVAL_SYNC", 1) != 0;
-    const int syncResult = params.vram && !evalSync ? 0 : g.cu.ctxSynchronize();
+    const int syncResult = params.vram && !evalSync ? 0 : synchronize_default_stream(*feature);
     const auto evaluated = ProfileClock::now();
     const bool gpuEventsRecorded = eventStartResult == 0 && eventEndResult == 0;
     float gpuEvalMs = -1.0f;
     if (gpuEventsRecorded && syncResult == 0 && (!params.vram || evalSync) &&
         g.cu.eventElapsedTime(&gpuEvalMs, feature->profileStart, feature->profileEnd) != 0)
         gpuEvalMs = -1.0f;
-    if (params.vram && !gpuWait)
-        slot.busy = false; // the input copies ran before the evaluation
+    // Keep the input slot owned until output completion, including linear reads.
     if (result != NGX_SUCCESS || syncResult != 0)
     {
         logf("frame %u: CUDA_EvaluateFeature -> 0x%08x, sync %d", frame, result, syncResult);
-        if (gpuWait)
-        {
-            g.cu.ctxSynchronize();
-            feature->gpuWaitFrame = 0;
-            slot.busy = false;
-        }
+        slot.busy = false;
         frame_retired(feature, frame);
         return;
     }
@@ -3886,11 +3961,7 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram
     if (params.vram)
     {
         publish_vram(feature, frame, params, timing, start);
-        if (gpuWait)
-        {
-            feature->gpuWaitFrame = 0;
-            slot.busy = false; // publish_vram waited for the whole frame, input copies included
-        }
+        slot.busy = false; // output synchronization completed all reads of this input slot
         return;
     }
 
@@ -3932,13 +4003,14 @@ static void prepare_inputs(Feature* feature, int slotIndex, uint32_t frame, Fram
 {
     InputSlot& slot = feature->inputs[slotIndex];
     const auto start = ProfileClock::now();
-    // Wait until the game's command list with this frame's input copies ran (with gpuWait the GPU waits).
-    while (!(params.vram && feature->gpuWait) && static_cast<int32_t>(*feature->markerValue - frame) < 0)
+    // Wait for actual input readiness without queueing an unbounded GPU-side wait.
+    while (static_cast<int32_t>(*feature->markerValue - frame) < 0)
     {
         std::atomic_thread_fence(std::memory_order_acquire);
-        if (std::chrono::steady_clock::now() - start > std::chrono::seconds(5))
+        if (feature->retiring.load() || std::chrono::steady_clock::now() - start > std::chrono::seconds(5))
         {
-            logf("frame %u: timed out waiting for GPU marker (at %u); dropping", frame, *feature->markerValue);
+            logf("frame %u: %s waiting for GPU marker (at %u); dropping", frame,
+                 feature->retiring.load() ? "feature retired while" : "timed out", *feature->markerValue);
             slot.busy = false;
             frame_retired(feature, frame);
             return;
@@ -3948,7 +4020,7 @@ static void prepare_inputs(Feature* feature, int slotIndex, uint32_t frame, Fram
         if (markerPollUs == 0)
             std::this_thread::yield();
         else
-            std::this_thread::sleep_for(std::chrono::microseconds(markerPollUs));
+            d4r_sleep_us(markerPollUs);
     }
     const auto ready = ProfileClock::now();
     // With VRAM interop the inputs are already on the GPU; host staging runs
@@ -4013,7 +4085,7 @@ static int claim_output_slot(Feature* feature)
         }
         if (std::chrono::steady_clock::now() - start > std::chrono::milliseconds(50))
             return -1;
-        std::this_thread::sleep_for(std::chrono::microseconds(200));
+        d4r_sleep_us(200);
     }
 }
 
@@ -4046,10 +4118,10 @@ static void finish_publish(Feature* feature, int target, uint32_t frame, const F
              "color_convert=%.3f depth_convert=%.3f motion_convert=%.3f exposure_convert=%.3f "
              "color_h2d=%.3f depth_h2d=%.3f motion_h2d=%.3f exposure_h2d=%.3f "
              "ngx_host=%.3f ctx_sync=%.3f gpu_eval=%.3f output_sync_wall=%.3f "
-             "output_sync_cpu=%.3f output_sync_blocking=%d d2h=%.3f output_convert=%.3f "
+             "output_sync_cpu=%.3f output_sync_blocking=%d output_sync_method=%s d2h=%.3f output_convert=%.3f "
              "h2d_total=%.3f d2h_issue=%.3f d2h_wait=%.3f "
              "prep=%.3f worker_wait=%.3f worker=%.3f finish_wait=%.3f finish=%.3f pipeline=%.3f "
-             "throughput_interval=%.3f presented=%u age=%d",
+             "throughput_interval=%.3f presented=%u age=%d motion_dilation=%u motion_dilation_wall=%.3f",
              frame, params.renderWidth, params.renderHeight, feature->outWidth, feature->outHeight,
              slot.host[0].width, slot.host[0].height, slot.host[1].width, slot.host[1].height,
              slot.host[2].width, slot.host[2].height, params.hasExposure, feature->quality, feature->flags,
@@ -4060,11 +4132,13 @@ static void finish_publish(Feature* feature, int target, uint32_t frame, const F
              timing.convert[0], timing.convert[1], timing.convert[2], timing.convert[3],
              timing.upload[0], timing.upload[1], timing.upload[2], timing.upload[3],
              timing.ngxHost, timing.ctxSync, timing.gpuEval, timing.outputSyncWall, timing.outputSyncCpu,
-             timing.outputSyncBlocking, timing.download, timing.outputConvert,
+             timing.outputSyncBlocking, timing.outputSyncBlocking ? "event_query_sleep" : "context",
+             timing.download, timing.outputConvert,
              timing.h2dTotal, timing.d2hIssue, timing.d2hWait,
              timing.prep, timing.workerWait, timing.worker, timing.finishWait, profile_ms(start, finished),
              profile_ms(timing.prepStart, finished), interval,
-             timing.presentedFrame, timing.presentedFrame != 0 ? static_cast<int>(frame - timing.presentedFrame) : -1);
+             timing.presentedFrame, timing.presentedFrame != 0 ? static_cast<int>(frame - timing.presentedFrame) : -1,
+             params.motionDilation, timing.motionDilationWall);
     }
     if (completed <= 5 || completed % 120 == 0)
         logf("frame %u done (slot %d)", frame, target);
@@ -4099,6 +4173,40 @@ static void drain_pipeline()
     g.finish.call([] { return 0; });
 }
 
+static void destroy_recorded_resources(Feature* feature);
+static void finish_shutdown();
+
+class DeferredResources final : public IUnknown
+{
+public:
+    explicit DeferredResources(Feature* feature) : feature_(feature) { ++g.resourceOwners; }
+    HRESULT STDMETHODCALLTYPE QueryInterface(REFIID iid, void** out) override
+    {
+        if (out == nullptr) return E_POINTER;
+        *out = nullptr;
+        if (iid != __uuidof(IUnknown)) return E_NOINTERFACE;
+        *out = static_cast<IUnknown*>(this);
+        AddRef();
+        return S_OK;
+    }
+    ULONG STDMETHODCALLTYPE AddRef() override { return ++refs_; }
+    ULONG STDMETHODCALLTYPE Release() override
+    {
+        const ULONG remaining = --refs_;
+        if (remaining == 0)
+            g.cleanup.post([this] {
+                destroy_recorded_resources(feature_);
+                delete this;
+                --g.resourceOwners;
+                finish_shutdown();
+            });
+        return remaining;
+    }
+private:
+    std::atomic<ULONG> refs_{1};
+    Feature* feature_;
+};
+
 // --- exported NGX D3D12 API ------------------------------------------------------
 
 #define D4R_EXPORT extern "C" __declspec(dllexport)
@@ -4284,8 +4392,16 @@ static float get_float_or(void* parameters, const char* name, float fallback)
 D4R_EXPORT NgxResult NVSDK_NGX_D3D12_CreateFeature(ID3D12GraphicsCommandList*, unsigned int featureId, void* parameters,
                                                    NgxHandle** handle)
 {
+    std::lock_guard<std::mutex> apiLock(g_featureApiMutex);
+    std::lock_guard<std::mutex> stateLock(g.mutex);
     if (!g.ngxInitialized)
         return NGX_FAIL_NOT_INITIALIZED;
+    std::call_once(g_vk.once, init_vram_interop);
+    if (g_vk.lifetime == nullptr)
+    {
+        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)
     {
         logf("NVSDK_NGX_D3D12_CreateFeature(feature=%u) unsupported", featureId);
@@ -4355,12 +4471,12 @@ D4R_EXPORT NgxResult NVSDK_NGX_D3D12_CreateFeature(ID3D12GraphicsCommandList*, u
         else
             feature->scratch = 0;
         const NgxResult created = g.ngx.createFeature(NGX_FEATURE_SUPER_SAMPLING, p, &feature->cudaHandle);
-        if (created == NGX_SUCCESS && env_uint("D4R_SHIM_BLOCKING_SYNC", 0) != 0 &&
+        if (created == NGX_SUCCESS && env_uint("D4R_SHIM_BLOCKING_SYNC", 1) != 0 &&
             g.cu.eventCreate != nullptr && g.cu.eventRecord != nullptr &&
-            g.cu.eventSynchronize != nullptr && g.cu.eventDestroy != nullptr)
+            g.cu.eventSynchronize != nullptr && g.cu.eventQuery != nullptr && g.cu.eventDestroy != nullptr)
         {
-            // All VRAM-path input waits, NGX kernels and output copies use the null stream.
-            // A blocking event yields the CPU while the GPU finishes that stream.
+            // Explicit query/sleep waits avoid HIP's spinning event-sync path.
+            // Timing is unnecessary for this completion event.
             constexpr unsigned int kBlockingSyncNoTiming = 0x1u | 0x2u;
             const int eventResult = g.cu.eventCreate(&feature->outputReadyEvent, kBlockingSyncNoTiming);
             if (eventResult != 0)
@@ -4369,7 +4485,7 @@ D4R_EXPORT NgxResult NVSDK_NGX_D3D12_CreateFeature(ID3D12GraphicsCommandList*, u
                 feature->outputReadyEvent = nullptr;
             }
             else
-                logf("blocking output wait enabled for this feature");
+                logf("output event query/sleep wait enabled for this feature");
         }
         if (created == NGX_SUCCESS && profile_enabled() && g.cu.eventCreate != nullptr &&
             g.cu.eventRecord != nullptr && g.cu.eventElapsedTime != nullptr && g.cu.eventDestroy != nullptr)
@@ -4402,6 +4518,7 @@ D4R_EXPORT NgxResult NVSDK_NGX_D3D12_CreateFeature(ID3D12GraphicsCommandList*, u
         return NGX_FAIL_PLATFORM_ERROR;
     }
     *feature->markerValue = 0;
+    feature->resources = new DeferredResources(feature);
     feature->handle.Id = g.nextHandleId++;
     {
         std::lock_guard<std::mutex> lock(g_featuresMutex);
@@ -4415,7 +4532,7 @@ static Feature* find_feature(const NgxHandle* handle)
 {
     std::lock_guard<std::mutex> lock(g_featuresMutex);
     for (Feature* feature : g_features)
-        if (&feature->handle == handle)
+        if (&feature->handle == handle && !feature->retiring.load())
             return feature;
     return nullptr;
 }
@@ -4433,6 +4550,7 @@ static ID3D12Resource* get_resource(void* parameters, const char* name)
 
 static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* handle, void* parameters)
 {
+    std::lock_guard<std::mutex> apiLock(g_featureApiMutex);
     FrameTiming timing;
     timing.enabled = profile_enabled();
     if (timing.enabled)
@@ -4543,11 +4661,18 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand
             logf("frame %u: input slot %d busy or %d frames in flight; skipping evaluation", frame, slotIndex, feature->inFlight.load());
             return NGX_SUCCESS;
         }
-        std::this_thread::sleep_for(std::chrono::microseconds(200));
+        d4r_sleep_us(200);
     }
     if (timing.enabled)
         timing.slotWait = profile_ms(waitStart, ProfileClock::now());
 
+    if (g_vk.lifetime == nullptr ||
+        FAILED(g_vk.lifetime->RetainExternalResources(list, feature->resources)))
+    {
+        logf("frame %u: unable to retain command-list resources; evaluation not recorded", frame);
+        return NGX_FAIL_PLATFORM_ERROR;
+    }
+
     // VRAM interop when every resource qualifies (decided once per feature).
     ID3D12Resource* const inputs[4] = {color, depth, motion, exposure};
     const Plane planes[4] = {Plane::Color, Plane::Depth, Plane::Motion, Plane::Exposure};
@@ -4565,17 +4690,13 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand
         feature->vram = p.vram;
         feature->split = p.vram && env_uint("D4R_SHIM_SPLIT_FRAME", 0) != 0 && g_vk.split != nullptr &&
                          create_split_semaphore(*feature);
-        if (feature->split && env_uint("D4R_SHIM_GPU_WAIT", 0) != 0 && g.cu.streamWaitValue32 != nullptr &&
-            g.cu.writeValue32 != nullptr && g_vk.fill != nullptr)
-            feature->gpuWait = ensure_vram_buffer(*feature, feature->gpuMarker, 256) &&
-                               g.worker.call([feature] { return g.cu.writeValue32(feature->gpuMarker.device, 0); }) == 0;
-        feature->linearInputs = feature->gpuWait && env_uint("D4R_SHIM_LINEAR_INPUTS", 0) != 0 &&
+        feature->linearInputs = feature->split && env_uint("D4R_SHIM_LINEAR_INPUTS", 0) != 0 &&
                                 g.cu.registerLinearTexture != nullptr;
         if (feature->linearInputs)
             logf("linear inputs: NGX samples the interop buffers directly");
         logf("VRAM interop %s for this feature%s%s", p.vram ? "on" : "off",
              feature->split ? ", presenting each frame's own result (split frames)" : "",
-             feature->gpuWait ? ", DLSS queued behind a GPU-side wait for the inputs" : "");
+             p.vram ? ", waiting for input readiness with CPU sleeps" : "");
     }
     else if (feature->vram && !p.vram)
     {
@@ -4586,7 +4707,6 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand
             feature->split = false;
         }
         feature->vram = false;
-        feature->gpuWait = false;
         feature->linearInputs = false;
         feature->latestOutput = -1;
         feature->outputFormat = DXGI_FORMAT_UNKNOWN; // recreates the output staging below
@@ -4640,7 +4760,7 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand
         feature->outputHeight = outputDesc.Height;
         for (OutputSlot& outputSlot : feature->outputs)
         {
-            outputSlot.staging.release();
+            outputSlot.staging.retire();
             if (p.vram ? !ensure_vram_buffer(*feature, outputSlot.vram, static_cast<size_t>(outputDesc.Width) * 8 *
                                                                             outputDesc.Height)
                        : !ensure_staging(outputSlot.staging, output, D3D12_HEAP_TYPE_UPLOAD))
@@ -4657,7 +4777,7 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand
         const D3D12_RESOURCE_STATES states[4] = {inputState, depthState, inputState, inputState};
         for (int index = 0; index < inputCount; ++index)
             transition(list, inputs[index], states[index], D3D12_RESOURCE_STATE_COPY_SOURCE);
-        const bool recorded = record_vram_inputs(*feature, list, slot, vramInputs, inputCount, frame);
+        const bool recorded = record_vram_inputs(*feature, list, slot, vramInputs, inputCount);
         for (int index = 0; index < inputCount; ++index)
             transition(list, inputs[index], D3D12_RESOURCE_STATE_COPY_SOURCE, states[index]);
         if (!recorded)
@@ -4779,9 +4899,15 @@ static void release_feature(Feature* feature)
         if (feature->cudaParams != nullptr)
             g.ngx.destroyParameters(feature->cudaParams);
         for (InputSlot& slot : feature->inputs)
+        {
             for (CudaObject& texture : slot.linearTexture)
                 if (texture != 0)
                     g.cu.texObjectDestroy(texture), texture = 0;
+            if (slot.dilatedTexture != 0)
+                g.cu.texObjectDestroy(slot.dilatedTexture), slot.dilatedTexture = 0;
+            if (slot.dilatedMotion != 0)
+                g.cu.memFree(slot.dilatedMotion), slot.dilatedMotion = 0;
+        }
         for (CudaImage* image : {&feature->color, &feature->depth, &feature->motion, &feature->exposure})
         {
             if (image->object != 0)
@@ -4813,38 +4939,40 @@ static void release_feature(Feature* feature)
             release_host(host);
         return 0;
     });
+    {
+        std::lock_guard<std::mutex> lock(g_featuresMutex);
+        g_features.erase(std::remove(g_features.begin(), g_features.end(), feature), g_features.end());
+    }
+    // Recorded lists retain this owner through their allocator. Reset/discard
+    // or GPU completion followed by allocator destruction releases those refs.
+    feature->resources->Release();
+}
+
+static void destroy_recorded_resources(Feature* feature)
+{
     for (InputSlot& slot : feature->inputs)
         for (Staging* staging : {&slot.color, &slot.depth, &slot.motion, &slot.exposure})
             staging->release();
     for (OutputSlot& slot : feature->outputs)
         slot.staging.release();
     if (g_vk.ready && feature->vramDecided)
-    {
-        // Queued command lists may still copy from or wait on these; the last
-        // marker lands just before the last output copy.
-        const auto start = std::chrono::steady_clock::now();
-        while (feature->markerValue != nullptr &&
-               static_cast<int32_t>(*feature->markerValue - feature->frame) < 0 &&
-               std::chrono::steady_clock::now() - start < std::chrono::seconds(2))
-            std::this_thread::sleep_for(std::chrono::milliseconds(1));
-        std::this_thread::sleep_for(std::chrono::milliseconds(50));
         release_vram(*feature);
-    }
     if (feature->marker != nullptr)
         feature->marker->Release();
+    logf("feature %u recorded resources retired after allocator completion", feature->handle.Id);
     delete feature;
 }
 
 D4R_EXPORT NgxResult NVSDK_NGX_D3D12_ReleaseFeature(NgxHandle* handle)
 {
+    std::lock_guard<std::mutex> apiLock(g_featureApiMutex);
     Feature* feature = nullptr;
     {
         std::lock_guard<std::mutex> lock(g_featuresMutex);
-        for (auto it = g_features.begin(); it != g_features.end(); ++it)
-            if (&(*it)->handle == handle)
+        for (Feature* candidate : g_features)
+            if (&candidate->handle == handle && !candidate->retiring.exchange(true))
             {
-                feature = *it;
-                g_features.erase(it);
+                feature = candidate;
                 break;
             }
     }
@@ -4856,30 +4984,51 @@ D4R_EXPORT NgxResult NVSDK_NGX_D3D12_ReleaseFeature(NgxHandle* handle)
     return NGX_SUCCESS;
 }
 
-D4R_EXPORT NgxResult NVSDK_NGX_D3D12_Shutdown()
+static void finish_shutdown()
 {
-    logf("NVSDK_NGX_D3D12_Shutdown");
-    std::vector<Feature*> features;
-    {
-        std::lock_guard<std::mutex> lock(g_featuresMutex);
-        features.swap(g_features);
-    }
-    for (Feature* feature : features)
-        release_feature(feature);
+    std::lock_guard<std::mutex> lock(g.mutex);
+    if (!g.shutdownPending || g.resourceOwners.load() != 0)
+        return;
     if (g_vk.ready)
         free_vram_pool();
-    std::lock_guard<std::mutex> lock(g.mutex);
     if (g.ngxInitialized)
     {
         drain_pipeline();
         g.worker.call([] { return g.ngx.shutdown(); });
     }
     g.ngxInitialized = false;
+    g.shutdownPending = false;
     if (g.device != nullptr)
     {
         g.device->Release();
         g.device = nullptr;
     }
+    logf("NGX shutdown completed after recorded resources retired");
+}
+
+D4R_EXPORT NgxResult NVSDK_NGX_D3D12_Shutdown()
+{
+    std::lock_guard<std::mutex> apiLock(g_featureApiMutex);
+    logf("NVSDK_NGX_D3D12_Shutdown");
+    std::vector<Feature*> features;
+    {
+        std::lock_guard<std::mutex> lock(g_featuresMutex);
+        features = g_features;
+        for (Feature* feature : features)
+            feature->retiring = true;
+    }
+    for (Feature* feature : features)
+        release_feature(feature);
+    {
+        std::lock_guard<std::mutex> lock(g.mutex);
+        if (!g.started)
+            return NGX_SUCCESS;
+        g.shutdownPending = true;
+    }
+    // Drain already queued cleanup before returning. If command allocators
+    // still own resources, finish_shutdown leaves NGX alive until their last
+    // reference retires; otherwise shutdown completes synchronously.
+    g.cleanup.call([] { finish_shutdown(); return 0; });
     return NGX_SUCCESS;
 }
 
diff --git a/tools/d4r_vkd3d_interop.h b/tools/d4r_vkd3d_interop.h
new file mode 100644
index 0000000..e636b78
--- /dev/null
+++ b/tools/d4r_vkd3d_interop.h
@@ -0,0 +1,59 @@
+#pragma once
+#include <d3d12.h>
+#include <vulkan/vulkan_core.h>
+
+MIDL_INTERFACE("39da4e09-bd1c-4198-9fae-86bbe3be41fd")
+ID3D12DXVKInteropDevice : public IUnknown
+{
+    virtual HRESULT STDMETHODCALLTYPE GetDXGIAdapter(REFIID iid, void** object) = 0;
+    virtual HRESULT STDMETHODCALLTYPE GetInstanceExtensions(UINT* count, const char** extensions) = 0;
+    virtual HRESULT STDMETHODCALLTYPE GetDeviceExtensions(UINT* count, const char** extensions) = 0;
+    virtual HRESULT STDMETHODCALLTYPE GetDeviceFeatures(const VkPhysicalDeviceFeatures2** features) = 0;
+    virtual HRESULT STDMETHODCALLTYPE GetVulkanHandles(VkInstance* instance, VkPhysicalDevice* physical,
+                                                       VkDevice* device) = 0;
+    virtual HRESULT STDMETHODCALLTYPE GetVulkanQueueInfo(ID3D12CommandQueue* queue, VkQueue* vkQueue,
+                                                         UINT32* family) = 0;
+    virtual void STDMETHODCALLTYPE GetVulkanImageLayout(ID3D12Resource* resource, D3D12_RESOURCE_STATES state,
+                                                        VkImageLayout* layout) = 0;
+    virtual HRESULT STDMETHODCALLTYPE GetVulkanResourceInfo(ID3D12Resource* resource, UINT64* handle,
+                                                            UINT64* offset) = 0;
+    virtual HRESULT STDMETHODCALLTYPE LockCommandQueue(ID3D12CommandQueue* queue) = 0;
+    virtual HRESULT STDMETHODCALLTYPE UnlockCommandQueue(ID3D12CommandQueue* queue) = 0;
+};
+
+MIDL_INTERFACE("902d8115-59eb-4406-9518-fe00f991ee65")
+ID3D12DXVKInteropDevice1 : public ID3D12DXVKInteropDevice
+{
+    virtual HRESULT STDMETHODCALLTYPE GetVulkanResourceInfo1(ID3D12Resource* resource, UINT64* handle, UINT64* offset,
+                                                             VkFormat* format) = 0;
+    virtual HRESULT STDMETHODCALLTYPE CreateInteropCommandQueue(const D3D12_COMMAND_QUEUE_DESC* desc, UINT32 family,
+                                                                ID3D12CommandQueue** queue) = 0;
+    virtual HRESULT STDMETHODCALLTYPE CreateInteropCommandAllocator(D3D12_COMMAND_LIST_TYPE type, UINT32 family,
+                                                                    ID3D12CommandAllocator** allocator) = 0;
+    virtual HRESULT STDMETHODCALLTYPE BeginVkCommandBufferInterop(ID3D12CommandList* list, VkCommandBuffer* buffer) = 0;
+    virtual HRESULT STDMETHODCALLTYPE EndVkCommandBufferInterop(ID3D12CommandList* list) = 0;
+};
+__CRT_UUID_DECL(ID3D12DXVKInteropDevice1, 0x902d8115, 0x59eb, 0x4406, 0x95, 0x18, 0xfe, 0x00, 0xf9, 0x91, 0xee, 0x65)
+
+// Added by the d4r vkd3d-proton patch (patches/vkd3d-proton/): ends the list's
+// current Vulkan command buffer; what is recorded afterwards is submitted
+// separately, with a queue-level wait for the timeline semaphore value.
+MIDL_INTERFACE("5a7c8b3e-2f61-4d0e-9c1a-7e3b52d4a901")
+ID3D12DXVKInteropDeviceD4R : public ID3D12DXVKInteropDevice1
+{
+    virtual HRESULT STDMETHODCALLTYPE LockVulkanQueue(ID3D12CommandQueue* queue) = 0;
+    virtual HRESULT STDMETHODCALLTYPE UnlockVulkanQueue(ID3D12CommandQueue* queue) = 0;
+    virtual HRESULT STDMETHODCALLTYPE GetVulkanHeapInfo(ID3D12Heap* heap, UINT64* memory, UINT64* offset,
+                                                        UINT32* type) = 0;
+    virtual HRESULT STDMETHODCALLTYPE SplitCommandListForExternalWait(ID3D12CommandList* list, UINT64 semaphore,
+                                                                      UINT64 value) = 0;
+};
+__CRT_UUID_DECL(ID3D12DXVKInteropDeviceD4R, 0x5a7c8b3e, 0x2f61, 0x4d0e, 0x9c, 0x1a, 0x7e, 0x3b, 0x52, 0xd4, 0xa9, 0x01)
+
+MIDL_INTERFACE("4f8c9462-8197-4dc1-a584-cf44d9706db2")
+ID3D12DXVKInteropDeviceD4R2 : public ID3D12DXVKInteropDeviceD4R
+{
+    virtual HRESULT STDMETHODCALLTYPE RetainExternalResources(ID3D12CommandList* list, IUnknown* resources) = 0;
+};
+__CRT_UUID_DECL(ID3D12DXVKInteropDeviceD4R2, 0x4f8c9462, 0x8197, 0x4dc1, 0xa5, 0x84, 0xcf, 0x44, 0xd9, 0x70, 0x6d, 0xb2)
+
diff --git a/tools/d4r_win32_wait.h b/tools/d4r_win32_wait.h
new file mode 100644
index 0000000..56a28d3
--- /dev/null
+++ b/tools/d4r_win32_wait.h
@@ -0,0 +1,23 @@
+#pragma once
+#include <windows.h>
+#include <cstdint>
+
+// MinGW's Win32 std::sleep_for can truncate submillisecond durations to zero.
+// NtDelayExecution retains 100 ns precision and actually sleeps under Wine.
+inline void d4r_sleep_us(unsigned int microseconds)
+{
+    using Delay = LONG(WINAPI*)(BOOLEAN, LARGE_INTEGER*);
+    static const auto delay = reinterpret_cast<Delay>(reinterpret_cast<void*>(
+        GetProcAddress(GetModuleHandleW(L"ntdll.dll"), "NtDelayExecution")));
+    if (microseconds == 0)
+    {
+        SwitchToThread();
+        return;
+    }
+    LARGE_INTEGER interval;
+    interval.QuadPart = -static_cast<LONGLONG>(microseconds) * 10;
+    if (delay != nullptr && delay(FALSE, &interval) >= 0)
+        return;
+    // A rounded-up millisecond wait still releases the CPU if ntdll is absent.
+    Sleep(static_cast<DWORD>((static_cast<uint64_t>(microseconds) + 999) / 1000));
+}
diff --git a/tools/wine_nvcuda_bridge.c b/tools/wine_nvcuda_bridge.c
index 30a9c87..11247fa 100644
--- a/tools/wine_nvcuda_bridge.c
+++ b/tools/wine_nvcuda_bridge.c
@@ -12,6 +12,10 @@
 #include <sys/stat.h>
 #include <unistd.h>
 #include "d4r_native_selection.h"
+#include "d4r_motion_dilation.h"
+#ifdef D4R_MOTION_KERNELS
+#include "d4r_motion_kernels.h"
+#endif
 
 typedef int CUresult;
 typedef int CUdevice;
@@ -457,6 +461,15 @@ static void prepare_native_kernels(const char* cache_home)
     }
     if (kind == 2)
     {
+        /* A flat developer set has no per-kernel manifest to filter. Fail safely to
+           translated kernels rather than silently retaining the disabled encoders. */
+        if (!d4r_native_swin_encoders_enabled())
+        {
+            tracef("NativeSwinEncoders off: developer set needs d4r-kernels.txt for selective fallback; "
+                   "using translated kernels for this set");
+            unsetenv("D4R_ZLUDA_NATIVE_DIR");
+            return;
+        }
         setenv("D4R_ZLUDA_NATIVE_DIR", source, 1);
         tracef("native kernels: %sdeveloper set in %s", accuracy ? "accuracy " : "", source);
         return;
@@ -476,6 +489,8 @@ static void prepare_native_kernels(const char* cache_home)
         unsigned long long hash = 0;
         if (line[0] == '#' || sscanf(line, "%127s %llx", name, &hash) != 2)
             continue;
+        if (!d4r_native_kernel_allowed(name))
+            continue;
         if (native_kernel_count == capacity)
         {
             capacity = capacity != 0 ? capacity * 2 : 32;
@@ -490,6 +505,9 @@ static void prepare_native_kernels(const char* cache_home)
         kernel->linked = 0;
     }
     fclose(manifest);
+    if (!d4r_native_swin_encoders_enabled())
+        tracef("NativeSwinEncoders off: rrlite_enc1_4x4 and rrlite_enc2_4x4 use original translated PTX; "
+               "other native kernels retained");
 
     /* <cache>/d4r-native/<pid>; directories of processes that no longer exist are removed */
     char base[1024];
@@ -843,6 +861,57 @@ static void* find_zluda_symbol(const char* name)
     return cuda_library != NULL ? dlsym(cuda_library, name) : NULL;
 }
 
+/* d4r-owned, opt-in preprocessing. Runs on HIP's legacy/default stream, shared
+   with CUDA's null stream; the shim records and waits for completion before NGX. */
+CUresult WINAPI d4rDilateMotionVectors(const D4rMotionDilationParams* parameters)
+{
+    if (parameters == NULL || parameters->depth == 0 || parameters->motion == 0 || parameters->output == 0 ||
+        parameters->radius < 1 || parameters->radius > 2 || parameters->render_width == 0 ||
+        parameters->render_height == 0 || parameters->motion_width == 0 || parameters->motion_height == 0 ||
+        parameters->depth_pitch < (uint64_t)parameters->render_width * 4 ||
+        parameters->motion_pitch < (uint64_t)parameters->motion_width * 4 ||
+        (uint64_t)parameters->motion_width * parameters->motion_height > 0x7fffffffull)
+        return CUDA_ERROR_INVALID_VALUE;
+#ifdef D4R_MOTION_KERNELS
+    typedef int(__attribute__((sysv_abi)) * HIP_LOAD)(void**, const void*);
+    typedef int(__attribute__((sysv_abi)) * HIP_FUNCTION)(void**, void*, const char*);
+    typedef int(__attribute__((sysv_abi)) * HIP_LAUNCH)(void*, unsigned, unsigned, unsigned,
+        unsigned, unsigned, unsigned, unsigned, void*, void**, void**);
+    static pthread_mutex_t lock = PTHREAD_MUTEX_INITIALIZER;
+    static void* module;
+    static void* function;
+    static HIP_LAUNCH launch;
+    static int initialized, result = CUDA_ERROR_NOT_SUPPORTED;
+    pthread_mutex_lock(&lock);
+    if (!initialized)
+    {
+        initialized = 1;
+        pthread_once(&load_once, load_zluda);
+        void* hip = dlopen("libamdhip64.so.7", RTLD_NOW | RTLD_LOCAL);
+        char architecture[32];
+        gpu_architecture(architecture, sizeof(architecture));
+        const void* image = d4r_motion_image(architecture);
+        HIP_LOAD load = hip != NULL ? (HIP_LOAD)dlsym(hip, "hipModuleLoadData") : NULL;
+        HIP_FUNCTION get = hip != NULL ? (HIP_FUNCTION)dlsym(hip, "hipModuleGetFunction") : NULL;
+        launch = hip != NULL ? (HIP_LAUNCH)dlsym(hip, "hipModuleLaunchKernel") : NULL;
+        if (image != NULL && load != NULL && get != NULL && launch != NULL &&
+            load(&module, image) == 0 && get(&function, module, "d4r_motion_dilate") == 0)
+            result = CUDA_SUCCESS;
+        tracef("motion-vector dilation: %s (%s)", result == 0 ? "GPU kernel ready" : "unavailable", architecture);
+    }
+    const int status = result;
+    pthread_mutex_unlock(&lock);
+    if (status != CUDA_SUCCESS)
+        return status;
+    D4rMotionDilationParams copy = *parameters;
+    void* arguments[] = {&copy};
+    const uint64_t pixels = (uint64_t)copy.motion_width * copy.motion_height;
+    return launch(function, (unsigned)((pixels + 127) / 128), 1, 1, 128, 1, 1, 0, NULL, arguments, NULL);
+#else
+    return CUDA_ERROR_NOT_SUPPORTED;
+#endif
+}
+
 static void tracef(const char* format, ...)
 {
     char line[2048];
@@ -2579,7 +2648,17 @@ CUresult WINAPI d4rEventSynchronize(D4rArg a0)
     TRACE_CALL(result, "d4rEventSynchronize result=%d", result);
     return result;
 }
-D4R_FORWARD(cuEventQuery, cuEventQuery, 1, (D4rArg a0), (a0))
+CUresult WINAPI cuEventQuery(D4rArg event)
+{
+    typedef CUresult(__attribute__((sysv_abi)) *function_type)(D4rArg);
+    function_type function = (function_type)find_zluda_symbol("cuEventQuery");
+    CUresult result = function != NULL ? function(event) : missing("cuEventQuery");
+    /* NOT_READY is the normal result while the shim sleeps between queries,
+       not an error to flush to the log on every poll. */
+    TRACE_CALL(result == CUDA_ERROR_NOT_READY ? CUDA_SUCCESS : result,
+               "cuEventQuery result=%d", result);
+    return result;
+}
 /* With elided event waits an event may still be pending: report 0 ms instead of CUDA_ERROR_NOT_READY. */
 CUresult WINAPI cuEventElapsedTime(D4rArg a0, D4rArg a1, D4rArg a2)
 {
diff --git a/tools/wine_nvcuda_bridge.spec b/tools/wine_nvcuda_bridge.spec
index 3673e25..5e92457 100644
--- a/tools/wine_nvcuda_bridge.spec
+++ b/tools/wine_nvcuda_bridge.spec
@@ -100,3 +100,4 @@
 @ stdcall d4rSetEnv(str str long) d4rSetEnv
 @ stdcall d4rLoadError() d4rLoadError
 @ stdcall d4rOutputKernelNative() d4rOutputKernelNative
+@ stdcall d4rDilateMotionVectors(ptr) d4rDilateMotionVectors