Enable experimental native Windows RDNA3 and RDNA4 targets

Share explicit target metadata across build/runtime/package guards, preserve native gfx11 WMMA with a Windows LDS lane exchange, select HIP devices by D3D12 LUID, and reject wrong-target/vkd3d installations before changing game files. Review all issue #10 logs and document compile-only coverage plus RX9070XT regressions.

Co-authored-by: realdody <[email protected]>

kernels/common/wave32_exchange.h

@@ -0,0 +1,31 @@
+#pragma once
+#include <stdint.h>
+
+// 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
+}

kernels/common/wmma_backend.h

@@ -1,5 +1,6 @@
 #pragma once
 #include <stdint.h>
+#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"

kernels/common/wmma_layout.h

@@ -18,6 +18,7 @@
 // Only clang builtins are used (the texture tails compile with -nogpuinc -nogpulib).
 #pragma once
 #include <stdint.h>
+#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)
 {

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)