This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm.git
The following commit(s) were added to refs/heads/main by this push:
new 009c117d04 [FIX][TIRx] Preserve boolean bitwise not semantics in C
codegen (#20461)
009c117d04 is described below
commit 009c117d0469163128fc7098ade201f2f47ea469
Author: Shushi Hong <[email protected]>
AuthorDate: Fri Sep 25 19:20:52 2026 -0400
[FIX][TIRx] Preserve boolean bitwise not semantics in C codegen (#20461)
The same as #20445 but for the main. This pr fixes 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 a2ef816053..b4b85f6f66 100644
--- a/tests/python/codegen/test_target_codegen_bool.py
+++ b/tests/python/codegen/test_target_codegen_bool.py
@@ -87,5 +87,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()
