diff options
Diffstat (limited to '')
8 files changed, 406 insertions, 325 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); |