[BUG] Nested data-dependent serial loops inside T.Parallel fail during LowerTileOp
- Dominant language
- Python
- Stars
- 7.4k
- Forks
- 745
- Avg merge
- 1d 5h
- Merged PRs (30d)
- 104
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.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
```python
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
```pytb
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](https://github.com/tile-ai/tilelang/issues/1442), fixed by [#1446](https://github.com/tile-ai/tilelang/pull/1446): a related analyzer-binding conflict around loop partitioning/vectorization; that fix snapshots the analyzer before visiting the loop body in layout inference.
- [#1472](https://github.com/tile-ai/tilelang/issues/1472), fixed by [#1649](https://github.com/tile-ai/tilelang/pull/1649): conflicting let-variable bindings in vectorization. This reproducer instead fails when binding a nested `For` variable's range.
- [#1728](https://github.com/tile-ai/tilelang/issues/1728), fixed by [#1735](https://github.com/tile-ai/tilelang/pull/1735): the same maximum-value assertion, but in `Simplify`, rather than this reproducer's `LowerTileOp` path.
Contributor guide
Research direction
Start with the provided Python reproducer and the tilelang.lower entry point, then trace the tl.LowerTileOp path to the assertion in 3rdparty/tvm/src/arith/int_set.cc. Compare the nested T.serial bounds inside T.Parallel with the related analyzer issues and verify that lowering succeeds for the example without changing its dynamic bounds behavior.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- python
- Domain
- compilers
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Active
- Clarity
- Mostly clear
- Newbie friendliness
- 55/100