[BUG][Fuzzer][ice-on-valid-code] `int32` expression whose constant terms fold past `INT32_MAX` aborts compilation instead of wrapping

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

### System information

- TileLang `0.1.13` (installed wheel; the affected code is the vendored TVM `arith` simplifier, `3rdparty/tvm` pinned at commit `8df8ebd61659505dd7005f4820b030e983e2c005`).
- Python 3.13, Linux.
- GPU: NVIDIA L40S (sm_89). The defect is in the host-side IR simplifier and is architecture-independent — no GPU is needed to hit the compile-time crash.

### Problem description

A valid `int32` kernel fails to **compile** when the simplifier reassociates two (or more) compile-time-constant addends whose exact integer sum exceeds `INT32_MAX`. `B[i] = (A[i] + 2000000000) + 2000000000` aborts at `tilelang.compile()` with:

```
InternalError: Check failed: value < 1LL << (dtype.bits() - 1) (4000000000 vs. 2147483648)
: Literal value 4000000000 exceeds maximum of int32
```

The two constants (`2_000_000_000 + 2_000_000_000 = 4_000_000_000`) are combined into a single `int32` literal, which the `IntImm` constructor rejects. Under C/CUDA two's-complement `int32` semantics the sum should wrap (`4_000_000_000 mod 2^32 = -294_967_296`), so this is a legal kernel that never gets to run.

The trigger is narrow: the crash only fires when the constant fold happens **inside a sum that also has a runtime term** (the `+ A[i]`). Two controls bound it (both run below):

- **Runtime addend, no fold → compiles + wraps correctly.** Replacing the second constant with a runtime buffer value (`(A[i] + 2000000000) + C[i]`, `C[i]==2000000000`) never triggers the fold; it compiles and returns the correctly-wrapped `[-294967296, …]`, matching numpy `int32`.
- **Pure-constant, no runtime term → compiles.** `B[0] = T.int32(1073741824) * T.int32(2)` (product `= 2^31`) with no runtime term folds fine — the plain constant-folding path wraps, so nothing out-of-range is materialized.

So the same arithmetic compiles on either side of the boundary; only the "constant fold reached through the summation path" case crashes. The multiplication form (`T.int32(1073741824) * T.int32(2) + A[0]`, constant factor `= 2^31`) crashes identically (`2147483648 vs 2147483648`), confirming the root is the shared constant-accumulator, not a specific operator.

No traceback beyond the `InternalError` above — the failure is a compile-time `ICHECK` abort, not a runtime error.

### Reproducible example code

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

# --- Trigger: two int32 constant addends whose sum (4_000_000_000) > INT32_MAX,
# plus a runtime term A[i]. Legal int32 arithmetic; should wrap. ---
@tilelang.jit(out_idx=[1])
def make(N):
@T.prim_func
def main(A: T.Tensor((N,), "int32"), B: T.Tensor((N,), "int32")):
with T.Kernel(1, threads=N) as bx:
i = T.get_thread_binding(0)
B[i] = (A[i] + T.int32(2000000000)) + T.int32(2000000000)
return main

a = torch.tensor([0, 1, -1, 100], dtype=torch.int32, device="cuda")
try:
b = make(4)(a).cpu()
print("TRIGGER got", b.tolist()) # expected once fixed: wrapped values
except Exception as e:
line = [l for l in str(e).splitlines() if "exceeds maximum" in l]
print("TRIGGER CRASH:", (line[0] if line else str(e).splitlines()[-1]).strip())
# -> Check failed: value < 1LL << (dtype.bits()-1) (4000000000 vs. 2147483648) : Literal value 4000000000 exceeds maximum of int32

# --- Control: same arithmetic, second addend from a runtime buffer -> no fold ->
# compiles and wraps correctly (int32 two's-complement). ---
@tilelang.jit(out_idx=[2])
def make_ctl(N):
@T.prim_func
def main(A: T.Tensor((N,), "int32"), C: T.Tensor((N,), "int32"), B: T.Tensor((N,), "int32")):
with T.Kernel(1, threads=N) as bx:
i = T.get_thread_binding(0)
B[i] = (A[i] + T.int32(2000000000)) + C[i]
return main

c = torch.full((4,), 2000000000, dtype=torch.int32, device="cuda")
b = make_ctl(4)(a, c).cpu()
aa = a.cpu().numpy().astype(np.int64)
s = (aa + 2000000000 + 2000000000) % (2**32)
exp = torch.from_numpy(np.where(s >= 2**31, s - 2**32, s).astype(np.int32))
print("CONTROL got", b.tolist(), "exp", exp.tolist(),
"->", "PASS" if torch.equal(b, exp) else "FAIL")
# -> CONTROL got [-294967296, -294967295, -294967297, -294967196] exp [...] -> PASS
```

### Traceback

```
No traceback beyond the InternalError — the kernel never finishes compiling:

[Fatal] InternalError: Check failed: value < 1LL << (dtype.bits() - 1)
(4000000000 vs. 2147483648) : Literal value 4000000000
exceeds maximum of int32
```

### Expected behavior

The kernel should compile and the constant sum should be folded modulo `2^32` (`4_000_000_000 → -294_967_296`), matching the two's-complement `int32` result the hardware produces when the fold is avoided (the runtime-addend control returns exactly that). TileLang/TVM already wraps `int32` constants elsewhere: the plain constant-folder wraps the same `2^31` product without complaint (pure-constant control above), and the runtime path wraps correctly at execution. Compiling `(x + 2000000000) + 2000000000` should not depend on whether an addend happens to be a compile-time constant.

### Additional context

**Root cause.** The `arith` canonical simplifier accumulates a summation's constant part in a `int64_t` field without ever reducing it modulo the expression's `int32` dtype, then rematerializes it as an `int32` literal, which is out of range. `SumExprNode` stores the constant term in [`int64_t base`](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/src/arith/canonical_simplify.cc#L245); [`AddToSelf`](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/src/arith/canonical_simplify.cc#L296) does `base += value` and [`MulToSelf`](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/src/arith/canonical_simplify.cc#L274) does `base *= scale`, neither masked to the dtype. When the sum has a non-constant arg it goes through `Normalize_`, which emits [`res + make_const(dtype, base)`](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/src/arith/canonical_simplify.cc#L510) with `base` still `4_000_000_000` (or `2^31`); `make_const(int32, …)` builds an `IntImm` whose [range check](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/src/ir/expr.cc#L73) then aborts.

The contrast with the *scalar* constant-folder is the key. The plain folder, [`TryConstFold`](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/src/arith/const_fold.h#L81), routes its result through [`GetFoldResultInt64Repr`](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/src/arith/const_fold.h#L81) which masks to `dtype.bits()` and sign-extends — so it *wraps* and never materializes an out-of-range literal. The `SumExpr` accumulator applies no such mask. That is why the all-constant expression (no runtime arg, or a constant sub-expression isolated inside a `min`/`max` — both verified below) is fine, and why only "the fold reached through the summation path with a runtime term" crashes.

**Only signed `int32` crashes (verified — dtype boundary is sharp).** The `SumExpr` canonical path is entered **only for index-typed dtypes**, [`IsIndexType`](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/src/arith/const_fold.h#L75) `= is_int() && (bits==32 || bits==64)`. `int8`/`int16`/`uint32` are *not* index types, so they never enter the un-masked accumulator — they fold through `GetFoldResultInt64Repr` and **wrap** (verified: `int8` sum→`-56`, `int16` sum→`-25536`, `uint32` sum→`1705032704`, all compile). `int64` *is* an index type and *does* enter the accumulator, but the accumulator itself wraps in C++ `int64_t` and the `IntImm` range check is [gated on `dtype.bits() < 64`](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/src/ir/expr.cc#L70), so it never fires (verified: `int64` sum→wraps, compiles). Signed `int32` is the *only* dtype that is both an index type (enters the un-masked path) and narrow enough that the range check is armed — so it is the sole crash face.

**Suggested fix.** Apply `GetFoldResultInt64Repr`-style masking to `base` before materializing it — wrap at the `make_const(dtype, base)` sites in `Normalize`/`Normalize_` (and correspondingly in `AddToSelf`/`MulToSelf`), so an in-`int32`-range literal is emitted, consistent with the wrapping the scalar const-folder and the hardware already perform.

**Provenance.** The unmasked `int64_t base` accumulator and `make_const(dtype, base)` rematerialization are original to the TVM CanonicalSimplifier ([apache/tvm #2891](https://github.com/apache/tvm/pull/2891), 2019; relocated into the `arith/` subfolder by #4722), inherited unchanged by the TileLang TVM fork — present since the simplifier first shipped, not a regression.

**Dedup.** I searched the open and closed tracker; the nearest report is [#2560](https://github.com/tile-ai/tilelang/issues/2560), the opposite face — a `>2^31` *constant* index that is **silently** truncated to `int32` (wrong address, no crash), where adding a runtime term *fixes* it via int64 promotion. Here a runtime term is what *triggers* the crash, and the value is in `int32` range once wrapped; different mechanism, different symptom.

**Reach.** The trigger is: a **signed `int32`** expression (verified: NOT narrower ints — `int8`/`int16`/`uint32`/`int64` all wrap, see §17) where the simplifier folds two or more compile-time-constant terms whose exact sum/product leaves `[-2^31, 2^31)`, mixed with at least one runtime term so the summation path (not the scalar folder) is taken. `int32` is the default index/offset dtype, and the fold also fires through **index arithmetic** — `B[i] = A[(i + 2000000000) + 2000000000]` crashes identically (verified §17), which is the more idiomatic `base_off + stride*K` route.

Large-constant `int32` code appears in the tree — `examples/dsa_hisa/pool_mqa_fp8.py:76-77` seeds reduction bounds with `cu_k_s_min = 2147483647` / `cu_k_e_max = -2147483648` (also `examples/deepseek_v32/fp8_lighting_indexer.py:140-141`; `git grep` of ≥10-digit decimal literals under `examples/`+`testing/` hits ~5 non-mask/non-hash sites). **Ran the example** (`pool_mqa_fp8.py`, smallest shipped config `(32768,64,128,128,256,1)`, tilelang 0.1.13, L40S): the kernel **compiles cleanly past the IR simplifier** ("TileLang completes to compile kernel") with **no fold crash** — then fails at an unrelated launch-time step (`InternalError: Failed to set the allowed dynamic shared memory size to 116736`, i.e. the L40S ~99 KB smem cap, a hardware-capacity limit, not this bug). It dodges the fold because each sentinel is a **single** constant fed to `T.min`/`T.max` (`cu_k_s_min = T.min(cu_k_s_min, ...)`), never two constants summed together with a runtime term — the isolated-constant / max-arg case that §17 confirms wraps rather than crashes. So no shipped example currently exercises the crash; ordinary offset accumulation with large literals would.

**Impact.** The trigger is narrow: a compile-time fold of two or more constant terms whose exact sum/product leaves `[-2^31, 2^31)`, mixed with a runtime term. When it fires it is a compile-time `ICHECK` abort — loud and immediate, refusing to build that one legal kernel; nothing is silently miscomputed and no wrong value can reach a running workload. Fixing it removes a compile-time disagreement between the summation path and the plain constant-folder/hardware (both of which already wrap `int32`), closing the "constant fold reached through the summation path" boundary class.

§17 Generalization record

**Two-level root.**
- **SOURCE-level (fragile implementation):** the `arith` canonical simplifier's `SumExprNode` constant accumulator — `int64_t base` ([canonical_simplify.cc:245](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/src/arith/canonical_simplify.cc#L245)) mutated by `AddToSelf`/`MulToSelf` and rematerialized by `make_const(dtype, base)` in `Normalize`/`Normalize_` ([L510](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/src/arith/canonical_simplify.cc#L510)) **without** the dtype-masking (`GetFoldResultInt64Repr`) that the scalar folder [`TryConstFold`](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/src/arith/const_fold.h#L81) applies. The `IntImm` range check ([expr.cc:73](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/src/ir/expr.cc#L73)) then aborts.
- **OPERATOR-level (what correlates):** any construct whose two-or-more constant terms are canonicalized into one `SumExpr` *together with a runtime term* — value stores, index/address arithmetic, `+`/`*`/`-` mixes. The single-constant idioms (`min`/`max` sentinels) and all-constant sub-expressions route through the scalar folder instead and wrap.

**All cells run in their own fresh process, fresh `TILELANG_CACHE_DIR`, tilelang 0.1.13, L40S. Observed, not asserted.**

| Axis | Cell tested (input) | Observed result | Same root? |
|---|---|---|---|
| trigger | `(A[i] + int32(2e9)) + int32(2e9)` | CRASH `4000000000 vs 2147483648` | — |
| related-operator | mult: `int32(1073741824) * int32(2) + A[i]` | CRASH `2147483648 vs 2147483648` | **same** (base via `MulToSelf`) |
| related-operator | negative: `(A[i] + int32(-2e9)) + int32(-2e9)` | CRASH `4000000000 vs 2147483648` (materialized as positive operand of a sub) | **same** |
| related-operator | isolated const in max-arg: `T.max(A[i], int32(2e9)+int32(2e9))` | COMPILES, `[0,1,-1,100]` (const sub-expr wraps via scalar folder → `-294967296`) | **same root, dodged** — needs the runtime term *inside* the sum |
| related-source | index arithmetic: `A[(i + 2e9) + 2e9]` (big src buffer) | CRASH `4000000000 vs 2147483648` | **same** — fold fires through address expr too (idiomatic offset accumulation) |
| related-type | int8: `(A[i] + int8(100)) + int8(100)` | COMPILES, `[-56,-55,-57,-46]` (wraps) | **distinct** — int8 is not an `IsIndexType`, never enters `SumExpr`; scalar folder masks |
| related-type | int16: `(A[i] + int16(20000)) + int16(20000)` | COMPILES, `[-25536,…]` (wraps); mult form `int16(16384)*int16(4)` → wraps too | **distinct** — same reason as int8 |
| related-type | uint32: `(A[i] + uint32(3e9)) + uint32(3e9)` | COMPILES, `[1705032704,…]` (wraps) | **distinct** — `uint` is not `is_int()`, not an index type |
| related-type | int64: `(A[i] + int64(5e18)) + int64(5e18)` | COMPILES, `[-8446744073709551616,…]` (wraps) | **distinct boundary** — index type, but base wraps in C++ int64 AND range check gated off at `bits==64` |
| related-type (boundary) | int32 to exactly 2^31: `(A[i] + int32(2147483647)) + int32(1)` | CRASH `2147483648 vs 2147483648` | **same** — crash starts one past `INT32_MAX` |

**Reframed finding — the dtype boundary is much sharper than first stated.** The earlier "int32 or narrower-int" reach was **refuted by running it**: the crash is **signed-`int32`-ONLY**. `int8`/`int16` are excluded because `SumExpr` is entered only for [`IsIndexType`](https://github.com/tile-ai/tvm/blob/8df8ebd61659505dd7005f4820b030e983e2c005/src/arith/const_fold.h#L75) (`is_int() && bits∈{32,64}`); `uint32` is excluded (not `is_int()`); `int64` enters the path but is masked off by the C++ int64 wrap + the `bits<64` range-check gate. `int32` is the unique dtype that (a) enters the un-masked accumulator and (b) has the range check armed. The class is therefore narrow and specific — kept specific in the title, not broadened.

**Same-root neighbors folded into this report:** the multiplication form, the negative-overflow form, and the index-arithmetic form all share the identical `SumExpr.base`→`make_const` sink and are the same bug (documented above, not separate reports).

**Distinct-adjacent, not filed:** the `int8`/`int16`/`uint32`/`int64` wrapping behaviors are *correct* (not bugs) and delimit the boundary. No NEW bug was found while sweeping.

**Example-run result (PART 1).** `examples/dsa_hisa/pool_mqa_fp8.py` (the cited large-`int32`-literal site) was fetched at v0.1.13 with its local deps and run at its smallest shipped config on 0.1.13/L40S: it **compiles past the IR simplifier with no fold crash** and then fails at an unrelated launch step (`Failed to set the allowed dynamic shared memory size to 116736` — L40S smem cap, not this bug). It dodges the fold because the `2147483647`/`-2147483648` sentinels are each a **single** constant inside `T.min`/`T.max`, matching the "isolated const in max-arg" cell above which wraps rather than crashes. Confirms: no shipped example exercises this crash.

**SAME MECHANISM as filed [#2982](https://github.com/tile-ai/tilelang/issues/2982), different sink — do NOT merge.** #2982 is the same general root — an `IntImm`/`FloatImm` range check aborting on a value not first narrowed to the result dtype — at a different construction site. Here the sink is specifically the `SumExpr` `int64_t base` accumulator rematerialized via `make_const(dtype, base)`; #2982's is a distinct materialization point. Shared general root, distinct sinks: keep as separate reports.

Run the reproducer in the issue, then inspect 3rdparty/tvm/src/arith/canonical_simplify.cc, especially SumExprNode, AddToSelf, MulToSelf, and Normalize_. Compare this path with GetFoldResultInt64Repr in 3rdparty/tvm/src/arith/const_fold.h; done means the signed int32 kernel compiles and produces the wrapped result without an InternalError.

Written by the indexing model from the issue text.