From 5198fd38b304432a2d092fc5e8ca94974150402f Mon Sep 17 00:00:00 2001 From: woutermans Date: Tue, 6 Oct 2026 23:25:14 +0200 Subject: [PATCH] Hardware e4m3 in the texture tails: enc0 tail and dec0 head On the gfx12 native-FP8 path the enc0 tail and the dec0 head encode their results with v_cvt_pk_fp8_f32 on the f32 copies of the f16 values (e4m3x4_f32 in tex_common.h) instead of ZLUDA's integer e4m3x2. The clamps keep the conversion in range: +-2pi for enc0 (the f16 bound hi[0], same in f32) and +-448 for dec0, both as a median of three, where a NaN gives the lower bound. The codes equal e4m3x2 for every non-NaN f16, checked over all 65536 patterns on an RX 9070 XT; a NaN in dec0 now encodes as -448 instead of the NaN code. Harness frames are byte-identical in both sets (LDR and HDR, motion vectors at render and display resolution). dec0: 0.21 -> 0.15 ms at 1280x720 render and 0.74 -> 0.53 ms at 2560x1440; enc0 changes by 1-2% (its time is ZLUDA's compile of the feature code). -DD4R_TEX_SOFT_E4M3 (D4R_TEX_CFLAGS) keeps e4m3x2. --- docs/native-kernels.md | 2 ++ kernels/tex/dec0_head.hip | 9 +++++++++ kernels/tex/enc0_tail.hip | 11 +++++++++++ kernels/tex/tex_common.h | 13 +++++++++++++ 4 files changed, 35 insertions(+) diff --git a/docs/native-kernels.md b/docs/native-kernels.md index 4409ceb..75c0663 100644 --- a/docs/native-kernels.md +++ b/docs/native-kernels.md @@ -76,6 +76,8 @@ In that build the GEMM activations live in LDS as e4m3 bytes (`SWIN_A8` in `swin - `-DSWIN_NO_TOKEN_LANES` builds these layers from `swin_block.h` again. enc2 and dec2 (96 channels, four heads) use the same template: byte-identical on the same recorded launches and harness sequences, 0.22 → 0.14 and 0.20 → 0.14 ms at 1280×720 render, 0.75 → 0.47 and 0.69 → 0.47 ms at 2560×1440. - Tried, no measurable gain: interleaving the WMMAs of independent k32 steps. Skipping windows that lie entirely outside the image was also exact but below the measurement noise (at most 2.4% of the windows of every other tube launch). Not tried: true16 code generation (the half-register conversions would save about 3% of the instructions). +**Texture tails on the native-FP8 path.** The enc0 tail and the dec0 head encode their results with `v_cvt_pk_fp8_f32` on the f32 copies of the f16 values (`e4m3x4_f32` in `tex_common.h`) instead of ZLUDA's integer `e4m3x2`. With the tails' clamps (±2π for enc0, ±448 for dec0, as a median of three) the codes are equal for every non-NaN f16, checked over all 65536 patterns on an RX 9070 XT; a NaN in dec0 now encodes as −448 instead of the NaN code. Harness frames are byte-identical in both sets (LDR and HDR, motion vectors at render and display resolution). dec0: 0.21 → 0.15 ms at 1280×720 render and 0.74 → 0.53 ms at 2560×1440; enc0 changes by 1–2% (its time is ZLUDA's compile of the feature code). `-DD4R_TEX_SOFT_E4M3` (`D4R_TEX_CFLAGS`) keeps `e4m3x2`. + RX 9070 XT, accuracy set, D3D12 harness at 3840×2160 output, median GPU ms per frame (`D4R_CUDA_KERNEL_PROFILE`, frames 81–200), previous build → this one: | Render | enc1 | dec1 | enc3 tube (6 launches) | whole evaluation | diff --git a/kernels/tex/dec0_head.hip b/kernels/tex/dec0_head.hip index 80f28ef..cdf0518 100644 --- a/kernels/tex/dec0_head.hip +++ b/kernels/tex/dec0_head.hip @@ -87,12 +87,21 @@ DEV void d4r_dec0_head(uint32_t smem, uint64_t wp, uint64_t inp, uint32_t dimx, for (int i = 0; i < 8; ++i) cf[i] = (float)acc[i]; } +#ifdef D4R_TEX_HW_E4M3 + // cf holds the f16 results as f32. A NaN becomes -448 here (e4m3x2: the NaN code). +#pragma unroll + for (int i = 0; i < 8; ++i) + cf[i] = __builtin_amdgcn_fmed3f(cf[i], -448.0f, 448.0f); + E[p][0] = e4m3x4_f32(cf[0], cf[1], cf[2], cf[3]); + E[p][1] = e4m3x4_f32(cf[4], cf[5], cf[6], cf[7]); +#else uint32_t codes[4]; #pragma unroll for (int k = 0; k < 4; ++k) codes[k] = e4m3x2(pack2(acc[2 * k], acc[2 * k + 1])); E[p][0] = codes[0] | (codes[1] << 16); E[p][1] = codes[2] | (codes[3] << 16); +#endif } // staging (layout as below): this lane holds tiles 4G + hf and 4G + 2 + hf (channels 16p + 8hf ..), so it diff --git a/kernels/tex/enc0_tail.hip b/kernels/tex/enc0_tail.hip index 886f830..46251ea 100644 --- a/kernels/tex/enc0_tail.hip +++ b/kernels/tex/enc0_tail.hip @@ -88,6 +88,16 @@ DEV void d4r_enc0_tail(uint32_t smem, uint64_t wp, uint64_t outp, uint32_t dimx, for (int i = 0; i < 8; ++i) cf[i] = (float)acc[i]; } +#ifdef D4R_TEX_HW_E4M3 + // cf holds the f16 results as f32; the clamp to the f16 +-2pi bounds is the same in f32 (median of + // three: a NaN gives the lower bound, as maximumnum / minimumnum do), and keeps the conversion in range + const float bound = (float)hi[0]; +#pragma unroll + for (int i = 0; i < 8; ++i) + cf[i] = __builtin_amdgcn_fmed3f(cf[i], -bound, bound); + E[p][0] = e4m3x4_f32(cf[0], cf[1], cf[2], cf[3]); + E[p][1] = e4m3x4_f32(cf[4], cf[5], cf[6], cf[7]); +#else uint32_t codes[4]; #pragma unroll for (int k = 0; k < 4; ++k) @@ -98,6 +108,7 @@ DEV void d4r_enc0_tail(uint32_t smem, uint64_t wp, uint64_t outp, uint32_t dimx, } E[p][0] = codes[0] | (codes[1] << 16); E[p][1] = codes[2] | (codes[3] << 16); +#endif } // plane pl, word 4 j + q = columns 2q, 2q + 1 of tile (2 pl, j) then of tile (2 pl + 1, j) (tile (p, j) = diff --git a/kernels/tex/tex_common.h b/kernels/tex/tex_common.h index 3910696..1ea8503 100644 --- a/kernels/tex/tex_common.h +++ b/kernels/tex/tex_common.h @@ -138,6 +138,19 @@ __attribute__((device)) static inline f8v k32_e4m3(f8v c, const kop& a0, const k } #endif +// RDNA4 native FP8: the tails encode with the hardware conversion. e4m3x4_f32 gives the codes of four f32 values +// that hold f16 values already clamped to the finite e4m3 range (the instruction encodes overflow as NaN), in the +// bytes of one word. With the clamp it equals e4m3x2 for every non-NaN f16 (all 65536 patterns checked on an +// RX 9070 XT). -DD4R_TEX_SOFT_E4M3 keeps e4m3x2. +#if defined(D4R_TEX_FP8) && defined(__GFX12__) && !defined(D4R_TEX_SOFT_E4M3) +#define D4R_TEX_HW_E4M3 1 +__attribute__((device)) static inline uint32_t e4m3x4_f32(float a, float b, float c, float d) +{ + const uint32_t w = (uint32_t)__builtin_amdgcn_cvt_pk_fp8_f32(a, b, 0, false); + return (uint32_t)__builtin_amdgcn_cvt_pk_fp8_f32(c, d, (int)w, true); +} +#endif + // ZLUDA's f16x2_to_e4m3x2_satfinite_bits: RNE satfinite e4m3 of both halves, low half -> low byte. __attribute__((device)) static inline uint32_t e4m3x2(uint32_t bits) {