diff --git a/CMakeLists.txt b/CMakeLists.txt index 9a04875..926cdae 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -99,6 +99,8 @@ install(FILES "${probe_object}" DESTINATION bin) install(FILES "${wmma_object}" DESTINATION bin) install(FILES "${k_enc1_object}" DESTINATION experimental/k) install(FILES scripts/windows/test-windows-rdna4.ps1 DESTINATION .) +install(FILES tools/windows/k_enc1_validate.py DESTINATION .) +install(FILES kernels/tools/pwin_model.py DESTINATION reference) install(FILES docs/windows-rdna4-port.md DESTINATION docs) include(CTest) add_test(NAME hip_gfx1201 COMMAND d4r_hip_gfx1201_probe @@ -107,9 +109,9 @@ set_tests_properties(hip_gfx1201 PROPERTIES TIMEOUT 90 LABELS "hardware;gfx1201" add_test(NAME gfx12_wmma COMMAND d4r_gfx12_wmma_probe --hip-root "${D4R_HIP_ROOT}" --module "${wmma_object}" --iterations 16) set_tests_properties(gfx12_wmma PROPERTIES TIMEOUT 90 LABELS "hardware;gfx1201;wmma") -add_test(NAME k_enc1_module COMMAND d4r_k_module_probe +add_test(NAME k_enc1_execution COMMAND d4r_k_module_probe --hip-root "${D4R_HIP_ROOT}" --module "${k_enc1_object}") -set_tests_properties(k_enc1_module PROPERTIES TIMEOUT 90 LABELS "hardware;gfx1201;k_compile") +set_tests_properties(k_enc1_execution PROPERTIES TIMEOUT 90 LABELS "hardware;gfx1201;k_execution") add_test(NAME d3d12_hip_interop COMMAND d4r_d3d12_hip_interop_probe --hip-root "${D4R_HIP_ROOT}" --module "${probe_object}" --iterations 32) set_tests_properties(d3d12_hip_interop PROPERTIES TIMEOUT 120 LABELS "hardware;interop") diff --git a/README.md b/README.md index 512e32f..8b9e3c8 100644 --- a/README.md +++ b/README.md @@ -8,8 +8,9 @@ It is an experiment. It has been tested on one GPU in a handful of games, it dep **Windows / RDNA4 development:** the `windows-rdna4` branch adds native Windows gfx1201 HIP/ZLUDA diagnostics, validated D3D12/HIP interop, an NGX init probe, -and a CPU-checked gfx12 WMMA backend. An experimental K `enc1` kernel compiles -and loads but has not passed transformer-output validation. +and a CPU-checked gfx12 WMMA backend. The first K `enc1` kernel executes on +gfx1201 and matches the existing numpy model on a synthetic nonzero fixture. +The remaining K/M kernels and Windows game runtime are still in progress. Build commands, hardware results and remaining porting gates are in [docs/windows-rdna4-port.md](docs/windows-rdna4-port.md). This branch does not yet provide a Windows DLSS runtime. diff --git a/docs/windows-rdna4-port.md b/docs/windows-rdna4-port.md index 3ef6c9f..f518d2a 100644 --- a/docs/windows-rdna4-port.md +++ b/docs/windows-rdna4-port.md @@ -7,8 +7,8 @@ This is a work log for an implementation in progress, not a claim of DLSS suppor ## Milestone and gates Current milestone: M1/M2/M4/M5 passed; NGX init + DLSS CUDA capability discovery -passed with the user's DLLs. The first K `enc1` code object compiles and loads; -its prep/transformer have not yet been executed or checked against replay. +passed with the user's DLLs. M6 is in progress: the first K `enc1` prep and +transformer run on gfx1201 and pass a synthetic nonzero numpy reference check. Do not integrate NGX until the integer PTX workload is stable on the real GPU. | Gate | Status | @@ -19,7 +19,7 @@ Do not integrate NGX until the integer PTX workload is stable on the real GPU. | M3: Windows NGX initialization and simple DLSS path | Init + SR capabilities PASS; Evaluate/transformer pending | | 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, 16 cases, max abs/relative error 0 | -| M6: K layers, full transformer and image validation | Pending | +| M6: K layers, full transformer and image validation | Partial: `enc1` executes, nonzero fixture matches reference; other layers/full network pending | | M7: M FP16-equivalent baseline and full transformer | Pending | | M8: standalone Windows D3D12 NGX harness | Pending | | M9: OptiScaler integration, profiling, installation | Pending | @@ -42,6 +42,9 @@ Do not integrate NGX until the integer PTX workload is stable on the real GPU. `9623b97ca511eaa176905075a504393eca6fde594222af10b55e4fb1edfe5a25`. No machine-wide runtime replacement is performed. * CMake 3.31.6, Ninja 1.11.1 (Python package 1.11.1.4). +* Synthetic K numerical reference: Python 3.11 and local NumPy `2.4.6` + ([PyPI release](https://pypi.org/project/numpy/2.4.6/)); installed only + into `.tools/python/vendor` by setup, outside the game runtime. * Windows host toolchain: LLVM-MinGW `20260922`, UCRT x64, SHA256 `e3ad77d117a4bea19a7a3b333341824d79a5a371004a10e25b8504e7b3047666`. This is a portable fallback on this machine, where MSVC is not installed. @@ -136,9 +139,14 @@ compiler uses `kernels/common/hip_device_minimal.h`, exposing only public Clang AMDGPU work-item/grid/barrier builtins and HIP-compatible qualifiers; the ordinary Linux HIP include path remains. CMake builds `enc1` for gfx1201 into `experimental/k/` and verifies that HIP loads both prep and transformer -entry points. This is a **compile/load smoke test**, not K layer correctness or -DLSS execution. The code object must stay out of ZLUDA's native override path -until numerical K replay validation passes. +entry points. The `k_module_probe` now launches prep and transformer. The +identity fixture verifies all 4096 full-resolution and 1024 merged FP16 values +bitwise. A second fixture enables nonzero V projections, position-only +attention, Wo, MLP/GELU and patch merge; it checks finite values and a patch +identity, then saves all inputs/outputs for `k_enc1_validate.py`. That script +runs the existing `pwin_model.py` and measures max absolute/relative error and +PSNR. The code object remains in `experimental/k/`, outside ZLUDA's override +path, until the other K layers and real captured weights/activations pass. ## External interop blockers @@ -245,14 +253,16 @@ This builds into `build/windows-rdna4-therock` and installs into `dist/windows-rdna4-therock`. The stable profile stays available for reproducing the runtime bug. It is expected to return failure at the lifetime gate. -One command adds the local NGX initialization gate after the seven hardware gates: +One command adds the local NGX initialization gate after eight diagnostics: ```powershell powershell -NoProfile -File scripts/windows/test-windows-rdna4.ps1 -RuntimeProfile therock -NgxCore "C:/Users/Administrator/d4r/_nvngx.dll" -DlssDll "C:/Users/Administrator/d4r/nvngx_dlss.dll" ``` -Success requires all eight tests and `PASS NGX_INIT ... sr_available=1`; -`transformer_executed=0` is explicit. On NGX failure, CUDA trace and NVAPI query +Success requires all nine tests, including `PASS K_REFERENCE` and +`PASS NGX_INIT ... sr_available=1`. The NGX probe still reports +`transformer_executed=0`: the standalone K launch has not been integrated into +DLSS. On NGX failure, CUDA trace and NVAPI query logs are collected automatically. NVIDIA DLLs remain local test inputs. The build accepts `-HipRoot`, `-ZludaRoot`, `-ToolchainRoot`, `-BuildDirectory` @@ -324,11 +334,33 @@ stable mapped-view lifetime bug still excludes its interop fast path. WMMA raw and legacy-layout adapter each match the scalar reference exactly over 16 changing input matrices, with both one-step and two-step accumulation. The `enc1` K module is built and loaded on gfx1201, with functions resolved; -`transformer_executed=0` is explicit in the module probe. Neither K nor M -transformer correctness is claimed yet. The same K source also compiles for +`transformer_executed=0` was explicit in the module probe at that milestone. +The same K source also compiles for `gfx1101` through the unchanged gfx11 builtin branch (compile-only regression; no RDNA3 GPU is present in this Windows machine). +2026-09-29 M6 first kernel: `enc1` runs on the RX 9070 XT with zero and nonzero +synthetic weights. The nonzero fixture exercises V, position-only attention, +Wo, MLP/GELU and patch merge. Compared with `pwin_model.py`, full output +max absolute error `0.000244140625`, PSNR `98.21 dB`; merged output max +absolute error `0.000244140625`, PSNR `97.63 dB`. No NaN/Inf, all 4096 full +and 1024 merged elements written, 3886 full elements changed. TheRock +diagnostic bundle `test-results/m6-k-enc1-final.zip` has 9/9 gates PASS; +CTest has 7/7 PASS. +Reference dependency: pinned local NumPy `2.4.6` on Python 3.11, installed by +`setup-windows-tools.ps1` into `.tools/python/vendor` (not the game runtime). +This is a synthetic one-layer check, not proof of complete K or DLSS output. + +Stable HIP SDK 7.2 compiler caveat: its `-O2` code object for `enc1` produces +NaN already on the identity fixture (first element `0xfe00` instead of +`0xbc00`). Its `-O1` object also miscomputes the identity fixture; `-O0` +passes but is unsuitable for the fast path. Cross-testing isolates the problem +to the stable compiler/code object: the stable-built `-O2` object fails under +TheRock runtime, while the TheRock-built `-O2` object passes under stable HIP +runtime. The independent WMMA probe passes under both compilers. Until a +smaller compiler reproducer or source-level workaround is found, use the +TheRock compiler for native K and the TheRock runtime for leak-free interop. + Audit found an upstream validation blocker: `swin_model.py` imports `model_enc3`, which is absent from the repository. M reference validation needs that dependency restored or an independently checked replacement. diff --git a/scripts/windows/setup-windows-tools.ps1 b/scripts/windows/setup-windows-tools.ps1 index 6aab3ae..5b0d70a 100644 --- a/scripts/windows/setup-windows-tools.ps1 +++ b/scripts/windows/setup-windows-tools.ps1 @@ -38,4 +38,9 @@ if (!(Test-Path (Join-Path $localTools 'python/cmake/data/bin/cmake.exe')) -or python -m pip install --target (Join-Path $localTools 'python') cmake==3.31.6 ninja==1.11.1.4 if ($LASTEXITCODE) { throw "pip install exit=$LASTEXITCODE" } } +if (!(Test-Path (Join-Path $localTools 'python/vendor/numpy/__init__.py'))) { + python -m pip install --disable-pip-version-check --only-binary=:all: ` + --target (Join-Path $localTools 'python/vendor') numpy==2.4.6 + if ($LASTEXITCODE) { throw "NumPy reference install exit=$LASTEXITCODE" } +} Write-Host 'Tools ready. Run scripts/windows/build-windows-rdna4.ps1.' diff --git a/scripts/windows/test-windows-rdna4.ps1 b/scripts/windows/test-windows-rdna4.ps1 index dea2b89..21645e5 100644 --- a/scripts/windows/test-windows-rdna4.ps1 +++ b/scripts/windows/test-windows-rdna4.ps1 @@ -11,8 +11,8 @@ param( [int]$TimeoutSeconds = 180 ) $ErrorActionPreference = 'Stop' +$repoRoot = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) if ($RuntimeProfile -eq 'therock') { - $repoRoot = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) # Works both from scripts/windows and the installed dist/ directory. if (!$PSBoundParameters.ContainsKey('HipRoot')) { $HipRoot = Join-Path $repoRoot '.tools/therock-10.2.0a20260929/_rocm_sdk_core' } if (!$PackageRoot) { @@ -42,6 +42,7 @@ $originalDiag = $env:D4R_DIAG_DIR $originalLog = $env:ZLUDA_LOG_DIR $originalCuda = $env:ZLUDA_CUDA_LIB $originalNvapi = $env:D4R_NVAPI_BACKEND +$originalPythonPath = $env:PYTHONPATH $exitStatus = 1 $ngxRuntimeDirectory = $null $summary = [ordered]@{utc=[DateTime]::UtcNow.ToString('o'); profile=$RuntimeProfile; hipRoot=$HipRoot; zludaRoot=$ZludaRoot; tests=@()} @@ -91,6 +92,10 @@ try { $env:ZLUDA_LOG_DIR = Join-Path $OutputDirectory 'zluda-trace' $env:ZLUDA_CUDA_LIB = Join-Path $ZludaRoot 'nvcuda.dll' $env:D4R_NVAPI_BACKEND = Join-Path $ZludaRoot 'nvapi64.dll' + $localPythonVendor = Join-Path $repoRoot '.tools/python/vendor' + if (Test-Path (Join-Path $localPythonVendor 'numpy/__init__.py')) { + $env:PYTHONPATH = "$localPythonVendor;$originalPythonPath" + } $comgr = Join-Path $HipRoot 'bin/amd_comgr_3.dll' if (!(Test-Path -LiteralPath $comgr)) { $comgr = Join-Path $HipRoot 'bin/amd_comgr.dll' } $files = @((Join-Path $HipRoot 'bin/amdhip64_7.dll'), $comgr, @@ -122,9 +127,16 @@ try { $wmmaOk = Invoke-Probe 'gfx12-wmma' (Join-Path $bin 'd4r_gfx12_wmma_probe.exe') @( '--hip-root', $HipRoot, '--module', (Join-Path $bin 'wmma_gfx1201.hsaco'), '--iterations', "$Iterations") if (!$wmmaOk) { throw 'gfx12 WMMA layout validation failed; refusing to mark K/M readiness.' } - $kModuleOk = Invoke-Probe 'k-enc1-module' (Join-Path $bin 'd4r_k_module_probe.exe') @( - '--hip-root', $HipRoot, '--module', (Join-Path $PackageRoot 'experimental/k/dltss_pwin_enc1_layer_gfx1201.hsaco')) - if (!$kModuleOk) { throw 'K enc1 module cannot load on gfx1201; transformer execution remains untested.' } + $kFixture = Join-Path $OutputDirectory 'k-enc1-fixture' + $kModuleOk = Invoke-Probe 'k-enc1-gpu' (Join-Path $bin 'd4r_k_module_probe.exe') @( + '--hip-root', $HipRoot, '--module', (Join-Path $PackageRoot 'experimental/k/dltss_pwin_enc1_layer_gfx1201.hsaco'), + '--fixture-dir', $kFixture, '--iterations', "$Iterations") + if (!$kModuleOk) { throw 'K enc1 GPU execution failed; see the k-enc1-gpu logs.' } + $pythonCommand = Get-Command python -ErrorAction SilentlyContinue + if (!$pythonCommand) { throw 'Python 3.11+ with NumPy 2.4.6 is needed for the K reference check; run setup-windows-tools.ps1.' } + $kReferenceOk = Invoke-Probe 'k-enc1-reference' $pythonCommand.Source @( + (Join-Path $PackageRoot 'k_enc1_validate.py'), '--fixture-dir', $kFixture) + if (!$kReferenceOk) { throw 'K enc1 output differs from pwin_model.py; see the k-enc1-reference logs.' } $cudaOk = $true foreach ($context in @('primary', 'created')) { $ok = Invoke-Probe "cuda-$context" (Join-Path $bin 'd4r_cuda_driver_probe.exe') @( @@ -200,6 +212,7 @@ try { $env:ZLUDA_LOG_DIR = $originalLog $env:ZLUDA_CUDA_LIB = $originalCuda $env:D4R_NVAPI_BACKEND = $originalNvapi + $env:PYTHONPATH = $originalPythonPath $archive = "$OutputDirectory.zip" Compress-Archive -LiteralPath $OutputDirectory -DestinationPath $archive -Force Write-Host "Diagnostic bundle: $archive" diff --git a/tools/windows/diagnostic.h b/tools/windows/diagnostic.h index 7dc06ee..79e4096 100644 --- a/tools/windows/diagnostic.h +++ b/tools/windows/diagnostic.h @@ -89,6 +89,7 @@ struct Args { std::string ngx_core, dlss_dll; std::string nvapi_dll; std::string ngx_abi = "driver"; + std::string fixture_dir; unsigned iterations = 32; int device = -1; Args(int argc, char** argv) { @@ -106,6 +107,7 @@ struct Args { else if (key == "--dlss-dll") dlss_dll = value; else if (key == "--nvapi-dll") nvapi_dll = value; else if (key == "--ngx-abi") ngx_abi = value; + else if (key == "--fixture-dir") fixture_dir = value; else if (key == "--iterations") { size_t end = 0; iterations = static_cast(std::stoul(value, &end)); diff --git a/tools/windows/k_enc1_validate.py b/tools/windows/k_enc1_validate.py new file mode 100644 index 0000000..81ce18d --- /dev/null +++ b/tools/windows/k_enc1_validate.py @@ -0,0 +1,82 @@ +#!/usr/bin/env python3 +"""Compare a real gfx1201 K enc1 launch with the existing numpy K reference.""" +import argparse +import math +import pathlib +import sys + +import numpy as np + +SCRIPT = pathlib.Path(__file__).resolve() +REFERENCE = SCRIPT.parent / "reference" +if not (REFERENCE / "pwin_model.py").exists(): + REFERENCE = SCRIPT.parents[2] / "kernels" / "tools" +sys.path.insert(0, str(REFERENCE)) +from pwin_model import pwin_encoder # noqa: E402 + + +class SyntheticParams: + def __init__(self, folder): + self.input = np.fromfile(folder / "input.bin", dtype=np.uint8) + self.weights = np.fromfile(folder / "weights.bin", dtype=np.uint8) + + def i32x2(self, offset): + return (8, 8) if offset == 0 else (0, 0) + + def ptr(self, offset): + if offset == 8: + return self.input, 0 + if offset == 64: + return self.weights, 0 + raise ValueError(f"unsupported synthetic pointer offset {offset}") + + def f16(self, offset, byte, count): + if offset != 64: + raise ValueError(f"unsupported synthetic half pointer offset {offset}") + return self.weights[byte:byte + 2 * count].view(np.float16).copy() + + +def compare(name, actual, expected): + if actual.shape != expected.shape: + raise RuntimeError(f"{name} shape mismatch: {actual.shape} != {expected.shape}") + if not np.isfinite(actual).all() or not np.isfinite(expected).all(): + raise RuntimeError(f"{name} contains NaN or Inf") + diff = np.abs(actual.astype(np.float32) - expected.astype(np.float32)) + maximum = float(diff.max()) + relative = float((diff / np.maximum(np.abs(expected.astype(np.float32)), 1e-3)).max()) + rms = float(np.sqrt(np.mean(np.square(diff.astype(np.float64))))) + peak = max(1.0, float(np.abs(expected.astype(np.float32)).max())) + psnr = math.inf if rms == 0 else 20.0 * math.log10(peak / rms) + mismatch = np.unravel_index(int(diff.argmax()), diff.shape) + print(f"K_REFERENCE {name} max_abs={maximum:.9g} max_rel={relative:.9g} " + f"rms={rms:.9g} psnr_db={psnr:.5g} worst={mismatch} " + f"expected={float(expected[mismatch]):.9g} actual={float(actual[mismatch]):.9g}") + return maximum, psnr + + +def main(): + parser = argparse.ArgumentParser() + parser.add_argument("--fixture-dir", required=True, type=pathlib.Path) + args = parser.parse_args() + folder = args.fixture_dir + params = SyntheticParams(folder) + reference_full, reference_merged = pwin_encoder( + params, 0, 0, H=2, C=64, COUT=64, posattn=True) + actual_full = np.fromfile(folder / "full.bin", np.float16).reshape(64, 64) + actual_merged = np.fromfile(folder / "merged.bin", np.float16).reshape(16, 64) + full_abs, full_psnr = compare("full", actual_full, reference_full) + merged_abs, merged_psnr = compare("merged", actual_merged, reference_merged) + # Initial acceptance for the synthetic nonzero FP16 fixture. Tighten this + # after each K stage has been compared on captured DLSS weights/activations. + if full_abs > 0.001 or full_psnr < 60.0 or merged_abs > 0.001 or merged_psnr < 60.0: + raise RuntimeError("K enc1 numerical mismatch against pwin_model.py") + print("PASS K_REFERENCE architecture=gfx1201 kernel=enc1 nonzero_weights=1 " + "transformer_executed=1 nan_inf=0") + + +if __name__ == "__main__": + try: + main() + except Exception as error: + print(f"FAIL K_REFERENCE {error}", file=sys.stderr) + raise SystemExit(5) diff --git a/tools/windows/k_module_probe.cpp b/tools/windows/k_module_probe.cpp index 76bb41a..46fb79f 100644 --- a/tools/windows/k_module_probe.cpp +++ b/tools/windows/k_module_probe.cpp @@ -1,4 +1,37 @@ #include "hip_api.h" +#include + +namespace { +struct PwinParams { + int width, height; + const uint8_t* input; + const uint8_t* skip; + uint8_t* merged; + uint8_t* full; + uint64_t pad40, pad48; + int sx, sy; + const uint8_t* weights; + uint8_t rest[104]; +}; +static_assert(sizeof(PwinParams) == 176, "K parameter ABI"); +static_assert(offsetof(PwinParams, weights) == 64, "K weight pointer ABI"); + +uint16_t half_quarter(int value) +{ + const uint16_t positive[] = {0x0000u, 0x3400u, 0x3800u, 0x3a00u, 0x3c00u}; + if (value < -4 || value > 4) throw std::runtime_error("bad half fixture"); + return static_cast(positive[std::abs(value)] | (value < 0 ? 0x8000u : 0u)); +} + +struct Allocation { + d4r::diag::HipApi& hip; + void* pointer = nullptr; + Allocation(d4r::diag::HipApi& api, size_t bytes) : hip(api) { + hip.check(hip.hipMalloc(&pointer, bytes), "hipMalloc(K fixture)"); + } + ~Allocation() { if (pointer) (void)hip.hipFree(pointer); } +}; +} int main(int argc, char** argv) { @@ -16,12 +49,160 @@ int main(int argc, char** argv) HipApi& hip; hipModule_t module; ~Cleanup() { if (module) (void)hip.hipModuleUnload(module); } } cleanup{hip, module}; - for (const char* name : {"dltss_pwin_enc1_layer_prep", "dltss_pwin_enc1_layer"}) { - hipFunction_t function = nullptr; - hip.check(hip.hipModuleGetFunction(&function, module, name), name); - if (!function) throw std::runtime_error(std::string("Null K function: ") + name); + hipFunction_t prep = nullptr, transformer = nullptr; + hip.check(hip.hipModuleGetFunction(&prep, module, "dltss_pwin_enc1_layer_prep"), "K enc1 prep"); + hip.check(hip.hipModuleGetFunction(&transformer, module, "dltss_pwin_enc1_layer"), "K enc1 transformer"); + if (!prep || !transformer) throw std::runtime_error("Null K entry point"); + + // Identity fixture: all projection, MLP and patch-merge weights are zero. + // G1/G2 are one. The full transformer must return its input through its + // residual path and the merged output must be zero. Every stage launches. + constexpr size_t full_elements = 8u * 8u * 64u; + constexpr size_t merged_elements = 4u * 4u * 64u; + constexpr size_t weight_bytes = 149504u; + constexpr size_t prep_items = 256u * 16u + 1024u; + std::vector weights(weight_bytes, 0); + auto set_half = [&](size_t byte, uint16_t bits) { + weights.at(byte) = static_cast(bits); + weights.at(byte + 1) = static_cast(bits >> 8); + }; + for (size_t c = 0; c < 64u; ++c) { + set_half(2u * c, 0x3c00u); // G1 + set_half(49408u + 2u * c, 0x3c00u); // G2 } - std::printf("PASS K_MODULE architecture=gfx1201 functions=prep,transformer transformer_executed=0\n"); + std::vector input(full_elements), full(full_elements, 0x7e00u); + std::vector merged(merged_elements, 0x7e00u); + for (size_t i = 0; i < input.size(); ++i) + input[i] = half_quarter(static_cast(i % 9u) - 4); + Allocation dev_input(hip, input.size() * sizeof(uint16_t)); + Allocation dev_weights(hip, weights.size()); + Allocation dev_full(hip, full.size() * sizeof(uint16_t)); + Allocation dev_merged(hip, merged.size() * sizeof(uint16_t)); + hip.check(hip.hipMemcpy(dev_input.pointer, input.data(), input.size() * sizeof(uint16_t), + hipMemcpyHostToDevice), "hipMemcpy(K input)"); + hip.check(hip.hipMemcpy(dev_weights.pointer, weights.data(), weights.size(), + hipMemcpyHostToDevice), "hipMemcpy(K weights)"); + PwinParams params{}; + params.width = params.height = 8; + params.input = static_cast(dev_input.pointer); + params.full = static_cast(dev_full.pointer); + params.merged = static_cast(dev_merged.pointer); + params.weights = static_cast(dev_weights.pointer); + void* arguments[] = {¶ms}; + hip.check(hip.hipModuleLaunchKernel(prep, static_cast((prep_items + 127u) / 128u), 1, 1, + 128, 1, 1, 0, nullptr, arguments, nullptr), "hipModuleLaunchKernel(K prep)"); + hip.check(hip.hipDeviceSynchronize(), "hipDeviceSynchronize(K prep)"); + for (unsigned iteration = 0; iteration < args.iterations; ++iteration) { + std::fill(full.begin(), full.end(), 0x7e00u); + std::fill(merged.begin(), merged.end(), 0x7e00u); + hip.check(hip.hipMemcpy(dev_full.pointer, full.data(), full.size() * sizeof(uint16_t), + hipMemcpyHostToDevice), "hipMemcpy(K full sentinel)"); + hip.check(hip.hipMemcpy(dev_merged.pointer, merged.data(), merged.size() * sizeof(uint16_t), + hipMemcpyHostToDevice), "hipMemcpy(K merged sentinel)"); + hip.check(hip.hipModuleLaunchKernel(transformer, 1, 1, 1, 32, 1, 4, 0, + nullptr, arguments, nullptr), "hipModuleLaunchKernel(K transformer)"); + hip.check(hip.hipDeviceSynchronize(), "hipDeviceSynchronize(K transformer)"); + hip.check(hip.hipMemcpy(full.data(), dev_full.pointer, full.size() * sizeof(uint16_t), + hipMemcpyDeviceToHost), "hipMemcpy(K full result)"); + hip.check(hip.hipMemcpy(merged.data(), dev_merged.pointer, merged.size() * sizeof(uint16_t), + hipMemcpyDeviceToHost), "hipMemcpy(K merged result)"); + for (size_t i = 0; i < full.size(); ++i) + if (full[i] != input[i]) { + std::fprintf(stderr, "FAIL K_IDENTITY iteration=%u full_index=%zu expected=0x%04x actual=0x%04x\n", + iteration, i, input[i], full[i]); + return 5; + } + for (size_t i = 0; i < merged.size(); ++i) + if (merged[i] != 0u) { + std::fprintf(stderr, "FAIL K_IDENTITY iteration=%u merged_index=%zu expected=0x0000 actual=0x%04x\n", + iteration, i, merged[i]); + return 5; + } + } + std::printf("PASS K_IDENTITY architecture=gfx1201 iterations=%u full_elements=%zu merged_elements=%zu transformer_executed=1 nonzero_weights=0\n", + args.iterations, full_elements, merged_elements); + + // Exercise nonzero V projection, position-only attention and Wo GEMMs. + // The generated data is saved for comparison with pwin_model.py. + auto fragment_offset = [](unsigned k, unsigned n) { + return 64u * (n & 7u) + 16u * ((k & 7u) >> 1) + 8u * (n >> 3) + + 4u * (k >> 3) + 2u * (k & 1u); + }; + auto set_matrix = [&](size_t base, unsigned k, unsigned n, unsigned nt_stride) { + set_half(base + 512u * (k >> 4) + nt_stride * (n >> 4) + + fragment_offset(k & 15u, n & 15u), 0x3c00u); + }; + constexpr size_t qkv = 128u, head = 12288u, bias = 24704u, wo = 41216u; + constexpr size_t w1 = 49664u, w2 = 82944u, pm = 115712u; + for (unsigned k = 0; k < 32u; ++k) { + set_matrix(qkv + 2048u * 2u, k, k, 1024u); // head 0 V <- input 0:32 + set_matrix(qkv + head + 2048u * 5u, k, k, 1024u); // head 1 V <- input 32:64 + set_matrix(wo, k, k, 1024u); // head 0 Wo -> output 0:32 + set_matrix(wo + 64u * 64u, k, k + 32u, 1024u); // head 1 Wo -> output 32:64 + set_matrix(w1, k, k, 2048u); // first MLP group input -> hidden + set_matrix(w2, k, k, 1024u); // first MLP group hidden -> output + } + for (unsigned k = 0; k < 64u; ++k) + set_matrix(pm, k, k, 8192u); // patch merge selects top-left token + for (unsigned h = 0; h < 2u; ++h) + for (unsigned row = 0; row < 64u; ++row) + for (unsigned col = 0; col < 64u; ++col) { + const unsigned r = row & 15u, c = col & 15u; + const size_t offset = bias + h * 8192u + + 512u * (4u * (col >> 4) + (row >> 4)) + + 64u * (r & 7u) + 16u * ((c & 7u) >> 1) + + 2u * (c & 1u) + 4u * (r >> 3) + 8u * (c >> 3); + set_half(offset, 0x2400u); // 1/64 position average + } + hip.check(hip.hipMemcpy(dev_weights.pointer, weights.data(), weights.size(), + hipMemcpyHostToDevice), "hipMemcpy(K nonzero weights)"); + hip.check(hip.hipModuleLaunchKernel(prep, static_cast((prep_items + 127u) / 128u), 1, 1, + 128, 1, 1, 0, nullptr, arguments, nullptr), "hipModuleLaunchKernel(K nonzero prep)"); + hip.check(hip.hipDeviceSynchronize(), "hipDeviceSynchronize(K nonzero prep)"); + std::fill(full.begin(), full.end(), 0x7e00u); + std::fill(merged.begin(), merged.end(), 0x7e00u); + hip.check(hip.hipMemcpy(dev_full.pointer, full.data(), full.size() * sizeof(uint16_t), + hipMemcpyHostToDevice), "hipMemcpy(K nonzero full sentinel)"); + hip.check(hip.hipMemcpy(dev_merged.pointer, merged.data(), merged.size() * sizeof(uint16_t), + hipMemcpyHostToDevice), "hipMemcpy(K nonzero merged sentinel)"); + hip.check(hip.hipModuleLaunchKernel(transformer, 1, 1, 1, 32, 1, 4, 0, + nullptr, arguments, nullptr), "hipModuleLaunchKernel(K nonzero transformer)"); + hip.check(hip.hipDeviceSynchronize(), "hipDeviceSynchronize(K nonzero transformer)"); + hip.check(hip.hipMemcpy(full.data(), dev_full.pointer, full.size() * sizeof(uint16_t), + hipMemcpyDeviceToHost), "hipMemcpy(K nonzero full result)"); + hip.check(hip.hipMemcpy(merged.data(), dev_merged.pointer, merged.size() * sizeof(uint16_t), + hipMemcpyDeviceToHost), "hipMemcpy(K nonzero merged result)"); + size_t changed = 0; + for (size_t i = 0; i < full.size(); ++i) { + if ((full[i] & 0x7c00u) == 0x7c00u) + throw std::runtime_error("K nonzero full output has NaN/Inf or unwritten sentinel at " + std::to_string(i)); + changed += full[i] != input[i]; + } + for (unsigned mt = 0; mt < 16u; ++mt) { + const unsigned my = mt >> 2, mx = mt & 3u; + const size_t full_base = (8u * (2u * my) + 2u * mx) * 64u; + for (unsigned channel = 0; channel < 64u; ++channel) { + const size_t i = mt * 64u + channel; + if (merged[i] != full[full_base + channel]) + throw std::runtime_error("K patch merge identity mismatch at " + std::to_string(i)); + } + } + if (!changed) throw std::runtime_error("K nonzero weights had no visible transformer effect"); + if (!args.fixture_dir.empty()) { + std::filesystem::create_directories(wide(args.fixture_dir)); + auto save = [&](const char* name, const void* data, size_t bytes) { + const auto path = std::filesystem::path(wide(args.fixture_dir)) / name; + std::ofstream file(path, std::ios::binary); + if (!file.write(static_cast(data), static_cast(bytes))) + throw std::runtime_error(std::string("Failed writing K fixture ") + name); + }; + save("input.bin", input.data(), input.size() * sizeof(uint16_t)); + save("weights.bin", weights.data(), weights.size()); + save("full.bin", full.data(), full.size() * sizeof(uint16_t)); + save("merged.bin", merged.data(), merged.size() * sizeof(uint16_t)); + } + std::printf("PASS K_NONZERO architecture=gfx1201 changed_elements=%zu full_elements=%zu merged_elements=%zu fixture_saved=%d numerical_reference_pending=1\n", + changed, full.size(), merged.size(), !args.fixture_dir.empty()); return 0; } catch (const std::exception& error) { std::fprintf(stderr, "FAIL K_MODULE %s\n", error.what());