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

@@ -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

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.

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).

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
             }
         }
     }

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)
 {