From 0c18df62536ccf4ad671ae9ce72791dccd38e42c Mon Sep 17 00:00:00 2001 From: Codex Date: Thu, 1 Oct 2026 21:16:28 +0300 Subject: [PATCH] Validate Windows native M FP8 and reduce operand packing cost --- README.md | 10 +- docs/windows-game.md | 12 ++- docs/windows-performance.md | 58 +++++++++++- docs/windows-rdna4-port.md | 55 +++++++++-- kernels/m/swin_common.h | 14 +++ scripts/windows/build-native-m-fp8.ps1 | 59 ++++++++++++ scripts/windows/build-native-texture.ps1 | 2 +- scripts/windows/build-zluda-windows.ps1 | 2 +- scripts/windows/package-windows-game.ps1 | 14 ++- scripts/windows/profile-windows-rdna4.ps1 | 2 +- scripts/windows/test-native-m-fp8.ps1 | 28 ++++++ scripts/windows/test-native-overrides.ps1 | 2 +- scripts/windows/test-native-texture.ps1 | 2 +- scripts/windows/test-windows-rdna4.ps1 | 2 +- tools/windows/diagnostic.h | 2 + tools/windows/fp8_pack_probe.hip | 9 ++ tools/windows/fp8_pack_validate.py | 46 +++++++++ tools/windows/m_native_replay_compare.py | 109 ++++++++++++++++++++++ tools/windows/native_replay_probe.cpp | 71 +++++++++++++- tools/windows/profile_report.py | 20 +++- 20 files changed, 486 insertions(+), 33 deletions(-) create mode 100644 scripts/windows/build-native-m-fp8.ps1 create mode 100644 scripts/windows/test-native-m-fp8.ps1 create mode 100644 tools/windows/fp8_pack_probe.hip create mode 100644 tools/windows/fp8_pack_validate.py create mode 100644 tools/windows/m_native_replay_compare.py diff --git a/README.md b/README.md index c4a375c..2c2d131 100644 --- a/README.md +++ b/README.md @@ -14,13 +14,17 @@ submission, and returns output in the same frame without CPU image copies. RGBA16/32F, R11G11B10, RGBA/BGRA8, depth, exposure and motion formats have independent D3D12 typed SRV/UAV conversion tests. Four-frame K/M runs with packed output match the FP16 baseline exactly after resource-format quantisation. -The build incorporates upstream RDNA4 commit `dbef4b2`. Patched OptiScaler now +The build incorporates upstream updates through `f0d1a65`. Patched OptiScaler now passes standalone K/M native-network tests with exact RGB agreement against the direct-d4r baseline; see [the source patch](patches/optiscaler/README.md). Silent Hill 2 now renders K through the native Windows backend; M also renders a saved gameplay level with GPU output NaN/Inf checks. Startup depth/stencil and swap-chain lifetime issues found in the game have regression checks. -Performance and wider resolution coverage remain in progress. See +Both K and M now pass 3840x2160 game output scans. K reaches a user-reported +49-51 FPS in the measured scene after native output stores and command-hook +lookup caching. Native M FP8 is separately replay/network validated; the +default keeps the FP16-equivalent baseline based on measured performance. +All sixteen CTest gates pass with the final native Windows ZLUDA runtime. See [the reversible game package](docs/windows-game.md) and [build instructions and hardware results](docs/windows-rdna4-port.md). @@ -53,7 +57,7 @@ On the same GPU DLSS starts at a disadvantage: its networks were designed for NV For users who prioritize fidelity over speed, set `[Kernels] PreferAccuracy = true` in `d4r/d4r.ini` and restart the game. It defaults to `false`. This selects accuracy variants of every native kernel and restores conservative translation settings. It aims to match NVIDIA's arithmetic; 1:1 image quality against RTX DLSS is not yet proven. See [native kernel numerics](docs/native-kernels.md#numerics). -## GPU support +## GPU support (Linux upstream) d4r builds for RDNA3 and RDNA4. A newly built release compiles native DLSS 4 and 4.5 network kernels for the targets below and selects the matching set at runtime. Only the RX 7700 XT has been tested by this project on a real GPU; an external video reports DLSS 4.5 running through d4r 0.1.2 on an RX 7900 XTX. RDNA4 runtime and performance remain unverified on hardware. diff --git a/docs/windows-game.md b/docs/windows-game.md index 3374229..fbbfcd8 100644 --- a/docs/windows-game.md +++ b/docs/windows-game.md @@ -3,21 +3,27 @@ Validated hardware: Windows 11 x64 / Radeon RX 9070 XT / gfx1201. Native K and M transformer kernels pass standalone D3D12 + OptiScaler checks. Silent Hill 2 K now renders real menu frames and M renders a saved gameplay level with -GPU NaN/Inf checks. K also passes 3840x2160 output and M passes 1920x1080. -Further performance work is in progress. This is a local +GPU NaN/Inf checks. Both K and M pass 3840x2160 output on this machine. +The user reports 49-51 FPS for K in the measured scene after the command-hook +fix and locally built output-store optimization. This is a local development package; no NVIDIA proprietary DLL is included. Build the shim/diagnostics, corrected ZLUDA, K/M native objects and patched OptiScaler using the scripts in `scripts/windows`, then run: ```powershell -.\scripts\windows\package-windows-game.ps1 +.\scripts\windows\package-windows-game.ps1 -ArchivePath "$PWD\dist\windows-rdna4-game.zip" ``` The package pins TheRock 10.2.0a20260929 because its Windows external-memory mapping lifetime passes the repeated import/release test. Stable HIP 7.2 has a reproduced mapped-view leak; see `docs/windows-rdna4-port.md` for the source fix, exact dependency versions, validation and complete build commands. +The ZIP contains only manifest-listed files, including a snapshot of committed +d4r source. Private NVIDIA-derived objects, caches and diagnostic captures are +excluded. Original FP8 M and the optional hardware-packing variant pass strict +replay/network tests but do not establish a full-network speedup, so the game +package retains its FP16-equivalent M baseline. Run a D3D12 game without anti-cheat using your locally supplied NVIDIA DLLs: diff --git a/docs/windows-performance.md b/docs/windows-performance.md index c13560b..501a2ba 100644 --- a/docs/windows-performance.md +++ b/docs/windows-performance.md @@ -69,9 +69,8 @@ memory, LDS and predicted blocks per multiprocessor. Add `--metadata-directory ` for native wave sizes and compiler spill counts. Reports and proprietary workload inputs remain local. -Native FP8 is not the validated baseline and has no performance claim here. -Replacing the exact per-MMA f16 rounding with f32 FP8 accumulation requires a -new layer/reference validation; it is not justified by the observed CPU stalls. +Native FP8 remains a separate optional experiment; the package uses the +validated FP16-equivalent baseline. Its later checks and timings are below. Further game coverage (HIP kernel profiling disabled, GPU output checks enabled): M at 1920x1080 completes 4833 frames, all finite. K at 3840x2160 completes @@ -172,5 +171,58 @@ each physical command-backend test and the software indirect-root regression. Use `-ProfileCommandHooks -ProfileStages` to collect samples; the sampler is disabled for ordinary play. +Native M FP8 was subsequently compiled and tested on this gfx1201. ISA confirms +`v_wmma_f32_16x16x16_fp8_fp8`, with packed e4m3 weight images; exact FP16 +attention/PV remains. All logical outputs and merges in forty recorded launches +match the previously NumPy-validated FP16 baseline byte-for-byte. Full DLSS RGB +also matches exactly: four 512x288 LDR frames and four 3840x2160 HDR frames, +all RGBA finite. This establishes correctness for those fixtures, not a speedup. + +The original software operand packing makes the FP8 network slower in the +four-frame 4K test: sum of five main-kernel event means (tube weighted six times) +is 27.789640 ms versus FP16's 23.550875 ms. A paired in-process replay confirms +the regression without relying solely on separate process timings. Both modules +and prepared weights remain resident, 512 launches warm the GPU, then sixty-four +AB/BA pairs time the same restored input using HIP events. Diagnostic CPU input +restoration is outside the timed event and is unrelated to game VRAM interop. + +Added opt-in `D4R_FP8_HW_PACK`: native packed conversion of already quantised +operands replaces software exponent/subnormal reconstruction. It preserves all +254 finite OCP e4m3 encodings, including signed zero, in an exhaustive GPU check. +All forty recorded launches still match exactly. Four complete 4K HDR frames +also match FP16 RGB exactly with finite RGBA. ISA confirms both packed converts +and FP8 WMMA. The baseline build's sixteen objects remain byte-identical. + +| Paired replay median, ms | FP16 (software-packing test) | FP8 software packing | FP16 (hardware-packing test) | FP8 hardware packing | +| --- | ---: | ---: | ---: | ---: | +| enc1 | 0.175450 | 0.217050 | 0.175100 | 0.177700 | +| enc2 | 0.141350 | 0.164550 | 0.132550 | 0.131550 | +| enc3 tube | 0.108000 | 0.112000 | 0.107600 | 0.099850 | +| dec2 | 0.129900 | 0.146000 | 0.143250 | 0.140850 | +| dec1 | 0.169650 | 0.194450 | 0.154250 | 0.156550 | + +These are two paired tests on the captured 512x288 workload. Medians avoid large +isolated scheduling outliers; raw per-pair ratios are retained. They identify a +substantial cost in software packing but do not establish an overall FPS gain. +The separately run four-frame 4K hardware-packing network mean is 24.603463 ms; +it still does not beat the measured FP16 total. Default packaging retains FP16. +Hardware counters for bandwidth or WMMA utilization were not measured. + +Reproduce the isolated build, encoding/replay checks and paired timings using +the existing private validated captures; no source editing is required: + +```powershell +.\scripts\windows\build-native-m-fp8.ps1 -DlssDll "$PWD\nvngx_dlss.dll" -HardwarePacking +.\scripts\windows\test-native-m-fp8.ps1 -ModuleDirectory "$PWD\build\native-m-fp8-hardware-pack" -CaptureDirectory "$PWD\test-results\ngx-m-exact-pv-capture\captures" -BaselineValidation "$PWD\test-results\m-native-exact-pv-temporal-reference\summary.json" -Benchmark +``` + +Omit `-HardwarePacking` to compile the original FP8 experiment. The replay +validator stops at the first mismatch, retains GPU stdout/stderr/exit code and +hashes every compared logical output. `profile_report.py` accepts replay logs +and reports per-pair GPU time and candidate/control ratios. This experiment is +not installed by the game runner or substituted silently for the baseline. +The conversion builtin is listed in the official +[Clang AMDGPU builtin reference](https://clang.llvm.org/docs/AMDGPUBuiltinReference.html#builtin-amdgcn-cvt-pk-fp8-f32). + References: [HIP events](https://rocmdocs.amd.com/projects/HIP/en/develop/doxygen/html/group___event.html), [HIP occupancy API](https://rocm.docs.amd.com/projects/HIP/en/latest/doxygen/html/group___occupancy.html). diff --git a/docs/windows-rdna4-port.md b/docs/windows-rdna4-port.md index 237765d..583fb69 100644 --- a/docs/windows-rdna4-port.md +++ b/docs/windows-rdna4-port.md @@ -2,10 +2,11 @@ Branch: `windows-rdna4`. Target: Windows 11 x64, RX 9070 XT, **gfx1201**. Priority: correct K, correct M, native Windows, same-frame output, then speed. -This is a work log for an implementation in progress. Standalone DLSS K/M +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 -validation coverage are recorded below. Performance remains in progress. +validation coverage are recorded below. Performance and game coverage have +explicit limits; the validated package keeps conservative arithmetic. ## Milestone and gates @@ -16,9 +17,44 @@ M attention baseline. Upstream's optional denormal override now drops only FP32 requirements, preserving FP16/FP64. All sixteen rebuilt native K/M objects have the same SHA256 as the previously validated objects. All thirteen ZLUDA patches apply to the pinned fresh base; all sixteen CTest gates pass. -The local ZLUDA source includes that denormal correction at `2378af7`; the -measured runtime remains the separately recorded `b0161a4` binary, with the -override unset. Package metadata records the binary build's actual commit. +Final ZLUDA source and binary are `2378af72e586e1730d15f75e0322daab674f0862`, +built cleanly into `dist/zluda-windows-final`; nvcuda SHA256 is +`ccc700974fa5d0b9a5f5c39e617c17ca22f34701d40f745f2f96f039a75b22f8`. +All sixteen CTest gates pass on it. Three K/4K HDR frames and four M baseline +frames remain bit-exact to their prior controls, all RGBA finite. Historical +performance runs below used `b0161a4` with the denormal override unset. +Package metadata records the binary build's actual commit. + +The final-runtime M game test completes 1940 4K frames and 1940 finite GPU +scans, with 19410 native launches (including a partial final frame), zero +backend failures/CPU image copies/previous-frame outputs. Native M FP8 also +passes all logical output/merge values from forty NumPy-validated captures, +plus four 512x288 and four 4K complete frames. Its original software packing +regresses performance; the opt-in hardware packing preserves all 254 finite +FP8 encodings and all forty captures, with exact RGB in four additional 4K +frames. Paired in-process event tests confirm reduced packing cost, but the +full-network measurement does not establish a win over FP16. Production keeps +the baseline. See [windows-performance.md](windows-performance.md) for timings +and one-command replay/paired profiling. + +Current native Windows build, after the pinned tool/source setup described +below (MSVC v143 / Windows SDK 10.0.26100 required for OptiScaler): + +```powershell +.\scripts\windows\build-zluda-windows.ps1 +.\scripts\windows\build-windows-rdna4.ps1 -RuntimeProfile therock -ZludaRoot "$PWD\dist\zluda-windows-final" -InstallDirectory "$PWD\dist\windows-rdna4-command-list" +.\scripts\windows\stage-native-k.ps1 -DlssDll "$PWD\nvngx_dlss.dll" -PackageRoot "$PWD\dist\windows-rdna4-command-list" +.\scripts\windows\stage-native-m.ps1 -DlssDll "$PWD\nvngx_dlss.dll" -PackageRoot "$PWD\dist\windows-rdna4-command-list" +.\scripts\windows\build-optiscaler-windows.ps1 +& .\.tools\python\cmake\data\bin\ctest.exe --test-dir build/windows-rdna4-therock --output-on-failure +.\scripts\windows\package-windows-game.ps1 -ArchivePath "$PWD\dist\windows-rdna4-game.zip" +``` + +The archive requires committed source. It includes a tracked-source snapshot, +all runtime file hashes/versions, dependency licenses and source patches, and +excludes local NVIDIA DLLs, PTX/weights/captures and private texture objects. +The optional texture optimization is built/validated locally and supplied with +`-LocalTextureKernels`; it is never folded into the public source snapshot. Additional production-style runs: M at 1920x1080 completes 4833 frames with 4833 finite-output GPU scans and no backend failures. K at 3840x2160 completes @@ -178,14 +214,17 @@ with `-EarlyIndirectProbe` in the standalone OptiScaler runner. This avoids the previously documented unknown-signature rejection for observed startup objects; signatures created before the frontend is loaded remain untracked. -Current milestone: M1/M2/M4/M5 passed; native NGX Create/Evaluate on the simple +Current milestone: M0-M9 functional gates pass on the tested Windows GPU/game. +The package uses FP16-equivalent K/M and conservative denormal handling. +Native NGX Create/Evaluate on the simple requested E path passes with the user's DLLs. All 11 K layers now pass synthetic and real-weight NumPy/PTX replay checks; native K Evaluate produces four finite frames on Windows. M now passes all five synthetic and real-weight layers against NumPy and independent PTX, and four native CUDA frames. The standalone D3D12 harness now produces four K and four M frames using imported VRAM and shared fences, with bit-exact RGB agreement against the corresponding CUDA harnesses. -Game integration remains in progress. The ordinary NGX D3D12 Evaluate API now +Silent Hill 2 K/M integration passes, including current-frame 4K VRAM interop. +The ordinary NGX D3D12 Evaluate API now records a boundary in an open command list; a native queue interception backend executes prefix, CUDA evaluation, then suffix in the same submitted frame. A full temporal M scan @@ -207,7 +246,7 @@ Do not integrate NGX until the integer PTX workload is stable on the real GPU. | M6: K layers, full transformer and image validation | PASS: 11 real-weight layers, 44 native launches, four D3D12 frames bit-exact against native CUDA, 81.85–92.91 dB versus translated K | | M7: M FP16-equivalent baseline and full transformer | PASS: all blocks of 40 native launches exactly match NumPy; targeted independent PTX regression matches; four D3D12 frames bit-exact against corrected translated M | | M8: standalone Windows D3D12 NGX harness | PASS for ordinary recorded EvaluateFeature and explicit boundaries: K/M, four frames each, no shim CPU image copies; both bit-exact against explicit boundary baselines | -| M9: OptiScaler integration, profiling, installation | Patched frontend and game K/M PASS, including K 4K and M 1080p finite output; native/translated GPU profiling and reversible install PASS; remaining FPS investigation and final archive in progress | +| M9: OptiScaler integration, profiling, installation | Patched frontend and game K/M PASS at 4K with finite output; native/translated/FP8 profiling, reversible install and manifest-only public archive PASS; K user reports 49-51 FPS | ## Public NGX identity and OptiScaler gate (2026-10-01) diff --git a/kernels/m/swin_common.h b/kernels/m/swin_common.h index 862f81a..88d78b1 100644 --- a/kernels/m/swin_common.h +++ b/kernels/m/swin_common.h @@ -273,11 +273,25 @@ __device__ __forceinline__ sop lds16(const half_t* p) // e4m3 bytes of an f16 operand whose values are e4m3 values (K order kept: byte j = element j) __device__ __forceinline__ gop to_fp8(sop v) { +#if defined(D4R_FP8_HW_PACK) && defined(__GFX12__) + // Inputs already contain exact e4m3 values. gfx12's native packed convert + // avoids the software subnormal/sign/exponent reconstruction below. Keep + // opt-in until all finite encodings and layer/network outputs pass replay. + gop result; +#pragma unroll + for (int word = 0; word < 2; ++word) { + int packed = __builtin_amdgcn_cvt_pk_fp8_f32((float)v[4 * word], (float)v[4 * word + 1], 0, false); + packed = __builtin_amdgcn_cvt_pk_fp8_f32((float)v[4 * word + 2], (float)v[4 * word + 3], packed, true); + result[word] = (uint32_t)packed; + } + return result; +#else uint32_t c[4]; #pragma unroll for (int j = 0; j < 4; ++j) c[j] = codes8x2((hv2){v[2 * j], v[2 * j + 1]}); return (u2v){__builtin_amdgcn_perm(c[1], c[0], 0x06040200u), __builtin_amdgcn_perm(c[3], c[2], 0x06040200u)}; +#endif } // f32 += A B over 16 K for e4m3 operands __device__ __forceinline__ f8v wmma8(gop a, gop b, f8v c) diff --git a/scripts/windows/build-native-m-fp8.ps1 b/scripts/windows/build-native-m-fp8.ps1 new file mode 100644 index 0000000..45d1168 --- /dev/null +++ b/scripts/windows/build-native-m-fp8.ps1 @@ -0,0 +1,59 @@ +[CmdletBinding()] +param([Parameter(Mandatory=$true)][string]$DlssDll, [string]$HipRoot, [string]$OutputDirectory, [switch]$HardwarePacking) +$ErrorActionPreference='Stop' +$repo=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) +if (!$HipRoot) { $HipRoot=Join-Path $repo '.tools/therock-10.2.0a20260929/_rocm_sdk_core' } +if (!$OutputDirectory) { $OutputDirectory=Join-Path $repo $(if ($HardwarePacking) { 'build/native-m-fp8-hardware-pack' } else { 'build/native-m-fp8-experimental' }) } +$OutputDirectory=[IO.Path]::GetFullPath($OutputDirectory) +if ($OutputDirectory -eq [IO.Path]::GetFullPath((Join-Path $repo 'build/native-m-gfx1201'))) { + throw 'Keep the validated FP16 baseline separate from the experimental FP8 build.' +} +$DlssDll=(Get-Item -LiteralPath $DlssDll -ErrorAction Stop).FullName +$compiler=Join-Path $HipRoot 'lib/llvm/bin/clang++.exe' +if (!(Test-Path -LiteralPath $compiler)) { $compiler=Join-Path $HipRoot 'bin/clang++.exe' } +if (!(Test-Path -LiteralPath $compiler)) { throw "HIP compiler missing under $HipRoot" } +New-Item -ItemType Directory -Force $OutputDirectory | Out-Null +$disassembler=Join-Path (Split-Path $compiler) 'llvm-objdump.exe' +if (!(Test-Path -LiteralPath $disassembler)) { throw 'HIP llvm-objdump is required to confirm native FP8 instructions.' } +foreach ($source in Get-ChildItem -LiteralPath (Join-Path $repo 'kernels/m') -Filter '*.hip' | Sort-Object Name) { + $flags=@() + if ((Get-Content -LiteralPath $source.FullName -Raw) -match '// d4r-build-flags: ([^\r\n]+)') { $flags=$Matches[1].Trim().Split(' ') } + if ($HardwarePacking) { $flags+='-DD4R_FP8_HW_PACK' } + $object=Join-Path $OutputDirectory ($source.BaseName+'.hsaco') + $ErrorActionPreference='Continue' + & $compiler -x hip -std=c++17 --offload-arch=gfx1201 --offload-device-only --no-gpu-bundle-output ` + -mno-wavefrontsize64 -nogpuinc -nogpulib -O3 -DD4R_DEVICE_ONLY_MINIMAL ` + -DSWIN_EXACT -DSWIN_EXACT_PV -DD4R_FP8_WMMA @flags $source.FullName ` + -o $object > (Join-Path $OutputDirectory ($source.BaseName+'.build.stdout.log')) ` + 2> (Join-Path $OutputDirectory ($source.BaseName+'.build.stderr.log')) + $ErrorActionPreference='Stop' + if ($LASTEXITCODE) { throw "Experimental FP8 compile failed: $($source.Name)" } + $isa=Join-Path $OutputDirectory ($source.BaseName+'.isa.txt') + & $disassembler --disassemble --mcpu=gfx1201 $object > $isa + if ($LASTEXITCODE -or !(Select-String -LiteralPath $isa -Pattern 'v_wmma_f32_16x16x16_fp8_fp8')) { + throw "Native FP8 WMMA is absent from $($source.Name); no FP16 fallback is accepted." + } + if ($HardwarePacking -and !(Select-String -LiteralPath $isa -Pattern 'v_cvt_pk_fp8_f32')) { throw 'Native FP8 operand packing is absent from ISA.' } + Write-Host "Compiled gfx1201 FP8 WMMA: $($source.BaseName)" +} +if ($HardwarePacking) { + $probeDirectory=Join-Path $OutputDirectory 'packing-probe' + New-Item -ItemType Directory -Force $probeDirectory | Out-Null + & $compiler -x hip -std=c++17 --offload-arch=gfx1201 --offload-device-only --no-gpu-bundle-output ` + -mno-wavefrontsize64 -nogpuinc -nogpulib -O3 -DD4R_DEVICE_ONLY_MINIMAL ` + -DD4R_FP8_WMMA -DD4R_FP8_HW_PACK -DSWIN_EXACT -DSWIN_EXACT_PV ` + (Join-Path $repo 'tools/windows/fp8_pack_probe.hip') -o (Join-Path $probeDirectory 'fp8_pack_probe.hsaco') + if ($LASTEXITCODE) { throw 'Native hardware packing diagnostic compile failed.' } +} +& python (Join-Path $repo 'kernels/tools/kernel_manifest.py') $OutputDirectory $DlssDll +if ($LASTEXITCODE) { throw 'Experimental FP8 manifest generation failed.' } +@{architecture='gfx1201'; family='M'; experimental=$true; nativeFp8=$true; accuracy=$true; + numerics='SWIN_EXACT; SWIN_EXACT_PV; FP8 WMMA for e4m3 GEMMs; exact FP16 attention'; + validation='not validated; do not install before layer/full-network checks'; + compiler=$compiler; compilerVersion=(& $compiler --version | Select-Object -First 1); + sourceCommit=(& git -C $repo rev-parse HEAD).Trim(); nativeFp8IsaVerified=$true; hardwarePacking=[bool]$HardwarePacking; + sourceWorkingTreeDirty=[bool](& git -C $repo status --porcelain); + dlssSha256=(Get-FileHash -LiteralPath $DlssDll -Algorithm SHA256).Hash; + modules=@(Get-ChildItem -LiteralPath $OutputDirectory -Filter '*.hsaco' | Get-FileHash -Algorithm SHA256)} | + ConvertTo-Json -Depth 5 | Set-Content -LiteralPath (Join-Path $OutputDirectory 'build-info.json') -Encoding UTF8 +Write-Host "Experimental native FP8 M objects: $OutputDirectory (baseline unchanged)" diff --git a/scripts/windows/build-native-texture.ps1 b/scripts/windows/build-native-texture.ps1 index d0d503f..55d3a06 100644 --- a/scripts/windows/build-native-texture.ps1 +++ b/scripts/windows/build-native-texture.ps1 @@ -8,7 +8,7 @@ param( ) $ErrorActionPreference='Stop' $repo=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) -if (!$ZludaRoot) { $ZludaRoot=Join-Path $repo 'dist/zluda-windows-profile' } +if (!$ZludaRoot) { $ZludaRoot=Join-Path $repo 'dist/zluda-windows-final' } if (!$OutputDirectory) { $OutputDirectory=Join-Path $repo 'build/private-textures-gfx1201' } & python (Join-Path $repo 'tools/windows/build_native_texture.py') --dlss-dll $DlssDll --hip-root $HipRoot ` --zluda-root $ZludaRoot --kernel $Kernel --output-directory $OutputDirectory diff --git a/scripts/windows/build-zluda-windows.ps1 b/scripts/windows/build-zluda-windows.ps1 index 862b271..fcc192f 100644 --- a/scripts/windows/build-zluda-windows.ps1 +++ b/scripts/windows/build-zluda-windows.ps1 @@ -4,7 +4,7 @@ $ErrorActionPreference = 'Stop' $repo = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) if (!$SourceRoot) { $SourceRoot = Join-Path $repo 'external/ZLUDA' } if (!$LlvmBuildRoot) { $LlvmBuildRoot = Join-Path $repo 'build/zluda-llvm-windows-gnu' } -if (!$InstallRoot) { $InstallRoot = Join-Path $repo 'dist/zluda-windows-native' } +if (!$InstallRoot) { $InstallRoot = Join-Path $repo 'dist/zluda-windows-final' } $rust = Join-Path $repo '.tools/rust-1.98.1-gnu' $toolchain = Join-Path $repo '.tools/llvm-mingw-20260922-ucrt-x86_64' $target = Join-Path $repo 'build/zluda-rust-windows-gnu' diff --git a/scripts/windows/package-windows-game.ps1 b/scripts/windows/package-windows-game.ps1 index 8fcace2..14e0258 100644 --- a/scripts/windows/package-windows-game.ps1 +++ b/scripts/windows/package-windows-game.ps1 @@ -7,13 +7,16 @@ param( $ErrorActionPreference = 'Stop' $repo = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) $defaults = @{PackageRoot='dist/windows-rdna4-game'; DiagnosticRoot='dist/windows-rdna4-command-list'; - OptiScalerRoot='dist/optiscaler-windows-d4r'; ZludaRoot='dist/zluda-windows-profile'; + OptiScalerRoot='dist/optiscaler-windows-d4r'; ZludaRoot='dist/zluda-windows-final'; HipRoot='.tools/therock-10.2.0a20260929/_rocm_sdk_core'; KernelKRoot='build/native-k-gfx1201'; KernelMRoot='build/native-m-gfx1201'} foreach ($name in $defaults.Keys) { if (!(Get-Variable -Name $name -ValueOnly)) { Set-Variable -Name $name -Value (Join-Path $repo $defaults[$name]) } } $PackageRoot = [IO.Path]::GetFullPath($PackageRoot) +$sourceCommit=(& git -C $repo rev-parse HEAD).Trim() +$sourceDirty=[bool](& git -C $repo status --porcelain) +if ($ArchivePath -and $sourceDirty) { throw 'Commit source changes before creating the distributable archive.' } New-Item -ItemType Directory -Force $PackageRoot | Out-Null $stagedFiles = [Collections.Generic.List[string]]::new() function Stage([string]$source, [string]$relative) { @@ -61,10 +64,15 @@ Stage (Join-Path $HipRoot 'share/doc/amd_comgr/LICENSE.txt') 'licenses/amd_comgr Stage (Join-Path $repo 'external/clr/LICENSE.md') 'licenses/HIP-CLR.txt' foreach ($patch in Get-ChildItem -LiteralPath (Join-Path $repo 'patches/optiscaler') -Filter '*.patch') { Stage $patch.FullName ('source-patches/optiscaler/' + $patch.Name) } foreach ($patch in Get-ChildItem -LiteralPath (Join-Path $repo 'patches/zluda') -Filter '*.patch') { Stage $patch.FullName ('source-patches/zluda/' + $patch.Name) } +$sourceArchive=Join-Path $repo 'build/d4r-windows-rdna4-source.zip' +New-Item -ItemType Directory -Force (Split-Path $sourceArchive) | Out-Null +& git -C $repo archive --format=zip --output $sourceArchive $sourceCommit +if ($LASTEXITCODE) { throw 'Cannot archive the committed d4r source.' } +Stage $sourceArchive 'source/d4r-windows-rdna4-source.zip' $metadata = [ordered]@{ architecture='gfx1201'; nativeKernels=16; dlssSha256='3975567b8943c53acce397f2b72380092f84f162d00b0d2c7d08a1025c563983'; - d4rCommit=(& git -C $repo rev-parse HEAD).Trim(); - d4rWorkingTreeDirty=[bool](& git -C $repo status --porcelain); + d4rCommit=$sourceCommit; + d4rWorkingTreeDirty=$sourceDirty; optiScaler=(Get-Content -LiteralPath (Join-Path $OptiScalerRoot 'build-info.json') -Raw | ConvertFrom-Json); zludaBuild=(Get-Content -LiteralPath (Join-Path $ZludaRoot 'build-info.json') -Raw | ConvertFrom-Json); hipRuntime='TheRock 10.2.0a20260929 / HIP 7.17.26386'; diff --git a/scripts/windows/profile-windows-rdna4.ps1 b/scripts/windows/profile-windows-rdna4.ps1 index c6a24d2..fc6c4d7 100644 --- a/scripts/windows/profile-windows-rdna4.ps1 +++ b/scripts/windows/profile-windows-rdna4.ps1 @@ -7,7 +7,7 @@ param( ) $ErrorActionPreference='Stop' $repo=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) -if (!$ZludaRoot) { $ZludaRoot=Join-Path $repo 'dist/zluda-windows-profile' } +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'; diff --git a/scripts/windows/test-native-m-fp8.ps1 b/scripts/windows/test-native-m-fp8.ps1 new file mode 100644 index 0000000..149e0f4 --- /dev/null +++ b/scripts/windows/test-native-m-fp8.ps1 @@ -0,0 +1,28 @@ +[CmdletBinding()] +param([Parameter(Mandatory=$true)][string]$CaptureDirectory, + [Parameter(Mandatory=$true)][string]$BaselineValidation, + [string]$ModuleDirectory, [string]$HipRoot, [string]$OutputDirectory, [switch]$Benchmark) +$ErrorActionPreference='Stop' +$repo=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) +if (!$ModuleDirectory) { $ModuleDirectory=Join-Path $repo 'build/native-m-fp8-experimental' } +if (!$HipRoot) { $HipRoot=Join-Path $repo '.tools/therock-10.2.0a20260929/_rocm_sdk_core' } +if (!$OutputDirectory) { $OutputDirectory=Join-Path $repo 'test-results/m-native-fp8-replay' } +$previousPython=$env:PYTHONPATH +try { + $env:PYTHONPATH="$(Join-Path $repo '.tools/python/vendor');$previousPython" + $metadata=Get-Content -LiteralPath (Join-Path $ModuleDirectory 'build-info.json') -Raw | ConvertFrom-Json + if ($metadata.hardwarePacking) { + & python -u (Join-Path $repo 'tools/windows/fp8_pack_validate.py') ` + --probe (Join-Path $repo 'dist/windows-rdna4-command-list/bin/d4r_native_replay_probe.exe') ` + --module (Join-Path $ModuleDirectory 'packing-probe/fp8_pack_probe.hsaco') ` + --hip-root $HipRoot --output-dir (Join-Path $OutputDirectory 'packing-check') + if ($LASTEXITCODE) { throw 'Hardware packing does not preserve every finite FP8 encoding; stopping before network replay.' } + } + $benchmarkArguments=@() + if ($Benchmark) { $benchmarkArguments=@('--benchmark-control-dir',(Join-Path $repo 'build/native-m-gfx1201')) } + & python -u (Join-Path $repo 'tools/windows/m_native_replay_compare.py') ` + --capture-dir $CaptureDirectory --module-dir $ModuleDirectory --hip-root $HipRoot ` + --probe (Join-Path $repo 'dist/windows-rdna4-command-list/bin/d4r_native_replay_probe.exe') ` + --baseline-validation $BaselineValidation --output-dir $OutputDirectory @benchmarkArguments + if ($LASTEXITCODE) { throw "Experimental FP8 validation failed; baseline unchanged. Logs: $OutputDirectory" } +} finally { $env:PYTHONPATH=$previousPython } diff --git a/scripts/windows/test-native-overrides.ps1 b/scripts/windows/test-native-overrides.ps1 index 8279696..5e8ca26 100644 --- a/scripts/windows/test-native-overrides.ps1 +++ b/scripts/windows/test-native-overrides.ps1 @@ -4,7 +4,7 @@ $ErrorActionPreference = 'Stop' $repo = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) if (!$HipRoot) { $HipRoot = Join-Path $repo '.tools/therock-10.2.0a20260929/_rocm_sdk_core' } if (!$PackageRoot) { $PackageRoot = Join-Path $repo 'dist/windows-rdna4-therock' } -if (!$ZludaRoot) { $ZludaRoot = Join-Path $repo 'dist/zluda-windows-native' } +if (!$ZludaRoot) { $ZludaRoot = Join-Path $repo 'dist/zluda-windows-final' } if (!$OutputDirectory) { $OutputDirectory = Join-Path $repo ('test-results/native-identity-' + (Get-Date -Format 'yyyyMMdd-HHmmss')) } New-Item -ItemType Directory -Force $OutputDirectory | Out-Null $previous = @{D4R_ZLUDA_NATIVE_DIR=$env:D4R_ZLUDA_NATIVE_DIR; D4R_ZLUDA_VERBOSE=$env:D4R_ZLUDA_VERBOSE; D4R_DIAG_DIR=$env:D4R_DIAG_DIR} diff --git a/scripts/windows/test-native-texture.ps1 b/scripts/windows/test-native-texture.ps1 index 58e3e69..6899023 100644 --- a/scripts/windows/test-native-texture.ps1 +++ b/scripts/windows/test-native-texture.ps1 @@ -22,7 +22,7 @@ try { foreach ($flags in @(0,11)) { foreach ($mode in @('control','candidate')) { $env:D4R_ZLUDA_NATIVE_DIR=if ($mode -eq 'control') { Join-Path $repo 'build/native-k-gfx1201' } else { $combined } - $arguments=@{RuntimeProfile='therock'; ZludaRoot=(Join-Path $repo 'dist/zluda-windows-profile'); + $arguments=@{RuntimeProfile='therock'; ZludaRoot=(Join-Path $repo 'dist/zluda-windows-final'); PackageRoot=(Join-Path $repo 'dist/windows-rdna4-command-list'); NgxCore=(Join-Path $repo '_nvngx.dll'); DlssDll=(Join-Path $repo 'nvngx_dlss.dll'); NgxMode='d3d12'; NgxAbi='project-legacy'; Preset=11; NgxCreateFlags=$flags; NgxOutputResolution=$OutputResolution; NgxOnly=$true; RequireNativeNetwork=$true; CommandListBackend=$true; diff --git a/scripts/windows/test-windows-rdna4.ps1 b/scripts/windows/test-windows-rdna4.ps1 index da2ed37..e65bf40 100644 --- a/scripts/windows/test-windows-rdna4.ps1 +++ b/scripts/windows/test-windows-rdna4.ps1 @@ -42,7 +42,7 @@ $PackageRoot = [IO.Path]::GetFullPath($PackageRoot) if (!$HipRoot) { $HipRoot = 'C:\Program Files\AMD\ROCm\7.2' } $HipRoot = [IO.Path]::GetFullPath($HipRoot) if (!$ZludaRoot) { - $builtZluda = Join-Path $repoRoot 'dist/zluda-windows-native' + $builtZluda = Join-Path $repoRoot 'dist/zluda-windows-final' $localZluda = [IO.Path]::GetFullPath((Join-Path $PackageRoot '../../.tools/zluda/zluda')) if (Test-Path (Join-Path $builtZluda 'build-info.json')) { $ZludaRoot = $builtZluda } elseif (Test-Path (Join-Path $localZluda 'nvcuda.dll')) { $ZludaRoot = $localZluda } diff --git a/tools/windows/diagnostic.h b/tools/windows/diagnostic.h index 6161a9f..19213b0 100644 --- a/tools/windows/diagnostic.h +++ b/tools/windows/diagnostic.h @@ -128,6 +128,7 @@ struct Args { unsigned ngx_create_flags = 0; unsigned output_width = 512, output_height = 288; std::string fixture_dir; + std::string benchmark_module; std::string output_dir; std::string kernel_name = "enc1"; unsigned iterations = 32; @@ -184,6 +185,7 @@ struct Args { throw std::runtime_error("preset must be 5 (E), 11 (K) or 13 (M)"); } else if (key == "--fixture-dir") fixture_dir = value; + else if (key == "--benchmark-module") benchmark_module = value; else if (key == "--output-dir") output_dir = value; else if (key == "--kernel-name") kernel_name = value; else if (key == "--iterations") { diff --git a/tools/windows/fp8_pack_probe.hip b/tools/windows/fp8_pack_probe.hip new file mode 100644 index 0000000..994ca48 --- /dev/null +++ b/tools/windows/fp8_pack_probe.hip @@ -0,0 +1,9 @@ +// Diagnostic only: all finite OCP e4m3 encodings, including signed zero. +#include "../../kernels/m/swin_common.h" +struct PackParams { const half_t* input; uint32_t* output; }; +extern "C" __global__ void fp8_pack_probe(PackParams p) { + const unsigned lane = blockIdx.x * 32 + threadIdx.x; + if (lane >= 32) return; + const auto packed = to_fp8(*(const sop*)(p.input + lane * 8)); + p.output[lane * 2] = packed[0]; p.output[lane * 2 + 1] = packed[1]; +} diff --git a/tools/windows/fp8_pack_validate.py b/tools/windows/fp8_pack_validate.py new file mode 100644 index 0000000..41f2e10 --- /dev/null +++ b/tools/windows/fp8_pack_validate.py @@ -0,0 +1,46 @@ +"""Check gfx1201 hardware packing of every finite OCP e4m3 value and signed zero.""" +import argparse +import hashlib +import json +import os +import pathlib +import struct +import subprocess + +import numpy as np +from m_replay_validate import me + + +def main(args): + root = args.output_dir.resolve() + root.mkdir(parents=True, exist_ok=True) + codes = np.arange(256, dtype=np.uint8) + codes[(codes & 127) == 127] = 0 # NaN inputs are outside this finite-value packing contract. + inputs = me.E4M3[codes].astype(np.float16).tobytes() + (root / 'alloc-0.bin').write_bytes(inputs) + (root / 'alloc-1.bin').write_bytes(bytes(256)) + (root / 'args.bin').write_bytes(struct.pack(' initial; Allocation(d4r::diag::HipApi& api, int identifier, const std::vector& bytes) - : hip(api), id(identifier), size(bytes.size()) { + : hip(api), id(identifier), size(bytes.size()), initial(bytes) { hip.check(hip.hipMalloc(&device, size), "hipMalloc(replay)"); try { hip.check(hip.hipMemcpy(device, bytes.data(), size, hipMemcpyHostToDevice), "hipMemcpy(replay upload)"); } catch (...) { (void)hip.hipFree(device); device = nullptr; throw; } @@ -90,14 +91,15 @@ int main(int argc, char** argv) { if (size != sizeof(coordinates)) throw std::runtime_error("Swin diagnostic coordinates ABI"); hip.check(hip.hipMemcpy(address, coordinates, size, hipMemcpyHostToDevice), "Set Swin diagnostic block"); } - auto scalar = [&](const char* name, unsigned fallback) { + auto module_scalar = [&](hipModule_t owner, const char* name, unsigned fallback) { hipDeviceptr_t address = nullptr; size_t size = 0; - if (hip.hipModuleGetGlobal(&address, &size, module, name) != hipSuccess) return fallback; + if (hip.hipModuleGetGlobal(&address, &size, owner, name) != hipSuccess) return fallback; if (size != 4) throw std::runtime_error(std::string("Invalid native metadata ") + name); unsigned value = 0; hip.check(hip.hipMemcpy(&value, address, 4, hipMemcpyDeviceToHost), name); return value ? value : fallback; }; + auto scalar = [&](const char* name, unsigned fallback) { return module_scalar(module, name, fallback); }; 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; } @@ -111,7 +113,65 @@ int main(int argc, char** argv) { } std::printf("REPLAY kernel=%s grid=%u,%u,%u block=%u,%u,%u allocations=%zu arguments=%zu\n", kernel.c_str(), grid[0], grid[1], grid[2], block[0], block[1], block[2], allocations.size(), parameters.size()); - for (unsigned iteration = 0; iteration < args.iterations; ++iteration) + if (!args.benchmark_module.empty()) { + 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"); + 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) { + const auto blocks = module_scalar(candidate, "d4r_prep_blocks", 0); + if (!blocks) throw std::runtime_error("Paired candidate prep metadata is missing"); + hip.check(hip.hipModuleLaunchKernel(other_prep, blocks, 1, 1, 128, 1, 1, 0, nullptr, + launchArguments, nullptr), "Paired candidate prep"); + } + auto create = hip.library.symbol("hipEventCreate"); + auto destroy = hip.library.symbol("hipEventDestroy"); + auto record = hip.library.symbol("hipEventRecord"); + auto synchronize = hip.library.symbol("hipEventSynchronize"); + auto elapsed = hip.library.symbol("hipEventElapsedTime"); + struct Event { + decltype(destroy) release; hipEvent_t value = nullptr; + ~Event() { if (value) (void)release(value); } + } begin{destroy}, end{destroy}; + 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], + shared, nullptr, launchArguments, nullptr), "Paired benchmark launch"); + }; + // Both modules and weight images stay resident on the same GPU. + // Warm continuous work first, then alternate AB/BA order. Restore + // captured inputs outside the event so neither module consumes + // modified temporal inputs from a preceding sample. + const auto verbose = hip.verbose; hip.verbose = false; + for (unsigned warm = 0; warm < 16; ++warm) + for (auto fn : {function, other}) for (unsigned batch = 0; batch < 16; ++batch) launch(fn); + hip.check(hip.hipDeviceSynchronize(), "Paired benchmark warmup"); + for (unsigned pair = 0; pair < args.iterations; ++pair) for (unsigned order = 0; order < 2; ++order) { + const unsigned variant = order ^ (pair & 1u); + 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"); + hip.check(synchronize(end.value), "Complete paired sample"); + 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)); + } + hip.verbose = verbose; + // Leave a single ordinary control replay for numerical checking. + for (const auto& allocation : allocations) + hip.check(hip.hipMemcpy(allocation->device, allocation->initial.data(), allocation->size, + hipMemcpyHostToDevice), "Restore final reference input"); + launch(function); + } + for (unsigned iteration = 0; args.benchmark_module.empty() && iteration < args.iterations; ++iteration) hip.check(hip.hipModuleLaunchKernel(function, grid[0], grid[1], grid[2], block[0], block[1], block[2], shared, nullptr, launchArguments, nullptr), "hipModuleLaunchKernel(replay)"); hip.check(hip.hipDeviceSynchronize(), "hipDeviceSynchronize(replay)"); @@ -132,7 +192,8 @@ int main(int argc, char** argv) { file.write(reinterpret_cast(bytes.data()), std::streamsize(bytes.size())); if (!file) throw std::runtime_error("Cannot save replay allocation"); } - std::printf("PASS NATIVE_REPLAY architecture=gfx1201 kernel=%s iterations=%u numerical_reference=pending\n", kernel.c_str(), args.iterations); + std::printf("PASS NATIVE_REPLAY architecture=gfx1201 kernel=%s iterations=%u benchmark_pairs=%u numerical_reference=pending\n", + kernel.c_str(), args.benchmark_module.empty() ? args.iterations : 1, args.benchmark_module.empty() ? 0 : args.iterations); return 0; } catch (const std::exception& error) { std::fprintf(stderr, "FAIL NATIVE_REPLAY %s\n", error.what()); loaded_modules(); return 4; diff --git a/tools/windows/profile_report.py b/tools/windows/profile_report.py index f01f8b0..a88beef 100644 --- a/tools/windows/profile_report.py +++ b/tools/windows/profile_report.py @@ -42,8 +42,13 @@ def report(directory, metadata_directory=None): commands = collections.defaultdict(list) kernels = collections.defaultdict(list) hook_threads = collections.defaultdict(list) + replay = collections.defaultdict(lambda: collections.defaultdict(dict)) for path in sorted(directory.glob('*.log')): - text = path.read_text(encoding='utf-8-sig', errors='replace') + 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 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): @@ -64,6 +69,7 @@ def report(directory, metadata_directory=None): result = dict(cpuStages={name: stats(values) for name, values in stages.items()}, commandStages={name: stats(values) for name, values in commands.items()}, commandHookSamples=[], + replayPairs=[], kernels=[], notes=[ 'CPU stage times include waits and host work; they are not isolated GPU timings.', 'Command record intervals are between NGX recording calls on one thread, not Present/FPS.', @@ -72,6 +78,13 @@ def report(directory, metadata_directory=None): 'HIP-event profiling synchronizes every sampled launch and changes scheduling.', 'Occupancy is an API prediction, not a measured hardware counter.', 'Bandwidth and WMMA utilization require additional supported hardware tooling.']) + for name, 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]))) for thread, rows in hook_threads.items(): sums = {name: sum(float(row[name]) for row in rows) for name in ('period_ms', 'calls', 'samples', 'marked', 'access_sum_ms', 'capture_sum_ms', 'driver_sum_ms')} @@ -105,7 +118,7 @@ def main(): parser.add_argument('--metadata-directory', type=pathlib.Path, help='llvm-readobj --notes output for native code objects') args = parser.parse_args() result = report(args.directory, args.metadata_directory) - if not result['cpuStages'] and not result['kernels']: + if not any(result[key] for key in ('cpuStages', 'kernels', 'commandStages', 'commandHookSamples', 'replayPairs')): parser.error('no complete stage or HIP-event records found') if args.output: args.output.write_text(json.dumps(result, indent=2) + '\n', encoding='utf-8') @@ -116,6 +129,9 @@ def main(): for row in result['commandHookSamples']: 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 " + 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")