From ad43ad7d8116a7b5182531e96215064f0d27d703 Mon Sep 17 00:00:00 2001
From: Isaac Marovitz <isaacryu@icloud.com>
Date: Sat, 22 Jun 2024 12:52:16 +0100
Subject: [PATCH] Fix simd_ballot

---
 .../CodeGen/Msl/Instructions/InstGen.cs       |  3 +++
 .../CodeGen/Msl/Instructions/InstGenBallot.cs | 21 +++++++++++++++++++
 .../CodeGen/Msl/Instructions/InstGenHelper.cs |  2 +-
 3 files changed, 25 insertions(+), 1 deletion(-)
 create mode 100644 src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenBallot.cs

diff --git a/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGen.cs b/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGen.cs
index 696564992d..7626f94aae 100644
--- a/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGen.cs
+++ b/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGen.cs
@@ -3,6 +3,7 @@ using Ryujinx.Graphics.Shader.StructuredIr;
 using Ryujinx.Graphics.Shader.Translation;
 using System;
 using System.Text;
+using static Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions.InstGenBallot;
 using static Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions.InstGenCall;
 using static Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions.InstGenHelper;
 using static Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions.InstGenMemory;
@@ -118,6 +119,8 @@ namespace Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions
             {
                 switch (inst & Instruction.Mask)
                 {
+                    case Instruction.Ballot:
+                        return Ballot(context, operation);
                     case Instruction.Barrier:
                         return "threadgroup_barrier(mem_flags::mem_threadgroup)";
                     case Instruction.Call:
diff --git a/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenBallot.cs b/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenBallot.cs
new file mode 100644
index 0000000000..1f53c74ed3
--- /dev/null
+++ b/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenBallot.cs
@@ -0,0 +1,21 @@
+using Ryujinx.Graphics.Shader.StructuredIr;
+using Ryujinx.Graphics.Shader.Translation;
+
+using static Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions.InstGenHelper;
+using static Ryujinx.Graphics.Shader.StructuredIr.InstructionInfo;
+
+namespace Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions
+{
+    static class InstGenBallot
+    {
+        public static string Ballot(CodeGenContext context, AstOperation operation)
+        {
+            AggregateType dstType = GetSrcVarType(operation.Inst, 0);
+
+            string arg = GetSourceExpr(context, operation.GetSource(0), dstType);
+            char component = "xyzw"[operation.Index];
+
+            return $"uint4(as_type<uint2>((simd_vote::vote_t)simd_ballot({arg})), 0, 0).{component}";
+        }
+    }
+}
diff --git a/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenHelper.cs b/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenHelper.cs
index 014d070efe..983441e59c 100644
--- a/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenHelper.cs
+++ b/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenHelper.cs
@@ -25,7 +25,7 @@ namespace Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions
             Add(Instruction.AtomicXor,                InstType.AtomicBinary,   "atomic_xor_explicit");
             Add(Instruction.Absolute,                 InstType.CallUnary,      "abs");
             Add(Instruction.Add,                      InstType.OpBinaryCom,    "+",  2);
-            Add(Instruction.Ballot,                   InstType.CallUnary,      "simd_ballot");
+            Add(Instruction.Ballot,                   InstType.Special);
             Add(Instruction.Barrier,                  InstType.Special);
             Add(Instruction.BitCount,                 InstType.CallUnary,      "popcount");
             Add(Instruction.BitfieldExtractS32,       InstType.CallTernary,    "extract_bits");