From 6760294f6079ca6a7215c30c33966287c2c90509 Mon Sep 17 00:00:00 2001 From: Codex Date: Tue, 29 Sep 2026 19:50:50 +0300 Subject: [PATCH] Validate gfx1201 WMMA layout and compile first K kernel --- CMakeLists.txt | 33 ++++++- README.md | 4 +- docs/windows-rdna4-port.md | 46 +++++++-- kernels/common/hip_device_minimal.h | 32 +++++++ kernels/common/wmma_backend.h | 57 +++++++++++ kernels/k/pwin_common.h | 7 +- scripts/windows/test-windows-rdna4.ps1 | 10 +- tools/windows/gfx12_wmma_probe.cpp | 127 +++++++++++++++++++++++++ tools/windows/k_module_probe.cpp | 31 ++++++ tools/windows/wmma_gfx1201.hip | 67 +++++++++++++ 10 files changed, 404 insertions(+), 10 deletions(-) create mode 100644 kernels/common/hip_device_minimal.h create mode 100644 kernels/common/wmma_backend.h create mode 100644 tools/windows/gfx12_wmma_probe.cpp create mode 100644 tools/windows/k_module_probe.cpp create mode 100644 tools/windows/wmma_gfx1201.hip diff --git a/CMakeLists.txt b/CMakeLists.txt index 762e278..9a04875 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -21,7 +21,28 @@ add_custom_command(OUTPUT "${probe_object}" -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 d3d12_hip_interop_probe hip_stream_lifecycle_probe ngx_cuda_init_probe) +set(wmma_object "${CMAKE_CURRENT_BINARY_DIR}/wmma_gfx1201.hsaco") +add_custom_command(OUTPUT "${wmma_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/wmma_gfx1201.hip" + -o "${wmma_object}" + DEPENDS tools/windows/wmma_gfx1201.hip kernels/common/wmma_backend.h VERBATIM) +add_custom_target(d4r_wmma_code_object ALL DEPENDS "${wmma_object}") +# Compile-only K port smoke test. This object is intentionally not installed as +# a native override until its transformer outputs pass the K reference/replay. +set(k_enc1_object "${CMAKE_CURRENT_BINARY_DIR}/dltss_pwin_enc1_layer_gfx1201.hsaco") +add_custom_command(OUTPUT "${k_enc1_object}" + COMMAND "${D4R_HIP_CLANG}" -x hip -std=c++17 --offload-arch=gfx1201 + --offload-device-only --no-gpu-bundle-output -mno-wavefrontsize64 + -nogpuinc -nogpulib -O2 -DD4R_DEVICE_ONLY_MINIMAL + "${CMAKE_CURRENT_SOURCE_DIR}/kernels/k/dltss_pwin_enc1_layer.hip" + -o "${k_enc1_object}" + DEPENDS kernels/k/dltss_pwin_enc1_layer.hip kernels/k/pwin_layer.h + kernels/k/pwin_common.h kernels/common/wmma_backend.h + kernels/common/hip_device_minimal.h VERBATIM) +add_custom_target(d4r_k_gfx1201_compile ALL DEPENDS "${k_enc1_object}") +foreach(probe hip_gfx1201_probe gfx12_wmma_probe k_module_probe cuda_driver_probe d3d12_hip_interop_probe hip_stream_lifecycle_probe ngx_cuda_init_probe) add_executable(d4r_${probe} tools/windows/${probe}.cpp) target_compile_definitions(d4r_${probe} PRIVATE WIN32_LEAN_AND_MEAN NOMINMAX __HIP_PLATFORM_AMD__) target_include_directories(d4r_${probe} PRIVATE "${D4R_HIP_ROOT}/include") @@ -72,13 +93,23 @@ endif() install(TARGETS d4r_nvapi_compat RUNTIME DESTINATION nvapi-compat) target_link_libraries(d4r_d3d12_hip_interop_probe PRIVATE d3d12 dxgi dxguid) add_dependencies(d4r_hip_gfx1201_probe d4r_probe_code_object) +add_dependencies(d4r_gfx12_wmma_probe d4r_wmma_code_object) +add_dependencies(d4r_k_module_probe d4r_k_gfx1201_compile) install(FILES "${probe_object}" DESTINATION bin) +install(FILES "${wmma_object}" DESTINATION bin) +install(FILES "${k_enc1_object}" DESTINATION experimental/k) 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") +add_test(NAME gfx12_wmma COMMAND d4r_gfx12_wmma_probe + --hip-root "${D4R_HIP_ROOT}" --module "${wmma_object}" --iterations 16) +set_tests_properties(gfx12_wmma PROPERTIES TIMEOUT 90 LABELS "hardware;gfx1201;wmma") +add_test(NAME k_enc1_module COMMAND d4r_k_module_probe + --hip-root "${D4R_HIP_ROOT}" --module "${k_enc1_object}") +set_tests_properties(k_enc1_module PROPERTIES TIMEOUT 90 LABELS "hardware;gfx1201;k_compile") add_test(NAME d3d12_hip_interop COMMAND d4r_d3d12_hip_interop_probe --hip-root "${D4R_HIP_ROOT}" --module "${probe_object}" --iterations 32) set_tests_properties(d3d12_hip_interop PROPERTIES TIMEOUT 120 LABELS "hardware;interop") diff --git a/README.md b/README.md index b85055c..512e32f 100644 --- a/README.md +++ b/README.md @@ -7,7 +7,9 @@ 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/ZLUDA diagnostics, validated D3D12/HIP interop and an NGX init probe. +gfx1201 HIP/ZLUDA diagnostics, validated D3D12/HIP interop, an NGX init probe, +and a CPU-checked gfx12 WMMA backend. An experimental K `enc1` kernel compiles +and loads but has not passed transformer-output validation. 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. diff --git a/docs/windows-rdna4-port.md b/docs/windows-rdna4-port.md index 24ce0b6..3ef6c9f 100644 --- a/docs/windows-rdna4-port.md +++ b/docs/windows-rdna4-port.md @@ -6,8 +6,9 @@ This is a work log for an implementation in progress, not a claim of DLSS suppor ## Milestone and gates -Current milestone: M1/M2/M4 passed; NGX init + DLSS CUDA capability discovery -passed with the user's DLLs. Next: validated gfx12 WMMA and native K kernels. +Current milestone: M1/M2/M4/M5 passed; NGX init + DLSS CUDA capability discovery +passed with the user's DLLs. The first K `enc1` code object compiles and loads; +its prep/transformer have not yet been executed or checked against replay. Do not integrate NGX until the integer PTX workload is stable on the real GPU. | Gate | Status | @@ -17,7 +18,7 @@ Do not integrate NGX until the integer PTX workload is stable on the real GPU. | M2: Windows ZLUDA integer PTX, primary and created contexts | PASS, 32 iterations each + guard verification | | M3: Windows NGX initialization and simple DLSS path | Init + SR capabilities PASS; Evaluate/transformer pending | | M4: D3D12 / HIP external memory and fence round trip | PASS with TheRock; stable 7.2 has mapped-view leak | -| M5: independently validated gfx12 WMMA backend | Pending | +| M5: independently validated gfx12 WMMA backend | PASS: raw + legacy adapter, 16 cases, max abs/relative error 0 | | 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 | @@ -118,6 +119,27 @@ 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. +`kernels/common/wmma_backend.h` now preserves the gfx11 call and maps d4r's +16-half operand / even-odd accumulator contract to gfx12's 8-half operand / +contiguous accumulator contract. The gfx12 mapping uses an unconditional +wave32 half-exchange for every source fragment element before selection. +Performing the exchange only on lanes selected by a group-dependent branch +returned channel 9 in place of channel 1 on this card: the opposite half of +the wave was inactive at that instruction. This was caught by the CPU reference. +`tools/windows/wmma_gfx1201.hip` tests raw gfx12 WMMA, the adapter, the +half-exchange and lane ID with nonzero accumulator and one/two K16 steps. +Both stable HIP 7.2 and TheRock 10.2 produced exact 16x16 outputs across 16 +different inputs; max absolute and relative errors were both zero. + +`kernels/k/pwin_common.h` now calls the tested adapter. The Windows device-only +compiler uses `kernels/common/hip_device_minimal.h`, exposing only public +Clang AMDGPU work-item/grid/barrier builtins and HIP-compatible qualifiers; +the ordinary Linux HIP include path remains. CMake builds `enc1` for gfx1201 +into `experimental/k/` and verifies that HIP loads both prep and transformer +entry points. This is a **compile/load smoke test**, not K layer correctness or +DLSS execution. The code object must stay out of ZLUDA's native override path +until numerical K replay validation passes. + ## External interop blockers The installed HIP headers define D3D12Resource/D3D12Heap external memory and @@ -223,13 +245,13 @@ This builds into `build/windows-rdna4-therock` and installs into `dist/windows-rdna4-therock`. The stable profile stays available for reproducing the runtime bug. It is expected to return failure at the lifetime gate. -One command adds the local NGX initialization gate after the five hardware gates: +One command adds the local NGX initialization gate after the seven hardware gates: ```powershell powershell -NoProfile -File scripts/windows/test-windows-rdna4.ps1 -RuntimeProfile therock -NgxCore "C:/Users/Administrator/d4r/_nvngx.dll" -DlssDll "C:/Users/Administrator/d4r/nvngx_dlss.dll" ``` -Success requires all six tests and `PASS NGX_INIT ... sr_available=1`; +Success requires all eight tests and `PASS NGX_INIT ... sr_available=1`; `transformer_executed=0` is explicit. On NGX failure, CUDA trace and NVAPI query logs are collected automatically. NVIDIA DLLs remain local test inputs. @@ -295,6 +317,18 @@ MSVC parameter roundtrip Width=640, checked parameter/context/NGX cleanup. Local logs: `test-results/ngx-init-private-runtime.zip`. This is only the first part of M3; feature execution, K/M correctness and the Windows shim remain open. +2026-09-29 M5: TheRock full diagnostic run including user-provided NGX DLLs: +8/8 gates PASS (`test-results/m5-k-module.zip`). CTest: 7/7 PASS. Stable HIP 7.2: +5/5 non-interop gates PASS (HIP, gfx12 WMMA, K module load, both CUDA contexts); +stable mapped-view lifetime bug still excludes its interop fast path. +WMMA raw and legacy-layout adapter each match the scalar reference exactly +over 16 changing input matrices, with both one-step and two-step accumulation. +The `enc1` K module is built and loaded on gfx1201, with functions resolved; +`transformer_executed=0` is explicit in the module probe. Neither K nor M +transformer correctness is claimed yet. The same K source also compiles for +`gfx1101` through the unchanged gfx11 builtin branch (compile-only regression; +no RDNA3 GPU is present in this Windows machine). + 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. @@ -304,4 +338,4 @@ 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 these diagnostic milestones. +No end-to-end K/M/DLSS/OptiScaler support is claimed by these diagnostic milestones. diff --git a/kernels/common/hip_device_minimal.h b/kernels/common/hip_device_minimal.h new file mode 100644 index 0000000..650d371 --- /dev/null +++ b/kernels/common/hip_device_minimal.h @@ -0,0 +1,32 @@ +#pragma once + +// Device-only Windows HIP build shim. The diagnostic/kernel code objects are +// compiled with -nogpuinc/-nogpulib; no host HIP or C++ standard library is +// needed. Definitions below mirror public HIP qualifiers and LLVM AMDGPU +// work-item/grid builtins. The normal Linux HIP include path is untouched. +#define __device__ __attribute__((device)) +#define __global__ __attribute__((global)) +#define __shared__ __attribute__((shared)) +#define __forceinline__ __attribute__((always_inline)) inline +#define __launch_bounds__(n) __attribute__((amdgpu_flat_work_group_size(1, n))) + +struct d4r_device_dim3 { unsigned x, y, z; }; +__device__ __forceinline__ d4r_device_dim3 d4r_thread_idx() { + return {__builtin_amdgcn_workitem_id_x(), __builtin_amdgcn_workitem_id_y(), + __builtin_amdgcn_workitem_id_z()}; +} +__device__ __forceinline__ d4r_device_dim3 d4r_block_idx() { + return {__builtin_amdgcn_workgroup_id_x(), __builtin_amdgcn_workgroup_id_y(), + __builtin_amdgcn_workgroup_id_z()}; +} +__device__ __forceinline__ d4r_device_dim3 d4r_grid_dim() { + return {__builtin_amdgcn_grid_size_x() / __builtin_amdgcn_workgroup_size_x(), + __builtin_amdgcn_grid_size_y() / __builtin_amdgcn_workgroup_size_y(), + __builtin_amdgcn_grid_size_z() / __builtin_amdgcn_workgroup_size_z()}; +} +#define threadIdx d4r_thread_idx() +#define blockIdx d4r_block_idx() +#define gridDim d4r_grid_dim() +__device__ __attribute__((convergent, always_inline)) inline void __syncthreads() { + __builtin_amdgcn_s_barrier(); +} diff --git a/kernels/common/wmma_backend.h b/kernels/common/wmma_backend.h new file mode 100644 index 0000000..3165deb --- /dev/null +++ b/kernels/common/wmma_backend.h @@ -0,0 +1,57 @@ +#pragma once +#include + +// 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. +// gfx12 consumes only 8 FP16 operands per lane and emits 8 consecutive rows +// per lane group. Keep the old interface for K/M/texture call sites until their +// surrounding fragments are ported and validated individually. +using d4r_wmma_h16 = _Float16 __attribute__((ext_vector_type(16))); +using d4r_wmma_h8 = _Float16 __attribute__((ext_vector_type(8))); +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); + return __builtin_bit_cast(float, swapped); +} + +__attribute__((device, always_inline)) static inline d4r_wmma_f8 +d4r_wmma_legacy_layout(d4r_wmma_h16 a, d4r_wmma_h16 b, d4r_wmma_f8 c) +{ +#if defined(__gfx1200__) || defined(__gfx1201__) + const uint32_t group = __builtin_amdgcn_mbcnt_lo(~0u, 0u) >> 4; + d4r_wmma_h8 a12, b12; + d4r_wmma_f8 c12; + d4r_wmma_f8 remote_c; +#pragma unroll + for (int i = 0; i < 8; ++i) + remote_c[i] = d4r_wmma_other_half(c[i]); +#pragma unroll + for (int i = 0; i < 8; ++i) { + a12[i] = a[8u * group + i]; + b12[i] = b[8u * group + i]; + const uint32_t row = 8u * group + i; + const uint32_t old_index = row >> 1; + c12[i] = (row & 1u) == group ? c[old_index] : remote_c[old_index]; + } + const d4r_wmma_f8 d12 = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(a12, b12, c12); + d4r_wmma_f8 result, remote_d; +#pragma unroll + for (int i = 0; i < 8; ++i) + remote_d[i] = d4r_wmma_other_half(d12[i]); +#pragma unroll + for (int i = 0; i < 8; ++i) { + const uint32_t row = 2u * i + group; + const uint32_t new_index = row & 7u; + result[i] = (row >> 3) == group ? d12[new_index] : remote_d[new_index]; + } + return result; +#elif defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || defined(__gfx1103__) + return __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(a, b, c); +#else +#error "d4r WMMA backend requires a validated gfx11 or gfx12 target" +#endif +} diff --git a/kernels/k/pwin_common.h b/kernels/k/pwin_common.h index c6a39b2..14f9ded 100644 --- a/kernels/k/pwin_common.h +++ b/kernels/k/pwin_common.h @@ -6,7 +6,11 @@ // and the activations as the B operand (lane l: token l & 15, 16 K values). The f32 result then has // lane l = token l & 15 and VGPR i = output channel 2i + (l >> 4) of the 16-channel tile ("D^T"). #pragma once +#ifdef D4R_DEVICE_ONLY_MINIMAL +#include "../common/hip_device_minimal.h" +#else #include +#endif #include #pragma clang fp contract(off) @@ -17,6 +21,7 @@ typedef _Float16 h16 __attribute__((ext_vector_type(16))); typedef float f8v __attribute__((ext_vector_type(8))); typedef uint32_t u8v __attribute__((ext_vector_type(8))); typedef uint32_t u4v __attribute__((ext_vector_type(4))); +#include "../common/wmma_backend.h" __device__ __forceinline__ uint32_t lane_id() { @@ -25,7 +30,7 @@ __device__ __forceinline__ uint32_t lane_id() __device__ __forceinline__ f8v wmma(u8v a, u8v b, f8v c) { - return __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(__builtin_bit_cast(h16, a), __builtin_bit_cast(h16, b), c); + return d4r_wmma_legacy_layout(__builtin_bit_cast(h16, a), __builtin_bit_cast(h16, b), c); } // one k16 step with the f16 accumulator of NVIDIA's f16 wmma (rounded after the step) diff --git a/scripts/windows/test-windows-rdna4.ps1 b/scripts/windows/test-windows-rdna4.ps1 index 45e5891..dea2b89 100644 --- a/scripts/windows/test-windows-rdna4.ps1 +++ b/scripts/windows/test-windows-rdna4.ps1 @@ -94,7 +94,9 @@ try { $comgr = Join-Path $HipRoot 'bin/amd_comgr_3.dll' if (!(Test-Path -LiteralPath $comgr)) { $comgr = Join-Path $HipRoot 'bin/amd_comgr.dll' } $files = @((Join-Path $HipRoot 'bin/amdhip64_7.dll'), $comgr, - (Join-Path $ZludaRoot 'nvcuda.dll'), (Join-Path $PackageRoot 'bin/probe_gfx1201.hsaco')) + (Join-Path $ZludaRoot 'nvcuda.dll'), (Join-Path $PackageRoot 'bin/probe_gfx1201.hsaco'), + (Join-Path $PackageRoot 'bin/wmma_gfx1201.hsaco'), + (Join-Path $PackageRoot 'experimental/k/dltss_pwin_enc1_layer_gfx1201.hsaco')) if ($NgxCore -or $DlssDll) { if (!$NgxCore -or !$DlssDll) { throw 'Supply both -NgxCore and -DlssDll absolute paths.' } if (![IO.Path]::IsPathRooted($NgxCore) -or ![IO.Path]::IsPathRooted($DlssDll)) { @@ -117,6 +119,12 @@ try { $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) { + $wmmaOk = Invoke-Probe 'gfx12-wmma' (Join-Path $bin 'd4r_gfx12_wmma_probe.exe') @( + '--hip-root', $HipRoot, '--module', (Join-Path $bin 'wmma_gfx1201.hsaco'), '--iterations', "$Iterations") + if (!$wmmaOk) { throw 'gfx12 WMMA layout validation failed; refusing to mark K/M readiness.' } + $kModuleOk = Invoke-Probe 'k-enc1-module' (Join-Path $bin 'd4r_k_module_probe.exe') @( + '--hip-root', $HipRoot, '--module', (Join-Path $PackageRoot 'experimental/k/dltss_pwin_enc1_layer_gfx1201.hsaco')) + if (!$kModuleOk) { throw 'K enc1 module cannot load on gfx1201; transformer execution remains untested.' } $cudaOk = $true foreach ($context in @('primary', 'created')) { $ok = Invoke-Probe "cuda-$context" (Join-Path $bin 'd4r_cuda_driver_probe.exe') @( diff --git a/tools/windows/gfx12_wmma_probe.cpp b/tools/windows/gfx12_wmma_probe.cpp new file mode 100644 index 0000000..629ef52 --- /dev/null +++ b/tools/windows/gfx12_wmma_probe.cpp @@ -0,0 +1,127 @@ +#include "hip_api.h" +#include +#include +#include + +namespace { +uint16_t half_integer(int value) +{ + const unsigned n = static_cast(std::abs(value)); + if (n > 3u) throw std::runtime_error("test half integer out of range"); + const uint16_t bits[] = {0u, 0x3c00u, 0x4000u, 0x4200u}; + return static_cast(bits[n] | (value < 0 ? 0x8000u : 0u)); +} +struct Allocation { + d4r::diag::HipApi& hip; + void* pointer = nullptr; + Allocation(d4r::diag::HipApi& api, size_t bytes) : hip(api) { + hip.check(hip.hipMalloc(&pointer, bytes), "hipMalloc(WMMA)"); + } + ~Allocation() { if (pointer) (void)hip.hipFree(pointer); } +}; +} + +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 hip(args.hip_root); + hipDeviceProp_t properties{}; + hip.select_gfx1201(args.device, properties); + hipModule_t module = nullptr; + hip.check(hip.hipModuleLoad(&module, args.module.c_str()), "hipModuleLoad(WMMA)"); + struct ModuleCleanup { + HipApi& hip; hipModule_t module; + ~ModuleCleanup() { if (module) (void)hip.hipModuleUnload(module); } + } cleanup{hip, module}; + hipFunction_t raw_kernel = nullptr, adapter_kernel = nullptr, swap_kernel = nullptr; + hip.check(hip.hipModuleGetFunction(&raw_kernel, module, "d4r_wmma_gfx1201"), + "hipModuleGetFunction(WMMA)"); + hip.check(hip.hipModuleGetFunction(&adapter_kernel, module, "d4r_wmma_legacy_contract"), + "hipModuleGetFunction(WMMA adapter)"); + hip.check(hip.hipModuleGetFunction(&swap_kernel, module, "d4r_wmma_swap_test"), + "hipModuleGetFunction(WMMA swap)"); + constexpr size_t elements = 16u * 16u; + std::array a{}, b{}; + std::array initial{}, output{}; + Allocation da(hip, sizeof(a)), db(hip, sizeof(b)); + Allocation dc(hip, sizeof(initial)), dout(hip, sizeof(output)); + void* swap_out = dout.pointer; + void* swap_params[] = {&swap_out}; + hip.check(hip.hipModuleLaunchKernel(swap_kernel, 1, 1, 1, 32, 1, 1, 0, + nullptr, swap_params, nullptr), "hipModuleLaunchKernel(WMMA swap)"); + hip.check(hip.hipDeviceSynchronize(), "hipDeviceSynchronize(WMMA swap)"); + hip.check(hip.hipMemcpy(output.data(), dout.pointer, sizeof(output), hipMemcpyDeviceToHost), "hipMemcpy(WMMA swap)"); + for (unsigned lane = 0; lane < 32u; ++lane) { + if (output[lane] != static_cast(lane ^ 16u)) { + std::fprintf(stderr, "FAIL WMMA swap lane=%u expected=%u actual=%g\n", + lane, lane ^ 16u, output[lane]); + return 5; + } + if (output[32u + lane] != static_cast(lane)) { + std::fprintf(stderr, "FAIL WMMA lane-id lane=%u mbcnt=%g\n", lane, output[32u + lane]); + return 5; + } + } + float max_absolute_error = 0.0f; + float max_relative_error = 0.0f; + for (unsigned iteration = 0; iteration < args.iterations; ++iteration) { + const unsigned steps = 1u + (iteration & 1u); + for (unsigned row = 0; row < 16u; ++row) { + for (unsigned k = 0; k < 16u; ++k) { + a[row * 16u + k] = half_integer(static_cast((row * 3u + k * 5u + iteration) % 7u) - 3); + b[row * 16u + k] = half_integer(static_cast((row * 5u + k * 2u + iteration * 3u) % 7u) - 3); + initial[row * 16u + k] = static_cast(static_cast((row * 11u + k * 3u + iteration) % 17u) - 8) * 0.125f; + } + } + hip.check(hip.hipMemcpy(da.pointer, a.data(), sizeof(a), hipMemcpyHostToDevice), "hipMemcpy(A)"); + hip.check(hip.hipMemcpy(db.pointer, b.data(), sizeof(b), hipMemcpyHostToDevice), "hipMemcpy(B)"); + hip.check(hip.hipMemcpy(dc.pointer, initial.data(), sizeof(initial), hipMemcpyHostToDevice), "hipMemcpy(C)"); + void* a_arg = da.pointer; + void* b_arg = db.pointer; + void* c_arg = dc.pointer; + void* output_arg = dout.pointer; + unsigned steps_arg = steps; + void* params[] = {&a_arg, &b_arg, &c_arg, &output_arg, &steps_arg}; + for (unsigned mode = 0; mode < 2u; ++mode) { + output.fill(std::numeric_limits::quiet_NaN()); + hip.check(hip.hipMemcpy(dout.pointer, output.data(), sizeof(output), hipMemcpyHostToDevice), "hipMemcpy(output sentinel)"); + hip.check(hip.hipModuleLaunchKernel(mode ? adapter_kernel : raw_kernel, 1, 1, 1, 32, 1, 1, 0, nullptr, + params, nullptr), mode ? "hipModuleLaunchKernel(WMMA adapter)" : "hipModuleLaunchKernel(WMMA raw)"); + hip.check(hip.hipDeviceSynchronize(), "hipDeviceSynchronize(WMMA)"); + hip.check(hip.hipMemcpy(output.data(), dout.pointer, sizeof(output), hipMemcpyDeviceToHost), "hipMemcpy(output)"); + for (unsigned row = 0; row < 16u; ++row) { + for (unsigned col = 0; col < 16u; ++col) { + int dot = 0; + for (unsigned k = 0; k < 16u; ++k) { + const int av = static_cast((row * 3u + k * 5u + iteration) % 7u) - 3; + const int bv = static_cast((col * 5u + k * 2u + iteration * 3u) % 7u) - 3; + dot += av * bv; + } + const float expected = initial[row * 16u + col] + static_cast(static_cast(steps) * dot); + const float got = output[row * 16u + col]; + const float absolute = std::fabs(got - expected); + const float relative = absolute / std::max(1.0f, std::fabs(expected)); + if (!std::isfinite(got) || absolute > 0.0001f) { + std::fprintf(stderr, "FAIL WMMA mode=%s iteration=%u row=%u col=%u steps=%u expected=%.9g actual=%.9g abs=%.9g\n", + mode ? "adapter" : "raw", iteration, row, col, steps, expected, got, absolute); + return 5; + } + max_absolute_error = std::max(max_absolute_error, absolute); + max_relative_error = std::max(max_relative_error, relative); + } + } + } + } + std::printf("PASS GFX12_WMMA architecture=gfx1201 modes=raw,legacy_adapter iterations=%u max_abs=%.9g max_rel=%.9g\n", + args.iterations, max_absolute_error, max_relative_error); + return 0; + } catch (const std::exception& error) { + std::fprintf(stderr, "FAIL GFX12_WMMA %s\n", error.what()); + loaded_modules(); + return 4; + } +} diff --git a/tools/windows/k_module_probe.cpp b/tools/windows/k_module_probe.cpp new file mode 100644 index 0000000..76bb41a --- /dev/null +++ b/tools/windows/k_module_probe.cpp @@ -0,0 +1,31 @@ +#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 hip(args.hip_root); + hipDeviceProp_t properties{}; + hip.select_gfx1201(args.device, properties); + hipModule_t module = nullptr; + hip.check(hip.hipModuleLoad(&module, args.module.c_str()), "hipModuleLoad(K enc1)"); + struct Cleanup { + HipApi& hip; hipModule_t module; + ~Cleanup() { if (module) (void)hip.hipModuleUnload(module); } + } cleanup{hip, module}; + for (const char* name : {"dltss_pwin_enc1_layer_prep", "dltss_pwin_enc1_layer"}) { + hipFunction_t function = nullptr; + hip.check(hip.hipModuleGetFunction(&function, module, name), name); + if (!function) throw std::runtime_error(std::string("Null K function: ") + name); + } + std::printf("PASS K_MODULE architecture=gfx1201 functions=prep,transformer transformer_executed=0\n"); + return 0; + } catch (const std::exception& error) { + std::fprintf(stderr, "FAIL K_MODULE %s\n", error.what()); + loaded_modules(); + return 4; + } +} diff --git a/tools/windows/wmma_gfx1201.hip b/tools/windows/wmma_gfx1201.hip new file mode 100644 index 0000000..0c40720 --- /dev/null +++ b/tools/windows/wmma_gfx1201.hip @@ -0,0 +1,67 @@ +// Independent gfx12 WMMA layout probe. This deliberately does not use d4r fragments. +// A and B are row-major [logical row/column][K]; the intrinsic receives B,A +// so its output is row=A, column=B, as in AMD's RDNA 4 WMMA guide. +using h8 = _Float16 __attribute__((ext_vector_type(8))); +using f8 = float __attribute__((ext_vector_type(8))); +#include "../../kernels/common/wmma_backend.h" + +extern "C" __attribute__((global)) void d4r_wmma_swap_test(float* output) +{ + const unsigned int lane = __builtin_amdgcn_workitem_id_x(); + if (lane < 32u) { + output[lane] = d4r_wmma_other_half(static_cast(lane)); + output[32u + lane] = static_cast(__builtin_amdgcn_mbcnt_lo(~0u, 0u)); + } +} + +extern "C" __attribute__((global)) void d4r_wmma_gfx1201( + const unsigned short* a_bits, const unsigned short* b_bits, + const float* initial_c, float* output, unsigned int steps) +{ + const unsigned int lane = __builtin_amdgcn_workitem_id_x(); + if (lane >= 32u) return; + const unsigned int row = lane & 15u; + const unsigned int group = lane >> 4; + const _Float16* a = reinterpret_cast(a_bits); + const _Float16* b = reinterpret_cast(b_bits); + h8 a_frag, b_frag; + f8 c_frag; + for (unsigned int i = 0; i < 8u; ++i) { + a_frag[i] = a[row * 16u + group * 8u + i]; + b_frag[i] = b[row * 16u + group * 8u + i]; + c_frag[i] = initial_c[row * 16u + group * 8u + i]; + } + f8 result = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12( + b_frag, a_frag, c_frag); + if (steps == 2u) + result = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12( + b_frag, a_frag, result); + for (unsigned int i = 0; i < 8u; ++i) + output[row * 16u + group * 8u + i] = result[i]; +} + +// d4r's original gfx11 fragment contract, exercised independently of K/M. +extern "C" __attribute__((global)) void d4r_wmma_legacy_contract( + const unsigned short* a_bits, const unsigned short* b_bits, + const float* initial_c, float* output, unsigned int steps) +{ + const unsigned int lane = __builtin_amdgcn_workitem_id_x(); + if (lane >= 32u) return; + const unsigned int token = lane & 15u; + const unsigned int group = lane >> 4; + const _Float16* a = reinterpret_cast(a_bits); + const _Float16* b = reinterpret_cast(b_bits); + d4r_wmma_h16 a_frag, b_frag; + d4r_wmma_f8 c_frag; + for (unsigned int i = 0; i < 16u; ++i) { + a_frag[i] = a[token * 16u + i]; + b_frag[i] = b[token * 16u + i]; + } + for (unsigned int i = 0; i < 8u; ++i) + c_frag[i] = initial_c[(2u * i + group) * 16u + token]; + d4r_wmma_f8 result = d4r_wmma_legacy_layout(a_frag, b_frag, c_frag); + if (steps == 2u) + result = d4r_wmma_legacy_layout(a_frag, b_frag, result); + for (unsigned int i = 0; i < 8u; ++i) + output[(2u * i + group) * 16u + token] = result[i]; +}