From 973eb5ecd2fb184051f157a9fd214b36ce7e45de Mon Sep 17 00:00:00 2001 From: xdfnx-dev Date: Fri, 2 Oct 2026 12:06:15 +0300 Subject: [PATCH] Batch Windows DLSS GPU inputs and add deferred kernel diagnostics --- CMakeLists.txt | 4 + docs/windows-performance.md | 87 +++++- docs/windows-rdna4-port.md | 48 ++- .../zluda/0016-deferred-kernel-profile.patch | 282 ++++++++++++++++++ .../0017-asynchronous-2d-array-copies.patch | 50 ++++ scripts/windows/profile-windows-rdna4.ps1 | 1 + scripts/windows/test-async-k.ps1 | 22 +- .../windows/test-deferred-kernel-profile.ps1 | 101 +++++++ scripts/windows/windows-game.ps1 | 17 +- tools/windows/cuda_driver_probe.cpp | 32 +- tools/windows/cuda_image_api.h | 10 +- tools/windows/cuda_image_probe.cpp | 20 +- tools/windows/ngx_windows_runtime.h | 17 +- tools/windows/pixel_conversion.h | 10 +- tools/windows/profile_report.py | 28 +- 15 files changed, 692 insertions(+), 37 deletions(-) create mode 100644 patches/zluda/0016-deferred-kernel-profile.patch create mode 100644 patches/zluda/0017-asynchronous-2d-array-copies.patch create mode 100644 scripts/windows/test-deferred-kernel-profile.ps1 diff --git a/CMakeLists.txt b/CMakeLists.txt index e5bfcb8..4049b66 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -245,6 +245,10 @@ 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 cuda_images_async COMMAND d4r_cuda_image_probe + --hip-root "${D4R_HIP_ROOT}" --cuda-dll "${D4R_ZLUDA_ROOT}/nvcuda.dll" --iterations 32) + set_tests_properties(cuda_images_async PROPERTIES TIMEOUT 90 LABELS "hardware;zluda;descriptors;async" + ENVIRONMENT "D4R_DIAG_ASYNC_ARRAY_COPY=1") 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") diff --git a/docs/windows-performance.md b/docs/windows-performance.md index cec2521..6190659 100644 --- a/docs/windows-performance.md +++ b/docs/windows-performance.md @@ -4,7 +4,7 @@ Hardware: RX 9070 XT / gfx1201, Windows 11, display driver 32.0.31041.1004. Runtime: TheRock 10.2.0a20260929, HIP 7.17.26386; ZLUDA b0161a4 with `D4R_ZLUDA_WMMA=1`, FP8 widening, native FP8 disabled, f16 reference rounding. -The optional profiler uses HIP events on the actual launch stream, including +The original serializing profiler uses HIP events on the actual launch stream, including separate measurements of one-time weight preparation. It synchronizes each sampled launch and changes scheduling: these numbers are not application FPS. HIP occupancy predictions are not measured hardware occupancy. The native @@ -192,6 +192,91 @@ without backend errors; ordinary recording intervals later return to about 16 ms. A stable foreground-scene capture is needed before optimizing from the ETW metrics. Raw files: `test-results/silent-hill2-k-presentmon-4k`. +The following 30-second capture starts after the user confirms the same +foreground scene is ready. All 1902 presents use Hardware Independent Flip +and SyncInterval 0; input/output resolution remains 2259x1271 -> 3840x2160. +The mean present interval is 15.730 ms (**63.574 FPS**). PresentMon estimates +mean GPU busy/wait at 14.504/1.110 ms, GPU time at 15.614 ms and CPU wait at +0.052 ms. These frame-attributed metrics do not equal a hardware utilization +sensor and do not establish a physical CPU limit. FSR4 has not been captured +in this same scene, so its previously reported 80+ FPS remains a user reading. +Raw files: `test-results/silent-hill2-k-steady-profile-4k`. + +The containing ten-minute diagnostic completes 36591 K frames, with zero +backend failures, CPU image copies, previous-frame outputs or recorded game +crashes. It deliberately stops at its time bound. Output scans and serializing +kernel events are off. Whole-run median boundary timestamps measure input +copies 0.063 ms, external span 5.399 ms and output copies 0.212 ms. The external +span includes HIP execution and host/driver scheduling, not only transformer +arithmetic. These medians cover the whole run; the PresentMon numbers above +come from the explicitly selected 30-second scene. + +### Deferred kernel events and batched inputs + +The new optional `-ProfileKernelsDeferred` samples every seventeenth launch +by default (`-KernelProfileEvery`). Its thread-local pool is bounded to 64 +event pairs. It queries ready events before selected launches and after the +application's existing context synchronization, reusing a pair only after +reading its old timestamps. It adds no explicit event/stream/device completion +wait; timestamps, queries and logging still add overhead. Stream capture and +legacy default-stream launches are skipped by default. `-ProfileLegacyStream` +allows the latter explicitly, retaining their legacy cross-stream ordering. +This opt-in is necessary to profile this DLSS K binary, whose measured launches +use the legacy default stream. `completion_ms` now denotes collection lag for +deferred records. The report separates serializing/deferred records and retains +invalid timestamps and skipped samples outside kernel timing statistics. + +HIP 7.17.26386 reports some negative elapsed times in the tiny queued PTX +diagnostic, including intervals exceeding one microsecond. These are logged as +`D4R_KERNEL_PROFILE_INVALID`; they are not clamped or used as valid GPU samples. +The cause is not established, and nonnegative samples alone cannot prove a +runtime's timestamp accuracy. The diagnostic verifies all 516 PTX launches +across primary/created contexts and individual/queued submission, checks final +device output and guards, accounts for every collected event and verifies safe +default-stream skipping. Results: `test-results/zluda-deferred-profile-final-ptx`. +The same diagnostic with stable HIP SDK 7.2 also verifies all 516 outputs and +reproduces negative intervals (`test-results/zluda-deferred-profile-stable-ptx`). +A native K dec4 sample is negative too; this is not confined to tiny PTX kernels. +Do not infer hardware/kernel utilization or a speedup from these event samples. +HIP event semantics: [AMD event management](https://rocm.docs.amd.com/projects/HIP/en/latest/reference/hip_runtime_api/modules/event_management.html). + +The opt-in `-BatchInputCopies` (`D4R_BATCH_INPUT_COPIES=1`) enqueues native +pixel decoding and device-to-array input copies on the shared HIP/ZLUDA legacy +default stream. One all-stream completion remains before NGX can use any input +on its own streams. The independent D3D12 input-fence check, NGX completion and +same-frame output dependency remain. Output transfers retain their existing +completion. Every image stays in VRAM. The ZLUDA implementation adds +`cuMemcpy2DAsync_v2` using `hipMemcpyParam2DAsync`; format-80 array emulation +is explicitly rejected because its synchronous implementation uses temporary +host storage. The Windows path supplies canonical FP16/FP32 arrays instead. +This experiment is disabled by default until the correctness and game timing +gates below pass; there is no claimed FPS improvement yet. + +```powershell +.\scripts\windows\test-deferred-kernel-profile.ps1 +.\scripts\windows\test-async-k.ps1 -BatchInputCopies -ProfileGpuBoundary -ZludaRoot "$PWD\dist\zluda-windows-deferred-profile" +``` + +The rebuilt Windows shim and new ZLUDA pass all seventeen CTest gates, +including the added asynchronous pitched FP16/FP32 array-copy test and +padding checks. Completed boundary timestamps are also read before reuse if +the copy fence advances between the first completion scan and slot selection; +this prevents diagnostic records being overwritten without adding a wait. + +The depth/stencil and packed-format burst/recreation gates pass 24 full K/4K +frames. All twelve candidate RGB outputs match fresh unbatched synchronous +controls exactly, with finite RGBA. The former also enables deferred legacy +events; invalid timestamps remain separate. These short correctness fixtures +include initialization and do not establish steady-state FPS or a copy-stage +speedup. Raw results: `test-results/k-batch-deferred-4k` and +`test-results/k-batch-packed-4k`. ZLUDA candidate source is `67127dd`, isolated +in `dist/zluda-windows-deferred-profile`; patches 0016/0017 and all preceding +patches apply to the clean pinned base. The public prerelease is unchanged. +After committing/rebuilding ZLUDA with a clean worktree, the twelve-frame +four-way K/4K gate passes again with exact candidate RGB and finite RGBA +(`test-results/k-batch-committed-source-4k`). This warms the committed runtime's +JIT cache before the game comparison. + Further game coverage (HIP kernel profiling disabled, GPU output checks enabled): M at 1920x1080 completes 4833 frames, all finite. K at 3840x2160 completes 4685 frames, all finite. Neither has a backend failure, CPU image copy or an diff --git a/docs/windows-rdna4-port.md b/docs/windows-rdna4-port.md index db8dd30..9862e13 100644 --- a/docs/windows-rdna4-port.md +++ b/docs/windows-rdna4-port.md @@ -2,11 +2,13 @@ Branch: `windows-rdna4`. Target: Windows 11 x64, RX 9070 XT, **gfx1201**. Priority: correct K, correct M, native Windows, same-frame output, then speed. -Upstream integration: [draft PR #11](https://github.com/countervolts/d4r/pull/11) -targets `countervolts/d4r:windows` from `xdfnx-dev/d4r:upstream-windows`. +Upstream integration: [PR #11](https://github.com/countervolts/d4r/pull/11) +is merged into `countervolts/d4r:windows` at `c522101`. The integration branch retains upstream's README with a Windows status section; -the fork's `main` keeps the dedicated English Windows README. The PR is open, -mergeable and allows maintainer edits. K performance remains a separate open gate. +the fork's `main` keeps the dedicated English Windows README. The maintainer +invitation is accepted; GitHub confirms `xdfnx-dev` has push access and the +Windows branch is unprotected. Further validated Windows changes can go there +directly. K performance remains a separate open gate. This file records current results followed by the dated development history. Standalone DLSS K/M works on the real Windows GPU, including the patched OptiScaler frontend. Silent Hill 2 now renders through the Windows backend; gameplay and output @@ -15,8 +17,9 @@ explicit limits; the validated package keeps conservative arithmetic. ## Milestone and gates -**Current milestone: K performance remains open.** The user reports 49–51 FPS -and about 53% GPU utilization in 4K Silent Hill 2, compared with roughly 80+ +**Current milestone: K performance remains open.** Async submission improves +the user's 4K Silent Hill 2 reading to 62 FPS / 63% GPU utilization. A subsequent +controlled 30-second foreground-scene capture records 63.574 FPS, compared with roughly 80+ FPS / 100% GPU utilization with FSR4 on this system. Functional K/M and the reproducible archive do not close this gate. Prioritize K queue scheduling and host synchronization; postpone further M/FP8 optimization. These user readings @@ -82,7 +85,38 @@ measures median input/output copies at 0.061/0.208 ms and external span at as HIP. PresentMon 2.6.0 capture/report tooling is added; its first trace contains changing presentation modes and is not a controlled capture of the user's 62 FPS scene. See `windows-performance.md` for exact metrics and -the stable foreground-scene gate that remains open. +the initial capture limitations. A subsequent user-confirmed foreground-scene +capture completes 1902 presents at 63.574 FPS, all Independent Flip / SyncInterval +0. PresentMon mean GPU busy/wait is 14.504/1.110 ms; this is not sensor +utilization. The containing ten-minute run completes 36591 K frames without +backend errors; median external span is 5.399 ms. Log bundle: +`test-results/silent-hill2-k-steady-profile-4k`. + +The next candidate adds bounded, deferred HIP-event sampling and opt-in batched +GPU input conversion/array copies. It retains one all-stream input completion +before NGX and every D3D12/output dependency. All seventeen CTest gates pass, +including asynchronous pitched FP16/FP32 array copies and guard checks. +The deferred profiler's tiny PTX diagnostic verifies 516 launches and safe +default-stream skipping; negative HIP elapsed times are retained as invalid +samples, with cause unresolved. This DLSS K binary uses the legacy default +stream, which needs explicit opt-in for timestamp sampling. The depth/stencil +and packed-format burst/recreation gates pass 24 full K/4K frames: all twelve +candidate RGB images match their fresh unbatched synchronous controls exactly, +with finite RGBA. Results: `test-results/k-batch-deferred-4k` and +`test-results/k-batch-packed-4k`. The game performance gate is pending. The candidate remains +disabled by default; see `windows-performance.md` for semantics and commands. + +Current candidate source is ZLUDA `67127dde599879ada022d6439383d682e09e0b36` +in `dist/zluda-windows-deferred-profile`, containing patches 0016/0017. +All sixteen patches (0002 through 0017) apply to clean pinned base `ee2f25a`. +The committed ZLUDA is rebuilt with a clean worktree and the four-way K/4K +gate is repeated: twelve frames pass, all six candidate RGB outputs match +their controls exactly, and all six async boundary timestamp records survive +slot reuse/release. Results: `test-results/k-batch-committed-source-4k`. +The same negative HIP event-timing issue reproduces with stable SDK 7.2 as +well as TheRock; PTX output/guards remain correct. It also affects a native K +dec4 timing sample, so valid-looking events alone are not a timing-accuracy +gate. No hardware utilization or kernel speedup is inferred from these samples. The public prerelease `windows-rdna4-dev-20261002-a346d76` is published to `xdfnx-dev/d4r`, containing source commit `a346d76` and ZLUDA `a1c506f`. diff --git a/patches/zluda/0016-deferred-kernel-profile.patch b/patches/zluda/0016-deferred-kernel-profile.patch new file mode 100644 index 0000000..9bfb224 --- /dev/null +++ b/patches/zluda/0016-deferred-kernel-profile.patch @@ -0,0 +1,282 @@ +From 26a43f7047ab933a4afb5ebdd26d5ba9e51e58a3 Mon Sep 17 00:00:00 2001 +From: xdfnx-dev +Date: Fri, 2 Oct 2026 11:50:45 +0300 +Subject: [PATCH 16/17] Sample HIP kernel events without per-launch completion + waits + +--- + zluda/src/impl/context.rs | 4 +- + zluda/src/impl/d4r_profile.rs | 204 ++++++++++++++++++++++++++++------ + zluda/src/impl/module.rs | 1 + + 3 files changed, 171 insertions(+), 38 deletions(-) + +diff --git a/zluda/src/impl/context.rs b/zluda/src/impl/context.rs +index 9ccb372..374d383 100644 +--- a/zluda/src/impl/context.rs ++++ b/zluda/src/impl/context.rs +@@ -165,7 +165,7 @@ pub(crate) fn set_cache_config(config: CUfunc_cache) -> CUresult { + } + + pub(crate) fn synchronize() -> hipError_t { +- unsafe { hipDeviceSynchronize() } ++ unsafe { hipDeviceSynchronize()?; super::d4r_profile::collect_completed() } + } + + pub(crate) fn synchronize_v2(ctx: &Context) -> hipError_t { +@@ -177,7 +177,7 @@ pub(crate) fn synchronize_v2(ctx: &Context) -> hipError_t { + } else { + None + }; +- unsafe { hipDeviceSynchronize() } ++ unsafe { hipDeviceSynchronize()?; super::d4r_profile::collect_completed() } + } + + struct SetDeviceOnDrop(i32); +diff --git a/zluda/src/impl/d4r_profile.rs b/zluda/src/impl/d4r_profile.rs +index 33ba5f6..1b9d248 100644 +--- a/zluda/src/impl/d4r_profile.rs ++++ b/zluda/src/impl/d4r_profile.rs +@@ -1,15 +1,29 @@ +-//! Optional, deliberately serializing HIP-event diagnostics. Not a production +-//! scheduler or a claim about application FPS / measured hardware occupancy. ++//! Optional HIP-event diagnostics. Deferred mode never waits for its events, ++//! skips the legacy default stream and bounds outstanding events per thread. ++//! Timestamp insertion still adds overhead; this is not a production scheduler. + use hip_runtime_sys::*; +-use std::{ffi::CStr, ptr, time::Instant}; ++use std::{cell::RefCell, ffi::{CStr, CString}, ptr, sync::OnceLock, time::Instant}; + +-pub(super) struct Launch { +- start: hipEvent_t, +- end: hipEvent_t, +- host: Instant, ++#[derive(Clone, Copy)] ++enum Mode { Off, Serial, Deferred(u64) } ++fn mode() -> Mode { ++ static MODE: OnceLock = OnceLock::new(); ++ *MODE.get_or_init(|| { ++ if std::env::var_os("D4R_ZLUDA_PROFILE_DEFERRED").is_some() { ++ let every = std::env::var("D4R_ZLUDA_PROFILE_EVERY").ok() ++ .and_then(|v| v.parse::().ok()).filter(|n| *n > 0).unwrap_or(17); ++ Mode::Deferred(every) ++ } else if std::env::var_os("D4R_ZLUDA_PROFILE").is_some() { Mode::Serial } ++ else { Mode::Off } ++ }) ++} ++fn allow_legacy() -> bool { ++ static ALLOW: OnceLock = OnceLock::new(); ++ *ALLOW.get_or_init(|| std::env::var_os("D4R_ZLUDA_PROFILE_ALLOW_LEGACY").is_some()) + } + +-impl Drop for Launch { ++struct Events { start: hipEvent_t, end: hipEvent_t } ++impl Drop for Events { + fn drop(&mut self) { + unsafe { + if !self.end.is_null() { let _ = hipEventDestroy(self.end); } +@@ -17,44 +31,162 @@ impl Drop for Launch { + } + } + } ++impl Events { ++ unsafe fn new() -> Result { ++ let mut result = Self { start: ptr::null_mut(), end: ptr::null_mut() }; ++ hipEventCreateWithFlags(&mut result.start, 0)?; ++ hipEventCreateWithFlags(&mut result.end, 0)?; ++ Ok(result) ++ } ++} ++ ++struct Description { ++ name: CString, backend: String, phase: String, grid: [u32;3], block: [u32;3], ++ shared: u32, lds: i32, regs: i32, private: i32, max_threads: i32, blocks_per_multiprocessor: i32, ++} ++impl Description { ++ unsafe fn new(function: hipFunction_t, name: &CStr, backend: &str, phase: &str, ++ grid: [u32;3], block: [u32;3], shared: u32) -> Result { ++ let mut d = Self { name: name.to_owned(), backend: backend.to_owned(), phase: phase.to_owned(), ++ grid, block, shared, lds: 0, regs: 0, private: 0, max_threads: 0, blocks_per_multiprocessor: 0 }; ++ // Copy metadata while the function/module is alive. Collection never ++ // dereferences a function handle that the application may have unloaded. ++ hipFuncGetAttribute(&mut d.lds, hipFunction_attribute::HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES, function)?; ++ hipFuncGetAttribute(&mut d.regs, hipFunction_attribute::HIP_FUNC_ATTRIBUTE_NUM_REGS, function)?; ++ hipFuncGetAttribute(&mut d.private, hipFunction_attribute::HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES, function)?; ++ hipFuncGetAttribute(&mut d.max_threads, hipFunction_attribute::HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, function)?; ++ let threads = block.iter().try_fold(1u32, |a, b| a.checked_mul(*b)) ++ .filter(|n| *n <= i32::MAX as u32).ok_or(hipErrorCode_t::InvalidValue)?; ++ hipModuleOccupancyMaxActiveBlocksPerMultiprocessor(&mut d.blocks_per_multiprocessor, ++ function, threads as i32, shared as usize)?; ++ Ok(d) ++ } ++ fn print(&self, gpu_ms: f32, enqueue_ms: f64, completion_ms: f64, serializing: bool, stream: usize) { ++ if !gpu_ms.is_finite() || gpu_ms < 0. { ++ eprintln!("D4R_KERNEL_PROFILE_INVALID kernel={:?} backend={} phase={} gpu_ms={gpu_ms:.6} reason=invalid_elapsed_time serializing={} deferred={} stream=0x{stream:x}", ++ self.name, self.backend, self.phase, u32::from(serializing), u32::from(!serializing)); ++ return; ++ } ++ eprintln!("D4R_KERNEL_PROFILE kernel={:?} backend={} phase={} gpu_ms={gpu_ms:.6} enqueue_ms={enqueue_ms:.6} completion_ms={completion_ms:.6} grid={},{},{} block={},{},{} static_lds_bytes={} dynamic_lds_bytes={} regs={} private_bytes={} max_threads={} predicted_blocks_per_multiprocessor={} serializing={} deferred={} stream=0x{stream:x} legacy_stream={}", ++ self.name, self.backend, self.phase, self.grid[0], self.grid[1], self.grid[2], ++ self.block[0], self.block[1], self.block[2], self.lds, self.shared, self.regs, ++ self.private, self.max_threads, self.blocks_per_multiprocessor, ++ u32::from(serializing), u32::from(!serializing), u32::from(stream <= 1)); ++ } ++} ++ ++struct Pending { launch: Launch, description: Description, enqueue_ms: f64, stream: usize } ++const CAPACITY: usize = 64; ++struct Deferred { ++ calls: u64, skipped: u64, pending: Vec, free: Vec<(i32, Events)>, ++} ++thread_local! { ++ static DEFERRED: RefCell = RefCell::new(Deferred { ++ calls: 0, skipped: 0, pending: Vec::new(), free: Vec::new(), ++ }); ++} ++ ++impl Deferred { ++ unsafe fn collect(&mut self, device: i32) -> hipError_t { ++ let mut index = 0; ++ while index < self.pending.len() { ++ let p = &self.pending[index]; ++ if p.launch.device != device { index += 1; continue; } ++ match hipEventQuery(p.launch.events.end) { ++ Err(e) if e == hipErrorCode_t::NotReady => { index += 1; continue; } ++ result => result?, ++ } ++ let mut gpu_ms = 0.; ++ hipEventElapsedTime(&mut gpu_ms, p.launch.events.start, p.launch.events.end)?; ++ p.description.print(gpu_ms, p.enqueue_ms, p.launch.host.elapsed().as_secs_f64() * 1000., false, p.stream); ++ let p = self.pending.swap_remove(index); ++ // Re-record only after the old timestamps have been collected. ++ self.free.push((device, p.launch.events)); ++ } ++ Ok(()) ++ } ++} ++ ++/// Called only after the application's existing context synchronization, or ++/// before a selected launch. No event/stream/device wait is added here. ++pub(super) unsafe fn collect_completed() -> hipError_t { ++ if !matches!(mode(), Mode::Deferred(_)) { return Ok(()); } ++ DEFERRED.with(|state| { ++ let mut state = state.borrow_mut(); ++ if state.pending.is_empty() { return Ok(()); } ++ let mut device = 0; hipGetDevice(&mut device)?; ++ state.collect(device) ++ }) ++} ++ ++pub(super) struct Launch { ++ events: Events, ++ host: Instant, ++ device: i32, ++ deferred: bool, ++} + + impl Launch { + pub(super) unsafe fn begin(stream: hipStream_t) -> Result, hipErrorCode_t> { +- if std::env::var_os("D4R_ZLUDA_PROFILE").is_none() { return Ok(None); } +- let mut result = Self { start: ptr::null_mut(), end: ptr::null_mut(), host: Instant::now() }; +- hipEventCreateWithFlags(&mut result.start, 0)?; +- hipEventCreateWithFlags(&mut result.end, 0)?; +- hipEventRecord(result.start, stream)?; +- result.host = Instant::now(); +- Ok(Some(result)) ++ let mode = mode(); ++ if matches!(mode, Mode::Off) { return Ok(None); } ++ if let Mode::Deferred(every) = mode { ++ return DEFERRED.with(|state| { ++ let mut state = state.borrow_mut(); ++ state.calls = state.calls.wrapping_add(1); ++ if (state.calls - 1) % every != 0 { return Ok(None); } ++ // HIP timing events on the legacy default stream can order ++ // other streams. Do not introduce that dependency in this mode. ++ if stream.0 as usize <= 1 && !allow_legacy() { ++ state.skipped += 1; ++ if state.skipped.is_power_of_two() { ++ eprintln!("D4R_KERNEL_PROFILE_SKIPPED reason=legacy_default_stream count={} serializing=0", state.skipped); ++ } ++ return Ok(None); ++ } ++ let mut device = 0; hipGetDevice(&mut device)?; ++ state.collect(device)?; ++ let mut capture = hipStreamCaptureStatus::hipStreamCaptureStatusNone; ++ hipStreamIsCapturing(stream, &mut capture)?; ++ if capture != hipStreamCaptureStatus::hipStreamCaptureStatusNone { ++ state.skipped += 1; ++ eprintln!("D4R_KERNEL_PROFILE_SKIPPED reason=stream_capture count={} serializing=0", state.skipped); ++ return Ok(None); ++ } ++ let events = if let Some(index) = state.free.iter().position(|(d, _)| *d == device) { ++ state.free.swap_remove(index).1 ++ } else if state.pending.len() + state.free.len() < CAPACITY { ++ Events::new()? ++ } else { ++ state.skipped += 1; ++ eprintln!("D4R_KERNEL_PROFILE_SKIPPED reason=capacity count={} capacity={CAPACITY} serializing=0", state.skipped); ++ return Ok(None); ++ }; ++ hipEventRecord(events.start, stream)?; ++ Ok(Some(Self { events, host: Instant::now(), device, deferred: true })) ++ }); ++ } ++ let events = Events::new()?; ++ hipEventRecord(events.start, stream)?; ++ Ok(Some(Self { events, host: Instant::now(), device: -1, deferred: false })) + } + + pub(super) unsafe fn finish(self, function: hipFunction_t, name: &CStr, backend: &str, + phase: &str, grid: [u32;3], block: [u32;3], shared: u32, stream: hipStream_t) -> hipError_t { + let enqueue_ms = self.host.elapsed().as_secs_f64() * 1000.; +- hipEventRecord(self.end, stream)?; +- hipEventSynchronize(self.end)?; ++ hipEventRecord(self.events.end, stream)?; ++ let description = Description::new(function, name, backend, phase, grid, block, shared)?; ++ if self.deferred { ++ DEFERRED.with(|state| state.borrow_mut().pending.push(Pending { ++ launch: self, description, enqueue_ms, stream: stream.0 as usize, ++ })); ++ return Ok(()); ++ } ++ hipEventSynchronize(self.events.end)?; + let completion_ms = self.host.elapsed().as_secs_f64() * 1000.; + let mut gpu_ms = 0.; +- hipEventElapsedTime(&mut gpu_ms, self.start, self.end)?; +- // Module function handles use the Driver API attribute query. The +- // plural Runtime API expects a registered host function address. +- let mut lds = 0; +- let mut regs = 0; +- let mut private = 0; +- let mut max_threads = 0; +- hipFuncGetAttribute(&mut lds, hipFunction_attribute::HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES, function)?; +- hipFuncGetAttribute(&mut regs, hipFunction_attribute::HIP_FUNC_ATTRIBUTE_NUM_REGS, function)?; +- hipFuncGetAttribute(&mut private, hipFunction_attribute::HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES, function)?; +- hipFuncGetAttribute(&mut max_threads, hipFunction_attribute::HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, function)?; +- let threads = block.iter().try_fold(1u32, |a, b| a.checked_mul(*b)) +- .filter(|n| *n <= i32::MAX as u32).ok_or(hipErrorCode_t::InvalidValue)?; +- let mut blocks_per_multiprocessor = 0; +- hipModuleOccupancyMaxActiveBlocksPerMultiprocessor(&mut blocks_per_multiprocessor, +- function, threads as i32, shared as usize)?; +- eprintln!("D4R_KERNEL_PROFILE kernel={name:?} backend={backend} phase={phase} gpu_ms={gpu_ms:.6} enqueue_ms={enqueue_ms:.6} completion_ms={completion_ms:.6} grid={},{},{} block={},{},{} static_lds_bytes={} dynamic_lds_bytes={shared} regs={} private_bytes={} max_threads={} predicted_blocks_per_multiprocessor={blocks_per_multiprocessor} serializing=1", +- grid[0], grid[1], grid[2], block[0], block[1], block[2], lds, +- regs, private, max_threads); ++ hipEventElapsedTime(&mut gpu_ms, self.events.start, self.events.end)?; ++ description.print(gpu_ms, enqueue_ms, completion_ms, true, stream.0 as usize); + Ok(()) + } + } +diff --git a/zluda/src/impl/module.rs b/zluda/src/impl/module.rs +index 232f7d4..8d43abf 100644 +--- a/zluda/src/impl/module.rs ++++ b/zluda/src/impl/module.rs +@@ -503,6 +503,7 @@ fn zluda_build_version() -> &'static str { + && !matches!( + key.as_str(), + "D4R_ZLUDA_NATIVE_DIR" | "D4R_ZLUDA_DIR" | "D4R_ZLUDA_DUMP_DIR" | "D4R_ZLUDA_DUMP_PTX_DIR" | "D4R_ZLUDA_LIBCUDA" | "D4R_ZLUDA_VERBOSE" | "D4R_ZLUDA_PROFILE" ++ | "D4R_ZLUDA_PROFILE_API" | "D4R_ZLUDA_PROFILE_DEFERRED" | "D4R_ZLUDA_PROFILE_EVERY" | "D4R_ZLUDA_PROFILE_ALLOW_LEGACY" + | "D4R_ZLUDA_WMMA_FP8_NATIVE" // keyed by effect (ExtraCacheAttributes) + ) + }) +-- +2.55.0.windows.5 + diff --git a/patches/zluda/0017-asynchronous-2d-array-copies.patch b/patches/zluda/0017-asynchronous-2d-array-copies.patch new file mode 100644 index 0000000..4b02d80 --- /dev/null +++ b/patches/zluda/0017-asynchronous-2d-array-copies.patch @@ -0,0 +1,50 @@ +From 67127dde599879ada022d6439383d682e09e0b36 Mon Sep 17 00:00:00 2001 +From: xdfnx-dev +Date: Fri, 2 Oct 2026 11:50:45 +0300 +Subject: [PATCH 17/17] Implement asynchronous 2D array copies for Windows + input batching + +--- + zluda/src/impl/memory.rs | 13 +++++++++++++ + zluda/src/lib.rs | 1 + + 2 files changed, 14 insertions(+) + +diff --git a/zluda/src/impl/memory.rs b/zluda/src/impl/memory.rs +index b724942..42facb3 100644 +--- a/zluda/src/impl/memory.rs ++++ b/zluda/src/impl/memory.rs +@@ -469,6 +469,19 @@ pub(crate) unsafe fn copy_2d_v2(memcpy: hip_Memcpy2D) -> CUresult { + Ok(()) + } + ++pub(crate) unsafe fn copy_2d_async_v2(memcpy: hip_Memcpy2D, stream: hipStream_t) -> CUresult { ++ // The format-80 emulation in copy_2d_v2 uses temporary CPU storage. ++ // Reject it explicitly here: an asynchronous copy cannot outlive those ++ // temporaries, and the Windows fast path uses canonical float arrays. ++ let src_packed = memcpy.srcMemoryType == hipMemoryType::hipMemoryTypeArray ++ && array::original_format(memcpy.srcArray)? == Some(80); ++ let dst_packed = memcpy.dstMemoryType == hipMemoryType::hipMemoryTypeArray ++ && array::original_format(memcpy.dstArray)? == Some(80); ++ if src_packed || dst_packed { return Err(CUerror::NOT_SUPPORTED); } ++ hipMemcpyParam2DAsync(&memcpy, stream)?; ++ Ok(()) ++} ++ + pub(crate) unsafe fn copy_2d_unaligned_v2(memcpy: hip_Memcpy2D) -> CUresult { + hipDrvMemcpy2DUnaligned(&memcpy)?; + Ok(()) +diff --git a/zluda/src/lib.rs b/zluda/src/lib.rs +index 31bfa42..939981e 100644 +--- a/zluda/src/lib.rs ++++ b/zluda/src/lib.rs +@@ -150,6 +150,7 @@ cuda_macros::cuda_function_declarations!( + cuMemAlloc_v2, + cuMemAllocPitch_v2, + cuMemcpy2D_v2, ++ cuMemcpy2DAsync_v2, + cuMemcpy2DUnaligned_v2, + cuMemcpy3D_v2, + cuMemcpyAsync, +-- +2.55.0.windows.5 + diff --git a/scripts/windows/profile-windows-rdna4.ps1 b/scripts/windows/profile-windows-rdna4.ps1 index fc6c4d7..4ed3746 100644 --- a/scripts/windows/profile-windows-rdna4.ps1 +++ b/scripts/windows/profile-windows-rdna4.ps1 @@ -11,6 +11,7 @@ if (!$ZludaRoot) { $ZludaRoot=Join-Path $repo 'dist/zluda-windows-final' } if (!$PackageRoot) { $PackageRoot=Join-Path $repo 'dist/windows-rdna4-command-list' } if (!$OutputDirectory) { $OutputDirectory=Join-Path $repo ('test-results/profile-' + $Preset + '-' + (Get-Date -Format 'yyyyMMdd-HHmmss-fff')) } $settings=@{D4R_ZLUDA_PROFILE='1'; D4R_PROFILE_STAGES='1'; D4R_VALIDATE_OUTPUT='1'; + D4R_ZLUDA_PROFILE_DEFERRED=$null; D4R_ZLUDA_PROFILE_ALLOW_LEGACY=$null; D4R_ZLUDA_PROFILE_EVERY=$null; D4R_ZLUDA_NATIVE_DIR=$(if ($Translated) { $null } elseif ($Preset -eq 11) { Join-Path $repo 'build/native-k-gfx1201' } else { Join-Path $repo 'build/native-m-gfx1201' })} $old=@{} try { diff --git a/scripts/windows/test-async-k.ps1 b/scripts/windows/test-async-k.ps1 index 0ead9a8..3e263a0 100644 --- a/scripts/windows/test-async-k.ps1 +++ b/scripts/windows/test-async-k.ps1 @@ -4,6 +4,9 @@ param( [string]$ZludaRoot, [string]$PackageRoot, [string]$Resolution = '3840x2160', [switch]$ProfileGpuBoundary, + [switch]$ProfileKernelsDeferred, + [switch]$BatchInputCopies, + [ValidateSet('baseline','packed','unorm','depth-stencil')][string]$PixelProfile = 'depth-stencil', [string]$OutputDirectory ) $ErrorActionPreference = 'Stop' @@ -18,8 +21,12 @@ if ($Resolution -notmatch '^([0-9]+)x([0-9]+)$') { throw 'Resolution must be WID $width=$Matches[1]; $height=$Matches[2] New-Item -ItemType Directory -Force $OutputDirectory | Out-Null $settings = @{D4R_DIAG_BURST='1'; D4R_DIAG_RECREATE=$null; D4R_ASYNC_INTEROP=$null; + D4R_BATCH_INPUT_COPIES=$null; D4R_ZLUDA_NATIVE_DIR=$NativeRoot; D4R_QUIET_API='1'; D4R_VALIDATE_OUTPUT='1'; D4R_PROFILE_STAGES='1'; D4R_ZLUDA_PROFILE=$null; D4R_ZLUDA_PROFILE_API=$null; + D4R_ZLUDA_PROFILE_DEFERRED=$(if ($ProfileKernelsDeferred) { '1' } else { $null }); + D4R_ZLUDA_PROFILE_ALLOW_LEGACY=$(if ($ProfileKernelsDeferred) { '1' } else { $null }); + D4R_ZLUDA_PROFILE_EVERY=$(if ($ProfileKernelsDeferred) { '1' } else { $null }); D4R_PROFILE_GPU_BOUNDARY=$(if ($ProfileGpuBoundary) { '1' } else { $null }); ZLUDA_CACHE_DIR=(Join-Path $repo 'build/zluda-cache-windows'); PYTHONPATH=(Join-Path $repo '.tools/python/vendor')} @@ -33,15 +40,25 @@ try { [Environment]::SetEnvironmentVariable('D4R_DIAG_RECREATE', $(if ($mode -eq 'recreate') { '1' } else { $null }), 'Process') foreach ($variant in @('sync','async')) { [Environment]::SetEnvironmentVariable('D4R_ASYNC_INTEROP', $(if ($variant -eq 'async') { '1' } else { $null }), 'Process') + [Environment]::SetEnvironmentVariable('D4R_BATCH_INPUT_COPIES', $(if ($BatchInputCopies -and $variant -eq 'async') { '1' } else { $null }), 'Process') $output=Join-Path $OutputDirectory "$variant-$mode" $arguments=@{RuntimeProfile='therock'; ZludaRoot=$ZludaRoot; PackageRoot=$PackageRoot; NgxCore=$NgxCore; DlssDll=$DlssDll; NgxOnly=$true; NgxMode='d3d12'; NgxAbi='project-legacy'; Preset=11; NgxCreateFlags=11; NgxOutputResolution=$Resolution; RequireNativeNetwork=$true; - CommandListBackend=$true; EarlyIndirectProbe=$true; PixelProfile='depth-stencil'; + CommandListBackend=$true; EarlyIndirectProbe=$true; PixelProfile=$PixelProfile; BarrierMode='inherited-legacy'; Iterations=3; CaptureExceptions=$true; TimeoutSeconds=600; OptiScalerDll=(Join-Path $repo 'dist/optiscaler-windows-d4r/OptiScaler.dll'); OutputDirectory=$output} & (Join-Path $PSScriptRoot 'test-windows-rdna4.ps1') @arguments if ($LASTEXITCODE) { throw "Failed $variant-$mode; logs: $output" } + if ($ProfileKernelsDeferred) { + $stderr=Get-Content -LiteralPath (Join-Path $output 'd3d12-evaluate-preset-11.stderr.log') -Raw + if ($stderr -match '(?m)^D4R_KERNEL_PROFILE .*serializing=1') { throw 'Unexpected serializing kernel profiler.' } + foreach ($layer in @('enc0','enc1','enc2','enc3','enc4','dec5','dec4','dec3','dec2','dec1','dec0')) { + $pattern='(?m)^D4R_KERNEL_PROFILE(?:_INVALID)? kernel="dltss_pwin_' + $layer + '_layer" backend=native phase=main ' + $events=[regex]::Matches($stderr,$pattern) + if ($events.Count -ne 3) { throw "Expected three collected $layer events for $variant-$mode, found $($events.Count)." } + } + } if ($ProfileGpuBoundary -and $variant -eq 'async') { $stdout=Get-Content -LiteralPath (Join-Path $output 'd3d12-evaluate-preset-11.stdout.log') -Raw $timings=[regex]::Matches($stdout, '(?m)^D4R_GPU_BOUNDARY [^\r\n]*diagnostics_cpu_bytes=32 serializing=0\r?$') @@ -53,7 +70,8 @@ try { if ($LASTEXITCODE) { throw "Async $mode output mismatch" } } @{passed=$true; architecture='gfx1201'; preset=11; resolution=$Resolution; - modes=@('burst','recreate'); framesPerMode=3; exactRgb=$true; gpuBoundaryTiming=[bool]$ProfileGpuBoundary} | + modes=@('burst','recreate'); framesPerMode=3; exactRgb=$true; gpuBoundaryTiming=[bool]$ProfileGpuBoundary; + deferredKernelTiming=[bool]$ProfileKernelsDeferred; batchInputs=[bool]$BatchInputCopies; pixelProfile=$PixelProfile} | ConvertTo-Json | Set-Content -LiteralPath (Join-Path $OutputDirectory 'validation.json') -Encoding UTF8 Write-Host "PASS asynchronous K burst/recreate: $OutputDirectory" } finally { diff --git a/scripts/windows/test-deferred-kernel-profile.ps1 b/scripts/windows/test-deferred-kernel-profile.ps1 new file mode 100644 index 0000000..b3a8aa6 --- /dev/null +++ b/scripts/windows/test-deferred-kernel-profile.ps1 @@ -0,0 +1,101 @@ +[CmdletBinding()] +param([string]$ZludaRoot, [string]$PackageRoot, [string]$HipRoot, [string]$OutputDirectory) +$ErrorActionPreference = 'Stop' +$repo = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) +if (!$ZludaRoot) { $ZludaRoot = Join-Path $repo 'dist/zluda-windows-deferred-profile' } +if (!$PackageRoot) { $PackageRoot = Join-Path $repo 'dist/windows-rdna4-command-list' } +if (!$HipRoot) { $HipRoot = Join-Path $repo '.tools/therock-10.2.0a20260929/_rocm_sdk_core' } +if (!$OutputDirectory) { $OutputDirectory = Join-Path $repo ('test-results/deferred-profile-' + (Get-Date -Format 'yyyyMMdd-HHmmss-fff')) } +New-Item -ItemType Directory -Force $OutputDirectory | Out-Null +$settings = @{ + D4R_ZLUDA_PROFILE=$null; D4R_ZLUDA_PROFILE_DEFERRED='1'; D4R_ZLUDA_PROFILE_EVERY='1'; + D4R_ZLUDA_PROFILE_ALLOW_LEGACY=$null; + D4R_DIAG_NONBLOCKING_STREAM='1'; D4R_DIAG_QUEUE_PTX=$null; + D4R_DIAG_DIR=$OutputDirectory; + ZLUDA_CACHE_DIR=(Join-Path $repo 'build/zluda-cache-windows'); + PATH=((Join-Path $HipRoot 'bin') + ';' + $ZludaRoot + ';' + $env:PATH) +} +$old=@{}; $cases=@() +function RunProbe([string[]]$Arguments, [string]$Stdout, [string]$Stderr) { + $info=[Diagnostics.ProcessStartInfo]::new() + $info.FileName=Join-Path $PackageRoot 'bin/d4r_cuda_driver_probe.exe' + $info.Arguments=$Arguments -join ' ' + $info.UseShellExecute=$false; $info.CreateNoWindow=$true + $info.RedirectStandardOutput=$true; $info.RedirectStandardError=$true + $p=[Diagnostics.Process]::new(); $p.StartInfo=$info + $outFile=$null; $errFile=$null; $outTask=$null; $errTask=$null + try { + if (!$p.Start()) { throw 'Could not launch the CUDA diagnostic.' } + $outFile=[IO.File]::Open($Stdout,[IO.FileMode]::Create,[IO.FileAccess]::Write,[IO.FileShare]::Read) + $errFile=[IO.File]::Open($Stderr,[IO.FileMode]::Create,[IO.FileAccess]::Write,[IO.FileShare]::Read) + $outTask=$p.StandardOutput.BaseStream.CopyToAsync($outFile) + $errTask=$p.StandardError.BaseStream.CopyToAsync($errFile) + $timer=[Diagnostics.Stopwatch]::StartNew() + while (!$p.WaitForExit(1000)) { + if ($timer.Elapsed.TotalSeconds -ge 120) { throw "CUDA diagnostic timed out; logs: $OutputDirectory" } + } + $p.WaitForExit() + [void]$outTask.GetAwaiter().GetResult(); [void]$errTask.GetAwaiter().GetResult() + return $p.ExitCode + } finally { + if (!$p.HasExited) { $p.Kill(); $p.WaitForExit() } + if ($outTask) { try { [void]$outTask.GetAwaiter().GetResult() } catch {} } + if ($errTask) { try { [void]$errTask.GetAwaiter().GetResult() } catch {} } + if ($outFile) { $outFile.Dispose() }; if ($errFile) { $errFile.Dispose() } + $p.Dispose() + } +} +try { + foreach ($key in $settings.Keys) { + $old[$key]=[Environment]::GetEnvironmentVariable($key,'Process') + [Environment]::SetEnvironmentVariable($key,$settings[$key],'Process') + } + foreach ($context in @('primary','created')) { + foreach ($queued in @($false,$true)) { + $name="$context-$(if ($queued) { 'queued' } else { 'single' })" + [Environment]::SetEnvironmentVariable('D4R_DIAG_QUEUE_PTX', $(if ($queued) { '1' } else { $null }), 'Process') + $iterations=129 + $stdout=Join-Path $OutputDirectory "$name.stdout.log" + $stderr=Join-Path $OutputDirectory "$name.stderr.log" + # Native Windows argument parsing: quote absolute paths explicitly. + $arguments=@('--hip-root', ('"' + $HipRoot + '"'), '--cuda-dll', ('"' + (Join-Path $ZludaRoot 'nvcuda.dll') + '"'), + '--context', $context, '--iterations', "$iterations") + $exitCode=RunProbe $arguments $stdout $stderr + if ($exitCode) { throw "$name failed ($exitCode); logs: $OutputDirectory" } + $out=Get-Content -LiteralPath $stdout -Raw + $err=Get-Content -LiteralPath $stderr -Raw + if ($out -notmatch 'PASS CUDA architecture=gfx1201' -or $out -notmatch 'guard_verified=1') { throw "$name did not verify PTX output" } + $records=[regex]::Matches($err, '(?m)^D4R_KERNEL_PROFILE kernel="d4r_ptx_pattern" [^\r\n]*serializing=0 deferred=1 [^\r\n]+') + $invalid=[regex]::Matches($err, '(?m)^D4R_KERNEL_PROFILE_INVALID kernel="d4r_ptx_pattern" [^\r\n]*reason=invalid_elapsed_time serializing=0 deferred=1 [^\r\n]+') + if ($records.Count + $invalid.Count -ne $iterations -or $err -match 'serializing=1|D4R_KERNEL_PROFILE_SKIPPED') { + throw "$name expected $iterations collected deferred records, found $($records.Count) valid / $($invalid.Count) invalid" + } + foreach ($record in $records) { + if ($record.Value -notmatch 'gpu_ms=([0-9.]+)') { + throw "$name contains an invalid GPU time" + } + $gpuMs=[double]::Parse($Matches[1], [Globalization.CultureInfo]::InvariantCulture) + if ([double]::IsNaN($gpuMs) -or [double]::IsInfinity($gpuMs)) { throw "$name contains an invalid GPU time" } + } + $cases += @{context=$context; queued=$queued; launches=$iterations; collected=($records.Count+$invalid.Count); + validTimings=$records.Count; invalidTimings=$invalid.Count; exitCode=$exitCode} + Write-Host "PASS ${name}: $($records.Count) valid / $($invalid.Count) invalid deferred timings; exact PTX output and guards" + } + } + # Legacy default-stream timing markers would add stream dependencies. + [Environment]::SetEnvironmentVariable('D4R_DIAG_NONBLOCKING_STREAM',$null,'Process') + [Environment]::SetEnvironmentVariable('D4R_DIAG_QUEUE_PTX',$null,'Process') + $stdout=Join-Path $OutputDirectory 'default-stream.stdout.log'; $stderr=Join-Path $OutputDirectory 'default-stream.stderr.log' + $arguments=@('--hip-root', ('"' + $HipRoot + '"'), '--cuda-dll', ('"' + (Join-Path $ZludaRoot 'nvcuda.dll') + '"'), '--iterations', '3') + $exitCode=RunProbe $arguments $stdout $stderr + $out=Get-Content -LiteralPath $stdout -Raw; $err=Get-Content -LiteralPath $stderr -Raw + if ($exitCode -or $out -notmatch 'PASS CUDA architecture=gfx1201' -or + $err -notmatch 'reason=legacy_default_stream' -or $err -match '(?m)^D4R_KERNEL_PROFILE ') { + throw 'Default-stream diagnostics did not safely skip timing events.' + } + @{passed=$true; architecture='gfx1201'; cases=$cases; defaultStreamSkipped=$true; addedCompletionWaits=0} | + ConvertTo-Json -Depth 5 | Set-Content -LiteralPath (Join-Path $OutputDirectory 'validation.json') -Encoding UTF8 + Write-Host "PASS deferred kernel profiling: $OutputDirectory" +} finally { + foreach ($key in $old.Keys) { [Environment]::SetEnvironmentVariable($key,$old[$key],'Process') } +} diff --git a/scripts/windows/windows-game.ps1 b/scripts/windows/windows-game.ps1 index f585bc0..7214f72 100644 --- a/scripts/windows/windows-game.ps1 +++ b/scripts/windows/windows-game.ps1 @@ -11,11 +11,15 @@ param( [switch]$ValidateOutput, [switch]$ProfileStages, [switch]$ProfileKernels, + [switch]$ProfileKernelsDeferred, + [switch]$ProfileLegacyStream, + [ValidateRange(1,1000000)][int]$KernelProfileEvery = 17, [switch]$ProfileCommandHooks, [switch]$ProfileCudaApi, [switch]$ProfileGpuBoundary, [switch]$UncachedInteropLists, [switch]$AsyncInterop, + [switch]$BatchInputCopies, [switch]$VerboseRuntime, [switch]$CaptureExceptions, [string[]]$GameArguments = @('-dx12'), @@ -23,6 +27,8 @@ param( ) $ErrorActionPreference = 'Stop' if ($ProfileGpuBoundary -and !$AsyncInterop) { throw '-ProfileGpuBoundary requires -AsyncInterop.' } +if ($ProfileKernels -and $ProfileKernelsDeferred) { throw 'Choose either serializing or deferred kernel profiling.' } +if ($ProfileLegacyStream -and !$ProfileKernelsDeferred) { throw '-ProfileLegacyStream requires -ProfileKernelsDeferred.' } $package = [IO.Path]::GetFullPath($PSScriptRoot) $GameExe = (Get-Item -LiteralPath $GameExe -ErrorAction Stop).FullName $game = Split-Path $GameExe @@ -182,11 +188,15 @@ $settings = @{ D4R_VALIDATE_OUTPUT=$(if ($ValidateOutput) { '1' } else { $null }); D4R_PROFILE_STAGES=$(if ($ProfileStages) { '1' } else { $null }); D4R_ZLUDA_PROFILE=$(if ($ProfileKernels) { '1' } else { $null }); + D4R_ZLUDA_PROFILE_DEFERRED=$(if ($ProfileKernelsDeferred) { '1' } else { $null }); + D4R_ZLUDA_PROFILE_ALLOW_LEGACY=$(if ($ProfileLegacyStream) { '1' } else { $null }); + D4R_ZLUDA_PROFILE_EVERY=$(if ($ProfileKernelsDeferred) { "$KernelProfileEvery" } else { $null }); D4R_PROFILE_COMMAND_HOOKS=$(if ($ProfileCommandHooks) { '1' } else { $null }); D4R_ZLUDA_PROFILE_API=$(if ($ProfileCudaApi) { '1' } else { $null }); D4R_PROFILE_GPU_BOUNDARY=$(if ($ProfileGpuBoundary) { '1' } else { $null }); D4R_DISABLE_INTEROP_LIST_CACHE=$(if ($UncachedInteropLists) { '1' } else { $null }); D4R_ASYNC_INTEROP=$(if ($AsyncInterop) { '1' } else { $null }); + D4R_BATCH_INPUT_COPIES=$(if ($BatchInputCopies) { '1' } else { $null }); D4R_DIAG_DIR=$OutputDirectory; ZLUDA_LOG_DIR=(Join-Path $OutputDirectory 'zluda-trace'); ZLUDA_CACHE_DIR=$CacheDirectory; PATH=((GamePath 'd4r/hip/bin') + ';' + (GamePath 'd4r/zluda') + ';' + $env:PATH) } @@ -222,7 +232,12 @@ try { Write-Host "Temporary diagnostic resolution: $DiagnosticResolution; original settings restore at exit." } foreach ($key in $settings.Keys) { $old[$key] = [Environment]::GetEnvironmentVariable($key,'Process'); [Environment]::SetEnvironmentVariable($key,$settings[$key],'Process') } - @{preset=$Preset; game=$GameExe; package=$metadata; attachedDebugger=[bool]$CaptureExceptions; driver=@(Get-CimInstance Win32_VideoController | Select-Object Name,DriverVersion,DriverDate); arguments=$GameArguments} | + @{preset=$Preset; game=$GameExe; package=$metadata; attachedDebugger=[bool]$CaptureExceptions; + diagnostics=@{asyncInterop=[bool]$AsyncInterop; batchInputCopies=[bool]$BatchInputCopies; validateOutput=[bool]$ValidateOutput; + profileGpuBoundary=[bool]$ProfileGpuBoundary; profileStages=[bool]$ProfileStages; profileCudaApi=[bool]$ProfileCudaApi; + profileKernels=[bool]$ProfileKernels; profileKernelsDeferred=[bool]$ProfileKernelsDeferred; + profileLegacyStream=[bool]$ProfileLegacyStream; kernelProfileEvery=$KernelProfileEvery}; + driver=@(Get-CimInstance Win32_VideoController | Select-Object Name,DriverVersion,DriverDate); arguments=$GameArguments} | ConvertTo-Json -Depth 12 | Set-Content -LiteralPath (Join-Path $OutputDirectory 'environment.json') -Encoding UTF8 function Quote([string]$value) { if ($value.Contains('"')) { throw 'A command argument contains an unsupported quote.' }; return '"' + $value + '"' } $arguments = @('--output-directory', (Quote $OutputDirectory)) diff --git a/tools/windows/cuda_driver_probe.cpp b/tools/windows/cuda_driver_probe.cpp index 2594c3d..7153975 100644 --- a/tools/windows/cuda_driver_probe.cpp +++ b/tools/windows/cuda_driver_probe.cpp @@ -80,6 +80,21 @@ int main(int argc, char** argv) if (primary) api.cuDevicePrimaryCtxRelease_v2(device); else api.cuCtxDestroy_v2(context); } } } cleanup{api, context, device, primary}; api.check(api.cuCtxSetCurrent(context), "cuCtxSetCurrent"); + void* stream = nullptr; + using CreateStream = int(WINAPI*)(void**, unsigned); + using DestroyStream = int(WINAPI*)(void*); + DestroyStream destroyStream = nullptr; + if (std::getenv("D4R_DIAG_NONBLOCKING_STREAM")) { + const auto createStream = api.library.symbol("cuStreamCreate"); + destroyStream = api.library.symbol("cuStreamDestroy_v2"); + api.check(createStream(&stream, 1), "cuStreamCreate(NON_BLOCKING)"); + } + struct StreamCleanup { + void*& stream; DestroyStream destroy; + ~StreamCleanup() { if (stream && destroy) destroy(stream); } + } streamCleanup{stream, destroyStream}; + const bool queued = std::getenv("D4R_DIAG_QUEUE_PTX") != nullptr; + if (queued && !stream) throw std::runtime_error("Queued PTX diagnostic requires D4R_DIAG_NONBLOCKING_STREAM=1"); CUmodule module = nullptr; std::printf("STAGE cuModuleLoadData (PTX JIT)\n"); api.check(api.cuModuleLoadData(&module, ptx), "cuModuleLoadData"); @@ -93,25 +108,30 @@ int main(int argc, char** argv) CUdeviceptr pointer = 0; api.check(api.cuMemAlloc_v2(&pointer, bytes), "cuMemAlloc_v2"); struct MemCleanup { CudaApi& api; CUdeviceptr p; ~MemCleanup() { if (p) api.cuMemFree_v2(p); } } ac{api, pointer}; + if (queued) api.check(api.cuMemcpyHtoD_v2(pointer, result.data(), bytes), "cuMemcpyHtoD_v2(queued sentinel)"); for (unsigned i = 0; i < args.iterations; ++i) { - uint32_t n = nmax - i % 129u, seed = 0x12345678u + i * 31u; - std::fill(result.begin(), result.end(), 0xdeadbeefu); - api.check(api.cuMemcpyHtoD_v2(pointer, result.data(), bytes), "cuMemcpyHtoD_v2(sentinel)"); + uint32_t n = queued ? nmax : nmax - i % 129u, seed = 0x12345678u + i * 31u; + if (!queued) { + std::fill(result.begin(), result.end(), 0xdeadbeefu); + api.check(api.cuMemcpyHtoD_v2(pointer, result.data(), bytes), "cuMemcpyHtoD_v2(sentinel)"); + } void* parameters[] = {&pointer, &n, &seed}; api.check(api.cuLaunchKernel(function, (n + 127u) / 128u, 1, 1, 128, 1, 1, - 0, nullptr, parameters, nullptr), "cuLaunchKernel"); + 0, stream, parameters, nullptr), "cuLaunchKernel"); + if (queued && i + 1 != args.iterations) continue; api.check(api.cuCtxSynchronize(), "cuCtxSynchronize"); api.check(api.cuMemcpyDtoH_v2(result.data(), pointer, bytes), "cuMemcpyDtoH_v2"); if (!verify(result, n, seed)) return 5; } api.check(api.cuMemFree_v2(pointer), "cuMemFree_v2"); ac.p = 0; api.check(api.cuModuleUnload(module), "cuModuleUnload"); mc.m = nullptr; + if (stream) { api.check(destroyStream(stream), "cuStreamDestroy_v2"); stream = nullptr; } api.check(api.cuCtxSetCurrent(nullptr), "cuCtxSetCurrent(NULL)"); if (primary) api.check(api.cuDevicePrimaryCtxRelease_v2(device), "cuDevicePrimaryCtxRelease_v2"); else api.check(api.cuCtxDestroy_v2(context), "cuCtxDestroy_v2"); cleanup.context = nullptr; - std::printf("PASS CUDA architecture=gfx1201 context=%s iterations=%u guard_verified=1\n", - args.context.c_str(), args.iterations); + std::printf("PASS CUDA architecture=gfx1201 context=%s iterations=%u guard_verified=1 queued=%u\n", + args.context.c_str(), args.iterations, unsigned(queued)); return 0; } catch (const std::exception& e) { std::fprintf(stderr, "FAIL CUDA %s\n", e.what()); diff --git a/tools/windows/cuda_image_api.h b/tools/windows/cuda_image_api.h index 16811bd..254ed12 100644 --- a/tools/windows/cuda_image_api.h +++ b/tools/windows/cuda_image_api.h @@ -47,6 +47,7 @@ struct ImageApi { D4R_IMAGE(cuArray3DGetDescriptor_v2, Array3DDescriptor*, Array); D4R_IMAGE(cuArrayDestroy, Array); D4R_IMAGE(cuMemcpy2D_v2, const Copy2D*); + D4R_IMAGE(cuMemcpy2DAsync_v2, const Copy2D*, void*); D4R_IMAGE(cuTexObjectCreate, Texture*, const ResourceDescriptor*, const TextureDescriptor*, const void*); D4R_IMAGE(cuTexObjectDestroy, Texture); D4R_IMAGE(cuTexObjectGetResourceDesc, ResourceDescriptor*, Texture); @@ -131,13 +132,18 @@ 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) { + void upload_device(CUdeviceptr source, size_t pitch, bool deferred = false) { 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)"); + if (deferred) { + // The Windows ZLUDA backend and HIP conversion use the same + // legacy default stream. Caller must complete the whole input + // batch before NGX may consume arrays on its own streams. + api_.cuda.check(api_.cuMemcpy2DAsync_v2(©, nullptr), "cuMemcpy2DAsync_v2(external VRAM to array)"); + } else 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"); diff --git a/tools/windows/cuda_image_probe.cpp b/tools/windows/cuda_image_probe.cpp index de524fa..f2ea4ce 100644 --- a/tools/windows/cuda_image_probe.cpp +++ b/tools/windows/cuda_image_probe.cpp @@ -30,22 +30,25 @@ int main(int argc, char** argv) { struct Cleanup { CudaApi& api; CUcontext context; ~Cleanup() { (void)api.cuCtxDestroy_v2(context); } } cleanup{cuda, context}; d4r::cuda::ImageApi images(cuda); for (unsigned iteration = 0; iteration < args.iterations; ++iteration) { - for (unsigned channels : {1u, 2u, 4u}) for (bool surface : {false, true}) { - d4r::cuda::Image image(images, 17, 9, 16, channels, surface); + for (unsigned format : {16u, 32u}) for (unsigned channels : {1u, 2u, 4u}) for (bool surface : {false, true}) { + const unsigned halvesPerPixel = channels * (format == 16 ? 1 : 2); + d4r::cuda::Image image(images, 17, 9, format, channels, surface); image.verify_wrong_object_kind(); - std::vector input(17 * 9 * channels), output(input.size()); + std::vector input(17 * 9 * halvesPerPixel), output(input.size()); for (size_t i = 0; i < input.size(); ++i) input[i] = uint16_t((i + iteration) % 8 * 0x100); 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); + const size_t rowBytes = 17 * halvesPerPixel * 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); + for (unsigned y = 0; y < 9; ++y) std::memcpy(padded.data() + y * pitch, input.data() + y * 17 * halvesPerPixel, rowBytes); cuda.check(cuda.cuMemcpyHtoD_v2(devicePixels, padded.data(), padded.size()), "cuMemcpyHtoD_v2(pitched pattern)"); - image.upload_device(devicePixels, pitch); + const bool deferred = std::getenv("D4R_DIAG_ASYNC_ARRAY_COPY") != nullptr; + image.upload_device(devicePixels, pitch, deferred); + if (deferred) cuda.check(cuda.cuCtxSynchronize(), "cuCtxSynchronize(input batch regression)"); 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); @@ -53,7 +56,7 @@ int main(int argc, char** argv) { 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)) + if (std::memcmp(padded.data() + y * pitch, input.data() + y * 17 * halvesPerPixel, 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"); @@ -61,7 +64,8 @@ int main(int argc, char** argv) { } } cuda.check(cuda.cuCtxSynchronize(), "cuCtxSynchronize"); - std::printf("PASS CUDA_IMAGES architecture=gfx1201 iterations=%u descriptor_and_storage=1 pitched_device_copies=1\n", args.iterations); + std::printf("PASS CUDA_IMAGES architecture=gfx1201 iterations=%u descriptor_and_storage=1 pitched_device_copies=1 fp16_fp32=1 async_array_copy=%u\n", + args.iterations, unsigned(std::getenv("D4R_DIAG_ASYNC_ARRAY_COPY") != nullptr)); return 0; } catch (const std::exception& e) { std::fprintf(stderr, "FAIL CUDA_IMAGES %s\n", e.what()); diff --git a/tools/windows/ngx_windows_runtime.h b/tools/windows/ngx_windows_runtime.h index 266ac24..22463f9 100644 --- a/tools/windows/ngx_windows_runtime.h +++ b/tools/windows/ngx_windows_runtime.h @@ -112,6 +112,8 @@ struct Runtime { cuda.check(cuda.cuDevicePrimaryCtxRetain(&context, ordinal), "cuDevicePrimaryCtxRetain"); current(); timeline = std::make_unique(external, d3d); + std::printf("D4R_INTEROP async=%u batch_inputs=%u cpu_image_copies=0\n", + unsigned(std::getenv("D4R_ASYNC_INTEROP") != nullptr), unsigned(std::getenv("D4R_BATCH_INPUT_COPIES") != nullptr)); // 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"))); @@ -386,6 +388,10 @@ private: report_completed_copies(); } if (selected) { + // The fence can advance between report_completed_copies() and + // slot selection. Collect that now-complete ticket before Reset + // and before submitted() overwrites its timing identity. + if (selected->timing) selected->timing->report(queue_.Get(), this); dx(selected->input_allocator->Reset(), "Async input allocator Reset"); dx(selected->output_allocator->Reset(), "Async output allocator Reset"); dx(selected->input->Reset(selected->input_allocator.Get(), nullptr), "Async input list Reset"); @@ -580,16 +586,23 @@ private: if (ticket) rt_->timeline->wait_ready(ticket->input); else rt_->timeline->wait_input(queue); rt_->cuda.check(rt_->cuda.cuCtxSynchronize(), "Interop input producer completion"); mark("input_fence_wait"); + const bool batchInputs = std::getenv("D4R_BATCH_INPUT_COPIES") != nullptr; for (unsigned i : {0u, 1u, 2u, 4u}) if (planes_[i].texture) { auto& plane = planes_[i]; void* data = plane.shared->mapped; uint64_t pitch = plane.shared->footprint.Footprint.RowPitch; if (plane.canonical) { - rt_->pixels->convert(false, data, pitch, plane.canonical->data, plane.canonical->pitch, + if (batchInputs) rt_->pixels->convert_deferred(false, data, pitch, plane.canonical->data, plane.canonical->pitch, + unsigned(plane.desc.Width), plane.desc.Height, plane.spec.storage, i); + else rt_->pixels->convert(false, data, pitch, plane.canonical->data, plane.canonical->pitch, unsigned(plane.desc.Width), plane.desc.Height, plane.spec.storage, i); data = plane.canonical->data; pitch = plane.canonical->pitch; } - plane.image->upload_device(reinterpret_cast(data), pitch); + plane.image->upload_device(reinterpret_cast(data), pitch, batchInputs); } + // HIP conversion and CUDA array copies are ordered on the shared + // legacy default stream. Retain an all-stream completion before NGX, + // including any internal nonblocking streams it might choose to use. + if (batchInputs) rt_->cuda.check(rt_->cuda.cuCtxSynchronize(), "Interop input batch completion"); mark("input_conversion_array_upload"); if (std::getenv("D4R_INTEROP_VERIFY")) { for (unsigned i : {0u, 1u, 2u, 4u}) if (planes_[i].texture) { diff --git a/tools/windows/pixel_conversion.h b/tools/windows/pixel_conversion.h index 3cc9e0d..b32ada4 100644 --- a/tools/windows/pixel_conversion.h +++ b/tools/windows/pixel_conversion.h @@ -77,15 +77,19 @@ public: hip_.check(hip_.hipMemcpy(result.data(), counters_, sizeof(result), hipMemcpyDeviceToHost), "Read validation counters"); return result; } - void convert(bool encode, const void* source, uint64_t sourcePitch, void* dest, uint64_t destPitch, + void convert_deferred(bool encode, const void* source, uint64_t sourcePitch, void* dest, uint64_t destPitch, unsigned width, unsigned height, pixel::Storage format, unsigned plane) { if (!width || !height || uint64_t(width) * height > UINT32_MAX) throw std::runtime_error("Invalid GPU conversion dimensions"); unsigned storage = unsigned(format); void* args[] = {&source, &sourcePitch, &dest, &destPitch, &width, &height, &storage, &plane}; hip_.check(hip_.hipModuleLaunchKernel(encode ? encode_ : decode_, unsigned((uint64_t(width) * height + 255) / 256), 1, 1, 256, 1, 1, 0, nullptr, args, nullptr), encode ? "GPU pixel encode" : "GPU pixel decode"); - // CUDA may use other nonblocking streams. Establish completion before - // giving the converted allocation to that API; images remain in VRAM. + } + void convert(bool encode, const void* source, uint64_t sourcePitch, void* dest, uint64_t destPitch, + unsigned width, unsigned height, pixel::Storage format, unsigned plane) { + convert_deferred(encode, source, sourcePitch, dest, destPitch, width, height, format, plane); + // CUDA may use other nonblocking streams. Standalone conversions must + // establish completion before handing the allocation to that API. hip_.check(hip_.hipDeviceSynchronize(), "GPU conversion completion"); } }; diff --git a/tools/windows/profile_report.py b/tools/windows/profile_report.py index 602de91..36e8acb 100644 --- a/tools/windows/profile_report.py +++ b/tools/windows/profile_report.py @@ -1,4 +1,4 @@ -"""Summarize native Windows d4r CPU stages and optional serializing HIP events. +"""Summarize native Windows d4r CPU stages and optional HIP event diagnostics. CPU completion time is not GPU execution time. HIP occupancy API predictions are not hardware counters. No bandwidth / WMMA utilization is inferred here. @@ -83,6 +83,8 @@ def report(directory, metadata_directory=None): kernels = collections.defaultdict(list) hook_threads = collections.defaultdict(list) apis = collections.defaultdict(list) + sample_skips = [] + invalid_samples = [] boundaries = collections.defaultdict(list) replay = collections.defaultdict(lambda: collections.defaultdict(dict)) for path in sorted(directory.glob('*.log')): @@ -115,13 +117,25 @@ def report(directory, metadata_directory=None): 'access_sum_ms', 'capture_sum_ms', 'driver_sum_ms')): hook_threads[fields['thread']].append(fields) if not line.startswith('D4R_KERNEL_PROFILE '): + if line.startswith('D4R_KERNEL_PROFILE_SKIPPED '): + sample_skips.append(dict(re.findall(r'(\w+)=(\S+)', line))) + if line.startswith('D4R_KERNEL_PROFILE_INVALID '): + invalid_samples.append(dict(re.findall(r'(\w+)=("[^"]+"|\S+)', line))) continue fields = dict(re.findall(r'(\w+)=("[^"]+"|\S+)', line)) - kernels[(fields['kernel'].strip('"'), fields['backend'], fields['phase'])].append(fields) + required = ('kernel', 'backend', 'phase', 'gpu_ms', 'enqueue_ms', 'completion_ms', + 'grid', 'block', 'static_lds_bytes', 'dynamic_lds_bytes', 'regs', 'private_bytes', + 'max_threads', 'predicted_blocks_per_multiprocessor', 'serializing') + if not all(key in fields for key in required): + continue # A bounded process stop can leave a partial final line. + kernels[(fields['kernel'].strip('"'), fields['backend'], fields['phase'], + fields['serializing'])].append(fields) result = dict(cpuStages={name: stats(values) for name, values in stages.items()}, commandStages={name: stats(values) for name, values in commands.items()}, commandHookSamples=[], cudaApi=[], + kernelSampleSkips=sample_skips, + invalidKernelSamples=invalid_samples, gpuBoundaryStages={}, presentMon=presentmon(directory), replayPairs=[], @@ -134,7 +148,8 @@ def report(directory, metadata_directory=None): 'Command record intervals are between NGX recording calls on one thread, not Present/FPS.', 'Hook estimates use random one-in-64 samples; driver timing covers generated forwarding methods only.', 'Hook access includes lock waits; summing threads does not measure serial frame latency or CPU execution time.', - 'HIP-event profiling synchronizes every sampled launch and changes scheduling.', + 'Serializing HIP-event profiles wait after every launch. Deferred profiles query ready events and never add completion waits; timing events still add overhead.', + 'Deferred completionMs is host collection lag, not kernel completion latency. Legacy-stream sampling requires an opt-in because its timing markers retain legacy cross-stream ordering.', 'Occupancy is an API prediction, not a measured hardware counter.', 'Bandwidth and WMMA utilization require additional supported hardware tooling.']) boundary_stages = collections.defaultdict(list) @@ -178,8 +193,11 @@ def report(directory, metadata_directory=None): accessMeanUs=1000 * sums['access_sum_ms'] / sums['samples'], captureMeanUsPerSample=1000 * sums['capture_sum_ms'] / sums['samples'], estimatedAccessCaptureMs=estimate, estimatedAccessCaptureMsPerSecond=1000 * estimate / sums['period_ms'])) - for (name, backend, phase), rows in kernels.items(): + for (name, backend, phase, serializing), rows in kernels.items(): item = dict(kernel=name, backend=backend, phase=phase, + serializing=serializing == '1', + legacyStream=any(row.get('legacy_stream') == '1' for row in rows), + completionMeaning='host_completion_wait' if serializing == '1' else 'host_collection_lag', gpuMs=stats([float(row['gpu_ms']) for row in rows]), enqueueMs=stats([float(row['enqueue_ms']) for row in rows]), completionMs=stats([float(row['completion_ms']) for row in rows]), @@ -224,7 +242,7 @@ def main(): f"candidate_median={row['candidateGpuMs']['median']:.6f} ms ratio_median={row['candidateToControl']['median']:.6f}") for row in result['kernels']: times = row['gpuMs'] - print(f"GPU {row['kernel']} {row['backend']} {row['phase']}: n={times['samples']} mean={times['mean']:.6f} ms total={times['total']:.3f} ms") + print(f"GPU {row['kernel']} {row['backend']} {row['phase']} serializing={int(row['serializing'])}: n={times['samples']} mean={times['mean']:.6f} ms total={times['total']:.3f} ms") if __name__ == '__main__':