a

fornwall · Sep 5, 2023 · e1bf6df · e1bf6df
1 parent c69e5c1
commit e1bf6df
Show file tree

Hide file tree

Showing 43 changed files with 90 additions and 477 deletions.
diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs
@@ -1720,7 +1720,7 @@ impl<W: Write> Writer<W> {
                             kind: crate::ScalarKind::Sint,
                             ..
                         }
-                        | crate::TypeInner::Vector {
+                        | &crate::TypeInner::Vector {
                             kind: crate::ScalarKind::Sint,
                             ..
                         } => NAGA_ISIGN_FUNCTION,
@@ -4118,21 +4118,12 @@ impl<W: Write> Writer<W> {
 
     fn write_polyfills(&mut self) -> Result<(), Error> {
         writeln!(self.out)?;
-        for size in 1..5 {
-            let tmp;
-            let type_name = if size == 1 {
-                "int"
-            } else {
-                tmp = format!("int{size}");
-                &tmp
-            };
-            writeln!(
-                self.out,
-                "{type_name} {NAGA_ISIGN_FUNCTION}({type_name} arg) {{
-    return {NAMESPACE}::select({NAMESPACE}::select({type_name}(-1), {type_name}(1), (arg > 0)), 0, (arg == 0));
+        writeln!(
+            self.out,
+            "template <typename T> inline T {NAGA_ISIGN_FUNCTION}(T arg) {{
+    return {NAMESPACE}::select({NAMESPACE}::select(T(-1), T(1), (arg > 0)), 0, (arg == 0));
 }}"
-            )?;
-        }
+        )?;
         Ok(())
     }
 

diff --git a/tests/out/msl/access.msl b/tests/out/msl/access.msl
@@ -58,17 +58,8 @@ struct type_26 {
     metal::float4 inner[2];
 };
 
-int naga_isign(int arg) {
-    return metal::select(metal::select(int(-1), int(1), (arg > 0)), 0, (arg == 0));
-}
-int2 naga_isign(int2 arg) {
-    return metal::select(metal::select(int2(-1), int2(1), (arg > 0)), 0, (arg == 0));
-}
-int3 naga_isign(int3 arg) {
-    return metal::select(metal::select(int3(-1), int3(1), (arg > 0)), 0, (arg == 0));
-}
-int4 naga_isign(int4 arg) {
-    return metal::select(metal::select(int4(-1), int4(1), (arg > 0)), 0, (arg == 0));
+template <typename T> inline T naga_isign(T arg) {
+    return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0));
 }
 
 void test_matrix_within_struct_accesses(

diff --git a/tests/out/msl/array-in-ctor.msl b/tests/out/msl/array-in-ctor.msl
@@ -11,17 +11,8 @@ struct Ah {
     type_1 inner;
 };
 
-int naga_isign(int arg) {
-    return metal::select(metal::select(int(-1), int(1), (arg > 0)), 0, (arg == 0));
-}
-int2 naga_isign(int2 arg) {
-    return metal::select(metal::select(int2(-1), int2(1), (arg > 0)), 0, (arg == 0));
-}
-int3 naga_isign(int3 arg) {
-    return metal::select(metal::select(int3(-1), int3(1), (arg > 0)), 0, (arg == 0));
-}
-int4 naga_isign(int4 arg) {
-    return metal::select(metal::select(int4(-1), int4(1), (arg > 0)), 0, (arg == 0));
+template <typename T> inline T naga_isign(T arg) {
+    return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0));
 }
 
 kernel void cs_main(

diff --git a/tests/out/msl/array-in-function-return-type.msl b/tests/out/msl/array-in-function-return-type.msl
@@ -8,17 +8,8 @@ struct type_1 {
     float inner[2];
 };
 
-int naga_isign(int arg) {
-    return metal::select(metal::select(int(-1), int(1), (arg > 0)), 0, (arg == 0));
-}
-int2 naga_isign(int2 arg) {
-    return metal::select(metal::select(int2(-1), int2(1), (arg > 0)), 0, (arg == 0));
-}
-int3 naga_isign(int3 arg) {
-    return metal::select(metal::select(int3(-1), int3(1), (arg > 0)), 0, (arg == 0));
-}
-int4 naga_isign(int4 arg) {
-    return metal::select(metal::select(int4(-1), int4(1), (arg > 0)), 0, (arg == 0));
+template <typename T> inline T naga_isign(T arg) {
+    return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0));
 }
 
 type_1 ret_array(

diff --git a/tests/out/msl/atomicOps.msl b/tests/out/msl/atomicOps.msl
@@ -12,17 +12,8 @@ struct Struct {
     type_2 atomic_arr;
 };
 
-int naga_isign(int arg) {
-    return metal::select(metal::select(int(-1), int(1), (arg > 0)), 0, (arg == 0));
-}
-int2 naga_isign(int2 arg) {
-    return metal::select(metal::select(int2(-1), int2(1), (arg > 0)), 0, (arg == 0));
-}
-int3 naga_isign(int3 arg) {
-    return metal::select(metal::select(int3(-1), int3(1), (arg > 0)), 0, (arg == 0));
-}
-int4 naga_isign(int4 arg) {
-    return metal::select(metal::select(int4(-1), int4(1), (arg > 0)), 0, (arg == 0));
+template <typename T> inline T naga_isign(T arg) {
+    return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0));
 }
 
 struct cs_mainInput {

diff --git a/tests/out/msl/binding-arrays.msl b/tests/out/msl/binding-arrays.msl
@@ -17,17 +17,8 @@ struct FragmentIn {
     uint index;
 };
 
-int naga_isign(int arg) {
-    return metal::select(metal::select(int(-1), int(1), (arg > 0)), 0, (arg == 0));
-}
-int2 naga_isign(int2 arg) {
-    return metal::select(metal::select(int2(-1), int2(1), (arg > 0)), 0, (arg == 0));
-}
-int3 naga_isign(int3 arg) {
-    return metal::select(metal::select(int3(-1), int3(1), (arg > 0)), 0, (arg == 0));
-}
-int4 naga_isign(int4 arg) {
-    return metal::select(metal::select(int4(-1), int4(1), (arg > 0)), 0, (arg == 0));
+template <typename T> inline T naga_isign(T arg) {
+    return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0));
 }
 
 struct main_Input {

diff --git a/tests/out/msl/bitcast.msl b/tests/out/msl/bitcast.msl
@@ -5,17 +5,8 @@
 using metal::uint;
 
 
-int naga_isign(int arg) {
-    return metal::select(metal::select(int(-1), int(1), (arg > 0)), 0, (arg == 0));
-}
-int2 naga_isign(int2 arg) {
-    return metal::select(metal::select(int2(-1), int2(1), (arg > 0)), 0, (arg == 0));
-}
-int3 naga_isign(int3 arg) {
-    return metal::select(metal::select(int3(-1), int3(1), (arg > 0)), 0, (arg == 0));
-}
-int4 naga_isign(int4 arg) {
-    return metal::select(metal::select(int4(-1), int4(1), (arg > 0)), 0, (arg == 0));
+template <typename T> inline T naga_isign(T arg) {
+    return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0));
 }
 
 kernel void main_(

diff --git a/tests/out/msl/bits.msl b/tests/out/msl/bits.msl
@@ -5,17 +5,8 @@
 using metal::uint;
 
 
-int naga_isign(int arg) {
-    return metal::select(metal::select(int(-1), int(1), (arg > 0)), 0, (arg == 0));
-}
-int2 naga_isign(int2 arg) {
-    return metal::select(metal::select(int2(-1), int2(1), (arg > 0)), 0, (arg == 0));
-}
-int3 naga_isign(int3 arg) {
-    return metal::select(metal::select(int3(-1), int3(1), (arg > 0)), 0, (arg == 0));
-}
-int4 naga_isign(int4 arg) {
-    return metal::select(metal::select(int4(-1), int4(1), (arg > 0)), 0, (arg == 0));
+template <typename T> inline T naga_isign(T arg) {
+    return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0));
 }
 
 kernel void main_(

diff --git a/tests/out/msl/boids.msl b/tests/out/msl/boids.msl
@@ -28,17 +28,8 @@ struct Particles {
 };
 constant uint NUM_PARTICLES = 1500u;
 
-int naga_isign(int arg) {
-    return metal::select(metal::select(int(-1), int(1), (arg > 0)), 0, (arg == 0));
-}
-int2 naga_isign(int2 arg) {
-    return metal::select(metal::select(int2(-1), int2(1), (arg > 0)), 0, (arg == 0));
-}
-int3 naga_isign(int3 arg) {
-    return metal::select(metal::select(int3(-1), int3(1), (arg > 0)), 0, (arg == 0));
-}
-int4 naga_isign(int4 arg) {
-    return metal::select(metal::select(int4(-1), int4(1), (arg > 0)), 0, (arg == 0));
+template <typename T> inline T naga_isign(T arg) {
+    return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0));
 }
 
 struct main_Input {

diff --git a/tests/out/msl/bounds-check-image-restrict.msl b/tests/out/msl/bounds-check-image-restrict.msl
@@ -5,17 +5,8 @@
 using metal::uint;
 
 
-int naga_isign(int arg) {
-    return metal::select(metal::select(int(-1), int(1), (arg > 0)), 0, (arg == 0));
-}
-int2 naga_isign(int2 arg) {
-    return metal::select(metal::select(int2(-1), int2(1), (arg > 0)), 0, (arg == 0));
-}
-int3 naga_isign(int3 arg) {
-    return metal::select(metal::select(int3(-1), int3(1), (arg > 0)), 0, (arg == 0));
-}
-int4 naga_isign(int4 arg) {
-    return metal::select(metal::select(int4(-1), int4(1), (arg > 0)), 0, (arg == 0));
+template <typename T> inline T naga_isign(T arg) {
+    return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0));
 }
 
 metal::float4 test_textureLoad_1d(

diff --git a/tests/out/msl/bounds-check-image-rzsw.msl b/tests/out/msl/bounds-check-image-rzsw.msl
@@ -11,17 +11,8 @@ struct DefaultConstructible {
 };
 
 
-int naga_isign(int arg) {
-    return metal::select(metal::select(int(-1), int(1), (arg > 0)), 0, (arg == 0));
-}
-int2 naga_isign(int2 arg) {
-    return metal::select(metal::select(int2(-1), int2(1), (arg > 0)), 0, (arg == 0));
-}
-int3 naga_isign(int3 arg) {
-    return metal::select(metal::select(int3(-1), int3(1), (arg > 0)), 0, (arg == 0));
-}
-int4 naga_isign(int4 arg) {
-    return metal::select(metal::select(int4(-1), int4(1), (arg > 0)), 0, (arg == 0));
+template <typename T> inline T naga_isign(T arg) {
+    return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0));
 }
 
 metal::float4 test_textureLoad_1d(

diff --git a/tests/out/msl/bounds-check-restrict.msl b/tests/out/msl/bounds-check-restrict.msl
@@ -20,17 +20,8 @@ struct Globals {
     type_4 d;
 };
 
-int naga_isign(int arg) {
-    return metal::select(metal::select(int(-1), int(1), (arg > 0)), 0, (arg == 0));
-}
-int2 naga_isign(int2 arg) {
-    return metal::select(metal::select(int2(-1), int2(1), (arg > 0)), 0, (arg == 0));
-}
-int3 naga_isign(int3 arg) {
-    return metal::select(metal::select(int3(-1), int3(1), (arg > 0)), 0, (arg == 0));
-}
-int4 naga_isign(int4 arg) {
-    return metal::select(metal::select(int4(-1), int4(1), (arg > 0)), 0, (arg == 0));
+template <typename T> inline T naga_isign(T arg) {
+    return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0));
 }
 
 float index_array(

diff --git a/tests/out/msl/bounds-check-zero-atomic.msl b/tests/out/msl/bounds-check-zero-atomic.msl
@@ -24,17 +24,8 @@ struct Globals {
     type_2 c;
 };
 
-int naga_isign(int arg) {
-    return metal::select(metal::select(int(-1), int(1), (arg > 0)), 0, (arg == 0));
-}
-int2 naga_isign(int2 arg) {
-    return metal::select(metal::select(int2(-1), int2(1), (arg > 0)), 0, (arg == 0));
-}
-int3 naga_isign(int3 arg) {
-    return metal::select(metal::select(int3(-1), int3(1), (arg > 0)), 0, (arg == 0));
-}
-int4 naga_isign(int4 arg) {
-    return metal::select(metal::select(int4(-1), int4(1), (arg > 0)), 0, (arg == 0));
+template <typename T> inline T naga_isign(T arg) {
+    return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0));
 }
 
 uint fetch_add_atomic(

diff --git a/tests/out/msl/bounds-check-zero.msl b/tests/out/msl/bounds-check-zero.msl
@@ -26,17 +26,8 @@ struct Globals {
     type_4 d;
 };
 
-int naga_isign(int arg) {
-    return metal::select(metal::select(int(-1), int(1), (arg > 0)), 0, (arg == 0));
-}
-int2 naga_isign(int2 arg) {
-    return metal::select(metal::select(int2(-1), int2(1), (arg > 0)), 0, (arg == 0));
-}
-int3 naga_isign(int3 arg) {
-    return metal::select(metal::select(int3(-1), int3(1), (arg > 0)), 0, (arg == 0));
-}
-int4 naga_isign(int4 arg) {
-    return metal::select(metal::select(int4(-1), int4(1), (arg > 0)), 0, (arg == 0));
+template <typename T> inline T naga_isign(T arg) {
+    return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0));
 }
 
 float index_array(

diff --git a/tests/out/msl/break-if.msl b/tests/out/msl/break-if.msl
@@ -5,17 +5,8 @@
 using metal::uint;
 
 
-int naga_isign(int arg) {
-    return metal::select(metal::select(int(-1), int(1), (arg > 0)), 0, (arg == 0));
-}
-int2 naga_isign(int2 arg) {
-    return metal::select(metal::select(int2(-1), int2(1), (arg > 0)), 0, (arg == 0));
-}
-int3 naga_isign(int3 arg) {
-    return metal::select(metal::select(int3(-1), int3(1), (arg > 0)), 0, (arg == 0));
-}
-int4 naga_isign(int4 arg) {
-    return metal::select(metal::select(int4(-1), int4(1), (arg > 0)), 0, (arg == 0));
+template <typename T> inline T naga_isign(T arg) {
+    return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0));
 }
 
 void breakIfEmpty(

diff --git a/tests/out/msl/collatz.msl b/tests/out/msl/collatz.msl
@@ -13,17 +13,8 @@ struct PrimeIndices {
     type_1 data;
 };
 
-int naga_isign(int arg) {
-    return metal::select(metal::select(int(-1), int(1), (arg > 0)), 0, (arg == 0));
-}
-int2 naga_isign(int2 arg) {
-    return metal::select(metal::select(int2(-1), int2(1), (arg > 0)), 0, (arg == 0));
-}
-int3 naga_isign(int3 arg) {
-    return metal::select(metal::select(int3(-1), int3(1), (arg > 0)), 0, (arg == 0));
-}
-int4 naga_isign(int4 arg) {
-    return metal::select(metal::select(int4(-1), int4(1), (arg > 0)), 0, (arg == 0));
+template <typename T> inline T naga_isign(T arg) {
+    return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0));
 }
 
 uint collatz_iterations(

diff --git a/tests/out/msl/constructors.msl b/tests/out/msl/constructors.msl
@@ -30,17 +30,8 @@ constant type_10 cz6_ = type_10 {};
 constant Foo cz7_ = Foo {};
 constant type_11 cp3_ = type_11 {0, 1, 2, 3};
 
-int naga_isign(int arg) {
-    return metal::select(metal::select(int(-1), int(1), (arg > 0)), 0, (arg == 0));
-}
-int2 naga_isign(int2 arg) {
-    return metal::select(metal::select(int2(-1), int2(1), (arg > 0)), 0, (arg == 0));
-}
-int3 naga_isign(int3 arg) {
-    return metal::select(metal::select(int3(-1), int3(1), (arg > 0)), 0, (arg == 0));
-}
-int4 naga_isign(int4 arg) {
-    return metal::select(metal::select(int4(-1), int4(1), (arg > 0)), 0, (arg == 0));
+template <typename T> inline T naga_isign(T arg) {
+    return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0));
 }
 
 kernel void main_(

diff --git a/tests/out/msl/control-flow.msl b/tests/out/msl/control-flow.msl
@@ -5,17 +5,8 @@
 using metal::uint;
 
 
-int naga_isign(int arg) {
-    return metal::select(metal::select(int(-1), int(1), (arg > 0)), 0, (arg == 0));
-}
-int2 naga_isign(int2 arg) {
-    return metal::select(metal::select(int2(-1), int2(1), (arg > 0)), 0, (arg == 0));
-}
-int3 naga_isign(int3 arg) {
-    return metal::select(metal::select(int3(-1), int3(1), (arg > 0)), 0, (arg == 0));
-}
-int4 naga_isign(int4 arg) {
-    return metal::select(metal::select(int4(-1), int4(1), (arg > 0)), 0, (arg == 0));
+template <typename T> inline T naga_isign(T arg) {
+    return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0));
 }
 
 void switch_default_break(

diff --git a/tests/out/msl/do-while.msl b/tests/out/msl/do-while.msl
@@ -5,17 +5,8 @@
 using metal::uint;
 
 
-int naga_isign(int arg) {
-    return metal::select(metal::select(int(-1), int(1), (arg > 0)), 0, (arg == 0));
-}
-int2 naga_isign(int2 arg) {
-    return metal::select(metal::select(int2(-1), int2(1), (arg > 0)), 0, (arg == 0));
-}
-int3 naga_isign(int3 arg) {
-    return metal::select(metal::select(int3(-1), int3(1), (arg > 0)), 0, (arg == 0));
-}
-int4 naga_isign(int4 arg) {
-    return metal::select(metal::select(int4(-1), int4(1), (arg > 0)), 0, (arg == 0));
+template <typename T> inline T naga_isign(T arg) {
+    return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0));
 }
 
 void fb1_(