[Bug] LLVM codegen aborts on arithmetic over "bool" buffers (Add/Sub/Mul) #20558

Thanks for participating in the TVM community! We use https://discuss.tvm.apache.org/ for any general usage questions and discussions. The issue tracker is used for actionable items such as feature proposals discussion, roadmaps, and bug tracking. You are always welcomed to post on the forum first 😸

Issues that are inactive for a period of time may get closed. We adopt this policy so that we won't lose track of actionable issues that may fall at the bottom of the pile. Feel free to reopen a new one if you feel there is an additional problem that needs attention when an old one gets closed.

Expected behavior

A PrimFunc that performs +, -, or * on buffers whose dtype is "bool" should either:

• build successfully; or

• be rejected earlier by a user-facing validation/type-checking pass with a proper diagnostic.

bool is a valid dtype in TVM's type system, and the arithmetic type relation accepts it. The module below parses, constructs, and passes IR-level checks without complaint.

The same module also builds successfully with target="c".

Actual behavior

tvm.tirx.build(mod, target="llvm") fails with an internal check:

tvm.error.InternalError: Check failed: (t.MatchesCode(DLDataTypeCode::kDLFloat)) is false:

The failing check is in DEFINE_CODEGEN_BINARY_OP in:

src/target/llvm/codegen_llvm.cc:1595

For Add, the relevant expansion is used at:

src/target/llvm/codegen_llvm.cc:1608

The macro dispatches on dtype and only handles three cases:

if (t.MatchesCode(DLDataTypeCode::kDLInt)) {
  ...
} else if (t.MatchesCode(DLDataTypeCode::kDLUInt)) {
  ...
} else {
  TVM_FFI_ICHECK(t.MatchesCode(DLDataTypeCode::kDLFloat));
  return builder_->CreateF##Op(a, b);
}

bool matches neither kDLInt nor kDLUInt, so execution falls into the else branch and triggers the assertion because it also does not match kDLFloat.

The same behavior occurs for Sub and Mul.

Tail of the traceback:

File "src/target/llvm/codegen_llvm.cc", line 2077, in VisitStmt_(BufferStoreNode const*)
  llvm::Value* value = MakeValue(op->value);
File "src/target/llvm/codegen_llvm.cc", line 1608, in VisitExpr_(prim::AddNode const*)

DEFINE_CODEGEN_BINARY_OP(Add);
File "src/target/llvm/codegen_llvm.cc", line 1608, in CreateAdd(PrimType, llvm::Value*, llvm::Value*)

DEFINE_CODEGEN_BINARY_OP(Add);
tvm.error.InternalError: Check failed: (t.MatchesCode(DLDataTypeCode::kDLFloat)) is false:

Dtype coverage

For:

C[i] = A[i] + B[i]

the following behavior is observed on target="llvm".

The same result occurs for - and *.

Environment

• TVM commit: 8312a17f8734ddfd56e5f3977cd5df25b83ec49f

• tvm.__version__: 0.19.0

• Build: built from source

• Target: llvm

• Comparison target: the same module builds with target="c"

• Python: 3.12.12

• OS: Linux

• Kernel: 5.15.0-139-generic

Steps to reproduce

import tvm
from tvm.script import tirx as T
@T.prim_func

def f(

A: T.Buffer((4,), "bool"),

B: T.Buffer((4,), "bool"),

C: T.Buffer((4,), "bool"),

):

for i in T.serial(4):

C[i] = A[i] + B[i]
mod = tvm.IRModule({"f": f})

tvm.tirx.build(mod, target="llvm")

Triage

• needs-triage

• type: bug