From 3d9ecbe99844c44074c26f2db4db376059f50534 Mon Sep 17 00:00:00 2001
From: ameerj <52414509+ameerj@users.noreply.github.com>
Date: Mon, 24 May 2021 18:35:37 -0400
Subject: [PATCH] glsl: Wip storage atomic ops
---
.../backend/glsl/emit_context.cpp | 42 ++-
.../backend/glsl/emit_context.h | 8 +-
.../backend/glsl/emit_glsl_atomic.cpp | 301 ++++++++++++++++++
.../backend/glsl/emit_glsl_instructions.h | 110 ++++---
.../backend/glsl/emit_glsl_memory.cpp | 8 +-
.../glsl/emit_glsl_not_implemented.cpp | 253 +--------------
.../backend/glsl/reg_alloc.cpp | 6 +
.../backend/glsl/reg_alloc.h | 3 +-
.../ir_opt/collect_shader_info_pass.cpp | 9 +-
src/shader_recompiler/shader_info.h | 1 +
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{};