[spirv-reader][ir] Correctly handle GLSL 450 UMin
The SPIR-V `UMin` method allows unsigned arguments and return types.
This is not permitted in WGSL. The SPIR-V spec states that the argument
is treated as a signed value, so bitcast the argument/result as needed.
Bug: 42250952
Change-Id: I52fe1f869b15ca8a89c68b57be03928cc80fd7f3
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/222755
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 78d0254..84521f1 100644
--- a/src/tint/lang/spirv/builtin_fn.cc
+++ b/src/tint/lang/spirv/builtin_fn.cc
@@ -124,6 +124,8 @@
return "sclamp";
case BuiltinFn::kUmax:
return "umax";
+ case BuiltinFn::kUmin:
+ return "umin";
case BuiltinFn::kSdot:
return "sdot";
case BuiltinFn::kUdot:
@@ -185,6 +187,7 @@
case BuiltinFn::kSmin:
case BuiltinFn::kSclamp:
case BuiltinFn::kUmax:
+ case BuiltinFn::kUmin:
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 c20bf32..35967f4 100644
--- a/src/tint/lang/spirv/builtin_fn.cc.tmpl
+++ b/src/tint/lang/spirv/builtin_fn.cc.tmpl
@@ -80,6 +80,7 @@
case BuiltinFn::kSmin:
case BuiltinFn::kSclamp:
case BuiltinFn::kUmax:
+ case BuiltinFn::kUmin:
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 c34ecac..50b4cc4 100644
--- a/src/tint/lang/spirv/builtin_fn.h
+++ b/src/tint/lang/spirv/builtin_fn.h
@@ -89,6 +89,7 @@
kSmin,
kSclamp,
kUmax,
+ kUmin,
kSdot,
kUdot,
kNone,
diff --git a/src/tint/lang/spirv/intrinsic/data.cc b/src/tint/lang/spirv/intrinsic/data.cc
index ff0dd41..4ed1781 100644
--- a/src/tint/lang/spirv/intrinsic/data.cc
+++ b/src/tint/lang/spirv/intrinsic/data.cc
@@ -5682,12 +5682,19 @@
},
{
/* [41] */
+ /* fn umin<R : iu32>[T : iu32, U : iu32](T, U) -> R */
+ /* fn umin<R : iu32>[T : iu32, U : iu32, N : num](vec<N, T>, vec<N, U>) -> vec<N, R> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(156),
+ },
+ {
+ /* [42] */
/* fn sdot(u32, u32, u32) -> i32 */
/* num overloads */ 1,
/* overloads */ OverloadIndex(171),
},
{
- /* [42] */
+ /* [43] */
/* 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 0d67a1e..0188299 100644
--- a/src/tint/lang/spirv/reader/import_glsl_std450_test.cc
+++ b/src/tint/lang/spirv/reader/import_glsl_std450_test.cc
@@ -646,23 +646,6 @@
)");
}
-TEST_F(SpirvReaderTest, DISABLED_RectifyOperandsAndResult_UMin) {
- EXPECT_IR(Preamble() + R"(
- %1 = OpExtInst %int %glsl UMin %int_30 %int_35
- %2 = OpExtInst %v2int %glsl UMin %v2int_30_40 %v2int_40_30
- OpReturn
- OpFunctionEnd
- )",
- R"(
-%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
- $B1: {
- let x_1 = bitcast<i32>(min(bitcast<u32>(i1), bitcast<u32>(i2)));
- let x_2 = bitcast<vec2i>(min(bitcast<vec2u>(v2i1), bitcast<vec2u>(v2i2)));
- }
-}
-)");
-}
-
TEST_F(SpirvReaderTest, DISABLED_RectifyOperandsAndResult_UClamp) {
EXPECT_IR(Preamble() + R"(
%1 = OpExtInst %int %glsl UClamp %int_30 %uint_15 %int_40
diff --git a/src/tint/lang/spirv/reader/lower/builtins.cc b/src/tint/lang/spirv/reader/lower/builtins.cc
index 3cd419e..63d290f 100644
--- a/src/tint/lang/spirv/reader/lower/builtins.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins.cc
@@ -85,6 +85,9 @@
case spirv::BuiltinFn::kUmax:
UMax(builtin);
break;
+ case spirv::BuiltinFn::kUmin:
+ UMin(builtin);
+ break;
default:
TINT_UNREACHABLE() << "unknown spirv builtin: " << builtin->Func();
}
@@ -170,6 +173,9 @@
void UMax(spirv::ir::BuiltinCall* call) {
WrapUnsignedSpirvMethods(call, core::BuiltinFn::kMax);
}
+ void UMin(spirv::ir::BuiltinCall* call) {
+ WrapUnsignedSpirvMethods(call, core::BuiltinFn::kMin);
+ }
void Normalize(spirv::ir::BuiltinCall* call) {
auto* arg = call->Args()[0];
diff --git a/src/tint/lang/spirv/reader/lower/builtins_test.cc b/src/tint/lang/spirv/reader/lower/builtins_test.cc
index 6e87fd1..e2c4efb 100644
--- a/src/tint/lang/spirv/reader/lower/builtins_test.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins_test.cc
@@ -1551,7 +1551,8 @@
INSTANTIATE_TEST_SUITE_P(SpirvReader,
SpirvParser_BuiltinsTest_TwoParamUnsigned,
- ::testing::Values(SpirvReaderParams{spirv::BuiltinFn::kUmax, "max"}));
+ ::testing::Values(SpirvReaderParams{spirv::BuiltinFn::kUmax, "max"},
+ SpirvReaderParams{spirv::BuiltinFn::kUmin, "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 32010c6..2a17571 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
@@ -374,7 +374,8 @@
GlslStd450TwoParamTest,
::testing::Values(GlslStd450TwoParams{"SMax", "smax"},
GlslStd450TwoParams{"SMin", "smin"},
- GlslStd450TwoParams{"UMax", "umax"}));
+ GlslStd450TwoParams{"UMax", "umax"},
+ GlslStd450TwoParams{"UMin", "umin"}));
TEST_F(SpirvParserTest, GlslStd450_SClamp_UnsignedToUnsigned) {
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 9258df4..a2c13ea 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -695,7 +695,6 @@
case GLSLstd450NClamp:
case GLSLstd450FClamp: // FClamp is less prescriptive about NaN operands
return core::BuiltinFn::kClamp;
- case GLSLstd450UMin:
case GLSLstd450NMin:
case GLSLstd450FMin: // FMin is less prescriptive about NaN operands
return core::BuiltinFn::kMin;
@@ -771,6 +770,8 @@
return spirv::BuiltinFn::kSclamp;
case GLSLstd450UMax:
return spirv::BuiltinFn::kUmax;
+ case GLSLstd450UMin:
+ return spirv::BuiltinFn::kUmin;
default:
break;
}
@@ -781,7 +782,8 @@
const core::type::Type* result_ty) {
if (ext_opcode == GLSLstd450SSign || ext_opcode == GLSLstd450SAbs ||
ext_opcode == GLSLstd450SMax || ext_opcode == GLSLstd450SMin ||
- ext_opcode == GLSLstd450SClamp || ext_opcode == GLSLstd450UMax) {
+ ext_opcode == GLSLstd450SClamp || ext_opcode == GLSLstd450UMax ||
+ ext_opcode == GLSLstd450UMin) {
return {result_ty->DeepestElement()};
}
return {};
diff --git a/src/tint/lang/spirv/spirv.def b/src/tint/lang/spirv/spirv.def
index 1202a8f..5f6f0fc 100644
--- a/src/tint/lang/spirv/spirv.def
+++ b/src/tint/lang/spirv/spirv.def
@@ -340,6 +340,8 @@
implicit(T: iu32, U: iu32) fn umax<R: iu32>(T, U) -> R
implicit(T: iu32, U: iu32, N: num) fn umax<R: iu32>(vec<N, T>, vec<N, U>) -> vec<N, R>
+implicit(T: iu32, U: iu32) fn umin<R: iu32>(T, U) -> R
+implicit(T: iu32, U: iu32, N: num) fn umin<R: iu32>(vec<N, T>, vec<N, U>) -> 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 e768da5..b2ddaad 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -1390,6 +1390,9 @@
case spirv::BuiltinFn::kUmax:
ext_inst(GLSLstd450UMax);
break;
+ case spirv::BuiltinFn::kUmin:
+ ext_inst(GLSLstd450UMin);
+ break;
case spirv::BuiltinFn::kNormalize:
ext_inst(GLSLstd450Normalize);
break;
