5 files changed, 78 insertions, 0 deletions

diff --git a/src/shader_recompiler/frontend/ir/ir_emitter.cpp b/src/shader_recompiler/frontend/ir/ir_emitter.cpp
index 6d41442ee..d6a1d8ec2 100644
--- a/src/shader_recompiler/frontend/ir/ir_emitter.cpp
+++ b/src/shader_recompiler/frontend/ir/ir_emitter.cpp
@@ -355,6 +355,52 @@ void IREmitter::WriteGlobal128(const U64& address, const IR::Value& vector) {
     Inst(Opcode::WriteGlobal128, address, vector);
 }
 
+U32 IREmitter::LoadLocal(const IR::U32& word_offset) {
+    return Inst<U32>(Opcode::LoadLocal, word_offset);
+}
+
+void IREmitter::WriteLocal(const IR::U32& word_offset, const IR::U32& value) {
+    Inst(Opcode::WriteLocal, word_offset, value);
+}
+
+Value IREmitter::LoadShared(int bit_size, bool is_signed, const IR::U32& offset) {
+    switch (bit_size) {
+    case 8:
+        return Inst(is_signed ? Opcode::LoadSharedS8 : Opcode::LoadSharedU8, offset);
+    case 16:
+        return Inst(is_signed ? Opcode::LoadSharedS16 : Opcode::LoadSharedU16, offset);
+    case 32:
+        return Inst(Opcode::LoadSharedU32, offset);
+    case 64:
+        return Inst(Opcode::LoadSharedU64, offset);
+    case 128:
+        return Inst(Opcode::LoadSharedU128, offset);
+    }
+    throw InvalidArgument("Invalid bit size {}", bit_size);
+}
+
+void IREmitter::WriteShared(int bit_size, const IR::U32& offset, const IR::Value& value) {
+    switch (bit_size) {
+    case 8:
+        Inst(Opcode::WriteSharedU8, offset, value);
+        break;
+    case 16:
+        Inst(Opcode::WriteSharedU16, offset, value);
+        break;
+    case 32:
+        Inst(Opcode::WriteSharedU32, offset, value);
+        break;
+    case 64:
+        Inst(Opcode::WriteSharedU64, offset, value);
+        break;
+    case 128:
+        Inst(Opcode::WriteSharedU128, offset, value);
+        break;
+    default:
+        throw InvalidArgument("Invalid bit size {}", bit_size);
+    }
+}
+
 U1 IREmitter::GetZeroFromOp(const Value& op) {
     return Inst<U1>(Opcode::GetZeroFromOp, op);
 }
diff --git a/src/shader_recompiler/frontend/ir/ir_emitter.h b/src/shader_recompiler/frontend/ir/ir_emitter.h
index 8d50aa607..842c2bdaf 100644
--- a/src/shader_recompiler/frontend/ir/ir_emitter.h
+++ b/src/shader_recompiler/frontend/ir/ir_emitter.h
@@ -99,6 +99,12 @@ public:
     void WriteGlobal64(const U64& address, const IR::Value& vector);
     void WriteGlobal128(const U64& address, const IR::Value& vector);
 
+    [[nodiscard]] U32 LoadLocal(const U32& word_offset);
+    void WriteLocal(const U32& word_offset, const U32& value);
+
+    [[nodiscard]] Value LoadShared(int bit_size, bool is_signed, const U32& offset);
+    void WriteShared(int bit_size, const U32& offset, const Value& value);
+
     [[nodiscard]] U1 GetZeroFromOp(const Value& op);
     [[nodiscard]] U1 GetSignFromOp(const Value& op);
     [[nodiscard]] U1 GetCarryFromOp(const Value& op);
diff --git a/src/shader_recompiler/frontend/ir/microinstruction.cpp b/src/shader_recompiler/frontend/ir/microinstruction.cpp
index be8eb4d4c..52a5e5034 100644
--- a/src/shader_recompiler/frontend/ir/microinstruction.cpp
+++ b/src/shader_recompiler/frontend/ir/microinstruction.cpp
@@ -76,6 +76,12 @@ bool Inst::MayHaveSideEffects() const noexcept {
     case Opcode::WriteStorage32:
     case Opcode::WriteStorage64:
     case Opcode::WriteStorage128:
+    case Opcode::WriteLocal:
+    case Opcode::WriteSharedU8:
+    case Opcode::WriteSharedU16:
+    case Opcode::WriteSharedU32:
+    case Opcode::WriteSharedU64:
+    case Opcode::WriteSharedU128:
         return true;
     default:
         return false;
diff --git a/src/shader_recompiler/frontend/ir/opcodes.inc b/src/shader_recompiler/frontend/ir/opcodes.inc
index 5d7462d76..c75658328 100644
--- a/src/shader_recompiler/frontend/ir/opcodes.inc
+++ b/src/shader_recompiler/frontend/ir/opcodes.inc
@@ -89,6 +89,24 @@ OPCODE(WriteStorage32,                                      Void,           U32,
 OPCODE(WriteStorage64,                                      Void,           U32,            U32,            U32x2,                                          )
 OPCODE(WriteStorage128,                                     Void,           U32,            U32,            U32x4,                                          )
 
+// Local memory operations
+OPCODE(LoadLocal,                                           U32,            U32,                                                                            )
+OPCODE(WriteLocal,                                          Void,           U32,            U32,                                                            )
+
+// Shared memory operations
+OPCODE(LoadSharedU8,                                        U32,            U32,                                                                            )
+OPCODE(LoadSharedS8,                                        U32,            U32,                                                                            )
+OPCODE(LoadSharedU16,                                       U32,            U32,                                                                            )
+OPCODE(LoadSharedS16,                                       U32,            U32,                                                                            )
+OPCODE(LoadSharedU32,                                       U32,            U32,                                                                            )
+OPCODE(LoadSharedU64,                                       U32x2,          U32,                                                                            )
+OPCODE(LoadSharedU128,                                      U32x4,          U32,                                                                            )
+OPCODE(WriteSharedU8,                                       Void,           U32,            U32,                                                            )
+OPCODE(WriteSharedU16,                                      Void,           U32,            U32,                                                            )
+OPCODE(WriteSharedU32,                                      Void,           U32,            U32,                                                            )
+OPCODE(WriteSharedU64,                                      Void,           U32,            U32x2,                                                          )
+OPCODE(WriteSharedU128,                                     Void,           U32,            U32x4,                                                          )
+
 // Vector utility
 OPCODE(CompositeConstructU32x2,                             U32x2,          U32,            U32,                                                            )
 OPCODE(CompositeConstructU32x3,                             U32x3,          U32,            U32,            U32,                                            )
diff --git a/src/shader_recompiler/frontend/ir/program.h b/src/shader_recompiler/frontend/ir/program.h
index 0162e919c..3a37b3ab9 100644
--- a/src/shader_recompiler/frontend/ir/program.h
+++ b/src/shader_recompiler/frontend/ir/program.h
@@ -21,6 +21,8 @@ struct Program {
     Info info;
     Stage stage{};
     std::array<u32, 3> workgroup_size{};
+    u32 local_memory_size{};
+    u32 shared_memory_size{};
 };
 
 [[nodiscard]] std::string DumpProgram(const Program& program);