[BUG][Fuzzer][ice-on-valid-code] A finite fp32 constant above a narrow float dtype's max (fp16/fp8/fp4) aborts in `FloatImm` instead of rounding like the runtime cast

### Required prerequisites

- [x] I have read the documentation .
- [x] I have searched the [Issue Tracker](https://github.com/tile-ai/tilelang/issues) and found no existing report of this defect.

### What version of TileLang are you using?

0.1.13 (tilelang `8001cc4`, bundled tvm `8df8ebd`). The same check is present on `main`; on this pin the assertion macro is `TVM_FFI_CHECK_LE` and the whole range check is now wrapped in an `if (!std::isinf(value) && !std::isnan(value))` guard, so only finite literals are checked.

### System information

TileLang 0.1.13 (tvm `8df8ebd`); the crash is at compile time (host-side constant folding), so it is architecture-independent. Reproduced on an NVIDIA L40S (sm_89), CUDA 12.

### Problem description

Casting a compile-time `float32` constant whose magnitude is above the fp16 *maximum finite value* (65504) but still finite as fp32 aborts compilation with an internal check, e.g. `T.cast(T.float32(65510.0), "float16")` raises `Check failed: value <= support::kMaxFloat16 (65510 vs. 65504)`. The identical cast performed at runtime on a device value compiles and returns the correctly-rounded result — `65510.0 → 65504.0` (the fp16 max), and any value `≥ 65520.0 → inf`, both matching `torch`/`numpy`. So a legal fp32→fp16 conversion crashes the compiler on the constant-folding path while succeeding on the runtime path.

The affected region is every *finite fp32* constant above `65504`: `(65504, 65520)` rounds to the finite fp16 max `65504`, and `[65520, ∞)` (still finite as fp32) rounds to `inf` — both are defined fp16 conversions that the compiler refuses to fold. The only value in that magnitude range that compiles is an actual `float('inf')` literal, because the check is now skipped for non-finite inputs (`!std::isinf`). This is a hard compile abort (`ValueError` from the FFI check), not a silent miscompile.

The same `FloatImm` check governs the *other narrow float dtypes* and *other constant sinks*, all verified below: fp8 (`e5m2` max `57344`, `e4m3`/`e4m3fn` max `448`) and fp4 (`e2m1fn` max `6`) abort identically on a finite fp32 constant above their max; the sink need not be a `T.cast` — a direct `T.float16(65510.0)` construction or a `T.fill(buf, T.float16(65510.0))` aborts at the same line. The negative side is symmetric (`T.cast(T.float32(-65510.0),"float16")` trips the `GE` check against `-65504`). `bfloat16` is the one narrow dtype with *no* analogous edge, because its max (`≈3.895e38`) exceeds `float32::max` (`≈3.403e38`), so a finite fp32 constant can never reach the bf16 bound (it overflows the fp32 literal check first).

### Reproducible example code

```python
import os, tempfile
os.environ["TILELANG_CACHE_DIR"] = tempfile.mkdtemp(prefix="tl_")
import tilelang, tilelang.language as T, torch

C = 65510.0 # in (65504, 65520): IEEE round-to-nearest maps it to the finite fp16 max 65504.0

# CONTROL: cast the SAME value from a device tensor at runtime
@tilelang.jit(out_idx=[1])
def krun(n):
@T.prim_func
def main(A: T.Tensor((n,), "float32"), B: T.Tensor((n,), "float16")):
with T.Kernel(1, threads=n) as bx:
for i in T.Parallel(n):
B[i] = T.cast(A[i], "float16")
return main

A = torch.full((2,), C, dtype=torch.float32, device="cuda")
ref = torch.full((2,), C, dtype=torch.float32).to(torch.float16)[0].item() # torch RNE -> 65504.0
print("CONTROL runtime cast :", krun(2)(A)[0].item(), " ref", ref) # -> 65504.0 ref 65504.0

# TRIGGER: cast the SAME value as a compile-time float32 CONSTANT (const-folded)
@tilelang.jit(out_idx=[0])
def kfold(n):
@T.prim_func
def main(B: T.Tensor((n,), "float16")):
with T.Kernel(1, threads=n) as bx:
for i in T.Parallel(n):
B[i] = T.cast(T.float32(C), "float16")
return main

print("TRIGGER const-fold :", kfold(2)()[0].item()) # -> InternalError before this prints
```

Boundary: runtime cast is correct across the whole region; const-fold aborts on every value > 65504

| fp32 constant | torch/numpy fp16 (RNE) | runtime `T.cast` (control) | const-fold `T.cast(T.float32(C),"float16")` |
|---|---|---|---|
| 65505.0 | 65504.0 | 65504.0 | `Check failed: value <= support::kMaxFloat16 (65505 vs. 65504)` |
| 65510.0 | 65504.0 | 65504.0 | `... (65510 vs. 65504)` |
| 65519.0 | 65504.0 | 65504.0 | `... (65519 vs. 65504)` |
| 65520.0 | inf | inf | `... (65520 vs. 65504)` |
| 66000.0 | inf | inf | `... (66000 vs. 65504)` |
| 1e30 | inf | inf | `... (1e+30 vs. 65504)` |
| `float('inf')` | inf | inf | **compiles → inf** (non-finite input skips the check) |

### Traceback

```
ValueError: Check failed: value <= support::kMaxFloat16 (65510 vs. 65504) : Literal value 65510 exceeds maximum of float16
File ".../3rdparty/tvm/src/ir/expr.cc", line 108, in tvm::FloatImm::FloatImm(tvm::DataType, double, tvm::Span)
<- tvm::tirx::MakeConstScalar(...) <- tvm::tirx::make_const(...) <- tvm::cast(...)
<- T.cast(T.float32(65510.0), "float16")
```

(The `TVM_FFI_CHECK_LE(..., ValueError)` macro raises a `ValueError`, surfaced through the FFI as a hard compile-time abort; it is not caught or rounded.)

### Expected behavior

The constant-fold path should perform the same IEEE-754 round-to-nearest conversion the runtime cast already performs: fold a constant in `(65504, 65520)` to the fp16 max `65504.0`, and a constant `≥ 65520` to `inf` — matching `torch`/`numpy` and the compiler's own device cast. A finite fp32 value should not make a `float16` cast uncompilable when the runtime path converts it without error.

### Additional context

**Root cause.** The `FloatImm` constructor treats the fp16 *maximum finite value* as the *maximum admissible* literal, so it rejects any finite value above it instead of allowing round-to-nearest to reach `65504` or `inf`. In [`src/ir/expr.cc`](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/src/ir/expr.cc#L99-L109), the range check is entered only for finite values (`if (!std::isinf(value) && !std::isnan(value))`, L99); the `is_float16()` branch (L105) then does `TVM_FFI_CHECK_LE(value, support::kMaxFloat16)` (L108) against [`kMaxFloat16 = 65504.0`](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/src/support/limits.h#L32). Any *finite* constant in `(65504, ∞)` fed to the fp16 `FloatImm` — as happens when `T.cast(, "float16")` is constant-folded — trips the check and aborts. The runtime cast lowers to a device `float2half` conversion that is unaffected, which is why the two paths disagree. (An already-`inf` literal skips the check via the `!std::isinf` guard and compiles.)

**Suggested fix.** In the `FloatImm` constructor (`src/ir/expr.cc`), round the incoming finite value into the target dtype's representable set rather than erroring. The same fix covers every narrow-float branch (fp16 / fp8 / fp4), since all use the identical *check-max-before-round* shape: map a finite value in `(dtype_max, overflow_threshold)` to `dtype_max` and a finite value `≥ overflow_threshold` to `inf` (or the dtype's saturating max for the `fn`/no-inf formats), matching the device cast and the already-accepted `inf` literal. (bf16 already can't reach its bound from an fp32 constant, so it needs no change but shares the same corrected logic.)

**Provenance.** The fp16 range check in `FloatImm` is inherited from the Apache TVM base the TileLang fork tracks; on the 0.1.13-bundled tvm (`8df8ebd`) it uses the `TVM_FFI_CHECK_LE` macro under the `!isinf && !isnan` guard. Exact introducing commit not verified — it predates the fork's own history for this file. Not a regression: the same abort occurs on 0.1.13 (and on 0.1.9, the earliest version tested).

**Dedup.** I searched the open and closed tracker and found no existing report of the fp16 const-fold *cast* defect. #2982 (`T.fill` of a boundary float32/float16 literal aborting on the same `FloatImm` range check) is the same root mechanism — a `FloatImm` built from a value not first rounded into the target dtype's representable set — at a **different sink** (`T.fill` materialization vs. constant-folded `T.cast`). Not merged: distinct trigger and repro path; see the generalization record below.

**Reach.** The trigger is a `float16` cast (or direct construction) of an fp32 constant in `(65504, 65520)` — an in-contract, round-to-nearest-defined conversion (fp32→fp16 casting is a core documented op, with the runtime path as the reference behavior). fp16 constants near the type's maximum arise wherever a kernel materializes a large fp16 literal (a saturation/clip ceiling, an initializer, a max-value sentinel).

*Shipped-example check (run on 0.1.13).* `git grep` over `examples/` and `testing/` in the v0.1.13 tree finds **no** site that casts or constructs a constant `> 65504` at fp16. The near-neighbours all dodge the region for a concrete reason: the large-magnitude literal casts in the flash-attention backward examples (`examples/flash_attention/example_gqa_bwd_tma_reduce.py:50`, `T.fill(scores_max, T.cast(-1e30, accum_dtype))`) target `accum_dtype = T.float32` (verified: `-1e30` fits fp32, so it builds a fp32 `FloatImm`, not fp16); the fp16 literal in `examples/dequantize_gemm/example_dequant_gemm_fine_grained.py:325` is `T.float16(0)` (in range); and every fp16 *cast* in the examples (`examples/deepseek_v32/topk_selector.py:11` `T.cast(x, T.float16)`, the `acc_s.to(torch.float16)` reductions) casts a **runtime** value, which lowers to the unaffected device `float2half` path — this is why CI is green. I confirmed the runtime path directly: `T.cast(A[i],"float16")` on a device tensor of `65510.0` returns `65504.0` (matches torch), and `65520.0`/`66000.0` return `inf` (match torch), while the const-fold path aborts on all three (table above).

**Impact.** The trigger is narrow: it needs a compile-time float *constant* (not a runtime value) whose magnitude falls above the target narrow dtype's max and reaches `FloatImm` via cast const-fold / direct construction / `T.fill` / reducer init. When it fires it is a hard compile-time abort (`ValueError` surfaced through the FFI), so it is caught immediately, blocks that one kernel from building, and cannot silently corrupt a result or reach production — no wrong value is ever produced. One fix in the `FloatImm` constructor closes the whole class (fp16 + fp8 e5m2/e4m3/e4m3fn + fp4 e2m1fn, both signs), removing a compile-vs-runtime disagreement on documented, round-to-nearest-defined narrowing conversions.

§17 Generalization record

**Two-level root.**
- **SOURCE-level root:** `tvm::FloatImm::FloatImm(DataType, double, Span)` in the bundled tvm [`src/ir/expr.cc` L99–L109](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/src/ir/expr.cc#L99-L109). Under the `if (!isinf && !isnan)` guard (L99), every narrow-float branch does `TVM_FFI_CHECK_LE(value, )` (fp16 L108, bf16 L113, fp8 L161, fp6 L169, fp4 L176) — a range check against the dtype's *exact finite max*, performed **before** the value is rounded into the dtype. Any finite value in `(dtype_max, ∞)` is rejected instead of allowed to round to the dtype max / `inf`. This is the fragile code shape.
- **OPERATOR-level root:** any front-end path that materialises a compile-time float constant at a narrow dtype builds this `FloatImm`. Confirmed sinks: `T.cast(, narrow)` (const-fold), direct `T.(const)` construction, `T.fill(buf, T.(const))`, and the `comm_reducer` init literal. The *runtime* cast path (`T.cast(, narrow)`) does **not** go through `FloatImm`; it lowers to a device conversion and is unaffected — this is the compile-vs-runtime disagreement.

**Class hypothesis (confirmed):** the **"check-max-before-round"** family — a `FloatImm` is range-checked against the target dtype's exact bounds *before* the value is rounded into that dtype's representable set, so a value that would legally round to the dtype max / to `inf` is rejected. Same root as #2982 (`T.fill` boundary literal) at a different sink.

**4-axis sweep** (each cell = fresh `/tmp/gen_29_*` dir, fresh cache, one kernel/process, `/tmp/tl113_venv` = 0.1.13):

| axis | cell tested | observed result | same-root? |
|---|---|---|---|
| found-first | `T.cast(T.float32(65510.0),"float16")` | CRASH `value <= support::kMaxFloat16 (65510 vs. 65504)`; runtime cast → `65504.0` | — |
| similar-logic | negative side: `T.cast(T.float32(-65510.0),"float16")` | CRASH `value >= -support::kMaxFloat16 (-65510 vs. -65504)`; torch → `-65504.0` | **same root** (mirror `GE` check, same line pair) |
| related-source | `≥65520` (→inf) rows: `65520.0`, `66000.0`, `1e30` cast to fp16 | all CRASH (`(65520 vs. 65504)` … `(1e+30 vs. 65504)`); runtime cast of `65520`/`66000` → `inf` (= torch) | **same root** — finite value hits LE check before rounding to inf |
| related-operator | direct `T.float16(65510.0)` (no cast) | CRASH at same `expr.cc:108`; this is the **minimal** form — bug is not cast-specific | **same root, broader** |
| related-operator | `T.fill(buf, T.float16(65510.0))` | CRASH `(65510 vs. 65504)` at same line | **same root** = #2982 sink |
| related-type | `T.cast(T.float32(60000.0),"float8_e5m2")` (max 57344) | CRASH `value <= bound (60000 vs. 57344)` | **same root** — fp8 branch, reachable from fp32 |
| related-type | `T.cast(T.float32(500.0),"float8_e4m3"/"float8_e4m3fn")` (max 448) | CRASH `(500 vs. 448)` | **same root** — fp8 branch |
| related-type | `T.cast(T.float32(10.0),"float4_e2m1fn")` (max 6) | CRASH `(10 vs. 6)` | **same root** — fp4 branch |
| related-type | `T.cast(T.float32(3.0e38),"bfloat16")` | **compiles → `3.004e38`** (in range) | n/a |
| related-type | `T.cast(T.float32(3.5e38),"bfloat16")` | CRASH but on the **fp32** check `(3.5e+38 vs. 3.40282e+38)` — overflows fp32 literal first | **NOT reachable** — bf16 max `≈3.895e38 > float32::max`, no analogous edge via fp32-const cast |
| boundary | `T.cast(T.float32(float('inf')),"float16")` | **compiles → `inf`** — `!isinf` guard skips the check | distinct (marks the exact defect boundary: only *finite* > max crashes) |

**Reframe.** The defect is a **class** across the narrow-float branches of `FloatImm` (fp16 + fp8 e5m2/e4m3/e4m3fn + fp4 e2m1fn) and across the constant sinks (cast / direct construction / `T.fill` / reducer init), all one fix in one function. Title/Problem widened to the class; fp16 kept as the lead example (most common in practice). bf16 is explicitly out-of-class (bound unreachable from fp32). The negative side and the `≥ overflow_threshold`(→inf) sub-region are the same root, not separate bugs. #2982 is the same root at the `T.fill` sink — **not merged** (distinct trigger/repro), noted as the same class here.

**No new distinct bug found while sweeping.** All crashing cells trace to the single `check-max-before-round` root; the two clean cells (bf16 in-range, `inf` literal) are correct behaviour, not new defects.

Start with the FloatImm constructor and narrow-float range checks in src/ir/expr.cc, then run the provided Python reproducer to compare constant folding with the runtime cast. Check the fp16, fp8, and fp4 cases described in the issue. Done means finite constants no longer abort and constant-folded results match the runtime conversion behavior.

Written by the indexing model from the issue text.