tile-ai / tile-ai/tilelang

[BUG] `makeBufferWithLayout` segfaults on a symbolic shared-buffer extent instead of hitting its own `ICHECK`

Closed Beginner friendly
#2,906 3 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

**Version:** v0.1.11 (`VERSION` = 0.1.11), commit `607a9144cf375e1176ad487dd5b0e3d87c2cb79a`
**Platform:** Linux, CUDA 13.0, NVIDIA RTX 4000 Ada Generation (sm_89), Python 3.10, torch 2.12.1+cu130
**Still present on `main`:** the two loops below are unchanged as of `af4753045895`
(source inspected; the reproduction itself was run on `607a9144`).

## Summary

`makeBufferWithLayout` in `src/transform/lower_tile_op.cc` dereferences the result of
`.as()` without a null check when computing a shared buffer's replication
factor. A symbolic (non-constant) extent therefore crashes the compiler with SIGSEGV.

The loop immediately after it guards the identical condition with an `ICHECK`, so the
crash appears to be an oversight rather than intended behaviour — the diagnostic that was
written for this case can never fire, because the unguarded loop runs first.

## Location

`src/transform/lower_tile_op.cc`, inside `if (IsSharedBuffer(buffer))`:

```cpp
Array buffer_shape = buffer->shape;
int buffer_extent = 1;
int layout_extent = 1;
for (size_t i = 0; i < buffer_shape.size(); i++) {
auto shape = buffer_shape[i].as();
buffer_extent *= shape->value; // <-- no null check
}
for (size_t i = 0; i < layout_shape.size(); i++) {
auto shape = layout_shape[i].as();
ICHECK(shape) << "Layout output shape must be constant integer, but got: "
<< layout_shape[i]; // <-- guarded
layout_extent *= shape->value;
}
```

If `buffer_shape[i]` is not an `IntImmNode`, `.as()` returns `nullptr` and
`shape->value` dereferences it.

## Reproduction

Run this **from a file**, not `python -c` — TVMScript's `@T.prim_func` parser calls
`inspect.getsourcelines`, which has no source under `-c`, so the parser fails first and the
crash never surfaces.

```python
import tilelang
import tilelang.language as T
from tilelang.layout import Layout

m = T.dynamic("m", "int32")
N = 256

@tilelang.jit(out_idx=[1])
def build():
@T.prim_func
def k(A: T.Tensor((m, N), "float16"), B: T.Tensor((m, N), "float16")):
with T.Kernel(1, threads=128) as bx:
s = T.alloc_shared([m, N], "float16")
# element counts match, so the layout.cc:678 product check passes and
# execution reaches makeBufferWithLayout with a symbolic shape
T.annotate_layout({s: Layout([m, N], lambda i, j: [i, j])})
T.copy(A[0:m, :], s)
T.copy(s, B[0:m, :])
return k

build()
```

**Actual:** process dies with SIGSEGV (exit `-11`), no diagnostic.

**Expected:** the existing message, i.e.
`Layout output shape must be constant integer, but got: m`

## Notes

A dynamic shared extent *without* an explicit layout annotation compiles fine — that path
never reaches `makeBufferWithLayout`, so the bug only surfaces when a layout annotation
forces the buffer through the remap path.

`T.dynamic("m", "int32")` and `T.symbolic("m")` behave identically here — both segfault with
the same stack. That is expected from `tilelang/language/symbolics.py`, where `symbolic` is a
`@deprecated(..., "v0.1.9")` alias that calls `dynamic(name, dtype)` and returns the same
`tirx.Var`. The reproduction uses `T.dynamic` as the non-deprecated spelling.

## Suggested fix

Guard the first loop the same way as the second, and report which buffer failed:

```cpp
for (size_t i = 0; i < buffer_shape.size(); i++) {
auto shape = buffer_shape[i].as();
ICHECK(shape) << "Shared buffer '" << buffer->name
<< "' extent must be a constant integer, but got: " << buffer_shape[i];
buffer_extent *= shape->value;
}
```

Happy to send a PR if that shape looks right.

Contributor guide

Open the contributing guide

Research direction

Read the shared-buffer handling in src/transform/lower_tile_op.cc, especially the first loop inside the IsSharedBuffer(buffer) branch and the guarded layout-shape loop that follows it. Run the provided TVMScript reproduction from a file, then verify that a symbolic extent reaches the intended ICHECK diagnostic instead of terminating with SIGSEGV.

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
72/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.