From 934fb486c771cf459dc7dd4e9dbf27faab259f46 Mon Sep 17 00:00:00 2001 From: woutermans Date: Thu, 8 Oct 2026 02:21:41 +0200 Subject: [PATCH] =?UTF-8?q?Clamp=20the=20patch=20merge=20to=20=C2=B12?= =?UTF-8?q?=CF=80=20in=20the=20native=20Swin=20encoders?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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 | 2 ++ kernels/m/swin_block.h | 5 +++-- kernels/m/swin_common.h | 12 ++++++++++++ kernels/m/swin_tok.h | 3 ++- kernels/tools/swin_check.py | 4 +++- 5 files changed, 22 insertions(+), 4 deletions(-) diff --git a/docs/native-kernels.md b/docs/native-kernels.md index 56bfe12..2c39b93 100644 --- a/docs/native-kernels.md +++ b/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. diff --git a/kernels/m/swin_block.h b/kernels/m/swin_block.h index 41f471d..b953b0c 100644 --- a/kernels/m/swin_block.h +++ b/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 } } diff --git a/kernels/m/swin_common.h b/kernels/m/swin_common.h index 698d764..286ad1b 100644 --- a/kernels/m/swin_common.h +++ b/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) { diff --git a/kernels/m/swin_tok.h b/kernels/m/swin_tok.h index ace809f..514f8aa 100644 --- a/kernels/m/swin_tok.h +++ b/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; } diff --git a/kernels/tools/swin_check.py b/kernels/tools/swin_check.py index 2d45bde..f1e1c63 100644 --- a/kernels/tools/swin_check.py +++ b/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