Reproduce gfx12 FP16 accumulator differences and preserve K baseline

kernels/common/wmma_layout.h

@@ -229,6 +229,20 @@ WM_FN wm_f8v wm_mma(const wm_op& a, const wm_op& b, wm_f8v c)
 {
     return __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(__builtin_bit_cast(wm_h8, a), __builtin_bit_cast(wm_h8, b), c);
 }
+#ifdef D4R_K_F16_WMMA
+// Rejected K experiment: direct packed F16 changes the arithmetic on gfx1201.
+// The public WMMA reproducer and enc1 replay are in tools/windows and docs.
+// Experimental gfx12 packed FP16 C/D. K's strict baseline enters each K16
+// step with an already FP16-rounded accumulator. Keep the existing logical
+// eight-row contract; only this instruction/register representation changes.
+WM_FN wm_f8v wm_mma_f16_step(const wm_op& a, const wm_op& b, wm_f8v c)
+{
+    const wm_h8 hc = __builtin_convertvector(c, wm_h8);
+    const wm_h8 hd = __builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12(
+        __builtin_bit_cast(wm_h8, a), __builtin_bit_cast(wm_h8, b), hc);
+    return __builtin_convertvector(hd, wm_f8v);
+}
+#endif
 #else
 WM_FN wm_f8v wm_mma11(const wm_u8v& a, const wm_u8v& b, wm_f8v c)
 {

kernels/k/pwin_common.h

@@ -80,6 +80,12 @@ __device__ __forceinline__ half_t l2_sum(const half_t* row)
 // PWIN_F32ACC: keep the accumulator in f32 through the chain (rounded to f16 where the values are used)
 __device__ __forceinline__ f8v mma16(const op_t& a, const op_t& b, f8v c)
 {
+#ifdef D4R_K_F16_WMMA
+#if !defined(__GFX12__) || D4R_WMMA_LAYOUT != 12 || !defined(D4R_K_FP16_BASELINE)
+#error "experimental packed F16 WMMA requires gfx12 native layout and strict K baseline"
+#endif
+    return wm_mma_f16_step(a, b, c);
+#else
     f8v d = wm_mma(a, b, c);
 #if !defined(PWIN_F32ACC) || defined(D4R_K_FP16_BASELINE)
 #pragma unroll
@@ -87,6 +93,7 @@ __device__ __forceinline__ f8v mma16(const op_t& a, const op_t& b, f8v c)
         d[i] = (float)(half_t)d[i];
 #endif
     return d;
+#endif
 }
 
 __device__ __forceinline__ f8v splat8(float v)