Skip to content

[naga wgsl-out] Represent min value i64 literals as -9223372036854775807li - 1li #7424

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 1 commit into from
Mar 25, 2025
Merged

[naga wgsl-out] Represent min value i64 literals as -9223372036854775807li - 1li #7424

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
9 changes: 5 additions & 4 deletions naga/src/back/wgsl/writer.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1147,11 +1147,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. The most negative `i64`
// value can only be expressed in WGSL using AbstractInt and
// a unary negation operator.
// `-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.
if value == i64::MIN {
write!(self.out, "i64({value})")?;
write!(self.out, "{}li - 1li", value + 1)?;
} else {
write!(self.out, "{value}li")?;
}
Expand Down
2 changes: 2 additions & 0 deletions naga/tests/in/wgsl/int64.wgsl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -55,6 +55,8 @@ fn int64_function(x: i64) -> i64 {
val += bitcast<vec2<i64>>(input_uniform.val_u64_2).y;
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;

// Reading/writing to a uniform/storage buffer
output.val_i64 = input_uniform.val_i64 + input_storage.val_i64;
Expand Down
62 changes: 32 additions & 30 deletions naga/tests/out/hlsl/int64.hlsl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -95,49 +95,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 _e80 = input_uniform.val_i64_;
int64_t _e83 = input_storage.Load<int64_t>(128);
output.Store(128, (_e80 + _e83));
int64_t2 _e89 = input_uniform.val_i64_2_;
int64_t2 _e92 = input_storage.Load<int64_t2>(144);
output.Store(144, (_e89 + _e92));
int64_t3 _e98 = input_uniform.val_i64_3_;
int64_t3 _e101 = input_storage.Load<int64_t3>(160);
output.Store(160, (_e98 + _e101));
int64_t4 _e107 = input_uniform.val_i64_4_;
int64_t4 _e110 = input_storage.Load<int64_t4>(192);
output.Store(192, (_e107 + _e110));
int64_t _e116[2] = Constructarray2_int64_t_(input_arrays.Load<int64_t>(16+0), input_arrays.Load<int64_t>(16+8));
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 _value2[2] = _e116;
int64_t _value2[2] = _e121;
output_arrays.Store(16+0, _value2[0]);
output_arrays.Store(16+8, _value2[1]);
}
int64_t _e117 = val;
int64_t _e119 = val;
val = (_e119 + abs(_e117));
int64_t _e121 = val;
int64_t _e122 = val;
int64_t _e123 = val;
int64_t _e125 = val;
val = (_e125 + clamp(_e121, _e122, _e123));
int64_t _e124 = val;
val = (_e124 + abs(_e122));
int64_t _e126 = val;
int64_t _e127 = val;
int64_t _e129 = val;
int64_t _e128 = val;
int64_t _e130 = val;
val = (_e130 + clamp(_e126, _e127, _e128));
int64_t _e132 = val;
val = (_e132 + dot((_e127).xx, (_e129).xx));
int64_t _e134 = val;
int64_t _e135 = val;
int64_t _e137 = val;
val = (_e137 + max(_e134, _e135));
val = (_e137 + dot((_e132).xx, (_e134).xx));
int64_t _e139 = val;
int64_t _e140 = val;
int64_t _e142 = val;
val = (_e142 + min(_e139, _e140));
val = (_e142 + max(_e139, _e140));
int64_t _e144 = val;
int64_t _e146 = val;
val = (_e146 + sign(_e144));
int64_t _e148 = val;
return _e148;
int64_t _e145 = val;
int64_t _e147 = val;
val = (_e147 + min(_e144, _e145));
int64_t _e149 = val;
int64_t _e151 = val;
val = (_e151 + sign(_e149));
int64_t _e153 = val;
return _e153;
}

typedef uint64_t ret_Constructarray2_uint64_t_[2];
Expand Down
66 changes: 34 additions & 32 deletions naga/tests/out/msl/int64.msl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -79,47 +79,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 _e80 = input_uniform.val_i64_;
long _e83 = input_storage.val_i64_;
output.val_i64_ = as_type<long>(as_type<ulong>(_e80) + as_type<ulong>(_e83));
metal::long2 _e89 = input_uniform.val_i64_2_;
metal::long2 _e92 = input_storage.val_i64_2_;
output.val_i64_2_ = as_type<metal::long2>(as_type<metal::ulong2>(_e89) + as_type<metal::ulong2>(_e92));
metal::long3 _e98 = input_uniform.val_i64_3_;
metal::long3 _e101 = input_storage.val_i64_3_;
output.val_i64_3_ = as_type<metal::long3>(as_type<metal::ulong3>(_e98) + as_type<metal::ulong3>(_e101));
metal::long4 _e107 = input_uniform.val_i64_4_;
metal::long4 _e110 = input_storage.val_i64_4_;
output.val_i64_4_ = as_type<metal::long4>(as_type<metal::ulong4>(_e107) + as_type<metal::ulong4>(_e110));
type_12 _e116 = input_arrays.val_i64_array_2_;
output_arrays.val_i64_array_2_ = _e116;
long _e117 = val;
long _e119 = val;
val = as_type<long>(as_type<ulong>(_e119) + as_type<ulong>(naga_abs(_e117)));
long _e121 = val;
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 _e122 = val;
long _e123 = val;
long _e125 = val;
val = as_type<long>(as_type<ulong>(_e125) + as_type<ulong>(metal::clamp(_e121, _e122, _e123)));
long _e124 = val;
val = as_type<long>(as_type<ulong>(_e124) + as_type<ulong>(naga_abs(_e122)));
long _e126 = val;
long _e127 = val;
metal::long2 _e128 = metal::long2(_e127);
long _e129 = val;
metal::long2 _e130 = metal::long2(_e129);
long _e128 = val;
long _e130 = val;
val = as_type<long>(as_type<ulong>(_e130) + as_type<ulong>(metal::clamp(_e126, _e127, _e128)));
long _e132 = val;
val = as_type<long>(as_type<ulong>(_e132) + as_type<ulong>(( + _e128.x * _e130.x + _e128.y * _e130.y)));
metal::long2 _e133 = metal::long2(_e132);
long _e134 = val;
long _e135 = val;
metal::long2 _e135 = metal::long2(_e134);
long _e137 = val;
val = as_type<long>(as_type<ulong>(_e137) + as_type<ulong>(metal::max(_e134, _e135)));
val = as_type<long>(as_type<ulong>(_e137) + as_type<ulong>(( + _e133.x * _e135.x + _e133.y * _e135.y)));
long _e139 = val;
long _e140 = val;
long _e142 = val;
val = as_type<long>(as_type<ulong>(_e142) + as_type<ulong>(metal::min(_e139, _e140)));
val = as_type<long>(as_type<ulong>(_e142) + as_type<ulong>(metal::max(_e139, _e140)));
long _e144 = val;
long _e146 = val;
val = as_type<long>(as_type<ulong>(_e146) + as_type<ulong>(metal::select(metal::select(long(-1), long(1), (_e144 > 0)), long(0), (_e144 == 0))));
long _e148 = val;
return _e148;
long _e145 = val;
long _e147 = val;
val = as_type<long>(as_type<ulong>(_e147) + as_type<ulong>(metal::min(_e144, _e145)));
long _e149 = val;
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;
}

ulong uint64_function(
Expand Down
Loading