[spirv-reader][ir] Correctly handle GLSL 450 FindSMsb
The SPIR-V `FindSMsb` method allows differing argument and return type.
The argument must also be interpreted as a signed value. So, inject
bitcasts as needed for the argument and return.
Bug: 391487629
Change-Id: If98081c52ddeebb2c176cc286771ce06bbeaf5e6
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/223055
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: David Neto <dneto@google.com>
diff --git a/src/tint/lang/spirv/builtin_fn.cc b/src/tint/lang/spirv/builtin_fn.cc
index e43c360..47afb9a 100644
--- a/src/tint/lang/spirv/builtin_fn.cc
+++ b/src/tint/lang/spirv/builtin_fn.cc
@@ -130,6 +130,8 @@
return "uclamp";
case BuiltinFn::kFindILsb:
return "findILsb";
+ case BuiltinFn::kFindSMsb:
+ return "findSMsb";
case BuiltinFn::kSdot:
return "sdot";
case BuiltinFn::kUdot:
@@ -194,6 +196,7 @@
case BuiltinFn::kUmin:
case BuiltinFn::kUclamp:
case BuiltinFn::kFindILsb:
+ case BuiltinFn::kFindSMsb:
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 30d7591..b8c68f1 100644
--- a/src/tint/lang/spirv/builtin_fn.cc.tmpl
+++ b/src/tint/lang/spirv/builtin_fn.cc.tmpl
@@ -83,6 +83,7 @@
case BuiltinFn::kUmin:
case BuiltinFn::kUclamp:
case BuiltinFn::kFindILsb:
+ case BuiltinFn::kFindSMsb:
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 1f62a6c..a5c85ef 100644
--- a/src/tint/lang/spirv/builtin_fn.h
+++ b/src/tint/lang/spirv/builtin_fn.h
@@ -92,6 +92,7 @@
kUmin,
kUclamp,
kFindILsb,
+ kFindSMsb,
kSdot,
kUdot,
kNone,
diff --git a/src/tint/lang/spirv/intrinsic/data.cc b/src/tint/lang/spirv/intrinsic/data.cc
index 3e1a0bd..e5b7421 100644
--- a/src/tint/lang/spirv/intrinsic/data.cc
+++ b/src/tint/lang/spirv/intrinsic/data.cc
@@ -5703,12 +5703,19 @@
},
{
/* [44] */
+ /* fn findSMsb<R : iu32>[T : iu32](T) -> R */
+ /* fn findSMsb<R : iu32>[T : iu32, N : num](vec<N, T>) -> vec<N, R> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(154),
+ },
+ {
+ /* [45] */
/* fn sdot(u32, u32, u32) -> i32 */
/* num overloads */ 1,
/* overloads */ OverloadIndex(171),
},
{
- /* [45] */
+ /* [46] */
/* fn udot(u32, u32, u32) -> u32 */
/* num overloads */ 1,
/* overloads */ OverloadIndex(172),
diff --git a/src/tint/lang/spirv/reader/import_glsl_std450_test.cc b/src/tint/lang/spirv/reader/import_glsl_std450_test.cc
index cdea0a3..146868d 100644
--- a/src/tint/lang/spirv/reader/import_glsl_std450_test.cc
+++ b/src/tint/lang/spirv/reader/import_glsl_std450_test.cc
@@ -649,54 +649,6 @@
)");
}
-TEST_F(SpirvReaderTest, DISABLED_RectifyOperandsAndResult_FindSMsb) {
- // Check signedness conversion of arguments and results.
- // SPIR-V signed arg -> keep it
- // signed result -> keep it
- // unsigned result -> cast result to unsigned
- //
- // SPIR-V unsigned arg -> cast it to signed
- // signed result -> keept it
- // unsigned result -> cast result to unsigned
- EXPECT_IR(Preamble() + R"(
- ; signed arg
- ; signed result
- %1 = OpExtInst %int %glsl FindSMsb %int_30
- %2 = OpExtInst %v2int %glsl FindSMsb %v2int_30_40
-
- ; signed arg
- ; unsigned result
- %3 = OpExtInst %uint %glsl FindSMsb %int_30
- %4 = OpExtInst %v2uint %glsl FindSMsb %v2int_30_40
-
- ; unsigned arg
- ; signed result
- %5 = OpExtInst %int %glsl FindSMsb %uint_10
- %6 = OpExtInst %v2int %glsl FindSMsb %v2uint_10_20
-
- ; unsigned arg
- ; unsigned result
- %7 = OpExtInst %uint %glsl FindSMsb %uint_10
- %8 = OpExtInst %v2uint %glsl FindSMsb %v2uint_10_20
- OpReturn
- OpFunctionEnd
- )",
- R"(
-%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
- $B1: {
- let x_1 = firstLeadingBit(i1);
- let x_2 = firstLeadingBit(v2i1);
- let x_3 = bitcast<u32>(firstLeadingBit(i1));
- let x_4 = bitcast<vec2u>(firstLeadingBit(v2i1));
- let x_5 = firstLeadingBit(bitcast<i32>(u1));
- let x_6 = firstLeadingBit(bitcast<vec2i>(v2u1));
- let x_7 = bitcast<u32>(firstLeadingBit(bitcast<i32>(u1)));
- let x_8 = bitcast<vec2u>(firstLeadingBit(bitcast<vec2i>(v2u1)));
- }
-}
-)");
-}
-
TEST_F(SpirvReaderTest, DISABLED_RectifyOperandsAndResult_FindUMsb) {
// Check signedness conversion of arguments and results.
// SPIR-V signed arg -> cast arg to unsigned
diff --git a/src/tint/lang/spirv/reader/lower/builtins.cc b/src/tint/lang/spirv/reader/lower/builtins.cc
index 85acacb..b4ed7b1 100644
--- a/src/tint/lang/spirv/reader/lower/builtins.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins.cc
@@ -94,6 +94,9 @@
case spirv::BuiltinFn::kFindILsb:
FindILsb(builtin);
break;
+ case spirv::BuiltinFn::kFindSMsb:
+ FindSMsb(builtin);
+ break;
default:
TINT_UNREACHABLE() << "unknown spirv builtin: " << builtin->Func();
}
@@ -137,6 +140,9 @@
WrapSignedSpirvMethods(call, core::BuiltinFn::kSign);
}
void Abs(spirv::ir::BuiltinCall* call) { WrapSignedSpirvMethods(call, core::BuiltinFn::kAbs); }
+ void FindSMsb(spirv::ir::BuiltinCall* call) {
+ WrapSignedSpirvMethods(call, core::BuiltinFn::kFirstLeadingBit);
+ }
void SMax(spirv::ir::BuiltinCall* call) { WrapSignedSpirvMethods(call, core::BuiltinFn::kMax); }
void SMin(spirv::ir::BuiltinCall* call) { WrapSignedSpirvMethods(call, core::BuiltinFn::kMin); }
void SClamp(spirv::ir::BuiltinCall* call) {
diff --git a/src/tint/lang/spirv/reader/lower/builtins_test.cc b/src/tint/lang/spirv/reader/lower/builtins_test.cc
index cda9c6e..e589218 100644
--- a/src/tint/lang/spirv/reader/lower/builtins_test.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins_test.cc
@@ -1005,23 +1005,40 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, SAbs_UnsignedToUnsigned) {
+struct SpirvReaderParams {
+ spirv::BuiltinFn fn;
+ std::string spv_name;
+ std::string wgsl_name;
+};
+[[maybe_unused]] inline std::ostream& operator<<(std::ostream& out, SpirvReaderParams c) {
+ out << c.spv_name;
+ return out;
+}
+
+using SpirvParser_BuiltinsTest_OneParamSigned =
+ core::ir::transform::TransformTestWithParam<SpirvReaderParams>;
+
+TEST_P(SpirvParser_BuiltinsTest_OneParamSigned, UnsignedToUnsigned) {
+ auto params = GetParam();
+
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
- b.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kAbs,
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), params.fn,
Vector<const core::type::Type*, 1>{ty.u32()}, 10_u);
- b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), spirv::BuiltinFn::kAbs,
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), params.fn,
Vector<const core::type::Type*, 1>{ty.u32()},
b.Splat(ty.vec2<u32>(), 10_u));
b.Return(ep);
});
- auto* src = R"(
+ auto src = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:u32 = spirv.abs<u32> 10u
- %3:vec2<u32> = spirv.abs<u32> vec2<u32>(10u)
+ %2:u32 = spirv.)" +
+ params.spv_name + R"(<u32> 10u
+ %3:vec2<u32> = spirv.)" +
+ params.spv_name + R"(<u32> vec2<u32>(10u)
ret
}
}
@@ -1030,14 +1047,16 @@
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 10u
- %3:i32 = abs %2
+ %3:i32 = )" + params.wgsl_name +
+ R"( %2
%4:u32 = bitcast %3
%5:vec2<i32> = bitcast vec2<u32>(10u)
- %6:vec2<i32> = abs %5
+ %6:vec2<i32> = )" +
+ params.wgsl_name + R"( %5
%7:vec2<u32> = bitcast %6
ret
}
@@ -1046,23 +1065,27 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, SAbs_UnsignedToSigned) {
+TEST_P(SpirvParser_BuiltinsTest_OneParamSigned, UnsignedToSigned) {
+ auto params = GetParam();
+
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
- b.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kAbs,
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), params.fn,
Vector<const core::type::Type*, 1>{ty.i32()}, 10_u);
- b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), spirv::BuiltinFn::kAbs,
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), params.fn,
Vector<const core::type::Type*, 1>{ty.i32()},
b.Splat(ty.vec2<u32>(), 10_u));
b.Return(ep);
});
- auto* src = R"(
+ auto src = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:i32 = spirv.abs<i32> 10u
- %3:vec2<i32> = spirv.abs<i32> vec2<u32>(10u)
+ %2:i32 = spirv.)" +
+ params.spv_name + R"(<i32> 10u
+ %3:vec2<i32> = spirv.)" +
+ params.spv_name + R"(<i32> vec2<u32>(10u)
ret
}
}
@@ -1071,13 +1094,15 @@
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 10u
- %3:i32 = abs %2
+ %3:i32 = )" + params.wgsl_name +
+ R"( %2
%4:vec2<i32> = bitcast vec2<u32>(10u)
- %5:vec2<i32> = abs %4
+ %5:vec2<i32> = )" +
+ params.wgsl_name + R"( %4
ret
}
}
@@ -1085,23 +1110,27 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, SAbs_SignedToSigned) {
+TEST_P(SpirvParser_BuiltinsTest_OneParamSigned, SignedToSigned) {
+ auto params = GetParam();
+
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
- b.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kAbs,
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), params.fn,
Vector<const core::type::Type*, 1>{ty.i32()}, 10_i);
- b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), spirv::BuiltinFn::kAbs,
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), params.fn,
Vector<const core::type::Type*, 1>{ty.i32()},
b.Splat(ty.vec2<i32>(), 10_i));
b.Return(ep);
});
- auto* src = R"(
+ auto src = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:i32 = spirv.abs<i32> 10i
- %3:vec2<i32> = spirv.abs<i32> vec2<i32>(10i)
+ %2:i32 = spirv.)" +
+ params.spv_name + R"(<i32> 10i
+ %3:vec2<i32> = spirv.)" +
+ params.spv_name + R"(<i32> vec2<i32>(10i)
ret
}
}
@@ -1110,11 +1139,13 @@
EXPECT_EQ(src, str());
Run(Builtins);
- auto* expect = R"(
+ auto expect = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:i32 = abs 10i
- %3:vec2<i32> = abs vec2<i32>(10i)
+ %2:i32 = )" + params.wgsl_name +
+ R"( 10i
+ %3:vec2<i32> = )" +
+ params.wgsl_name + R"( vec2<i32>(10i)
ret
}
}
@@ -1122,23 +1153,27 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvParser_BuiltinsTest, SAbs_SignedToUnsigned) {
+TEST_P(SpirvParser_BuiltinsTest_OneParamSigned, SignedToUnsigned) {
+ auto params = GetParam();
+
auto* ep = b.ComputeFunction("foo");
b.Append(ep->Block(), [&] { //
- b.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kAbs,
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), params.fn,
Vector<const core::type::Type*, 1>{ty.u32()}, 10_i);
- b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), spirv::BuiltinFn::kAbs,
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), params.fn,
Vector<const core::type::Type*, 1>{ty.u32()},
b.Splat(ty.vec2<i32>(), 10_i));
b.Return(ep);
});
- auto* src = R"(
+ auto src = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:u32 = spirv.abs<u32> 10i
- %3:vec2<u32> = spirv.abs<u32> vec2<i32>(10i)
+ %2:u32 = spirv.)" +
+ params.spv_name + R"(<u32> 10i
+ %3:vec2<u32> = spirv.)" +
+ params.spv_name + R"(<u32> vec2<i32>(10i)
ret
}
}
@@ -1147,12 +1182,14 @@
EXPECT_EQ(src, str());
Run(Builtins);
- auto* expect = R"(
+ auto expect = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:i32 = abs 10i
+ %2:i32 = )" + params.wgsl_name +
+ R"( 10i
%3:u32 = bitcast %2
- %4:vec2<i32> = abs vec2<i32>(10i)
+ %4:vec2<i32> = )" +
+ params.wgsl_name + R"( vec2<i32>(10i)
%5:vec2<u32> = bitcast %4
ret
}
@@ -1161,14 +1198,11 @@
EXPECT_EQ(expect, str());
}
-struct SpirvReaderParams {
- spirv::BuiltinFn fn;
- std::string name;
-};
-[[maybe_unused]] inline std::ostream& operator<<(std::ostream& out, SpirvReaderParams c) {
- out << c.name;
- return out;
-}
+INSTANTIATE_TEST_SUITE_P(SpirvReader,
+ SpirvParser_BuiltinsTest_OneParamSigned,
+ ::testing::Values(SpirvReaderParams{spirv::BuiltinFn::kAbs, "abs", "abs"},
+ SpirvReaderParams{spirv::BuiltinFn::kFindSMsb,
+ "findSMsb", "firstLeadingBit"}));
using SpirvParser_BuiltinsTest_TwoParamSigned =
core::ir::transform::TransformTestWithParam<SpirvReaderParams>;
@@ -1191,9 +1225,9 @@
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:u32 = spirv.s)" +
- params.name + R"(<u32> 10u, 15u
+ params.spv_name + R"(<u32> 10u, 15u
%3:vec2<u32> = spirv.s)" +
- params.name + R"(<u32> vec2<u32>(10u), vec2<u32>(15u)
+ params.spv_name + R"(<u32> vec2<u32>(10u), vec2<u32>(15u)
ret
}
}
@@ -1207,13 +1241,13 @@
$B1: {
%2:i32 = bitcast 10u
%3:i32 = bitcast 15u
- %4:i32 = )" + params.name +
+ %4:i32 = )" + params.wgsl_name +
R"( %2, %3
%5:u32 = bitcast %4
%6:vec2<i32> = bitcast vec2<u32>(10u)
%7:vec2<i32> = bitcast vec2<u32>(15u)
%8:vec2<i32> = )" +
- params.name + R"( %6, %7
+ params.wgsl_name + R"( %6, %7
%9:vec2<u32> = bitcast %8
ret
}
@@ -1240,9 +1274,9 @@
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:i32 = spirv.s)" +
- params.name + R"(<i32> 10i, 15i
+ params.spv_name + R"(<i32> 10i, 15i
%3:vec2<i32> = spirv.s)" +
- params.name + R"(<i32> vec2<i32>(10i), vec2<i32>(15i)
+ params.spv_name + R"(<i32> vec2<i32>(10i), vec2<i32>(15i)
ret
}
}
@@ -1254,10 +1288,10 @@
auto expect = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:i32 = )" + params.name +
+ %2:i32 = )" + params.wgsl_name +
R"( 10i, 15i
%3:vec2<i32> = )" +
- params.name + R"( vec2<i32>(10i), vec2<i32>(15i)
+ params.wgsl_name + R"( vec2<i32>(10i), vec2<i32>(15i)
ret
}
}
@@ -1283,9 +1317,9 @@
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:u32 = spirv.s)" +
- params.name + R"(<u32> 10i, 10u
+ params.spv_name + R"(<u32> 10i, 10u
%3:vec2<u32> = spirv.s)" +
- params.name + R"(<u32> vec2<i32>(10i), vec2<u32>(10u)
+ params.spv_name + R"(<u32> vec2<i32>(10i), vec2<u32>(10u)
ret
}
}
@@ -1298,12 +1332,12 @@
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:i32 = bitcast 10u
- %3:i32 = )" + params.name +
+ %3:i32 = )" + params.wgsl_name +
R"( 10i, %2
%4:u32 = bitcast %3
%5:vec2<i32> = bitcast vec2<u32>(10u)
%6:vec2<i32> = )" +
- params.name + R"( vec2<i32>(10i), %5
+ params.wgsl_name + R"( vec2<i32>(10i), %5
%7:vec2<u32> = bitcast %6
ret
}
@@ -1330,9 +1364,9 @@
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:i32 = spirv.s)" +
- params.name + R"(<i32> 10u, 10i
+ params.spv_name + R"(<i32> 10u, 10i
%3:vec2<i32> = spirv.s)" +
- params.name + R"(<i32> vec2<u32>(10u), vec2<i32>(10i)
+ params.spv_name + R"(<i32> vec2<u32>(10u), vec2<i32>(10i)
ret
}
}
@@ -1345,11 +1379,11 @@
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:i32 = bitcast 10u
- %3:i32 = )" + params.name +
+ %3:i32 = )" + params.wgsl_name +
R"( %2, 10i
%4:vec2<i32> = bitcast vec2<u32>(10u)
%5:vec2<i32> = )" +
- params.name + R"( %4, vec2<i32>(10i)
+ params.wgsl_name + R"( %4, vec2<i32>(10i)
ret
}
}
@@ -1359,8 +1393,9 @@
INSTANTIATE_TEST_SUITE_P(SpirvReader,
SpirvParser_BuiltinsTest_TwoParamSigned,
- ::testing::Values(SpirvReaderParams{spirv::BuiltinFn::kSmax, "max"},
- SpirvReaderParams{spirv::BuiltinFn::kSmin, "min"}));
+ ::testing::Values(SpirvReaderParams{spirv::BuiltinFn::kSmax, "max", "max"},
+ SpirvReaderParams{spirv::BuiltinFn::kSmin, "min",
+ "min"}));
using SpirvParser_BuiltinsTest_TwoParamUnsigned =
core::ir::transform::TransformTestWithParam<SpirvReaderParams>;
@@ -1383,9 +1418,9 @@
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:u32 = spirv.u)" +
- params.name + R"(<u32> 10u, 15u
+ params.spv_name + R"(<u32> 10u, 15u
%3:vec2<u32> = spirv.u)" +
- params.name + R"(<u32> vec2<u32>(10u), vec2<u32>(15u)
+ params.spv_name + R"(<u32> vec2<u32>(10u), vec2<u32>(15u)
ret
}
}
@@ -1397,10 +1432,10 @@
auto expect = R"(
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- %2:u32 = )" + params.name +
+ %2:u32 = )" + params.wgsl_name +
R"( 10u, 15u
%3:vec2<u32> = )" +
- params.name + R"( vec2<u32>(10u), vec2<u32>(15u)
+ params.wgsl_name + R"( vec2<u32>(10u), vec2<u32>(15u)
ret
}
}
@@ -1426,9 +1461,9 @@
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:i32 = spirv.u)" +
- params.name + R"(<i32> 10i, 15i
+ params.spv_name + R"(<i32> 10i, 15i
%3:vec2<i32> = spirv.u)" +
- params.name + R"(<i32> vec2<i32>(10i), vec2<i32>(15i)
+ params.spv_name + R"(<i32> vec2<i32>(10i), vec2<i32>(15i)
ret
}
}
@@ -1442,13 +1477,13 @@
$B1: {
%2:u32 = bitcast 10i
%3:u32 = bitcast 15i
- %4:u32 = )" + params.name +
+ %4:u32 = )" + params.wgsl_name +
R"( %2, %3
%5:i32 = bitcast %4
%6:vec2<u32> = bitcast vec2<i32>(10i)
%7:vec2<u32> = bitcast vec2<i32>(15i)
%8:vec2<u32> = )" +
- params.name + R"( %6, %7
+ params.wgsl_name + R"( %6, %7
%9:vec2<i32> = bitcast %8
ret
}
@@ -1475,9 +1510,9 @@
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:u32 = spirv.u)" +
- params.name + R"(<u32> 10i, 10u
+ params.spv_name + R"(<u32> 10i, 10u
%3:vec2<u32> = spirv.u)" +
- params.name + R"(<u32> vec2<i32>(10i), vec2<u32>(10u)
+ params.spv_name + R"(<u32> vec2<i32>(10i), vec2<u32>(10u)
ret
}
}
@@ -1490,11 +1525,11 @@
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:u32 = bitcast 10i
- %3:u32 = )" + params.name +
+ %3:u32 = )" + params.wgsl_name +
R"( %2, 10u
%4:vec2<u32> = bitcast vec2<i32>(10i)
%5:vec2<u32> = )" +
- params.name + R"( %4, vec2<u32>(10u)
+ params.wgsl_name + R"( %4, vec2<u32>(10u)
ret
}
}
@@ -1520,9 +1555,9 @@
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:i32 = spirv.u)" +
- params.name + R"(<i32> 10u, 10i
+ params.spv_name + R"(<i32> 10u, 10i
%3:vec2<i32> = spirv.u)" +
- params.name + R"(<i32> vec2<u32>(10u), vec2<i32>(10i)
+ params.spv_name + R"(<i32> vec2<u32>(10u), vec2<i32>(10i)
ret
}
}
@@ -1535,12 +1570,12 @@
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
%2:u32 = bitcast 10i
- %3:u32 = )" + params.name +
+ %3:u32 = )" + params.wgsl_name +
R"( 10u, %2
%4:i32 = bitcast %3
%5:vec2<u32> = bitcast vec2<i32>(10i)
%6:vec2<u32> = )" +
- params.name + R"( vec2<u32>(10u), %5
+ params.wgsl_name + R"( vec2<u32>(10u), %5
%7:vec2<i32> = bitcast %6
ret
}
@@ -1551,8 +1586,9 @@
INSTANTIATE_TEST_SUITE_P(SpirvReader,
SpirvParser_BuiltinsTest_TwoParamUnsigned,
- ::testing::Values(SpirvReaderParams{spirv::BuiltinFn::kUmax, "max"},
- SpirvReaderParams{spirv::BuiltinFn::kUmin, "min"}));
+ ::testing::Values(SpirvReaderParams{spirv::BuiltinFn::kUmax, "max", "max"},
+ SpirvReaderParams{spirv::BuiltinFn::kUmin, "min",
+ "min"}));
TEST_F(SpirvParser_BuiltinsTest, SClamp_UnsignedToUnsigned) {
auto* ep = b.ComputeFunction("foo");
diff --git a/src/tint/lang/spirv/reader/parser/import_glsl_std450_test.cc b/src/tint/lang/spirv/reader/parser/import_glsl_std450_test.cc
index f87d7af..54020d8 100644
--- a/src/tint/lang/spirv/reader/parser/import_glsl_std450_test.cc
+++ b/src/tint/lang/spirv/reader/parser/import_glsl_std450_test.cc
@@ -159,104 +159,131 @@
)");
}
-TEST_F(SpirvParserTest, GlslStd450_SAbs_UnsignedToUnsigned) {
- EXPECT_IR(Preamble() + R"(
- %1 = OpExtInst %uint %glsl SAbs %uint_10
- %2 = OpExtInst %v2uint %glsl SAbs %v2uint_10_20
- %3 = OpCopyObject %uint %1
- %4 = OpCopyObject %v2uint %2
- OpReturn
- OpFunctionEnd
- )",
- R"(
-%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
- $B1: {
- %2:u32 = spirv.abs<u32> 10u
- %3:vec2<u32> = spirv.abs<u32> vec2<u32>(10u, 20u)
- %4:u32 = let %2
- %5:vec2<u32> = let %3
- ret
- }
-}
-)");
-}
-
-TEST_F(SpirvParserTest, GlslStd450_SAbs_UnsignedToSigned) {
- EXPECT_IR(Preamble() + R"(
- %1 = OpExtInst %int %glsl SAbs %uint_10
- %2 = OpExtInst %v2int %glsl SAbs %v2uint_10_20
- %3 = OpCopyObject %int %1
- %4 = OpCopyObject %v2int %2
- OpReturn
- OpFunctionEnd
- )",
- R"(
-%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
- $B1: {
- %2:i32 = spirv.abs<i32> 10u
- %3:vec2<i32> = spirv.abs<i32> vec2<u32>(10u, 20u)
- %4:i32 = let %2
- %5:vec2<i32> = let %3
- ret
- }
-}
-)");
-}
-
-TEST_F(SpirvParserTest, GlslStd450_SAbs_SignedToUnsigned) {
- EXPECT_IR(Preamble() + R"(
- %1 = OpExtInst %uint %glsl SAbs %int_10
- %2 = OpExtInst %v2uint %glsl SAbs %v2int_10_20
- %3 = OpCopyObject %uint %1
- %4 = OpCopyObject %v2uint %2
- OpReturn
- OpFunctionEnd
- )",
- R"(
-%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
- $B1: {
- %2:u32 = spirv.abs<u32> 10i
- %3:vec2<u32> = spirv.abs<u32> vec2<i32>(10i, 20i)
- %4:u32 = let %2
- %5:vec2<u32> = let %3
- ret
- }
-}
-)");
-}
-
-TEST_F(SpirvParserTest, GlslStd450_SAbs_SignedToSigned) {
- EXPECT_IR(Preamble() + R"(
- %1 = OpExtInst %int %glsl SAbs %int_10
- %2 = OpExtInst %v2int %glsl SAbs %v2int_10_20
- %3 = OpCopyObject %int %1
- %4 = OpCopyObject %v2int %2
- OpReturn
- OpFunctionEnd
- )",
- R"(
-%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
- $B1: {
- %2:i32 = spirv.abs<i32> 10i
- %3:vec2<i32> = spirv.abs<i32> vec2<i32>(10i, 20i)
- %4:i32 = let %2
- %5:vec2<i32> = let %3
- ret
- }
-}
-)");
-}
-
-struct GlslStd450TwoParams {
+struct GlslStd450Params {
std::string spv_name;
std::string ir_name;
};
-[[maybe_unused]] inline std::ostream& operator<<(std::ostream& out, GlslStd450TwoParams c) {
+[[maybe_unused]] inline std::ostream& operator<<(std::ostream& out, GlslStd450Params c) {
out << c.spv_name;
return out;
}
-using GlslStd450TwoParamTest = SpirvParserTestWithParam<GlslStd450TwoParams>;
+using GlslStd450OneParamTest = SpirvParserTestWithParam<GlslStd450Params>;
+
+TEST_P(GlslStd450OneParamTest, UnsignedToUnsigned) {
+ auto params = GetParam();
+ EXPECT_IR(Preamble() + R"(
+ %1 = OpExtInst %uint %glsl )" +
+ params.spv_name + R"( %uint_10
+ %2 = OpExtInst %v2uint %glsl )" +
+ params.spv_name + R"( %v2uint_10_20
+ %3 = OpCopyObject %uint %1
+ %4 = OpCopyObject %v2uint %2
+ OpReturn
+ OpFunctionEnd
+ )",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.)" +
+ params.ir_name + R"(<u32> 10u
+ %3:vec2<u32> = spirv.)" +
+ params.ir_name + R"(<u32> vec2<u32>(10u, 20u)
+ %4:u32 = let %2
+ %5:vec2<u32> = let %3
+ ret
+ }
+}
+)");
+}
+
+TEST_P(GlslStd450OneParamTest, UnsignedToSigned) {
+ auto params = GetParam();
+ EXPECT_IR(Preamble() + R"(
+ %1 = OpExtInst %int %glsl )" +
+ params.spv_name + R"( %uint_10
+ %2 = OpExtInst %v2int %glsl )" +
+ params.spv_name + R"( %v2uint_10_20
+ %3 = OpCopyObject %int %1
+ %4 = OpCopyObject %v2int %2
+ OpReturn
+ OpFunctionEnd
+ )",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.)" +
+ params.ir_name + R"(<i32> 10u
+ %3:vec2<i32> = spirv.)" +
+ params.ir_name + R"(<i32> vec2<u32>(10u, 20u)
+ %4:i32 = let %2
+ %5:vec2<i32> = let %3
+ ret
+ }
+}
+)");
+}
+
+TEST_P(GlslStd450OneParamTest, SignedToUnsigned) {
+ auto params = GetParam();
+ EXPECT_IR(Preamble() + R"(
+ %1 = OpExtInst %uint %glsl )" +
+ params.spv_name + R"( %int_10
+ %2 = OpExtInst %v2uint %glsl )" +
+ params.spv_name + R"( %v2int_10_20
+ %3 = OpCopyObject %uint %1
+ %4 = OpCopyObject %v2uint %2
+ OpReturn
+ OpFunctionEnd
+ )",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.)" +
+ params.ir_name + R"(<u32> 10i
+ %3:vec2<u32> = spirv.)" +
+ params.ir_name + R"(<u32> vec2<i32>(10i, 20i)
+ %4:u32 = let %2
+ %5:vec2<u32> = let %3
+ ret
+ }
+}
+)");
+}
+
+TEST_P(GlslStd450OneParamTest, SignedToSigned) {
+ auto params = GetParam();
+ EXPECT_IR(Preamble() + R"(
+ %1 = OpExtInst %int %glsl )" +
+ params.spv_name + R"( %int_10
+ %2 = OpExtInst %v2int %glsl )" +
+ params.spv_name + R"( %v2int_10_20
+ %3 = OpCopyObject %int %1
+ %4 = OpCopyObject %v2int %2
+ OpReturn
+ OpFunctionEnd
+ )",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.)" +
+ params.ir_name + R"(<i32> 10i
+ %3:vec2<i32> = spirv.)" +
+ params.ir_name + R"(<i32> vec2<i32>(10i, 20i)
+ %4:i32 = let %2
+ %5:vec2<i32> = let %3
+ ret
+ }
+}
+)");
+}
+
+INSTANTIATE_TEST_SUITE_P(SpirvParser,
+ GlslStd450OneParamTest,
+ ::testing::Values(GlslStd450Params{"SAbs", "abs"},
+ GlslStd450Params{"FindSMsb", "findSMsb"}));
+
+using GlslStd450TwoParamTest = SpirvParserTestWithParam<GlslStd450Params>;
TEST_P(GlslStd450TwoParamTest, UnsignedToUnsigned) {
auto params = GetParam();
@@ -372,12 +399,12 @@
INSTANTIATE_TEST_SUITE_P(SpirvParser,
GlslStd450TwoParamTest,
- ::testing::Values(GlslStd450TwoParams{"SMax", "smax"},
- GlslStd450TwoParams{"SMin", "smin"},
- GlslStd450TwoParams{"UMax", "umax"},
- GlslStd450TwoParams{"UMin", "umin"}));
+ ::testing::Values(GlslStd450Params{"SMax", "smax"},
+ GlslStd450Params{"SMin", "smin"},
+ GlslStd450Params{"UMax", "umax"},
+ GlslStd450Params{"UMin", "umin"}));
-using GlslStd450ThreeParamTest = SpirvParserTestWithParam<GlslStd450TwoParams>;
+using GlslStd450ThreeParamTest = SpirvParserTestWithParam<GlslStd450Params>;
TEST_P(GlslStd450ThreeParamTest, UnsignedToUnsigned) {
auto params = GetParam();
@@ -493,8 +520,8 @@
INSTANTIATE_TEST_SUITE_P(SpirvParser,
GlslStd450ThreeParamTest,
- ::testing::Values(GlslStd450TwoParams{"SClamp", "sclamp"},
- GlslStd450TwoParams{"UClamp", "uclamp"}));
+ ::testing::Values(GlslStd450Params{"SClamp", "sclamp"},
+ GlslStd450Params{"UClamp", "uclamp"}));
TEST_F(SpirvParserTest, FindILsb) {
EXPECT_IR(Preamble() + R"(
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index 4ada393..5db0870 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -675,8 +675,6 @@
return core::BuiltinFn::kAbs;
case GLSLstd450FSign:
return core::BuiltinFn::kSign;
- case GLSLstd450FindSMsb:
- return core::BuiltinFn::kFirstLeadingBit;
case GLSLstd450FindUMsb:
return core::BuiltinFn::kFirstLeadingBit;
case GLSLstd450Floor:
@@ -779,6 +777,8 @@
return spirv::BuiltinFn::kUclamp;
case GLSLstd450FindILsb:
return spirv::BuiltinFn::kFindILsb;
+ case GLSLstd450FindSMsb:
+ return spirv::BuiltinFn::kFindSMsb;
default:
break;
}
@@ -791,7 +791,7 @@
ext_opcode == GLSLstd450SMax || ext_opcode == GLSLstd450SMin ||
ext_opcode == GLSLstd450SClamp || ext_opcode == GLSLstd450UMax ||
ext_opcode == GLSLstd450UMin || ext_opcode == GLSLstd450UClamp ||
- ext_opcode == GLSLstd450FindILsb) {
+ ext_opcode == GLSLstd450FindILsb || ext_opcode == GLSLstd450FindSMsb) {
return {result_ty->DeepestElement()};
}
return {};
diff --git a/src/tint/lang/spirv/spirv.def b/src/tint/lang/spirv/spirv.def
index 3e687c3..3724f02 100644
--- a/src/tint/lang/spirv/spirv.def
+++ b/src/tint/lang/spirv/spirv.def
@@ -347,6 +347,8 @@
implicit(T: iu32) fn findILsb<R: iu32>(T) -> R
implicit(T: iu32, N: num) fn findILsb<R: iu32>(vec<N, T>) -> vec<N, R>
+implicit(T: iu32) fn findSMsb<R: iu32>(T) -> R
+implicit(T: iu32, N: num) fn findSMsb<R: iu32>(vec<N, T>) -> 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 566055f..4814b7f 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -1420,6 +1420,9 @@
case spirv::BuiltinFn::kFindILsb:
ext_inst(GLSLstd450FindILsb);
break;
+ case spirv::BuiltinFn::kFindSMsb:
+ ext_inst(GLSLstd450FindSMsb);
+ break;
case spirv::BuiltinFn::kUdot:
module_.PushExtension("SPV_KHR_integer_dot_product");
module_.PushCapability(SpvCapabilityDotProductKHR);
