[spirv-reader][ir] Convert OpBitwiseAnd
Add conversion for OpBitwiseAnd, converting the arguments such that they
match the type of the first parameter.
Bug: 391487132
Change-Id: I036d29907103282d939396d52971f42b7e3f17ac
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/226854
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/spirv/builtin_fn.cc b/src/tint/lang/spirv/builtin_fn.cc
index 60cb6b9..ebb1b3b 100644
--- a/src/tint/lang/spirv/builtin_fn.cc
+++ b/src/tint/lang/spirv/builtin_fn.cc
@@ -170,6 +170,8 @@
return "convertSToF";
case BuiltinFn::kConvertUToF:
return "convertUToF";
+ case BuiltinFn::kBitwiseAnd:
+ return "bitwiseAnd";
case BuiltinFn::kSdot:
return "sdot";
case BuiltinFn::kUdot:
@@ -263,6 +265,7 @@
case BuiltinFn::kConvertFToS:
case BuiltinFn::kConvertSToF:
case BuiltinFn::kConvertUToF:
+ case BuiltinFn::kBitwiseAnd:
break;
}
return core::ir::Instruction::Accesses{};
diff --git a/src/tint/lang/spirv/builtin_fn.cc.tmpl b/src/tint/lang/spirv/builtin_fn.cc.tmpl
index 8203aa9..87210f2 100644
--- a/src/tint/lang/spirv/builtin_fn.cc.tmpl
+++ b/src/tint/lang/spirv/builtin_fn.cc.tmpl
@@ -106,6 +106,7 @@
case BuiltinFn::kConvertFToS:
case BuiltinFn::kConvertSToF:
case BuiltinFn::kConvertUToF:
+ case BuiltinFn::kBitwiseAnd:
break;
}
return core::ir::Instruction::Accesses{};
diff --git a/src/tint/lang/spirv/builtin_fn.h b/src/tint/lang/spirv/builtin_fn.h
index 687eaa6..43e52cc 100644
--- a/src/tint/lang/spirv/builtin_fn.h
+++ b/src/tint/lang/spirv/builtin_fn.h
@@ -112,6 +112,7 @@
kConvertFToS,
kConvertSToF,
kConvertUToF,
+ kBitwiseAnd,
kSdot,
kUdot,
kCooperativeMatrixLoad,
diff --git a/src/tint/lang/spirv/intrinsic/data.cc b/src/tint/lang/spirv/intrinsic/data.cc
index 15f7fee..ac8af96 100644
--- a/src/tint/lang/spirv/intrinsic/data.cc
+++ b/src/tint/lang/spirv/intrinsic/data.cc
@@ -6718,30 +6718,37 @@
},
{
/* [64] */
+ /* fn bitwiseAnd<R : iu32>[A : iu32, B : iu32](A, B) -> R */
+ /* fn bitwiseAnd<R : iu32>[A : iu32, B : iu32, N : num](vec<N, A>, vec<N, B>) -> vec<N, R> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(176),
+ },
+ {
+ /* [65] */
/* fn sdot(u32, u32, u32) -> i32 */
/* num overloads */ 1,
/* overloads */ OverloadIndex(193),
},
{
- /* [65] */
+ /* [66] */
/* fn udot(u32, u32, u32) -> u32 */
/* num overloads */ 1,
/* overloads */ OverloadIndex(194),
},
{
- /* [66] */
+ /* [67] */
/* fn cooperative_matrix_load<T : subgroup_matrix<K, S, C, R>>[K : subgroup_matrix_kind, S : fiu32_f16, C : num, R : num](ptr<workgroup_or_storage, S, readable>, u32, u32, u32) -> T */
/* num overloads */ 1,
/* overloads */ OverloadIndex(195),
},
{
- /* [67] */
+ /* [68] */
/* fn cooperative_matrix_store[K : subgroup_matrix_kind, S : fiu32_f16, C : num, R : num](ptr<workgroup_or_storage, S, writable>, subgroup_matrix<K, S, C, R>, u32, u32, u32) */
/* num overloads */ 1,
/* overloads */ OverloadIndex(196),
},
{
- /* [68] */
+ /* [69] */
/* fn cooperative_matrix_mul_add[T : subgroup_matrix_elements, TR : subgroup_matrix_elements, C : num, R : num, K : num](subgroup_matrix<subgroup_matrix_kind_left, T, K, R>, subgroup_matrix<subgroup_matrix_kind_right, T, C, K>, subgroup_matrix<subgroup_matrix_kind_result, TR, C, R>, u32) -> subgroup_matrix<subgroup_matrix_kind_result, TR, C, R> */
/* num overloads */ 1,
/* overloads */ OverloadIndex(197),
diff --git a/src/tint/lang/spirv/reader/lower/builtins.cc b/src/tint/lang/spirv/reader/lower/builtins.cc
index fc2b6cc..e7c7b88 100644
--- a/src/tint/lang/spirv/reader/lower/builtins.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins.cc
@@ -155,6 +155,9 @@
case spirv::BuiltinFn::kConvertUToF:
ConvertUToF(builtin);
break;
+ case spirv::BuiltinFn::kBitwiseAnd:
+ BitwiseAnd(builtin);
+ break;
default:
TINT_UNREACHABLE() << "unknown spirv builtin: " << builtin->Func();
}
@@ -225,6 +228,10 @@
call->Destroy();
}
+ void BitwiseAnd(spirv::ir::BuiltinCall* call) {
+ EmitBinaryWrappedAsFirstArg(call, core::BinaryOp::kAnd);
+ }
+
void Add(spirv::ir::BuiltinCall* call) {
EmitBinaryWrappedAsFirstArg(call, core::BinaryOp::kAdd);
}
diff --git a/src/tint/lang/spirv/reader/lower/builtins_test.cc b/src/tint/lang/spirv/reader/lower/builtins_test.cc
index 71ba45e..34f4beb 100644
--- a/src/tint/lang/spirv/reader/lower/builtins_test.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins_test.cc
@@ -5467,5 +5467,525 @@
EXPECT_EQ(expect, str());
}
+TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Scalar_SignedSigned_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kBitwiseAnd,
+ Vector{ty.i32()}, 1_i, 2_i);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.bitwiseAnd<i32> 1i, 2i
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = and 1i, 2i
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Scalar_SignedUnsigned_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kBitwiseAnd,
+ Vector{ty.i32()}, 1_i, 8_u);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.bitwiseAnd<i32> 1i, 8u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = bitcast 8u
+ %3:i32 = and 1i, %2
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Scalar_UnsignedSigned_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kBitwiseAnd,
+ Vector{ty.i32()}, 8_u, 1_i);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.bitwiseAnd<i32> 8u, 1i
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = bitcast 1i
+ %3:u32 = and 8u, %2
+ %4:i32 = bitcast %3
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Scalar_UnsignedUnsigned_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kBitwiseAnd,
+ Vector{ty.i32()}, 8_u, 9_u);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.bitwiseAnd<i32> 8u, 9u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = and 8u, 9u
+ %3:i32 = bitcast %2
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Scalar_UnsignedUnsigned_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kBitwiseAnd,
+ Vector{ty.u32()}, 8_u, 9_u);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.bitwiseAnd<u32> 8u, 9u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = and 8u, 9u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Scalar_UnsignedSigned_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kBitwiseAnd,
+ Vector{ty.u32()}, 8_u, 1_i);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.bitwiseAnd<u32> 8u, 1i
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = bitcast 1i
+ %3:u32 = and 8u, %2
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Scalar_SignedUnsigned_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kBitwiseAnd,
+ Vector{ty.u32()}, 1_i, 8_u);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.bitwiseAnd<u32> 1i, 8u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = bitcast 8u
+ %3:i32 = and 1i, %2
+ %4:u32 = bitcast %3
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Scalar_SignedSigned_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kBitwiseAnd,
+ Vector{ty.u32()}, 1_i, 2_i);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.bitwiseAnd<u32> 1i, 2i
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = and 1i, 2i
+ %3:u32 = bitcast %2
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Vector_SignedSigned_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), spirv::BuiltinFn::kBitwiseAnd,
+ Vector{ty.i32()}, b.Splat<vec2<i32>>(1_i),
+ b.Splat<vec2<i32>>(2_i));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = spirv.bitwiseAnd<i32> vec2<i32>(1i), vec2<i32>(2i)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = and vec2<i32>(1i), vec2<i32>(2i)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Vector_SignedUnsigned_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), spirv::BuiltinFn::kBitwiseAnd,
+ Vector{ty.i32()}, b.Splat<vec2<i32>>(1_i),
+ b.Splat<vec2<u32>>(8_u));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = spirv.bitwiseAnd<i32> vec2<i32>(1i), vec2<u32>(8u)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = bitcast vec2<u32>(8u)
+ %3:vec2<i32> = and vec2<i32>(1i), %2
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Vector_UnsignedSigned_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), spirv::BuiltinFn::kBitwiseAnd,
+ Vector{ty.i32()}, b.Splat<vec2<u32>>(8_u),
+ b.Splat<vec2<i32>>(1_i));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = spirv.bitwiseAnd<i32> vec2<u32>(8u), vec2<i32>(1i)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = bitcast vec2<i32>(1i)
+ %3:vec2<u32> = and vec2<u32>(8u), %2
+ %4:vec2<i32> = bitcast %3
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Vector_UnsignedUnsigned_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), spirv::BuiltinFn::kBitwiseAnd,
+ Vector{ty.i32()}, b.Splat<vec2<u32>>(8_u),
+ b.Splat<vec2<u32>>(9_u));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = spirv.bitwiseAnd<i32> vec2<u32>(8u), vec2<u32>(9u)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = and vec2<u32>(8u), vec2<u32>(9u)
+ %3:vec2<i32> = bitcast %2
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Vector_UnsignedUnsigned_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), spirv::BuiltinFn::kBitwiseAnd,
+ Vector{ty.u32()}, b.Splat<vec2<u32>>(8_u),
+ b.Splat<vec2<u32>>(9_u));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = spirv.bitwiseAnd<u32> vec2<u32>(8u), vec2<u32>(9u)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = and vec2<u32>(8u), vec2<u32>(9u)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Vector_UnsignedSigned_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), spirv::BuiltinFn::kBitwiseAnd,
+ Vector{ty.u32()}, b.Splat<vec2<u32>>(8_u),
+ b.Splat<vec2<i32>>(1_i));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = spirv.bitwiseAnd<u32> vec2<u32>(8u), vec2<i32>(1i)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = bitcast vec2<i32>(1i)
+ %3:vec2<u32> = and vec2<u32>(8u), %2
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Vector_SignedUnsigned_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), spirv::BuiltinFn::kBitwiseAnd,
+ Vector{ty.u32()}, b.Splat<vec2<i32>>(1_i),
+ b.Splat<vec2<u32>>(8_u));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = spirv.bitwiseAnd<u32> vec2<i32>(1i), vec2<u32>(8u)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = bitcast vec2<u32>(8u)
+ %3:vec2<i32> = and vec2<i32>(1i), %2
+ %4:vec2<u32> = bitcast %3
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Vector_SignedSigned_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), spirv::BuiltinFn::kBitwiseAnd,
+ Vector{ty.u32()}, b.Splat<vec2<i32>>(1_i),
+ b.Splat<vec2<i32>>(2_i));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = spirv.bitwiseAnd<u32> vec2<i32>(1i), vec2<i32>(2i)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = and vec2<i32>(1i), vec2<i32>(2i)
+ %3:vec2<u32> = bitcast %2
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
} // namespace
} // namespace tint::spirv::reader::lower
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.bazel b/src/tint/lang/spirv/reader/parser/BUILD.bazel
index fffb95a..c426687 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.bazel
+++ b/src/tint/lang/spirv/reader/parser/BUILD.bazel
@@ -82,6 +82,7 @@
alwayslink = True,
srcs = [
"binary_test.cc",
+ "bit_test.cc",
"branch_test.cc",
"builtin_test.cc",
"composite_test.cc",
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.cmake b/src/tint/lang/spirv/reader/parser/BUILD.cmake
index 51c33bd..330e5c1 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.cmake
+++ b/src/tint/lang/spirv/reader/parser/BUILD.cmake
@@ -91,6 +91,7 @@
################################################################################
tint_add_target(tint_lang_spirv_reader_parser_test test
lang/spirv/reader/parser/binary_test.cc
+ lang/spirv/reader/parser/bit_test.cc
lang/spirv/reader/parser/branch_test.cc
lang/spirv/reader/parser/builtin_test.cc
lang/spirv/reader/parser/composite_test.cc
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.gn b/src/tint/lang/spirv/reader/parser/BUILD.gn
index 382792d..13fa322 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.gn
+++ b/src/tint/lang/spirv/reader/parser/BUILD.gn
@@ -90,6 +90,7 @@
tint_unittests_source_set("unittests") {
sources = [
"binary_test.cc",
+ "bit_test.cc",
"branch_test.cc",
"builtin_test.cc",
"composite_test.cc",
diff --git a/src/tint/lang/spirv/reader/parser/bit_test.cc b/src/tint/lang/spirv/reader/parser/bit_test.cc
new file mode 100644
index 0000000..2e6de9b
--- /dev/null
+++ b/src/tint/lang/spirv/reader/parser/bit_test.cc
@@ -0,0 +1,575 @@
+// Copyright 2025 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/spirv/reader/parser/helper_test.h"
+
+namespace tint::spirv::reader {
+namespace {
+
+TEST_F(SpirvParserTest, BitwiseAnd_Scalar_SignedSigned_Signed) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %int = OpTypeInt 32 1
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpBitwiseAnd %int %one %two
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.bitwiseAnd<i32> 1i, 2i
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, BitwiseAnd_Scalar_SignedUnsigned_Signed) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %int = OpTypeInt 32 1
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpBitwiseAnd %int %one %eight
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.bitwiseAnd<i32> 1i, 8u
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, BitwiseAnd_Scalar_UnsignedSigned_Signed) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %int = OpTypeInt 32 1
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpBitwiseAnd %int %eight %one
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.bitwiseAnd<i32> 8u, 1i
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, BitwiseAnd_Scalar_UnsignedUnsigned_Signed) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %int = OpTypeInt 32 1
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpBitwiseAnd %int %eight %nine
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.bitwiseAnd<i32> 8u, 9u
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, BitwiseAnd_Scalar_UnsignedUnsigned_Unsigned) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %int = OpTypeInt 32 1
+ %uint = OpTypeInt 32 0
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpBitwiseAnd %uint %eight %nine
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.bitwiseAnd<u32> 8u, 9u
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, BitwiseAnd_Scalar_UnsignedSigned_Unsigned) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %int = OpTypeInt 32 1
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpBitwiseAnd %uint %eight %one
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.bitwiseAnd<u32> 8u, 1i
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, BitwiseAnd_Scalar_SignedUnsigned_Unsigned) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %int = OpTypeInt 32 1
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpBitwiseAnd %uint %one %eight
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.bitwiseAnd<u32> 1i, 8u
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, BitwiseAnd_Scalar_SignedSigned_Unsigned) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %int = OpTypeInt 32 1
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpBitwiseAnd %uint %one %two
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.bitwiseAnd<u32> 1i, 2i
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, BitwiseAnd_Vector_SignedSigned_Signed) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %int = OpTypeInt 32 1
+ %v2uint = OpTypeVector %uint 2
+ %v2int = OpTypeVector %int 2
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2two = OpConstantComposite %v2int %two %two
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %v2nine = OpConstantComposite %v2uint %nine %nine
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpBitwiseAnd %v2int %v2one %v2two
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = spirv.bitwiseAnd<i32> vec2<i32>(1i), vec2<i32>(2i)
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, BitwiseAnd_Vector_SignedUnsigned_Signed) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %int = OpTypeInt 32 1
+ %v2uint = OpTypeVector %uint 2
+ %v2int = OpTypeVector %int 2
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2two = OpConstantComposite %v2int %two %two
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %v2nine = OpConstantComposite %v2uint %nine %nine
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpBitwiseAnd %v2int %v2one %v2eight
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = spirv.bitwiseAnd<i32> vec2<i32>(1i), vec2<u32>(8u)
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, BitwiseAnd_Vector_UnsignedSigned_Signed) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %int = OpTypeInt 32 1
+ %v2uint = OpTypeVector %uint 2
+ %v2int = OpTypeVector %int 2
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2two = OpConstantComposite %v2int %two %two
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %v2nine = OpConstantComposite %v2uint %nine %nine
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpBitwiseAnd %v2int %v2eight %v2one
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = spirv.bitwiseAnd<i32> vec2<u32>(8u), vec2<i32>(1i)
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, BitwiseAnd_Vector_UnsignedUnsigned_Signed) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %int = OpTypeInt 32 1
+ %v2uint = OpTypeVector %uint 2
+ %v2int = OpTypeVector %int 2
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2two = OpConstantComposite %v2int %two %two
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %v2nine = OpConstantComposite %v2uint %nine %nine
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpBitwiseAnd %v2int %v2eight %v2nine
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = spirv.bitwiseAnd<i32> vec2<u32>(8u), vec2<u32>(9u)
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, BitwiseAnd_Vector_UnsignedUnsigned_Unsigned) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %int = OpTypeInt 32 1
+ %v2uint = OpTypeVector %uint 2
+ %v2int = OpTypeVector %int 2
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2two = OpConstantComposite %v2int %two %two
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %v2nine = OpConstantComposite %v2uint %nine %nine
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpBitwiseAnd %v2uint %v2eight %v2nine
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = spirv.bitwiseAnd<u32> vec2<u32>(8u), vec2<u32>(9u)
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, BitwiseAnd_Vector_UnsignedSigned_Unsigned) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %int = OpTypeInt 32 1
+ %v2uint = OpTypeVector %uint 2
+ %v2int = OpTypeVector %int 2
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2two = OpConstantComposite %v2int %two %two
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %v2nine = OpConstantComposite %v2uint %nine %nine
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpBitwiseAnd %v2uint %v2eight %v2one
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = spirv.bitwiseAnd<u32> vec2<u32>(8u), vec2<i32>(1i)
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, BitwiseAnd_Vector_SignedUnsigned_Unsigned) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %int = OpTypeInt 32 1
+ %v2uint = OpTypeVector %uint 2
+ %v2int = OpTypeVector %int 2
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2two = OpConstantComposite %v2int %two %two
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %v2nine = OpConstantComposite %v2uint %nine %nine
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpBitwiseAnd %v2uint %v2one %v2eight
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = spirv.bitwiseAnd<u32> vec2<i32>(1i), vec2<u32>(8u)
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, BitwiseAnd_Vector_SignedSigned_Unsigned) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %int = OpTypeInt 32 1
+ %v2uint = OpTypeVector %uint 2
+ %v2int = OpTypeVector %int 2
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2two = OpConstantComposite %v2int %two %two
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %v2nine = OpConstantComposite %v2uint %nine %nine
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpBitwiseAnd %v2uint %v2one %v2two
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = spirv.bitwiseAnd<u32> vec2<i32>(1i), vec2<i32>(2i)
+ ret
+ }
+}
+)");
+}
+
+} // namespace
+} // namespace tint::spirv::reader
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index 736ca0e..80d054f 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -668,6 +668,9 @@
case spv::Op::OpConvertUToF:
EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kConvertUToF);
break;
+ case spv::Op::OpBitwiseAnd:
+ EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kBitwiseAnd);
+ break;
case spv::Op::OpAccessChain:
case spv::Op::OpInBoundsAccessChain:
EmitAccess(inst);
diff --git a/src/tint/lang/spirv/spirv.def b/src/tint/lang/spirv/spirv.def
index 3cf4ad5..c8c879b 100644
--- a/src/tint/lang/spirv/spirv.def
+++ b/src/tint/lang/spirv/spirv.def
@@ -421,6 +421,9 @@
implicit(T: iu32) fn convertUToF<R: f32_f16>(T) -> R
implicit(T: iu32, N: num) fn convertUToF<R: f32_f16>(vec<N, T>) -> vec<N, R>
+implicit(A: iu32, B: iu32) fn bitwiseAnd<R: iu32>(A, B) -> R
+implicit(A: iu32, B: iu32, N: num) fn bitwiseAnd<R: iu32>(vec<N, A>, vec<N, B>) -> vec<N, R>
+
////////////////////////////////////////////////////////////////////////////////
// SPV_KHR_integer_dot_product instructions
////////////////////////////////////////////////////////////////////////////////
diff --git a/src/tint/lang/spirv/writer/printer/printer.cc b/src/tint/lang/spirv/writer/printer/printer.cc
index eaed98a..e5b4f59 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -1533,6 +1533,9 @@
case BuiltinFn::kConvertUToF:
op = spv::Op::OpConvertUToF;
break;
+ case BuiltinFn::kBitwiseAnd:
+ op = spv::Op::OpBitwiseAnd;
+ break;
case spirv::BuiltinFn::kNone:
TINT_ICE() << "undefined spirv ir function";
}
