[spirv-reader][ir] Convert OpBitwiseOr
Add conversion for OpBitwiseOr, converting the arguments such that they
match the type of the first parameter.
Bug: 391486220
Change-Id: Ib5b01b1ac75df58557b994d994a58cbb341001de
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/226895
Auto-Submit: dan sinclair <dsinclair@chromium.org>
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 ebb1b3b..6bb3d17 100644
--- a/src/tint/lang/spirv/builtin_fn.cc
+++ b/src/tint/lang/spirv/builtin_fn.cc
@@ -172,6 +172,8 @@
return "convertUToF";
case BuiltinFn::kBitwiseAnd:
return "bitwiseAnd";
+ case BuiltinFn::kBitwiseOr:
+ return "bitwiseOr";
case BuiltinFn::kSdot:
return "sdot";
case BuiltinFn::kUdot:
@@ -266,6 +268,7 @@
case BuiltinFn::kConvertSToF:
case BuiltinFn::kConvertUToF:
case BuiltinFn::kBitwiseAnd:
+ case BuiltinFn::kBitwiseOr:
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 87210f2..0862f38 100644
--- a/src/tint/lang/spirv/builtin_fn.cc.tmpl
+++ b/src/tint/lang/spirv/builtin_fn.cc.tmpl
@@ -107,6 +107,7 @@
case BuiltinFn::kConvertSToF:
case BuiltinFn::kConvertUToF:
case BuiltinFn::kBitwiseAnd:
+ case BuiltinFn::kBitwiseOr:
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 43e52cc..fccebad 100644
--- a/src/tint/lang/spirv/builtin_fn.h
+++ b/src/tint/lang/spirv/builtin_fn.h
@@ -113,6 +113,7 @@
kConvertSToF,
kConvertUToF,
kBitwiseAnd,
+ kBitwiseOr,
kSdot,
kUdot,
kCooperativeMatrixLoad,
diff --git a/src/tint/lang/spirv/intrinsic/data.cc b/src/tint/lang/spirv/intrinsic/data.cc
index ac8af96..efb10fb 100644
--- a/src/tint/lang/spirv/intrinsic/data.cc
+++ b/src/tint/lang/spirv/intrinsic/data.cc
@@ -6725,30 +6725,37 @@
},
{
/* [65] */
+ /* fn bitwiseOr<R : iu32>[A : iu32, B : iu32](A, B) -> R */
+ /* fn bitwiseOr<R : iu32>[A : iu32, B : iu32, N : num](vec<N, A>, vec<N, B>) -> vec<N, R> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(176),
+ },
+ {
+ /* [66] */
/* fn sdot(u32, u32, u32) -> i32 */
/* num overloads */ 1,
/* overloads */ OverloadIndex(193),
},
{
- /* [66] */
+ /* [67] */
/* fn udot(u32, u32, u32) -> u32 */
/* num overloads */ 1,
/* overloads */ OverloadIndex(194),
},
{
- /* [67] */
+ /* [68] */
/* 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),
},
{
- /* [68] */
+ /* [69] */
/* 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),
},
{
- /* [69] */
+ /* [70] */
/* 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 e7c7b88..40095c3 100644
--- a/src/tint/lang/spirv/reader/lower/builtins.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins.cc
@@ -158,6 +158,9 @@
case spirv::BuiltinFn::kBitwiseAnd:
BitwiseAnd(builtin);
break;
+ case spirv::BuiltinFn::kBitwiseOr:
+ BitwiseOr(builtin);
+ break;
default:
TINT_UNREACHABLE() << "unknown spirv builtin: " << builtin->Func();
}
@@ -231,6 +234,9 @@
void BitwiseAnd(spirv::ir::BuiltinCall* call) {
EmitBinaryWrappedAsFirstArg(call, core::BinaryOp::kAnd);
}
+ void BitwiseOr(spirv::ir::BuiltinCall* call) {
+ EmitBinaryWrappedAsFirstArg(call, core::BinaryOp::kOr);
+ }
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 34f4beb..b0da81d 100644
--- a/src/tint/lang/spirv/reader/lower/builtins_test.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins_test.cc
@@ -38,9 +38,9 @@
using namespace tint::core::fluent_types; // NOLINT
using namespace tint::core::number_suffixes; // NOLINT
-using SpirvParser_BuiltinsTest = core::ir::transform::TransformTest;
+using SpirvReader_BuiltinsTest = core::ir::transform::TransformTest;
-TEST_F(SpirvParser_BuiltinsTest, Normalize_Scalar) {
+TEST_F(SpirvReader_BuiltinsTest, Normalize_Scalar) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -70,7 +70,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Normalize_Vector) {
+TEST_F(SpirvReader_BuiltinsTest, Normalize_Vector) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -101,7 +101,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Inverse_Mat2x2f) {
+TEST_F(SpirvReader_BuiltinsTest, Inverse_Mat2x2f) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -148,7 +148,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Inverse_Mat2x2h) {
+TEST_F(SpirvReader_BuiltinsTest, Inverse_Mat2x2h) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -195,7 +195,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Inverse_Mat3x3f) {
+TEST_F(SpirvReader_BuiltinsTest, Inverse_Mat3x3f) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -272,7 +272,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Inverse_Mat3x3h) {
+TEST_F(SpirvReader_BuiltinsTest, Inverse_Mat3x3h) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -349,7 +349,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Inverse_Mat4x4f) {
+TEST_F(SpirvReader_BuiltinsTest, Inverse_Mat4x4f) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -545,7 +545,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Inverse_Mat4x4h) {
+TEST_F(SpirvReader_BuiltinsTest, Inverse_Mat4x4h) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -741,7 +741,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, SSign_Scalar_UnsignedArg) {
+TEST_F(SpirvReader_BuiltinsTest, SSign_Scalar_UnsignedArg) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -774,7 +774,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, SSign_Scalar_UnsignedResult) {
+TEST_F(SpirvReader_BuiltinsTest, SSign_Scalar_UnsignedResult) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -807,7 +807,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, SSign_Scalar_UnsignedArgAndResult) {
+TEST_F(SpirvReader_BuiltinsTest, SSign_Scalar_UnsignedArgAndResult) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -841,7 +841,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, SSign_Scalar_SignedArgAndResult) {
+TEST_F(SpirvReader_BuiltinsTest, SSign_Scalar_SignedArgAndResult) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -873,7 +873,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, SSign_Vector_UnsignedArg) {
+TEST_F(SpirvReader_BuiltinsTest, SSign_Vector_UnsignedArg) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -906,7 +906,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, SSign_Vector_UnsignedResult) {
+TEST_F(SpirvReader_BuiltinsTest, SSign_Vector_UnsignedResult) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -939,7 +939,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, SSign_Vector_UnsignedArgAndResult) {
+TEST_F(SpirvReader_BuiltinsTest, SSign_Vector_UnsignedArgAndResult) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -973,7 +973,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, SSign_Vector_SignedArgAndResult) {
+TEST_F(SpirvReader_BuiltinsTest, SSign_Vector_SignedArgAndResult) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -1015,10 +1015,10 @@
return out;
}
-using SpirvParser_BuiltinsTest_OneParamSigned =
+using SpirvReader_BuiltinsTest_OneParamSigned =
core::ir::transform::TransformTestWithParam<SpirvReaderParams>;
-TEST_P(SpirvParser_BuiltinsTest_OneParamSigned, UnsignedToUnsigned) {
+TEST_P(SpirvReader_BuiltinsTest_OneParamSigned, UnsignedToUnsigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -1065,7 +1065,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsTest_OneParamSigned, UnsignedToSigned) {
+TEST_P(SpirvReader_BuiltinsTest_OneParamSigned, UnsignedToSigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -1110,7 +1110,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsTest_OneParamSigned, SignedToSigned) {
+TEST_P(SpirvReader_BuiltinsTest_OneParamSigned, SignedToSigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -1153,7 +1153,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsTest_OneParamSigned, SignedToUnsigned) {
+TEST_P(SpirvReader_BuiltinsTest_OneParamSigned, SignedToUnsigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -1199,15 +1199,15 @@
}
INSTANTIATE_TEST_SUITE_P(SpirvReader,
- SpirvParser_BuiltinsTest_OneParamSigned,
+ SpirvReader_BuiltinsTest_OneParamSigned,
::testing::Values(SpirvReaderParams{spirv::BuiltinFn::kAbs, "abs", "abs"},
SpirvReaderParams{spirv::BuiltinFn::kFindSMsb,
"findSMsb", "firstLeadingBit"}));
-using SpirvParser_BuiltinsTest_OneParamUnsigned =
+using SpirvReader_BuiltinsTest_OneParamUnsigned =
core::ir::transform::TransformTestWithParam<SpirvReaderParams>;
-TEST_P(SpirvParser_BuiltinsTest_OneParamUnsigned, UnsignedToUnsigned) {
+TEST_P(SpirvReader_BuiltinsTest_OneParamUnsigned, UnsignedToUnsigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -1250,7 +1250,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsTest_OneParamUnsigned, UnsignedToSigned) {
+TEST_P(SpirvReader_BuiltinsTest_OneParamUnsigned, UnsignedToSigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -1295,7 +1295,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsTest_OneParamUnsigned, SignedToSigned) {
+TEST_P(SpirvReader_BuiltinsTest_OneParamUnsigned, SignedToSigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -1342,7 +1342,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsTest_OneParamUnsigned, SignedToUnsigned) {
+TEST_P(SpirvReader_BuiltinsTest_OneParamUnsigned, SignedToUnsigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -1388,14 +1388,14 @@
}
INSTANTIATE_TEST_SUITE_P(SpirvReader,
- SpirvParser_BuiltinsTest_OneParamUnsigned,
+ SpirvReader_BuiltinsTest_OneParamUnsigned,
::testing::Values(SpirvReaderParams{spirv::BuiltinFn::kFindUMsb,
"findUMsb", "firstLeadingBit"}));
-using SpirvParser_BuiltinsTest_TwoParamSigned =
+using SpirvReader_BuiltinsTest_TwoParamSigned =
core::ir::transform::TransformTestWithParam<SpirvReaderParams>;
-TEST_P(SpirvParser_BuiltinsTest_TwoParamSigned, UnsignedToUnsigned) {
+TEST_P(SpirvReader_BuiltinsTest_TwoParamSigned, UnsignedToUnsigned) {
auto& params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -1444,7 +1444,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsTest_TwoParamSigned, SignedToSigned) {
+TEST_P(SpirvReader_BuiltinsTest_TwoParamSigned, SignedToSigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -1487,7 +1487,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsTest_TwoParamSigned, MixedToUnsigned) {
+TEST_P(SpirvReader_BuiltinsTest_TwoParamSigned, MixedToUnsigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -1534,7 +1534,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsTest_TwoParamSigned, MixedToSigned) {
+TEST_P(SpirvReader_BuiltinsTest_TwoParamSigned, MixedToSigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -1580,15 +1580,15 @@
}
INSTANTIATE_TEST_SUITE_P(SpirvReader,
- SpirvParser_BuiltinsTest_TwoParamSigned,
+ SpirvReader_BuiltinsTest_TwoParamSigned,
::testing::Values(SpirvReaderParams{spirv::BuiltinFn::kSmax, "max", "max"},
SpirvReaderParams{spirv::BuiltinFn::kSmin, "min",
"min"}));
-using SpirvParser_BuiltinsTest_TwoParamUnsigned =
+using SpirvReader_BuiltinsTest_TwoParamUnsigned =
core::ir::transform::TransformTestWithParam<SpirvReaderParams>;
-TEST_P(SpirvParser_BuiltinsTest_TwoParamUnsigned, UnsignedToUnsigned) {
+TEST_P(SpirvReader_BuiltinsTest_TwoParamUnsigned, UnsignedToUnsigned) {
auto& params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -1631,7 +1631,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsTest_TwoParamUnsigned, SignedToSigned) {
+TEST_P(SpirvReader_BuiltinsTest_TwoParamUnsigned, SignedToSigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -1680,7 +1680,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsTest_TwoParamUnsigned, MixedToUnsigned) {
+TEST_P(SpirvReader_BuiltinsTest_TwoParamUnsigned, MixedToUnsigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -1725,7 +1725,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsTest_TwoParamUnsigned, MixedToSigned) {
+TEST_P(SpirvReader_BuiltinsTest_TwoParamUnsigned, MixedToSigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -1773,12 +1773,12 @@
}
INSTANTIATE_TEST_SUITE_P(SpirvReader,
- SpirvParser_BuiltinsTest_TwoParamUnsigned,
+ SpirvReader_BuiltinsTest_TwoParamUnsigned,
::testing::Values(SpirvReaderParams{spirv::BuiltinFn::kUmax, "max", "max"},
SpirvReaderParams{spirv::BuiltinFn::kUmin, "min",
"min"}));
-TEST_F(SpirvParser_BuiltinsTest, SClamp_UnsignedToUnsigned) {
+TEST_F(SpirvReader_BuiltinsTest, SClamp_UnsignedToUnsigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -1825,7 +1825,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, SClamp_SignedToSigned) {
+TEST_F(SpirvReader_BuiltinsTest, SClamp_SignedToSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -1864,7 +1864,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, SClamp_MixedToUnsigned) {
+TEST_F(SpirvReader_BuiltinsTest, SClamp_MixedToUnsigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -1907,7 +1907,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, SClamp_MixedToSigned) {
+TEST_F(SpirvReader_BuiltinsTest, SClamp_MixedToSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -1950,7 +1950,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, UClamp_UnsignedToUnsigned) {
+TEST_F(SpirvReader_BuiltinsTest, UClamp_UnsignedToUnsigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -1989,7 +1989,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, UClamp_SignedToSigned) {
+TEST_F(SpirvReader_BuiltinsTest, UClamp_SignedToSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2036,7 +2036,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, UClamp_MixedToUnsigned) {
+TEST_F(SpirvReader_BuiltinsTest, UClamp_MixedToUnsigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2079,7 +2079,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, UClamp_MixedToSigned) {
+TEST_F(SpirvReader_BuiltinsTest, UClamp_MixedToSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2122,7 +2122,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, FindILsb_SignedToSigned) {
+TEST_F(SpirvReader_BuiltinsTest, FindILsb_SignedToSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2159,7 +2159,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, FindILsb_UnsignedToUnsigned) {
+TEST_F(SpirvReader_BuiltinsTest, FindILsb_UnsignedToUnsigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2196,7 +2196,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, FindILsb_SignedToUnsigned) {
+TEST_F(SpirvReader_BuiltinsTest, FindILsb_SignedToUnsigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2235,7 +2235,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, FindILsb_UnsignedToSigned) {
+TEST_F(SpirvReader_BuiltinsTest, FindILsb_UnsignedToSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2274,7 +2274,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Refract_Scalar) {
+TEST_F(SpirvReader_BuiltinsTest, Refract_Scalar) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2309,7 +2309,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Refract_Vector) {
+TEST_F(SpirvReader_BuiltinsTest, Refract_Vector) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2342,7 +2342,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, FaceForward_Scalar) {
+TEST_F(SpirvReader_BuiltinsTest, FaceForward_Scalar) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2377,7 +2377,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, FaceForward_Vector) {
+TEST_F(SpirvReader_BuiltinsTest, FaceForward_Vector) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2410,7 +2410,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Reflect_Scalar) {
+TEST_F(SpirvReader_BuiltinsTest, Reflect_Scalar) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2445,7 +2445,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Reflect_Vector) {
+TEST_F(SpirvReader_BuiltinsTest, Reflect_Vector) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2478,7 +2478,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Ldexp_ScalarSigned) {
+TEST_F(SpirvReader_BuiltinsTest, Ldexp_ScalarSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2509,7 +2509,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Ldexp_ScalarUnsigned) {
+TEST_F(SpirvReader_BuiltinsTest, Ldexp_ScalarUnsigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2541,7 +2541,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Ldexp_VectorUnsigned) {
+TEST_F(SpirvReader_BuiltinsTest, Ldexp_VectorUnsigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2574,7 +2574,7 @@
)";
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Ldexp_VectorSigned) {
+TEST_F(SpirvReader_BuiltinsTest, Ldexp_VectorSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2607,7 +2607,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Modf_Scalar) {
+TEST_F(SpirvReader_BuiltinsTest, Modf_Scalar) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2654,7 +2654,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Modf_Vector) {
+TEST_F(SpirvReader_BuiltinsTest, Modf_Vector) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2702,7 +2702,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Frexp_ScalarSigned) {
+TEST_F(SpirvReader_BuiltinsTest, Frexp_ScalarSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2749,7 +2749,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Frexp_ScalarUnSigned) {
+TEST_F(SpirvReader_BuiltinsTest, Frexp_ScalarUnSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2797,7 +2797,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Frexp_VectorSigned) {
+TEST_F(SpirvReader_BuiltinsTest, Frexp_VectorSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2845,7 +2845,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, Frexp_VectorUnsigned) {
+TEST_F(SpirvReader_BuiltinsTest, Frexp_VectorUnsigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2894,7 +2894,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitCount_Scalar_UnsignedToUnsigned) {
+TEST_F(SpirvReader_BuiltinsTest, BitCount_Scalar_UnsignedToUnsigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2926,7 +2926,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitCount_Scalar_UnsignedToSigned) {
+TEST_F(SpirvReader_BuiltinsTest, BitCount_Scalar_UnsignedToSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2959,7 +2959,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitCount_Scalar_SignedToUnsigned) {
+TEST_F(SpirvReader_BuiltinsTest, BitCount_Scalar_SignedToUnsigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -2992,7 +2992,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitCount_Scalar_SignedToSigned) {
+TEST_F(SpirvReader_BuiltinsTest, BitCount_Scalar_SignedToSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3024,7 +3024,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitCount_Vector_UnsignedToUnsigned) {
+TEST_F(SpirvReader_BuiltinsTest, BitCount_Vector_UnsignedToUnsigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3057,7 +3057,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitCount_Vector_UnsignedToSigned) {
+TEST_F(SpirvReader_BuiltinsTest, BitCount_Vector_UnsignedToSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3091,7 +3091,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitCount_Vector_SignedToUnsigned) {
+TEST_F(SpirvReader_BuiltinsTest, BitCount_Vector_SignedToUnsigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3125,7 +3125,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitCount_Vector_SignedToSigned) {
+TEST_F(SpirvReader_BuiltinsTest, BitCount_Vector_SignedToSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3158,7 +3158,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldInsert_Int_UnsignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldInsert_Int_UnsignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3190,7 +3190,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldInsert_Int_SignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldInsert_Int_SignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3224,7 +3224,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldInsert_IntVector_UnsignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldInsert_IntVector_UnsignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3257,7 +3257,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldInsert_IntVector_SignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldInsert_IntVector_SignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3292,7 +3292,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldInsert_Uint_UnsignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldInsert_Uint_UnsignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3324,7 +3324,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldInsert_Uint_SignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldInsert_Uint_SignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3358,7 +3358,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldInsert_UintVector_UnsignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldInsert_UintVector_UnsignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3391,7 +3391,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldInsert_UintVector_SignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldInsert_UintVector_SignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3426,7 +3426,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldInsert_Uint_SignedOffsetAndUnsignedCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldInsert_Uint_SignedOffsetAndUnsignedCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3459,7 +3459,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldSExtract_Int_UnsignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldSExtract_Int_UnsignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3491,7 +3491,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldSExtract_Int_SignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldSExtract_Int_SignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3525,7 +3525,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldSExtract_IntVector_UnsignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldSExtract_IntVector_UnsignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3557,7 +3557,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldSExtract_IntVector_SignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldSExtract_IntVector_SignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3591,7 +3591,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldSExtract_IntVector_SignedOffsetAndUnsignedCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldSExtract_IntVector_SignedOffsetAndUnsignedCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3624,7 +3624,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldSExtract_Uint_UnsignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldSExtract_Uint_UnsignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3658,7 +3658,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldSExtract_Uint_SignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldSExtract_Uint_SignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3694,7 +3694,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldSExtract_UintVector_UnsignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldSExtract_UintVector_UnsignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3728,7 +3728,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldSExtract_UintVector_SignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldSExtract_UintVector_SignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3764,7 +3764,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldSExtract_UintVector_SignedOffsetAndUnsignedCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldSExtract_UintVector_SignedOffsetAndUnsignedCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3799,7 +3799,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldUExtract_Uint_UnsignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldUExtract_Uint_UnsignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3831,7 +3831,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldUExtract_Uint_SignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldUExtract_Uint_SignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3865,7 +3865,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldUExtract_UintVector_UnsignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldUExtract_UintVector_UnsignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3897,7 +3897,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldUExtract_UintVector_SignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldUExtract_UintVector_SignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3931,7 +3931,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldUExtract_UintVector_UnsignedOffsetAndSignedCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldUExtract_UintVector_UnsignedOffsetAndSignedCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3964,7 +3964,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldUExtract_Int_UnsignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldUExtract_Int_UnsignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -3998,7 +3998,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldUExtract_Int_SignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldUExtract_Int_SignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -4034,7 +4034,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldUExtract_IntVector_UnsignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldUExtract_IntVector_UnsignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -4068,7 +4068,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldUExtract_IntVector_SignedOffsetAndCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldUExtract_IntVector_SignedOffsetAndCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -4104,7 +4104,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitFieldUExtract_IntVector_UnsignedOffsetAndSignedCount) {
+TEST_F(SpirvReader_BuiltinsTest, BitFieldUExtract_IntVector_UnsignedOffsetAndSignedCount) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -4144,9 +4144,9 @@
std::string ir;
};
-using SpirvParser_BuiltinsMixedSignTest = core::ir::transform::TransformTestWithParam<BinaryCase>;
+using SpirvReader_BuiltinsMixedSignTest = core::ir::transform::TransformTestWithParam<BinaryCase>;
-TEST_P(SpirvParser_BuiltinsMixedSignTest, Scalar_Signed_SignedUnsigned) {
+TEST_P(SpirvReader_BuiltinsMixedSignTest, Scalar_Signed_SignedUnsigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4183,7 +4183,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsMixedSignTest, Scalar_Signed_UnsignedSigned) {
+TEST_P(SpirvReader_BuiltinsMixedSignTest, Scalar_Signed_UnsignedSigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4221,7 +4221,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsMixedSignTest, Scalar_Signed_UnsignedUnsigned) {
+TEST_P(SpirvReader_BuiltinsMixedSignTest, Scalar_Signed_UnsignedUnsigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4258,7 +4258,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsMixedSignTest, Scalar_Unsigned_SignedUnsigned) {
+TEST_P(SpirvReader_BuiltinsMixedSignTest, Scalar_Unsigned_SignedUnsigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4296,7 +4296,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsMixedSignTest, Scalar_Unsigned_UnsignedSigned) {
+TEST_P(SpirvReader_BuiltinsMixedSignTest, Scalar_Unsigned_UnsignedSigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4333,7 +4333,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsMixedSignTest, Scalar_Unsigned_SignedSigned) {
+TEST_P(SpirvReader_BuiltinsMixedSignTest, Scalar_Unsigned_SignedSigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4370,7 +4370,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsMixedSignTest, Vector_Signed_SignedUnsigned) {
+TEST_P(SpirvReader_BuiltinsMixedSignTest, Vector_Signed_SignedUnsigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4409,7 +4409,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsMixedSignTest, Vector_Signed_UnsignedSigned) {
+TEST_P(SpirvReader_BuiltinsMixedSignTest, Vector_Signed_UnsignedSigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4449,7 +4449,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsMixedSignTest, Vector_Signed_UnsignedUnsigned) {
+TEST_P(SpirvReader_BuiltinsMixedSignTest, Vector_Signed_UnsignedUnsigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4488,7 +4488,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsMixedSignTest, Vector_Unsigned_SignedUnsigned) {
+TEST_P(SpirvReader_BuiltinsMixedSignTest, Vector_Unsigned_SignedUnsigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4528,7 +4528,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsMixedSignTest, Vector_Unsigned_UnsignedSigned) {
+TEST_P(SpirvReader_BuiltinsMixedSignTest, Vector_Unsigned_UnsignedSigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4567,7 +4567,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsMixedSignTest, Vector_Unsigned_SignedSigned) {
+TEST_P(SpirvReader_BuiltinsMixedSignTest, Vector_Unsigned_SignedSigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4607,7 +4607,7 @@
}
INSTANTIATE_TEST_SUITE_P(SpirvReader,
- SpirvParser_BuiltinsMixedSignTest,
+ SpirvReader_BuiltinsMixedSignTest,
testing::Values(BinaryCase{spirv::BuiltinFn::kAdd, "add"},
BinaryCase{spirv::BuiltinFn::kSub, "sub"},
BinaryCase{spirv::BuiltinFn::kMul, "mul"}));
@@ -4618,10 +4618,10 @@
std::string wgsl;
};
-using SpirvParser_BuiltinsSignedTest =
+using SpirvReader_BuiltinsSignedTest =
core::ir::transform::TransformTestWithParam<SignedBinaryCase>;
-TEST_P(SpirvParser_BuiltinsSignedTest, Scalar_Signed_SignedUnsigned) {
+TEST_P(SpirvReader_BuiltinsSignedTest, Scalar_Signed_SignedUnsigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4658,7 +4658,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsSignedTest, Scalar_Signed_UnsignedSigned) {
+TEST_P(SpirvReader_BuiltinsSignedTest, Scalar_Signed_UnsignedSigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4695,7 +4695,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsSignedTest, Scalar_Signed_UnsignedUnsigned) {
+TEST_P(SpirvReader_BuiltinsSignedTest, Scalar_Signed_UnsignedUnsigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4733,7 +4733,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsSignedTest, Scalar_Unsigned_SignedUnsigned) {
+TEST_P(SpirvReader_BuiltinsSignedTest, Scalar_Unsigned_SignedUnsigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4771,7 +4771,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsSignedTest, Scalar_Unsigned_UnsignedSigned) {
+TEST_P(SpirvReader_BuiltinsSignedTest, Scalar_Unsigned_UnsignedSigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4809,7 +4809,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsSignedTest, Scalar_Unsigned_SignedSigned) {
+TEST_P(SpirvReader_BuiltinsSignedTest, Scalar_Unsigned_SignedSigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4846,7 +4846,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsSignedTest, Vector_Signed_SignedUnsigned) {
+TEST_P(SpirvReader_BuiltinsSignedTest, Vector_Signed_SignedUnsigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4885,7 +4885,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsSignedTest, Vector_Signed_UnsignedSigned) {
+TEST_P(SpirvReader_BuiltinsSignedTest, Vector_Signed_UnsignedSigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4924,7 +4924,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsSignedTest, Vector_Signed_UnsignedUnsigned) {
+TEST_P(SpirvReader_BuiltinsSignedTest, Vector_Signed_UnsignedUnsigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -4964,7 +4964,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsSignedTest, Vector_Unsigned_SignedUnsigned) {
+TEST_P(SpirvReader_BuiltinsSignedTest, Vector_Unsigned_SignedUnsigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -5004,7 +5004,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsSignedTest, Vector_Unsigned_UnsignedSigned) {
+TEST_P(SpirvReader_BuiltinsSignedTest, Vector_Unsigned_UnsignedSigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -5044,7 +5044,7 @@
EXPECT_EQ(expect, str());
}
-TEST_P(SpirvParser_BuiltinsSignedTest, Vector_Unsigned_SignedSigned) {
+TEST_P(SpirvReader_BuiltinsSignedTest, Vector_Unsigned_SignedSigned) {
auto params = GetParam();
auto* ep = b.ComputeFunction("foo");
@@ -5084,12 +5084,12 @@
}
INSTANTIATE_TEST_SUITE_P(SpirvReader,
- SpirvParser_BuiltinsSignedTest,
+ SpirvReader_BuiltinsSignedTest,
testing::Values(SignedBinaryCase{spirv::BuiltinFn::kSDiv, "s_div", "div"},
SignedBinaryCase{spirv::BuiltinFn::kSMod, "s_mod",
"mod"}));
-TEST_F(SpirvParser_BuiltinsTest, ConvertFToS_ScalarSigned) {
+TEST_F(SpirvReader_BuiltinsTest, ConvertFToS_ScalarSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -5120,7 +5120,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, ConvertFToS_ScalarUnsigned) {
+TEST_F(SpirvReader_BuiltinsTest, ConvertFToS_ScalarUnsigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -5152,7 +5152,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, ConvertFToS_VectorSigned) {
+TEST_F(SpirvReader_BuiltinsTest, ConvertFToS_VectorSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -5183,7 +5183,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, ConvertFToS_VectorUnsigned) {
+TEST_F(SpirvReader_BuiltinsTest, ConvertFToS_VectorUnsigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -5215,7 +5215,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, ConvertSToF_ScalarSigned) {
+TEST_F(SpirvReader_BuiltinsTest, ConvertSToF_ScalarSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -5246,7 +5246,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, ConvertSToF_ScalarUnsigned) {
+TEST_F(SpirvReader_BuiltinsTest, ConvertSToF_ScalarUnsigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -5278,7 +5278,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, ConvertSToF_VectorSigned) {
+TEST_F(SpirvReader_BuiltinsTest, ConvertSToF_VectorSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -5309,7 +5309,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, ConvertSToF_VectorUnsigned) {
+TEST_F(SpirvReader_BuiltinsTest, ConvertSToF_VectorUnsigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -5341,7 +5341,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, ConvertUToF_ScalarSigned) {
+TEST_F(SpirvReader_BuiltinsTest, ConvertUToF_ScalarSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -5373,7 +5373,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, ConvertUToF_ScalarUnsigned) {
+TEST_F(SpirvReader_BuiltinsTest, ConvertUToF_ScalarUnsigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -5404,7 +5404,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, ConvertUToF_VectorSigned) {
+TEST_F(SpirvReader_BuiltinsTest, ConvertUToF_VectorSigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -5436,7 +5436,7 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, ConvertUToF_VectorUnsigned) {
+TEST_F(SpirvReader_BuiltinsTest, ConvertUToF_VectorUnsigned) {
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
@@ -5467,19 +5467,22 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Scalar_SignedSigned_Signed) {
+using SpirvReader_BitwiseTest = core::ir::transform::TransformTestWithParam<SpirvReaderParams>;
+
+TEST_P(SpirvReader_BitwiseTest, Scalar_SignedSigned_Signed) {
+ auto& params = GetParam();
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.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), params.fn, Vector{ty.i32()}, 1_i, 2_i);
b.Return(ep);
});
- auto* src = R"(
+ auto src = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:i32 = spirv.bitwiseAnd<i32> 1i, 2i
+ %2:i32 = spirv.)" +
+ params.spv_name + R"(<i32> 1i, 2i
ret
}
}
@@ -5487,10 +5490,11 @@
EXPECT_EQ(src, str());
Run(Builtins);
- auto* expect = R"(
+ auto expect = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:i32 = and 1i, 2i
+ %2:i32 = )" + params.wgsl_name +
+ R"( 1i, 2i
ret
}
}
@@ -5498,19 +5502,20 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Scalar_SignedUnsigned_Signed) {
+TEST_P(SpirvReader_BitwiseTest, Scalar_SignedUnsigned_Signed) {
+ auto& params = GetParam();
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.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), params.fn, Vector{ty.i32()}, 1_i, 8_u);
b.Return(ep);
});
- auto* src = R"(
+ auto src = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:i32 = spirv.bitwiseAnd<i32> 1i, 8u
+ %2:i32 = spirv.)" +
+ params.spv_name + R"(<i32> 1i, 8u
ret
}
}
@@ -5518,11 +5523,12 @@
EXPECT_EQ(src, str());
Run(Builtins);
- auto* expect = R"(
+ auto expect = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:i32 = bitcast 8u
- %3:i32 = and 1i, %2
+ %3:i32 = )" + params.wgsl_name +
+ R"( 1i, %2
ret
}
}
@@ -5530,19 +5536,20 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Scalar_UnsignedSigned_Signed) {
+TEST_P(SpirvReader_BitwiseTest, Scalar_UnsignedSigned_Signed) {
+ auto& params = GetParam();
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.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), params.fn, Vector{ty.i32()}, 8_u, 1_i);
b.Return(ep);
});
- auto* src = R"(
+ auto src = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:i32 = spirv.bitwiseAnd<i32> 8u, 1i
+ %2:i32 = spirv.)" +
+ params.spv_name + R"(<i32> 8u, 1i
ret
}
}
@@ -5550,11 +5557,12 @@
EXPECT_EQ(src, str());
Run(Builtins);
- auto* expect = R"(
+ auto expect = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:u32 = bitcast 1i
- %3:u32 = and 8u, %2
+ %3:u32 = )" + params.wgsl_name +
+ R"( 8u, %2
%4:i32 = bitcast %3
ret
}
@@ -5563,19 +5571,20 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Scalar_UnsignedUnsigned_Signed) {
+TEST_P(SpirvReader_BitwiseTest, Scalar_UnsignedUnsigned_Signed) {
+ auto& params = GetParam();
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.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), params.fn, Vector{ty.i32()}, 8_u, 9_u);
b.Return(ep);
});
- auto* src = R"(
+ auto src = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:i32 = spirv.bitwiseAnd<i32> 8u, 9u
+ %2:i32 = spirv.)" +
+ params.spv_name + R"(<i32> 8u, 9u
ret
}
}
@@ -5583,10 +5592,11 @@
EXPECT_EQ(src, str());
Run(Builtins);
- auto* expect = R"(
+ auto expect = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:u32 = and 8u, 9u
+ %2:u32 = )" + params.wgsl_name +
+ R"( 8u, 9u
%3:i32 = bitcast %2
ret
}
@@ -5595,19 +5605,20 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Scalar_UnsignedUnsigned_Unsigned) {
+TEST_P(SpirvReader_BitwiseTest, Scalar_UnsignedUnsigned_Unsigned) {
+ auto& params = GetParam();
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.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), params.fn, Vector{ty.u32()}, 8_u, 9_u);
b.Return(ep);
});
- auto* src = R"(
+ auto src = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:u32 = spirv.bitwiseAnd<u32> 8u, 9u
+ %2:u32 = spirv.)" +
+ params.spv_name + R"(<u32> 8u, 9u
ret
}
}
@@ -5615,10 +5626,11 @@
EXPECT_EQ(src, str());
Run(Builtins);
- auto* expect = R"(
+ auto expect = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:u32 = and 8u, 9u
+ %2:u32 = )" + params.wgsl_name +
+ R"( 8u, 9u
ret
}
}
@@ -5626,19 +5638,20 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Scalar_UnsignedSigned_Unsigned) {
+TEST_P(SpirvReader_BitwiseTest, Scalar_UnsignedSigned_Unsigned) {
+ auto& params = GetParam();
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.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), params.fn, Vector{ty.u32()}, 8_u, 1_i);
b.Return(ep);
});
- auto* src = R"(
+ auto src = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:u32 = spirv.bitwiseAnd<u32> 8u, 1i
+ %2:u32 = spirv.)" +
+ params.spv_name + R"(<u32> 8u, 1i
ret
}
}
@@ -5646,11 +5659,12 @@
EXPECT_EQ(src, str());
Run(Builtins);
- auto* expect = R"(
+ auto expect = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:u32 = bitcast 1i
- %3:u32 = and 8u, %2
+ %3:u32 = )" + params.wgsl_name +
+ R"( 8u, %2
ret
}
}
@@ -5658,19 +5672,20 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Scalar_SignedUnsigned_Unsigned) {
+TEST_P(SpirvReader_BitwiseTest, Scalar_SignedUnsigned_Unsigned) {
+ auto& params = GetParam();
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.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), params.fn, Vector{ty.u32()}, 1_i, 8_u);
b.Return(ep);
});
- auto* src = R"(
+ auto src = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:u32 = spirv.bitwiseAnd<u32> 1i, 8u
+ %2:u32 = spirv.)" +
+ params.spv_name + R"(<u32> 1i, 8u
ret
}
}
@@ -5678,11 +5693,12 @@
EXPECT_EQ(src, str());
Run(Builtins);
- auto* expect = R"(
+ auto expect = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:i32 = bitcast 8u
- %3:i32 = and 1i, %2
+ %3:i32 = )" + params.wgsl_name +
+ R"( 1i, %2
%4:u32 = bitcast %3
ret
}
@@ -5691,19 +5707,20 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Scalar_SignedSigned_Unsigned) {
+TEST_P(SpirvReader_BitwiseTest, Scalar_SignedSigned_Unsigned) {
+ auto& params = GetParam();
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.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), params.fn, Vector{ty.u32()}, 1_i, 2_i);
b.Return(ep);
});
- auto* src = R"(
+ auto src = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:u32 = spirv.bitwiseAnd<u32> 1i, 2i
+ %2:u32 = spirv.)" +
+ params.spv_name + R"(<u32> 1i, 2i
ret
}
}
@@ -5711,10 +5728,11 @@
EXPECT_EQ(src, str());
Run(Builtins);
- auto* expect = R"(
+ auto expect = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:i32 = and 1i, 2i
+ %2:i32 = )" + params.wgsl_name +
+ R"( 1i, 2i
%3:u32 = bitcast %2
ret
}
@@ -5723,20 +5741,21 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Vector_SignedSigned_Signed) {
+TEST_P(SpirvReader_BitwiseTest, Vector_SignedSigned_Signed) {
+ auto& params = GetParam();
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.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), params.fn, Vector{ty.i32()},
+ b.Splat<vec2<i32>>(1_i), b.Splat<vec2<i32>>(2_i));
b.Return(ep);
});
- auto* src = R"(
+ 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)
+ %2:vec2<i32> = spirv.)" +
+ params.spv_name + R"(<i32> vec2<i32>(1i), vec2<i32>(2i)
ret
}
}
@@ -5744,10 +5763,11 @@
EXPECT_EQ(src, str());
Run(Builtins);
- auto* expect = R"(
+ auto expect = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:vec2<i32> = and vec2<i32>(1i), vec2<i32>(2i)
+ %2:vec2<i32> = )" +
+ params.wgsl_name + R"( vec2<i32>(1i), vec2<i32>(2i)
ret
}
}
@@ -5755,20 +5775,21 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Vector_SignedUnsigned_Signed) {
+TEST_P(SpirvReader_BitwiseTest, Vector_SignedUnsigned_Signed) {
+ auto& params = GetParam();
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.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), params.fn, Vector{ty.i32()},
+ b.Splat<vec2<i32>>(1_i), b.Splat<vec2<u32>>(8_u));
b.Return(ep);
});
- auto* src = R"(
+ 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)
+ %2:vec2<i32> = spirv.)" +
+ params.spv_name + R"(<i32> vec2<i32>(1i), vec2<u32>(8u)
ret
}
}
@@ -5776,11 +5797,12 @@
EXPECT_EQ(src, str());
Run(Builtins);
- auto* expect = R"(
+ 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
+ %3:vec2<i32> = )" +
+ params.wgsl_name + R"( vec2<i32>(1i), %2
ret
}
}
@@ -5788,20 +5810,21 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Vector_UnsignedSigned_Signed) {
+TEST_P(SpirvReader_BitwiseTest, Vector_UnsignedSigned_Signed) {
+ auto& params = GetParam();
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.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), params.fn, Vector{ty.i32()},
+ b.Splat<vec2<u32>>(8_u), b.Splat<vec2<i32>>(1_i));
b.Return(ep);
});
- auto* src = R"(
+ 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)
+ %2:vec2<i32> = spirv.)" +
+ params.spv_name + R"(<i32> vec2<u32>(8u), vec2<i32>(1i)
ret
}
}
@@ -5809,11 +5832,12 @@
EXPECT_EQ(src, str());
Run(Builtins);
- auto* expect = R"(
+ 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
+ %3:vec2<u32> = )" +
+ params.wgsl_name + R"( vec2<u32>(8u), %2
%4:vec2<i32> = bitcast %3
ret
}
@@ -5822,20 +5846,21 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Vector_UnsignedUnsigned_Signed) {
+TEST_P(SpirvReader_BitwiseTest, Vector_UnsignedUnsigned_Signed) {
+ auto& params = GetParam();
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.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), params.fn, Vector{ty.i32()},
+ b.Splat<vec2<u32>>(8_u), b.Splat<vec2<u32>>(9_u));
b.Return(ep);
});
- auto* src = R"(
+ 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)
+ %2:vec2<i32> = spirv.)" +
+ params.spv_name + R"(<i32> vec2<u32>(8u), vec2<u32>(9u)
ret
}
}
@@ -5843,10 +5868,11 @@
EXPECT_EQ(src, str());
Run(Builtins);
- auto* expect = R"(
+ auto expect = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:vec2<u32> = and vec2<u32>(8u), vec2<u32>(9u)
+ %2:vec2<u32> = )" +
+ params.wgsl_name + R"( vec2<u32>(8u), vec2<u32>(9u)
%3:vec2<i32> = bitcast %2
ret
}
@@ -5855,20 +5881,21 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Vector_UnsignedUnsigned_Unsigned) {
+TEST_P(SpirvReader_BitwiseTest, Vector_UnsignedUnsigned_Unsigned) {
+ auto& params = GetParam();
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.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), params.fn, Vector{ty.u32()},
+ b.Splat<vec2<u32>>(8_u), b.Splat<vec2<u32>>(9_u));
b.Return(ep);
});
- auto* src = R"(
+ 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)
+ %2:vec2<u32> = spirv.)" +
+ params.spv_name + R"(<u32> vec2<u32>(8u), vec2<u32>(9u)
ret
}
}
@@ -5876,10 +5903,11 @@
EXPECT_EQ(src, str());
Run(Builtins);
- auto* expect = R"(
+ auto expect = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:vec2<u32> = and vec2<u32>(8u), vec2<u32>(9u)
+ %2:vec2<u32> = )" +
+ params.wgsl_name + R"( vec2<u32>(8u), vec2<u32>(9u)
ret
}
}
@@ -5887,20 +5915,21 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Vector_UnsignedSigned_Unsigned) {
+TEST_P(SpirvReader_BitwiseTest, Vector_UnsignedSigned_Unsigned) {
+ auto& params = GetParam();
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.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), params.fn, Vector{ty.u32()},
+ b.Splat<vec2<u32>>(8_u), b.Splat<vec2<i32>>(1_i));
b.Return(ep);
});
- auto* src = R"(
+ 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)
+ %2:vec2<u32> = spirv.)" +
+ params.spv_name + R"(<u32> vec2<u32>(8u), vec2<i32>(1i)
ret
}
}
@@ -5908,11 +5937,12 @@
EXPECT_EQ(src, str());
Run(Builtins);
- auto* expect = R"(
+ 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
+ %3:vec2<u32> = )" +
+ params.wgsl_name + R"( vec2<u32>(8u), %2
ret
}
}
@@ -5920,20 +5950,21 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Vector_SignedUnsigned_Unsigned) {
+TEST_P(SpirvReader_BitwiseTest, Vector_SignedUnsigned_Unsigned) {
+ auto& params = GetParam();
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.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), params.fn, Vector{ty.u32()},
+ b.Splat<vec2<i32>>(1_i), b.Splat<vec2<u32>>(8_u));
b.Return(ep);
});
- auto* src = R"(
+ 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)
+ %2:vec2<u32> = spirv.)" +
+ params.spv_name + R"(<u32> vec2<i32>(1i), vec2<u32>(8u)
ret
}
}
@@ -5941,11 +5972,12 @@
EXPECT_EQ(src, str());
Run(Builtins);
- auto* expect = R"(
+ 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
+ %3:vec2<i32> = )" +
+ params.wgsl_name + R"( vec2<i32>(1i), %2
%4:vec2<u32> = bitcast %3
ret
}
@@ -5954,20 +5986,21 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, BitwiseAnd_Vector_SignedSigned_Unsigned) {
+TEST_P(SpirvReader_BitwiseTest, Vector_SignedSigned_Unsigned) {
+ auto& params = GetParam();
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.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), params.fn, Vector{ty.u32()},
+ b.Splat<vec2<i32>>(1_i), b.Splat<vec2<i32>>(2_i));
b.Return(ep);
});
- auto* src = R"(
+ 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)
+ %2:vec2<u32> = spirv.)" +
+ params.spv_name + R"(<u32> vec2<i32>(1i), vec2<i32>(2i)
ret
}
}
@@ -5975,10 +6008,11 @@
EXPECT_EQ(src, str());
Run(Builtins);
- auto* expect = R"(
+ auto expect = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:vec2<i32> = and vec2<i32>(1i), vec2<i32>(2i)
+ %2:vec2<i32> = )" +
+ params.wgsl_name + R"( vec2<i32>(1i), vec2<i32>(2i)
%3:vec2<u32> = bitcast %2
ret
}
@@ -5986,6 +6020,11 @@
)";
EXPECT_EQ(expect, str());
}
+INSTANTIATE_TEST_SUITE_P(
+ SpirvReader,
+ SpirvReader_BitwiseTest,
+ ::testing::Values(SpirvReaderParams{spirv::BuiltinFn::kBitwiseAnd, "bitwiseAnd", "and"},
+ SpirvReaderParams{spirv::BuiltinFn::kBitwiseOr, "bitwiseOr", "or"}));
} // 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 2e6de9b..9ff1b5e 100644
--- a/src/tint/lang/spirv/reader/parser/bit_test.cc
+++ b/src/tint/lang/spirv/reader/parser/bit_test.cc
@@ -30,7 +30,18 @@
namespace tint::spirv::reader {
namespace {
-TEST_F(SpirvParserTest, BitwiseAnd_Scalar_SignedSigned_Signed) {
+struct SpirvBitwiseParam {
+ std::string name;
+};
+[[maybe_unused]] inline std::ostream& operator<<(std::ostream& out, SpirvBitwiseParam c) {
+ out << c.name;
+ return out;
+}
+
+using SpirvParser_BitwiseTest = SpirvParserTestWithParam<SpirvBitwiseParam>;
+
+TEST_P(SpirvParser_BitwiseTest, Scalar_SignedSigned_Signed) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -44,21 +55,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpBitwiseAnd %int %one %two
+ %1 = OpBitwise)" +
+ params.name + R"( %int %one %two
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:i32 = spirv.bitwiseAnd<i32> 1i, 2i
+ %2:i32 = spirv.bitwise)" +
+ params.name + R"(<i32> 1i, 2i
ret
}
}
)");
}
-TEST_F(SpirvParserTest, BitwiseAnd_Scalar_SignedUnsigned_Signed) {
+TEST_P(SpirvParser_BitwiseTest, Scalar_SignedUnsigned_Signed) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -75,21 +89,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpBitwiseAnd %int %one %eight
+ %1 = OpBitwise)" +
+ params.name + R"( %int %one %eight
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:i32 = spirv.bitwiseAnd<i32> 1i, 8u
+ %2:i32 = spirv.bitwise)" +
+ params.name + R"(<i32> 1i, 8u
ret
}
}
)");
}
-TEST_F(SpirvParserTest, BitwiseAnd_Scalar_UnsignedSigned_Signed) {
+TEST_P(SpirvParser_BitwiseTest, Scalar_UnsignedSigned_Signed) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -106,21 +123,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpBitwiseAnd %int %eight %one
+ %1 = OpBitwise)" +
+ params.name + R"( %int %eight %one
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:i32 = spirv.bitwiseAnd<i32> 8u, 1i
+ %2:i32 = spirv.bitwise)" +
+ params.name + R"(<i32> 8u, 1i
ret
}
}
)");
}
-TEST_F(SpirvParserTest, BitwiseAnd_Scalar_UnsignedUnsigned_Signed) {
+TEST_P(SpirvParser_BitwiseTest, Scalar_UnsignedUnsigned_Signed) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -137,21 +157,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpBitwiseAnd %int %eight %nine
+ %1 = OpBitwise)" +
+ params.name + R"( %int %eight %nine
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:i32 = spirv.bitwiseAnd<i32> 8u, 9u
+ %2:i32 = spirv.bitwise)" +
+ params.name + R"(<i32> 8u, 9u
ret
}
}
)");
}
-TEST_F(SpirvParserTest, BitwiseAnd_Scalar_UnsignedUnsigned_Unsigned) {
+TEST_P(SpirvParser_BitwiseTest, Scalar_UnsignedUnsigned_Unsigned) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -168,21 +191,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpBitwiseAnd %uint %eight %nine
+ %1 = OpBitwise)" +
+ params.name + R"( %uint %eight %nine
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:u32 = spirv.bitwiseAnd<u32> 8u, 9u
+ %2:u32 = spirv.bitwise)" +
+ params.name + R"(<u32> 8u, 9u
ret
}
}
)");
}
-TEST_F(SpirvParserTest, BitwiseAnd_Scalar_UnsignedSigned_Unsigned) {
+TEST_P(SpirvParser_BitwiseTest, Scalar_UnsignedSigned_Unsigned) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -199,21 +225,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpBitwiseAnd %uint %eight %one
+ %1 = OpBitwise)" +
+ params.name + R"( %uint %eight %one
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:u32 = spirv.bitwiseAnd<u32> 8u, 1i
+ %2:u32 = spirv.bitwise)" +
+ params.name + R"(<u32> 8u, 1i
ret
}
}
)");
}
-TEST_F(SpirvParserTest, BitwiseAnd_Scalar_SignedUnsigned_Unsigned) {
+TEST_P(SpirvParser_BitwiseTest, Scalar_SignedUnsigned_Unsigned) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -230,21 +259,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpBitwiseAnd %uint %one %eight
+ %1 = OpBitwise)" +
+ params.name + R"( %uint %one %eight
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:u32 = spirv.bitwiseAnd<u32> 1i, 8u
+ %2:u32 = spirv.bitwise)" +
+ params.name + R"(<u32> 1i, 8u
ret
}
}
)");
}
-TEST_F(SpirvParserTest, BitwiseAnd_Scalar_SignedSigned_Unsigned) {
+TEST_P(SpirvParser_BitwiseTest, Scalar_SignedSigned_Unsigned) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -261,21 +293,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpBitwiseAnd %uint %one %two
+ %1 = OpBitwise)" +
+ params.name + R"( %uint %one %two
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:u32 = spirv.bitwiseAnd<u32> 1i, 2i
+ %2:u32 = spirv.bitwise)" +
+ params.name + R"(<u32> 1i, 2i
ret
}
}
)");
}
-TEST_F(SpirvParserTest, BitwiseAnd_Vector_SignedSigned_Signed) {
+TEST_P(SpirvParser_BitwiseTest, Vector_SignedSigned_Signed) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -298,21 +333,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpBitwiseAnd %v2int %v2one %v2two
+ %1 = OpBitwise)" +
+ params.name + R"( %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)
+ %2:vec2<i32> = spirv.bitwise)" +
+ params.name + R"(<i32> vec2<i32>(1i), vec2<i32>(2i)
ret
}
}
)");
}
-TEST_F(SpirvParserTest, BitwiseAnd_Vector_SignedUnsigned_Signed) {
+TEST_P(SpirvParser_BitwiseTest, Vector_SignedUnsigned_Signed) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -335,21 +373,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpBitwiseAnd %v2int %v2one %v2eight
+ %1 = OpBitwise)" +
+ params.name + R"( %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)
+ %2:vec2<i32> = spirv.bitwise)" +
+ params.name + R"(<i32> vec2<i32>(1i), vec2<u32>(8u)
ret
}
}
)");
}
-TEST_F(SpirvParserTest, BitwiseAnd_Vector_UnsignedSigned_Signed) {
+TEST_P(SpirvParser_BitwiseTest, Vector_UnsignedSigned_Signed) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -372,21 +413,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpBitwiseAnd %v2int %v2eight %v2one
+ %1 = OpBitwise)" +
+ params.name + R"( %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)
+ %2:vec2<i32> = spirv.bitwise)" +
+ params.name + R"(<i32> vec2<u32>(8u), vec2<i32>(1i)
ret
}
}
)");
}
-TEST_F(SpirvParserTest, BitwiseAnd_Vector_UnsignedUnsigned_Signed) {
+TEST_P(SpirvParser_BitwiseTest, Vector_UnsignedUnsigned_Signed) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -409,21 +453,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpBitwiseAnd %v2int %v2eight %v2nine
+ %1 = OpBitwise)" +
+ params.name + R"( %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)
+ %2:vec2<i32> = spirv.bitwise)" +
+ params.name + R"(<i32> vec2<u32>(8u), vec2<u32>(9u)
ret
}
}
)");
}
-TEST_F(SpirvParserTest, BitwiseAnd_Vector_UnsignedUnsigned_Unsigned) {
+TEST_P(SpirvParser_BitwiseTest, Vector_UnsignedUnsigned_Unsigned) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -446,21 +493,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpBitwiseAnd %v2uint %v2eight %v2nine
+ %1 = OpBitwise)" +
+ params.name + R"( %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)
+ %2:vec2<u32> = spirv.bitwise)" +
+ params.name + R"(<u32> vec2<u32>(8u), vec2<u32>(9u)
ret
}
}
)");
}
-TEST_F(SpirvParserTest, BitwiseAnd_Vector_UnsignedSigned_Unsigned) {
+TEST_P(SpirvParser_BitwiseTest, Vector_UnsignedSigned_Unsigned) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -483,21 +533,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpBitwiseAnd %v2uint %v2eight %v2one
+ %1 = OpBitwise)" +
+ params.name + R"( %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)
+ %2:vec2<u32> = spirv.bitwise)" +
+ params.name + R"(<u32> vec2<u32>(8u), vec2<i32>(1i)
ret
}
}
)");
}
-TEST_F(SpirvParserTest, BitwiseAnd_Vector_SignedUnsigned_Unsigned) {
+TEST_P(SpirvParser_BitwiseTest, Vector_SignedUnsigned_Unsigned) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -520,21 +573,24 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpBitwiseAnd %v2uint %v2one %v2eight
+ %1 = OpBitwise)" +
+ params.name + R"( %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)
+ %2:vec2<u32> = spirv.bitwise)" +
+ params.name + R"(<u32> vec2<i32>(1i), vec2<u32>(8u)
ret
}
}
)");
}
-TEST_F(SpirvParserTest, BitwiseAnd_Vector_SignedSigned_Unsigned) {
+TEST_P(SpirvParser_BitwiseTest, Vector_SignedSigned_Unsigned) {
+ auto& params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -557,19 +613,26 @@
%main = OpFunction %void None %void_fn
%main_start = OpLabel
- %1 = OpBitwiseAnd %v2uint %v2one %v2two
+ %1 = OpBitwise)" +
+ params.name + R"( %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)
+ %2:vec2<u32> = spirv.bitwise)" +
+ params.name + R"(<u32> vec2<i32>(1i), vec2<i32>(2i)
ret
}
}
)");
}
+INSTANTIATE_TEST_SUITE_P(SpirvParser,
+ SpirvParser_BitwiseTest,
+ testing::Values(SpirvBitwiseParam{"And"}, //
+ SpirvBitwiseParam{"Or"}));
+
} // 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 80d054f..836a058 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -671,6 +671,9 @@
case spv::Op::OpBitwiseAnd:
EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kBitwiseAnd);
break;
+ case spv::Op::OpBitwiseOr:
+ EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kBitwiseOr);
+ 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 c8c879b..6950e70 100644
--- a/src/tint/lang/spirv/spirv.def
+++ b/src/tint/lang/spirv/spirv.def
@@ -423,6 +423,8 @@
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>
+implicit(A: iu32, B: iu32) fn bitwiseOr<R: iu32>(A, B) -> R
+implicit(A: iu32, B: iu32, N: num) fn bitwiseOr<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 e5b4f59..a1c31f4 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -1536,6 +1536,9 @@
case BuiltinFn::kBitwiseAnd:
op = spv::Op::OpBitwiseAnd;
break;
+ case BuiltinFn::kBitwiseOr:
+ op = spv::Op::OpBitwiseOr;
+ break;
case spirv::BuiltinFn::kNone:
TINT_ICE() << "undefined spirv ir function";
}
