diff --git a/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGen.cs b/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGen.cs index 05fc3b2c8..f9601f62c 100644 --- a/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGen.cs +++ b/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGen.cs @@ -131,7 +131,8 @@ namespace Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions case Instruction.FSIEnd: return "|| FSI END ||"; case Instruction.GroupMemoryBarrier: - return "|| FIND GROUP MEMORY BARRIER ||"; + case Instruction.MemoryBarrier: + return MemoryBarrier(context, operation); case Instruction.ImageLoad: case Instruction.ImageStore: case Instruction.ImageAtomic: @@ -140,8 +141,6 @@ namespace Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions return Load(context, operation); case Instruction.Lod: return Lod(context, operation); - case Instruction.MemoryBarrier: - return "|| MEMORY BARRIER ||"; case Instruction.Store: return Store(context, operation); case Instruction.SwizzleAdd: diff --git a/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenMemory.cs b/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenMemory.cs index d13300e05..04f887d56 100644 --- a/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenMemory.cs +++ b/src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenMemory.cs @@ -600,6 +600,13 @@ namespace Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions return $"float2(as_type({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)}";