From 2930dccecc933d6748772e9f51a5724fe1e6771b Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Mon, 8 Feb 2021 02:54:35 -0300 Subject: spirv: Initial SPIR-V support --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 134 +++++++++++++++++++++ 1 file changed, 134 insertions(+) create mode 100644 src/shader_recompiler/backend/spirv/emit_spirv.cpp (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp new file mode 100644 index 000000000..7c4269fad --- /dev/null +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -0,0 +1,134 @@ +// Copyright 2021 yuzu Emulator Project +// Licensed under GPLv2 or any later version +// Refer to the license.txt file included. + +#include +#include + +#include "shader_recompiler/backend/spirv/emit_spirv.h" +#include "shader_recompiler/frontend/ir/basic_block.h" +#include "shader_recompiler/frontend/ir/function.h" +#include "shader_recompiler/frontend/ir/microinstruction.h" +#include "shader_recompiler/frontend/ir/program.h" + +namespace Shader::Backend::SPIRV { + +EmitContext::EmitContext(IR::Program& program) { + AddCapability(spv::Capability::Shader); + AddCapability(spv::Capability::Float16); + AddCapability(spv::Capability::Float64); + void_id = TypeVoid(); + + u1 = Name(TypeBool(), "u1"); + f32.Define(*this, TypeFloat(32), "f32"); + u32.Define(*this, TypeInt(32, false), "u32"); + f16.Define(*this, TypeFloat(16), "f16"); + f64.Define(*this, TypeFloat(64), "f64"); + + for (const IR::Function& function : program.functions) { + for (IR::Block* const block : function.blocks) { + block_label_map.emplace_back(block, OpLabel()); + } + } + std::ranges::sort(block_label_map, {}, &std::pair::first); +} + +EmitContext::~EmitContext() = default; + +EmitSPIRV::EmitSPIRV(IR::Program& program) { + EmitContext ctx{program}; + const Id void_function{ctx.TypeFunction(ctx.void_id)}; + // FIXME: Forward declare functions (needs sirit support) + Id func{}; + for (IR::Function& function : program.functions) { + func = ctx.OpFunction(ctx.void_id, spv::FunctionControlMask::MaskNone, void_function); + for (IR::Block* const block : function.blocks) { + ctx.AddLabel(ctx.BlockLabel(block)); + for (IR::Inst& inst : block->Instructions()) { + EmitInst(ctx, &inst); + } + } + ctx.OpFunctionEnd(); + } + ctx.AddEntryPoint(spv::ExecutionModel::GLCompute, func, "main"); + + std::vector result{ctx.Assemble()}; + std::FILE* file{std::fopen("shader.spv", "wb")}; + std::fwrite(result.data(), sizeof(u32), result.size(), file); + std::fclose(file); + std::system("spirv-dis shader.spv"); + std::system("spirv-val shader.spv"); +} + +template +static void Invoke(EmitSPIRV& emit, EmitContext& ctx, IR::Inst* inst) { + using M = decltype(method); + using std::is_invocable_r_v; + if constexpr (is_invocable_r_v) { + ctx.Define(inst, (emit.*method)(ctx)); + } else if constexpr (is_invocable_r_v) { + ctx.Define(inst, (emit.*method)(ctx, ctx.Def(inst->Arg(0)))); + } else if constexpr (is_invocable_r_v) { + ctx.Define(inst, (emit.*method)(ctx, ctx.Def(inst->Arg(0)), ctx.Def(inst->Arg(1)))); + } else if constexpr (is_invocable_r_v) { + ctx.Define(inst, (emit.*method)(ctx, ctx.Def(inst->Arg(0)), ctx.Def(inst->Arg(1)), + ctx.Def(inst->Arg(2)))); + } else if constexpr (is_invocable_r_v) { + ctx.Define(inst, (emit.*method)(ctx, inst, ctx.Def(inst->Arg(0)), ctx.Def(inst->Arg(1)))); + } else if constexpr (is_invocable_r_v) { + ctx.Define(inst, (emit.*method)(ctx, inst, ctx.Def(inst->Arg(0)), ctx.Def(inst->Arg(1)), + ctx.Def(inst->Arg(2)))); + } else if constexpr (is_invocable_r_v) { + ctx.Define(inst, (emit.*method)(ctx, ctx.Def(inst->Arg(0)), inst->Arg(1).U32())); + } else if constexpr (is_invocable_r_v) { + ctx.Define(inst, (emit.*method)(ctx, inst->Arg(0))); + } else if constexpr (is_invocable_r_v) { + ctx.Define(inst, (emit.*method)(ctx, inst->Arg(0), inst->Arg(1))); + } else if constexpr (is_invocable_r_v) { + (emit.*method)(ctx, inst); + } else if constexpr (is_invocable_r_v) { + (emit.*method)(ctx); + } else { + static_assert(false, "Bad format"); + } +} + +void EmitSPIRV::EmitInst(EmitContext& ctx, IR::Inst* inst) { + switch (inst->Opcode()) { +#define OPCODE(name, result_type, ...) \ + case IR::Opcode::name: \ + return Invoke<&EmitSPIRV::Emit##name>(*this, ctx, inst); +#include "shader_recompiler/frontend/ir/opcodes.inc" +#undef OPCODE + } + throw LogicError("Invalid opcode {}", inst->Opcode()); +} + +void EmitSPIRV::EmitPhi(EmitContext&) { + throw NotImplementedException("SPIR-V Instruction"); +} + +void EmitSPIRV::EmitVoid(EmitContext&) {} + +void EmitSPIRV::EmitIdentity(EmitContext&) { + throw NotImplementedException("SPIR-V Instruction"); +} + +void EmitSPIRV::EmitGetZeroFromOp(EmitContext&) { + throw LogicError("Unreachable instruction"); +} + +void EmitSPIRV::EmitGetSignFromOp(EmitContext&) { + throw LogicError("Unreachable instruction"); +} + +void EmitSPIRV::EmitGetCarryFromOp(EmitContext&) { + throw LogicError("Unreachable instruction"); +} + +void EmitSPIRV::EmitGetOverflowFromOp(EmitContext&) { + throw LogicError("Unreachable instruction"); +} + +} // namespace Shader::Backend::SPIRV -- cgit v1.2.3 From 9170200a11715d131645d1ffb92e86e6ef0d7e88 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Thu, 11 Feb 2021 16:39:06 -0300 Subject: shader: Initial implementation of an AST --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 45 +++++++++++++++++++++- 1 file changed, 43 insertions(+), 2 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 7c4269fad..5022b5159 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -105,8 +105,26 @@ void EmitSPIRV::EmitInst(EmitContext& ctx, IR::Inst* inst) { throw LogicError("Invalid opcode {}", inst->Opcode()); } -void EmitSPIRV::EmitPhi(EmitContext&) { - throw NotImplementedException("SPIR-V Instruction"); +static Id TypeId(const EmitContext& ctx, IR::Type type) { + switch (type) { + case IR::Type::U1: + return ctx.u1; + default: + throw NotImplementedException("Phi node type {}", type); + } +} + +Id EmitSPIRV::EmitPhi(EmitContext& ctx, IR::Inst* inst) { + const size_t num_args{inst->NumArgs()}; + boost::container::small_vector operands; + operands.reserve(num_args * 2); + for (size_t index = 0; index < num_args; ++index) { + IR::Block* const phi_block{inst->PhiBlock(index)}; + operands.push_back(ctx.Def(inst->Arg(index))); + operands.push_back(ctx.BlockLabel(phi_block)); + } + const Id result_type{TypeId(ctx, inst->Arg(0).Type())}; + return ctx.OpPhi(result_type, std::span(operands.data(), operands.size())); } void EmitSPIRV::EmitVoid(EmitContext&) {} @@ -115,6 +133,29 @@ void EmitSPIRV::EmitIdentity(EmitContext&) { throw NotImplementedException("SPIR-V Instruction"); } +// FIXME: Move to its own file +void EmitSPIRV::EmitBranch(EmitContext& ctx, IR::Inst* inst) { + ctx.OpBranch(ctx.BlockLabel(inst->Arg(0).Label())); +} + +void EmitSPIRV::EmitBranchConditional(EmitContext& ctx, IR::Inst* inst) { + ctx.OpBranchConditional(ctx.Def(inst->Arg(0)), ctx.BlockLabel(inst->Arg(1).Label()), + ctx.BlockLabel(inst->Arg(2).Label())); +} + +void EmitSPIRV::EmitLoopMerge(EmitContext& ctx, IR::Inst* inst) { + ctx.OpLoopMerge(ctx.BlockLabel(inst->Arg(0).Label()), ctx.BlockLabel(inst->Arg(1).Label()), + spv::LoopControlMask::MaskNone); +} + +void EmitSPIRV::EmitSelectionMerge(EmitContext& ctx, IR::Inst* inst) { + ctx.OpSelectionMerge(ctx.BlockLabel(inst->Arg(0).Label()), spv::SelectionControlMask::MaskNone); +} + +void EmitSPIRV::EmitReturn(EmitContext& ctx) { + ctx.OpReturn(); +} + void EmitSPIRV::EmitGetZeroFromOp(EmitContext&) { throw LogicError("Unreachable instruction"); } -- cgit v1.2.3 From 8af9297f0972d0aaa8306369c5d04926b886a89e Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Sun, 14 Feb 2021 01:24:32 -0300 Subject: shader: Misc fixes --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 6 ++++++ 1 file changed, 6 insertions(+) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 5022b5159..e29e448c7 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -25,6 +25,9 @@ EmitContext::EmitContext(IR::Program& program) { f16.Define(*this, TypeFloat(16), "f16"); f64.Define(*this, TypeFloat(64), "f64"); + true_value = ConstantTrue(u1); + false_value = ConstantFalse(u1); + for (const IR::Function& function : program.functions) { for (IR::Block* const block : function.blocks) { block_label_map.emplace_back(block, OpLabel()); @@ -58,6 +61,7 @@ EmitSPIRV::EmitSPIRV(IR::Program& program) { std::fclose(file); std::system("spirv-dis shader.spv"); std::system("spirv-val shader.spv"); + std::system("spirv-cross shader.spv"); } template @@ -109,6 +113,8 @@ static Id TypeId(const EmitContext& ctx, IR::Type type) { switch (type) { case IR::Type::U1: return ctx.u1; + case IR::Type::U32: + return ctx.u32[1]; default: throw NotImplementedException("Phi node type {}", type); } -- cgit v1.2.3 From 1b0cf2309c760c1cb97a230a1572f8e87f84444a Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Sun, 14 Feb 2021 22:46:40 -0300 Subject: shader: Add support for forward declarations --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 61 +++++++++++++++++----- 1 file changed, 48 insertions(+), 13 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index e29e448c7..0895414b4 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -64,31 +64,49 @@ EmitSPIRV::EmitSPIRV(IR::Program& program) { std::system("spirv-cross shader.spv"); } +template +static void SetDefinition(EmitSPIRV& emit, EmitContext& ctx, IR::Inst* inst, Args... args) { + const Id forward_id{inst->Definition()}; + const bool has_forward_id{Sirit::ValidId(forward_id)}; + Id current_id{}; + if (has_forward_id) { + current_id = ctx.ExchangeCurrentId(forward_id); + } + const Id new_id{(emit.*method)(ctx, std::forward(args)...)}; + if (has_forward_id) { + ctx.ExchangeCurrentId(current_id); + } else { + inst->SetDefinition(new_id); + } +} + template static void Invoke(EmitSPIRV& emit, EmitContext& ctx, IR::Inst* inst) { using M = decltype(method); using std::is_invocable_r_v; if constexpr (is_invocable_r_v) { - ctx.Define(inst, (emit.*method)(ctx)); + SetDefinition(emit, ctx, inst); } else if constexpr (is_invocable_r_v) { - ctx.Define(inst, (emit.*method)(ctx, ctx.Def(inst->Arg(0)))); + SetDefinition(emit, ctx, inst, ctx.Def(inst->Arg(0))); } else if constexpr (is_invocable_r_v) { - ctx.Define(inst, (emit.*method)(ctx, ctx.Def(inst->Arg(0)), ctx.Def(inst->Arg(1)))); + SetDefinition(emit, ctx, inst, ctx.Def(inst->Arg(0)), ctx.Def(inst->Arg(1))); } else if constexpr (is_invocable_r_v) { - ctx.Define(inst, (emit.*method)(ctx, ctx.Def(inst->Arg(0)), ctx.Def(inst->Arg(1)), - ctx.Def(inst->Arg(2)))); + SetDefinition(emit, ctx, inst, ctx.Def(inst->Arg(0)), ctx.Def(inst->Arg(1)), + ctx.Def(inst->Arg(2))); + } else if constexpr (is_invocable_r_v) { + SetDefinition(emit, ctx, inst, inst); } else if constexpr (is_invocable_r_v) { - ctx.Define(inst, (emit.*method)(ctx, inst, ctx.Def(inst->Arg(0)), ctx.Def(inst->Arg(1)))); + SetDefinition(emit, ctx, inst, inst, ctx.Def(inst->Arg(0)), ctx.Def(inst->Arg(1))); } else if constexpr (is_invocable_r_v) { - ctx.Define(inst, (emit.*method)(ctx, inst, ctx.Def(inst->Arg(0)), ctx.Def(inst->Arg(1)), - ctx.Def(inst->Arg(2)))); + SetDefinition(emit, ctx, inst, inst, ctx.Def(inst->Arg(0)), ctx.Def(inst->Arg(1)), + ctx.Def(inst->Arg(2))); } else if constexpr (is_invocable_r_v) { - ctx.Define(inst, (emit.*method)(ctx, ctx.Def(inst->Arg(0)), inst->Arg(1).U32())); + SetDefinition(emit, ctx, inst, ctx.Def(inst->Arg(0)), inst->Arg(1).U32()); } else if constexpr (is_invocable_r_v) { - ctx.Define(inst, (emit.*method)(ctx, inst->Arg(0))); + SetDefinition(emit, ctx, inst, inst->Arg(0)); } else if constexpr (is_invocable_r_v) { - ctx.Define(inst, (emit.*method)(ctx, inst->Arg(0), inst->Arg(1))); + SetDefinition(emit, ctx, inst, inst->Arg(0), inst->Arg(1)); } else if constexpr (is_invocable_r_v) { (emit.*method)(ctx, inst); } else if constexpr (is_invocable_r_v) { @@ -122,11 +140,28 @@ static Id TypeId(const EmitContext& ctx, IR::Type type) { Id EmitSPIRV::EmitPhi(EmitContext& ctx, IR::Inst* inst) { const size_t num_args{inst->NumArgs()}; - boost::container::small_vector operands; + boost::container::small_vector operands; operands.reserve(num_args * 2); for (size_t index = 0; index < num_args; ++index) { + // Phi nodes can have forward declarations, if an argument is not defined provide a forward + // declaration of it. Invoke will take care of giving it the right definition when it's + // actually defined. + const IR::Value arg{inst->Arg(index)}; + Id def{}; + if (arg.IsImmediate()) { + // Let the context handle immediate definitions, as it already knows how + def = ctx.Def(arg); + } else { + IR::Inst* const arg_inst{arg.Inst()}; + def = arg_inst->Definition(); + if (!Sirit::ValidId(def)) { + // If it hasn't been defined, get a forward declaration + def = ctx.ForwardDeclarationId(); + arg_inst->SetDefinition(def); + } + } IR::Block* const phi_block{inst->PhiBlock(index)}; - operands.push_back(ctx.Def(inst->Arg(index))); + operands.push_back(def); operands.push_back(ctx.BlockLabel(phi_block)); } const Id result_type{TypeId(ctx, inst->Arg(0).Type())}; -- cgit v1.2.3 From b5d7279d878211654b4abb165d94af763a365f47 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Tue, 16 Feb 2021 04:10:22 -0300 Subject: spirv: Initial bindings support --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 189 ++++++++++----------- 1 file changed, 88 insertions(+), 101 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 0895414b4..c79c09774 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -12,31 +12,83 @@ #include "shader_recompiler/frontend/ir/program.h" namespace Shader::Backend::SPIRV { +namespace { +template +struct FuncTraits : FuncTraits {}; -EmitContext::EmitContext(IR::Program& program) { - AddCapability(spv::Capability::Shader); - AddCapability(spv::Capability::Float16); - AddCapability(spv::Capability::Float64); - void_id = TypeVoid(); +template +struct FuncTraits { + using ReturnType = ReturnType_; - u1 = Name(TypeBool(), "u1"); - f32.Define(*this, TypeFloat(32), "f32"); - u32.Define(*this, TypeInt(32, false), "u32"); - f16.Define(*this, TypeFloat(16), "f16"); - f64.Define(*this, TypeFloat(64), "f64"); + static constexpr size_t NUM_ARGS = sizeof...(Args); - true_value = ConstantTrue(u1); - false_value = ConstantFalse(u1); + template + using ArgType = std::tuple_element_t>; +}; - for (const IR::Function& function : program.functions) { - for (IR::Block* const block : function.blocks) { - block_label_map.emplace_back(block, OpLabel()); +template +void SetDefinition(EmitSPIRV& emit, EmitContext& ctx, IR::Inst* inst, Args... args) { + const Id forward_id{inst->Definition()}; + const bool has_forward_id{Sirit::ValidId(forward_id)}; + Id current_id{}; + if (has_forward_id) { + current_id = ctx.ExchangeCurrentId(forward_id); + } + const Id new_id{(emit.*method)(ctx, std::forward(args)...)}; + if (has_forward_id) { + ctx.ExchangeCurrentId(current_id); + } else { + inst->SetDefinition(new_id); + } +} + +template +ArgType Arg(EmitContext& ctx, const IR::Value& arg) { + if constexpr (std::is_same_v) { + return ctx.Def(arg); + } else if constexpr (std::is_same_v) { + return arg; + } else if constexpr (std::is_same_v) { + return arg.U32(); + } else if constexpr (std::is_same_v) { + return arg.Label(); + } +} + +template +void Invoke(EmitSPIRV& emit, EmitContext& ctx, IR::Inst* inst, std::index_sequence) { + using Traits = FuncTraits; + if constexpr (std::is_same_v) { + if constexpr (is_first_arg_inst) { + SetDefinition(emit, ctx, inst, inst, + Arg>(ctx, inst->Arg(I))...); + } else { + SetDefinition(emit, ctx, inst, + Arg>(ctx, inst->Arg(I))...); + } + } else { + if constexpr (is_first_arg_inst) { + (emit.*method)(ctx, inst, Arg>(ctx, inst->Arg(I))...); + } else { + (emit.*method)(ctx, Arg>(ctx, inst->Arg(I))...); } } - std::ranges::sort(block_label_map, {}, &std::pair::first); } -EmitContext::~EmitContext() = default; +template +void Invoke(EmitSPIRV& emit, EmitContext& ctx, IR::Inst* inst) { + using Traits = FuncTraits; + static_assert(Traits::NUM_ARGS >= 1, "Insufficient arguments"); + if constexpr (Traits::NUM_ARGS == 1) { + Invoke(emit, ctx, inst, std::make_index_sequence<0>{}); + } else { + using FirstArgType = typename Traits::template ArgType<1>; + static constexpr bool is_first_arg_inst = std::is_same_v; + using Indices = std::make_index_sequence; + Invoke(emit, ctx, inst, Indices{}); + } +} +} // Anonymous namespace EmitSPIRV::EmitSPIRV(IR::Program& program) { EmitContext ctx{program}; @@ -46,74 +98,32 @@ EmitSPIRV::EmitSPIRV(IR::Program& program) { for (IR::Function& function : program.functions) { func = ctx.OpFunction(ctx.void_id, spv::FunctionControlMask::MaskNone, void_function); for (IR::Block* const block : function.blocks) { - ctx.AddLabel(ctx.BlockLabel(block)); + ctx.AddLabel(block->Definition()); for (IR::Inst& inst : block->Instructions()) { EmitInst(ctx, &inst); } } ctx.OpFunctionEnd(); } - ctx.AddEntryPoint(spv::ExecutionModel::GLCompute, func, "main"); + boost::container::small_vector interfaces; + if (program.info.uses_workgroup_id) { + interfaces.push_back(ctx.workgroup_id); + } + if (program.info.uses_local_invocation_id) { + interfaces.push_back(ctx.local_invocation_id); + } + + const std::span interfaces_span(interfaces.data(), interfaces.size()); + ctx.AddEntryPoint(spv::ExecutionModel::Fragment, func, "main", interfaces_span); + ctx.AddExecutionMode(func, spv::ExecutionMode::OriginUpperLeft); std::vector result{ctx.Assemble()}; - std::FILE* file{std::fopen("shader.spv", "wb")}; + std::FILE* file{std::fopen("D:\\shader.spv", "wb")}; std::fwrite(result.data(), sizeof(u32), result.size(), file); std::fclose(file); - std::system("spirv-dis shader.spv"); - std::system("spirv-val shader.spv"); - std::system("spirv-cross shader.spv"); -} - -template -static void SetDefinition(EmitSPIRV& emit, EmitContext& ctx, IR::Inst* inst, Args... args) { - const Id forward_id{inst->Definition()}; - const bool has_forward_id{Sirit::ValidId(forward_id)}; - Id current_id{}; - if (has_forward_id) { - current_id = ctx.ExchangeCurrentId(forward_id); - } - const Id new_id{(emit.*method)(ctx, std::forward(args)...)}; - if (has_forward_id) { - ctx.ExchangeCurrentId(current_id); - } else { - inst->SetDefinition(new_id); - } -} - -template -static void Invoke(EmitSPIRV& emit, EmitContext& ctx, IR::Inst* inst) { - using M = decltype(method); - using std::is_invocable_r_v; - if constexpr (is_invocable_r_v) { - SetDefinition(emit, ctx, inst); - } else if constexpr (is_invocable_r_v) { - SetDefinition(emit, ctx, inst, ctx.Def(inst->Arg(0))); - } else if constexpr (is_invocable_r_v) { - SetDefinition(emit, ctx, inst, ctx.Def(inst->Arg(0)), ctx.Def(inst->Arg(1))); - } else if constexpr (is_invocable_r_v) { - SetDefinition(emit, ctx, inst, ctx.Def(inst->Arg(0)), ctx.Def(inst->Arg(1)), - ctx.Def(inst->Arg(2))); - } else if constexpr (is_invocable_r_v) { - SetDefinition(emit, ctx, inst, inst); - } else if constexpr (is_invocable_r_v) { - SetDefinition(emit, ctx, inst, inst, ctx.Def(inst->Arg(0)), ctx.Def(inst->Arg(1))); - } else if constexpr (is_invocable_r_v) { - SetDefinition(emit, ctx, inst, inst, ctx.Def(inst->Arg(0)), ctx.Def(inst->Arg(1)), - ctx.Def(inst->Arg(2))); - } else if constexpr (is_invocable_r_v) { - SetDefinition(emit, ctx, inst, ctx.Def(inst->Arg(0)), inst->Arg(1).U32()); - } else if constexpr (is_invocable_r_v) { - SetDefinition(emit, ctx, inst, inst->Arg(0)); - } else if constexpr (is_invocable_r_v) { - SetDefinition(emit, ctx, inst, inst->Arg(0), inst->Arg(1)); - } else if constexpr (is_invocable_r_v) { - (emit.*method)(ctx, inst); - } else if constexpr (is_invocable_r_v) { - (emit.*method)(ctx); - } else { - static_assert(false, "Bad format"); - } + std::system("spirv-dis D:\\shader.spv") == 0 && + std::system("spirv-val --uniform-buffer-standard-layout D:\\shader.spv") == 0 && + std::system("spirv-cross -V D:\\shader.spv") == 0; } void EmitSPIRV::EmitInst(EmitContext& ctx, IR::Inst* inst) { @@ -130,9 +140,9 @@ void EmitSPIRV::EmitInst(EmitContext& ctx, IR::Inst* inst) { static Id TypeId(const EmitContext& ctx, IR::Type type) { switch (type) { case IR::Type::U1: - return ctx.u1; + return ctx.U1; case IR::Type::U32: - return ctx.u32[1]; + return ctx.U32[1]; default: throw NotImplementedException("Phi node type {}", type); } @@ -162,7 +172,7 @@ Id EmitSPIRV::EmitPhi(EmitContext& ctx, IR::Inst* inst) { } IR::Block* const phi_block{inst->PhiBlock(index)}; operands.push_back(def); - operands.push_back(ctx.BlockLabel(phi_block)); + operands.push_back(phi_block->Definition()); } const Id result_type{TypeId(ctx, inst->Arg(0).Type())}; return ctx.OpPhi(result_type, std::span(operands.data(), operands.size())); @@ -174,29 +184,6 @@ void EmitSPIRV::EmitIdentity(EmitContext&) { throw NotImplementedException("SPIR-V Instruction"); } -// FIXME: Move to its own file -void EmitSPIRV::EmitBranch(EmitContext& ctx, IR::Inst* inst) { - ctx.OpBranch(ctx.BlockLabel(inst->Arg(0).Label())); -} - -void EmitSPIRV::EmitBranchConditional(EmitContext& ctx, IR::Inst* inst) { - ctx.OpBranchConditional(ctx.Def(inst->Arg(0)), ctx.BlockLabel(inst->Arg(1).Label()), - ctx.BlockLabel(inst->Arg(2).Label())); -} - -void EmitSPIRV::EmitLoopMerge(EmitContext& ctx, IR::Inst* inst) { - ctx.OpLoopMerge(ctx.BlockLabel(inst->Arg(0).Label()), ctx.BlockLabel(inst->Arg(1).Label()), - spv::LoopControlMask::MaskNone); -} - -void EmitSPIRV::EmitSelectionMerge(EmitContext& ctx, IR::Inst* inst) { - ctx.OpSelectionMerge(ctx.BlockLabel(inst->Arg(0).Label()), spv::SelectionControlMask::MaskNone); -} - -void EmitSPIRV::EmitReturn(EmitContext& ctx) { - ctx.OpReturn(); -} - void EmitSPIRV::EmitGetZeroFromOp(EmitContext&) { throw LogicError("Unreachable instruction"); } -- cgit v1.2.3 From 3a59fffaa16838985f9f953f30d1af4aa0f86252 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Tue, 16 Feb 2021 19:48:58 -0300 Subject: spirv: Implement EmitIdentity --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index c79c09774..55018332e 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -180,8 +180,8 @@ Id EmitSPIRV::EmitPhi(EmitContext& ctx, IR::Inst* inst) { void EmitSPIRV::EmitVoid(EmitContext&) {} -void EmitSPIRV::EmitIdentity(EmitContext&) { - throw NotImplementedException("SPIR-V Instruction"); +Id EmitSPIRV::EmitIdentity(EmitContext& ctx, const IR::Value& value) { + return ctx.Def(value); } void EmitSPIRV::EmitGetZeroFromOp(EmitContext&) { -- cgit v1.2.3 From 85cce78583bc2232428a8fb39e43182877c8d5ad Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Wed, 17 Feb 2021 00:59:28 -0300 Subject: shader: Primitive Vulkan integration --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 117 ++++++++++----------- 1 file changed, 57 insertions(+), 60 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 55018332e..d59718435 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -2,8 +2,11 @@ // Licensed under GPLv2 or any later version // Refer to the license.txt file included. -#include +#include +#include #include +#include +#include #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/frontend/ir/basic_block.h" @@ -14,10 +17,10 @@ namespace Shader::Backend::SPIRV { namespace { template -struct FuncTraits : FuncTraits {}; +struct FuncTraits : FuncTraits {}; -template -struct FuncTraits { +template +struct FuncTraits { using ReturnType = ReturnType_; static constexpr size_t NUM_ARGS = sizeof...(Args); @@ -26,15 +29,15 @@ struct FuncTraits { using ArgType = std::tuple_element_t>; }; -template -void SetDefinition(EmitSPIRV& emit, EmitContext& ctx, IR::Inst* inst, Args... args) { +template +void SetDefinition(EmitContext& ctx, IR::Inst* inst, Args... args) { const Id forward_id{inst->Definition()}; const bool has_forward_id{Sirit::ValidId(forward_id)}; Id current_id{}; if (has_forward_id) { current_id = ctx.ExchangeCurrentId(forward_id); } - const Id new_id{(emit.*method)(ctx, std::forward(args)...)}; + const Id new_id{func(ctx, std::forward(args)...)}; if (has_forward_id) { ctx.ExchangeCurrentId(current_id); } else { @@ -55,42 +58,62 @@ ArgType Arg(EmitContext& ctx, const IR::Value& arg) { } } -template -void Invoke(EmitSPIRV& emit, EmitContext& ctx, IR::Inst* inst, std::index_sequence) { - using Traits = FuncTraits; +template +void Invoke(EmitContext& ctx, IR::Inst* inst, std::index_sequence) { + using Traits = FuncTraits; if constexpr (std::is_same_v) { if constexpr (is_first_arg_inst) { - SetDefinition(emit, ctx, inst, inst, - Arg>(ctx, inst->Arg(I))...); + SetDefinition(ctx, inst, inst, Arg>(ctx, inst->Arg(I))...); } else { - SetDefinition(emit, ctx, inst, - Arg>(ctx, inst->Arg(I))...); + SetDefinition(ctx, inst, Arg>(ctx, inst->Arg(I))...); } } else { if constexpr (is_first_arg_inst) { - (emit.*method)(ctx, inst, Arg>(ctx, inst->Arg(I))...); + func(ctx, inst, Arg>(ctx, inst->Arg(I))...); } else { - (emit.*method)(ctx, Arg>(ctx, inst->Arg(I))...); + func(ctx, Arg>(ctx, inst->Arg(I))...); } } } -template -void Invoke(EmitSPIRV& emit, EmitContext& ctx, IR::Inst* inst) { - using Traits = FuncTraits; +template +void Invoke(EmitContext& ctx, IR::Inst* inst) { + using Traits = FuncTraits; static_assert(Traits::NUM_ARGS >= 1, "Insufficient arguments"); if constexpr (Traits::NUM_ARGS == 1) { - Invoke(emit, ctx, inst, std::make_index_sequence<0>{}); + Invoke(ctx, inst, std::make_index_sequence<0>{}); } else { using FirstArgType = typename Traits::template ArgType<1>; static constexpr bool is_first_arg_inst = std::is_same_v; using Indices = std::make_index_sequence; - Invoke(emit, ctx, inst, Indices{}); + Invoke(ctx, inst, Indices{}); + } +} + +void EmitInst(EmitContext& ctx, IR::Inst* inst) { + switch (inst->Opcode()) { +#define OPCODE(name, result_type, ...) \ + case IR::Opcode::name: \ + return Invoke<&Emit##name>(ctx, inst); +#include "shader_recompiler/frontend/ir/opcodes.inc" +#undef OPCODE + } + throw LogicError("Invalid opcode {}", inst->Opcode()); +} + +Id TypeId(const EmitContext& ctx, IR::Type type) { + switch (type) { + case IR::Type::U1: + return ctx.U1; + case IR::Type::U32: + return ctx.U32[1]; + default: + throw NotImplementedException("Phi node type {}", type); } } } // Anonymous namespace -EmitSPIRV::EmitSPIRV(IR::Program& program) { +std::vector EmitSPIRV(Environment& env, IR::Program& program) { EmitContext ctx{program}; const Id void_function{ctx.TypeFunction(ctx.void_id)}; // FIXME: Forward declare functions (needs sirit support) @@ -112,43 +135,17 @@ EmitSPIRV::EmitSPIRV(IR::Program& program) { if (program.info.uses_local_invocation_id) { interfaces.push_back(ctx.local_invocation_id); } - const std::span interfaces_span(interfaces.data(), interfaces.size()); - ctx.AddEntryPoint(spv::ExecutionModel::Fragment, func, "main", interfaces_span); - ctx.AddExecutionMode(func, spv::ExecutionMode::OriginUpperLeft); - - std::vector result{ctx.Assemble()}; - std::FILE* file{std::fopen("D:\\shader.spv", "wb")}; - std::fwrite(result.data(), sizeof(u32), result.size(), file); - std::fclose(file); - std::system("spirv-dis D:\\shader.spv") == 0 && - std::system("spirv-val --uniform-buffer-standard-layout D:\\shader.spv") == 0 && - std::system("spirv-cross -V D:\\shader.spv") == 0; -} + ctx.AddEntryPoint(spv::ExecutionModel::GLCompute, func, "main", interfaces_span); -void EmitSPIRV::EmitInst(EmitContext& ctx, IR::Inst* inst) { - switch (inst->Opcode()) { -#define OPCODE(name, result_type, ...) \ - case IR::Opcode::name: \ - return Invoke<&EmitSPIRV::Emit##name>(*this, ctx, inst); -#include "shader_recompiler/frontend/ir/opcodes.inc" -#undef OPCODE - } - throw LogicError("Invalid opcode {}", inst->Opcode()); -} + const std::array workgroup_size{env.WorkgroupSize()}; + ctx.AddExecutionMode(func, spv::ExecutionMode::LocalSize, workgroup_size[0], workgroup_size[1], + workgroup_size[2]); -static Id TypeId(const EmitContext& ctx, IR::Type type) { - switch (type) { - case IR::Type::U1: - return ctx.U1; - case IR::Type::U32: - return ctx.U32[1]; - default: - throw NotImplementedException("Phi node type {}", type); - } + return ctx.Assemble(); } -Id EmitSPIRV::EmitPhi(EmitContext& ctx, IR::Inst* inst) { +Id EmitPhi(EmitContext& ctx, IR::Inst* inst) { const size_t num_args{inst->NumArgs()}; boost::container::small_vector operands; operands.reserve(num_args * 2); @@ -178,25 +175,25 @@ Id EmitSPIRV::EmitPhi(EmitContext& ctx, IR::Inst* inst) { return ctx.OpPhi(result_type, std::span(operands.data(), operands.size())); } -void EmitSPIRV::EmitVoid(EmitContext&) {} +void EmitVoid(EmitContext&) {} -Id EmitSPIRV::EmitIdentity(EmitContext& ctx, const IR::Value& value) { +Id EmitIdentity(EmitContext& ctx, const IR::Value& value) { return ctx.Def(value); } -void EmitSPIRV::EmitGetZeroFromOp(EmitContext&) { +void EmitGetZeroFromOp(EmitContext&) { throw LogicError("Unreachable instruction"); } -void EmitSPIRV::EmitGetSignFromOp(EmitContext&) { +void EmitGetSignFromOp(EmitContext&) { throw LogicError("Unreachable instruction"); } -void EmitSPIRV::EmitGetCarryFromOp(EmitContext&) { +void EmitGetCarryFromOp(EmitContext&) { throw LogicError("Unreachable instruction"); } -void EmitSPIRV::EmitGetOverflowFromOp(EmitContext&) { +void EmitGetOverflowFromOp(EmitContext&) { throw LogicError("Unreachable instruction"); } -- cgit v1.2.3 From 6db69990da9f232e6d982cdcb69c2e27d93075cf Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Fri, 19 Feb 2021 18:10:18 -0300 Subject: spirv: Add lower fp16 to fp32 pass --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 2 ++ 1 file changed, 2 insertions(+) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index d59718435..4ce07c281 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -14,6 +14,8 @@ #include "shader_recompiler/frontend/ir/microinstruction.h" #include "shader_recompiler/frontend/ir/program.h" +#pragma optimize("", off) + namespace Shader::Backend::SPIRV { namespace { template -- cgit v1.2.3 From e2bc05b17d91854cbb9c0ce3647141bf7d33143e Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Sat, 20 Feb 2021 03:30:13 -0300 Subject: shader: Add denorm flush support --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 63 ++++++++++++++++++++-- 1 file changed, 58 insertions(+), 5 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 4ce07c281..2519e446a 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -14,8 +14,6 @@ #include "shader_recompiler/frontend/ir/microinstruction.h" #include "shader_recompiler/frontend/ir/program.h" -#pragma optimize("", off) - namespace Shader::Backend::SPIRV { namespace { template @@ -113,9 +111,61 @@ Id TypeId(const EmitContext& ctx, IR::Type type) { throw NotImplementedException("Phi node type {}", type); } } + +void SetupDenormControl(const Profile& profile, const IR::Program& program, EmitContext& ctx, + Id main_func) { + if (!profile.support_float_controls) { + return; + } + const Info& info{program.info}; + if (!info.uses_fp32_denorms_flush && !info.uses_fp32_denorms_preserve && + !info.uses_fp16_denorms_flush && !info.uses_fp16_denorms_preserve) { + return; + } + ctx.AddExtension("SPV_KHR_float_controls"); + + if (info.uses_fp32_denorms_flush && info.uses_fp32_denorms_preserve) { + // LOG_ERROR(HW_GPU, "Fp32 denorm flush and preserve on the same shader"); + } else if (info.uses_fp32_denorms_flush) { + if (profile.support_fp32_denorm_flush) { + ctx.AddCapability(spv::Capability::DenormFlushToZero); + ctx.AddExecutionMode(main_func, spv::ExecutionMode::DenormFlushToZero, 32U); + } else { + // Drivers will most likely flush denorms by default, no need to warn + } + } else if (info.uses_fp32_denorms_preserve) { + if (profile.support_fp32_denorm_preserve) { + ctx.AddCapability(spv::Capability::DenormPreserve); + ctx.AddExecutionMode(main_func, spv::ExecutionMode::DenormPreserve, 32U); + } else { + // LOG_WARNING(HW_GPU, "Fp32 denorm preserve used in shader without host support"); + } + } + if (!profile.support_separate_denorm_behavior) { + // No separate denorm behavior + return; + } + if (info.uses_fp16_denorms_flush && info.uses_fp16_denorms_preserve) { + // LOG_ERROR(HW_GPU, "Fp16 denorm flush and preserve on the same shader"); + } else if (info.uses_fp16_denorms_flush) { + if (profile.support_fp16_denorm_flush) { + ctx.AddCapability(spv::Capability::DenormFlushToZero); + ctx.AddExecutionMode(main_func, spv::ExecutionMode::DenormPreserve, 16U); + } else { + // Same as fp32, no need to warn as most drivers will flush by default + } + } else if (info.uses_fp32_denorms_preserve) { + if (profile.support_fp16_denorm_preserve) { + ctx.AddCapability(spv::Capability::DenormPreserve); + ctx.AddExecutionMode(main_func, spv::ExecutionMode::DenormPreserve, 16U); + } else { + // LOG_WARNING(HW_GPU, "Fp16 denorm preserve used in shader without host support"); + } + } +} } // Anonymous namespace -std::vector EmitSPIRV(Environment& env, IR::Program& program) { +std::vector EmitSPIRV(const Profile& profile, Environment& env, IR::Program& program) { EmitContext ctx{program}; const Id void_function{ctx.TypeFunction(ctx.void_id)}; // FIXME: Forward declare functions (needs sirit support) @@ -131,10 +181,11 @@ std::vector EmitSPIRV(Environment& env, IR::Program& program) { ctx.OpFunctionEnd(); } boost::container::small_vector interfaces; - if (program.info.uses_workgroup_id) { + const Info& info{program.info}; + if (info.uses_workgroup_id) { interfaces.push_back(ctx.workgroup_id); } - if (program.info.uses_local_invocation_id) { + if (info.uses_local_invocation_id) { interfaces.push_back(ctx.local_invocation_id); } const std::span interfaces_span(interfaces.data(), interfaces.size()); @@ -144,6 +195,8 @@ std::vector EmitSPIRV(Environment& env, IR::Program& program) { ctx.AddExecutionMode(func, spv::ExecutionMode::LocalSize, workgroup_size[0], workgroup_size[1], workgroup_size[2]); + SetupDenormControl(profile, program, ctx, func); + return ctx.Assemble(); } -- cgit v1.2.3 From 274897dfd59b4d08029ab7e93be4f84654abcdc8 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Sun, 21 Feb 2021 23:42:38 -0300 Subject: spirv: Fixes and Intel specific workarounds --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 2519e446a..f3aca90d0 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -150,11 +150,11 @@ void SetupDenormControl(const Profile& profile, const IR::Program& program, Emit } else if (info.uses_fp16_denorms_flush) { if (profile.support_fp16_denorm_flush) { ctx.AddCapability(spv::Capability::DenormFlushToZero); - ctx.AddExecutionMode(main_func, spv::ExecutionMode::DenormPreserve, 16U); + ctx.AddExecutionMode(main_func, spv::ExecutionMode::DenormFlushToZero, 16U); } else { // Same as fp32, no need to warn as most drivers will flush by default } - } else if (info.uses_fp32_denorms_preserve) { + } else if (info.uses_fp16_denorms_preserve) { if (profile.support_fp16_denorm_preserve) { ctx.AddCapability(spv::Capability::DenormPreserve); ctx.AddExecutionMode(main_func, spv::ExecutionMode::DenormPreserve, 16U); @@ -166,7 +166,7 @@ void SetupDenormControl(const Profile& profile, const IR::Program& program, Emit } // Anonymous namespace std::vector EmitSPIRV(const Profile& profile, Environment& env, IR::Program& program) { - EmitContext ctx{program}; + EmitContext ctx{profile, program}; const Id void_function{ctx.TypeFunction(ctx.void_id)}; // FIXME: Forward declare functions (needs sirit support) Id func{}; -- cgit v1.2.3 From 7496bbf7584049fb9c99cf705e9cc16aee61a55a Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Wed, 24 Feb 2021 18:31:32 -0300 Subject: spirv: Add support for self-referencing phi nodes --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 13 ++++++++++--- 1 file changed, 10 insertions(+), 3 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index f3aca90d0..bcd6bda28 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -217,9 +217,16 @@ Id EmitPhi(EmitContext& ctx, IR::Inst* inst) { IR::Inst* const arg_inst{arg.Inst()}; def = arg_inst->Definition(); if (!Sirit::ValidId(def)) { - // If it hasn't been defined, get a forward declaration - def = ctx.ForwardDeclarationId(); - arg_inst->SetDefinition(def); + if (arg_inst == inst) { + // This is a self referencing phi node + def = ctx.CurrentId(); + // Self-referencing definition will be set by the caller + } else { + // If it hasn't been defined and it's not a self reference, + // get a forward declaration + def = ctx.ForwardDeclarationId(); + arg_inst->SetDefinition(def); + } } } IR::Block* const phi_block{inst->PhiBlock(index)}; -- cgit v1.2.3 From 726625cf5057157fb5e4c9c210676930ff520dd2 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Wed, 24 Feb 2021 18:37:47 -0300 Subject: spirv: Move phi arguments emit to a separate function --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 54 +++++++++++----------- 1 file changed, 27 insertions(+), 27 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index bcd6bda28..8097fe82d 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -163,6 +163,31 @@ void SetupDenormControl(const Profile& profile, const IR::Program& program, Emit } } } + +Id PhiArgDef(EmitContext& ctx, IR::Inst* inst, size_t index) { + // Phi nodes can have forward declarations, if an argument is not defined provide a forward + // declaration of it. Invoke will take care of giving it the right definition when it's + // actually defined. + const IR::Value arg{inst->Arg(index)}; + if (arg.IsImmediate()) { + // Let the context handle immediate definitions, as it already knows how + return ctx.Def(arg); + } + IR::Inst* const arg_inst{arg.Inst()}; + if (const Id def{arg_inst->Definition()}; Sirit::ValidId(def)) { + // Return the current definition if it exists + return def; + } + if (arg_inst == inst) { + // This is a self referencing phi node + // Self-referencing definition will be set by the caller, so just grab the current id + return ctx.CurrentId(); + } + // If it hasn't been defined and it's not a self reference, get a forward declaration + const Id def{ctx.ForwardDeclarationId()}; + arg_inst->SetDefinition(def); + return def; +} } // Anonymous namespace std::vector EmitSPIRV(const Profile& profile, Environment& env, IR::Program& program) { @@ -205,33 +230,8 @@ Id EmitPhi(EmitContext& ctx, IR::Inst* inst) { boost::container::small_vector operands; operands.reserve(num_args * 2); for (size_t index = 0; index < num_args; ++index) { - // Phi nodes can have forward declarations, if an argument is not defined provide a forward - // declaration of it. Invoke will take care of giving it the right definition when it's - // actually defined. - const IR::Value arg{inst->Arg(index)}; - Id def{}; - if (arg.IsImmediate()) { - // Let the context handle immediate definitions, as it already knows how - def = ctx.Def(arg); - } else { - IR::Inst* const arg_inst{arg.Inst()}; - def = arg_inst->Definition(); - if (!Sirit::ValidId(def)) { - if (arg_inst == inst) { - // This is a self referencing phi node - def = ctx.CurrentId(); - // Self-referencing definition will be set by the caller - } else { - // If it hasn't been defined and it's not a self reference, - // get a forward declaration - def = ctx.ForwardDeclarationId(); - arg_inst->SetDefinition(def); - } - } - } - IR::Block* const phi_block{inst->PhiBlock(index)}; - operands.push_back(def); - operands.push_back(phi_block->Definition()); + operands.push_back(PhiArgDef(ctx, inst, index)); + operands.push_back(inst->PhiBlock(index)->Definition()); } const Id result_type{TypeId(ctx, inst->Arg(0).Type())}; return ctx.OpPhi(result_type, std::span(operands.data(), operands.size())); -- cgit v1.2.3 From ab463712474de5f99eec137a9c6233e55fe184f0 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Mon, 8 Mar 2021 18:31:53 -0300 Subject: shader: Initial support for textures and TEX --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 12 ++++++++++++ 1 file changed, 12 insertions(+) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 8097fe82d..a94e9cb2d 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -221,6 +221,14 @@ std::vector EmitSPIRV(const Profile& profile, Environment& env, IR::Program workgroup_size[2]); SetupDenormControl(profile, program, ctx, func); + if (info.uses_sampled_1d) { + ctx.AddCapability(spv::Capability::Sampled1D); + } + if (info.uses_sparse_residency) { + ctx.AddCapability(spv::Capability::SparseResidency); + } + // TODO: Track this usage + ctx.AddCapability(spv::Capability::ImageGatherExtended); return ctx.Assemble(); } @@ -259,4 +267,8 @@ void EmitGetOverflowFromOp(EmitContext&) { throw LogicError("Unreachable instruction"); } +void EmitGetSparseFromOp(EmitContext&) { + throw LogicError("Unreachable instruction"); +} + } // namespace Shader::Backend::SPIRV -- cgit v1.2.3 From b9f7bf4472b8e0a5aad1aec3a5ff5bb56470bfff Mon Sep 17 00:00:00 2001 From: ameerj <52414509+ameerj@users.noreply.github.com> Date: Sun, 14 Mar 2021 01:51:40 -0500 Subject: spirv: Add SignedZeroInfNanPreserve logic --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 6 ++++++ 1 file changed, 6 insertions(+) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index a94e9cb2d..c7cba6279 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -124,6 +124,12 @@ void SetupDenormControl(const Profile& profile, const IR::Program& program, Emit } ctx.AddExtension("SPV_KHR_float_controls"); + if (info.uses_fp16 && profile.support_fp16_signed_zero_nan_preserve) { + ctx.AddExecutionMode(main_func, spv::ExecutionMode::SignedZeroInfNanPreserve); + } + if (profile.support_fp32_signed_zero_nan_preserve) { + ctx.AddExecutionMode(main_func, spv::ExecutionMode::SignedZeroInfNanPreserve); + } if (info.uses_fp32_denorms_flush && info.uses_fp32_denorms_preserve) { // LOG_ERROR(HW_GPU, "Fp32 denorm flush and preserve on the same shader"); } else if (info.uses_fp32_denorms_flush) { -- cgit v1.2.3 From 71f96fa6366dc6dd306a953bca1b958fb32bc55a Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Sun, 14 Mar 2021 03:41:05 -0300 Subject: shader: Implement CAL inlining function calls --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 17 ++++++----------- 1 file changed, 6 insertions(+), 11 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index c7cba6279..7e7db9161 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -10,7 +10,6 @@ #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/frontend/ir/basic_block.h" -#include "shader_recompiler/frontend/ir/function.h" #include "shader_recompiler/frontend/ir/microinstruction.h" #include "shader_recompiler/frontend/ir/program.h" @@ -199,18 +198,14 @@ Id PhiArgDef(EmitContext& ctx, IR::Inst* inst, size_t index) { std::vector EmitSPIRV(const Profile& profile, Environment& env, IR::Program& program) { EmitContext ctx{profile, program}; const Id void_function{ctx.TypeFunction(ctx.void_id)}; - // FIXME: Forward declare functions (needs sirit support) - Id func{}; - for (IR::Function& function : program.functions) { - func = ctx.OpFunction(ctx.void_id, spv::FunctionControlMask::MaskNone, void_function); - for (IR::Block* const block : function.blocks) { - ctx.AddLabel(block->Definition()); - for (IR::Inst& inst : block->Instructions()) { - EmitInst(ctx, &inst); - } + const Id func{ctx.OpFunction(ctx.void_id, spv::FunctionControlMask::MaskNone, void_function)}; + for (IR::Block* const block : program.blocks) { + ctx.AddLabel(block->Definition()); + for (IR::Inst& inst : block->Instructions()) { + EmitInst(ctx, &inst); } - ctx.OpFunctionEnd(); } + ctx.OpFunctionEnd(); boost::container::small_vector interfaces; const Info& info{program.info}; if (info.uses_workgroup_id) { -- cgit v1.2.3 From fa2f6e38f4d465ba6e5efe6c6bd23d8ef39b080d Mon Sep 17 00:00:00 2001 From: ameerj <52414509+ameerj@users.noreply.github.com> Date: Tue, 16 Mar 2021 00:57:07 -0400 Subject: shader: Implement FSET and FSETP Also fix oversight with adding SignedZeroInfNanPreserve execution mode. --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 7e7db9161..50c0f7243 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -124,10 +124,12 @@ void SetupDenormControl(const Profile& profile, const IR::Program& program, Emit ctx.AddExtension("SPV_KHR_float_controls"); if (info.uses_fp16 && profile.support_fp16_signed_zero_nan_preserve) { - ctx.AddExecutionMode(main_func, spv::ExecutionMode::SignedZeroInfNanPreserve); + ctx.AddCapability(spv::Capability::SignedZeroInfNanPreserve); + ctx.AddExecutionMode(main_func, spv::ExecutionMode::SignedZeroInfNanPreserve, 16U); } if (profile.support_fp32_signed_zero_nan_preserve) { - ctx.AddExecutionMode(main_func, spv::ExecutionMode::SignedZeroInfNanPreserve); + ctx.AddCapability(spv::Capability::SignedZeroInfNanPreserve); + ctx.AddExecutionMode(main_func, spv::ExecutionMode::SignedZeroInfNanPreserve, 32U); } if (info.uses_fp32_denorms_flush && info.uses_fp32_denorms_preserve) { // LOG_ERROR(HW_GPU, "Fp32 denorm flush and preserve on the same shader"); -- cgit v1.2.3 From 260743f371236f7c57b01334b1c3474b15a47c39 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Fri, 19 Mar 2021 19:28:31 -0300 Subject: shader: Add partial rasterizer integration --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 44 +++++++++++++++------- 1 file changed, 30 insertions(+), 14 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 50c0f7243..b8978b94a 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -54,6 +54,8 @@ ArgType Arg(EmitContext& ctx, const IR::Value& arg) { return arg.U32(); } else if constexpr (std::is_same_v) { return arg.Label(); + } else if constexpr (std::is_same_v) { + return arg.Attribute(); } } @@ -197,8 +199,9 @@ Id PhiArgDef(EmitContext& ctx, IR::Inst* inst, size_t index) { } } // Anonymous namespace -std::vector EmitSPIRV(const Profile& profile, Environment& env, IR::Program& program) { - EmitContext ctx{profile, program}; +std::vector EmitSPIRV(const Profile& profile, Environment& env, IR::Program& program, + u32& binding) { + EmitContext ctx{profile, program, binding}; const Id void_function{ctx.TypeFunction(ctx.void_id)}; const Id func{ctx.OpFunction(ctx.void_id, spv::FunctionControlMask::MaskNone, void_function)}; for (IR::Block* const block : program.blocks) { @@ -208,28 +211,41 @@ std::vector EmitSPIRV(const Profile& profile, Environment& env, IR::Program } } ctx.OpFunctionEnd(); - boost::container::small_vector interfaces; - const Info& info{program.info}; - if (info.uses_workgroup_id) { - interfaces.push_back(ctx.workgroup_id); + + const std::span interfaces(ctx.interfaces.data(), ctx.interfaces.size()); + spv::ExecutionModel execution_model{}; + switch (env.ShaderStage()) { + case Shader::Stage::Compute: { + const std::array workgroup_size{env.WorkgroupSize()}; + execution_model = spv::ExecutionModel::GLCompute; + ctx.AddExecutionMode(func, spv::ExecutionMode::LocalSize, workgroup_size[0], + workgroup_size[1], workgroup_size[2]); + break; } - if (info.uses_local_invocation_id) { - interfaces.push_back(ctx.local_invocation_id); + case Shader::Stage::VertexB: + execution_model = spv::ExecutionModel::Vertex; + break; + case Shader::Stage::Fragment: + execution_model = spv::ExecutionModel::Fragment; + ctx.AddExecutionMode(func, spv::ExecutionMode::OriginUpperLeft); + break; + default: + throw NotImplementedException("Stage {}", env.ShaderStage()); } - const std::span interfaces_span(interfaces.data(), interfaces.size()); - ctx.AddEntryPoint(spv::ExecutionModel::GLCompute, func, "main", interfaces_span); - - const std::array workgroup_size{env.WorkgroupSize()}; - ctx.AddExecutionMode(func, spv::ExecutionMode::LocalSize, workgroup_size[0], workgroup_size[1], - workgroup_size[2]); + ctx.AddEntryPoint(execution_model, func, "main", interfaces); SetupDenormControl(profile, program, ctx, func); + const Info& info{program.info}; if (info.uses_sampled_1d) { ctx.AddCapability(spv::Capability::Sampled1D); } if (info.uses_sparse_residency) { ctx.AddCapability(spv::Capability::SparseResidency); } + if (info.uses_demote_to_helper_invocation) { + ctx.AddExtension("SPV_EXT_demote_to_helper_invocation"); + ctx.AddCapability(spv::Capability::DemoteToHelperInvocationEXT); + } // TODO: Track this usage ctx.AddCapability(spv::Capability::ImageGatherExtended); -- cgit v1.2.3 From 76c8a962ac4eae77e71d66a72c448930240339f9 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Sat, 20 Mar 2021 19:11:56 -0300 Subject: spirv: Implement VertexId and InstanceId, refactor code --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 107 ++++++++++++--------- 1 file changed, 60 insertions(+), 47 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index b8978b94a..efd0b70b7 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -113,6 +113,43 @@ Id TypeId(const EmitContext& ctx, IR::Type type) { } } +Id DefineMain(EmitContext& ctx, IR::Program& program) { + const Id void_function{ctx.TypeFunction(ctx.void_id)}; + const Id main{ctx.OpFunction(ctx.void_id, spv::FunctionControlMask::MaskNone, void_function)}; + for (IR::Block* const block : program.blocks) { + ctx.AddLabel(block->Definition()); + for (IR::Inst& inst : block->Instructions()) { + EmitInst(ctx, &inst); + } + } + ctx.OpFunctionEnd(); + return main; +} + +void DefineEntryPoint(Environment& env, EmitContext& ctx, Id main) { + const std::span interfaces(ctx.interfaces.data(), ctx.interfaces.size()); + spv::ExecutionModel execution_model{}; + switch (env.ShaderStage()) { + case Shader::Stage::Compute: { + const std::array workgroup_size{env.WorkgroupSize()}; + execution_model = spv::ExecutionModel::GLCompute; + ctx.AddExecutionMode(main, spv::ExecutionMode::LocalSize, workgroup_size[0], + workgroup_size[1], workgroup_size[2]); + break; + } + case Shader::Stage::VertexB: + execution_model = spv::ExecutionModel::Vertex; + break; + case Shader::Stage::Fragment: + execution_model = spv::ExecutionModel::Fragment; + ctx.AddExecutionMode(main, spv::ExecutionMode::OriginUpperLeft); + break; + default: + throw NotImplementedException("Stage {}", env.ShaderStage()); + } + ctx.AddEntryPoint(execution_model, main, "main", interfaces); +} + void SetupDenormControl(const Profile& profile, const IR::Program& program, EmitContext& ctx, Id main_func) { if (!profile.support_float_controls) { @@ -173,6 +210,25 @@ void SetupDenormControl(const Profile& profile, const IR::Program& program, Emit } } +void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ctx) { + if (info.uses_sampled_1d) { + ctx.AddCapability(spv::Capability::Sampled1D); + } + if (info.uses_sparse_residency) { + ctx.AddCapability(spv::Capability::SparseResidency); + } + if (info.uses_demote_to_helper_invocation) { + ctx.AddExtension("SPV_EXT_demote_to_helper_invocation"); + ctx.AddCapability(spv::Capability::DemoteToHelperInvocationEXT); + } + if (!profile.support_vertex_instance_id && (info.loads_instance_id || info.loads_vertex_id)) { + ctx.AddExtension("SPV_KHR_shader_draw_parameters"); + ctx.AddCapability(spv::Capability::DrawParameters); + } + // TODO: Track this usage + ctx.AddCapability(spv::Capability::ImageGatherExtended); +} + Id PhiArgDef(EmitContext& ctx, IR::Inst* inst, size_t index) { // Phi nodes can have forward declarations, if an argument is not defined provide a forward // declaration of it. Invoke will take care of giving it the right definition when it's @@ -202,53 +258,10 @@ Id PhiArgDef(EmitContext& ctx, IR::Inst* inst, size_t index) { std::vector EmitSPIRV(const Profile& profile, Environment& env, IR::Program& program, u32& binding) { EmitContext ctx{profile, program, binding}; - const Id void_function{ctx.TypeFunction(ctx.void_id)}; - const Id func{ctx.OpFunction(ctx.void_id, spv::FunctionControlMask::MaskNone, void_function)}; - for (IR::Block* const block : program.blocks) { - ctx.AddLabel(block->Definition()); - for (IR::Inst& inst : block->Instructions()) { - EmitInst(ctx, &inst); - } - } - ctx.OpFunctionEnd(); - - const std::span interfaces(ctx.interfaces.data(), ctx.interfaces.size()); - spv::ExecutionModel execution_model{}; - switch (env.ShaderStage()) { - case Shader::Stage::Compute: { - const std::array workgroup_size{env.WorkgroupSize()}; - execution_model = spv::ExecutionModel::GLCompute; - ctx.AddExecutionMode(func, spv::ExecutionMode::LocalSize, workgroup_size[0], - workgroup_size[1], workgroup_size[2]); - break; - } - case Shader::Stage::VertexB: - execution_model = spv::ExecutionModel::Vertex; - break; - case Shader::Stage::Fragment: - execution_model = spv::ExecutionModel::Fragment; - ctx.AddExecutionMode(func, spv::ExecutionMode::OriginUpperLeft); - break; - default: - throw NotImplementedException("Stage {}", env.ShaderStage()); - } - ctx.AddEntryPoint(execution_model, func, "main", interfaces); - - SetupDenormControl(profile, program, ctx, func); - const Info& info{program.info}; - if (info.uses_sampled_1d) { - ctx.AddCapability(spv::Capability::Sampled1D); - } - if (info.uses_sparse_residency) { - ctx.AddCapability(spv::Capability::SparseResidency); - } - if (info.uses_demote_to_helper_invocation) { - ctx.AddExtension("SPV_EXT_demote_to_helper_invocation"); - ctx.AddCapability(spv::Capability::DemoteToHelperInvocationEXT); - } - // TODO: Track this usage - ctx.AddCapability(spv::Capability::ImageGatherExtended); - + const Id main{DefineMain(ctx, program)}; + DefineEntryPoint(env, ctx, main); + SetupDenormControl(profile, program, ctx, main); + SetupCapabilities(profile, program.info, ctx); return ctx.Assemble(); } -- cgit v1.2.3 From e4e1cc11b8f7649171fe922b2899e57120bfba53 Mon Sep 17 00:00:00 2001 From: ameerj <52414509+ameerj@users.noreply.github.com> Date: Sun, 21 Mar 2021 19:28:37 -0400 Subject: shader: Implement DMNMX, DSET, DSETP --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 39 ++++++++++++---------- 1 file changed, 21 insertions(+), 18 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index efd0b70b7..93e851133 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -152,24 +152,7 @@ void DefineEntryPoint(Environment& env, EmitContext& ctx, Id main) { void SetupDenormControl(const Profile& profile, const IR::Program& program, EmitContext& ctx, Id main_func) { - if (!profile.support_float_controls) { - return; - } const Info& info{program.info}; - if (!info.uses_fp32_denorms_flush && !info.uses_fp32_denorms_preserve && - !info.uses_fp16_denorms_flush && !info.uses_fp16_denorms_preserve) { - return; - } - ctx.AddExtension("SPV_KHR_float_controls"); - - if (info.uses_fp16 && profile.support_fp16_signed_zero_nan_preserve) { - ctx.AddCapability(spv::Capability::SignedZeroInfNanPreserve); - ctx.AddExecutionMode(main_func, spv::ExecutionMode::SignedZeroInfNanPreserve, 16U); - } - if (profile.support_fp32_signed_zero_nan_preserve) { - ctx.AddCapability(spv::Capability::SignedZeroInfNanPreserve); - ctx.AddExecutionMode(main_func, spv::ExecutionMode::SignedZeroInfNanPreserve, 32U); - } if (info.uses_fp32_denorms_flush && info.uses_fp32_denorms_preserve) { // LOG_ERROR(HW_GPU, "Fp32 denorm flush and preserve on the same shader"); } else if (info.uses_fp32_denorms_flush) { @@ -210,6 +193,22 @@ void SetupDenormControl(const Profile& profile, const IR::Program& program, Emit } } +void SetupSignedNanCapabilities(const Profile& profile, const IR::Program& program, + EmitContext& ctx, Id main_func) { + if (program.info.uses_fp16 && profile.support_fp16_signed_zero_nan_preserve) { + ctx.AddCapability(spv::Capability::SignedZeroInfNanPreserve); + ctx.AddExecutionMode(main_func, spv::ExecutionMode::SignedZeroInfNanPreserve, 16U); + } + if (profile.support_fp32_signed_zero_nan_preserve) { + ctx.AddCapability(spv::Capability::SignedZeroInfNanPreserve); + ctx.AddExecutionMode(main_func, spv::ExecutionMode::SignedZeroInfNanPreserve, 32U); + } + if (program.info.uses_fp64 && profile.support_fp64_signed_zero_nan_preserve) { + ctx.AddCapability(spv::Capability::SignedZeroInfNanPreserve); + ctx.AddExecutionMode(main_func, spv::ExecutionMode::SignedZeroInfNanPreserve, 64U); + } +} + void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ctx) { if (info.uses_sampled_1d) { ctx.AddCapability(spv::Capability::Sampled1D); @@ -260,7 +259,11 @@ std::vector EmitSPIRV(const Profile& profile, Environment& env, IR::Program EmitContext ctx{profile, program, binding}; const Id main{DefineMain(ctx, program)}; DefineEntryPoint(env, ctx, main); - SetupDenormControl(profile, program, ctx, main); + if (profile.support_float_controls) { + ctx.AddExtension("SPV_KHR_float_controls"); + SetupDenormControl(profile, program, ctx, main); + SetupSignedNanCapabilities(profile, program, ctx, main); + } SetupCapabilities(profile, program.info, ctx); return ctx.Assemble(); } -- cgit v1.2.3 From 3d07cef009cf9e287744c7771c67166ef5761ce8 Mon Sep 17 00:00:00 2001 From: ameerj <52414509+ameerj@users.noreply.github.com> Date: Tue, 23 Mar 2021 20:27:17 -0400 Subject: shader: Implement VOTE --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 9 +++++++++ 1 file changed, 9 insertions(+) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 93e851133..107403912 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -224,6 +224,15 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct ctx.AddExtension("SPV_KHR_shader_draw_parameters"); ctx.AddCapability(spv::Capability::DrawParameters); } + if (info.uses_subgroup_vote && profile.support_vote) { + ctx.AddExtension("SPV_KHR_shader_ballot"); + ctx.AddCapability(spv::Capability::SubgroupBallotKHR); + if (!profile.warp_size_potentially_larger_than_guest) { + // vote ops are only used when not taking the long path + ctx.AddExtension("SPV_KHR_subgroup_vote"); + ctx.AddCapability(spv::Capability::SubgroupVoteKHR); + } + } // TODO: Track this usage ctx.AddCapability(spv::Capability::ImageGatherExtended); } -- cgit v1.2.3 From 32c5483beb2f79f5d55eb2906f2bfdfa1698bca3 Mon Sep 17 00:00:00 2001 From: ameerj <52414509+ameerj@users.noreply.github.com> Date: Thu, 25 Mar 2021 11:31:37 -0400 Subject: shader: Implement SHFL --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 107403912..cee72f50d 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -224,7 +224,7 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct ctx.AddExtension("SPV_KHR_shader_draw_parameters"); ctx.AddCapability(spv::Capability::DrawParameters); } - if (info.uses_subgroup_vote && profile.support_vote) { + if ((info.uses_subgroup_vote || info.uses_subgroup_invocation_id) && profile.support_vote) { ctx.AddExtension("SPV_KHR_shader_ballot"); ctx.AddCapability(spv::Capability::SubgroupBallotKHR); if (!profile.warp_size_potentially_larger_than_guest) { @@ -315,4 +315,8 @@ void EmitGetSparseFromOp(EmitContext&) { throw LogicError("Unreachable instruction"); } +void EmitGetInBoundsFromOp(EmitContext&) { + throw LogicError("Unreachable instruction"); +} + } // namespace Shader::Backend::SPIRV -- cgit v1.2.3 From 17063d16a3cfe6542e74265739191e1d018fc456 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Fri, 26 Mar 2021 18:45:38 -0300 Subject: shader: Implement TXQ and fix FragDepth --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 10 +++++++--- 1 file changed, 7 insertions(+), 3 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index cee72f50d..4bed16e7b 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -126,10 +126,10 @@ Id DefineMain(EmitContext& ctx, IR::Program& program) { return main; } -void DefineEntryPoint(Environment& env, EmitContext& ctx, Id main) { +void DefineEntryPoint(Environment& env, const IR::Program& program, EmitContext& ctx, Id main) { const std::span interfaces(ctx.interfaces.data(), ctx.interfaces.size()); spv::ExecutionModel execution_model{}; - switch (env.ShaderStage()) { + switch (program.stage) { case Shader::Stage::Compute: { const std::array workgroup_size{env.WorkgroupSize()}; execution_model = spv::ExecutionModel::GLCompute; @@ -143,6 +143,9 @@ void DefineEntryPoint(Environment& env, EmitContext& ctx, Id main) { case Shader::Stage::Fragment: execution_model = spv::ExecutionModel::Fragment; ctx.AddExecutionMode(main, spv::ExecutionMode::OriginUpperLeft); + if (program.info.stores_frag_depth) { + ctx.AddExecutionMode(main, spv::ExecutionMode::DepthReplacing); + } break; default: throw NotImplementedException("Stage {}", env.ShaderStage()); @@ -235,6 +238,7 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct } // TODO: Track this usage ctx.AddCapability(spv::Capability::ImageGatherExtended); + ctx.AddCapability(spv::Capability::ImageQuery); } Id PhiArgDef(EmitContext& ctx, IR::Inst* inst, size_t index) { @@ -267,7 +271,7 @@ std::vector EmitSPIRV(const Profile& profile, Environment& env, IR::Program u32& binding) { EmitContext ctx{profile, program, binding}; const Id main{DefineMain(ctx, program)}; - DefineEntryPoint(env, ctx, main); + DefineEntryPoint(env, program, ctx, main); if (profile.support_float_controls) { ctx.AddExtension("SPV_KHR_float_controls"); SetupDenormControl(profile, program, ctx, main); -- cgit v1.2.3 From 675a82416d7775dc7a252a5d8f5b704e6b8f2326 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Sat, 27 Mar 2021 03:08:31 -0300 Subject: spirv: Remove dependencies on Environment when generating SPIR-V --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 11 +++++------ 1 file changed, 5 insertions(+), 6 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 4bed16e7b..2e7e6bb0c 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -126,12 +126,12 @@ Id DefineMain(EmitContext& ctx, IR::Program& program) { return main; } -void DefineEntryPoint(Environment& env, const IR::Program& program, EmitContext& ctx, Id main) { +void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) { const std::span interfaces(ctx.interfaces.data(), ctx.interfaces.size()); spv::ExecutionModel execution_model{}; switch (program.stage) { case Shader::Stage::Compute: { - const std::array workgroup_size{env.WorkgroupSize()}; + const std::array workgroup_size{program.workgroup_size}; execution_model = spv::ExecutionModel::GLCompute; ctx.AddExecutionMode(main, spv::ExecutionMode::LocalSize, workgroup_size[0], workgroup_size[1], workgroup_size[2]); @@ -148,7 +148,7 @@ void DefineEntryPoint(Environment& env, const IR::Program& program, EmitContext& } break; default: - throw NotImplementedException("Stage {}", env.ShaderStage()); + throw NotImplementedException("Stage {}", program.stage); } ctx.AddEntryPoint(execution_model, main, "main", interfaces); } @@ -267,11 +267,10 @@ Id PhiArgDef(EmitContext& ctx, IR::Inst* inst, size_t index) { } } // Anonymous namespace -std::vector EmitSPIRV(const Profile& profile, Environment& env, IR::Program& program, - u32& binding) { +std::vector EmitSPIRV(const Profile& profile, IR::Program& program, u32& binding) { EmitContext ctx{profile, program, binding}; const Id main{DefineMain(ctx, program)}; - DefineEntryPoint(env, program, ctx, main); + DefineEntryPoint(program, ctx, main); if (profile.support_float_controls) { ctx.AddExtension("SPV_KHR_float_controls"); SetupDenormControl(profile, program, ctx, main); -- cgit v1.2.3 From 514a6b07eedace58b4a0c95282bdfc729623d1d9 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Tue, 30 Mar 2021 03:19:50 -0300 Subject: shader: Store type of phi nodes in flags This is needed because pseudo-instructions where invalidated. --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 2e7e6bb0c..6389d80bf 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -288,7 +288,8 @@ Id EmitPhi(EmitContext& ctx, IR::Inst* inst) { operands.push_back(PhiArgDef(ctx, inst, index)); operands.push_back(inst->PhiBlock(index)->Definition()); } - const Id result_type{TypeId(ctx, inst->Arg(0).Type())}; + // The type of a phi instruction is stored in its flags + const Id result_type{TypeId(ctx, inst->Flags())}; return ctx.OpPhi(result_type, std::span(operands.data(), operands.size())); } -- cgit v1.2.3 From f1dd743731bd0e7b7f1ef172882971bcd15eb5bc Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Thu, 1 Apr 2021 01:07:51 -0300 Subject: shader: Fix dependency on identity removal pass --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 6389d80bf..9dc769307 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -250,7 +250,7 @@ Id PhiArgDef(EmitContext& ctx, IR::Inst* inst, size_t index) { // Let the context handle immediate definitions, as it already knows how return ctx.Def(arg); } - IR::Inst* const arg_inst{arg.Inst()}; + IR::Inst* const arg_inst{arg.InstRecursive()}; if (const Id def{arg_inst->Definition()}; Sirit::ValidId(def)) { // Return the current definition if it exists return def; @@ -296,7 +296,12 @@ Id EmitPhi(EmitContext& ctx, IR::Inst* inst) { void EmitVoid(EmitContext&) {} Id EmitIdentity(EmitContext& ctx, const IR::Value& value) { - return ctx.Def(value); + if (const Id id = ctx.Def(value); Sirit::ValidId(id)) { + return id; + } + const Id def{ctx.ForwardDeclarationId()}; + value.InstRecursive()->SetDefinition(def); + return def; } void EmitGetZeroFromOp(EmitContext&) { -- cgit v1.2.3 From b4a5e767d0a60d44c77460bd3a4062c5f69fb6c7 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Fri, 2 Apr 2021 01:17:47 -0300 Subject: shader: Fix branches to visited virtual blocks --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 2 ++ 1 file changed, 2 insertions(+) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 9dc769307..b8e3b8527 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -56,6 +56,8 @@ ArgType Arg(EmitContext& ctx, const IR::Value& arg) { return arg.Label(); } else if constexpr (std::is_same_v) { return arg.Attribute(); + } else if constexpr (std::is_same_v) { + return arg.Reg(); } } -- cgit v1.2.3 From d819ba4489b90955286341c739083e638173b938 Mon Sep 17 00:00:00 2001 From: FernandoS27 Date: Thu, 1 Apr 2021 08:34:45 +0200 Subject: shader: Implement ViewportIndex --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 10 ++++++++++ 1 file changed, 10 insertions(+) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index b8e3b8527..cc6b98f7e 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -225,6 +225,16 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct ctx.AddExtension("SPV_EXT_demote_to_helper_invocation"); ctx.AddCapability(spv::Capability::DemoteToHelperInvocationEXT); } + if (info.stores_viewport_index) { + ctx.AddCapability(spv::Capability::MultiViewport); + if (profile.support_viewport_index_layer_non_geometry && + ctx.stage == Shader::Stage::VertexB) { + ctx.AddExtension("SPV_EXT_shader_viewport_index_layer"); + ctx.AddCapability(spv::Capability::ShaderViewportIndexLayerEXT); + } else { + ctx.ignore_viewport_layer = true; + } + } if (!profile.support_vertex_instance_id && (info.loads_instance_id || info.loads_vertex_id)) { ctx.AddExtension("SPV_KHR_shader_draw_parameters"); ctx.AddCapability(spv::Capability::DrawParameters); -- cgit v1.2.3 From baec84247fe815199595d9e8077b71f3b5c8317e Mon Sep 17 00:00:00 2001 From: FernandoS27 Date: Sat, 3 Apr 2021 01:48:39 +0200 Subject: shader: Address Feedback --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index cc6b98f7e..191380db0 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -228,11 +228,9 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct if (info.stores_viewport_index) { ctx.AddCapability(spv::Capability::MultiViewport); if (profile.support_viewport_index_layer_non_geometry && - ctx.stage == Shader::Stage::VertexB) { + ctx.stage != Shader::Stage::Geometry) { ctx.AddExtension("SPV_EXT_shader_viewport_index_layer"); ctx.AddCapability(spv::Capability::ShaderViewportIndexLayerEXT); - } else { - ctx.ignore_viewport_layer = true; } } if (!profile.support_vertex_instance_id && (info.loads_instance_id || info.loads_vertex_id)) { -- cgit v1.2.3 From 1f3eb601acdcdfa4c119cffbf36b5792147b893f Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Tue, 6 Apr 2021 02:56:15 -0300 Subject: shader: Implement texture buffers --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 1 + 1 file changed, 1 insertion(+) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 191380db0..32512a0e5 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -249,6 +249,7 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct // TODO: Track this usage ctx.AddCapability(spv::Capability::ImageGatherExtended); ctx.AddCapability(spv::Capability::ImageQuery); + ctx.AddCapability(spv::Capability::SampledBuffer); } Id PhiArgDef(EmitContext& ctx, IR::Inst* inst, size_t index) { -- cgit v1.2.3 From 0bb85f6a753c769266c95c4ba146b25b9eaaaffd Mon Sep 17 00:00:00 2001 From: lat9nq <22451773+lat9nq@users.noreply.github.com> Date: Mon, 5 Apr 2021 22:25:22 -0400 Subject: shader_recompiler,video_core: Cleanup some GCC and Clang errors Mostly fixing unused *, implicit conversion, braced scalar init, fpermissive, and some others. Some Clang errors likely remain in video_core, and std::ranges is still a pertinent issue in shader_recompiler shader_recompiler: cmake: Force bracket depth to 1024 on Clang Increases the maximum fold expression depth thread_worker: Include condition_variable Don't use list initializers in control flow Co-authored-by: ReinUsesLisp --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 19 +++++++++++-------- 1 file changed, 11 insertions(+), 8 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 32512a0e5..355cf0ca8 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -16,7 +16,7 @@ namespace Shader::Backend::SPIRV { namespace { template -struct FuncTraits : FuncTraits {}; +struct FuncTraits {}; template struct FuncTraits { @@ -64,17 +64,20 @@ ArgType Arg(EmitContext& ctx, const IR::Value& arg) { template void Invoke(EmitContext& ctx, IR::Inst* inst, std::index_sequence) { using Traits = FuncTraits; - if constexpr (std::is_same_v) { + if constexpr (std::is_same_v) { if constexpr (is_first_arg_inst) { - SetDefinition(ctx, inst, inst, Arg>(ctx, inst->Arg(I))...); + SetDefinition( + ctx, inst, inst, + Arg>(ctx, inst->Arg(I))...); } else { - SetDefinition(ctx, inst, Arg>(ctx, inst->Arg(I))...); + SetDefinition( + ctx, inst, Arg>(ctx, inst->Arg(I))...); } } else { if constexpr (is_first_arg_inst) { - func(ctx, inst, Arg>(ctx, inst->Arg(I))...); + func(ctx, inst, Arg>(ctx, inst->Arg(I))...); } else { - func(ctx, Arg>(ctx, inst->Arg(I))...); + func(ctx, Arg>(ctx, inst->Arg(I))...); } } } @@ -94,14 +97,14 @@ void Invoke(EmitContext& ctx, IR::Inst* inst) { } void EmitInst(EmitContext& ctx, IR::Inst* inst) { - switch (inst->Opcode()) { + switch (inst->GetOpcode()) { #define OPCODE(name, result_type, ...) \ case IR::Opcode::name: \ return Invoke<&Emit##name>(ctx, inst); #include "shader_recompiler/frontend/ir/opcodes.inc" #undef OPCODE } - throw LogicError("Invalid opcode {}", inst->Opcode()); + throw LogicError("Invalid opcode {}", inst->GetOpcode()); } Id TypeId(const EmitContext& ctx, IR::Type type) { -- cgit v1.2.3 From 7cb2ab358517d95ebcd35c94c72b9e91762906c3 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Fri, 9 Apr 2021 01:45:39 -0300 Subject: shader: Implement SULD and SUST --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 1 + 1 file changed, 1 insertion(+) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 355cf0ca8..ecd0fac5c 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -253,6 +253,7 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct ctx.AddCapability(spv::Capability::ImageGatherExtended); ctx.AddCapability(spv::Capability::ImageQuery); ctx.AddCapability(spv::Capability::SampledBuffer); + ctx.AddCapability(spv::Capability::StorageImageReadWithoutFormat); } Id PhiArgDef(EmitContext& ctx, IR::Inst* inst, size_t index) { -- cgit v1.2.3 From 2e71e4c5c02d133e3e85597b7eb52b88084a31fe Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Sun, 11 Apr 2021 02:08:15 -0300 Subject: spirv: Fix forward declarations on phi nodes --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 72 ++++++++-------------- 1 file changed, 25 insertions(+), 47 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index ecd0fac5c..63ed92a5d 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -30,18 +30,7 @@ struct FuncTraits { template void SetDefinition(EmitContext& ctx, IR::Inst* inst, Args... args) { - const Id forward_id{inst->Definition()}; - const bool has_forward_id{Sirit::ValidId(forward_id)}; - Id current_id{}; - if (has_forward_id) { - current_id = ctx.ExchangeCurrentId(forward_id); - } - const Id new_id{func(ctx, std::forward(args)...)}; - if (has_forward_id) { - ctx.ExchangeCurrentId(current_id); - } else { - inst->SetDefinition(new_id); - } + inst->SetDefinition(func(ctx, std::forward(args)...)); } template @@ -255,31 +244,6 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct ctx.AddCapability(spv::Capability::SampledBuffer); ctx.AddCapability(spv::Capability::StorageImageReadWithoutFormat); } - -Id PhiArgDef(EmitContext& ctx, IR::Inst* inst, size_t index) { - // Phi nodes can have forward declarations, if an argument is not defined provide a forward - // declaration of it. Invoke will take care of giving it the right definition when it's - // actually defined. - const IR::Value arg{inst->Arg(index)}; - if (arg.IsImmediate()) { - // Let the context handle immediate definitions, as it already knows how - return ctx.Def(arg); - } - IR::Inst* const arg_inst{arg.InstRecursive()}; - if (const Id def{arg_inst->Definition()}; Sirit::ValidId(def)) { - // Return the current definition if it exists - return def; - } - if (arg_inst == inst) { - // This is a self referencing phi node - // Self-referencing definition will be set by the caller, so just grab the current id - return ctx.CurrentId(); - } - // If it hasn't been defined and it's not a self reference, get a forward declaration - const Id def{ctx.ForwardDeclarationId()}; - arg_inst->SetDefinition(def); - return def; -} } // Anonymous namespace std::vector EmitSPIRV(const Profile& profile, IR::Program& program, u32& binding) { @@ -292,31 +256,45 @@ std::vector EmitSPIRV(const Profile& profile, IR::Program& program, u32& bi SetupSignedNanCapabilities(profile, program, ctx, main); } SetupCapabilities(profile, program.info, ctx); + + auto inst{program.blocks.front()->begin()}; + size_t block_index{}; + ctx.PatchDeferredPhi([&](size_t phi_arg) { + if (phi_arg == 0) { + ++inst; + if (inst == program.blocks[block_index]->end() || + inst->GetOpcode() != IR::Opcode::Phi) { + do { + ++block_index; + inst = program.blocks[block_index]->begin(); + } while (inst->GetOpcode() != IR::Opcode::Phi); + } + } + return ctx.Def(inst->Arg(phi_arg)); + }); return ctx.Assemble(); } Id EmitPhi(EmitContext& ctx, IR::Inst* inst) { const size_t num_args{inst->NumArgs()}; - boost::container::small_vector operands; - operands.reserve(num_args * 2); + boost::container::small_vector blocks; + blocks.reserve(num_args); for (size_t index = 0; index < num_args; ++index) { - operands.push_back(PhiArgDef(ctx, inst, index)); - operands.push_back(inst->PhiBlock(index)->Definition()); + blocks.push_back(inst->PhiBlock(index)->Definition()); } // The type of a phi instruction is stored in its flags const Id result_type{TypeId(ctx, inst->Flags())}; - return ctx.OpPhi(result_type, std::span(operands.data(), operands.size())); + return ctx.DeferredOpPhi(result_type, std::span(blocks.data(), blocks.size())); } void EmitVoid(EmitContext&) {} Id EmitIdentity(EmitContext& ctx, const IR::Value& value) { - if (const Id id = ctx.Def(value); Sirit::ValidId(id)) { - return id; + const Id id{ctx.Def(value)}; + if (!Sirit::ValidId(id)) { + throw NotImplementedException("Forward identity declaration"); } - const Id def{ctx.ForwardDeclarationId()}; - value.InstRecursive()->SetDefinition(def); - return def; + return id; } void EmitGetZeroFromOp(EmitContext&) { -- cgit v1.2.3 From ab543f18213133b3076b81f30df386d5cb470e49 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Sun, 11 Apr 2021 02:37:03 -0300 Subject: spirv: Guard against typeless image reads on unsupported devices --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 63ed92a5d..db7b3f1b2 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -238,11 +238,13 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct ctx.AddCapability(spv::Capability::SubgroupVoteKHR); } } + if (info.uses_typeless_image_reads && profile.support_typeless_image_loads) { + ctx.AddCapability(spv::Capability::StorageImageReadWithoutFormat); + } // TODO: Track this usage ctx.AddCapability(spv::Capability::ImageGatherExtended); ctx.AddCapability(spv::Capability::ImageQuery); ctx.AddCapability(spv::Capability::SampledBuffer); - ctx.AddCapability(spv::Capability::StorageImageReadWithoutFormat); } } // Anonymous namespace -- cgit v1.2.3 From 106764a6d51566b050dfe7124f82cf9a4de468c1 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Sun, 11 Apr 2021 02:46:51 -0300 Subject: spirv: Move phi node patching to a separate function --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 29 ++++++++++++---------- 1 file changed, 16 insertions(+), 13 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index db7b3f1b2..5a1ffd61c 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -246,21 +246,10 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct ctx.AddCapability(spv::Capability::ImageQuery); ctx.AddCapability(spv::Capability::SampledBuffer); } -} // Anonymous namespace - -std::vector EmitSPIRV(const Profile& profile, IR::Program& program, u32& binding) { - EmitContext ctx{profile, program, binding}; - const Id main{DefineMain(ctx, program)}; - DefineEntryPoint(program, ctx, main); - if (profile.support_float_controls) { - ctx.AddExtension("SPV_KHR_float_controls"); - SetupDenormControl(profile, program, ctx, main); - SetupSignedNanCapabilities(profile, program, ctx, main); - } - SetupCapabilities(profile, program.info, ctx); +void PatchPhiNodes(IR::Program& program, EmitContext& ctx) { auto inst{program.blocks.front()->begin()}; - size_t block_index{}; + size_t block_index{0}; ctx.PatchDeferredPhi([&](size_t phi_arg) { if (phi_arg == 0) { ++inst; @@ -274,6 +263,20 @@ std::vector EmitSPIRV(const Profile& profile, IR::Program& program, u32& bi } return ctx.Def(inst->Arg(phi_arg)); }); +} +} // Anonymous namespace + +std::vector EmitSPIRV(const Profile& profile, IR::Program& program, u32& binding) { + EmitContext ctx{profile, program, binding}; + const Id main{DefineMain(ctx, program)}; + DefineEntryPoint(program, ctx, main); + if (profile.support_float_controls) { + ctx.AddExtension("SPV_KHR_float_controls"); + SetupDenormControl(profile, program, ctx, main); + SetupSignedNanCapabilities(profile, program, ctx, main); + } + SetupCapabilities(profile, program.info, ctx); + PatchPhiNodes(program, ctx); return ctx.Assemble(); } -- cgit v1.2.3 From 3db2b3effa953ae66457b7a19b419fc4db2c4801 Mon Sep 17 00:00:00 2001 From: ameerj <52414509+ameerj@users.noreply.github.com> Date: Sun, 11 Apr 2021 02:07:02 -0400 Subject: shader: Implement ATOM/S and RED --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 3 +++ 1 file changed, 3 insertions(+) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 5a1ffd61c..9248bd78b 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -238,6 +238,9 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct ctx.AddCapability(spv::Capability::SubgroupVoteKHR); } } + if (info.uses_64_bit_atomics && profile.support_int64_atomics) { + ctx.AddCapability(spv::Capability::Int64Atomics); + } if (info.uses_typeless_image_reads && profile.support_typeless_image_loads) { ctx.AddCapability(spv::Capability::StorageImageReadWithoutFormat); } -- cgit v1.2.3 From 23b87147321d02abf47868f231f00f29b0d3b87d Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Sun, 11 Apr 2021 21:02:44 -0300 Subject: spirv: Define StorageImageWriteWithoutFormat capability when used --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 3 +++ 1 file changed, 3 insertions(+) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 9248bd78b..3258b0cf8 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -244,6 +244,9 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct if (info.uses_typeless_image_reads && profile.support_typeless_image_loads) { ctx.AddCapability(spv::Capability::StorageImageReadWithoutFormat); } + if (info.uses_typeless_image_writes) { + ctx.AddCapability(spv::Capability::StorageImageWriteWithoutFormat); + } // TODO: Track this usage ctx.AddCapability(spv::Capability::ImageGatherExtended); ctx.AddCapability(spv::Capability::ImageQuery); -- cgit v1.2.3 From f263760c5a3aff771123b32b15677e1f7a089640 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Mon, 12 Apr 2021 19:41:22 -0300 Subject: shader: Implement geometry shaders --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 38 ++++++++++++++++++++++ 1 file changed, 38 insertions(+) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 3258b0cf8..d7c5890ab 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -134,6 +134,44 @@ void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) { case Shader::Stage::VertexB: execution_model = spv::ExecutionModel::Vertex; break; + case Shader::Stage::Geometry: + execution_model = spv::ExecutionModel::Geometry; + ctx.AddCapability(spv::Capability::Geometry); + ctx.AddCapability(spv::Capability::GeometryStreams); + switch (ctx.profile.input_topology) { + case InputTopology::Points: + ctx.AddExecutionMode(main, spv::ExecutionMode::InputPoints); + break; + case InputTopology::Lines: + ctx.AddExecutionMode(main, spv::ExecutionMode::InputLines); + break; + case InputTopology::LinesAdjacency: + ctx.AddExecutionMode(main, spv::ExecutionMode::InputLinesAdjacency); + break; + case InputTopology::Triangles: + ctx.AddExecutionMode(main, spv::ExecutionMode::Triangles); + break; + case InputTopology::TrianglesAdjacency: + ctx.AddExecutionMode(main, spv::ExecutionMode::InputTrianglesAdjacency); + break; + } + switch (program.output_topology) { + case OutputTopology::PointList: + ctx.AddExecutionMode(main, spv::ExecutionMode::OutputPoints); + break; + case OutputTopology::LineStrip: + ctx.AddExecutionMode(main, spv::ExecutionMode::OutputLineStrip); + break; + case OutputTopology::TriangleStrip: + ctx.AddExecutionMode(main, spv::ExecutionMode::OutputTriangleStrip); + break; + } + if (program.info.stores_point_size) { + ctx.AddCapability(spv::Capability::GeometryPointSize); + } + ctx.AddExecutionMode(main, spv::ExecutionMode::OutputVertices, program.output_vertices); + ctx.AddExecutionMode(main, spv::ExecutionMode::Invocations, program.invocations); + break; case Shader::Stage::Fragment: execution_model = spv::ExecutionModel::Fragment; ctx.AddExecutionMode(main, spv::ExecutionMode::OriginUpperLeft); -- cgit v1.2.3 From fa75b9b0626c8e118e27207dd1e82e2f415fc0bc Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Tue, 13 Apr 2021 05:32:21 -0300 Subject: spirv: Rework storage buffers and shader memory --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index d7c5890ab..61a2018d7 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -276,7 +276,7 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct ctx.AddCapability(spv::Capability::SubgroupVoteKHR); } } - if (info.uses_64_bit_atomics && profile.support_int64_atomics) { + if (info.uses_int64_bit_atomics && profile.support_int64_atomics) { ctx.AddCapability(spv::Capability::Int64Atomics); } if (info.uses_typeless_image_reads && profile.support_typeless_image_loads) { -- cgit v1.2.3 From a83579b50a167ab9483e5058fd1c748018ef6d7c Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Tue, 13 Apr 2021 16:56:22 -0300 Subject: shader: Implement early Z tests --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 3 +++ 1 file changed, 3 insertions(+) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 61a2018d7..7ad00c434 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -178,6 +178,9 @@ void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) { if (program.info.stores_frag_depth) { ctx.AddExecutionMode(main, spv::ExecutionMode::DepthReplacing); } + if (ctx.profile.force_early_z) { + ctx.AddExecutionMode(main, spv::ExecutionMode::EarlyFragmentTests); + } break; default: throw NotImplementedException("Stage {}", program.stage); -- cgit v1.2.3 From b126987c59964d81ae3705ad7ad6c0ace8714e19 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Wed, 14 Apr 2021 01:04:59 -0300 Subject: shader: Implement transform feedbacks and define file format --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 3 +++ 1 file changed, 3 insertions(+) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 7ad00c434..444ba276f 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -288,6 +288,9 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct if (info.uses_typeless_image_writes) { ctx.AddCapability(spv::Capability::StorageImageWriteWithoutFormat); } + if (!ctx.profile.xfb_varyings.empty()) { + ctx.AddCapability(spv::Capability::TransformFeedback); + } // TODO: Track this usage ctx.AddCapability(spv::Capability::ImageGatherExtended); ctx.AddCapability(spv::Capability::ImageQuery); -- cgit v1.2.3 From d8ec99dadaa033aa440671572ed38e2614815e11 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Wed, 14 Apr 2021 18:09:18 -0300 Subject: spirv: Implement Layer stores --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 16 ++++++++++------ 1 file changed, 10 insertions(+), 6 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 444ba276f..3bf4c6a9e 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -124,17 +124,17 @@ void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) { const std::span interfaces(ctx.interfaces.data(), ctx.interfaces.size()); spv::ExecutionModel execution_model{}; switch (program.stage) { - case Shader::Stage::Compute: { + case Stage::Compute: { const std::array workgroup_size{program.workgroup_size}; execution_model = spv::ExecutionModel::GLCompute; ctx.AddExecutionMode(main, spv::ExecutionMode::LocalSize, workgroup_size[0], workgroup_size[1], workgroup_size[2]); break; } - case Shader::Stage::VertexB: + case Stage::VertexB: execution_model = spv::ExecutionModel::Vertex; break; - case Shader::Stage::Geometry: + case Stage::Geometry: execution_model = spv::ExecutionModel::Geometry; ctx.AddCapability(spv::Capability::Geometry); ctx.AddCapability(spv::Capability::GeometryStreams); @@ -172,7 +172,7 @@ void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) { ctx.AddExecutionMode(main, spv::ExecutionMode::OutputVertices, program.output_vertices); ctx.AddExecutionMode(main, spv::ExecutionMode::Invocations, program.invocations); break; - case Shader::Stage::Fragment: + case Stage::Fragment: execution_model = spv::ExecutionModel::Fragment; ctx.AddExecutionMode(main, spv::ExecutionMode::OriginUpperLeft); if (program.info.stores_frag_depth) { @@ -258,10 +258,14 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct ctx.AddExtension("SPV_EXT_demote_to_helper_invocation"); ctx.AddCapability(spv::Capability::DemoteToHelperInvocationEXT); } + if (info.stores_layer) { + ctx.AddCapability(spv::Capability::ShaderLayer); + } if (info.stores_viewport_index) { ctx.AddCapability(spv::Capability::MultiViewport); - if (profile.support_viewport_index_layer_non_geometry && - ctx.stage != Shader::Stage::Geometry) { + } + if (info.stores_layer || info.stores_viewport_index) { + if (profile.support_viewport_index_layer_non_geometry && ctx.stage != Stage::Geometry) { ctx.AddExtension("SPV_EXT_shader_viewport_index_layer"); ctx.AddCapability(spv::Capability::ShaderViewportIndexLayerEXT); } -- cgit v1.2.3 From 183855e396cc6918d36fbf3e38ea426e934b4e3e Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Thu, 15 Apr 2021 22:46:11 -0300 Subject: shader: Implement tessellation shaders, polygon mode and invocation id --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 39 ++++++++++++++++++++++ 1 file changed, 39 insertions(+) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 3bf4c6a9e..105602ccf 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -45,6 +45,8 @@ ArgType Arg(EmitContext& ctx, const IR::Value& arg) { return arg.Label(); } else if constexpr (std::is_same_v) { return arg.Attribute(); + } else if constexpr (std::is_same_v) { + return arg.Patch(); } else if constexpr (std::is_same_v) { return arg.Reg(); } @@ -120,6 +122,30 @@ Id DefineMain(EmitContext& ctx, IR::Program& program) { return main; } +spv::ExecutionMode ExecutionMode(TessPrimitive primitive) { + switch (primitive) { + case TessPrimitive::Isolines: + return spv::ExecutionMode::Isolines; + case TessPrimitive::Triangles: + return spv::ExecutionMode::Triangles; + case TessPrimitive::Quads: + return spv::ExecutionMode::Quads; + } + throw InvalidArgument("Tessellation primitive {}", primitive); +} + +spv::ExecutionMode ExecutionMode(TessSpacing spacing) { + switch (spacing) { + case TessSpacing::Equal: + return spv::ExecutionMode::SpacingEqual; + case TessSpacing::FractionalOdd: + return spv::ExecutionMode::SpacingFractionalOdd; + case TessSpacing::FractionalEven: + return spv::ExecutionMode::SpacingFractionalEven; + } + throw InvalidArgument("Tessellation spacing {}", spacing); +} + void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) { const std::span interfaces(ctx.interfaces.data(), ctx.interfaces.size()); spv::ExecutionModel execution_model{}; @@ -134,6 +160,19 @@ void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) { case Stage::VertexB: execution_model = spv::ExecutionModel::Vertex; break; + case Stage::TessellationControl: + execution_model = spv::ExecutionModel::TessellationControl; + ctx.AddCapability(spv::Capability::Tessellation); + ctx.AddExecutionMode(main, spv::ExecutionMode::OutputVertices, program.invocations); + break; + case Stage::TessellationEval: + execution_model = spv::ExecutionModel::TessellationEvaluation; + ctx.AddCapability(spv::Capability::Tessellation); + ctx.AddExecutionMode(main, ExecutionMode(ctx.profile.tess_primitive)); + ctx.AddExecutionMode(main, ExecutionMode(ctx.profile.tess_spacing)); + ctx.AddExecutionMode(main, ctx.profile.tess_clockwise ? spv::ExecutionMode::VertexOrderCw + : spv::ExecutionMode::VertexOrderCcw); + break; case Stage::Geometry: execution_model = spv::ExecutionModel::Geometry; ctx.AddCapability(spv::Capability::Geometry); -- cgit v1.2.3 From e3514bcd6b09f623da14c4f3c4ffd988e75577ed Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Fri, 16 Apr 2021 16:31:15 -0300 Subject: spirv: Implement ViewportMask with NV_viewport_array2 --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 4 ++++ 1 file changed, 4 insertions(+) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 105602ccf..90c4833a8 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -303,6 +303,10 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct if (info.stores_viewport_index) { ctx.AddCapability(spv::Capability::MultiViewport); } + if (info.stores_viewport_mask && profile.support_viewport_mask) { + ctx.AddExtension("SPV_NV_viewport_array2"); + ctx.AddCapability(spv::Capability::ShaderViewportMaskNV); + } if (info.stores_layer || info.stores_viewport_index) { if (profile.support_viewport_index_layer_non_geometry && ctx.stage != Stage::Geometry) { ctx.AddExtension("SPV_EXT_shader_viewport_index_layer"); -- cgit v1.2.3 From 95815a3883d708f71db5119f42243e183f32f9a2 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Fri, 16 Apr 2021 17:22:59 -0300 Subject: shader: Implement PIXLD.MY_INDEX --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 3 +++ 1 file changed, 3 insertions(+) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 90c4833a8..9ec970706 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -335,6 +335,9 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct if (info.uses_typeless_image_writes) { ctx.AddCapability(spv::Capability::StorageImageWriteWithoutFormat); } + if (info.uses_sample_id) { + ctx.AddCapability(spv::Capability::SampleRateShading); + } if (!ctx.profile.xfb_varyings.empty()) { ctx.AddCapability(spv::Capability::TransformFeedback); } -- cgit v1.2.3 From 04c459fc8d99b41fa8a03c49523599e9bf797f9d Mon Sep 17 00:00:00 2001 From: FernandoS27 Date: Sat, 17 Apr 2021 11:56:45 +0200 Subject: shader: Implement fine derivates constant propagation --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 3 +++ 1 file changed, 3 insertions(+) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 9ec970706..c4d5874ca 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -341,6 +341,9 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct if (!ctx.profile.xfb_varyings.empty()) { ctx.AddCapability(spv::Capability::TransformFeedback); } + if (info.uses_derivates) { + ctx.AddCapability(spv::Capability::DerivativeControl); + } // TODO: Track this usage ctx.AddCapability(spv::Capability::ImageGatherExtended); ctx.AddCapability(spv::Capability::ImageQuery); -- cgit v1.2.3 From f69d0b91ffad7d9ab827f55a9297b8f6da815cc9 Mon Sep 17 00:00:00 2001 From: FernandoS27 Date: Sun, 18 Apr 2021 09:07:48 +0200 Subject: shader: Address feedback --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index c4d5874ca..5d6fdeb65 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -341,7 +341,7 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct if (!ctx.profile.xfb_varyings.empty()) { ctx.AddCapability(spv::Capability::TransformFeedback); } - if (info.uses_derivates) { + if (info.uses_derivatives) { ctx.AddCapability(spv::Capability::DerivativeControl); } // TODO: Track this usage -- cgit v1.2.3 From 050e81500c002f304d581f28700de549b828a2bc Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Wed, 21 Apr 2021 00:35:47 -0300 Subject: shader: Move microinstruction header to the value header --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 1 - 1 file changed, 1 deletion(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 5d6fdeb65..815b3cd95 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -10,7 +10,6 @@ #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/frontend/ir/basic_block.h" -#include "shader_recompiler/frontend/ir/microinstruction.h" #include "shader_recompiler/frontend/ir/program.h" namespace Shader::Backend::SPIRV { -- cgit v1.2.3 From 48a17298d76cd8ed3bf2b53aca1e1ac097693976 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Sun, 23 May 2021 03:58:11 -0300 Subject: spirv: Support OpenGL uniform buffers and change bindings --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 815b3cd95..0cb075670 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -368,7 +368,7 @@ void PatchPhiNodes(IR::Program& program, EmitContext& ctx) { } } // Anonymous namespace -std::vector EmitSPIRV(const Profile& profile, IR::Program& program, u32& binding) { +std::vector EmitSPIRV(const Profile& profile, IR::Program& program, Bindings& binding) { EmitContext ctx{profile, program, binding}; const Id main{DefineMain(ctx, program)}; DefineEntryPoint(program, ctx, main); -- cgit v1.2.3 From cfd873275d705f124efff6ceae33efc8994e64fa Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Sun, 23 May 2021 04:18:22 -0300 Subject: spirv: Use OriginLowerLeft when requested --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 0cb075670..10de612cd 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -212,7 +212,11 @@ void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) { break; case Stage::Fragment: execution_model = spv::ExecutionModel::Fragment; - ctx.AddExecutionMode(main, spv::ExecutionMode::OriginUpperLeft); + if (ctx.profile.lower_left_origin_mode) { + ctx.AddExecutionMode(main, spv::ExecutionMode::OriginLowerLeft); + } else { + ctx.AddExecutionMode(main, spv::ExecutionMode::OriginUpperLeft); + } if (program.info.stores_frag_depth) { ctx.AddExecutionMode(main, spv::ExecutionMode::DepthReplacing); } -- cgit v1.2.3 From 2b434b74af7be8c6ed2e96970e4e3965e351edbd Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Sun, 23 May 2021 04:18:55 -0300 Subject: spirv: Enable DemoteToHelperInvocationEXT only when supported --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 10de612cd..eb192e3c9 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -296,7 +296,7 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct if (info.uses_sparse_residency) { ctx.AddCapability(spv::Capability::SparseResidency); } - if (info.uses_demote_to_helper_invocation) { + if (info.uses_demote_to_helper_invocation && profile.support_demote_to_helper_invocation) { ctx.AddExtension("SPV_EXT_demote_to_helper_invocation"); ctx.AddCapability(spv::Capability::DemoteToHelperInvocationEXT); } -- cgit v1.2.3 From d2a0f9d7ad89184294e4d3f05ae0843e4ff4e6be Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Sun, 23 May 2021 04:19:27 -0300 Subject: spirv: Do not enable ShaderLayer This is enabled by an extension instead of the capability. --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 3 --- 1 file changed, 3 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index eb192e3c9..745a834e3 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -300,9 +300,6 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct ctx.AddExtension("SPV_EXT_demote_to_helper_invocation"); ctx.AddCapability(spv::Capability::DemoteToHelperInvocationEXT); } - if (info.stores_layer) { - ctx.AddCapability(spv::Capability::ShaderLayer); - } if (info.stores_viewport_index) { ctx.AddCapability(spv::Capability::MultiViewport); } -- cgit v1.2.3 From 7ecc6de56ae01602b25408db8b6658d7a41a419a Mon Sep 17 00:00:00 2001 From: ameerj <52414509+ameerj@users.noreply.github.com> Date: Fri, 23 Apr 2021 17:47:54 -0400 Subject: shader: Implement Int32 SUATOM/SURED --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 3 +++ 1 file changed, 3 insertions(+) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 745a834e3..3f9adc902 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -335,6 +335,9 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct if (info.uses_typeless_image_writes) { ctx.AddCapability(spv::Capability::StorageImageWriteWithoutFormat); } + if (info.uses_image_buffers) { + ctx.AddCapability(spv::Capability::ImageBuffer); + } if (info.uses_sample_id) { ctx.AddCapability(spv::Capability::SampleRateShading); } -- cgit v1.2.3 From bed090807afd3364ed6ef18a031a0ffd95a1b89b Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Mon, 3 May 2021 20:53:00 -0300 Subject: Move SPIR-V emission functions to their own header --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 1 + 1 file changed, 1 insertion(+) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 3f9adc902..0681dfd16 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -9,6 +9,7 @@ #include #include "shader_recompiler/backend/spirv/emit_spirv.h" +#include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" #include "shader_recompiler/frontend/ir/basic_block.h" #include "shader_recompiler/frontend/ir/program.h" -- cgit v1.2.3 From c4fd6b55bc9acd06b2fc89f84fd175d78e14110a Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Mon, 10 May 2021 18:21:28 -0300 Subject: glasm: Implement shuffle and vote instructions on GLASM --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 0681dfd16..2dad87e87 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -318,7 +318,9 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct ctx.AddExtension("SPV_KHR_shader_draw_parameters"); ctx.AddCapability(spv::Capability::DrawParameters); } - if ((info.uses_subgroup_vote || info.uses_subgroup_invocation_id) && profile.support_vote) { + if ((info.uses_subgroup_vote || info.uses_subgroup_invocation_id || + info.uses_subgroup_shuffles) && + profile.support_vote) { ctx.AddExtension("SPV_KHR_shader_ballot"); ctx.AddCapability(spv::Capability::SubgroupBallotKHR); if (!profile.warp_size_potentially_larger_than_guest) { -- cgit v1.2.3 From d54d7de40e7295827b0e4e4026441b53d3fc9569 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Fri, 14 May 2021 00:40:54 -0300 Subject: glasm: Rework control flow introducing a syntax list This commit regresses VertexA shaders, their transformation pass has to be adapted to the new control flow. --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 71 ++++++++++++++++++++-- 1 file changed, 65 insertions(+), 6 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 2dad87e87..c22edfec2 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -41,8 +41,6 @@ ArgType Arg(EmitContext& ctx, const IR::Value& arg) { return arg; } else if constexpr (std::is_same_v) { return arg.U32(); - } else if constexpr (std::is_same_v) { - return arg.Label(); } else if constexpr (std::is_same_v) { return arg.Attribute(); } else if constexpr (std::is_same_v) { @@ -109,15 +107,74 @@ Id TypeId(const EmitContext& ctx, IR::Type type) { } } +void Traverse(EmitContext& ctx, IR::Program& program) { + IR::Block* current_block{}; + for (const IR::AbstractSyntaxNode& node : program.syntax_list) { + switch (node.type) { + case IR::AbstractSyntaxNode::Type::Block: + const Id label{node.block->Definition()}; + if (current_block) { + ctx.OpBranch(label); + } + current_block = node.block; + ctx.AddLabel(label); + for (IR::Inst& inst : node.block->Instructions()) { + EmitInst(ctx, &inst); + } + break; + case IR::AbstractSyntaxNode::Type::If: { + const Id if_label{node.if_node.body->Definition()}; + const Id endif_label{node.if_node.merge->Definition()}; + ctx.OpSelectionMerge(endif_label, spv::SelectionControlMask::MaskNone); + ctx.OpBranchConditional(ctx.Def(node.if_node.cond), if_label, endif_label); + break; + } + case IR::AbstractSyntaxNode::Type::Loop: { + const Id body_label{node.loop.body->Definition()}; + const Id continue_label{node.loop.continue_block->Definition()}; + const Id endloop_label{node.loop.merge->Definition()}; + + ctx.OpLoopMerge(endloop_label, continue_label, spv::LoopControlMask::MaskNone); + ctx.OpBranch(node.loop.body->Definition()); + break; + } + case IR::AbstractSyntaxNode::Type::Break: { + const Id break_label{node.break_node.merge->Definition()}; + const Id skip_label{node.break_node.skip->Definition()}; + ctx.OpBranchConditional(ctx.Def(node.break_node.cond), break_label, skip_label); + break; + } + case IR::AbstractSyntaxNode::Type::EndIf: + if (current_block) { + ctx.OpBranch(node.end_if.merge->Definition()); + } + break; + case IR::AbstractSyntaxNode::Type::Repeat: { + const Id loop_header_label{node.repeat.loop_header->Definition()}; + const Id merge_label{node.repeat.merge->Definition()}; + ctx.OpBranchConditional(ctx.Def(node.repeat.cond), loop_header_label, merge_label); + break; + } + case IR::AbstractSyntaxNode::Type::Return: + ctx.OpReturn(); + break; + case IR::AbstractSyntaxNode::Type::Unreachable: + ctx.OpUnreachable(); + break; + } + if (node.type != IR::AbstractSyntaxNode::Type::Block) { + current_block = nullptr; + } + } +} + Id DefineMain(EmitContext& ctx, IR::Program& program) { const Id void_function{ctx.TypeFunction(ctx.void_id)}; const Id main{ctx.OpFunction(ctx.void_id, spv::FunctionControlMask::MaskNone, void_function)}; for (IR::Block* const block : program.blocks) { - ctx.AddLabel(block->Definition()); - for (IR::Inst& inst : block->Instructions()) { - EmitInst(ctx, &inst); - } + block->SetDefinition(ctx.OpLabel()); } + Traverse(ctx, program); ctx.OpFunctionEnd(); return main; } @@ -411,6 +468,8 @@ Id EmitIdentity(EmitContext& ctx, const IR::Value& value) { return id; } +void EmitBranchConditionRef(EmitContext&) {} + void EmitGetZeroFromOp(EmitContext&) { throw LogicError("Unreachable instruction"); } -- cgit v1.2.3 From bf5e48ffe4bd48ea681f2a01c8919c97125e88df Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Fri, 14 May 2021 04:48:46 -0300 Subject: glasm: Initial implementation of phi nodes on GLASM --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index c22edfec2..7bf8c78de 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -468,7 +468,11 @@ Id EmitIdentity(EmitContext& ctx, const IR::Value& value) { return id; } -void EmitBranchConditionRef(EmitContext&) {} +void EmitDummyReference(EmitContext&) {} + +void EmitPhiMove(EmitContext&) { + throw LogicError("Unreachable instruction"); +} void EmitGetZeroFromOp(EmitContext&) { throw LogicError("Unreachable instruction"); -- cgit v1.2.3 From ab5dbe7c299d904e36099d209b7290e538a84745 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Fri, 14 May 2021 22:01:01 -0300 Subject: emit_spirv: Add missing block in case --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 7bf8c78de..266ac690c 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -111,7 +111,7 @@ void Traverse(EmitContext& ctx, IR::Program& program) { IR::Block* current_block{}; for (const IR::AbstractSyntaxNode& node : program.syntax_list) { switch (node.type) { - case IR::AbstractSyntaxNode::Type::Block: + case IR::AbstractSyntaxNode::Type::Block: { const Id label{node.block->Definition()}; if (current_block) { ctx.OpBranch(label); @@ -122,6 +122,7 @@ void Traverse(EmitContext& ctx, IR::Program& program) { EmitInst(ctx, &inst); } break; + } case IR::AbstractSyntaxNode::Type::If: { const Id if_label{node.if_node.body->Definition()}; const Id endif_label{node.if_node.merge->Definition()}; -- cgit v1.2.3 From 38e7b8c80552a453580db641c20eba9e86cdf106 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Sat, 15 May 2021 18:17:40 -0300 Subject: emit_spirv: Jump to loop body with local variable Silence unused variable warning --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 266ac690c..881a5dc4c 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -136,7 +136,7 @@ void Traverse(EmitContext& ctx, IR::Program& program) { const Id endloop_label{node.loop.merge->Definition()}; ctx.OpLoopMerge(endloop_label, continue_label, spv::LoopControlMask::MaskNone); - ctx.OpBranch(node.loop.body->Definition()); + ctx.OpBranch(body_label); break; } case IR::AbstractSyntaxNode::Type::Break: { -- cgit v1.2.3 From f7a2340205b4fa2db32403f20d7b7afe32b15f33 Mon Sep 17 00:00:00 2001 From: lat9nq <22451773+lat9nq@users.noreply.github.com> Date: Sun, 16 May 2021 17:06:13 -0400 Subject: shader_recompiler: GCC fixes Fixes members of unnamed union not being accessible, and one function without a declaration. --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 32 +++++++++++----------- 1 file changed, 16 insertions(+), 16 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 881a5dc4c..9ed2af991 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -112,48 +112,48 @@ void Traverse(EmitContext& ctx, IR::Program& program) { for (const IR::AbstractSyntaxNode& node : program.syntax_list) { switch (node.type) { case IR::AbstractSyntaxNode::Type::Block: { - const Id label{node.block->Definition()}; + const Id label{node.data.block->Definition()}; if (current_block) { ctx.OpBranch(label); } - current_block = node.block; + current_block = node.data.block; ctx.AddLabel(label); - for (IR::Inst& inst : node.block->Instructions()) { + for (IR::Inst& inst : node.data.block->Instructions()) { EmitInst(ctx, &inst); } break; } case IR::AbstractSyntaxNode::Type::If: { - const Id if_label{node.if_node.body->Definition()}; - const Id endif_label{node.if_node.merge->Definition()}; + const Id if_label{node.data.if_node.body->Definition()}; + const Id endif_label{node.data.if_node.merge->Definition()}; ctx.OpSelectionMerge(endif_label, spv::SelectionControlMask::MaskNone); - ctx.OpBranchConditional(ctx.Def(node.if_node.cond), if_label, endif_label); + ctx.OpBranchConditional(ctx.Def(node.data.if_node.cond), if_label, endif_label); break; } case IR::AbstractSyntaxNode::Type::Loop: { - const Id body_label{node.loop.body->Definition()}; - const Id continue_label{node.loop.continue_block->Definition()}; - const Id endloop_label{node.loop.merge->Definition()}; + const Id body_label{node.data.loop.body->Definition()}; + const Id continue_label{node.data.loop.continue_block->Definition()}; + const Id endloop_label{node.data.loop.merge->Definition()}; ctx.OpLoopMerge(endloop_label, continue_label, spv::LoopControlMask::MaskNone); ctx.OpBranch(body_label); break; } case IR::AbstractSyntaxNode::Type::Break: { - const Id break_label{node.break_node.merge->Definition()}; - const Id skip_label{node.break_node.skip->Definition()}; - ctx.OpBranchConditional(ctx.Def(node.break_node.cond), break_label, skip_label); + const Id break_label{node.data.break_node.merge->Definition()}; + const Id skip_label{node.data.break_node.skip->Definition()}; + ctx.OpBranchConditional(ctx.Def(node.data.break_node.cond), break_label, skip_label); break; } case IR::AbstractSyntaxNode::Type::EndIf: if (current_block) { - ctx.OpBranch(node.end_if.merge->Definition()); + ctx.OpBranch(node.data.end_if.merge->Definition()); } break; case IR::AbstractSyntaxNode::Type::Repeat: { - const Id loop_header_label{node.repeat.loop_header->Definition()}; - const Id merge_label{node.repeat.merge->Definition()}; - ctx.OpBranchConditional(ctx.Def(node.repeat.cond), loop_header_label, merge_label); + const Id loop_header_label{node.data.repeat.loop_header->Definition()}; + const Id merge_label{node.data.repeat.merge->Definition()}; + ctx.OpBranchConditional(ctx.Def(node.data.repeat.cond), loop_header_label, merge_label); break; } case IR::AbstractSyntaxNode::Type::Return: -- cgit v1.2.3 From 9bb3e008c9f4bbdd35c095b506c3a3312d17e383 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Tue, 18 May 2021 02:04:22 -0300 Subject: shader: Read branch conditions from an instruction Fixes the identity removal pass. --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 10 +++++++++- 1 file changed, 9 insertions(+), 1 deletion(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 9ed2af991..3e20ac3b9 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -469,7 +469,15 @@ Id EmitIdentity(EmitContext& ctx, const IR::Value& value) { return id; } -void EmitDummyReference(EmitContext&) {} +Id EmitConditionRef(EmitContext& ctx, const IR::Value& value) { + const Id id{ctx.Def(value)}; + if (!Sirit::ValidId(id)) { + throw NotImplementedException("Forward identity declaration"); + } + return id; +} + +void EmitReference(EmitContext&) {} void EmitPhiMove(EmitContext&) { throw LogicError("Unreachable instruction"); -- cgit v1.2.3 From 9e7b6622c25aa858b96bf0f1c7f94223a2f449a2 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Fri, 21 May 2021 02:12:32 -0300 Subject: shader: Split profile and runtime information in separate structs --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 20 +++++++++++--------- 1 file changed, 11 insertions(+), 9 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 3e20ac3b9..cba420cda 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -226,16 +226,17 @@ void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) { case Stage::TessellationEval: execution_model = spv::ExecutionModel::TessellationEvaluation; ctx.AddCapability(spv::Capability::Tessellation); - ctx.AddExecutionMode(main, ExecutionMode(ctx.profile.tess_primitive)); - ctx.AddExecutionMode(main, ExecutionMode(ctx.profile.tess_spacing)); - ctx.AddExecutionMode(main, ctx.profile.tess_clockwise ? spv::ExecutionMode::VertexOrderCw - : spv::ExecutionMode::VertexOrderCcw); + ctx.AddExecutionMode(main, ExecutionMode(ctx.runtime_info.tess_primitive)); + ctx.AddExecutionMode(main, ExecutionMode(ctx.runtime_info.tess_spacing)); + ctx.AddExecutionMode(main, ctx.runtime_info.tess_clockwise + ? spv::ExecutionMode::VertexOrderCw + : spv::ExecutionMode::VertexOrderCcw); break; case Stage::Geometry: execution_model = spv::ExecutionModel::Geometry; ctx.AddCapability(spv::Capability::Geometry); ctx.AddCapability(spv::Capability::GeometryStreams); - switch (ctx.profile.input_topology) { + switch (ctx.runtime_info.input_topology) { case InputTopology::Points: ctx.AddExecutionMode(main, spv::ExecutionMode::InputPoints); break; @@ -279,7 +280,7 @@ void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) { if (program.info.stores_frag_depth) { ctx.AddExecutionMode(main, spv::ExecutionMode::DepthReplacing); } - if (ctx.profile.force_early_z) { + if (ctx.runtime_info.force_early_z) { ctx.AddExecutionMode(main, spv::ExecutionMode::EarlyFragmentTests); } break; @@ -402,7 +403,7 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct if (info.uses_sample_id) { ctx.AddCapability(spv::Capability::SampleRateShading); } - if (!ctx.profile.xfb_varyings.empty()) { + if (!ctx.runtime_info.xfb_varyings.empty()) { ctx.AddCapability(spv::Capability::TransformFeedback); } if (info.uses_derivatives) { @@ -433,8 +434,9 @@ void PatchPhiNodes(IR::Program& program, EmitContext& ctx) { } } // Anonymous namespace -std::vector EmitSPIRV(const Profile& profile, IR::Program& program, Bindings& binding) { - EmitContext ctx{profile, program, binding}; +std::vector EmitSPIRV(const Profile& profile, const RuntimeInfo& runtime_info, + IR::Program& program, Bindings& bindings) { + EmitContext ctx{profile, runtime_info, program, bindings}; const Id main{DefineMain(ctx, program)}; DefineEntryPoint(program, ctx, main); if (profile.support_float_controls) { -- cgit v1.2.3 From 61cd7dd30128633b656ce3264da74bef1ba00bb5 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Mon, 14 Jun 2021 02:27:49 -0300 Subject: shader: Add logging --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index cba420cda..14a99750d 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -294,7 +294,7 @@ void SetupDenormControl(const Profile& profile, const IR::Program& program, Emit Id main_func) { const Info& info{program.info}; if (info.uses_fp32_denorms_flush && info.uses_fp32_denorms_preserve) { - // LOG_ERROR(HW_GPU, "Fp32 denorm flush and preserve on the same shader"); + LOG_ERROR(Shader_SPIRV, "Fp32 denorm flush and preserve on the same shader"); } else if (info.uses_fp32_denorms_flush) { if (profile.support_fp32_denorm_flush) { ctx.AddCapability(spv::Capability::DenormFlushToZero); @@ -307,7 +307,7 @@ void SetupDenormControl(const Profile& profile, const IR::Program& program, Emit ctx.AddCapability(spv::Capability::DenormPreserve); ctx.AddExecutionMode(main_func, spv::ExecutionMode::DenormPreserve, 32U); } else { - // LOG_WARNING(HW_GPU, "Fp32 denorm preserve used in shader without host support"); + LOG_WARNING(Shader_SPIRV, "Fp32 denorm preserve used in shader without host support"); } } if (!profile.support_separate_denorm_behavior) { @@ -315,7 +315,7 @@ void SetupDenormControl(const Profile& profile, const IR::Program& program, Emit return; } if (info.uses_fp16_denorms_flush && info.uses_fp16_denorms_preserve) { - // LOG_ERROR(HW_GPU, "Fp16 denorm flush and preserve on the same shader"); + LOG_ERROR(Shader_SPIRV, "Fp16 denorm flush and preserve on the same shader"); } else if (info.uses_fp16_denorms_flush) { if (profile.support_fp16_denorm_flush) { ctx.AddCapability(spv::Capability::DenormFlushToZero); @@ -328,7 +328,7 @@ void SetupDenormControl(const Profile& profile, const IR::Program& program, Emit ctx.AddCapability(spv::Capability::DenormPreserve); ctx.AddExecutionMode(main_func, spv::ExecutionMode::DenormPreserve, 16U); } else { - // LOG_WARNING(HW_GPU, "Fp16 denorm preserve used in shader without host support"); + LOG_WARNING(Shader_SPIRV, "Fp16 denorm preserve used in shader without host support"); } } } -- cgit v1.2.3 From 3d822faea1af9cab2e58fcd9edcec09940e290a4 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Tue, 15 Jun 2021 17:49:33 -0300 Subject: spirv: Reduce log severity of mismatching denorm rules --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 14a99750d..fd59b4d0a 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -294,7 +294,7 @@ void SetupDenormControl(const Profile& profile, const IR::Program& program, Emit Id main_func) { const Info& info{program.info}; if (info.uses_fp32_denorms_flush && info.uses_fp32_denorms_preserve) { - LOG_ERROR(Shader_SPIRV, "Fp32 denorm flush and preserve on the same shader"); + LOG_WARNING(Shader_SPIRV, "Fp32 denorm flush and preserve on the same shader"); } else if (info.uses_fp32_denorms_flush) { if (profile.support_fp32_denorm_flush) { ctx.AddCapability(spv::Capability::DenormFlushToZero); @@ -315,7 +315,7 @@ void SetupDenormControl(const Profile& profile, const IR::Program& program, Emit return; } if (info.uses_fp16_denorms_flush && info.uses_fp16_denorms_preserve) { - LOG_ERROR(Shader_SPIRV, "Fp16 denorm flush and preserve on the same shader"); + LOG_WARNING(Shader_SPIRV, "Fp16 denorm flush and preserve on the same shader"); } else if (info.uses_fp16_denorms_flush) { if (profile.support_fp16_denorm_flush) { ctx.AddCapability(spv::Capability::DenormFlushToZero); -- cgit v1.2.3 From 808ef97a086e7cc58a3ceded1de516ad6a6be5d3 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Mon, 21 Jun 2021 01:07:10 -0300 Subject: shader: Move loop safety tests to code emission --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 19 ++++++++++++++++++- 1 file changed, 18 insertions(+), 1 deletion(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index fd59b4d0a..278c262f8 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -8,6 +8,7 @@ #include #include +#include "common/settings.h" #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" #include "shader_recompiler/frontend/ir/basic_block.h" @@ -151,9 +152,25 @@ void Traverse(EmitContext& ctx, IR::Program& program) { } break; case IR::AbstractSyntaxNode::Type::Repeat: { + Id cond{ctx.Def(node.data.repeat.cond)}; + if (!Settings::values.disable_shader_loop_safety_checks) { + const Id pointer_type{ctx.TypePointer(spv::StorageClass::Private, ctx.U32[1])}; + const Id safety_counter{ctx.AddGlobalVariable( + pointer_type, spv::StorageClass::Private, ctx.Const(0x2000u))}; + if (ctx.profile.supported_spirv >= 0x00010400) { + ctx.interfaces.push_back(safety_counter); + } + const Id old_counter{ctx.OpLoad(ctx.U32[1], safety_counter)}; + const Id new_counter{ctx.OpISub(ctx.U32[1], old_counter, ctx.Const(1u))}; + ctx.OpStore(safety_counter, new_counter); + + const Id safety_cond{ + ctx.OpSGreaterThanEqual(ctx.U1, new_counter, ctx.u32_zero_value)}; + cond = ctx.OpLogicalAnd(ctx.U1, cond, safety_cond); + } const Id loop_header_label{node.data.repeat.loop_header->Definition()}; const Id merge_label{node.data.repeat.merge->Definition()}; - ctx.OpBranchConditional(ctx.Def(node.data.repeat.cond), loop_header_label, merge_label); + ctx.OpBranchConditional(cond, loop_header_label, merge_label); break; } case IR::AbstractSyntaxNode::Type::Return: -- cgit v1.2.3 From 7dafa96ab59892b7f1fbffdb61e4326e6443955f Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Thu, 24 Jun 2021 02:41:09 -0300 Subject: shader: Rework varyings and implement passthrough geometry shaders Put all varyings into a single std::bitset with helpers to access it. Implement passthrough geometry shaders using host's. --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 19 ++++++++++++++----- 1 file changed, 14 insertions(+), 5 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 278c262f8..ddb86d070 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -281,11 +281,19 @@ void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) { ctx.AddExecutionMode(main, spv::ExecutionMode::OutputTriangleStrip); break; } - if (program.info.stores_point_size) { + if (program.info.stores[IR::Attribute::PointSize]) { ctx.AddCapability(spv::Capability::GeometryPointSize); } ctx.AddExecutionMode(main, spv::ExecutionMode::OutputVertices, program.output_vertices); ctx.AddExecutionMode(main, spv::ExecutionMode::Invocations, program.invocations); + if (program.is_geometry_passthrough) { + if (ctx.profile.support_geometry_shader_passthrough) { + ctx.AddExtension("SPV_NV_geometry_shader_passthrough"); + ctx.AddCapability(spv::Capability::GeometryShaderPassthroughNV); + } else { + LOG_WARNING(Shader_SPIRV, "Geometry shader passthrough used with no support"); + } + } break; case Stage::Fragment: execution_model = spv::ExecutionModel::Fragment; @@ -377,20 +385,21 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct ctx.AddExtension("SPV_EXT_demote_to_helper_invocation"); ctx.AddCapability(spv::Capability::DemoteToHelperInvocationEXT); } - if (info.stores_viewport_index) { + if (info.stores[IR::Attribute::ViewportIndex]) { ctx.AddCapability(spv::Capability::MultiViewport); } - if (info.stores_viewport_mask && profile.support_viewport_mask) { + if (info.stores[IR::Attribute::ViewportMask] && profile.support_viewport_mask) { ctx.AddExtension("SPV_NV_viewport_array2"); ctx.AddCapability(spv::Capability::ShaderViewportMaskNV); } - if (info.stores_layer || info.stores_viewport_index) { + if (info.stores[IR::Attribute::Layer] || info.stores[IR::Attribute::ViewportIndex]) { if (profile.support_viewport_index_layer_non_geometry && ctx.stage != Stage::Geometry) { ctx.AddExtension("SPV_EXT_shader_viewport_index_layer"); ctx.AddCapability(spv::Capability::ShaderViewportIndexLayerEXT); } } - if (!profile.support_vertex_instance_id && (info.loads_instance_id || info.loads_vertex_id)) { + if (!profile.support_vertex_instance_id && + (info.loads[IR::Attribute::InstanceId] || info.loads[IR::Attribute::VertexId])) { ctx.AddExtension("SPV_KHR_shader_draw_parameters"); ctx.AddCapability(spv::Capability::DrawParameters); } -- cgit v1.2.3 From 8722668b3c027f0132d0be07e867247debd08d30 Mon Sep 17 00:00:00 2001 From: ReinUsesLisp Date: Tue, 29 Jun 2021 18:42:17 -0300 Subject: emit_spirv: Workaround VK_KHR_shader_float_controls on fp16 Nvidia Fix regression on Fire Emblem: Three Houses when using native fp16. --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index ddb86d070..d7a86e270 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -319,7 +319,7 @@ void SetupDenormControl(const Profile& profile, const IR::Program& program, Emit Id main_func) { const Info& info{program.info}; if (info.uses_fp32_denorms_flush && info.uses_fp32_denorms_preserve) { - LOG_WARNING(Shader_SPIRV, "Fp32 denorm flush and preserve on the same shader"); + LOG_DEBUG(Shader_SPIRV, "Fp32 denorm flush and preserve on the same shader"); } else if (info.uses_fp32_denorms_flush) { if (profile.support_fp32_denorm_flush) { ctx.AddCapability(spv::Capability::DenormFlushToZero); @@ -332,15 +332,15 @@ void SetupDenormControl(const Profile& profile, const IR::Program& program, Emit ctx.AddCapability(spv::Capability::DenormPreserve); ctx.AddExecutionMode(main_func, spv::ExecutionMode::DenormPreserve, 32U); } else { - LOG_WARNING(Shader_SPIRV, "Fp32 denorm preserve used in shader without host support"); + LOG_DEBUG(Shader_SPIRV, "Fp32 denorm preserve used in shader without host support"); } } - if (!profile.support_separate_denorm_behavior) { + if (!profile.support_separate_denorm_behavior || profile.has_broken_fp16_float_controls) { // No separate denorm behavior return; } if (info.uses_fp16_denorms_flush && info.uses_fp16_denorms_preserve) { - LOG_WARNING(Shader_SPIRV, "Fp16 denorm flush and preserve on the same shader"); + LOG_DEBUG(Shader_SPIRV, "Fp16 denorm flush and preserve on the same shader"); } else if (info.uses_fp16_denorms_flush) { if (profile.support_fp16_denorm_flush) { ctx.AddCapability(spv::Capability::DenormFlushToZero); @@ -353,13 +353,16 @@ void SetupDenormControl(const Profile& profile, const IR::Program& program, Emit ctx.AddCapability(spv::Capability::DenormPreserve); ctx.AddExecutionMode(main_func, spv::ExecutionMode::DenormPreserve, 16U); } else { - LOG_WARNING(Shader_SPIRV, "Fp16 denorm preserve used in shader without host support"); + LOG_DEBUG(Shader_SPIRV, "Fp16 denorm preserve used in shader without host support"); } } } void SetupSignedNanCapabilities(const Profile& profile, const IR::Program& program, EmitContext& ctx, Id main_func) { + if (profile.has_broken_fp16_float_controls && program.info.uses_fp16) { + return; + } if (program.info.uses_fp16 && profile.support_fp16_signed_zero_nan_preserve) { ctx.AddCapability(spv::Capability::SignedZeroInfNanPreserve); ctx.AddExecutionMode(main_func, spv::ExecutionMode::SignedZeroInfNanPreserve, 16U); -- cgit v1.2.3