Fix incompatible type issue with MSL output

gfx-rs · Nov 30, 2021 · 1550ba5 · 1550ba5
1 parent fb98408
commit 1550ba5
Show file tree

Hide file tree

Showing 2 changed files with 20 additions and 20 deletions.
diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs
@@ -1207,13 +1207,13 @@ 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, "(((int(1) + {}::ctz(", 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, "(((int(1) + {}::clz(", 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 = (((int(1) + metal::ctz(_e120)) % 33) - 1);
     metal::int2 _e122 = i2_;
-    i2_ = (((metal::ctz(_e122) + 1) % 33) - 1);
+    i2_ = (((int(1) + metal::ctz(_e122)) % 33) - 1);
     metal::int3 _e124 = i3_;
-    i3_ = (((metal::ctz(_e124) + 1) % 33) - 1);
+    i3_ = (((int(1) + metal::ctz(_e124)) % 33) - 1);
     metal::int4 _e126 = i4_;
-    i4_ = (((metal::ctz(_e126) + 1) % 33) - 1);
+    i4_ = (((int(1) + metal::ctz(_e126)) % 33) - 1);
     metal::uint _e128 = u;
-    i = (((metal::ctz(_e128) + 1) % 33) - 1);
+    i = (((int(1) + metal::ctz(_e128)) % 33) - 1);
     metal::uint2 _e130 = u2_;
-    i2_ = (((metal::ctz(_e130) + 1) % 33) - 1);
+    i2_ = (((int(1) + metal::ctz(_e130)) % 33) - 1);
     metal::uint3 _e132 = u3_;
-    i3_ = (((metal::ctz(_e132) + 1) % 33) - 1);
+    i3_ = (((int(1) + metal::ctz(_e132)) % 33) - 1);
     metal::uint4 _e134 = u4_;
-    i4_ = (((metal::ctz(_e134) + 1) % 33) - 1);
+    i4_ = (((int(1) + metal::ctz(_e134)) % 33) - 1);
     int _e136 = i;
-    i = (((metal::clz(_e136) + 1) % 33) - 1);
+    i = (((int(1) + metal::clz(_e136)) % 33) - 1);
     metal::int2 _e138 = i2_;
-    i2_ = (((metal::clz(_e138) + 1) % 33) - 1);
+    i2_ = (((int(1) + metal::clz(_e138)) % 33) - 1);
     metal::int3 _e140 = i3_;
-    i3_ = (((metal::clz(_e140) + 1) % 33) - 1);
+    i3_ = (((int(1) + metal::clz(_e140)) % 33) - 1);
     metal::int4 _e142 = i4_;
-    i4_ = (((metal::clz(_e142) + 1) % 33) - 1);
+    i4_ = (((int(1) + metal::clz(_e142)) % 33) - 1);
     metal::uint _e144 = u;
-    i = (((metal::clz(_e144) + 1) % 33) - 1);
+    i = (((int(1) + metal::clz(_e144)) % 33) - 1);
     metal::uint2 _e146 = u2_;
-    i2_ = (((metal::clz(_e146) + 1) % 33) - 1);
+    i2_ = (((int(1) + metal::clz(_e146)) % 33) - 1);
     metal::uint3 _e148 = u3_;
-    i3_ = (((metal::clz(_e148) + 1) % 33) - 1);
+    i3_ = (((int(1) + metal::clz(_e148)) % 33) - 1);
     metal::uint4 _e150 = u4_;
-    i4_ = (((metal::clz(_e150) + 1) % 33) - 1);
+    i4_ = (((int(1) + metal::clz(_e150)) % 33) - 1);
     return;
 }