From 9bdf2dcc6bbda9b727f8b365f5bc081c7b2150ef Mon Sep 17 00:00:00 2001 From: Antonio Maiorano Date: Fri, 30 Jul 2021 18:56:26 +0000 Subject: [PATCH] MSL: fix i32 INT_MIN literal emitted as `long` instead of `int` Bug: tint:124 Change-Id: Ie632b78cd67948b65e823f0a3c52fda7ef7343f3 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60440 Kokoro: Kokoro Commit-Queue: Antonio Maiorano Reviewed-by: Ben Clayton Reviewed-by: David Neto --- src/writer/msl/generator_impl.cc | 13 +++++++++++- src/writer/msl/generator_impl_cast_test.cc | 11 ++++++++++ .../msl/generator_impl_unary_op_test.cc | 12 +++++++++++ .../bitcast/scalar/i32min-u32.wgsl | 4 ++++ .../scalar/i32min-u32.wgsl.expected.hlsl | 5 +++++ .../scalar/i32min-u32.wgsl.expected.msl | 8 ++++++++ .../scalar/i32min-u32.wgsl.expected.spvasm | 20 +++++++++++++++++++ .../scalar/i32min-u32.wgsl.expected.wgsl | 4 ++++ .../0-opt.spvasm.expected.msl | 2 +- .../0-opt.wgsl.expected.msl | 2 +- .../0-opt.spvasm.expected.msl | 2 +- .../0-opt.wgsl.expected.msl | 2 +- 12 files changed, 80 insertions(+), 5 deletions(-) create mode 100644 test/expressions/bitcast/scalar/i32min-u32.wgsl create mode 100644 test/expressions/bitcast/scalar/i32min-u32.wgsl.expected.hlsl create mode 100644 test/expressions/bitcast/scalar/i32min-u32.wgsl.expected.msl create mode 100644 test/expressions/bitcast/scalar/i32min-u32.wgsl.expected.spvasm create mode 100644 test/expressions/bitcast/scalar/i32min-u32.wgsl.expected.wgsl diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc index c5220a00a2..1307f2051f 100644 --- a/src/writer/msl/generator_impl.cc +++ b/src/writer/msl/generator_impl.cc @@ -17,6 +17,7 @@ #include #include #include +#include #include #include @@ -1246,7 +1247,17 @@ bool GeneratorImpl::EmitLiteral(std::ostream& out, ast::Literal* lit) { out << FloatToString(fl->value()) << "f"; } } else if (auto* sl = lit->As()) { - out << sl->value(); + // MSL (and C++) parse `-2147483648` as a `long` because it parses unary + // minus and `2147483648` as separate tokens, and the latter doesn't + // fit into an (32-bit) `int`. WGSL, OTOH, parses this as an `i32`. To avoid + // issues with `long` to `int` casts, emit `(2147483647 - 1)` instead, which + // ensures the expression type is `int`. + const auto int_min = std::numeric_limits::min(); + if (sl->value_as_i32() == int_min) { + out << "(" << int_min + 1 << " - 1)"; + } else { + out << sl->value(); + } } else if (auto* ul = lit->As()) { out << ul->value() << "u"; } else { diff --git a/src/writer/msl/generator_impl_cast_test.cc b/src/writer/msl/generator_impl_cast_test.cc index 6f66478ad2..c871382b49 100644 --- a/src/writer/msl/generator_impl_cast_test.cc +++ b/src/writer/msl/generator_impl_cast_test.cc @@ -43,6 +43,17 @@ TEST_F(MslGeneratorImplTest, EmitExpression_Cast_Vector) { EXPECT_EQ(out.str(), "float3(int3(1, 2, 3))"); } +TEST_F(MslGeneratorImplTest, EmitExpression_Cast_IntMin) { + auto* cast = Construct(std::numeric_limits::min()); + WrapInFunction(cast); + + GeneratorImpl& gen = Build(); + + std::stringstream out; + ASSERT_TRUE(gen.EmitExpression(out, cast)) << gen.error(); + EXPECT_EQ(out.str(), "uint((-2147483647 - 1))"); +} + } // namespace } // namespace msl } // namespace writer diff --git a/src/writer/msl/generator_impl_unary_op_test.cc b/src/writer/msl/generator_impl_unary_op_test.cc index 9254494084..31eac9c664 100644 --- a/src/writer/msl/generator_impl_unary_op_test.cc +++ b/src/writer/msl/generator_impl_unary_op_test.cc @@ -88,6 +88,18 @@ TEST_F(MslUnaryOpTest, Negation) { EXPECT_EQ(out.str(), "-(expr)"); } +TEST_F(MslUnaryOpTest, NegationOfIntMin) { + auto* op = create( + ast::UnaryOp::kNegation, Expr(std::numeric_limits::min())); + WrapInFunction(op); + + GeneratorImpl& gen = Build(); + + std::stringstream out; + ASSERT_TRUE(gen.EmitExpression(out, op)) << gen.error(); + EXPECT_EQ(out.str(), "-((-2147483647 - 1))"); +} + } // namespace } // namespace msl } // namespace writer diff --git a/test/expressions/bitcast/scalar/i32min-u32.wgsl b/test/expressions/bitcast/scalar/i32min-u32.wgsl new file mode 100644 index 0000000000..9812c83dbd --- /dev/null +++ b/test/expressions/bitcast/scalar/i32min-u32.wgsl @@ -0,0 +1,4 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let b : u32 = bitcast(-2147483648); +} diff --git a/test/expressions/bitcast/scalar/i32min-u32.wgsl.expected.hlsl b/test/expressions/bitcast/scalar/i32min-u32.wgsl.expected.hlsl new file mode 100644 index 0000000000..6044ed9a5c --- /dev/null +++ b/test/expressions/bitcast/scalar/i32min-u32.wgsl.expected.hlsl @@ -0,0 +1,5 @@ +[numthreads(1, 1, 1)] +void f() { + const uint b = asuint(-2147483648); + return; +} diff --git a/test/expressions/bitcast/scalar/i32min-u32.wgsl.expected.msl b/test/expressions/bitcast/scalar/i32min-u32.wgsl.expected.msl new file mode 100644 index 0000000000..5ad270a067 --- /dev/null +++ b/test/expressions/bitcast/scalar/i32min-u32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +kernel void f() { + uint const b = as_type((-2147483647 - 1)); + return; +} + diff --git a/test/expressions/bitcast/scalar/i32min-u32.wgsl.expected.spvasm b/test/expressions/bitcast/scalar/i32min-u32.wgsl.expected.spvasm new file mode 100644 index 0000000000..ffd7f0c451 --- /dev/null +++ b/test/expressions/bitcast/scalar/i32min-u32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 9 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %int = OpTypeInt 32 1 +%int_n2147483648 = OpConstant %int -2147483648 + %f = OpFunction %void None %1 + %4 = OpLabel + %5 = OpBitcast %uint %int_n2147483648 + OpReturn + OpFunctionEnd diff --git a/test/expressions/bitcast/scalar/i32min-u32.wgsl.expected.wgsl b/test/expressions/bitcast/scalar/i32min-u32.wgsl.expected.wgsl new file mode 100644 index 0000000000..9812c83dbd --- /dev/null +++ b/test/expressions/bitcast/scalar/i32min-u32.wgsl.expected.wgsl @@ -0,0 +1,4 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let b : u32 = bitcast(-2147483648); +} diff --git a/test/vk-gl-cts/graphicsfuzz/cov-fold-negate-min-int-value/0-opt.spvasm.expected.msl b/test/vk-gl-cts/graphicsfuzz/cov-fold-negate-min-int-value/0-opt.spvasm.expected.msl index bc8e92f6d8..6b6a5dd547 100644 --- a/test/vk-gl-cts/graphicsfuzz/cov-fold-negate-min-int-value/0-opt.spvasm.expected.msl +++ b/test/vk-gl-cts/graphicsfuzz/cov-fold-negate-min-int-value/0-opt.spvasm.expected.msl @@ -14,7 +14,7 @@ struct tint_symbol_1 { void main_1(constant buf0& x_7, thread float4* const tint_symbol_4) { int minValue = 0; int negMinValue = 0; - minValue = -2147483648; + minValue = (-2147483647 - 1); int const x_25 = minValue; negMinValue = -(x_25); int const x_27 = negMinValue; diff --git a/test/vk-gl-cts/graphicsfuzz/cov-fold-negate-min-int-value/0-opt.wgsl.expected.msl b/test/vk-gl-cts/graphicsfuzz/cov-fold-negate-min-int-value/0-opt.wgsl.expected.msl index bc8e92f6d8..6b6a5dd547 100644 --- a/test/vk-gl-cts/graphicsfuzz/cov-fold-negate-min-int-value/0-opt.wgsl.expected.msl +++ b/test/vk-gl-cts/graphicsfuzz/cov-fold-negate-min-int-value/0-opt.wgsl.expected.msl @@ -14,7 +14,7 @@ struct tint_symbol_1 { void main_1(constant buf0& x_7, thread float4* const tint_symbol_4) { int minValue = 0; int negMinValue = 0; - minValue = -2147483648; + minValue = (-2147483647 - 1); int const x_25 = minValue; negMinValue = -(x_25); int const x_27 = negMinValue; diff --git a/test/vk-gl-cts/graphicsfuzz/cov-x86-isel-lowering-negative-left-shift/0-opt.spvasm.expected.msl b/test/vk-gl-cts/graphicsfuzz/cov-x86-isel-lowering-negative-left-shift/0-opt.spvasm.expected.msl index a7a8dbb87d..933626f679 100644 --- a/test/vk-gl-cts/graphicsfuzz/cov-x86-isel-lowering-negative-left-shift/0-opt.spvasm.expected.msl +++ b/test/vk-gl-cts/graphicsfuzz/cov-x86-isel-lowering-negative-left-shift/0-opt.spvasm.expected.msl @@ -69,7 +69,7 @@ void main_1(constant buf1& x_6, constant buf0& x_9, thread float4* const tint_sy } case 0: { int const x_70 = i; - if ((-2147483648 < x_70)) { + if (((-2147483647 - 1) < x_70)) { { int const x_82 = j; j = (x_82 + 1); diff --git a/test/vk-gl-cts/graphicsfuzz/cov-x86-isel-lowering-negative-left-shift/0-opt.wgsl.expected.msl b/test/vk-gl-cts/graphicsfuzz/cov-x86-isel-lowering-negative-left-shift/0-opt.wgsl.expected.msl index a7a8dbb87d..933626f679 100644 --- a/test/vk-gl-cts/graphicsfuzz/cov-x86-isel-lowering-negative-left-shift/0-opt.wgsl.expected.msl +++ b/test/vk-gl-cts/graphicsfuzz/cov-x86-isel-lowering-negative-left-shift/0-opt.wgsl.expected.msl @@ -69,7 +69,7 @@ void main_1(constant buf1& x_6, constant buf0& x_9, thread float4* const tint_sy } case 0: { int const x_70 = i; - if ((-2147483648 < x_70)) { + if (((-2147483647 - 1) < x_70)) { { int const x_82 = j; j = (x_82 + 1);