tile-ai / tile-ai/tilelang

[BUG] Nested data-dependent serial loops inside T.Parallel fail during LowerTileOp

Open
#3,225 0 comments 0 reactions 0 assignees View on GitHub
bug
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

Open the contributing 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.