Required prerequisites
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] - 1
Expected 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
For variable's range.
- #1728, fixed by #1735: the same maximum-value assertion, but in
Simplify, rather than this reproducer's LowerTileOp path.
Required prerequisites
What version of TileLang are you using?
0.1.14+cuda.git030556de
System information
030556de976356b21061d413843b4a29f75b20a8.Problem description
Two nested
T.serialloops with tensor-dependent bounds insideT.Parallelfail with an internal analyzer assertion intl.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
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
Forvariable's range.Simplify, rather than this reproducer'sLowerTileOppath.