Queue native Windows DLSS asynchronously and drain feature consumers

docs/windows-game.md

@@ -61,11 +61,23 @@ The recording interval measures NGX calls on the recording thread, not Present.
 frame scheduling; use it separately from a production FPS measurement.
 `-ProfileCommandHooks` samples one in 64 D3D12 interceptions to distinguish
 tracking/lock waits from driver calls. It is disabled in ordinary runs.
+`-ProfileCudaApi` aggregates existing CUDA API host durations without adding
+HIP events or synchronization. `-UncachedInteropLists` restores per-frame
+copy-list allocation for comparison with the normal safely reused lists.
 The runner drains stdout/stderr concurrently as raw bytes; the earlier
 line-oriented PowerShell collector imposed large per-frame delays. The optional
 ZLUDA kernel profiler is for diagnostics only. Measurements and reproduction
 commands are in [windows-performance.md](windows-performance.md).
 
+`-AsyncInterop` enables the experimental K scheduling path. Submission queues
+the current frame's input/output dependencies and returns while a CUDA worker
+executes DLSS. Owned command memory is retained until GPU completion; feature
+release and resource reconfiguration wait for pending consumers. It passes
+the standalone queued-frame and Release/CreateFeature exact-output regression,
+but remains opt-in while game stability and performance are checked. It
+currently requires one D3D12 command queue. This switch is independent of
+profiling and output scans; omit those diagnostics for a production FPS run.
+
 Ordinary runs use OptiScaler info logging and suppress successful HIP/CUDA API
 messages. Errors, GPU identity, imports, frame records and native/translated
 kernel launches remain logged. `-VerboseRuntime` enables all API and frontend

docs/windows-performance.md

@@ -72,6 +72,64 @@ and compiler spill counts. Reports and proprietary workload inputs remain local.
 Native FP8 remains a separate optional experiment; the package uses the
 validated FP16-equivalent baseline. Its later checks and timings are below.
 
+## K host synchronization and asynchronous submission
+
+Current priority is K performance. The user's synchronous 4K reading is
+49–51 FPS / 53% GPU utilization, versus approximately 80+ FPS / 100% with
+FSR4. These readings do not identify a CPU hardware limit or provide a
+controlled same-scene comparison. M/FP8 optimization is deferred.
+
+ZLUDA candidate `a1c506f` adds `D4R_ZLUDA_PROFILE_API`, enabled by the game
+runner's `-ProfileCudaApi`. It measures existing API host durations, aggregates
+per thread and logs every two seconds without HIP events or additional waits.
+The synchronous control shows five `cuCtxSynchronize` calls per frame with
+about 3.9 ms aggregate host time. CUDA launch host overhead is approximately
+0.46 ms per frame. Cached copy-list setup is only 0.011 ms median. These host
+measurements include GPU waits and must not be added to GPU kernel times.
+
+| K/4K command submission, CPU median | Synchronous | First async | Async without CUDA mutex |
+| --- | ---: | ---: | ---: |
+| Entire submission, ms | 10.512 | 2.676 | 0.181 |
+| DLSS boundary, ms | 10.083 | 2.604 | 0.124 |
+| Consumer suffix completion, ms | 0.350 | 0.005 | 0.004 |
+
+The first async game run completes 9481 finite frames; the user reports
+60 FPS / 63% GPU. The mutex-free run crashes during feature recreation after
+366 finite frames, so its submission timing is not a successful gameplay
+benchmark. The standalone burst/recreate reproducer subsequently identifies
+caller-owned GPU backing freed after ReleaseFeature returns. Waiting for
+pending D3D12 consumers in ReleaseFeature fixes the reproducer: all twelve
+sync/async test frames pass, and corresponding RGB outputs are exact. Game
+stability and FPS must be checked separately on the fixed build.
+
+The fixed game retest on 2026-10-02 completes 16941 K/4K frames, all finite,
+including feature recreation and color-format reimports, with zero backend
+failures. The user reports 62 FPS / 63% GPU. Median CPU submission is
+0.197 ms, NGX host completion on the worker is 4.400 ms and output validation
+is 0.380 ms. The 10.162 ms median worker input wait includes queued game
+rendering. These results establish a stable five-minute test of this fix,
+not a conclusion about the remaining GPU idle time or FSR4 performance.
+
+`-AsyncInterop` is opt-in. It queues input/output fences on the submission
+thread, executes CUDA on a worker and retires owned command memory by GPU
+completion. Three copy slots bound outstanding ownership. The queue consumes
+the output of the same frame, with no CPU image copy. The worker's input wait
+includes its lead behind submitted rendering and is not a CPU submission
+stall. Neither that duration nor the NGX recording interval measures Present
+latency. Reconfiguration waits for all consumers; ordinary submission does not
+hold the CUDA mutex. The experimental runtime currently supports one queue.
+
+```powershell
+.\scripts\windows\test-async-k.ps1
+```
+
+The test generates fresh synchronous controls, submits three frames without
+intermediate CPU completion, and repeats while releasing/recreating the DLSS
+feature with work pending. Results include exact RGB comparisons, finite GPU
+scans, stdout/stderr, exceptions and runtime provenance. To include the local
+validated output-store optimization, pass its combined native directory with
+`-NativeRoot`; proprietary objects remain local.
+
 Further game coverage (HIP kernel profiling disabled, GPU output checks enabled):
 M at 1920x1080 completes 4833 frames, all finite. K at 3840x2160 completes
 4685 frames, all finite. Neither has a backend failure, CPU image copy or an

docs/windows-rdna4-port.md

@@ -10,6 +10,87 @@ explicit limits; the validated package keeps conservative arithmetic.
 
 ## Milestone and gates
 
+**Current milestone: K performance remains open.** The user reports 49–51 FPS
+and about 53% GPU utilization in 4K Silent Hill 2, compared with roughly 80+
+FPS / 100% GPU utilization with FSR4 on this system. Functional K/M and the
+reproducible archive do not close this gate. Prioritize K queue scheduling and
+host synchronization; postpone further M/FP8 optimization. These user readings
+are not a controlled same-scene benchmark. Do not infer a hardware CPU limit
+from utilization alone.
+
+The next diagnostics add `-ProfileCudaApi` (ZLUDA patch 0015): aggregate host
+CUDA API durations without HIP-event waits. Boundary stages now separate
+prepare, copy-list setup, input/output submission and output-copy drain.
+Feature-owned input/output command lists and allocators are reused only after
+the existing completion fence; `-UncachedInteropLists` restores per-frame
+allocation for an A/B control. Submission drain fences/events are queue-owned
+and reused. No completion dependency or current-frame output is removed.
+
+The opt-in `-AsyncInterop` experiment reserves D3D12 input/output dependencies
+at submission, runs CUDA on a feature worker, and retires owned suffix command
+memory on a separate fence-completion service. Worker order is shared across
+features using the context; this experimental runtime requires one D3D12 queue.
+Three reusable copy contexts bound outstanding ownership. Every consumer waits
+for **its own** output; no previous-image selection or CPU image staging occurs.
+Steady-state submission uses immutable allocation metadata and per-frame
+resource/parameter snapshots, avoiding the CUDA mutex. Resize/reimport and
+feature creation/release drain D3D12 outside that mutex, keeping workers able to signal
+already-published waits. Worker failures poison the context and are logged.
+
+First async game run: 9481 K/4K frames and finite GPU scans, zero backend failures;
+the user reports **60 FPS / 63% GPU**. Median command submission drops from
+10.512 ms (synchronous control) to 2.676 ms. The next mutex-free candidate
+submits in median 0.181 ms, but its initial game test exits with UE
+`DXGI_ERROR_DEVICE_HUNG` during resource/feature recreation after 366 finite
+frames. Its crash context and dump are in
+`test-results/silent-hill2-k-async-nolock-4k/ue-crash`. Do not classify this as
+success or enable the experiment by default until the game reconfiguration gate
+passes. The new standalone burst + Release/CreateFeature mode reproduces the
+GPU hang without the game. Retaining the d4r Feature through completion does
+not retain OptiScaler's caller-owned root-CBV buffers: OptiScaler frees its
+shader backing immediately after NGX ReleaseFeature returns. ReleaseFeature
+now waits for pending workers, copy contexts and all submitted D3D12 consumers
+before returning, outside the CUDA mutex. The fixed four-way regression
+(sync/async x burst/recreate) completes twelve K/4K frames; all six async RGB
+outputs exactly match their corresponding fresh synchronous controls, with
+no NaN/Inf or GPU hang. Results:
+`test-results/k-async-release-drain-4k/validation.json`. The fixed game retest
+on 2026-10-02 completes **16941 K/4K frames**, 16941 finite GPU scans and
+203292 native launches, with zero backend failures/CPU image copies/previous
+frame outputs. It includes two feature creations and color-format reimports;
+the previous recreation crash does not recur in this five-minute run. The
+user reports **62 FPS / 63% GPU** in the comparison scene. Median command
+submission is 0.197 ms; DLSS worker NGX completion is 4.400 ms and output
+scan is 0.380 ms. Logs:
+`test-results/silent-hill2-k-async-release-fix-4k`. The runner deliberately
+stops at its bound (`diagnostic_timeout`); there is no recorded game crash.
+The K performance gate remains open.
+
+Three ordinary K/4K async frames and three queued together without intermediate
+CPU completion all match their corresponding prior synchronous RGB exactly;
+the temporal controls differ between frames. One burst attempt times out during
+NGX creation before submitting a frame after the game GPU crash; HIP/D3D12
+sanity probes pass and the 600-second retry completes and matches all frames.
+`test-async-k.ps1` generates fresh synchronous controls and checks burst and
+Release/CreateFeature with pending work, rather than depending on saved fixtures:
+
+```powershell
+.\scripts\windows\test-async-k.ps1
+```
+
+New game crashes are automatically copied from the current Silent Hill 2 UE
+crash directory into the diagnostic bundle, including its error reason and
+minidump; proprietary/private artifacts remain excluded from public packaging.
+
+The profiling candidate is ZLUDA
+`a1c506fc956cfc347c56c44311669700c0d500ac` in
+`dist/zluda-windows-api-profile`. Patch 0015 adds optional per-thread host API
+aggregates without extra synchronization; ordinary runs leave it disabled.
+All fourteen patches (0002 through 0015) apply to pinned base `ee2f25a`.
+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
 [`f0d1a65`](https://github.com/countervolts/d4r/commit/f0d1a65e27aff6cbe7eeaa13227c80facb9056ab)
 and artifact fix `65dfc9f`. Windows retains strict f16 rounding and the exact

patches/zluda/0015-nonserializing-cuda-api-profile.patch

@@ -0,0 +1,83 @@
+From a1c506fc956cfc347c56c44311669700c0d500ac Mon Sep 17 00:00:00 2001
+From: Codex <codex@localhost>
+Date: Thu, 1 Oct 2026 22:16:19 +0300
+Subject: [PATCH] Profile CUDA host API duration without serializing GPU
+ launches
+
+---
+ zluda/src/d4r_api_profile.rs | 34 ++++++++++++++++++++++++++++++++++
+ zluda/src/lib.rs             |  3 +++
+ 2 files changed, 37 insertions(+)
+ create mode 100644 zluda/src/d4r_api_profile.rs
+
+diff --git a/zluda/src/d4r_api_profile.rs b/zluda/src/d4r_api_profile.rs
+new file mode 100644
+index 0000000..ef2da1b
+--- /dev/null
++++ b/zluda/src/d4r_api_profile.rs
+@@ -0,0 +1,34 @@
++//! Host API timings only: no HIP events, stream waits, or altered scheduling.
++use std::{cell::RefCell, collections::BTreeMap, sync::OnceLock, time::Instant};
++
++#[derive(Default)]
++struct Row { calls: u64, total_ms: f64, max_ms: f64 }
++struct Samples { since: Instant, rows: BTreeMap<&'static str, Row> }
++thread_local! {
++    static SAMPLES: RefCell<Samples> = RefCell::new(Samples { since: Instant::now(), rows: BTreeMap::new() });
++}
++pub(crate) struct Call { name: &'static str, start: Instant }
++impl Call {
++    pub(crate) fn begin(name: &'static str) -> Option<Self> {
++        static ENABLED: OnceLock<bool> = OnceLock::new();
++        if !*ENABLED.get_or_init(|| std::env::var_os("D4R_ZLUDA_PROFILE_API").is_some()) { return None; }
++        Some(Self { name, start: Instant::now() })
++    }
++}
++impl Drop for Call {
++    fn drop(&mut self) {
++        let ms = self.start.elapsed().as_secs_f64() * 1000.;
++        SAMPLES.with(|samples| {
++            let mut samples = samples.borrow_mut();
++            let row = samples.rows.entry(self.name).or_default();
++            row.calls += 1; row.total_ms += ms; row.max_ms = row.max_ms.max(ms);
++            let period_ms = samples.since.elapsed().as_secs_f64() * 1000.;
++            if period_ms < 2000. { return; }
++            for (name, row) in &samples.rows {
++                eprintln!("D4R_CUDA_API_PROFILE thread={:?} api={name} calls={} total_ms={:.6} max_ms={:.6} period_ms={period_ms:.6} serializing=0",
++                    std::thread::current().id(), row.calls, row.total_ms, row.max_ms);
++            }
++            samples.rows.clear(); samples.since = Instant::now();
++        });
++    }
++}
+diff --git a/zluda/src/lib.rs b/zluda/src/lib.rs
+index ecdeb3e..31bfa42 100644
+--- a/zluda/src/lib.rs
++++ b/zluda/src/lib.rs
+@@ -2,6 +2,7 @@ use cuda_types::cuda::CUerror;
+ use std::sync::atomic::{AtomicBool, Ordering};
+ 
+ pub(crate) mod r#impl;
++mod d4r_api_profile;
+ #[cfg_attr(windows, path = "os_win.rs")]
+ #[cfg_attr(not(windows), path = "os_unix.rs")]
+ mod os;
+@@ -40,6 +41,7 @@ macro_rules! implemented {
+                 if !initialized() {
+                     return Err(CUerror::DEINITIALIZED);
+                 }
++                let _profile = d4r_api_profile::Call::begin(stringify!($fn_name));
+                 cuda_macros::cuda_normalize_fn!( crate::r#impl::$fn_name ) ($(zluda_common::FromCuda::<_, CUerror>::from_cuda(&$arg_id)?),*)?;
+                 Ok(())
+             }
+@@ -57,6 +59,7 @@ macro_rules! implemented_in_function {
+                 if !initialized() {
+                     return Err(CUerror::DEINITIALIZED);
+                 }
++                let _profile = d4r_api_profile::Call::begin(stringify!($fn_name));
+                 cuda_macros::cuda_normalize_fn!( crate::r#impl::function::$fn_name ) ($(zluda_common::FromCuda::<_, CUerror>::from_cuda(&$arg_id)?),*)?;
+                 Ok(())
+             }
+-- 
+2.55.0.windows.5
+

scripts/windows/test-async-k.ps1

@@ -0,0 +1,54 @@
+[CmdletBinding()]
+param(
+    [string]$NgxCore, [string]$DlssDll, [string]$NativeRoot,
+    [string]$ZludaRoot, [string]$PackageRoot,
+    [string]$Resolution = '3840x2160',
+    [string]$OutputDirectory
+)
+$ErrorActionPreference = 'Stop'
+$repo = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..'))
+if (!$NgxCore) { $NgxCore = Join-Path $repo '_nvngx.dll' }
+if (!$DlssDll) { $DlssDll = Join-Path $repo 'nvngx_dlss.dll' }
+if (!$NativeRoot) { $NativeRoot = Join-Path $repo 'build/native-k-gfx1201' }
+if (!$ZludaRoot) { $ZludaRoot = Join-Path $repo 'dist/zluda-windows-api-profile' }
+if (!$PackageRoot) { $PackageRoot = Join-Path $repo 'dist/windows-rdna4-command-list' }
+if (!$OutputDirectory) { $OutputDirectory = Join-Path $repo ('test-results/async-k-' + (Get-Date -Format 'yyyyMMdd-HHmmss-fff')) }
+if ($Resolution -notmatch '^([0-9]+)x([0-9]+)$') { throw 'Resolution must be WIDTHxHEIGHT' }
+$width=$Matches[1]; $height=$Matches[2]
+New-Item -ItemType Directory -Force $OutputDirectory | Out-Null
+$settings = @{D4R_DIAG_BURST='1'; D4R_DIAG_RECREATE=$null; D4R_ASYNC_INTEROP=$null;
+    D4R_ZLUDA_NATIVE_DIR=$NativeRoot; D4R_QUIET_API='1'; D4R_VALIDATE_OUTPUT='1';
+    D4R_PROFILE_STAGES='1'; D4R_ZLUDA_PROFILE=$null; D4R_ZLUDA_PROFILE_API=$null;
+    ZLUDA_CACHE_DIR=(Join-Path $repo 'build/zluda-cache-windows');
+    PYTHONPATH=(Join-Path $repo '.tools/python/vendor')}
+$old=@{}
+try {
+    foreach ($key in $settings.Keys) {
+        $old[$key]=[Environment]::GetEnvironmentVariable($key,'Process')
+        [Environment]::SetEnvironmentVariable($key,$settings[$key],'Process')
+    }
+    foreach ($mode in @('burst','recreate')) {
+        [Environment]::SetEnvironmentVariable('D4R_DIAG_RECREATE', $(if ($mode -eq 'recreate') { '1' } else { $null }), 'Process')
+        foreach ($variant in @('sync','async')) {
+            [Environment]::SetEnvironmentVariable('D4R_ASYNC_INTEROP', $(if ($variant -eq 'async') { '1' } else { $null }), 'Process')
+            $output=Join-Path $OutputDirectory "$variant-$mode"
+            $arguments=@{RuntimeProfile='therock'; ZludaRoot=$ZludaRoot; PackageRoot=$PackageRoot;
+                NgxCore=$NgxCore; DlssDll=$DlssDll; NgxOnly=$true; NgxMode='d3d12'; NgxAbi='project-legacy';
+                Preset=11; NgxCreateFlags=11; NgxOutputResolution=$Resolution; RequireNativeNetwork=$true;
+                CommandListBackend=$true; EarlyIndirectProbe=$true; PixelProfile='depth-stencil';
+                BarrierMode='inherited-legacy'; Iterations=3; CaptureExceptions=$true; TimeoutSeconds=600;
+                OptiScalerDll=(Join-Path $repo 'dist/optiscaler-windows-d4r/OptiScaler.dll'); OutputDirectory=$output}
+            & (Join-Path $PSScriptRoot 'test-windows-rdna4.ps1') @arguments
+            if ($LASTEXITCODE) { throw "Failed $variant-$mode; logs: $output" }
+        }
+        & python.exe (Join-Path $repo 'tools/windows/frame_compare.py') --reference (Join-Path $OutputDirectory "sync-$mode") --actual (Join-Path $OutputDirectory "async-$mode") --width $width --height $height --exact |
+            Tee-Object -FilePath (Join-Path $OutputDirectory "$mode-comparison.log")
+        if ($LASTEXITCODE) { throw "Async $mode output mismatch" }
+    }
+    @{passed=$true; architecture='gfx1201'; preset=11; resolution=$Resolution;
+        modes=@('burst','recreate'); framesPerMode=3; exactRgb=$true} |
+        ConvertTo-Json | Set-Content -LiteralPath (Join-Path $OutputDirectory 'validation.json') -Encoding UTF8
+    Write-Host "PASS asynchronous K burst/recreate: $OutputDirectory"
+} finally {
+    foreach ($key in $old.Keys) { [Environment]::SetEnvironmentVariable($key,$old[$key],'Process') }
+}

scripts/windows/windows-game.ps1

@@ -12,6 +12,9 @@ param(
     [switch]$ProfileStages,
     [switch]$ProfileKernels,
     [switch]$ProfileCommandHooks,
+    [switch]$ProfileCudaApi,
+    [switch]$UncachedInteropLists,
+    [switch]$AsyncInterop,
     [switch]$VerboseRuntime,
     [switch]$CaptureExceptions,
     [string[]]$GameArguments = @('-dx12'),
@@ -178,6 +181,9 @@ $settings = @{
     D4R_PROFILE_STAGES=$(if ($ProfileStages) { '1' } else { $null });
     D4R_ZLUDA_PROFILE=$(if ($ProfileKernels) { '1' } else { $null });
     D4R_PROFILE_COMMAND_HOOKS=$(if ($ProfileCommandHooks) { '1' } else { $null });
+    D4R_ZLUDA_PROFILE_API=$(if ($ProfileCudaApi) { '1' } else { $null });
+    D4R_DISABLE_INTEROP_LIST_CACHE=$(if ($UncachedInteropLists) { '1' } else { $null });
+    D4R_ASYNC_INTEROP=$(if ($AsyncInterop) { '1' } else { $null });
     D4R_DIAG_DIR=$OutputDirectory; ZLUDA_LOG_DIR=(Join-Path $OutputDirectory 'zluda-trace');
     ZLUDA_CACHE_DIR=$CacheDirectory; PATH=((GamePath 'd4r/hip/bin') + ';' + (GamePath 'd4r/zluda') + ';' + $env:PATH)
 }
@@ -246,6 +252,19 @@ try {
     $stdout = Get-Content -LiteralPath (Join-Path $OutputDirectory 'd4r.stdout.log') -Raw
     $debugger = Get-Content -LiteralPath (Join-Path $OutputDirectory 'debugger.log') -Raw
     $gameExit = [regex]::Match($debugger, 'EXIT code=0x([0-9a-f]+)')
+    $gameCrash = $null
+    if ([IO.Path]::GetFileName($GameExe) -eq 'SHProto-Win64-Shipping.exe' -and $gameExit.Success -and $gameExit.Groups[1].Value -ne '0') {
+        $crashRoot = Join-Path $env:LOCALAPPDATA 'SilentHill2/Saved/Crashes'
+        $latest = Get-ChildItem -LiteralPath $crashRoot -Directory -ErrorAction SilentlyContinue |
+            Where-Object { $_.LastWriteTime -ge (Get-Date).AddSeconds(-$timer.Elapsed.TotalSeconds-5) } |
+            Sort-Object LastWriteTime -Descending | Select-Object -First 1
+        if ($latest) {
+            Copy-Item -LiteralPath $latest.FullName -Destination (Join-Path $OutputDirectory 'ue-crash') -Recurse
+            [xml]$context = Get-Content -LiteralPath (Join-Path $latest.FullName 'CrashContext.runtime-xml') -Raw
+            $gameCrash = @{type=[string]$context.FGenericCrashContext.RuntimeProperties.CrashType;
+                reason=[string]$context.FGenericCrashContext.RuntimeProperties.ErrorMessage}
+        }
+    }
     # A bounded stop can interrupt the last printf; count only complete records.
     $frames = [regex]::Matches($stdout, '(?m)^D4R_FRAME cpu_frame_copies=([0-9]+) frame_age=([0-9]+) interop_ngx_ms=([0-9.]+)\r?$')
     $checks = [regex]::Matches($stdout, '(?m)^D4R_OUTPUT_VALIDATION elements=([0-9]+) nan=([0-9]+) inf=([0-9]+) diagnostics_cpu_bytes=8\r?$')
@@ -254,7 +273,7 @@ try {
     $kernels = @($launches | ForEach-Object { $_.Groups[1].Value + ' ' + $_.Groups[2].Value } | Group-Object | ForEach-Object {
         $parts=$_.Name.Split(' '); @{kernel=$parts[0]; backend=$parts[1]; launches=$_.Count}
     })
-    @{exitCode=$(if ($gameExit.Success) { '0x' + $gameExit.Groups[1].Value } else { 'diagnostic_timeout' }); runSeconds=$timer.Elapsed.TotalSeconds; preset=$Preset;
+    @{exitCode=$(if ($gameExit.Success) { '0x' + $gameExit.Groups[1].Value } else { 'diagnostic_timeout' }); runSeconds=$timer.Elapsed.TotalSeconds; preset=$Preset; gameCrash=$gameCrash;
         nativeLaunches=([regex]::Matches($stderr, '\[d4r-launch\].*backend=native')).Count;
         completedFrames=$frames.Count; outputValidationRequested=[bool]$ValidateOutput;
         outputGpuChecks=$checks.Count; nonfiniteOutputs=$nonfinite; kernels=$kernels;

tools/windows/d3d12_command_hooks.cpp

@@ -1,5 +1,6 @@
 #include "d3d12_command_hooks.h"
 #include "d3d12_newer_commands.h"
+#include "d3d12_retirement.h"
 #include <MinHook.h>
 #include <atomic>
 #include <memory>
@@ -406,6 +407,9 @@ std::mutex queue_metadata_mutex;
 struct SubmissionLock final : IUnknown {
     std::atomic<ULONG> references{1};
     std::recursive_mutex mutex;
+    ComPtr<ID3D12Fence> completion_fence;
+    Handle completion_event;
+    uint64_t completion_value = 0;
     HRESULT STDMETHODCALLTYPE QueryInterface(REFIID iid, void** out) override {
         if (!out) return E_POINTER; *out = nullptr;
         if (iid != IID_IUnknown) return E_NOINTERFACE;
@@ -442,15 +446,31 @@ HRESULT STDMETHODCALLTYPE hook_queue_wait(ID3D12CommandQueue* queue, ID3D12Fence
     return queue_fence(original_queue_wait, queue, fence, value);
 }
 void complete(ID3D12CommandQueue* queue) {
-    ComPtr<ID3D12Device> device; dx(queue->GetDevice(IID_PPV_ARGS(device.GetAddressOf())), "Queue device(submit drain)");
+    auto owner = submission_lock(queue);
-    ComPtr<ID3D12Fence> fence; dx(device->CreateFence(0, D3D12_FENCE_FLAG_NONE, IID_PPV_ARGS(fence.GetAddressOf())), "Queue fence(submit drain)");
+    std::lock_guard<std::recursive_mutex> ordered(owner->mutex);
-    Handle event; event.value = CreateEventW(nullptr, FALSE, FALSE, nullptr);
+    if (!owner->completion_fence) {
-    if (!event.value) throw std::runtime_error("CreateEvent(submit drain)");
+        ComPtr<ID3D12Device> device; dx(queue->GetDevice(IID_PPV_ARGS(device.GetAddressOf())), "Queue device(submit drain)");
-    dx(queue->Signal(fence.Get(), 1), "Queue Signal(submit drain)");
+        dx(device->CreateFence(0, D3D12_FENCE_FLAG_NONE, IID_PPV_ARGS(owner->completion_fence.GetAddressOf())), "Queue fence(submit drain)");
-    if (fence->GetCompletedValue() < 1) {
-        dx(fence->SetEventOnCompletion(1, event.value), "Queue event(submit drain)");
-        if (WaitForSingleObject(event.value, 30000) != WAIT_OBJECT_0) throw std::runtime_error("GPU submit boundary timed out");
     }
+    if (!owner->completion_event.value) owner->completion_event.value = CreateEventW(nullptr, FALSE, FALSE, nullptr);
+    if (!owner->completion_event.value) throw std::runtime_error("CreateEvent(submit drain)");
+    const auto value = ++owner->completion_value;
+    dx(queue->Signal(owner->completion_fence.Get(), value), "Queue Signal(submit drain)");
+    if (owner->completion_fence->GetCompletedValue() < value) {
+        dx(owner->completion_fence->SetEventOnCompletion(value, owner->completion_event.value), "Queue event(submit drain)");
+        if (WaitForSingleObject(owner->completion_event.value, 30000) != WAIT_OBJECT_0) throw std::runtime_error("GPU submit boundary timed out");
+    }
+}
+void retire(ID3D12CommandQueue* queue, std::vector<std::shared_ptr<Recording>> owners) {
+    auto owner = submission_lock(queue);
+    if (!owner->completion_fence) {
+        ComPtr<ID3D12Device> device; dx(queue->GetDevice(IID_PPV_ARGS(device.GetAddressOf())), "Retirement device");
+        dx(device->CreateFence(0, D3D12_FENCE_FLAG_NONE, IID_PPV_ARGS(owner->completion_fence.GetAddressOf())), "Retirement fence");
+    }
+    const auto value = ++owner->completion_value;
+    dx(queue->Signal(owner->completion_fence.Get(), value), "Queue Signal(retirement)");
+    GpuRetirement::instance().submit(owner->completion_fence.Get(), value,
+        std::make_shared<std::vector<std::shared_ptr<Recording>>>(std::move(owners)));
 }
 void STDMETHODCALLTYPE hook_execute(ID3D12CommandQueue* queue, UINT count, ID3D12CommandList* const* lists) {
     if (internal) { original_execute(queue, count, lists); return; }
@@ -498,7 +518,9 @@ void STDMETHODCALLTYPE hook_execute(ID3D12CommandQueue* queue, UINT count, ID3D1
         }
         flush();
         if (boundaries) {
-            const auto drainStart = Clock::now(); complete(queue); drainMs = milliseconds(drainStart);
+            const auto drainStart = Clock::now();
+            if (std::getenv("D4R_ASYNC_INTEROP")) retire(queue, std::move(owners)); else complete(queue);
+            drainMs = milliseconds(drainStart);
             if (profile_commands()) std::printf("D4R_COMMAND_SUBMIT total_ms=%.3f boundary_ms=%.3f publish_ms=%.3f suffix_drain_ms=%.3f\n",
                 milliseconds(submitStart), boundaryMs, publishMs, drainMs);
         }

tools/windows/d3d12_dlss_probe.cpp

@@ -281,7 +281,63 @@ int main(int argc, char** argv) {
                 IID_PPV_ARGS(list.ReleaseAndGetAddressOf())), "Create game-style command list");
             dx(list->Close(), "Close initial game-style list");
         }
+        const bool burst = std::getenv("D4R_DIAG_BURST") != nullptr;
+        const bool recreate = std::getenv("D4R_DIAG_RECREATE") != nullptr;
+        if (burst && (!commandBackend || args.iterations > 8)) throw std::runtime_error("Burst requires command-list backend and at most 8 frames");
+        struct SubmittedFrame {
+            ComPtr<ID3D12CommandAllocator> allocator;
+            ComPtr<ID3D12GraphicsCommandList> list;
+            ComPtr<ID3D12Resource> readback;
+        };
+        std::vector<SubmittedFrame> submitted;
+        std::vector<std::pair<ComPtr<ID3D12CommandAllocator>, ComPtr<ID3D12GraphicsCommandList>>> creation_owners;
+        auto validate_output = [&](ID3D12Resource* image_readback, unsigned frame) {
+            void* raw = nullptr; D3D12_RANGE range{0, size_t(textures[3].bytes)};
+            dx(image_readback->Map(0, &range, &raw), "Map(verification only)");
+            std::vector<uint16_t> pixels(size_t(outWidth) * outHeight * 4);
+            const auto format = resource_desc(textures[3].image.Get()).Format;
+            if (format == DXGI_FORMAT_R16G16B16A16_FLOAT) {
+                for (unsigned y = 0; y < outHeight; ++y) std::memcpy(pixels.data() + size_t(y) * outWidth * 4,
+                    static_cast<uint8_t*>(raw) + y * textures[3].footprint.Footprint.RowPitch, outWidth * 8);
+            } else {
+                std::ofstream native(directory / ("output-" + std::to_string(frame) + (args.pixel_profile == "packed" ? ".r11g11b10f" : ".bgra8")), std::ios::binary);
+                // Independently decode the diagnostic readback for previews and
+                // frame comparisons. No CPU image copies are in the backend.
+                static const auto unormHalf = [] {
+                    std::array<uint16_t, 256> lut{};
+                    for (unsigned b = 0; b < 256; ++b) { double best = INFINITY; for (unsigned h = 0; h <= 0x3c00; ++h) { double error = std::abs(double(half_value(uint16_t(h))) - double(b)/255.); if (error < best || (error == best && !(h & 1))) { best = error; lut[b] = uint16_t(h); } } }
+                    return lut;
+                }();
+                for (unsigned y = 0; y < outHeight; ++y) {
+                    const auto* row = static_cast<uint8_t*>(raw) + y * textures[3].footprint.Footprint.RowPitch;
+                    native.write(reinterpret_cast<const char*>(row), outWidth * 4);
+                    for (unsigned x = 0; x < outWidth; ++x) {
+                        auto* p = pixels.data() + (size_t(y) * outWidth + x) * 4;
+                        if (format == DXGI_FORMAT_R11G11B10_FLOAT) { uint32_t value = 0; std::memcpy(&value,row + x * 4,4); p[0] = uint16_t((value & 2047) << 4); p[1] = uint16_t(((value >> 11) & 2047) << 4); p[2] = uint16_t((value >> 22) << 5); p[3] = 0x3c00; }
+                        else for (unsigned c = 0; c < 4; ++c) p[c] = unormHalf[row[x * 4 + (c == 0 ? 2 : c == 2 ? 0 : c)]];
+                    }
+                }
+                if (!native) throw std::runtime_error("Saving native output format failed");
+            }
+            D3D12_RANGE noWrite{}; image_readback->Unmap(0, &noWrite);
+            double sum = 0, square = 0;
+            for (size_t i = 0; i < pixels.size(); ++i) {
+                const float value = half_value(pixels[i]);
+                if (!std::isfinite(value)) throw std::runtime_error("D3D12 output contains NaN/Inf or unwritten pixels");
+                if ((i & 3) < 3) { sum += value; square += double(value) * value; }
+            }
+            const double count = outWidth * outHeight * 3, variance = square / count - (sum / count) * (sum / count);
+            save_output(directory, pixels, outWidth, outHeight, frame);
+            std::printf("D3D12_OUTPUT preset=%u frame=%u finite=1 mean=%.9g variance=%.9g frame_age=0\n", args.preset, frame, sum/count, variance);
+            if (variance < 1e-6) throw std::runtime_error("D3D12 output lost the synthetic pattern");
+        };
         for (unsigned frame = 0; frame < args.iterations; ++frame) {
+            if (burst) {
+                dx(device->CreateCommandAllocator(D3D12_COMMAND_LIST_TYPE_DIRECT, IID_PPV_ARGS(allocator.ReleaseAndGetAddressOf())), "Burst allocator");
+                dx(device->CreateCommandList(0, D3D12_COMMAND_LIST_TYPE_DIRECT, allocator.Get(), nullptr, IID_PPV_ARGS(list.ReleaseAndGetAddressOf())), "Burst list");
+                dx(list->Close(), "Burst initial Close");
+                readback = make_buffer(device.Get(), textures[3].bytes, D3D12_HEAP_TYPE_READBACK, D3D12_HEAP_FLAG_NONE, D3D12_RESOURCE_STATE_COPY_DEST);
+            }
             d4r_ngx_set_int(parameters, "Reset", frame == 0 ? 1 : 0);
             if (commandBackend) {
                 dx(allocator->Reset(), "Reset(game-style allocator)"); dx(list->Reset(allocator.Get(), nullptr), "Reset(game-style list)");
@@ -312,46 +368,33 @@ int main(int argc, char** argv) {
             if (commandBackend && frame == 0) {
                 ID3D12CommandList* batch[] = {predecessor.Get(), list.Get()}; queue->ExecuteCommandLists(2, batch);
             } else { ID3D12CommandList* verify[] = {list.Get()}; queue->ExecuteCommandLists(1, verify); }
-            drain();
+            if (!burst) drain(); else submitted.push_back({allocator, list, readback});
             if (commandBackend) d4r_ngx_set_d3d12_resource(parameters, "Color", textures[0].image.Get());
-            void* raw = nullptr; D3D12_RANGE range{0, size_t(textures[3].bytes)};
+            if (!burst) validate_output(readback.Get(), frame);
-            dx(readback->Map(0, &range, &raw), "Map(verification only)");
+            if (recreate && frame + 1 < args.iterations) {
-            std::vector<uint16_t> pixels(size_t(outWidth) * outHeight * 4);
+                // Release/create without an explicit harness drain. In burst
-            const auto format = resource_desc(textures[3].image.Get()).Format;
+                // mode earlier frames can still be using the old feature.
-            if (format == DXGI_FORMAT_R16G16B16A16_FLOAT) {
+                check(release(handle), "Recreate ReleaseFeature"); featureCleanup.h = nullptr;
-                for (unsigned y = 0; y < outHeight; ++y) std::memcpy(pixels.data() + size_t(y) * outWidth * 4,
+                dx(creationAllocator->Reset(), "Recreate allocator Reset");
-                    static_cast<uint8_t*>(raw) + y * textures[3].footprint.Footprint.RowPitch, outWidth * 8);
+                dx(creationList->Reset(creationAllocator.Get(), nullptr), "Recreate list Reset");
-            } else {
+                handle = nullptr;
-                std::ofstream native(directory / ("output-" + std::to_string(frame) + (args.pixel_profile == "packed" ? ".r11g11b10f" : ".bgra8")), std::ios::binary);
+                check(shim.symbol<Create>("NVSDK_NGX_D3D12_CreateFeature")(creationList.Get(), 1, parameters, &handle), "Recreate CreateFeature");
-                // Independently decode the diagnostic readback for previews and
+                featureCleanup.h = handle;
-                // frame comparisons. No CPU image copies are in the backend.
+                dx(creationList->Close(), "Recreate list Close");
-                static const auto unormHalf = [] {
+                ID3D12CommandList* created[] = {creationList.Get()}; queue->ExecuteCommandLists(1, created);
-                    std::array<uint16_t, 256> lut{};
+                creation_owners.emplace_back(creationAllocator, creationList);
-                    for (unsigned b = 0; b < 256; ++b) { double best = INFINITY; for (unsigned h = 0; h <= 0x3c00; ++h) { double error = std::abs(double(half_value(uint16_t(h))) - double(b)/255.); if (error < best || (error == best && !(h & 1))) { best = error; lut[b] = uint16_t(h); } } }
+                // This empty creation list owns no work; the next recreation
-                    return lut;
+                // allocates independent command memory without waiting on GPU.
-                }();
+                dx(device->CreateCommandAllocator(D3D12_COMMAND_LIST_TYPE_DIRECT, IID_PPV_ARGS(creationAllocator.ReleaseAndGetAddressOf())), "Next recreate allocator");
-                for (unsigned y = 0; y < outHeight; ++y) {
+                dx(device->CreateCommandList(0, D3D12_COMMAND_LIST_TYPE_DIRECT, creationAllocator.Get(), nullptr, IID_PPV_ARGS(creationList.ReleaseAndGetAddressOf())), "Next recreate list");
-                    const auto* row = static_cast<uint8_t*>(raw) + y * textures[3].footprint.Footprint.RowPitch;
+                dx(creationList->Close(), "Next recreate initial Close");
-                    native.write(reinterpret_cast<const char*>(row), outWidth * 4);
+                std::printf("D4R_RECREATE_VALIDATION submitted_frame=%u next_feature_created=1\n", frame);
-                    for (unsigned x = 0; x < outWidth; ++x) {
-                        auto* p = pixels.data() + (size_t(y) * outWidth + x) * 4;
-                        if (format == DXGI_FORMAT_R11G11B10_FLOAT) { uint32_t value = 0; std::memcpy(&value,row + x * 4,4); p[0] = uint16_t((value & 2047) << 4); p[1] = uint16_t(((value >> 11) & 2047) << 4); p[2] = uint16_t((value >> 22) << 5); p[3] = 0x3c00; }
-                        else for (unsigned c = 0; c < 4; ++c) p[c] = unormHalf[row[x * 4 + (c == 0 ? 2 : c == 2 ? 0 : c)]];
-                    }
-                }
-                if (!native) throw std::runtime_error("Saving native output format failed");
             }
-            D3D12_RANGE noWrite{}; readback->Unmap(0, &noWrite);
+        }
-            double sum = 0, square = 0;
+        if (burst) {
-            for (size_t i = 0; i < pixels.size(); ++i) {
+            drain();
-                const float value = half_value(pixels[i]);
+            for (unsigned frame = 0; frame < submitted.size(); ++frame) validate_output(submitted[frame].readback.Get(), frame);
-                if (!std::isfinite(value)) throw std::runtime_error("D3D12 output contains NaN/Inf or unwritten pixels");
+            std::printf("D4R_BURST_VALIDATION frames=%zu cpu_readback_after_all_submissions=1\n", submitted.size());
-                if ((i & 3) < 3) { sum += value; square += double(value) * value; }
-            }
-            const double count = outWidth * outHeight * 3, variance = square / count - (sum / count) * (sum / count);
-            save_output(directory, pixels, outWidth, outHeight, frame);
-            std::printf("D3D12_OUTPUT preset=%u frame=%u finite=1 mean=%.9g variance=%.9g frame_age=0\n", args.preset, frame, sum/count, variance);
-            if (variance < 1e-6) throw std::runtime_error("D3D12 output lost the synthetic pattern");
         }
         loaded_modules();
         std::printf("PASS D3D12_DLSS architecture=gfx1201 preset=%u frames=%u fast_path_cpu_copies=0 frame_age=0 enhanced=%u queue_integration=%s\n",

tools/windows/d3d12_external.h

@@ -178,6 +178,7 @@ class SharedTimeline {
     Handle event_;
     uint64_t value_ = 0;
 public:
+    struct Ticket { uint64_t input, output; };
     ComPtr<ID3D12Fence> fence;
     SharedTimeline(ExternalApi& api, ID3D12Device* device) : api_(api) {
         dx(device->CreateFence(0, D3D12_FENCE_FLAG_SHARED, IID_PPV_ARGS(fence.GetAddressOf())), "CreateFence(D3D12/HIP)");
@@ -200,6 +201,17 @@ public:
     void wait_input(ID3D12CommandQueue* queue) {
         const uint64_t ready = ++value_;
         dx(queue->Signal(fence.Get(), ready), "Queue Signal(input)");
+        wait_ready(ready);
+    }
+    // Reserve both dependencies on the submission thread. The CUDA worker
+    // never submits to the D3D12 queue, so later game Signal/Wait calls retain
+    // their original ordering without waiting on a CPU callback.
+    Ticket reserve() { return {++value_, ++value_}; }
+    void enqueue(ID3D12CommandQueue* queue, Ticket ticket) {
+        dx(queue->Signal(fence.Get(), ticket.input), "Queue Signal(async input)");
+        dx(queue->Wait(fence.Get(), ticket.output), "Queue Wait(async output)");
+    }
+    void wait_ready(uint64_t ready) {
         hipExternalSemaphoreWaitParams wait{}; wait.params.fence.value = ready;
         api_.hip.check(api_.hipWaitExternalSemaphoresAsync(&semaphore_, &wait, 1, stream_), "HIP wait(D3D12 input)");
         api_.hip.check(api_.hip.hipStreamSynchronize(stream_), "HIP input wait completion");
@@ -214,9 +226,12 @@ public:
     }
     void signal_output(ID3D12CommandQueue* queue) {
         const uint64_t done = ++value_;
+        signal_ready(done);
+        dx(queue->Wait(fence.Get(), done), "Queue Wait(HIP output)");
+    }
+    void signal_ready(uint64_t done) {
         hipExternalSemaphoreSignalParams signal{}; signal.params.fence.value = done;
         api_.hip.check(api_.hipSignalExternalSemaphoresAsync(&semaphore_, &signal, 1, stream_), "HIP signal(output)");
-        dx(queue->Wait(fence.Get(), done), "Queue Wait(HIP output)");
     }
     void drain(ID3D12CommandQueue* queue) {
         const uint64_t done = ++value_;

tools/windows/d3d12_retirement.h

@@ -0,0 +1,46 @@
+#pragma once
+#include "d3d12_external.h"
+#include <condition_variable>
+#include <deque>
+#include <memory>
+#include <mutex>
+#include <thread>
+
+namespace d4r::win {
+// GPU-completion ownership, independent of command-list Reset and queue COM
+// lifetime. The command backend pins its DLL. One process-lifetime thread
+// releases completed batches outside its lock, including the final batch;
+// queue -> recording -> feature -> queue cannot become a reference cycle.
+class GpuRetirement {
+    struct Batch { ComPtr<ID3D12Fence> fence; uint64_t value; std::shared_ptr<void> owner; };
+    std::mutex mutex_;
+    std::condition_variable changed_;
+    std::deque<Batch> pending_;
+    GpuRetirement() {
+        std::thread([this] {
+            Handle event; event.value = CreateEventW(nullptr, FALSE, FALSE, nullptr);
+            for (;;) {
+                Batch batch;
+                { std::unique_lock<std::mutex> lock(mutex_);
+                  changed_.wait(lock, [&] { return !pending_.empty(); });
+                  batch = std::move(pending_.front()); pending_.pop_front(); }
+                if (batch.fence->GetCompletedValue() < batch.value) {
+                    if (!event.value || FAILED(batch.fence->SetEventOnCompletion(batch.value, event.value)) ||
+                        WaitForSingleObject(event.value, 30000) != WAIT_OBJECT_0) {
+                        std::fprintf(stderr, "D4R_RETIREMENT_FAILURE GPU completion timeout; retaining resources\n");
+                        // A hung GPU must not make us free live command memory.
+                        auto* retained = new Batch(std::move(batch)); (void)retained;
+                    }
+                }
+                batch.owner.reset();
+            }
+        }).detach();
+    }
+public:
+    static GpuRetirement& instance() { static auto* service = new GpuRetirement; return *service; }
+    void submit(ID3D12Fence* fence, uint64_t value, std::shared_ptr<void> owner) {
+        { std::lock_guard<std::mutex> lock(mutex_); pending_.push_back({fence, value, std::move(owner)}); }
+        changed_.notify_one();
+    }
+};
+}

tools/windows/d4r_nvngx_windows.cpp

@@ -184,11 +184,14 @@ API unsigned NVSDK_NGX_D3D12_ReleaseFeature(Handle* handle) {
     { std::lock_guard<std::mutex> lock(apiMutex);
       auto it = features.find(handle); if (it == features.end()) return invalid;
       feature = std::move(it->second); features.erase(it); delete handle; }
-    return call([&] { feature.reset(); return 1u; });
+    // OptiScaler frees its root-CBV/descriptor backing after Release returns.
+    // A retained C++ Feature alone does not keep that caller-owned memory alive.
+    return call([&] { feature->drain(); feature.reset(); return 1u; });
 }
 API unsigned NVSDK_NGX_D3D12_Shutdown() {
     std::lock_guard<std::mutex> lock(apiMutex);
     return call([&] {
+        if (runtime && std::getenv("D4R_ASYNC_INTEROP")) runtime->drain_async_queue();
         for (auto& pair : features) delete pair.first;
         features.clear();
         for (auto& pair : parameters) d4r_ngx_parameters_destroy(pair.first);

tools/windows/ngx_windows_runtime.h

@@ -5,6 +5,10 @@
 #include "ngx_parameters.h"
 #include "ngx_public_abi.h"
 #include "d3d12_command_hooks.h"
+#include "d3d12_retirement.h"
+#include <condition_variable>
+#include <deque>
+#include <thread>
 #include <chrono>
 #include <memory>
 #include <mutex>
@@ -55,6 +59,14 @@ struct Runtime {
     std::unique_ptr<SharedTimeline> timeline;
     std::unique_ptr<PixelProgram> pixels;
     std::mutex mutex;
+    // Features share one CUDA context. Follow submission order even when
+    // separate feature workers wake in a different order; otherwise a worker
+    // could hold the context while waiting for an earlier worker's output.
+    std::mutex async_mutex;
+    std::condition_variable async_changed;
+    std::deque<uint64_t> async_tickets;
+    std::string async_failure;
+    ComPtr<ID3D12CommandQueue> async_queue;
     ngx::ProjectIdentity project;
     ngx::FeatureCommonInfo common{};
     std::vector<std::wstring> featurePaths;
@@ -146,6 +158,24 @@ struct Runtime {
     Runtime(const Runtime&) = delete;
     ~Runtime() { cleanup(); }
     void current() { cuda.check(cuda.cuCtxSetCurrent(context), "cuCtxSetCurrent(runtime)"); }
+    void drain_async_queue() {
+        ComPtr<ID3D12CommandQueue> queue;
+        { std::lock_guard<std::mutex> lock(async_mutex); queue = async_queue; }
+        if (!queue) return;
+        // Creation/free/import may synchronize the AMD device internally.
+        // Drain queued D3D12 consumers before taking the CUDA mutex: workers
+        // must remain able to satisfy any already-published output waits.
+        ComPtr<ID3D12Fence> completion;
+        dx(device->CreateFence(0, D3D12_FENCE_FLAG_NONE, IID_PPV_ARGS(completion.GetAddressOf())), "Reconfigure fence");
+        Handle event; event.value = CreateEventW(nullptr, FALSE, FALSE, nullptr);
+        if (!event.value) throw std::runtime_error("Reconfigure event");
+        dx(queue->Signal(completion.Get(), 1), "Reconfigure queue Signal");
+        if (completion->GetCompletedValue() < 1) {
+            dx(completion->SetEventOnCompletion(1, event.value), "Reconfigure completion event");
+            if (WaitForSingleObject(event.value, 30000) != WAIT_OBJECT_0) throw std::runtime_error("Reconfigure GPU timeout");
+        }
+        dx(device->GetDeviceRemovedReason(), "Reconfigure device status");
+    }
     void cleanup() noexcept {
         if (context) (void)cuda.cuCtxSetCurrent(context);
         if (initialized) { (void)shutdown(); initialized = false; }
@@ -164,6 +194,32 @@ class Feature {
     void* handle_ = nullptr;
     CUdeviceptr scratch_ = 0;
     ComPtr<ID3D12CommandQueue> queue_;
+    // A boundary drains both copies before returning. Reuse their allocators
+    // only after that completion; no game allocator is retained here.
+    ComPtr<ID3D12CommandAllocator> input_allocator_, output_allocator_;
+    ComPtr<ID3D12GraphicsCommandList> input_list_, output_list_;
+    struct CopySlot {
+        ComPtr<ID3D12CommandAllocator> input_allocator, output_allocator;
+        ComPtr<ID3D12GraphicsCommandList> input, output;
+        uint64_t done = 0;
+    };
+    struct AsyncFrame {
+        void* parameters = d4r_ngx_parameters_create();
+        std::vector<ComPtr<ID3D12Resource>> resources;
+        SharedTimeline::Ticket ticket{};
+        ~AsyncFrame() { if (parameters) d4r_ngx_parameters_destroy(parameters); }
+    };
+    std::vector<std::unique_ptr<CopySlot>> copy_slots_;
+    ComPtr<ID3D12Fence> copy_completion_;
+    Handle copy_event_;
+    uint64_t copy_value_ = 0;
+    std::mutex jobs_mutex_;
+    std::mutex submit_mutex_;
+    std::condition_variable jobs_changed_;
+    std::deque<std::shared_ptr<AsyncFrame>> jobs_;
+    std::thread worker_;
+    bool worker_stop_ = false, worker_active_ = false;
+    std::string worker_failure_;
     struct Plane {
         ComPtr<ID3D12Resource> texture;
         D3D12_RESOURCE_DESC desc{};
@@ -185,6 +241,7 @@ class Feature {
 public:
     unsigned width = 0, height = 0;
     explicit Feature(std::shared_ptr<Runtime> runtime, void* parameters) : rt_(std::move(runtime)) {
+        if (std::getenv("D4R_ASYNC_INTEROP")) rt_->drain_async_queue();
         std::lock_guard<std::mutex> lock(rt_->mutex);
         rt_->current();
         try {
@@ -203,12 +260,38 @@ public:
             ngx_check(rt_->create(1, parameters_, &handle_), "CUDA CreateFeature");
         } catch (...) { cleanup(); throw; }
     }
-    ~Feature() { std::lock_guard<std::mutex> lock(rt_->mutex); cleanup(); }
+    ~Feature() {
+        { std::lock_guard<std::mutex> lock(jobs_mutex_); worker_stop_ = true; }
+        jobs_changed_.notify_all();
+        if (worker_.joinable()) worker_.join();
+        if (copy_completion_) {
+            try { wait_copies(copy_value_); }
+            catch (const std::exception& e) { std::fprintf(stderr, "D4R_ASYNC_CLEANUP_FAILURE %s\n", e.what()); }
+        }
+        std::lock_guard<std::mutex> lock(rt_->mutex); cleanup();
+    }
     Feature(const Feature&) = delete;
+    void drain() {
+        if (!std::getenv("D4R_ASYNC_INTEROP")) return;
+        wait_worker();
+        if (copy_value_) wait_copies(copy_value_);
+        rt_->drain_async_queue();
+    }
 
     // Explicit harness boundary: all earlier producers must already be queued.
     // The forthcoming command-list backend calls the same stages at queue submit.
     void evaluate_boundary(ID3D12CommandQueue* queue, void* parameters) {
+        if (std::getenv("D4R_ASYNC_INTEROP")) { evaluate_async(queue, parameters); return; }
+        const auto boundaryBegin = std::chrono::steady_clock::now();
+        auto boundaryStage = boundaryBegin;
+        const bool profile = std::getenv("D4R_PROFILE_STAGES") != nullptr;
+        auto mark = [&](const char* name) {
+            if (!profile) return;
+            const auto now = std::chrono::steady_clock::now();
+            std::printf("D4R_STAGE name=%s cpu_ms=%.6f\n", name,
+                std::chrono::duration<double, std::milli>(now - boundaryStage).count());
+            boundaryStage = std::chrono::steady_clock::now();
+        };
         commands::InternalScope internal;
         std::lock_guard<std::mutex> lock(rt_->mutex);
         if (!queue) throw std::runtime_error("Null D3D12 queue");
@@ -222,6 +305,7 @@ public:
         if (queueIdentity.Get() != runtimeIdentity.Get()) throw std::runtime_error("Queue belongs to a different D3D12 device");
         rt_->current();
         prepare(parameters);
+        mark("boundary_prepare");
         ComPtr<ID3D12CommandAllocator> inputAllocator, outputAllocator;
         ComPtr<ID3D12GraphicsCommandList> inputList, outputList;
         auto list = [&](ComPtr<ID3D12CommandAllocator>& allocator, ComPtr<ID3D12GraphicsCommandList>& commandList) {
@@ -229,20 +313,191 @@ public:
             dx(rt_->device->CreateCommandList(0, D3D12_COMMAND_LIST_TYPE_DIRECT, allocator.Get(), nullptr,
                 IID_PPV_ARGS(commandList.GetAddressOf())), "CreateCommandList(interop)");
         };
-        list(inputAllocator, inputList); list(outputAllocator, outputList);
+        const bool cached = !std::getenv("D4R_DISABLE_INTEROP_LIST_CACHE");
+        if (cached) {
+            if (!input_list_) {
+                // Publish the pair only after both creations succeed.
+                list(inputAllocator, inputList); list(outputAllocator, outputList);
+                input_allocator_ = inputAllocator; output_allocator_ = outputAllocator;
+                input_list_ = inputList; output_list_ = outputList;
+            } else {
+                dx(input_allocator_->Reset(), "Reset(input interop allocator)");
+                dx(output_allocator_->Reset(), "Reset(output interop allocator)");
+                dx(input_list_->Reset(input_allocator_.Get(), nullptr), "Reset(input interop list)");
+                dx(output_list_->Reset(output_allocator_.Get(), nullptr), "Reset(output interop list)");
+            }
+            inputAllocator = input_allocator_; outputAllocator = output_allocator_;
+            inputList = input_list_; outputList = output_list_;
+        } else { list(inputAllocator, inputList); list(outputAllocator, outputList); }
+        mark("boundary_copy_list_setup");
         record_inputs(inputList.Get(), parameters); dx(inputList->Close(), "Close(input copies)");
         ID3D12CommandList* before[] = {inputList.Get()}; queue->ExecuteCommandLists(1, before);
+        mark("boundary_input_copy_submit");
         try {
             execute_cuda(queue);
+            mark("boundary_cuda_transaction");
             record_output(outputList.Get(), parameters); dx(outputList->Close(), "Close(output copy)");
             ID3D12CommandList* after[] = {outputList.Get()}; queue->ExecuteCommandLists(1, after);
+            mark("boundary_output_copy_submit");
             rt_->timeline->drain(queue);
+            mark("boundary_output_copy_drain");
         } catch (...) {
             // Keep all recorded resources alive until submitted work is complete.
             rt_->timeline->drain(queue); throw;
         }
     }
 private:
+    void wait_copies(uint64_t value) {
+        if (copy_completion_->GetCompletedValue() >= value) return;
+        dx(copy_completion_->SetEventOnCompletion(value, copy_event_.value), "Async copy completion event");
+        if (WaitForSingleObject(copy_event_.value, 30000) != WAIT_OBJECT_0)
+            throw std::runtime_error("Async interop copy completion timeout");
+    }
+    void wait_worker() {
+        std::unique_lock<std::mutex> lock(jobs_mutex_);
+        jobs_changed_.wait(lock, [&] { return jobs_.empty() && !worker_active_; });
+        if (!worker_failure_.empty()) throw std::runtime_error(worker_failure_);
+    }
+    CopySlot& copy_slot() {
+        if (!copy_completion_) {
+            dx(rt_->device->CreateFence(0, D3D12_FENCE_FLAG_NONE, IID_PPV_ARGS(copy_completion_.GetAddressOf())), "Async copy fence");
+            copy_event_.value = CreateEventW(nullptr, FALSE, FALSE, nullptr);
+            if (!copy_event_.value) throw std::runtime_error("Async copy event");
+        }
+        CopySlot* selected = nullptr;
+        const auto completed = copy_completion_->GetCompletedValue();
+        for (auto& slot : copy_slots_) if (slot->done <= completed) { selected = slot.get(); break; }
+        if (!selected && copy_slots_.size() == 3) {
+            selected = copy_slots_.front().get();
+            for (auto& slot : copy_slots_) if (slot->done < selected->done) selected = slot.get();
+            // Bounded ownership; this wait is outside the runtime mutex.
+            wait_copies(selected->done);
+        }
+        if (selected) {
+            dx(selected->input_allocator->Reset(), "Async input allocator Reset");
+            dx(selected->output_allocator->Reset(), "Async output allocator Reset");
+            dx(selected->input->Reset(selected->input_allocator.Get(), nullptr), "Async input list Reset");
+            dx(selected->output->Reset(selected->output_allocator.Get(), nullptr), "Async output list Reset");
+            return *selected;
+        }
+        auto slot = std::make_unique<CopySlot>();
+        auto create = [&](auto& allocator, auto& list) {
+            dx(rt_->device->CreateCommandAllocator(D3D12_COMMAND_LIST_TYPE_DIRECT, IID_PPV_ARGS(allocator.GetAddressOf())), "Async copy allocator");
+            dx(rt_->device->CreateCommandList(0, D3D12_COMMAND_LIST_TYPE_DIRECT, allocator.Get(), nullptr, IID_PPV_ARGS(list.GetAddressOf())), "Async copy list");
+        };
+        create(slot->input_allocator, slot->input); create(slot->output_allocator, slot->output);
+        copy_slots_.push_back(std::move(slot)); return *copy_slots_.back();
+    }
+    bool resize_needed(void* parameters) {
+        const char* names[] = {"Color", "Depth", "MotionVectors", "Output", "ExposureTexture"};
+        for (unsigned i = 0; i < 5; ++i) {
+            ID3D12Resource* texture = nullptr; (void)d4r_ngx_get_d3d12_resource(parameters, names[i], &texture);
+            if (!texture) {
+                if (i < 4) throw std::runtime_error(std::string("Missing async resource ") + names[i]);
+                continue;
+            }
+            const auto desc = resource_desc(texture);
+            if (desc.Dimension != D3D12_RESOURCE_DIMENSION_TEXTURE2D || desc.SampleDesc.Count != 1 || desc.DepthOrArraySize != 1)
+                throw std::runtime_error("Async interop requires single-sample 2D textures");
+            if (!planes_[i].shared || planes_[i].desc.Width != desc.Width || planes_[i].desc.Height != desc.Height || planes_[i].desc.Format != desc.Format) return true;
+        }
+        return false;
+    }
+    void evaluate_async(ID3D12CommandQueue* queue, void* parameters) {
+        commands::InternalScope internal;
+        std::lock_guard<std::mutex> submitted(submit_mutex_);
+        if (!queue || (queue_ && queue_.Get() != queue)) throw std::runtime_error("Async interop requires one feature queue");
+        ComPtr<ID3D12Device> queueDevice; dx(queue->GetDevice(IID_PPV_ARGS(queueDevice.GetAddressOf())), "Async queue device");
+        ComPtr<IUnknown> a, b;
+        dx(queueDevice->QueryInterface(IID_PPV_ARGS(a.GetAddressOf())), "Async queue identity");
+        dx(rt_->device->QueryInterface(IID_PPV_ARGS(b.GetAddressOf())), "Async runtime identity");
+        if (a.Get() != b.Get()) throw std::runtime_error("Async queue device mismatch");
+        { std::lock_guard<std::mutex> lock(jobs_mutex_);
+          if (!worker_failure_.empty()) throw std::runtime_error(worker_failure_); }
+        { std::lock_guard<std::mutex> lock(rt_->async_mutex);
+          if (!rt_->async_failure.empty()) throw std::runtime_error(rt_->async_failure);
+          if (rt_->async_queue && rt_->async_queue.Get() != queue) throw std::runtime_error("Async runtime features must share one D3D12 queue");
+          rt_->async_queue = queue; }
+        auto& slot = copy_slot();
+        if (resize_needed(parameters)) {
+            wait_worker(); if (copy_value_) wait_copies(copy_value_);
+            rt_->drain_async_queue();
+            std::lock_guard<std::mutex> lock(rt_->mutex);
+            rt_->current(); prepare(parameters);
+        }
+        // Steady-state recording touches immutable allocation metadata and the
+        // caller's resource snapshot. It does not contend with CUDA or mutate
+        // the worker's NGX parameters / texture references.
+        queue_ = queue;
+        auto frame = std::make_shared<AsyncFrame>();
+        if (!frame->parameters) throw std::runtime_error("Async frame parameters allocation");
+        ngx::copy_create(parameters, frame->parameters); ngx::copy_frame(parameters, frame->parameters);
+        const char* names[] = {"Color", "Depth", "MotionVectors", "Output", "ExposureTexture"};
+        for (unsigned i = 0; i < 5; ++i) {
+            ID3D12Resource* resource = nullptr; (void)d4r_ngx_get_d3d12_resource(parameters, names[i], &resource);
+            frame->resources.emplace_back(resource);
+            d4r_ngx_set_d3d12_resource(frame->parameters, names[i], resource);
+        }
+        record_inputs(slot.input.Get(), parameters); dx(slot.input->Close(), "Async input Close");
+        record_output(slot.output.Get(), parameters); dx(slot.output->Close(), "Async output Close");
+        if (!worker_.joinable()) worker_ = std::thread([this] { work(); });
+        ID3D12CommandList* input[] = {slot.input.Get()}, *output[] = {slot.output.Get()};
+        // Allocate the job before any GPU wait is published. Worker cannot
+        // acquire the runtime mutex until every queue dependency is enqueued.
+        std::lock_guard<std::mutex> pending(jobs_mutex_);
+        { std::lock_guard<std::mutex> order(rt_->async_mutex);
+          frame->ticket = rt_->timeline->reserve(); rt_->async_tickets.push_back(frame->ticket.output); }
+        try { jobs_.push_back(frame); }
+        catch (...) {
+            std::lock_guard<std::mutex> order(rt_->async_mutex); rt_->async_tickets.pop_back(); throw;
+        }
+        try {
+            queue->ExecuteCommandLists(1, input);
+            rt_->timeline->enqueue(queue, frame->ticket);
+            queue->ExecuteCommandLists(1, output);
+            slot.done = ++copy_value_;
+            dx(queue->Signal(copy_completion_.Get(), slot.done), "Async output-copy retirement Signal");
+        } catch (const std::exception& e) {
+            jobs_.pop_back(); worker_failure_ = e.what();
+            { std::lock_guard<std::mutex> order(rt_->async_mutex);
+              rt_->async_tickets.pop_back(); rt_->async_failure = e.what(); }
+            (void)rt_->timeline->fence->Signal(frame->ticket.output);
+            jobs_changed_.notify_all(); rt_->async_changed.notify_all(); throw;
+        }
+        jobs_changed_.notify_one();
+        std::printf("D4R_ASYNC_SUBMIT ticket=%llu copy_slots=%zu cpu_frame_copies=0 frame_age=0\n",
+            (unsigned long long)frame->ticket.output, copy_slots_.size());
+    }
+    void work() noexcept {
+        commands::InternalScope internal;
+        for (;;) {
+            std::shared_ptr<AsyncFrame> frame;
+            { std::unique_lock<std::mutex> lock(jobs_mutex_);
+              jobs_changed_.wait(lock, [&] { return worker_stop_ || !jobs_.empty(); });
+              if (jobs_.empty()) return;
+              frame = std::move(jobs_.front()); jobs_.pop_front(); worker_active_ = true; }
+            { std::unique_lock<std::mutex> order(rt_->async_mutex);
+              rt_->async_changed.wait(order, [&] { return !rt_->async_tickets.empty() && rt_->async_tickets.front() == frame->ticket.output; }); }
+            try {
+                std::lock_guard<std::mutex> lock(rt_->mutex);
+                { std::lock_guard<std::mutex> pending(jobs_mutex_);
+                  if (!worker_failure_.empty()) throw std::runtime_error(worker_failure_); }
+                rt_->current(); prepare(frame->parameters); execute_cuda(nullptr, &frame->ticket);
+            } catch (const std::exception& e) {
+                { std::lock_guard<std::mutex> lock(jobs_mutex_); worker_failure_ = e.what(); }
+                { std::lock_guard<std::mutex> order(rt_->async_mutex); rt_->async_failure = e.what(); }
+                std::fprintf(stderr, "D4R_ASYNC_INTEROP_FAILURE %s\n", e.what());
+                // Poison the feature and release the queue for error handling;
+                // subsequent evaluates fail explicitly, never reuse an old frame.
+                (void)rt_->timeline->fence->Signal(frame->ticket.output);
+            }
+            { std::lock_guard<std::mutex> order(rt_->async_mutex); rt_->async_tickets.pop_front(); }
+            rt_->async_changed.notify_all();
+            frame.reset();
+            { std::lock_guard<std::mutex> lock(jobs_mutex_); worker_active_ = false; }
+            jobs_changed_.notify_all();
+        }
+    }
     void prepare(void* parameters) {
         static const char* names[] = {"Color", "Depth", "MotionVectors", "Output", "ExposureTexture"};
         for (unsigned i = 0; i < 5; ++i) {
@@ -283,12 +538,15 @@ private:
         for (const char* name : {"TransparencyMask", "DLSS.Input.Bias.Current.Color.Mask"}) d4r_ngx_set_void(parameters_, name, nullptr);
     }
     void record_inputs(ID3D12GraphicsCommandList* list, void* parameters) {
-        for (unsigned i : {0u, 1u, 2u, 4u}) if (planes_[i].texture) {
+        const char* names[] = {"Color", "Depth", "MotionVectors", "Output", "ExposureTexture"};
+        for (unsigned i : {0u, 1u, 2u, 4u}) {
+            ID3D12Resource* texture = nullptr; (void)d4r_ngx_get_d3d12_resource(parameters, names[i], &texture);
+            if (!texture) continue;
             static const char* states[] = {"D4R.Color.State", "D4R.Depth.State", "D4R.Motion.State", "D4R.Output.State", "D4R.Exposure.State"};
-            planes_[i].shared->copy_input(list, planes_[i].texture.Get(), ngx_resource_access(parameters, states[i], D3D12_RESOURCE_STATE_NON_PIXEL_SHADER_RESOURCE));
+            planes_[i].shared->copy_input(list, texture, ngx_resource_access(parameters, states[i], D3D12_RESOURCE_STATE_NON_PIXEL_SHADER_RESOURCE));
         }
     }
-    void execute_cuda(ID3D12CommandQueue* queue) {
+    void execute_cuda(ID3D12CommandQueue* queue, const SharedTimeline::Ticket* ticket = nullptr) {
         const auto begin = std::chrono::steady_clock::now();
         auto stage = begin;
         const bool profileStages = std::getenv("D4R_PROFILE_STAGES") != nullptr;
@@ -299,7 +557,7 @@ private:
                 std::chrono::duration<double, std::milli>(now - stage).count());
             stage = std::chrono::steady_clock::now();
         };
-        rt_->timeline->wait_input(queue);
+        if (ticket) rt_->timeline->wait_ready(ticket->input); else rt_->timeline->wait_input(queue);
         rt_->cuda.check(rt_->cuda.cuCtxSynchronize(), "Interop input producer completion");
         mark("input_fence_wait");
         for (unsigned i : {0u, 1u, 2u, 4u}) if (planes_[i].texture) {
@@ -347,13 +605,15 @@ private:
         mark("output_diagnostic");
         if (output.canonical) rt_->pixels->convert(true, output.canonical->data, output.canonical->pitch, output.shared->mapped,
             output.shared->footprint.Footprint.RowPitch, unsigned(output.desc.Width), output.desc.Height, output.spec.storage, 3);
-        rt_->timeline->signal_output(queue);
+        if (ticket) rt_->timeline->signal_ready(ticket->output); else rt_->timeline->signal_output(queue);
         mark("output_conversion_fence_signal");
         std::printf("D4R_FRAME cpu_frame_copies=%u frame_age=0 interop_ngx_ms=%.3f\n", std::getenv("D4R_INTEROP_VERIFY") ? 1u : 0u,
             std::chrono::duration<double, std::milli>(std::chrono::steady_clock::now() - begin).count());
     }
     void record_output(ID3D12GraphicsCommandList* list, void* parameters) {
-        planes_[3].shared->copy_output(list, planes_[3].texture.Get(), ngx_resource_access(parameters, "D4R.Output.State", D3D12_RESOURCE_STATE_UNORDERED_ACCESS));
+        ID3D12Resource* texture = nullptr; (void)d4r_ngx_get_d3d12_resource(parameters, "Output", &texture);
+        if (!texture) throw std::runtime_error("Missing output resource");
+        planes_[3].shared->copy_output(list, texture, ngx_resource_access(parameters, "D4R.Output.State", D3D12_RESOURCE_STATE_UNORDERED_ACCESS));
     }
 };
 }

tools/windows/profile_report.py

@@ -42,6 +42,7 @@ def report(directory, metadata_directory=None):
     commands = collections.defaultdict(list)
     kernels = collections.defaultdict(list)
     hook_threads = collections.defaultdict(list)
+    apis = collections.defaultdict(list)
     replay = collections.defaultdict(lambda: collections.defaultdict(dict))
     for path in sorted(directory.glob('*.log')):
         data = path.read_bytes()
@@ -57,6 +58,10 @@ def report(directory, metadata_directory=None):
                 if name != 'interval_ms' or number > 0:
                     commands[kind.lower() + '.' + name].append(number)
         for line in text.splitlines():
+            if line.startswith('D4R_CUDA_API_PROFILE '):
+                fields = dict(re.findall(r'(\w+)=(\S+)', line))
+                if all(key in fields for key in ('api', 'calls', 'total_ms', 'max_ms', 'period_ms', 'thread')):
+                    apis[(fields['thread'], fields['api'])].append(fields)
             if line.startswith('D4R_COMMAND_HOOK_SAMPLE '):
                 fields = dict(re.findall(r'(\w+)=(\S+)', line))
                 if all(name in fields for name in ('thread', 'period_ms', 'calls', 'samples', 'marked',
@@ -69,9 +74,11 @@ def report(directory, metadata_directory=None):
     result = dict(cpuStages={name: stats(values) for name, values in stages.items()},
                   commandStages={name: stats(values) for name, values in commands.items()},
                   commandHookSamples=[],
+                  cudaApi=[],
                   replayPairs=[],
                   kernels=[], notes=[
                       'CPU stage times include waits and host work; they are not isolated GPU timings.',
+                      'CUDA API profiles measure host call duration without GPU events or added synchronization; they still include existing waits.',
                       'Command record intervals are between NGX recording calls on one thread, not Present/FPS.',
                       'Hook estimates use random one-in-64 samples; driver timing covers generated forwarding methods only.',
                       'Hook access includes lock waits; summing threads does not measure serial frame latency or CPU execution time.',
@@ -85,6 +92,13 @@ def report(directory, metadata_directory=None):
                 controlGpuMs=stats([row['control'] for row in complete]),
                 candidateGpuMs=stats([row['candidate'] for row in complete]),
                 candidateToControl=stats([row['candidate'] / row['control'] for row in complete if row['control'] > 0])))
+    for (thread, api), rows in apis.items():
+        calls = sum(int(row['calls']) for row in rows)
+        total = sum(float(row['total_ms']) for row in rows)
+        result['cudaApi'].append(dict(thread=thread, api=api, calls=calls, totalMs=total,
+            meanMs=total / calls, maxMs=max(float(row['max_ms']) for row in rows),
+            periodMs=sum(float(row['period_ms']) for row in rows)))
+    result['cudaApi'].sort(key=lambda row: row['totalMs'], reverse=True)
     for thread, rows in hook_threads.items():
         sums = {name: sum(float(row[name]) for row in rows) for name in
                 ('period_ms', 'calls', 'samples', 'marked', 'access_sum_ms', 'capture_sum_ms', 'driver_sum_ms')}
@@ -118,7 +132,7 @@ def main():
     parser.add_argument('--metadata-directory', type=pathlib.Path, help='llvm-readobj --notes output for native code objects')
     args = parser.parse_args()
     result = report(args.directory, args.metadata_directory)
-    if not any(result[key] for key in ('cpuStages', 'kernels', 'commandStages', 'commandHookSamples', 'replayPairs')):
+    if not any(result[key] for key in ('cpuStages', 'kernels', 'commandStages', 'commandHookSamples', 'replayPairs', 'cudaApi')):
         parser.error('no complete stage or HIP-event records found')
     if args.output:
         args.output.write_text(json.dumps(result, indent=2) + '\n', encoding='utf-8')
@@ -126,6 +140,8 @@ def main():
         print(f"CPU {name}: n={row['samples']} mean={row['mean']:.3f} ms median={row['median']:.3f} ms p95={row['p95']:.3f} ms")
     for name, row in result['commandStages'].items():
         print(f"COMMAND {name}: n={row['samples']} mean={row['mean']:.3f} ms median={row['median']:.3f} ms p95={row['p95']:.3f} ms")
+    for row in result['cudaApi'][:20]:
+        print(f"CUDA API {row['api']} {row['thread']}: n={row['calls']} mean={row['meanMs']:.6f} ms total={row['totalMs']:.3f} ms max={row['maxMs']:.3f} ms")
     for row in result['commandHookSamples']:
         print(f"HOOK thread={row['thread']} samples={int(row['samples'])} access={row['accessMeanUs']:.3f} us "
               f"capture={row['captureMeanUsPerSample']:.3f} us estimated_cpu={row['estimatedAccessCaptureMsPerSecond']:.3f} ms/s")