yunligou711-commits opened a new issue, #20399: URL: https://github.com/apache/tvm/issues/20399
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 :smile_cat: 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`: ```cpp 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. -- This is an automated message from the Apache Git Service. To respond to the message, please log on to GitHub and use the URL above to go to the specific comment. To unsubscribe, e-mail: [email protected] For queries about this service, please contact Infrastructure at: [email protected] --------------------------------------------------------------------- To unsubscribe, e-mail: [email protected] For additional commands, e-mail: [email protected]
