[BUG] `makeBufferWithLayout` segfaults on a symbolic shared-buffer extent instead of hitting its own `ICHECK`
- 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
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