Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Apr 9, 2025
Merged

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
c3121eb
[naga wgsl-out] Improve representation of minimum int64 literal
jamienicol Mar 27, 2025
903f812
[naga msl-out] invalid code emitted for minimum i32/i64 value literal.
LegNeato Mar 26, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
22 changes: 20 additions & 2 deletions naga/src/back/msl/writer.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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}")?;
Expand Down
10 changes: 5 additions & 5 deletions naga/src/back/wgsl/writer.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1139,12 +1139,12 @@ impl<W: Write> Writer<W> {
crate::Literal::Bool(value) => write!(self.out, "{value}")?,
crate::Literal::F64(value) => write!(self.out, "{value:?}lf")?,
crate::Literal::I64(value) => {
// `-9223372036854775808li` is not valid WGSL. Nor can we use the AbstractInt
// trick above, as AbstractInt also cannot represent `9223372036854775808`.
// The most negative `i64` value can only be expressed in WGSL using
// subtracting 1 from the second most negative value.
// `-9223372036854775808li` is not valid WGSL. Nor can we simply use the
// AbstractInt trick above, as AbstractInt also cannot represent
// `9223372036854775808`. Instead construct the second most negative
// AbstractInt, subtract one from it, then cast to i64.
if value == i64::MIN {
write!(self.out, "{}li - 1li", value + 1)?;
write!(self.out, "i64({} - 1)", value + 1)?;
} else {
write!(self.out, "{value}li")?;
}
Expand Down
2 changes: 1 addition & 1 deletion naga/tests/in/wgsl/int64.wgsl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -56,7 +56,7 @@ fn int64_function(x: i64) -> i64 {
val += bitcast<vec3<i64>>(input_uniform.val_u64_3).z;
val += bitcast<vec4<i64>>(input_uniform.val_u64_4).w;
// Most negative i64
val += -9223372036854775807li - 1li;
val += i64(-9223372036854775807 - 1);

// Reading/writing to a uniform/storage buffer
output.val_i64 = input_uniform.val_i64 + input_storage.val_i64;
Expand Down
8 changes: 8 additions & 0 deletions naga/tests/in/wgsl/operators.wgsl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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() {
Expand Down
4 changes: 4 additions & 0 deletions naga/tests/out/glsl/wgsl-operators.main.Compute.glsl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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));
Expand Down Expand Up @@ -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;
}

Expand Down
56 changes: 28 additions & 28 deletions naga/tests/out/hlsl/wgsl-int64.hlsl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -99,51 +99,51 @@ int64_t int64_function(int64_t x)
uint64_t4 _e71 = input_uniform.val_u64_4_;
int64_t _e74 = val;
val = (_e74 + _e71.w);
int64_t _e79 = val;
val = (_e79 + (-9223372036854775807L - 1L));
int64_t _e85 = input_uniform.val_i64_;
int64_t _e88 = input_storage.Load<int64_t>(128);
output.Store(128, (_e85 + _e88));
int64_t2 _e94 = input_uniform.val_i64_2_;
int64_t2 _e97 = input_storage.Load<int64_t2>(144);
output.Store(144, (_e94 + _e97));
int64_t3 _e103 = input_uniform.val_i64_3_;
int64_t3 _e106 = input_storage.Load<int64_t3>(160);
output.Store(160, (_e103 + _e106));
int64_t4 _e112 = input_uniform.val_i64_4_;
int64_t4 _e115 = input_storage.Load<int64_t4>(192);
output.Store(192, (_e112 + _e115));
int64_t _e121[2] = Constructarray2_int64_t_(input_arrays.Load<int64_t>(16+0), input_arrays.Load<int64_t>(16+8));
int64_t _e77 = val;
val = (_e77 + -9223372036854775808L);
int64_t _e83 = input_uniform.val_i64_;
int64_t _e86 = input_storage.Load<int64_t>(128);
output.Store(128, (_e83 + _e86));
int64_t2 _e92 = input_uniform.val_i64_2_;
int64_t2 _e95 = input_storage.Load<int64_t2>(144);
output.Store(144, (_e92 + _e95));
int64_t3 _e101 = input_uniform.val_i64_3_;
int64_t3 _e104 = input_storage.Load<int64_t3>(160);
output.Store(160, (_e101 + _e104));
int64_t4 _e110 = input_uniform.val_i64_4_;
int64_t4 _e113 = input_storage.Load<int64_t4>(192);
output.Store(192, (_e110 + _e113));
int64_t _e119[2] = Constructarray2_int64_t_(input_arrays.Load<int64_t>(16+0), input_arrays.Load<int64_t>(16+8));
{
int64_t _value2[2] = _e121;
int64_t _value2[2] = _e119;
output_arrays.Store(16+0, _value2[0]);
output_arrays.Store(16+8, _value2[1]);
}
int64_t _e120 = val;
int64_t _e122 = val;
val = (_e122 + abs(_e120));
int64_t _e124 = val;
val = (_e124 + abs(_e122));
int64_t _e125 = val;
int64_t _e126 = val;
int64_t _e127 = val;
int64_t _e128 = val;
val = (_e128 + clamp(_e124, _e125, _e126));
int64_t _e130 = val;
val = (_e130 + clamp(_e126, _e127, _e128));
int64_t _e132 = val;
int64_t _e134 = val;
int64_t _e135 = val;
val = (_e135 + dot((_e130).xx, (_e132).xx));
int64_t _e137 = val;
val = (_e137 + dot((_e132).xx, (_e134).xx));
int64_t _e139 = val;
int64_t _e138 = val;
int64_t _e140 = val;
val = (_e140 + max(_e137, _e138));
int64_t _e142 = val;
val = (_e142 + max(_e139, _e140));
int64_t _e144 = val;
int64_t _e143 = val;
int64_t _e145 = val;
val = (_e145 + min(_e142, _e143));
int64_t _e147 = val;
val = (_e147 + min(_e144, _e145));
int64_t _e149 = val;
val = (_e149 + sign(_e147));
int64_t _e151 = val;
val = (_e151 + sign(_e149));
int64_t _e153 = val;
return _e153;
return _e151;
}

uint64_t naga_f2u64(float value) {
Expand Down
5 changes: 5 additions & 0 deletions naga/tests/out/hlsl/wgsl-operators.hlsl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down Expand Up @@ -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;
}

Expand Down
4 changes: 2 additions & 2 deletions naga/tests/out/msl/wgsl-abstract-types-operators.msl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand All @@ -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(
) {
Expand Down
8 changes: 4 additions & 4 deletions naga/tests/out/msl/wgsl-conversion-float-to-int-no-f64.msl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
58 changes: 29 additions & 29 deletions naga/tests/out/msl/wgsl-int64.msl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -83,49 +83,49 @@ long int64_function(
metal::ulong4 _e71 = input_uniform.val_u64_4_;
long _e74 = val;
val = as_type<long>(as_type<ulong>(_e74) + as_type<ulong>(as_type<metal::long4>(_e71).w));
long _e79 = val;
val = as_type<long>(as_type<ulong>(_e79) + as_type<ulong>(as_type<long>(as_type<ulong>(-9223372036854775807L) - as_type<ulong>(1L))));
long _e85 = input_uniform.val_i64_;
long _e88 = input_storage.val_i64_;
output.val_i64_ = as_type<long>(as_type<ulong>(_e85) + as_type<ulong>(_e88));
metal::long2 _e94 = input_uniform.val_i64_2_;
metal::long2 _e97 = input_storage.val_i64_2_;
output.val_i64_2_ = as_type<metal::long2>(as_type<metal::ulong2>(_e94) + as_type<metal::ulong2>(_e97));
metal::long3 _e103 = input_uniform.val_i64_3_;
metal::long3 _e106 = input_storage.val_i64_3_;
output.val_i64_3_ = as_type<metal::long3>(as_type<metal::ulong3>(_e103) + as_type<metal::ulong3>(_e106));
metal::long4 _e112 = input_uniform.val_i64_4_;
metal::long4 _e115 = input_storage.val_i64_4_;
output.val_i64_4_ = as_type<metal::long4>(as_type<metal::ulong4>(_e112) + as_type<metal::ulong4>(_e115));
type_12 _e121 = input_arrays.val_i64_array_2_;
output_arrays.val_i64_array_2_ = _e121;
long _e77 = val;
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));
metal::long2 _e92 = input_uniform.val_i64_2_;
metal::long2 _e95 = input_storage.val_i64_2_;
output.val_i64_2_ = as_type<metal::long2>(as_type<metal::ulong2>(_e92) + as_type<metal::ulong2>(_e95));
metal::long3 _e101 = input_uniform.val_i64_3_;
metal::long3 _e104 = input_storage.val_i64_3_;
output.val_i64_3_ = as_type<metal::long3>(as_type<metal::ulong3>(_e101) + as_type<metal::ulong3>(_e104));
metal::long4 _e110 = input_uniform.val_i64_4_;
metal::long4 _e113 = input_storage.val_i64_4_;
output.val_i64_4_ = as_type<metal::long4>(as_type<metal::ulong4>(_e110) + as_type<metal::ulong4>(_e113));
type_12 _e119 = input_arrays.val_i64_array_2_;
output_arrays.val_i64_array_2_ = _e119;
long _e120 = val;
long _e122 = val;
val = as_type<long>(as_type<ulong>(_e122) + as_type<ulong>(naga_abs(_e120)));
long _e124 = val;
val = as_type<long>(as_type<ulong>(_e124) + as_type<ulong>(naga_abs(_e122)));
long _e125 = val;
long _e126 = val;
long _e127 = val;
long _e128 = val;
val = as_type<long>(as_type<ulong>(_e128) + as_type<ulong>(metal::clamp(_e124, _e125, _e126)));
long _e130 = val;
val = as_type<long>(as_type<ulong>(_e130) + as_type<ulong>(metal::clamp(_e126, _e127, _e128)));
metal::long2 _e131 = metal::long2(_e130);
long _e132 = val;
metal::long2 _e133 = metal::long2(_e132);
long _e134 = val;
metal::long2 _e135 = metal::long2(_e134);
long _e135 = val;
val = as_type<long>(as_type<ulong>(_e135) + as_type<ulong>(( + _e131.x * _e133.x + _e131.y * _e133.y)));
long _e137 = val;
val = as_type<long>(as_type<ulong>(_e137) + as_type<ulong>(( + _e133.x * _e135.x + _e133.y * _e135.y)));
long _e139 = val;
long _e138 = val;
long _e140 = val;
val = as_type<long>(as_type<ulong>(_e140) + as_type<ulong>(metal::max(_e137, _e138)));
long _e142 = val;
val = as_type<long>(as_type<ulong>(_e142) + as_type<ulong>(metal::max(_e139, _e140)));
long _e144 = val;
long _e143 = val;
long _e145 = val;
val = as_type<long>(as_type<ulong>(_e145) + as_type<ulong>(metal::min(_e142, _e143)));
long _e147 = val;
val = as_type<long>(as_type<ulong>(_e147) + as_type<ulong>(metal::min(_e144, _e145)));
long _e149 = val;
val = as_type<long>(as_type<ulong>(_e149) + as_type<ulong>(metal::select(metal::select(long(-1), long(1), (_e147 > 0)), long(0), (_e147 == 0))));
long _e151 = val;
val = as_type<long>(as_type<ulong>(_e151) + as_type<ulong>(metal::select(metal::select(long(-1), long(1), (_e149 > 0)), long(0), (_e149 == 0))));
long _e153 = val;
return _e153;
return _e151;
}

ulong naga_f2u64(float value) {
Expand Down
4 changes: 4 additions & 0 deletions naga/tests/out/msl/wgsl-operators.msl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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));
Expand Down Expand Up @@ -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;
}

Expand Down
Loading