Batch Windows DLSS GPU inputs and add deferred kernel diagnostics

2026-10-02 12:06:15 +03:00
1 parent 4f3b3a9fe6
commit 973eb5ecd2
@@ -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")
@@ -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
@@ -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)
+Upstream integration: [PR #11](https://github.com/countervolts/d4r/pull/11)
-targets `countervolts/d4r:windows` from `xdfnx-dev/d4r:upstream-windows`.
+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,
+the fork's `main` keeps the dedicated English Windows README. The maintainer
-mergeable and allows maintainer edits. K performance remains a separate open gate.
+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
+**Current milestone: K performance remains open.** Async submission improves
-and about 53% GPU utilization in 4K Silent Hill 2, compared with roughly 80+
+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`.
@@ -0,0 +1,282 @@
 From 26a43f7047ab933a4afb5ebdd26d5ba9e51e58a3 Mon Sep 17 00:00:00 2001
 From: xdfnx-dev <[email protected]>
 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<Mode> = 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::<u64>().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<bool> = 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<Self, hipErrorCode_t> {
 +        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<Self, hipErrorCode_t> {
 +        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<Pending>, free: Vec<(i32, Events)>,
 +}
 +thread_local! {
 +    static DEFERRED: RefCell<Deferred> = 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<Option<Self>, 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
@@ -0,0 +1,50 @@
 From 67127dde599879ada022d6439383d682e09e0b36 Mon Sep 17 00:00:00 2001
 From: xdfnx-dev <[email protected]>
 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
@@ -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 {
@@ -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 {
@@ -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') }
 }
@@ -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))
@@ -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<CreateStream>("cuStreamCreate");
            destroyStream = api.library.symbol<DestroyStream>("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;
+            uint32_t n = queued ? nmax : nmax - i % 129u, seed = 0x12345678u + i * 31u;
-            std::fill(result.begin(), result.end(), 0xdeadbeefu);
+            if (!queued) {
-            api.check(api.cuMemcpyHtoD_v2(pointer, result.data(), bytes), "cuMemcpyHtoD_v2(sentinel)");
+                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",
+        std::printf("PASS CUDA architecture=gfx1201 context=%s iterations=%u guard_verified=1 queued=%u\n",
-            args.context.c_str(), args.iterations);
+            args.context.c_str(), args.iterations, unsigned(queued));
        return 0;
    } catch (const std::exception& e) {
        std::fprintf(stderr, "FAIL CUDA %s\n", e.what());
@@ -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(&copy), "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(&copy), "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(&copy, nullptr), "cuMemcpy2DAsync_v2(external VRAM to array)");
        } else api_.cuda.check(api_.cuMemcpy2D_v2(&copy), "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");
@@ -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}) {
+            for (unsigned format : {16u, 32u}) for (unsigned channels : {1u, 2u, 4u}) for (bool surface : {false, true}) {
-                d4r::cuda::Image image(images, 17, 9, 16, channels, surface);
+                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<uint16_t> input(17 * 9 * channels), output(input.size());
+                std::vector<uint16_t> 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<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);
+                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());
@@ -112,6 +112,8 @@ struct Runtime {
            cuda.check(cuda.cuDevicePrimaryCtxRetain(&context, ordinal), "cuDevicePrimaryCtxRetain");
            current();
            timeline = std::make_unique<SharedTimeline>(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::Library>(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<uintptr_t>(data), pitch);
+            plane.image->upload_device(reinterpret_cast<uintptr_t>(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) {
@@ -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");
    }
 };
@@ -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__':