Skip to content

Commit e3828fc

Browse files
jamienicolErichDonGubler
jamienicol
authored and
ErichDonGubler
committed
[naga wgsl-out] Improve representation of minimum int64 literal
In #7424 we fixed a bug where the representation of the minimum int64 literal generated by naga was invalid WGSL. It changed us from expressing it as `-9223372036854775808` which was invalid, to `-9223372036854775807li - 1li`. This is valid WGSL. However, as the values are concrete i64 types if the shader is parsed again by naga the expression does not get const evaluated away, leading to suboptimal code generated by the backends. This patch makes us perform the subtraction using abstract integers before casting to i64, solving this problem. Additionally the input WGSL test is updated to use the same construct.
1 parent 4761e27 commit e3828fc

File tree

7 files changed

+236
-236
lines changed

7 files changed

+236
-236
lines changed

‎naga/src/back/wgsl/writer.rs‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1139,12 +1139,12 @@ impl<W: Write> Writer<W> {
11391139
crate::Literal::Bool(value) => write!(self.out, "{value}")?,
11401140
crate::Literal::F64(value) => write!(self.out, "{value:?}lf")?,
11411141
crate::Literal::I64(value) => {
1142-
// `-9223372036854775808li` is not valid WGSL. Nor can we use the AbstractInt
1143-
// trick above, as AbstractInt also cannot represent `9223372036854775808`.
1144-
// The most negative `i64` value can only be expressed in WGSL using
1145-
// subtracting 1 from the second most negative value.
1142+
// `-9223372036854775808li` is not valid WGSL. Nor can we simply use the
1143+
// AbstractInt trick above, as AbstractInt also cannot represent
1144+
// `9223372036854775808`. Instead construct the second most negative
1145+
// AbstractInt, subtract one from it, then cast to i64.
11461146
if value == i64::MIN {
1147-
write!(self.out, "{}li - 1li", value + 1)?;
1147+
write!(self.out, "i64({} - 1)", value + 1)?;
11481148
} else {
11491149
write!(self.out, "{value}li")?;
11501150
}

‎naga/tests/in/wgsl/int64.wgsl‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -56,7 +56,7 @@ fn int64_function(x: i64) -> i64 {
5656
val += bitcast<vec3<i64>>(input_uniform.val_u64_3).z;
5757
val += bitcast<vec4<i64>>(input_uniform.val_u64_4).w;
5858
// Most negative i64
59-
val += -9223372036854775807li - 1li;
59+
val += i64(-9223372036854775807 - 1);
6060

6161
// Reading/writing to a uniform/storage buffer
6262
output.val_i64 = input_uniform.val_i64 + input_storage.val_i64;

‎naga/tests/out/hlsl/wgsl-int64.hlsl‎

Lines changed: 28 additions & 28 deletions
Original file line numberDiff line numberDiff line change
@@ -99,51 +99,51 @@ int64_t int64_function(int64_t x)
9999
uint64_t4 _e71 = input_uniform.val_u64_4_;
100100
int64_t _e74 = val;
101101
val = (_e74 + _e71.w);
102-
int64_t _e79 = val;
103-
val = (_e79 + (-9223372036854775807L - 1L));
104-
int64_t _e85 = input_uniform.val_i64_;
105-
int64_t _e88 = input_storage.Load<int64_t>(128);
106-
output.Store(128, (_e85 + _e88));
107-
int64_t2 _e94 = input_uniform.val_i64_2_;
108-
int64_t2 _e97 = input_storage.Load<int64_t2>(144);
109-
output.Store(144, (_e94 + _e97));
110-
int64_t3 _e103 = input_uniform.val_i64_3_;
111-
int64_t3 _e106 = input_storage.Load<int64_t3>(160);
112-
output.Store(160, (_e103 + _e106));
113-
int64_t4 _e112 = input_uniform.val_i64_4_;
114-
int64_t4 _e115 = input_storage.Load<int64_t4>(192);
115-
output.Store(192, (_e112 + _e115));
116-
int64_t _e121[2] = Constructarray2_int64_t_(input_arrays.Load<int64_t>(16+0), input_arrays.Load<int64_t>(16+8));
102+
int64_t _e77 = val;
103+
val = (_e77 + -9223372036854775808L);
104+
int64_t _e83 = input_uniform.val_i64_;
105+
int64_t _e86 = input_storage.Load<int64_t>(128);
106+
output.Store(128, (_e83 + _e86));
107+
int64_t2 _e92 = input_uniform.val_i64_2_;
108+
int64_t2 _e95 = input_storage.Load<int64_t2>(144);
109+
output.Store(144, (_e92 + _e95));
110+
int64_t3 _e101 = input_uniform.val_i64_3_;
111+
int64_t3 _e104 = input_storage.Load<int64_t3>(160);
112+
output.Store(160, (_e101 + _e104));
113+
int64_t4 _e110 = input_uniform.val_i64_4_;
114+
int64_t4 _e113 = input_storage.Load<int64_t4>(192);
115+
output.Store(192, (_e110 + _e113));
116+
int64_t _e119[2] = Constructarray2_int64_t_(input_arrays.Load<int64_t>(16+0), input_arrays.Load<int64_t>(16+8));
117117
{
118-
int64_t _value2[2] = _e121;
118+
int64_t _value2[2] = _e119;
119119
output_arrays.Store(16+0, _value2[0]);
120120
output_arrays.Store(16+8, _value2[1]);
121121
}
122+
int64_t _e120 = val;
122123
int64_t _e122 = val;
124+
val = (_e122 + abs(_e120));
123125
int64_t _e124 = val;
124-
val = (_e124 + abs(_e122));
126+
int64_t _e125 = val;
125127
int64_t _e126 = val;
126-
int64_t _e127 = val;
127128
int64_t _e128 = val;
129+
val = (_e128 + clamp(_e124, _e125, _e126));
128130
int64_t _e130 = val;
129-
val = (_e130 + clamp(_e126, _e127, _e128));
130131
int64_t _e132 = val;
131-
int64_t _e134 = val;
132+
int64_t _e135 = val;
133+
val = (_e135 + dot((_e130).xx, (_e132).xx));
132134
int64_t _e137 = val;
133-
val = (_e137 + dot((_e132).xx, (_e134).xx));
134-
int64_t _e139 = val;
135+
int64_t _e138 = val;
135136
int64_t _e140 = val;
137+
val = (_e140 + max(_e137, _e138));
136138
int64_t _e142 = val;
137-
val = (_e142 + max(_e139, _e140));
138-
int64_t _e144 = val;
139+
int64_t _e143 = val;
139140
int64_t _e145 = val;
141+
val = (_e145 + min(_e142, _e143));
140142
int64_t _e147 = val;
141-
val = (_e147 + min(_e144, _e145));
142143
int64_t _e149 = val;
144+
val = (_e149 + sign(_e147));
143145
int64_t _e151 = val;
144-
val = (_e151 + sign(_e149));
145-
int64_t _e153 = val;
146-
return _e153;
146+
return _e151;
147147
}
148148

149149
uint64_t naga_f2u64(float value) {

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

Lines changed: 29 additions & 29 deletions
Original file line numberDiff line numberDiff line change
@@ -83,49 +83,49 @@ long int64_function(
8383
metal::ulong4 _e71 = input_uniform.val_u64_4_;
8484
long _e74 = val;
8585
val = as_type<long>(as_type<ulong>(_e74) + as_type<ulong>(as_type<metal::long4>(_e71).w));
86-
long _e79 = val;
87-
val = as_type<long>(as_type<ulong>(_e79) + as_type<ulong>(as_type<long>(as_type<ulong>(-9223372036854775807L) - as_type<ulong>(1L))));
88-
long _e85 = input_uniform.val_i64_;
89-
long _e88 = input_storage.val_i64_;
90-
output.val_i64_ = as_type<long>(as_type<ulong>(_e85) + as_type<ulong>(_e88));
91-
metal::long2 _e94 = input_uniform.val_i64_2_;
92-
metal::long2 _e97 = input_storage.val_i64_2_;
93-
output.val_i64_2_ = as_type<metal::long2>(as_type<metal::ulong2>(_e94) + as_type<metal::ulong2>(_e97));
94-
metal::long3 _e103 = input_uniform.val_i64_3_;
95-
metal::long3 _e106 = input_storage.val_i64_3_;
96-
output.val_i64_3_ = as_type<metal::long3>(as_type<metal::ulong3>(_e103) + as_type<metal::ulong3>(_e106));
97-
metal::long4 _e112 = input_uniform.val_i64_4_;
98-
metal::long4 _e115 = input_storage.val_i64_4_;
99-
output.val_i64_4_ = as_type<metal::long4>(as_type<metal::ulong4>(_e112) + as_type<metal::ulong4>(_e115));
100-
type_12 _e121 = input_arrays.val_i64_array_2_;
101-
output_arrays.val_i64_array_2_ = _e121;
86+
long _e77 = val;
87+
val = as_type<long>(as_type<ulong>(_e77) + as_type<ulong>(-9223372036854775808L));
88+
long _e83 = input_uniform.val_i64_;
89+
long _e86 = input_storage.val_i64_;
90+
output.val_i64_ = as_type<long>(as_type<ulong>(_e83) + as_type<ulong>(_e86));
91+
metal::long2 _e92 = input_uniform.val_i64_2_;
92+
metal::long2 _e95 = input_storage.val_i64_2_;
93+
output.val_i64_2_ = as_type<metal::long2>(as_type<metal::ulong2>(_e92) + as_type<metal::ulong2>(_e95));
94+
metal::long3 _e101 = input_uniform.val_i64_3_;
95+
metal::long3 _e104 = input_storage.val_i64_3_;
96+
output.val_i64_3_ = as_type<metal::long3>(as_type<metal::ulong3>(_e101) + as_type<metal::ulong3>(_e104));
97+
metal::long4 _e110 = input_uniform.val_i64_4_;
98+
metal::long4 _e113 = input_storage.val_i64_4_;
99+
output.val_i64_4_ = as_type<metal::long4>(as_type<metal::ulong4>(_e110) + as_type<metal::ulong4>(_e113));
100+
type_12 _e119 = input_arrays.val_i64_array_2_;
101+
output_arrays.val_i64_array_2_ = _e119;
102+
long _e120 = val;
102103
long _e122 = val;
104+
val = as_type<long>(as_type<ulong>(_e122) + as_type<ulong>(naga_abs(_e120)));
103105
long _e124 = val;
104-
val = as_type<long>(as_type<ulong>(_e124) + as_type<ulong>(naga_abs(_e122)));
106+
long _e125 = val;
105107
long _e126 = val;
106-
long _e127 = val;
107108
long _e128 = val;
109+
val = as_type<long>(as_type<ulong>(_e128) + as_type<ulong>(metal::clamp(_e124, _e125, _e126)));
108110
long _e130 = val;
109-
val = as_type<long>(as_type<ulong>(_e130) + as_type<ulong>(metal::clamp(_e126, _e127, _e128)));
111+
metal::long2 _e131 = metal::long2(_e130);
110112
long _e132 = val;
111113
metal::long2 _e133 = metal::long2(_e132);
112-
long _e134 = val;
113-
metal::long2 _e135 = metal::long2(_e134);
114+
long _e135 = val;
115+
val = as_type<long>(as_type<ulong>(_e135) + as_type<ulong>(( + _e131.x * _e133.x + _e131.y * _e133.y)));
114116
long _e137 = val;
115-
val = as_type<long>(as_type<ulong>(_e137) + as_type<ulong>(( + _e133.x * _e135.x + _e133.y * _e135.y)));
116-
long _e139 = val;
117+
long _e138 = val;
117118
long _e140 = val;
119+
val = as_type<long>(as_type<ulong>(_e140) + as_type<ulong>(metal::max(_e137, _e138)));
118120
long _e142 = val;
119-
val = as_type<long>(as_type<ulong>(_e142) + as_type<ulong>(metal::max(_e139, _e140)));
120-
long _e144 = val;
121+
long _e143 = val;
121122
long _e145 = val;
123+
val = as_type<long>(as_type<ulong>(_e145) + as_type<ulong>(metal::min(_e142, _e143)));
122124
long _e147 = val;
123-
val = as_type<long>(as_type<ulong>(_e147) + as_type<ulong>(metal::min(_e144, _e145)));
124125
long _e149 = val;
126+
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))));
125127
long _e151 = val;
126-
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))));
127-
long _e153 = val;
128-
return _e153;
128+
return _e151;
129129
}
130130

131131
ulong naga_f2u64(float value) {

0 commit comments

Comments
 (0)