owen/Ryujinx
1
0
Fork 0
mirror of https://github.com/Ryubing/Ryujinx.git synced 2025-11-21 18:40:56 -05:00
Code Issues Packages Projects Releases Wiki Activity

Instruction.Barrier

Browse Source 
Whoops

Fix inline functions in compute stage

Fix regression

Declare SharedMemories + Only Declare Memories on Main Func

Lowecase struct

Avoid magic strings

Make function signatures readable

Change how unsized arrays are indexed

Use string builder

Fix shuffle instructions

Cleanup NumberFormater

Bunch of Subgroup I/O Vars

Will probably need further refinement

Fix point_coord type

Fix support buffer declaration

Fix point_coord
This commit is contained in:
Isaac Marovitz
2024-06-21 10:31:21 +01:00
committed by Evan Husted
parent d0e4adac36
commit becf828d0a
10 changed files with 110 additions and 77 deletions
Download Patch File Download Diff File Expand all files Collapse all files

25
src/Ryujinx.Graphics.Shader/CodeGen/Msl/Instructions/InstGen.cs
View File

@@ -2,7 +2,7 @@ using Ryujinx.Graphics.Shader.IntermediateRepresentation;
using Ryujinx.Graphics.Shader.StructuredIr;
using Ryujinx.Graphics.Shader.Translation;
using System;
using System.Text;
using static Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions.InstGenCall;
using static Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions.InstGenHelper;
using static Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions.InstGenMemory;
@@ -39,11 +39,20 @@ namespace Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions
int arity = (int)(info.Type & InstType.ArityMask);
string args = string.Empty;
StringBuilder builder = new();
if (atomic)
if (atomic && (operation.StorageKind == StorageKind.StorageBuffer || operation.StorageKind == StorageKind.SharedMemory))
{
// Hell
builder.Append(GenerateLoadOrStore(context, operation, isStore: false));
AggregateType dstType = operation.Inst == Instruction.AtomicMaxS32 || operation.Inst == Instruction.AtomicMinS32
? AggregateType.S32
: AggregateType.U32;
for (int argIndex = operation.SourcesCount - arity + 2; argIndex < operation.SourcesCount; argIndex++)
{
builder.Append($", {GetSourceExpr(context, operation.GetSource(argIndex), dstType)}");
}
}
else
{
@@ -51,16 +60,16 @@ namespace Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions
{
if (argIndex != 0)
{
args += ", ";
builder.Append(", ");
}
AggregateType dstType = GetSrcVarType(inst, argIndex);
args += GetSourceExpr(context, operation.GetSource(argIndex), dstType);
builder.Append(GetSourceExpr(context, operation.GetSource(argIndex), dstType));
}
}
return info.OpName + '(' + args + ')';
return $"{info.OpName}({builder})";
}
else if ((info.Type & InstType.Op) != 0)
{
@@ -110,7 +119,7 @@ namespace Ryujinx.Graphics.Shader.CodeGen.Msl.Instructions
switch (inst & Instruction.Mask)
{
case Instruction.Barrier:
return "|| BARRIER ||";
return "threadgroup_barrier(mem_flags::mem_threadgroup)";
case Instruction.Call:
return Call(context, operation);
case Instruction.FSIBegin: