Validate all eleven native gfx1201 K layers with Windows replay

CMakeLists.txt

@@ -59,7 +59,7 @@ foreach(family k m)
   endforeach()
 endforeach()
 add_custom_target(d4r_native_gfx1201_compile ALL DEPENDS ${native_objects})
-foreach(probe hip_gfx1201_probe gfx12_wmma_probe k_module_probe cuda_driver_probe cuda_image_probe d3d12_hip_interop_probe hip_stream_lifecycle_probe ngx_cuda_init_probe)
+foreach(probe hip_gfx1201_probe gfx12_wmma_probe k_module_probe native_replay_probe cuda_driver_probe cuda_image_probe d3d12_hip_interop_probe hip_stream_lifecycle_probe ngx_cuda_init_probe)
   add_executable(d4r_${probe} tools/windows/${probe}.cpp)
   target_compile_definitions(d4r_${probe} PRIVATE WIN32_LEAN_AND_MEAN NOMINMAX __HIP_PLATFORM_AMD__)
   target_include_directories(d4r_${probe} PRIVATE "${D4R_HIP_ROOT}/include")
@@ -115,8 +115,10 @@ add_dependencies(d4r_k_module_probe d4r_native_gfx1201_compile)
 install(FILES "${probe_object}" DESTINATION bin)
 install(FILES "${wmma_object}" DESTINATION bin)
 install(FILES scripts/windows/test-windows-rdna4.ps1 DESTINATION .)
+install(FILES scripts/windows/test-k-replay.ps1 DESTINATION .)
 install(FILES tools/windows/k_layer_validate.py DESTINATION .)
-install(FILES kernels/tools/pwin_model.py DESTINATION reference)
+install(FILES tools/windows/k_replay_validate.py DESTINATION .)
+install(FILES kernels/tools/pwin_model.py kernels/tools/pwin_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

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; native NGX Create/Evaluate on the simple
-requested E path passes with the user's DLLs. M6 is in progress: K `enc1` and `enc2` prep and
+requested E path passes with the user's DLLs. M6 is in progress: all 11 K layers
-transformers run on gfx1201 and pass synthetic nonzero numpy reference checks.
+run on gfx1201 and pass synthetic nonzero NumPy reference checks.
 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, Create/Evaluate PASS; four finite synthetic frames; K/M 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 + upstream layout, max abs/relative error 0 |
-| M6: K layers, full transformer and image validation | Partial: `enc1`/`enc2` execute and match nonzero references; other layers/full network pending |
+| M6: K layers, full transformer and image validation | All 11 synthetic layer replays PASS; real weights, full network and K image validation pending |
 | M7: M FP16-equivalent baseline and full transformer | Pending |
 | M8: standalone Windows D3D12 NGX harness | Pending |
 | M9: OptiScaler integration, profiling, installation | Pending |
@@ -95,7 +95,7 @@ Windows CMake compiles all **11 K + 5 M** source modules with gfx1201 wave32,
 `-O3` and each source's upstream `d4r-build-flags` (including `-mcumode`). Both
 stable HIP 7.2 and isolated TheRock build successfully. The objects install to
 `experimental/k` and `experimental/m`, outside the ZLUDA override directory.
-Only `enc1`/`enc2` currently have GPU layer correctness checks.
+All 11 K layers now have GPU synthetic replay/reference checks (results below).
 
 2026-09-30 integration results on RX 9070 XT, Windows 11:
 
@@ -116,13 +116,53 @@ fix/TheRock; passing K does not resolve that leak. Build logs remain under each
 `build/windows-rdna4*` directory; stable GPU logs/fixtures remain under
 `test-results/upstream-stable-*-gpu.log` and `test-results/upstream-rdna4-stable`.
 
-Next gate: validate the remaining K layers against references/captures, then
+Next gate: validate K against real DLSS captures with the built Windows ZLUDA,
-build the d4r-patched Windows ZLUDA and implement same-frame Windows NGX Evaluate
+then implement same-frame Windows NGX Evaluate
 with explicit D3D12 queue submission ownership. Full K precedes M and native
 FP8/performance work. No complete DLSS or OptiScaler game result is claimed.
 
 ## Native Windows ZLUDA build and Evaluate work (2026-09-30)
 
+### All K layers: native Windows synthetic replay (2026-09-30)
+
+`native_replay_probe` reads upstream `manifest.txt`/`args.bin`/`alloc-N.bin`,
+relocates allocation pointers, runs the module's prep and main kernels, and
+saves GPU allocations for the existing `pwin_check`/`pwin_model` references.
+No NVIDIA DLL, PTX or weights are needed for these synthetic fixtures.
+Each fixture enables nonzero projections, attention, MLP and the layer's
+embedding/merge/head operations; it covers one 8x8 window with sparse weights.
+It does not replace validation with real network weights, multiple windows,
+shifted borders, textures or complete DLSS output.
+
+Run from the repository:
+
+```powershell
+powershell -NoProfile -ExecutionPolicy Bypass -File scripts/windows/test-k-replay.ps1
+```
+
+Hardware result: `test-results/k-all-layers-synthetic.zip`, **11/11 PASS**,
+32 native launches per layer, all output elements finite. The acceptance limit
+is max absolute error <= 0.001 and PSNR >= 60 dB; relative error uses a 0.001
+denominator floor. Validation stops on the first mismatch.
+
+| Layer | Max absolute error | Max relative error | PSNR (dB) |
+| --- | --- | --- | --- |
+| enc0 | 0.00048828125 | 0.0011682243 | 88.75 |
+| enc1 | 0 | 0 | exact |
+| enc2 | 0.0000305175781 | 0.0009319664 | 114.40 |
+| enc3 | 0 | 0 | exact |
+| enc4 | 0 | 0 | exact |
+| dec5 | 0 | 0 | exact |
+| dec4 | 0 | 0 | exact |
+| dec3 | 0 | 0 | exact |
+| dec2 | 0.0000305175781 | 0.0008741259 | 111.42 |
+| dec1 | 0.00048828125 | 0.0011086475 | 99.33 |
+| dec0 | 0.000122070312 | 0.0133333333 | 100.84 |
+
+The next K gate is capture/reference validation with actual DLSS weights and
+proof of native transformer launches during Evaluate. Native overrides must
+check the originating PTX identity before replacing a function.
+
 The local fork now builds natively under Windows, including d4r patches
 `0002`–`0007`. New `0008-native-windows-build.patch` preserves MSVC/Linux builds
 and adds LLVM-MinGW support: prebuilt LLVM selection, GNU delay-load flags,

scripts/windows/test-k-replay.ps1

@@ -0,0 +1,53 @@
+[CmdletBinding()]
+param(
+    [string]$HipRoot,
+    [string]$PackageRoot,
+    [string]$OutputDirectory,
+    [string[]]$Layers = @('enc0','enc1','enc2','enc3','enc4','dec5','dec4','dec3','dec2','dec1','dec0'),
+    [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/k-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 {
+    if ($Iterations -lt 1 -or $Iterations -gt 10000) { throw 'Iterations must be 1..10000' }
+    foreach ($layer in $Layers) {
+        if ($layer -notmatch '^(enc[0-4]|dec[0-5])$') { throw "Invalid K 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 'k_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 "K $layer fixture generation failed; see $folder" }
+        & (Join-Path $PackageRoot 'bin/d4r_native_replay_probe.exe') --hip-root $HipRoot `
+            --module (Join-Path $PackageRoot "experimental/k/dltss_pwin_${layer}_layer_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 "K $layer GPU replay failed; see $folder" }
+        & python (Join-Path $PackageRoot 'k_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 "K $layer reference mismatch; stopping before the next layer. See $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 "K replay bundle: $OutputDirectory.zip"
+}
+if (!$summary.passed) { exit 1 }

tools/windows/diagnostic.h

@@ -116,6 +116,7 @@ struct Args {
     std::string ngx_mode = "init";
     unsigned preset = 11;
     std::string fixture_dir;
+    std::string output_dir;
     std::string kernel_name = "enc1";
     unsigned iterations = 32;
     int device = -1;
@@ -142,6 +143,7 @@ struct Args {
                     throw std::runtime_error("preset must be 5 (E), 11 (K) or 13 (M)");
             }
             else if (key == "--fixture-dir") fixture_dir = value;
+            else if (key == "--output-dir") output_dir = value;
             else if (key == "--kernel-name") kernel_name = value;
             else if (key == "--iterations") {
                 size_t end = 0;

tools/windows/hip_api.h

@@ -25,6 +25,7 @@ struct HipApi {
     D4R_HIP_FUNCTION(hipModuleLoad);
     D4R_HIP_FUNCTION(hipModuleUnload);
     D4R_HIP_FUNCTION(hipModuleGetFunction);
+    D4R_HIP_FUNCTION(hipModuleGetGlobal);
     D4R_HIP_FUNCTION(hipModuleLaunchKernel);
     D4R_HIP_FUNCTION(hipDeviceSynchronize);
     D4R_HIP_FUNCTION(hipStreamCreateWithFlags);

tools/windows/k_replay_validate.py

@@ -0,0 +1,152 @@
+#!/usr/bin/env python3
+"""Generate public synthetic K replay fixtures and validate native outputs.
+
+Uses the upstream replay format and NumPy model for all eleven K layers.
+No NVIDIA weights, PTX or DLLs are used by this fixture generator.
+"""
+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 / 'pwin_check.py').exists():
+    REFERENCE = ROOT / 'kernels/tools'
+sys.path.insert(0, str(REFERENCE))
+import pwin_model as pm
+from pwin_check import LAYERS, model_window, Outputs, stored
+
+
+def generate(layer, directory):
+    kind, heads, channels, other, extra = LAYERS[layer]
+    cin = 16 if kind == 'enc0' else channels
+    low = kind in ('dec', 'dec0')
+    cout = other if kind in ('enc', 'enc0') else channels
+    layout = pm.Layout(heads, channels, cout)
+    core_base = 2 * channels + 2 * 16 * channels if kind == 'enc0' else 0
+    if low:
+        core_base = 2 * other * 4 * channels + 8 * channels
+    weight_bytes = core_base + layout.PMB + 2 * layout.NPA * heads
+    weights = np.zeros(weight_bytes, np.uint8)
+
+    def half(offset, value):
+        weights[offset:offset + 2] = np.array([value], np.float16).view(np.uint8)
+
+    def matrix(base, k, n, stride, value=1.0):
+        half(base + 512 * (k // 16) + stride * (n // 16) + pm.frag_offset(k % 16, n % 16), value)
+
+    if kind == 'enc0':
+        for k in range(16):
+            matrix(2 * channels, k, k, 512)
+            matrix(2 * channels, k, k + 16, 512, .5)
+    if low:
+        for q in range(4):
+            for k in range(min(other, channels)):
+                matrix(0, k, q * channels + k, 512 * (other // 16), .25)
+    for c in range(channels):
+        half(core_base + layout.G1 + 2 * c, 1)
+        half(core_base + layout.G2 + 2 * c, 1)
+        half(core_base + layout.BO + 2 * c, 1 / 64)
+        half(core_base + layout.B2 + 2 * c, 1 / 64)
+    position_only = kind in ('enc0', 'dec0') or bool(extra)
+    for head in range(heads):
+        source_group = head % (channels // 32)
+        for j in range(3):
+            for k in range(32):
+                matrix(core_base + layout.QKV + layout.HEAD * head + 2048 * (3 * source_group + j),
+                       k, k, 1024, .5 if j == 2 else .25)
+        for k in range(32):
+            matrix(core_base + layout.WO + 64 * channels * head, k,
+                   32 * source_group + k, 1024, .5)
+        if position_only:
+            for row in range(64):
+                for col in range(64):
+                    r, c = row % 16, col % 16
+                    offset = core_base + layout.BIAS + head * 8192 + 512 * (4 * (col // 16) + row // 16)
+                    offset += 64 * (r & 7) + 16 * ((c & 7) >> 1) + 2 * (c & 1) + 4 * (r >> 3) + 8 * (c >> 3)
+                    half(offset, 1 / 64)
+    for group in range(channels // 32):
+        for k in range(32):
+            matrix(core_base + layout.W1 + 64 * channels * group, 32 * group + k, k, 32 * channels, .5)
+            matrix(core_base + layout.W2 + 64 * channels * group, k, 32 * group + k, 1024, .5)
+    if kind in ('enc', 'enc0'):
+        for n in range(cout):
+            matrix(core_base + layout.PM, n % channels + channels * ((n // channels) % 4),
+                   n, 128 * channels, .5)
+    if kind == 'dec0':
+        for n in range(40):
+            matrix(core_base + layout.PM + 2 * 48, n % channels, n, 1024, .5)
+            half(core_base + layout.PM + 2 * n, 1 / 64)
+
+    input_elements = 16 * other if low else 64 * cin
+    x = ((np.arange(input_elements) % 9 - 4) / 4).astype(np.float16)
+    skip = ((np.arange(64 * channels) % 7 - 3) / 8).astype(np.float16)
+    output_channels = 40 if kind == 'dec0' else channels
+    allocations = {0: x.view(np.uint8), 1: skip.view(np.uint8), 2: weights}
+    if kind in ('enc0', 'enc'):
+        allocations[3] = np.full(16 * cout, np.nan, np.float16).view(np.uint8)
+        allocations[4] = np.full(64 * channels, np.nan, np.float16).view(np.uint8)
+    else:
+        allocations[3] = np.full(64 * output_channels, np.nan, np.float16).view(np.uint8)
+    parameters = bytearray(176)
+    struct.pack_into('<ii', parameters, 0, 8, 8)
+    pointers = {8: 0, 64: 2, 24: 3}
+    if low:
+        pointers[16] = 1
+    if kind in ('enc0', 'enc'):
+        pointers[32] = 4
+    bases = {i: 0x100000000 + i * 0x1000000 for i in allocations}
+    for offset, identifier in pointers.items():
+        struct.pack_into('<Q', parameters, offset, bases[identifier])
+    directory.mkdir(parents=True, exist_ok=True)
+    for identifier, data in allocations.items():
+        data.tofile(directory / f'alloc-{identifier}.bin')
+    (directory / 'args.bin').write_bytes(parameters)
+    lines = [f'kernel dltss_pwin_{layer}_layer', 'launch 1 1 1 32 1 4 0']
+    lines += [f'alloc {i} {bases[i]:016x} {data.size}' for i, data in allocations.items()]
+    lines += [f'pointer {offset} {i} 0' for offset, i in pointers.items()]
+    (directory / 'manifest.txt').write_text('\n'.join(lines) + '\n', encoding='utf-8')
+    print(f'PASS K_FIXTURE kernel={layer} weights={weight_bytes} public_synthetic=1')
+
+
+def validate(layer, directory, output):
+    dump = pm.Dump(str(directory))
+    model, extra = model_window(dump, layer, 0, 0)
+    expected, actual = stored(dump, Outputs(dump, str(output)), layer, 0, 0, model, extra)
+    if not np.isfinite(expected).all() or not np.isfinite(actual).all():
+        raise RuntimeError('K replay contains NaN/Inf or unwritten output')
+    difference = np.abs(actual.astype(np.float64) - expected.astype(np.float64))
+    maximum = float(difference.max())
+    relative = float((difference / np.maximum(np.abs(expected.astype(np.float64)), 1e-3)).max())
+    rms = float(np.sqrt(np.mean(difference ** 2)))
+    peak = max(1., float(np.abs(expected.astype(np.float64)).max()))
+    psnr = math.inf if not rms else 20 * math.log10(peak / rms)
+    print(f'K_REPLAY_REFERENCE kernel={layer} elements={actual.size} max_abs={maximum:.9g} '
+          f'max_rel={relative:.9g} rms={rms:.9g} psnr_db={psnr:.6g}')
+    if maximum > .001 or psnr < 60:
+        worst = int(difference.argmax())
+        raise RuntimeError(f'K mismatch at {worst}: reference={expected[worst]}, gpu={actual[worst]}')
+    print(f'PASS K_REPLAY_REFERENCE kernel={layer} architecture=gfx1201 nan_inf=0')
+
+
+if __name__ == '__main__':
+    parser = argparse.ArgumentParser()
+    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)
+    args = parser.parse_args()
+    try:
+        if args.mode == 'generate':
+            generate(args.layer, args.fixture_dir)
+        else:
+            if not args.output_dir:
+                parser.error('validate requires --output-dir')
+            validate(args.layer, args.fixture_dir, args.output_dir)
+    except Exception as error:
+        print(f'FAIL K_REPLAY {error}', file=sys.stderr)
+        raise SystemExit(5)

tools/windows/native_replay_probe.cpp

@@ -0,0 +1,123 @@
+#include "hip_api.h"
+#include <fstream>
+#include <sstream>
+#include <memory>
+
+namespace {
+std::vector<uint8_t> read_file(const std::filesystem::path& path) {
+    std::ifstream file(path, std::ios::binary);
+    if (!file) throw std::runtime_error("Cannot read replay file " + d4r::diag::utf8(path.c_str()));
+    return {std::istreambuf_iterator<char>(file), std::istreambuf_iterator<char>()};
+}
+struct Allocation {
+    d4r::diag::HipApi& hip;
+    int id;
+    size_t size;
+    void* device = nullptr;
+    Allocation(d4r::diag::HipApi& api, int identifier, const std::vector<uint8_t>& bytes)
+        : hip(api), id(identifier), size(bytes.size()) {
+        hip.check(hip.hipMalloc(&device, size), "hipMalloc(replay)");
+        try { hip.check(hip.hipMemcpy(device, bytes.data(), size, hipMemcpyHostToDevice), "hipMemcpy(replay upload)"); }
+        catch (...) { (void)hip.hipFree(device); device = nullptr; throw; }
+    }
+    ~Allocation() { if (device) (void)hip.hipFree(device); }
+};
+}
+
+int main(int argc, char** argv) {
+    using namespace d4r::diag;
+    start();
+    try {
+        Args args(argc, argv);
+        if (args.module.empty() || args.fixture_dir.empty() || args.output_dir.empty())
+            throw std::runtime_error("--module, --fixture-dir and --output-dir are required");
+        HipApi hip(args.hip_root);
+        hipDeviceProp_t props{};
+        hip.select_gfx1201(args.device, props);
+        const std::filesystem::path input(wide(args.fixture_dir)), output(wide(args.output_dir));
+        auto parameters = read_file(input / L"args.bin");
+        if (parameters.empty() || parameters.size() > 4096) throw std::runtime_error("Invalid replay parameter size");
+        std::ifstream manifest(input / L"manifest.txt");
+        if (!manifest) throw std::runtime_error("Replay manifest is missing");
+        std::vector<std::string> lines;
+        std::string line, kernel;
+        unsigned grid[3] = {1, 1, 1}, block[3] = {32, 1, 4}, shared = 0;
+        std::vector<std::unique_ptr<Allocation>> allocations;
+        while (std::getline(manifest, line)) {
+            lines.push_back(line);
+            std::istringstream stream(line);
+            std::string kind;
+            stream >> kind;
+            if (kind == "kernel") stream >> kernel;
+            else if (kind == "launch") stream >> grid[0] >> grid[1] >> grid[2] >> block[0] >> block[1] >> block[2] >> shared;
+            else if (kind == "texture" || kind == "surface") throw std::runtime_error("This replay requires a texture/surface recreation backend");
+            else if (kind == "alloc") {
+                int id = -1; std::string base; size_t size = 0;
+                stream >> id >> base >> size;
+                if (!stream || id < 0 || size == 0 || size > props.totalGlobalMem / 2)
+                    throw std::runtime_error("Invalid replay allocation record");
+                for (const auto& allocation : allocations)
+                    if (allocation->id == id) throw std::runtime_error("Duplicate replay allocation ID");
+                auto bytes = read_file(input / ("alloc-" + std::to_string(id) + ".bin"));
+                if (bytes.size() != size) throw std::runtime_error("Replay allocation size mismatch");
+                allocations.push_back(std::make_unique<Allocation>(hip, id, bytes));
+            }
+        }
+        if (kernel.empty() || allocations.empty() || !grid[0] || !grid[1] || !grid[2])
+            throw std::runtime_error("Incomplete replay manifest");
+        for (const auto& record : lines) {
+            std::istringstream stream(record); std::string kind;
+            stream >> kind;
+            if (kind != "pointer") continue;
+            size_t offset = 0, delta = 0; int id = -1;
+            stream >> offset >> id >> delta;
+            if (!stream || offset > parameters.size() || parameters.size() - offset < 8)
+                throw std::runtime_error("Invalid replay pointer offset");
+            auto allocation = std::find_if(allocations.begin(), allocations.end(), [&](const auto& a) { return a->id == id; });
+            if (allocation == allocations.end() || delta >= (*allocation)->size)
+                throw std::runtime_error("Invalid replay pointer relocation");
+            const uint64_t pointer = reinterpret_cast<uintptr_t>((*allocation)->device) + delta;
+            std::memcpy(parameters.data() + offset, &pointer, 8);
+        }
+        hipModule_t module = nullptr;
+        hip.check(hip.hipModuleLoad(&module, args.module.c_str()), "hipModuleLoad(replay gfx1201)");
+        struct Cleanup { HipApi& hip; hipModule_t module; ~Cleanup() { (void)hip.hipModuleUnload(module); } } cleanup{hip, module};
+        auto scalar = [&](const char* name, unsigned fallback) {
+            hipDeviceptr_t address = nullptr; size_t size = 0;
+            if (hip.hipModuleGetGlobal(&address, &size, module, name) != hipSuccess) return fallback;
+            if (size != 4) throw std::runtime_error(std::string("Invalid native metadata ") + name);
+            unsigned value = 0;
+            hip.check(hip.hipMemcpy(&value, address, 4, hipMemcpyDeviceToHost), name);
+            return value ? value : fallback;
+        };
+        block[2] = scalar("d4r_block_z", block[2]);
+        const unsigned persistentGrid = scalar("d4r_grid_x", 0);
+        if (persistentGrid) { grid[0] = persistentGrid; grid[1] = grid[2] = 1; }
+        hipFunction_t function = nullptr, prep = nullptr;
+        hip.check(hip.hipModuleGetFunction(&function, module, kernel.c_str()), "hipModuleGetFunction(replay)");
+        void* launchArguments[] = {parameters.data()};
+        if (hip.hipModuleGetFunction(&prep, module, (kernel + "_prep").c_str()) == hipSuccess) {
+            const unsigned prepBlocks = scalar("d4r_prep_blocks", 0);
+            if (!prepBlocks) throw std::runtime_error("Native prep has no valid d4r_prep_blocks");
+            hip.check(hip.hipModuleLaunchKernel(prep, prepBlocks, 1, 1, 128, 1, 1, 0, nullptr, launchArguments, nullptr), "hipModuleLaunchKernel(replay prep)");
+        }
+        std::printf("REPLAY kernel=%s grid=%u,%u,%u block=%u,%u,%u allocations=%zu arguments=%zu\n",
+            kernel.c_str(), grid[0], grid[1], grid[2], block[0], block[1], block[2], allocations.size(), parameters.size());
+        for (unsigned iteration = 0; iteration < args.iterations; ++iteration)
+            hip.check(hip.hipModuleLaunchKernel(function, grid[0], grid[1], grid[2], block[0], block[1], block[2], shared,
+                nullptr, launchArguments, nullptr), "hipModuleLaunchKernel(replay)");
+        hip.check(hip.hipDeviceSynchronize(), "hipDeviceSynchronize(replay)");
+        std::filesystem::create_directories(output);
+        for (const auto& allocation : allocations) {
+            std::vector<uint8_t> bytes(allocation->size);
+            hip.check(hip.hipMemcpy(bytes.data(), allocation->device, bytes.size(), hipMemcpyDeviceToHost), "hipMemcpy(replay result)");
+            std::ofstream file(output / ("alloc-" + std::to_string(allocation->id) + ".bin"), std::ios::binary);
+            file.write(reinterpret_cast<const char*>(bytes.data()), std::streamsize(bytes.size()));
+            if (!file) throw std::runtime_error("Cannot save replay allocation");
+        }
+        std::printf("PASS NATIVE_REPLAY architecture=gfx1201 kernel=%s iterations=%u numerical_reference=pending\n", kernel.c_str(), args.iterations);
+        return 0;
+    } catch (const std::exception& error) {
+        std::fprintf(stderr, "FAIL NATIVE_REPLAY %s\n", error.what()); loaded_modules(); return 4;
+    }
+}