nordic-dev.net/wgpu
nordic-dev.net/wgpu
2
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2025-04-16 22:16:36 +00:00
Code Issues Packages Projects Releases Wiki Activity

[naga msl-out hlsl-out] Ensure wrapper functions for integer division and modulo avoid using minimum value literals

Browse Source 
As we know that minimum value integer literals can cause problems for
some compilers. (See #7437)

Make the code which generates these functions call
msl::Writer::put_literal() and hlsl::Writer::write_literal()
respectively to output the minimum value integer literals instead of
just writing them directly, ensuring we only have to handle this
workaround in a single location (per backend).
This commit is contained in:
Jamie Nicol 2025-04-10 09:32:10 +01:00 committed by Erich Gubler
parent 6dc3b7186c
commit 08d3bb5196
8 changed files with 61 additions and 22 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

24
naga/src/back/hlsl/help.rs
View File

@ -1283,8 +1283,16 @@ impl<W: Write> super::Writer<'_, W> {
let level = crate::back::Level(1);
match scalar.kind {
ScalarKind::Sint => {
let min = -1i64 << (scalar.width as u32 * 8 - 1);
writeln!(self.out, "{level}return lhs / (((lhs == {min} & rhs == -1) | (rhs == 0)) ? 1 : rhs);")?
let min_val = match scalar.width {
4 => crate::Literal::I32(i32::MIN),
8 => crate::Literal::I64(i64::MIN),
_ => {
return Err(super::Error::UnsupportedScalar(scalar));
}
};
write!(self.out, "{level}return lhs / (((lhs == ")?;
self.write_literal(min_val)?;
writeln!(self.out, " & rhs == -1) | (rhs == 0)) ? 1 : rhs);")?
}
ScalarKind::Uint => {
writeln!(self.out, "{level}return lhs / (rhs == 0u ? 1u : rhs);")?
@ -1339,10 +1347,18 @@ impl<W: Write> super::Writer<'_, W> {
let level = crate::back::Level(1);
match scalar.kind {
ScalarKind::Sint => {
let min = -1i64 << (scalar.width as u32 * 8 - 1);
let min_val = match scalar.width {
4 => crate::Literal::I32(i32::MIN),
8 => crate::Literal::I64(i64::MIN),
_ => {
return Err(super::Error::UnsupportedScalar(scalar));
}
};
write!(self.out, "{level}")?;
self.write_value_type(module, right_ty)?;
writeln!(self.out, " divisor = ((lhs == {min} & rhs == -1) | (rhs == 0)) ? 1 : rhs;")?;
write!(self.out, " divisor = ((lhs == ")?;
self.write_literal(min_val)?;
writeln!(self.out, " & rhs == -1) | (rhs == 0)) ? 1 : rhs;")?;
writeln!(
self.out,
"{level}return lhs - (lhs / divisor) * divisor;"

31
naga/src/back/msl/writer.rs
View File

@ -5346,8 +5346,21 @@ template <typename A>
let level = back::Level(1);
match scalar.kind {
crate::ScalarKind::Sint => {
let min = -1i64 << (scalar.width as u32 * 8 - 1);
writeln!(self.out, "{level}return lhs / metal::select(rhs, 1, (lhs == {min} & rhs == -1) | (rhs == 0));")?
let min_val = match scalar.width {
4 => crate::Literal::I32(i32::MIN),
8 => crate::Literal::I64(i64::MIN),
_ => {
return Err(Error::GenericValidation(format!(
"Unexpected width for scalar {scalar:?}"
)));
}
};
write!(
self.out,
"{level}return lhs / metal::select(rhs, 1, (lhs == "
)?;
self.put_literal(min_val)?;
writeln!(self.out, " & rhs == -1) | (rhs == 0));")?
}
crate::ScalarKind::Uint => writeln!(
self.out,
@ -5415,8 +5428,18 @@ template <typename A>
let level = back::Level(1);
match scalar.kind {
crate::ScalarKind::Sint => {
let min = -1i64 << (scalar.width as u32 * 8 - 1);
writeln!(self.out, "{level}{rhs_type_name} divisor = metal::select(rhs, 1, (lhs == {min} & rhs == -1) | (rhs == 0));")?;
let min_val = match scalar.width {
4 => crate::Literal::I32(i32::MIN),
8 => crate::Literal::I64(i64::MIN),
_ => {
return Err(Error::GenericValidation(format!(
"Unexpected width for scalar {scalar:?}"
)));
}
};
write!(self.out, "{level}{rhs_type_name} divisor = metal::select(rhs, 1, (lhs == ")?;
self.put_literal(min_val)?;
writeln!(self.out, " & rhs == -1) | (rhs == 0));")?;
writeln!(
self.out,
"{level}return lhs - (lhs / divisor) * divisor;"

2
naga/tests/out/hlsl/wgsl-image.hlsl
View File

@ -25,7 +25,7 @@ Texture2DArray<float> image_2d_array_depth : register(t3, space1);
TextureCube<float> image_cube_depth : register(t4, space1);
int2 naga_mod(int2 lhs, int2 rhs) {
int2 divisor = ((lhs == -2147483648 & rhs == -1) | (rhs == 0)) ? 1 : rhs;
int2 divisor = ((lhs == int(-2147483647 - 1) & rhs == -1) | (rhs == 0)) ? 1 : rhs;
return lhs - (lhs / divisor) * divisor;
}

10
naga/tests/out/hlsl/wgsl-operators.hlsl
View File

@ -17,7 +17,7 @@ float4 builtins()
}
int4 naga_mod(int4 lhs, int4 rhs) {
int4 divisor = ((lhs == -2147483648 & rhs == -1) | (rhs == 0)) ? 1 : rhs;
int4 divisor = ((lhs == int(-2147483647 - 1) & rhs == -1) | (rhs == 0)) ? 1 : rhs;
return lhs - (lhs / divisor) * divisor;
}
@ -66,7 +66,7 @@ int2 naga_neg(int2 val) {
}
int naga_div(int lhs, int rhs) {
return lhs / (((lhs == -2147483648 & rhs == -1) | (rhs == 0)) ? 1 : rhs);
return lhs / (((lhs == int(-2147483647 - 1) & rhs == -1) | (rhs == 0)) ? 1 : rhs);
}
uint naga_div(uint lhs, uint rhs) {
@ -74,7 +74,7 @@ uint naga_div(uint lhs, uint rhs) {
}
int2 naga_div(int2 lhs, int2 rhs) {
return lhs / (((lhs == -2147483648 & rhs == -1) | (rhs == 0)) ? 1 : rhs);
return lhs / (((lhs == int(-2147483647 - 1) & rhs == -1) | (rhs == 0)) ? 1 : rhs);
}
uint3 naga_div(uint3 lhs, uint3 rhs) {
@ -82,7 +82,7 @@ uint3 naga_div(uint3 lhs, uint3 rhs) {
}
int naga_mod(int lhs, int rhs) {
int divisor = ((lhs == -2147483648 & rhs == -1) | (rhs == 0)) ? 1 : rhs;
int divisor = ((lhs == int(-2147483647 - 1) & rhs == -1) | (rhs == 0)) ? 1 : rhs;
return lhs - (lhs / divisor) * divisor;
}
@ -91,7 +91,7 @@ uint naga_mod(uint lhs, uint rhs) {
}
int2 naga_mod(int2 lhs, int2 rhs) {
int2 divisor = ((lhs == -2147483648 & rhs == -1) | (rhs == 0)) ? 1 : rhs;
int2 divisor = ((lhs == int(-2147483647 - 1) & rhs == -1) | (rhs == 0)) ? 1 : rhs;
return lhs - (lhs / divisor) * divisor;
}

2
naga/tests/out/hlsl/wgsl-skybox.hlsl
View File

@ -33,7 +33,7 @@ struct FragmentInput_fs_main {
};
int naga_div(int lhs, int rhs) {
return lhs / (((lhs == -2147483648 & rhs == -1) | (rhs == 0)) ? 1 : rhs);
return lhs / (((lhs == int(-2147483647 - 1) & rhs == -1) | (rhs == 0)) ? 1 : rhs);
}
VertexOutput ConstructVertexOutput(float4 arg0, float3 arg1) {

2
naga/tests/out/msl/wgsl-image.msl
View File

@ -5,7 +5,7 @@
using metal::uint;
metal::int2 naga_mod(metal::int2 lhs, metal::int2 rhs) {
metal::int2 divisor = metal::select(rhs, 1, (lhs == -2147483648 & rhs == -1) | (rhs == 0));
metal::int2 divisor = metal::select(rhs, 1, (lhs == (-2147483647 - 1) & rhs == -1) | (rhs == 0));
return lhs - (lhs / divisor) * divisor;
}

10
naga/tests/out/msl/wgsl-operators.msl
View File

@ -23,7 +23,7 @@ metal::float4 builtins(
}
metal::int4 naga_mod(metal::int4 lhs, metal::int4 rhs) {
metal::int4 divisor = metal::select(rhs, 1, (lhs == -2147483648 & rhs == -1) | (rhs == 0));
metal::int4 divisor = metal::select(rhs, 1, (lhs == (-2147483647 - 1) & rhs == -1) | (rhs == 0));
return lhs - (lhs / divisor) * divisor;
}
@ -74,7 +74,7 @@ metal::int2 naga_neg(metal::int2 val) {
}
int naga_div(int lhs, int rhs) {
return lhs / metal::select(rhs, 1, (lhs == -2147483648 & rhs == -1) | (rhs == 0));
return lhs / metal::select(rhs, 1, (lhs == (-2147483647 - 1) & rhs == -1) | (rhs == 0));
}
uint naga_div(uint lhs, uint rhs) {
@ -82,7 +82,7 @@ uint naga_div(uint lhs, uint rhs) {
}
metal::int2 naga_div(metal::int2 lhs, metal::int2 rhs) {
return lhs / metal::select(rhs, 1, (lhs == -2147483648 & rhs == -1) | (rhs == 0));
return lhs / metal::select(rhs, 1, (lhs == (-2147483647 - 1) & rhs == -1) | (rhs == 0));
}
metal::uint3 naga_div(metal::uint3 lhs, metal::uint3 rhs) {
@ -90,7 +90,7 @@ metal::uint3 naga_div(metal::uint3 lhs, metal::uint3 rhs) {
}
int naga_mod(int lhs, int rhs) {
int divisor = metal::select(rhs, 1, (lhs == -2147483648 & rhs == -1) | (rhs == 0));
int divisor = metal::select(rhs, 1, (lhs == (-2147483647 - 1) & rhs == -1) | (rhs == 0));
return lhs - (lhs / divisor) * divisor;
}
@ -99,7 +99,7 @@ uint naga_mod(uint lhs, uint rhs) {
}
metal::int2 naga_mod(metal::int2 lhs, metal::int2 rhs) {
metal::int2 divisor = metal::select(rhs, 1, (lhs == -2147483648 & rhs == -1) | (rhs == 0));
metal::int2 divisor = metal::select(rhs, 1, (lhs == (-2147483647 - 1) & rhs == -1) | (rhs == 0));
return lhs - (lhs / divisor) * divisor;
}

2
naga/tests/out/msl/wgsl-skybox.msl
View File

@ -13,7 +13,7 @@ struct Data {
metal::float4x4 view;
};
int naga_div(int lhs, int rhs) {
return lhs / metal::select(rhs, 1, (lhs == -2147483648 & rhs == -1) | (rhs == 0));
return lhs / metal::select(rhs, 1, (lhs == (-2147483647 - 1) & rhs == -1) | (rhs == 0));
}