Execute native Windows K DLSS and validate real transformer weights

kernels/common/hip_device_minimal.h

@@ -35,5 +35,9 @@ __device__ __forceinline__ d4r_device_dim3 d4r_grid_dim() {
 #define blockIdx d4r_block_idx()
 #define gridDim d4r_grid_dim()
 __device__ __attribute__((convergent, always_inline)) inline void __syncthreads() {
+    // Match HIP __work_group_barrier: execution rendezvous alone does not
+    // publish LDS stores or prevent loads moving across the barrier.
+    __builtin_amdgcn_fence(__ATOMIC_RELEASE, "workgroup");
     __builtin_amdgcn_s_barrier();
+    __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup");
 }

kernels/k/pwin_common.h

@@ -30,12 +30,55 @@ __device__ __forceinline__ uint32_t lane_id()
 
 typedef wm_op op_t; // WMMA operand: u8v (gfx11) or u4v (gfx12)
 
+// The CUDA fragment reduction rounds products and every tree node to f16.
+// A float sum of the squared halves changes rsqrt and can amplify differences
+// through attention. The strict baseline preserves its channel grouping.
+template <int C>
+__device__ __forceinline__ half_t l2_sum(const half_t* row)
+{
+#ifdef D4R_K_FP16_BASELINE
+    static_assert(C % 32 == 0 && C <= 160, "audited K channel counts");
+    hv2 partial[4];
+#pragma unroll
+    for (int t = 0; t < 4; ++t)
+    {
+        hv2 halves[2];
+#pragma unroll
+        for (int parity = 0; parity < 2; ++parity)
+        {
+            hv2 pairs[C / 32];
+#pragma unroll
+            for (int group = 0; group < C / 32; ++group)
+            {
+                const hv2 a = *(const hv2*)(row + 2 * t + 8 * parity + 32 * group);
+                const hv2 b = *(const hv2*)(row + 2 * t + 8 * parity + 32 * group + 16);
+                pairs[group] = (a * a) + (b * b);
+            }
+            hv2 value = pairs[0];
+            if constexpr (C >= 64) value = value + pairs[1];
+            if constexpr (C == 96) value = value + pairs[2];
+            if constexpr (C >= 128) value = value + (pairs[2] + pairs[3]);
+            if constexpr (C == 160) value = value + pairs[4];
+            halves[parity] = value;
+        }
+        partial[t] = halves[0] + halves[1];
+    }
+    const hv2 sum = (partial[0] + partial[1]) + (partial[2] + partial[3]);
+    return sum[0] + sum[1];
+#else
+    float sum = 0;
+#pragma unroll
+    for (int c = 0; c < C; ++c) sum += (float)(half_t)(row[c] * row[c]);
+    return (half_t)sum;
+#endif
+}
+
 // one k16 step with the f16 accumulator of NVIDIA's f16 wmma (rounded after the step)
 // 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)
 {
     f8v d = wm_mma(a, b, c);
-#ifndef PWIN_F32ACC
+#if !defined(PWIN_F32ACC) || defined(D4R_K_FP16_BASELINE)
 #pragma unroll
     for (int i = 0; i < 8; ++i)
         d[i] = (float)(half_t)d[i];

kernels/k/pwin_layer.h

@@ -259,17 +259,7 @@ __device__ void pwin_core(const PwinParams& p, const u8v* __restrict__ img, cons
 
     STAMP(1);
     // ---- L2 norm: h1 = x0 * (rsqrt(sum x0^2) * gamma1) -> hb
-    float ss = 0.0f;
-    for (int kt = 0; kt < KT; ++kt)
-    {
-        const u8v xv = lds_row16(arow + 16 * kt);
-#pragma unroll
-        for (int c = 0; c < 16; ++c)
-        {
-            const half_t v = op_get(xv, c);
-            ss += (float)(half_t)(v * v);
-        }
-    }
+    const half_t ss = l2_sum<C>(arow);
     const half_t r16 = (half_t)__builtin_amdgcn_rsqf((float)(half_t)ss);
     for (int kt = (int)hf; kt < KT; kt += 2)
     {

kernels/k/pwin_pos.h

@@ -19,18 +19,7 @@ __device__ void pos_core(const PwinParams& p, const u8v* __restrict__ img, const
     {
         const half_t* arow = act[16 * (wv * MT + mi) + m];
         half_t* hrow = hb[16 * (wv * MT + mi) + m];
-        float ss = 0.0f;
-#pragma unroll
-        for (int kt = 0; kt < KT; ++kt)
-        {
-            const u8v xv = lds_row16(arow + 16 * kt);
-#pragma unroll
-            for (int c = 0; c < 16; ++c)
-            {
-                const half_t v = op_get(xv, c);
-                ss += (float)(half_t)(v * v);
-            }
-        }
+        const half_t ss = l2_sum<C>(arow);
         const half_t r16 = (half_t)__builtin_amdgcn_rsqf((float)(half_t)ss);
         for (int kt = (int)hf; kt < KT; kt += 2)
         {

kernels/k/pwin_wide.h

@@ -38,17 +38,7 @@ __device__ void wide_core(const PwinParams& p, const u8v* __restrict__ img, cons
     half_t* hrow = hb[tok];
 
     // ---- L2 norm: h1 = x0 * (rsqrt(sum x0^2) * gamma1) -> hb (all groups of a token tile compute the sum)
-    float ss = 0.0f;
-    for (int kt = 0; kt < KT; ++kt)
-    {
-        const u8v xv = lds_row16(arow + 16 * kt);
-#pragma unroll
-        for (int c = 0; c < 16; ++c)
-        {
-            const half_t v = op_get(xv, c);
-            ss += (float)(half_t)(v * v);
-        }
-    }
+    const half_t ss = l2_sum<C>(arow);
     const half_t r16 = (half_t)__builtin_amdgcn_rsqf((float)(half_t)ss);
     for (int kt = 2 * g + (int)hf; kt < KT; kt += 2 * NG)
     {

kernels/tools/pwin_model.py

@@ -6,6 +6,38 @@ import numpy as np
 
 F16, F32, F64 = np.float16, np.float32, np.float64
 
+# Legacy replay reports retain their original approximate norm by default.
+# Windows strict K validation selects the PTX fragment's rounded reduction.
+NORM_FRAGMENT_ORDER = False
+
+
+def norm_rows(x):
+    if not NORM_FRAGMENT_ORDER:
+        return f16(np.sum(x.astype(F32).astype(F64) ** 2, axis=1))
+    tokens, channels = x.shape
+    assert channels in (32, 64, 96, 128, 160)
+    square = f16(x.astype(F64) ** 2)
+    partial = []
+    for lane in range(4):
+        halves = []
+        for parity in range(2):
+            pairs = [add16(square[:, 2 * lane + 8 * parity + 32 * g:2 * lane + 8 * parity + 32 * g + 2],
+                           square[:, 2 * lane + 8 * parity + 32 * g + 16:2 * lane + 8 * parity + 32 * g + 18])
+                     for g in range(channels // 32)]
+            value = pairs[0]
+            if channels >= 64:
+                value = add16(value, pairs[1])
+            if channels == 96:
+                value = add16(value, pairs[2])
+            if channels >= 128:
+                value = add16(value, add16(pairs[2], pairs[3]))
+            if channels == 160:
+                value = add16(value, pairs[4])
+            halves.append(value)
+        partial.append(add16(halves[0], halves[1]))
+    total = add16(add16(partial[0], partial[1]), add16(partial[2], partial[3]))
+    return add16(total[:, 0], total[:, 1])
+
 
 class Dump:
     def __init__(self, d):
@@ -181,7 +213,7 @@ def pwin_enc2(P, bx, by, pm_order=((0, 0), (0, 1), (1, 0), (1, 1)), b2_in_c0=Tru
     """logical enc2 layer (H 2, C 64 -> merged 96) for one window; returns (y [64 x 64], merged [16 x 96])"""
     C, H, Dh = 64, 2, 32
     x0 = load_window(P, bx, by, C)
-    s = f16(np.sum(x0.astype(F32).astype(F64) ** 2, axis=1))
+    s = norm_rows(x0)
     r = rsqrt16(s)
     h1 = mul16(x0, mul16(r[:, None], vec(P, 0, C)[None, :]))
     Os = []
@@ -248,7 +280,7 @@ def swin_core(P, L, x0, dbg=None, posattn=False):
     """norm -> attention -> Wo + residual -> MLP (+ residual): the block output y [64 x C]"""
     C, H = L.C, L.H
     NCH = C // 32
-    s = f16(np.sum(x0.astype(F32).astype(F64) ** 2, axis=1))
+    s = norm_rows(x0)
     r = rsqrt16(s)
     h1 = mul16(x0, mul16(r[:, None], vec(P, L.G1, C)[None, :]))
     Os = []