summaryrefslogtreecommitdiff
path: root/src/shader_recompiler/backend/spirv/emit_spirv.cpp
diff options
context:
space:
mode:
Diffstat (limited to 'src/shader_recompiler/backend/spirv/emit_spirv.cpp')
-rw-r--r--src/shader_recompiler/backend/spirv/emit_spirv.cpp117
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 @@
14namespace Shader::Backend::SPIRV { 17namespace Shader::Backend::SPIRV {
15namespace { 18namespace {
16template <class Func> 19template <class Func>
17struct FuncTraits : FuncTraits<decltype(&Func::operator())> {}; 20struct FuncTraits : FuncTraits<Func> {};
18 21
19template <class ClassType, class ReturnType_, class... Args> 22template <class ReturnType_, class... Args>
20struct FuncTraits<ReturnType_ (ClassType::*)(Args...)> { 23struct 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
29template <auto method, typename... Args> 32template <auto func, typename... Args>
30void SetDefinition(EmitSPIRV& emit, EmitContext& ctx, IR::Inst* inst, Args... args) { 33void 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
58template <auto method, bool is_first_arg_inst, size_t... I> 61template <auto func, bool is_first_arg_inst, size_t... I>
59void Invoke(EmitSPIRV& emit, EmitContext& ctx, IR::Inst* inst, std::index_sequence<I...>) { 62void 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
78template <auto method> 79template <auto func>
79void Invoke(EmitSPIRV& emit, EmitContext& ctx, IR::Inst* inst) { 80void 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
93void 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
104Id 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
93EmitSPIRV::EmitSPIRV(IR::Program& program) { 116std::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
129void 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
140static 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
151Id EmitSPIRV::EmitPhi(EmitContext& ctx, IR::Inst* inst) { 148Id 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
181void EmitSPIRV::EmitVoid(EmitContext&) {} 178void EmitVoid(EmitContext&) {}
182 179
183Id EmitSPIRV::EmitIdentity(EmitContext& ctx, const IR::Value& value) { 180Id EmitIdentity(EmitContext& ctx, const IR::Value& value) {
184 return ctx.Def(value); 181 return ctx.Def(value);
185} 182}
186 183
187void EmitSPIRV::EmitGetZeroFromOp(EmitContext&) { 184void EmitGetZeroFromOp(EmitContext&) {
188 throw LogicError("Unreachable instruction"); 185 throw LogicError("Unreachable instruction");
189} 186}
190 187
191void EmitSPIRV::EmitGetSignFromOp(EmitContext&) { 188void EmitGetSignFromOp(EmitContext&) {
192 throw LogicError("Unreachable instruction"); 189 throw LogicError("Unreachable instruction");
193} 190}
194 191
195void EmitSPIRV::EmitGetCarryFromOp(EmitContext&) { 192void EmitGetCarryFromOp(EmitContext&) {
196 throw LogicError("Unreachable instruction"); 193 throw LogicError("Unreachable instruction");
197} 194}
198 195
199void EmitSPIRV::EmitGetOverflowFromOp(EmitContext&) { 196void EmitGetOverflowFromOp(EmitContext&) {
200 throw LogicError("Unreachable instruction"); 197 throw LogicError("Unreachable instruction");
201} 198}
202 199