mirrors/Ryujinx
1
0
Fork 0
mirror of https://github.com/Ryujinx/Ryujinx.git synced 2024-12-18 21:02:00 +00:00
Code Issues Releases Activity

Consolodate barriers

Browse source
This commit is contained in:
Isaac Marovitz 2024-07-24 15:35:50 +01:00 committed by Isaac Marovitz
parent a31e461db8
commit 558752594c
3 changed files with 6 additions and 15 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

5
src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGen.cs
View file

@ -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:

9
src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenBarrier.cs
View file

@ -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")})";
}
}
}

7
src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGenMemory.cs
View file

@ -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)}";