This is an automated email from the ASF dual-hosted git repository.
MasterJH5574 pushed a commit to branch v0.27.0
in repository https://gitbox.apache.org/repos/asf/tvm.git
The following commit(s) were added to refs/heads/v0.27.0 by this push:
new 4d0ac168c7 [FIX][TIRx] Fix boolean bitwise-not codegen for C-like
targets (#20445)
4d0ac168c7 is described below
commit 4d0ac168c78c1e20bf05a6ec5926760bb1c52b05
Author: Shushi Hong <[email protected]>
AuthorDate: Thu Sep 24 23:01:11 2026 -0400
[FIX][TIRx] Fix boolean bitwise-not codegen for C-like targets (#20445)
Fix incorrect code generation for boolean `BitwiseNot`.
For boolean values, bitwise NOT should flip `false` to `true` and `true`
to `false`. The current C code generator emits `~value`, but C/C++ first
converts the boolean to an integer:
- `~false` becomes `~0`, which is `-1`.
- `~true` becomes `~1`, which is `-2`.
Both results are nonzero, so converting them back to boolean produces
`true` in both cases. CUDA boolean vectors have a related problem: their
generated vector types do not support the `~` operator, causing
compilation to fail.
This change emits `value == false` for boolean operands. It produces the
correct result for scalar booleans and uses the existing vector
comparison code to apply the operation to each lane. Integer operands
continue to use `~value`..
---
src/target/source/codegen_c.cc | 6 ++++++
tests/python/codegen/test_target_codegen_bool.py | 17 +++++++++++++++++
2 files changed, 23 insertions(+)
diff --git a/src/target/source/codegen_c.cc b/src/target/source/codegen_c.cc
index a496cee1b2..dd59944d85 100644
--- a/src/target/source/codegen_c.cc
+++ b/src/target/source/codegen_c.cc
@@ -667,6 +667,12 @@ void CodeGenC::Dispatch_(const prim::BitwiseXorNode* op,
std::ostream& os) { //
PrintBinaryExpr(op, "^", os, this);
}
void CodeGenC::Dispatch_(const prim::BitwiseNotNode* op, std::ostream& os) {
// NOLINT(*)
+ if (op->a.ty().MatchesCode(DLDataTypeCode::kDLBool)) {
+ // C promotes ~bool to int. Comparing against false preserves the boolean
+ // complement and uses the target's elementwise path for vector operands.
+ PrintExpr(prim::EQ(op->a, prim::MakeConst(op->a.ty(), false)), os);
+ return;
+ }
os << "(~";
PrintExpr(op->a, os);
os << ')';
diff --git a/tests/python/codegen/test_target_codegen_bool.py
b/tests/python/codegen/test_target_codegen_bool.py
index c47e469223..3069256ffe 100644
--- a/tests/python/codegen/test_target_codegen_bool.py
+++ b/tests/python/codegen/test_target_codegen_bool.py
@@ -105,5 +105,22 @@ def test_cmp_load_store(target):
run_and_check()
+def test_bitwise_not_c(tmp_path):
+ @T.prim_func
+ def complement(values: T.Buffer((5,), "int32"), output: T.Buffer((5,),
"bool")):
+ for i in range(5):
+ output[i] = T.bitwise_not(values[i] != 0)
+
+ built = tvm.compile(complement, target="c")
+ library = str(tmp_path / "complement.so")
+ built.export_library(library)
+ built = tvm.runtime.load_module(library)
+
+ values = np.array([-3, -1, 0, 1, 2], dtype="int32")
+ output = tvm.runtime.tensor(np.zeros(5, dtype="bool"))
+ built(tvm.runtime.tensor(values), output)
+ np.testing.assert_array_equal(output.numpy(), np.logical_not(values))
+
+
if __name__ == "__main__":
tvm.testing.main()
