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

jamienicol · Vecvec · commit 158b9bdc44f8 · 2025-04-11T10:00:53.000+12:00
As we know that minimum value integer literals can cause problems for some compilers. (See gfx-rs#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).
diff --git 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
@@ -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
@@ -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
@@ -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,23 +66,23 @@ 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) {
     return lhs / (rhs == 0u ? 1u : 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) {
     return lhs / (rhs == 0u ? 1u : 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
@@ -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
@@ -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
@@ -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,23 +74,23 @@ 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) {
     return lhs / metal::select(rhs, 1u, rhs == 0u);
 }
 
 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) {
     return lhs / metal::select(rhs, 1u, rhs == 0u);
 }
 
 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
@@ -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));
 }
 
 

Original file line number	Diff line number	Diff line change
`@@ -25,7 +25,7 @@ Texture2DArray<float> image_2d_array_depth : register(t3, space1);`
`25`	`25`	`TextureCube<float> image_cube_depth : register(t4, space1);`
`26`	`26`
`27`	`27`	`int2 naga_mod(int2 lhs, int2 rhs) {`
`28`		`- int2 divisor = ((lhs == -2147483648 & rhs == -1) \| (rhs == 0)) ? 1 : rhs;`
	`28`	`+ int2 divisor = ((lhs == int(-2147483647 - 1) & rhs == -1) \| (rhs == 0)) ? 1 : rhs;`
`29`	`29`	`return lhs - (lhs / divisor) * divisor;`
`30`	`30`	`}`
`31`	`31`
Original file line number	Diff line number	Diff line change
`@@ -17,7 +17,7 @@ float4 builtins()`
`17`	`17`	`}`
`18`	`18`
`19`	`19`	`int4 naga_mod(int4 lhs, int4 rhs) {`
`20`		`- int4 divisor = ((lhs == -2147483648 & rhs == -1) \| (rhs == 0)) ? 1 : rhs;`
	`20`	`+ int4 divisor = ((lhs == int(-2147483647 - 1) & rhs == -1) \| (rhs == 0)) ? 1 : rhs;`
`21`	`21`	`return lhs - (lhs / divisor) * divisor;`
`22`	`22`	`}`
`23`	`23`
`@@ -66,23 +66,23 @@ int2 naga_neg(int2 val) {`
`66`	`66`	`}`
`67`	`67`
`68`	`68`	`int naga_div(int lhs, int rhs) {`
`69`		`- return lhs / (((lhs == -2147483648 & rhs == -1) \| (rhs == 0)) ? 1 : rhs);`
	`69`	`+ return lhs / (((lhs == int(-2147483647 - 1) & rhs == -1) \| (rhs == 0)) ? 1 : rhs);`
`70`	`70`	`}`
`71`	`71`
`72`	`72`	`uint naga_div(uint lhs, uint rhs) {`
`73`	`73`	`return lhs / (rhs == 0u ? 1u : rhs);`
`74`	`74`	`}`
`75`	`75`
`76`	`76`	`int2 naga_div(int2 lhs, int2 rhs) {`
`77`		`- return lhs / (((lhs == -2147483648 & rhs == -1) \| (rhs == 0)) ? 1 : rhs);`
	`77`	`+ return lhs / (((lhs == int(-2147483647 - 1) & rhs == -1) \| (rhs == 0)) ? 1 : rhs);`
`78`	`78`	`}`
`79`	`79`
`80`	`80`	`uint3 naga_div(uint3 lhs, uint3 rhs) {`
`81`	`81`	`return lhs / (rhs == 0u ? 1u : rhs);`
`82`	`82`	`}`
`83`	`83`
`84`	`84`	`int naga_mod(int lhs, int rhs) {`
`85`		`- int divisor = ((lhs == -2147483648 & rhs == -1) \| (rhs == 0)) ? 1 : rhs;`
	`85`	`+ int divisor = ((lhs == int(-2147483647 - 1) & rhs == -1) \| (rhs == 0)) ? 1 : rhs;`
`86`	`86`	`return lhs - (lhs / divisor) * divisor;`
`87`	`87`	`}`
`88`	`88`
`@@ -91,7 +91,7 @@ uint naga_mod(uint lhs, uint rhs) {`
`91`	`91`	`}`
`92`	`92`
`93`	`93`	`int2 naga_mod(int2 lhs, int2 rhs) {`
`94`		`- int2 divisor = ((lhs == -2147483648 & rhs == -1) \| (rhs == 0)) ? 1 : rhs;`
	`94`	`+ int2 divisor = ((lhs == int(-2147483647 - 1) & rhs == -1) \| (rhs == 0)) ? 1 : rhs;`
`95`	`95`	`return lhs - (lhs / divisor) * divisor;`
`96`	`96`	`}`
`97`	`97`
Original file line number	Diff line number	Diff line change
`@@ -33,7 +33,7 @@ struct FragmentInput_fs_main {`
`33`	`33`	`};`
`34`	`34`
`35`	`35`	`int naga_div(int lhs, int rhs) {`
`36`		`- return lhs / (((lhs == -2147483648 & rhs == -1) \| (rhs == 0)) ? 1 : rhs);`
	`36`	`+ return lhs / (((lhs == int(-2147483647 - 1) & rhs == -1) \| (rhs == 0)) ? 1 : rhs);`
`37`	`37`	`}`
`38`	`38`
`39`	`39`	`VertexOutput ConstructVertexOutput(float4 arg0, float3 arg1) {`
Original file line number	Diff line number	Diff line change
`@@ -5,7 +5,7 @@`
`5`	`5`	`using metal::uint;`
`6`	`6`
`7`	`7`	`metal::int2 naga_mod(metal::int2 lhs, metal::int2 rhs) {`
`8`		`- metal::int2 divisor = metal::select(rhs, 1, (lhs == -2147483648 & rhs == -1) \| (rhs == 0));`
	`8`	`+ metal::int2 divisor = metal::select(rhs, 1, (lhs == (-2147483647 - 1) & rhs == -1) \| (rhs == 0));`
`9`	`9`	`return lhs - (lhs / divisor) * divisor;`
`10`	`10`	`}`
`11`	`11`
Original file line number	Diff line number	Diff line change
`@@ -23,7 +23,7 @@ metal::float4 builtins(`
`23`	`23`	`}`
`24`	`24`
`25`	`25`	`metal::int4 naga_mod(metal::int4 lhs, metal::int4 rhs) {`
`26`		`- metal::int4 divisor = metal::select(rhs, 1, (lhs == -2147483648 & rhs == -1) \| (rhs == 0));`
	`26`	`+ metal::int4 divisor = metal::select(rhs, 1, (lhs == (-2147483647 - 1) & rhs == -1) \| (rhs == 0));`
`27`	`27`	`return lhs - (lhs / divisor) * divisor;`
`28`	`28`	`}`
`29`	`29`
`@@ -74,23 +74,23 @@ metal::int2 naga_neg(metal::int2 val) {`
`74`	`74`	`}`
`75`	`75`
`76`	`76`	`int naga_div(int lhs, int rhs) {`
`77`		`- return lhs / metal::select(rhs, 1, (lhs == -2147483648 & rhs == -1) \| (rhs == 0));`
	`77`	`+ return lhs / metal::select(rhs, 1, (lhs == (-2147483647 - 1) & rhs == -1) \| (rhs == 0));`
`78`	`78`	`}`
`79`	`79`
`80`	`80`	`uint naga_div(uint lhs, uint rhs) {`
`81`	`81`	`return lhs / metal::select(rhs, 1u, rhs == 0u);`
`82`	`82`	`}`
`83`	`83`
`84`	`84`	`metal::int2 naga_div(metal::int2 lhs, metal::int2 rhs) {`
`85`		`- return lhs / metal::select(rhs, 1, (lhs == -2147483648 & rhs == -1) \| (rhs == 0));`
	`85`	`+ return lhs / metal::select(rhs, 1, (lhs == (-2147483647 - 1) & rhs == -1) \| (rhs == 0));`
`86`	`86`	`}`
`87`	`87`
`88`	`88`	`metal::uint3 naga_div(metal::uint3 lhs, metal::uint3 rhs) {`
`89`	`89`	`return lhs / metal::select(rhs, 1u, rhs == 0u);`
`90`	`90`	`}`
`91`	`91`
`92`	`92`	`int naga_mod(int lhs, int rhs) {`
`93`		`- int divisor = metal::select(rhs, 1, (lhs == -2147483648 & rhs == -1) \| (rhs == 0));`
	`93`	`+ int divisor = metal::select(rhs, 1, (lhs == (-2147483647 - 1) & rhs == -1) \| (rhs == 0));`
`94`	`94`	`return lhs - (lhs / divisor) * divisor;`
`95`	`95`	`}`
`96`	`96`
`@@ -99,7 +99,7 @@ uint naga_mod(uint lhs, uint rhs) {`
`99`	`99`	`}`
`100`	`100`
`101`	`101`	`metal::int2 naga_mod(metal::int2 lhs, metal::int2 rhs) {`
`102`		`- metal::int2 divisor = metal::select(rhs, 1, (lhs == -2147483648 & rhs == -1) \| (rhs == 0));`
	`102`	`+ metal::int2 divisor = metal::select(rhs, 1, (lhs == (-2147483647 - 1) & rhs == -1) \| (rhs == 0));`
`103`	`103`	`return lhs - (lhs / divisor) * divisor;`
`104`	`104`	`}`
`105`	`105`
Original file line number	Diff line number	Diff line change
`@@ -13,7 +13,7 @@ struct Data {`
`13`	`13`	`metal::float4x4 view;`
`14`	`14`	`};`
`15`	`15`	`int naga_div(int lhs, int rhs) {`
`16`		`- return lhs / metal::select(rhs, 1, (lhs == -2147483648 & rhs == -1) \| (rhs == 0));`
	`16`	`+ return lhs / metal::select(rhs, 1, (lhs == (-2147483647 - 1) & rhs == -1) \| (rhs == 0));`
`17`	`17`	`}`
`18`	`18`
`19`	`19`