10 files changed, 414 insertions, 327 deletions

diff --git a/src/shader_recompiler/backend/glsl/emit_context.cpp b/src/shader_recompiler/backend/glsl/emit_context.cpp
index 67772c46d..3c610a08a 100644
--- a/src/shader_recompiler/backend/glsl/emit_context.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_context.cpp
@@ -19,8 +19,10 @@ EmitContext::EmitContext(IR::Program& program, [[maybe_unused]] Bindings& bindin
                               program.workgroup_size[2]);
     }
     code += header;
+
     DefineConstantBuffers();
     DefineStorageBuffers();
+    DefineHelperFunctions();
     code += "void main(){\n";
 }
 
@@ -28,6 +30,15 @@ void EmitContext::SetupExtensions(std::string& header) {
     if (info.uses_int64) {
         header += "#extension GL_ARB_gpu_shader_int64 : enable\n";
     }
+    if (info.uses_int64_bit_atomics) {
+        header += "#extension GL_NV_shader_atomic_int64 : enable\n";
+    }
+    if (info.uses_atomic_f32_add) {
+        header += "#extension GL_NV_shader_atomic_float : enable\n";
+    }
+    if (info.uses_atomic_f16x2_add || info.uses_atomic_f16x2_min || info.uses_atomic_f16x2_max) {
+        header += "#extension NV_shader_atomic_fp16_vector : enable\n";
+    }
 }
 
 void EmitContext::DefineConstantBuffers() {
@@ -48,18 +59,39 @@ void EmitContext::DefineStorageBuffers() {
     }
     u32 binding{};
     for (const auto& desc : info.storage_buffers_descriptors) {
-        if (True(info.used_storage_buffer_types & IR::Type::U32) ||
-            True(info.used_storage_buffer_types & IR::Type::F32)) {
+        if (info.uses_s32_atomics) {
+            Add("layout(std430,binding={}) buffer ssbo_{}_s32{{int ssbo{}_s32[];}};", binding,
+                binding, desc.cbuf_index, desc.count);
+        }
+        if (True(info.used_storage_buffer_types & IR::Type::U32)) {
             Add("layout(std430,binding={}) buffer ssbo_{}_u32{{uint ssbo{}_u32[];}};", binding,
                 binding, desc.cbuf_index, desc.count);
         }
-        if (True(info.used_storage_buffer_types & IR::Type::U32x2) ||
-            True(info.used_storage_buffer_types & IR::Type::F32x2)) {
-            Add("layout(std430,binding={}) buffer ssbo_{}_u64{{uvec2 ssbo{}_u64[];}};", binding,
+        if (True(info.used_storage_buffer_types & IR::Type::F32)) {
+            Add("layout(std430,binding={}) buffer ssbo_{}_f32{{float ssbo{}_f32[];}};", binding,
+                binding, desc.cbuf_index, desc.count);
+        }
+        if (True(info.used_storage_buffer_types & IR::Type::U32x2)) {
+            Add("layout(std430,binding={}) buffer ssbo_{}_u32x2{{uvec2 ssbo{}_u32x2[];}};", binding,
+                binding, desc.cbuf_index, desc.count);
+        }
+        if (True(info.used_storage_buffer_types & IR::Type::U64) ||
+            True(info.used_storage_buffer_types & IR::Type::F64)) {
+            Add("layout(std430,binding={}) buffer ssbo_{}_u64{{uint64_t ssbo{}_u64[];}};", binding,
                 binding, desc.cbuf_index, desc.count);
         }
         ++binding;
     }
 }
 
+void EmitContext::DefineHelperFunctions() {
+    if (info.uses_global_increment) {
+        code += "uint CasIncrement(uint op_a,uint op_b){return(op_a>=op_b)?0u:(op_a+1u);}\n";
+    }
+    if (info.uses_global_decrement) {
+        code +=
+            "uint CasDecrement(uint op_a,uint op_b){return(op_a==0||op_a>op_b)?op_b:(op_a-1u);}\n";
+    }
+}
+
 } // namespace Shader::Backend::GLSL
diff --git a/src/shader_recompiler/backend/glsl/emit_context.h b/src/shader_recompiler/backend/glsl/emit_context.h
index 9472f71b9..ca5657fe7 100644
--- a/src/shader_recompiler/backend/glsl/emit_context.h
+++ b/src/shader_recompiler/backend/glsl/emit_context.h
@@ -31,13 +31,6 @@ class EmitContext {
 public:
     explicit EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_);
 
-    // template <typename... Args>
-    // void Add(const char* format_str, IR::Inst& inst, Args&&... args) {
-    //     code += fmt::format(format_str, reg_alloc.Define(inst), std::forward<Args>(args)...);
-    //     // TODO: Remove this
-    //     code += '\n';
-    // }
-
     template <Type type, typename... Args>
     void Add(const char* format_str, IR::Inst& inst, Args&&... args) {
         code += fmt::format(format_str, reg_alloc.Define(inst, type), std::forward<Args>(args)...);
@@ -106,6 +99,7 @@ private:
     void SetupExtensions(std::string& header);
     void DefineConstantBuffers();
     void DefineStorageBuffers();
+    void DefineHelperFunctions();
 };
 
 } // namespace Shader::Backend::GLSL
diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_atomic.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_atomic.cpp
index e69de29bb..f3ef37873 100644
--- a/src/shader_recompiler/backend/glsl/emit_glsl_atomic.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_atomic.cpp
@@ -0,0 +1,301 @@
+
+// Copyright 2021 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+#include <string_view>
+
+#include "shader_recompiler/backend/glsl/emit_context.h"
+#include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
+#include "shader_recompiler/frontend/ir/value.h"
+#include "shader_recompiler/profile.h"
+
+namespace Shader::Backend::GLSL {
+namespace {
+static constexpr std::string_view cas_loop{R"(
+{} {};
+for (;;){{
+    {} old_value={};
+    {} = atomicCompSwap({},old_value,{}({},{}));
+    if ({}==old_value){{break;}}
+}})"};
+
+void CasFunction(EmitContext& ctx, IR::Inst& inst, std::string_view ssbo, std::string_view value,
+                 std::string_view type, std::string_view function) {
+    const auto ret{ctx.reg_alloc.Define(inst)};
+    ctx.Add(cas_loop.data(), type, ret, type, ssbo, ret, ssbo, function, ssbo, value, ret);
+}
+} // namespace
+
+void EmitStorageAtomicIAdd32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value) {
+    ctx.AddU32("{}=atomicAdd(ssbo{}_u32[{}],{});", inst, binding.U32(), offset.U32(), value);
+}
+
+void EmitStorageAtomicSMin32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value) {
+    ctx.AddS32("{}=atomicMin(ssbo{}_s32[{}],int({}));", inst, binding.U32(), offset.U32(), value);
+}
+
+void EmitStorageAtomicUMin32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value) {
+    ctx.AddU32("{}=atomicMin(ssbo{}_u32[{}],{});", inst, binding.U32(), offset.U32(), value);
+}
+
+void EmitStorageAtomicSMax32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value) {
+    ctx.AddS32("{}=atomicMax(ssbo{}_s32[{}],int({}));", inst, binding.U32(), offset.U32(), value);
+}
+
+void EmitStorageAtomicUMax32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value) {
+    ctx.AddU32("{}=atomicMax(ssbo{}_u32[{}],{});", inst, binding.U32(), offset.U32(), value);
+}
+
+void EmitStorageAtomicInc32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                            [[maybe_unused]] const IR::Value& offset, std::string_view value) {
+    // const auto ret{ctx.reg_alloc.Define(inst)};
+    // const auto type{"uint"};
+    // ctx.Add(cas_loop.data(), type, ret, type, ssbo, ret, ssbo, "CasIncrement", ssbo, value, ret);
+    const std::string ssbo{fmt::format("ssbo{}_u32[{}]", binding.U32(), offset.U32())};
+    CasFunction(ctx, inst, ssbo, value, "uint", "CasIncrement");
+}
+
+void EmitStorageAtomicDec32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                            const IR::Value& offset, std::string_view value) {
+    const std::string ssbo{fmt::format("ssbo{}_u32[{}]", binding.U32(), offset.U32())};
+    CasFunction(ctx, inst, ssbo, value, "uint", "CasDecrement");
+}
+
+void EmitStorageAtomicAnd32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                            const IR::Value& offset, std::string_view value) {
+    ctx.AddU32("{}=atomicAnd(ssbo{}_u32[{}],{});", inst, binding.U32(), offset.U32(), value);
+}
+
+void EmitStorageAtomicOr32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                           const IR::Value& offset, std::string_view value) {
+    ctx.AddU32("{}=atomicOr(ssbo{}_u32[{}],{});", inst, binding.U32(), offset.U32(), value);
+}
+
+void EmitStorageAtomicXor32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                            const IR::Value& offset, std::string_view value) {
+    ctx.AddU32("{}=atomicXor(ssbo{}_u32[{}],{});", inst, binding.U32(), offset.U32(), value);
+}
+
+void EmitStorageAtomicExchange32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                                 const IR::Value& offset, std::string_view value) {
+    ctx.AddU32("{}=atomicExchange(ssbo{}_u32[{}],{});", inst, binding.U32(), offset.U32(), value);
+}
+
+void EmitStorageAtomicIAdd64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value) {
+    // ctx.AddU64("{}=atomicAdd(ssbo{}_u64[{}],{});", inst, binding.U32(), offset.U32(), value);
+    ctx.AddU64("{}=ssbo{}_u64[{}];", inst, binding.U32(), offset.U32());
+    ctx.Add("ssbo{}_u64[{}]+={};", binding.U32(), offset.U32(), value);
+}
+
+void EmitStorageAtomicSMin64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value) {
+    ctx.AddS64("{}=atomicMin(int64_t(ssbo{}_u64[{}]),int64_t({}));", inst, binding.U32(),
+               offset.U32(), value);
+}
+
+void EmitStorageAtomicUMin64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value) {
+    ctx.AddU64("{}=atomicMin(ssbo{}_u64[{}],{});", inst, binding.U32(), offset.U32(), value);
+}
+
+void EmitStorageAtomicSMax64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value) {
+    ctx.AddS64("{}=atomicMax(int64_t(ssbo{}_u64[{}]),int64_t({}));", inst, binding.U32(),
+               offset.U32(), value);
+}
+
+void EmitStorageAtomicUMax64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value) {
+    ctx.AddU64("{}=atomicMax(ssbo{}_u64[{}],{});", inst, binding.U32(), offset.U32(), value);
+}
+
+void EmitStorageAtomicAnd64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                            const IR::Value& offset, std::string_view value) {
+    ctx.AddU64("{}=atomicAnd(ssbo{}_u64[{}],{});", inst, binding.U32(), offset.U32(), value);
+}
+
+void EmitStorageAtomicOr64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                           const IR::Value& offset, std::string_view value) {
+    ctx.AddU64("{}=atomicOr(ssbo{}_u64[{}],{});", inst, binding.U32(), offset.U32(), value);
+}
+
+void EmitStorageAtomicXor64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                            const IR::Value& offset, std::string_view value) {
+    ctx.AddU64("{}=atomicXor(ssbo{}_u64[{}],{});", inst, binding.U32(), offset.U32(), value);
+}
+
+void EmitStorageAtomicExchange64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                                 const IR::Value& offset, std::string_view value) {
+    ctx.AddU64("{}=atomicExchange(ssbo{}_u64[{}],{});", inst, binding.U32(), offset.U32(), value);
+}
+
+void EmitStorageAtomicAddF32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value) {
+    ctx.AddF32("{}=atomicAdd(ssbo{}_u32[{}],{});", inst, binding.U32(), offset.U32(), value);
+}
+
+void EmitStorageAtomicAddF16x2([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst,
+                               [[maybe_unused]] const IR::Value& binding,
+                               [[maybe_unused]] const IR::Value& offset,
+                               [[maybe_unused]] std::string_view value) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitStorageAtomicAddF32x2([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst,
+                               [[maybe_unused]] const IR::Value& binding,
+                               [[maybe_unused]] const IR::Value& offset,
+                               [[maybe_unused]] std::string_view value) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitStorageAtomicMinF16x2([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst,
+                               [[maybe_unused]] const IR::Value& binding,
+                               [[maybe_unused]] const IR::Value& offset,
+                               [[maybe_unused]] std::string_view value) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitStorageAtomicMinF32x2([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst,
+                               [[maybe_unused]] const IR::Value& binding,
+                               [[maybe_unused]] const IR::Value& offset,
+                               [[maybe_unused]] std::string_view value) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitStorageAtomicMaxF16x2([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst,
+                               [[maybe_unused]] const IR::Value& binding,
+                               [[maybe_unused]] const IR::Value& offset,
+                               [[maybe_unused]] std::string_view value) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitStorageAtomicMaxF32x2([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst,
+                               [[maybe_unused]] const IR::Value& binding,
+                               [[maybe_unused]] const IR::Value& offset,
+                               [[maybe_unused]] std::string_view value) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicIAdd32(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicSMin32(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicUMin32(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicSMax32(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicUMax32(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicInc32(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicDec32(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicAnd32(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicOr32(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicXor32(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicExchange32(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicIAdd64(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicSMin64(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicUMin64(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicSMax64(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicUMax64(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicInc64(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicDec64(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicAnd64(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicOr64(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicXor64(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicExchange64(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicAddF32(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicAddF16x2(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicAddF32x2(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicMinF16x2(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicMinF32x2(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicMaxF16x2(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+
+void EmitGlobalAtomicMaxF32x2(EmitContext&) {
+    throw NotImplementedException("GLSL Instrucion");
+}
+} // namespace Shader::Backend::GLSL
diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_instructions.h b/src/shader_recompiler/backend/glsl/emit_glsl_instructions.h
index 4e0487543..c2836898f 100644
--- a/src/shader_recompiler/backend/glsl/emit_glsl_instructions.h
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_instructions.h
@@ -15,10 +15,7 @@ class Inst;
 class Value;
 } // namespace Shader::IR
 
-#pragma optimize("", off)
-
 namespace Shader::Backend::GLSL {
-
 class EmitContext;
 
 inline void EmitSetLoopSafetyVariable(EmitContext&) {}
@@ -114,7 +111,8 @@ void EmitLoadStorageU8(EmitContext& ctx, const IR::Value& binding, const IR::Val
 void EmitLoadStorageS8(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
 void EmitLoadStorageU16(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
 void EmitLoadStorageS16(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
-void EmitLoadStorage32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
+void EmitLoadStorage32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                       const IR::Value& offset);
 void EmitLoadStorage64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
 void EmitLoadStorage128(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
 void EmitWriteStorageU8(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
@@ -431,60 +429,60 @@ void EmitSharedAtomicExchange32(EmitContext& ctx, std::string_view pointer_offse
                                 std::string_view value);
 void EmitSharedAtomicExchange64(EmitContext& ctx, std::string_view pointer_offset,
                                 std::string_view value);
-void EmitStorageAtomicIAdd32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value);
-void EmitStorageAtomicSMin32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value);
-void EmitStorageAtomicUMin32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value);
-void EmitStorageAtomicSMax32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value);
-void EmitStorageAtomicUMax32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value);
-void EmitStorageAtomicInc32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                            std::string_view value);
-void EmitStorageAtomicDec32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                            std::string_view value);
-void EmitStorageAtomicAnd32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                            std::string_view value);
-void EmitStorageAtomicOr32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                           std::string_view value);
-void EmitStorageAtomicXor32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                            std::string_view value);
-void EmitStorageAtomicExchange32(EmitContext& ctx, const IR::Value& binding,
+void EmitStorageAtomicIAdd32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicSMin32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicUMin32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicSMax32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicUMax32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicInc32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                            const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicDec32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                            const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicAnd32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                            const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicOr32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                           const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicXor32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                            const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicExchange32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
                                  const IR::Value& offset, std::string_view value);
-void EmitStorageAtomicIAdd64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value);
-void EmitStorageAtomicSMin64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value);
-void EmitStorageAtomicUMin64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value);
-void EmitStorageAtomicSMax64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value);
-void EmitStorageAtomicUMax64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value);
-void EmitStorageAtomicAnd64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                            std::string_view value);
-void EmitStorageAtomicOr64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                           std::string_view value);
-void EmitStorageAtomicXor64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                            std::string_view value);
-void EmitStorageAtomicExchange64(EmitContext& ctx, const IR::Value& binding,
+void EmitStorageAtomicIAdd64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicSMin64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicUMin64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicSMax64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicUMax64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicAnd64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                            const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicOr64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                           const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicXor64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                            const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicExchange64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
                                  const IR::Value& offset, std::string_view value);
-void EmitStorageAtomicAddF32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value);
-void EmitStorageAtomicAddF16x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                               std::string_view value);
-void EmitStorageAtomicAddF32x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                               std::string_view value);
-void EmitStorageAtomicMinF16x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                               std::string_view value);
-void EmitStorageAtomicMinF32x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                               std::string_view value);
-void EmitStorageAtomicMaxF16x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                               std::string_view value);
-void EmitStorageAtomicMaxF32x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                               std::string_view value);
+void EmitStorageAtomicAddF32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                             const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicAddF16x2(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                               const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicAddF32x2(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                               const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicMinF16x2(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                               const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicMinF32x2(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                               const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicMaxF16x2(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                               const IR::Value& offset, std::string_view value);
+void EmitStorageAtomicMaxF32x2(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                               const IR::Value& offset, std::string_view value);
 void EmitGlobalAtomicIAdd32(EmitContext& ctx);
 void EmitGlobalAtomicSMin32(EmitContext& ctx);
 void EmitGlobalAtomicUMin32(EmitContext& ctx);
diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_memory.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_memory.cpp
index b042ae853..d1e6f074d 100644
--- a/src/shader_recompiler/backend/glsl/emit_glsl_memory.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_memory.cpp
@@ -32,9 +32,9 @@ void EmitLoadStorageS16([[maybe_unused]] EmitContext& ctx,
     throw NotImplementedException("GLSL Instrucion");
 }
 
-void EmitLoadStorage32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] const IR::Value& binding,
-                       [[maybe_unused]] const IR::Value& offset) {
-    throw NotImplementedException("GLSL Instrucion");
+void EmitLoadStorage32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
+                       const IR::Value& offset) {
+    ctx.AddU32("{}=ssbo{}_u32[{}];", inst, binding.U32(), offset.U32());
 }
 
 void EmitLoadStorage64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] const IR::Value& binding,
@@ -83,7 +83,7 @@ void EmitWriteStorage32(EmitContext& ctx, const IR::Value& binding, const IR::Va
 
 void EmitWriteStorage64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
                         std::string_view value) {
-    ctx.Add("ssbo{}_u64[{}]={};", binding.U32(), offset.U32(), value);
+    ctx.Add("ssbo{}_u32x2[{}]={};", binding.U32(), offset.U32(), value);
 }
 
 void EmitWriteStorage128([[maybe_unused]] EmitContext& ctx,
diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_not_implemented.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_not_implemented.cpp
index cb6562ebf..65eccaece 100644
--- a/src/shader_recompiler/backend/glsl/emit_glsl_not_implemented.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_not_implemented.cpp
@@ -20,7 +20,7 @@ static void NotImplemented() {
 }
 
 void EmitPhi(EmitContext& ctx, IR::Inst& inst) {
-    NotImplemented();
+    // NotImplemented();
 }
 
 void EmitVoid(EmitContext& ctx) {
@@ -439,257 +439,6 @@ void EmitSharedAtomicExchange64(EmitContext& ctx, std::string_view pointer_offse
     NotImplemented();
 }
 
-void EmitStorageAtomicIAdd32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicSMin32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicUMin32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicSMax32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicUMax32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicInc32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                            std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicDec32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                            std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicAnd32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                            std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicOr32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                           std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicXor32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                            std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicExchange32(EmitContext& ctx, const IR::Value& binding,
-                                 const IR::Value& offset, std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicIAdd64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicSMin64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicUMin64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicSMax64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicUMax64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicAnd64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                            std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicOr64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                           std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicXor64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                            std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicExchange64(EmitContext& ctx, const IR::Value& binding,
-                                 const IR::Value& offset, std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicAddF32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                             std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicAddF16x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                               std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicAddF32x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                               std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicMinF16x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                               std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicMinF32x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                               std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicMaxF16x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                               std::string_view value) {
-    NotImplemented();
-}
-
-void EmitStorageAtomicMaxF32x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
-                               std::string_view value) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicIAdd32(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicSMin32(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicUMin32(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicSMax32(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicUMax32(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicInc32(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicDec32(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicAnd32(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicOr32(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicXor32(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicExchange32(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicIAdd64(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicSMin64(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicUMin64(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicSMax64(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicUMax64(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicInc64(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicDec64(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicAnd64(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicOr64(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicXor64(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicExchange64(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicAddF32(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicAddF16x2(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicAddF32x2(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicMinF16x2(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicMinF32x2(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicMaxF16x2(EmitContext& ctx) {
-    NotImplemented();
-}
-
-void EmitGlobalAtomicMaxF32x2(EmitContext& ctx) {
-    NotImplemented();
-}
-
 void EmitBindlessImageSampleImplicitLod(EmitContext&) {
     NotImplemented();
 }
diff --git a/src/shader_recompiler/backend/glsl/reg_alloc.cpp b/src/shader_recompiler/backend/glsl/reg_alloc.cpp
index 9f529c358..8db1391fd 100644
--- a/src/shader_recompiler/backend/glsl/reg_alloc.cpp
+++ b/src/shader_recompiler/backend/glsl/reg_alloc.cpp
@@ -61,6 +61,12 @@ std::string MakeImm(const IR::Value& value) {
 }
 } // Anonymous namespace
 
+std::string RegAlloc::Define(IR::Inst& inst) {
+    const Id id{Alloc()};
+    inst.SetDefinition<Id>(id);
+    return Representation(id);
+}
+
 std::string RegAlloc::Define(IR::Inst& inst, Type type) {
     const Id id{Alloc()};
     const auto type_str{GetType(type, id.index)};
diff --git a/src/shader_recompiler/backend/glsl/reg_alloc.h b/src/shader_recompiler/backend/glsl/reg_alloc.h
index 63c940d3a..7891c30e0 100644
--- a/src/shader_recompiler/backend/glsl/reg_alloc.h
+++ b/src/shader_recompiler/backend/glsl/reg_alloc.h
@@ -48,7 +48,8 @@ static_assert(sizeof(Id) == sizeof(u32));
 
 class RegAlloc {
 public:
-    std::string Define(IR::Inst& inst, Type type = Type::Void);
+    std::string Define(IR::Inst& inst);
+    std::string Define(IR::Inst& inst, Type type);
 
     std::string Consume(const IR::Value& value);
 
diff --git a/src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp b/src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp
index fb2031fc8..884ade004 100644
--- a/src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp
+++ b/src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp
@@ -687,9 +687,7 @@ void VisitUsages(Info& info, IR::Inst& inst) {
     case IR::Opcode::LoadStorage32:
     case IR::Opcode::WriteStorage32:
     case IR::Opcode::StorageAtomicIAdd32:
-    case IR::Opcode::StorageAtomicSMin32:
     case IR::Opcode::StorageAtomicUMin32:
-    case IR::Opcode::StorageAtomicSMax32:
     case IR::Opcode::StorageAtomicUMax32:
     case IR::Opcode::StorageAtomicAnd32:
     case IR::Opcode::StorageAtomicOr32:
@@ -714,6 +712,13 @@ void VisitUsages(Info& info, IR::Inst& inst) {
     case IR::Opcode::SharedAtomicExchange64:
         info.uses_int64_bit_atomics = true;
         break;
+    case IR::Opcode::GlobalAtomicSMin32:
+    case IR::Opcode::GlobalAtomicSMax32:
+    case IR::Opcode::StorageAtomicSMin32:
+    case IR::Opcode::StorageAtomicSMax32:
+        info.used_storage_buffer_types |= IR::Type::U32;
+        info.uses_s32_atomics = true;
+        break;
     case IR::Opcode::GlobalAtomicInc32:
     case IR::Opcode::StorageAtomicInc32:
         info.used_storage_buffer_types |= IR::Type::U32;
diff --git a/src/shader_recompiler/shader_info.h b/src/shader_recompiler/shader_info.h
index 32f8a50ea..be05eafcf 100644
--- a/src/shader_recompiler/shader_info.h
+++ b/src/shader_recompiler/shader_info.h
@@ -190,6 +190,7 @@ struct Info {
     bool uses_atomic_f32x2_min{};
     bool uses_atomic_f32x2_max{};
     bool uses_int64_bit_atomics{};
+    bool uses_s32_atomics{};
     bool uses_global_memory{};
     bool uses_atomic_image_u32{};