#20315·tvm

[Bug][TIRx][FP8] DTypeConversion mis-encodes float8_e5m2 zero/subnormal values when converting to/from float32

Author: LittlehamsterXuCreated Sep 11, 2026Updated Sep 16, 2026

Labelstype: bugneeds-triage

Summary

On targets without native FP8 support, tirx.transform.FP8ComputeLegalize uses tvm::tirx::DTypeConversion for FP8 promotion and narrowing. DTypeConversion mis-encodes valid float8_e5m2 zero and subnormal values when converting to/from float32.

Environment

Target: llvm on x86_64 Linux
Dependencies: numpy, ml_dtypes

The runtime reproductions used TVM v0.25.0.post1 (b3e249b7d75f8f3bc7cbee48188d3c80ae323437) and a source build based on upstream main commit 141a187f35e9b6263db6665bc99eb03146d9b23f. On 2026-09-11, upstream main was 8f328e802cfe5e41fcc8f5c17e7582b1c28bfce4, 15 commits ahead; the current source was inspected at that commit, and none of those intervening commits modified dtype_conversion.cc or unsupported_dtype_legalize.cc. Both runtime reproductions produced the same result below.

Reproduction

Save the following as reproduce-tvm-fp8-zero-subnormal-float32.py and run it in an environment containing TVM 0.25.0.post1:

bash

python reproduce-tvm-fp8-zero-subnormal-float32.py

python

import ml_dtypes
import numpy as np
import tvm
from tvm import tirx
from tvm.script import tirx as T


# +0, -0, smallest subnormal, middle subnormal, largest subnormal, smallest normal.
FP8_TO_F32_INPUT_BITS = np.array([0x00, 0x80, 0x01, 0x02, 0x03, 0x04], dtype=np.uint8)
FP8_TO_F32_EXPECTED_BITS = [0x00000000, 0x80000000, 0x37800000, 0x38000000, 0x38400000, 0x38800000]
F32_TO_FP8_INPUT_BITS = np.array(
    [0x00000000, 0x80000000, 0x37800000, 0x38000000, 0x38400000, 0x38800000],
    dtype=np.uint32,
)
F32_TO_FP8_EXPECTED_BITS = [0x00, 0x80, 0x01, 0x02, 0x03, 0x04]


@T.prim_func
def fp8_to_f32(A: T.Buffer((6,), "float8_e5m2"), B: T.Buffer((6,), "float32")):
    for i in T.serial(6):
        B[i] = A[i]


@T.prim_func
def f32_to_fp8(A: T.Buffer((6,), "float32"), B: T.Buffer((6,), "float8_e5m2")):
    for i in T.serial(6):
        B[i] = A[i]


fp8_to_f32_module = tirx.build(fp8_to_f32, target="llvm")
fp8_to_f32_src = tvm.runtime.tensor(FP8_TO_F32_INPUT_BITS.view(ml_dtypes.float8_e5m2))
fp8_to_f32_dst = tvm.runtime.tensor(np.zeros(6, dtype=np.float32))
fp8_to_f32_module["main"](fp8_to_f32_src, fp8_to_f32_dst)
fp8_to_f32_observed = [int(value) for value in fp8_to_f32_dst.numpy().view(np.uint32)]

f32_to_fp8_module = tirx.build(f32_to_fp8, target="llvm")
f32_to_fp8_src = tvm.runtime.tensor(F32_TO_FP8_INPUT_BITS.view(np.float32))
f32_to_fp8_dst = tvm.runtime.tensor(np.zeros(6, dtype=ml_dtypes.float8_e5m2))
f32_to_fp8_module["main"](f32_to_fp8_src, f32_to_fp8_dst)
f32_to_fp8_observed = [int(value) for value in f32_to_fp8_dst.numpy().view(np.uint8)]

print("float8_e5m2 -> float32 observed:", [f"0x{x:08X}" for x in fp8_to_f32_observed])
print("float8_e5m2 -> float32 expected:", [f"0x{x:08X}" for x in FP8_TO_F32_EXPECTED_BITS])
print("float32 -> float8_e5m2 observed:", [f"0x{x:02X}" for x in f32_to_fp8_observed])
print("float32 -> float8_e5m2 expected:", [f"0x{x:02X}" for x in F32_TO_FP8_EXPECTED_BITS])

Expected behavior

The six inputs are, in order, +0, -0, the smallest subnormal, the middle subnormal, the largest subnormal, and the smallest normal:

For float8_e5m2 (bias = 15, mantissa bits = 2), an exponent field of zero denotes signed zero when the mantissa is zero and a subnormal otherwise. For a subnormal with mantissa M > 0, the value is (-1)^S × M × 2^-16; normal values use (-1)^S × (1 + M / 4) × 2^(E - 15). Therefore, 0x00 and 0x80 are +0 and -0, while 0x01, 0x02, and 0x03 are 2^-16, 2^-15, and 3 × 2^-16.

value	`+0`	`-0`	smallest subnormal	middle subnormal	largest subnormal	smallest normal
`float8_e5m2` bits	`00`	`80`	`01`	`02`	`03`	`04`
`float32` bits	`00000000`	`80000000`	`37800000`	`38000000`	`38400000`	`38800000`

float8_e5m2 -> float32 should produce the listed float32 words, and float32 -> float8_e5m2 should produce the listed FP8 bytes.

Actual behavior

The reproducer prints:

float8_e5m2 -> float32 observed: 0x38000000 0xB8000000 0x38200000 0x38400000 0x38600000 0x38800000
float8_e5m2 -> float32 expected: 0x00000000 0x80000000 0x37800000 0x38000000 0x38400000 0x38800000
float32 -> float8_e5m2 observed: 0x00 0x00 0x00 0x00 0x02 0x04
float32 -> float8_e5m2 expected: 0x00 0x80 0x01 0x02 0x03 0x04

The smallest-normal control (0x04, 2^-14) is correct in both directions, while all tested exponent-zero cases are mis-encoded.

The same output was observed with the rebuilt main-derived source checkout described above.

Suspected cause

The relevant code is src/tirx/transform/dtype_conversion.cc. The two relevant branches are:

cpp

PrimExpr exponent_before_delta = ((src_uint_value << 1) >> (src_fp.mantissa + 1));
if (bias_delta >= 0) {
  PrimExpr ret_exponent =
      (bias_delta > 0) ? (cast(tgt_uint, exponent_before_delta + bias_delta) << tgt_fp.mantissa)
                       : (cast(tgt_uint, exponent_before_delta) << tgt_fp.mantissa);
  return reinterpret(tgt_dtype, ret_mantissa | ret_exponent | ret_sign);
} else {
  PrimExpr round_to_zero = exponent_before_delta < (-bias_delta);
  PrimExpr ret_exponent = cast(tgt_uint, exponent_before_delta - (-bias_delta))
                          << tgt_fp.mantissa;
  return reinterpret(tgt_dtype, if_then_else(
      round_to_zero, MakeConst(tgt_uint, 0), ret_mantissa | ret_exponent | ret_sign));
}

Widening: when the source exponent is 0, the code still applies the normal-value bias delta (112 for float8_e5m2 -> float32). For this input, exponent_before_delta is 0, so zero becomes 2^-15 instead of zero.
Narrowing: when exponent_before_delta < -bias_delta, the underflow branch directly returns zero and does not construct a destination FP8 subnormal. It also loses the sign of -0.0.

TVM declares float8_e5m2 as an IEEE-style format in dtype_conversion.h. The smallest-normal control above is correct; the exponent-zero cases are not handled. FP8StorageLegalize only remaps the storage representation and is not implicated in this numeric conversion.

Impact

This is a silent numerical correctness issue in the software FP8 legalization path. A regular float8_e5m2 zero value is sufficient to trigger the mismatch.

This affects CPU builds and any target that falls back to the software FP8 conversion, including FP8 dequantization at storage/compute boundaries.

Suggested regression test

Regression coverage for +0, -0, subnormals, and the smallest normal in both directions would catch this case.

Source: apache/tvm

View original on GitHub View discussion on GitHub