ShuriZma
/
Ryujinx
mirror of https://github.com/Ryujinx/Ryujinx.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

Shader Memory Barriers 

Fixes some of the shader generation failures in Sonic Frontiers
This commit is contained in:
Isaac Marovitz 2024-07-24 15:23:16 +01:00
parent fdf5d49985
commit 61c7d3bb9c
No known key found for this signature in database
GPG Key ID: 97250B2B09A132E1
2 changed files with 9 additions and 3 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

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

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

@ -600,6 +600,13 @@ 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)}";