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 <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index c5220a0..1307f20 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -17,6 +17,7 @@
#include <algorithm>
#include <cmath>
#include <iomanip>
+#include <limits>
#include <utility>
#include <vector>
@@ -1246,7 +1247,17 @@
out << FloatToString(fl->value()) << "f";
}
} else if (auto* sl = lit->As<ast::SintLiteral>()) {
- 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<int32_t>::min();
+ if (sl->value_as_i32() == int_min) {
+ out << "(" << int_min + 1 << " - 1)";
+ } else {
+ out << sl->value();
+ }
} else if (auto* ul = lit->As<ast::UintLiteral>()) {
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 6f66478..c871382 100644
--- a/src/writer/msl/generator_impl_cast_test.cc
+++ b/src/writer/msl/generator_impl_cast_test.cc
@@ -43,6 +43,17 @@
EXPECT_EQ(out.str(), "float3(int3(1, 2, 3))");
}
+TEST_F(MslGeneratorImplTest, EmitExpression_Cast_IntMin) {
+ auto* cast = Construct<u32>(std::numeric_limits<int32_t>::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 9254494..31eac9c 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 @@
EXPECT_EQ(out.str(), "-(expr)");
}
+TEST_F(MslUnaryOpTest, NegationOfIntMin) {
+ auto* op = create<ast::UnaryOpExpression>(
+ ast::UnaryOp::kNegation, Expr(std::numeric_limits<int32_t>::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 0000000..9812c83
--- /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<u32>(-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 0000000..6044ed9
--- /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 0000000..5ad270a
--- /dev/null
+++ b/test/expressions/bitcast/scalar/i32min-u32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ uint const b = as_type<uint>((-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 0000000..ffd7f0c
--- /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 0000000..9812c83
--- /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<u32>(-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 bc8e92f..6b6a5dd 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 @@
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 bc8e92f..6b6a5dd 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 @@
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 a7a8dbb..933626f 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 @@
}
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 a7a8dbb..933626f 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 @@
}
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);
