From 28832eef6c4f58f1476e80b13e210b1eeef37ada Mon Sep 17 00:00:00 2001 From: xdfnx-dev Date: Fri, 2 Oct 2026 15:01:25 +0300 Subject: [PATCH] Guard paired GPU replay timings and normalize Windows diagnostic environments --- scripts/windows/benchmark-native-replay.ps1 | 17 ++++ scripts/windows/test-async-k.ps1 | 3 + scripts/windows/test-windows-rdna4.ps1 | 17 +++- scripts/windows/windows-game.ps1 | 10 ++ tools/windows/benchmark_native_replay.py | 104 ++++++++++++++++++++ tools/windows/d3d12_dlss_probe.cpp | 3 +- tools/windows/diagnostic.h | 28 ++++++ tools/windows/native_replay_probe.cpp | 66 +++++++++++-- tools/windows/profile_report.py | 37 +++++-- 9 files changed, 265 insertions(+), 20 deletions(-) create mode 100644 scripts/windows/benchmark-native-replay.ps1 create mode 100644 tools/windows/benchmark_native_replay.py diff --git a/scripts/windows/benchmark-native-replay.ps1 b/scripts/windows/benchmark-native-replay.ps1 new file mode 100644 index 0000000..148ff81 --- /dev/null +++ b/scripts/windows/benchmark-native-replay.ps1 @@ -0,0 +1,17 @@ +[CmdletBinding()] +param([Parameter(Mandatory=$true)][string]$Module, + [Parameter(Mandatory=$true)][string]$Candidate, + [Parameter(Mandatory=$true)][string]$FixtureDirectory, + [string]$HipRoot, [string]$Probe, [string]$OutputDirectory, + [ValidateRange(1,10000)][int]$Pairs=32, + [ValidateRange(1,256)][int[]]$Batches=@(1,64), + [ValidateSet('events','dispatch')][string[]]$Timings=@('events','dispatch')) +$ErrorActionPreference='Stop' +$repo=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) +if (!$HipRoot) { $HipRoot=Join-Path $repo '.tools/therock-10.2.0a20260929/_rocm_sdk_core' } +if (!$Probe) { $Probe=Join-Path $repo 'build/windows-rdna4-therock/d4r_native_replay_probe.exe' } +if (!$OutputDirectory) { $OutputDirectory=Join-Path $repo ('test-results/native-benchmark-'+(Get-Date -Format 'yyyyMMdd-HHmmss-fff')) } +& python.exe (Join-Path $repo 'tools/windows/benchmark_native_replay.py') --probe $Probe --hip-root $HipRoot ` + --module $Module --candidate $Candidate --fixture-dir $FixtureDirectory --output-dir $OutputDirectory ` + --pairs $Pairs --batches @Batches --timings @Timings +if ($LASTEXITCODE) { throw "Native benchmark failed. Logs: $OutputDirectory" } diff --git a/scripts/windows/test-async-k.ps1 b/scripts/windows/test-async-k.ps1 index 3e263a0..0ff58c3 100644 --- a/scripts/windows/test-async-k.ps1 +++ b/scripts/windows/test-async-k.ps1 @@ -3,6 +3,7 @@ param( [string]$NgxCore, [string]$DlssDll, [string]$NativeRoot, [string]$ZludaRoot, [string]$PackageRoot, [string]$Resolution = '3840x2160', + [string]$InputResolution, [switch]$ProfileGpuBoundary, [switch]$ProfileKernelsDeferred, [switch]$BatchInputCopies, @@ -45,6 +46,7 @@ try { $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; + NgxInputResolution=$InputResolution; 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} @@ -70,6 +72,7 @@ try { if ($LASTEXITCODE) { throw "Async $mode output mismatch" } } @{passed=$true; architecture='gfx1201'; preset=11; resolution=$Resolution; + inputResolution=$InputResolution; 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 diff --git a/scripts/windows/test-windows-rdna4.ps1 b/scripts/windows/test-windows-rdna4.ps1 index e65bf40..c343d50 100644 --- a/scripts/windows/test-windows-rdna4.ps1 +++ b/scripts/windows/test-windows-rdna4.ps1 @@ -13,6 +13,7 @@ param( [ValidateSet(5, 11, 13)][int]$Preset = 11, [ValidateSet(0,11)][int]$NgxCreateFlags = 0, [string]$NgxOutputResolution = '512x288', + [string]$NgxInputResolution, [switch]$NgxOnly, [switch]$Trace, [switch]$RequireNativeNetwork, @@ -90,6 +91,19 @@ function Invoke-Probe([string]$Name, [string]$Exe, [string[]]$Arguments) { $info.CreateNoWindow = $true $info.RedirectStandardOutput = $true $info.RedirectStandardError = $true + # Some PowerShell/.NET hosts restore a removed variable as an empty value. + # These task switches/paths use an empty value to mean unset, whereas Rust + # presence-based diagnostic flags would otherwise enable profiling. + foreach ($key in @($info.EnvironmentVariables.Keys)) { + if ($key.StartsWith('D4R_') -and [string]::IsNullOrWhiteSpace($info.EnvironmentVariables[$key])) { + $info.EnvironmentVariables.Remove($key) + } + } + # Disabled WGP has the ordinary CU semantics. Omit a disabled/empty switch + # from child environments so older ZLUDA builds keep their existing cache. + if ([Environment]::GetEnvironmentVariable('D4R_ZLUDA_WGP','Process') -ne '1') { + $info.EnvironmentVariables.Remove('D4R_ZLUDA_WGP') + } $process = New-Object System.Diagnostics.Process $process.StartInfo = $info if (!$process.Start()) { throw "Process start failed: $Exe" } @@ -118,7 +132,7 @@ try { if (!$env:ZLUDA_CACHE_DIR) { $env:ZLUDA_CACHE_DIR = Join-Path $repoRoot 'build/zluda-cache-windows' } $summary.zludaCacheDirectory = $env:ZLUDA_CACHE_DIR foreach ($setting in @('D4R_ZLUDA_WMMA','D4R_ZLUDA_WMMA_FP8','D4R_ZLUDA_WMMA_FP8_NATIVE','D4R_ZLUDA_WMMA_F16_REFERENCE')) { - if ($null -eq $originalCodegen[$setting]) { + if ([string]::IsNullOrWhiteSpace($originalCodegen[$setting])) { [Environment]::SetEnvironmentVariable($setting, $(if ($setting -eq 'D4R_ZLUDA_WMMA_FP8_NATIVE') { '0' } else { '1' }), 'Process') } } @@ -315,6 +329,7 @@ DisableSplash=true } $ngxArguments += @('--module', $localShim, '--pixel-profile', $PixelProfile, '--barrier-mode', $BarrierMode, '--ngx-create-flags', "$NgxCreateFlags", '--ngx-output-resolution', $NgxOutputResolution) + if ($NgxInputResolution) { $ngxArguments += @('--ngx-input-resolution', $NgxInputResolution) } if ($EarlyIndirectProbe) { if (!$CommandListBackend) { throw '-EarlyIndirectProbe requires -CommandListBackend' } $ngxArguments += @('--early-indirect', '1') diff --git a/scripts/windows/windows-game.ps1 b/scripts/windows/windows-game.ps1 index 7214f72..21d8b7a 100644 --- a/scripts/windows/windows-game.ps1 +++ b/scripts/windows/windows-game.ps1 @@ -251,6 +251,16 @@ try { $info.WorkingDirectory = $game $info.UseShellExecute = $false; $info.CreateNoWindow = $true $info.RedirectStandardOutput = $true; $info.RedirectStandardError = $true + # Empty task switches mean unset. Some PowerShell/.NET hosts preserve empty + # values, which would enable Rust's presence-based diagnostic flags. + foreach ($key in @($info.EnvironmentVariables.Keys)) { + if ($key.StartsWith('D4R_') -and [string]::IsNullOrWhiteSpace($info.EnvironmentVariables[$key])) { + $info.EnvironmentVariables.Remove($key) + } + } + if ([Environment]::GetEnvironmentVariable('D4R_ZLUDA_WGP','Process') -ne '1') { + $info.EnvironmentVariables.Remove('D4R_ZLUDA_WGP') + } $process = [Diagnostics.Process]::new(); $process.StartInfo = $info if (!$process.Start()) { throw 'Game diagnostic launch failed.' } $outFile = [IO.File]::Open((Join-Path $OutputDirectory 'd4r.stdout.log'), [IO.FileMode]::Create, [IO.FileAccess]::Write, [IO.FileShare]::Read) diff --git a/tools/windows/benchmark_native_replay.py b/tools/windows/benchmark_native_replay.py new file mode 100644 index 0000000..e5c7220 --- /dev/null +++ b/tools/windows/benchmark_native_replay.py @@ -0,0 +1,104 @@ +"""Compare resident native modules with exact replay and paired batch timings. + +Fixtures can contain private weights; this utility never packages them. +""" +import argparse +import hashlib +import json +import os +import pathlib +import subprocess +import sys + +from profile_report import report + + +def sha(path): + return hashlib.sha256(path.read_bytes()).hexdigest() + + +def run(args): + args.output_dir.mkdir(parents=True, exist_ok=False) + summary = dict(passed=False, numericalExact=False, fixture=str(args.fixture_dir), + module=str(args.module), candidate=str(args.candidate), probe=str(args.probe), + pairs=args.pairs, batches=args.batches, + timingSources=args.timings, executions=[]) + environment = os.environ.copy() + environment['D4R_DIAG_DIR'] = str(args.output_dir) + environment['PATH'] = str(args.hip_root / 'bin') + os.pathsep + environment.get('PATH', '') + try: + summary.update(moduleSha256=sha(args.module), candidateSha256=sha(args.candidate), + probeSha256=sha(args.probe)) + def execute(name, module, options=()): + command = [str(args.probe), '--hip-root', str(args.hip_root), '--module', str(module), + '--fixture-dir', str(args.fixture_dir), '--output-dir', str(args.output_dir / name), + '--iterations', '1', *options] + with (args.output_dir / (name + '.stdout.log')).open('wb') as out, \ + (args.output_dir / (name + '.stderr.log')).open('wb') as err: + try: + completed = subprocess.run(command, stdout=out, stderr=err, env=environment, timeout=180) + except subprocess.TimeoutExpired: + summary['executions'].append(dict(name=name, command=command, timedOut=True)) + raise + summary['executions'].append(dict(name=name, command=command, exitCode=completed.returncode)) + if completed.returncode: + raise RuntimeError(f'{name} failed ({completed.returncode}); inspect preserved logs/dump') + return args.output_dir / name + + control = execute('control-reference', args.module) + candidate = execute('candidate-reference', args.candidate) + reference = {p.name: sha(p) for p in control.glob('alloc-*.bin')} + if not reference: + raise RuntimeError('Replay produced no allocations') + + def check(folder): + actual = {p.name: sha(p) for p in folder.glob('alloc-*.bin')} + if actual != reference: + changed = sorted(set(actual) ^ set(reference) | + {name for name in reference if actual.get(name) != reference[name]}) + raise RuntimeError(f'{folder.name}: allocations differ from fresh control: {changed}') + + # Stop before timing a candidate that changes the numerical result. + check(candidate) + summary['numericalExact'] = True + summary['allocationSha256'] = reference + for timing in args.timings: + for batch in args.batches: + folder = execute(f'{timing}-{batch}', args.module, + ['--benchmark-module', str(args.candidate), '--benchmark-timing', timing, + '--benchmark-batch', str(batch), '--iterations', str(args.pairs)]) + check(folder) + measured = report(args.output_dir) + summary['profile'] = measured + expected = len(args.timings) * len(args.batches) + if len(measured['replayPairs']) != expected or any( + row['controlGpuMs']['samples'] != args.pairs for row in measured['replayPairs']): + raise RuntimeError('Incomplete/invalid timing pairs; this run is not accepted for performance comparison') + summary['passed'] = True + for row in measured['replayPairs']: + print(f"PAIRED {row['kernel']} {row['timingSource']} batch={row['launchesPerSample']} " + f"control={row['controlGpuMs']['median']:.6f} candidate={row['candidateGpuMs']['median']:.6f} ms " + f"ratio={row['candidateToControl']['median']:.6f}") + print('PASS NATIVE_BENCHMARK exact_allocations=1 timings_valid=1 ' + 'isolated_replay=1 game_speedup=unmeasured') + except Exception as error: + summary['error'] = str(error) + print(f'FAIL NATIVE_BENCHMARK {error}', file=sys.stderr) + finally: + (args.output_dir / 'summary.json').write_text(json.dumps(summary, indent=2) + '\n', encoding='utf-8') + return 0 if summary['passed'] else 5 + + +if __name__ == '__main__': + parser = argparse.ArgumentParser(description=__doc__) + for name in ('probe', 'hip-root', 'module', 'candidate', 'fixture-dir', 'output-dir'): + parser.add_argument('--' + name, required=True, type=lambda s: pathlib.Path(s).resolve()) + parser.add_argument('--pairs', type=int, default=32) + parser.add_argument('--batches', type=int, nargs='+', default=[1, 64]) + parser.add_argument('--timings', choices=['events', 'dispatch'], nargs='+', default=['events', 'dispatch']) + args = parser.parse_args() + if not 1 <= args.pairs <= 10000 or any(not 1 <= batch <= 256 for batch in args.batches): + parser.error('pairs must be 1..10000 and batches 1..256') + if len(set(args.batches)) != len(args.batches) or len(set(args.timings)) != len(args.timings): + parser.error('duplicate batches/timing sources are not allowed') + raise SystemExit(run(args)) diff --git a/tools/windows/d3d12_dlss_probe.cpp b/tools/windows/d3d12_dlss_probe.cpp index 75bf48c..8d37b60 100644 --- a/tools/windows/d3d12_dlss_probe.cpp +++ b/tools/windows/d3d12_dlss_probe.cpp @@ -74,7 +74,8 @@ int main(int argc, char** argv) { dx(device->CreateCommandAllocator(D3D12_COMMAND_LIST_TYPE_DIRECT, IID_PPV_ARGS(allocator.GetAddressOf())), "CreateCommandAllocator"); dx(device->CreateCommandList(0, D3D12_COMMAND_LIST_TYPE_DIRECT, allocator.Get(), nullptr, IID_PPV_ARGS(list.GetAddressOf())), "CreateCommandList"); const unsigned outWidth = args.output_width, outHeight = args.output_height; - const unsigned width = outWidth / 2, height = outHeight / 2; + const unsigned width = args.input_width ? args.input_width : outWidth / 2; + const unsigned height = args.input_height ? args.input_height : outHeight / 2; std::printf("D3D12_NGX_DIMENSIONS input=%ux%u output=%ux%u\n", width, height, outWidth, outHeight); struct Texture { ComPtr image, upload; D3D12_PLACED_SUBRESOURCE_FOOTPRINT footprint{}; UINT64 bytes; } textures[5]; for (unsigned i = 0; i < 5; ++i) { diff --git a/tools/windows/diagnostic.h b/tools/windows/diagnostic.h index 19213b0..ad17731 100644 --- a/tools/windows/diagnostic.h +++ b/tools/windows/diagnostic.h @@ -127,8 +127,11 @@ struct Args { unsigned preset = 11; unsigned ngx_create_flags = 0; unsigned output_width = 512, output_height = 288; + unsigned input_width = 0, input_height = 0; std::string fixture_dir; std::string benchmark_module; + std::string benchmark_timing = "events"; + unsigned benchmark_batch = 1; std::string output_dir; std::string kernel_name = "enc1"; unsigned iterations = 32; @@ -178,6 +181,17 @@ struct Args { output_height < 128 || output_height > 2160 || output_width % 2 || output_height % 2) throw std::runtime_error("Output resolution must be even and within 128x128..4096x2160"); } + else if (key == "--ngx-input-resolution") { + const auto separator = value.find('x'); + if (separator == std::string::npos) throw std::runtime_error("Input resolution must be WIDTHxHEIGHT"); + size_t wend = 0, hend = 0; + const auto width = std::stoul(value.substr(0, separator), &wend); + const auto height = std::stoul(value.substr(separator + 1), &hend); + if (wend != separator || hend != value.size() - separator - 1 || width < 64 || width > 4096 || + height < 64 || height > 2160) + throw std::runtime_error("Input resolution must be within 64x64..4096x2160"); + input_width = static_cast(width); input_height = static_cast(height); + } else if (key == "--preset") { size_t end = 0; preset = static_cast(std::stoul(value, &end)); @@ -186,6 +200,18 @@ struct Args { } else if (key == "--fixture-dir") fixture_dir = value; else if (key == "--benchmark-module") benchmark_module = value; + else if (key == "--benchmark-timing") { + if (value != "events" && value != "dispatch") + throw std::runtime_error("benchmark-timing must be events or dispatch"); + benchmark_timing = value; + } + else if (key == "--benchmark-batch") { + size_t end = 0; + const auto parsed = std::stoul(value, &end); + if (end != value.size() || parsed < 1 || parsed > 256) + throw std::runtime_error("benchmark-batch must be 1..256"); + benchmark_batch = static_cast(parsed); + } else if (key == "--output-dir") output_dir = value; else if (key == "--kernel-name") kernel_name = value; else if (key == "--iterations") { @@ -200,6 +226,8 @@ struct Args { } else throw std::runtime_error("Unknown option " + key); } if (hip_root.empty()) throw std::runtime_error("--hip-root is required"); + if (input_width > output_width || input_height > output_height) + throw std::runtime_error("Input resolution cannot exceed output resolution"); } }; class Library { diff --git a/tools/windows/native_replay_probe.cpp b/tools/windows/native_replay_probe.cpp index fecc8ac..aae1a38 100644 --- a/tools/windows/native_replay_probe.cpp +++ b/tools/windows/native_replay_probe.cpp @@ -2,6 +2,9 @@ #include #include #include +#include +#include +#include namespace { std::vector read_file(const std::filesystem::path& path) { @@ -100,6 +103,8 @@ int main(int argc, char** argv) { return value ? value : fallback; }; auto scalar = [&](const char* name, unsigned fallback) { return module_scalar(module, name, fallback); }; + const std::vector captured_grid{grid[0], grid[1], grid[2]}; + const unsigned captured_block_z = block[2]; block[2] = scalar("d4r_block_z", block[2]); const unsigned persistentGrid = scalar("d4r_grid_x", 0); if (persistentGrid) { grid[0] = persistentGrid; grid[1] = grid[2] = 1; } @@ -117,9 +122,16 @@ int main(int argc, char** argv) { hipModule_t candidate = nullptr; hip.check(hip.hipModuleLoad(&candidate, args.benchmark_module.c_str()), "Load paired benchmark candidate"); Cleanup candidate_cleanup{hip, candidate}; - if (module_scalar(candidate, "d4r_block_z", block[2]) != block[2] || - module_scalar(candidate, "d4r_grid_x", 0) != persistentGrid) - throw std::runtime_error("Paired modules require identical launch geometry"); + unsigned grids[2][3] = {{grid[0], grid[1], grid[2]}, + {captured_grid[0], captured_grid[1], captured_grid[2]}}; + unsigned blocks[2][3] = {{block[0], block[1], block[2]}, + {block[0], block[1], module_scalar(candidate, "d4r_block_z", captured_block_z)}}; + const unsigned candidate_persistent = module_scalar(candidate, "d4r_grid_x", 0); + if (candidate_persistent) { grids[1][0] = candidate_persistent; grids[1][1] = grids[1][2] = 1; } + for (unsigned variant = 0; variant < 2; ++variant) + std::printf("D4R_REPLAY_MODULE variant=%s grid=%u,%u,%u block=%u,%u,%u geometry_source=native_metadata\n", + variant ? "candidate" : "control", grids[variant][0], grids[variant][1], grids[variant][2], + blocks[variant][0], blocks[variant][1], blocks[variant][2]); hipFunction_t other = nullptr, other_prep = nullptr; hip.check(hip.hipModuleGetFunction(&other, candidate, kernel.c_str()), "Find paired benchmark function"); if (hip.hipModuleGetFunction(&other_prep, candidate, (kernel + "_prep").c_str()) == hipSuccess) { @@ -133,6 +145,22 @@ int main(int argc, char** argv) { auto record = hip.library.symbol("hipEventRecord"); auto synchronize = hip.library.symbol("hipEventSynchronize"); auto elapsed = hip.library.symbol("hipEventElapsedTime"); + // Documented hip/hip_ext.h ABI. Resolve dynamically like the other + // HIP APIs so the host probe does not depend on HIP host linking. + using DispatchFn = hipError_t (*)(hipFunction_t, uint32_t, uint32_t, uint32_t, + uint32_t, uint32_t, uint32_t, size_t, hipStream_t, void**, void**, + hipEvent_t, hipEvent_t, uint32_t); + DispatchFn dispatch = nullptr; + uint32_t global[2][3]{}; + if (args.benchmark_timing == "dispatch") { + dispatch = hip.library.symbol("hipExtModuleLaunchKernel"); + for (unsigned variant = 0; variant < 2; ++variant) for (unsigned axis = 0; axis < 3; ++axis) { + const auto items = uint64_t(grids[variant][axis]) * blocks[variant][axis]; + if (!items || items > std::numeric_limits::max()) + throw std::runtime_error("Dispatch profiling work-item dimension overflow"); + global[variant][axis] = static_cast(items); + } + } struct Event { decltype(destroy) release; hipEvent_t value = nullptr; ~Event() { if (value) (void)release(value); } @@ -140,7 +168,9 @@ int main(int argc, char** argv) { hip.check(create(&begin.value), "Create paired start event"); hip.check(create(&end.value), "Create paired end event"); auto launch = [&](hipFunction_t fn) { - hip.check(hip.hipModuleLaunchKernel(fn, grid[0], grid[1], grid[2], block[0], block[1], block[2], + const unsigned variant = fn == other; + hip.check(hip.hipModuleLaunchKernel(fn, grids[variant][0], grids[variant][1], grids[variant][2], + blocks[variant][0], blocks[variant][1], blocks[variant][2], shared, nullptr, launchArguments, nullptr), "Paired benchmark launch"); }; // Both modules and weight images stay resident on the same GPU. @@ -156,13 +186,31 @@ int main(int argc, char** argv) { for (const auto& allocation : allocations) hip.check(hip.hipMemcpy(allocation->device, allocation->initial.data(), allocation->size, hipMemcpyHostToDevice), "Restore paired captured input"); - hip.check(record(begin.value, nullptr), "Record paired start"); - launch(variant ? other : function); - hip.check(record(end.value, nullptr), "Record paired end"); + const auto host_begin = std::chrono::steady_clock::now(); + if (!dispatch) hip.check(record(begin.value, nullptr), "Record paired start"); + for (unsigned batch = 0; batch < args.benchmark_batch; ++batch) { + if (dispatch) { + // Flags zero preserve stream order. Global sizes are + // work-items, unlike hipModuleLaunchKernel's blocks. + hip.check(dispatch(variant ? other : function, global[variant][0], global[variant][1], global[variant][2], + blocks[variant][0], blocks[variant][1], blocks[variant][2], shared, nullptr, launchArguments, nullptr, + batch == 0 ? begin.value : nullptr, + batch + 1 == args.benchmark_batch ? end.value : nullptr, 0), + "Paired dispatch-profiled launch"); + } else launch(variant ? other : function); + } + if (!dispatch) hip.check(record(end.value, nullptr), "Record paired end"); hip.check(synchronize(end.value), "Complete paired sample"); + const auto host_ms = std::chrono::duration( + std::chrono::steady_clock::now() - host_begin).count(); float ms = 0; hip.check(elapsed(&ms, begin.value, end.value), "Paired GPU elapsed time"); - std::printf("D4R_REPLAY_PROFILE kernel=%s variant=%s pair=%u gpu_ms=%.6f warmup_launches=512 paired=1 input_restore_outside_event=1\n", - kernel.c_str(), variant ? "candidate" : "control", pair, double(ms)); + const bool valid = std::isfinite(ms) && ms > 0 && double(ms) <= host_ms; + std::printf("D4R_REPLAY_PROFILE%s kernel=%s variant=%s pair=%u gpu_ms=%.6f " + "host_ms=%.6f timing_source=%s launches=%u warmup_launches=512 paired=1 " + "input_restore_outside_event=1\n", valid ? "" : "_INVALID", + kernel.c_str(), variant ? "candidate" : "control", pair, + double(ms) / args.benchmark_batch, host_ms / args.benchmark_batch, + args.benchmark_timing.c_str(), args.benchmark_batch); } hip.verbose = verbose; // Leave a single ordinary control replay for numerical checking. diff --git a/tools/windows/profile_report.py b/tools/windows/profile_report.py index 9d86a6c..1d012ed 100644 --- a/tools/windows/profile_report.py +++ b/tools/windows/profile_report.py @@ -119,12 +119,24 @@ def report(directory, metadata_directory=None): invalid_samples = [] boundaries = collections.defaultdict(list) replay = collections.defaultdict(lambda: collections.defaultdict(dict)) + invalid_replay = [] for path in sorted(directory.glob('*.log')): data = path.read_bytes() text = data.decode('utf-16' if data.startswith((b'\xff\xfe', b'\xfe\xff')) else 'utf-8-sig', errors='replace') - for name, variant, pair, value in re.findall( - r'D4R_REPLAY_PROFILE kernel=(\S+) variant=(control|candidate) pair=([0-9]+) gpu_ms=([0-9.]+)', text): - replay[name][int(pair)][variant] = float(value) + for line in text.splitlines(): + if line.startswith('D4R_REPLAY_PROFILE_INVALID '): + invalid_replay.append(dict(file=path.name, **dict(re.findall(r'(\w+)=(\S+)', line)))) + if not line.startswith('D4R_REPLAY_PROFILE '): + continue + fields = dict(re.findall(r'(\w+)=(\S+)', line)) + if not all(key in fields for key in ('kernel', 'variant', 'pair', 'gpu_ms')): + continue + value = float(fields['gpu_ms']) + if fields['variant'] not in ('control', 'candidate') or not math.isfinite(value) or value <= 0: + invalid_replay.append(dict(file=path.name, **fields)) + continue + key = (path.name, fields['kernel'], fields.get('timing_source', 'events'), int(fields.get('launches', 1))) + replay[key][int(fields['pair'])][fields['variant']] = fields for name, value in re.findall(r'D4R_STAGE name=(\S+) cpu_ms=([0-9.]+)', text): stages[name].append(float(value)) for kind, fields in re.findall(r'(?m)^D4R_COMMAND_(RECORD|SUBMIT) ([^\r\n]+)', text): @@ -172,6 +184,7 @@ def report(directory, metadata_directory=None): presentMon=presentmon(directory), gpuTelemetry=gpu_telemetry(directory), replayPairs=[], + invalidReplaySamples=invalid_replay, kernels=[], notes=[ 'CPU stage times include waits and host work; they are not isolated GPU timings.', 'CUDA API profiles measure host call duration without GPU events or added synchronization; they still include existing waits.', @@ -179,6 +192,7 @@ def report(directory, metadata_directory=None): 'Boundary timestamps are read after existing completion fences; only 32 timing bytes are read, with no extra completion wait. GPU clock idle behavior can affect intervals.', 'PresentMon ETW GPU busy/wait are per-process frame estimates, not hardware sensor utilization. HWS and cross-API context attribution can affect them.', 'ADLX telemetry is driver-reported device-wide utilization/clocks/power, sampled in a separate read-only process. It does not isolate DLSS or WMMA occupancy.', + 'Replay GPU/host intervals are normalized per launch; multi-launch batches include queue/dispatch gaps and are not isolated instruction time or application FPS.', '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.', @@ -202,13 +216,18 @@ def report(directory, metadata_directory=None): boundary_stages['between_boundaries_ms'].append(before * 1000. / frequency) boundary_stages['input_interval_ms'].append(interval * 1000. / frequency) result['gpuBoundaryStages'] = {name: stats(values) for name, values in boundary_stages.items()} - for name, pairs in replay.items(): + for (file, name, timing, launches), pairs in replay.items(): complete = [row for row in pairs.values() if set(row) == {'control', 'candidate'}] if complete: - result['replayPairs'].append(dict(kernel=name, - controlGpuMs=stats([row['control'] for row in complete]), - candidateGpuMs=stats([row['candidate'] for row in complete]), - candidateToControl=stats([row['candidate'] / row['control'] for row in complete if row['control'] > 0]))) + values = {variant: [float(row[variant]['gpu_ms']) for row in complete] for variant in ('control', 'candidate')} + item = dict(file=file, kernel=name, timingSource=timing, launchesPerSample=launches, + controlGpuMs=stats(values['control']), candidateGpuMs=stats(values['candidate']), + candidateToControl=stats([candidate / control for control, candidate in zip(values['control'], values['candidate'])])) + for variant in ('control', 'candidate'): + host = [float(row[variant]['host_ms']) for row in complete if 'host_ms' in row[variant]] + if host: + item[variant + 'HostCompletionMs'] = stats(host) + result['replayPairs'].append(item) for (thread, api), rows in apis.items(): calls = sum(int(row['calls']) for row in rows) total = sum(float(row['total_ms']) for row in rows) @@ -276,7 +295,7 @@ def main(): print(f"HOOK thread={row['thread']} samples={int(row['samples'])} access={row['accessMeanUs']:.3f} us " f"capture={row['captureMeanUsPerSample']:.3f} us estimated_cpu={row['estimatedAccessCaptureMsPerSecond']:.3f} ms/s") for row in result['replayPairs']: - print(f"PAIRED {row['kernel']}: n={row['controlGpuMs']['samples']} control_median={row['controlGpuMs']['median']:.6f} ms " + print(f"PAIRED {row['kernel']} timing={row['timingSource']} batch={row['launchesPerSample']}: n={row['controlGpuMs']['samples']} control_median={row['controlGpuMs']['median']:.6f} ms " f"candidate_median={row['candidateGpuMs']['median']:.6f} ms ratio_median={row['candidateToControl']['median']:.6f}") for row in result['kernels']: times = row['gpuMs']