Skip to content

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

Description

@sepcnt

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.

Activity

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Labels

    bugSomething isn't working

    Type

    No type

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions