emu-mirrors/Ryujinx
0
0
Fork 0
mirror of https://github.com/GreemDev/Ryujinx.git synced 2024-12-23 05:45:47 +00:00
Code Wiki Activity

Shader Memory Barriers

Browse source 
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 committed by Isaac Marovitz
parent 650f309b58
commit a31e461db8
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)}";