[spirv-reader][ir] Add OpShiftRightLogical support
Add support for converting an OpShiftRightLogical to a bitwise `>>`
operator.
Bug: 391487016
Change-Id: I607d2d0639ffc895c0f53a9789a4a49a3ae6d62a
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/227635
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/spirv/builtin_fn.cc b/src/tint/lang/spirv/builtin_fn.cc
index 7a705a1..236544e 100644
--- a/src/tint/lang/spirv/builtin_fn.cc
+++ b/src/tint/lang/spirv/builtin_fn.cc
@@ -198,6 +198,8 @@
return "not_equal";
case BuiltinFn::kShiftLeftLogical:
return "shift_left_logical";
+ case BuiltinFn::kShiftRightLogical:
+ return "shift_right_logical";
case BuiltinFn::kSdot:
return "sdot";
case BuiltinFn::kUdot:
@@ -305,6 +307,7 @@
case BuiltinFn::kULessThan:
case BuiltinFn::kULessThanEqual:
case BuiltinFn::kShiftLeftLogical:
+ case BuiltinFn::kShiftRightLogical:
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 e3b3344..3497e81 100644
--- a/src/tint/lang/spirv/builtin_fn.cc.tmpl
+++ b/src/tint/lang/spirv/builtin_fn.cc.tmpl
@@ -120,6 +120,7 @@
case BuiltinFn::kULessThan:
case BuiltinFn::kULessThanEqual:
case BuiltinFn::kShiftLeftLogical:
+ case BuiltinFn::kShiftRightLogical:
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 93b6e62..76139ae 100644
--- a/src/tint/lang/spirv/builtin_fn.h
+++ b/src/tint/lang/spirv/builtin_fn.h
@@ -126,6 +126,7 @@
kEqual,
kNotEqual,
kShiftLeftLogical,
+ kShiftRightLogical,
kSdot,
kUdot,
kCooperativeMatrixLoad,
diff --git a/src/tint/lang/spirv/intrinsic/data.cc b/src/tint/lang/spirv/intrinsic/data.cc
index fcd2f46..a91ff42 100644
--- a/src/tint/lang/spirv/intrinsic/data.cc
+++ b/src/tint/lang/spirv/intrinsic/data.cc
@@ -6846,30 +6846,37 @@
},
{
/* [78] */
+ /* fn shift_right_logical<R : iu32>[A : iu32, B : iu32](A, B) -> R */
+ /* fn shift_right_logical<R : iu32>[A : iu32, B : iu32, N : num](vec<N, A>, vec<N, B>) -> vec<N, R> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(176),
+ },
+ {
+ /* [79] */
/* fn sdot(u32, u32, u32) -> i32 */
/* num overloads */ 1,
/* overloads */ OverloadIndex(195),
},
{
- /* [79] */
+ /* [80] */
/* fn udot(u32, u32, u32) -> u32 */
/* num overloads */ 1,
/* overloads */ OverloadIndex(196),
},
{
- /* [80] */
+ /* [81] */
/* 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(197),
},
{
- /* [81] */
+ /* [82] */
/* 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(198),
},
{
- /* [82] */
+ /* [83] */
/* 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(199),
diff --git a/src/tint/lang/spirv/reader/lower/builtins.cc b/src/tint/lang/spirv/reader/lower/builtins.cc
index 77cb35c..28e4293 100644
--- a/src/tint/lang/spirv/reader/lower/builtins.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins.cc
@@ -197,6 +197,9 @@
case spirv::BuiltinFn::kShiftLeftLogical:
ShiftLeftLogical(builtin);
break;
+ case spirv::BuiltinFn::kShiftRightLogical:
+ ShiftRightLogical(builtin);
+ break;
default:
TINT_UNREACHABLE() << "unknown spirv builtin: " << builtin->Func();
}
@@ -750,6 +753,30 @@
call->Destroy();
}
+ void ShiftRightLogical(spirv::ir::BuiltinCall* call) {
+ const auto& args = call->Args();
+
+ b.InsertBefore(call, [&] {
+ auto* base = args[0];
+ auto* shift = args[1];
+
+ auto* u_ty = ty.MatchWidth(ty.u32(), base->Type());
+ if (!base->Type()->IsUnsignedIntegerScalarOrVector()) {
+ base = b.Bitcast(u_ty, base)->Result(0);
+ }
+ if (!shift->Type()->IsUnsignedIntegerScalarOrVector()) {
+ shift = b.Bitcast(u_ty, shift)->Result(0);
+ }
+
+ auto* bin = b.Binary(core::BinaryOp::kShiftRight, u_ty, base, shift)->Result(0);
+ if (u_ty != call->Result(0)->Type()) {
+ bin = b.Bitcast(call->Result(0)->Type(), bin)->Result(0);
+ }
+ call->Result(0)->ReplaceAllUsesWith(bin);
+ });
+ call->Destroy();
+ }
+
void Inverse(spirv::ir::BuiltinCall* call) {
auto* arg = call->Args()[0];
auto* mat_ty = arg->Type()->As<core::type::Matrix>();
diff --git a/src/tint/lang/spirv/reader/lower/builtins_test.cc b/src/tint/lang/spirv/reader/lower/builtins_test.cc
index 7d43f1d..cbb44c5 100644
--- a/src/tint/lang/spirv/reader/lower/builtins_test.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins_test.cc
@@ -7413,5 +7413,549 @@
EXPECT_EQ(expect, str());
}
+TEST_F(SpirvReader_BuiltinsTest, ShiftRightLogical_Scalar_UnsignedUnsigned_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kShiftRightLogical,
+ 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.shift_right_logical<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 = shr 8u, 9u
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, ShiftRightLogical_Scalar_UnsignedSigned_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kShiftRightLogical,
+ 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.shift_right_logical<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 = shr 8u, %2
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, ShiftRightLogical_Scalar_SignedUnsigned_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kShiftRightLogical,
+ Vector{ty.u32()}, 1_i, 9_u);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.shift_right_logical<u32> 1i, 9u
+ 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 = shr %2, 9u
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, ShiftRightLogical_Scalar_SignedSigned_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kShiftRightLogical,
+ 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.shift_right_logical<u32> 1i, 2i
+ 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 = bitcast 2i
+ %4:u32 = shr %2, %3
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, ShiftRightLogical_Scalar_UnsignedUnsigned_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kShiftRightLogical,
+ 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.shift_right_logical<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 = shr 8u, 9u
+ %3:i32 = bitcast %2
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, ShiftRightLogical_Scalar_UnsignedSigned_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kShiftRightLogical,
+ 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.shift_right_logical<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 = shr 8u, %2
+ %4:i32 = bitcast %3
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, ShiftRightLogical_Scalar_SignedUnsigned_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kShiftRightLogical,
+ Vector{ty.i32()}, 1_i, 9_u);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.shift_right_logical<i32> 1i, 9u
+ 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 = shr %2, 9u
+ %4:i32 = bitcast %3
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, ShiftRightLogical_Scalar_SignedSigned_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kShiftRightLogical,
+ 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.shift_right_logical<i32> 1i, 2i
+ 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 = bitcast 2i
+ %4:u32 = shr %2, %3
+ %5:i32 = bitcast %4
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, ShiftRightLogical_Vector_UnsignedUnsigned_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), spirv::BuiltinFn::kShiftRightLogical,
+ 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.shift_right_logical<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> = shr vec2<u32>(8u), vec2<u32>(9u)
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, ShiftRightLogical_Vector_UnsignedSigned_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), spirv::BuiltinFn::kShiftRightLogical,
+ 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.shift_right_logical<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> = shr vec2<u32>(8u), %2
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, ShiftRightLogical_Vector_SignedUnsigned_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), spirv::BuiltinFn::kShiftRightLogical,
+ Vector{ty.u32()}, b.Splat<vec2<i32>>(1_i),
+ 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.shift_right_logical<u32> vec2<i32>(1i), 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> = bitcast vec2<i32>(1i)
+ %3:vec2<u32> = shr %2, vec2<u32>(9u)
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, ShiftRightLogical_Vector_SignedSigned_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), spirv::BuiltinFn::kShiftRightLogical,
+ 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.shift_right_logical<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<u32> = bitcast vec2<i32>(1i)
+ %3:vec2<u32> = bitcast vec2<i32>(2i)
+ %4:vec2<u32> = shr %2, %3
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, ShiftRightLogical_Vector_UnsignedUnsigned_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), spirv::BuiltinFn::kShiftRightLogical,
+ 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.shift_right_logical<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> = shr vec2<u32>(8u), vec2<u32>(9u)
+ %3:vec2<i32> = bitcast %2
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, ShiftRightLogical_Vector_UnsignedSigned_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), spirv::BuiltinFn::kShiftRightLogical,
+ 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.shift_right_logical<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> = shr vec2<u32>(8u), %2
+ %4:vec2<i32> = bitcast %3
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, ShiftRightLogical_Vector_SignedUnsigned_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), spirv::BuiltinFn::kShiftRightLogical,
+ Vector{ty.i32()}, b.Splat<vec2<i32>>(1_i),
+ 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.shift_right_logical<i32> vec2<i32>(1i), 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> = bitcast vec2<i32>(1i)
+ %3:vec2<u32> = shr %2, vec2<u32>(9u)
+ %4:vec2<i32> = bitcast %3
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, ShiftRightLogical_Vector_SignedSigned_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), spirv::BuiltinFn::kShiftRightLogical,
+ 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.shift_right_logical<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<u32> = bitcast vec2<i32>(1i)
+ %3:vec2<u32> = bitcast vec2<i32>(2i)
+ %4:vec2<u32> = shr %2, %3
+ %5:vec2<i32> = bitcast %4
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
} // namespace
} // namespace tint::spirv::reader::lower
diff --git a/src/tint/lang/spirv/reader/parser/bit_test.cc b/src/tint/lang/spirv/reader/parser/bit_test.cc
index d8503bb..f472caf 100644
--- a/src/tint/lang/spirv/reader/parser/bit_test.cc
+++ b/src/tint/lang/spirv/reader/parser/bit_test.cc
@@ -30,15 +30,16 @@
namespace tint::spirv::reader {
namespace {
-struct SpirvBitwiseParam {
- std::string name;
+struct SpirvBitParam {
+ std::string spv_name;
+ std::string ir_name;
};
-[[maybe_unused]] inline std::ostream& operator<<(std::ostream& out, SpirvBitwiseParam c) {
- out << c.name;
+[[maybe_unused]] inline std::ostream& operator<<(std::ostream& out, SpirvBitParam c) {
+ out << c.spv_name;
return out;
}
-using SpirvParser_BitwiseTest = SpirvParserTestWithParam<SpirvBitwiseParam>;
+using SpirvParser_BitwiseTest = SpirvParserTestWithParam<SpirvBitParam>;
TEST_P(SpirvParser_BitwiseTest, Scalar_SignedSigned_Signed) {
auto& params = GetParam();
@@ -56,7 +57,7 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
%1 = OpBitwise)" +
- params.name + R"( %int %one %two
+ params.spv_name + R"( %int %one %two
OpReturn
OpFunctionEnd
)",
@@ -64,7 +65,7 @@
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:i32 = spirv.bitwise)" +
- params.name + R"(<i32> 1i, 2i
+ params.ir_name + R"(<i32> 1i, 2i
ret
}
}
@@ -90,7 +91,7 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
%1 = OpBitwise)" +
- params.name + R"( %int %one %eight
+ params.spv_name + R"( %int %one %eight
OpReturn
OpFunctionEnd
)",
@@ -98,7 +99,7 @@
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:i32 = spirv.bitwise)" +
- params.name + R"(<i32> 1i, 8u
+ params.ir_name + R"(<i32> 1i, 8u
ret
}
}
@@ -124,7 +125,7 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
%1 = OpBitwise)" +
- params.name + R"( %int %eight %one
+ params.spv_name + R"( %int %eight %one
OpReturn
OpFunctionEnd
)",
@@ -132,7 +133,7 @@
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:i32 = spirv.bitwise)" +
- params.name + R"(<i32> 8u, 1i
+ params.ir_name + R"(<i32> 8u, 1i
ret
}
}
@@ -158,7 +159,7 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
%1 = OpBitwise)" +
- params.name + R"( %int %eight %nine
+ params.spv_name + R"( %int %eight %nine
OpReturn
OpFunctionEnd
)",
@@ -166,7 +167,7 @@
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:i32 = spirv.bitwise)" +
- params.name + R"(<i32> 8u, 9u
+ params.ir_name + R"(<i32> 8u, 9u
ret
}
}
@@ -192,7 +193,7 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
%1 = OpBitwise)" +
- params.name + R"( %uint %eight %nine
+ params.spv_name + R"( %uint %eight %nine
OpReturn
OpFunctionEnd
)",
@@ -200,7 +201,7 @@
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:u32 = spirv.bitwise)" +
- params.name + R"(<u32> 8u, 9u
+ params.ir_name + R"(<u32> 8u, 9u
ret
}
}
@@ -226,7 +227,7 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
%1 = OpBitwise)" +
- params.name + R"( %uint %eight %one
+ params.spv_name + R"( %uint %eight %one
OpReturn
OpFunctionEnd
)",
@@ -234,7 +235,7 @@
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:u32 = spirv.bitwise)" +
- params.name + R"(<u32> 8u, 1i
+ params.ir_name + R"(<u32> 8u, 1i
ret
}
}
@@ -260,7 +261,7 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
%1 = OpBitwise)" +
- params.name + R"( %uint %one %eight
+ params.spv_name + R"( %uint %one %eight
OpReturn
OpFunctionEnd
)",
@@ -268,7 +269,7 @@
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:u32 = spirv.bitwise)" +
- params.name + R"(<u32> 1i, 8u
+ params.ir_name + R"(<u32> 1i, 8u
ret
}
}
@@ -294,7 +295,7 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
%1 = OpBitwise)" +
- params.name + R"( %uint %one %two
+ params.spv_name + R"( %uint %one %two
OpReturn
OpFunctionEnd
)",
@@ -302,7 +303,7 @@
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:u32 = spirv.bitwise)" +
- params.name + R"(<u32> 1i, 2i
+ params.ir_name + R"(<u32> 1i, 2i
ret
}
}
@@ -334,7 +335,7 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
%1 = OpBitwise)" +
- params.name + R"( %v2int %v2one %v2two
+ params.spv_name + R"( %v2int %v2one %v2two
OpReturn
OpFunctionEnd
)",
@@ -342,7 +343,7 @@
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:vec2<i32> = spirv.bitwise)" +
- params.name + R"(<i32> vec2<i32>(1i), vec2<i32>(2i)
+ params.ir_name + R"(<i32> vec2<i32>(1i), vec2<i32>(2i)
ret
}
}
@@ -374,7 +375,7 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
%1 = OpBitwise)" +
- params.name + R"( %v2int %v2one %v2eight
+ params.spv_name + R"( %v2int %v2one %v2eight
OpReturn
OpFunctionEnd
)",
@@ -382,7 +383,7 @@
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:vec2<i32> = spirv.bitwise)" +
- params.name + R"(<i32> vec2<i32>(1i), vec2<u32>(8u)
+ params.ir_name + R"(<i32> vec2<i32>(1i), vec2<u32>(8u)
ret
}
}
@@ -414,7 +415,7 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
%1 = OpBitwise)" +
- params.name + R"( %v2int %v2eight %v2one
+ params.spv_name + R"( %v2int %v2eight %v2one
OpReturn
OpFunctionEnd
)",
@@ -422,7 +423,7 @@
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:vec2<i32> = spirv.bitwise)" +
- params.name + R"(<i32> vec2<u32>(8u), vec2<i32>(1i)
+ params.ir_name + R"(<i32> vec2<u32>(8u), vec2<i32>(1i)
ret
}
}
@@ -454,7 +455,7 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
%1 = OpBitwise)" +
- params.name + R"( %v2int %v2eight %v2nine
+ params.spv_name + R"( %v2int %v2eight %v2nine
OpReturn
OpFunctionEnd
)",
@@ -462,7 +463,7 @@
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:vec2<i32> = spirv.bitwise)" +
- params.name + R"(<i32> vec2<u32>(8u), vec2<u32>(9u)
+ params.ir_name + R"(<i32> vec2<u32>(8u), vec2<u32>(9u)
ret
}
}
@@ -494,7 +495,7 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
%1 = OpBitwise)" +
- params.name + R"( %v2uint %v2eight %v2nine
+ params.spv_name + R"( %v2uint %v2eight %v2nine
OpReturn
OpFunctionEnd
)",
@@ -502,7 +503,7 @@
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:vec2<u32> = spirv.bitwise)" +
- params.name + R"(<u32> vec2<u32>(8u), vec2<u32>(9u)
+ params.ir_name + R"(<u32> vec2<u32>(8u), vec2<u32>(9u)
ret
}
}
@@ -534,7 +535,7 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
%1 = OpBitwise)" +
- params.name + R"( %v2uint %v2eight %v2one
+ params.spv_name + R"( %v2uint %v2eight %v2one
OpReturn
OpFunctionEnd
)",
@@ -542,7 +543,7 @@
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:vec2<u32> = spirv.bitwise)" +
- params.name + R"(<u32> vec2<u32>(8u), vec2<i32>(1i)
+ params.ir_name + R"(<u32> vec2<u32>(8u), vec2<i32>(1i)
ret
}
}
@@ -574,7 +575,7 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
%1 = OpBitwise)" +
- params.name + R"( %v2uint %v2one %v2eight
+ params.spv_name + R"( %v2uint %v2one %v2eight
OpReturn
OpFunctionEnd
)",
@@ -582,7 +583,7 @@
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:vec2<u32> = spirv.bitwise)" +
- params.name + R"(<u32> vec2<i32>(1i), vec2<u32>(8u)
+ params.ir_name + R"(<u32> vec2<i32>(1i), vec2<u32>(8u)
ret
}
}
@@ -614,7 +615,7 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
%1 = OpBitwise)" +
- params.name + R"( %v2uint %v2one %v2two
+ params.spv_name + R"( %v2uint %v2one %v2two
OpReturn
OpFunctionEnd
)",
@@ -622,7 +623,7 @@
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:vec2<u32> = spirv.bitwise)" +
- params.name + R"(<u32> vec2<i32>(1i), vec2<i32>(2i)
+ params.ir_name + R"(<u32> vec2<i32>(1i), vec2<i32>(2i)
ret
}
}
@@ -631,11 +632,13 @@
INSTANTIATE_TEST_SUITE_P(SpirvParser,
SpirvParser_BitwiseTest,
- testing::Values(SpirvBitwiseParam{"And"}, //
- SpirvBitwiseParam{"Or"}, //
- SpirvBitwiseParam{"Xor"}));
+ testing::Values(SpirvBitParam{"And", "And"}, //
+ SpirvBitParam{"Or", "Or"}, //
+ SpirvBitParam{"Xor", "Xor"}));
-TEST_F(SpirvParserTest, ShiftLeftLogical_Scalar_UnsignedUnsigned_Unsigned) {
+using SpirvParser_ShiftTest = SpirvParserTestWithParam<SpirvBitParam>;
+TEST_P(SpirvParser_ShiftTest, Scalar_UnsignedUnsigned_Unsigned) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -658,21 +661,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpShiftLeftLogical %uint %eight %nine
+ %1 = Op)" +
+ params.spv_name + R"( %uint %eight %nine
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:u32 = spirv.shift_left_logical<u32> 8u, 9u
+ %2:u32 = spirv.)" +
+ params.ir_name + R"(<u32> 8u, 9u
ret
}
}
)");
}
-TEST_F(SpirvParserTest, ShiftLeftLogical_Scalar_UnsignedSigned_Unsigned) {
+TEST_P(SpirvParser_ShiftTest, Scalar_UnsignedSigned_Unsigned) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -695,21 +701,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpShiftLeftLogical %uint %eight %one
+ %1 = Op)" +
+ params.spv_name + R"( %uint %eight %one
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:u32 = spirv.shift_left_logical<u32> 8u, 1i
+ %2:u32 = spirv.)" +
+ params.ir_name + R"(<u32> 8u, 1i
ret
}
}
)");
}
-TEST_F(SpirvParserTest, ShiftLeftLogical_Scalar_SignedUnsigned_Unsigned) {
+TEST_P(SpirvParser_ShiftTest, Scalar_SignedUnsigned_Unsigned) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -732,21 +741,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpShiftLeftLogical %uint %one %nine
+ %1 = Op)" +
+ params.spv_name + R"( %uint %one %nine
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:u32 = spirv.shift_left_logical<u32> 1i, 9u
+ %2:u32 = spirv.)" +
+ params.ir_name + R"(<u32> 1i, 9u
ret
}
}
)");
}
-TEST_F(SpirvParserTest, ShiftLeftLogical_Scalar_SignedSigned_Unsigned) {
+TEST_P(SpirvParser_ShiftTest, Scalar_SignedSigned_Unsigned) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -769,21 +781,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpShiftLeftLogical %uint %one %two
+ %1 = Op)" +
+ params.spv_name + R"( %uint %one %two
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:u32 = spirv.shift_left_logical<u32> 1i, 2i
+ %2:u32 = spirv.)" +
+ params.ir_name + R"(<u32> 1i, 2i
ret
}
}
)");
}
-TEST_F(SpirvParserTest, ShiftLeftLogical_Scalar_UnsignedUnsigned_Signed) {
+TEST_P(SpirvParser_ShiftTest, Scalar_UnsignedUnsigned_Signed) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -806,21 +821,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpShiftLeftLogical %int %eight %nine
+ %1 = Op)" +
+ params.spv_name + R"( %int %eight %nine
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:i32 = spirv.shift_left_logical<i32> 8u, 9u
+ %2:i32 = spirv.)" +
+ params.ir_name + R"(<i32> 8u, 9u
ret
}
}
)");
}
-TEST_F(SpirvParserTest, ShiftLeftLogical_Scalar_UnsignedSigned_Signed) {
+TEST_P(SpirvParser_ShiftTest, Scalar_UnsignedSigned_Signed) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -843,21 +861,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpShiftLeftLogical %int %eight %one
+ %1 = Op)" +
+ params.spv_name + R"( %int %eight %one
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:i32 = spirv.shift_left_logical<i32> 8u, 1i
+ %2:i32 = spirv.)" +
+ params.ir_name + R"(<i32> 8u, 1i
ret
}
}
)");
}
-TEST_F(SpirvParserTest, ShiftLeftLogical_Scalar_SignedUnsigned_Signed) {
+TEST_P(SpirvParser_ShiftTest, Scalar_SignedUnsigned_Signed) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -880,21 +901,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpShiftLeftLogical %int %one %nine
+ %1 = Op)" +
+ params.spv_name + R"( %int %one %nine
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:i32 = spirv.shift_left_logical<i32> 1i, 9u
+ %2:i32 = spirv.)" +
+ params.ir_name + R"(<i32> 1i, 9u
ret
}
}
)");
}
-TEST_F(SpirvParserTest, ShiftLeftLogical_Scalar_SignedSigned_Signed) {
+TEST_P(SpirvParser_ShiftTest, Scalar_SignedSigned_Signed) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -917,21 +941,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpShiftLeftLogical %int %one %two
+ %1 = Op)" +
+ params.spv_name + R"( %int %one %two
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:i32 = spirv.shift_left_logical<i32> 1i, 2i
+ %2:i32 = spirv.)" +
+ params.ir_name + R"(<i32> 1i, 2i
ret
}
}
)");
}
-TEST_F(SpirvParserTest, ShiftLeftLogical_Vector_UnsignedUnsigned_Unsigned) {
+TEST_P(SpirvParser_ShiftTest, Vector_UnsignedUnsigned_Unsigned) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -954,21 +981,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpShiftLeftLogical %v2uint %v2eight %v2nine
+ %1 = Op)" +
+ params.spv_name + R"( %v2uint %v2eight %v2nine
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:vec2<u32> = spirv.shift_left_logical<u32> vec2<u32>(8u), vec2<u32>(9u)
+ %2:vec2<u32> = spirv.)" +
+ params.ir_name + R"(<u32> vec2<u32>(8u), vec2<u32>(9u)
ret
}
}
)");
}
-TEST_F(SpirvParserTest, ShiftLeftLogical_Vector_UnsignedSigned_Unsigned) {
+TEST_P(SpirvParser_ShiftTest, Vector_UnsignedSigned_Unsigned) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -991,21 +1021,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpShiftLeftLogical %v2uint %v2eight %v2one
+ %1 = Op)" +
+ params.spv_name + R"( %v2uint %v2eight %v2one
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:vec2<u32> = spirv.shift_left_logical<u32> vec2<u32>(8u), vec2<i32>(1i)
+ %2:vec2<u32> = spirv.)" +
+ params.ir_name + R"(<u32> vec2<u32>(8u), vec2<i32>(1i)
ret
}
}
)");
}
-TEST_F(SpirvParserTest, ShiftLeftLogical_Vector_SignedUnsigned_Unsigned) {
+TEST_P(SpirvParser_ShiftTest, Vector_SignedUnsigned_Unsigned) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -1028,21 +1061,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpShiftLeftLogical %v2uint %v2one %v2nine
+ %1 = Op)" +
+ params.spv_name + R"( %v2uint %v2one %v2nine
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:vec2<u32> = spirv.shift_left_logical<u32> vec2<i32>(1i), vec2<u32>(9u)
+ %2:vec2<u32> = spirv.)" +
+ params.ir_name + R"(<u32> vec2<i32>(1i), vec2<u32>(9u)
ret
}
}
)");
}
-TEST_F(SpirvParserTest, ShiftLeftLogical_Vector_SignedSigned_Unsigned) {
+TEST_P(SpirvParser_ShiftTest, Vector_SignedSigned_Unsigned) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -1065,21 +1101,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpShiftLeftLogical %v2uint %v2one %v2two
+ %1 = Op)" +
+ params.spv_name + R"( %v2uint %v2one %v2two
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:vec2<u32> = spirv.shift_left_logical<u32> vec2<i32>(1i), vec2<i32>(2i)
+ %2:vec2<u32> = spirv.)" +
+ params.ir_name + R"(<u32> vec2<i32>(1i), vec2<i32>(2i)
ret
}
}
)");
}
-TEST_F(SpirvParserTest, ShiftLeftLogical_Vector_UnsignedUnsigned_Signed) {
+TEST_P(SpirvParser_ShiftTest, Vector_UnsignedUnsigned_Signed) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -1102,21 +1141,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpShiftLeftLogical %v2int %v2eight %v2nine
+ %1 = Op)" +
+ params.spv_name + R"( %v2int %v2eight %v2nine
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:vec2<i32> = spirv.shift_left_logical<i32> vec2<u32>(8u), vec2<u32>(9u)
+ %2:vec2<i32> = spirv.)" +
+ params.ir_name + R"(<i32> vec2<u32>(8u), vec2<u32>(9u)
ret
}
}
)");
}
-TEST_F(SpirvParserTest, ShiftLeftLogical_Vector_UnsignedSigned_Signed) {
+TEST_P(SpirvParser_ShiftTest, Vector_UnsignedSigned_Signed) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -1139,21 +1181,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpShiftLeftLogical %v2int %v2eight %v2one
+ %1 = Op)" +
+ params.spv_name + R"( %v2int %v2eight %v2one
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:vec2<i32> = spirv.shift_left_logical<i32> vec2<u32>(8u), vec2<i32>(1i)
+ %2:vec2<i32> = spirv.)" +
+ params.ir_name + R"(<i32> vec2<u32>(8u), vec2<i32>(1i)
ret
}
}
)");
}
-TEST_F(SpirvParserTest, ShiftLeftLogical_Vector_SignedUnsigned_Signed) {
+TEST_P(SpirvParser_ShiftTest, Vector_SignedUnsigned_Signed) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -1176,21 +1221,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpShiftLeftLogical %v2int %v2one %v2nine
+ %1 = Op)" +
+ params.spv_name + R"( %v2int %v2one %v2nine
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:vec2<i32> = spirv.shift_left_logical<i32> vec2<i32>(1i), vec2<u32>(9u)
+ %2:vec2<i32> = spirv.)" +
+ params.ir_name + R"(<i32> vec2<i32>(1i), vec2<u32>(9u)
ret
}
}
)");
}
-TEST_F(SpirvParserTest, ShiftLeftLogical_Vector_SignedSigned_Signed) {
+TEST_P(SpirvParser_ShiftTest, Vector_SignedSigned_Signed) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -1213,19 +1261,27 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpShiftLeftLogical %v2int %v2one %v2two
+ %1 = Op)" +
+ params.spv_name + R"( %v2int %v2one %v2two
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:vec2<i32> = spirv.shift_left_logical<i32> vec2<i32>(1i), vec2<i32>(2i)
+ %2:vec2<i32> = spirv.)" +
+ params.ir_name + R"(<i32> vec2<i32>(1i), vec2<i32>(2i)
ret
}
}
)");
}
+INSTANTIATE_TEST_SUITE_P(
+ SpirvParser,
+ SpirvParser_ShiftTest,
+ testing::Values(SpirvBitParam{"ShiftLeftLogical", "shift_left_logical"}, //
+ SpirvBitParam{"ShiftRightLogical", "shift_right_logical"}));
+
} // 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 d4f7078..9e98306 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -886,6 +886,9 @@
case spv::Op::OpShiftLeftLogical:
EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kShiftLeftLogical);
break;
+ case spv::Op::OpShiftRightLogical:
+ EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kShiftRightLogical);
+ break;
default:
TINT_UNIMPLEMENTED()
<< "unhandled SPIR-V instruction: " << static_cast<uint32_t>(inst.opcode());
diff --git a/src/tint/lang/spirv/spirv.def b/src/tint/lang/spirv/spirv.def
index 04c7ac3c..70186ab 100644
--- a/src/tint/lang/spirv/spirv.def
+++ b/src/tint/lang/spirv/spirv.def
@@ -453,6 +453,8 @@
implicit(A: iu32, B: iu32) fn shift_left_logical<R: iu32>(A, B) -> R
implicit(A: iu32, B: iu32, N: num) fn shift_left_logical<R: iu32>(vec<N, A>, vec<N, B>) -> vec<N, R>
+implicit(A: iu32, B: iu32) fn shift_right_logical<R: iu32>(A, B) -> R
+implicit(A: iu32, B: iu32, N: num) fn shift_right_logical<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 c300b1c..50011b0 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -1575,6 +1575,9 @@
case BuiltinFn::kShiftLeftLogical:
op = spv::Op::OpShiftLeftLogical;
break;
+ case BuiltinFn::kShiftRightLogical:
+ op = spv::Op::OpShiftRightLogical;
+ break;
case spirv::BuiltinFn::kNone:
TINT_ICE() << "undefined spirv ir function";
}
