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

Start with src/target/llvm/codegen_llvm.cc, especially CodeGenLLVM::VisitExpr_ for DivNode and the adjacent ModNode visitor, then run the provided FuseTIR and relax.build reproduction with an LLVM target. Done means boolean division and modulo no longer trigger an uncaught ICHECK; they should instead produce a diagnostic or a meaningful, documented lowering.

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

Either of these is 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 / documented as unsupported.

The compiler should never hard-abort on well-formed, successfully parsed user IR.

Actual behavior

The process aborts with an internal-check failure, with no source location and no user-actionable message:

InternalError: Check failed: dtype.MatchesCode(DLDataTypeCode::kDLFloat)
(UInt(8 bits, lanes=1) vs. kDLFloat) :

The abort 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));   // <-- aborts 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. Note the ModNode visitor immediately below has the identical
structure, so x % y on bool aborts the same way.

The check is meant as a "this can never happen" assertion, but it is reachable from
ordinary user code, so it turns a type error into a compiler crash.

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, USE_LLVM=ON
Steps to reproduce
The prim_func must be inlined into the Relax function (i.e. it has to go through
FuseTIR) before relax.build, otherwise the TIR body is never sent to the LLVM
backend.


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")     # <-- aborts here
Output (no Python traceback, the C++ ICHECK kills the process):


InternalError: Check failed: dtype.MatchesCode(DLDataTypeCode::kDLFloat)
 (UInt(8 bits, lanes=1) vs. kDLFloat) :
How this was found: we are running a differential-testing fuzzer over Relax programs
and this abort fires on ~100 distinct generated modules out of ~33k. The generated IR
is well-typed and parses cleanly; only codegen rejects it.

Triage
needs-triage
Bug
target: llvm
Please follow this repository's security policy and code of conduct.