[Bug] CodeGenLLVM hits an ICHECK instead of reporting an error when dividing boolean values

[yunligou711-commits]

Expected behavior

Either of these would be acceptable:

1. T.div on bool — a dtype code that is neither kDLInt, kDLUInt, nor kDLFloat — is rejected early by the TIR/Relax type checker with a proper DiagnosticError pointing at the offending expression; or
2. it is lowered to something meaningful, or documented as unsupported.

Codegen should not raise an internal check ("this should be impossible") on IR that parsed and type-checked successfully.

---

Actual behavior

relax.build fails with a tvm.error.InternalError raised from an internal check in the LLVM backend.

The message carries no source location and no user-actionable guidance:

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

The message ends there — it does not even report the offending dtype.

The exception is raised as an ordinary catchable Python exception; the process exits with code 1, it is not aborted.

The check comes from CodeGenLLVM::VisitExpr_(const prim::DivNode*) in src/target/llvm/codegen_llvm.cc:

llvm::Value* CodeGenLLVM::VisitExpr_(const prim::DivNode* op) {
  llvm::Value* a = MakeValue(op->a);
  llvm::Value* b = MakeValue(op->b);
  PrimType dtype(op->ty.as_or_throw<PrimType>()->dtype);

  if (dtype.MatchesCode(DLDataTypeCode::kDLInt)) {
    return builder_->CreateSDiv(a, b);
  } else if (dtype.MatchesCode(DLDataTypeCode::kDLUInt)) {
    return builder_->CreateUDiv(a, b);
  } else {
    TVM_FFI_ICHECK(dtype.MatchesCode(DLDataTypeCode::kDLFloat));  // <-- fires for bool
    return builder_->CreateFDiv(a, b);
  }
}

bool has DLDataTypeCode::kDLBool (3rdparty/dlpack/include/dlpack/dlpack.h:161), which matches neither kDLInt nor kDLUInt, so control falls into the else branch and the ICHECK fires.

By inspection, the ModNode visitor immediately below has the identical structure, so x % y on bool should hit the same check.

---

Environment

• TVM version: 0.26.dev0

• Commit: 8312a17f8734ddfd56e5f3977cd5df25b83ec49f

  • [REFACTOR][Arith] Evaluate iterator domains through Var maps (#20354)
  • 2026-09-15

• OS: Ubuntu 20.04.6 LTS

• Kernel: 5.15.0-139-generic

• Compiler: GCC 13.1.0

• Target: LLVM (host)

• CPU: Intel Xeon Gold 6326

• Build: Standard CMake build with USE_LLVM=ON

---

Steps to reproduce

The prim_func has to be inlined into the Relax function (i.e. it must go through FuseTIR) before relax.build.

Otherwise, the TIR body is never handed to the LLVM backend and no failure occurs.

import numpy as np
import tvm
from tvm import relax
from tvm.script import ir as I
from tvm.script import tirx as T
from tvm.script import relax as R


@I.ir_module
class Module:
    @T.prim_func(private=True, s_tir=True)
    def bool_div(
        x: T.Buffer((T.int64(4),), "bool"),
        y: T.Buffer((T.int64(4),), "bool"),
        out: T.Buffer((T.int64(4),), "bool"),
    ):
        T.func_attr({"tir.noalias": T.bool(True)})
        for ax0 in range(T.int64(4)):
            with T.sblock("bool_div"):
                v_ax0 = T.axis.spatial(T.int64(4), ax0)
                T.reads(x[v_ax0], y[v_ax0])
                T.writes(out[v_ax0])
                out[v_ax0] = x[v_ax0] / y[v_ax0]  # <-- bool / bool

    @R.function
    def main(
        x: R.Tensor((4,), dtype="bool"),
        y: R.Tensor((4,), dtype="bool"),
    ) -> R.Tensor((4,), dtype="bool"):
        cls = Module
        with R.dataflow():
            gv = R.call_tir(
                cls.bool_div,
                (x, y),
                out_ty=R.Tensor((4,), dtype="bool"),
            )
            R.output(gv)
        return gv


mod = Module

with tvm.transform.PassContext(disabled_pass=["RemoveUnusedParameters"]):
    mod = relax.transform.FuseTIR()(mod)

mod = relax.transform.LegalizeOps()(mod)
mod = relax.transform.LambdaLift()(mod)

with tvm.transform.PassContext(opt_level=4):
    ex = relax.build(mod, target="llvm")  # <-- fails here

---

Output

Exit code: 1

Tail of the traceback:

File ".../src/target/llvm/codegen_llvm.cc", line 1641, in virtual llvm::Value* tvm::codegen::CodeGenLLVM::VisitExpr_(const tvm::prim::DivNode*)
    TVM_FFI_ICHECK(dtype.MatchesCode(DLDataTypeCode::kDLFloat));

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

---

How this was found

The IR parses and type-checks cleanly; only the LLVM backend rejects it.

---

Triage

• needs-triage
• Bug
• target: llvm

[Nanmur]

I opened #20412 to address this.

The patch rejects boolean operands in TIRx div and truncmod before LLVM code generation, so bool / bool and bool % bool now raise a clear TypeError instead of reaching the CodeGenLLVM internal check. It also adds a regression test for both operations.