apache / apache/tvm

[Bug] InternalError "Check failed: indices.size() == 1 (2 vs. 1): CodeGenLLVM requires all buffers to be flat 1-d buffers"

Open
#17,478 0 comments 0 reactions 0 assignees View on GitHub
needs-triage type: bug
Dominant language
Python
Stars
13.7k
Forks
4k
Avg merge
2d 19h
Merged PRs (30d)
111

Description

When building a Relax module with a TIR function that uses axis separators in the buffer layout, TVM throws an InternalError during the compilation process. The error message suggests that TVM's LLVM code generation requires all buffers to be flat 1-dimensional, but the provided buffer has more than one index.

The issue appears to be related to how the `axis_separators` attribute is handled in match_buffer, which results in a mismatch during LLVM code generation.

### Steps to reproduce
```python
import tvm
from tvm import relax

from tvm.script import ir as I
from tvm.script import tir as T
from tvm.script import relax as R

@I.ir_module
class Module:
@T.prim_func(private=True)
def te_layout_transform_axis_separator(x: T.Buffer((T.int64(16),), "float32"), var_te_layout_transform_axis_separator: T.handle):
T.func_attr({"tir.noalias": T.bool(True)})
te_layout_transform_axis_separator = T.match_buffer(var_te_layout_transform_axis_separator, (T.int64(4), T.int64(4)), axis_separators=[1])
# with T.block("root"):
for self in range(T.int64(16)):
with T.block("te_layout_transform_axis_separator"):
v_self = T.axis.spatial(T.int64(16), self)
T.reads(x[v_self])
T.writes(te_layout_transform_axis_separator[v_self // T.int64(4), v_self % T.int64(4)])
te_layout_transform_axis_separator[v_self // T.int64(4), v_self % T.int64(4)] = x[v_self]

@R.function
def main(x: R.Tensor((16,), dtype="float32")) -> R.Tensor((4, 4), dtype="float32"):
cls = Module
with R.dataflow():
lv = R.call_tir(cls.te_layout_transform_axis_separator, (x,), out_sinfo=R.Tensor((4, 4), dtype="float32"))
R.output(lv)
return lv

mod = Module
ex = relax.build(mod, target='llvm')
```

### Actual behavior

```
Traceback (most recent call last):
File "path/to/script.py", line 31, in
ex = relax.build(mod, target='llvm')
...
tvm.error.InternalError: Traceback (most recent call last):
...
File "/software/tvm/src/target/llvm/codegen_llvm.cc", line 960
InternalError: Check failed: indices.size() == 1 (2 vs. 1) : CodeGenLLVM requires all buffers to be flat 1-d buffers.
```

Could you please help confirm if this is a bug in TVM or an issue with my usage?

Contributor guide

No contributing guide indexed for this repository

Research direction

Start by running the provided Relax module with the LLVM target and inspect src/target/llvm/codegen_llvm.cc at line 960. Trace how axis_separators on T.match_buffer reaches LLVM code generation; the issue is complete when this reproduction no longer raises the reported InternalError or the supported usage is clearly established.

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
Stale
Clarity
Needs clarification
Newbie friendliness
35/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.