[BUG][Fuzzer][ice-on-valid-code] Casting an out-of-range integer constant to a narrow int crashes the compiler instead of wrapping
- Dominant language
- Python
- Stars
- 7.4k
- Forks
- 742
- Avg merge
- 1d 2h
- Merged PRs (30d)
- 108
Description
### Required prerequisites
- [x] I have read the documentation .
- [x] I have searched the [Issue Tracker](https://github.com/tile-ai/tilelang/issues) that this hasn't already been reported. (comment there if it has.)
### What version of TileLang are you using?
0.1.13 (tag `v0.1.13` = `8001cc4ccf6149382d2019654a19f59c1d4d0482`, bundling the `tile-ai/tvm` submodule at `8df8ebd61659505dd7005f4820b030e983e2c005`).
### System information
- NVIDIA L40S (`sm_89`), CUDA 12.8, Python 3.13, torch 2.8, tilelang 0.1.13.
- **Architecture-independent — no GPU required.** The crash fires host-side while the arithmetic simplifier folds the store's value expression at compile time, before any target codegen. It was reproduced with `target="cuda"`, but the abort happens while lowering the `@T.prim_func`, before a device is touched.
### Problem description
Writing a compile-time integer constant that does not fit the destination's narrow integer type crashes compilation with an internal `Check failed` abort. `T.Cast("int8", 300)` aborts with `Check failed: value < 1LL << (dtype.bits() - 1) (300 vs. 128) : Literal value 300 exceeds maximum of int8`, even though narrowing `300` to `int8` is a defined conversion that wraps to `44` (`300 - 256`). The same narrowing computed at runtime returns `44` — only the compile-time-constant path crashes.
The trigger is a constant whose value exceeds the target's range, cast to any integer type narrower than 64 bits (int8/int16/int32/uint*; only 64-bit types skip the range check):
Trigger / control matrix (run this session, L40S, tilelang 0.1.13)
| cast | expected (C narrowing / numpy) | outcome |
|---|---|---|
| `T.Cast("int8", 300)` | `44` | **CRASH** `value < 1LL << (dtype.bits() - 1) (300 vs. 128) : Literal value 300 exceeds maximum of int8` |
| `T.Cast("int8", 200)` | `-56` | **CRASH** `(200 vs. 128) : ... exceeds maximum of int8` |
| `T.Cast("int8", -200)` | `56` | **CRASH** `value >= -(1LL << (dtype.bits() - 1)) (-200 vs. -128) : ... exceeds minimum of int8` (GE check) |
| `T.Cast("int8", 128)` | `-128` | **CRASH** `(128 vs. 128) : ... exceeds maximum of int8` (128 is one past `int8` max) |
| `T.Cast("int16", 70000)` | `4464` | **CRASH** `(70000 vs. 32768) : ... exceeds maximum of int16` |
| `T.Cast("uint8", 300)` | `44` | **CRASH** `value < 1LL << dtype.bits() (300 vs. 256) : ... exceeds maximum of uint8` (uint LT check) |
| `T.Cast("int32", 3000000000)` | `-1294967296` | **CRASH** `(3000000000 vs. 2147483648) : ... exceeds maximum of int32` — int32 is NOT wide enough to be safe |
| `T.Cast("int64", 5000000000)` (bits==64) | `5000000000` | OK — 64-bit types skip the range check, so no crash |
| `T.Cast("int8", 44)` (in range) | `44` | OK — compiles and returns `44` |
| runtime `int32 -> int8` of `300` (value from a buffer) | `44` | OK — wraps, returns `44` |
The runtime narrowing cast and the in-range constant both work; only an **out-of-range constant** cast to an integer type narrower than 64 bits aborts. This is not a regression — the fold shortcut is long-standing.
### Reproducible example code
```python
import torch, numpy as np, tilelang, tilelang.language as T
N = 8
# (A) BUG: a compile-time integer constant that overflows the narrow target type
# crashes at compile time. 300 as int8 is a defined narrowing conversion
# (300 - 256 = 44), yet the fold aborts.
@tilelang.jit(out_idx=[0])
def const_cast(dd, val):
@T.prim_func
def main(B: T.Tensor((N,), dd)):
with T.Kernel(1, threads=1):
for i in T.serial(N):
B[i] = T.Cast(dd, val) # val folded at compile time
return main
# (B) CONTROL: the SAME narrowing, value supplied at runtime -> computes 44 (wraps).
@tilelang.jit(out_idx=[1])
def runtime_cast(sd, dd):
@T.prim_func
def main(A: T.Tensor((N,), sd), B: T.Tensor((N,), dd)):
with T.Kernel(1, threads=1):
for i in T.serial(N):
B[i] = T.Cast(dd, A[i])
return main
for dd, val in [("int8", 300), ("int16", 70000), ("uint8", 300)]:
ref = int(np.array([val]).astype({"int8":np.int8,"int16":np.int16,"uint8":np.uint8}[dd])[0])
try:
got = int(const_cast(dd, val)()[0].item())
print(f"{dd} <- {val}: got={got} expected={ref}")
except Exception as e:
print(f"{dd} <- {val}: CRASH -> {str(e).splitlines()[-1][:70]} (expected {ref})")
# int8 <- 300 : CRASH -> ... 300 exceeds maximum of int8 (expected 44)
# int16 <- 70000: CRASH -> ... 70000 exceeds maximum of int16 (expected 4464)
# uint8 <- 300 : CRASH -> ... 300 exceeds maximum of uint8 (expected 44)
a = torch.tensor([300, 208, 200, 128, 44, -1, 127, 255], dtype=torch.int32, device="cuda")
b = runtime_cast("int32", "int8")(a).cpu().numpy()
print("runtime int32->int8:", b, "(wraps: 300 -> 44)") # -> [44 -48 -56 -128 44 -1 127 -1]
```
### Traceback
```
tvm.error.InternalError: Check failed: value < 1LL << (dtype.bits() - 1) (300 vs. 128) :
Literal value 300 exceeds maximum of int8
...
in tvm::arith::RewriteSimplifier::Impl::VisitExpr_(tvm::tir::CastNode const*)
in tvm::cast(tvm::runtime::DataType const&, tvm::PrimExpr, tvm::Span)
in tvm::PrimExpr tvm::tir::MakeConstScalar(tvm::runtime::DataType, long, tvm::Span)
in tvm::IntImm::IntImm(tvm::runtime::DataType, long, tvm::Span)
```
### Expected behavior
Casting an integer constant to a narrower integer type should produce the same value the runtime cast produces — the low `bits()` of the source, interpreted in the target type (`300 -> int8 == 44`, matching C 2's-complement narrowing and `numpy.int8(300)`). The runtime path already computes exactly this over the same inputs, so the constant path returning a different result — a hard compiler abort — is the inconsistency.
### Additional context
**Root cause.** The arithmetic simplifier's constant-cast shortcut narrows an integer constant by re-using the source value unmasked, so it hands an out-of-range value to `IntImm`, whose range assert aborts. When `tvm::cast` sees a `Cast` of an `IntImm`, it folds via [`make_const(t, op->value, op->span)`](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/src/tirx/op/op.cc#L445) ([comment "const fold IntImm as they are used in index computations"](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/src/tirx/op/op.cc#L442)). That routes to [`MakeConstScalar`](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/include/tvm/tirx/op.h#L978), which for an int/bool type does `IntImm(t, static_cast(value), span)` with no truncation to `t.bits()`. `IntImm`'s constructor then hits its [range check](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/src/ir/expr.cc#L73) and aborts. A wrapping `value & ((1 << bits) - 1)` (sign-extended for signed types) before constructing the `IntImm` — the operation the emitted device code already performs — would fold to the correct in-range constant.
**Suggested fix.** In the integer branch of `MakeConstScalar` (or in the `cast`/`make_const` fold shortcut), truncate the value to the destination width and sign/zero-extend per signedness before building the `IntImm`, matching the runtime narrowing. That keeps the constant fold and returns the same value the non-folded path computes, instead of aborting.
**Provenance.** The `tvm::cast` IntImm const-fold shortcut is long-standing upstream TVM code inherited by the `tile-ai/tvm` fork; it is present at the v0.1.13 submodule pin (`8df8ebd6`) with the "const fold IntImm as they are used in index computations" comment. Origin predates the fork; the exact upstream commit is unverified.
**Dedup.** I searched the open and closed tracker (keywords "constant fold", "cast", "exceeds maximum", "narrowing", "IntImm") and found no report of this narrowing-cast crash. It shares its general mechanism with the already-filed #2982 — that report is the same "range-check on an un-narrowed constant" defect, but on the **FloatImm** branch reached via `T.fill(float32, ...)` (`expr.cc:100-107`, the float bounds checked before rounding). This one is the **IntImm** branch reached via `T.Cast` to a narrow int (`expr.cc:71-73`, the int bounds checked before truncation), fed by a different sink (`tvm::cast` -> `MakeConstScalar`, which does no mask step). Same root family, distinct sink — do NOT merge. Also distinct from #2582, a `uint64` `Add`/`Mul` fold whose *result* lands in `[2^63, 2^64)`: that path (`GetFoldResultInt64Repr`) masks to `bits()` for narrow types, so it is a separate mechanism.
**Reach.** The trigger is typed and documented (`T.Cast` accepts any dtype and value pair, and the runtime path defines the result), and it is natural usage: writing a constant clamp/bias/fill into a narrow-int buffer is the natural way to express a quantization step; 32 example files under `examples/` use `int8`/`int16` dtypes. The **fragile narrowing-cast path IS exercised by shipped examples** — the dequantize kernels build narrow-int constants through this exact fold via `tirx.const(..., T.uint8/uint16)` and `expr.astype(T.uint8/uint16)` (e.g. `examples/dequantize_gemm/example_dequant_gemm_w4a8.py`, `..._bf16_mxfp4_hopper.py`). Those examples **dodge the crash only because every constant they narrow is in range** (masks `(1<<4)-1=15`, bit-shift amounts `1..8`, bias `126` — all `< 256`). I ran the shipped `example_dequant_gemm_w4a8.py` verbatim on 0.1.13 (`--m 128 --n 256 --k 256`) and it prints `All checks pass.` — it uses the same `tirx.const(..., T.uint8)` narrowing fold this bug lives in, but every folded constant is in-range so the range check never trips. No shipped example uses an *out-of-range* narrow-int constant, which is why CI is green. The one ingredient that fires it is a constant outside the target's `[min, max]`; any narrow-int destination trips it, an in-range one does not.
**Impact.** The trigger is narrow: a compile-time integer constant, statically outside the destination integer type's `[min, max]` (a runtime value or an in-range constant does not fire). When it fires it is a hard host-side compile abort — loud, deterministic, caught immediately, and it corrupts no data and cannot slip silently into a running workload; the cost is that a valid kernel writing an out-of-range constant into a narrow-int buffer refuses to build. Fixing it removes the const-fold-vs-runtime disagreement (the folded path would then return the same wrapped value the runtime cast already produces) and closes this out-of-range narrowing boundary.
§17 Generalization record
**Two-level root.**
- **SOURCE-level (fragile implementation):** the `IntImm` constructor's range check (`src/ir/expr.cc`, the `is_uint` LT check ~L64, the signed lower-bound GE ~L71 and upper-bound LT ~L73, and the `bits()==1`/bool `value==0||value==1` branch ~L69) is fed by `tvm::cast`'s IntImm const-fold shortcut (`src/tirx/op/op.cc:445`, `make_const(t, op->value, ...)`) → `MakeConstScalar` (`include/tvm/tirx/op.h:979`, `IntImm(t, static_cast(value), span)`) which passes the source value through with **no mask/truncate to `t.bits()`**. The abort is the range check firing on a value the destination could legally hold post-narrowing.
- **OPERATOR-level (what correlates):** ANY frontend path that folds an out-of-range integer constant into a narrow-int `IntImm` — every alias of "cast/construct a constant into a narrow integer" routes through the same `cast`→`MakeConstScalar`→`IntImm` sink and shares this root.
**Class.** "Un-narrowed constant construct" family: a fold builds an `IntImm`/`FloatImm` from the source value WITHOUT masking/rounding to the destination's `bits()`. General root shared with the already-filed **#2982**, which is the **FloatImm** sibling sink (`T.fill(float32,...)`, float bounds checked before rounding); this draft is the **IntImm** sink (`cast`→`MakeConstScalar`, no mask step). Same family, distinct sink — do NOT merge. Also distinct from **#2582** (`uint64` `Add`/`Mul` fold whose *result* is masked to `bits()` in `GetFoldResultInt64Repr` — separate mechanism; confirmed below `int64<-5e9` does not fire).
**4-axis sweep (every cell RUN, fresh process + fresh cache, L40S, 0.1.13):**
| axis | cell tested | input → observed result | same-root? |
|---|---|---|---|
| found-first | `T.Cast("int8", 300)` | CRASH `value < 1LL<<(bits-1) (300 vs. 128) : exceeds maximum of int8` (wants `44`) | — (root) |
| related-operator | `tirx.const(300, "int8")` (the helper examples use) | CRASH `(300 vs. 128) : exceeds maximum of int8` | SAME |
| related-operator | `T.Cast("int32", 300).astype("int8")` (the `.astype` example idiom) | CRASH `(300 vs. 128) : exceeds maximum of int8` | SAME |
| related-operator | `T.Cast("int8", 100+250)` (constant *arithmetic* folds to 350) | CRASH `(350 vs. 128) : exceeds maximum of int8` | SAME — need not be a literal |
| related-source | (FloatImm sink, `T.fill(float32,...)`) = **#2982** | filed separately | same family, distinct sink |
| related-source | (`uint64` result-mask fold) = **#2582** | masks, no crash | distinct mechanism |
| related-operator | `T.Cast("uint8", 300)` | CRASH `value < 1LL<= -(1LL<<(bits-1)) (-200 vs. -128) : exceeds minimum of int8` (wants `56`) | SAME, signed lower-bound GE (expr.cc~L71) |
| related-type (boundary) | `T.Cast("int8", 128)` | CRASH `(128 vs. 128)` (wants `-128`) | SAME — max is exclusive |
| related-type (width) | `T.Cast("int16", 70000)` | CRASH `(70000 vs. 32768) : exceeds maximum of int16` (wants `4464`) | SAME |
| related-type (width) | `T.Cast("int32", 3000000000)` | CRASH `(3000000000 vs. 2147483648) : exceeds maximum of int32` (wants `-1294967296`) | SAME — int32 NOT safe |
| related-type (width) | `T.Cast("int64", 5000000000)` | OK, returns `5000000000` | DISTINCT boundary — `bits()==64` skips the range check |
| control | `T.Cast("int8", 44)` (in range) | OK, returns `44` | — proves boundary |
| control | runtime `int32→int8` of `300` (buffer value) | OK, returns `44` (wraps) | — the non-fold path is correct |
**What the sweep establishes:**
- The class is **"any integer type narrower than 64 bits, out-of-range constant, via any cast/const-construct idiom"** — not "sub-int32" (int32 fires) and not "only `T.Cast`" (`tirx.const`, `.astype`, and constant *arithmetic* all fire the same sink).
- `bool` (`bits()==1`) hits a *different* assert inside the same `IntImm` ctor — same root, one more sub-check, so it belongs to the class.
- 64-bit is the sole safe integer width (range check gated on `bits() < 64`).
**PART 1 — shipped-example run.** No `examples/` file spells `T.Cast("int8"|"int16"|"uint8"` literally (`grep`: 0 sites), but the narrowing-cast fold **is** exercised by the dequantize examples via `tirx.const(..., T.uint8/uint16)` and `.astype(T.uint8/uint16)`. I ran `examples/dequantize_gemm/example_dequant_gemm_w4a8.py --m 128 --n 256 --k 256` verbatim on 0.1.13: it prints `All checks pass.` It **dodges** the crash because every constant it narrows is in range (`(1<<4)-1=15`, shift amounts, `T.int8` reinterpret masks — all `< 256`). So a shipped example touches the exact fragile path and passes only by never handing it an out-of-range value. The bug is real (all matrix cells above reproduce), the example does not trigger it, and that is why CI is green.
Contributor guide
Research direction
Start at src/tirx/op/op.cc around the IntImm constant-fold shortcut, then follow MakeConstScalar in include/tvm/tirx/op.h and the IntImm checks in src/ir/expr.cc. Run the reproducible T.Cast cases, including the runtime control, and verify that out-of-range narrow integer constants compile and match wrapped results without aborting while in-range behavior remains intact.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp, python
- Domain
- compilers
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Quiet
- Clarity
- Clearly specified
- Newbie friendliness
- 62/100