diff --git a/naga/src/back/hlsl/help.rs b/naga/src/back/hlsl/help.rs
index 058d39692..782f73175 100644
--- a/naga/src/back/hlsl/help.rs
+++ b/naga/src/back/hlsl/help.rs
@@ -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;"
diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs
index 2903ad80e..7e037c43e 100644
--- a/naga/src/back/msl/writer.rs
+++ b/naga/src/back/msl/writer.rs
@@ -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;"
diff --git a/naga/tests/out/hlsl/wgsl-image.hlsl b/naga/tests/out/hlsl/wgsl-image.hlsl
index d41a5d945..a014d57e9 100644
--- a/naga/tests/out/hlsl/wgsl-image.hlsl
+++ b/naga/tests/out/hlsl/wgsl-image.hlsl
@@ -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;
}
diff --git a/naga/tests/out/hlsl/wgsl-operators.hlsl b/naga/tests/out/hlsl/wgsl-operators.hlsl
index 3dbec4345..c4e3578c2 100644
--- a/naga/tests/out/hlsl/wgsl-operators.hlsl
+++ b/naga/tests/out/hlsl/wgsl-operators.hlsl
@@ -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;
}
diff --git a/naga/tests/out/hlsl/wgsl-skybox.hlsl b/naga/tests/out/hlsl/wgsl-skybox.hlsl
index 9c890953b..9bcc7debf 100644
--- a/naga/tests/out/hlsl/wgsl-skybox.hlsl
+++ b/naga/tests/out/hlsl/wgsl-skybox.hlsl
@@ -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) {
diff --git a/naga/tests/out/msl/wgsl-image.msl b/naga/tests/out/msl/wgsl-image.msl
index da4c41063..bd55ec3f1 100644
--- a/naga/tests/out/msl/wgsl-image.msl
+++ b/naga/tests/out/msl/wgsl-image.msl
@@ -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;
}
diff --git a/naga/tests/out/msl/wgsl-operators.msl b/naga/tests/out/msl/wgsl-operators.msl
index 383c65713..4bb4e66b1 100644
--- a/naga/tests/out/msl/wgsl-operators.msl
+++ b/naga/tests/out/msl/wgsl-operators.msl
@@ -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;
}
diff --git a/naga/tests/out/msl/wgsl-skybox.msl b/naga/tests/out/msl/wgsl-skybox.msl
index b5f51c946..d0dbe5e07 100644
--- a/naga/tests/out/msl/wgsl-skybox.msl
+++ b/naga/tests/out/msl/wgsl-skybox.msl
@@ -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));
}