Fix incompatible type issue with MSL output

gfx-rs · Dec 18, 2021 · b9162e4 · b9162e4
1 parent fb98408
commit b9162e4
Show file tree

Hide file tree

Showing 2 changed files with 43 additions and 20 deletions.
diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs
@@ -1121,6 +1121,21 @@ impl<W: Write> Writer<W> {
                     crate::TypeInner::Scalar { .. } => true,
                     _ => false,
                 };
+                let argument_size_suffix = match *context.resolve_type(arg) {
+                    crate::TypeInner::Vector {
+                        size: crate::VectorSize::Bi,
+                        ..
+                    } => "2",
+                    crate::TypeInner::Vector {
+                        size: crate::VectorSize::Tri,
+                        ..
+                    } => "3",
+                    crate::TypeInner::Vector {
+                        size: crate::VectorSize::Quad,
+                        ..
+                    } => "4",
+                    _ => "",
+                };
 
                 let fun_name = match fun {
                     // comparison
@@ -1207,13 +1222,21 @@ impl<W: Write> Writer<W> {
                     self.put_expression(arg1.unwrap(), context, false)?;
                     write!(self.out, ")")?;
                 } else if fun == Mf::FindLsb {
-                    write!(self.out, "((({}::ctz(", NAMESPACE)?;
+                    write!(
+                        self.out,
+                        "(((1 + int{}({}::ctz(",
+                        argument_size_suffix, NAMESPACE
+                    )?;
                     self.put_expression(arg, context, false)?;
-                    write!(self.out, ") + 1) % 33) - 1)")?;
+                    write!(self.out, "))) % 33) - 1)")?;
                 } else if fun == Mf::FindMsb {
-                    write!(self.out, "((({}::clz(", NAMESPACE)?;
+                    write!(
+                        self.out,
+                        "(((1 + int{}({}::clz(",
+                        argument_size_suffix, NAMESPACE
+                    )?;
                     self.put_expression(arg, context, false)?;
-                    write!(self.out, ") + 1) % 33) - 1)")?;
+                    write!(self.out, "))) % 33) - 1)")?;
                 } else if fun == Mf::Unpack2x16float {
                     write!(self.out, "float2(as_type<half2>(")?;
                     self.put_expression(arg, context, false)?;

diff --git a/tests/out/msl/bits.msl b/tests/out/msl/bits.msl
@@ -84,36 +84,36 @@ kernel void main_(
     metal::uint4 _e116 = u4_;
     u4_ = metal::extract_bits(_e116, 5u, 10u);
     int _e120 = i;
-    i = (((metal::ctz(_e120) + 1) % 33) - 1);
+    i = (((1 + int(metal::ctz(_e120))) % 33) - 1);
     metal::int2 _e122 = i2_;
-    i2_ = (((metal::ctz(_e122) + 1) % 33) - 1);
+    i2_ = (((1 + int2(metal::ctz(_e122))) % 33) - 1);
     metal::int3 _e124 = i3_;
-    i3_ = (((metal::ctz(_e124) + 1) % 33) - 1);
+    i3_ = (((1 + int3(metal::ctz(_e124))) % 33) - 1);
     metal::int4 _e126 = i4_;
-    i4_ = (((metal::ctz(_e126) + 1) % 33) - 1);
+    i4_ = (((1 + int4(metal::ctz(_e126))) % 33) - 1);
     metal::uint _e128 = u;
-    i = (((metal::ctz(_e128) + 1) % 33) - 1);
+    i = (((1 + int(metal::ctz(_e128))) % 33) - 1);
     metal::uint2 _e130 = u2_;
-    i2_ = (((metal::ctz(_e130) + 1) % 33) - 1);
+    i2_ = (((1 + int2(metal::ctz(_e130))) % 33) - 1);
     metal::uint3 _e132 = u3_;
-    i3_ = (((metal::ctz(_e132) + 1) % 33) - 1);
+    i3_ = (((1 + int3(metal::ctz(_e132))) % 33) - 1);
     metal::uint4 _e134 = u4_;
-    i4_ = (((metal::ctz(_e134) + 1) % 33) - 1);
+    i4_ = (((1 + int4(metal::ctz(_e134))) % 33) - 1);
     int _e136 = i;
-    i = (((metal::clz(_e136) + 1) % 33) - 1);
+    i = (((1 + int(metal::clz(_e136))) % 33) - 1);
     metal::int2 _e138 = i2_;
-    i2_ = (((metal::clz(_e138) + 1) % 33) - 1);
+    i2_ = (((1 + int2(metal::clz(_e138))) % 33) - 1);
     metal::int3 _e140 = i3_;
-    i3_ = (((metal::clz(_e140) + 1) % 33) - 1);
+    i3_ = (((1 + int3(metal::clz(_e140))) % 33) - 1);
     metal::int4 _e142 = i4_;
-    i4_ = (((metal::clz(_e142) + 1) % 33) - 1);
+    i4_ = (((1 + int4(metal::clz(_e142))) % 33) - 1);
     metal::uint _e144 = u;
-    i = (((metal::clz(_e144) + 1) % 33) - 1);
+    i = (((1 + int(metal::clz(_e144))) % 33) - 1);
     metal::uint2 _e146 = u2_;
-    i2_ = (((metal::clz(_e146) + 1) % 33) - 1);
+    i2_ = (((1 + int2(metal::clz(_e146))) % 33) - 1);
     metal::uint3 _e148 = u3_;
-    i3_ = (((metal::clz(_e148) + 1) % 33) - 1);
+    i3_ = (((1 + int3(metal::clz(_e148))) % 33) - 1);
     metal::uint4 _e150 = u4_;
-    i4_ = (((metal::clz(_e150) + 1) % 33) - 1);
+    i4_ = (((1 + int4(metal::clz(_e150))) % 33) - 1);
     return;
 }