From e95e061e813603f867a7ad3d194be6978f372467 Mon Sep 17 00:00:00 2001 From: Codex Date: Tue, 29 Sep 2026 18:50:09 +0300 Subject: [PATCH] Add native Windows gfx1201 HIP and ZLUDA diagnostics --- .gitignore | 5 +- CMakeLists.txt | 56 +++++++ README.md | 5 + docs/windows-rdna4-port.md | 209 ++++++++++++++++++++++++ scripts/windows/build-windows-rdna4.ps1 | 56 +++++++ scripts/windows/setup-windows-tools.ps1 | 31 ++++ scripts/windows/test-windows-rdna4.ps1 | 127 ++++++++++++++ tools/windows/cuda_driver_probe.cpp | 159 ++++++++++++++++++ tools/windows/diagnostic.h | 156 ++++++++++++++++++ tools/windows/hip_api.h | 68 ++++++++ tools/windows/hip_gfx1201_probe.cpp | 51 ++++++ tools/windows/probe_kernel.hip | 9 + 12 files changed, 931 insertions(+), 1 deletion(-) create mode 100644 CMakeLists.txt create mode 100644 docs/windows-rdna4-port.md create mode 100644 scripts/windows/build-windows-rdna4.ps1 create mode 100644 scripts/windows/setup-windows-tools.ps1 create mode 100644 scripts/windows/test-windows-rdna4.ps1 create mode 100644 tools/windows/cuda_driver_probe.cpp create mode 100644 tools/windows/diagnostic.h create mode 100644 tools/windows/hip_api.h create mode 100644 tools/windows/hip_gfx1201_probe.cpp create mode 100644 tools/windows/probe_kernel.hip diff --git a/.gitignore b/.gitignore index 104ea9d..4a6c178 100644 --- a/.gitignore +++ b/.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/ \ No newline at end of file +/kernels/extracted/ diff --git a/CMakeLists.txt b/CMakeLists.txt new file mode 100644 index 0000000..87ef6bc --- /dev/null +++ b/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$<$: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() diff --git a/README.md b/README.md index a3573df..d3183f1 100644 --- a/README.md +++ b/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.** diff --git a/docs/windows-rdna4-port.md b/docs/windows-rdna4-port.md new file mode 100644 index 0000000..35dd9f1 --- /dev/null +++ b/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. diff --git a/scripts/windows/build-windows-rdna4.ps1 b/scripts/windows/build-windows-rdna4.ps1 new file mode 100644 index 0000000..3f7ea05 --- /dev/null +++ b/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 } diff --git a/scripts/windows/setup-windows-tools.ps1 b/scripts/windows/setup-windows-tools.ps1 new file mode 100644 index 0000000..b8a1c04 --- /dev/null +++ b/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.' diff --git a/scripts/windows/test-windows-rdna4.ps1 b/scripts/windows/test-windows-rdna4.ps1 new file mode 100644 index 0000000..b94caf4 --- /dev/null +++ b/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 diff --git a/tools/windows/cuda_driver_probe.cpp b/tools/windows/cuda_driver_probe.cpp new file mode 100644 index 0000000..582bdba --- /dev/null +++ b/tools/windows/cuda_driver_probe.cpp @@ -0,0 +1,159 @@ +#include "hip_api.h" +#include + +// 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) + 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 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; + } +} diff --git a/tools/windows/diagnostic.h b/tools/windows/diagnostic.h new file mode 100644 index 0000000..c7b6a1b --- /dev/null +++ b/tools/windows/diagnostic.h @@ -0,0 +1,156 @@ +#pragma once +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +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 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(&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(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 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& 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(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; +} +} diff --git a/tools/windows/hip_api.h b/tools/windows/hip_api.h new file mode 100644 index 0000000..da53175 --- /dev/null +++ b/tools/windows/hip_api.h @@ -0,0 +1,68 @@ +#pragma once +#include "diagnostic.h" +#include + +namespace d4r::diag { +struct HipApi { + SearchDirectory search; + Library library; +#define D4R_HIP_FUNCTION(name) decltype(&::name) name = library.symbol(#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("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("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(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; + } +}; +} diff --git a/tools/windows/hip_gfx1201_probe.cpp b/tools/windows/hip_gfx1201_probe.cpp new file mode 100644 index 0000000..3cd5b9c --- /dev/null +++ b/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 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; + } +} diff --git a/tools/windows/probe_kernel.hip b/tools/windows/probe_kernel.hip new file mode 100644 index 0000000..9ab0c08 --- /dev/null +++ b/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; +}