From 2db93adb1ac43b4232c7cd4f638da148c6211d84 Mon Sep 17 00:00:00 2001 From: Codex Date: Wed, 30 Sep 2026 19:01:18 +0300 Subject: [PATCH] Validate native Windows M against real PTX and correct patch-merge bounds --- CMakeLists.txt | 9 + docs/windows-rdna4-port.md | 66 ++++++- kernels/m/swin_block.h | 11 +- kernels/tools/ptxsim.py | 22 ++- kernels/tools/swin_check.py | 10 +- .../0011-audited-k-m-windows-replay.patch | 93 +++++++++ scripts/windows/stage-native-m.ps1 | 24 +++ scripts/windows/test-m-captures.ps1 | 49 +++++ scripts/windows/test-m-replay.ps1 | 45 +++++ scripts/windows/test-windows-rdna4.ps1 | 13 +- tools/windows/m_ptx_validate.py | 62 ++++++ tools/windows/m_replay_validate.py | 178 ++++++++++++++++++ tools/windows/test_ptx_reference.py | 46 +++++ 13 files changed, 611 insertions(+), 17 deletions(-) create mode 100644 patches/zluda/0011-audited-k-m-windows-replay.patch create mode 100644 scripts/windows/stage-native-m.ps1 create mode 100644 scripts/windows/test-m-captures.ps1 create mode 100644 scripts/windows/test-m-replay.ps1 create mode 100644 tools/windows/m_ptx_validate.py create mode 100644 tools/windows/m_replay_validate.py create mode 100644 tools/windows/test_ptx_reference.py diff --git a/CMakeLists.txt b/CMakeLists.txt index 676e738..8a4862d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -35,6 +35,7 @@ add_custom_target(d4r_wmma_code_object ALL DEPENDS "${wmma_object}") # M uses the existing FP16 widening baseline; native FP8 is not enabled here. set(native_objects) option(D4R_K_FP16_BASELINE "Round every K16 accumulation step like NVIDIA f16 WMMA" ON) +option(D4R_M_FP16_BASELINE "Preserve M e4m3 quantisation and f16 accumulation per K32" ON) foreach(family k m) file(GLOB native_sources CONFIGURE_DEPENDS "${CMAKE_CURRENT_SOURCE_DIR}/kernels/${family}/*.hip") file(GLOB native_headers CONFIGURE_DEPENDS "${CMAKE_CURRENT_SOURCE_DIR}/kernels/${family}/*.h") @@ -48,6 +49,9 @@ foreach(family k m) if(family STREQUAL "k" AND D4R_K_FP16_BASELINE) list(APPEND native_flags -DD4R_K_FP16_BASELINE) endif() + if(family STREQUAL "m" AND D4R_M_FP16_BASELINE) + list(APPEND native_flags -DSWIN_EXACT) + endif() if(flags_line) separate_arguments(source_flags NATIVE_COMMAND "${CMAKE_MATCH_1}") list(APPEND native_flags ${source_flags}) @@ -127,8 +131,13 @@ install(FILES tools/windows/k_layer_validate.py DESTINATION .) install(FILES tools/windows/k_replay_validate.py DESTINATION .) install(FILES tools/windows/k_ptx_validate.py DESTINATION .) install(FILES tools/windows/frame_compare.py DESTINATION .) +install(FILES tools/windows/m_replay_validate.py DESTINATION .) +install(FILES tools/windows/m_ptx_validate.py DESTINATION .) +install(FILES scripts/windows/test-m-replay.ps1 DESTINATION .) +install(FILES scripts/windows/test-m-captures.ps1 DESTINATION .) install(FILES kernels/tools/pwin_model.py kernels/tools/pwin_check.py kernels/tools/ptxsim.py kernels/tools/extract_dlss_ptx.py kernels/tools/kernel_manifest.py DESTINATION reference) +install(FILES kernels/tools/model_enc3.py kernels/tools/swin_model.py kernels/tools/swin_check.py DESTINATION reference) install(FILES docs/windows-rdna4-port.md DESTINATION docs) include(CTest) add_test(NAME hip_gfx1201 COMMAND d4r_hip_gfx1201_probe diff --git a/docs/windows-rdna4-port.md b/docs/windows-rdna4-port.md index b24eba4..17af852 100644 --- a/docs/windows-rdna4-port.md +++ b/docs/windows-rdna4-port.md @@ -9,7 +9,9 @@ This is a work log for an implementation in progress, not a claim of DLSS suppor Current milestone: M1/M2/M4/M5 passed; native NGX Create/Evaluate on the simple requested E path passes with the user's DLLs. All 11 K layers now pass synthetic and real-weight NumPy/PTX replay checks; native K Evaluate produces four finite -frames on Windows. D3D12 shim integration and M remain in progress. +frames on Windows. M now passes all five synthetic and real-weight layers against +NumPy and independent PTX, and four native CUDA frames. D3D12 integration remains +in progress; the fully translated M comparison currently exceeds its JIT timeout. Do not integrate NGX until the integer PTX workload is stable on the real GPU. | Gate | Status | @@ -21,7 +23,7 @@ Do not integrate NGX until the integer PTX workload is stable on the real GPU. | M4: D3D12 / HIP external memory and fence round trip | PASS with TheRock; stable 7.2 has mapped-view leak | | M5: independently validated gfx12 WMMA backend | PASS: raw + legacy adapter + upstream layout, max abs/relative error 0 | | M6: K layers, full transformer and image validation | CUDA harness PASS: 11 real-weight layers, native launches for all layers, four finite frames, 81.85–92.91 dB versus translated K; D3D12 integration pending | -| M7: M FP16-equivalent baseline and full transformer | Pending | +| M7: M FP16-equivalent baseline and full transformer | CUDA harness PASS: 5 real-weight layers exactly match NumPy/PTX, 4 finite frames, 40 native Swin launches; D3D12 and full translated M comparison pending | | M8: standalone Windows D3D12 NGX harness | Pending | | M9: OptiScaler integration, profiling, installation | Pending | @@ -124,6 +126,66 @@ FP8/performance work. No complete DLSS or OptiScaler game result is claimed. ## Native Windows ZLUDA build and Evaluate work (2026-09-30) +### M strict FP16 baseline and independent reference (2026-09-30) + +`D4R_M_FP16_BASELINE=ON` is the Windows CMake default and defines `SWIN_EXACT`. +Activations retain e4m3 requantisation; matrix accumulation rounds to f16 after +each K32 step. Gfx1201 WMMA widens FP8 operands to FP16. Native FP8 is disabled +until this baseline also passes the D3D12 harness. + +All five `rrlite_{enc1,enc2,enc3_tube,dec2,dec1}_4x4` modules pass 32 iterations +on public synthetic 16x16 fixtures, including nonzero Q/V/WO/MLP weights, +decoder expansion/skip and patch merge. Every stored e4m3 code exactly matches +the NumPy model; max absolute and relative error are zero. The fixture's patch +merge deliberately uses large weights to exercise its output bound. + +An independent CPU PTX check exposed an upstream omission: DLSS 310.9's patch +merge clamps f16 results to +/- f16(2*pi), before e4m3 rounding. Without this +bound, real `enc2` merge outputs reached +/-9 or 10 where the PTX stored +/-6.5. +The correction is limited to `SWIN_EXACT`; upstream fast arithmetic stays +unchanged. The Windows model enables the same audited bound explicitly; the +legacy Linux oracle remains the default. Both implementations now agree with +the independent PTX interpreter, rather than merely agreeing with each other. + +`0011-audited-k-m-windows-replay.patch` adds exact M capture ABI sizes (56/88 +bytes) and restricts K captures to the eleven audited names. The public CPU PTX +interpreter now handles divergent finite backward loops by suspending completed +lanes at the loop fallthrough. `tools/windows/test_ptx_reference.py` verifies +lane-specific loop counts and reconvergence with a public synthetic program. + +With the user's 310.9 DLL, `test-results/ngx-m-bounded-capture.zip` passes four +frames with native counts enc1=4, enc2=4, enc3_tube=24, dec2=4, dec1=4. No NaN/Inf +or missing transformer is accepted. `test-results/m-real-final-dual-reference` +passes 12 NumPy blocks and two independent CPU PTX blocks per layer; all sampled +codes agree exactly (max abs/relative error 0, PSNR infinite). Capture files, +weights and extracted PTX remain private and must not be redistributed. + +Reproduce the layer gates: + +```powershell +& scripts/windows/test-m-replay.ps1 +& scripts/windows/stage-native-m.ps1 -DlssDll "$PWD/nvngx_dlss.dll" +$env:D4R_ZLUDA_NATIVE_DIR="$PWD/build/native-m-gfx1201" +$env:D4R_WMMA='1'; $env:D4R_FP8_WMMA='1'; $env:D4R_FP8_NATIVE='0' +& scripts/windows/test-windows-rdna4.ps1 -RuntimeProfile therock -NgxOnly -NgxMode evaluate -Preset 13 -Iterations 4 -RequireNativeNetwork -NgxCore "$PWD/_nvngx.dll" -DlssDll "$PWD/nvngx_dlss.dll" +``` + +The fully translated M feature creation currently exceeds 600 seconds; its +automatic trace localises this to loading the transformer modules. Native M +does not have this JIT delay. A completed translated image comparison is not +claimed. The user selected Silent Hill 2 at +`D:\Games\SILENT HILL 2\SHProto\Binaries\Win64\SHProto-Win64-Shipping.exe` +for the later game gate. No game files have been changed yet. + +OptiScaler upstream inspected at `45a2001303ddff632e279f77aef85ceede5832cb`, +2026-09-30. Its DLSS backend forwards an open command list into Evaluate; +its swapchain hooks retain the game queue. A correct same-frame Windows path +must preserve submission order when earlier command lists have not yet been +submitted. Closing/submitting the current list immediately inside Evaluate +is insufficient for a general game integration. The next backend gate is a +D3D12 harness with explicit queue boundaries; command-list/queue integration +must be validated before the game gate. + ### K real-weight native Evaluate and correctness baseline The locally supplied DLSS 310.9.1.0 selects all eleven `dltss_pwin_*` functions diff --git a/kernels/m/swin_block.h b/kernels/m/swin_block.h index d114a29..148f0a6 100644 --- a/kernels/m/swin_block.h +++ b/kernels/m/swin_block.h @@ -718,7 +718,16 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara if (MX < 0 || MX >= W2 || MY < 0 || MY >= H2) continue; p.p48[((size_t)((ch >> 5) * H2 + MY) * W2 + MX) * 32 + (ch & 31)] = - (uint8_t)enc8(y[j].get(i)); + (uint8_t)enc8( +#ifdef SWIN_EXACT + // The audited 310.9 patch-merge path bounds the f16 + // result to +/- 2*pi before e4m3 quantisation. + __builtin_elementwise_max(__builtin_elementwise_min(y[j].get(i), (half_t)6.283185307179586f), + (half_t)-6.283185307179586f) +#else + y[j].get(i) +#endif + ); } } } diff --git a/kernels/tools/ptxsim.py b/kernels/tools/ptxsim.py index 478a95b..2ea05da 100644 --- a/kernels/tools/ptxsim.py +++ b/kernels/tools/ptxsim.py @@ -96,7 +96,8 @@ def split_args(s): def parse_kernel(path, name): - src = open(path).read() + with open(path) as file: + src = file.read() i = src.index('.entry ' + name + '(') j = src.find('\n.visible', i + 10) j = len(src) if j < 0 else j @@ -340,12 +341,11 @@ class Block: steps = 0 while pc < len(stmts): st = stmts[pc] - for l in st.labels: - if l in pending: - active |= pending.pop(l) + if pc in pending: + active |= pending.pop(pc) if not active.any(): - # jump to the nearest label someone waits at - targets = sorted(self.labels[l] for l in pending) + # Reconverge at either a forward target or a loop fallthrough. + targets = sorted(pending) if not targets: break pc = targets[0] @@ -365,10 +365,16 @@ class Block: if not taking.any(): pc += 1 continue - raise NotImplementedError('divergent backward branch at ' + st.text) + # Lanes finishing a loop wait at its fallthrough while the + # other lanes execute their remaining iterations. Their + # registers and memory writes stay masked until reconvergence. + pending[pc + 1] = pending.get(pc + 1, np.zeros(self.T, bool)) | (active & ~taking) + active = taking.copy() + pc = target + continue active = active & ~taking if taking.any(): - pending[st.args[0]] = pending.get(st.args[0], np.zeros(self.T, bool)) | taking + pending[target] = pending.get(target, np.zeros(self.T, bool)) | taking pc += 1 continue if mask.any() or st.op.startswith(('mma', 'shfl', 'movmatrix', 'bar')): diff --git a/kernels/tools/swin_check.py b/kernels/tools/swin_check.py index 2d45bde..fa389c6 100644 --- a/kernels/tools/swin_check.py +++ b/kernels/tools/swin_check.py @@ -23,6 +23,10 @@ import swin_model as sm # noqa: E402 LAYERS = {'enc1': (64, 2, 96, 0), 'enc2': (96, 4, 128, 0), 'enc3_tube': (128, 4, 0, 0), 'dec2': (96, 4, 0, 128), 'dec1': (64, 2, 0, 96)} +# Preserve the existing Linux oracle by default; Windows strict replay enables +# the audited 310.9 f16 patch-merge bound before final e4m3 quantisation. +CLAMP_PATCH_MERGE = False + def blocks(grid, n, seed=1): gx, gy = grid[0], grid[1] @@ -47,7 +51,11 @@ def model_block(P, layer, bx, by): Wg = me.E4M3[P.wbuf[sm.woff_table(PM0 + g * 128 * C, 512, 64 * C, 4 * C, 32)]] bias = P.f16vec(PM0 + 4 * C * NPM + 64 * g, 32) cols.append(me.mma_chain(A, Wg, C0=np.broadcast_to(bias, (16, 32)).copy())) - merged = me.q8(sm.pair_to_nat(np.concatenate(cols, axis=1))) + merged = sm.pair_to_nat(np.concatenate(cols, axis=1)) + if CLAMP_PATCH_MERGE: + limit = me.f16(2 * np.pi) + merged = np.clip(merged, -limit, limit) + merged = me.q8(merged) return out, merged diff --git a/patches/zluda/0011-audited-k-m-windows-replay.patch b/patches/zluda/0011-audited-k-m-windows-replay.patch new file mode 100644 index 0000000..0742db1 --- /dev/null +++ b/patches/zluda/0011-audited-k-m-windows-replay.patch @@ -0,0 +1,93 @@ +diff --git a/zluda/src/impl/d4r_replay.rs b/zluda/src/impl/d4r_replay.rs +index 875b2e3..d491ae2 100644 +--- a/zluda/src/impl/d4r_replay.rs ++++ b/zluda/src/impl/d4r_replay.rs +@@ -7,6 +7,18 @@ struct Config { directory: PathBuf, filter: String, limit: u32 } + static CONFIG: OnceLock> = OnceLock::new(); + static COUNTS: OnceLock>> = OnceLock::new(); + ++fn parameter_size(name: &str) -> Option { ++ match name { ++ "dltss_pwin_enc0_layer" | "dltss_pwin_enc1_layer" | "dltss_pwin_enc2_layer" | ++ "dltss_pwin_enc3_layer" | "dltss_pwin_enc4_layer" | "dltss_pwin_dec5_layer" | ++ "dltss_pwin_dec4_layer" | "dltss_pwin_dec3_layer" | "dltss_pwin_dec2_layer" | ++ "dltss_pwin_dec1_layer" | "dltss_pwin_dec0_layer" => Some(176), ++ "rrlite_enc1_4x4" | "rrlite_enc2_4x4" | "rrlite_dec2_4x4" | "rrlite_dec1_4x4" => Some(56), ++ "rrlite_enc3_tube_4x4" => Some(88), ++ _ => None, ++ } ++} ++ + pub(crate) struct Capture { directory: PathBuf, allocations: Vec<(usize, usize)> } + + fn io_error(error: impl std::fmt::Display) -> hipErrorCode_t { +@@ -43,8 +55,10 @@ pub(crate) unsafe fn begin( + limit: std::env::var("D4R_CUDA_REPLAY_DUMP_LIMIT").ok().and_then(|s| s.parse().ok()).unwrap_or(1), + })) else { return Ok(None) }; + let name = name.to_str().map_err(io_error)?; +- // Only the audited, single 176-byte K argument ABI is captured here. +- if !name.starts_with("dltss_pwin_") || !name.ends_with("_layer") || !name.contains(&config.filter) { ++ // Only audited single-by-value K/M argument blocks. Never infer the size ++ // from a prefix: reading a different kernel's parameters can overrun host memory. ++ let Some(parameter_bytes) = parameter_size(name) else { return Ok(None); }; ++ if !name.contains(&config.filter) { + return Ok(None); + } + let mut block = std::ptr::null::(); +@@ -60,10 +74,10 @@ pub(crate) unsafe fn begin( + } + } else if !parameters.is_null() { + block = (*parameters).cast(); +- bytes = Some(176); ++ bytes = Some(parameter_bytes); + } +- if block.is_null() || bytes != Some(176) { +- return Err(io_error(format!("unexpected K parameter ABI for {name}: {bytes:?}"))); ++ if block.is_null() || bytes != Some(parameter_bytes) { ++ return Err(io_error(format!("unexpected parameter ABI for {name}: {bytes:?}, expected {parameter_bytes}"))); + } + let sequence = { + let mut counts = COUNTS.get_or_init(|| Mutex::new(HashMap::new())).lock().map_err(io_error)?; +@@ -73,14 +87,14 @@ pub(crate) unsafe fn begin( + *count + }; + hipStreamSynchronize(stream)?; +- let arguments = std::slice::from_raw_parts(block, 176); ++ let arguments = std::slice::from_raw_parts(block, parameter_bytes); + let directory = config.directory.join(format!("replay-{sequence:06}-{name}")); + fs::create_dir_all(&directory).map_err(io_error)?; + fs::write(directory.join("args.bin"), arguments).map_err(io_error)?; +- let mut manifest = format!("kernel {name}\nlaunch {} {} {} {} {} {} {}\nargs 176\n", ++ let mut manifest = format!("kernel {name}\nlaunch {} {} {} {} {} {} {}\nargs {parameter_bytes}\n", + launch[0], launch[1], launch[2], launch[3], launch[4], launch[5], launch[6]); + let mut allocations = Vec::<(usize, usize)>::new(); +- for offset in (0..176).step_by(8) { ++ for offset in (0..parameter_bytes).step_by(8) { + let pointer = u64::from_le_bytes(arguments[offset..offset + 8].try_into().unwrap()) as usize; + if pointer == 0 { continue; } + let mut base = hipDeviceptr_t(std::ptr::null_mut()); +@@ -100,10 +114,23 @@ pub(crate) unsafe fn begin( + }; + manifest.push_str(&format!("pointer {offset} {id} {}\n", pointer - base)); + } +- if allocations.is_empty() { return Err(io_error("no K allocation pointers resolved")); } ++ if allocations.is_empty() { return Err(io_error("no allocation pointers resolved")); } + fs::write(directory.join("manifest.txt"), manifest).map_err(io_error)?; + let capture = Capture { directory, allocations }; + capture.save(&capture.directory)?; + eprintln!("[d4r-replay] begin kernel={name} allocations={} path={}", capture.allocations.len(), capture.directory.display()); + Ok(Some(capture)) + } ++ ++#[cfg(test)] ++mod tests { ++ #[test] ++ fn only_audited_argument_sizes() { ++ assert_eq!(super::parameter_size("dltss_pwin_enc0_layer"), Some(176)); ++ assert_eq!(super::parameter_size("rrlite_dec2_4x4"), Some(56)); ++ assert_eq!(super::parameter_size("rrlite_enc3_tube_4x4"), Some(88)); ++ assert_eq!(super::parameter_size("dltss_pwin_enc0_layer_prep"), None); ++ assert_eq!(super::parameter_size("rrlite_enc3_tube_4x4_prep"), None); ++ assert_eq!(super::parameter_size("dltss_pwin_unknown_layer"), None); ++ } ++} diff --git a/scripts/windows/stage-native-m.ps1 b/scripts/windows/stage-native-m.ps1 new file mode 100644 index 0000000..2f75986 --- /dev/null +++ b/scripts/windows/stage-native-m.ps1 @@ -0,0 +1,24 @@ +[CmdletBinding()] +param([Parameter(Mandatory=$true)][string]$DlssDll, [string]$PackageRoot, [string]$OutputDirectory) +$ErrorActionPreference = 'Stop' +$repo = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) +if (!$PackageRoot) { $PackageRoot = Join-Path $repo 'dist/windows-rdna4-therock' } +if (!$OutputDirectory) { $OutputDirectory = Join-Path $repo 'build/native-m-gfx1201' } +if (![IO.Path]::IsPathRooted($DlssDll) -or !(Test-Path -LiteralPath $DlssDll)) { + throw 'Supply the absolute path to your local nvngx_dlss.dll.' +} +New-Item -ItemType Directory -Force $OutputDirectory | Out-Null +foreach ($layer in @('enc1','enc2','enc3_tube','dec2','dec1')) { + $name = "rrlite_${layer}_4x4" + $source = Join-Path $PackageRoot "experimental/m/${name}_gfx1201.hsaco" + if (!(Test-Path -LiteralPath $source)) { throw "Native M module missing: $source" } + Copy-Item -LiteralPath $source -Destination (Join-Path $OutputDirectory "$name.hsaco") -Force +} +& python (Join-Path $repo 'kernels/tools/kernel_manifest.py') $OutputDirectory $DlssDll +if ($LASTEXITCODE) { throw "Native manifest generation failed ($LASTEXITCODE)" } +@{architecture='gfx1201'; family='M'; numerics='SWIN_EXACT; FP16 widening; native FP8 disabled'; + dlssSha256=(Get-FileHash -LiteralPath $DlssDll -Algorithm SHA256).Hash; + validation='synthetic layers passed; actual DLSS capture/reference validation required'; + 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 M directory: $OutputDirectory" diff --git a/scripts/windows/test-m-captures.ps1 b/scripts/windows/test-m-captures.ps1 new file mode 100644 index 0000000..1396c7e --- /dev/null +++ b/scripts/windows/test-m-captures.ps1 @@ -0,0 +1,49 @@ +[CmdletBinding()] +param([Parameter(Mandatory=$true)][string]$CaptureDirectory, [string]$HipRoot, + [string]$PackageRoot, [string]$OutputDirectory, [int]$Blocks = 12, + [Parameter(Mandatory=$true)][string]$DlssDll, [int]$PtxBlocks = 2) +$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 (!$OutputDirectory) { $OutputDirectory = Join-Path $repo ('test-results/m-real-reference-' + (Get-Date -Format 'yyyyMMdd-HHmmss')) } +New-Item -ItemType Directory -Force $OutputDirectory | Out-Null +$previousPython=$env:PYTHONPATH; $previousDiag=$env:D4R_DIAG_DIR +$env:PYTHONPATH = "$(Join-Path $repo '.tools/python/vendor');$previousPython" +$summary = [ordered]@{utc=[DateTime]::UtcNow.ToString('o'); captureDirectory=$CaptureDirectory; layers=@(); passed=$false} +try { + foreach ($layer in @('enc1','enc2','enc3_tube','dec2','dec1')) { + $name = "rrlite_${layer}_4x4" + $captures = @(Get-ChildItem -LiteralPath $CaptureDirectory -Directory | Where-Object Name -Match ("^replay-[0-9]+-" + $name + '$')) + if (!$captures.Count) { throw "Missing actual M capture: $name" } + $capture = $captures | Sort-Object Name | Select-Object -First 1 + $folder = Join-Path $OutputDirectory $layer + $output = Join-Path $folder 'output' + New-Item -ItemType Directory -Force $folder | Out-Null + $env:D4R_DIAG_DIR = $folder + & (Join-Path $PackageRoot 'bin/d4r_native_replay_probe.exe') --hip-root $HipRoot ` + --module (Join-Path $PackageRoot "experimental/m/${name}_gfx1201.hsaco") ` + --fixture-dir $capture.FullName --output-dir $output --iterations 1 ` + > (Join-Path $folder 'gpu.stdout.log') 2> (Join-Path $folder 'gpu.stderr.log') + if ($LASTEXITCODE) { throw "M $layer actual GPU replay failed: $folder" } + & python (Join-Path $PackageRoot 'm_replay_validate.py') validate --layer $layer ` + --fixture-dir $capture.FullName --output-dir $output --blocks $Blocks --real-capture ` + > (Join-Path $folder 'reference.stdout.log') 2> (Join-Path $folder 'reference.stderr.log') + if ($LASTEXITCODE) { throw "M $layer mismatch; stopping before the next layer: $folder" } + Get-Content -LiteralPath (Join-Path $folder 'reference.stdout.log') | Write-Host + & python -u (Join-Path $PackageRoot 'm_ptx_validate.py') --layer $layer ` + --capture-dir $capture.FullName --native-dir $output --dlss-dll $DlssDll ` + --output-dir (Join-Path $folder 'private-ptx') --blocks $PtxBlocks ` + > (Join-Path $folder 'ptx-reference.stdout.log') 2> (Join-Path $folder 'ptx-reference.stderr.log') + if ($LASTEXITCODE) { throw "M $layer independent PTX mismatch; stopping: $folder" } + Get-Content -LiteralPath (Join-Path $folder 'ptx-reference.stdout.log') | Write-Host + $summary.layers += @{layer=$layer; passed=$true} + } + $summary.passed=$true +} catch { $summary.error=$_.Exception.Message; Write-Warning $summary.error } +finally { + $env:PYTHONPATH=$previousPython; $env:D4R_DIAG_DIR=$previousDiag + $summary | ConvertTo-Json -Depth 5 | Set-Content -LiteralPath (Join-Path $OutputDirectory 'summary.json') -Encoding UTF8 + Write-Host "Private real-weight replay results: $OutputDirectory (do not redistribute captures/allocations)" +} +if (!$summary.passed) { exit 1 } diff --git a/scripts/windows/test-m-replay.ps1 b/scripts/windows/test-m-replay.ps1 new file mode 100644 index 0000000..9f9f0af --- /dev/null +++ b/scripts/windows/test-m-replay.ps1 @@ -0,0 +1,45 @@ +[CmdletBinding()] +param([string]$HipRoot, [string]$PackageRoot, [string]$OutputDirectory, + [string[]]$Layers = @('enc1','enc2','enc3_tube','dec2','dec1'), + [ValidateRange(1,10000)][int]$Iterations = 32) +$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 (!$OutputDirectory) { $OutputDirectory = Join-Path $repo ('test-results/m-replay-' + (Get-Date -Format 'yyyyMMdd-HHmmss')) } +New-Item -ItemType Directory -Force $OutputDirectory | Out-Null +$previousPython = $env:PYTHONPATH +$previousDiag = $env:D4R_DIAG_DIR +$env:PYTHONPATH = "$(Join-Path $repo '.tools/python/vendor');$previousPython" +$summary = [ordered]@{utc=[DateTime]::UtcNow.ToString('o'); hipRoot=$HipRoot; layers=@(); passed=$false} +try { + foreach ($layer in $Layers) { + if ($layer -notmatch '^(enc[12]|enc3_tube|dec[12])$') { throw "Invalid M layer $layer" } + $folder = Join-Path $OutputDirectory $layer + $fixture = Join-Path $folder 'fixture' + $output = Join-Path $folder 'output' + New-Item -ItemType Directory -Force $folder | Out-Null + $env:D4R_DIAG_DIR = $folder + & python (Join-Path $PackageRoot 'm_replay_validate.py') generate --layer $layer --fixture-dir $fixture ` + > (Join-Path $folder 'fixture.stdout.log') 2> (Join-Path $folder 'fixture.stderr.log') + if ($LASTEXITCODE) { throw "M $layer fixture generation failed: $folder" } + & (Join-Path $PackageRoot 'bin/d4r_native_replay_probe.exe') --hip-root $HipRoot ` + --module (Join-Path $PackageRoot "experimental/m/rrlite_${layer}_4x4_gfx1201.hsaco") ` + --fixture-dir $fixture --output-dir $output --iterations $Iterations ` + > (Join-Path $folder 'gpu.stdout.log') 2> (Join-Path $folder 'gpu.stderr.log') + if ($LASTEXITCODE) { throw "M $layer GPU replay failed: $folder" } + & python (Join-Path $PackageRoot 'm_replay_validate.py') validate --layer $layer --fixture-dir $fixture --output-dir $output ` + > (Join-Path $folder 'reference.stdout.log') 2> (Join-Path $folder 'reference.stderr.log') + if ($LASTEXITCODE) { throw "M $layer reference mismatch; stopping before the next layer: $folder" } + Get-Content -LiteralPath (Join-Path $folder 'reference.stdout.log') | Write-Host + $summary.layers += @{layer=$layer; passed=$true; iterations=$Iterations} + } + $summary.passed = $true +} catch { $summary.error=$_.Exception.Message; Write-Warning $summary.error } +finally { + $env:PYTHONPATH=$previousPython; $env:D4R_DIAG_DIR=$previousDiag + $summary | ConvertTo-Json -Depth 5 | Set-Content -LiteralPath (Join-Path $OutputDirectory 'summary.json') -Encoding UTF8 + Compress-Archive -LiteralPath $OutputDirectory -DestinationPath "$OutputDirectory.zip" -Force + Write-Host "M replay bundle: $OutputDirectory.zip" +} +if (!$summary.passed) { exit 1 } diff --git a/scripts/windows/test-windows-rdna4.ps1 b/scripts/windows/test-windows-rdna4.ps1 index fe5c078..29e9e77 100644 --- a/scripts/windows/test-windows-rdna4.ps1 +++ b/scripts/windows/test-windows-rdna4.ps1 @@ -98,8 +98,8 @@ function Invoke-Probe([string]$Name, [string]$Exe, [string[]]$Arguments) { try { if ($NgxOnly -and (!$NgxCore -or !$DlssDll)) { throw '-NgxOnly requires both locally supplied NVIDIA DLL paths.' } if ($RequireNativeNetwork) { - if ($NgxMode -ne 'evaluate' -or $Preset -ne 11 -or !$env:D4R_ZLUDA_NATIVE_DIR) { - throw '-RequireNativeNetwork currently requires K Evaluate and D4R_ZLUDA_NATIVE_DIR.' + if ($NgxMode -ne 'evaluate' -or $Preset -notin @(11,13) -or !$env:D4R_ZLUDA_NATIVE_DIR) { + throw '-RequireNativeNetwork requires K/M Evaluate and D4R_ZLUDA_NATIVE_DIR.' } $env:D4R_ZLUDA_VERBOSE = '1' } @@ -223,15 +223,18 @@ try { $ngxOk = Invoke-Probe $ngxName $ngxExe $ngxArguments if ($ngxOk -and $RequireNativeNetwork) { $hits = @() - foreach ($layer in @('enc0','enc1','enc2','enc3','enc4','dec5','dec4','dec3','dec2','dec1','dec0')) { - $pattern = '\[d4r-launch\] kernel="dltss_pwin_' + $layer + '_layer" backend=native' + $nativeLayers = if ($Preset -eq 11) { @('enc0','enc1','enc2','enc3','enc4','dec5','dec4','dec3','dec2','dec1','dec0') } + else { @('enc1','enc2','enc3_tube','dec2','dec1') } + foreach ($layer in $nativeLayers) { + $name = if ($Preset -eq 11) { "dltss_pwin_${layer}_layer" } else { "rrlite_${layer}_4x4" } + $pattern = '\[d4r-launch\] kernel="' + $name + '" backend=native' $count = @(Select-String -LiteralPath (Join-Path $OutputDirectory "$ngxName.stderr.log") -Pattern $pattern).Count $hits += @{layer=$layer; nativeLaunches=$count; expectedFrames=$Iterations} if ($count -lt $Iterations) { $ngxOk=$false } } $summary.nativeTransformer = $hits $summary.nativeTransformerPassed = $ngxOk - Write-Host "Native K transformer launches validated=$ngxOk" + Write-Host "Native preset $Preset transformer launches validated=$ngxOk" } if (!$ngxOk) { $exitStatus = 1 diff --git a/tools/windows/m_ptx_validate.py b/tools/windows/m_ptx_validate.py new file mode 100644 index 0000000..d4c373e --- /dev/null +++ b/tools/windows/m_ptx_validate.py @@ -0,0 +1,62 @@ +"""Independent PTX interpreter validation of real M captures; keep outputs private.""" +import argparse +import pathlib +import re +import sys + +import numpy as np + +from m_replay_validate import me, LAYERS, blocks, model_block, stored, metrics +import ptxsim +from extract_dlss_ptx import fatbins, ptx_entries + + +def validate(args): + name = f'rrlite_{args.layer}_4x4' + parameters = me.Params(str(args.capture_dir)) + if parameters.kernel != name: + raise RuntimeError('Capture kernel does not match requested M layer') + native = {i: np.fromfile(args.native_dir / f'alloc-{i}.bin', np.uint8) for i in parameters.allocs} + candidates = [] + for bx, by in blocks(parameters.grid, 12): + model = model_block(parameters, args.layer, bx, by) + expected, actual = stored(parameters, args.layer, bx, by, native, model) + difference = float(np.abs(expected - actual).max()) + candidates.append((difference, bx, by, model, expected, actual)) + selected = sorted(candidates, key=lambda c: (-c[0], c[1], c[2]))[:args.blocks] + pattern = re.compile(rb'\.entry\s+' + name.encode() + rb'\s*\(') + sources = [source.rstrip(b'\0') for _, fatbin in fatbins(args.dlss_dll.read_bytes()) + for source in ptx_entries(fatbin) if pattern.search(source)] + if not sources: + raise RuntimeError('Local DLSS DLL does not contain this M kernel') + args.output_dir.mkdir(parents=True, exist_ok=True) + path = args.output_dir / (name + '.ptx') + path.write_bytes(sources[-1]) + for _, bx, by, model, numpy_expected, actual in selected: + _, memory, steps = ptxsim.run_block(str(path), name, str(args.capture_dir), (bx, by, 0)) + allocation_data = dict(zip(memory.ids, memory.data)) + _, expected = stored(parameters, args.layer, bx, by, allocation_data, model) + print(f'PTX_REFERENCE layer={args.layer} block={bx},{by} statements={steps}', flush=True) + _, _, psnr, fp8_steps = metrics(expected, actual, 'PTX_VS_NATIVE') + metrics(expected, numpy_expected, 'PTX_VS_NUMPY') + if fp8_steps.max() > 1 or psnr < 50: + raise RuntimeError('Native M differs from independent PTX reference') + print(f'PASS M_PTX_REFERENCE layer={args.layer} blocks={len(selected)} nan_inf=0', flush=True) + + +if __name__ == '__main__': + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument('--layer', choices=list(LAYERS), required=True) + parser.add_argument('--capture-dir', type=pathlib.Path, required=True) + parser.add_argument('--native-dir', type=pathlib.Path, required=True) + parser.add_argument('--dlss-dll', type=pathlib.Path, required=True) + parser.add_argument('--output-dir', type=pathlib.Path, required=True) + parser.add_argument('--blocks', type=int, default=2) + args = parser.parse_args() + if args.blocks < 1: + parser.error('--blocks must be positive') + try: + validate(args) + except Exception as error: + print(f'FAIL M_PTX_REFERENCE {error}', file=sys.stderr) + raise SystemExit(5) diff --git a/tools/windows/m_replay_validate.py b/tools/windows/m_replay_validate.py new file mode 100644 index 0000000..504e807 --- /dev/null +++ b/tools/windows/m_replay_validate.py @@ -0,0 +1,178 @@ +"""Public synthetic M fixtures and strict FP16-widening replay validation. + +No NVIDIA binaries, weights or extracted PTX are needed to generate fixtures. +The wire format is upstream's replay format, also accepted by native_replay_probe. +""" +import argparse +import math +import pathlib +import struct +import sys + +import numpy as np + +ROOT = pathlib.Path(__file__).resolve().parents[2] +REFERENCE = pathlib.Path(__file__).resolve().parent / 'reference' +if not (REFERENCE / 'swin_check.py').exists(): + REFERENCE = ROOT / 'kernels/tools' +sys.path.insert(0, str(REFERENCE)) +import model_enc3 as me +import swin_model as sm +import swin_check +swin_check.CLAMP_PATCH_MERGE = True +from swin_check import LAYERS, blocks, model_block, codes_at + + +def encode(values): + """Encode already quantised finite values without depending on vendor FP8.""" + values = me.q8(values) + table = me.E4M3[:127] + codes = np.searchsorted(table, np.abs(values)).astype(np.uint8) + return codes | ((values < 0).astype(np.uint8) << 7) + + +def generate(layer, directory, size): + C, heads, NPM, CIN = LAYERS[layer] + pre = 4 * C * CIN + 8 * C if CIN else 0 + layout = sm.layout(C, heads, pre) + weights = np.zeros(layout['end'] + 4 * C * NPM + 2 * NPM, np.uint8) + + def half(offset, values): + data = np.asarray(values, np.float16).reshape(-1).view(np.uint8) + weights[offset:offset + data.size] = data + + def matrix(base, ks, ns, k, n, source, dest, value=.25): + offsets = sm.woff_table(base, ks, ns, k, n) + weights[offsets[source, dest]] = encode(np.array(value)) + + if CIN: + cols = 4 * C // heads + for head in range(heads): + for n in range(cols): + matrix(head * CIN * cols, 512, 16 * CIN, CIN, cols, + (head * cols + n) % CIN, n) + half(4 * C * CIN, np.full(4 * C, 1 / 64)) + half(layout['g1'], np.ones(C)) + half(layout['g2'], np.ones(C)) + half(layout['bo'], np.full(C, 1 / 64)) + half(layout['b2'], np.full(C, 1 / 64)) + for head, base in enumerate(layout['head']): + for n in range(32): + source = (32 * head + n) % C + matrix(base, 1024, 512, C, 32, source, n) + matrix(base + 32 * C, 1024, 512, C, 32, source, n, .5) + matrix(base + 64 * C + 512, 0, 512, 32, C, n, source, .5) + for group, base in enumerate(layout['b1']): + half(base + 32 * C, np.full(32, 1 / 64)) + for n in range(32): + source = (32 * group + n) % C + matrix(base, 512, 16 * C, C, 32, source, n, .5) + matrix(base + 32 * C + 64, 0, 512, 32, C, n, source, .25) + for group in range(NPM // 32): + for n in range(32): + matrix(layout['end'] + group * 128 * C, 512, 64 * C, + 4 * C, 32, (group * 32 + n) % (4 * C), n, 8.) + half(layout['end'] + 4 * C * NPM + 64 * group, np.full(32, 1 / 64)) + + input_count = size * size * C if not CIN else (size // 2) ** 2 * CIN + input_data = encode((np.arange(input_count) % 9 - 4) / 4) + skip = encode((np.arange(size * size * C) % 7 - 3) / 8) if CIN else np.empty(0, np.uint8) + allocations = {0: weights, 1: np.concatenate([input_data, skip]), + 2: np.full(size * size * C, 0x7f, np.uint8)} + if NPM: + allocations[3] = np.full((size // 2) ** 2 * NPM, 0x7f, np.uint8) + bases = {i: 0x100000000 + i * 0x1000000 for i in allocations} + parameters = bytearray(88 if layer == 'enc3_tube' else 56) + struct.pack_into('= 0) & (x < w) & (y >= 0) & (y < h) + codes = codes_at(allocations[identifier], base, ptr, w, h, x[valid], y[valid], channels) + return me.E4M3[codes].ravel(), valid + + x, y = parameters.token_xy(bx, by) + actual, valid = read(parameters.out, parameters.W, parameters.H, x, y, C) + got, want = [actual], [model[0][valid].ravel()] + if NPM: + r = np.arange(16) + mx = (8 * bx - parameters.sx) // 2 + (r & 3) + my = (8 * by - parameters.sy) // 2 + (r >> 2) + actual, valid = read(parameters.p48, parameters.W // 2, parameters.H // 2, mx, my, NPM) + got.append(actual) + want.append(model[1][valid].ravel()) + return np.concatenate(want), np.concatenate(got) + + +def metrics(expected, actual, label): + if not np.isfinite(expected).all() or not np.isfinite(actual).all(): + raise RuntimeError(f'{label}: NaN/Inf or unwritten output') + error = np.abs(actual.astype(np.float64) - expected.astype(np.float64)) + peak = max(1., float(np.abs(expected).max())) + maximum = float(error.max()) + relative = float((error / np.maximum(np.abs(expected), .001)).max()) + rms = float(np.sqrt(np.mean(error ** 2))) + psnr = math.inf if rms == 0 else 20 * math.log10(peak / rms) + rank = lambda v: np.searchsorted(me._POS, np.abs(v)) * np.sign(v) + steps = np.abs(rank(actual) - rank(expected)) + print(f'{label} elements={actual.size} max_abs={maximum:.9g} max_rel={relative:.9g} ' + f'rms={rms:.9g} psnr_db={psnr:.6g} identical={np.mean(error == 0):.9g} ' + f'max_fp8_steps={int(steps.max())}', flush=True) + return maximum, peak, psnr, steps + + +def validate(layer, directory, output, block_count, real_capture): + parameters = me.Params(str(directory)) + if parameters.kernel != f'rrlite_{layer}_4x4': + raise RuntimeError('Replay kernel does not match requested M layer') + allocations = {i: np.fromfile(output / f'alloc-{i}.bin', np.uint8) for i in parameters.allocs} + pairs = [stored(parameters, layer, bx, by, allocations, model_block(parameters, layer, bx, by)) + for bx, by in blocks(parameters.grid, block_count)] + expected, actual = (np.concatenate(values) for values in zip(*pairs)) + maximum, peak, psnr, steps = metrics(expected, actual, f'M_REPLAY_REFERENCE kernel={layer}') + if (not real_capture and maximum != 0) or (real_capture and (steps.max() > 1 or psnr < 50)): + raise RuntimeError(f'M mismatch: exact synthetic / <=1 FP8 step and >=50 dB real gate; peak={peak}') + print(f'PASS M_REPLAY_REFERENCE kernel={layer} blocks={len(pairs)} architecture=gfx1201 nan_inf=0') + + +if __name__ == '__main__': + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument('mode', choices=['generate', 'validate']) + parser.add_argument('--layer', choices=list(LAYERS), required=True) + parser.add_argument('--fixture-dir', type=pathlib.Path, required=True) + parser.add_argument('--output-dir', type=pathlib.Path) + parser.add_argument('--size', type=int, choices=[8, 16, 24, 32], default=16) + parser.add_argument('--blocks', type=int, default=12) + parser.add_argument('--real-capture', action='store_true') + args = parser.parse_args() + try: + if args.mode == 'generate': + generate(args.layer, args.fixture_dir, args.size) + else: + if not args.output_dir or args.blocks < 1: + parser.error('validate requires --output-dir and positive --blocks') + validate(args.layer, args.fixture_dir, args.output_dir, args.blocks, args.real_capture) + except Exception as error: + print(f'FAIL M_REPLAY {error}', file=sys.stderr) + raise SystemExit(5) diff --git a/tools/windows/test_ptx_reference.py b/tools/windows/test_ptx_reference.py new file mode 100644 index 0000000..585e7e7 --- /dev/null +++ b/tools/windows/test_ptx_reference.py @@ -0,0 +1,46 @@ +"""Regression for divergent finite loops in the public CPU PTX reference.""" +import pathlib +import sys +import tempfile +import unittest + +import numpy as np + +ROOT = pathlib.Path(__file__).resolve().parents[2] +sys.path.insert(0, str(ROOT / 'kernels/tools')) +import ptxsim + + +class ControlFlow(unittest.TestCase): + def test_divergent_backward_loop_reconverges(self): + source = '''.visible .entry pattern(.param .align 8 .b8 pattern_param_0[8]) + .maxntid 32, 1, 1 + { + .reg .b32 %r<4>; + .reg .pred %p<1>; + mov.u32 %r0, %tid.x; + and.b32 %r0, %r0, 3; + add.u32 %r0, %r0, 1; + mov.u32 %r1, 0; + $loop: + add.u32 %r1, %r1, %r0; + sub.u32 %r0, %r0, 1; + setp.ne.u32 %p0, %r0, 0; + @%p0 bra $loop; + add.u32 %r2, %r1, 7; + ret; + } + ''' + with tempfile.TemporaryDirectory() as directory: + path = pathlib.Path(directory) / 'pattern.ptx' + path.write_text(source) + statements, types, shared = ptxsim.parse_kernel(str(path), 'pattern') + block = ptxsim.Block(statements, types, shared, None, 'pattern_param_0', (32, 1, 1), (1, 1, 1), (0, 0, 0)) + block.run(limit=1000) + n = np.arange(32) % 4 + 1 + np.testing.assert_array_equal(block.regs['%r1'], n * (n + 1) // 2) + np.testing.assert_array_equal(block.regs['%r2'], n * (n + 1) // 2 + 7) + + +if __name__ == '__main__': + unittest.main()