[naga msl-out] invalid code emitted for minimum i32/i64 value literal.

The literal `-2147483648` is parsed by Metal as negation of positive
2147483648. As 2147483648 is too large for a int, the expression is
silently promoted to a long. Sometimes this does not matter as it will
often be implicitly converted back to an int after the negation.
However, if the expression is used in a bitcast then we hit a compiler
error due to mismatched bitwidths.

Similarily for `-9223372036854775808`, as 9223372036854775808 is too
large for a long, metal emits a `-Wconstant-conversion` warning and
changes the value to -9223372036854775808. This would then be negated
again, possibly causing undefined behaviour.

In both cases we can avoid the issue by expressing the literals as the
second most negative value expressible by the type, minus one.

eg `-2147483647 - 1` and `-9223372036854775807L - 1L`.

We have added a test which uses the most negative i32 literal in an
addition. Because we bitcast addition operands to unsigned in metal,
this would cause a validation error without this fix. For the i64 case
existing tests already make use of the minimum literal value. Passing
the flag `-Werror=constant-conversion` to Metal during validation will
therefore catch this issue.

Author: LegNeato

naga/src/back/msl/writer.rs

@@ -1494,13 +1494,31 @@ impl<W: Write> Writer<W> {
                 write!(self.out, "{value}u")?;
             }
             crate::Literal::I32(value) => {
-                write!(self.out, "{value}")?;
+                // `-2147483648` is parsed as unary negation of positive 2147483648.
+                // 2147483648 is too large for int32_t meaning the expression gets
+                // promoted to a int64_t which is not our intention. Avoid this by instead
+                // using `-2147483647 - 1`.
+                if value == i32::MIN {
+                    write!(self.out, "({} - 1)", value + 1)?;
+                } else {
+                    write!(self.out, "{value}")?;
+                }
             }
             crate::Literal::U64(value) => {
                 write!(self.out, "{value}uL")?;
             }
             crate::Literal::I64(value) => {
-                write!(self.out, "{value}L")?;
+                // `-9223372036854775808` is parsed as unary negation of positive
+                // 9223372036854775808. 9223372036854775808 is too large for int64_t
+                // causing Metal to emit a `-Wconstant-conversion` warning, and change the
+                // value to `-9223372036854775808`. Which would then be negated, possibly
+                // causing undefined behaviour. Avoid this by instead using
+                // `-9223372036854775808L - 1L`.
+                if value == i64::MIN {
+                    write!(self.out, "({}L - 1L)", value + 1)?;
+                } else {
+                    write!(self.out, "{value}L")?;
+                }
             }
             crate::Literal::Bool(value) => {
                 write!(self.out, "{value}")?;

naga/tests/in/wgsl/operators.wgsl

@@ -160,6 +160,14 @@ fn arithmetic() {
     let mul_vector1 = vec3f(two_f) * mat4x3f();
 
     let mul = mat4x3<f32>() * mat3x4<f32>();
+
+    // Arithmetic involving the minimum value i32 literal. What we're really testing here
+    // is how this literal is expressed by Naga backends. eg in Metal, `-2147483648` is
+    // silently promoted to a `long` which we don't want. The addition ensures this would
+    // be caught as a compiler error, as we bitcast the operands to unsigned which fails
+    // if the expression's type has an unexpected width.
+    var prevent_const_eval: i32;
+    var wgpu_7437 = prevent_const_eval + -2147483648;
 }
 
 fn bit() {

naga/tests/out/glsl/wgsl-operators.main.Compute.glsl

@@ -59,6 +59,8 @@ void logical() {
 }
 
 void arithmetic() {
+    int prevent_const_eval = 0;
+    int wgpu_7437_ = 0;
     float neg0_1 = -(1.0);
     ivec2 neg1_1 = -(ivec2(1));
     vec2 neg2_ = -(vec2(1.0));
@@ -131,6 +133,8 @@ void arithmetic() {
     vec3 mul_vector0_ = (mat4x3(0.0) * vec4(1.0));
     vec4 mul_vector1_ = (vec3(2.0) * mat4x3(0.0));
     mat3x3 mul = (mat4x3(0.0) * mat3x4(0.0));
+    int _e175 = prevent_const_eval;
+    wgpu_7437_ = (_e175 + -2147483648);
     return;
 }
 

naga/tests/out/hlsl/wgsl-operators.hlsl

@@ -121,6 +121,9 @@ float3x4 ZeroValuefloat3x4() {
 
 void arithmetic()
 {
+    int prevent_const_eval = (int)0;
+    int wgpu_7437_ = (int)0;
+
     float neg0_1 = -(1.0);
     int2 neg1_1 = naga_neg((int(1)).xx);
     float2 neg2_ = -((1.0).xx);
@@ -193,6 +196,8 @@ void arithmetic()
     float3 mul_vector0_ = mul((1.0).xxxx, ZeroValuefloat4x3());
     float4 mul_vector1_ = mul(ZeroValuefloat4x3(), (2.0).xxx);
     float3x3 mul_ = mul(ZeroValuefloat3x4(), ZeroValuefloat4x3());
+    int _e175 = prevent_const_eval;
+    wgpu_7437_ = asint(asuint(_e175) + asuint(int(-2147483648)));
     return;
 }
 

naga/tests/out/msl/wgsl-abstract-types-operators.msl

@@ -26,7 +26,7 @@ constant uint plus_u_uai_1 = 3u;
 constant uint plus_u_u_u_1 = 3u;
 constant uint bitflip_u_u = 0u;
 constant uint bitflip_uai = 0u;
-constant int least_i32_ = -2147483648;
+constant int least_i32_ = (-2147483647 - 1);
 constant float least_f32_ = -340282350000000000000000000000000000000.0;
 constant int shl_iaiai = 4;
 constant int shl_iai_u_1 = 4;
@@ -36,7 +36,7 @@ constant int shr_iaiai = 0;
 constant int shr_iai_u_1 = 0;
 constant uint shr_uaiai = 0u;
 constant uint shr_uai_u = 0u;
-constant int wgpu_4492_ = -2147483648;
+constant int wgpu_4492_ = (-2147483647 - 1);
 
 void runtime_values(
 ) {

naga/tests/out/msl/wgsl-conversion-float-to-int-no-f64.msl

@@ -19,19 +19,19 @@ void test_const_eval(
     long max_f16_to_i64_ = 65504L;
     ulong min_f16_to_u64_ = 0uL;
     ulong max_f16_to_u64_ = 65504uL;
-    int min_f32_to_i32_ = -2147483648;
+    int min_f32_to_i32_ = (-2147483647 - 1);
     int max_f32_to_i32_ = 2147483520;
     uint min_f32_to_u32_ = 0u;
     uint max_f32_to_u32_ = 4294967040u;
-    long min_f32_to_i64_ = -9223372036854775808L;
+    long min_f32_to_i64_ = (-9223372036854775807L - 1L);
     long max_f32_to_i64_ = 9223371487098961920L;
     ulong min_f32_to_u64_ = 0uL;
     ulong max_f32_to_u64_ = 18446742974197923840uL;
-    int min_abstract_float_to_i32_ = -2147483648;
+    int min_abstract_float_to_i32_ = (-2147483647 - 1);
     int max_abstract_float_to_i32_ = 2147483647;
     uint min_abstract_float_to_u32_ = 0u;
     uint max_abstract_float_to_u32_ = 4294967295u;
-    long min_abstract_float_to_i64_ = -9223372036854775808L;
+    long min_abstract_float_to_i64_ = (-9223372036854775807L - 1L);
     long max_abstract_float_to_i64_ = 9223372036854774784L;
     ulong min_abstract_float_to_u64_ = 0uL;
     ulong max_abstract_float_to_u64_ = 18446744073709549568uL;

naga/tests/out/msl/wgsl-int64.msl

@@ -84,7 +84,7 @@ long int64_function(
     long _e74 = val;
     val = as_type<long>(as_type<ulong>(_e74) + as_type<ulong>(as_type<metal::long4>(_e71).w));
     long _e77 = val;
-    val = as_type<long>(as_type<ulong>(_e77) + as_type<ulong>(-9223372036854775808L));
+    val = as_type<long>(as_type<ulong>(_e77) + as_type<ulong>((-9223372036854775807L - 1L)));
     long _e83 = input_uniform.val_i64_;
     long _e86 = input_storage.val_i64_;
     output.val_i64_ = as_type<long>(as_type<ulong>(_e83) + as_type<ulong>(_e86));

naga/tests/out/msl/wgsl-operators.msl

@@ -117,6 +117,8 @@ metal::uint2 naga_mod(metal::uint2 lhs, metal::uint2 rhs) {
 
 void arithmetic(
 ) {
+    int prevent_const_eval = {};
+    int wgpu_7437_ = {};
     float neg0_1 = -(1.0);
     metal::int2 neg1_1 = naga_neg(metal::int2(1));
     metal::float2 neg2_ = -(metal::float2(1.0));
@@ -189,6 +191,8 @@ void arithmetic(
     metal::float3 mul_vector0_ = metal::float4x3 {} * metal::float4(1.0);
     metal::float4 mul_vector1_ = metal::float3(2.0) * metal::float4x3 {};
     metal::float3x3 mul = metal::float4x3 {} * metal::float3x4 {};
+    int _e175 = prevent_const_eval;
+    wgpu_7437_ = as_type<int>(as_type<uint>(_e175) + as_type<uint>((-2147483647 - 1)));
     return;
 }