diff options
Diffstat (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp')
| -rw-r--r-- | src/shader_recompiler/backend/spirv/emit_spirv.cpp | 117 |
1 files changed, 57 insertions, 60 deletions
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 @@ | |||
| 2 | // Licensed under GPLv2 or any later version | 2 | // Licensed under GPLv2 or any later version |
| 3 | // Refer to the license.txt file included. | 3 | // Refer to the license.txt file included. |
| 4 | 4 | ||
| 5 | #include <numeric> | 5 | #include <span> |
| 6 | #include <tuple> | ||
| 6 | #include <type_traits> | 7 | #include <type_traits> |
| 8 | #include <utility> | ||
| 9 | #include <vector> | ||
| 7 | 10 | ||
| 8 | #include "shader_recompiler/backend/spirv/emit_spirv.h" | 11 | #include "shader_recompiler/backend/spirv/emit_spirv.h" |
| 9 | #include "shader_recompiler/frontend/ir/basic_block.h" | 12 | #include "shader_recompiler/frontend/ir/basic_block.h" |
| @@ -14,10 +17,10 @@ | |||
| 14 | namespace Shader::Backend::SPIRV { | 17 | namespace Shader::Backend::SPIRV { |
| 15 | namespace { | 18 | namespace { |
| 16 | template <class Func> | 19 | template <class Func> |
| 17 | struct FuncTraits : FuncTraits<decltype(&Func::operator())> {}; | 20 | struct FuncTraits : FuncTraits<Func> {}; |
| 18 | 21 | ||
| 19 | template <class ClassType, class ReturnType_, class... Args> | 22 | template <class ReturnType_, class... Args> |
| 20 | struct FuncTraits<ReturnType_ (ClassType::*)(Args...)> { | 23 | struct FuncTraits<ReturnType_ (*)(Args...)> { |
| 21 | using ReturnType = ReturnType_; | 24 | using ReturnType = ReturnType_; |
| 22 | 25 | ||
| 23 | static constexpr size_t NUM_ARGS = sizeof...(Args); | 26 | static constexpr size_t NUM_ARGS = sizeof...(Args); |
| @@ -26,15 +29,15 @@ struct FuncTraits<ReturnType_ (ClassType::*)(Args...)> { | |||
| 26 | using ArgType = std::tuple_element_t<I, std::tuple<Args...>>; | 29 | using ArgType = std::tuple_element_t<I, std::tuple<Args...>>; |
| 27 | }; | 30 | }; |
| 28 | 31 | ||
| 29 | template <auto method, typename... Args> | 32 | template <auto func, typename... Args> |
| 30 | void SetDefinition(EmitSPIRV& emit, EmitContext& ctx, IR::Inst* inst, Args... args) { | 33 | void SetDefinition(EmitContext& ctx, IR::Inst* inst, Args... args) { |
| 31 | const Id forward_id{inst->Definition<Id>()}; | 34 | const Id forward_id{inst->Definition<Id>()}; |
| 32 | const bool has_forward_id{Sirit::ValidId(forward_id)}; | 35 | const bool has_forward_id{Sirit::ValidId(forward_id)}; |
| 33 | Id current_id{}; | 36 | Id current_id{}; |
| 34 | if (has_forward_id) { | 37 | if (has_forward_id) { |
| 35 | current_id = ctx.ExchangeCurrentId(forward_id); | 38 | current_id = ctx.ExchangeCurrentId(forward_id); |
| 36 | } | 39 | } |
| 37 | const Id new_id{(emit.*method)(ctx, std::forward<Args>(args)...)}; | 40 | const Id new_id{func(ctx, std::forward<Args>(args)...)}; |
| 38 | if (has_forward_id) { | 41 | if (has_forward_id) { |
| 39 | ctx.ExchangeCurrentId(current_id); | 42 | ctx.ExchangeCurrentId(current_id); |
| 40 | } else { | 43 | } else { |
| @@ -55,42 +58,62 @@ ArgType Arg(EmitContext& ctx, const IR::Value& arg) { | |||
| 55 | } | 58 | } |
| 56 | } | 59 | } |
| 57 | 60 | ||
| 58 | template <auto method, bool is_first_arg_inst, size_t... I> | 61 | template <auto func, bool is_first_arg_inst, size_t... I> |
| 59 | void Invoke(EmitSPIRV& emit, EmitContext& ctx, IR::Inst* inst, std::index_sequence<I...>) { | 62 | void Invoke(EmitContext& ctx, IR::Inst* inst, std::index_sequence<I...>) { |
| 60 | using Traits = FuncTraits<decltype(method)>; | 63 | using Traits = FuncTraits<decltype(func)>; |
| 61 | if constexpr (std::is_same_v<Traits::ReturnType, Id>) { | 64 | if constexpr (std::is_same_v<Traits::ReturnType, Id>) { |
| 62 | if constexpr (is_first_arg_inst) { | 65 | if constexpr (is_first_arg_inst) { |
| 63 | SetDefinition<method>(emit, ctx, inst, inst, | 66 | SetDefinition<func>(ctx, inst, inst, Arg<Traits::ArgType<I + 2>>(ctx, inst->Arg(I))...); |
| 64 | Arg<Traits::ArgType<I + 2>>(ctx, inst->Arg(I))...); | ||
| 65 | } else { | 67 | } else { |
| 66 | SetDefinition<method>(emit, ctx, inst, | 68 | SetDefinition<func>(ctx, inst, Arg<Traits::ArgType<I + 1>>(ctx, inst->Arg(I))...); |
| 67 | Arg<Traits::ArgType<I + 1>>(ctx, inst->Arg(I))...); | ||
| 68 | } | 69 | } |
| 69 | } else { | 70 | } else { |
| 70 | if constexpr (is_first_arg_inst) { | 71 | if constexpr (is_first_arg_inst) { |
| 71 | (emit.*method)(ctx, inst, Arg<Traits::ArgType<I + 2>>(ctx, inst->Arg(I))...); | 72 | func(ctx, inst, Arg<Traits::ArgType<I + 2>>(ctx, inst->Arg(I))...); |
| 72 | } else { | 73 | } else { |
| 73 | (emit.*method)(ctx, Arg<Traits::ArgType<I + 1>>(ctx, inst->Arg(I))...); | 74 | func(ctx, Arg<Traits::ArgType<I + 1>>(ctx, inst->Arg(I))...); |
| 74 | } | 75 | } |
| 75 | } | 76 | } |
| 76 | } | 77 | } |
| 77 | 78 | ||
| 78 | template <auto method> | 79 | template <auto func> |
| 79 | void Invoke(EmitSPIRV& emit, EmitContext& ctx, IR::Inst* inst) { | 80 | void Invoke(EmitContext& ctx, IR::Inst* inst) { |
| 80 | using Traits = FuncTraits<decltype(method)>; | 81 | using Traits = FuncTraits<decltype(func)>; |
| 81 | static_assert(Traits::NUM_ARGS >= 1, "Insufficient arguments"); | 82 | static_assert(Traits::NUM_ARGS >= 1, "Insufficient arguments"); |
| 82 | if constexpr (Traits::NUM_ARGS == 1) { | 83 | if constexpr (Traits::NUM_ARGS == 1) { |
| 83 | Invoke<method, false>(emit, ctx, inst, std::make_index_sequence<0>{}); | 84 | Invoke<func, false>(ctx, inst, std::make_index_sequence<0>{}); |
| 84 | } else { | 85 | } else { |
| 85 | using FirstArgType = typename Traits::template ArgType<1>; | 86 | using FirstArgType = typename Traits::template ArgType<1>; |
| 86 | static constexpr bool is_first_arg_inst = std::is_same_v<FirstArgType, IR::Inst*>; | 87 | static constexpr bool is_first_arg_inst = std::is_same_v<FirstArgType, IR::Inst*>; |
| 87 | using Indices = std::make_index_sequence<Traits::NUM_ARGS - (is_first_arg_inst ? 2 : 1)>; | 88 | using Indices = std::make_index_sequence<Traits::NUM_ARGS - (is_first_arg_inst ? 2 : 1)>; |
| 88 | Invoke<method, is_first_arg_inst>(emit, ctx, inst, Indices{}); | 89 | Invoke<func, is_first_arg_inst>(ctx, inst, Indices{}); |
| 90 | } | ||
| 91 | } | ||
| 92 | |||
| 93 | void EmitInst(EmitContext& ctx, IR::Inst* inst) { | ||
| 94 | switch (inst->Opcode()) { | ||
| 95 | #define OPCODE(name, result_type, ...) \ | ||
| 96 | case IR::Opcode::name: \ | ||
| 97 | return Invoke<&Emit##name>(ctx, inst); | ||
| 98 | #include "shader_recompiler/frontend/ir/opcodes.inc" | ||
| 99 | #undef OPCODE | ||
| 100 | } | ||
| 101 | throw LogicError("Invalid opcode {}", inst->Opcode()); | ||
| 102 | } | ||
| 103 | |||
| 104 | Id TypeId(const EmitContext& ctx, IR::Type type) { | ||
| 105 | switch (type) { | ||
| 106 | case IR::Type::U1: | ||
| 107 | return ctx.U1; | ||
| 108 | case IR::Type::U32: | ||
| 109 | return ctx.U32[1]; | ||
| 110 | default: | ||
| 111 | throw NotImplementedException("Phi node type {}", type); | ||
| 89 | } | 112 | } |
| 90 | } | 113 | } |
| 91 | } // Anonymous namespace | 114 | } // Anonymous namespace |
| 92 | 115 | ||
| 93 | EmitSPIRV::EmitSPIRV(IR::Program& program) { | 116 | std::vector<u32> EmitSPIRV(Environment& env, IR::Program& program) { |
| 94 | EmitContext ctx{program}; | 117 | EmitContext ctx{program}; |
| 95 | const Id void_function{ctx.TypeFunction(ctx.void_id)}; | 118 | const Id void_function{ctx.TypeFunction(ctx.void_id)}; |
| 96 | // FIXME: Forward declare functions (needs sirit support) | 119 | // FIXME: Forward declare functions (needs sirit support) |
| @@ -112,43 +135,17 @@ EmitSPIRV::EmitSPIRV(IR::Program& program) { | |||
| 112 | if (program.info.uses_local_invocation_id) { | 135 | if (program.info.uses_local_invocation_id) { |
| 113 | interfaces.push_back(ctx.local_invocation_id); | 136 | interfaces.push_back(ctx.local_invocation_id); |
| 114 | } | 137 | } |
| 115 | |||
| 116 | const std::span interfaces_span(interfaces.data(), interfaces.size()); | 138 | const std::span interfaces_span(interfaces.data(), interfaces.size()); |
| 117 | ctx.AddEntryPoint(spv::ExecutionModel::Fragment, func, "main", interfaces_span); | 139 | ctx.AddEntryPoint(spv::ExecutionModel::GLCompute, func, "main", interfaces_span); |
| 118 | ctx.AddExecutionMode(func, spv::ExecutionMode::OriginUpperLeft); | ||
| 119 | |||
| 120 | std::vector<u32> result{ctx.Assemble()}; | ||
| 121 | std::FILE* file{std::fopen("D:\\shader.spv", "wb")}; | ||
| 122 | std::fwrite(result.data(), sizeof(u32), result.size(), file); | ||
| 123 | std::fclose(file); | ||
| 124 | std::system("spirv-dis D:\\shader.spv") == 0 && | ||
| 125 | std::system("spirv-val --uniform-buffer-standard-layout D:\\shader.spv") == 0 && | ||
| 126 | std::system("spirv-cross -V D:\\shader.spv") == 0; | ||
| 127 | } | ||
| 128 | 140 | ||
| 129 | void EmitSPIRV::EmitInst(EmitContext& ctx, IR::Inst* inst) { | 141 | const std::array<u32, 3> workgroup_size{env.WorkgroupSize()}; |
| 130 | switch (inst->Opcode()) { | 142 | ctx.AddExecutionMode(func, spv::ExecutionMode::LocalSize, workgroup_size[0], workgroup_size[1], |
| 131 | #define OPCODE(name, result_type, ...) \ | 143 | workgroup_size[2]); |
| 132 | case IR::Opcode::name: \ | ||
| 133 | return Invoke<&EmitSPIRV::Emit##name>(*this, ctx, inst); | ||
| 134 | #include "shader_recompiler/frontend/ir/opcodes.inc" | ||
| 135 | #undef OPCODE | ||
| 136 | } | ||
| 137 | throw LogicError("Invalid opcode {}", inst->Opcode()); | ||
| 138 | } | ||
| 139 | 144 | ||
| 140 | static Id TypeId(const EmitContext& ctx, IR::Type type) { | 145 | return ctx.Assemble(); |
| 141 | switch (type) { | ||
| 142 | case IR::Type::U1: | ||
| 143 | return ctx.U1; | ||
| 144 | case IR::Type::U32: | ||
| 145 | return ctx.U32[1]; | ||
| 146 | default: | ||
| 147 | throw NotImplementedException("Phi node type {}", type); | ||
| 148 | } | ||
| 149 | } | 146 | } |
| 150 | 147 | ||
| 151 | Id EmitSPIRV::EmitPhi(EmitContext& ctx, IR::Inst* inst) { | 148 | Id EmitPhi(EmitContext& ctx, IR::Inst* inst) { |
| 152 | const size_t num_args{inst->NumArgs()}; | 149 | const size_t num_args{inst->NumArgs()}; |
| 153 | boost::container::small_vector<Id, 32> operands; | 150 | boost::container::small_vector<Id, 32> operands; |
| 154 | operands.reserve(num_args * 2); | 151 | operands.reserve(num_args * 2); |
| @@ -178,25 +175,25 @@ Id EmitSPIRV::EmitPhi(EmitContext& ctx, IR::Inst* inst) { | |||
| 178 | return ctx.OpPhi(result_type, std::span(operands.data(), operands.size())); | 175 | return ctx.OpPhi(result_type, std::span(operands.data(), operands.size())); |
| 179 | } | 176 | } |
| 180 | 177 | ||
| 181 | void EmitSPIRV::EmitVoid(EmitContext&) {} | 178 | void EmitVoid(EmitContext&) {} |
| 182 | 179 | ||
| 183 | Id EmitSPIRV::EmitIdentity(EmitContext& ctx, const IR::Value& value) { | 180 | Id EmitIdentity(EmitContext& ctx, const IR::Value& value) { |
| 184 | return ctx.Def(value); | 181 | return ctx.Def(value); |
| 185 | } | 182 | } |
| 186 | 183 | ||
| 187 | void EmitSPIRV::EmitGetZeroFromOp(EmitContext&) { | 184 | void EmitGetZeroFromOp(EmitContext&) { |
| 188 | throw LogicError("Unreachable instruction"); | 185 | throw LogicError("Unreachable instruction"); |
| 189 | } | 186 | } |
| 190 | 187 | ||
| 191 | void EmitSPIRV::EmitGetSignFromOp(EmitContext&) { | 188 | void EmitGetSignFromOp(EmitContext&) { |
| 192 | throw LogicError("Unreachable instruction"); | 189 | throw LogicError("Unreachable instruction"); |
| 193 | } | 190 | } |
| 194 | 191 | ||
| 195 | void EmitSPIRV::EmitGetCarryFromOp(EmitContext&) { | 192 | void EmitGetCarryFromOp(EmitContext&) { |
| 196 | throw LogicError("Unreachable instruction"); | 193 | throw LogicError("Unreachable instruction"); |
| 197 | } | 194 | } |
| 198 | 195 | ||
| 199 | void EmitSPIRV::EmitGetOverflowFromOp(EmitContext&) { | 196 | void EmitGetOverflowFromOp(EmitContext&) { |
| 200 | throw LogicError("Unreachable instruction"); | 197 | throw LogicError("Unreachable instruction"); |
| 201 | } | 198 | } |
| 202 | 199 | ||