[BUG][Fuzzer][wrong-code] `T.ceil`/`T.floor`/`T.trunc` of a narrow-float literal (`float16`/`bfloat16`/`float32`) silently folds to the wrong integer instead of the runtime result

### Required prerequisites

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

### What version of TileLang are you using?

0.1.13 (latest release)

### System information

TileLang 0.1.13; CUDA 13.0 (nvcc V13.0.88); PyTorch 2.13.0+cu130; NVIDIA L40S (sm_89). The defect is a compile-time constant fold and is target-independent.

### Problem description

`T.ceil` (and `T.floor`, `T.trunc`) applied to a compile-time narrow-float literal folds to a different integer than the same op computes at runtime on the same value. `T.ceil(T.float32(8388608.5))` folds to `8388609.0`, but the runtime device `T.ceil` on a buffer holding that same `float32` value returns `8388608.0` — and `8388608.0` is correct: `float32(8388608.5)` rounds (round-half-to-even) to `8388608.0`, whose ceiling is itself. The result is silently wrong; the kernel compiles and runs with no error.

The literal `8388608.5` is representable only in `double`; assigned to a `float32` it becomes `8388608.0`. The fold rounds the wrong way because it applies `ceil` to the un-narrowed `double` value the `FloatImm` still carries, not the narrowed value the operand actually holds.

**This is not specific to `float32`, nor to `ceil`.** The same fold error hits every float dtype narrower than `double` — `float16` and `bfloat16` too — because each carries the un-narrowed `double` in its `FloatImm`. The affected operation depends on which value the narrowing produces:

- **`.5`-type literal** (rounds to an integer under narrowing, e.g. `float16(1024.5) = 1024.0`): only `ceil` misfolds (it rounds the phantom `.5` up); `floor`/`trunc` happen to agree.
- **odd-integer-past-precision literal** (rounds to an even integer under narrowing, e.g. `float16(2049.0) = 2048.0`, `bfloat16(257.0) = 256.0`): **all three of `ceil`/`floor`/`trunc` misfold** — they operate on `2049`/`257` and return `2049`/`257`, while the runtime operand is `2048`/`256`.

Measured folds vs the narrowed value's correct result:

```
dtype op literal narrowed fold correct
float32 ceil 8388608.5 8388608.0 8388609 8388608 FAIL
float16 ceil 1024.5 1024.0 1025 1024 FAIL
bfloat16 ceil 128.5 128.0 129 128 FAIL
float16 ceil 2049.0 2048.0 2049 2048 FAIL (floor/trunc also FAIL: 2049 vs 2048)
bfloat16 ceil 257.0 256.0 257 256 FAIL (floor/trunc also FAIL: 257 vs 256)
```

Three faces, and non-boundary controls that pass

Same kernel/value, folded (compile-time) vs runtime (device), against `numpy` on the `float32` value:

```
ceil (8388608.5): float32 = 8388608.0 folded = 8388609.0 runtime = 8388608.0 numpy = 8388608.0 -> FAIL
floor(8388609.5): float32 = 8388610.0 folded = 8388609.0 runtime = 8388610.0 numpy = 8388610.0 -> FAIL
trunc(8388609.5): float32 = 8388610.0 folded = 8388609.0 runtime = 8388610.0 numpy = 8388610.0 -> FAIL
ceil (3.2): float32 = 3.2000000 folded = 4.0 runtime = 4.0 numpy = 4.0 -> PASS
floor(3.7): float32 = 3.7000000 folded = 3.0 runtime = 3.0 numpy = 3.0 -> PASS
```

The runtime device path is correct in every case; only the compile-time fold disagrees, and only for literals whose `double`→`float32` narrowing crosses an integer boundary that the op would otherwise not cross.

Not a regression — see Provenance.

### Reproducible example code

```python
import numpy as np, torch
import tilelang, tilelang.language as T

N = 8
LIT = 8388608.5 # 2^23 + 0.5; float32(LIT) rounds (half-to-even) to 8388608.0

# Const-fold path: T.ceil applied to a compile-time float32 literal
@T.prim_func
def fold(C: T.Tensor((N,), "float32")):
with T.Kernel(1, threads=N) as bx:
for i in T.Parallel(N):
C[i] = T.ceil(T.float32(LIT))

kf = tilelang.compile(fold, target="cuda")
cf = torch.empty(N, dtype=torch.float32, device="cuda"); kf(cf)
folded = cf[0].item()

# Runtime control: same op, same value, arriving through a buffer (holds float32(LIT))
@T.prim_func
def rt(A: T.Tensor((N,), "float32"), C: T.Tensor((N,), "float32")):
with T.Kernel(1, threads=N) as bx:
for i in T.Parallel(N):
C[i] = T.ceil(A[i])

kr = tilelang.compile(rt, target="cuda")
a = torch.full((N,), LIT, dtype=torch.float32, device="cuda") # == float32(8388608.5) == 8388608.0
cr = torch.empty(N, dtype=torch.float32, device="cuda"); kr(a, cr)
runtime = cr[0].item()

ref = float(np.ceil(np.float32(LIT))) # numpy on the float32 value
print("const-folded T.ceil :", folded) # -> 8388609.0
print("runtime device T.ceil:", runtime) # -> 8388608.0
print("numpy ceil(float32) :", ref) # -> 8388608.0
print("ORACLE:", "PASS" if folded == ref else f"FAIL (folded {folded} != {ref}; runtime gives {runtime})")
# -> ORACLE: FAIL (folded 8388609.0 != 8388608.0; runtime gives 8388608.0)
```

### Traceback

No traceback — the kernel compiles and runs to completion; the folded constant is silently wrong and deterministic.

### Expected behavior

`T.ceil`/`T.floor`/`T.trunc` of a narrow-float literal should fold to the value the operation produces on the operand in its **declared** dtype — i.e. the same value the runtime `tir.Call` path (and IEEE-754 / `numpy` on that narrowed value) produce. Rounding the literal into its declared dtype (`float16`/`bfloat16`/`float32`) before applying `ceil`/`floor`/`trunc` makes the compile-time and runtime paths agree.

### Additional context

**Root cause.** The `FloatImm` const-fold for `floor`/`ceil`/`trunc` rounds the *un-narrowed* stored value rather than the value in the expression's declared dtype. `T.float32(8388608.5)` produces a `FloatImm` whose `dtype` is `float32` but whose `value` field still holds the `double` `8388608.5`; the fold calls `std::ceil`/`std::floor` on that `double` and re-wraps the result as `float32`, so it sees `8388608.5` where the device only ever sees `float32(8388608.5) == 8388608.0`.

Mechanism and source

- [`floor`, `op.cc:1039`](https://github.com/TileLang/tvm/blob/0e15b274bce8b46f971abf5ac390e844aa6acee5/src/tirx/op/op.cc#L1039): `if (fx) return FloatImm(x.dtype(), std::floor(fx->value), fx->span);`
- [`ceil`, `op.cc:1053`](https://github.com/TileLang/tvm/blob/0e15b274bce8b46f971abf5ac390e844aa6acee5/src/tirx/op/op.cc#L1053): `if (fx) return FloatImm(x.dtype(), std::ceil(fx->value), fx->span);`
- [`trunc`, `op.cc:1096`](https://github.com/TileLang/tvm/blob/0e15b274bce8b46f971abf5ac390e844aa6acee5/src/tirx/op/op.cc#L1096): `FloatImm(x.dtype(), (fx->value < 0 ? std::ceil(fx->value) : std::floor(fx->value)), ...)`

`fx->value` is a `double`; `x.dtype()` is `float32`. Rounding the `double` and stamping it `float32` is where the paths diverge — narrowing `fx->value` to `x.dtype()` before the `ceil`/`floor` (when the dtype is narrower than `double`) would make the fold match the runtime and `numpy`. The un-narrowed value is confirmed directly: `T.float32(8388608.5).value == 8388608.5`.

The runtime `tir.Call` path (`floorf`/`ceilf`/`truncf` via the CUDA intrinsic lowering) is unaffected — it operates on the already-narrowed `float32`, which is why it is correct and serves as the control above.

**Suggested fix.** In the `floor`/`ceil`/`trunc` `FloatImm` fold (`op.cc:1039/1044/1087`), round `fx->value` into `x.dtype()`'s precision before applying `std::ceil`/`std::floor`/`std::trunc` (for a `float32` dtype, cast through `float`), so the fold matches the value the operand holds at runtime. The fix is in vendored C++ (not verified end-to-end).

**Provenance.** Vendored from apache/tvm via the `TileLang/tvm` submodule (pinned at `0e15b274` for the 0.1.9 tag); the `floor`/`ceil`/`trunc` `FloatImm` fold has computed on the un-narrowed value since these functions existed upstream. Not a TileLang regression. Same vendored-const-fold lineage as the previously reported float `x*0→0` inf/NaN drop (#2638), the `max`/`min` NaN-order const-fold (#2882), and the `uint64` const-fold crash (#2582), but a distinct mechanism (double-vs-declared-dtype rounding of `floor`/`ceil`/`trunc`).

**Dedup.** I searched the open and closed tracker and found no existing report of this defect. It is unrelated to #2565 (fp16/bf16 `T.floor` compile crash via the `hfloor` dtype dispatch) — that is a compile failure on a narrow dtype; this is a silent wrong value on `float32` literals through the constant folder.

**Reach.** `T.ceil`/`T.floor`/`T.trunc` are documented, exported ops in real use: they appear in one example (`examples/dequantize_gemm/quantize/nvfp4.py:361`, `T.ceil(x[0] * 512.0)`) and in shipped tests (`testing/python/math/test_math_fast_math.py:288-290` and `testing/python/fastmath/test_mathops_fastmath.py:274-276`). The bug needs the operand to be a *compile-time float literal* (a constant expression the folder reaches) *and* that literal to sit at a `double`→narrow-dtype narrowing boundary that flips the rounded integer — near `2^23` for `float32`, `2^10` for `float16`, `2^7` for `bfloat16` (the smaller mantissas make `float16`/`bfloat16` far *easier* to trip, e.g. `float16(2049.0)`); a runtime operand takes the `tir.Call` codegen path, which is correct. Those tests apply the ops to random runtime buffer values, and the `nvfp4` example's argument is runtime, so all shipped sites take the runtime path and none trip the fold — that is why CI is green.

**Impact.** Low-consequence: the trigger is narrow (a compile-time constant on a narrowing boundary; runtime
values take the correct path), and even when it fires the result is one constant off by one, computed
deterministically — no data corruption, no crash, and, being a literal, wrong the same way every run rather
than poisoning a workload mid-computation. The value of fixing it is removing a compile-time/runtime
disagreement that should not exist, and closing the whole boundary class (all three ops, all narrow float
dtypes) at once.

Read the FloatImm folding paths in src/tirx/op/op.cc around lines 1039, 1053, and 1096, then inspect the existing math tests in testing/python/math/test_math_fast_math.py and testing/python/fastmath/test_mathops_fastmath.py. Reproduce the listed narrow-float boundary cases and add coverage for ceil, floor, and trunc across the affected dtypes. Done means compile-time results agree with runtime results and the regression tests pass.

Written by the indexing model from the issue text.