diff --git a/CMakeLists.txt b/CMakeLists.txt index 1685a4f..55a9211 100644 --- a/CMakeLists.txt +++ b/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 diff --git a/docs/windows-rdna4-port.md b/docs/windows-rdna4-port.md index 24a76eb..d59b49d 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; 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 -transformers run on gfx1201 and pass synthetic nonzero numpy reference checks. +requested E path passes with the user's DLLs. M6 is in progress: all 11 K layers +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 -build the d4r-patched Windows ZLUDA and implement same-frame Windows NGX Evaluate +Next gate: validate K against real DLSS captures with the built Windows ZLUDA, +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, diff --git a/scripts/windows/test-k-replay.ps1 b/scripts/windows/test-k-replay.ps1 new file mode 100644 index 0000000..2fcfef6 --- /dev/null +++ b/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 } diff --git a/tools/windows/diagnostic.h b/tools/windows/diagnostic.h index 3188b1c..312bfb3 100644 --- a/tools/windows/diagnostic.h +++ b/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; diff --git a/tools/windows/hip_api.h b/tools/windows/hip_api.h index 12ce53c..a465da5 100644 --- a/tools/windows/hip_api.h +++ b/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); diff --git a/tools/windows/k_replay_validate.py b/tools/windows/k_replay_validate.py new file mode 100644 index 0000000..8c8b3c4 --- /dev/null +++ b/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(' .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) diff --git a/tools/windows/native_replay_probe.cpp b/tools/windows/native_replay_probe.cpp new file mode 100644 index 0000000..183b56c --- /dev/null +++ b/tools/windows/native_replay_probe.cpp @@ -0,0 +1,123 @@ +#include "hip_api.h" +#include +#include +#include + +namespace { +std::vector 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(file), std::istreambuf_iterator()}; +} +struct Allocation { + d4r::diag::HipApi& hip; + int id; + size_t size; + void* device = nullptr; + Allocation(d4r::diag::HipApi& api, int identifier, const std::vector& 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 lines; + std::string line, kernel; + unsigned grid[3] = {1, 1, 1}, block[3] = {32, 1, 4}, shared = 0; + std::vector> 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(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((*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 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(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; + } +}