diff --git a/docs/windows-game.md b/docs/windows-game.md index fbbfcd8..84da57b 100644 --- a/docs/windows-game.md +++ b/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 diff --git a/docs/windows-performance.md b/docs/windows-performance.md index 501a2ba..9343deb 100644 --- a/docs/windows-performance.md +++ b/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 diff --git a/docs/windows-rdna4-port.md b/docs/windows-rdna4-port.md index a8a94d0..510b79d 100644 --- a/docs/windows-rdna4-port.md +++ b/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 diff --git a/patches/zluda/0015-nonserializing-cuda-api-profile.patch b/patches/zluda/0015-nonserializing-cuda-api-profile.patch new file mode 100644 index 0000000..d33259d --- /dev/null +++ b/patches/zluda/0015-nonserializing-cuda-api-profile.patch @@ -0,0 +1,83 @@ +From a1c506fc956cfc347c56c44311669700c0d500ac Mon Sep 17 00:00:00 2001 +From: Codex +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 = 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 { ++ static ENABLED: OnceLock = 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 + diff --git a/scripts/windows/test-async-k.ps1 b/scripts/windows/test-async-k.ps1 new file mode 100644 index 0000000..fbff164 --- /dev/null +++ b/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') } +} diff --git a/scripts/windows/windows-game.ps1 b/scripts/windows/windows-game.ps1 index b1414bf..18a6e3d 100644 --- a/scripts/windows/windows-game.ps1 +++ b/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; diff --git a/tools/windows/d3d12_command_hooks.cpp b/tools/windows/d3d12_command_hooks.cpp index 65d65cc..e8a7cf2 100644 --- a/tools/windows/d3d12_command_hooks.cpp +++ b/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 #include #include @@ -406,6 +407,9 @@ std::mutex queue_metadata_mutex; struct SubmissionLock final : IUnknown { std::atomic references{1}; std::recursive_mutex mutex; + ComPtr 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 device; dx(queue->GetDevice(IID_PPV_ARGS(device.GetAddressOf())), "Queue device(submit drain)"); - ComPtr fence; dx(device->CreateFence(0, D3D12_FENCE_FLAG_NONE, IID_PPV_ARGS(fence.GetAddressOf())), "Queue fence(submit drain)"); - Handle event; event.value = CreateEventW(nullptr, FALSE, FALSE, nullptr); - if (!event.value) throw std::runtime_error("CreateEvent(submit drain)"); - dx(queue->Signal(fence.Get(), 1), "Queue Signal(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"); + auto owner = submission_lock(queue); + std::lock_guard ordered(owner->mutex); + if (!owner->completion_fence) { + ComPtr device; dx(queue->GetDevice(IID_PPV_ARGS(device.GetAddressOf())), "Queue device(submit drain)"); + dx(device->CreateFence(0, D3D12_FENCE_FLAG_NONE, IID_PPV_ARGS(owner->completion_fence.GetAddressOf())), "Queue fence(submit drain)"); } + 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> owners) { + auto owner = submission_lock(queue); + if (!owner->completion_fence) { + ComPtr 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::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); } diff --git a/tools/windows/d3d12_dlss_probe.cpp b/tools/windows/d3d12_dlss_probe.cpp index fd2cb23..75bf48c 100644 --- a/tools/windows/d3d12_dlss_probe.cpp +++ b/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 allocator; + ComPtr list; + ComPtr readback; + }; + std::vector submitted; + std::vector, ComPtr>> 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 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(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 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(raw) + y * textures[3].footprint.Footprint.RowPitch; + native.write(reinterpret_cast(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)}; - dx(readback->Map(0, &range, &raw), "Map(verification only)"); - std::vector 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(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 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(raw) + y * textures[3].footprint.Footprint.RowPitch; - native.write(reinterpret_cast(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"); + if (!burst) validate_output(readback.Get(), frame); + if (recreate && frame + 1 < args.iterations) { + // Release/create without an explicit harness drain. In burst + // mode earlier frames can still be using the old feature. + check(release(handle), "Recreate ReleaseFeature"); featureCleanup.h = nullptr; + dx(creationAllocator->Reset(), "Recreate allocator Reset"); + dx(creationList->Reset(creationAllocator.Get(), nullptr), "Recreate list Reset"); + handle = nullptr; + check(shim.symbol("NVSDK_NGX_D3D12_CreateFeature")(creationList.Get(), 1, parameters, &handle), "Recreate CreateFeature"); + featureCleanup.h = handle; + dx(creationList->Close(), "Recreate list Close"); + ID3D12CommandList* created[] = {creationList.Get()}; queue->ExecuteCommandLists(1, created); + creation_owners.emplace_back(creationAllocator, creationList); + // This empty creation list owns no work; the next recreation + // allocates independent command memory without waiting on GPU. + dx(device->CreateCommandAllocator(D3D12_COMMAND_LIST_TYPE_DIRECT, IID_PPV_ARGS(creationAllocator.ReleaseAndGetAddressOf())), "Next recreate allocator"); + dx(device->CreateCommandList(0, D3D12_COMMAND_LIST_TYPE_DIRECT, creationAllocator.Get(), nullptr, IID_PPV_ARGS(creationList.ReleaseAndGetAddressOf())), "Next recreate list"); + dx(creationList->Close(), "Next recreate initial Close"); + std::printf("D4R_RECREATE_VALIDATION submitted_frame=%u next_feature_created=1\n", frame); } - D3D12_RANGE noWrite{}; 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"); + } + if (burst) { + drain(); + for (unsigned frame = 0; frame < submitted.size(); ++frame) validate_output(submitted[frame].readback.Get(), frame); + std::printf("D4R_BURST_VALIDATION frames=%zu cpu_readback_after_all_submissions=1\n", submitted.size()); } 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", diff --git a/tools/windows/d3d12_external.h b/tools/windows/d3d12_external.h index e24928a..80dabab 100644 --- a/tools/windows/d3d12_external.h +++ b/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 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_; diff --git a/tools/windows/d3d12_retirement.h b/tools/windows/d3d12_retirement.h new file mode 100644 index 0000000..500ac81 --- /dev/null +++ b/tools/windows/d3d12_retirement.h @@ -0,0 +1,46 @@ +#pragma once +#include "d3d12_external.h" +#include +#include +#include +#include +#include + +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 fence; uint64_t value; std::shared_ptr owner; }; + std::mutex mutex_; + std::condition_variable changed_; + std::deque pending_; + GpuRetirement() { + std::thread([this] { + Handle event; event.value = CreateEventW(nullptr, FALSE, FALSE, nullptr); + for (;;) { + Batch batch; + { std::unique_lock 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 owner) { + { std::lock_guard lock(mutex_); pending_.push_back({fence, value, std::move(owner)}); } + changed_.notify_one(); + } +}; +} diff --git a/tools/windows/d4r_nvngx_windows.cpp b/tools/windows/d4r_nvngx_windows.cpp index b455002..157fc3e 100644 --- a/tools/windows/d4r_nvngx_windows.cpp +++ b/tools/windows/d4r_nvngx_windows.cpp @@ -184,11 +184,14 @@ API unsigned NVSDK_NGX_D3D12_ReleaseFeature(Handle* handle) { { std::lock_guard 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 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); diff --git a/tools/windows/ngx_windows_runtime.h b/tools/windows/ngx_windows_runtime.h index 366eaab..11ffd8e 100644 --- a/tools/windows/ngx_windows_runtime.h +++ b/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 +#include +#include #include #include #include @@ -55,6 +59,14 @@ struct Runtime { std::unique_ptr timeline; std::unique_ptr 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 async_tickets; + std::string async_failure; + ComPtr async_queue; ngx::ProjectIdentity project; ngx::FeatureCommonInfo common{}; std::vector 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 queue; + { std::lock_guard 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 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 queue_; + // A boundary drains both copies before returning. Reuse their allocators + // only after that completion; no game allocator is retained here. + ComPtr input_allocator_, output_allocator_; + ComPtr input_list_, output_list_; + struct CopySlot { + ComPtr input_allocator, output_allocator; + ComPtr input, output; + uint64_t done = 0; + }; + struct AsyncFrame { + void* parameters = d4r_ngx_parameters_create(); + std::vector> resources; + SharedTimeline::Ticket ticket{}; + ~AsyncFrame() { if (parameters) d4r_ngx_parameters_destroy(parameters); } + }; + std::vector> copy_slots_; + ComPtr 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> jobs_; + std::thread worker_; + bool worker_stop_ = false, worker_active_ = false; + std::string worker_failure_; struct Plane { ComPtr texture; D3D12_RESOURCE_DESC desc{}; @@ -185,6 +241,7 @@ class Feature { public: unsigned width = 0, height = 0; explicit Feature(std::shared_ptr runtime, void* parameters) : rt_(std::move(runtime)) { + if (std::getenv("D4R_ASYNC_INTEROP")) rt_->drain_async_queue(); std::lock_guard 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 lock(rt_->mutex); cleanup(); } + ~Feature() { + { std::lock_guard 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 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(now - boundaryStage).count()); + boundaryStage = std::chrono::steady_clock::now(); + }; commands::InternalScope internal; std::lock_guard 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 inputAllocator, outputAllocator; ComPtr inputList, outputList; auto list = [&](ComPtr& allocator, ComPtr& 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 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(); + 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 submitted(submit_mutex_); + if (!queue || (queue_ && queue_.Get() != queue)) throw std::runtime_error("Async interop requires one feature queue"); + ComPtr queueDevice; dx(queue->GetDevice(IID_PPV_ARGS(queueDevice.GetAddressOf())), "Async queue device"); + ComPtr 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 lock(jobs_mutex_); + if (!worker_failure_.empty()) throw std::runtime_error(worker_failure_); } + { std::lock_guard 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 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(); + 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 pending(jobs_mutex_); + { std::lock_guard 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 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 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 frame; + { std::unique_lock 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 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 lock(rt_->mutex); + { std::lock_guard 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 lock(jobs_mutex_); worker_failure_ = e.what(); } + { std::lock_guard 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 order(rt_->async_mutex); rt_->async_tickets.pop_front(); } + rt_->async_changed.notify_all(); + frame.reset(); + { std::lock_guard 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(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(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)); } }; } diff --git a/tools/windows/profile_report.py b/tools/windows/profile_report.py index a88beef..de088ec 100644 --- a/tools/windows/profile_report.py +++ b/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")