Clamp the patch merge to ±2π in the native Swin encoders

NVIDIA's enc1 and enc2 clamp the patch merge's result to ±2π (as f16) before
its e4m3 encoding (max.f16x2 / min.f16x2 with 0fC0C90FDB / 0f40C90FDB after the
merge's last MMA). The native layers did not, in either template or set:
merge_clamp (swin_common.h) now does, at the three sites that encode the merged
tokens (swin_block.h's e4m3-activation and f16 paths, swin_tok.h), and
swin_check.py's model clamps the same way.

Measured on an RX 9070 XT against a run with no native kernels (preset L,
accuracy set, the flags of Ghost of Tsushima, 1280x720 -> 3840x2160, one
frame): native enc1 alone 34 dB and enc2 alone 31 dB from the translated
kernels before, byte-identical after; the whole native set 28 -> 80 dB (the
rest is dec1, 0.09% of the output values). Preset M's fast set against
translated M: 30.8 -> 35.1 dB after 30 frames. The gfx11 path has the same
change and was not run on hardware.

The repo's tests pass (36, 3 skipped).

docs/native-kernels.md

@@ -55,6 +55,8 @@ The release is built this way, with hashes from DLSS 310.7.0 and 310.9.1. The PT
 
 **Preset L texture variants.** L uses unfolded `rrlite_enc0_4x4_*`, `rrlite_dec0_4x4`, and `rrlite_post_3_*` kernels. The L build keeps their full neural arithmetic and replaces surface stores with the existing native format-conversion functions. M's folded enc0 tail and dec0 head cannot be used for these variants. All four input flag combinations and all eight post variants are built, along with the shared downsample kernels. `all` and `tex` include these variants; `l` builds the shared Swin layers plus only the L texture set. L remains experimental; these replacements do not establish RTX parity or game performance.
 
+**Patch-merge clamp.** NVIDIA's enc1 and enc2 clamp the patch merge's result to ±2π (as f16) before its e4m3 encoding; the native layers did not, in either template or set (`merge_clamp` in `swin_common.h` now does). Measured on an RX 9070 XT against a run with no native kernels (preset L, accuracy set, the flags of Ghost of Tsushima, 1280×720 → 3840×2160, one frame): native enc1 alone 34 dB and enc2 alone 31 dB from the translated kernels before, byte-identical after; the whole native set 28 dB before and 80 dB after (the rest is dec1, 0.09% of the output values). Preset M's fast set against translated M: 30.8 → 35.1 dB after 30 frames. This is very likely what the switch below worked around; it has not been re-evaluated in a game. The gfx11 path has the same one-line change and was not run on hardware.
+
 **Motion mitigation.** `[Kernels] NativeSwinEncoders = false` (`D4R_NATIVE_SWIN_ENCODERS=0`) bypasses only `rrlite_enc1_4x4` and `rrlite_enc2_4x4` in the native manifest. L/M then use the original translated PTX for those encoders, retaining native tube, decoder and texture kernels. Keep `PreferAccuracy = true` for the tested conservative path. It reduced the extra building-edge trails in a 64-frame moving Townfall capture at 853×480 → 2560×1440 on gfx1101; it does not recover detail absent from the low-resolution input or establish RTX parity. The default remains `true` for speed. Flat developer sets without manifests fall back entirely when this switch is off; add a manifest to keep selective acceleration.
 
 **L's enc0 and dec0 on RDNA4 native FP8 (`kernels/tex/enc0l_tail.hip`, `dec0l_block.hip`).** The unfolded kernels contain a whole Swin block at pixel resolution (32 channels, 2 heads, 4×4 windows) that M's folded variants do not have: enc0 is the feature code, a two-layer embedding (32 → 32 → 32 with a ReLU), the block, and the 2×2 patch merge to 64 channels; dec0 is the patch expand of the half-resolution tokens, the skip sum, the block, and output code that writes four surfaces. On the `-fp8` build `make_ptx.py` keeps the texture code (enc0's feature staging, dec0's output code) as ZLUDA's compile and replaces the rest with the token-lane block of `swin_tok.h` (`swin_tok_core`), one window after the other in the kernel's single wave per 8×8-pixel block.

kernels/m/swin_block.h

@@ -975,7 +975,8 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
                 continue;
             const int ch = ginv(16 * (wv * NTP + j) + l16);
 #ifdef SWIN_A8
-            const uint32_t code[2] = {codes4(y[j].pair(0), y[j].pair(1)), codes4(y[j].pair(2), y[j].pair(3))};
+            const uint32_t code[2] = {codes4(merge_clamp(y[j].pair(0)), merge_clamp(y[j].pair(1))),
+                                      codes4(merge_clamp(y[j].pair(2)), merge_clamp(y[j].pair(3)))};
 #endif
 #pragma unroll
             for (int i = 0; i < 8; ++i)
@@ -988,7 +989,7 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
 #ifdef SWIN_A8
                     (uint8_t)(code[i >> 2] >> (8 * (i & 3)));
 #else
-                    (uint8_t)enc8(y[j].get(i));
+                    (uint8_t)enc8(merge_clamp(y[j].get(i)));
 #endif
             }
         }

kernels/m/swin_common.h

@@ -274,6 +274,18 @@ __device__ __forceinline__ half_t recip16(half_t s)
 #endif
 }
 
+// NVIDIA's encoders clamp the patch merge's result to +-2 pi (as f16) before its e4m3 encoding
+// (max.f16x2 / min.f16x2 with 0fC0C90FDB / 0f40C90FDB after the merge's last MMA).
+__device__ __forceinline__ hv2 merge_clamp(hv2 v)
+{
+    const hv2 lim = {(half_t)6.2831855f, (half_t)6.2831855f};
+    return __builtin_elementwise_minimumnum(__builtin_elementwise_maximumnum(v, -lim), lim);
+}
+__device__ __forceinline__ half_t merge_clamp(half_t v)
+{
+    return merge_clamp((hv2){v, v})[0];
+}
+
 // natural channel c -> pair-order column, and inverse (per 32-channel group)
 __device__ __forceinline__ int gperm(int c)
 {

kernels/m/swin_tok.h

@@ -738,7 +738,8 @@ __device__ __forceinline__ void swin_tok_block(const CommonParams& p, const Tube
 #pragma unroll
                     for (int kc = 0; kc < C / 8; ++kc)
                         k32(y, wt(K::PM, NTP, kc, 0, u), as[kc][0], wt(K::PM, NTP, kc, 1, u), as[kc][1]);
-                    const u2v code = tok_codes(y);
+                    const u2v code = {codes4(merge_clamp(y.pair(0)), merge_clamp(y.pair(1))),
+                                      codes4(merge_clamp(y.pair(2)), merge_clamp(y.pair(3)))};
                     if (inside)
                         *(u2v*)(p.p48 + ((size_t)((u >> 1) * H2 + MY) * W2 + MX) * 32 + 16 * (u & 1) + 8 * hf) = code;
                 }

kernels/tools/swin_check.py

@@ -47,7 +47,9 @@ def model_block(P, layer, bx, by):
             Wg = me.E4M3[P.wbuf[sm.woff_table(PM0 + g * 128 * C, 512, 64 * C, 4 * C, 32)]]
             bias = P.f16vec(PM0 + 4 * C * NPM + 64 * g, 32)
             cols.append(me.mma_chain(A, Wg, C0=np.broadcast_to(bias, (16, 32)).copy()))
-        merged = me.q8(sm.pair_to_nat(np.concatenate(cols, axis=1)))
+        # the merged tokens are clamped to +-2 pi (f16) before the e4m3 encoding
+        lim = float(np.float16(6.2831855))
+        merged = me.q8(np.clip(sm.pair_to_nat(np.concatenate(cols, axis=1)), -lim, lim))
     return out, merged