diff --git a/CMakeLists.txt b/CMakeLists.txt index 8a4862d..1b7a158 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -68,7 +68,7 @@ foreach(family k m) endforeach() endforeach() add_custom_target(d4r_native_gfx1201_compile ALL DEPENDS ${native_objects}) -foreach(probe hip_gfx1201_probe gfx12_wmma_probe k_module_probe native_replay_probe cuda_driver_probe cuda_image_probe d3d12_hip_interop_probe hip_stream_lifecycle_probe ngx_cuda_init_probe) +foreach(probe hip_gfx1201_probe gfx12_wmma_probe k_module_probe native_replay_probe cuda_driver_probe cuda_image_probe d3d12_hip_interop_probe hip_stream_lifecycle_probe ngx_cuda_init_probe d3d12_dlss_probe) add_executable(d4r_${probe} tools/windows/${probe}.cpp) target_compile_definitions(d4r_${probe} PRIVATE WIN32_LEAN_AND_MEAN NOMINMAX __HIP_PLATFORM_AMD__) target_include_directories(d4r_${probe} PRIVATE "${D4R_HIP_ROOT}/include") @@ -84,8 +84,19 @@ foreach(probe hip_gfx1201_probe gfx12_wmma_probe k_module_probe native_replay_pr endif() install(TARGETS d4r_${probe} RUNTIME DESTINATION bin) endforeach() +add_library(d4r_nvngx_windows SHARED tools/windows/d4r_nvngx_windows.cpp) +set_target_properties(d4r_nvngx_windows PROPERTIES PREFIX "" OUTPUT_NAME "d4r_nvngx") +target_compile_definitions(d4r_nvngx_windows PRIVATE WIN32_LEAN_AND_MEAN NOMINMAX __HIP_PLATFORM_AMD__) +target_include_directories(d4r_nvngx_windows PRIVATE "${D4R_HIP_ROOT}/include") +target_link_libraries(d4r_nvngx_windows PRIVATE version dbghelp d3d12 dxguid) +if(MINGW) + target_link_options(d4r_nvngx_windows PRIVATE -static) +endif() +install(TARGETS d4r_nvngx_windows RUNTIME DESTINATION bin) if(MSVC) target_sources(d4r_ngx_cuda_init_probe PRIVATE tools/d4r_ngx_param_msvc.cpp tools/d4r_ngx_param_host.cpp) + target_sources(d4r_d3d12_dlss_probe PRIVATE tools/d4r_ngx_param_msvc.cpp tools/d4r_ngx_param_host.cpp) + target_sources(d4r_nvngx_windows PRIVATE tools/d4r_ngx_param_msvc.cpp tools/d4r_ngx_param_host.cpp) else() find_program(D4R_CLANG_CL NAMES clang-cl.exe PATHS "${D4R_HIP_ROOT}/bin" "${D4R_HIP_ROOT}/lib/llvm/bin" NO_DEFAULT_PATH REQUIRED) set(ngx_parameter_object "${CMAKE_CURRENT_BINARY_DIR}/ngx_parameter_msvc.obj") @@ -95,6 +106,8 @@ else() DEPENDS tools/d4r_ngx_param_msvc.cpp VERBATIM) set_source_files_properties("${ngx_parameter_object}" PROPERTIES GENERATED TRUE EXTERNAL_OBJECT TRUE) target_sources(d4r_ngx_cuda_init_probe PRIVATE "${ngx_parameter_object}" tools/d4r_ngx_param_host.cpp) + target_sources(d4r_d3d12_dlss_probe PRIVATE "${ngx_parameter_object}" tools/d4r_ngx_param_host.cpp) + target_sources(d4r_nvngx_windows PRIVATE "${ngx_parameter_object}" tools/d4r_ngx_param_host.cpp) endif() add_library(d4r_nvapi_trace SHARED tools/windows/nvapi_trace.cpp) set_target_properties(d4r_nvapi_trace PROPERTIES PREFIX "" OUTPUT_NAME "nvapi64" @@ -118,6 +131,7 @@ if(MINGW) endif() install(TARGETS d4r_nvapi_compat RUNTIME DESTINATION nvapi-compat) target_link_libraries(d4r_d3d12_hip_interop_probe PRIVATE d3d12 dxgi dxguid) +target_link_libraries(d4r_d3d12_dlss_probe PRIVATE d3d12 dxgi dxguid) add_dependencies(d4r_hip_gfx1201_probe d4r_probe_code_object) add_dependencies(d4r_gfx12_wmma_probe d4r_wmma_code_object) add_dependencies(d4r_k_module_probe d4r_native_gfx1201_compile) @@ -171,5 +185,8 @@ if(EXISTS "${D4R_ZLUDA_ROOT}/nvcuda.dll") add_test(NAME cuda_images COMMAND d4r_cuda_image_probe --hip-root "${D4R_HIP_ROOT}" --cuda-dll "${D4R_ZLUDA_ROOT}/nvcuda.dll" --iterations 32) set_tests_properties(cuda_images PROPERTIES TIMEOUT 90 LABELS "hardware;zluda;descriptors") + add_test(NAME d3d12_cuda_images COMMAND d4r_d3d12_dlss_probe + --hip-root "${D4R_HIP_ROOT}" --cuda-dll "${D4R_ZLUDA_ROOT}/nvcuda.dll" --interop-mode images --iterations 1) + set_tests_properties(d3d12_cuda_images PROPERTIES TIMEOUT 90 LABELS "hardware;interop;zluda;descriptors") endif() endif() diff --git a/docs/windows-rdna4-port.md b/docs/windows-rdna4-port.md index 17af852..3bd3587 100644 --- a/docs/windows-rdna4-port.md +++ b/docs/windows-rdna4-port.md @@ -10,8 +10,11 @@ Current milestone: M1/M2/M4/M5 passed; native NGX Create/Evaluate on the simple requested E path passes with the user's DLLs. All 11 K layers now pass synthetic and real-weight NumPy/PTX replay checks; native K Evaluate produces four finite frames on Windows. M now passes all five synthetic and real-weight layers against -NumPy and independent PTX, and four native CUDA frames. D3D12 integration remains -in progress; the fully translated M comparison currently exceeds its JIT timeout. +NumPy and independent PTX, and four native CUDA frames. The standalone D3D12 +harness now produces four K and four M frames using imported VRAM and shared +fences, with bit-exact RGB agreement against the corresponding CUDA harnesses. +Game command-list/queue integration remains in progress. The fully translated M +comparison is being rerun with the correct WMMA options and a writable JIT cache. Do not integrate NGX until the integer PTX workload is stable on the real GPU. | Gate | Status | @@ -19,12 +22,12 @@ Do not integrate NGX until the integer PTX workload is stable on the real GPU. | M0: branch, audit, Windows build and diagnostics | Initial implementation complete | | M1: native HIP gfx1201 allocation, kernel, CPU verification | PASS, 32 iterations + guard verification | | M2: Windows ZLUDA integer PTX, primary and created contexts | PASS, 32 iterations each + guard verification | -| M3: Windows NGX initialization and simple DLSS path | Init, SR capabilities, Create/Evaluate PASS; four finite synthetic frames; K/M pending | +| M3: Windows NGX initialization and simple DLSS path | Init, SR capabilities, Create/Evaluate PASS; E, K and M CUDA harnesses execute | | M4: D3D12 / HIP external memory and fence round trip | PASS with TheRock; stable 7.2 has mapped-view leak | | M5: independently validated gfx12 WMMA backend | PASS: raw + legacy adapter + upstream layout, max abs/relative error 0 | | M6: K layers, full transformer and image validation | CUDA harness PASS: 11 real-weight layers, native launches for all layers, four finite frames, 81.85–92.91 dB versus translated K; D3D12 integration pending | | M7: M FP16-equivalent baseline and full transformer | CUDA harness PASS: 5 real-weight layers exactly match NumPy/PTX, 4 finite frames, 40 native Swin launches; D3D12 and full translated M comparison pending | -| M8: standalone Windows D3D12 NGX harness | Pending | +| M8: standalone Windows D3D12 NGX harness | PASS at explicit submitted queue boundaries: K/M, four frames each, no shim CPU image copies, bit-exact RGB against CUDA harness | | M9: OptiScaler integration, profiling, installation | Pending | ## Source baseline (checked 2026-09-29) @@ -166,14 +169,18 @@ Reproduce the layer gates: & scripts/windows/test-m-replay.ps1 & scripts/windows/stage-native-m.ps1 -DlssDll "$PWD/nvngx_dlss.dll" $env:D4R_ZLUDA_NATIVE_DIR="$PWD/build/native-m-gfx1201" -$env:D4R_WMMA='1'; $env:D4R_FP8_WMMA='1'; $env:D4R_FP8_NATIVE='0' +$env:D4R_ZLUDA_WMMA='1'; $env:D4R_ZLUDA_WMMA_FP8='1'; $env:D4R_ZLUDA_WMMA_FP8_NATIVE='0' & scripts/windows/test-windows-rdna4.ps1 -RuntimeProfile therock -NgxOnly -NgxMode evaluate -Preset 13 -Iterations 4 -RequireNativeNetwork -NgxCore "$PWD/_nvngx.dll" -DlssDll "$PWD/nvngx_dlss.dll" ``` -The fully translated M feature creation currently exceeds 600 seconds; its -automatic trace localises this to loading the transformer modules. Native M -does not have this JIT delay. A completed translated image comparison is not -claimed. The user selected Silent Hill 2 at +Earlier translated M attempts used incorrect environment names and an +unwritable default JIT cache. The runner now defaults to the actual +`D4R_ZLUDA_WMMA`, `D4R_ZLUDA_WMMA_FP8`, `D4R_ZLUDA_WMMA_FP8_NATIVE` options +(1, 1, 0), records them, and restores the caller's environment. Patch 0012 adds +absolute `ZLUDA_CACHE_DIR` and reports directory/open failures. The diagnostic +default is `build/zluda-cache-windows`; five Rust cache tests and all eleven +patches applied to the pinned clean upstream pass. A completed translated M +image comparison is not yet claimed. The user selected Silent Hill 2 at `D:\Games\SILENT HILL 2\SHProto\Binaries\Win64\SHProto-Win64-Shipping.exe` for the later game gate. No game files have been changed yet. @@ -186,6 +193,49 @@ is insufficient for a general game integration. The next backend gate is a D3D12 harness with explicit queue boundaries; command-list/queue integration must be validated before the game gate. +### Native D3D12 harness milestone (2026-09-30) + +The new `d4r_nvngx.dll` has a Windows backend independent of the Wine shim. +`d3d12_external.h` owns Win32 handles, shared DEFAULT buffers, HIP mappings and +the D3D12 fence/semaphore timeline. Canonical color/output RGBA16F, motion RG16F, +and depth/exposure R32F remain in VRAM. Texture-to-buffer copies, CUDA array +transfers and output copies execute on the GPU. Harness initialization uploads +synthetic inputs; final readback is confined to its verifier/output exporter. + +The first complete attempt exposed a synchronization race: waiting on a HIP +external semaphore on its nonblocking stream and synchronizing that stream did +not reliably establish completion for work subsequently submitted through other +CUDA streams. Samples read back by the diagnostic could hide the race. The +baseline additionally waits for the producer's actual D3D12 fence event before +CUDA submission. It waits for completion, without transferring image data or +using another frame. `cuCtxSynchronize` covers NGX's internal streams, followed +by HIP fence signaling and a D3D12 queue wait before copying the output texture. +Reducing these host completion waits is a later profiling/ordering task. + +Private results: `test-results/d3d12-k-direct-fence.zip` and +`test-results/d3d12-m-direct-fence.zip`, four finite, nonconstant frames each, +44 native K and 40 native M launches. Every RGB value agrees bit-for-bit with +`ngx-k-native-verified` / `ngx-m-bounded-capture`: max absolute/relative error 0, +PSNR infinite. CTest: 10/10 PASS including a five-plane D3D12-to-CUDA array +roundtrip and pitched CUDA device/array copies with padding sentinels. + +Build/package: `dist/windows-rdna4-d3d12`. Reproduce K (use preset 13 and the +staged `native-m-gfx1201` directory for M): + +```powershell +$env:D4R_ZLUDA_NATIVE_DIR="$PWD/build/native-k-gfx1201" +& scripts/windows/test-windows-rdna4.ps1 -RuntimeProfile therock -PackageRoot "$PWD/dist/windows-rdna4-d3d12" -NgxOnly -NgxMode d3d12 -Preset 11 -Iterations 4 -RequireNativeNetwork -NgxCore "$PWD/_nvngx.dll" -DlssDll "$PWD/nvngx_dlss.dll" +``` + +The runner places the shim beside the two private NVIDIA DLLs in an owned +temporary directory, matching NGX's caller-relative feature lookup. They are +never installed into the public package. `D4R_INTEROP_VERIFY` explicitly enables +extra CPU diagnostics and changes the logged copy flag; keep it unset for the +normal path. The open-command-list game export still rejects execution until +submission ordering is implemented. The explicit harness export requires that +all preceding producers have already been submitted. Other texture formats, +dynamic resolution and real-game resource-state tracking remain unvalidated. + ### K real-weight native Evaluate and correctness baseline The locally supplied DLSS 310.9.1.0 selects all eleven `dltss_pwin_*` functions diff --git a/patches/zluda/0012-configurable-windows-jit-cache.patch b/patches/zluda/0012-configurable-windows-jit-cache.patch new file mode 100644 index 0000000..e72945e --- /dev/null +++ b/patches/zluda/0012-configurable-windows-jit-cache.patch @@ -0,0 +1,65 @@ +diff --git a/zluda_cache/src/lib.rs b/zluda_cache/src/lib.rs +index acd262e..c3ac1be 100644 +--- a/zluda_cache/src/lib.rs ++++ b/zluda_cache/src/lib.rs +@@ -22,18 +22,36 @@ pub struct ModuleCache(SqliteConnection); + + impl ModuleCache { + pub fn create_cache_dir_and_get_path() -> Option { +- let mut cache_dir = dirs::cache_dir()?; +- cache_dir.extend(["zluda", "ComputeCache"]); ++ let override_dir = std::env::var_os("ZLUDA_CACHE_DIR"); ++ let cache_dir = match override_dir { ++ Some(path) => { ++ let path = std::path::PathBuf::from(path); ++ if !path.is_absolute() { ++ eprintln!("[zluda-cache] ZLUDA_CACHE_DIR must be absolute; cache disabled"); ++ return None; ++ } ++ path ++ } ++ None => dirs::cache_dir()?.join("zluda").join("ComputeCache"), ++ }; ++ Self::cache_path(&cache_dir) ++ } ++ ++ fn cache_path(cache_dir: &std::path::Path) -> Option { + // We ensure that the cache directory exists +- std::fs::create_dir_all(&cache_dir).ok()?; ++ if let Err(error) = std::fs::create_dir_all(cache_dir) { ++ eprintln!("[zluda-cache] cannot create {}: {error}; cache disabled", cache_dir.display()); ++ return None; ++ } + // No need to create the file, it will be created by SQLite on first access + // zluda.db might be in use by older versions of ZLUDA +- cache_dir.push("zluda2.db"); +- Some(cache_dir.to_string_lossy().into()) ++ Some(cache_dir.join("zluda2.db").to_string_lossy().into()) + } + + pub fn open(file_path: &str) -> Option { +- let mut conn = SqliteConnection::establish(file_path).ok()?; ++ let mut conn = SqliteConnection::establish(file_path).map_err(|error| { ++ eprintln!("[zluda-cache] cannot open {file_path}: {error}; cache disabled"); ++ }).ok()?; + conn.batch_execute("PRAGMA journal_mode = WAL; PRAGMA synchronous = NORMAL;") + .ok()?; + conn.run_pending_migrations(MIGRATIONS).ok()?; +@@ -86,6 +104,17 @@ mod tests { + use arrayvec::ArrayString; + use diesel::prelude::*; + ++ #[test] ++ fn explicit_directory_creates_persistent_database() { ++ let directory = std::env::temp_dir().join(format!("zluda-cache-test-{}-{}", std::process::id(), ModuleCache::time_now())); ++ let path = ModuleCache::cache_path(&directory).unwrap(); ++ assert_eq!(std::path::Path::new(&path), directory.join("zluda2.db")); ++ let cache = ModuleCache::open(&path).unwrap(); ++ drop(cache); ++ assert!(std::path::Path::new(&path).is_file()); ++ std::fs::remove_dir_all(directory).unwrap(); ++ } ++ + #[derive(Queryable, Selectable)] + #[diesel(table_name = crate::schema::modules)] + #[diesel(check_for_backend(diesel::sqlite::Sqlite))] diff --git a/scripts/windows/build-zluda-windows.ps1 b/scripts/windows/build-zluda-windows.ps1 index 1aa98f9..862b271 100644 --- a/scripts/windows/build-zluda-windows.ps1 +++ b/scripts/windows/build-zluda-windows.ps1 @@ -1,5 +1,5 @@ [CmdletBinding()] -param([string]$SourceRoot, [string]$LlvmBuildRoot, [string]$InstallRoot, [switch]$FetchOnly, [int]$Jobs = 4) +param([string]$SourceRoot, [string]$LlvmBuildRoot, [string]$InstallRoot, [switch]$FetchOnly, [switch]$CacheTestsOnly, [int]$Jobs = 4) $ErrorActionPreference = 'Stop' $repo = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) if (!$SourceRoot) { $SourceRoot = Join-Path $repo 'external/ZLUDA' } @@ -33,6 +33,14 @@ Push-Location $SourceRoot try { $cargo = Join-Path $rust 'bin/cargo.exe' if (!(Test-Path $cargo)) { throw 'Run setup-rust-toolchain.ps1 first.' } + if ($CacheTestsOnly) { + $ErrorActionPreference = 'Continue' + & $cargo test --locked --release --target x86_64-pc-windows-gnu -p zluda_cache ` + 2>&1 | Tee-Object (Join-Path $target 'cache-tests.log') + $ErrorActionPreference = 'Stop' + if ($LASTEXITCODE) { throw "ZLUDA cache regression tests failed ($LASTEXITCODE)" } + return + } if ($FetchOnly) { $ErrorActionPreference = 'Continue' & $cargo fetch --locked --target x86_64-pc-windows-gnu 2>&1 | Tee-Object (Join-Path $target 'fetch.log') diff --git a/scripts/windows/test-windows-rdna4.ps1 b/scripts/windows/test-windows-rdna4.ps1 index 29e9e77..4139f91 100644 --- a/scripts/windows/test-windows-rdna4.ps1 +++ b/scripts/windows/test-windows-rdna4.ps1 @@ -7,7 +7,7 @@ param( [string]$OutputDirectory, [string]$NgxCore = $env:D4R_NGX_CORE, [string]$DlssDll = $env:D4R_DLSS_DLL, - [ValidateSet('init', 'evaluate')][string]$NgxMode = 'init', + [ValidateSet('init', 'evaluate', 'd3d12')][string]$NgxMode = 'init', [ValidateSet(5, 11, 13)][int]$Preset = 11, [switch]$NgxOnly, [switch]$Trace, @@ -51,6 +51,11 @@ $originalCuda = $env:ZLUDA_CUDA_LIB $originalNvapi = $env:D4R_NVAPI_BACKEND $originalPythonPath = $env:PYTHONPATH $originalVerbose = $env:D4R_ZLUDA_VERBOSE +$originalCache = $env:ZLUDA_CACHE_DIR +$originalCodegen = @{} +foreach ($setting in @('D4R_ZLUDA_WMMA','D4R_ZLUDA_WMMA_FP8','D4R_ZLUDA_WMMA_FP8_NATIVE')) { + $originalCodegen[$setting] = [Environment]::GetEnvironmentVariable($setting,'Process') +} $exitStatus = 1 $ngxRuntimeDirectory = $null $summary = [ordered]@{utc=[DateTime]::UtcNow.ToString('o'); profile=$RuntimeProfile; hipRoot=$HipRoot; zludaRoot=$ZludaRoot; tests=@()} @@ -96,9 +101,19 @@ function Invoke-Probe([string]$Name, [string]$Exe, [string[]]$Arguments) { return $pass } try { + if (!$env:ZLUDA_CACHE_DIR) { $env:ZLUDA_CACHE_DIR = Join-Path $repoRoot 'build/zluda-cache-windows' } + $summary.zludaCacheDirectory = $env:ZLUDA_CACHE_DIR + foreach ($setting in @('D4R_ZLUDA_WMMA','D4R_ZLUDA_WMMA_FP8','D4R_ZLUDA_WMMA_FP8_NATIVE')) { + if ($null -eq $originalCodegen[$setting]) { + [Environment]::SetEnvironmentVariable($setting, $(if ($setting -eq 'D4R_ZLUDA_WMMA_FP8_NATIVE') { '0' } else { '1' }), 'Process') + } + } + $summary.codegen = @{} + foreach ($setting in $originalCodegen.Keys) { $summary.codegen[$setting] = [Environment]::GetEnvironmentVariable($setting,'Process') } + Write-Host "ZLUDA codegen: WMMA=$env:D4R_ZLUDA_WMMA FP8 widening=$env:D4R_ZLUDA_WMMA_FP8 native FP8=$env:D4R_ZLUDA_WMMA_FP8_NATIVE" if ($NgxOnly -and (!$NgxCore -or !$DlssDll)) { throw '-NgxOnly requires both locally supplied NVIDIA DLL paths.' } if ($RequireNativeNetwork) { - if ($NgxMode -ne 'evaluate' -or $Preset -notin @(11,13) -or !$env:D4R_ZLUDA_NATIVE_DIR) { + if ($NgxMode -notin @('evaluate','d3d12') -or $Preset -notin @(11,13) -or !$env:D4R_ZLUDA_NATIVE_DIR) { throw '-RequireNativeNetwork requires K/M Evaluate and D4R_ZLUDA_NATIVE_DIR.' } $env:D4R_ZLUDA_VERBOSE = '1' @@ -209,7 +224,8 @@ try { $ngxExe = Join-Path $ngxRuntimeDirectory 'd4r_ngx_cuda_init_probe.exe' $localCore = Join-Path $ngxRuntimeDirectory '_nvngx.dll' $localDlss = Join-Path $ngxRuntimeDirectory 'nvngx_dlss.dll' - Copy-Item -LiteralPath (Join-Path $bin 'd4r_ngx_cuda_init_probe.exe') -Destination $ngxExe + $probeName = if ($NgxMode -eq 'd3d12') { 'd4r_d3d12_dlss_probe.exe' } else { 'd4r_ngx_cuda_init_probe.exe' } + Copy-Item -LiteralPath (Join-Path $bin $probeName) -Destination $ngxExe Copy-Item -LiteralPath $NgxCore -Destination $localCore Copy-Item -LiteralPath $DlssDll -Destination $localDlss $selectedCuda = if ($Trace) { Join-Path $ZludaRoot 'trace/nvcuda.dll' } else { Join-Path $ZludaRoot 'nvcuda.dll' } @@ -219,7 +235,15 @@ try { '--ngx-core', $localCore, '--dlss-dll', $localDlss, '--nvapi-dll', (Join-Path $PackageRoot 'nvapi-compat/nvapi64.dll'), '--ngx-mode', $NgxMode, '--preset', "$Preset", '--iterations', "$Iterations") - $ngxName = if ($NgxMode -eq 'evaluate') { "ngx-evaluate-preset-$Preset" } else { 'ngx-init' } + if ($NgxMode -eq 'd3d12') { + # The driver NGX core resolves feature DLLs relative to its + # caller module as well as the executable. Keep the shim + # alongside the private user-supplied DLLs for this test. + $localShim = Join-Path $ngxRuntimeDirectory 'd4r_nvngx.dll' + Copy-Item -LiteralPath (Join-Path $bin 'd4r_nvngx.dll') -Destination $localShim + $ngxArguments += @('--module', $localShim) + } + $ngxName = if ($NgxMode -eq 'evaluate') { "ngx-evaluate-preset-$Preset" } elseif ($NgxMode -eq 'd3d12') { "d3d12-evaluate-preset-$Preset" } else { 'ngx-init' } $ngxOk = Invoke-Probe $ngxName $ngxExe $ngxArguments if ($ngxOk -and $RequireNativeNetwork) { $hits = @() @@ -238,12 +262,14 @@ try { } if (!$ngxOk) { $exitStatus = 1 - if (Test-Path (Join-Path $ZludaRoot 'trace/nvcuda.dll')) { - Invoke-Probe "$ngxName-trace" $ngxExe @( + if (!$Trace -and (Test-Path (Join-Path $ZludaRoot 'trace/nvcuda.dll'))) { + $traceArguments = @( '--hip-root', $HipRoot, '--cuda-dll', (Join-Path $ZludaRoot 'trace/nvcuda.dll'), '--ngx-core', $localCore, '--dlss-dll', $localDlss, '--nvapi-dll', (Join-Path $PackageRoot 'nvapi-compat/nvapi64.dll'), - '--ngx-mode', $NgxMode, '--preset', "$Preset", '--iterations', '1') | Out-Null + '--ngx-mode', $NgxMode, '--preset', "$Preset", '--iterations', '1') + if ($NgxMode -eq 'd3d12') { $traceArguments += @('--module', $localShim) } + Invoke-Probe "$ngxName-trace" $ngxExe $traceArguments | Out-Null } } } @@ -272,6 +298,8 @@ try { $env:D4R_NVAPI_BACKEND = $originalNvapi $env:PYTHONPATH = $originalPythonPath $env:D4R_ZLUDA_VERBOSE = $originalVerbose + $env:ZLUDA_CACHE_DIR = $originalCache + foreach ($setting in $originalCodegen.Keys) { [Environment]::SetEnvironmentVariable($setting, $originalCodegen[$setting], 'Process') } $archive = "$OutputDirectory.zip" Compress-Archive -LiteralPath $OutputDirectory -DestinationPath $archive -Force Write-Host "Diagnostic bundle: $archive" diff --git a/tools/windows/cuda_image_api.h b/tools/windows/cuda_image_api.h index 441779f..16811bd 100644 --- a/tools/windows/cuda_image_api.h +++ b/tools/windows/cuda_image_api.h @@ -131,5 +131,21 @@ public: copy.widthBytes = rowBytes_; copy.height = height_; api_.cuda.check(api_.cuMemcpy2D_v2(©), "cuMemcpy2D_v2(diagnostic array to host)"); } + void upload_device(CUdeviceptr source, size_t pitch) { + if (pitch < rowBytes_) throw std::runtime_error("Device input pitch is smaller than its row"); + Copy2D copy{}; + copy.srcType = 2; copy.srcDevice = source; copy.srcPitch = pitch; + copy.dstType = 3; copy.dstArray = array_; + copy.widthBytes = rowBytes_; copy.height = height_; + api_.cuda.check(api_.cuMemcpy2D_v2(©), "cuMemcpy2D_v2(external VRAM to array)"); + } + void download_device(CUdeviceptr destination, size_t pitch) { + if (pitch < rowBytes_) throw std::runtime_error("Device output pitch is smaller than its row"); + Copy2D copy{}; + copy.srcType = 3; copy.srcArray = array_; + copy.dstType = 2; copy.dstDevice = destination; copy.dstPitch = pitch; + copy.widthBytes = rowBytes_; copy.height = height_; + api_.cuda.check(api_.cuMemcpy2D_v2(©), "cuMemcpy2D_v2(array to external VRAM)"); + } }; } diff --git a/tools/windows/cuda_image_probe.cpp b/tools/windows/cuda_image_probe.cpp index 25a972c..de524fa 100644 --- a/tools/windows/cuda_image_probe.cpp +++ b/tools/windows/cuda_image_probe.cpp @@ -38,10 +38,30 @@ int main(int argc, char** argv) { image.upload(input.data()); image.download(output.data()); if (input != output) throw std::runtime_error("CUDA array storage roundtrip mismatch"); + const size_t rowBytes = 17 * channels * sizeof(uint16_t), pitch = (rowBytes + 255) & ~size_t(255); + CUdeviceptr devicePixels = 0; + cuda.check(cuda.cuMemAlloc_v2(&devicePixels, pitch * 9), "cuMemAlloc_v2(pitched image regression)"); + struct PixelsCleanup { CudaApi& cuda; CUdeviceptr pointer; ~PixelsCleanup() { (void)cuda.cuMemFree_v2(pointer); } } pixelsCleanup{cuda, devicePixels}; + std::vector padded(pitch * 9, 0xcd); + for (unsigned y = 0; y < 9; ++y) std::memcpy(padded.data() + y * pitch, input.data() + y * 17 * channels, rowBytes); + cuda.check(cuda.cuMemcpyHtoD_v2(devicePixels, padded.data(), padded.size()), "cuMemcpyHtoD_v2(pitched pattern)"); + image.upload_device(devicePixels, pitch); + image.download(output.data()); + if (input != output) throw std::runtime_error("CUDA device-to-array pitched copy mismatch"); + std::fill(padded.begin(), padded.end(), 0xee); + cuda.check(cuda.cuMemcpyHtoD_v2(devicePixels, padded.data(), padded.size()), "cuMemcpyHtoD_v2(output sentinel)"); + image.download_device(devicePixels, pitch); + cuda.check(cuda.cuMemcpyDtoH_v2(padded.data(), devicePixels, padded.size()), "cuMemcpyDtoH_v2(device array output)"); + for (unsigned y = 0; y < 9; ++y) { + if (std::memcmp(padded.data() + y * pitch, input.data() + y * 17 * channels, rowBytes)) + throw std::runtime_error("CUDA array-to-device pitched copy mismatch"); + if (!std::all_of(padded.begin() + y * pitch + rowBytes, padded.begin() + (y + 1) * pitch, [](uint8_t value) { return value == 0xee; })) + throw std::runtime_error("CUDA array-to-device copy overwrote pitch padding"); + } } } cuda.check(cuda.cuCtxSynchronize(), "cuCtxSynchronize"); - std::printf("PASS CUDA_IMAGES architecture=gfx1201 iterations=%u descriptor_and_storage=1\n", args.iterations); + std::printf("PASS CUDA_IMAGES architecture=gfx1201 iterations=%u descriptor_and_storage=1 pitched_device_copies=1\n", args.iterations); return 0; } catch (const std::exception& e) { std::fprintf(stderr, "FAIL CUDA_IMAGES %s\n", e.what()); diff --git a/tools/windows/d3d12_dlss_probe.cpp b/tools/windows/d3d12_dlss_probe.cpp new file mode 100644 index 0000000..71117ae --- /dev/null +++ b/tools/windows/d3d12_dlss_probe.cpp @@ -0,0 +1,199 @@ +#include "d3d12_external.h" +#include "ngx_parameters.h" +#include "ngx_cuda_evaluate.h" +#include +#include + +int main(int argc, char** argv) { + using namespace d4r::diag; + using namespace d4r::win; + start(); + try { + Args args(argc, argv); + if (args.cuda_dll.empty() || (args.interop_mode != "images" && (args.module.empty() || args.ngx_core.empty() || args.dlss_dll.empty()))) + throw std::runtime_error("--module shim DLL, --cuda-dll, --ngx-core and --dlss-dll required"); + const char* diagDir = std::getenv("D4R_DIAG_DIR"); + if (!diagDir && args.interop_mode != "images") throw std::runtime_error("Set D4R_DIAG_DIR"); + const std::filesystem::path directory(diagDir ? wide(diagDir) : std::filesystem::current_path().wstring()); + std::filesystem::create_directories(directory); + for (const auto& pair : {std::make_pair("D4R_HIP_ROOT", args.hip_root), + std::make_pair("D4R_NVCUDA_DLL", args.cuda_dll), std::make_pair("D4R_NVAPI_DLL", args.nvapi_dll), + std::make_pair("D4R_NGX_CORE", args.ngx_core), std::make_pair("D4R_DLSS_DLL", args.dlss_dll)}) + if (_putenv_s(pair.first, pair.second.c_str())) throw std::runtime_error("Cannot set runtime environment"); + HipApi hip(args.hip_root); hipDeviceProp_t props{}; hip.select_gfx1201(args.device, props); + ComPtr factory; dx(CreateDXGIFactory1(IID_PPV_ARGS(factory.GetAddressOf())), "CreateDXGIFactory1"); + ComPtr adapter; + bool found = false; + for (UINT i = 0; ; ++i) { + const HRESULT result = factory->EnumAdapters1(i, adapter.ReleaseAndGetAddressOf()); + if (result == DXGI_ERROR_NOT_FOUND) break; + dx(result, "EnumAdapters"); DXGI_ADAPTER_DESC1 desc{}; dx(adapter->GetDesc1(&desc), "GetDesc1"); + if (!std::memcmp(&desc.AdapterLuid, props.luid, sizeof(LUID))) { found = true; break; } + } + if (!found) throw std::runtime_error("HIP adapter missing from DXGI"); + ComPtr device; + dx(D3D12CreateDevice(adapter.Get(), D3D_FEATURE_LEVEL_12_0, IID_PPV_ARGS(device.GetAddressOf())), "D3D12CreateDevice"); + ComPtr queue; + D3D12_COMMAND_QUEUE_DESC qd{}; qd.Type = D3D12_COMMAND_LIST_TYPE_DIRECT; + dx(device->CreateCommandQueue(&qd, IID_PPV_ARGS(queue.GetAddressOf())), "CreateCommandQueue"); + ComPtr completion; + dx(device->CreateFence(0, D3D12_FENCE_FLAG_NONE, IID_PPV_ARGS(completion.GetAddressOf())), "CreateFence(harness)"); + Handle event; event.value = CreateEventW(nullptr, FALSE, FALSE, nullptr); + if (!event.value) throw std::runtime_error("CreateEvent(harness)"); + uint64_t fenceValue = 0; + auto drain = [&] { + dx(queue->Signal(completion.Get(), ++fenceValue), "Queue Signal(harness)"); + if (completion->GetCompletedValue() < fenceValue) { + dx(completion->SetEventOnCompletion(fenceValue, event.value), "SetEventOnCompletion(harness)"); + if (WaitForSingleObject(event.value, 30000) != WAIT_OBJECT_0) throw std::runtime_error("Harness GPU timeout"); + } + }; + ComPtr allocator; + ComPtr list; + dx(device->CreateCommandAllocator(D3D12_COMMAND_LIST_TYPE_DIRECT, IID_PPV_ARGS(allocator.GetAddressOf())), "CreateCommandAllocator"); + dx(device->CreateCommandList(0, D3D12_COMMAND_LIST_TYPE_DIRECT, allocator.Get(), nullptr, IID_PPV_ARGS(list.GetAddressOf())), "CreateCommandList"); + constexpr unsigned width = 256, height = 144, outWidth = 512, outHeight = 288; + struct Texture { ComPtr image, upload; D3D12_PLACED_SUBRESOURCE_FOOTPRINT footprint{}; UINT64 bytes; } textures[5]; + for (unsigned i = 0; i < 5; ++i) { + auto& texture = textures[i]; + D3D12_RESOURCE_DESC desc{}; desc.Dimension = D3D12_RESOURCE_DIMENSION_TEXTURE2D; + desc.Width = i == 4 ? 1 : i == 3 ? outWidth : width; + desc.Height = i == 4 ? 1 : i == 3 ? outHeight : height; + desc.DepthOrArraySize = desc.MipLevels = desc.SampleDesc.Count = 1; + desc.Format = i == 0 || i == 3 ? DXGI_FORMAT_R16G16B16A16_FLOAT : i == 2 ? DXGI_FORMAT_R16G16_FLOAT : DXGI_FORMAT_R32_FLOAT; + if (i == 3) desc.Flags = D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS; + D3D12_HEAP_PROPERTIES heap{}; heap.Type = D3D12_HEAP_TYPE_DEFAULT; + heap.CreationNodeMask = heap.VisibleNodeMask = 1; + dx(device->CreateCommittedResource(&heap, D3D12_HEAP_FLAG_NONE, &desc, D3D12_RESOURCE_STATE_COPY_DEST, + nullptr, IID_PPV_ARGS(texture.image.GetAddressOf())), "CreateCommittedResource(texture)"); + UINT rows; UINT64 rowBytes; + device->GetCopyableFootprints(&desc, 0, 1, 0, &texture.footprint, &rows, &rowBytes, &texture.bytes); + texture.upload = make_buffer(device.Get(), texture.bytes, D3D12_HEAP_TYPE_UPLOAD, D3D12_HEAP_FLAG_NONE, D3D12_RESOURCE_STATE_GENERIC_READ); + void* raw = nullptr; D3D12_RANGE noRead{}; + dx(texture.upload->Map(0, &noRead, &raw), "Map(synthetic initialization)"); + std::memset(raw, 0, size_t(texture.bytes)); + for (unsigned y = 0; y < desc.Height; ++y) for (unsigned x = 0; x < desc.Width; ++x) { + auto* row = static_cast(raw) + y * texture.footprint.Footprint.RowPitch; + if (i == 0) { + uint16_t pixel[] = {uint16_t(((x / 16 + y / 16) & 1) ? 0x3a00 : 0x3400), + uint16_t((x & 32) ? 0x3800 : 0x3400), uint16_t((y & 32) ? 0x3a00 : 0x3800), 0x3c00}; + std::memcpy(row + x * 8, pixel, sizeof(pixel)); + } else if (i == 1 || i == 4) { + const float value = i == 1 ? .5f : 1.f; std::memcpy(row + x * 4, &value, 4); + } else if (i == 3) { + const uint16_t pixel[] = {0x7e00, 0x7e00, 0x7e00, 0x7e00}; std::memcpy(row + x * 8, pixel, 8); + } + } + texture.upload->Unmap(0, nullptr); + D3D12_TEXTURE_COPY_LOCATION src{}, dst{}; + src.pResource = texture.upload.Get(); src.Type = D3D12_TEXTURE_COPY_TYPE_PLACED_FOOTPRINT; src.PlacedFootprint = texture.footprint; + dst.pResource = texture.image.Get(); dst.Type = D3D12_TEXTURE_COPY_TYPE_SUBRESOURCE_INDEX; + list->CopyTextureRegion(&dst, 0, 0, 0, &src, nullptr); + transition(list.Get(), texture.image.Get(), D3D12_RESOURCE_STATE_COPY_DEST, + i == 3 ? D3D12_RESOURCE_STATE_UNORDERED_ACCESS : D3D12_RESOURCE_STATE_NON_PIXEL_SHADER_RESOURCE); + } + dx(list->Close(), "Close(synthetic producers)"); ID3D12CommandList* producer[] = {list.Get()}; + queue->ExecuteCommandLists(1, producer); drain(); + if (args.interop_mode == "images") { + CudaApi cuda(args.cuda_dll); cuda.check(cuda.cuInit(0), "cuInit(image interop)"); + CUdevice ordinal = 0; CUcontext context = nullptr; + cuda.check(cuda.cuDeviceGet(&ordinal, 0), "cuDeviceGet(image interop)"); + cuda.check(cuda.cuDevicePrimaryCtxRetain(&context, ordinal), "cuDevicePrimaryCtxRetain(image interop)"); + struct ContextCleanup { CudaApi& cuda; CUdevice ordinal; ~ContextCleanup() { (void)cuda.cuCtxSetCurrent(nullptr); (void)cuda.cuDevicePrimaryCtxRelease_v2(ordinal); } } contextCleanup{cuda, ordinal}; + cuda.check(cuda.cuCtxSetCurrent(context), "cuCtxSetCurrent(image interop)"); + ExternalApi external{hip}; SharedTimeline timeline(external, device.Get()); + d4r::cuda::ImageApi images(cuda); + std::unique_ptr planes[5]; + std::unique_ptr arrays[5]; + dx(allocator->Reset(), "Reset(image interop allocator)"); dx(list->Reset(allocator.Get(), nullptr), "Reset(image interop list)"); + for (unsigned i = 0; i < 5; ++i) { + const auto desc = resource_desc(textures[i].image.Get()); + const unsigned channels = i == 0 || i == 3 ? 4 : i == 2 ? 2 : 1; + const bool half = i == 0 || i == 2 || i == 3; + planes[i] = std::make_unique(external, device.Get(), desc); + arrays[i] = std::make_unique(images, unsigned(desc.Width), desc.Height, half ? 16 : 32, channels, i == 3, i == 0 ? 1 : 0); + planes[i]->copy_input(list.Get(), textures[i].image.Get(), i == 3 ? D3D12_RESOURCE_STATE_UNORDERED_ACCESS : D3D12_RESOURCE_STATE_NON_PIXEL_SHADER_RESOURCE); + } + dx(list->Close(), "Close(image interop input)"); ID3D12CommandList* copies[] = {list.Get()}; queue->ExecuteCommandLists(1, copies); + timeline.wait_input(queue.Get()); + for (unsigned i = 0; i < 5; ++i) arrays[i]->upload_device(reinterpret_cast(planes[i]->mapped), planes[i]->footprint.Footprint.RowPitch); + for (unsigned i = 0; i < 5; ++i) { + const auto desc = resource_desc(textures[i].image.Get()); + const unsigned channels = i == 0 || i == 3 ? 4 : i == 2 ? 2 : 1; + const bool half = i == 0 || i == 2 || i == 3; + const size_t rowBytes = size_t(desc.Width) * channels * (half ? 2 : 4); + std::vector actual(rowBytes * desc.Height); + arrays[i]->download(actual.data()); + void* expected = nullptr; D3D12_RANGE read{0, size_t(textures[i].bytes)}; + dx(textures[i].upload->Map(0, &read, &expected), "Map(expected diagnostic image)"); + bool matches = true; + for (unsigned y = 0; y < desc.Height; ++y) if (std::memcmp(actual.data() + y * rowBytes, + static_cast(expected) + y * textures[i].footprint.Footprint.RowPitch, rowBytes)) { matches = false; break; } + D3D12_RANGE noWrite{}; textures[i].upload->Unmap(0, &noWrite); + std::printf("D3D12_ARRAY_COPY plane=%u format=%u matches=%u first=%02x%02x%02x%02x\n", i, desc.Format, matches, actual[0], actual[1], actual[2], actual[3]); + if (!matches) throw std::runtime_error("Imported D3D12 VRAM to CUDA array mismatch on plane " + std::to_string(i)); + timeline.drain(queue.Get()); + } + std::printf("PASS D3D12_ARRAY_COPY architecture=gfx1201 planes=5 cpu_copies_between_apis=0\n"); + return 0; + } + Library shim(wide(args.module)); + using Init = unsigned(*)(unsigned long long, const wchar_t*, ID3D12Device*, unsigned, const void*); + using Allocate = unsigned(*)(void**); using Destroy = unsigned(*)(void*); + using Create = unsigned(*)(ID3D12GraphicsCommandList*, unsigned, void*, void**); + using Evaluate = unsigned(*)(ID3D12CommandQueue*, void*, void*); + using Release = unsigned(*)(void*); using Shutdown = unsigned(*)(); + auto check = [](unsigned result, const char* call) { if (result != 1) throw std::runtime_error(std::string(call) + " NGX=" + std::to_string(result)); }; + check(shim.symbol("NVSDK_NGX_D3D12_Init_Ext")(241534723ull, directory.c_str(), device.Get(), 0x15, nullptr), "D3D12 Init"); + auto shutdown = shim.symbol("NVSDK_NGX_D3D12_Shutdown"); + struct ShutdownCleanup { Shutdown fn; ~ShutdownCleanup() { (void)fn(); } } shutdownCleanup{shutdown}; + void* parameters = nullptr; check(shim.symbol("NVSDK_NGX_D3D12_AllocateParameters")(¶meters), "AllocateParameters"); + auto destroy = shim.symbol("NVSDK_NGX_D3D12_DestroyParameters"); + struct ParameterCleanup { Destroy fn; void* p; ~ParameterCleanup() { (void)fn(p); } } parameterCleanup{destroy, parameters}; + for (const auto& pair : {std::make_pair("Width", width), std::make_pair("Height", height), + std::make_pair("OutWidth", outWidth), std::make_pair("OutHeight", outHeight), + std::make_pair("CreationNodeMask", 1u), std::make_pair("VisibilityNodeMask", 1u)}) d4r_ngx_set_uint(parameters, pair.first, pair.second); + d4r_ngx_set_int(parameters, "PerfQualityValue", 2); d4r_ngx_set_int(parameters, "DLSS.Feature.Create.Flags", 0); + for (const char* name : {"DLSS.Hint.Render.Preset.DLAA", "DLSS.Hint.Render.Preset.Quality", "DLSS.Hint.Render.Preset.Balanced", + "DLSS.Hint.Render.Preset.Performance", "DLSS.Hint.Render.Preset.UltraPerformance", "DLSS.Hint.Render.Preset.UltraQuality"}) d4r_ngx_set_uint(parameters, name, args.preset); + void* handle = nullptr; + check(shim.symbol("NVSDK_NGX_D3D12_CreateFeature")(nullptr, 1, parameters, &handle), "D3D12 CreateFeature"); + auto release = shim.symbol("NVSDK_NGX_D3D12_ReleaseFeature"); + struct FeatureCleanup { Release fn; void* h; ~FeatureCleanup() { (void)fn(h); } } featureCleanup{release, handle}; + static const char* names[] = {"Color", "Depth", "MotionVectors", "Output", "ExposureTexture"}; + for (unsigned i = 0; i < 5; ++i) d4r_ngx_set_d3d12_resource(parameters, names[i], textures[i].image.Get()); + d4r_ngx_set_int(parameters, "Disable.Watermark", 1); + auto readback = make_buffer(device.Get(), textures[3].bytes, D3D12_HEAP_TYPE_READBACK, D3D12_HEAP_FLAG_NONE, D3D12_RESOURCE_STATE_COPY_DEST); + for (unsigned frame = 0; frame < args.iterations; ++frame) { + d4r_ngx_set_int(parameters, "Reset", frame == 0 ? 1 : 0); + check(shim.symbol("d4r_D3D12_EvaluateAtBoundary")(queue.Get(), handle, parameters), "D3D12 same-frame Evaluate"); + dx(allocator->Reset(), "Reset(harness allocator)"); dx(list->Reset(allocator.Get(), nullptr), "Reset(harness list)"); + transition(list.Get(), textures[3].image.Get(), D3D12_RESOURCE_STATE_UNORDERED_ACCESS, D3D12_RESOURCE_STATE_COPY_SOURCE); + D3D12_TEXTURE_COPY_LOCATION src{}, dst{}; + src.pResource = textures[3].image.Get(); src.Type = D3D12_TEXTURE_COPY_TYPE_SUBRESOURCE_INDEX; + dst.pResource = readback.Get(); dst.Type = D3D12_TEXTURE_COPY_TYPE_PLACED_FOOTPRINT; dst.PlacedFootprint = textures[3].footprint; + list->CopyTextureRegion(&dst, 0, 0, 0, &src, nullptr); + transition(list.Get(), textures[3].image.Get(), D3D12_RESOURCE_STATE_COPY_SOURCE, D3D12_RESOURCE_STATE_UNORDERED_ACCESS); + dx(list->Close(), "Close(verification)"); ID3D12CommandList* verify[] = {list.Get()}; queue->ExecuteCommandLists(1, verify); drain(); + 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); + 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); + 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"); + } + loaded_modules(); + std::printf("PASS D3D12_DLSS architecture=gfx1201 preset=%u frames=%u fast_path_cpu_copies=0 frame_age=0 queue_integration=explicit_harness\n", args.preset, args.iterations); + return 0; + } catch (const std::exception& error) { std::fprintf(stderr, "FAIL D3D12_DLSS %s\n", error.what()); loaded_modules(); return 4; } +} diff --git a/tools/windows/d3d12_external.h b/tools/windows/d3d12_external.h new file mode 100644 index 0000000..c83db3b --- /dev/null +++ b/tools/windows/d3d12_external.h @@ -0,0 +1,190 @@ +#pragma once +#if defined(__MINGW32__) +#define WIDL_EXPLICIT_AGGREGATE_RETURNS +#endif +#include "hip_api.h" +#include +#include + +namespace d4r::win { +using Microsoft::WRL::ComPtr; +using diag::HipApi; +inline void dx(HRESULT result, const char* operation) { + if (FAILED(result)) throw std::runtime_error(std::string(operation) + " HRESULT=" + std::to_string(result)); +} +struct Handle { + HANDLE value = nullptr; + Handle() = default; + Handle(const Handle&) = delete; + ~Handle() { if (value) CloseHandle(value); } +}; +struct ExternalApi { + HipApi& hip; +#define D4R_EXTERNAL(name) decltype(&::name) name = hip.library.symbol(#name) + D4R_EXTERNAL(hipImportExternalMemory); + D4R_EXTERNAL(hipExternalMemoryGetMappedBuffer); + D4R_EXTERNAL(hipDestroyExternalMemory); + D4R_EXTERNAL(hipImportExternalSemaphore); + D4R_EXTERNAL(hipDestroyExternalSemaphore); + D4R_EXTERNAL(hipWaitExternalSemaphoresAsync); + D4R_EXTERNAL(hipSignalExternalSemaphoresAsync); +#undef D4R_EXTERNAL +}; +inline D3D12_RESOURCE_DESC resource_desc(ID3D12Resource* resource) { +#if defined(__MINGW32__) + D3D12_RESOURCE_DESC desc{}; resource->GetDesc(&desc); return desc; +#else + return resource->GetDesc(); +#endif +} +inline LUID adapter_luid(ID3D12Device* device) { +#if defined(__MINGW32__) + LUID luid{}; device->GetAdapterLuid(&luid); return luid; +#else + return device->GetAdapterLuid(); +#endif +} +inline ComPtr make_buffer(ID3D12Device* device, uint64_t bytes, + D3D12_HEAP_TYPE heapType, D3D12_HEAP_FLAGS flags, D3D12_RESOURCE_STATES state) { + D3D12_HEAP_PROPERTIES heap{}; heap.Type = heapType; + heap.CreationNodeMask = heap.VisibleNodeMask = 1; + D3D12_RESOURCE_DESC desc{}; + desc.Dimension = D3D12_RESOURCE_DIMENSION_BUFFER; desc.Width = bytes; + desc.Height = desc.DepthOrArraySize = desc.MipLevels = desc.SampleDesc.Count = 1; + desc.Layout = D3D12_TEXTURE_LAYOUT_ROW_MAJOR; + ComPtr resource; + dx(device->CreateCommittedResource(&heap, flags, &desc, state, nullptr, + IID_PPV_ARGS(resource.GetAddressOf())), "CreateCommittedResource(buffer)"); + return resource; +} +inline void transition(ID3D12GraphicsCommandList* list, ID3D12Resource* resource, + D3D12_RESOURCE_STATES before, D3D12_RESOURCE_STATES after) { + if (before == after) return; + D3D12_RESOURCE_BARRIER barrier{}; + barrier.Type = D3D12_RESOURCE_BARRIER_TYPE_TRANSITION; + barrier.Transition = {resource, D3D12_RESOURCE_BARRIER_ALL_SUBRESOURCES, before, after}; + list->ResourceBarrier(1, &barrier); +} + +class SharedPlane { + ExternalApi& api_; + hipExternalMemory_t memory_ = nullptr; + void release() noexcept { + if (mapped) { (void)api_.hip.hipFree(mapped); mapped = nullptr; } + if (memory_) { (void)api_.hipDestroyExternalMemory(memory_); memory_ = nullptr; } + } +public: + ComPtr buffer; + D3D12_PLACED_SUBRESOURCE_FOOTPRINT footprint{}; + UINT64 bytes = 0; + void* mapped = nullptr; + SharedPlane(ExternalApi& api, ID3D12Device* device, const D3D12_RESOURCE_DESC& texture) : api_(api) { + UINT rows = 0; UINT64 rowBytes = 0; + device->GetCopyableFootprints(&texture, 0, 1, 0, &footprint, &rows, &rowBytes, &bytes); + if (!bytes || bytes == UINT64_MAX || !footprint.Footprint.RowPitch) + throw std::runtime_error("Invalid shared texture footprint"); + buffer = make_buffer(device, bytes, D3D12_HEAP_TYPE_DEFAULT, D3D12_HEAP_FLAG_SHARED, D3D12_RESOURCE_STATE_COMMON); + const auto desc = resource_desc(buffer.Get()); + D3D12_RESOURCE_ALLOCATION_INFO allocation{}; +#if defined(__MINGW32__) + device->GetResourceAllocationInfo(&allocation, 0, 1, &desc); +#else + allocation = device->GetResourceAllocationInfo(0, 1, &desc); +#endif + Handle handle; + dx(device->CreateSharedHandle(buffer.Get(), nullptr, GENERIC_ALL, nullptr, &handle.value), "CreateSharedHandle(VRAM)"); + try { + hipExternalMemoryHandleDesc import{}; + import.type = hipExternalMemoryHandleTypeD3D12Resource; + import.handle.win32.handle = handle.value; + import.size = allocation.SizeInBytes; import.flags = hipExternalMemoryDedicated; + api.hip.check(api.hipImportExternalMemory(&memory_, &import), "hipImportExternalMemory(D3D12)"); + hipExternalMemoryBufferDesc map{}; map.size = bytes; + api.hip.check(api.hipExternalMemoryGetMappedBuffer(&mapped, memory_, &map), "hipExternalMemoryGetMappedBuffer"); + } catch (...) { release(); throw; } + // Win32 imports retain their own reference; the exporter closes its HANDLE. + std::printf("D4R_IMPORT width=%u height=%u pitch=%u allocation=%llu cpu_copy=0\n", + footprint.Footprint.Width, footprint.Footprint.Height, footprint.Footprint.RowPitch, allocation.SizeInBytes); + } + SharedPlane(const SharedPlane&) = delete; + ~SharedPlane() { release(); } + void copy_input(ID3D12GraphicsCommandList* list, ID3D12Resource* texture, D3D12_RESOURCE_STATES state) { + transition(list, texture, state, D3D12_RESOURCE_STATE_COPY_SOURCE); + transition(list, buffer.Get(), D3D12_RESOURCE_STATE_COMMON, D3D12_RESOURCE_STATE_COPY_DEST); + D3D12_TEXTURE_COPY_LOCATION source{}, dest{}; + source.pResource = texture; source.Type = D3D12_TEXTURE_COPY_TYPE_SUBRESOURCE_INDEX; + dest.pResource = buffer.Get(); dest.Type = D3D12_TEXTURE_COPY_TYPE_PLACED_FOOTPRINT; dest.PlacedFootprint = footprint; + list->CopyTextureRegion(&dest, 0, 0, 0, &source, nullptr); + transition(list, buffer.Get(), D3D12_RESOURCE_STATE_COPY_DEST, D3D12_RESOURCE_STATE_COMMON); + transition(list, texture, D3D12_RESOURCE_STATE_COPY_SOURCE, state); + } + void copy_output(ID3D12GraphicsCommandList* list, ID3D12Resource* texture, D3D12_RESOURCE_STATES state) { + transition(list, texture, state, D3D12_RESOURCE_STATE_COPY_DEST); + transition(list, buffer.Get(), D3D12_RESOURCE_STATE_COMMON, D3D12_RESOURCE_STATE_COPY_SOURCE); + D3D12_TEXTURE_COPY_LOCATION source{}, dest{}; + source.pResource = buffer.Get(); source.Type = D3D12_TEXTURE_COPY_TYPE_PLACED_FOOTPRINT; source.PlacedFootprint = footprint; + dest.pResource = texture; dest.Type = D3D12_TEXTURE_COPY_TYPE_SUBRESOURCE_INDEX; + list->CopyTextureRegion(&dest, 0, 0, 0, &source, nullptr); + transition(list, buffer.Get(), D3D12_RESOURCE_STATE_COPY_SOURCE, D3D12_RESOURCE_STATE_COMMON); + transition(list, texture, D3D12_RESOURCE_STATE_COPY_DEST, state); + } +}; + +class SharedTimeline { + ExternalApi& api_; + hipExternalSemaphore_t semaphore_ = nullptr; + hipStream_t stream_ = nullptr; + Handle event_; + uint64_t value_ = 0; +public: + 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)"); + Handle handle; + dx(device->CreateSharedHandle(fence.Get(), nullptr, GENERIC_ALL, nullptr, &handle.value), "CreateSharedHandle(fence)"); + event_.value = CreateEventW(nullptr, FALSE, FALSE, nullptr); + if (!event_.value) throw std::runtime_error("CreateEvent(fence) failed"); + hipExternalSemaphoreHandleDesc import{}; + import.type = hipExternalSemaphoreHandleTypeD3D12Fence; + import.handle.win32.handle = handle.value; + api.hip.check(api.hipImportExternalSemaphore(&semaphore_, &import), "hipImportExternalSemaphore(D3D12Fence)"); + try { api.hip.check(api.hip.hipStreamCreateWithFlags(&stream_, hipStreamNonBlocking), "hipStreamCreate(interop)"); } + catch (...) { (void)api.hipDestroyExternalSemaphore(semaphore_); semaphore_ = nullptr; throw; } + } + ~SharedTimeline() { + if (stream_) { (void)api_.hip.hipStreamSynchronize(stream_); (void)api_.hip.hipStreamDestroy(stream_); } + if (semaphore_) (void)api_.hipDestroyExternalSemaphore(semaphore_); + } + // Caller must serialize transactions and submit the producer prefix first. + void wait_input(ID3D12CommandQueue* queue) { + const uint64_t ready = ++value_; + dx(queue->Signal(fence.Get(), ready), "Queue Signal(input)"); + 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"); + // Establish host-side completion as well: the Windows HIP semaphore + // wait currently needs a separately verified completion boundary before + // work is launched on CUDA's other streams. No image crosses the CPU. + if (fence->GetCompletedValue() < ready) { + dx(fence->SetEventOnCompletion(ready, event_.value), "SetEventOnCompletion(input)"); + if (WaitForSingleObject(event_.value, 30000) != WAIT_OBJECT_0) + throw std::runtime_error("D3D12 input completion timeout"); + } + } + void signal_output(ID3D12CommandQueue* queue) { + const uint64_t done = ++value_; + 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_; + dx(queue->Signal(fence.Get(), done), "Queue Signal(completion)"); + if (fence->GetCompletedValue() < done) { + dx(fence->SetEventOnCompletion(done, event_.value), "SetEventOnCompletion"); + if (WaitForSingleObject(event_.value, 30000) != WAIT_OBJECT_0) + throw std::runtime_error("D3D12 completion timeout; resources must remain alive"); + } + } +}; +} diff --git a/tools/windows/d4r_nvngx_windows.cpp b/tools/windows/d4r_nvngx_windows.cpp new file mode 100644 index 0000000..51e8aaf --- /dev/null +++ b/tools/windows/d4r_nvngx_windows.cpp @@ -0,0 +1,111 @@ +// Native Windows backend. The Linux/Wine implementation remains separate. +#include "ngx_windows_runtime.h" +#include +#include + +namespace { +std::mutex apiMutex; +std::shared_ptr runtime; +struct Handle { unsigned Id; }; +unsigned nextId = 1; +std::unordered_map> features; +std::unordered_map parameters; +constexpr unsigned failure = 0xBAD00002, invalid = 0xBAD00005; +unsigned call(const std::function& fn) noexcept { + try { return fn(); } + catch (const std::exception& error) { + std::fprintf(stderr, "D4R_WINDOWS_FAILURE %s\n", error.what()); return failure; + } +} +unsigned allocate(void** out, bool caps) { + if (!runtime || !out) return invalid; + void* value = d4r_ngx_parameters_create(); + if (!value) return failure; + parameters.emplace(value, true); *out = value; + if (caps) { + std::lock_guard lock(runtime->mutex); + runtime->current(); void* official = nullptr; + d4r::win::ngx_check(runtime->capabilities(&official), "CUDA capabilities"); + struct Cleanup { d4r::win::Runtime& rt; void* p; ~Cleanup() { (void)rt.destroy(p); } } cleanup{*runtime, official}; + for (const char* name : {"SuperSampling.Available", "SuperSampling.NeedsUpdatedDriver", "SuperSampling.FeatureInitResult"}) + d4r_ngx_set_int(value, name, d4r::ngx::int_value(official, name)); + for (const char* name : {"SuperSampling.MinDriverVersionMajor", "SuperSampling.MinDriverVersionMinor"}) + d4r_ngx_set_uint(value, name, d4r::ngx::uint_value(official, name)); + for (const char* name : {"DLSSOptimalSettingsCallback", "DLSSGetStatsCallback"}) { + void* callback = nullptr; + if (d4r_ngx_get_void(official, name, &callback) == 1) d4r_ngx_set_void(value, name, callback); + } + } + return 1; +} +} +#define API extern "C" __declspec(dllexport) +API unsigned NVSDK_NGX_D3D12_Init_Ext(unsigned long long app, const wchar_t* data, ID3D12Device* device, unsigned sdk, const void*) { + std::lock_guard lock(apiMutex); + return call([&] { + if (!runtime) runtime = std::make_shared(device, app, data, sdk); + return 1u; + }); +} +API unsigned NVSDK_NGX_D3D12_Init(unsigned long long app, const wchar_t* data, ID3D12Device* device, const void* info, unsigned sdk) { + return NVSDK_NGX_D3D12_Init_Ext(app, data, device, sdk, info); +} +API unsigned NVSDK_NGX_D3D12_Init_ProjectID(const char*, int, const char*, const wchar_t* data, ID3D12Device* device, unsigned sdk, const void* info) { + // Driver core's CUDA ABI uses the numerical application identity. + return NVSDK_NGX_D3D12_Init_Ext(241534723ull, data, device, sdk, info); +} +API unsigned NVSDK_NGX_D3D12_AllocateParameters(void** p) { + std::lock_guard lock(apiMutex); return call([&] { return allocate(p, false); }); +} +API unsigned NVSDK_NGX_D3D12_GetCapabilityParameters(void** p) { + std::lock_guard lock(apiMutex); return call([&] { return allocate(p, true); }); +} +API unsigned NVSDK_NGX_D3D12_GetParameters(void** p) { return NVSDK_NGX_D3D12_GetCapabilityParameters(p); } +API unsigned NVSDK_NGX_D3D12_DestroyParameters(void* p) { + std::lock_guard lock(apiMutex); + if (!parameters.erase(p)) return invalid; + d4r_ngx_parameters_destroy(p); return 1; +} +API unsigned NVSDK_NGX_D3D12_GetScratchBufferSize(unsigned, const void*, size_t* size) { + if (!size) return invalid; *size = 0; return 1; +} +API unsigned NVSDK_NGX_D3D12_CreateFeature(ID3D12GraphicsCommandList*, unsigned feature, void* p, Handle** out) { + std::lock_guard lock(apiMutex); + return call([&] { + if (!runtime || feature != 1 || !p || !out) return invalid; + auto value = std::make_shared(runtime, p); + auto handle = std::make_unique(); handle->Id = nextId++; + features.emplace(handle.get(), std::move(value)); *out = handle.release(); return 1u; + }); +} +API unsigned d4r_D3D12_EvaluateAtBoundary(ID3D12CommandQueue* queue, Handle* handle, void* p) { + std::shared_ptr feature; + { std::lock_guard lock(apiMutex); + auto it = features.find(handle); if (it == features.end() || !p) return invalid; feature = it->second; } + return call([&] { feature->evaluate_boundary(queue, p); return 1u; }); +} +API unsigned NVSDK_NGX_D3D12_EvaluateFeature(ID3D12GraphicsCommandList*, const Handle*, void*, void*) { + // Explicitly refuse an unsubmitted game list until queue integration exists. + std::fprintf(stderr, "D4R_WINDOWS_FAILURE command-list queue backend pending; no frame-age fallback\n"); + return failure; +} +API unsigned NVSDK_NGX_D3D12_EvaluateFeature_C(ID3D12GraphicsCommandList* list, const Handle* handle, void* p, void* callback) { + return NVSDK_NGX_D3D12_EvaluateFeature(list, handle, p, callback); +} +API unsigned NVSDK_NGX_D3D12_ReleaseFeature(Handle* handle) { + std::shared_ptr feature; + { 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; }); +} +API unsigned NVSDK_NGX_D3D12_Shutdown() { + std::lock_guard lock(apiMutex); + return call([&] { + for (auto& pair : features) delete pair.first; + features.clear(); + for (auto& pair : parameters) d4r_ngx_parameters_destroy(pair.first); + parameters.clear(); runtime.reset(); return 1u; + }); +} +API unsigned NVSDK_NGX_D3D12_Shutdown1(ID3D12Device*) { return NVSDK_NGX_D3D12_Shutdown(); } diff --git a/tools/windows/ngx_parameters.h b/tools/windows/ngx_parameters.h new file mode 100644 index 0000000..b5e6d0f --- /dev/null +++ b/tools/windows/ngx_parameters.h @@ -0,0 +1,56 @@ +#pragma once +#include +struct ID3D12Resource; +extern "C" { +void d4r_ngx_set_ull(void*, const char*, unsigned long long); +void d4r_ngx_set_float(void*, const char*, float); +void d4r_ngx_set_uint(void*, const char*, unsigned); +void d4r_ngx_set_int(void*, const char*, int); +void d4r_ngx_set_void(void*, const char*, void*); +void d4r_ngx_set_d3d12_resource(void*, const char*, ID3D12Resource*); +unsigned d4r_ngx_get_ull(void*, const char*, unsigned long long*); +unsigned d4r_ngx_get_float(void*, const char*, float*); +unsigned d4r_ngx_get_uint(void*, const char*, unsigned*); +unsigned d4r_ngx_get_int(void*, const char*, int*); +unsigned d4r_ngx_get_void(void*, const char*, void**); +unsigned d4r_ngx_get_d3d12_resource(void*, const char*, ID3D12Resource**); +void* d4r_ngx_parameters_create(); +void d4r_ngx_parameters_destroy(void*); +} +namespace d4r::ngx { +inline unsigned uint_value(void* p, const char* name, unsigned fallback = 0) { + unsigned value = fallback; (void)d4r_ngx_get_uint(p, name, &value); return value; +} +inline int int_value(void* p, const char* name, int fallback = 0) { + int value = fallback; (void)d4r_ngx_get_int(p, name, &value); return value; +} +inline float float_value(void* p, const char* name, float fallback = 0) { + float value = fallback; (void)d4r_ngx_get_float(p, name, &value); return value; +} +inline void copy_create(void* source, void* dest) { + for (const char* name : {"Width", "Height", "OutWidth", "OutHeight", "CreationNodeMask", "VisibilityNodeMask", + "DLSS.Hint.Render.Preset.DLAA", "DLSS.Hint.Render.Preset.Quality", "DLSS.Hint.Render.Preset.Balanced", + "DLSS.Hint.Render.Preset.Performance", "DLSS.Hint.Render.Preset.UltraPerformance", "DLSS.Hint.Render.Preset.UltraQuality"}) { + unsigned value = 0; + if (d4r_ngx_get_uint(source, name, &value) == 1) d4r_ngx_set_uint(dest, name, value); + } + for (const char* name : {"PerfQualityValue", "DLSS.Feature.Create.Flags", "DLSS.Enable.Output.Subrects"}) { + int value = 0; + if (d4r_ngx_get_int(source, name, &value) == 1) d4r_ngx_set_int(dest, name, value); + } +} +inline void copy_frame(void* source, void* dest) { + for (const char* name : {"Jitter.Offset.X", "Jitter.Offset.Y", "Sharpness", "MV.Scale.X", "MV.Scale.Y", + "DLSS.Pre.Exposure", "DLSS.Exposure.Scale", "FrameTimeDeltaInMsec"}) { + const float fallback = name[0] == 'M' || name[0] == 'D' ? 1.f : name[0] == 'F' ? 16.666667f : 0.f; + d4r_ngx_set_float(dest, name, float_value(source, name, fallback)); + } + for (const char* name : {"Reset", "Disable.Watermark", "DLSS.Indicator.Invert.X.Axis", "DLSS.Indicator.Invert.Y.Axis"}) + d4r_ngx_set_int(dest, name, int_value(source, name)); + for (const char* name : {"DLSS.Render.Subrect.Dimensions.Width", "DLSS.Render.Subrect.Dimensions.Height", + "DLSS.Input.Color.Subrect.Base.X", "DLSS.Input.Color.Subrect.Base.Y", "DLSS.Input.Depth.Subrect.Base.X", + "DLSS.Input.Depth.Subrect.Base.Y", "DLSS.Input.MV.Subrect.Base.X", "DLSS.Input.MV.Subrect.Base.Y", + "DLSS.Output.Subrect.Base.X", "DLSS.Output.Subrect.Base.Y"}) + d4r_ngx_set_uint(dest, name, uint_value(source, name)); +} +} diff --git a/tools/windows/ngx_windows_runtime.h b/tools/windows/ngx_windows_runtime.h new file mode 100644 index 0000000..29e5691 --- /dev/null +++ b/tools/windows/ngx_windows_runtime.h @@ -0,0 +1,270 @@ +#pragma once +#include "d3d12_external.h" +#include "cuda_image_api.h" +#include "ngx_parameters.h" +#include +#include +#include + +namespace d4r::win { +inline std::string env_path(const char* name) { + const char* value = std::getenv(name); + if (!value || !*value) throw std::runtime_error(std::string("Set the absolute path in ") + name); + if (!std::filesystem::path(diag::wide(value)).is_absolute()) throw std::runtime_error(std::string(name) + " must be absolute"); + return value; +} +inline void ngx_check(unsigned result, const char* operation) { + std::printf("D4R_NGX %s result=0x%08x\n", operation, result); + if (result != 1) throw std::runtime_error(std::string(operation) + " NGX=" + std::to_string(result)); +} + +struct Runtime { + diag::HipApi hip; + CudaApi cuda; + std::unique_ptr nvapi; + diag::SearchDirectory featureSearch; + std::unique_ptr dlss, core; + ExternalApi external; + cuda::ImageApi images; + ComPtr device; + CUdevice ordinal = -1; + CUcontext context = nullptr; + bool initialized = false; + std::unique_ptr timeline; + std::mutex mutex; + using Init = unsigned(*)(unsigned long long, const wchar_t*, unsigned); + using Shutdown = unsigned(*)(); + using Allocate = unsigned(*)(void**); + using Destroy = unsigned(*)(void*); + using Create = unsigned(*)(unsigned, void*, void**); + using Evaluate = unsigned(*)(void*, void*, void*); + using Release = unsigned(*)(void*); + Allocate allocate = nullptr, capabilities = nullptr; + Destroy destroy = nullptr; + Create create = nullptr; + Evaluate evaluate = nullptr; + Release release = nullptr; + Shutdown shutdown = nullptr; + + Runtime(ID3D12Device* d3d, unsigned long long app, const wchar_t* data, unsigned sdk) + : hip(env_path("D4R_HIP_ROOT")), cuda(env_path("D4R_NVCUDA_DLL")), + featureSearch(std::filesystem::path(diag::wide(env_path("D4R_DLSS_DLL"))).parent_path()), + external{hip}, images(cuda), device(d3d) { + if (!d3d) throw std::runtime_error("Null D3D12 device"); + try { + hipDeviceProp_t props{}; + hip.select_gfx1201(-1, props); + const auto luid = adapter_luid(d3d); + if (std::memcmp(&luid, props.luid, sizeof(luid))) throw std::runtime_error("D3D12/HIP LUID mismatch"); + cuda.check(cuda.cuInit(0), "cuInit(Windows runtime)"); + int count = 0; cuda.check(cuda.cuDeviceGetCount(&count), "cuDeviceGetCount"); + for (int i = 0; i < count; ++i) { + CUdevice candidate = -1; cuda.check(cuda.cuDeviceGet(&candidate, i), "cuDeviceGet"); + int bus = -1, pciDevice = -1, domain = -1; + cuda.check(cuda.cuDeviceGetAttribute(&bus, 33, candidate), "PCI bus"); + cuda.check(cuda.cuDeviceGetAttribute(&pciDevice, 34, candidate), "PCI device"); + cuda.check(cuda.cuDeviceGetAttribute(&domain, 50, candidate), "PCI domain"); + if (bus == props.pciBusID && pciDevice == props.pciDeviceID && domain == props.pciDomainID) ordinal = candidate; + } + if (ordinal < 0) throw std::runtime_error("CUDA/HIP PCI identity mismatch"); + cuda.check(cuda.cuDevicePrimaryCtxRetain(&context, ordinal), "cuDevicePrimaryCtxRetain"); + current(); + timeline = std::make_unique(external, d3d); + // NGX's loader can probe CUDA during DLL initialization: establish + // the correct primary context before loading either NVIDIA DLL. + nvapi = std::make_unique(diag::wide(env_path("D4R_NVAPI_DLL"))); + dlss = std::make_unique(diag::wide(env_path("D4R_DLSS_DLL"))); + core = std::make_unique(diag::wide(env_path("D4R_NGX_CORE"))); + allocate = core->symbol("NVSDK_NGX_CUDA_AllocateParameters"); + capabilities = core->symbol("NVSDK_NGX_CUDA_GetCapabilityParameters"); + destroy = core->symbol("NVSDK_NGX_CUDA_DestroyParameters"); + create = core->symbol("NVSDK_NGX_CUDA_CreateFeature"); + evaluate = core->symbol("NVSDK_NGX_CUDA_EvaluateFeature"); + release = core->symbol("NVSDK_NGX_CUDA_ReleaseFeature"); + shutdown = core->symbol("NVSDK_NGX_CUDA_Shutdown"); + ngx_check(core->symbol("NVSDK_NGX_CUDA_Init")(app, data, sdk), "CUDA Init(driver ABI)"); + initialized = true; + void* caps = nullptr; + ngx_check(capabilities(&caps), "CUDA feature capabilities"); + if (!caps) throw std::runtime_error("Null CUDA capability parameters"); + const int available = ngx::int_value(caps, "SuperSampling.Available", -1); + const int initResult = ngx::int_value(caps, "SuperSampling.FeatureInitResult", -1); + const int needsDriver = ngx::int_value(caps, "SuperSampling.NeedsUpdatedDriver", -1); + std::printf("D4R_CAPABILITIES available=%d feature_init=0x%08x needs_driver=%d\n", available, unsigned(initResult), needsDriver); + ngx_check(destroy(caps), "Destroy capability parameters"); + if (available != 1 || initResult != 1 || needsDriver != 0) + throw std::runtime_error("CUDA DLSS is unavailable after capability initialization"); + std::printf("D4R_RUNTIME platform=Windows architecture=gfx1201 cpu_frame_copies=0 frame_age=0\n"); + } catch (...) { cleanup(); throw; } + } + Runtime(const Runtime&) = delete; + ~Runtime() { cleanup(); } + void current() { cuda.check(cuda.cuCtxSetCurrent(context), "cuCtxSetCurrent(runtime)"); } + void cleanup() noexcept { + if (context) (void)cuda.cuCtxSetCurrent(context); + if (initialized) { (void)shutdown(); initialized = false; } + timeline.reset(); + if (context) { + (void)cuda.cuCtxSynchronize(); (void)cuda.cuCtxSetCurrent(nullptr); + (void)cuda.cuDevicePrimaryCtxRelease_v2(ordinal); context = nullptr; + } + } +}; + +class Feature { + std::shared_ptr rt_; + void* parameters_ = nullptr; + void* handle_ = nullptr; + CUdeviceptr scratch_ = 0; + ComPtr queue_; + struct Plane { + ComPtr texture; + D3D12_RESOURCE_DESC desc{}; + std::unique_ptr shared; + std::unique_ptr image; + } planes_[5]; + void cleanup() noexcept { + (void)rt_->cuda.cuCtxSetCurrent(rt_->context); + (void)rt_->cuda.cuCtxSynchronize(); + if (handle_) { (void)rt_->release(handle_); handle_ = nullptr; } + for (auto& plane : planes_) { + plane.image.reset(); plane.shared.reset(); plane.texture.Reset(); + } + if (scratch_) { (void)rt_->cuda.cuMemFree_v2(scratch_); scratch_ = 0; } + if (parameters_) { (void)rt_->destroy(parameters_); parameters_ = nullptr; } + } +public: + unsigned width = 0, height = 0; + explicit Feature(std::shared_ptr runtime, void* parameters) : rt_(std::move(runtime)) { + std::lock_guard lock(rt_->mutex); + rt_->current(); + try { + ngx_check(rt_->allocate(¶meters_), "CUDA AllocateParameters"); + ngx::copy_create(parameters, parameters_); + d4r_ngx_set_int(parameters_, "DLSS.Enable.Output.Subrects", ngx::int_value(parameters, "DLSS.Enable.Output.Subrects")); + width = ngx::uint_value(parameters, "Width"); height = ngx::uint_value(parameters, "Height"); + std::printf("D4R_CREATE source=%ux%u->%ux%u cuda=%ux%u->%ux%u\n", width, height, + ngx::uint_value(parameters, "OutWidth"), ngx::uint_value(parameters, "OutHeight"), + ngx::uint_value(parameters_, "Width"), ngx::uint_value(parameters_, "Height"), + ngx::uint_value(parameters_, "OutWidth"), ngx::uint_value(parameters_, "OutHeight")); + if (!width || !height) throw std::runtime_error("Zero NGX input dimensions"); + rt_->cuda.check(rt_->cuda.cuMemAlloc_v2(&scratch_, 64ull * 1024 * 1024), "NGX VRAM scratch"); + d4r_ngx_set_void(parameters_, "Scratch", reinterpret_cast(uintptr_t(scratch_))); + d4r_ngx_set_ull(parameters_, "Scratch.SizeInBytes", 64ull * 1024 * 1024); + ngx_check(rt_->create(1, parameters_, &handle_), "CUDA CreateFeature"); + } catch (...) { cleanup(); throw; } + } + ~Feature() { std::lock_guard lock(rt_->mutex); cleanup(); } + Feature(const Feature&) = delete; + + // 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) { + std::lock_guard lock(rt_->mutex); + if (!queue) throw std::runtime_error("Null D3D12 queue"); + if (queue_ && queue_.Get() != queue) throw std::runtime_error("A feature cannot switch queues without an explicit drain"); + queue_ = queue; + ComPtr queueDevice; + dx(queue->GetDevice(IID_PPV_ARGS(queueDevice.GetAddressOf())), "GetDevice(queue)"); + ComPtr queueIdentity, runtimeIdentity; + dx(queueDevice->QueryInterface(IID_PPV_ARGS(queueIdentity.GetAddressOf())), "Query queue device identity"); + dx(rt_->device->QueryInterface(IID_PPV_ARGS(runtimeIdentity.GetAddressOf())), "Query runtime device identity"); + if (queueIdentity.Get() != runtimeIdentity.Get()) throw std::runtime_error("Queue belongs to a different D3D12 device"); + rt_->current(); + prepare(parameters); + ComPtr inputAllocator, outputAllocator; + ComPtr inputList, outputList; + auto list = [&](ComPtr& allocator, ComPtr& commandList) { + dx(rt_->device->CreateCommandAllocator(D3D12_COMMAND_LIST_TYPE_DIRECT, IID_PPV_ARGS(allocator.GetAddressOf())), "CreateCommandAllocator(interop)"); + 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); + record_inputs(inputList.Get(), parameters); dx(inputList->Close(), "Close(input copies)"); + ID3D12CommandList* before[] = {inputList.Get()}; queue->ExecuteCommandLists(1, before); + try { + execute_cuda(queue); + record_output(outputList.Get(), parameters); dx(outputList->Close(), "Close(output copy)"); + ID3D12CommandList* after[] = {outputList.Get()}; queue->ExecuteCommandLists(1, after); + rt_->timeline->drain(queue); + } catch (...) { + // Keep all recorded resources alive until submitted work is complete. + rt_->timeline->drain(queue); throw; + } + } +private: + void prepare(void* parameters) { + static 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 resource ") + names[i]); + planes_[i].texture.Reset(); d4r_ngx_set_void(parameters_, names[i], nullptr); continue; + } + auto desc = resource_desc(texture); + if (desc.Dimension != D3D12_RESOURCE_DIMENSION_TEXTURE2D || desc.SampleDesc.Count != 1 || desc.DepthOrArraySize != 1) + throw std::runtime_error("Interop baseline requires single-sample 2D textures"); + const bool half = i == 0 || i == 2 || i == 3; + const unsigned channels = i == 0 || i == 3 ? 4 : i == 2 ? 2 : 1; + const DXGI_FORMAT format = channels == 4 ? DXGI_FORMAT_R16G16B16A16_FLOAT : + channels == 2 ? DXGI_FORMAT_R16G16_FLOAT : DXGI_FORMAT_R32_FLOAT; + if (desc.Format != format && !(i == 1 && (desc.Format == DXGI_FORMAT_R32_TYPELESS || desc.Format == DXGI_FORMAT_D32_FLOAT))) + throw std::runtime_error(std::string("GPU conversion not yet implemented for ") + names[i] + " format=" + std::to_string(desc.Format)); + auto& plane = planes_[i]; + if (!plane.shared || plane.desc.Width != desc.Width || plane.desc.Height != desc.Height || plane.desc.Format != desc.Format) { + rt_->cuda.check(rt_->cuda.cuCtxSynchronize(), "Synchronize(resize interop)"); + plane.image.reset(); plane.shared.reset(); + plane.shared = std::make_unique(rt_->external, rt_->device.Get(), desc); + plane.image = std::make_unique(rt_->images, unsigned(desc.Width), desc.Height, half ? 16 : 32, channels, i == 3, i == 0 ? 1 : 0); + plane.desc = desc; + } + plane.texture = texture; + d4r_ngx_set_void(parameters_, names[i], &plane.image->object); + } + ngx::copy_frame(parameters, parameters_); + if (!ngx::uint_value(parameters_, "DLSS.Render.Subrect.Dimensions.Width")) d4r_ngx_set_uint(parameters_, "DLSS.Render.Subrect.Dimensions.Width", width); + if (!ngx::uint_value(parameters_, "DLSS.Render.Subrect.Dimensions.Height")) d4r_ngx_set_uint(parameters_, "DLSS.Render.Subrect.Dimensions.Height", height); + 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) { + 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(), D3D12_RESOURCE_STATES(ngx::uint_value(parameters, states[i], D3D12_RESOURCE_STATE_NON_PIXEL_SHADER_RESOURCE))); + } + } + void execute_cuda(ID3D12CommandQueue* queue) { + const auto begin = std::chrono::steady_clock::now(); + rt_->timeline->wait_input(queue); + rt_->cuda.check(rt_->cuda.cuCtxSynchronize(), "Interop input producer completion"); + for (unsigned i : {0u, 1u, 2u, 4u}) if (planes_[i].texture) { + planes_[i].image->upload_device(reinterpret_cast(planes_[i].shared->mapped), planes_[i].shared->footprint.Footprint.RowPitch); + } + if (std::getenv("D4R_INTEROP_VERIFY")) { + for (unsigned i : {0u, 1u, 2u, 4u}) if (planes_[i].texture) { + uint32_t sample[8]{}; + rt_->hip.check(rt_->hip.hipMemcpy(sample, planes_[i].shared->mapped, + std::min(sizeof(sample), planes_[i].shared->bytes), hipMemcpyDeviceToHost), "DIAGNOSTIC external-memory sample"); + std::printf("D4R_VRAM_VERIFY plane=%u words=%08x,%08x,%08x,%08x\n", i, sample[0], sample[1], sample[2], sample[3]); + const unsigned bytesPerPixel = i == 0 ? 8 : i == 2 ? 4 : 4; + std::vector array(size_t(planes_[i].desc.Width) * planes_[i].desc.Height * bytesPerPixel); + planes_[i].image->download(array.data()); + uint32_t first = 0; std::memcpy(&first, array.data(), 4); + std::printf("D4R_ARRAY_VERIFY plane=%u first=%08x\n", i, first); + } + } + ngx_check(rt_->evaluate(handle_, parameters_, nullptr), "CUDA EvaluateFeature"); + // NGX can use internal nonblocking streams. The default stream alone + // does not establish completion of those streams. + rt_->cuda.check(rt_->cuda.cuCtxSynchronize(), "NGX all-stream completion"); + planes_[3].image->download_device(reinterpret_cast(planes_[3].shared->mapped), planes_[3].shared->footprint.Footprint.RowPitch); + rt_->cuda.check(rt_->cuda.cuCtxSynchronize(), "VRAM output completion"); + rt_->timeline->signal_output(queue); + 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(), D3D12_RESOURCE_STATES(ngx::uint_value(parameters, "D4R.Output.State", D3D12_RESOURCE_STATE_UNORDERED_ACCESS))); + } +}; +}