diff --git a/README.md b/README.md index 758f5cd..6f9695d 100644 --- a/README.md +++ b/README.md @@ -17,8 +17,12 @@ output match the FP16 baseline exactly after resource-format quantisation. The build incorporates upstream RDNA4 commit `dbef4b2`. Patched OptiScaler now passes standalone K/M native-network tests with exact RGB agreement against the direct-d4r baseline; see [the source patch](patches/optiscaler/README.md). -Game integration and performance profiling remain in progress; no game support is -claimed yet. See [build instructions and hardware results](docs/windows-rdna4-port.md). +Silent Hill 2 now renders K through the native Windows backend; M also renders +a saved gameplay level with GPU output NaN/Inf checks. Startup depth/stencil +and swap-chain lifetime issues found in the game have regression checks. +Performance and wider resolution coverage remain in progress. See +[the reversible game package](docs/windows-game.md) and +[build instructions and hardware results](docs/windows-rdna4-port.md). See [supported games](SUPPORTED_GAMES.md) for the tested games and DLSS models. diff --git a/docs/windows-game.md b/docs/windows-game.md index a51ea1f..64b7c57 100644 --- a/docs/windows-game.md +++ b/docs/windows-game.md @@ -1,8 +1,10 @@ # Native Windows RDNA4 game development package Validated hardware: Windows 11 x64 / Radeon RX 9070 XT / gfx1201. Native K and -M transformer kernels pass standalone D3D12 + OptiScaler checks. Actual game -compatibility and performance are still being validated. This is a local +M transformer kernels pass standalone D3D12 + OptiScaler checks. Silent Hill 2 +K now renders real menu frames and M renders a saved gameplay level with +GPU NaN/Inf checks. Resolution coverage and performance are still being +validated. This is a local development package; no NVIDIA proprietary DLL is included. Build the shim/diagnostics, corrected ZLUDA, K/M native objects and patched @@ -39,6 +41,11 @@ original bytes after the test. It does not change progression saves. First-use PTX compilation can be slow; `-CacheDirectory ` reuses modules with matching GPU, codegen switches and ZLUDA binary fingerprint. +Add `-ValidateOutput` to scan the entire RGBA16F DLSS output for NaN/Inf on +the GPU before returning it to D3D12. Only two diagnostic counters (8 bytes) +are read by the CPU; image data remains in VRAM. Validation failures are +reported explicitly. Disable this diagnostic when measuring production speed. + Original files are backed up before replacement. Restore with: ```powershell diff --git a/docs/windows-rdna4-port.md b/docs/windows-rdna4-port.md index ade9be6..925fe95 100644 --- a/docs/windows-rdna4-port.md +++ b/docs/windows-rdna4-port.md @@ -4,11 +4,40 @@ Branch: `windows-rdna4`. Target: Windows 11 x64, RX 9070 XT, **gfx1201**. Priority: correct K, correct M, native Windows, same-frame output, then speed. This is a work log for an implementation in progress. Standalone DLSS K/M works on the real Windows GPU, including the patched OptiScaler frontend. -An actual game session remains unverified. +Silent Hill 2 now renders through the Windows backend; gameplay and output +validation coverage are recorded below. Performance remains in progress. ## Milestone and gates -First actual game runs (2026-10-01): Silent Hill 2 1.2, XeSS input -> OptiScaler +Actual K rendering: `silent-hill2-k-depth-plane.zip` completes **1794 DLSS +frames / 19734 native K launches / zero d4r failures** at 753x424 -> 1280x720 +with the game's depth-inverted HDR path and DXGI-19 depth plane. The captured +menu is rendered correctly; no CPU image copies or frame-age fallback are +used. Its restored graphics settings match the original SHA256. + +GPU output validation is now available via `D4R_VALIDATE_OUTPUT=1` or the +game runner's `-ValidateOutput`. A native GPU scan classifies all canonical +RGBA16F components, returning only two 32-bit NaN/Inf counters. A negative +fixture with one NaN and two infinities passes. All four K/M standalone +validated frames have zero NaN/Inf and bit-exact RGB against their control +frames (`optiscaler-k-output-validation.zip`, `optiscaler-m-output-validation.zip`). +The pixel format tests including the negative fixture pass legacy/enhanced +barriers (`output-validation-formats-ctest.log`). + +`silent-hill2-m-output-validation` now renders the actual saved gameplay +level with all five native M transformer layers. A screenshot shows the +character and street scene; 1349 full 1280x720 output scans have zero +NaN/Inf in all 3686400 components. There are 1348 complete frame log records, +13490 native launches and zero d4r failures; the bounded stop interrupts the +last frame log after its successful scan. The runner now counts only complete +records and reports validation counts and per-kernel native/translated hits. +The current-frame backend remains slow: +observed interop + NGX CPU completion is roughly 74–86 ms in these runs. +These are initial stage measurements, not isolated kernel GPU timings or +a controlled K/M performance comparison. Profile before optimizing. + +First actual game runs (2026-10-01): Silent Hill 2 v1.1.258834 (the GSA settings +schema reports 1.2.0.0), XeSS input -> OptiScaler DLSS K. `silent-hill2-k-first.zip` found the old Project ID wrapper initializing CUDA NGX with app 4919 instead of the caller's identity, and back-buffer tracking references blocking `ResizeBuffers` (DXGI_ERROR_INVALID_CALL). @@ -53,7 +82,7 @@ restores original DLL/config sentinel bytes exactly. For Silent Hill 2, files; this avoids modifying progression saves. Original settings after the stopped failing run have an identical SHA256. `-CacheDirectory` can reuse an already validated local JIT cache. Development package usage is documented in -`docs/windows-game.md`; in-game success is still a gate. +`docs/windows-game.md`; K menu and M saved-level results are recorded above. Startup coverage (2026-10-01): OptiScaler patch `0002` installs d4r at device discovery, before NGX initialization. `optiscaler-k-preinit-signature.zip` and @@ -95,7 +124,7 @@ Do not integrate NGX until the integer PTX workload is stable on the real GPU. | M6: K layers, full transformer and image validation | PASS: 11 real-weight layers, 44 native launches, four D3D12 frames bit-exact against native CUDA, 81.85–92.91 dB versus translated K | | M7: M FP16-equivalent baseline and full transformer | PASS: all blocks of 40 native launches exactly match NumPy; targeted independent PTX regression matches; four D3D12 frames bit-exact against corrected translated M | | M8: standalone Windows D3D12 NGX harness | PASS for ordinary recorded EvaluateFeature and explicit boundaries: K/M, four frames each, no shim CPU image copies; both bit-exact against explicit boundary baselines | -| M9: OptiScaler integration, profiling, installation | Patched frontend K/M PASS; game, profiling and final package pending | +| M9: OptiScaler integration, profiling, installation | Patched frontend K/M PASS; K game menu and M saved level PASS; profiling and final package pending | ## Public NGX identity and OptiScaler gate (2026-10-01) @@ -137,7 +166,8 @@ NVIDIA architecture as a proxy. GetFeatureRequirements spoofing is bypassed. Explicit external DLSS selection returns errors instead of silently using FSR. Microsoft Build Tools 17.14 / MSVC 14.44.35207 and SDK 10.0.26100 were installed from the signature-verified official bootstrapper without restarting Windows. -No game files have been modified. +At this historical frontend milestone no game files had been modified; the +later game installer and reversible file backups are described above. The OptiScaler fork commit is `33bbac2` on `windows-rdna4-d4r`. Its reproducible source patch and build instructions are in [patches/optiscaler](../patches/optiscaler/README.md). diff --git a/scripts/windows/capture-game-window.ps1 b/scripts/windows/capture-game-window.ps1 index e99dee5..d09d1e9 100644 --- a/scripts/windows/capture-game-window.ps1 +++ b/scripts/windows/capture-game-window.ps1 @@ -15,10 +15,13 @@ public static class D4RWindowCapture { [DllImport("user32.dll")] public static extern bool GetWindowRect(IntPtr window, out Rect rect); [DllImport("user32.dll",CharSet=CharSet.Unicode)] public static extern int GetWindowText(IntPtr window, StringBuilder text, int count); [DllImport("user32.dll")] public static extern bool IsWindowVisible(IntPtr window); + [DllImport("user32.dll")] public static extern IntPtr SetThreadDpiAwarenessContext(IntPtr context); } '@ } $windows = [Collections.Generic.List[object]]::new() +$previousDpi=[D4RWindowCapture]::SetThreadDpiAwarenessContext([IntPtr](-4)) +try { $callback = [D4RWindowCapture+Callback]{ param($window,$data) [uint32]$owner = 0 [void][D4RWindowCapture]::GetWindowThreadProcessId($window,[ref]$owner) @@ -42,5 +45,10 @@ $graphics = [Drawing.Graphics]::FromImage($bitmap) try { $graphics.CopyFromScreen($rect.Left,$rect.Top,0,0,$bitmap.Size) $bitmap.Save([IO.Path]::GetFullPath($OutputPath),[Drawing.Imaging.ImageFormat]::Png) -} finally { $graphics.Dispose(); $bitmap.Dispose() } +} finally { + $graphics.Dispose(); $bitmap.Dispose() +} Write-Host "Captured game window: $OutputPath" +} finally { + if ($previousDpi -ne [IntPtr]::Zero) { [void][D4RWindowCapture]::SetThreadDpiAwarenessContext($previousDpi) } +} diff --git a/scripts/windows/windows-game.ps1 b/scripts/windows/windows-game.ps1 index 0b0db2f..6b44f91 100644 --- a/scripts/windows/windows-game.ps1 +++ b/scripts/windows/windows-game.ps1 @@ -7,6 +7,7 @@ param( [int]$RunSeconds = 0, [string]$DiagnosticResolution, [string]$CacheDirectory = $env:ZLUDA_CACHE_DIR, + [switch]$ValidateOutput, [string[]]$GameArguments = @('-dx12','-windowed','-ResX=1280','-ResY=720'), [string]$OutputDirectory ) @@ -132,6 +133,7 @@ $settings = @{ ZLUDA_CACHE_DIR=$CacheDirectory; PATH=((GamePath 'd4r/hip/bin') + ';' + (GamePath 'd4r/zluda') + ';' + $env:PATH) } $old = @{}; $process = $null; $originalSettings = @() +if ($ValidateOutput) { $settings['D4R_VALIDATE_OUTPUT']='1' } try { if ($DiagnosticResolution) { if ([IO.Path]::GetFileName($GameExe) -ne 'SHProto-Win64-Shipping.exe' -or $DiagnosticResolution -notmatch '^([0-9]{3,4})x([0-9]{3,4})$') { throw 'DiagnosticResolution currently supports Silent Hill 2 only; use e.g. 1280x720.' } @@ -179,11 +181,22 @@ 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]+)') + # 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?$') + $nonfinite = @($checks | Where-Object { $_.Groups[2].Value -ne '0' -or $_.Groups[3].Value -ne '0' }).Count + $launches = [regex]::Matches($stderr, '\[d4r-launch\] kernel="([^"]+)" backend=(native|translated)') + $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; nativeLaunches=([regex]::Matches($stderr, '\[d4r-launch\].*backend=native')).Count; - completedFrames=([regex]::Matches($stdout, 'D4R_FRAME ')).Count; + completedFrames=$frames.Count; outputValidationRequested=[bool]$ValidateOutput; + outputGpuChecks=$checks.Count; nonfiniteOutputs=$nonfinite; kernels=$kernels; + previousFrameOutputs=@($frames | Where-Object { $_.Groups[2].Value -ne '0' }).Count; + cpuImageCopyFrames=@($frames | Where-Object { $_.Groups[1].Value -ne '0' }).Count; failures=([regex]::Matches($stderr, 'D4R_WINDOWS_FAILURE|D4R_COMMAND_FAILURE')).Count} | - ConvertTo-Json | Set-Content -LiteralPath (Join-Path $OutputDirectory 'summary.json') -Encoding UTF8 + ConvertTo-Json -Depth 6 | Set-Content -LiteralPath (Join-Path $OutputDirectory 'summary.json') -Encoding UTF8 Compress-Archive -LiteralPath $OutputDirectory -DestinationPath ($OutputDirectory + '.zip') -Force Write-Host "Game diagnostic bundle: $OutputDirectory.zip" } finally { diff --git a/tools/windows/d3d12_pixel_probe.cpp b/tools/windows/d3d12_pixel_probe.cpp index 2fca777..6870999 100644 --- a/tools/windows/d3d12_pixel_probe.cpp +++ b/tools/windows/d3d12_pixel_probe.cpp @@ -172,6 +172,14 @@ int main(int argc, char** argv) { try { Args args(argc, argv); HipApi hip(args.hip_root); hipDeviceProp_t props{}; hip.select_gfx1201(args.device, props); ExternalApi external{hip}; PixelProgram program(hip, wide(args.module)); + { + PixelAllocation diagnostic(hip,3,1,8); + const uint16_t bad[] = {0x7e00,0x7c00,0xfc00,0x3c00,0,1,0x3ff,0x7bff,0x8000,0x8001,0xbc00,0x400}; + hip.check(hip.hipMemcpy(diagnostic.data,bad,sizeof(bad),hipMemcpyHostToDevice), "Upload nonfinite validation fixture"); + const auto counts = program.validate(diagnostic.data,diagnostic.pitch,3,1); + if (counts[0] != 1 || counts[1] != 2) throw std::runtime_error("GPU NaN/Inf validation counters mismatch"); + std::printf("PASS OUTPUT_VALIDATION_FIXTURE nan=1 inf=2 finite=9\n"); + } ComPtr factory; dx(CreateDXGIFactory1(IID_PPV_ARGS(factory.GetAddressOf())), "Pixel DXGI"); ComPtr adapter; for (UINT i = 0; ; ++i) { dx(factory->EnumAdapters1(i, adapter.ReleaseAndGetAddressOf()), "Pixel adapter"); DXGI_ADAPTER_DESC1 d{}; dx(adapter->GetDesc1(&d), "Pixel adapter desc"); if (!std::memcmp(&d.AdapterLuid, props.luid, sizeof(LUID))) break; } diff --git a/tools/windows/ngx_windows_runtime.h b/tools/windows/ngx_windows_runtime.h index 169d926..6cce4f3 100644 --- a/tools/windows/ngx_windows_runtime.h +++ b/tools/windows/ngx_windows_runtime.h @@ -322,6 +322,13 @@ private: output.image->download_device(reinterpret_cast(output.canonical ? output.canonical->data : output.shared->mapped), output.canonical ? output.canonical->pitch : output.shared->footprint.Footprint.RowPitch); rt_->cuda.check(rt_->cuda.cuCtxSynchronize(), "VRAM output completion"); + if (std::getenv("D4R_VALIDATE_OUTPUT")) { + if (!rt_->pixels) rt_->pixels = std::make_unique(rt_->hip, pixel_module_path()); + const auto counts = rt_->pixels->validate(output.canonical ? output.canonical->data : output.shared->mapped, + output.canonical ? output.canonical->pitch : output.shared->footprint.Footprint.RowPitch, unsigned(output.desc.Width), output.desc.Height); + std::printf("D4R_OUTPUT_VALIDATION elements=%llu nan=%u inf=%u diagnostics_cpu_bytes=8\n", output.desc.Width * output.desc.Height * 4, counts[0], counts[1]); + if (counts[0] || counts[1]) throw std::runtime_error("DLSS output contains NaN/Inf"); + } 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); diff --git a/tools/windows/pixel_conversion.h b/tools/windows/pixel_conversion.h index df2fc6f..3cc9e0d 100644 --- a/tools/windows/pixel_conversion.h +++ b/tools/windows/pixel_conversion.h @@ -1,6 +1,7 @@ #pragma once #include "d3d12_external.h" #include "pixel_format.h" +#include namespace d4r::win { struct PixelSpec { @@ -52,6 +53,8 @@ class PixelProgram { diag::HipApi& hip_; hipModule_t module_ = nullptr; hipFunction_t decode_ = nullptr, encode_ = nullptr; + hipFunction_t validate_ = nullptr; + void* counters_ = nullptr; public: PixelProgram(diag::HipApi& hip, const std::filesystem::path& module) : hip_(hip) { try { @@ -60,8 +63,20 @@ public: hip.check(hip.hipModuleGetFunction(&encode_, module_, "d4r_pixel_encode"), "Resolve GPU pixel encode"); } catch (...) { if (module_) (void)hip.hipModuleUnload(module_); throw; } } - ~PixelProgram() { if (module_) (void)hip_.hipModuleUnload(module_); } + ~PixelProgram() { if (counters_) (void)hip_.hipFree(counters_); if (module_) (void)hip_.hipModuleUnload(module_); } PixelProgram(const PixelProgram&) = delete; + std::array validate(const void* source, uint64_t pitch, unsigned width, unsigned height) { + if (!width || !height || uint64_t(width) * height > UINT32_MAX / 4) throw std::runtime_error("Invalid output validation dimensions"); + if (!validate_) hip_.check(hip_.hipModuleGetFunction(&validate_, module_, "d4r_pixel_validate"), "Resolve output validation kernel"); + if (!counters_) hip_.check(hip_.hipMalloc(&counters_, 2 * sizeof(unsigned)), "Allocate validation counters"); + std::array result{}; + hip_.check(hip_.hipMemcpy(counters_, result.data(), sizeof(result), hipMemcpyHostToDevice), "Clear validation counters"); + void* args[] = {&source, &pitch, &width, &height, &counters_}; + hip_.check(hip_.hipModuleLaunchKernel(validate_, unsigned((uint64_t(width) * height + 255) / 256),1,1,256,1,1,0,nullptr,args,nullptr), "Validate output on GPU"); + hip_.check(hip_.hipDeviceSynchronize(), "Output validation completion"); + hip_.check(hip_.hipMemcpy(result.data(), counters_, sizeof(result), hipMemcpyDeviceToHost), "Read validation counters"); + return result; + } void convert(bool encode, const void* source, uint64_t sourcePitch, void* dest, uint64_t destPitch, unsigned width, unsigned height, pixel::Storage format, unsigned plane) { if (!width || !height || uint64_t(width) * height > UINT32_MAX) throw std::runtime_error("Invalid GPU conversion dimensions"); diff --git a/tools/windows/pixel_convert_gfx1201.hip b/tools/windows/pixel_convert_gfx1201.hip index b4d8907..a72d67f 100644 --- a/tools/windows/pixel_convert_gfx1201.hip +++ b/tools/windows/pixel_convert_gfx1201.hip @@ -131,3 +131,17 @@ extern "C" __global__ void d4r_pixel_encode(const u8* source, u64 source_pitch, float value[4]; for (unsigned c = 0; c < 4; ++c) value[c] = from_half(halves[c]); encode(dest + u64(y) * dest_pitch + u64(x) * pixel_bytes(Storage(format)), Storage(format), value); } + +// Diagnostic scan of the canonical RGBA16F output; no image readback. +extern "C" __global__ void d4r_pixel_validate(const u8* source, u64 pitch, + unsigned width, unsigned height, unsigned* counters) { + unsigned i = blockIdx.x * 256u + threadIdx.x; + if (u64(i) >= u64(width) * height) return; + const auto* values = reinterpret_cast(source + u64(i / width) * pitch) + (i % width) * 4; + unsigned nans = 0, infinities = 0; + for (unsigned c = 0; c < 4; ++c) if ((values[c] & 0x7c00u) == 0x7c00u) { + if (values[c] & 0x3ffu) ++nans; else ++infinities; + } + if (nans) __atomic_fetch_add(counters, nans, __ATOMIC_RELAXED); + if (infinities) __atomic_fetch_add(counters + 1, infinities, __ATOMIC_RELAXED); +}