From ba77946a4ad5b179c38e97dd2700ad9595a70508 Mon Sep 17 00:00:00 2001
From: woutermans <woutertje554@gmail.com>
Date: Tue, 6 Oct 2026 14:11:19 +0200
Subject: [PATCH 1/4] Speed up preset M on RDNA4: e4m3 byte activations in the
 FP8 Swin layers

On the gfx12 native-FP8 path the Swin layers kept GEMM activations as f16
in LDS and encoded them to e4m3 again at every operand load. They are now
encoded once where they are produced, with v_cvt_pk_fp8_f32, and stored as
bytes (SWIN_A8); the patch merge reads the stored output codes. The M
downsample kernels are built as wave64 on gfx12.

RX 9070 XT, 1280x720 -> 3840x2160, D3D12 harness, DLSS 310.7: whole
evaluation 5.9 ms -> 2.8 ms of GPU time with the GPU otherwise idle, 2.4 ms
at 80-98% GPU utilisation. Harness output is byte-identical to the previous
kernels; the accuracy build matches swin_model.py on sampled blocks.

Builds without D4R_FP8_WMMA (RDNA3, and RDNA4 with NativeFp8 off) compile to
the same instructions as before.
---
 docs/native-kernels.md  |   6 +-
 docs/performance.md     |  22 +++++-
 kernels/build.sh        |   5 +-
 kernels/m/swin_block.h  | 168 +++++++++++++++++++++++++++++++++++++---
 kernels/m/swin_common.h |  43 ++++++++++
 5 files changed, 232 insertions(+), 12 deletions(-)

diff --git a/docs/native-kernels.md b/docs/native-kernels.md
index e09bd98..f01bc99 100644
--- a/docs/native-kernels.md
+++ b/docs/native-kernels.md
@@ -65,6 +65,10 @@ 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.
+
+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 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.
 
 For ZLUDA-compiled code, compare rocjitsu runs with other rocjitsu runs: the emulator does not reproduce the denormal-flush mode of those code objects, so a GPU byte comparison is not a valid arithmetic check. Captures with texture or surface objects must not be sent to the replay tools; the standalone `tail_check.hip` and the full D3D12 harness cover those tails.
@@ -77,7 +81,7 @@ For ZLUDA-compiled code, compare rocjitsu runs with other rocjitsu runs: the emu
   - for M's enc0 and dec0, a range of the kernel is replaced by a native tail or head.
 - **Build:** ZLUDA's `d4r_emit` compiles the edited PTX for `D4R_GPU_ARCH` with the HIP bitcode linked in (`D4R_ZLUDA_EXTRA_BC`), and the resulting code object is saved. This is an offline per-target build.
 - **Every flag combination:** DLSS picks kernel variants by the game's settings (motion vectors at render or display resolution, HDR or LDR input, inverted or regular depth, ...). `kernels/build.sh` builds all of them: 24 K output kernels, 8 M post, 4 M downsample and 4 M enc0 kernels. enc0's native tail is spliced into each variant at its own line, with the registers read off the variant's PTX (`make_ptx.py` checks that the code after the cut matches the reference variant line for line).
-- **Wave64:** kernels without MMAs can be compiled as wave64 (`D4R_ZLUDA_WAVE64=1`, ZLUDA patch 0006): each wave runs two CUDA warps, and RDNA3 issues FP32 work for all 64 lanes at once. The post and K output kernels are built this way (post 0.89 → 0.77 ms), with bit-identical results.
+- **Wave64:** kernels without MMAs can be compiled as wave64 (`D4R_ZLUDA_WAVE64=1`, ZLUDA patch 0006): each wave runs two CUDA warps, and RDNA3 issues FP32 work for all 64 lanes at once. The post and K output kernels are built this way (post 0.89 → 0.77 ms), with bit-identical results. The M downsample kernels are wave64 on gfx12 only (RX 9070 XT: 0.37 → 0.34 ms at 3840×2160, identical output); on RDNA3 they were slower that way.
 - **Output redirect:** the kernel that writes a preset's result can write it straight into the shim's output buffer instead of a CUDA array: `hiluma_engine_output_*` for K, `rrlite_downsample_kernel_*` for M (M's post kernel works at 3840×2160; the downsample produces the output). See `D4R_SHIM_OUTPUT_DIRECT` in [architecture.md](architecture.md).
 
 ## Numerics
diff --git a/docs/performance.md b/docs/performance.md
index 8b861c1..2304b6d 100644
--- a/docs/performance.md
+++ b/docs/performance.md
@@ -1,6 +1,6 @@
 # Performance
 
-All frame rates and kernel timings on this page were measured on an RX 7700 XT (gfx1101). RDNA4 (gfx1200/gfx1201) has compile and emulator validation only; no RDNA4 frame rate, latency, or kernel timing has been measured. Native FP8 may change both arithmetic and cost, so the gfx11 results cannot be projected to it.
+Unless stated otherwise, the frame rates and kernel timings on this page were measured on an RX 7700 XT (gfx1101). Native FP8 changes both arithmetic and cost, so the gfx11 results cannot be projected to RDNA4; the one RDNA4 measurement is in [RDNA4](#rdna4). No RDNA4 game frame rate has been measured.
 
 ## Method
 
@@ -75,3 +75,23 @@ For DLSS 4 at Quality the GPU spends about 2.8 ms per frame in DLSS kernels:
 - The input kernel, exposure and miscellaneous kernels: about 0.35 ms.
 
 The game-side input and output copies add about 0.3 ms. [native-kernels.md](native-kernels.md) lists every kernel.
+
+## RDNA4
+
+DLSS 4.5 (M) on an RX 9070 XT (gfx1201, native FP8, fast set), 1280×720 → 3840×2160 in the D3D12 harness with DLSS 310.7. Median GPU time per frame in milliseconds, from `D4R_CUDA_KERNEL_PROFILE` (frames 21–40); "before" is the same texture set with the previous Swin layers.
+
+| Kernel | before | now |
+|---|---|---|
+| enc3 tube (6 launches) | 1.50 | 0.52 |
+| enc1 | 1.32 | 0.33 |
+| dec1 | 0.94 | 0.33 |
+| post (texture kernel) | 0.67 | 0.68 |
+| enc2 | 0.51 | 0.19 |
+| dec2 | 0.50 | 0.17 |
+| downsample (texture kernel, now wave64) | 0.36 | 0.34 |
+| dec0, enc0 (texture kernels) | 0.18, 0.13 | 0.18, 0.13 |
+| **total** | **6.10** | **2.86** |
+
+The shim's whole-evaluation GPU time (`D4R_PROFILE`, no per-kernel synchronisation) went from 5.9–6.0 ms to 2.80 ms with the harness's 200 ms pause between frames, and to 2.39–2.43 ms with frames 0–4 ms apart (2000 frames at 4 ms spacing: median 2.39, maximum 2.56). The difference is the driver's clock governor: with the 200 ms pause the GPU is about 9% busy and runs at about 2790 MHz, with frames 0–4 ms apart it is 80–98% busy and runs at 3180–3300 MHz, as it does under a game's rendering load. Other GPU users on the desktop add occasional slower frames to any of these figures; the previous kernels show the same disturbances.
+
+The change is in [native-kernels.md](native-kernels.md#rdna4): the FP8 layers keep their activations as e4m3 bytes, and the downsample kernel is built as wave64 on gfx12. The output image is byte-identical before and after.
diff --git a/kernels/build.sh b/kernels/build.sh
index efbc1be..d844c55 100755
--- a/kernels/build.sh
+++ b/kernels/build.sh
@@ -118,8 +118,11 @@ if [[ "$WHAT" == all || "$WHAT" == tex || "$WHAT" == l ]]; then
             for v in 3_1 3_2; do specs+=("rrlite_post_${v}_${mv}_${range}:sust_only:w64"); done
         done
     done
+    # downsample: wave64 is slower on RDNA3 and faster on RDNA4 (RX 9070 XT: 0.37 -> 0.34 ms at 4K, same output)
+    downsample_mode=""
+    [[ "$ARCH" == gfx12* ]] && downsample_mode=":w64"
     for mode in static dynamic; do
-        for range in hdr ldr; do specs+=("rrlite_downsample_kernel_${mode}_${range}:sust_only"); done
+        for range in hdr ldr; do specs+=("rrlite_downsample_kernel_${mode}_${range}:sust_only${downsample_mode}"); done
     done
     # The fast set serves two K kernels from native code translated from the PTX (kernels/native): these builds
     # replace the sust_only build of the same name. The accuracy set keeps ZLUDA's compile (denormals preserved).
diff --git a/kernels/m/swin_block.h b/kernels/m/swin_block.h
index 29cf769..a3b2505 100644
--- a/kernels/m/swin_block.h
+++ b/kernels/m/swin_block.h
@@ -97,6 +97,25 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
     half_t* QP = (half_t*)U;                   // stage 2: per-wave Q / P scratch [NW][16][QST]
     half_t* PA = QP;                           // stage 0: low-res tokens f16 [16][PAST] (apos order)
     half_t* G = (half_t*)U;                    // stage 4: fc1 outputs [R][64][GST]
+#ifdef SWIN_A8
+    // e4m3 byte images of the GEMM operands: the same rows and columns as the f16 images, one byte per value.
+    // A token's bytes start where its f16 row starts, so they stay inside that row.
+    constexpr int ARB = 2 * AST;
+    uint8_t* const A8 = (uint8_t*)A;
+    uint8_t* const G8 = U;
+    // codes of one column in four rows
+    auto put22 = [](uint8_t* p0, uint8_t* p1, uint8_t* p2, uint8_t* p3, hv2 a, hv2 b) {
+        const uint32_t c = codes4(a, b);
+        *p0 = (uint8_t)c;
+        *p1 = (uint8_t)(c >> 8);
+        *p2 = (uint8_t)(c >> 16);
+        *p3 = (uint8_t)(c >> 24);
+    };
+#else
+    constexpr int ARB = AST;
+    half_t* const A8 = A;
+    half_t* const G8 = G;
+#endif
 
     const int lane = threadIdx.x, wv = threadIdx.z, tid = lane + 32 * wv;
     const int l16 = lane & 15, hi = lane >> 4;
@@ -169,7 +188,9 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
             const int MX = mirror((8 * bx - p.sx) / 2 + (r & 3), W2), MY = mirror((8 * by - p.sy) / 2 + (r >> 2), H2);
             const uint32_t bytes =
                 input_word((const uint32_t*)(p.in + ((size_t)((c4 >> 5) * H2 + MY) * W2 + MX) * 32 + (c4 & 31)));
-#ifdef SWIN_PACKED_INPUT
+#ifdef SWIN_A8
+            *(uint32_t*)((uint8_t*)PA + r * PAST + apos(c4)) = bytes; // apos(c4 + b) = apos(c4) + b
+#elif defined(SWIN_PACKED_INPUT)
             *(h4*)(PA + r * PAST + apos(c4)) = decode4(bytes);
 #else
 #pragma unroll
@@ -187,7 +208,11 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
 #pragma unroll
         for (int kc = 0; kc < CIN / 32; ++kc)
         {
+#ifdef SWIN_A8
+            const gop a0 = lda((const uint8_t*)PA + l16 * PAST + 32 * kc), a1 = lda((const uint8_t*)PA + l16 * PAST + 32 * kc + 16);
+#else
             const gop a0 = lda(PA + l16 * PAST + 32 * kc), a1 = lda(PA + l16 * PAST + 32 * kc + 16);
+#endif
 #pragma unroll
             for (int j = 0; j < NTEW; ++j)
             {
@@ -318,6 +343,20 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
             const half_t rn = rsqrt16(a0 + a1);
 #endif
             const hv2 rn2 = {rn, rn};
+#ifdef SWIN_A8
+#pragma unroll
+            for (int pl = 0; pl < NPL; ++pl)
+#pragma unroll
+                for (int ww = 0; ww < 2; ++ww)
+                {
+                    const int c = 32 * pl + 4 * (t + 4 * ww);
+                    const hv2 g0 = __builtin_bit_cast(hv2, *(const uint32_t*)(W + L::G1 + 2 * gperm(c)));
+                    const hv2 g1 = __builtin_bit_cast(hv2, *(const uint32_t*)(W + L::G1 + 2 * gperm(c + 2)));
+                    // apos(c + b) = apos(c) + b for b < 4: the four codes are one word
+                    *(uint32_t*)(A8 + T * ARB + apos(c)) = codes4((hv2){v[pl][ww][0], v[pl][ww][1]} * (hv2)(rn2 * g0),
+                                                                  (hv2){v[pl][ww][2], v[pl][ww][3]} * (hv2)(rn2 * g1));
+                }
+#else
 #pragma unroll
             for (int pl = 0; pl < NPL; ++pl)
 #pragma unroll
@@ -332,6 +371,7 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
                         // apos(c + 1) = apos(c) + 1 for even c
                         *(hv2*)(A + T * AST + apos(c)) = hq;
                     }
+#endif
         }
     }
     __syncthreads();
@@ -400,8 +440,8 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
 #pragma unroll
             for (int m = 0; m < MW; ++m)
             {
-                const gop a0 = lda(A + (16 * (m0 + m) + l16) * AST + 32 * kc);
-                const gop a1 = lda(A + (16 * (m0 + m) + l16) * AST + 32 * kc + 16);
+                const gop a0 = lda(A8 + (16 * (m0 + m) + l16) * ARB + 32 * kc);
+                const gop a1 = lda(A8 + (16 * (m0 + m) + l16) * ARB + 32 * kc + 16);
 #pragma unroll
                 for (int q = 0; q < 4; ++q)
                     k32(acc[m][q], a0, b[0][q], a1, b[1][q]);
@@ -415,7 +455,11 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
 #pragma unroll
                 for (int i = 0; i < 8; i += 2)
                 {
+#ifdef SWIN_A8
+                    qp[m][q][i / 2] = acc[m][q].pair(i / 2); // encoded when staged
+#else
                     qp[m][q][i / 2] = q8x2(acc[m][q].pair(i / 2));
+#endif
                     vp[m][q][i / 2] = acc[m][2 + q].pair(i / 2);
                 }
 #endif
@@ -437,6 +481,18 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
 #pragma unroll
         for (int m = 0; m < MW; ++m)
         {
+#ifdef SWIN_A8
+            uint8_t* const Q8 = (uint8_t*)Qs;
+#pragma unroll
+            for (int q = 0; q < 2; ++q)
+#pragma unroll
+                for (int i = 0; i < 8; i += 4)
+                    put22(Q8 + mrow(i) * QST + 16 * q + l16, Q8 + mrow(i + 1) * QST + 16 * q + l16,
+                          Q8 + mrow(i + 2) * QST + 16 * q + l16, Q8 + mrow(i + 3) * QST + 16 * q + l16, qp[m][q][i / 2],
+                          qp[m][q][i / 2 + 1]);
+            wave_sync();
+            const gop a0 = lda(Q8 + l16 * QST), a1 = lda(Q8 + l16 * QST + 16);
+#else
 #pragma unroll
             for (int q = 0; q < 2; ++q)
 #pragma unroll
@@ -447,6 +503,7 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
                 }
             wave_sync();
             const gop a0 = lda(Qs + l16 * QST), a1 = lda(Qs + l16 * QST + 16);
+#endif
             T16 S = bias;
             k32(S, a0, a0, a1, a1);
             // rows mrow(2k), mrow(2k+1) of this lane's column, packed: scale, clamp and the cubic
@@ -509,6 +566,15 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
 #endif
                 const f8v o = wmma(pa, vb, splat(0.0f));
                 const int ocol = apos(32 * h + ginv(16 * nt + l16));
+#ifdef SWIN_A8
+#pragma unroll
+                for (int i = 0; i < 8; i += 4)
+                {
+                    uint8_t* const oc = A8 + 16 * (m0 + m) * ARB + ocol;
+                    put22(oc + mrow(i) * ARB, oc + mrow(i + 1) * ARB, oc + mrow(i + 2) * ARB, oc + mrow(i + 3) * ARB,
+                          (hv2){(half_t)o[i], (half_t)o[i + 1]}, (hv2){(half_t)o[i + 2], (half_t)o[i + 3]});
+                }
+#else
 #pragma unroll
                 for (int i = 0; i < 8; i += 2)
                 {
@@ -516,6 +582,7 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
                     A[(16 * (m0 + m) + mrow(i)) * AST + ocol] = oq[0];
                     A[(16 * (m0 + m) + mrow(i + 1)) * AST + ocol] = oq[1];
                 }
+#endif
             }
             wave_sync();
         }
@@ -545,8 +612,8 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
 #pragma unroll
         for (int mm = 0; mm < MT; ++mm)
         {
-            const gop a0 = lda(A + (16 * (mg * MT + mm) + l16) * AST + 32 * kc);
-            const gop a1 = lda(A + (16 * (mg * MT + mm) + l16) * AST + 32 * kc + 16);
+            const gop a0 = lda(A8 + (16 * (mg * MT + mm) + l16) * ARB + 32 * kc);
+            const gop a1 = lda(A8 + (16 * (mg * MT + mm) + l16) * ARB + 32 * kc + 16);
 #pragma unroll
             for (int j = 0; j < NTW; ++j)
                 k32(x[mm][j], a0, b[0][j], a1, b[1][j]);
@@ -647,6 +714,26 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
         const int col = apos(ginv(n));
 #pragma unroll
         for (int mm = 0; mm < MT; ++mm)
+#ifdef SWIN_A8
+#pragma unroll
+            for (int i = 0; i < 8; i += 4)
+            {
+                hv2 h[2];
+                uint8_t* at[4];
+#pragma unroll
+                for (int e = 0; e < 2; ++e)
+                {
+                    const int row = 16 * (mg * MT + mm) + mrow(i + 2 * e), row2 = 16 * (mg * MT + mm) + mrow(i + 2 * e + 1);
+                    const hv2 v = x[mm][j].pair(i / 2 + e);
+                    const hv2 r2 = {R2[row], R2[row2]};
+                    h[e] = v * (hv2)(r2 * (hv2){g2, g2});
+                    at[2 * e] = A8 + row * ARB + col;
+                    at[2 * e + 1] = A8 + row2 * ARB + col;
+                    x[mm][j].set_pair(i / 2 + e, v + (hv2){b2, b2});
+                }
+                put22(at[0], at[1], at[2], at[3], h[0], h[1]);
+            }
+#else
 #pragma unroll
             for (int i = 0; i < 8; i += 2)
             {
@@ -658,6 +745,7 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
                 A[row2 * AST + col] = hq[1];
                 x[mm][j].set_pair(i / 2, v + (hv2){b2, b2});
             }
+#endif
     }
     __syncthreads();
 
@@ -695,8 +783,8 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
 #pragma unroll
                 for (int mm = 0; mm < MPW; ++mm)
                 {
-                    const gop a0 = lda(A + (16 * (mb + mm) + l16) * AST + 32 * kc);
-                    const gop a1 = lda(A + (16 * (mb + mm) + l16) * AST + 32 * kc + 16);
+                    const gop a0 = lda(A8 + (16 * (mb + mm) + l16) * ARB + 32 * kc);
+                    const gop a1 = lda(A8 + (16 * (mb + mm) + l16) * ARB + 32 * kc + 16);
                     k32(z[mm][0], a0, b[0][0], a1, b[1][0]);
                     k32(z[mm][1], a0, b[0][1], a1, b[1][1]);
                 }
@@ -708,6 +796,27 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
                 for (int nt = 0; nt < 2; ++nt)
                 {
                     const int col = apos(ginv(16 * nt + l16));
+#ifdef SWIN_A8
+#pragma unroll
+                    for (int i = 0; i < 8; i += 4)
+                    {
+                        hv2 gl[2];
+#pragma unroll
+                        for (int e = 0; e < 2; ++e)
+                        {
+                            const hv2 zv = z[mm][nt].pair(i / 2 + e);
+                            const hv2 cz = __builtin_elementwise_max(__builtin_elementwise_min(zv, (hv2){2.0f16, 2.0f16}),
+                                                                     (hv2){-2.0f16, -2.0f16});
+                            const hv2 az = __builtin_bit_cast(hv2, (u16x2)(__builtin_bit_cast(u16x2, cz) & (unsigned short)0x7fff));
+                            const hv2 inner = (hv2){f16(0.41216981f), f16(0.41216981f)} -
+                                              (hv2)((hv2){f16(0.08108133f), f16(0.08108133f)} * az);
+                            gl[e] = zv * (hv2)((hv2){0.5f16, 0.5f16} + (hv2)(cz * inner));
+                        }
+                        uint8_t* const gc = G8 + (cw * 64 + 16 * (mb + mm)) * GST + col;
+                        put22(gc + mrow(i) * GST, gc + mrow(i + 1) * GST, gc + mrow(i + 2) * GST, gc + mrow(i + 3) * GST, gl[0],
+                              gl[1]);
+                    }
+#else
 #pragma unroll
                     for (int i = 0; i < 8; i += 2)
                     {
@@ -722,6 +831,7 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
                         Gc[(16 * (mb + mm) + mrow(i)) * GST + col] = gq[0];
                         Gc[(16 * (mb + mm) + mrow(i + 1)) * GST + col] = gq[1];
                     }
+#endif
                 }
             __syncthreads();
 #pragma unroll
@@ -737,8 +847,8 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
 #pragma unroll
                 for (int mm = 0; mm < MT; ++mm)
                 {
-                    const gop a0 = lda(G + (cc * 64 + 16 * (mg * MT + mm) + l16) * GST);
-                    const gop a1 = lda(G + (cc * 64 + 16 * (mg * MT + mm) + l16) * GST + 16);
+                    const gop a0 = lda(G8 + (cc * 64 + 16 * (mg * MT + mm) + l16) * GST);
+                    const gop a1 = lda(G8 + (cc * 64 + 16 * (mg * MT + mm) + l16) * GST + 16);
 #pragma unroll
                     for (int j = 0; j < NTW; ++j)
                         k32(x[mm][j], a0, b[0][j], a1, b[1][j]);
@@ -752,6 +862,9 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
     // XO: [64 tokens][C channels] natural order; the row padding keeps the patch merge's 8-byte reads of 16
     // tokens on distinct LDS banks (see xo_row)
     half_t* XO = A;
+#ifdef SWIN_A8
+    constexpr int XB = C + 4; // byte row stride of the patch merge's code image (rows stay 4-byte aligned)
+#endif
     static_assert(xo_row<C>(63) + C <= 64 * AST, "XO fits in A");
 #pragma unroll
     for (int j = 0; j < NTW; ++j)
@@ -775,6 +888,9 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
         for (int k = 0; k < 4; ++k)
         {
             const uint2 pr = xs[k];
+#ifdef SWIN_A8
+            w[k] = codes4(__builtin_bit_cast(hv2, pr.x), __builtin_bit_cast(hv2, pr.y));
+#else
             hv2 qa, qb;
             uint32_t ca, cb;
 #ifdef SWIN_DIRECT_CODES
@@ -791,8 +907,20 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
             // the patch merge consumes decode(code) = q8_exact(x), written back over x (this thread's own values)
             if constexpr (NPM > 0)
                 xs[k] = (uint2){__builtin_bit_cast(uint32_t, qa), __builtin_bit_cast(uint32_t, qb)};
+#endif
         }
         const uint4 v = {w[0], w[1], w[2], w[3]};
+#ifdef SWIN_A8
+        // the patch merge's operands are these codes: a byte image [64][XB] in U (G is dead)
+        if constexpr (NPM > 0)
+        {
+            static_assert(64 * XB <= USIZE, "merge byte image fits in U");
+            uint32_t* xb = (uint32_t*)(U + T * XB + c16);
+#pragma unroll
+            for (int k = 0; k < 4; ++k)
+                xb[k] = w[k];
+        }
+#endif
         const int X = 8 * bx - p.sx + tok_x(T), Y = 8 * by - p.sy + tok_y(T);
         if (X >= 0 && X < p.tw && Y >= 0 && Y < p.th)
             *(uint4*)(p.out + ((size_t)((c16 >> 5) * p.th + Y) * p.tw + X) * 32 + (c16 & 31)) = v;
@@ -822,6 +950,20 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
 #pragma unroll
             for (int s = 0; s < 2; ++s)
             {
+#ifdef SWIN_A8
+                // this half's two runs of four K slots, as stored codes
+                uint32_t run4[2];
+#pragma unroll
+                for (int rr = 0; rr < 2; ++rr)
+                {
+                    const int k = 32 * kc + kslot(s, 4 * (2 * hi + rr));
+                    const int q = k / C, ch = k % C;
+                    const int tx = 2 * mx + (q & 1), ty = 2 * my + (q >> 1);
+                    const int T = (tx & 3) + 4 * (ty & 3) + 16 * (tx >> 2) + 32 * (ty >> 2);
+                    run4[rr] = *(const uint32_t*)(U + T * XB + ch);
+                }
+                as[s] = (u2v){run4[0], run4[1]};
+#else
                 sop a;
                 // the operand's K slots in runs of 4 (gfx12: this half's slots 8 hi .. 8 hi + 7)
                 constexpr int RUNS = sizeof(sop) / sizeof(half_t) / 4;
@@ -840,6 +982,7 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
                         a[4 * rr + b] = v[b];
                 }
                 as[s] = as_gop(a);
+#endif
             }
 #pragma unroll
             for (int j = 0; j < NTP; ++j)
@@ -857,6 +1000,9 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
             if (!PM_EVEN && wv * NTP + j >= NPM / 16)
                 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))};
+#endif
 #pragma unroll
             for (int i = 0; i < 8; ++i)
             {
@@ -865,7 +1011,11 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara
                 if (MX < 0 || MX >= W2 || MY < 0 || MY >= H2)
                     continue;
                 p.p48[((size_t)((ch >> 5) * H2 + MY) * W2 + MX) * 32 + (ch & 31)] =
+#ifdef SWIN_A8
+                    (uint8_t)(code[i >> 2] >> (8 * (i & 3)));
+#else
                     (uint8_t)enc8(y[j].get(i));
+#endif
             }
         }
     }
diff --git a/kernels/m/swin_common.h b/kernels/m/swin_common.h
index f07c28f..110c758 100644
--- a/kernels/m/swin_common.h
+++ b/kernels/m/swin_common.h
@@ -28,6 +28,13 @@
 #undef SWIN_F32ACC
 #endif
 
+// SWIN_A8 (RDNA4 native FP8): the e4m3 GEMMs' activations live in LDS as e4m3 bytes, encoded once where they
+// are produced (v_cvt_pk_fp8_f32) instead of as f16 values that every operand load encodes again. The bytes
+// are the codes of the same q8 values, so results are unchanged.
+#if defined(D4R_FP8_WMMA) && defined(__GFX12__) && !defined(SWIN_NO_A8)
+#define SWIN_A8 1
+#endif
+
 #pragma clang fp contract(off)
 
 typedef _Float16 half_t;
@@ -199,6 +206,35 @@ __device__ __forceinline__ uint32_t codes8x2(hv2 q)
     return __builtin_bit_cast(uint32_t, (u16x2)(((qb >> 8) & (unsigned short)0x80) | code));
 }
 
+#ifdef SWIN_A8
+// e4m3 codes of two halves (RNE, satfinite, NaN -> 0x7f | sign: equal to half_to_e4m3 for all inputs) in bytes 0
+// and 1. The instruction rounds like NVIDIA's conversion but encodes overflow as 0x7f, so magnitudes are
+// clamped to 448 first and NaNs enter as infinities.
+__device__ __forceinline__ uint32_t codes2(hv2 h)
+{
+    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
+    return codes2(a) | (codes2(b) << 16);
+#else
+    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 w - ((((w & 0x7f7f7f7fu) + 0x01010101u) >> 7) & 0x01010101u);
+#endif
+}
+#endif
+
 // four e4m3 codes (a, b from enc8x2) as the bytes of one word in element order
 __device__ __forceinline__ uint32_t pack_codes(uint32_t a, uint32_t b)
 {
@@ -336,6 +372,13 @@ __device__ __forceinline__ gop lda(const half_t* p)
 {
     return to_fp8(lds16(p));
 }
+#ifdef SWIN_A8
+// GEMM A operand from an e4m3 byte image in LDS: this half's 8 K values
+__device__ __forceinline__ gop lda(const uint8_t* p)
+{
+    return *(const u2v*)(p + 8 * m_half());
+}
+#endif
 #else
 __device__ __forceinline__ f8v wmma8(gop a, gop b, f8v c)
 {

From ffa686e079926fbb355134566579dabef83ae060 Mon Sep 17 00:00:00 2001
From: woutermans <woutertje554@gmail.com>
Date: Tue, 6 Oct 2026 14:29:21 +0200
Subject: [PATCH 2/4] 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   |  4 ++--
 kernels/build.sh         | 12 ++++++++++--
 kernels/m/swin_common.h  | 38 +++++++++++++++-----------------------
 kernels/tex/build_tex.sh |  5 +++--
 4 files changed, 30 insertions(+), 29 deletions(-)

diff --git a/docs/native-kernels.md b/docs/native-kernels.md
index f01bc99..fc8e202 100644
--- a/docs/native-kernels.md
+++ b/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.
 
diff --git a/kernels/build.sh b/kernels/build.sh
index d844c55..2fae61a 100755
--- a/kernels/build.sh
+++ b/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" \
-            D4R_ZLUDA_WAVE64="$([[ "$mode" == w64 && "$ACCURACY" == 0 ]] && echo 1 || echo 0)" D4R_ROCM_DIR="$ROCM" D4R_GPU_ARCH="$ARCH" \
+        # accuracy sets are wave32, except M's post and downsample kernels on RDNA4: their wave64 builds keep the
+        # 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
diff --git a/kernels/m/swin_common.h b/kernels/m/swin_common.h
index 110c758..34335d6 100644
--- a/kernels/m/swin_common.h
+++ b/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
-// and 1. The instruction rounds like NVIDIA's conversion but encodes overflow as 0x7f, so magnitudes are
-// clamped to 448 first and NaNs enter as infinities.
-__device__ __forceinline__ uint32_t codes2(hv2 h)
-{
-    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.
+// e4m3 codes of four halves (RNE, satfinite) in the bytes of one word. The instruction rounds like NVIDIA's
+// conversion but encodes overflow as 0x7f, so the values are clamped to +-448 first, with the NaN-propagating
+// minimum / maximum. Equal to half_to_e4m3 for every non-NaN input; a NaN encodes as the NaN code 0xff (the
+// reference keeps its sign: 0x7f or 0xff).
 __device__ __forceinline__ uint32_t codes4(hv2 a, hv2 b)
 {
-#ifdef SWIN_EXACT
-    return codes2(a) | (codes2(b) << 16);
-#else
-    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 w - ((((w & 0x7f7f7f7fu) + 0x01010101u) >> 7) & 0x01010101u);
-#endif
+    const hv2 lim = {(half_t)448.0f, (half_t)448.0f};
+    a = __builtin_elementwise_maximum(__builtin_elementwise_minimum(a, lim), -lim);
+    b = __builtin_elementwise_maximum(__builtin_elementwise_minimum(b, lim), -lim);
+    const uint32_t w = (uint32_t)__builtin_amdgcn_cvt_pk_fp8_f32((float)a[0], (float)a[1], 0, false);
+    return (uint32_t)__builtin_amdgcn_cvt_pk_fp8_f32((float)b[0], (float)b[1], (int)w, true);
 }
 #endif
 
diff --git a/kernels/tex/build_tex.sh b/kernels/tex/build_tex.sh
index bf61e30..e8163e3 100755
--- a/kernels/tex/build_tex.sh
+++ b/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' ')"

From 11c84193e6eba2d942f101368386c443c10ff3ac Mon Sep 17 00:00:00 2001
From: woutermans <woutertje554@gmail.com>
Date: Tue, 6 Oct 2026 14:59:14 +0200
Subject: [PATCH 3/4] docs: compare the RDNA4 FP8 set with the 16-bit path

The 2x figure is against the previous native-FP8 layers. With NativeFp8 off
the RX 9070 XT was already at 3.37 ms for 1280x720 -> 3840x2160; the FP8 set
is 2.88 ms there.
---
 docs/performance.md | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/docs/performance.md b/docs/performance.md
index 2304b6d..2e6cb55 100644
--- a/docs/performance.md
+++ b/docs/performance.md
@@ -92,6 +92,8 @@ DLSS 4.5 (M) on an RX 9070 XT (gfx1201, native FP8, fast set), 1280×720 → 384
 | dec0, enc0 (texture kernels) | 0.18, 0.13 | 0.18, 0.13 |
 | **total** | **6.10** | **2.86** |
 
+With `NativeFp8` off the same card runs the layers as 16-bit math, which never re-encoded its operands: that set takes 3.37 ms here (5.94 ms at 1920×1080 → 3840×2160, against 4.63 ms for the FP8 set). So against a 9070 XT that was already running with `NativeFp8` off, the FP8 set is about 15% faster at 1280×720 and 22% at 1920×1080, not twice as fast. In Ghost of Tsushima at 1920×1080 → 3840×2160 the upscaler time shown in game went from 7.0 ms (16-bit set) to 5.8 ms (this FP8 set).
+
 The shim's whole-evaluation GPU time (`D4R_PROFILE`, no per-kernel synchronisation) went from 5.9–6.0 ms to 2.80 ms with the harness's 200 ms pause between frames, and to 2.39–2.43 ms with frames 0–4 ms apart (2000 frames at 4 ms spacing: median 2.39, maximum 2.56). The difference is the driver's clock governor: with the 200 ms pause the GPU is about 9% busy and runs at about 2790 MHz, with frames 0–4 ms apart it is 80–98% busy and runs at 3180–3300 MHz, as it does under a game's rendering load. Other GPU users on the desktop add occasional slower frames to any of these figures; the previous kernels show the same disturbances.
 
 The change is in [native-kernels.md](native-kernels.md#rdna4): the FP8 layers keep their activations as e4m3 bytes, and the downsample kernel is built as wave64 on gfx12. The output image is byte-identical before and after.

From cf741c2e3ddfa0b94c497c8a4993316324d484e3 Mon Sep 17 00:00:00 2001
From: woutermans <woutertje554@gmail.com>
Date: Tue, 6 Oct 2026 15:18:11 +0200
Subject: [PATCH 4/4] docs: RDNA4 fast and accuracy sets, FP8 and 16-bit,
 before and after

---
 docs/native-kernels.md |  2 +-
 docs/performance.md    | 11 ++++++++++-
 2 files changed, 11 insertions(+), 2 deletions(-)

diff --git a/docs/native-kernels.md b/docs/native-kernels.md
index fc8e202..7effe5e 100644
--- a/docs/native-kernels.md
+++ b/docs/native-kernels.md
@@ -67,7 +67,7 @@ The gfx12 M variant can use native e4m3 FP8 WMMA with `D4R_NATIVE_FP8=1`. Its ac
 
 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. 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.
+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.06 ms per evaluation in the same harness (2.59 ms under load) against 2.77 ms (2.38 ms) for the fast set; with the previous FP8 layers it cost 6.49 ms (6.32 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.
 
diff --git a/docs/performance.md b/docs/performance.md
index 2e6cb55..8d45382 100644
--- a/docs/performance.md
+++ b/docs/performance.md
@@ -92,7 +92,16 @@ DLSS 4.5 (M) on an RX 9070 XT (gfx1201, native FP8, fast set), 1280×720 → 384
 | dec0, enc0 (texture kernels) | 0.18, 0.13 | 0.18, 0.13 |
 | **total** | **6.10** | **2.86** |
 
-With `NativeFp8` off the same card runs the layers as 16-bit math, which never re-encoded its operands: that set takes 3.37 ms here (5.94 ms at 1920×1080 → 3840×2160, against 4.63 ms for the FP8 set). So against a 9070 XT that was already running with `NativeFp8` off, the FP8 set is about 15% faster at 1280×720 and 22% at 1920×1080, not twice as fast. In Ghost of Tsushima at 1920×1080 → 3840×2160 the upscaler time shown in game went from 7.0 ms (16-bit set) to 5.8 ms (this FP8 set).
+Fast and accuracy sets, on the native-FP8 path and on the 16-bit path (`NativeFp8` off), before and after the change, same harness and resolution. Median GPU time per frame in milliseconds from `D4R_CUDA_KERNEL_PROFILE`; the bracketed figures are the whole evaluation (`D4R_PROFILE`) with the GPU otherwise idle and at about 80% utilisation.
+
+| Set | before | now |
+|---|---|---|
+| native FP8, fast | 6.11 (6.01 / 5.89) | 2.88 (2.77 / 2.38) |
+| native FP8, accuracy | 6.38 (6.49 / 6.32) | 3.15 (3.06 / 2.59) |
+| 16-bit, fast | 3.48 | 3.40 |
+| 16-bit, accuracy | 4.38 | 4.20 |
+
+The 16-bit path never re-encoded its operands, so it was already faster than the previous FP8 layers; its small gain here is the wave64 downsample (and post, in the accuracy set). Against a 9070 XT that was running with `NativeFp8` off, the FP8 fast set is about 17% faster at 1280×720 and 22% at 1920×1080 → 3840×2160 (4.63 ms against 5.94 ms, the latter with a 16-bit set from an earlier build), not twice as fast. In Ghost of Tsushima at 1920×1080 → 3840×2160 the upscaler time shown in game went from 7.0 ms (16-bit fast set) to 5.8 ms (this FP8 fast set).
 
 The shim's whole-evaluation GPU time (`D4R_PROFILE`, no per-kernel synchronisation) went from 5.9–6.0 ms to 2.80 ms with the harness's 200 ms pause between frames, and to 2.39–2.43 ms with frames 0–4 ms apart (2000 frames at 4 ms spacing: median 2.39, maximum 2.56). The difference is the driver's clock governor: with the 200 ms pause the GPU is about 9% busy and runs at about 2790 MHz, with frames 0–4 ms apart it is 80–98% busy and runs at 3180–3300 MHz, as it does under a game's rendering load. Other GPU users on the desktop add occasional slower frames to any of these figures; the previous kernels show the same disturbances.