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

## 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:

```text
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`](https://github.com/apache/tvm/blob/b3e249b7d75f8f3bc7cbee48188d3c80ae323437/src/tirx/transform/dtype_conversion.cc#L63-L98). 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`](https://github.com/apache/tvm/blob/b3e249b7d75f8f3bc7cbee48188d3c80ae323437/src/tirx/transform/dtype_conversion.h#L140-L148). 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.

Start with src/tirx/transform/dtype_conversion.cc and the format definitions in src/tirx/transform/dtype_conversion.h, then run the supplied Python reproducer with numpy and ml_dtypes. Check both float8_e5m2 conversion directions for signed zero, subnormals, and the smallest normal. Done means the expected bit patterns are produced and regression coverage exists for these cases.

Written by the indexing model from the issue text.