diff --git a/CMakeLists.txt b/CMakeLists.txt index 1909c96..55103e2 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -15,15 +15,36 @@ endif() find_program(D4R_HIP_CLANG NAMES clang++.exe PATHS "${D4R_HIP_ROOT}/bin" "${D4R_HIP_ROOT}/lib/llvm/bin" NO_DEFAULT_PATH REQUIRED) # Architecture parameterization adapted from realdody/d4r ae530b0. # Keep explicit subtargets: no gfx11 substitution or untested future gfx12 ISA. +file(READ "${CMAKE_CURRENT_SOURCE_DIR}/tools/windows/gpu-targets.json" d4r_target_json) +set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS tools/windows/gpu-targets.json) +string(JSON d4r_target_count LENGTH "${d4r_target_json}" targets) +math(EXPR d4r_target_last "${d4r_target_count} - 1") +set(d4r_supported_targets) +set(d4r_target_cpp_entries "") +foreach(index RANGE ${d4r_target_last}) + foreach(field architecture elfMachine isaMajor family) + string(JSON target_${field} GET "${d4r_target_json}" targets ${index} ${field}) + endforeach() + list(APPEND d4r_supported_targets "${target_architecture}") + string(APPEND d4r_target_cpp_entries " GpuTarget{\"${target_architecture}\", ${target_elfMachine}, ${target_isaMajor}, \"${target_family}\"},\n") +endforeach() +configure_file(tools/windows/gpu_targets.generated.h.in gpu_targets.generated.h @ONLY) +include_directories("${CMAKE_CURRENT_BINARY_DIR}") set(D4R_GPU_ARCH "gfx1201" CACHE STRING "Native Windows package GPU target") -set_property(CACHE D4R_GPU_ARCH PROPERTY STRINGS gfx1201 gfx1200) -set(d4r_supported_targets gfx1200 gfx1201) +set_property(CACHE D4R_GPU_ARCH PROPERTY STRINGS ${d4r_supported_targets}) if(NOT D4R_GPU_ARCH IN_LIST d4r_supported_targets) - message(FATAL_ERROR "D4R_GPU_ARCH must be gfx1200 or gfx1201") + message(FATAL_ERROR "D4R_GPU_ARCH must be one of: ${d4r_supported_targets}") endif() add_compile_definitions(D4R_TARGET_ARCH="${D4R_GPU_ARCH}") configure_file(tools/windows/gpu-target.json.in gpu-target.json @ONLY) install(FILES "${CMAKE_CURRENT_BINARY_DIR}/gpu-target.json" DESTINATION .) +install(FILES tools/windows/gpu-targets.json DESTINATION .) +set(d4r_device_flags) +if(D4R_GPU_ARCH MATCHES "^gfx11") + # Windows-only workaround informed by realdody/d4r 1897b4a. Keep the real + # gfx11 layout; do not route RDNA3 through gfx12 fragments. + list(APPEND d4r_device_flags -DD4R_WINDOWS_GFX11_LANE_EXCHANGE) +endif() set(probe_object "${CMAKE_CURRENT_BINARY_DIR}/probe_${D4R_GPU_ARCH}.hsaco") add_custom_command(OUTPUT "${probe_object}" COMMAND "${D4R_HIP_CLANG}" -x hip --offload-arch=${D4R_GPU_ARCH} @@ -44,10 +65,10 @@ set(wmma_object "${CMAKE_CURRENT_BINARY_DIR}/wmma_${D4R_GPU_ARCH}.hsaco") add_custom_command(OUTPUT "${wmma_object}" COMMAND "${D4R_HIP_CLANG}" -x hip --offload-arch=${D4R_GPU_ARCH} --offload-device-only --no-gpu-bundle-output -mno-wavefrontsize64 - -nogpuinc -nogpulib -O2 "${CMAKE_CURRENT_SOURCE_DIR}/tools/windows/wmma_gfx1201.hip" + -nogpuinc -nogpulib -O2 ${d4r_device_flags} "${CMAKE_CURRENT_SOURCE_DIR}/tools/windows/wmma_gfx1201.hip" -o "${wmma_object}" DEPENDS tools/windows/wmma_gfx1201.hip kernels/common/wmma_backend.h - kernels/common/wmma_layout.h VERBATIM) + kernels/common/wmma_layout.h kernels/common/wave32_exchange.h VERBATIM) add_custom_target(d4r_wmma_code_object ALL DEPENDS "${wmma_object}") # K/M objects stay in experimental/, outside the ZLUDA native override # directory, until each layer and then the complete network pass replay. @@ -78,9 +99,9 @@ foreach(family k m) add_custom_command(OUTPUT "${native_object}" COMMAND "${D4R_HIP_CLANG}" -x hip -std=c++17 --offload-arch=${D4R_GPU_ARCH} --offload-device-only --no-gpu-bundle-output -mno-wavefrontsize64 - -nogpuinc -nogpulib -O3 -DD4R_DEVICE_ONLY_MINIMAL ${native_flags} + -nogpuinc -nogpulib -O3 -DD4R_DEVICE_ONLY_MINIMAL ${d4r_device_flags} ${native_flags} "${source}" -o "${native_object}" - DEPENDS "${source}" ${native_headers} kernels/common/wmma_layout.h + DEPENDS "${source}" ${native_headers} kernels/common/wmma_layout.h kernels/common/wave32_exchange.h kernels/common/hip_device_minimal.h VERBATIM) list(APPEND native_objects "${native_object}") install(FILES "${native_object}" DESTINATION "experimental/${family}") @@ -89,7 +110,7 @@ endforeach() add_custom_target(d4r_native_compile ALL DEPENDS ${native_objects}) # Retain the old manual build target for existing scripts. add_custom_target(d4r_native_gfx1201_compile DEPENDS d4r_native_compile) -foreach(probe hip_gfx1201_probe gfx12_wmma_probe k_module_probe native_replay_probe cuda_driver_probe cuda_mma_probe cuda_image_probe d3d12_command_probe d3d12_pixel_probe d3d12_hip_interop_probe hip_stream_lifecycle_probe ngx_cuda_init_probe d3d12_dlss_probe) +foreach(probe gpu_inventory hip_gfx1201_probe gfx12_wmma_probe k_module_probe native_replay_probe cuda_driver_probe cuda_mma_probe cuda_image_probe d3d12_command_probe d3d12_pixel_probe d3d12_hip_interop_probe hip_stream_lifecycle_probe ngx_cuda_init_probe d3d12_dlss_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") @@ -200,6 +221,7 @@ install(FILES "${pixel_object}" DESTINATION bin) install(FILES "${wmma_object}" DESTINATION bin) install(FILES scripts/windows/test-windows-rdna4.ps1 DESTINATION .) install(FILES scripts/windows/gpu-target.ps1 DESTINATION .) +install(FILES scripts/windows/gpu-preflight.ps1 DESTINATION .) install(FILES scripts/windows/test-gpu-target-contract.ps1 DESTINATION .) install(FILES scripts/windows/test-k-replay.ps1 DESTINATION .) install(FILES scripts/windows/test-k-captures.ps1 DESTINATION .) diff --git a/docs/windows-game.md b/docs/windows-game.md index eb5d3bd..bc7b4f2 100644 --- a/docs/windows-game.md +++ b/docs/windows-game.md @@ -1,21 +1,29 @@ -# Native Windows RDNA4 game development package +# Native Windows RDNA3 / RDNA4 game development package Validated hardware: Windows 11 x64 / Radeon RX 9070 XT / gfx1201. Native K and M transformer kernels pass standalone D3D12 + OptiScaler checks. Silent Hill 2 K now renders real menu frames and M renders a saved gameplay level with GPU NaN/Inf checks. Both K and M pass 3840x2160 output on this machine. -The user reports 49-51 FPS for K in the measured scene after the command-hook -fix and locally built output-store optimization. This is a local +Performance measurements and their controls are recorded in +[windows-performance.md](windows-performance.md). This is a development package; no NVIDIA proprietary DLL is included. -Builds also accept gfx1200 (for example RX 9060/9060 XT). This target is -**compile-tested only**; no physical gfx1200 validation has been performed. +Experimental builds also accept gfx1100..gfx1103 (RDNA3), gfx1150..gfx1154 +(RDNA3.5) and gfx1200 (for example RX 9060/9060 XT). These targets are +**compile-tested only**; no physical validation here has been performed. RX 9070 and other gfx1201 cards also remain untested. Use a package built for the GPU's actual HIP architecture: its `gpu-target.json`, host shim and ELF code objects must agree. Installation/runtime reject mixed targets. See [windows-gpu-support.md](windows-gpu-support.md) for separate-target build commands, the fork review and the first hardware diagnostic gate. +Before changing game files, the runner inventories HIP devices without +executing kernels. It rejects another architecture's package and a local +D3D12 proxy identifying itself as vkd3d/Wine, producing one diagnostic ZIP. +RX 9060 XT needs gfx1200: the old gfx1201-only package fails before DLSS +initialization on it. Use native D3D12. On an untested GPU, start with +synchronous K and output/exception diagnostics. + Build the shim/diagnostics, corrected ZLUDA, K/M native objects and patched OptiScaler using the scripts in `scripts/windows`, then run: diff --git a/docs/windows-gpu-support.md b/docs/windows-gpu-support.md index 1ab4c9d..b9b30d4 100644 --- a/docs/windows-gpu-support.md +++ b/docs/windows-gpu-support.md @@ -1,40 +1,80 @@ -# Windows GPU targets +# Experimental Windows RDNA3 / RDNA4 coverage -The native Windows build accepts **gfx1200** and **gfx1201**. It builds a -separate shim, diagnostics and code objects for each target. The default stays -gfx1201. GPU selection uses HIP's real `gcnArchName`, followed by D3D12/HIP -adapter LUID matching; it does not depend on a marketing name or overridden -DXGI vendor ID. +The shim, diagnostics, native K/M kernels and game installer accept the +targets below, with one package per actual ISA. **Only RX 9070 XT has physical +validation here.** All other targets are experimental and compile-tested; +working DLSS, interop and performance on them remain unverified. -| Target | Examples in LLVM's processor table | This project's validation | +| Family | Target | Examples / scope | | --- | --- | --- | -| gfx1201 | RX 9070, RX 9070 XT | RX 9070 XT: hardware, K/M and game tests. Other cards untested. | -| gfx1200 | RX 9060, RX 9060 XT | Host/device compilation and software tests only. No physical gfx1200 available. | +| RDNA3 | gfx1100 | RX 7900 family; same-target workstation GPUs | +| RDNA3 | gfx1101 | RX 7700 / 7700 XT / 7800 XT; same-target mobile/workstation GPUs | +| RDNA3 | gfx1102 | RX 7600 / 7600 XT; same-target mobile/workstation GPUs | +| RDNA3 | gfx1103 | RDNA3 integrated graphics | +| RDNA3.5 | gfx1150, gfx1151, gfx1152, gfx1153, gfx1154 | Corresponding integrated graphics targets | +| RDNA4 | gfx1200 | RX 9060 / 9060 XT; same-target GPUs | +| RDNA4 | gfx1201 | RX 9070 / 9070 XT / 9070 GRE; same-target GPUs | -Sources: [LLVM AMDGPU processors](https://llvm.org/docs/AMDGPUUsage.html#processors) -and [AMD's Windows compatibility matrix](https://rocm.docs.amd.com/projects/radeon-ryzen/en/latest/docs/compatibility/compatibilityrad/windows/windows_compatibility.html). -AMD's matrix is dependency information, not evidence that d4r's K/M or external -memory path works on an untested card. Other RDNA4 SKUs are eligible only when -HIP reports one of these exact targets and exposes the required Windows APIs. -Device-specific drivers, performance and memory capacity still need testing. -gfx1250/gfx1251 are a separate ISA group and are not enabled by this change. +Names and IDs follow [LLVM processors](https://llvm.org/docs/AMDGPUUsage.html#processors) +and [ELF ABI](https://llvm.org/docs/AMDGPUUsage.html#amdgpu-elf-header-e-flags). +`tools/windows/gpu-targets.json` supplies CMake-generated C++, PowerShell and +the private texture builder. Selection uses HIP's real `gcnArchName` and the +D3D12 adapter LUID, including systems with several GPUs sharing a target. +Marketing names do not control eligibility. -Do not rename a gfx1201 package to gfx1200. `gpu-target.json`, the compiled -host target and each ELF code-object target must agree. Runtime selection -rejects another architecture before launching kernels. Staging, packaging and -installation inspect the documented AMDGPU ELF machine ID and reject mixed -objects. They never rewrite a binary target or set `HSA_OVERRIDE_GFX_VERSION`. -Legacy packages without target metadata remain gfx1201 only. +All listed RDNA3/RDNA3.5/RDNA4 targets with existing gfx11/gfx12 WMMA contracts +are included. TBD RDNA4m gfx1170..gfx1172 and gfx12.5 gfx1250/gfx1251 are +different ISA groups and excluded. Unknown future targets need explicit review. +Check [AMD's Windows matrix](https://rocm.docs.amd.com/projects/radeon-ryzen/en/latest/docs/compatibility/compatibilityrad/windows/windows_compatibility.html) +for driver/runtime eligibility. Compiler support cannot supply missing Windows +driver support or external-memory APIs, especially on APUs. -## Build and package +## Installation and diagnostics -Prepare the pinned dependencies as described in `windows-rdna4-port.md`. -Select a target explicitly; this example creates the experimental gfx1200 -package. Use gfx1201 for a 9070/9070 XT. Keep separate output directories when -building both targets. +Use the actual target's archive: **gfx1200 for RX 9060 XT**, for example. +Metadata, host target and every ELF target must agree. Runtime/scripts reject +mixed targets and never rename targets, rewrite ELF IDs or set +`HSA_OVERRIDE_GFX_VERSION`. Legacy metadata-free packages remain gfx1201 only. + +Before changing game files, `windows-game.ps1` now runs a read-only HIP +inventory and checks the target and wave32. Inventory allocates no GPU buffers +and runs no kernels; success is discovery, not validation. Failure prints one +ZIP containing stdout/stderr, exit code, runtime/DLL and driver details. A local +`d3d12.dll` identifying itself as vkd3d/Wine is rejected before install. Restore +the game's native D3D12 installation; the script never loads or removes that +proxy itself. Begin synchronously on an untested GPU: ```powershell -$arch = 'gfx1200' +.\windows-game.ps1 -GameExe "D:\Games\Example\Game.exe" -NgxCore "C:\LocalDLLs\_nvngx.dll" -DlssDll "C:\LocalDLLs\nvngx_dlss.dll" -Preset 11 -ValidateOutput -CaptureExceptions +``` + +The local DLLs must match the documented NGX version and DLSS 310.9.1 SHA256 +in `package.json`. NVIDIA DLLs are never distributed. Retain the printed ZIP +on failure. Async K and performance are subsequent gates; gfx1201-specific +FP8, arithmetic and tuning experiments remain separate. + +## Issue #10 review, 2026-10-02 + +Reviewed [comment 5952822115](https://github.com/countervolts/d4r/issues/10#issuecomment-5952822115), +including all 32 attached text logs from eight attempts. HIP discovers +**RX 9060 XT / gfx1200** successfully. Every attempt fails with the old shim's +**"No selected gfx1201 device"** rejection; OptiScaler reports `support=false` +/ `BAD00002`. None reaches `D4R_RUNTIME` initialization or native transformer +execution. These logs do not demonstrate a K/M kernel crash. Use the gfx1200 +package instead of overriding architecture. Three attempts also load +vkd3d-proton 3.1.0 from a local D3D12 proxy, outside this Windows backend. + +The new packages/preflight address the identified failures. They do not prove +those games work on RX 9060 XT; a new native-D3D12 run on that GPU is required. +Raw community logs and personal directory paths stay local. + +## Build + +Prepare dependencies following `windows-rdna4-port.md`. Default: gfx1201. +Select another target explicitly and keep separate directories: + +```powershell +$arch = 'gfx1101' # for example RX 7800 XT $diag = "$PWD\dist\windows-$arch-diagnostics" .\scripts\windows\build-windows-rdna4.ps1 -RuntimeProfile therock -GpuArch $arch -ZludaRoot "$PWD\dist\zluda-windows-final" -BuildDirectory "$PWD\build\windows-$arch" -InstallDirectory $diag .\scripts\windows\stage-native-k.ps1 -DlssDll "$PWD\nvngx_dlss.dll" -PackageRoot $diag @@ -42,89 +82,60 @@ $diag = "$PWD\dist\windows-$arch-diagnostics" .\scripts\windows\package-windows-game.ps1 -DiagnosticRoot $diag -GpuArch $arch -PackageRoot "$PWD\dist\windows-$arch-game" -ArchivePath "$PWD\dist\windows-$arch-game.zip" ``` -ZLUDA, HIP and OptiScaler must already be built/prepared. The package contains -the strict K FP16 and M FP16-equivalent baseline; native FP8 stays disabled. -The architecture does not change arithmetic or the existing WMMA layout. -Private NVIDIA-derived texture overrides remain optional local experiments; -their builder accepts `-GpuArch`, and their validation must match the target -package before installation. Existing gfx1201 performance measurements are -not measurements of gfx1200. -The separate experimental FP8/F16 arithmetic and tuning recipes remain -gfx1201-specific; this change enables the conservative production baseline -on both build targets. +Rebuild diagnostics for the new inventory executable. The package keeps strict +K FP16 and M FP16-equivalent baselines; native M FP8 remains disabled. CMake: +`-DD4R_GPU_ARCH=gfx1101`. Some source/executable names retain gfx1201/gfx12 for +compatibility; actual targets follow the selected architecture. -Equivalent CMake selection is `-DD4R_GPU_ARCH=gfx1200`. Unknown targets fail -configuration. Some diagnostic executable names and source filenames retain -`gfx1201` for compatibility; the logged architecture and compiled code objects -use the selected target. - -## First hardware gate on another card - -From the repository root, with the matching diagnostic package above and the -pinned dependencies, run this single command: +One command compiles all eleven targets without GPU workloads: ```powershell -.\scripts\windows\test-windows-rdna4.ps1 -RuntimeProfile therock -PackageRoot "$PWD\dist\windows-gfx1200-diagnostics" -ZludaRoot "$PWD\dist\zluda-windows-final" +.\scripts\windows\build-windows-gpu-matrix.ps1 -ZludaRoot "$PWD\dist\zluda-windows-final" ``` -Success requires exit 0, all selected tests passed in `summary.json`, the -correct actual GPU/target in stdout, exact HIP/PTX/WMMA/pixel-format checks and -stable interop lifetime. On failure, retain the single printed diagnostic ZIP; -it includes stdout/stderr, exit code, runtime/driver inventory and DLL identity. -No NVIDIA DLL is required for this first gate. Full K/M reference, temporal, -standalone DLSS and game validation are subsequent gates on that hardware; -passing the first gate alone does not close them. +Add `-PackageGames -DlssDll "$PWD\nvngx_dlss.dll" -ArchiveDirectory "$PWD\dist\windows-gpu-archives"` +to stage/archive all packages after committing source and preparing OptiScaler. +Use `-GpuArchitectures gfx1100,gfx1101,gfx1102,gfx1200,gfx1201` for a subset. +Per-target logs and `matrix-build.json` distinguish compilation from execution. +The supplied DLL generates manifests locally and is never copied into archives. -Without that GPU, run `ctest --test-dir build/windows-gfx1200 -L software ---output-on-failure`. These CPU/WARP checks do not execute gfx1200 kernels. +## Fork adaptation -## Fork review +Reviewed [realdody's branch](https://github.com/realdody/d4r/tree/windows-rdna3) +at `ae530b0f53ed5c0bb87558dd7c20f619e014e4fd` and the newer +[`1897b4a` correction](https://github.com/realdody/d4r/commit/1897b4afa79832bd879234cfb37e7c1468db5cde). +Earlier MSVC binary16 and generated `BOOL` fixes remain. The earlier forced +gfx12 fragment layout is not used. -Reviewed [realdody's windows-rdna3 branch](https://github.com/realdody/d4r/tree/windows-rdna3) -at `ae530b0f53ed5c0bb87558dd7c20f619e014e4fd`, on 2026-10-02. Adapted: +The newer fork reports Windows gfx1101 `permlanex16` returning its own lane, +duplicating even K values into odd slots. This port adapts that finding into +`kernels/common/wave32_exchange.h`: a Windows/gfx11-only per-wave LDS exchange, +with native gfx11 layout, generic x/y/z indexing, 2 KiB scratch (512 threads +maximum), volatile accesses and wave fences/barriers protecting scratch reuse. +It covers K/M fragment conversions and other K half-wave exchanges. gfx12 and +Linux keep the original instruction path. The raw gfx11 probe uses its actual +intrinsic/layout, without a forced gfx12 shim. -- Per-target build configuration and actual-target diagnostics, narrowed here - to the two documented RDNA4 subtargets, with package/ELF guards added. -- MSVC binary16 conversion in the synthetic CUDA MMA reference. An exhaustive - CPU test covers 63,490 finite/infinity encodings and 31,743 rounding - boundaries; LLVM also checks it against native `_Float16` conversion. -- `WINBOOL` to `BOOL` normalization in generated D3D12 hooks, so the generated - header works with the Microsoft SDK as well as MinGW headers. +The fork reports RX 7800 XT enc1/enc2 identity/nonzero/reference checks. Those +are **the author's results**, not validation of this adaptation or full K/M +here. LDS exchange, full network, translated CUDA, accuracy, interop and speed +still need physical gfx11 testing. Credit: realdody . -The fork's RDNA3 implementation forces `D4R_WMMA_LAYOUT=12` for gfx11. That -kernel-layout change, the raw gfx11 probe and the expanded gfx11 runtime -allowlist are not imported. Its reported RX 7800 XT/K result has not been -independently reproduced here; its M/translated-path limitations still require -separate checks. This change does not advertise native Windows RDNA3 readiness. -The existing upstream Linux/RDNA3 implementation is unchanged. Credit for -the adapted portability fixes: realdody . +## Results and next hardware gate -## Local results, 2026-10-02 +All eleven targets compile the complete MinGW host build and 19 device objects +with TheRock `10.2.0a20260929`, HIP `7.17.26386`, Clang 24. Software target +contracts inspect actual ELF objects without other-target GPU launches. +RX 9070 XT passes all 20 CTest gates after these changes. K/M full-frame and +package checks are recorded in `windows-rdna4-port.md`. Local logs are in +`build/windows-gpu-coverage` and `test-results/windows-gpu-coverage`. -- Both targets: complete MinGW host build and 19 device code objects using - TheRock `10.2.0a20260929`, HIP `7.17.26386`, Clang 24. The runtime, driver, - patched ZLUDA and OptiScaler versions are unchanged from the validated port. -- gfx1201/RX 9070 XT: 20/20 CTest gates, including 16 physical-GPU gates. -- gfx1200: 4/4 software gates; no gfx1200 kernel execution. -- MSVC v143 `14.44.35207`: binary16 CPU test passes; CUDA MMA probe and D3D12 - hook translation units compile. This is not a complete MSVC build claim. -- K: 12 finite 4K frames at 2259x1271 input, burst and feature recreation; - all six async RGB images exactly match their synchronous controls. -- M: eight finite 4K control/candidate frames at the same input dimensions; - all four new-build RGB images exactly match the prior validated build. - Both presets log every required native layer, zero CPU image copies and - current-frame output. -- Original gfx1201 objects remain byte-identical on disk. In fresh builds, - `.text` and `.rodata*` match the prior gfx1201 objects for all 19 kernels; - fresh ELF files have different overall hashes. gfx1200 and gfx1201 have - distinct correct ELF targets. Matching instruction sections are compile - evidence, not a substitute for testing gfx1200 hardware. -- A gfx1200 HIP probe on the physical gfx1201 fails cleanly with exit 4 before - any GPU allocation/kernel execution. This tests refusal, not gfx1200 support. -- Both target packages pass install/restore checks in isolated workspace - fixtures; no game is launched. Only manifest-listed runtime files install. - Packaging rejects a mismatched destination or native kernel directory before - writing files. Every packaged code object has the expected ELF target. +The first hardware gate on another GPU needs no NVIDIA DLL: -Logs are local in `test-results/rdna4-targets-*` and -`test-results/rdna4-target-objects.json`; they contain no public NVIDIA DLLs. +```powershell +.\scripts\windows\test-windows-rdna4.ps1 -RuntimeProfile therock -PackageRoot "$PWD\dist\windows-gfx1101-diagnostics" -ZludaRoot "$PWD\dist\zluda-windows-final" +``` + +Success requires exit 0, correct actual target, exact HIP/PTX/WMMA/pixel checks +and stable interop lifetime. Failure prints one ZIP. Full K/M reference, +temporal, DLSS and game validation remain subsequent gates. diff --git a/docs/windows-rdna4-port.md b/docs/windows-rdna4-port.md index 0d1a2c5..a315ec3 100644 --- a/docs/windows-rdna4-port.md +++ b/docs/windows-rdna4-port.md @@ -1,6 +1,7 @@ # Windows / RDNA4 port -Branch: `windows-rdna4`. Native Windows build targets: **gfx1200 / gfx1201**. +Branch: `windows-rdna4`. Experimental Windows build targets: **gfx1100..gfx1103, +gfx1150..gfx1154, gfx1200 / gfx1201**. Validated physical hardware: Windows 11 x64, RX 9070 XT, **gfx1201**. Priority: correct K, correct M, native Windows, same-frame output, then speed. Upstream integration: [PR #11](https://github.com/countervolts/d4r/pull/11) @@ -18,15 +19,38 @@ explicit limits; the validated package keeps conservative arithmetic. ## Milestone and gates -2026-10-02 target coverage: gfx1200 and gfx1201 now have separate host/device -builds and guarded packages. Both compile; gfx1201 passes 20 CTest gates on -RX 9070 XT, and gfx1200 passes four software-only gates. No physical gfx1200 -test is available. K's new host build passes twelve finite 4K burst/recreation -frames with all six async RGB images exact. M's four new-build 4K RGB images -exactly match the previous build. The fork audit, adopted MSVC fixes, -per-target commands, validation limits and first hardware gate are maintained -in [windows-gpu-support.md](windows-gpu-support.md). RDNA3 runtime/layout changes -from the reviewed fork are not imported. K performance remains open. +2026-10-02 coverage: all eleven listed RDNA3 / RDNA3.5 / RDNA4 targets compile +with separate host/device objects and guarded packages. Only RX 9070 XT has +physical validation here; no other-target GPU workloads are executed. The +newest realdody correction informs a Windows-only gfx11 LDS half-wave exchange +preserving native gfx11 WMMA layout. gfx12 and Linux keep their original +instruction path. A shared registry drives build, runtime, ELF checks and +scripts. HIP selection matches the D3D12 LUID before selecting an ordinal, +including same-target multi-GPU systems. + +Issue #10 comment 5952822115 contains 32 logs / eight attempts: all discover +RX 9060 XT / gfx1200 but reject the old gfx1201 package before NGX runtime +initialization. Three also load a vkd3d-proton D3D12 proxy. The installer now +catches these conditions before changing game files and prints one diagnostic +ZIP. The next community gate is the gfx1200 package on native D3D12; those logs +do not establish a native K/M crash. Coverage, sources, fork credit, commands +and hardware limits are in [windows-gpu-support.md](windows-gpu-support.md). + +The rebuilt gfx1201 passes 20/20 CTest gates on RX 9070 XT. K completes twelve +finite 4K burst/recreation frames at 2259x1271 input, with all six async RGB +images exactly matching sync. M completes eight finite control/candidate 4K +frames, with all four candidate RGB images matching the prior build exactly. +Both log every required native layer, zero CPU frame copies and frame age 0. +All 209 compiled objects have the correct ELF target. For all nineteen gfx1201 +objects, instruction/constant sections match the previous validated build. + +All eleven game directories stage strict K/M baselines without NVIDIA DLLs. +An isolated gfx1201 install/restore fixture preserves original bytes; wrong +gfx1200-on-gfx1201 and dummy vkd3d-proxy fixtures refuse before game changes +and produce one diagnostic ZIP each. The real game installation is untouched. +Results: `test-results/windows-gpu-coverage`; build logs and the compile-only +matrix report: `build/windows-gpu-coverage`. No physical multi-GPU, APU or +other-target validation is claimed. K performance remains open. **Current milestone: K performance remains open.** The latest unprofiled 4K Silent Hill 2 scene capture records 65.250 FPS. Functional K/M and the diff --git a/kernels/common/wave32_exchange.h b/kernels/common/wave32_exchange.h new file mode 100644 index 0000000..16205db --- /dev/null +++ b/kernels/common/wave32_exchange.h @@ -0,0 +1,31 @@ +#pragma once +#include + +// Windows gfx11 portability workaround informed by realdody/d4r 1897b4a. +// The reported permlanex16 failure duplicates the even K value into the odd +// slot. Retain the native gfx11 layout and exchange through per-wave LDS. +// Other platforms and gfx12 keep their original instruction sequence. +__attribute__((device, always_inline)) static inline uint32_t d4r_wave32_other_half(uint32_t value) +{ +#if defined(D4R_WINDOWS_GFX11_LANE_EXCHANGE) && defined(__GFX11__) + __attribute__((shared)) volatile uint32_t scratch[512]; + const uint32_t tid = __builtin_amdgcn_workitem_id_x() + __builtin_amdgcn_workgroup_size_x() * + (__builtin_amdgcn_workitem_id_y() + __builtin_amdgcn_workgroup_size_y() * __builtin_amdgcn_workitem_id_z()); + // All native K/M launch bounds are <=512. Do not silently corrupt LDS if + // this helper is later reused by a larger workgroup. + if (tid >= 512) __builtin_trap(); + scratch[tid] = value; + __builtin_amdgcn_fence(__ATOMIC_RELEASE, "wavefront"); + __builtin_amdgcn_wave_barrier(); + __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "wavefront"); + const uint32_t result = scratch[tid ^ 16u]; + // Protect scratch reuse by the next exchange in this wave. No block-wide + // rendezvous: different waves can reach GEMMs in a different order. + __builtin_amdgcn_fence(__ATOMIC_RELEASE, "wavefront"); + __builtin_amdgcn_wave_barrier(); + __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "wavefront"); + return result; +#else + return __builtin_amdgcn_permlanex16(value, value, 0x76543210u, 0xfedcba98u, false, false); +#endif +} diff --git a/kernels/common/wmma_backend.h b/kernels/common/wmma_backend.h index 3165deb..21a3ee6 100644 --- a/kernels/common/wmma_backend.h +++ b/kernels/common/wmma_backend.h @@ -1,5 +1,6 @@ #pragma once #include +#include "wave32_exchange.h" // d4r's gfx11 interface: each lane supplies 16 FP16 K operands. The 8 FP32 // outputs have logical row 2*i + lane/16 and logical column lane%16. @@ -13,8 +14,7 @@ using d4r_wmma_f8 = float __attribute__((ext_vector_type(8))); __attribute__((device, always_inline)) static inline float d4r_wmma_other_half(float value) { const uint32_t bits = __builtin_bit_cast(uint32_t, value); - const uint32_t swapped = __builtin_amdgcn_permlanex16( - bits, bits, 0x76543210u, 0xfedcba98u, false, false); + const uint32_t swapped = d4r_wave32_other_half(bits); return __builtin_bit_cast(float, swapped); } @@ -49,7 +49,7 @@ d4r_wmma_legacy_layout(d4r_wmma_h16 a, d4r_wmma_h16 b, d4r_wmma_f8 c) result[i] = (row >> 3) == group ? d12[new_index] : remote_d[new_index]; } return result; -#elif defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || defined(__gfx1103__) +#elif defined(__GFX11__) return __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(a, b, c); #else #error "d4r WMMA backend requires a validated gfx11 or gfx12 target" diff --git a/kernels/common/wmma_layout.h b/kernels/common/wmma_layout.h index 6335874..215ce7d 100644 --- a/kernels/common/wmma_layout.h +++ b/kernels/common/wmma_layout.h @@ -18,6 +18,7 @@ // Only clang builtins are used (the texture tails compile with -nogpuinc -nogpulib). #pragma once #include +#include "wave32_exchange.h" #ifndef D4R_WMMA_LAYOUT #if defined(__GFX12__) @@ -56,7 +57,7 @@ WM_FN uint32_t wm_half() // value of the lane with the same index in the other half of the wave WM_FN uint32_t wm_other_half(uint32_t v) { - return __builtin_amdgcn_permlanex16(v, v, 0x76543210u, 0xfedcba98u, false, false); + return d4r_wave32_other_half(v); } WM_FN float wm_other_half_f(float v) { diff --git a/kernels/k/pwin_common.h b/kernels/k/pwin_common.h index 1bccb17..4b174c8 100644 --- a/kernels/k/pwin_common.h +++ b/kernels/k/pwin_common.h @@ -122,7 +122,7 @@ __device__ __forceinline__ acc8v splat8(float v) __device__ __forceinline__ uint32_t other_half(uint32_t v) { - return __builtin_amdgcn_permlanex16(v, v, 0x76543210u, 0xfedcba98u, false, false); + return wm_other_half(v); } __device__ __forceinline__ uint32_t pack2(half_t lo, half_t hi) diff --git a/scripts/windows/build-native-texture.ps1 b/scripts/windows/build-native-texture.ps1 index 20ae134..1b9e4cd 100644 --- a/scripts/windows/build-native-texture.ps1 +++ b/scripts/windows/build-native-texture.ps1 @@ -5,10 +5,12 @@ param( [string]$HipRoot = 'C:\Program Files\AMD\ROCm\7.2', [string]$ZludaRoot, [ValidateSet('cu','wgp')][string]$ShaderMode='cu', - [ValidateSet('gfx1200','gfx1201')][string]$GpuArch='gfx1201', + [string]$GpuArch='gfx1201', [string]$OutputDirectory ) $ErrorActionPreference='Stop' +. (Join-Path $PSScriptRoot 'gpu-target.ps1') +[void](Get-D4RGpuTargetInfo $GpuArch) $repo=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) if (!$ZludaRoot) { $ZludaRoot=Join-Path $repo 'dist/zluda-windows-final' } if (!$OutputDirectory) { $OutputDirectory=Join-Path $repo "build/private-textures-$GpuArch" } diff --git a/scripts/windows/build-windows-gpu-matrix.ps1 b/scripts/windows/build-windows-gpu-matrix.ps1 new file mode 100644 index 0000000..d71b938 --- /dev/null +++ b/scripts/windows/build-windows-gpu-matrix.ps1 @@ -0,0 +1,44 @@ +[CmdletBinding()] +param( + [string[]]$GpuArchitectures, + [ValidateSet('stable','therock')][string]$RuntimeProfile='therock', + [string]$ZludaRoot, [string]$DlssDll, + [switch]$PackageGames, + [string]$BuildRoot, [string]$OutputRoot, [string]$ArchiveDirectory +) +$ErrorActionPreference='Stop' +. (Join-Path $PSScriptRoot 'gpu-target.ps1') +$repo=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) +if (!$GpuArchitectures) { $GpuArchitectures=@(Get-D4RGpuTargets | ForEach-Object { $_.architecture }) } +foreach ($arch in $GpuArchitectures) { [void](Get-D4RGpuTargetInfo $arch) } +if (@($GpuArchitectures | Select-Object -Unique).Count -ne $GpuArchitectures.Count) { throw 'Duplicate GPU target.' } +if (!$BuildRoot) { $BuildRoot=Join-Path $repo 'build/windows-gpu-coverage' } +if (!$OutputRoot) { $OutputRoot=Join-Path $repo 'dist/windows-gpu-coverage' } +if (!$ZludaRoot) { $ZludaRoot=Join-Path $repo 'dist/zluda-windows-final' } +if ($ArchiveDirectory -and !$PackageGames) { throw 'ArchiveDirectory requires PackageGames.' } +if ($PackageGames -and (!$DlssDll -or ![IO.Path]::IsPathRooted($DlssDll) -or !(Test-Path -LiteralPath $DlssDll))) { + throw 'PackageGames requires an absolute path to your local nvngx_dlss.dll; it will not be redistributed.' +} +New-Item -ItemType Directory -Force $BuildRoot | Out-Null +if ($ArchiveDirectory) { New-Item -ItemType Directory -Force $ArchiveDirectory | Out-Null } +$results=[Collections.Generic.List[object]]::new() +foreach ($arch in $GpuArchitectures) { + $build="$BuildRoot-$arch"; $diag="$OutputRoot-$arch" + $log=Join-Path $BuildRoot "$arch-build.log" + Write-Host "Building $arch (compile/package only; no GPU workloads). Log: $log" + & (Join-Path $PSScriptRoot 'build-windows-rdna4.ps1') -RuntimeProfile $RuntimeProfile -GpuArch $arch ` + -ZludaRoot $ZludaRoot -BuildDirectory $build -InstallDirectory $diag *> $log + $row=[ordered]@{architecture=$arch; compiled=$true; hardwareTestExecuted=$false; buildDirectory=$build; diagnosticRoot=$diag; log=$log} + if ($PackageGames) { + $k="$BuildRoot-native-k-$arch"; $m="$BuildRoot-native-m-$arch"; $game="$OutputRoot-game-$arch" + & (Join-Path $PSScriptRoot 'stage-native-k.ps1') -DlssDll $DlssDll -PackageRoot $diag -GpuArch $arch -OutputDirectory $k + & (Join-Path $PSScriptRoot 'stage-native-m.ps1') -DlssDll $DlssDll -PackageRoot $diag -GpuArch $arch -OutputDirectory $m + $packageArgs=@{DiagnosticRoot=$diag; GpuArch=$arch; PackageRoot=$game; ZludaRoot=$ZludaRoot; KernelKRoot=$k; KernelMRoot=$m} + if ($ArchiveDirectory) { $packageArgs.ArchivePath=Join-Path $ArchiveDirectory "d4r-windows-$arch.zip" } + & (Join-Path $PSScriptRoot 'package-windows-game.ps1') @packageArgs + $row.gamePackage=$game + } + $results.Add($row) + $results | ConvertTo-Json -Depth 5 | Set-Content (Join-Path $BuildRoot 'matrix-build.json') -Encoding UTF8 + Write-Host "Compiled $arch; hardware validation remains separate." +} diff --git a/scripts/windows/build-windows-rdna4.ps1 b/scripts/windows/build-windows-rdna4.ps1 index a60a752..80667b7 100644 --- a/scripts/windows/build-windows-rdna4.ps1 +++ b/scripts/windows/build-windows-rdna4.ps1 @@ -2,13 +2,15 @@ param( [string]$HipRoot = $env:HIP_PATH, [ValidateSet('stable', 'therock')][string]$RuntimeProfile = 'stable', - [ValidateSet('gfx1200','gfx1201')][string]$GpuArch = 'gfx1201', + [string]$GpuArch = 'gfx1201', [string]$ZludaRoot, [string]$ToolchainRoot, [string]$BuildDirectory, [string]$InstallDirectory ) $ErrorActionPreference = 'Stop' +. (Join-Path $PSScriptRoot 'gpu-target.ps1') +[void](Get-D4RGpuTargetInfo $GpuArch) $repo = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) if ($RuntimeProfile -eq 'therock') { if (!$PSBoundParameters.ContainsKey('HipRoot')) { $HipRoot = Join-Path $repo '.tools/therock-10.2.0a20260929/_rocm_sdk_core' } diff --git a/scripts/windows/gpu-preflight.ps1 b/scripts/windows/gpu-preflight.ps1 new file mode 100644 index 0000000..c432109 --- /dev/null +++ b/scripts/windows/gpu-preflight.ps1 @@ -0,0 +1,58 @@ +# Read-only inventory before touching game files. It allocates no GPU memory +# and executes no kernels; successful discovery is not hardware validation. +function Invoke-D4RGpuPreflight([string]$PackageRoot, [string]$GameDirectory, [string]$Target, [string]$OutputDirectory) { + New-Item -ItemType Directory -Force $OutputDirectory | Out-Null + $result=[ordered]@{passed=$false; target=$Target; gpuWorkExecuted=$false; failure=$null} + try { + $proxy=Join-Path $GameDirectory 'd3d12.dll' + if (Test-Path -LiteralPath $proxy) { + $item=Get-Item -LiteralPath $proxy + $result.localD3D12=@{path=$proxy; sha256=(Get-FileHash -LiteralPath $proxy).Hash; product=$item.VersionInfo.ProductName} + # Identify known Vulkan/Wine proxies by their own strings. Do not + # load, remove or replace a third-party game DLL. + if ($item.Length -le 64MB -and [Text.Encoding]::ASCII.GetString([IO.File]::ReadAllBytes($proxy)) -match '(?i)vkd3d|wined3d') { + throw 'Local d3d12.dll identifies as vkd3d/Wine. Restore the game native D3D12 installation before using this Windows backend.' + } + } + $info=[Diagnostics.ProcessStartInfo]::new() + $info.FileName=Join-Path $PackageRoot 'd4r/d4r_gpu_inventory.exe' + $hip=Join-Path $PackageRoot 'd4r/hip' + foreach ($value in @($hip,$OutputDirectory)) { if ($value.Contains('"')) { throw 'Unsupported quote in inventory path.' } } + $info.Arguments='--hip-root "' + $hip + '" --output-dir "' + $OutputDirectory + '"' + $info.UseShellExecute=$false; $info.CreateNoWindow=$true + $info.EnvironmentVariables['D4R_DIAG_DIR']=$OutputDirectory + $info.RedirectStandardOutput=$true; $info.RedirectStandardError=$true + $process=[Diagnostics.Process]::new(); $process.StartInfo=$info + try { + if (!$process.Start()) { throw 'Cannot start GPU inventory.' } + $stdout=$process.StandardOutput.ReadToEndAsync(); $stderr=$process.StandardError.ReadToEndAsync() + $timedOut=!$process.WaitForExit(30000) + if ($timedOut) { $process.Kill(); $process.WaitForExit() } + $stdout.GetAwaiter().GetResult() | Set-Content (Join-Path $OutputDirectory 'gpu-inventory.stdout.log') -Encoding UTF8 + $stderr.GetAwaiter().GetResult() | Set-Content (Join-Path $OutputDirectory 'gpu-inventory.stderr.log') -Encoding UTF8 + $result.exitCode=$process.ExitCode + $result.exceptionOrExitHex='0x{0:X8}' -f [BitConverter]::ToUInt32([BitConverter]::GetBytes([int]$process.ExitCode),0) + if ($timedOut) { $result.timedOut=$true; throw 'HIP inventory timed out after 30 seconds.' } + if ($process.ExitCode) { throw "HIP inventory failed (exit $($process.ExitCode))." } + } finally { $process.Dispose() } + $inventory=Get-Content (Join-Path $OutputDirectory 'gpu-inventory.json') -Raw | ConvertFrom-Json + $result.inventory=$inventory + if ($inventory.compiledTarget -ne $Target) { throw "Inventory executable targets $($inventory.compiledTarget); package requires $Target." } + $matching=@($inventory.devices | Where-Object { $_.architecture -eq $Target -and $_.wavefrontSize -eq 32 }) + if (!$matching.Count) { + $actual=@($inventory.devices | ForEach-Object { "$($_.architecture) ($($_.name))" }) -join ', ' + throw "GPU target mismatch: detected $actual; package requires $Target. Use the package for the actual HIP architecture; do not rename code objects." + } + $result.passed=$true + Write-Host "D4R_PREFLIGHT target=$Target devices=$($matching.Count) gpu_work_executed=0" + } catch { + $result.failure=$_.Exception.Message + $result | ConvertTo-Json -Depth 8 | Set-Content (Join-Path $OutputDirectory 'preflight.json') -Encoding UTF8 + try { Get-CimInstance Win32_VideoController | Select-Object Name,DriverVersion,DriverDate | ConvertTo-Json | + Set-Content (Join-Path $OutputDirectory 'driver.json') -Encoding UTF8 } catch {} + $bundle=$OutputDirectory+'.zip' + Compress-Archive -Path (Join-Path $OutputDirectory '*') -DestinationPath $bundle -Force + throw "$($result.failure) Diagnostic bundle: $bundle" + } + $result | ConvertTo-Json -Depth 8 | Set-Content (Join-Path $OutputDirectory 'preflight.json') -Encoding UTF8 +} diff --git a/scripts/windows/gpu-target.ps1 b/scripts/windows/gpu-target.ps1 index e05ab38..3656690 100644 --- a/scripts/windows/gpu-target.ps1 +++ b/scripts/windows/gpu-target.ps1 @@ -1,19 +1,31 @@ # Shared by the build, diagnostics, staging and game package scripts. +function Get-D4RGpuTargets { + $registry=Join-Path $PSScriptRoot 'gpu-targets.json' + if (!(Test-Path -LiteralPath $registry)) { $registry=Join-Path $PSScriptRoot '../../tools/windows/gpu-targets.json' } + $data=Get-Content -LiteralPath $registry -Raw | ConvertFrom-Json + if ($data.schema -ne 1) { throw 'Unsupported GPU target registry schema.' } + return $data.targets +} +function Get-D4RGpuTargetInfo([string]$Target) { + $info=@(Get-D4RGpuTargets | Where-Object { $_.architecture -eq $Target }) + if ($info.Count -ne 1) { throw "Unsupported GPU target: $Target" } + return $info[0] +} function Get-D4RGpuTarget([string]$Root, [string]$Requested) { $configured=$null $targetFile=Join-Path $Root 'gpu-target.json' if (Test-Path -LiteralPath $targetFile) { $configured=(Get-Content -LiteralPath $targetFile -Raw | ConvertFrom-Json).architecture - if ($configured -notin @('gfx1200','gfx1201')) { throw "Unsupported package GPU target: $configured" } + [void](Get-D4RGpuTargetInfo $configured) } $target=if ($Requested) { $Requested } elseif ($configured) { $configured } else { 'gfx1201' } - if ($target -notin @('gfx1200','gfx1201')) { throw "Unsupported GPU target: $target" } - if (!$configured -and $target -ne 'gfx1201') { throw 'A gfx1200 build requires gpu-target.json; rebuild with -GpuArch gfx1200.' } + [void](Get-D4RGpuTargetInfo $target) + if (!$configured -and $target -ne 'gfx1201') { throw "A $target build requires gpu-target.json; rebuild with -GpuArch $target." } if ($configured -and $target -ne $configured) { throw "Requested $target differs from package target $configured" } return $target } function Assert-D4RCodeObjectTarget([string]$Path, [string]$Target) { - if ($Target -notin @('gfx1200','gfx1201')) { throw "Unsupported GPU target: $Target" } + $info=Get-D4RGpuTargetInfo $Target $reader=[IO.BinaryReader]::new([IO.File]::OpenRead($Path)) try { $header=$reader.ReadBytes(64) } finally { $reader.Dispose() } if ($header.Length -ne 64 -or $header[0] -ne 127 -or $header[1] -ne 69 -or $header[2] -ne 76 -or $header[3] -ne 70 -or @@ -21,6 +33,6 @@ function Assert-D4RCodeObjectTarget([string]$Path, [string]$Target) { throw "Invalid AMDGPU ELF64 code object: $Path" } # LLVM AMDGPU e_flags machine IDs; inspect, never patch or spoof them. - $expected=if ($Target -eq 'gfx1200') { 0x48 } else { 0x4e } + $expected=$info.elfMachine if ($header[48] -ne $expected) { throw "Code object does not target ${Target}: $Path (ELF machine=$($header[48]))" } } diff --git a/scripts/windows/package-windows-game.ps1 b/scripts/windows/package-windows-game.ps1 index 8c33dc7..98b2aac 100644 --- a/scripts/windows/package-windows-game.ps1 +++ b/scripts/windows/package-windows-game.ps1 @@ -2,7 +2,7 @@ param( [string]$PackageRoot, [string]$DiagnosticRoot, [string]$OptiScalerRoot, [string]$ZludaRoot, [string]$HipRoot, [string]$KernelKRoot, [string]$KernelMRoot, - [ValidateSet('gfx1200','gfx1201')][string]$GpuArch, + [string]$GpuArch, [string]$ArchivePath ) $ErrorActionPreference = 'Stop' @@ -15,7 +15,7 @@ foreach ($name in $defaults.Keys) { if (!(Get-Variable -Name $name -ValueOnly)) { Set-Variable -Name $name -Value (Join-Path $repo $defaults[$name]) } } $GpuArch=Get-D4RGpuTarget $DiagnosticRoot $GpuArch -if ($GpuArch -eq 'gfx1200' -and !$PSBoundParameters.ContainsKey('PackageRoot')) { $PackageRoot += '-gfx1200' } +if ($GpuArch -ne 'gfx1201' -and !$PSBoundParameters.ContainsKey('PackageRoot')) { $PackageRoot += "-$GpuArch" } if (!$KernelKRoot) { $KernelKRoot=Join-Path $repo "build/native-k-$GpuArch" } if (!$KernelMRoot) { $KernelMRoot=Join-Path $repo "build/native-m-$GpuArch" } foreach ($directory in @($KernelKRoot,$KernelMRoot)) { @@ -46,6 +46,7 @@ Stage (Join-Path $DiagnosticRoot 'bin/d4r_nvngx.dll') 'd4r/_nvngx.dll' Stage (Join-Path $DiagnosticRoot "bin/pixel_convert_${GpuArch}.hsaco") "d4r/pixel_convert_${GpuArch}.hsaco" Stage (Join-Path $DiagnosticRoot 'nvapi-compat/nvapi64.dll') 'd4r/nvapi/nvapi64.dll' Stage (Join-Path $DiagnosticRoot 'bin/d4r_debug_launcher.exe') 'd4r/d4r_debug_launcher.exe' +Stage (Join-Path $DiagnosticRoot 'bin/d4r_gpu_inventory.exe') 'd4r/d4r_gpu_inventory.exe' foreach ($file in @('nvcuda.dll','nvapi64.dll','libc++.dll','libunwind.dll')) { Stage (Join-Path $ZludaRoot $file) "d4r/zluda/$file" } foreach ($file in @('amdhip64_7.dll','amd_comgr.dll','hiprtc0717.dll','hiprtc-builtins0717.dll','rocm_kpack.dll','AMD.ROCM.Comgr.MANIFEST','.hipVersion')) { Stage (Join-Path $HipRoot "bin/$file") "d4r/hip/bin/$file" @@ -66,6 +67,8 @@ if ($entries.Count -ne 16) { throw "Expected 11 K and 5 M native objects; found $stagedFiles.Add('d4r/native/d4r-kernels.txt') Stage (Join-Path $repo 'scripts/windows/windows-game.ps1') 'windows-game.ps1' Stage (Join-Path $repo 'scripts/windows/gpu-target.ps1') 'gpu-target.ps1' +Stage (Join-Path $repo 'scripts/windows/gpu-preflight.ps1') 'gpu-preflight.ps1' +Stage (Join-Path $repo 'tools/windows/gpu-targets.json') 'gpu-targets.json' if (Test-Path -LiteralPath (Join-Path $DiagnosticRoot 'gpu-target.json')) { Stage (Join-Path $DiagnosticRoot 'gpu-target.json') 'gpu-target.json' } else { diff --git a/scripts/windows/stage-native-k.ps1 b/scripts/windows/stage-native-k.ps1 index f470c50..12c102b 100644 --- a/scripts/windows/stage-native-k.ps1 +++ b/scripts/windows/stage-native-k.ps1 @@ -1,5 +1,5 @@ [CmdletBinding()] -param([Parameter(Mandatory=$true)][string]$DlssDll, [string]$PackageRoot, [string]$OutputDirectory, [ValidateSet('gfx1200','gfx1201')][string]$GpuArch) +param([Parameter(Mandatory=$true)][string]$DlssDll, [string]$PackageRoot, [string]$OutputDirectory, [string]$GpuArch) $ErrorActionPreference = 'Stop' . (Join-Path $PSScriptRoot 'gpu-target.ps1') $repo = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) diff --git a/scripts/windows/stage-native-m.ps1 b/scripts/windows/stage-native-m.ps1 index c882b25..d856f8c 100644 --- a/scripts/windows/stage-native-m.ps1 +++ b/scripts/windows/stage-native-m.ps1 @@ -1,5 +1,5 @@ [CmdletBinding()] -param([Parameter(Mandatory=$true)][string]$DlssDll, [string]$PackageRoot, [string]$OutputDirectory, [ValidateSet('gfx1200','gfx1201')][string]$GpuArch) +param([Parameter(Mandatory=$true)][string]$DlssDll, [string]$PackageRoot, [string]$OutputDirectory, [string]$GpuArch) $ErrorActionPreference = 'Stop' . (Join-Path $PSScriptRoot 'gpu-target.ps1') $repo = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) diff --git a/scripts/windows/test-windows-rdna4.ps1 b/scripts/windows/test-windows-rdna4.ps1 index 6f71bfd..9bea7e2 100644 --- a/scripts/windows/test-windows-rdna4.ps1 +++ b/scripts/windows/test-windows-rdna4.ps1 @@ -4,7 +4,7 @@ param( [ValidateSet('stable', 'therock')][string]$RuntimeProfile = 'stable', [string]$ZludaRoot, [string]$PackageRoot, - [ValidateSet('gfx1200','gfx1201')][string]$GpuArch, + [string]$GpuArch, [string]$OutputDirectory, [string]$NgxCore = $env:D4R_NGX_CORE, [string]$DlssDll = $env:D4R_DLSS_DLL, diff --git a/scripts/windows/windows-game.ps1 b/scripts/windows/windows-game.ps1 index 38f83d4..83dbd75 100644 --- a/scripts/windows/windows-game.ps1 +++ b/scripts/windows/windows-game.ps1 @@ -27,6 +27,7 @@ param( ) $ErrorActionPreference = 'Stop' . (Join-Path $PSScriptRoot 'gpu-target.ps1') +. (Join-Path $PSScriptRoot 'gpu-preflight.ps1') if ($ProfileGpuBoundary -and !$AsyncInterop) { throw '-ProfileGpuBoundary requires -AsyncInterop.' } if ($ProfileKernels -and $ProfileKernelsDeferred) { throw 'Choose either serializing or deferred kernel profiling.' } if ($ProfileLegacyStream -and !$ProfileKernelsDeferred) { throw '-ProfileLegacyStream requires -ProfileKernelsDeferred.' } @@ -74,6 +75,7 @@ foreach ($file in $metadata.files) { if ((Get-FileHash -LiteralPath (Join-Path $package $file.path) -Algorithm SHA256).Hash -ne $file.sha256) { throw "Package file changed; rebuild package: $($file.path)" } } if ((Get-FileHash -LiteralPath $DlssDll -Algorithm SHA256).Hash -ne $metadata.dlssSha256) { throw 'DLSS DLL does not match the validated 310.9.1 native manifest. Rebuild and validate native manifests for this DLL first.' } +Invoke-D4RGpuPreflight $package $game $GpuArch (Join-Path $package ('results/preflight-' + (Get-Date -Format 'yyyyMMdd-HHmmss-fff'))) $textureFiles=@(); $textureManifest=@() if ($LocalTextureKernels) { $LocalTextureKernels=(Get-Item -LiteralPath $LocalTextureKernels -ErrorAction Stop).FullName @@ -193,6 +195,7 @@ $settings = @{ # false have the same numerical effect; unset reuses the validated cache. D4R_ZLUDA_IGNORE_DENORMAL=$null; D4R_ZLUDA_FAST_MATH=$null; D4R_ZLUDA_WMMA_F32ACC=$null; D4R_ZLUDA_WAVE64=$null; + D4R_ZLUDA_WMMA_LAYOUT=$null; D4R_QUIET_API=$(if ($VerboseRuntime) { $null } else { '1' }); D4R_VALIDATE_OUTPUT=$(if ($ValidateOutput) { '1' } else { $null }); D4R_PROFILE_STAGES=$(if ($ProfileStages) { '1' } else { $null }); diff --git a/tools/windows/build_native_texture.py b/tools/windows/build_native_texture.py index 67ff26c..0b0c11f 100644 --- a/tools/windows/build_native_texture.py +++ b/tools/windows/build_native_texture.py @@ -1,4 +1,4 @@ -"""Build a private RDNA4 texture override from the user's local DLSS DLL. +"""Build a private target-specific texture override from the user's local DLSS DLL. The resulting PTX/code object contains NVIDIA code and must remain local. No Wine, shell script, GPU, CUDA SDK or system installation is required. @@ -13,13 +13,14 @@ import re import shutil import subprocess import sys +from gpu_targets import TARGETS def require_code_object_target(path, target): with path.open('rb') as file: header = file.read(64) if (len(header) != 64 or header[:6] != b'\x7fELF\x02\x01' or - header[18:20] != b'\xe0\x00' or header[48] != {'gfx1200': 0x48, 'gfx1201': 0x4e}[target]): + header[18:20] != b'\xe0\x00' or header[48] != TARGETS[target]['elfMachine']): raise RuntimeError(f'Code object does not target {target}: {path}') @@ -30,7 +31,7 @@ def main(): parser.add_argument('--zluda-root', type=pathlib.Path, required=True) parser.add_argument('--kernel', required=True) parser.add_argument('--output-directory', type=pathlib.Path, required=True) - parser.add_argument('--gpu-arch', choices=['gfx1200', 'gfx1201'], default='gfx1201') + parser.add_argument('--gpu-arch', choices=tuple(TARGETS), default='gfx1201') parser.add_argument('--shader-mode', choices=['cu', 'wgp'], default='cu') args = parser.parse_args() if not re.fullmatch(r'hiluma_engine_(?:output_depth(?:inv|reg)_mv(?:hi|lo)_(?:hdr|ldr)(?:_max)?|' diff --git a/tools/windows/d4r_nvngx_windows.cpp b/tools/windows/d4r_nvngx_windows.cpp index bf99083..2f3696a 100644 --- a/tools/windows/d4r_nvngx_windows.cpp +++ b/tools/windows/d4r_nvngx_windows.cpp @@ -87,7 +87,7 @@ API unsigned NVSDK_NGX_D3D12_GetFeatureRequirements(IDXGIAdapter* adapter, const DXGI_ADAPTER_DESC desc{}; d4r::win::dx(adapter->GetDesc(&desc), "NGX requirements adapter"); // OptiScaler can override the DXGI vendor ID. The physical HIP // architecture and full adapter LUID are the authoritative identity. - d4r::diag::HipApi hip(d4r::win::env_path("D4R_HIP_ROOT")); hipDeviceProp_t props{}; hip.select_architecture(-1,props); + d4r::diag::HipApi hip(d4r::win::env_path("D4R_HIP_ROOT")); hipDeviceProp_t props{}; hip.select_architecture(-1,props,&desc.AdapterLuid); requirements->FeatureSupported = std::memcmp(&desc.AdapterLuid,props.luid,sizeof(LUID)) ? 4u : 0u; std::strcpy(requirements->MinOSVersion,"10.0.22000.0"); std::printf("D4R_REQUIREMENTS feature=%u support=%u architecture=%s\n",info->FeatureID,requirements->FeatureSupported,d4r::diag::HipApi::target_arch()); diff --git a/tools/windows/gpu-targets.json b/tools/windows/gpu-targets.json new file mode 100644 index 0000000..62ec491 --- /dev/null +++ b/tools/windows/gpu-targets.json @@ -0,0 +1,16 @@ +{ + "schema": 1, + "targets": [ + {"architecture":"gfx1100","elfMachine":65,"isaMajor":11,"family":"RDNA3"}, + {"architecture":"gfx1101","elfMachine":70,"isaMajor":11,"family":"RDNA3"}, + {"architecture":"gfx1102","elfMachine":71,"isaMajor":11,"family":"RDNA3"}, + {"architecture":"gfx1103","elfMachine":68,"isaMajor":11,"family":"RDNA3"}, + {"architecture":"gfx1150","elfMachine":67,"isaMajor":11,"family":"RDNA3.5"}, + {"architecture":"gfx1151","elfMachine":74,"isaMajor":11,"family":"RDNA3.5"}, + {"architecture":"gfx1152","elfMachine":85,"isaMajor":11,"family":"RDNA3.5"}, + {"architecture":"gfx1153","elfMachine":88,"isaMajor":11,"family":"RDNA3.5"}, + {"architecture":"gfx1154","elfMachine":87,"isaMajor":11,"family":"RDNA3.5"}, + {"architecture":"gfx1200","elfMachine":72,"isaMajor":12,"family":"RDNA4"}, + {"architecture":"gfx1201","elfMachine":78,"isaMajor":12,"family":"RDNA4"} + ] +} diff --git a/tools/windows/gpu_inventory.cpp b/tools/windows/gpu_inventory.cpp new file mode 100644 index 0000000..580111f --- /dev/null +++ b/tools/windows/gpu_inventory.cpp @@ -0,0 +1,57 @@ +#include "hip_api.h" +#include +#include +#include + +static std::string json_string(const std::string& value) { + std::ostringstream out; out << '"'; + for (unsigned char c : value) { + if (c == '"' || c == '\\') out << '\\' << char(c); + else if (c < 32) out << "\\u" << std::hex << std::setw(4) << std::setfill('0') << unsigned(c); + else out << char(c); + } + out << '"'; return out.str(); +} +int main(int argc, char** argv) { + using namespace d4r::diag; + start(); + try { + Args args(argc, argv); + if (args.output_dir.empty()) throw std::runtime_error("--output-dir is required"); + 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); + HipApi hip(args.hip_root); + hip.check(hip.hipInit(0), "hipInit(inventory)"); + int count=0, version=0, driver=0; + hip.check(hip.hipRuntimeGetVersion(&version), "hipRuntimeGetVersion"); + hip.check(hip.hipDriverGetVersion(&driver), "hipDriverGetVersion"); + std::printf("GPU_INVENTORY_RUNTIME 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)); + hip.check(hip.hipGetDeviceCount(&count), "hipGetDeviceCount"); + std::ostringstream json; + json << "{\"schema\":1,\"runtimeVersion\":" << version << ",\"driverApiVersion\":" << driver << ",\"compiledTarget\":" + << json_string(D4R_TARGET_ARCH) << ",\"gpuWorkExecuted\":false,\"devices\":["; + for (int i=0; i #include #include +#include "gpu_targets.generated.h" #ifndef D4R_TARGET_ARCH #define D4R_TARGET_ARCH "gfx1201" @@ -11,7 +12,8 @@ namespace d4r::diag { inline bool supported_gpu_target(const std::string& target) { - return target == "gfx1200" || target == "gfx1201"; + for (const auto& row : gpu_targets) if (target == row.architecture) return true; + return false; } inline std::string gpu_architecture(const std::string& name) { return name.substr(0, name.find(':')); @@ -29,11 +31,8 @@ inline std::string code_object_target(const std::filesystem::path& path) { header[0] != 0x7f || header[1] != 'E' || header[2] != 'L' || header[3] != 'F' || header[4] != 2 || header[5] != 1 || header[18] != 0xe0 || header[19] != 0) throw std::runtime_error("Invalid AMDGPU ELF64 code object: " + path.string()); - switch (header[48]) { - case 0x48: return "gfx1200"; - case 0x4e: return "gfx1201"; - default: throw std::runtime_error("Unsupported AMDGPU code-object target: " + path.string()); - } + for (const auto& row : gpu_targets) if (header[48] == row.elf_machine) return row.architecture; + throw std::runtime_error("Unsupported AMDGPU code-object target: " + path.string()); } inline void require_code_object_target(const std::filesystem::path& path, const std::string& expected) { require_gpu_target(code_object_target(path), expected); diff --git a/tools/windows/gpu_target_probe.cpp b/tools/windows/gpu_target_probe.cpp index ffcdf6c..e3f60e9 100644 --- a/tools/windows/gpu_target_probe.cpp +++ b/tools/windows/gpu_target_probe.cpp @@ -5,10 +5,11 @@ int main(int argc, char** argv) { using namespace d4r::diag; try { unsigned rejected = 0, objects = 0; - for (const char* target : {"gfx1200", "gfx1201"}) { + for (const auto& info : gpu_targets) { + const char* target = info.architecture; require_gpu_target(target, target); require_gpu_target(std::string(target) + ":xnack-:sramecc-", target); - for (const char* invalid : {"gfx1101", "gfx12000", "gfx1202", "gfx1250", ""}) { + for (const char* invalid : {"gfx1030", "gfx1104", "gfx1155", "gfx1170", "gfx12000", "gfx1202", "gfx1250", ""}) { try { require_gpu_target(invalid, target); } catch (const std::exception&) { ++rejected; continue; } throw std::runtime_error("Invalid GPU target accepted"); diff --git a/tools/windows/gpu_targets.generated.h.in b/tools/windows/gpu_targets.generated.h.in new file mode 100644 index 0000000..ad0b2b5 --- /dev/null +++ b/tools/windows/gpu_targets.generated.h.in @@ -0,0 +1,7 @@ +#pragma once +#include +namespace d4r::diag { +struct GpuTarget { const char* architecture; unsigned elf_machine; unsigned isa_major; const char* family; }; +inline constexpr std::array gpu_targets = {{ +@d4r_target_cpp_entries@}}; +} diff --git a/tools/windows/gpu_targets.py b/tools/windows/gpu_targets.py new file mode 100644 index 0000000..f7f5a0d --- /dev/null +++ b/tools/windows/gpu_targets.py @@ -0,0 +1,10 @@ +"""Canonical Windows build target registry, shared with CMake and PowerShell.""" +import json +import pathlib + +data = json.loads(pathlib.Path(__file__).with_name('gpu-targets.json').read_text()) +if data['schema'] != 1: + raise RuntimeError('Unsupported GPU target registry schema') +TARGETS = {row['architecture']: row for row in data['targets']} +if len(TARGETS) != len(data['targets']): + raise RuntimeError('Duplicate GPU target') diff --git a/tools/windows/hip_api.h b/tools/windows/hip_api.h index bc32088..c96e539 100644 --- a/tools/windows/hip_api.h +++ b/tools/windows/hip_api.h @@ -42,7 +42,7 @@ struct HipApi { if (result != hipSuccess) throw std::runtime_error(std::string(call) + ": " + hipGetErrorString(result)); } static const char* target_arch() { return D4R_TARGET_ARCH; } - int select_architecture(int requested, hipDeviceProp_t& selected) const { + int select_architecture(int requested, hipDeviceProp_t& selected, const void* required_luid = nullptr) 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); @@ -61,12 +61,14 @@ struct HipApi { i, props.name, props.gcnArchName, props.warpSize, props.totalGlobalMem, props.pciDomainID, props.pciBusID, props.pciDeviceID); const std::string arch = props.gcnArchName; - if (gpu_architecture(arch) == target_arch() && (requested < 0 || requested == i) && found < 0) { + if (gpu_architecture(arch) == target_arch() && (requested < 0 || requested == i) && found < 0 && + (!required_luid || std::memcmp(props.luid, required_luid, sizeof(props.luid)) == 0)) { found = i; selected = props; } } - if (found < 0) throw std::runtime_error(std::string("No selected ") + target_arch() + " device; refusing architecture substitution"); + if (found < 0) throw std::runtime_error(std::string("No selected ") + target_arch() + + " device for this adapter; use the package matching the actual HIP architecture (see GPU log)"); require_gpu_target(selected.gcnArchName, target_arch()); check(hipSetDevice(found), "hipSetDevice"); std::printf("SELECTED HIP ordinal=%d architecture=%s validation=%s\n", found, target_arch(), diff --git a/tools/windows/k_layer_validate.py b/tools/windows/k_layer_validate.py index 88d6c32..a35d864 100644 --- a/tools/windows/k_layer_validate.py +++ b/tools/windows/k_layer_validate.py @@ -1,5 +1,5 @@ #!/usr/bin/env python3 -"""Compare RDNA4 K enc1/enc2 launches with the existing numpy reference.""" +"""Compare native K enc1/enc2 launches with the existing numpy reference.""" import argparse import math import pathlib @@ -73,7 +73,7 @@ def main(): # after each K stage has been compared on captured DLSS weights/activations. if full_abs > 0.001 or full_psnr < 60.0 or merged_abs > 0.001 or merged_psnr < 60.0: raise RuntimeError(f"K {args.kernel_name} numerical mismatch against pwin_model.py") - print(f"PASS K_REFERENCE reference_layout=gfx12 kernel={args.kernel_name} " + print(f"PASS K_REFERENCE reference=logical kernel={args.kernel_name} " "nonzero_weights=1 transformer_executed=1 nan_inf=0") diff --git a/tools/windows/k_replay_validate.py b/tools/windows/k_replay_validate.py index 2b6563a..6c0cbf1 100644 --- a/tools/windows/k_replay_validate.py +++ b/tools/windows/k_replay_validate.py @@ -139,7 +139,7 @@ def validate(layer, directory, output, window_count=1, real_capture=False): if maximum > absolute_limit or psnr < 60: worst = int(difference.argmax()) raise RuntimeError(f'K mismatch at {worst}: reference={expected[worst]}, gpu={actual[worst]}') - print(f'PASS K_REPLAY_REFERENCE kernel={layer} reference_layout=gfx12 nan_inf=0') + print(f'PASS K_REPLAY_REFERENCE kernel={layer} reference=logical nan_inf=0') if __name__ == '__main__': diff --git a/tools/windows/m_replay_validate.py b/tools/windows/m_replay_validate.py index 991d786..0c9c569 100644 --- a/tools/windows/m_replay_validate.py +++ b/tools/windows/m_replay_validate.py @@ -153,7 +153,7 @@ def validate(layer, directory, output, block_count, real_capture, exact=False): maximum, peak, psnr, steps = metrics(expected, actual, f'M_REPLAY_REFERENCE kernel={layer}') if ((exact or not real_capture) and maximum != 0) or (real_capture and (steps.max() > 1 or psnr < 50)): raise RuntimeError(f'M mismatch: exact synthetic / <=1 FP8 step and >=50 dB real gate; peak={peak}') - print(f'PASS M_REPLAY_REFERENCE kernel={layer} blocks={len(pairs)} reference_layout=gfx12 nan_inf=0') + print(f'PASS M_REPLAY_REFERENCE kernel={layer} blocks={len(pairs)} reference=logical nan_inf=0') if __name__ == '__main__': diff --git a/tools/windows/ngx_windows_runtime.h b/tools/windows/ngx_windows_runtime.h index d6acced..6c74e0a 100644 --- a/tools/windows/ngx_windows_runtime.h +++ b/tools/windows/ngx_windows_runtime.h @@ -94,9 +94,9 @@ struct Runtime { if (!d3d) throw std::runtime_error("Null D3D12 device"); try { hipDeviceProp_t props{}; - hip.select_architecture(-1, props); - hip.verbose = cuda.verbose = !std::getenv("D4R_QUIET_API"); const auto luid = adapter_luid(d3d); + hip.select_architecture(-1, props, &luid); + hip.verbose = cuda.verbose = !std::getenv("D4R_QUIET_API"); if (std::memcmp(&luid, props.luid, sizeof(luid))) throw std::runtime_error("D3D12/HIP LUID mismatch"); cuda.check(cuda.cuInit(0), "cuInit(Windows runtime)"); int count = 0; cuda.check(cuda.cuDeviceGetCount(&count), "cuDeviceGetCount"); diff --git a/tools/windows/wmma_gfx1201.hip b/tools/windows/wmma_gfx1201.hip index 42dc911..b61781b 100644 --- a/tools/windows/wmma_gfx1201.hip +++ b/tools/windows/wmma_gfx1201.hip @@ -25,6 +25,7 @@ extern "C" __attribute__((global)) void d4r_wmma_gfx1201( const unsigned int group = lane >> 4; const _Float16* a = reinterpret_cast(a_bits); const _Float16* b = reinterpret_cast(b_bits); +#if defined(__GFX12__) h8 a_frag, b_frag; f8 c_frag; for (unsigned int i = 0; i < 8u; ++i) { @@ -39,6 +40,23 @@ extern "C" __attribute__((global)) void d4r_wmma_gfx1201( b_frag, a_frag, result); for (unsigned int i = 0; i < 8u; ++i) output[row * 16u + group * 8u + i] = result[i]; +#else + // Adapted from realdody/d4r ae530b0: real gfx11 operands/accumulators, + // independent of the shared layout abstraction. + d4r_wmma_h16 a_frag, b_frag; + d4r_wmma_f8 c_frag; + for (unsigned int i = 0; i < 16u; ++i) { + a_frag[i] = a[row * 16u + i]; + b_frag[i] = b[row * 16u + i]; + } + for (unsigned int i = 0; i < 8u; ++i) + c_frag[i] = initial_c[(2u * i + group) * 16u + row]; + d4r_wmma_f8 result = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(a_frag, b_frag, c_frag); + if (steps == 2u) + result = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(a_frag, b_frag, result); + for (unsigned int i = 0; i < 8u; ++i) + output[(2u * i + group) * 16u + row] = result[i]; +#endif } // d4r's original gfx11 fragment contract, exercised independently of K/M.