[BUG] Nested data-dependent serial loops inside T.Parallel fail during LowerTileOp
Author: sepcntCreated Sep 14, 2026Updated Sep 14, 2026
Labelsbug
Required prerequisites
- I have read the documentation https://tilelang.com.
- I have searched the Issue Tracker that this hasn't already been reported. (comment there if it has.)
What version of TileLang are you using?
0.1.14+cuda.git030556de
System information
- Windows x86-64; freshly rebuilt from commit
030556de976356b21061d413843b4a29f75b20a8. - Python 3.12.13, PyTorch 2.11.0+cu130, apache-tvm-ffi 0.1.12.
- RTX 5060 Laptop GPU; CUDA toolkit 13.2.78.
Problem description
Two nested T.serial loops with tensor-dependent bounds inside T.Parallel fail with an internal analyzer assertion in tl.LowerTileOp.
The example has fixed tensor shapes and only performs lowering: no GPU execution, NVCC compilation, or CUDA Graph is involved. Assume input bounds are in [0, 4]; each parallel iteration writes a distinct output slice.
Reproducible example code
import tilelang
import tilelang.language as T
@T.prim_func
def main(bounds: T.Tensor[(32, 2), "int32"],
out: T.Tensor[(32, 4, 4), "int32"]):
with T.Kernel(1, threads=32):
for p in T.Parallel(32):
for y in T.serial(bounds[p, 0]):
for x in T.serial(bounds[p, 1]):
out[p, y, x] = p * 100 + y * 10 + x
target = tilelang.tvm.target.Target({"kind": "cuda", "arch": "sm_120"})
with target:
tilelang.lower(main, target=target,
enable_host_codegen=False, enable_device_compile=False)Traceback
mod = tilelang.transform.LowerTileOp()(mod)
...
File "..\3rdparty\tvm\src\arith\int_set.cc", line 701,
in tvm::arith::IntSetAnalyzer::Impl::Update(...)
tvm.error.InternalError: Check failed: (ExprDeepEqual()(old_info.max(), info.max())) is false:
Trying to update var 'y' with a different maximum value:
original=bounds[p, 0] - 1, new=bounds[tx, 0] - 1Expected behavior
Successful lowering, allowing each thread to execute its own dynamically bounded serial loops.
Additional context
- #1442, fixed by #1446: a related analyzer-binding conflict around loop partitioning/vectorization; that fix snapshots the analyzer before visiting the loop body in layout inference.
- #1472, fixed by #1649: conflicting let-variable bindings in vectorization. This reproducer instead fails when binding a nested
Forvariable's range. - #1728, fixed by #1735: the same maximum-value assertion, but in
Simplify, rather than this reproducer'sLowerTileOppath.
Source: tile-ai/tilelang