Narrow the RDNA4 accuracy set's cost for preset M

The FP8 Swin layers encode with a clamp to +-448 (NaN-propagating packed
minimum/maximum) followed by v_cvt_pk_fp8_f32, in both the fast and the
accuracy build: equal to half_to_e4m3 for every non-NaN input, NaN encodes
as 0xff. The accuracy build rounds each k32 step with v_fma_mix (K32_MIX).
On gfx12 the accuracy builds of M's post and downsample kernels are wave64
with denormal handling kept.

RX 9070 XT, 1280x720 -> 3840x2160, D3D12 harness: the accuracy set goes
from 3.54 ms to 3.05 ms per evaluation (3.01 -> 2.59 ms under load). Its
Swin layers still equal swin_model.py on sampled blocks and its image is
byte-identical to the previous accuracy set over 8 temporal frames.

docs/native-kernels.md

@@ -65,9 +65,9 @@ Initial validation: a gfx1101 fast L build completed two Ultra Performance harne
 
 The gfx12 M variant can use native e4m3 FP8 WMMA with `D4R_NATIVE_FP8=1`. Its activations are requantized to e4m3 at operand load, and its prepared weights remain bytes. The matching ZLUDA and texture-tail variant uses `D4R_ZLUDA_WMMA_FP8_NATIVE=1` / `D4R_TEX_FP8=1`; the release bridge chooses the `-fp8` kernel folder. `NativeFp8` in d4r.ini controls that choice and defaults to on. The f16 widening path remains available.
 
-In that build the GEMM activations live in LDS as e4m3 bytes (`SWIN_A8` in `swin_common.h`). Every value is encoded once where it is produced, four at a time with `v_cvt_pk_fp8_f32`, and an operand load is one 8-byte LDS read; the earlier build kept f16 values and encoded them again at every operand load (up to 16 times per value in the MLP). The instruction rounds like NVIDIA's conversion but encodes overflow as `0x7f`: the accuracy set clamps and maps NaNs first (`codes2`, equal to `half_to_e4m3` for all 65536 f16 inputs on an RX 9070 XT), the fast set fixes the overflow code afterwards (`codes4`; a NaN encodes as `0xfe`). The patch merge reads the stored output codes from a byte image. `-DSWIN_NO_A8` restores the previous operand path. Builds without `D4R_FP8_WMMA` (RDNA3, and RDNA4 with `NativeFp8` off) compile to the same code as before.
+In that build the GEMM activations live in LDS as e4m3 bytes (`SWIN_A8` in `swin_common.h`). Every value is encoded once where it is produced, four at a time with `v_cvt_pk_fp8_f32`, and an operand load is one 8-byte LDS read; the earlier build kept f16 values and encoded them again at every operand load (up to 16 times per value in the MLP). The instruction rounds like NVIDIA's conversion but encodes overflow as `0x7f`, so values are clamped to ±448 first with the NaN-propagating packed minimum and maximum (`codes4`). It equals `half_to_e4m3` for every non-NaN f16 input on an RX 9070 XT; a NaN encodes as the NaN code `0xff`, where the reference keeps its sign (`0x7f` or `0xff`). The accuracy build of these layers rounds each k32 step with `v_fma_mix` (`K32_MIX`; `-DK32_NO_MIX` restores the convert-and-round form) and its stored codes still equal `swin_model.py`. The patch merge reads the stored output codes from a byte image. `-DSWIN_NO_A8` restores the previous operand path. Builds without `D4R_FP8_WMMA` (RDNA3, and RDNA4 with `NativeFp8` off) compile to the same code as before.
 
-RX 9070 XT (gfx1201), 1280×720 → 3840×2160, D3D12 harness, DLSS 310.7: the five Swin layers went from 3.99 ms to 1.37 ms per frame in replay (enc1 0.97 → 0.28, enc2 0.42 → 0.17, enc3 6 × 0.22 → 6 × 0.08, dec2 0.41 → 0.15, dec1 0.89 → 0.31), and the whole evaluation from 5.9 ms to 2.8 ms of GPU time (2.4 ms when the GPU is not left idle between frames) ([performance.md](performance.md#rdna4)). The layer outputs are byte-identical to the previous build on recorded launches, and the harness output is byte-identical over a 12-frame temporal sequence. The accuracy build's stored codes equal `swin_model.py` on sampled blocks of all five layers.
+RX 9070 XT (gfx1201), 1280×720 → 3840×2160, D3D12 harness, DLSS 310.7: the five Swin layers went from 3.99 ms to 1.37 ms per frame in replay (enc1 0.97 → 0.28, enc2 0.42 → 0.17, enc3 6 × 0.22 → 6 × 0.08, dec2 0.41 → 0.15, dec1 0.89 → 0.31), and the whole evaluation from 5.9 ms to 2.8 ms of GPU time (2.4 ms when the GPU is not left idle between frames) ([performance.md](performance.md#rdna4)). The layer outputs are byte-identical to the previous build on recorded launches, and the harness output is byte-identical over a 12-frame temporal sequence. The accuracy build's stored codes equal `swin_model.py` on sampled blocks of all five layers. The complete accuracy set costs 3.05 ms per evaluation in the same harness (2.59 ms under load) against 2.77 ms (2.37 ms) for the fast set; before the encoder, rounding and wave64 changes it cost 3.54 ms (3.01 ms). On gfx12 the accuracy builds of M's post and downsample kernels are wave64 with denormal handling kept; the image was identical to the wave32 accuracy set over an 8-frame temporal sequence.
 
 The gfx11 refactor and gfx12-layout shim matched the recorded gfx11 replays and 88-frame K/M harness images byte for byte on the RX 7700 XT. rocjitsu gfx1201 replay comparisons put the native K layers within 96–109 dB PSNR of gfx11, and M differed in 0.001–0.3% of e4m3 bytes. These are emulator results, not RDNA4 hardware results. rocjitsu's FP8 arithmetic is not hardware-checked; bit-exactness, speed and driver stability on an RDNA4 GPU remain unverified.
 

kernels/build.sh

@@ -18,6 +18,7 @@
 #                   which compiles PTX offline for D4R_GPU_ARCH (no GPU of that kind needed)
 #   D4R_DLSS_PTX_DIR optional pre-extracted PTX directory for parallel per-target builds
 set -euo pipefail
+shopt -s extglob
 
 HERE="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
 WHAT="${1:-all}"
@@ -131,8 +132,15 @@ if [[ "$WHAT" == all || "$WHAT" == tex || "$WHAT" == l ]]; then
     for spec in "${specs[@]}"; do
         IFS=: read -r kernel src mode <<< "$spec"
         [[ "$ACCURACY" == 0 && "$kernel" == "$native_out" ]] && continue
-        D4R_PREFER_ACCURACY="$ACCURACY" \
+        # accuracy sets are wave32, except M's post and downsample kernels on RDNA4: their wave64 builds keep the
-            D4R_ZLUDA_WAVE64="$([[ "$mode" == w64 && "$ACCURACY" == 0 ]] && echo 1 || echo 0)" D4R_ROCM_DIR="$ROCM" D4R_GPU_ARCH="$ARCH" \
+        # denormal handling and gave the same image on an RX 9070 XT (post 0.79 -> 0.67 ms, downsample 0.43 -> 0.39 ms)
+        w64=0
+        if [[ "$mode" == w64 ]]; then
+            [[ "$ACCURACY" == 0 ]] && w64=1
+            [[ "$ARCH" == gfx12* && "$kernel" == rrlite_@(post|downsample)_* ]] && w64=1
+        fi
+        D4R_PREFER_ACCURACY="$ACCURACY" D4R_TEX_ACCURACY_WAVE64="$w64" \
+            D4R_ZLUDA_WAVE64="$w64" D4R_ROCM_DIR="$ROCM" D4R_GPU_ARCH="$ARCH" \
             D4R_TEX_FP8="$FP8" D4R_DLSS_PTX_DIR="$PTX_DIR" "$HERE/tex/build_tex.sh" "$kernel" "$src" "$OUT"
     done
     if [[ "$WHAT" != l && "$ACCURACY" == 0 ]]; then

kernels/m/swin_common.h

@@ -34,6 +34,12 @@
 #if defined(D4R_FP8_WMMA) && defined(__GFX12__) && !defined(SWIN_NO_A8)
 #define SWIN_A8 1
 #endif
+// The exact SWIN_A8 build rounds each k32 step with v_fma_mix (f16(P + C) in one rounding) instead of
+// converting the accumulator to f32 and back: the same codes as the conversion form on the recorded launches
+// and as tools/swin_model.py, at about half the cost. K32_NO_MIX keeps the conversion form.
+#if defined(SWIN_A8) && defined(SWIN_EXACT) && !defined(K32_NO_MIX) && !defined(K32_MIX)
+#define K32_MIX
+#endif
 
 #pragma clang fp contract(off)
 
@@ -207,31 +213,17 @@ __device__ __forceinline__ uint32_t codes8x2(hv2 q)
 }
 
 #ifdef SWIN_A8
-// e4m3 codes of two halves (RNE, satfinite, NaN -> 0x7f | sign: equal to half_to_e4m3 for all inputs) in bytes 0
+// e4m3 codes of four halves (RNE, satfinite) in the bytes of one word. The instruction rounds like NVIDIA's
-// and 1. The instruction rounds like NVIDIA's conversion but encodes overflow as 0x7f, so magnitudes are
+// conversion but encodes overflow as 0x7f, so the values are clamped to +-448 first, with the NaN-propagating
-// clamped to 448 first and NaNs enter as infinities.
+// minimum / maximum. Equal to half_to_e4m3 for every non-NaN input; a NaN encodes as the NaN code 0xff (the
-__device__ __forceinline__ uint32_t codes2(hv2 h)
+// reference keeps its sign: 0x7f or 0xff).
-{
-    const u16x2 u = __builtin_bit_cast(u16x2, h);
-    const u16x2 m = u & (unsigned short)0x7fff;
-    const u16x2 nan = __builtin_bit_cast(u16x2, (i16x2)(__builtin_bit_cast(i16x2, (u16x2)((unsigned short)0x7c00 - m)) >> 15));
-    const u16x2 k = __builtin_elementwise_max(__builtin_elementwise_min(m, (u16x2){0x5f00, 0x5f00}), (u16x2)(nan & (unsigned short)0x7c00));
-    const hv2 c = __builtin_bit_cast(hv2, (u16x2)(k | (u & (unsigned short)0x8000)));
-    return (uint32_t)__builtin_amdgcn_cvt_pk_fp8_f32((float)c[0], (float)c[1], 0, false) & 0xffffu;
-}
-// e4m3 codes of four halves in the bytes of one word. Exact builds use codes2. The fast build converts
-// first and then turns the instruction's overflow code into the saturated one in all four bytes at once
-// (low seven bits 0x7f -> 0x7e); it equals half_to_e4m3 for every non-NaN input, and a NaN encodes as the
-// negative saturated value 0xfe instead of 0x7f | sign.
 __device__ __forceinline__ uint32_t codes4(hv2 a, hv2 b)
 {
-#ifdef SWIN_EXACT
+    const hv2 lim = {(half_t)448.0f, (half_t)448.0f};
-    return codes2(a) | (codes2(b) << 16);
+    a = __builtin_elementwise_maximum(__builtin_elementwise_minimum(a, lim), -lim);
-#else
+    b = __builtin_elementwise_maximum(__builtin_elementwise_minimum(b, lim), -lim);
-    uint32_t w = (uint32_t)__builtin_amdgcn_cvt_pk_fp8_f32((float)a[0], (float)a[1], 0, false);
+    const uint32_t w = (uint32_t)__builtin_amdgcn_cvt_pk_fp8_f32((float)a[0], (float)a[1], 0, false);
-    w = (uint32_t)__builtin_amdgcn_cvt_pk_fp8_f32((float)b[0], (float)b[1], (int)w, true);
+    return (uint32_t)__builtin_amdgcn_cvt_pk_fp8_f32((float)b[0], (float)b[1], (int)w, true);
-    return w - ((((w & 0x7f7f7f7fu) + 0x01010101u) >> 7) & 0x01010101u);
-#endif
 }
 #endif
 

kernels/tex/build_tex.sh

@@ -3,7 +3,8 @@
 # part of it calls functions from SRC.hip, is compiled by ZLUDA with SRC's bitcode linked in
 # (D4R_ZLUDA_EXTRA_BC) and the resulting code object is saved as OUT_DIR/KERNEL.hsaco.
 # Called by kernels/build.sh, which sets D4R_ROCM_DIR, D4R_GPU_ARCH, D4R_DLSS_PTX_DIR, D4R_ZLUDA_EMIT,
-# D4R_ZLUDA_WAVE64 (1 = wave64 code object) and D4R_TEX_FP8 (1 = the gfx12 native-FP8 variant).
+# D4R_ZLUDA_WAVE64 (1 = wave64 code object; accuracy builds only with D4R_TEX_ACCURACY_WAVE64=1) and
+# D4R_TEX_FP8 (1 = the gfx12 native-FP8 variant).
 # ZLUDA compiles offline for D4R_GPU_ARCH with its d4r_emit tool, so any target builds on any machine.
 set -euo pipefail
 K=$1 SRC=$2 OUT=$3
@@ -31,7 +32,7 @@ D4R_PREFER_ACCURACY="$ACCURACY" python3 "$D/make_ptx.py" "$K" "$B/$K.ptx" >/dev/
 env LD_LIBRARY_PATH="$R/lib" D4R_ZLUDA_EXTRA_BC="$B/extra.bc" D4R_ZLUDA_WMMA=1 D4R_ZLUDA_WMMA_FP8=1 \
     D4R_ZLUDA_WMMA_FP8_NATIVE="$FP8" D4R_ZLUDA_IGNORE_DENORMAL="$((1 - ACCURACY))" \
     D4R_ZLUDA_WMMA_F32ACC=0 D4R_ZLUDA_FAST_MATH=0 D4R_ZLUDA_IMPLICIT_MAX_BLOCK=256 \
-    D4R_ZLUDA_WAVE64="$([[ "$ACCURACY" == 1 ]] && echo 0 || echo "${D4R_ZLUDA_WAVE64:-0}")" \
+    D4R_ZLUDA_WAVE64="$([[ "$ACCURACY" == 1 && "${D4R_TEX_ACCURACY_WAVE64:-0}" != 1 ]] && echo 0 || echo "${D4R_ZLUDA_WAVE64:-0}")" \
     "$D4R_ZLUDA_EMIT" "$B/$K.ptx" "$B/dump" "$ARCH" >/dev/null
 cp "$B/dump/module.hsaco" "$OUT/$K.hsaco"
 printf '%-52s %s\n' "$K" "$(grep -E '^\s+\.(vgpr_count|private_segment_fixed_size|wavefront_size):' "$B/dump/asm.s" | tr -s ' ' | paste -sd' ')"