Add native Windows gfx1201 HIP and ZLUDA diagnostics

.gitignore

@@ -15,6 +15,9 @@
 /traces/
 /build/
 /dist/
+/.tools/
+/external/
+/test-results/
 *.trace
 *.dump
 *.log
@@ -33,4 +36,4 @@ videos/
 /kernels/out/
 *.hsaco
 *.ptx
-/kernels/extracted/
+/kernels/extracted/

CMakeLists.txt

@@ -0,0 +1,56 @@
+cmake_minimum_required(VERSION 3.24)
+project(d4r_windows_diagnostics LANGUAGES CXX)
+if(NOT WIN32 OR NOT CMAKE_SIZEOF_VOID_P EQUAL 8)
+  message(FATAL_ERROR "These diagnostics require native Windows x64")
+endif()
+set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+set(D4R_HIP_ROOT "$ENV{HIP_PATH}" CACHE PATH "One stable HIP SDK or TheRock installation")
+if(NOT D4R_HIP_ROOT AND EXISTS "C:/Program Files/AMD/ROCm/7.2")
+  set(D4R_HIP_ROOT "C:/Program Files/AMD/ROCm/7.2")
+endif()
+if(NOT EXISTS "${D4R_HIP_ROOT}/include/hip/hip_runtime_api.h")
+  message(FATAL_ERROR "Set D4R_HIP_ROOT to a Windows HIP SDK (headers and bin/clang++.exe required)")
+endif()
+find_program(D4R_HIP_CLANG NAMES clang++.exe PATHS "${D4R_HIP_ROOT}/bin" NO_DEFAULT_PATH REQUIRED)
+set(probe_object "${CMAKE_CURRENT_BINARY_DIR}/probe_gfx1201.hsaco")
+add_custom_command(OUTPUT "${probe_object}"
+  COMMAND "${D4R_HIP_CLANG}" -x hip --offload-arch=gfx1201
+    --offload-device-only --no-gpu-bundle-output -mno-wavefrontsize64
+    -nogpuinc -nogpulib -O2 "${CMAKE_CURRENT_SOURCE_DIR}/tools/windows/probe_kernel.hip"
+    -o "${probe_object}"
+  DEPENDS tools/windows/probe_kernel.hip VERBATIM)
+add_custom_target(d4r_probe_code_object ALL DEPENDS "${probe_object}")
+foreach(probe hip_gfx1201_probe cuda_driver_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")
+  target_link_libraries(d4r_${probe} PRIVATE version dbghelp)
+  if(MINGW)
+    target_link_options(d4r_${probe} PRIVATE -static)
+  endif()
+  if(MSVC)
+    target_compile_options(d4r_${probe} PRIVATE /W4 /EHsc)
+    set_property(TARGET d4r_${probe} PROPERTY MSVC_RUNTIME_LIBRARY "MultiThreaded$<$<CONFIG:Debug>:Debug>")
+  else()
+    target_compile_options(d4r_${probe} PRIVATE -Wall -Wextra)
+  endif()
+  install(TARGETS d4r_${probe} RUNTIME DESTINATION bin)
+endforeach()
+add_dependencies(d4r_hip_gfx1201_probe d4r_probe_code_object)
+install(FILES "${probe_object}" DESTINATION bin)
+install(FILES scripts/windows/test-windows-rdna4.ps1 DESTINATION .)
+install(FILES docs/windows-rdna4-port.md DESTINATION docs)
+include(CTest)
+add_test(NAME hip_gfx1201 COMMAND d4r_hip_gfx1201_probe
+  --hip-root "${D4R_HIP_ROOT}" --module "${probe_object}")
+set_tests_properties(hip_gfx1201 PROPERTIES TIMEOUT 90 LABELS "hardware;gfx1201")
+set(D4R_ZLUDA_ROOT "" CACHE PATH "Local Windows ZLUDA directory (nvcuda.dll)")
+if(EXISTS "${D4R_ZLUDA_ROOT}/nvcuda.dll")
+  foreach(context primary created)
+    add_test(NAME cuda_${context} COMMAND d4r_cuda_driver_probe
+      --hip-root "${D4R_HIP_ROOT}" --cuda-dll "${D4R_ZLUDA_ROOT}/nvcuda.dll"
+      --context ${context} --iterations 32)
+    set_tests_properties(cuda_${context} PROPERTIES TIMEOUT 180 LABELS "hardware;zluda")
+  endforeach()
+endif()

README.md

@@ -6,6 +6,11 @@ d4r runs NVIDIA's official DLSS Super Resolution library (`nvngx_dlss.dll`) in W
 
 It is an experiment. It has been tested on one GPU in a handful of games, it depends on unreleased patches to ZLUDA and vkd3d-proton.
 
+**Windows / RDNA4 development:** the `windows-rdna4` branch adds native Windows
+gfx1201 HIP and ZLUDA diagnostics. 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.
+
 See [supported games](SUPPORTED_GAMES.md) for the tested games and DLSS models.
 
 > **Not affiliated with NVIDIA or AMD.**

docs/windows-rdna4-port.md

@@ -0,0 +1,209 @@
+# Windows / RDNA4 port
+
+Branch: `windows-rdna4`. Target: Windows 11 x64, RX 9070 XT, **gfx1201**.
+Priority: correct K, correct M, native Windows, same-frame output, then speed.
+This is a work log for an implementation in progress, not a claim of DLSS support.
+
+## Milestone and gates
+
+Current milestone: M1/M2 passed on the physical RX 9070 XT; next: independent
+Windows D3D12/HIP interop and validated gfx12 WMMA, before NGX integration.
+Do not integrate NGX until the integer PTX workload is stable on the real GPU.
+
+| Gate | Status |
+| --- | --- |
+| M0: branch, audit, Windows build and diagnostics | Initial implementation complete |
+| M1: native HIP gfx1201 allocation, kernel, CPU verification | PASS, 32 iterations + guard verification |
+| M2: Windows ZLUDA integer PTX, primary and created contexts | PASS, 32 iterations each + guard verification |
+| M3: Windows NGX initialization and simple DLSS path | Pending |
+| M4: D3D12 / HIP external memory and fence round trip | Pending |
+| M5: independently validated gfx12 WMMA backend | Pending |
+| M6: K layers, full transformer and 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 |
+
+## Source baseline (checked 2026-09-29)
+
+* d4r: `ed1ab60`, main, 14 commits; https://github.com/countervolts/d4r.
+* ZLUDA master: `ee2f25a180099fa42f36b2346732e1f2470a03ad`,
+  2026-09-22. This is also the base named by d4r's existing patches.
+* Windows binary baseline: ZLUDA `v7-preview.11`,
+  `zluda-windows-ee2f25a.zip`, SHA256
+  `7788c1ed43385e62f0cc5f4d1aead24c92b671c02f1d685f63182540dea16d74`.
+* Local HIP SDK: `C:\Program Files\AMD\ROCm\7.2`, headers report
+  `7.2.60201`, commit `38d754472`. Compiler: AMD clang 21,
+  LLVM commit `590b9320a5be90e40268759c6203c01fde121e68`.
+* Keep stable SDK and TheRock installations separate. Select one directory per
+  build and test run; never mix headers, import libraries, runtime or COMGR.
+* CMake 3.31.6, Ninja 1.11.1 (Python package 1.11.1.4).
+* Windows host toolchain: LLVM-MinGW `20260922`, UCRT x64, SHA256
+  `e3ad77d117a4bea19a7a3b333341824d79a5a371004a10e25b8504e7b3047666`.
+  This is a portable fallback on this machine, where MSVC is not installed.
+  The device compiler remains the selected AMD HIP clang.
+
+## Architecture and platform boundaries
+
+Existing path:
+`D3D12 -> OptiScaler -> d4r shim -> NGX CUDA -> Wine nvcuda -> ZLUDA -> HIP`.
+The shim records Vulkan copies through vkd3d-proton interop. A patched vkd3d
+splits recording into submissions and gates the second part on a semaphore.
+The Wine bridge translates Windows/SysV ABIs, loads Linux libraries, chooses
+native kernels from KFD topology and performs fd-based Vulkan memory imports.
+
+Planned Windows path:
+`D3D12 -> OptiScaler -> d4r shim -> NGX CUDA -> Windows ZLUDA -> HIP gfx1201`.
+Platform-independent NGX parameter handling, formats, references, replay format
+and kernel identity checks should be retained. A Windows backend must use
+Win32 DLL loading, HIP device properties, shared D3D12 allocations and fences.
+Wine builtins, KFD, file descriptors, Proton and vkd3d are excluded from runtime.
+
+An NGX Evaluate call receives a **recording** command list, not a queue that can
+be submitted halfway through by the shim. Native D3D12 has no equivalent to the
+vkd3d split extension. Same-frame synchronization therefore needs an explicit
+submission boundary supplied by the harness and then OptiScaler/queue integration.
+A HIP wait inserted behind an unsplit D3D12 submission can deadlock. Merely
+mapping VRAM does not solve command recording/submission ownership.
+
+## Audit map and decisions
+
+* `README.md`, `docs/{architecture,building,native-kernels,performance}.md`:
+  Linux-tested baseline; upstream K translation can produce nonfinite results.
+  A successful NGX return alone is not transformer validation.
+* `tools/d4r_nvngx_shim.cpp`: Windows API/NGX entry points and C++ worker are
+  reusable; Vulkan resource conversion, markers and split-frame functions form
+  the platform boundary. Portable ini paths currently assume Wine translation.
+* `tools/wine_nvcuda_bridge.c`: loading (`dlopen`, pthread_once), KFD target
+  selection, descriptor bookkeeping, PTX manifests, replay/profile and CUDA
+  forwarding are interleaved. Preserve identities/instrumentation, replace
+  platform services rather than copy the Wine bridge into Windows.
+* `patches/zluda/0002`: DLSS arrays/textures/surfaces, PTX instructions, dumps;
+  `0003`: bounds and gfx11 MMA; `0004`: native override/prep, FP8, extra bitcode,
+  handoffs; `0005`: null texture reads. The gfx11 MMA optimization is explicitly
+  architecture-gated. It must not be enabled for gfx12 by changing that gate.
+* `patches/vkd3d-proton/0001`: submission split, not a portable D3D12 interface.
+* `kernels/k`: `pwin_common.h`, layer/position/wide templates and eleven entries.
+  Operands are replicated 16-half rows; accumulator row is `2*i + (lane>>4)`;
+  `operand_from_dt`, `wtile`, weight prep and stores depend on that convention.
+* `kernels/m`: `swin_common.h`, `swin_block.h`, five standalone entries;
+  FP8 weights expanded to FP16, `kslot`, `woff`, `bload`, `k32`, accumulator
+  reconstruction and xor shuffles need layout audit. Preserve rounding per k32.
+* `kernels/tex`: enc0 tail/dec0 head also use gfx11 WMMA and output shuffles;
+  changing only K/M would leave part of M incorrect. Generated PTX is private.
+* `kernels/build.sh`, texture build and `scripts/build_zluda_ptx_helpers.sh`:
+  Linux tool paths, shell, bitcode cleanup and target assumptions need a Windows
+  build path. Keep existing Linux scripts operational.
+* `kernels/tools/{pwin_model,swin_model,ptxsim,psnr,kernel_manifest}.py` and
+  `dump_runner.cpp`, `tools/replay_dlss_*`: use existing reference semantics
+  and capture relocation format; add finite checks and quantitative comparison.
+* `scripts/package_release.sh`, install/launch scripts and `packaging/`:
+  Linux release bundles Proton components. Windows diagnostics/package must be
+  separate and must never include NVIDIA DLLs or extracted NVIDIA PTX/weights.
+* Upstream ZLUDA already has Windows delay-loading for HIP 7 and fallback HIP 6.
+  Its Rust binding calls versioned `hipGetDevicePropertiesR0600`; a reported
+  gfx1201 crash is not evidence of an ABI bug without a reproducer/trace here.
+
+## gfx11 versus gfx12 WMMA
+
+For wave32, gfx11 FP16 A/B are 16 halves per lane with replication across the
+two 16-lane groups. gfx12 uses 8 halves per lane without replication. The gfx12
+builtin has a `_gfx12` suffix; its FP32 accumulator is still eight floats, but
+rows are contiguous within the lane group, rather than gfx11 even/odd rows.
+The backend must explicitly convert operands **and** accumulator distribution,
+or update all dependent loads/stores and intermediate fragments together.
+Start with FP16 / FP32 accumulate. M native FP8 is a later measured optimization.
+Compile success is not proof of lane layout or numerical correctness.
+
+## External interop blockers
+
+The installed HIP headers define D3D12Resource/D3D12Heap external memory and
+D3D12Fence external semaphore types. Their presence is not proof the Windows
+runtime implements them. Validate imports, mapping, GPU waits/signals and
+cleanup using an independent round trip before modifying the shim.
+Use HIP/DXGI device identity to prevent importing allocations across adapters.
+NT HANDLE ownership must be explicit; Win32 import is not fd ownership transfer.
+Test repeated create/destroy cycles as upstream noted a mapped-buffer leak on
+Linux ROCm. CPU verification copies are allowed in tests, never in the fast path.
+
+## Authoritative references
+
+* [AMD HIP SDK Windows support matrix](https://rocm.docs.amd.com/projects/install-on-windows/en/latest/reference/system-requirements.html)
+  lists RX 9070 XT as gfx1201 with runtime and SDK support, release 7.2.0.
+* [ZLUDA Windows SDK setup](https://github.com/vosen/ZLUDA/blob/ee2f25a180099fa42f36b2346732e1f2470a03ad/docs/src/hip_sdk.md)
+  distinguishes official SDK and TheRock nightlies.
+* [ZLUDA build](https://github.com/vosen/ZLUDA/blob/ee2f25a180099fa42f36b2346732e1f2470a03ad/docs/src/building.md),
+  [trace instructions](https://github.com/vosen/ZLUDA/blob/ee2f25a180099fa42f36b2346732e1f2470a03ad/docs/src/troubleshooting.md).
+* [LLVM AMDGPU target guide](https://llvm.org/docs/AMDGPUUsage.html),
+  [Clang AMDGPU builtins](https://github.com/llvm/llvm-project/blob/main/clang/include/clang/Basic/BuiltinsAMDGPU.td).
+* [AMD RDNA4 matrix cores](https://gpuopen.com/learn/using_matrix_core_amd_rdna4/),
+  [RDNA4 WMMA guide](https://gpuopen.com/learn/wmma-guide-amd-rdna-4-gpus-part-1/),
+  [RDNA4 ISA](https://www.amd.com/content/dam/amd/en/documents/radeon-tech-docs/instruction-set-architectures/rdna4-instruction-set-architecture.pdf).
+* [TheRock Windows development](https://github.com/ROCm/TheRock/blob/main/docs/development/windows_support.md).
+
+## Build, test, package
+
+Run from the repository root in PowerShell. The setup installs tools only inside
+`.tools`, verifies downloaded SHA256 values, and does not install a driver/SDK.
+
+```powershell
+powershell -NoProfile -File scripts/windows/setup-windows-tools.ps1
+powershell -NoProfile -File scripts/windows/build-windows-rdna4.ps1
+powershell -NoProfile -File scripts/windows/test-windows-rdna4.ps1
+```
+
+The build accepts `-HipRoot`, `-ZludaRoot`, `-ToolchainRoot`, `-BuildDirectory`
+and `-InstallDirectory`. MSVC users can invoke CMake from a Developer Shell
+without the portable MinGW compiler. Equivalent direct build:
+
+```powershell
+cmake -S . -B build/windows-rdna4 -G Ninja -DCMAKE_BUILD_TYPE=RelWithDebInfo "-DD4R_HIP_ROOT=C:/Program Files/AMD/ROCm/7.2" "-DD4R_ZLUDA_ROOT=C:/path/to/zluda"
+cmake --build build/windows-rdna4
+ctest --test-dir build/windows-rdna4 --output-on-failure
+cmake --install build/windows-rdna4 --prefix dist/windows-rdna4-diagnostics
+```
+
+The one-command diagnostic runner is also installed at
+`dist/windows-rdna4-diagnostics/test-windows-rdna4.ps1`. It requires the selected
+HIP SDK and ZLUDA directory; the developer setup auto-discovers local defaults.
+Success: exit 0, all three tests passed in `summary.json`, `PASS HIP` and
+`PASS CUDA` in stdout. Failure: send the **single printed ZIP path**; it contains
+stdout/stderr, timeout/exception exit code, DLL paths/versions/hashes, driver/OS
+inventory and a minidump for an unhandled exception. Upstream trace is attempted
+automatically on CUDA failure when `zluda.exe` is available.
+
+The HIP module compiler uses the selected SDK, fixed `--offload-arch=gfx1201`,
+wave32 and no host/device library dependency. Host code uses the installed HIP
+header ABI and loads that SDK's `amdhip64_7.dll` by absolute path. CUDA loads the
+provided ZLUDA DLL by absolute path, matches the adapter against HIP PCI identity,
+tests two context lifecycles, JITs embedded PTX and verifies an integer pattern
+and out-of-range sentinels. No CUDA SDK is required.
+
+No NVIDIA DLL is needed for M1/M2. Later NGX tests accept user-supplied absolute
+paths; NVIDIA software is never downloaded or added to git by these scripts.
+
+## Known limitations and results
+
+2026-09-29: installed `hipInfo.exe` passed device enumeration in this workspace:
+RX 9070 XT, gfx1201, warpSize 32, 15.92 GiB total, 15.77 GiB free.
+2026-09-29 hardware results: Windows 11 Pro 10.0.26200, AMD display driver
+`32.0.31041.1004`, HIP runtime/driver API `70260201`. HIP struct size 1472 bytes;
+physical device PCI `0000:03:00`. Native HIP module: 32 successful iterations,
+4099 elements with changing launch counts and seeds, guard verified. CUDA PTX:
+32 iterations each in primary/created contexts, same pattern/guards verified,
+all context/module/memory teardown calls successful. ZLUDA binary SHA256
+`51dd32dc116a6c14c7a4bf5acf620bba0546ca7ec5f8a6e2148dd68f1714b165`.
+Logs: local `test-results/m1-m2.zip`; ignored by git.
+
+Local proprietary inputs supplied by the user: `nvngx_dlss.dll` 310.9.1.0,
+`_nvngx.dll` 32.0.16.1714. Neither is packaged or tracked.
+
+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.
+
+Environment-specific issues resolved: sandbox disallows writes to `.git` and
+launching the MinGW child compiler, so branch/commit/build require scoped tool
+approval in this managed environment. HIP_PATH ends in a backslash; build script
+normalizes it before PowerShell 5 argument quoting. Test runner uses .NET Process
+to retain reliable exit codes and concurrently drain both diagnostic streams.
+No K/M/DLSS/OptiScaler support is claimed by this first milestone.

scripts/windows/build-windows-rdna4.ps1

@@ -0,0 +1,56 @@
+[CmdletBinding()]
+param(
+    [string]$HipRoot = $env:HIP_PATH,
+    [string]$ZludaRoot,
+    [string]$ToolchainRoot,
+    [string]$BuildDirectory,
+    [string]$InstallDirectory
+)
+$ErrorActionPreference = 'Stop'
+$repo = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..'))
+if (!$HipRoot) { $HipRoot = 'C:\Program Files\AMD\ROCm\7.2' }
+if (!$BuildDirectory) { $BuildDirectory = Join-Path $repo 'build/windows-rdna4' }
+if (!$InstallDirectory) { $InstallDirectory = Join-Path $repo 'dist/windows-rdna4-diagnostics' }
+$HipRoot = [IO.Path]::GetFullPath($HipRoot).TrimEnd('\', '/')
+$BuildDirectory = [IO.Path]::GetFullPath($BuildDirectory).TrimEnd('\', '/')
+$InstallDirectory = [IO.Path]::GetFullPath($InstallDirectory).TrimEnd('\', '/')
+if (!$ZludaRoot -and (Test-Path (Join-Path $repo '.tools/zluda/zluda/nvcuda.dll'))) {
+    $ZludaRoot = Join-Path $repo '.tools/zluda/zluda'
+}
+if (!$ToolchainRoot -and (Test-Path (Join-Path $repo '.tools/llvm-mingw-20260922-ucrt-x86_64'))) {
+    $ToolchainRoot = Join-Path $repo '.tools/llvm-mingw-20260922-ucrt-x86_64'
+}
+$cmakeCommand = Get-Command cmake -ErrorAction SilentlyContinue
+$cmake = if ($cmakeCommand) { $cmakeCommand.Source } else { Join-Path $repo '.tools/python/cmake/data/bin/cmake.exe' }
+$ninjaCommand = Get-Command ninja -ErrorAction SilentlyContinue
+$ninja = if ($ninjaCommand) { $ninjaCommand.Source } else { Join-Path $repo '.tools/python/ninja/data/bin/ninja.exe' }
+if (!(Test-Path $ninja)) { $ninja = Join-Path $repo '.tools/python/bin/ninja.exe' }
+if (!(Test-Path $cmake) -or !(Test-Path $ninja)) {
+    throw 'CMake >=3.24 and Ninja required. See docs/windows-rdna4-port.md for local tool setup.'
+}
+New-Item -ItemType Directory -Force $BuildDirectory | Out-Null
+$originalPath = $env:PATH
+try {
+    $configure = @('--fresh', '-S', $repo, '-B', $BuildDirectory, '-G', 'Ninja', '-DCMAKE_BUILD_TYPE=RelWithDebInfo',
+        "-DCMAKE_MAKE_PROGRAM=$ninja", "-DD4R_HIP_ROOT=$HipRoot", "-DCMAKE_INSTALL_PREFIX=$InstallDirectory")
+    if ($ToolchainRoot) {
+        $compiler = Join-Path $ToolchainRoot 'bin/x86_64-w64-mingw32-clang++.exe'
+        if (!(Test-Path $compiler)) { throw "Compiler missing: $compiler" }
+        $env:PATH = "$(Join-Path $ToolchainRoot 'bin');$originalPath"
+        $configure += "-DCMAKE_CXX_COMPILER=$compiler"
+    }
+    if ($ZludaRoot) { $configure += "-DD4R_ZLUDA_ROOT=$ZludaRoot" }
+    $ErrorActionPreference = 'Continue'
+    & $cmake @configure 2>&1 | Tee-Object (Join-Path $BuildDirectory 'configure.log')
+    $ErrorActionPreference = 'Stop'
+    if ($LASTEXITCODE) { throw "CMake configure exit=$LASTEXITCODE" }
+    $ErrorActionPreference = 'Continue'
+    & $cmake --build $BuildDirectory --parallel 4 2>&1 | Tee-Object (Join-Path $BuildDirectory 'build.log')
+    $ErrorActionPreference = 'Stop'
+    if ($LASTEXITCODE) { throw "Build exit=$LASTEXITCODE" }
+    $ErrorActionPreference = 'Continue'
+    & $cmake --install $BuildDirectory 2>&1 | Tee-Object (Join-Path $BuildDirectory 'install.log')
+    $ErrorActionPreference = 'Stop'
+    if ($LASTEXITCODE) { throw "Install exit=$LASTEXITCODE" }
+    Write-Host "Built diagnostics: $InstallDirectory"
+} finally { $env:PATH = $originalPath }

scripts/windows/setup-windows-tools.ps1

@@ -0,0 +1,31 @@
+[CmdletBinding()]
+param()
+$ErrorActionPreference = 'Stop'
+$repo = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..'))
+$localTools = Join-Path $repo '.tools'
+$downloadDirectory = Join-Path $localTools 'downloads'
+New-Item -ItemType Directory -Force $downloadDirectory | Out-Null
+function Get-VerifiedArchive([string]$Url, [string]$Name, [string]$Hash, [string]$Destination, [string]$Marker) {
+    if (Test-Path -LiteralPath $Marker) { Write-Host "Already present: $Marker"; return }
+    $archive = Join-Path $downloadDirectory $Name
+    if (!(Test-Path -LiteralPath $archive)) {
+        Invoke-WebRequest -UseBasicParsing -Uri $Url -OutFile $archive
+    }
+    if ((Get-FileHash -LiteralPath $archive -Algorithm SHA256).Hash -ne $Hash) {
+        throw "SHA256 mismatch: $archive (remove this download and retry)"
+    }
+    Expand-Archive -LiteralPath $archive -DestinationPath $Destination -Force
+    if (!(Test-Path -LiteralPath $Marker)) { throw "Archive layout changed: $Marker" }
+}
+Get-VerifiedArchive 'https://github.com/mstorsjo/llvm-mingw/releases/download/20260922/llvm-mingw-20260922-ucrt-x86_64.zip' `
+    'llvm-mingw.zip' 'e3ad77d117a4bea19a7a3b333341824d79a5a371004a10e25b8504e7b3047666' $localTools `
+    (Join-Path $localTools 'llvm-mingw-20260922-ucrt-x86_64/bin/x86_64-w64-mingw32-clang++.exe')
+Get-VerifiedArchive 'https://github.com/vosen/ZLUDA/releases/download/v7-preview.11/zluda-windows-ee2f25a.zip' `
+    'zluda.zip' '7788c1ed43385e62f0cc5f4d1aead24c92b671c02f1d685f63182540dea16d74' (Join-Path $localTools 'zluda') `
+    (Join-Path $localTools 'zluda/zluda/nvcuda.dll')
+if (!(Test-Path (Join-Path $localTools 'python/cmake/data/bin/cmake.exe')) -or
+    !(Test-Path (Join-Path $localTools 'python/bin/ninja.exe'))) {
+    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" }
+}
+Write-Host 'Tools ready. Run scripts/windows/build-windows-rdna4.ps1.'

scripts/windows/test-windows-rdna4.ps1

@@ -0,0 +1,127 @@
+[CmdletBinding()]
+param(
+    [string]$HipRoot = $env:HIP_PATH,
+    [string]$ZludaRoot,
+    [string]$PackageRoot,
+    [string]$OutputDirectory,
+    [int]$Iterations = 32,
+    [int]$TimeoutSeconds = 180
+)
+$ErrorActionPreference = 'Stop'
+if (!$PackageRoot) {
+    if (Test-Path (Join-Path $PSScriptRoot 'bin/d4r_hip_gfx1201_probe.exe')) { $PackageRoot = $PSScriptRoot }
+    else { $PackageRoot = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../../dist/windows-rdna4-diagnostics')) }
+}
+$PackageRoot = [IO.Path]::GetFullPath($PackageRoot)
+if (!$HipRoot) { $HipRoot = 'C:\Program Files\AMD\ROCm\7.2' }
+$HipRoot = [IO.Path]::GetFullPath($HipRoot)
+if (!$ZludaRoot) {
+    $localZluda = [IO.Path]::GetFullPath((Join-Path $PackageRoot '../../.tools/zluda/zluda'))
+    if (Test-Path (Join-Path $localZluda 'nvcuda.dll')) { $ZludaRoot = $localZluda }
+    else { throw 'Pass -ZludaRoot with the directory containing the Windows ZLUDA nvcuda.dll.' }
+}
+$ZludaRoot = [IO.Path]::GetFullPath($ZludaRoot)
+if (!$OutputDirectory) { $OutputDirectory = Join-Path $PackageRoot ('results/' + (Get-Date -Format 'yyyyMMdd-HHmmss-fff')) }
+$OutputDirectory = [IO.Path]::GetFullPath($OutputDirectory)
+New-Item -ItemType Directory -Force $OutputDirectory | Out-Null
+$originalPath = $env:PATH
+$originalHip = $env:HIP_PATH
+$originalDiag = $env:D4R_DIAG_DIR
+$originalLog = $env:ZLUDA_LOG_DIR
+$originalCuda = $env:ZLUDA_CUDA_LIB
+$exitStatus = 1
+$summary = [ordered]@{utc=[DateTime]::UtcNow.ToString('o'); hipRoot=$HipRoot; zludaRoot=$ZludaRoot; tests=@()}
+function Quote-Argument([string]$Value) {
+    # All generated paths are absolute file/directory paths, with no trailing backslash.
+    if ($Value.Contains('"')) { throw 'Double quotes are not allowed in arguments' }
+    return '"' + $Value.TrimEnd('\') + '"'
+}
+function Invoke-Probe([string]$Name, [string]$Exe, [string[]]$Arguments) {
+    if (!(Test-Path -LiteralPath $Exe)) { throw "Executable missing: $Exe" }
+    $stdout = Join-Path $OutputDirectory "$Name.stdout.log"
+    $stderr = Join-Path $OutputDirectory "$Name.stderr.log"
+    $argumentLine = ($Arguments | ForEach-Object { Quote-Argument $_ }) -join ' '
+    $info = New-Object System.Diagnostics.ProcessStartInfo
+    $info.FileName = $Exe
+    $info.Arguments = $argumentLine
+    $info.UseShellExecute = $false
+    $info.CreateNoWindow = $true
+    $info.RedirectStandardOutput = $true
+    $info.RedirectStandardError = $true
+    $process = New-Object System.Diagnostics.Process
+    $process.StartInfo = $info
+    if (!$process.Start()) { throw "Process start failed: $Exe" }
+    # Consume both streams concurrently; a verbose JIT trace must not block its child.
+    $outTask = $process.StandardOutput.ReadToEndAsync()
+    $errTask = $process.StandardError.ReadToEndAsync()
+    $finished = $process.WaitForExit($TimeoutSeconds * 1000)
+    if (!$finished) { $process.Kill(); $process.WaitForExit() }
+    $code = $process.ExitCode
+    [IO.File]::WriteAllText($stdout, $outTask.GetAwaiter().GetResult())
+    [IO.File]::WriteAllText($stderr, $errTask.GetAwaiter().GetResult())
+    $process.Dispose()
+    $codeHex = if ($null -eq $code) { 'unknown' } else {
+        '0x{0:x8}' -f [BitConverter]::ToUInt32([BitConverter]::GetBytes([int]$code), 0)
+    }
+    $pass = $finished -and $code -eq 0 -and (Select-String -LiteralPath $stdout -Pattern '^PASS ' -Quiet)
+    $summary.tests += [ordered]@{name=$Name; passed=$pass; timedOut=(!$finished); exitCode=$code; exceptionOrExitHex=$codeHex}
+    Write-Host "$Name passed=$pass exit=$codeHex timeout=$(!$finished)"
+    return $pass
+}
+try {
+    if ($Iterations -lt 1 -or $Iterations -gt 10000) { throw 'Iterations must be 1..10000' }
+    if ($TimeoutSeconds -lt 1) { throw 'TimeoutSeconds must be positive' }
+    $env:HIP_PATH = $HipRoot
+    $env:PATH = "$(Join-Path $HipRoot 'bin');$ZludaRoot;$originalPath"
+    $env:D4R_DIAG_DIR = $OutputDirectory
+    $env:ZLUDA_LOG_DIR = Join-Path $OutputDirectory 'zluda-trace'
+    $env:ZLUDA_CUDA_LIB = Join-Path $ZludaRoot 'nvcuda.dll'
+    $files = @((Join-Path $HipRoot 'bin/amdhip64_7.dll'), (Join-Path $HipRoot 'bin/amd_comgr_3.dll'),
+        (Join-Path $ZludaRoot 'nvcuda.dll'), (Join-Path $PackageRoot 'bin/probe_gfx1201.hsaco'))
+    $summary.files = @($files | ForEach-Object {
+        if (!(Test-Path -LiteralPath $_)) { throw "Required file missing: $_" }
+        $file = Get-Item -LiteralPath $_
+        [ordered]@{path=$file.FullName; version=$file.VersionInfo.FileVersion;
+            sha256=(Get-FileHash -LiteralPath $_ -Algorithm SHA256).Hash}
+    })
+    try {
+        $summary.displayDrivers = @(Get-CimInstance Win32_VideoController |
+            Select-Object Name,DriverVersion,DriverDate,PNPDeviceID)
+        $summary.os = Get-CimInstance Win32_OperatingSystem | Select-Object Caption,Version,BuildNumber
+    } catch { $summary.inventoryError = $_.Exception.Message }
+    $bin = Join-Path $PackageRoot 'bin'
+    $hipOk = Invoke-Probe 'hip' (Join-Path $bin 'd4r_hip_gfx1201_probe.exe') @(
+        '--hip-root', $HipRoot, '--module', (Join-Path $bin 'probe_gfx1201.hsaco'), '--iterations', "$Iterations")
+    if ($hipOk) {
+        $cudaOk = $true
+        foreach ($context in @('primary', 'created')) {
+            $ok = Invoke-Probe "cuda-$context" (Join-Path $bin 'd4r_cuda_driver_probe.exe') @(
+                '--hip-root', $HipRoot, '--cuda-dll', (Join-Path $ZludaRoot 'nvcuda.dll'),
+                '--context', $context, '--iterations', "$Iterations")
+            $cudaOk = $cudaOk -and $ok
+            # On failure collect the upstream trace automatically if available.
+            if (!$ok -and (Test-Path (Join-Path $ZludaRoot 'zluda.exe'))) {
+                Invoke-Probe "cuda-$context-trace" (Join-Path $ZludaRoot 'zluda.exe') @(
+                    '--zluda-trace', '--', (Join-Path $bin 'd4r_cuda_driver_probe.exe'),
+                    '--hip-root', $HipRoot, '--cuda-dll', (Join-Path $ZludaRoot 'nvcuda.dll'),
+                    '--context', $context, '--iterations', '1') | Out-Null
+            }
+        }
+        if ($cudaOk) { $exitStatus = 0 }
+    }
+} catch {
+    $summary.error = $_.Exception.Message
+    Write-Warning $summary.error
+} finally {
+    $summary.passed = $exitStatus -eq 0
+    $summary | ConvertTo-Json -Depth 8 | Set-Content -LiteralPath (Join-Path $OutputDirectory 'summary.json') -Encoding UTF8
+    $env:PATH = $originalPath
+    $env:HIP_PATH = $originalHip
+    $env:D4R_DIAG_DIR = $originalDiag
+    $env:ZLUDA_LOG_DIR = $originalLog
+    $env:ZLUDA_CUDA_LIB = $originalCuda
+    $archive = "$OutputDirectory.zip"
+    Compress-Archive -LiteralPath $OutputDirectory -DestinationPath $archive -Force
+    Write-Host "Diagnostic bundle: $archive"
+}
+exit $exitStatus

tools/windows/cuda_driver_probe.cpp

@@ -0,0 +1,159 @@
+#include "hip_api.h"
+#include <algorithm>
+
+// CUDA Driver API ABI subset. No CUDA SDK, CUDA runtime or NVIDIA binary is needed.
+using CUdevice = int;
+using CUdeviceptr = uint64_t;
+using CUcontext = void*;
+using CUmodule = void*;
+using CUfunction = void*;
+struct CudaApi {
+    d4r::diag::Library library;
+#define D4R_CU(name, ret, ...) using name##_fn = ret(WINAPI*)(__VA_ARGS__); name##_fn name = library.symbol<name##_fn>(#name)
+    D4R_CU(cuInit, int, unsigned);
+    D4R_CU(cuDriverGetVersion, int, int*);
+    D4R_CU(cuDeviceGetCount, int, int*);
+    D4R_CU(cuDeviceGet, int, CUdevice*, int);
+    D4R_CU(cuDeviceGetName, int, char*, int, CUdevice);
+    D4R_CU(cuDeviceGetAttribute, int, int*, int, CUdevice);
+    D4R_CU(cuDevicePrimaryCtxRetain, int, CUcontext*, CUdevice);
+    D4R_CU(cuDevicePrimaryCtxRelease_v2, int, CUdevice);
+    D4R_CU(cuCtxSetCurrent, int, CUcontext);
+    D4R_CU(cuCtxCreate_v2, int, CUcontext*, unsigned, CUdevice);
+    D4R_CU(cuCtxDestroy_v2, int, CUcontext);
+    D4R_CU(cuCtxSynchronize, int);
+    D4R_CU(cuMemAlloc_v2, int, CUdeviceptr*, size_t);
+    D4R_CU(cuMemFree_v2, int, CUdeviceptr);
+    D4R_CU(cuMemcpyHtoD_v2, int, CUdeviceptr, const void*, size_t);
+    D4R_CU(cuMemcpyDtoH_v2, int, void*, CUdeviceptr, size_t);
+    D4R_CU(cuModuleLoadData, int, CUmodule*, const void*);
+    D4R_CU(cuModuleUnload, int, CUmodule);
+    D4R_CU(cuModuleGetFunction, int, CUfunction*, CUmodule, const char*);
+    D4R_CU(cuLaunchKernel, int, CUfunction, unsigned, unsigned, unsigned, unsigned, unsigned, unsigned,
+        unsigned, void*, void**, void**);
+    D4R_CU(cuGetErrorName, int, int, const char**);
+    D4R_CU(cuGetErrorString, int, int, const char**);
+#undef D4R_CU
+    explicit CudaApi(const std::string& path) : library(d4r::diag::wide(path)) {}
+    void check(int result, const char* call) {
+        const char* name = "unknown", *message = "unknown";
+        cuGetErrorName(result, &name);
+        cuGetErrorString(result, &message);
+        std::printf("CUDA %s -> %d (%s)\n", call, result, name ? name : "unknown");
+        if (result) throw std::runtime_error(std::string(call) + ": " + (message ? message : "unknown"));
+    }
+};
+
+static constexpr char ptx[] = R"PTX(
+.version 7.0
+.target sm_75
+.address_size 64
+.visible .entry d4r_ptx_pattern(
+    .param .u64 output, .param .u32 count, .param .u32 seed)
+{
+    .reg .pred %p;
+    .reg .b32 %r<8>;
+    .reg .b64 %rd<4>;
+    ld.param.u64 %rd0, [output];
+    ld.param.u32 %r0, [count];
+    ld.param.u32 %r1, [seed];
+    mov.u32 %r2, %ctaid.x;
+    mov.u32 %r3, %ntid.x;
+    mov.u32 %r4, %tid.x;
+    mad.lo.u32 %r5, %r2, %r3, %r4;
+    setp.ge.u32 %p, %r5, %r0;
+    @%p bra done;
+    mad.lo.u32 %r6, %r5, 1664525, %r1;
+    xor.b32 %r7, %r6, 0xa5a55a5a;
+    mul.wide.u32 %rd1, %r5, 4;
+    add.u64 %rd2, %rd0, %rd1;
+    st.global.u32 [%rd2], %r7;
+done:
+    ret;
+}
+)PTX";
+
+int main(int argc, char** argv)
+{
+    using namespace d4r::diag;
+    start();
+    try {
+        Args args(argc, argv);
+        if (args.cuda_dll.empty()) throw std::runtime_error("--cuda-dll is required (absolute ZLUDA nvcuda.dll path)");
+        if (args.context != "primary" && args.context != "created") throw std::runtime_error("context must be primary or created");
+        // Require the real HIP architecture, then match the CUDA device by PCI identity.
+        HipApi hip(args.hip_root);
+        hipDeviceProp_t props{};
+        hip.select_gfx1201(args.device, props);
+        CudaApi api(args.cuda_dll);
+        loaded_modules();
+        std::printf("STAGE cuInit\n");
+        api.check(api.cuInit(0), "cuInit");
+        int driver = 0, count = 0;
+        api.check(api.cuDriverGetVersion(&driver), "cuDriverGetVersion");
+        api.check(api.cuDeviceGetCount(&count), "cuDeviceGetCount");
+        std::printf("CUDA driver_api=%d devices=%d\n", driver, count);
+        CUdevice device = -1;
+        for (int ordinal = 0; ordinal < count; ++ordinal) {
+            CUdevice candidate;
+            api.check(api.cuDeviceGet(&candidate, ordinal), "cuDeviceGet");
+            char name[256]{};
+            int bus = -1, pci_device = -1, domain = -1;
+            api.check(api.cuDeviceGetName(name, sizeof(name), candidate), "cuDeviceGetName");
+            api.check(api.cuDeviceGetAttribute(&bus, 33, candidate), "cuDeviceGetAttribute(PCI_BUS_ID)");
+            api.check(api.cuDeviceGetAttribute(&pci_device, 34, candidate), "cuDeviceGetAttribute(PCI_DEVICE_ID)");
+            api.check(api.cuDeviceGetAttribute(&domain, 50, candidate), "cuDeviceGetAttribute(PCI_DOMAIN_ID)");
+            std::printf("CUDA GPU ordinal=%d name=%s pci=%04x:%02x:%02x\n", ordinal, name, domain, bus, pci_device);
+            if (bus == props.pciBusID && pci_device == props.pciDeviceID && domain == props.pciDomainID) device = candidate;
+        }
+        if (device < 0) throw std::runtime_error("CUDA/HIP PCI identity mismatch");
+        CUcontext context = nullptr;
+        const bool primary = args.context == "primary";
+        if (primary) {
+            api.check(api.cuDevicePrimaryCtxRetain(&context, device), "cuDevicePrimaryCtxRetain");
+        } else api.check(api.cuCtxCreate_v2(&context, 0, device), "cuCtxCreate_v2");
+        struct ContextCleanup {
+            CudaApi& api; CUcontext context; CUdevice device; bool primary;
+            ~ContextCleanup() { if (context) { api.cuCtxSetCurrent(nullptr);
+                if (primary) api.cuDevicePrimaryCtxRelease_v2(device); else api.cuCtxDestroy_v2(context); } }
+        } cleanup{api, context, device, primary};
+        api.check(api.cuCtxSetCurrent(context), "cuCtxSetCurrent");
+        CUmodule module = nullptr;
+        std::printf("STAGE cuModuleLoadData (PTX JIT)\n");
+        api.check(api.cuModuleLoadData(&module, ptx), "cuModuleLoadData");
+        struct ModuleCleanup { CudaApi& api; CUmodule m; ~ModuleCleanup() { if (m) api.cuModuleUnload(m); } } mc{api, module};
+        loaded_modules();
+        CUfunction function = nullptr;
+        api.check(api.cuModuleGetFunction(&function, module, "d4r_ptx_pattern"), "cuModuleGetFunction");
+        constexpr uint32_t nmax = 4099;
+        std::vector<uint32_t> result(nmax + 61, 0xdeadbeefu);
+        const size_t bytes = result.size() * sizeof(uint32_t);
+        CUdeviceptr pointer = 0;
+        api.check(api.cuMemAlloc_v2(&pointer, bytes), "cuMemAlloc_v2");
+        struct MemCleanup { CudaApi& api; CUdeviceptr p; ~MemCleanup() { if (p) api.cuMemFree_v2(p); } } ac{api, pointer};
+        for (unsigned i = 0; i < args.iterations; ++i) {
+            uint32_t n = nmax - i % 129u, seed = 0x12345678u + i * 31u;
+            std::fill(result.begin(), result.end(), 0xdeadbeefu);
+            api.check(api.cuMemcpyHtoD_v2(pointer, result.data(), bytes), "cuMemcpyHtoD_v2(sentinel)");
+            void* parameters[] = {&pointer, &n, &seed};
+            api.check(api.cuLaunchKernel(function, (n + 127u) / 128u, 1, 1, 128, 1, 1,
+                0, nullptr, parameters, nullptr), "cuLaunchKernel");
+            api.check(api.cuCtxSynchronize(), "cuCtxSynchronize");
+            api.check(api.cuMemcpyDtoH_v2(result.data(), pointer, bytes), "cuMemcpyDtoH_v2");
+            if (!verify(result, n, seed)) return 5;
+        }
+        api.check(api.cuMemFree_v2(pointer), "cuMemFree_v2"); ac.p = 0;
+        api.check(api.cuModuleUnload(module), "cuModuleUnload"); mc.m = nullptr;
+        api.check(api.cuCtxSetCurrent(nullptr), "cuCtxSetCurrent(NULL)");
+        if (primary) api.check(api.cuDevicePrimaryCtxRelease_v2(device), "cuDevicePrimaryCtxRelease_v2");
+        else api.check(api.cuCtxDestroy_v2(context), "cuCtxDestroy_v2");
+        cleanup.context = nullptr;
+        std::printf("PASS CUDA architecture=gfx1201 context=%s iterations=%u guard_verified=1\n",
+            args.context.c_str(), args.iterations);
+        return 0;
+    } catch (const std::exception& e) {
+        std::fprintf(stderr, "FAIL CUDA %s\n", e.what());
+        loaded_modules();
+        return 4;
+    }
+}

tools/windows/diagnostic.h

@@ -0,0 +1,156 @@
+#pragma once
+#include <windows.h>
+#include <tlhelp32.h>
+#include <dbghelp.h>
+#include <cstdint>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <filesystem>
+#include <stdexcept>
+#include <string>
+#include <vector>
+#include <algorithm>
+
+namespace d4r::diag {
+inline std::wstring wide(const std::string& s)
+{
+    const int n = MultiByteToWideChar(CP_UTF8, MB_ERR_INVALID_CHARS, s.c_str(), -1, nullptr, 0);
+    if (!n) throw std::runtime_error("Invalid UTF-8 path");
+    std::wstring out(n, L'\0');
+    MultiByteToWideChar(CP_UTF8, MB_ERR_INVALID_CHARS, s.c_str(), -1, out.data(), n);
+    out.pop_back();
+    return out;
+}
+inline std::string utf8(const wchar_t* s)
+{
+    const int n = WideCharToMultiByte(CP_UTF8, 0, s, -1, nullptr, 0, nullptr, nullptr);
+    std::string out(n, '\0');
+    WideCharToMultiByte(CP_UTF8, 0, s, -1, out.data(), n, nullptr, nullptr);
+    if (n) out.pop_back();
+    return out;
+}
+inline void loaded_modules()
+{
+    HANDLE snapshot = CreateToolhelp32Snapshot(TH32CS_SNAPMODULE, GetCurrentProcessId());
+    if (snapshot == INVALID_HANDLE_VALUE) return;
+    MODULEENTRY32W entry{};
+    entry.dwSize = sizeof(entry);
+    if (Module32FirstW(snapshot, &entry)) do {
+        DWORD unused = 0;
+        DWORD size = GetFileVersionInfoSizeW(entry.szExePath, &unused);
+        std::vector<unsigned char> bytes(size);
+        VS_FIXEDFILEINFO* version = nullptr;
+        UINT length = 0;
+        std::printf("DLL path=%s", utf8(entry.szExePath).c_str());
+        if (size && GetFileVersionInfoW(entry.szExePath, 0, size, bytes.data()) &&
+            VerQueryValueW(bytes.data(), L"\\", reinterpret_cast<void**>(&version), &length) &&
+            length >= sizeof(*version))
+            std::printf(" version=%u.%u.%u.%u", HIWORD(version->dwFileVersionMS),
+                LOWORD(version->dwFileVersionMS), HIWORD(version->dwFileVersionLS),
+                LOWORD(version->dwFileVersionLS));
+        std::printf("\n");
+    } while (Module32NextW(snapshot, &entry));
+    CloseHandle(snapshot);
+}
+inline LONG WINAPI unhandled(EXCEPTION_POINTERS* info)
+{
+    std::fprintf(stderr, "EXCEPTION code=0x%08lx address=%p\n",
+        info->ExceptionRecord->ExceptionCode, info->ExceptionRecord->ExceptionAddress);
+    loaded_modules();
+    wchar_t directory[32768]{};
+    const DWORD n = GetEnvironmentVariableW(L"D4R_DIAG_DIR", directory, 32768);
+    if (n && n < 32768) {
+        const std::wstring path = std::wstring(directory) + L"\\crash-" +
+            std::to_wstring(GetCurrentProcessId()) + L".dmp";
+        HANDLE file = CreateFileW(path.c_str(), GENERIC_WRITE, 0, nullptr, CREATE_ALWAYS, 0, nullptr);
+        if (file != INVALID_HANDLE_VALUE) {
+            MINIDUMP_EXCEPTION_INFORMATION exception{GetCurrentThreadId(), info, FALSE};
+            const BOOL ok = MiniDumpWriteDump(GetCurrentProcess(), GetCurrentProcessId(), file,
+                MiniDumpNormal, &exception, nullptr, nullptr);
+            std::fprintf(stderr, "MINIDUMP success=%d error=%lu\n", ok, ok ? 0 : GetLastError());
+            CloseHandle(file);
+        }
+    }
+    return EXCEPTION_EXECUTE_HANDLER;
+}
+inline void start()
+{
+    std::setvbuf(stdout, nullptr, _IONBF, 0);
+    std::setvbuf(stderr, nullptr, _IONBF, 0);
+    SetErrorMode(SEM_FAILCRITICALERRORS | SEM_NOGPFAULTERRORBOX);
+    SetUnhandledExceptionFilter(unhandled);
+    std::printf("d4r native Windows diagnostic x64 pid=%lu\n", GetCurrentProcessId());
+}
+struct Args {
+    std::string hip_root, cuda_dll, module, context = "primary";
+    unsigned iterations = 32;
+    int device = -1;
+    Args(int argc, char** argv) {
+        for (int i = 1; i < argc; ++i) {
+            const std::string key = argv[i];
+            if (i + 1 == argc) throw std::runtime_error("Missing value for " + key);
+            const std::string value = argv[++i];
+            if (key == "--hip-root") hip_root = value;
+            else if (key == "--cuda-dll") cuda_dll = value;
+            else if (key == "--module") module = value;
+            else if (key == "--context") context = value;
+            else if (key == "--iterations") {
+                size_t end = 0;
+                iterations = static_cast<unsigned>(std::stoul(value, &end));
+                if (end != value.size() || iterations < 1 || iterations > 10000)
+                    throw std::runtime_error("iterations must be 1..10000");
+            } else if (key == "--device") {
+                size_t end = 0;
+                device = std::stoi(value, &end);
+                if (end != value.size() || device < 0) throw std::runtime_error("Invalid device ordinal");
+            } else throw std::runtime_error("Unknown option " + key);
+        }
+        if (hip_root.empty()) throw std::runtime_error("--hip-root is required");
+    }
+};
+class Library {
+    HMODULE handle_ = nullptr;
+public:
+    explicit Library(const std::filesystem::path& path) {
+        if (!path.is_absolute()) throw std::runtime_error("DLL path must be absolute");
+        handle_ = LoadLibraryExW(path.c_str(), nullptr,
+            LOAD_LIBRARY_SEARCH_DLL_LOAD_DIR | LOAD_LIBRARY_SEARCH_DEFAULT_DIRS | LOAD_LIBRARY_SEARCH_USER_DIRS);
+        if (!handle_) throw std::runtime_error("LoadLibrary " + utf8(path.c_str()) +
+            " Win32=" + std::to_string(GetLastError()));
+        std::printf("LOAD requested=%s\n", utf8(path.c_str()).c_str());
+    }
+    Library(const Library&) = delete;
+    Library& operator=(const Library&) = delete;
+    ~Library() { if (handle_) FreeLibrary(handle_); }
+    template<class T> T symbol(const char* name) const {
+        FARPROC raw = GetProcAddress(handle_, name);
+        if (!raw) throw std::runtime_error(std::string("Missing export ") + name);
+        static_assert(sizeof(T) == sizeof(raw));
+        T out;
+        std::memcpy(&out, &raw, sizeof(out));
+        return out;
+    }
+};
+class SearchDirectory {
+    DLL_DIRECTORY_COOKIE cookie_ = nullptr;
+public:
+    explicit SearchDirectory(const std::filesystem::path& path) {
+        if (!path.is_absolute()) throw std::runtime_error("HIP root must be absolute");
+        cookie_ = AddDllDirectory(path.c_str());
+        if (!cookie_) throw std::runtime_error("AddDllDirectory failed");
+    }
+    ~SearchDirectory() { if (cookie_) RemoveDllDirectory(cookie_); }
+};
+inline uint32_t pattern(uint32_t i, uint32_t seed) { return (i * 1664525u + seed) ^ 0xa5a55a5au; }
+inline bool verify(const std::vector<uint32_t>& result, uint32_t count, uint32_t seed) {
+    for (size_t i = 0; i < result.size(); ++i) {
+        const uint32_t expected = i < count ? pattern(static_cast<uint32_t>(i), seed) : 0xdeadbeefu;
+        if (result[i] != expected) {
+            std::fprintf(stderr, "MISMATCH index=%zu expected=0x%08x actual=0x%08x\n", i, expected, result[i]);
+            return false;
+        }
+    }
+    return true;
+}
+}

tools/windows/hip_api.h

@@ -0,0 +1,68 @@
+#pragma once
+#include "diagnostic.h"
+#include <hip/hip_runtime_api.h>
+
+namespace d4r::diag {
+struct HipApi {
+    SearchDirectory search;
+    Library library;
+#define D4R_HIP_FUNCTION(name) decltype(&::name) name = library.symbol<decltype(&::name)>(#name)
+    D4R_HIP_FUNCTION(hipInit);
+    D4R_HIP_FUNCTION(hipGetDeviceCount);
+    // hipGetDeviceProperties is an SDK macro; the ABI-versioned export is intentional.
+    decltype(&::hipGetDeviceProperties) hipGetDeviceProperties =
+        library.symbol<decltype(&::hipGetDeviceProperties)>("hipGetDevicePropertiesR0600");
+    D4R_HIP_FUNCTION(hipSetDevice);
+    D4R_HIP_FUNCTION(hipRuntimeGetVersion);
+    D4R_HIP_FUNCTION(hipDriverGetVersion);
+    D4R_HIP_FUNCTION(hipGetErrorName);
+    D4R_HIP_FUNCTION(hipGetErrorString);
+    using MallocFn = hipError_t (*)(void**, size_t);
+    MallocFn hipMalloc = library.symbol<MallocFn>("hipMalloc");
+    D4R_HIP_FUNCTION(hipFree);
+    D4R_HIP_FUNCTION(hipMemcpy);
+    D4R_HIP_FUNCTION(hipMemGetInfo);
+    D4R_HIP_FUNCTION(hipModuleLoad);
+    D4R_HIP_FUNCTION(hipModuleUnload);
+    D4R_HIP_FUNCTION(hipModuleGetFunction);
+    D4R_HIP_FUNCTION(hipModuleLaunchKernel);
+    D4R_HIP_FUNCTION(hipDeviceSynchronize);
+#undef D4R_HIP_FUNCTION
+    explicit HipApi(const std::string& root) :
+        search(std::filesystem::path(wide(root)) / L"bin"),
+        library(std::filesystem::path(wide(root)) / L"bin" / L"amdhip64_7.dll") {}
+    void check(hipError_t result, const char* call) const {
+        std::printf("HIP %s -> %d (%s)\n", call, static_cast<int>(result), hipGetErrorName(result));
+        if (result != hipSuccess) throw std::runtime_error(std::string(call) + ": " + hipGetErrorString(result));
+    }
+    int select_gfx1201(int requested, hipDeviceProp_t& selected) const {
+        // Architecture spoofing invalidates a gfx12 correctness test.
+        for (const char* name : {"HSA_OVERRIDE_GFX_VERSION", "HSA_OVERRIDE_GFX_VERSION_0"})
+            if (std::getenv(name)) throw std::runtime_error(std::string("Unset architecture override ") + name);
+        check(hipInit(0), "hipInit");
+        int version = 0, driver = 0, count = 0;
+        check(hipRuntimeGetVersion(&version), "hipRuntimeGetVersion");
+        check(hipDriverGetVersion(&driver), "hipDriverGetVersion");
+        std::printf("HIP runtime=%d driver_api=%d headers=%d.%d.%d sizeof(properties)=%zu\n",
+            version, driver, HIP_VERSION_MAJOR, HIP_VERSION_MINOR, HIP_VERSION_PATCH, sizeof(hipDeviceProp_t));
+        check(hipGetDeviceCount(&count), "hipGetDeviceCount");
+        int found = -1;
+        for (int i = 0; i < count; ++i) {
+            hipDeviceProp_t props{};
+            check(hipGetDeviceProperties(&props, i), "hipGetDevicePropertiesR0600");
+            std::printf("GPU ordinal=%d name=%s gcnArchName=%s wave=%d vram=%zu pci=%04x:%02x:%02x\n",
+                i, props.name, props.gcnArchName, props.warpSize, props.totalGlobalMem,
+                props.pciDomainID, props.pciBusID, props.pciDeviceID);
+            const std::string arch = props.gcnArchName;
+            if (arch.substr(0, arch.find(':')) == "gfx1201" && (requested < 0 || requested == i) && found < 0) {
+                found = i;
+                selected = props;
+            }
+        }
+        if (found < 0) throw std::runtime_error("No selected gfx1201 device; never substitute gfx110x");
+        check(hipSetDevice(found), "hipSetDevice");
+        std::printf("SELECTED HIP ordinal=%d architecture=gfx1201\n", found);
+        return found;
+    }
+};
+}

tools/windows/hip_gfx1201_probe.cpp

@@ -0,0 +1,51 @@
+#include "hip_api.h"
+
+int main(int argc, char** argv)
+{
+    using namespace d4r::diag;
+    start();
+    try {
+        Args args(argc, argv);
+        if (args.module.empty()) throw std::runtime_error("--module is required");
+        HipApi api(args.hip_root);
+        hipDeviceProp_t props{};
+        api.select_gfx1201(args.device, props);
+        loaded_modules();
+        size_t free = 0, total = 0;
+        api.check(api.hipMemGetInfo(&free, &total), "hipMemGetInfo");
+        std::printf("VRAM free=%zu total=%zu\n", free, total);
+        hipModule_t module = nullptr;
+        api.check(api.hipModuleLoad(&module, args.module.c_str()), "hipModuleLoad(gfx1201)");
+        struct ModuleCleanup { HipApi& api; hipModule_t m; ~ModuleCleanup() { if (m) (void)api.hipModuleUnload(m); } } cleanup{api, module};
+        hipFunction_t kernel = nullptr;
+        api.check(api.hipModuleGetFunction(&kernel, module, "d4r_probe_pattern"), "hipModuleGetFunction");
+        constexpr uint32_t count = 4099, guard = 61;
+        std::vector<uint32_t> result(count + guard, 0xdeadbeefu);
+        const size_t bytes = result.size() * sizeof(uint32_t);
+        void* allocation = nullptr;
+        api.check(api.hipMalloc(&allocation, bytes), "hipMalloc(VRAM)");
+        struct AllocCleanup { HipApi& api; void* p; ~AllocCleanup() { if (p) (void)api.hipFree(p); } } allocation_cleanup{api, allocation};
+        for (unsigned iteration = 0; iteration < args.iterations; ++iteration) {
+            uint32_t n = count - iteration % 129u, seed = 0x12345678u + iteration * 31u;
+            std::fill(result.begin(), result.end(), 0xdeadbeefu);
+            api.check(api.hipMemcpy(allocation, result.data(), bytes, hipMemcpyHostToDevice), "hipMemcpy(HtoD sentinel)");
+            void* parameters[] = {&allocation, &n, &seed};
+            api.check(api.hipModuleLaunchKernel(kernel, (n + 127u) / 128u, 1, 1, 128, 1, 1,
+                0, nullptr, parameters, nullptr), "hipModuleLaunchKernel");
+            api.check(api.hipDeviceSynchronize(), "hipDeviceSynchronize");
+            api.check(api.hipMemcpy(result.data(), allocation, bytes, hipMemcpyDeviceToHost), "hipMemcpy(DtoH)");
+            if (!verify(result, n, seed)) return 5;
+        }
+        // Explicitly validate teardown too; disable RAII cleanup after success.
+        api.check(api.hipFree(allocation), "hipFree");
+        allocation_cleanup.p = nullptr;
+        api.check(api.hipModuleUnload(module), "hipModuleUnload");
+        cleanup.m = nullptr;
+        std::printf("PASS HIP architecture=gfx1201 iterations=%u guard_verified=1\n", args.iterations);
+        return 0;
+    } catch (const std::exception& e) {
+        std::fprintf(stderr, "FAIL HIP %s\n", e.what());
+        loaded_modules();
+        return 4;
+    }
+}

tools/windows/probe_kernel.hip

@@ -0,0 +1,9 @@
+// No host SDK, CRT, or device library dependency. Compiled only for gfx1201.
+extern "C" __attribute__((global)) void d4r_probe_pattern(
+    unsigned int* output, unsigned int count, unsigned int seed)
+{
+    const unsigned int i = __builtin_amdgcn_workgroup_id_x() * 128u +
+                           __builtin_amdgcn_workitem_id_x();
+    if (i < count)
+        output[i] = (i * 1664525u + seed) ^ 0xa5a55a5au;
+}