tile-ai / tile-ai/tilelang

[BUG][Fuzzer][ice-on-valid-code] `T.Parallel(coalesced_width=<int>)` aborts compilation instead of accepting the documented int

Open Beginner friendly
#3,013 0 comments 0 reactions 0 assignees View on GitHub
Dominant language
Python
Stars
7.4k
Forks
745
Avg merge
1d 1h
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) and [Discussions](https://github.com/tile-ai/tilelang/discussions) that this hasn't already been reported.

### What version of TileLang are you using?

0.1.13

### System information

TileLang 0.1.13 / CUDA 12.8 / PyTorch 2.8 ; run on NVIDIA L40S (sm_89). The failure is in `LayoutInference` (host-side layout planning), before any device code is generated, so it is architecture-independent.

### Problem description

Passing an `int` to `T.Parallel`'s documented `coalesced_width` keyword aborts compilation:

```
tvm.error.InternalError: coalesced_width should be an IntImmNode.
```

`coalesced_width` is a public keyword of `T.Parallel`, typed `int | None` in the signature ([loop.py#L15](https://github.com/tile-ai/tilelang/blob/8001cc4ccf6149382d2019654a19f59c1d4d0482/tilelang/language/loop.py#L13-L15)) and documented as `Optional[int]` ([loop.py#L28](https://github.com/tile-ai/tilelang/blob/8001cc4ccf6149382d2019654a19f59c1d4d0482/tilelang/language/loop.py#L28-L29)), so a plain Python int is exactly what the signature promises. It never reaches codegen — the crash fires during `LayoutInference`. The same int works everywhere else it is accepted: passing it via `T.copy(..., coalesced_width=)` compiles fine, and wrapping the int in `tir.IntImm` before handing it to `T.Parallel` also compiles fine (both verified below), so the value itself is legal.

Not a regression — see Provenance.

### Reproducible example code

```python
import tilelang
import tilelang.language as T
import torch
# 0.1.13: IntImm moved from tvm.tir to tvm.tirx; tilelang re-exports it as T.IntImm.

N, blk = 4096, 256
A = torch.arange(N, device="cuda", dtype=torch.float32)

# --- FAIL: documented int coalesced_width aborts compilation ---
@T.prim_func
def bug(A: T.Tensor((N,), "float32"), B: T.Tensor((N,), "float32")):
with T.Kernel(N // blk, threads=64) as bx:
for i in T.Parallel(blk, coalesced_width=4): # int, per the signature
B[bx * blk + i] = A[bx * blk + i]

tilelang.compile(bug, out_idx=[1], target="cuda")
# -> tvm.error.InternalError: coalesced_width should be an IntImmNode.

# --- CONTROL: identical kernel, coalesced_width wrapped in IntImm -> compiles & correct ---
@T.prim_func
def ok(A: T.Tensor((N,), "float32"), B: T.Tensor((N,), "float32")):
with T.Kernel(N // blk, threads=64) as bx:
for i in T.Parallel(blk, coalesced_width=T.IntImm("int32", 4)):
B[bx * blk + i] = A[bx * blk + i]

k = tilelang.compile(ok, out_idx=[1], target="cuda")
print("control match:", torch.equal(k(A), A)) # -> True
```

### Traceback

```
File ".../tilelang/cuda/pipeline.py", line 117, in CUDAPassPipelineBodyPrologue
mod = tilelang.transform.LayoutInference()(mod)
...
tvm::tl::ParallelOpNode::ComputePlanCandidate(...) const
tvm.error.InternalError: coalesced_width should be an IntImmNode.
```

### Expected behavior

`T.Parallel(..., coalesced_width=4)` should compile and use the requested coalesced width, matching its own `int | None` signature and the sibling `T.copy(..., coalesced_width=)` path, which accepts a bare int over the same value. At minimum a documented, correctly-typed argument should not abort compilation.

### Additional context

Guard boundary tested (this session, 0.1.13, L40S) — the bare-int path aborts unconditionally

Each a separate compile of the same kernel, varying only how `coalesced_width` is passed to `T.Parallel`:

| `coalesced_width=` | result |
|---|---|
| `1` (bare int, smallest legal) | CRASH `coalesced_width should be an IntImmNode.` |
| `2` (bare int) | CRASH `coalesced_width should be an IntImmNode.` |
| `T.IntImm("int32", 2)` (wrapped) | compiles, output matches |
| `T.IntImm("int32", 4)` (wrapped, CONTROL) | compiles, output matches |

So the abort is unconditional on the bare-`int` path (independent of the value) and disappears the moment the same value is wrapped in `IntImm` — a **missing-case / type-coercion gap**, not a value problem. It is a *distinct* branch from the sibling divisibility abort ("Vector size N is not divisible by coalesced width M", `parallel.cc:804`): that one honors the int type and rejects on geometry; this one rejects the int type itself (`parallel.cc:808`). Same function `ComputePlanCandidate`, adjacent code, different root — reported separately.

**Root cause.** `T.Parallel` attaches the `coalesced_width` argument to the loop-annotation map as a raw Python int, but the C++ layout planner reads it back as an `IntImmNode` and hard-fails when it is not one. In [`loop.py#L79-L87`](https://github.com/tile-ai/tilelang/blob/8001cc4ccf6149382d2019654a19f59c1d4d0482/tilelang/language/loop.py#L79-L87) the value is placed in `merged_annotations` unconverted (`merged_annotations["coalesced_width"] = coalesced_width`) and handed straight to `_ffi_api.Parallel(extents, merged_annotations)`, which attaches the dict to the `For` node's annotation Map; the value arrives as a runtime int, not a `tir.IntImm`. Then in [`parallel.cc#L798-L809`](https://github.com/tile-ai/tilelang/blob/8001cc4ccf6149382d2019654a19f59c1d4d0482/src/op/parallel.cc#L798-L809), `ParallelOpNode::ComputePlanCandidate` does `coalesced_width->as()`, which returns null, and takes the `LOG(FATAL)` branch (`parallel.cc:808`).

The `T.copy` family stores the argument the same raw way ([`copy_op.py#L124-L125`](https://github.com/tile-ai/tilelang/blob/8001cc4ccf6149382d2019654a19f59c1d4d0482/tilelang/language/copy_op.py#L124-L125)), and its lowering ([`copy.cc#L538-L540`](https://github.com/tile-ai/tilelang/blob/8001cc4ccf6149382d2019654a19f59c1d4d0482/src/op/copy.cc#L538-L540)) *does* re-attach `kCoalescedWidth` onto a `ForKind::kParallel` loop that reaches the **same** `ComputePlanCandidate` / `as()` read. The reason copy does not crash is not that it skips that read — it is that copy passes its annotations through `call_intrin(..., annotations=ann)`, which builds a TVM `Map` and coerces the Python int to a `tir.IntImm` (a proper `IntImmNode`) as part of node construction, whereas `_ffi_api.Parallel` attaches the raw Python dict directly and no such coercion happens. So the divergence is a coercion gap at the FFI boundary of `T.Parallel`, not a difference in what the C++ planner reads. (Verified this session: `T.copy(..., coalesced_width=4)` compiles and returns the correct result on CUDA; see Reach.)

**Suggested fix.** Wrapping the value at the Python boundary — e.g. `merged_annotations["coalesced_width"] = IntImm("int32", coalesced_width)` in `T.Parallel` — makes the annotation an `IntImmNode` and lets the existing `parallel.cc` branch run; the CONTROL above (which does exactly this wrap at the call site) compiles and returns the correct result, so this is the minimal change. Alternatively `ComputePlanCandidate` could coerce the annotation via `Downcast`/`arith` before the `as()` read.

**Provenance.** The `coalesced_width` keyword on `T.Parallel` was added by [#1887](https://github.com/tile-ai/tilelang/pull/1887) (merged 2026-03-06), which introduced the raw-int passthrough at `loop.py:80`; that is when the crash first became reachable. The `as()` FATAL in `parallel.cc` predates it (present since the codebase migration #10, 2025-01-11), so the two halves never agreed on the type — the argument crashes on every release that exposes the keyword, not a regression.

**Dedup.** Searched the open and closed tracker (`coalesced_width`, `IntImmNode`, `should be an IntImmNode`) and found no existing report of this defect.

**Reach.** The trigger is documented and typed: `coalesced_width` is declared `int | None` and documented `Optional[int]`, so an int is the type the API asks for. `coalesced_width` appears in shipped code (`git grep coalesced_width examples/ testing/` at v0.1.13): every site routes it through `T.copy` / `T.async_copy`, e.g. `examples/amd/example_amd_flash_attn_bwd.py:163,172-173` and `examples/amd/example_amd_flash_attn_fwd.py:183,192-193` (`coalesced_width=vec_size`), and `testing/python/amd/test_tilelang_gfx950_copy_async.py`. No example or test passes it to `T.Parallel` (`git grep 'Parallel(' examples/ testing/ | grep coalesced_width` at v0.1.13 is empty), so CI never exercises the crashing path. These shipped sites are all AMD gfx950-targeted and were not runnable on this session's hardware (L40S, sm_89), but the load-bearing fact they establish — that a **bare int** through the copy family is a legal `coalesced_width` — was verified directly on CUDA this session: `T.copy(dst, src, coalesced_width=4)` compiled and returned `torch.equal(...) == True`. Given the keyword on `T.Parallel`, any value trips the crash — the failure is unconditional on that int path (the CONTROL only differs by wrapping the same value in `IntImm`), and it also fires when the int is supplied via the generic `annotations={"coalesced_width": }` dict, so it is the FFI-attach path of `T.Parallel`, not the specific keyword.

**Impact.** The narrow ingredient is passing `coalesced_width` as a bare int through `T.Parallel` specifically (the sibling `T.copy` path and the `IntImm`-wrapped form both compile). When it fires it is a loud compile-time abort during `LayoutInference` — no device code is produced, so nothing is silently miscompiled and no output is corrupted; the failure is deterministic and caught immediately at build time. Fixing it removes a type disagreement between the Python signature (`int | None`) and the C++ `as()` read, letting a documented, correctly-typed argument build the kernel it should.

Generalization (root analysis + 4-axis sweep, all cells run this session, 0.1.13 / L40S)

**Two-level root.**
- *Source-level:* `_ffi_api.Parallel(extents, merged_annotations)` ([loop.py#L79-L87](https://github.com/tile-ai/tilelang/blob/8001cc4ccf6149382d2019654a19f59c1d4d0482/tilelang/language/loop.py#L79-L87)) attaches the raw Python annotation dict to the `For` node with no int→`IntImm` coercion, while `ParallelOpNode::ComputePlanCandidate` reads the value through `as()` and `LOG(FATAL)`s on a miss ([parallel.cc#L798-L809](https://github.com/tile-ai/tilelang/blob/8001cc4ccf6149382d2019654a19f59c1d4d0482/src/op/parallel.cc#L798-L809)). The fragile code is this uncoerced FFI-attach on the `T.Parallel` frame.
- *Operator-level:* the `coalesced_width` (`kCoalescedWidth`) annotation is consumed only in `ComputePlanCandidate`. Every op that produces a parallel loop carrying that annotation — `T.Parallel`, and via re-attach `T.copy` / `T.async_copy` ([copy.cc#L538-L540](https://github.com/tile-ai/tilelang/blob/8001cc4ccf6149382d2019654a19f59c1d4d0482/src/op/copy.cc#L538-L540)) and the atomic reductions (`atomic_add.cc:200-202`, `atomic_reduce.h:129-131`) — funnels through that single `as()` read. The copy/atomic ops are safe only because their annotations are laundered into an `IntImm` by `call_intrin`'s `Map` construction; `T.Parallel` is the one entry point that attaches raw.

| axis | cell tested | result | same-root? |
|---|---|---|---|
| repro | `T.Parallel(..., coalesced_width=4)` (bare int) | CRASH `coalesced_width should be an IntImmNode.` | — (the bug) |
| control | `T.Parallel(..., coalesced_width=T.IntImm("int32",4))` | COMPILES, `torch.equal==True` | — (control) |
| similar-logic | same int via generic `T.Parallel(..., annotations={"coalesced_width":4})` (bypass the typed kwarg) | CRASH `... should be an IntImmNode.` | **same root** — proves it is the FFI-attach path, not the keyword |
| related-operator | `T.copy(dst, src, coalesced_width=4)` (bare int) | COMPILES, `torch.equal==True` | distinct outcome — same `as()` read, but `call_intrin` coerces int→`IntImm`, so no crash (documents WHY sibling dodges) |
| related-operator | `T.copy(dst, src, annotations={"coalesced_width":4})` (bare int via dict) | COMPILES | same as above — copy launders the dict through `Map` construction |
| related-source | divisibility sibling branch: `coalesced_width=T.IntImm("int32",3)` (vector_size=4) | CRASH `Vector size 4 is not divisible by coalesced width 3` (parallel.cc:804) | **distinct** — honors the int type, rejects on geometry; not a type-coercion gap. Not filed. |
| related-type | `coalesced_width=T.IntImm("int64",4)` (width neighbor) | COMPILES, `torch.equal==True` | any `IntImmNode` width passes — confirms the check is type-tag only |
| related-type | `coalesced_width=np.int32(4)` (common user int, non-`IntImm` runtime scalar) | CRASH `... should be an IntImmNode.` | **same root** — any non-`IntImm` scalar on the `T.Parallel` path aborts |

**PART-1 example run.** The shipped `coalesced_width` sites (AMD flash-attn fwd/bwd examples, gfx950 copy-async test) were located in the v0.1.13 tree; all route the value through `T.copy`/`T.async_copy`, none through `T.Parallel`, so none exercise this crash — the draft's Reach is confirmed by source, not merely inherited. They are AMD-targeted and were not run on this L40S session; the sibling-path claim they support (bare int is a legal `coalesced_width`) was instead verified directly on CUDA (row 3/4 above).

**Framing decision.** The two crashing cells (typed kwarg, generic annotations dict) share one root, so the write-up is kept at the `T.Parallel`-int-path level (title/Problem unchanged) — the class is exactly "a bare/non-`IntImm` `coalesced_width` reaching `ComputePlanCandidate` via the raw `T.Parallel` FFI attach." The divisibility branch (parallel.cc:804) is a **distinct** adjacent check and is *not filed*. No new bug was found while sweeping.

Contributor guide

Open the contributing guide

Research direction

Start in tilelang/language/loop.py at the T.Parallel annotation handling, then compare it with src/op/parallel.cc around ComputePlanCandidate and its coalesced_width check. Reproduce the bare-int and IntImm control cases, and consider the issue resolved when the documented bare int compiles successfully with the requested width.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp, python
Domain
compilers
Issue type
Bug
Difficulty
2/5
Estimated time
1-3 hours
Activity status
Quiet
Clarity
Clearly specified
Newbie friendliness
78/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.