Hardware e4m3 decode in dec0's tail on the native-FP8 build

make_ptx.py turns the 32 cvt.rn.f16x2.e4m3x2 of rrlite_dec0_4x4_folded into calls of d4r_dec_e4m3x2 (tex_common.h) when D4R_TEX_FP8=1: v_cvt_pk_f32_fp8 and two conversions to f16 instead of ZLUDA's integer lowering (about 25 instructions each; the kernel goes from 4058 to 3418 instructions). The conversion equals ZLUDA's e4m3_to_f16_bits for every code except the NaN codes 0x7f / 0xff, checked over all 65536 code pairs on an RX 9070 XT, and the head never stages those codes (it clamps to +-448 before encoding). Harness frames are byte-identical in the accuracy and the fast set, over a 16-frame moving sequence in the accuracy set. dec0 at 2227x1253 render: 0.43 -> 0.38 ms (accuracy), 0.36 -> 0.30 ms (fast). -DD4R_TEX_SOFT_E4M3 gives ZLUDA's lowering; builds without D4R_TEX_FP8 keep the PTX instruction.

docs/native-kernels.md

@@ -78,6 +78,8 @@ In that build the GEMM activations live in LDS as e4m3 bytes (`SWIN_A8` in `swin
 
 **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`.
 
+dec0's tail decodes the staged codes with the hardware conversion too. On this build `make_ptx.py` turns its 32 `cvt.rn.f16x2.e4m3x2` into calls of `d4r_dec_e4m3x2` (`tex_common.h`: `v_cvt_pk_f32_fp8` and two conversions to f16), where ZLUDA lowers each one to about 25 integer instructions. The conversion equals ZLUDA's `e4m3_to_f16_bits` for every code except the NaN codes `0x7f` / `0xff` (all 65536 code pairs checked on an RX 9070 XT), and the head never stages those: it clamps to ±448 before it encodes. Harness frames are byte-identical in both sets, over a 16-frame moving sequence in the accuracy set. dec0 at 2227×1253 render: 0.43 → 0.38 ms in the accuracy set and 0.36 → 0.30 ms in the fast set. With `-DD4R_TEX_SOFT_E4M3` the function is ZLUDA's lowering; builds without `D4R_TEX_FP8` keep the PTX instruction.
+
 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 |

kernels/tex/build_tex.sh

@@ -27,7 +27,7 @@ read -r -a FLAGS <<< "${D4R_TEX_CFLAGS:-}"
 "$LLVM/llvm-dis" "$B/raw.bc" -o "$B/raw.ll"
 sed -E -e '/@llvm.used/d' -e '/wchar_size/d' -e '/llvm.module.flags/d' -e '/__hip_cuid/d' -e 's/optnone//g' \
     -e "s/\"target-cpu\"=\"$ARCH\"//g" -e 's/"target-features"="[^"]+"//g' "$B/raw.ll" | "$LLVM/llvm-as" -o "$B/extra.bc" -
-D4R_PREFER_ACCURACY="$ACCURACY" python3 "$D/make_ptx.py" "$K" "$B/$K.ptx" >/dev/null
+D4R_PREFER_ACCURACY="$ACCURACY" D4R_TEX_FP8="$FP8" python3 "$D/make_ptx.py" "$K" "$B/$K.ptx" >/dev/null
 # the code-generation switches of the runtime (d4r.ini [Kernels]), which ZLUDA's module cache keys on
 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))" \

kernels/tex/make_ptx.py

@@ -56,6 +56,7 @@ KERNELS = {
         delete=(139, 587),
         delete_check=("mov.u32 %r17, %laneid;", "st.shared.v2.u16 [%r570+1168], {%rs63, %rs64};"),
         sust=True,
+        dec8=32,
         extern=""".extern .func d4r_dec0_head
 (
 	.param .b32 d4r_a0, .param .b64 d4r_a1, .param .b64 d4r_a2, .param .b32 d4r_a3,
@@ -176,6 +177,36 @@ def replace_sust(lines):
     return out, n
 
 
+# dec0 on the native-FP8 build (D4R_TEX_FP8=1): the tail's e4m3 decodes call d4r_dec_e4m3x2 (tex_common.h, the
+# hardware conversion) instead of ZLUDA's integer lowering of cvt.rn.f16x2.e4m3x2
+DEC8_EXTERN = """.extern .func (.param .b32 d4r_dr) d4r_dec_e4m3x2
+(
+	.param .b16 d4r_dc
+)
+;
+"""
+DEC8_RE = re.compile(r"^cvt\.rn\.f16x2\.e4m3x2 (%r\d+), (%rs\d+);$")
+
+
+def replace_dec8(lines):
+    out, n = [], 0
+    for line in lines:
+        m = DEC8_RE.match(line.strip())
+        if not m:
+            out.append(line)
+            continue
+        dst, src = m.groups()
+        out.append(f"""{{
+	.param .b16 c0;
+	st.param.b16 [c0+0], {src};
+	.param .b32 d0;
+	call (d0), d4r_dec_e4m3x2, (c0);
+	ld.param.b32 {dst}, [d0+0];
+}}""")
+        n += 1
+    return out, n
+
+
 RZ_RE = re.compile(r"^add\.rz\.ftz\.f32 (%f\d+), (%f\d+), (%f\d+);$")
 CVT_RE = re.compile(r"^cvt\.rzi\.f32\.f32 (%f\d+), (%f\d+);$")
 
@@ -318,6 +349,11 @@ def main():
         at = next(i for i, l in enumerate(body) if l.startswith(".reg"))
         body = body[:at] + [f".reg .f32 %rq<{4 * nrz + 1}>;", f".reg .pred %pq<{nrz + 1}>;"] + body[at:]
         n["rz"] = nrz
+    if k.get("dec8") and os.environ.get("D4R_TEX_FP8", "0") == "1":
+        body, n["dec8"] = replace_dec8(body)
+        if n["dec8"] != k["dec8"]:
+            sys.exit(f"{name}: {n['dec8']} e4m3 decodes, expected {k['dec8']}")
+        k = dict(k, extern=k["extern"] + DEC8_EXTERN)
     extern = k["extern"] + (SUST_EXTERN if n["b16"] else "") + (SUSTP_EXTERN if n["p"] else "") + (SUSTB32_EXTERN if n["b32"] else "")
     text = "\n".join(head[:entry]) + "\n" + extern + "\n".join(body + tail)
     out.write_text(text)

kernels/tex/tex_common.h

@@ -151,6 +151,32 @@ __attribute__((device)) static inline uint32_t e4m3x4_f32(float a, float b, floa
 }
 #endif
 
+// cvt.rn.f16x2.e4m3x2 for the kernels whose PTX make_ptx.py points here (dec0's tail on the native-FP8 build):
+// two e4m3 codes -> an f16 pair. The hardware conversion equals ZLUDA's e4m3_to_f16_bits for every code except
+// the NaN codes 0x7f / 0xff (a different NaN), which the head never stages: it clamps to +-448 before encoding.
+DEV __attribute__((always_inline)) uint32_t d4r_dec_e4m3x2(uint16_t codes)
+{
+#ifdef D4R_TEX_HW_E4M3
+    typedef float dec_f2v __attribute__((ext_vector_type(2)));
+    typedef _Float16 dec_h2v __attribute__((ext_vector_type(2)));
+    const dec_f2v f = __builtin_amdgcn_cvt_pk_f32_fp8((int)codes, false);
+    return __builtin_bit_cast(uint32_t, (dec_h2v){(_Float16)f[0], (_Float16)f[1]});
+#else
+    uint32_t out = 0;
+#pragma unroll
+    for (int i = 0; i < 2; ++i)
+    {
+        const uint32_t code = (codes >> (8 * i)) & 0xffu, magnitude = code & 0x7fu;
+        const uint32_t normal = (magnitude << 7) + 0x2000u;
+        const uint32_t subnormal = __builtin_bit_cast(uint16_t, (_Float16)((float)(magnitude & 7u) * 0x1p-9f));
+        uint32_t half = magnitude >= 8u ? normal : subnormal;
+        half = magnitude == 0x7fu ? 0x7fffu : half;
+        out |= (((code & 0x80u) << 8) | half) << (16 * i);
+    }
+    return out;
+#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)
 {