summaryrefslogtreecommitdiff
path: root/src/shader_recompiler
diff options
context:
space:
mode:
Diffstat (limited to 'src/shader_recompiler')
-rw-r--r--src/shader_recompiler/CMakeLists.txt4
-rw-r--r--src/shader_recompiler/backend/spirv/emit_context.cpp64
-rw-r--r--src/shader_recompiler/backend/spirv/emit_context.h18
-rw-r--r--src/shader_recompiler/backend/spirv/emit_spirv.cpp44
-rw-r--r--src/shader_recompiler/backend/spirv/emit_spirv.h18
-rw-r--r--src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp55
-rw-r--r--src/shader_recompiler/backend/spirv/emit_spirv_control_flow.cpp23
-rw-r--r--src/shader_recompiler/environment.h14
-rw-r--r--src/shader_recompiler/frontend/ir/attribute.cpp2
-rw-r--r--src/shader_recompiler/frontend/ir/attribute.h2
-rw-r--r--src/shader_recompiler/frontend/ir/ir_emitter.cpp14
-rw-r--r--src/shader_recompiler/frontend/ir/ir_emitter.h4
-rw-r--r--src/shader_recompiler/frontend/ir/microinstruction.cpp3
-rw-r--r--src/shader_recompiler/frontend/ir/opcodes.inc11
-rw-r--r--src/shader_recompiler/frontend/ir/program.h2
-rw-r--r--src/shader_recompiler/frontend/ir/reg.h4
-rw-r--r--src/shader_recompiler/frontend/maxwell/control_flow.cpp31
-rw-r--r--src/shader_recompiler/frontend/maxwell/control_flow.h3
-rw-r--r--src/shader_recompiler/frontend/maxwell/program.cpp1
-rw-r--r--src/shader_recompiler/frontend/maxwell/structured_control_flow.cpp18
-rw-r--r--src/shader_recompiler/frontend/maxwell/translate/impl/exit.cpp15
-rw-r--r--src/shader_recompiler/frontend/maxwell/translate/impl/exit_program.cpp43
-rw-r--r--src/shader_recompiler/frontend/maxwell/translate/impl/impl.h4
-rw-r--r--src/shader_recompiler/frontend/maxwell/translate/impl/load_store_attribute.cpp86
-rw-r--r--src/shader_recompiler/frontend/maxwell/translate/impl/not_implemented.cpp16
-rw-r--r--src/shader_recompiler/frontend/maxwell/translate/impl/texture_fetch.cpp2
-rw-r--r--src/shader_recompiler/frontend/maxwell/translate/impl/texture_fetch_swizzled.cpp2
-rw-r--r--src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp60
-rw-r--r--src/shader_recompiler/ir_opt/ssa_rewrite_pass.cpp2
-rw-r--r--src/shader_recompiler/program_header.h143
-rw-r--r--src/shader_recompiler/recompiler.cpp28
-rw-r--r--src/shader_recompiler/recompiler.h20
-rw-r--r--src/shader_recompiler/shader_info.h10
-rw-r--r--src/shader_recompiler/stage.h19
34 files changed, 629 insertions, 156 deletions
diff --git a/src/shader_recompiler/CMakeLists.txt b/src/shader_recompiler/CMakeLists.txt
index b870e9937..31c394106 100644
--- a/src/shader_recompiler/CMakeLists.txt
+++ b/src/shader_recompiler/CMakeLists.txt
@@ -65,6 +65,7 @@ add_library(shader_recompiler STATIC
65 frontend/maxwell/translate/impl/common_funcs.h 65 frontend/maxwell/translate/impl/common_funcs.h
66 frontend/maxwell/translate/impl/condition_code_set.cpp 66 frontend/maxwell/translate/impl/condition_code_set.cpp
67 frontend/maxwell/translate/impl/double_add.cpp 67 frontend/maxwell/translate/impl/double_add.cpp
68 frontend/maxwell/translate/impl/exit_program.cpp
68 frontend/maxwell/translate/impl/find_leading_one.cpp 69 frontend/maxwell/translate/impl/find_leading_one.cpp
69 frontend/maxwell/translate/impl/floating_point_add.cpp 70 frontend/maxwell/translate/impl/floating_point_add.cpp
70 frontend/maxwell/translate/impl/floating_point_compare.cpp 71 frontend/maxwell/translate/impl/floating_point_compare.cpp
@@ -121,9 +122,8 @@ add_library(shader_recompiler STATIC
121 ir_opt/texture_pass.cpp 122 ir_opt/texture_pass.cpp
122 ir_opt/verification_pass.cpp 123 ir_opt/verification_pass.cpp
123 object_pool.h 124 object_pool.h
125 program_header.h
124 profile.h 126 profile.h
125 recompiler.cpp
126 recompiler.h
127 shader_info.h 127 shader_info.h
128) 128)
129 129
diff --git a/src/shader_recompiler/backend/spirv/emit_context.cpp b/src/shader_recompiler/backend/spirv/emit_context.cpp
index 204389d74..6c79b611b 100644
--- a/src/shader_recompiler/backend/spirv/emit_context.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_context.cpp
@@ -62,18 +62,15 @@ void VectorTypes::Define(Sirit::Module& sirit_ctx, Id base_type, std::string_vie
62 } 62 }
63} 63}
64 64
65EmitContext::EmitContext(const Profile& profile_, IR::Program& program) 65EmitContext::EmitContext(const Profile& profile_, IR::Program& program, u32& binding)
66 : Sirit::Module(0x00010000), profile{profile_} { 66 : Sirit::Module(0x00010000), profile{profile_} {
67 AddCapability(spv::Capability::Shader); 67 AddCapability(spv::Capability::Shader);
68 DefineCommonTypes(program.info); 68 DefineCommonTypes(program.info);
69 DefineCommonConstants(); 69 DefineCommonConstants();
70 DefineSpecialVariables(program.info); 70 DefineInterfaces(program.info, program.stage);
71
72 u32 binding{};
73 DefineConstantBuffers(program.info, binding); 71 DefineConstantBuffers(program.info, binding);
74 DefineStorageBuffers(program.info, binding); 72 DefineStorageBuffers(program.info, binding);
75 DefineTextures(program.info, binding); 73 DefineTextures(program.info, binding);
76
77 DefineLabels(program); 74 DefineLabels(program);
78} 75}
79 76
@@ -96,6 +93,8 @@ Id EmitContext::Def(const IR::Value& value) {
96 return Constant(F32[1], value.F32()); 93 return Constant(F32[1], value.F32());
97 case IR::Type::F64: 94 case IR::Type::F64:
98 return Constant(F64[1], value.F64()); 95 return Constant(F64[1], value.F64());
96 case IR::Type::Label:
97 return value.Label()->Definition<Id>();
99 default: 98 default:
100 throw NotImplementedException("Immediate type {}", value.Type()); 99 throw NotImplementedException("Immediate type {}", value.Type());
101 } 100 }
@@ -109,6 +108,9 @@ void EmitContext::DefineCommonTypes(const Info& info) {
109 F32.Define(*this, TypeFloat(32), "f32"); 108 F32.Define(*this, TypeFloat(32), "f32");
110 U32.Define(*this, TypeInt(32, false), "u32"); 109 U32.Define(*this, TypeInt(32, false), "u32");
111 110
111 input_f32 = Name(TypePointer(spv::StorageClass::Input, F32[1]), "input_f32");
112 output_f32 = Name(TypePointer(spv::StorageClass::Output, F32[1]), "output_f32");
113
112 if (info.uses_int8) { 114 if (info.uses_int8) {
113 AddCapability(spv::Capability::Int8); 115 AddCapability(spv::Capability::Int8);
114 U8 = Name(TypeInt(8, false), "u8"); 116 U8 = Name(TypeInt(8, false), "u8");
@@ -139,15 +141,20 @@ void EmitContext::DefineCommonConstants() {
139 u32_zero_value = Constant(U32[1], 0U); 141 u32_zero_value = Constant(U32[1], 0U);
140} 142}
141 143
142void EmitContext::DefineSpecialVariables(const Info& info) { 144void EmitContext::DefineInterfaces(const Info& info, Stage stage) {
143 const auto define{[this](Id type, spv::BuiltIn builtin, spv::StorageClass storage_class) { 145 const auto define{
144 const Id pointer_type{TypePointer(storage_class, type)}; 146 [this](Id type, std::optional<spv::BuiltIn> builtin, spv::StorageClass storage_class) {
145 const Id id{AddGlobalVariable(pointer_type, spv::StorageClass::Input)}; 147 const Id pointer_type{TypePointer(storage_class, type)};
146 Decorate(id, spv::Decoration::BuiltIn, builtin); 148 const Id id{AddGlobalVariable(pointer_type, storage_class)};
147 return id; 149 if (builtin) {
148 }}; 150 Decorate(id, spv::Decoration::BuiltIn, *builtin);
151 }
152 interfaces.push_back(id);
153 return id;
154 }};
149 using namespace std::placeholders; 155 using namespace std::placeholders;
150 const auto define_input{std::bind(define, _1, _2, spv::StorageClass::Input)}; 156 const auto define_input{std::bind(define, _1, _2, spv::StorageClass::Input)};
157 const auto define_output{std::bind(define, _1, _2, spv::StorageClass::Output)};
151 158
152 if (info.uses_workgroup_id) { 159 if (info.uses_workgroup_id) {
153 workgroup_id = define_input(U32[3], spv::BuiltIn::WorkgroupId); 160 workgroup_id = define_input(U32[3], spv::BuiltIn::WorkgroupId);
@@ -155,6 +162,39 @@ void EmitContext::DefineSpecialVariables(const Info& info) {
155 if (info.uses_local_invocation_id) { 162 if (info.uses_local_invocation_id) {
156 local_invocation_id = define_input(U32[3], spv::BuiltIn::LocalInvocationId); 163 local_invocation_id = define_input(U32[3], spv::BuiltIn::LocalInvocationId);
157 } 164 }
165 if (info.loads_position) {
166 const bool is_fragment{stage != Stage::Fragment};
167 const spv::BuiltIn built_in{is_fragment ? spv::BuiltIn::Position : spv::BuiltIn::FragCoord};
168 input_position = define_input(F32[4], built_in);
169 }
170 for (size_t i = 0; i < info.loads_generics.size(); ++i) {
171 if (info.loads_generics[i]) {
172 // FIXME: Declare size from input
173 input_generics[i] = define_input(F32[4], std::nullopt);
174 Decorate(input_generics[i], spv::Decoration::Location, static_cast<u32>(i));
175 Name(input_generics[i], fmt::format("in_attr{}", i));
176 }
177 }
178 if (info.stores_position) {
179 output_position = define_output(F32[4], spv::BuiltIn::Position);
180 }
181 for (size_t i = 0; i < info.stores_generics.size(); ++i) {
182 if (info.stores_generics[i]) {
183 output_generics[i] = define_output(F32[4], std::nullopt);
184 Decorate(output_generics[i], spv::Decoration::Location, static_cast<u32>(i));
185 Name(output_generics[i], fmt::format("out_attr{}", i));
186 }
187 }
188 if (stage == Stage::Fragment) {
189 for (size_t i = 0; i < 8; ++i) {
190 if (!info.stores_frag_color[i]) {
191 continue;
192 }
193 frag_color[i] = define_output(F32[4], std::nullopt);
194 Decorate(frag_color[i], spv::Decoration::Location, static_cast<u32>(i));
195 Name(frag_color[i], fmt::format("frag_color{}", i));
196 }
197 }
158} 198}
159 199
160void EmitContext::DefineConstantBuffers(const Info& info, u32& binding) { 200void EmitContext::DefineConstantBuffers(const Info& info, u32& binding) {
diff --git a/src/shader_recompiler/backend/spirv/emit_context.h b/src/shader_recompiler/backend/spirv/emit_context.h
index 35eca258a..2d7961ac3 100644
--- a/src/shader_recompiler/backend/spirv/emit_context.h
+++ b/src/shader_recompiler/backend/spirv/emit_context.h
@@ -46,7 +46,7 @@ struct UniformDefinitions {
46 46
47class EmitContext final : public Sirit::Module { 47class EmitContext final : public Sirit::Module {
48public: 48public:
49 explicit EmitContext(const Profile& profile, IR::Program& program); 49 explicit EmitContext(const Profile& profile, IR::Program& program, u32& binding);
50 ~EmitContext(); 50 ~EmitContext();
51 51
52 [[nodiscard]] Id Def(const IR::Value& value); 52 [[nodiscard]] Id Def(const IR::Value& value);
@@ -71,6 +71,9 @@ public:
71 71
72 UniformDefinitions uniform_types; 72 UniformDefinitions uniform_types;
73 73
74 Id input_f32{};
75 Id output_f32{};
76
74 Id storage_u32{}; 77 Id storage_u32{};
75 78
76 std::array<UniformDefinitions, Info::MAX_CBUFS> cbufs{}; 79 std::array<UniformDefinitions, Info::MAX_CBUFS> cbufs{};
@@ -80,10 +83,21 @@ public:
80 Id workgroup_id{}; 83 Id workgroup_id{};
81 Id local_invocation_id{}; 84 Id local_invocation_id{};
82 85
86 Id input_position{};
87 std::array<Id, 32> input_generics{};
88
89 Id output_position{};
90 std::array<Id, 32> output_generics{};
91
92 std::array<Id, 8> frag_color{};
93 Id frag_depth {};
94
95 std::vector<Id> interfaces;
96
83private: 97private:
84 void DefineCommonTypes(const Info& info); 98 void DefineCommonTypes(const Info& info);
85 void DefineCommonConstants(); 99 void DefineCommonConstants();
86 void DefineSpecialVariables(const Info& info); 100 void DefineInterfaces(const Info& info, Stage stage);
87 void DefineConstantBuffers(const Info& info, u32& binding); 101 void DefineConstantBuffers(const Info& info, u32& binding);
88 void DefineConstantBuffers(const Info& info, Id UniformDefinitions::*member_type, u32 binding, 102 void DefineConstantBuffers(const Info& info, Id UniformDefinitions::*member_type, u32 binding,
89 Id type, char type_char, u32 element_size); 103 Id type, char type_char, u32 element_size);
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) {
54 return arg.U32(); 54 return arg.U32();
55 } else if constexpr (std::is_same_v<ArgType, IR::Block*>) { 55 } else if constexpr (std::is_same_v<ArgType, IR::Block*>) {
56 return arg.Label(); 56 return arg.Label();
57 } else if constexpr (std::is_same_v<ArgType, IR::Attribute>) {
58 return arg.Attribute();
57 } 59 }
58} 60}
59 61
@@ -197,8 +199,9 @@ Id PhiArgDef(EmitContext& ctx, IR::Inst* inst, size_t index) {
197} 199}
198} // Anonymous namespace 200} // Anonymous namespace
199 201
200std::vector<u32> EmitSPIRV(const Profile& profile, Environment& env, IR::Program& program) { 202std::vector<u32> EmitSPIRV(const Profile& profile, Environment& env, IR::Program& program,
201 EmitContext ctx{profile, program}; 203 u32& binding) {
204 EmitContext ctx{profile, program, binding};
202 const Id void_function{ctx.TypeFunction(ctx.void_id)}; 205 const Id void_function{ctx.TypeFunction(ctx.void_id)};
203 const Id func{ctx.OpFunction(ctx.void_id, spv::FunctionControlMask::MaskNone, void_function)}; 206 const Id func{ctx.OpFunction(ctx.void_id, spv::FunctionControlMask::MaskNone, void_function)};
204 for (IR::Block* const block : program.blocks) { 207 for (IR::Block* const block : program.blocks) {
@@ -208,28 +211,41 @@ std::vector<u32> EmitSPIRV(const Profile& profile, Environment& env, IR::Program
208 } 211 }
209 } 212 }
210 ctx.OpFunctionEnd(); 213 ctx.OpFunctionEnd();
211 boost::container::small_vector<Id, 32> interfaces; 214
212 const Info& info{program.info}; 215 const std::span interfaces(ctx.interfaces.data(), ctx.interfaces.size());
213 if (info.uses_workgroup_id) { 216 spv::ExecutionModel execution_model{};
214 interfaces.push_back(ctx.workgroup_id); 217 switch (env.ShaderStage()) {
218 case Shader::Stage::Compute: {
219 const std::array<u32, 3> workgroup_size{env.WorkgroupSize()};
220 execution_model = spv::ExecutionModel::GLCompute;
221 ctx.AddExecutionMode(func, spv::ExecutionMode::LocalSize, workgroup_size[0],
222 workgroup_size[1], workgroup_size[2]);
223 break;
215 } 224 }
216 if (info.uses_local_invocation_id) { 225 case Shader::Stage::VertexB:
217 interfaces.push_back(ctx.local_invocation_id); 226 execution_model = spv::ExecutionModel::Vertex;
227 break;
228 case Shader::Stage::Fragment:
229 execution_model = spv::ExecutionModel::Fragment;
230 ctx.AddExecutionMode(func, spv::ExecutionMode::OriginUpperLeft);
231 break;
232 default:
233 throw NotImplementedException("Stage {}", env.ShaderStage());
218 } 234 }
219 const std::span interfaces_span(interfaces.data(), interfaces.size()); 235 ctx.AddEntryPoint(execution_model, func, "main", interfaces);
220 ctx.AddEntryPoint(spv::ExecutionModel::GLCompute, func, "main", interfaces_span);
221
222 const std::array<u32, 3> workgroup_size{env.WorkgroupSize()};
223 ctx.AddExecutionMode(func, spv::ExecutionMode::LocalSize, workgroup_size[0], workgroup_size[1],
224 workgroup_size[2]);
225 236
226 SetupDenormControl(profile, program, ctx, func); 237 SetupDenormControl(profile, program, ctx, func);
238 const Info& info{program.info};
227 if (info.uses_sampled_1d) { 239 if (info.uses_sampled_1d) {
228 ctx.AddCapability(spv::Capability::Sampled1D); 240 ctx.AddCapability(spv::Capability::Sampled1D);
229 } 241 }
230 if (info.uses_sparse_residency) { 242 if (info.uses_sparse_residency) {
231 ctx.AddCapability(spv::Capability::SparseResidency); 243 ctx.AddCapability(spv::Capability::SparseResidency);
232 } 244 }
245 if (info.uses_demote_to_helper_invocation) {
246 ctx.AddExtension("SPV_EXT_demote_to_helper_invocation");
247 ctx.AddCapability(spv::Capability::DemoteToHelperInvocationEXT);
248 }
233 // TODO: Track this usage 249 // TODO: Track this usage
234 ctx.AddCapability(spv::Capability::ImageGatherExtended); 250 ctx.AddCapability(spv::Capability::ImageGatherExtended);
235 251
diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.h b/src/shader_recompiler/backend/spirv/emit_spirv.h
index 89566c83d..ae121f534 100644
--- a/src/shader_recompiler/backend/spirv/emit_spirv.h
+++ b/src/shader_recompiler/backend/spirv/emit_spirv.h
@@ -16,18 +16,18 @@
16namespace Shader::Backend::SPIRV { 16namespace Shader::Backend::SPIRV {
17 17
18[[nodiscard]] std::vector<u32> EmitSPIRV(const Profile& profile, Environment& env, 18[[nodiscard]] std::vector<u32> EmitSPIRV(const Profile& profile, Environment& env,
19 IR::Program& program); 19 IR::Program& program, u32& binding);
20 20
21// Microinstruction emitters 21// Microinstruction emitters
22Id EmitPhi(EmitContext& ctx, IR::Inst* inst); 22Id EmitPhi(EmitContext& ctx, IR::Inst* inst);
23void EmitVoid(EmitContext& ctx); 23void EmitVoid(EmitContext& ctx);
24Id EmitIdentity(EmitContext& ctx, const IR::Value& value); 24Id EmitIdentity(EmitContext& ctx, const IR::Value& value);
25void EmitBranch(EmitContext& ctx, IR::Block* label); 25void EmitBranch(EmitContext& ctx, Id label);
26void EmitBranchConditional(EmitContext& ctx, Id condition, IR::Block* true_label, 26void EmitBranchConditional(EmitContext& ctx, Id condition, Id true_label, Id false_label);
27 IR::Block* false_label); 27void EmitLoopMerge(EmitContext& ctx, Id merge_label, Id continue_label);
28void EmitLoopMerge(EmitContext& ctx, IR::Block* merge_label, IR::Block* continue_label); 28void EmitSelectionMerge(EmitContext& ctx, Id merge_label);
29void EmitSelectionMerge(EmitContext& ctx, IR::Block* merge_label);
30void EmitReturn(EmitContext& ctx); 29void EmitReturn(EmitContext& ctx);
30void EmitDemoteToHelperInvocation(EmitContext& ctx, Id continue_label);
31void EmitGetRegister(EmitContext& ctx); 31void EmitGetRegister(EmitContext& ctx);
32void EmitSetRegister(EmitContext& ctx); 32void EmitSetRegister(EmitContext& ctx);
33void EmitGetPred(EmitContext& ctx); 33void EmitGetPred(EmitContext& ctx);
@@ -41,10 +41,12 @@ Id EmitGetCbufS16(EmitContext& ctx, const IR::Value& binding, const IR::Value& o
41Id EmitGetCbufU32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset); 41Id EmitGetCbufU32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
42Id EmitGetCbufF32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset); 42Id EmitGetCbufF32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
43Id EmitGetCbufU64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset); 43Id EmitGetCbufU64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
44void EmitGetAttribute(EmitContext& ctx); 44Id EmitGetAttribute(EmitContext& ctx, IR::Attribute attr);
45void EmitSetAttribute(EmitContext& ctx); 45void EmitSetAttribute(EmitContext& ctx, IR::Attribute attr, Id value);
46void EmitGetAttributeIndexed(EmitContext& ctx); 46void EmitGetAttributeIndexed(EmitContext& ctx);
47void EmitSetAttributeIndexed(EmitContext& ctx); 47void EmitSetAttributeIndexed(EmitContext& ctx);
48void EmitSetFragColor(EmitContext& ctx, u32 index, u32 component, Id value);
49void EmitSetFragDepth(EmitContext& ctx, Id value);
48void EmitGetZFlag(EmitContext& ctx); 50void EmitGetZFlag(EmitContext& ctx);
49void EmitGetSFlag(EmitContext& ctx); 51void EmitGetSFlag(EmitContext& ctx);
50void EmitGetCFlag(EmitContext& ctx); 52void EmitGetCFlag(EmitContext& ctx);
diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp
index 125b58cf7..02d115740 100644
--- a/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp
@@ -5,6 +5,43 @@
5#include "shader_recompiler/backend/spirv/emit_spirv.h" 5#include "shader_recompiler/backend/spirv/emit_spirv.h"
6 6
7namespace Shader::Backend::SPIRV { 7namespace Shader::Backend::SPIRV {
8namespace {
9Id InputAttrPointer(EmitContext& ctx, IR::Attribute attr) {
10 const u32 element{static_cast<u32>(attr) % 4};
11 const auto element_id{[&] { return ctx.Constant(ctx.U32[1], element); }};
12 if (IR::IsGeneric(attr)) {
13 const u32 index{IR::GenericAttributeIndex(attr)};
14 return ctx.OpAccessChain(ctx.input_f32, ctx.input_generics.at(index), element_id());
15 }
16 switch (attr) {
17 case IR::Attribute::PositionX:
18 case IR::Attribute::PositionY:
19 case IR::Attribute::PositionZ:
20 case IR::Attribute::PositionW:
21 return ctx.OpAccessChain(ctx.input_f32, ctx.input_position, element_id());
22 default:
23 throw NotImplementedException("Read attribute {}", attr);
24 }
25}
26
27Id OutputAttrPointer(EmitContext& ctx, IR::Attribute attr) {
28 const u32 element{static_cast<u32>(attr) % 4};
29 const auto element_id{[&] { return ctx.Constant(ctx.U32[1], element); }};
30 if (IR::IsGeneric(attr)) {
31 const u32 index{IR::GenericAttributeIndex(attr)};
32 return ctx.OpAccessChain(ctx.output_f32, ctx.output_generics.at(index), element_id());
33 }
34 switch (attr) {
35 case IR::Attribute::PositionX:
36 case IR::Attribute::PositionY:
37 case IR::Attribute::PositionZ:
38 case IR::Attribute::PositionW:
39 return ctx.OpAccessChain(ctx.output_f32, ctx.output_position, element_id());
40 default:
41 throw NotImplementedException("Read attribute {}", attr);
42 }
43}
44} // Anonymous namespace
8 45
9void EmitGetRegister(EmitContext&) { 46void EmitGetRegister(EmitContext&) {
10 throw NotImplementedException("SPIR-V Instruction"); 47 throw NotImplementedException("SPIR-V Instruction");
@@ -87,12 +124,12 @@ Id EmitGetCbufU64(EmitContext& ctx, const IR::Value& binding, const IR::Value& o
87 return GetCbuf(ctx, ctx.U64, &UniformDefinitions::U64, sizeof(u64), binding, offset); 124 return GetCbuf(ctx, ctx.U64, &UniformDefinitions::U64, sizeof(u64), binding, offset);
88} 125}
89 126
90void EmitGetAttribute(EmitContext&) { 127Id EmitGetAttribute(EmitContext& ctx, IR::Attribute attr) {
91 throw NotImplementedException("SPIR-V Instruction"); 128 return ctx.OpLoad(ctx.F32[1], InputAttrPointer(ctx, attr));
92} 129}
93 130
94void EmitSetAttribute(EmitContext&) { 131void EmitSetAttribute(EmitContext& ctx, IR::Attribute attr, Id value) {
95 throw NotImplementedException("SPIR-V Instruction"); 132 ctx.OpStore(OutputAttrPointer(ctx, attr), value);
96} 133}
97 134
98void EmitGetAttributeIndexed(EmitContext&) { 135void EmitGetAttributeIndexed(EmitContext&) {
@@ -103,6 +140,16 @@ void EmitSetAttributeIndexed(EmitContext&) {
103 throw NotImplementedException("SPIR-V Instruction"); 140 throw NotImplementedException("SPIR-V Instruction");
104} 141}
105 142
143void EmitSetFragColor(EmitContext& ctx, u32 index, u32 component, Id value) {
144 const Id component_id{ctx.Constant(ctx.U32[1], component)};
145 const Id pointer{ctx.OpAccessChain(ctx.output_f32, ctx.frag_color.at(index), component_id)};
146 ctx.OpStore(pointer, value);
147}
148
149void EmitSetFragDepth(EmitContext& ctx, Id value) {
150 ctx.OpStore(ctx.frag_depth, value);
151}
152
106void EmitGetZFlag(EmitContext&) { 153void EmitGetZFlag(EmitContext&) {
107 throw NotImplementedException("SPIR-V Instruction"); 154 throw NotImplementedException("SPIR-V Instruction");
108} 155}
diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_control_flow.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_control_flow.cpp
index 48755b827..6b81f0169 100644
--- a/src/shader_recompiler/backend/spirv/emit_spirv_control_flow.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_control_flow.cpp
@@ -6,26 +6,29 @@
6 6
7namespace Shader::Backend::SPIRV { 7namespace Shader::Backend::SPIRV {
8 8
9void EmitBranch(EmitContext& ctx, IR::Block* label) { 9void EmitBranch(EmitContext& ctx, Id label) {
10 ctx.OpBranch(label->Definition<Id>()); 10 ctx.OpBranch(label);
11} 11}
12 12
13void EmitBranchConditional(EmitContext& ctx, Id condition, IR::Block* true_label, 13void EmitBranchConditional(EmitContext& ctx, Id condition, Id true_label, Id false_label) {
14 IR::Block* false_label) { 14 ctx.OpBranchConditional(condition, true_label, false_label);
15 ctx.OpBranchConditional(condition, true_label->Definition<Id>(), false_label->Definition<Id>());
16} 15}
17 16
18void EmitLoopMerge(EmitContext& ctx, IR::Block* merge_label, IR::Block* continue_label) { 17void EmitLoopMerge(EmitContext& ctx, Id merge_label, Id continue_label) {
19 ctx.OpLoopMerge(merge_label->Definition<Id>(), continue_label->Definition<Id>(), 18 ctx.OpLoopMerge(merge_label, continue_label, spv::LoopControlMask::MaskNone);
20 spv::LoopControlMask::MaskNone);
21} 19}
22 20
23void EmitSelectionMerge(EmitContext& ctx, IR::Block* merge_label) { 21void EmitSelectionMerge(EmitContext& ctx, Id merge_label) {
24 ctx.OpSelectionMerge(merge_label->Definition<Id>(), spv::SelectionControlMask::MaskNone); 22 ctx.OpSelectionMerge(merge_label, spv::SelectionControlMask::MaskNone);
25} 23}
26 24
27void EmitReturn(EmitContext& ctx) { 25void EmitReturn(EmitContext& ctx) {
28 ctx.OpReturn(); 26 ctx.OpReturn();
29} 27}
30 28
29void EmitDemoteToHelperInvocation(EmitContext& ctx, Id continue_label) {
30 ctx.OpDemoteToHelperInvocationEXT();
31 ctx.OpBranch(continue_label);
32}
33
31} // namespace Shader::Backend::SPIRV 34} // namespace Shader::Backend::SPIRV
diff --git a/src/shader_recompiler/environment.h b/src/shader_recompiler/environment.h
index 0fcb68050..1fcaa56dd 100644
--- a/src/shader_recompiler/environment.h
+++ b/src/shader_recompiler/environment.h
@@ -3,6 +3,8 @@
3#include <array> 3#include <array>
4 4
5#include "common/common_types.h" 5#include "common/common_types.h"
6#include "shader_recompiler/stage.h"
7#include "shader_recompiler/program_header.h"
6 8
7namespace Shader { 9namespace Shader {
8 10
@@ -15,6 +17,18 @@ public:
15 [[nodiscard]] virtual u32 TextureBoundBuffer() = 0; 17 [[nodiscard]] virtual u32 TextureBoundBuffer() = 0;
16 18
17 [[nodiscard]] virtual std::array<u32, 3> WorkgroupSize() = 0; 19 [[nodiscard]] virtual std::array<u32, 3> WorkgroupSize() = 0;
20
21 [[nodiscard]] const ProgramHeader& SPH() const noexcept {
22 return sph;
23 }
24
25 [[nodiscard]] Stage ShaderStage() const noexcept {
26 return stage;
27 }
28
29protected:
30 ProgramHeader sph{};
31 Stage stage{};
18}; 32};
19 33
20} // namespace Shader 34} // namespace Shader
diff --git a/src/shader_recompiler/frontend/ir/attribute.cpp b/src/shader_recompiler/frontend/ir/attribute.cpp
index 2fb7d576f..4811242ea 100644
--- a/src/shader_recompiler/frontend/ir/attribute.cpp
+++ b/src/shader_recompiler/frontend/ir/attribute.cpp
@@ -13,7 +13,7 @@ bool IsGeneric(Attribute attribute) noexcept {
13 return attribute >= Attribute::Generic0X && attribute <= Attribute::Generic31X; 13 return attribute >= Attribute::Generic0X && attribute <= Attribute::Generic31X;
14} 14}
15 15
16int GenericAttributeIndex(Attribute attribute) { 16u32 GenericAttributeIndex(Attribute attribute) {
17 if (!IsGeneric(attribute)) { 17 if (!IsGeneric(attribute)) {
18 throw InvalidArgument("Attribute is not generic {}", attribute); 18 throw InvalidArgument("Attribute is not generic {}", attribute);
19 } 19 }
diff --git a/src/shader_recompiler/frontend/ir/attribute.h b/src/shader_recompiler/frontend/ir/attribute.h
index bb2cad6af..34ec7e0cd 100644
--- a/src/shader_recompiler/frontend/ir/attribute.h
+++ b/src/shader_recompiler/frontend/ir/attribute.h
@@ -224,7 +224,7 @@ enum class Attribute : u64 {
224 224
225[[nodiscard]] bool IsGeneric(Attribute attribute) noexcept; 225[[nodiscard]] bool IsGeneric(Attribute attribute) noexcept;
226 226
227[[nodiscard]] int GenericAttributeIndex(Attribute attribute); 227[[nodiscard]] u32 GenericAttributeIndex(Attribute attribute);
228 228
229[[nodiscard]] std::string NameOf(Attribute attribute); 229[[nodiscard]] std::string NameOf(Attribute attribute);
230 230
diff --git a/src/shader_recompiler/frontend/ir/ir_emitter.cpp b/src/shader_recompiler/frontend/ir/ir_emitter.cpp
index 958282160..672836c0b 100644
--- a/src/shader_recompiler/frontend/ir/ir_emitter.cpp
+++ b/src/shader_recompiler/frontend/ir/ir_emitter.cpp
@@ -82,6 +82,12 @@ void IREmitter::Return() {
82 Inst(Opcode::Return); 82 Inst(Opcode::Return);
83} 83}
84 84
85void IREmitter::DemoteToHelperInvocation(Block* continue_label) {
86 block->SetBranch(continue_label);
87 continue_label->AddImmediatePredecessor(block);
88 Inst(Opcode::DemoteToHelperInvocation, continue_label);
89}
90
85U32 IREmitter::GetReg(IR::Reg reg) { 91U32 IREmitter::GetReg(IR::Reg reg) {
86 return Inst<U32>(Opcode::GetRegister, reg); 92 return Inst<U32>(Opcode::GetRegister, reg);
87} 93}
@@ -248,6 +254,14 @@ void IREmitter::SetAttribute(IR::Attribute attribute, const F32& value) {
248 Inst(Opcode::SetAttribute, attribute, value); 254 Inst(Opcode::SetAttribute, attribute, value);
249} 255}
250 256
257void IREmitter::SetFragColor(u32 index, u32 component, const F32& value) {
258 Inst(Opcode::SetFragColor, Imm32(index), Imm32(component), value);
259}
260
261void IREmitter::SetFragDepth(const F32& value) {
262 Inst(Opcode::SetFragDepth, value);
263}
264
251U32 IREmitter::WorkgroupIdX() { 265U32 IREmitter::WorkgroupIdX() {
252 return U32{CompositeExtract(Inst(Opcode::WorkgroupId), 0)}; 266 return U32{CompositeExtract(Inst(Opcode::WorkgroupId), 0)};
253} 267}
diff --git a/src/shader_recompiler/frontend/ir/ir_emitter.h b/src/shader_recompiler/frontend/ir/ir_emitter.h
index 05263fe8b..72af5db37 100644
--- a/src/shader_recompiler/frontend/ir/ir_emitter.h
+++ b/src/shader_recompiler/frontend/ir/ir_emitter.h
@@ -36,6 +36,7 @@ public:
36 void LoopMerge(Block* merge_block, Block* continue_target); 36 void LoopMerge(Block* merge_block, Block* continue_target);
37 void SelectionMerge(Block* merge_block); 37 void SelectionMerge(Block* merge_block);
38 void Return(); 38 void Return();
39 void DemoteToHelperInvocation(Block* continue_label);
39 40
40 [[nodiscard]] U32 GetReg(IR::Reg reg); 41 [[nodiscard]] U32 GetReg(IR::Reg reg);
41 void SetReg(IR::Reg reg, const U32& value); 42 void SetReg(IR::Reg reg, const U32& value);
@@ -67,6 +68,9 @@ public:
67 [[nodiscard]] F32 GetAttribute(IR::Attribute attribute); 68 [[nodiscard]] F32 GetAttribute(IR::Attribute attribute);
68 void SetAttribute(IR::Attribute attribute, const F32& value); 69 void SetAttribute(IR::Attribute attribute, const F32& value);
69 70
71 void SetFragColor(u32 index, u32 component, const F32& value);
72 void SetFragDepth(const F32& value);
73
70 [[nodiscard]] U32 WorkgroupIdX(); 74 [[nodiscard]] U32 WorkgroupIdX();
71 [[nodiscard]] U32 WorkgroupIdY(); 75 [[nodiscard]] U32 WorkgroupIdY();
72 [[nodiscard]] U32 WorkgroupIdZ(); 76 [[nodiscard]] U32 WorkgroupIdZ();
diff --git a/src/shader_recompiler/frontend/ir/microinstruction.cpp b/src/shader_recompiler/frontend/ir/microinstruction.cpp
index 5946105d2..21b7d8a9f 100644
--- a/src/shader_recompiler/frontend/ir/microinstruction.cpp
+++ b/src/shader_recompiler/frontend/ir/microinstruction.cpp
@@ -55,8 +55,11 @@ bool Inst::MayHaveSideEffects() const noexcept {
55 case Opcode::LoopMerge: 55 case Opcode::LoopMerge:
56 case Opcode::SelectionMerge: 56 case Opcode::SelectionMerge:
57 case Opcode::Return: 57 case Opcode::Return:
58 case Opcode::DemoteToHelperInvocation:
58 case Opcode::SetAttribute: 59 case Opcode::SetAttribute:
59 case Opcode::SetAttributeIndexed: 60 case Opcode::SetAttributeIndexed:
61 case Opcode::SetFragColor:
62 case Opcode::SetFragDepth:
60 case Opcode::WriteGlobalU8: 63 case Opcode::WriteGlobalU8:
61 case Opcode::WriteGlobalS8: 64 case Opcode::WriteGlobalS8:
62 case Opcode::WriteGlobalU16: 65 case Opcode::WriteGlobalU16:
diff --git a/src/shader_recompiler/frontend/ir/opcodes.inc b/src/shader_recompiler/frontend/ir/opcodes.inc
index 9052a4903..593faca52 100644
--- a/src/shader_recompiler/frontend/ir/opcodes.inc
+++ b/src/shader_recompiler/frontend/ir/opcodes.inc
@@ -13,6 +13,7 @@ OPCODE(BranchConditional, Void, U1,
13OPCODE(LoopMerge, Void, Label, Label, ) 13OPCODE(LoopMerge, Void, Label, Label, )
14OPCODE(SelectionMerge, Void, Label, ) 14OPCODE(SelectionMerge, Void, Label, )
15OPCODE(Return, Void, ) 15OPCODE(Return, Void, )
16OPCODE(DemoteToHelperInvocation, Void, Label, )
16 17
17// Context getters/setters 18// Context getters/setters
18OPCODE(GetRegister, U32, Reg, ) 19OPCODE(GetRegister, U32, Reg, )
@@ -28,10 +29,12 @@ OPCODE(GetCbufS16, U32, U32,
28OPCODE(GetCbufU32, U32, U32, U32, ) 29OPCODE(GetCbufU32, U32, U32, U32, )
29OPCODE(GetCbufF32, F32, U32, U32, ) 30OPCODE(GetCbufF32, F32, U32, U32, )
30OPCODE(GetCbufU64, U64, U32, U32, ) 31OPCODE(GetCbufU64, U64, U32, U32, )
31OPCODE(GetAttribute, U32, Attribute, ) 32OPCODE(GetAttribute, F32, Attribute, )
32OPCODE(SetAttribute, Void, Attribute, U32, ) 33OPCODE(SetAttribute, Void, Attribute, F32, )
33OPCODE(GetAttributeIndexed, U32, U32, ) 34OPCODE(GetAttributeIndexed, F32, U32, )
34OPCODE(SetAttributeIndexed, Void, U32, U32, ) 35OPCODE(SetAttributeIndexed, Void, U32, F32, )
36OPCODE(SetFragColor, Void, U32, U32, F32, )
37OPCODE(SetFragDepth, Void, F32, )
35OPCODE(GetZFlag, U1, Void, ) 38OPCODE(GetZFlag, U1, Void, )
36OPCODE(GetSFlag, U1, Void, ) 39OPCODE(GetSFlag, U1, Void, )
37OPCODE(GetCFlag, U1, Void, ) 40OPCODE(GetCFlag, U1, Void, )
diff --git a/src/shader_recompiler/frontend/ir/program.h b/src/shader_recompiler/frontend/ir/program.h
index bce8b19b3..733513c8b 100644
--- a/src/shader_recompiler/frontend/ir/program.h
+++ b/src/shader_recompiler/frontend/ir/program.h
@@ -10,6 +10,7 @@
10 10
11#include "shader_recompiler/frontend/ir/basic_block.h" 11#include "shader_recompiler/frontend/ir/basic_block.h"
12#include "shader_recompiler/shader_info.h" 12#include "shader_recompiler/shader_info.h"
13#include "shader_recompiler/stage.h"
13 14
14namespace Shader::IR { 15namespace Shader::IR {
15 16
@@ -17,6 +18,7 @@ struct Program {
17 BlockList blocks; 18 BlockList blocks;
18 BlockList post_order_blocks; 19 BlockList post_order_blocks;
19 Info info; 20 Info info;
21 Stage stage{};
20}; 22};
21 23
22[[nodiscard]] std::string DumpProgram(const Program& program); 24[[nodiscard]] std::string DumpProgram(const Program& program);
diff --git a/src/shader_recompiler/frontend/ir/reg.h b/src/shader_recompiler/frontend/ir/reg.h
index 8fea05f7b..3845ec5fb 100644
--- a/src/shader_recompiler/frontend/ir/reg.h
+++ b/src/shader_recompiler/frontend/ir/reg.h
@@ -293,12 +293,12 @@ constexpr size_t NUM_REGS = 256;
293 return reg + (-num); 293 return reg + (-num);
294} 294}
295 295
296[[nodiscard]] constexpr Reg operator++(Reg& reg) { 296constexpr Reg operator++(Reg& reg) {
297 reg = reg + 1; 297 reg = reg + 1;
298 return reg; 298 return reg;
299} 299}
300 300
301[[nodiscard]] constexpr Reg operator++(Reg& reg, int) { 301constexpr Reg operator++(Reg& reg, int) {
302 const Reg copy{reg}; 302 const Reg copy{reg};
303 reg = reg + 1; 303 reg = reg + 1;
304 return copy; 304 return copy;
diff --git a/src/shader_recompiler/frontend/maxwell/control_flow.cpp b/src/shader_recompiler/frontend/maxwell/control_flow.cpp
index 715c0e92d..4f6707fae 100644
--- a/src/shader_recompiler/frontend/maxwell/control_flow.cpp
+++ b/src/shader_recompiler/frontend/maxwell/control_flow.cpp
@@ -104,6 +104,7 @@ bool HasFlowTest(Opcode opcode) {
104 case Opcode::EXIT: 104 case Opcode::EXIT:
105 case Opcode::JMP: 105 case Opcode::JMP:
106 case Opcode::JMX: 106 case Opcode::JMX:
107 case Opcode::KIL:
107 case Opcode::BRK: 108 case Opcode::BRK:
108 case Opcode::CONT: 109 case Opcode::CONT:
109 case Opcode::LONGJMP: 110 case Opcode::LONGJMP:
@@ -287,6 +288,13 @@ CFG::AnalysisState CFG::AnalyzeInst(Block* block, FunctionId function_id, Locati
287 block->end = pc; 288 block->end = pc;
288 return AnalysisState::Branch; 289 return AnalysisState::Branch;
289 } 290 }
291 case Opcode::KIL: {
292 const Predicate pred{inst.Pred()};
293 const auto ir_pred{static_cast<IR::Pred>(pred.index)};
294 const IR::Condition cond{inst.branch.flow_test, ir_pred, pred.negated};
295 AnalyzeCondInst(block, function_id, pc, EndClass::Kill, cond);
296 return AnalysisState::Branch;
297 }
290 case Opcode::PBK: 298 case Opcode::PBK:
291 case Opcode::PCNT: 299 case Opcode::PCNT:
292 case Opcode::PEXIT: 300 case Opcode::PEXIT:
@@ -324,13 +332,12 @@ CFG::AnalysisState CFG::AnalyzeInst(Block* block, FunctionId function_id, Locati
324 return AnalysisState::Continue; 332 return AnalysisState::Continue;
325 } 333 }
326 const IR::Condition cond{static_cast<IR::Pred>(pred.index), pred.negated}; 334 const IR::Condition cond{static_cast<IR::Pred>(pred.index), pred.negated};
327 AnalyzeCondInst(block, function_id, pc, EndClass::Branch, cond, true); 335 AnalyzeCondInst(block, function_id, pc, EndClass::Branch, cond);
328 return AnalysisState::Branch; 336 return AnalysisState::Branch;
329} 337}
330 338
331void CFG::AnalyzeCondInst(Block* block, FunctionId function_id, Location pc, 339void CFG::AnalyzeCondInst(Block* block, FunctionId function_id, Location pc,
332 EndClass insn_end_class, IR::Condition cond, 340 EndClass insn_end_class, IR::Condition cond) {
333 bool visit_conditional_inst) {
334 if (block->begin != pc) { 341 if (block->begin != pc) {
335 // If the block doesn't start in the conditional instruction 342 // If the block doesn't start in the conditional instruction
336 // mark it as a label to visit it later 343 // mark it as a label to visit it later
@@ -356,14 +363,16 @@ void CFG::AnalyzeCondInst(Block* block, FunctionId function_id, Location pc,
356 // Impersonate the visited block with a virtual block 363 // Impersonate the visited block with a virtual block
357 *block = std::move(virtual_block); 364 *block = std::move(virtual_block);
358 // Set the end properties of the conditional instruction 365 // Set the end properties of the conditional instruction
359 conditional_block->end = visit_conditional_inst ? (pc + 1) : pc; 366 conditional_block->end = pc + 1;
360 conditional_block->end_class = insn_end_class; 367 conditional_block->end_class = insn_end_class;
361 // Add a label to the instruction after the conditional instruction 368 // Add a label to the instruction after the conditional instruction
362 Block* const endif_block{AddLabel(conditional_block, block->stack, pc + 1, function_id)}; 369 Block* const endif_block{AddLabel(conditional_block, block->stack, pc + 1, function_id)};
363 // Branch to the next instruction from the virtual block 370 // Branch to the next instruction from the virtual block
364 block->branch_false = endif_block; 371 block->branch_false = endif_block;
365 // And branch to it from the conditional instruction if it is a branch 372 // And branch to it from the conditional instruction if it is a branch or a kill instruction
366 if (insn_end_class == EndClass::Branch) { 373 // Kill instructions are considered a branch because they demote to a helper invocation and
374 // execution may continue.
375 if (insn_end_class == EndClass::Branch || insn_end_class == EndClass::Kill) {
367 conditional_block->cond = IR::Condition{true}; 376 conditional_block->cond = IR::Condition{true};
368 conditional_block->branch_true = endif_block; 377 conditional_block->branch_true = endif_block;
369 conditional_block->branch_false = nullptr; 378 conditional_block->branch_false = nullptr;
@@ -415,7 +424,7 @@ CFG::AnalysisState CFG::AnalyzeEXIT(Block* block, FunctionId function_id, Locati
415 throw NotImplementedException("Conditional EXIT with PEXIT token"); 424 throw NotImplementedException("Conditional EXIT with PEXIT token");
416 } 425 }
417 const IR::Condition cond{flow_test, static_cast<IR::Pred>(pred.index), pred.negated}; 426 const IR::Condition cond{flow_test, static_cast<IR::Pred>(pred.index), pred.negated};
418 AnalyzeCondInst(block, function_id, pc, EndClass::Exit, cond, false); 427 AnalyzeCondInst(block, function_id, pc, EndClass::Exit, cond);
419 return AnalysisState::Branch; 428 return AnalysisState::Branch;
420 } 429 }
421 if (const std::optional<Location> exit_pc{block->stack.Peek(Token::PEXIT)}) { 430 if (const std::optional<Location> exit_pc{block->stack.Peek(Token::PEXIT)}) {
@@ -425,7 +434,7 @@ CFG::AnalysisState CFG::AnalyzeEXIT(Block* block, FunctionId function_id, Locati
425 block->branch_false = nullptr; 434 block->branch_false = nullptr;
426 return AnalysisState::Branch; 435 return AnalysisState::Branch;
427 } 436 }
428 block->end = pc; 437 block->end = pc + 1;
429 block->end_class = EndClass::Exit; 438 block->end_class = EndClass::Exit;
430 return AnalysisState::Branch; 439 return AnalysisState::Branch;
431} 440}
@@ -505,6 +514,12 @@ std::string CFG::Dot() const {
505 node_uid); 514 node_uid);
506 ++node_uid; 515 ++node_uid;
507 break; 516 break;
517 case EndClass::Kill:
518 dot += fmt::format("\t\t{}->N{};\n", name, node_uid);
519 dot += fmt::format("\t\tN{} [label=\"Kill\"][shape=square][style=stripped];\n",
520 node_uid);
521 ++node_uid;
522 break;
508 } 523 }
509 } 524 }
510 if (function.entrypoint == 8) { 525 if (function.entrypoint == 8) {
diff --git a/src/shader_recompiler/frontend/maxwell/control_flow.h b/src/shader_recompiler/frontend/maxwell/control_flow.h
index fe74f210f..22f134194 100644
--- a/src/shader_recompiler/frontend/maxwell/control_flow.h
+++ b/src/shader_recompiler/frontend/maxwell/control_flow.h
@@ -29,6 +29,7 @@ enum class EndClass {
29 Call, 29 Call,
30 Exit, 30 Exit,
31 Return, 31 Return,
32 Kill,
32}; 33};
33 34
34enum class Token { 35enum class Token {
@@ -130,7 +131,7 @@ private:
130 AnalysisState AnalyzeInst(Block* block, FunctionId function_id, Location pc); 131 AnalysisState AnalyzeInst(Block* block, FunctionId function_id, Location pc);
131 132
132 void AnalyzeCondInst(Block* block, FunctionId function_id, Location pc, EndClass insn_end_class, 133 void AnalyzeCondInst(Block* block, FunctionId function_id, Location pc, EndClass insn_end_class,
133 IR::Condition cond, bool visit_conditional_inst); 134 IR::Condition cond);
134 135
135 /// Return true when the branch instruction is confirmed to be a branch 136 /// Return true when the branch instruction is confirmed to be a branch
136 bool AnalyzeBranch(Block* block, FunctionId function_id, Location pc, Instruction inst, 137 bool AnalyzeBranch(Block* block, FunctionId function_id, Location pc, Instruction inst,
diff --git a/src/shader_recompiler/frontend/maxwell/program.cpp b/src/shader_recompiler/frontend/maxwell/program.cpp
index 8bfa64326..0074eb89b 100644
--- a/src/shader_recompiler/frontend/maxwell/program.cpp
+++ b/src/shader_recompiler/frontend/maxwell/program.cpp
@@ -32,6 +32,7 @@ IR::Program TranslateProgram(ObjectPool<IR::Inst>& inst_pool, ObjectPool<IR::Blo
32 IR::Program program; 32 IR::Program program;
33 program.blocks = VisitAST(inst_pool, block_pool, env, cfg); 33 program.blocks = VisitAST(inst_pool, block_pool, env, cfg);
34 program.post_order_blocks = PostOrder(program.blocks); 34 program.post_order_blocks = PostOrder(program.blocks);
35 program.stage = env.ShaderStage();
35 RemoveUnreachableBlocks(program); 36 RemoveUnreachableBlocks(program);
36 37
37 // Replace instructions before the SSA rewrite 38 // Replace instructions before the SSA rewrite
diff --git a/src/shader_recompiler/frontend/maxwell/structured_control_flow.cpp b/src/shader_recompiler/frontend/maxwell/structured_control_flow.cpp
index 5f5d9cf17..cec03e73e 100644
--- a/src/shader_recompiler/frontend/maxwell/structured_control_flow.cpp
+++ b/src/shader_recompiler/frontend/maxwell/structured_control_flow.cpp
@@ -45,6 +45,7 @@ enum class StatementType {
45 Loop, 45 Loop,
46 Break, 46 Break,
47 Return, 47 Return,
48 Kill,
48 Function, 49 Function,
49 Identity, 50 Identity,
50 Not, 51 Not,
@@ -70,6 +71,7 @@ struct If {};
70struct Loop {}; 71struct Loop {};
71struct Break {}; 72struct Break {};
72struct Return {}; 73struct Return {};
74struct Kill {};
73struct FunctionTag {}; 75struct FunctionTag {};
74struct Identity {}; 76struct Identity {};
75struct Not {}; 77struct Not {};
@@ -93,6 +95,7 @@ struct Statement : ListBaseHook {
93 Statement(Break, Statement* cond_, Statement* up_) 95 Statement(Break, Statement* cond_, Statement* up_)
94 : cond{cond_}, up{up_}, type{StatementType::Break} {} 96 : cond{cond_}, up{up_}, type{StatementType::Break} {}
95 Statement(Return) : type{StatementType::Return} {} 97 Statement(Return) : type{StatementType::Return} {}
98 Statement(Kill) : type{StatementType::Kill} {}
96 Statement(FunctionTag) : children{}, type{StatementType::Function} {} 99 Statement(FunctionTag) : children{}, type{StatementType::Function} {}
97 Statement(Identity, IR::Condition cond_) : guest_cond{cond_}, type{StatementType::Identity} {} 100 Statement(Identity, IR::Condition cond_) : guest_cond{cond_}, type{StatementType::Identity} {}
98 Statement(Not, Statement* op_) : op{op_}, type{StatementType::Not} {} 101 Statement(Not, Statement* op_) : op{op_}, type{StatementType::Not} {}
@@ -174,6 +177,9 @@ std::string DumpTree(const Tree& tree, u32 indentation = 0) {
174 case StatementType::Return: 177 case StatementType::Return:
175 ret += fmt::format("{} return;\n", indent); 178 ret += fmt::format("{} return;\n", indent);
176 break; 179 break;
180 case StatementType::Kill:
181 ret += fmt::format("{} kill;\n", indent);
182 break;
177 case StatementType::SetVariable: 183 case StatementType::SetVariable:
178 ret += fmt::format("{} goto_L{} = {};\n", indent, stmt->id, DumpExpr(stmt->op)); 184 ret += fmt::format("{} goto_L{} = {};\n", indent, stmt->id, DumpExpr(stmt->op));
179 break; 185 break;
@@ -424,6 +430,9 @@ private:
424 gotos.push_back(root.insert(ip, *goto_stmt)); 430 gotos.push_back(root.insert(ip, *goto_stmt));
425 break; 431 break;
426 } 432 }
433 case Flow::EndClass::Kill:
434 root.insert(ip, *pool.Create(Kill{}));
435 break;
427 } 436 }
428 } 437 }
429 } 438 }
@@ -729,6 +738,15 @@ private:
729 current_block = nullptr; 738 current_block = nullptr;
730 break; 739 break;
731 } 740 }
741 case StatementType::Kill: {
742 if (!current_block) {
743 current_block = block_pool.Create(inst_pool);
744 block_list.push_back(current_block);
745 }
746 IR::IREmitter{*current_block}.DemoteToHelperInvocation(continue_block);
747 current_block = nullptr;
748 break;
749 }
732 default: 750 default:
733 throw NotImplementedException("Statement type {}", stmt.type); 751 throw NotImplementedException("Statement type {}", stmt.type);
734 } 752 }
diff --git a/src/shader_recompiler/frontend/maxwell/translate/impl/exit.cpp b/src/shader_recompiler/frontend/maxwell/translate/impl/exit.cpp
deleted file mode 100644
index e98bbd0d1..000000000
--- a/src/shader_recompiler/frontend/maxwell/translate/impl/exit.cpp
+++ /dev/null
@@ -1,15 +0,0 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#include "common/common_types.h"
6#include "shader_recompiler/exception.h"
7#include "shader_recompiler/frontend/maxwell/translate/impl/impl.h"
8
9namespace Shader::Maxwell {
10
11void TranslatorVisitor::EXIT(u64) {
12 ir.Exit();
13}
14
15} // namespace Shader::Maxwell
diff --git a/src/shader_recompiler/frontend/maxwell/translate/impl/exit_program.cpp b/src/shader_recompiler/frontend/maxwell/translate/impl/exit_program.cpp
new file mode 100644
index 000000000..ea9b33da9
--- /dev/null
+++ b/src/shader_recompiler/frontend/maxwell/translate/impl/exit_program.cpp
@@ -0,0 +1,43 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#include "common/common_types.h"
6#include "shader_recompiler/exception.h"
7#include "shader_recompiler/frontend/maxwell/translate/impl/impl.h"
8
9namespace Shader::Maxwell {
10namespace {
11void ExitFragment(TranslatorVisitor& v) {
12 const ProgramHeader sph{v.env.SPH()};
13 IR::Reg src_reg{IR::Reg::R0};
14 for (u32 render_target = 0; render_target < 8; ++render_target) {
15 const std::array<bool, 4> mask{sph.ps.EnabledOutputComponents(render_target)};
16 for (u32 component = 0; component < 4; ++component) {
17 if (!mask[component]) {
18 continue;
19 }
20 v.ir.SetFragColor(render_target, component, v.F(src_reg));
21 ++src_reg;
22 }
23 }
24 if (sph.ps.omap.sample_mask != 0) {
25 throw NotImplementedException("Sample mask");
26 }
27 if (sph.ps.omap.depth != 0) {
28 throw NotImplementedException("Fragment depth");
29 }
30}
31} // Anonymous namespace
32
33void TranslatorVisitor::EXIT() {
34 switch (env.ShaderStage()) {
35 case Stage::Fragment:
36 ExitFragment(*this);
37 break;
38 default:
39 break;
40 }
41}
42
43} // namespace Shader::Maxwell
diff --git a/src/shader_recompiler/frontend/maxwell/translate/impl/impl.h b/src/shader_recompiler/frontend/maxwell/translate/impl/impl.h
index e3e298c3b..ed81d9c36 100644
--- a/src/shader_recompiler/frontend/maxwell/translate/impl/impl.h
+++ b/src/shader_recompiler/frontend/maxwell/translate/impl/impl.h
@@ -108,7 +108,7 @@ public:
108 void DSETP_reg(u64 insn); 108 void DSETP_reg(u64 insn);
109 void DSETP_cbuf(u64 insn); 109 void DSETP_cbuf(u64 insn);
110 void DSETP_imm(u64 insn); 110 void DSETP_imm(u64 insn);
111 void EXIT(u64 insn); 111 void EXIT();
112 void F2F_reg(u64 insn); 112 void F2F_reg(u64 insn);
113 void F2F_cbuf(u64 insn); 113 void F2F_cbuf(u64 insn);
114 void F2F_imm(u64 insn); 114 void F2F_imm(u64 insn);
@@ -220,7 +220,7 @@ public:
220 void JCAL(u64 insn); 220 void JCAL(u64 insn);
221 void JMP(u64 insn); 221 void JMP(u64 insn);
222 void JMX(u64 insn); 222 void JMX(u64 insn);
223 void KIL(u64 insn); 223 void KIL();
224 void LD(u64 insn); 224 void LD(u64 insn);
225 void LDC(u64 insn); 225 void LDC(u64 insn);
226 void LDG(u64 insn); 226 void LDG(u64 insn);
diff --git a/src/shader_recompiler/frontend/maxwell/translate/impl/load_store_attribute.cpp b/src/shader_recompiler/frontend/maxwell/translate/impl/load_store_attribute.cpp
index ad97786d4..2922145ee 100644
--- a/src/shader_recompiler/frontend/maxwell/translate/impl/load_store_attribute.cpp
+++ b/src/shader_recompiler/frontend/maxwell/translate/impl/load_store_attribute.cpp
@@ -11,6 +11,13 @@
11 11
12namespace Shader::Maxwell { 12namespace Shader::Maxwell {
13namespace { 13namespace {
14enum class Size : u64 {
15 B32,
16 B64,
17 B96,
18 B128,
19};
20
14enum class InterpolationMode : u64 { 21enum class InterpolationMode : u64 {
15 Pass, 22 Pass,
16 Multiply, 23 Multiply,
@@ -23,8 +30,85 @@ enum class SampleMode : u64 {
23 Centroid, 30 Centroid,
24 Offset, 31 Offset,
25}; 32};
33
34int NumElements(Size size) {
35 switch (size) {
36 case Size::B32:
37 return 1;
38 case Size::B64:
39 return 2;
40 case Size::B96:
41 return 3;
42 case Size::B128:
43 return 4;
44 }
45 throw InvalidArgument("Invalid size {}", size);
46}
26} // Anonymous namespace 47} // Anonymous namespace
27 48
49void TranslatorVisitor::ALD(u64 insn) {
50 union {
51 u64 raw;
52 BitField<0, 8, IR::Reg> dest_reg;
53 BitField<8, 8, IR::Reg> index_reg;
54 BitField<20, 10, u64> absolute_offset;
55 BitField<20, 11, s64> relative_offset;
56 BitField<39, 8, IR::Reg> stream_reg;
57 BitField<32, 1, u64> o;
58 BitField<31, 1, u64> patch;
59 BitField<47, 2, Size> size;
60 } const ald{insn};
61
62 if (ald.o != 0) {
63 throw NotImplementedException("O");
64 }
65 if (ald.patch != 0) {
66 throw NotImplementedException("P");
67 }
68 if (ald.index_reg != IR::Reg::RZ) {
69 throw NotImplementedException("Indexed");
70 }
71 const u64 offset{ald.absolute_offset.Value()};
72 if (offset % 4 != 0) {
73 throw NotImplementedException("Unaligned absolute offset {}", offset);
74 }
75 const int num_elements{NumElements(ald.size)};
76 for (int element = 0; element < num_elements; ++element) {
77 F(ald.dest_reg + element, ir.GetAttribute(IR::Attribute{offset / 4 + element}));
78 }
79}
80
81void TranslatorVisitor::AST(u64 insn) {
82 union {
83 u64 raw;
84 BitField<0, 8, IR::Reg> src_reg;
85 BitField<8, 8, IR::Reg> index_reg;
86 BitField<20, 10, u64> absolute_offset;
87 BitField<20, 11, s64> relative_offset;
88 BitField<31, 1, u64> patch;
89 BitField<39, 8, IR::Reg> stream_reg;
90 BitField<47, 2, Size> size;
91 } const ast{insn};
92
93 if (ast.patch != 0) {
94 throw NotImplementedException("P");
95 }
96 if (ast.stream_reg != IR::Reg::RZ) {
97 throw NotImplementedException("Stream store");
98 }
99 if (ast.index_reg != IR::Reg::RZ) {
100 throw NotImplementedException("Indexed store");
101 }
102 const u64 offset{ast.absolute_offset.Value()};
103 if (offset % 4 != 0) {
104 throw NotImplementedException("Unaligned absolute offset {}", offset);
105 }
106 const int num_elements{NumElements(ast.size)};
107 for (int element = 0; element < num_elements; ++element) {
108 ir.SetAttribute(IR::Attribute{offset / 4 + element}, F(ast.src_reg + element));
109 }
110}
111
28void TranslatorVisitor::IPA(u64 insn) { 112void TranslatorVisitor::IPA(u64 insn) {
29 // IPA is the instruction used to read varyings from a fragment shader. 113 // IPA is the instruction used to read varyings from a fragment shader.
30 // gl_FragCoord is mapped to the gl_Position attribute. 114 // gl_FragCoord is mapped to the gl_Position attribute.
@@ -51,7 +135,7 @@ void TranslatorVisitor::IPA(u64 insn) {
51 // } 135 // }
52 const bool is_indexed{ipa.idx != 0 && ipa.index_reg != IR::Reg::RZ}; 136 const bool is_indexed{ipa.idx != 0 && ipa.index_reg != IR::Reg::RZ};
53 if (is_indexed) { 137 if (is_indexed) {
54 throw NotImplementedException("IPA.IDX"); 138 throw NotImplementedException("IDX");
55 } 139 }
56 140
57 const IR::Attribute attribute{ipa.attribute}; 141 const IR::Attribute attribute{ipa.attribute};
diff --git a/src/shader_recompiler/frontend/maxwell/translate/impl/not_implemented.cpp b/src/shader_recompiler/frontend/maxwell/translate/impl/not_implemented.cpp
index 9675cef54..59252bcc5 100644
--- a/src/shader_recompiler/frontend/maxwell/translate/impl/not_implemented.cpp
+++ b/src/shader_recompiler/frontend/maxwell/translate/impl/not_implemented.cpp
@@ -17,14 +17,6 @@ void TranslatorVisitor::AL2P(u64) {
17 ThrowNotImplemented(Opcode::AL2P); 17 ThrowNotImplemented(Opcode::AL2P);
18} 18}
19 19
20void TranslatorVisitor::ALD(u64) {
21 ThrowNotImplemented(Opcode::ALD);
22}
23
24void TranslatorVisitor::AST(u64) {
25 ThrowNotImplemented(Opcode::AST);
26}
27
28void TranslatorVisitor::ATOM_cas(u64) { 20void TranslatorVisitor::ATOM_cas(u64) {
29 ThrowNotImplemented(Opcode::ATOM_cas); 21 ThrowNotImplemented(Opcode::ATOM_cas);
30} 22}
@@ -153,10 +145,6 @@ void TranslatorVisitor::DSETP_imm(u64) {
153 ThrowNotImplemented(Opcode::DSETP_imm); 145 ThrowNotImplemented(Opcode::DSETP_imm);
154} 146}
155 147
156void TranslatorVisitor::EXIT(u64) {
157 throw LogicError("Visting EXIT instruction");
158}
159
160void TranslatorVisitor::F2F_reg(u64) { 148void TranslatorVisitor::F2F_reg(u64) {
161 ThrowNotImplemented(Opcode::F2F_reg); 149 ThrowNotImplemented(Opcode::F2F_reg);
162} 150}
@@ -345,8 +333,8 @@ void TranslatorVisitor::JMX(u64) {
345 ThrowNotImplemented(Opcode::JMX); 333 ThrowNotImplemented(Opcode::JMX);
346} 334}
347 335
348void TranslatorVisitor::KIL(u64) { 336void TranslatorVisitor::KIL() {
349 ThrowNotImplemented(Opcode::KIL); 337 // KIL is a no-op
350} 338}
351 339
352void TranslatorVisitor::LD(u64) { 340void TranslatorVisitor::LD(u64) {
diff --git a/src/shader_recompiler/frontend/maxwell/translate/impl/texture_fetch.cpp b/src/shader_recompiler/frontend/maxwell/translate/impl/texture_fetch.cpp
index 98d9f4c64..0fbb87ec4 100644
--- a/src/shader_recompiler/frontend/maxwell/translate/impl/texture_fetch.cpp
+++ b/src/shader_recompiler/frontend/maxwell/translate/impl/texture_fetch.cpp
@@ -215,7 +215,7 @@ void TranslatorVisitor::TEX(u64 insn) {
215 BitField<36, 13, u64> cbuf_offset; 215 BitField<36, 13, u64> cbuf_offset;
216 } const tex{insn}; 216 } const tex{insn};
217 217
218 Impl(*this, insn, tex.aoffi != 0, tex.blod, tex.lc != 0, static_cast<u32>(tex.cbuf_offset)); 218 Impl(*this, insn, tex.aoffi != 0, tex.blod, tex.lc != 0, static_cast<u32>(tex.cbuf_offset * 4));
219} 219}
220 220
221void TranslatorVisitor::TEX_b(u64 insn) { 221void TranslatorVisitor::TEX_b(u64 insn) {
diff --git a/src/shader_recompiler/frontend/maxwell/translate/impl/texture_fetch_swizzled.cpp b/src/shader_recompiler/frontend/maxwell/translate/impl/texture_fetch_swizzled.cpp
index ac1615b00..54f0df754 100644
--- a/src/shader_recompiler/frontend/maxwell/translate/impl/texture_fetch_swizzled.cpp
+++ b/src/shader_recompiler/frontend/maxwell/translate/impl/texture_fetch_swizzled.cpp
@@ -70,7 +70,7 @@ IR::F32 ReadArray(TranslatorVisitor& v, const IR::U32& value) {
70 70
71IR::Value Sample(TranslatorVisitor& v, u64 insn) { 71IR::Value Sample(TranslatorVisitor& v, u64 insn) {
72 const Encoding texs{insn}; 72 const Encoding texs{insn};
73 const IR::U32 handle{v.ir.Imm32(static_cast<u32>(texs.cbuf_offset))}; 73 const IR::U32 handle{v.ir.Imm32(static_cast<u32>(texs.cbuf_offset * 4))};
74 const IR::F32 zero{v.ir.Imm32(0.0f)}; 74 const IR::F32 zero{v.ir.Imm32(0.0f)};
75 const IR::Reg reg_a{texs.src_reg_a}; 75 const IR::Reg reg_a{texs.src_reg_a};
76 const IR::Reg reg_b{texs.src_reg_b}; 76 const IR::Reg reg_b{texs.src_reg_b};
diff --git a/src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp b/src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp
index 708b6b267..fbbe28632 100644
--- a/src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp
+++ b/src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp
@@ -17,10 +17,47 @@ void AddConstantBufferDescriptor(Info& info, u32 index, u32 count) {
17 return; 17 return;
18 } 18 }
19 info.constant_buffer_mask |= 1U << index; 19 info.constant_buffer_mask |= 1U << index;
20 info.constant_buffer_descriptors.push_back({ 20
21 .index{index}, 21 auto& cbufs{info.constant_buffer_descriptors};
22 .count{1}, 22 cbufs.insert(std::ranges::lower_bound(cbufs, index, {}, &ConstantBufferDescriptor::index),
23 }); 23 ConstantBufferDescriptor{
24 .index{index},
25 .count{1},
26 });
27}
28
29void GetAttribute(Info& info, IR::Attribute attribute) {
30 if (IR::IsGeneric(attribute)) {
31 info.loads_generics.at(IR::GenericAttributeIndex(attribute)) = true;
32 return;
33 }
34 switch (attribute) {
35 case IR::Attribute::PositionX:
36 case IR::Attribute::PositionY:
37 case IR::Attribute::PositionZ:
38 case IR::Attribute::PositionW:
39 info.loads_position = true;
40 break;
41 default:
42 throw NotImplementedException("Get attribute {}", attribute);
43 }
44}
45
46void SetAttribute(Info& info, IR::Attribute attribute) {
47 if (IR::IsGeneric(attribute)) {
48 info.stores_generics.at(IR::GenericAttributeIndex(attribute)) = true;
49 return;
50 }
51 switch (attribute) {
52 case IR::Attribute::PositionX:
53 case IR::Attribute::PositionY:
54 case IR::Attribute::PositionZ:
55 case IR::Attribute::PositionW:
56 info.stores_position = true;
57 break;
58 default:
59 throw NotImplementedException("Set attribute {}", attribute);
60 }
24} 61}
25 62
26void VisitUsages(Info& info, IR::Inst& inst) { 63void VisitUsages(Info& info, IR::Inst& inst) {
@@ -162,6 +199,21 @@ void VisitUsages(Info& info, IR::Inst& inst) {
162 break; 199 break;
163 } 200 }
164 switch (inst.Opcode()) { 201 switch (inst.Opcode()) {
202 case IR::Opcode::DemoteToHelperInvocation:
203 info.uses_demote_to_helper_invocation = true;
204 break;
205 case IR::Opcode::GetAttribute:
206 GetAttribute(info, inst.Arg(0).Attribute());
207 break;
208 case IR::Opcode::SetAttribute:
209 SetAttribute(info, inst.Arg(0).Attribute());
210 break;
211 case IR::Opcode::SetFragColor:
212 info.stores_frag_color[inst.Arg(0).U32()] = true;
213 break;
214 case IR::Opcode::SetFragDepth:
215 info.stores_frag_depth = true;
216 break;
165 case IR::Opcode::WorkgroupId: 217 case IR::Opcode::WorkgroupId:
166 info.uses_workgroup_id = true; 218 info.uses_workgroup_id = true;
167 break; 219 break;
diff --git a/src/shader_recompiler/ir_opt/ssa_rewrite_pass.cpp b/src/shader_recompiler/ir_opt/ssa_rewrite_pass.cpp
index d09bcec36..bab7ca186 100644
--- a/src/shader_recompiler/ir_opt/ssa_rewrite_pass.cpp
+++ b/src/shader_recompiler/ir_opt/ssa_rewrite_pass.cpp
@@ -169,7 +169,7 @@ private:
169 const size_t num_args{phi.NumArgs()}; 169 const size_t num_args{phi.NumArgs()};
170 for (size_t arg_index = 0; arg_index < num_args; ++arg_index) { 170 for (size_t arg_index = 0; arg_index < num_args; ++arg_index) {
171 const IR::Value& op{phi.Arg(arg_index)}; 171 const IR::Value& op{phi.Arg(arg_index)};
172 if (op == same || op == IR::Value{&phi}) { 172 if (op.Resolve() == same.Resolve() || op == IR::Value{&phi}) {
173 // Unique value or self-reference 173 // Unique value or self-reference
174 continue; 174 continue;
175 } 175 }
diff --git a/src/shader_recompiler/program_header.h b/src/shader_recompiler/program_header.h
new file mode 100644
index 000000000..1544bfa42
--- /dev/null
+++ b/src/shader_recompiler/program_header.h
@@ -0,0 +1,143 @@
1// Copyright 2018 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include <array>
8#include <optional>
9
10#include "common/bit_field.h"
11#include "common/common_funcs.h"
12#include "common/common_types.h"
13
14namespace Shader {
15
16enum class OutputTopology : u32 {
17 PointList = 1,
18 LineStrip = 6,
19 TriangleStrip = 7,
20};
21
22enum class PixelImap : u8 {
23 Unused = 0,
24 Constant = 1,
25 Perspective = 2,
26 ScreenLinear = 3,
27};
28
29// Documentation in:
30// http://download.nvidia.com/open-gpu-doc/Shader-Program-Header/1/Shader-Program-Header.html
31struct ProgramHeader {
32 union {
33 BitField<0, 5, u32> sph_type;
34 BitField<5, 5, u32> version;
35 BitField<10, 4, u32> shader_type;
36 BitField<14, 1, u32> mrt_enable;
37 BitField<15, 1, u32> kills_pixels;
38 BitField<16, 1, u32> does_global_store;
39 BitField<17, 4, u32> sass_version;
40 BitField<21, 5, u32> reserved;
41 BitField<26, 1, u32> does_load_or_store;
42 BitField<27, 1, u32> does_fp64;
43 BitField<28, 4, u32> stream_out_mask;
44 } common0;
45
46 union {
47 BitField<0, 24, u32> shader_local_memory_low_size;
48 BitField<24, 8, u32> per_patch_attribute_count;
49 } common1;
50
51 union {
52 BitField<0, 24, u32> shader_local_memory_high_size;
53 BitField<24, 8, u32> threads_per_input_primitive;
54 } common2;
55
56 union {
57 BitField<0, 24, u32> shader_local_memory_crs_size;
58 BitField<24, 4, OutputTopology> output_topology;
59 BitField<28, 4, u32> reserved;
60 } common3;
61
62 union {
63 BitField<0, 12, u32> max_output_vertices;
64 BitField<12, 8, u32> store_req_start; // NOTE: not used by geometry shaders.
65 BitField<20, 4, u32> reserved;
66 BitField<24, 8, u32> store_req_end; // NOTE: not used by geometry shaders.
67 } common4;
68
69 union {
70 struct {
71 INSERT_PADDING_BYTES_NOINIT(3); // ImapSystemValuesA
72 INSERT_PADDING_BYTES_NOINIT(1); // ImapSystemValuesB
73 INSERT_PADDING_BYTES_NOINIT(16); // ImapGenericVector[32]
74 INSERT_PADDING_BYTES_NOINIT(2); // ImapColor
75 union {
76 BitField<0, 8, u16> clip_distances;
77 BitField<8, 1, u16> point_sprite_s;
78 BitField<9, 1, u16> point_sprite_t;
79 BitField<10, 1, u16> fog_coordinate;
80 BitField<12, 1, u16> tessellation_eval_point_u;
81 BitField<13, 1, u16> tessellation_eval_point_v;
82 BitField<14, 1, u16> instance_id;
83 BitField<15, 1, u16> vertex_id;
84 };
85 INSERT_PADDING_BYTES_NOINIT(5); // ImapFixedFncTexture[10]
86 INSERT_PADDING_BYTES_NOINIT(1); // ImapReserved
87 INSERT_PADDING_BYTES_NOINIT(3); // OmapSystemValuesA
88 INSERT_PADDING_BYTES_NOINIT(1); // OmapSystemValuesB
89 INSERT_PADDING_BYTES_NOINIT(16); // OmapGenericVector[32]
90 INSERT_PADDING_BYTES_NOINIT(2); // OmapColor
91 INSERT_PADDING_BYTES_NOINIT(2); // OmapSystemValuesC
92 INSERT_PADDING_BYTES_NOINIT(5); // OmapFixedFncTexture[10]
93 INSERT_PADDING_BYTES_NOINIT(1); // OmapReserved
94 } vtg;
95
96 struct {
97 INSERT_PADDING_BYTES_NOINIT(3); // ImapSystemValuesA
98 INSERT_PADDING_BYTES_NOINIT(1); // ImapSystemValuesB
99
100 union {
101 BitField<0, 2, PixelImap> x;
102 BitField<2, 2, PixelImap> y;
103 BitField<4, 2, PixelImap> z;
104 BitField<6, 2, PixelImap> w;
105 u8 raw;
106 } imap_generic_vector[32];
107
108 INSERT_PADDING_BYTES_NOINIT(2); // ImapColor
109 INSERT_PADDING_BYTES_NOINIT(2); // ImapSystemValuesC
110 INSERT_PADDING_BYTES_NOINIT(10); // ImapFixedFncTexture[10]
111 INSERT_PADDING_BYTES_NOINIT(2); // ImapReserved
112
113 struct {
114 u32 target;
115 union {
116 BitField<0, 1, u32> sample_mask;
117 BitField<1, 1, u32> depth;
118 BitField<2, 30, u32> reserved;
119 };
120 } omap;
121
122 [[nodiscard]] std::array<bool, 4> EnabledOutputComponents(u32 rt) const noexcept {
123 const u32 bits{omap.target >> (rt * 4)};
124 return {(bits & 1) != 0, (bits & 2) != 0, (bits & 4) != 0, (bits & 8) != 0};
125 }
126
127 [[nodiscard]] std::array<PixelImap, 4> GenericInputMap(u32 attribute) const {
128 const auto& vector{imap_generic_vector[attribute]};
129 return {vector.x, vector.y, vector.z, vector.w};
130 }
131 } ps;
132
133 std::array<u32, 0xf> raw;
134 };
135
136 [[nodiscard]] u64 LocalMemorySize() const noexcept {
137 return (common1.shader_local_memory_low_size |
138 (common2.shader_local_memory_high_size << 24));
139 }
140};
141static_assert(sizeof(ProgramHeader) == 0x50, "Incorrect structure size");
142
143} // namespace Shader
diff --git a/src/shader_recompiler/recompiler.cpp b/src/shader_recompiler/recompiler.cpp
deleted file mode 100644
index 527e19c27..000000000
--- a/src/shader_recompiler/recompiler.cpp
+++ /dev/null
@@ -1,28 +0,0 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#include <vector>
6
7#include "common/common_types.h"
8#include "shader_recompiler/backend/spirv/emit_spirv.h"
9#include "shader_recompiler/environment.h"
10#include "shader_recompiler/frontend/maxwell/control_flow.h"
11#include "shader_recompiler/frontend/maxwell/program.h"
12#include "shader_recompiler/object_pool.h"
13#include "shader_recompiler/recompiler.h"
14
15namespace Shader {
16
17std::pair<Info, std::vector<u32>> RecompileSPIRV(const Profile& profile, Environment& env,
18 u32 start_address) {
19 ObjectPool<Maxwell::Flow::Block> flow_block_pool;
20 ObjectPool<IR::Inst> inst_pool;
21 ObjectPool<IR::Block> block_pool;
22
23 Maxwell::Flow::CFG cfg{env, flow_block_pool, start_address};
24 IR::Program program{Maxwell::TranslateProgram(inst_pool, block_pool, env, cfg)};
25 return {std::move(program.info), Backend::SPIRV::EmitSPIRV(profile, env, program)};
26}
27
28} // namespace Shader
diff --git a/src/shader_recompiler/recompiler.h b/src/shader_recompiler/recompiler.h
deleted file mode 100644
index 2529463ae..000000000
--- a/src/shader_recompiler/recompiler.h
+++ /dev/null
@@ -1,20 +0,0 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include <utility>
8#include <vector>
9
10#include "common/common_types.h"
11#include "shader_recompiler/environment.h"
12#include "shader_recompiler/profile.h"
13#include "shader_recompiler/shader_info.h"
14
15namespace Shader {
16
17[[nodiscard]] std::pair<Info, std::vector<u32>> RecompileSPIRV(const Profile& profile,
18 Environment& env, u32 start_address);
19
20} // namespace Shader
diff --git a/src/shader_recompiler/shader_info.h b/src/shader_recompiler/shader_info.h
index adc1d9a64..6eff762e2 100644
--- a/src/shader_recompiler/shader_info.h
+++ b/src/shader_recompiler/shader_info.h
@@ -56,6 +56,15 @@ struct Info {
56 56
57 bool uses_workgroup_id{}; 57 bool uses_workgroup_id{};
58 bool uses_local_invocation_id{}; 58 bool uses_local_invocation_id{};
59
60 std::array<bool, 32> loads_generics{};
61 bool loads_position{};
62
63 std::array<bool, 8> stores_frag_color{};
64 bool stores_frag_depth{};
65 std::array<bool, 32> stores_generics{};
66 bool stores_position{};
67
59 bool uses_fp16{}; 68 bool uses_fp16{};
60 bool uses_fp64{}; 69 bool uses_fp64{};
61 bool uses_fp16_denorms_flush{}; 70 bool uses_fp16_denorms_flush{};
@@ -68,6 +77,7 @@ struct Info {
68 bool uses_image_1d{}; 77 bool uses_image_1d{};
69 bool uses_sampled_1d{}; 78 bool uses_sampled_1d{};
70 bool uses_sparse_residency{}; 79 bool uses_sparse_residency{};
80 bool uses_demote_to_helper_invocation{};
71 81
72 IR::Type used_constant_buffer_types{}; 82 IR::Type used_constant_buffer_types{};
73 83
diff --git a/src/shader_recompiler/stage.h b/src/shader_recompiler/stage.h
new file mode 100644
index 000000000..fc6ce6043
--- /dev/null
+++ b/src/shader_recompiler/stage.h
@@ -0,0 +1,19 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7namespace Shader {
8
9enum class Stage {
10 Compute,
11 VertexA,
12 VertexB,
13 TessellationControl,
14 TessellationEval,
15 Geometry,
16 Fragment,
17};
18
19} // namespace Shader
HIC SVNT DRACONES