Execute gfx1201 K enc1 and validate against numpy reference

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")

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.

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.

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.'

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/<profile> 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"

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<unsigned>(std::stoul(value, &end));

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)

tools/windows/k_module_probe.cpp

@@ -1,4 +1,37 @@
 #include "hip_api.h"
+#include <fstream>
+
+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<uint16_t>(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<uint8_t> weights(weight_bytes, 0);
+        auto set_half = [&](size_t byte, uint16_t bits) {
+            weights.at(byte) = static_cast<uint8_t>(bits);
+            weights.at(byte + 1) = static_cast<uint8_t>(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<uint16_t> input(full_elements), full(full_elements, 0x7e00u);
+        std::vector<uint16_t> merged(merged_elements, 0x7e00u);
+        for (size_t i = 0; i < input.size(); ++i)
+            input[i] = half_quarter(static_cast<int>(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<const uint8_t*>(dev_input.pointer);
+        params.full = static_cast<uint8_t*>(dev_full.pointer);
+        params.merged = static_cast<uint8_t*>(dev_merged.pointer);
+        params.weights = static_cast<const uint8_t*>(dev_weights.pointer);
+        void* arguments[] = {&params};
+        hip.check(hip.hipModuleLaunchKernel(prep, static_cast<unsigned>((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<unsigned>((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<const char*>(data), static_cast<std::streamsize>(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());