Run native Windows D3D12 DLSS K and M through shared VRAM
This commit is contained in:
Codex committed
1 parent
2db93adb1a
commit
e15f6a73a6
12 files changed
+1049
-19
No files matched your search
+18
-1
@@ -68,7 +68,7 @@ foreach(family k m)
|
|||||||
endforeach()
|
endforeach()
|
||||||
endforeach()
|
endforeach()
|
||||||
add_custom_target(d4r_native_gfx1201_compile ALL DEPENDS ${native_objects})
|
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)
|
add_executable(d4r_${probe} tools/windows/${probe}.cpp)
|
||||||
target_compile_definitions(d4r_${probe} PRIVATE WIN32_LEAN_AND_MEAN NOMINMAX __HIP_PLATFORM_AMD__)
|
target_compile_definitions(d4r_${probe} PRIVATE WIN32_LEAN_AND_MEAN NOMINMAX __HIP_PLATFORM_AMD__)
|
||||||
target_include_directories(d4r_${probe} PRIVATE "${D4R_HIP_ROOT}/include")
|
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()
|
endif()
|
||||||
install(TARGETS d4r_${probe} RUNTIME DESTINATION bin)
|
install(TARGETS d4r_${probe} RUNTIME DESTINATION bin)
|
||||||
endforeach()
|
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)
|
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_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()
|
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)
|
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")
|
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)
|
DEPENDS tools/d4r_ngx_param_msvc.cpp VERBATIM)
|
||||||
set_source_files_properties("${ngx_parameter_object}" PROPERTIES GENERATED TRUE EXTERNAL_OBJECT TRUE)
|
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_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()
|
endif()
|
||||||
add_library(d4r_nvapi_trace SHARED tools/windows/nvapi_trace.cpp)
|
add_library(d4r_nvapi_trace SHARED tools/windows/nvapi_trace.cpp)
|
||||||
set_target_properties(d4r_nvapi_trace PROPERTIES PREFIX "" OUTPUT_NAME "nvapi64"
|
set_target_properties(d4r_nvapi_trace PROPERTIES PREFIX "" OUTPUT_NAME "nvapi64"
|
||||||
@@ -118,6 +131,7 @@ if(MINGW)
|
|||||||
endif()
|
endif()
|
||||||
install(TARGETS d4r_nvapi_compat RUNTIME DESTINATION nvapi-compat)
|
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_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_hip_gfx1201_probe d4r_probe_code_object)
|
||||||
add_dependencies(d4r_gfx12_wmma_probe d4r_wmma_code_object)
|
add_dependencies(d4r_gfx12_wmma_probe d4r_wmma_code_object)
|
||||||
add_dependencies(d4r_k_module_probe d4r_native_gfx1201_compile)
|
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
|
add_test(NAME cuda_images COMMAND d4r_cuda_image_probe
|
||||||
--hip-root "${D4R_HIP_ROOT}" --cuda-dll "${D4R_ZLUDA_ROOT}/nvcuda.dll" --iterations 32)
|
--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")
|
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()
|
||||||
endif()
|
endif()
|
||||||
@@ -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
|
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
|
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
|
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
|
NumPy and independent PTX, and four native CUDA frames. The standalone D3D12
|
||||||
in progress; the fully translated M comparison currently exceeds its JIT timeout.
|
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.
|
Do not integrate NGX until the integer PTX workload is stable on the real GPU.
|
||||||
|
|
||||||
| Gate | Status |
|
| 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 |
|
| M0: branch, audit, Windows build and diagnostics | Initial implementation complete |
|
||||||
| M1: native HIP gfx1201 allocation, kernel, CPU verification | PASS, 32 iterations + guard verification |
|
| 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 |
|
| 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 |
|
| 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 |
|
| 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 |
|
| 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 |
|
| 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 |
|
| M9: OptiScaler integration, profiling, installation | Pending |
|
||||||
|
|
||||||
## Source baseline (checked 2026-09-29)
|
## Source baseline (checked 2026-09-29)
|
||||||
@@ -166,14 +169,18 @@ Reproduce the layer gates:
|
|||||||
& scripts/windows/test-m-replay.ps1
|
& scripts/windows/test-m-replay.ps1
|
||||||
& scripts/windows/stage-native-m.ps1 -DlssDll "$PWD/nvngx_dlss.dll"
|
& scripts/windows/stage-native-m.ps1 -DlssDll "$PWD/nvngx_dlss.dll"
|
||||||
$env:D4R_ZLUDA_NATIVE_DIR="$PWD/build/native-m-gfx1201"
|
$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"
|
& 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
|
Earlier translated M attempts used incorrect environment names and an
|
||||||
automatic trace localises this to loading the transformer modules. Native M
|
unwritable default JIT cache. The runner now defaults to the actual
|
||||||
does not have this JIT delay. A completed translated image comparison is not
|
`D4R_ZLUDA_WMMA`, `D4R_ZLUDA_WMMA_FP8`, `D4R_ZLUDA_WMMA_FP8_NATIVE` options
|
||||||
claimed. The user selected Silent Hill 2 at
|
(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`
|
`D:\Games\SILENT HILL 2\SHProto\Binaries\Win64\SHProto-Win64-Shipping.exe`
|
||||||
for the later game gate. No game files have been changed yet.
|
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
|
D3D12 harness with explicit queue boundaries; command-list/queue integration
|
||||||
must be validated before the game gate.
|
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
|
### K real-weight native Evaluate and correctness baseline
|
||||||
|
|
||||||
The locally supplied DLSS 310.9.1.0 selects all eleven `dltss_pwin_*` functions
|
The locally supplied DLSS 310.9.1.0 selects all eleven `dltss_pwin_*` functions
|
||||||
|
|||||||
@@ -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<String> {
|
||||||
|
- 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<String> {
|
||||||
|
// 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<Self> {
|
||||||
|
- 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))]
|
||||||
@@ -1,5 +1,5 @@
|
|||||||
[CmdletBinding()]
|
[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'
|
$ErrorActionPreference = 'Stop'
|
||||||
$repo = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..'))
|
$repo = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..'))
|
||||||
if (!$SourceRoot) { $SourceRoot = Join-Path $repo 'external/ZLUDA' }
|
if (!$SourceRoot) { $SourceRoot = Join-Path $repo 'external/ZLUDA' }
|
||||||
@@ -33,6 +33,14 @@ Push-Location $SourceRoot
|
|||||||
try {
|
try {
|
||||||
$cargo = Join-Path $rust 'bin/cargo.exe'
|
$cargo = Join-Path $rust 'bin/cargo.exe'
|
||||||
if (!(Test-Path $cargo)) { throw 'Run setup-rust-toolchain.ps1 first.' }
|
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) {
|
if ($FetchOnly) {
|
||||||
$ErrorActionPreference = 'Continue'
|
$ErrorActionPreference = 'Continue'
|
||||||
& $cargo fetch --locked --target x86_64-pc-windows-gnu 2>&1 | Tee-Object (Join-Path $target 'fetch.log')
|
& $cargo fetch --locked --target x86_64-pc-windows-gnu 2>&1 | Tee-Object (Join-Path $target 'fetch.log')
|
||||||
|
|||||||
@@ -7,7 +7,7 @@ param(
|
|||||||
[string]$OutputDirectory,
|
[string]$OutputDirectory,
|
||||||
[string]$NgxCore = $env:D4R_NGX_CORE,
|
[string]$NgxCore = $env:D4R_NGX_CORE,
|
||||||
[string]$DlssDll = $env:D4R_DLSS_DLL,
|
[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,
|
[ValidateSet(5, 11, 13)][int]$Preset = 11,
|
||||||
[switch]$NgxOnly,
|
[switch]$NgxOnly,
|
||||||
[switch]$Trace,
|
[switch]$Trace,
|
||||||
@@ -51,6 +51,11 @@ $originalCuda = $env:ZLUDA_CUDA_LIB
|
|||||||
$originalNvapi = $env:D4R_NVAPI_BACKEND
|
$originalNvapi = $env:D4R_NVAPI_BACKEND
|
||||||
$originalPythonPath = $env:PYTHONPATH
|
$originalPythonPath = $env:PYTHONPATH
|
||||||
$originalVerbose = $env:D4R_ZLUDA_VERBOSE
|
$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
|
$exitStatus = 1
|
||||||
$ngxRuntimeDirectory = $null
|
$ngxRuntimeDirectory = $null
|
||||||
$summary = [ordered]@{utc=[DateTime]::UtcNow.ToString('o'); profile=$RuntimeProfile; hipRoot=$HipRoot; zludaRoot=$ZludaRoot; tests=@()}
|
$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
|
return $pass
|
||||||
}
|
}
|
||||||
try {
|
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 ($NgxOnly -and (!$NgxCore -or !$DlssDll)) { throw '-NgxOnly requires both locally supplied NVIDIA DLL paths.' }
|
||||||
if ($RequireNativeNetwork) {
|
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.'
|
throw '-RequireNativeNetwork requires K/M Evaluate and D4R_ZLUDA_NATIVE_DIR.'
|
||||||
}
|
}
|
||||||
$env:D4R_ZLUDA_VERBOSE = '1'
|
$env:D4R_ZLUDA_VERBOSE = '1'
|
||||||
@@ -209,7 +224,8 @@ try {
|
|||||||
$ngxExe = Join-Path $ngxRuntimeDirectory 'd4r_ngx_cuda_init_probe.exe'
|
$ngxExe = Join-Path $ngxRuntimeDirectory 'd4r_ngx_cuda_init_probe.exe'
|
||||||
$localCore = Join-Path $ngxRuntimeDirectory '_nvngx.dll'
|
$localCore = Join-Path $ngxRuntimeDirectory '_nvngx.dll'
|
||||||
$localDlss = Join-Path $ngxRuntimeDirectory 'nvngx_dlss.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 $NgxCore -Destination $localCore
|
||||||
Copy-Item -LiteralPath $DlssDll -Destination $localDlss
|
Copy-Item -LiteralPath $DlssDll -Destination $localDlss
|
||||||
$selectedCuda = if ($Trace) { Join-Path $ZludaRoot 'trace/nvcuda.dll' } else { Join-Path $ZludaRoot 'nvcuda.dll' }
|
$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,
|
'--ngx-core', $localCore, '--dlss-dll', $localDlss,
|
||||||
'--nvapi-dll', (Join-Path $PackageRoot 'nvapi-compat/nvapi64.dll'),
|
'--nvapi-dll', (Join-Path $PackageRoot 'nvapi-compat/nvapi64.dll'),
|
||||||
'--ngx-mode', $NgxMode, '--preset', "$Preset", '--iterations', "$Iterations")
|
'--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
|
$ngxOk = Invoke-Probe $ngxName $ngxExe $ngxArguments
|
||||||
if ($ngxOk -and $RequireNativeNetwork) {
|
if ($ngxOk -and $RequireNativeNetwork) {
|
||||||
$hits = @()
|
$hits = @()
|
||||||
@@ -238,12 +262,14 @@ try {
|
|||||||
}
|
}
|
||||||
if (!$ngxOk) {
|
if (!$ngxOk) {
|
||||||
$exitStatus = 1
|
$exitStatus = 1
|
||||||
if (Test-Path (Join-Path $ZludaRoot 'trace/nvcuda.dll')) {
|
if (!$Trace -and (Test-Path (Join-Path $ZludaRoot 'trace/nvcuda.dll'))) {
|
||||||
Invoke-Probe "$ngxName-trace" $ngxExe @(
|
$traceArguments = @(
|
||||||
'--hip-root', $HipRoot, '--cuda-dll', (Join-Path $ZludaRoot 'trace/nvcuda.dll'),
|
'--hip-root', $HipRoot, '--cuda-dll', (Join-Path $ZludaRoot 'trace/nvcuda.dll'),
|
||||||
'--ngx-core', $localCore, '--dlss-dll', $localDlss,
|
'--ngx-core', $localCore, '--dlss-dll', $localDlss,
|
||||||
'--nvapi-dll', (Join-Path $PackageRoot 'nvapi-compat/nvapi64.dll'),
|
'--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:D4R_NVAPI_BACKEND = $originalNvapi
|
||||||
$env:PYTHONPATH = $originalPythonPath
|
$env:PYTHONPATH = $originalPythonPath
|
||||||
$env:D4R_ZLUDA_VERBOSE = $originalVerbose
|
$env:D4R_ZLUDA_VERBOSE = $originalVerbose
|
||||||
|
$env:ZLUDA_CACHE_DIR = $originalCache
|
||||||
|
foreach ($setting in $originalCodegen.Keys) { [Environment]::SetEnvironmentVariable($setting, $originalCodegen[$setting], 'Process') }
|
||||||
$archive = "$OutputDirectory.zip"
|
$archive = "$OutputDirectory.zip"
|
||||||
Compress-Archive -LiteralPath $OutputDirectory -DestinationPath $archive -Force
|
Compress-Archive -LiteralPath $OutputDirectory -DestinationPath $archive -Force
|
||||||
Write-Host "Diagnostic bundle: $archive"
|
Write-Host "Diagnostic bundle: $archive"
|
||||||
|
|||||||
@@ -131,5 +131,21 @@ public:
|
|||||||
copy.widthBytes = rowBytes_; copy.height = height_;
|
copy.widthBytes = rowBytes_; copy.height = height_;
|
||||||
api_.cuda.check(api_.cuMemcpy2D_v2(©), "cuMemcpy2D_v2(diagnostic array to host)");
|
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)");
|
||||||
|
}
|
||||||
};
|
};
|
||||||
}
|
}
|
||||||
@@ -38,10 +38,30 @@ int main(int argc, char** argv) {
|
|||||||
image.upload(input.data());
|
image.upload(input.data());
|
||||||
image.download(output.data());
|
image.download(output.data());
|
||||||
if (input != output) throw std::runtime_error("CUDA array storage roundtrip mismatch");
|
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<uint8_t> 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");
|
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;
|
return 0;
|
||||||
} catch (const std::exception& e) {
|
} catch (const std::exception& e) {
|
||||||
std::fprintf(stderr, "FAIL CUDA_IMAGES %s\n", e.what());
|
std::fprintf(stderr, "FAIL CUDA_IMAGES %s\n", e.what());
|
||||||
|
|||||||
@@ -0,0 +1,199 @@
|
|||||||
|
#include "d3d12_external.h"
|
||||||
|
#include "ngx_parameters.h"
|
||||||
|
#include "ngx_cuda_evaluate.h"
|
||||||
|
#include <dxgi1_6.h>
|
||||||
|
#include <memory>
|
||||||
|
|
||||||
|
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<IDXGIFactory4> factory; dx(CreateDXGIFactory1(IID_PPV_ARGS(factory.GetAddressOf())), "CreateDXGIFactory1");
|
||||||
|
ComPtr<IDXGIAdapter1> 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<ID3D12Device> device;
|
||||||
|
dx(D3D12CreateDevice(adapter.Get(), D3D_FEATURE_LEVEL_12_0, IID_PPV_ARGS(device.GetAddressOf())), "D3D12CreateDevice");
|
||||||
|
ComPtr<ID3D12CommandQueue> queue;
|
||||||
|
D3D12_COMMAND_QUEUE_DESC qd{}; qd.Type = D3D12_COMMAND_LIST_TYPE_DIRECT;
|
||||||
|
dx(device->CreateCommandQueue(&qd, IID_PPV_ARGS(queue.GetAddressOf())), "CreateCommandQueue");
|
||||||
|
ComPtr<ID3D12Fence> 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<ID3D12CommandAllocator> allocator;
|
||||||
|
ComPtr<ID3D12GraphicsCommandList> 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<ID3D12Resource> 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<uint8_t*>(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<SharedPlane> planes[5];
|
||||||
|
std::unique_ptr<d4r::cuda::Image> 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<SharedPlane>(external, device.Get(), desc);
|
||||||
|
arrays[i] = std::make_unique<d4r::cuda::Image>(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<uintptr_t>(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<uint8_t> 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<uint8_t*>(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<Init>("NVSDK_NGX_D3D12_Init_Ext")(241534723ull, directory.c_str(), device.Get(), 0x15, nullptr), "D3D12 Init");
|
||||||
|
auto shutdown = shim.symbol<Shutdown>("NVSDK_NGX_D3D12_Shutdown");
|
||||||
|
struct ShutdownCleanup { Shutdown fn; ~ShutdownCleanup() { (void)fn(); } } shutdownCleanup{shutdown};
|
||||||
|
void* parameters = nullptr; check(shim.symbol<Allocate>("NVSDK_NGX_D3D12_AllocateParameters")(¶meters), "AllocateParameters");
|
||||||
|
auto destroy = shim.symbol<Destroy>("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<Create>("NVSDK_NGX_D3D12_CreateFeature")(nullptr, 1, parameters, &handle), "D3D12 CreateFeature");
|
||||||
|
auto release = shim.symbol<Release>("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<Evaluate>("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<uint16_t> pixels(size_t(outWidth) * outHeight * 4);
|
||||||
|
for (unsigned y = 0; y < outHeight; ++y) std::memcpy(pixels.data() + size_t(y) * outWidth * 4,
|
||||||
|
static_cast<uint8_t*>(raw) + y * textures[3].footprint.Footprint.RowPitch, outWidth * 8);
|
||||||
|
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; }
|
||||||
|
}
|
||||||
@@ -0,0 +1,190 @@
|
|||||||
|
#pragma once
|
||||||
|
#if defined(__MINGW32__)
|
||||||
|
#define WIDL_EXPLICIT_AGGREGATE_RETURNS
|
||||||
|
#endif
|
||||||
|
#include "hip_api.h"
|
||||||
|
#include <d3d12.h>
|
||||||
|
#include <wrl/client.h>
|
||||||
|
|
||||||
|
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<decltype(&::name)>(#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<ID3D12Resource> 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<ID3D12Resource> 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<ID3D12Resource> 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<ID3D12Fence> fence;
|
||||||
|
SharedTimeline(ExternalApi& api, ID3D12Device* device) : api_(api) {
|
||||||
|
dx(device->CreateFence(0, D3D12_FENCE_FLAG_SHARED, IID_PPV_ARGS(fence.GetAddressOf())), "CreateFence(D3D12/HIP)");
|
||||||
|
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");
|
||||||
|
}
|
||||||
|
}
|
||||||
|
};
|
||||||
|
}
|
||||||
@@ -0,0 +1,111 @@
|
|||||||
|
// Native Windows backend. The Linux/Wine implementation remains separate.
|
||||||
|
#include "ngx_windows_runtime.h"
|
||||||
|
#include <unordered_map>
|
||||||
|
#include <functional>
|
||||||
|
|
||||||
|
namespace {
|
||||||
|
std::mutex apiMutex;
|
||||||
|
std::shared_ptr<d4r::win::Runtime> runtime;
|
||||||
|
struct Handle { unsigned Id; };
|
||||||
|
unsigned nextId = 1;
|
||||||
|
std::unordered_map<Handle*, std::shared_ptr<d4r::win::Feature>> features;
|
||||||
|
std::unordered_map<void*, bool> parameters;
|
||||||
|
constexpr unsigned failure = 0xBAD00002, invalid = 0xBAD00005;
|
||||||
|
unsigned call(const std::function<unsigned()>& 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<std::mutex> 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<std::mutex> lock(apiMutex);
|
||||||
|
return call([&] {
|
||||||
|
if (!runtime) runtime = std::make_shared<d4r::win::Runtime>(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<std::mutex> lock(apiMutex); return call([&] { return allocate(p, false); });
|
||||||
|
}
|
||||||
|
API unsigned NVSDK_NGX_D3D12_GetCapabilityParameters(void** p) {
|
||||||
|
std::lock_guard<std::mutex> 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<std::mutex> 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<std::mutex> lock(apiMutex);
|
||||||
|
return call([&] {
|
||||||
|
if (!runtime || feature != 1 || !p || !out) return invalid;
|
||||||
|
auto value = std::make_shared<d4r::win::Feature>(runtime, p);
|
||||||
|
auto handle = std::make_unique<Handle>(); 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<d4r::win::Feature> feature;
|
||||||
|
{ std::lock_guard<std::mutex> 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<d4r::win::Feature> feature;
|
||||||
|
{ std::lock_guard<std::mutex> lock(apiMutex);
|
||||||
|
auto it = features.find(handle); if (it == features.end()) return invalid;
|
||||||
|
feature = std::move(it->second); features.erase(it); delete handle; }
|
||||||
|
return call([&] { feature.reset(); return 1u; });
|
||||||
|
}
|
||||||
|
API unsigned NVSDK_NGX_D3D12_Shutdown() {
|
||||||
|
std::lock_guard<std::mutex> 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(); }
|
||||||
@@ -0,0 +1,56 @@
|
|||||||
|
#pragma once
|
||||||
|
#include <cstdint>
|
||||||
|
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));
|
||||||
|
}
|
||||||
|
}
|
||||||
@@ -0,0 +1,270 @@
|
|||||||
|
#pragma once
|
||||||
|
#include "d3d12_external.h"
|
||||||
|
#include "cuda_image_api.h"
|
||||||
|
#include "ngx_parameters.h"
|
||||||
|
#include <chrono>
|
||||||
|
#include <memory>
|
||||||
|
#include <mutex>
|
||||||
|
|
||||||
|
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<diag::Library> nvapi;
|
||||||
|
diag::SearchDirectory featureSearch;
|
||||||
|
std::unique_ptr<diag::Library> dlss, core;
|
||||||
|
ExternalApi external;
|
||||||
|
cuda::ImageApi images;
|
||||||
|
ComPtr<ID3D12Device> device;
|
||||||
|
CUdevice ordinal = -1;
|
||||||
|
CUcontext context = nullptr;
|
||||||
|
bool initialized = false;
|
||||||
|
std::unique_ptr<SharedTimeline> 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<SharedTimeline>(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::Library>(diag::wide(env_path("D4R_NVAPI_DLL")));
|
||||||
|
dlss = std::make_unique<diag::Library>(diag::wide(env_path("D4R_DLSS_DLL")));
|
||||||
|
core = std::make_unique<diag::Library>(diag::wide(env_path("D4R_NGX_CORE")));
|
||||||
|
allocate = core->symbol<Allocate>("NVSDK_NGX_CUDA_AllocateParameters");
|
||||||
|
capabilities = core->symbol<Allocate>("NVSDK_NGX_CUDA_GetCapabilityParameters");
|
||||||
|
destroy = core->symbol<Destroy>("NVSDK_NGX_CUDA_DestroyParameters");
|
||||||
|
create = core->symbol<Create>("NVSDK_NGX_CUDA_CreateFeature");
|
||||||
|
evaluate = core->symbol<Evaluate>("NVSDK_NGX_CUDA_EvaluateFeature");
|
||||||
|
release = core->symbol<Release>("NVSDK_NGX_CUDA_ReleaseFeature");
|
||||||
|
shutdown = core->symbol<Shutdown>("NVSDK_NGX_CUDA_Shutdown");
|
||||||
|
ngx_check(core->symbol<Init>("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<Runtime> rt_;
|
||||||
|
void* parameters_ = nullptr;
|
||||||
|
void* handle_ = nullptr;
|
||||||
|
CUdeviceptr scratch_ = 0;
|
||||||
|
ComPtr<ID3D12CommandQueue> queue_;
|
||||||
|
struct Plane {
|
||||||
|
ComPtr<ID3D12Resource> texture;
|
||||||
|
D3D12_RESOURCE_DESC desc{};
|
||||||
|
std::unique_ptr<SharedPlane> shared;
|
||||||
|
std::unique_ptr<cuda::Image> 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> runtime, void* parameters) : rt_(std::move(runtime)) {
|
||||||
|
std::lock_guard<std::mutex> 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<void*>(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<std::mutex> 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<std::mutex> 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<ID3D12Device> queueDevice;
|
||||||
|
dx(queue->GetDevice(IID_PPV_ARGS(queueDevice.GetAddressOf())), "GetDevice(queue)");
|
||||||
|
ComPtr<IUnknown> 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<ID3D12CommandAllocator> inputAllocator, outputAllocator;
|
||||||
|
ComPtr<ID3D12GraphicsCommandList> inputList, outputList;
|
||||||
|
auto list = [&](ComPtr<ID3D12CommandAllocator>& allocator, ComPtr<ID3D12GraphicsCommandList>& 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<SharedPlane>(rt_->external, rt_->device.Get(), desc);
|
||||||
|
plane.image = std::make_unique<cuda::Image>(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<uintptr_t>(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<size_t>(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<uint8_t> 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<uintptr_t>(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<double, std::milli>(std::chrono::steady_clock::now() - begin).count());
|
||||||
|
}
|
||||||
|
void record_output(ID3D12GraphicsCommandList* list, void* parameters) {
|
||||||
|
planes_[3].shared->copy_output(list, planes_[3].texture.Get(), D3D12_RESOURCE_STATES(ngx::uint_value(parameters, "D4R.Output.State", D3D12_RESOURCE_STATE_UNORDERED_ACCESS)));
|
||||||
|
}
|
||||||
|
};
|
||||||
|
}
|
||||||
Reference in new issue
Block a user