diff --git a/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGen.cs b/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGen.cs
index f9601f62c..887dfae4f 100644
--- a/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGen.cs
+++ b/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGen.cs
@@ -122,8 +122,6 @@ namespace Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions
                 {
                     case Instruction.Ballot:
                         return Ballot(context, operation);
-                    case Instruction.Barrier:
-                        return Barrier(context, operation);
                     case Instruction.Call:
                         return Call(context, operation);
                     case Instruction.FSIBegin:
@@ -132,7 +130,8 @@ namespace Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions
                         return "|| FSI END ||";
                     case Instruction.GroupMemoryBarrier:
                     case Instruction.MemoryBarrier:
-                        return MemoryBarrier(context, operation);
+                    case Instruction.Barrier:
+                        return Barrier(context, operation);
                     case Instruction.ImageLoad:
                     case Instruction.ImageStore:
                     case Instruction.ImageAtomic:
diff --git a/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenBarrier.cs b/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenBarrier.cs
index 7d681de26..77f05defc 100644
--- a/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenBarrier.cs
+++ b/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenBarrier.cs
@@ -1,8 +1,5 @@
+using Ryujinx.Graphics.Shader.IntermediateRepresentation;
 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
 {
@@ -10,7 +7,9 @@ namespace Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions
     {
         public static string Barrier(CodeGenContext context, AstOperation operation)
         {
-            return "threadgroup_barrier(mem_flags::mem_threadgroup)";
+            var device = (operation.Inst & Instruction.Mask) == Instruction.MemoryBarrier;
+
+            return $"threadgroup_barrier(mem_flags::mem_{(device ? "device" : "threadgroup")})";
         }
     }
 }
diff --git a/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenMemory.cs b/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenMemory.cs
index 04f887d56..d13300e05 100644
--- a/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenMemory.cs
+++ b/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenMemory.cs
@@ -600,13 +600,6 @@ namespace Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions
             return $"float2(as_type<half2>({srcExpr})){GetMask(operation.Index)}";
         }
 
-        public static string MemoryBarrier(CodeGenContext context, AstOperation operation)
-        {
-            var grouped = (operation.Inst & Instruction.Mask) == Instruction.GroupMemoryBarrier;
-
-            return $"threadgroup_barrier(mem_flags::mem_{(grouped ? "threadgroup" : "device")})";
-        }
-
         private static string GetMask(int index)
         {
             return $".{"xy".AsSpan(index, 1)}";