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
     frontend/maxwell/translate/impl/common_funcs.h
     frontend/maxwell/translate/impl/condition_code_set.cpp
     frontend/maxwell/translate/impl/double_add.cpp
+    frontend/maxwell/translate/impl/exit_program.cpp
     frontend/maxwell/translate/impl/find_leading_one.cpp
     frontend/maxwell/translate/impl/floating_point_add.cpp
     frontend/maxwell/translate/impl/floating_point_compare.cpp
@@ -121,9 +122,8 @@ add_library(shader_recompiler STATIC
     ir_opt/texture_pass.cpp
     ir_opt/verification_pass.cpp
     object_pool.h
+    program_header.h
     profile.h
-    recompiler.cpp
-    recompiler.h
     shader_info.h
 )
 
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
     }
 }
 
-EmitContext::EmitContext(const Profile& profile_, IR::Program& program)
+EmitContext::EmitContext(const Profile& profile_, IR::Program& program, u32& binding)
     : Sirit::Module(0x00010000), profile{profile_} {
     AddCapability(spv::Capability::Shader);
     DefineCommonTypes(program.info);
     DefineCommonConstants();
-    DefineSpecialVariables(program.info);
-
-    u32 binding{};
+    DefineInterfaces(program.info, program.stage);
     DefineConstantBuffers(program.info, binding);
     DefineStorageBuffers(program.info, binding);
     DefineTextures(program.info, binding);
-
     DefineLabels(program);
 }
 
@@ -96,6 +93,8 @@ Id EmitContext::Def(const IR::Value& value) {
         return Constant(F32[1], value.F32());
     case IR::Type::F64:
         return Constant(F64[1], value.F64());
+    case IR::Type::Label:
+        return value.Label()->Definition<Id>();
     default:
         throw NotImplementedException("Immediate type {}", value.Type());
     }
@@ -109,6 +108,9 @@ void EmitContext::DefineCommonTypes(const Info& info) {
     F32.Define(*this, TypeFloat(32), "f32");
     U32.Define(*this, TypeInt(32, false), "u32");
 
+    input_f32 = Name(TypePointer(spv::StorageClass::Input, F32[1]), "input_f32");
+    output_f32 = Name(TypePointer(spv::StorageClass::Output, F32[1]), "output_f32");
+
     if (info.uses_int8) {
         AddCapability(spv::Capability::Int8);
         U8 = Name(TypeInt(8, false), "u8");
@@ -139,15 +141,20 @@ void EmitContext::DefineCommonConstants() {
     u32_zero_value = Constant(U32[1], 0U);
 }
 
-void EmitContext::DefineSpecialVariables(const Info& info) {
-    const auto define{[this](Id type, spv::BuiltIn builtin, spv::StorageClass storage_class) {
-        const Id pointer_type{TypePointer(storage_class, type)};
-        const Id id{AddGlobalVariable(pointer_type, spv::StorageClass::Input)};
-        Decorate(id, spv::Decoration::BuiltIn, builtin);
-        return id;
-    }};
+void EmitContext::DefineInterfaces(const Info& info, Stage stage) {
+    const auto define{
+        [this](Id type, std::optional<spv::BuiltIn> builtin, spv::StorageClass storage_class) {
+            const Id pointer_type{TypePointer(storage_class, type)};
+            const Id id{AddGlobalVariable(pointer_type, storage_class)};
+            if (builtin) {
+                Decorate(id, spv::Decoration::BuiltIn, *builtin);
+            }
+            interfaces.push_back(id);
+            return id;
+        }};
     using namespace std::placeholders;
     const auto define_input{std::bind(define, _1, _2, spv::StorageClass::Input)};
+    const auto define_output{std::bind(define, _1, _2, spv::StorageClass::Output)};
 
     if (info.uses_workgroup_id) {
         workgroup_id = define_input(U32[3], spv::BuiltIn::WorkgroupId);
@@ -155,6 +162,39 @@ void EmitContext::DefineSpecialVariables(const Info& info) {
     if (info.uses_local_invocation_id) {
         local_invocation_id = define_input(U32[3], spv::BuiltIn::LocalInvocationId);
     }
+    if (info.loads_position) {
+        const bool is_fragment{stage != Stage::Fragment};
+        const spv::BuiltIn built_in{is_fragment ? spv::BuiltIn::Position : spv::BuiltIn::FragCoord};
+        input_position = define_input(F32[4], built_in);
+    }
+    for (size_t i = 0; i < info.loads_generics.size(); ++i) {
+        if (info.loads_generics[i]) {
+            // FIXME: Declare size from input
+            input_generics[i] = define_input(F32[4], std::nullopt);
+            Decorate(input_generics[i], spv::Decoration::Location, static_cast<u32>(i));
+            Name(input_generics[i], fmt::format("in_attr{}", i));
+        }
+    }
+    if (info.stores_position) {
+        output_position = define_output(F32[4], spv::BuiltIn::Position);
+    }
+    for (size_t i = 0; i < info.stores_generics.size(); ++i) {
+        if (info.stores_generics[i]) {
+            output_generics[i] = define_output(F32[4], std::nullopt);
+            Decorate(output_generics[i], spv::Decoration::Location, static_cast<u32>(i));
+            Name(output_generics[i], fmt::format("out_attr{}", i));
+        }
+    }
+    if (stage == Stage::Fragment) {
+        for (size_t i = 0; i < 8; ++i) {
+            if (!info.stores_frag_color[i]) {
+                continue;
+            }
+            frag_color[i] = define_output(F32[4], std::nullopt);
+            Decorate(frag_color[i], spv::Decoration::Location, static_cast<u32>(i));
+            Name(frag_color[i], fmt::format("frag_color{}", i));
+        }
+    }
 }
 
 void 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 {
 
 class EmitContext final : public Sirit::Module {
 public:
-    explicit EmitContext(const Profile& profile, IR::Program& program);
+    explicit EmitContext(const Profile& profile, IR::Program& program, u32& binding);
     ~EmitContext();
 
     [[nodiscard]] Id Def(const IR::Value& value);
@@ -71,6 +71,9 @@ public:
 
     UniformDefinitions uniform_types;
 
+    Id input_f32{};
+    Id output_f32{};
+
     Id storage_u32{};
 
     std::array<UniformDefinitions, Info::MAX_CBUFS> cbufs{};
@@ -80,10 +83,21 @@ public:
     Id workgroup_id{};
     Id local_invocation_id{};
 
+    Id input_position{};
+    std::array<Id, 32> input_generics{};
+
+    Id output_position{};
+    std::array<Id, 32> output_generics{};
+
+    std::array<Id, 8> frag_color{};
+    Id frag_depth {};
+
+    std::vector<Id> interfaces;
+
 private:
     void DefineCommonTypes(const Info& info);
     void DefineCommonConstants();
-    void DefineSpecialVariables(const Info& info);
+    void DefineInterfaces(const Info& info, Stage stage);
     void DefineConstantBuffers(const Info& info, u32& binding);
     void DefineConstantBuffers(const Info& info, Id UniformDefinitions::*member_type, u32 binding,
                                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) {
         return arg.U32();
     } else if constexpr (std::is_same_v<ArgType, IR::Block*>) {
         return arg.Label();
+    } else if constexpr (std::is_same_v<ArgType, IR::Attribute>) {
+        return arg.Attribute();
     }
 }
 
@@ -197,8 +199,9 @@ Id PhiArgDef(EmitContext& ctx, IR::Inst* inst, size_t index) {
 }
 } // Anonymous namespace
 
-std::vector<u32> EmitSPIRV(const Profile& profile, Environment& env, IR::Program& program) {
-    EmitContext ctx{profile, program};
+std::vector<u32> 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<u32> EmitSPIRV(const Profile& profile, Environment& env, IR::Program
         }
     }
     ctx.OpFunctionEnd();
-    boost::container::small_vector<Id, 32> 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<u32, 3> 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<u32, 3> 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);
 
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 @@
 namespace Shader::Backend::SPIRV {
 
 [[nodiscard]] std::vector<u32> EmitSPIRV(const Profile& profile, Environment& env,
-                                         IR::Program& program);
+                                         IR::Program& program, u32& binding);
 
 // Microinstruction emitters
 Id EmitPhi(EmitContext& ctx, IR::Inst* inst);
 void EmitVoid(EmitContext& ctx);
 Id EmitIdentity(EmitContext& ctx, const IR::Value& value);
-void EmitBranch(EmitContext& ctx, IR::Block* label);
-void EmitBranchConditional(EmitContext& ctx, Id condition, IR::Block* true_label,
-                           IR::Block* false_label);
-void EmitLoopMerge(EmitContext& ctx, IR::Block* merge_label, IR::Block* continue_label);
-void EmitSelectionMerge(EmitContext& ctx, IR::Block* merge_label);
+void EmitBranch(EmitContext& ctx, Id label);
+void EmitBranchConditional(EmitContext& ctx, Id condition, Id true_label, Id false_label);
+void EmitLoopMerge(EmitContext& ctx, Id merge_label, Id continue_label);
+void EmitSelectionMerge(EmitContext& ctx, Id merge_label);
 void EmitReturn(EmitContext& ctx);
+void EmitDemoteToHelperInvocation(EmitContext& ctx, Id continue_label);
 void EmitGetRegister(EmitContext& ctx);
 void EmitSetRegister(EmitContext& ctx);
 void EmitGetPred(EmitContext& ctx);
@@ -41,10 +41,12 @@ Id EmitGetCbufS16(EmitContext& ctx, const IR::Value& binding, const IR::Value& o
 Id EmitGetCbufU32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
 Id EmitGetCbufF32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
 Id EmitGetCbufU64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
-void EmitGetAttribute(EmitContext& ctx);
-void EmitSetAttribute(EmitContext& ctx);
+Id EmitGetAttribute(EmitContext& ctx, IR::Attribute attr);
+void EmitSetAttribute(EmitContext& ctx, IR::Attribute attr, Id value);
 void EmitGetAttributeIndexed(EmitContext& ctx);
 void EmitSetAttributeIndexed(EmitContext& ctx);
+void EmitSetFragColor(EmitContext& ctx, u32 index, u32 component, Id value);
+void EmitSetFragDepth(EmitContext& ctx, Id value);
 void EmitGetZFlag(EmitContext& ctx);
 void EmitGetSFlag(EmitContext& ctx);
 void 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 @@
 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 
 namespace Shader::Backend::SPIRV {
+namespace {
+Id InputAttrPointer(EmitContext& ctx, IR::Attribute attr) {
+    const u32 element{static_cast<u32>(attr) % 4};
+    const auto element_id{[&] { return ctx.Constant(ctx.U32[1], element); }};
+    if (IR::IsGeneric(attr)) {
+        const u32 index{IR::GenericAttributeIndex(attr)};
+        return ctx.OpAccessChain(ctx.input_f32, ctx.input_generics.at(index), element_id());
+    }
+    switch (attr) {
+    case IR::Attribute::PositionX:
+    case IR::Attribute::PositionY:
+    case IR::Attribute::PositionZ:
+    case IR::Attribute::PositionW:
+        return ctx.OpAccessChain(ctx.input_f32, ctx.input_position, element_id());
+    default:
+        throw NotImplementedException("Read attribute {}", attr);
+    }
+}
+
+Id OutputAttrPointer(EmitContext& ctx, IR::Attribute attr) {
+    const u32 element{static_cast<u32>(attr) % 4};
+    const auto element_id{[&] { return ctx.Constant(ctx.U32[1], element); }};
+    if (IR::IsGeneric(attr)) {
+        const u32 index{IR::GenericAttributeIndex(attr)};
+        return ctx.OpAccessChain(ctx.output_f32, ctx.output_generics.at(index), element_id());
+    }
+    switch (attr) {
+    case IR::Attribute::PositionX:
+    case IR::Attribute::PositionY:
+    case IR::Attribute::PositionZ:
+    case IR::Attribute::PositionW:
+        return ctx.OpAccessChain(ctx.output_f32, ctx.output_position, element_id());
+    default:
+        throw NotImplementedException("Read attribute {}", attr);
+    }
+}
+} // Anonymous namespace
 
 void EmitGetRegister(EmitContext&) {
     throw NotImplementedException("SPIR-V Instruction");
@@ -87,12 +124,12 @@ Id EmitGetCbufU64(EmitContext& ctx, const IR::Value& binding, const IR::Value& o
     return GetCbuf(ctx, ctx.U64, &UniformDefinitions::U64, sizeof(u64), binding, offset);
 }
 
-void EmitGetAttribute(EmitContext&) {
-    throw NotImplementedException("SPIR-V Instruction");
+Id EmitGetAttribute(EmitContext& ctx, IR::Attribute attr) {
+    return ctx.OpLoad(ctx.F32[1], InputAttrPointer(ctx, attr));
 }
 
-void EmitSetAttribute(EmitContext&) {
-    throw NotImplementedException("SPIR-V Instruction");
+void EmitSetAttribute(EmitContext& ctx, IR::Attribute attr, Id value) {
+    ctx.OpStore(OutputAttrPointer(ctx, attr), value);
 }
 
 void EmitGetAttributeIndexed(EmitContext&) {
@@ -103,6 +140,16 @@ void EmitSetAttributeIndexed(EmitContext&) {
     throw NotImplementedException("SPIR-V Instruction");
 }
 
+void EmitSetFragColor(EmitContext& ctx, u32 index, u32 component, Id value) {
+    const Id component_id{ctx.Constant(ctx.U32[1], component)};
+    const Id pointer{ctx.OpAccessChain(ctx.output_f32, ctx.frag_color.at(index), component_id)};
+    ctx.OpStore(pointer, value);
+}
+
+void EmitSetFragDepth(EmitContext& ctx, Id value) {
+    ctx.OpStore(ctx.frag_depth, value);
+}
+
 void EmitGetZFlag(EmitContext&) {
     throw NotImplementedException("SPIR-V Instruction");
 }
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 @@
 
 namespace Shader::Backend::SPIRV {
 
-void EmitBranch(EmitContext& ctx, IR::Block* label) {
-    ctx.OpBranch(label->Definition<Id>());
+void EmitBranch(EmitContext& ctx, Id label) {
+    ctx.OpBranch(label);
 }
 
-void EmitBranchConditional(EmitContext& ctx, Id condition, IR::Block* true_label,
-                           IR::Block* false_label) {
-    ctx.OpBranchConditional(condition, true_label->Definition<Id>(), false_label->Definition<Id>());
+void EmitBranchConditional(EmitContext& ctx, Id condition, Id true_label, Id false_label) {
+    ctx.OpBranchConditional(condition, true_label, false_label);
 }
 
-void EmitLoopMerge(EmitContext& ctx, IR::Block* merge_label, IR::Block* continue_label) {
-    ctx.OpLoopMerge(merge_label->Definition<Id>(), continue_label->Definition<Id>(),
-                    spv::LoopControlMask::MaskNone);
+void EmitLoopMerge(EmitContext& ctx, Id merge_label, Id continue_label) {
+    ctx.OpLoopMerge(merge_label, continue_label, spv::LoopControlMask::MaskNone);
 }
 
-void EmitSelectionMerge(EmitContext& ctx, IR::Block* merge_label) {
-    ctx.OpSelectionMerge(merge_label->Definition<Id>(), spv::SelectionControlMask::MaskNone);
+void EmitSelectionMerge(EmitContext& ctx, Id merge_label) {
+    ctx.OpSelectionMerge(merge_label, spv::SelectionControlMask::MaskNone);
 }
 
 void EmitReturn(EmitContext& ctx) {
     ctx.OpReturn();
 }
 
+void EmitDemoteToHelperInvocation(EmitContext& ctx, Id continue_label) {
+    ctx.OpDemoteToHelperInvocationEXT();
+    ctx.OpBranch(continue_label);
+}
+
 } // 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 @@
 #include <array>
 
 #include "common/common_types.h"
+#include "shader_recompiler/stage.h"
+#include "shader_recompiler/program_header.h"
 
 namespace Shader {
 
@@ -15,6 +17,18 @@ public:
     [[nodiscard]] virtual u32 TextureBoundBuffer() = 0;
 
     [[nodiscard]] virtual std::array<u32, 3> WorkgroupSize() = 0;
+
+    [[nodiscard]] const ProgramHeader& SPH() const noexcept {
+        return sph;
+    }
+
+    [[nodiscard]] Stage ShaderStage() const noexcept {
+        return stage;
+    }
+
+protected:
+    ProgramHeader sph{};
+    Stage stage{};
 };
 
 } // 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 {
     return attribute >= Attribute::Generic0X && attribute <= Attribute::Generic31X;
 }
 
-int GenericAttributeIndex(Attribute attribute) {
+u32 GenericAttributeIndex(Attribute attribute) {
     if (!IsGeneric(attribute)) {
         throw InvalidArgument("Attribute is not generic {}", attribute);
     }
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 {
 
 [[nodiscard]] bool IsGeneric(Attribute attribute) noexcept;
 
-[[nodiscard]] int GenericAttributeIndex(Attribute attribute);
+[[nodiscard]] u32 GenericAttributeIndex(Attribute attribute);
 
 [[nodiscard]] std::string NameOf(Attribute attribute);
 
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() {
     Inst(Opcode::Return);
 }
 
+void IREmitter::DemoteToHelperInvocation(Block* continue_label) {
+    block->SetBranch(continue_label);
+    continue_label->AddImmediatePredecessor(block);
+    Inst(Opcode::DemoteToHelperInvocation, continue_label);
+}
+
 U32 IREmitter::GetReg(IR::Reg reg) {
     return Inst<U32>(Opcode::GetRegister, reg);
 }
@@ -248,6 +254,14 @@ void IREmitter::SetAttribute(IR::Attribute attribute, const F32& value) {
     Inst(Opcode::SetAttribute, attribute, value);
 }
 
+void IREmitter::SetFragColor(u32 index, u32 component, const F32& value) {
+    Inst(Opcode::SetFragColor, Imm32(index), Imm32(component), value);
+}
+
+void IREmitter::SetFragDepth(const F32& value) {
+    Inst(Opcode::SetFragDepth, value);
+}
+
 U32 IREmitter::WorkgroupIdX() {
     return U32{CompositeExtract(Inst(Opcode::WorkgroupId), 0)};
 }
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:
     void LoopMerge(Block* merge_block, Block* continue_target);
     void SelectionMerge(Block* merge_block);
     void Return();
+    void DemoteToHelperInvocation(Block* continue_label);
 
     [[nodiscard]] U32 GetReg(IR::Reg reg);
     void SetReg(IR::Reg reg, const U32& value);
@@ -67,6 +68,9 @@ public:
     [[nodiscard]] F32 GetAttribute(IR::Attribute attribute);
     void SetAttribute(IR::Attribute attribute, const F32& value);
 
+    void SetFragColor(u32 index, u32 component, const F32& value);
+    void SetFragDepth(const F32& value);
+
     [[nodiscard]] U32 WorkgroupIdX();
     [[nodiscard]] U32 WorkgroupIdY();
     [[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 {
     case Opcode::LoopMerge:
     case Opcode::SelectionMerge:
     case Opcode::Return:
+    case Opcode::DemoteToHelperInvocation:
     case Opcode::SetAttribute:
     case Opcode::SetAttributeIndexed:
+    case Opcode::SetFragColor:
+    case Opcode::SetFragDepth:
     case Opcode::WriteGlobalU8:
     case Opcode::WriteGlobalS8:
     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,
 OPCODE(LoopMerge,                                           Void,           Label,          Label,                                                          )
 OPCODE(SelectionMerge,                                      Void,           Label,                                                                          )
 OPCODE(Return,                                              Void,                                                                                           )
+OPCODE(DemoteToHelperInvocation,                            Void,           Label,                                                                          )
 
 // Context getters/setters
 OPCODE(GetRegister,                                         U32,            Reg,                                                                            )
@@ -28,10 +29,12 @@ OPCODE(GetCbufS16,                                          U32,            U32,
 OPCODE(GetCbufU32,                                          U32,            U32,            U32,                                                            )
 OPCODE(GetCbufF32,                                          F32,            U32,            U32,                                                            )
 OPCODE(GetCbufU64,                                          U64,            U32,            U32,                                                            )
-OPCODE(GetAttribute,                                        U32,            Attribute,                                                                      )
-OPCODE(SetAttribute,                                        Void,           Attribute,      U32,                                                            )
-OPCODE(GetAttributeIndexed,                                 U32,            U32,                                                                            )
-OPCODE(SetAttributeIndexed,                                 Void,           U32,            U32,                                                            )
+OPCODE(GetAttribute,                                        F32,            Attribute,                                                                      )
+OPCODE(SetAttribute,                                        Void,           Attribute,      F32,                                                            )
+OPCODE(GetAttributeIndexed,                                 F32,            U32,                                                                            )
+OPCODE(SetAttributeIndexed,                                 Void,           U32,            F32,                                                            )
+OPCODE(SetFragColor,                                        Void,           U32,            U32,            F32,                                            )
+OPCODE(SetFragDepth,                                        Void,           F32,                                                                            )
 OPCODE(GetZFlag,                                            U1,             Void,                                                                           )
 OPCODE(GetSFlag,                                            U1,             Void,                                                                           )
 OPCODE(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 @@
 
 #include "shader_recompiler/frontend/ir/basic_block.h"
 #include "shader_recompiler/shader_info.h"
+#include "shader_recompiler/stage.h"
 
 namespace Shader::IR {
 
@@ -17,6 +18,7 @@ struct Program {
     BlockList blocks;
     BlockList post_order_blocks;
     Info info;
+    Stage stage{};
 };
 
 [[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;
     return reg + (-num);
 }
 
-[[nodiscard]] constexpr Reg operator++(Reg& reg) {
+constexpr Reg operator++(Reg& reg) {
     reg = reg + 1;
     return reg;
 }
 
-[[nodiscard]] constexpr Reg operator++(Reg& reg, int) {
+constexpr Reg operator++(Reg& reg, int) {
     const Reg copy{reg};
     reg = reg + 1;
     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) {
     case Opcode::EXIT:
     case Opcode::JMP:
     case Opcode::JMX:
+    case Opcode::KIL:
     case Opcode::BRK:
     case Opcode::CONT:
     case Opcode::LONGJMP:
@@ -287,6 +288,13 @@ CFG::AnalysisState CFG::AnalyzeInst(Block* block, FunctionId function_id, Locati
         block->end = pc;
         return AnalysisState::Branch;
     }
+    case Opcode::KIL: {
+        const Predicate pred{inst.Pred()};
+        const auto ir_pred{static_cast<IR::Pred>(pred.index)};
+        const IR::Condition cond{inst.branch.flow_test, ir_pred, pred.negated};
+        AnalyzeCondInst(block, function_id, pc, EndClass::Kill, cond);
+        return AnalysisState::Branch;
+    }
     case Opcode::PBK:
     case Opcode::PCNT:
     case Opcode::PEXIT:
@@ -324,13 +332,12 @@ CFG::AnalysisState CFG::AnalyzeInst(Block* block, FunctionId function_id, Locati
         return AnalysisState::Continue;
     }
     const IR::Condition cond{static_cast<IR::Pred>(pred.index), pred.negated};
-    AnalyzeCondInst(block, function_id, pc, EndClass::Branch, cond, true);
+    AnalyzeCondInst(block, function_id, pc, EndClass::Branch, cond);
     return AnalysisState::Branch;
 }
 
 void CFG::AnalyzeCondInst(Block* block, FunctionId function_id, Location pc,
-                          EndClass insn_end_class, IR::Condition cond,
-                          bool visit_conditional_inst) {
+                          EndClass insn_end_class, IR::Condition cond) {
     if (block->begin != pc) {
         // If the block doesn't start in the conditional instruction
         // mark it as a label to visit it later
@@ -356,14 +363,16 @@ void CFG::AnalyzeCondInst(Block* block, FunctionId function_id, Location pc,
     // Impersonate the visited block with a virtual block
     *block = std::move(virtual_block);
     // Set the end properties of the conditional instruction
-    conditional_block->end = visit_conditional_inst ? (pc + 1) : pc;
+    conditional_block->end = pc + 1;
     conditional_block->end_class = insn_end_class;
     // Add a label to the instruction after the conditional instruction
     Block* const endif_block{AddLabel(conditional_block, block->stack, pc + 1, function_id)};
     // Branch to the next instruction from the virtual block
     block->branch_false = endif_block;
-    // And branch to it from the conditional instruction if it is a branch
-    if (insn_end_class == EndClass::Branch) {
+    // And branch to it from the conditional instruction if it is a branch or a kill instruction
+    // Kill instructions are considered a branch because they demote to a helper invocation and
+    // execution may continue.
+    if (insn_end_class == EndClass::Branch || insn_end_class == EndClass::Kill) {
         conditional_block->cond = IR::Condition{true};
         conditional_block->branch_true = endif_block;
         conditional_block->branch_false = nullptr;
@@ -415,7 +424,7 @@ CFG::AnalysisState CFG::AnalyzeEXIT(Block* block, FunctionId function_id, Locati
             throw NotImplementedException("Conditional EXIT with PEXIT token");
         }
         const IR::Condition cond{flow_test, static_cast<IR::Pred>(pred.index), pred.negated};
-        AnalyzeCondInst(block, function_id, pc, EndClass::Exit, cond, false);
+        AnalyzeCondInst(block, function_id, pc, EndClass::Exit, cond);
         return AnalysisState::Branch;
     }
     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
         block->branch_false = nullptr;
         return AnalysisState::Branch;
     }
-    block->end = pc;
+    block->end = pc + 1;
     block->end_class = EndClass::Exit;
     return AnalysisState::Branch;
 }
@@ -505,6 +514,12 @@ std::string CFG::Dot() const {
                                    node_uid);
                 ++node_uid;
                 break;
+            case EndClass::Kill:
+                dot += fmt::format("\t\t{}->N{};\n", name, node_uid);
+                dot += fmt::format("\t\tN{} [label=\"Kill\"][shape=square][style=stripped];\n",
+                                   node_uid);
+                ++node_uid;
+                break;
             }
         }
         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 {
     Call,
     Exit,
     Return,
+    Kill,
 };
 
 enum class Token {
@@ -130,7 +131,7 @@ private:
     AnalysisState AnalyzeInst(Block* block, FunctionId function_id, Location pc);
 
     void AnalyzeCondInst(Block* block, FunctionId function_id, Location pc, EndClass insn_end_class,
-                         IR::Condition cond, bool visit_conditional_inst);
+                         IR::Condition cond);
 
     /// Return true when the branch instruction is confirmed to be a branch
     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
     IR::Program program;
     program.blocks = VisitAST(inst_pool, block_pool, env, cfg);
     program.post_order_blocks = PostOrder(program.blocks);
+    program.stage = env.ShaderStage();
     RemoveUnreachableBlocks(program);
 
     // 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 {
     Loop,
     Break,
     Return,
+    Kill,
     Function,
     Identity,
     Not,
@@ -70,6 +71,7 @@ struct If {};
 struct Loop {};
 struct Break {};
 struct Return {};
+struct Kill {};
 struct FunctionTag {};
 struct Identity {};
 struct Not {};
@@ -93,6 +95,7 @@ struct Statement : ListBaseHook {
     Statement(Break, Statement* cond_, Statement* up_)
         : cond{cond_}, up{up_}, type{StatementType::Break} {}
     Statement(Return) : type{StatementType::Return} {}
+    Statement(Kill) : type{StatementType::Kill} {}
     Statement(FunctionTag) : children{}, type{StatementType::Function} {}
     Statement(Identity, IR::Condition cond_) : guest_cond{cond_}, type{StatementType::Identity} {}
     Statement(Not, Statement* op_) : op{op_}, type{StatementType::Not} {}
@@ -174,6 +177,9 @@ std::string DumpTree(const Tree& tree, u32 indentation = 0) {
         case StatementType::Return:
             ret += fmt::format("{}    return;\n", indent);
             break;
+        case StatementType::Kill:
+            ret += fmt::format("{}    kill;\n", indent);
+            break;
         case StatementType::SetVariable:
             ret += fmt::format("{}    goto_L{} = {};\n", indent, stmt->id, DumpExpr(stmt->op));
             break;
@@ -424,6 +430,9 @@ private:
                 gotos.push_back(root.insert(ip, *goto_stmt));
                 break;
             }
+            case Flow::EndClass::Kill:
+                root.insert(ip, *pool.Create(Kill{}));
+                break;
             }
         }
     }
@@ -729,6 +738,15 @@ private:
                 current_block = nullptr;
                 break;
             }
+            case StatementType::Kill: {
+                if (!current_block) {
+                    current_block = block_pool.Create(inst_pool);
+                    block_list.push_back(current_block);
+                }
+                IR::IREmitter{*current_block}.DemoteToHelperInvocation(continue_block);
+                current_block = nullptr;
+                break;
+            }
             default:
                 throw NotImplementedException("Statement type {}", stmt.type);
             }
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 @@
-// Copyright 2021 yuzu Emulator Project
-// Licensed under GPLv2 or any later version
-// Refer to the license.txt file included.
-
-#include "common/common_types.h"
-#include "shader_recompiler/exception.h"
-#include "shader_recompiler/frontend/maxwell/translate/impl/impl.h"
-
-namespace Shader::Maxwell {
-
-void TranslatorVisitor::EXIT(u64) {
-    ir.Exit();
-}
-
-} // 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 @@
+// Copyright 2021 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+#include "common/common_types.h"
+#include "shader_recompiler/exception.h"
+#include "shader_recompiler/frontend/maxwell/translate/impl/impl.h"
+
+namespace Shader::Maxwell {
+namespace {
+void ExitFragment(TranslatorVisitor& v) {
+    const ProgramHeader sph{v.env.SPH()};
+    IR::Reg src_reg{IR::Reg::R0};
+    for (u32 render_target = 0; render_target < 8; ++render_target) {
+        const std::array<bool, 4> mask{sph.ps.EnabledOutputComponents(render_target)};
+        for (u32 component = 0; component < 4; ++component) {
+            if (!mask[component]) {
+                continue;
+            }
+            v.ir.SetFragColor(render_target, component, v.F(src_reg));
+            ++src_reg;
+        }
+    }
+    if (sph.ps.omap.sample_mask != 0) {
+        throw NotImplementedException("Sample mask");
+    }
+    if (sph.ps.omap.depth != 0) {
+        throw NotImplementedException("Fragment depth");
+    }
+}
+} // Anonymous namespace
+
+void TranslatorVisitor::EXIT() {
+    switch (env.ShaderStage()) {
+    case Stage::Fragment:
+        ExitFragment(*this);
+        break;
+    default:
+        break;
+    }
+}
+
+} // 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:
     void DSETP_reg(u64 insn);
     void DSETP_cbuf(u64 insn);
     void DSETP_imm(u64 insn);
-    void EXIT(u64 insn);
+    void EXIT();
     void F2F_reg(u64 insn);
     void F2F_cbuf(u64 insn);
     void F2F_imm(u64 insn);
@@ -220,7 +220,7 @@ public:
     void JCAL(u64 insn);
     void JMP(u64 insn);
     void JMX(u64 insn);
-    void KIL(u64 insn);
+    void KIL();
     void LD(u64 insn);
     void LDC(u64 insn);
     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 @@
 
 namespace Shader::Maxwell {
 namespace {
+enum class Size : u64 {
+    B32,
+    B64,
+    B96,
+    B128,
+};
+
 enum class InterpolationMode : u64 {
     Pass,
     Multiply,
@@ -23,8 +30,85 @@ enum class SampleMode : u64 {
     Centroid,
     Offset,
 };
+
+int NumElements(Size size) {
+    switch (size) {
+    case Size::B32:
+        return 1;
+    case Size::B64:
+        return 2;
+    case Size::B96:
+        return 3;
+    case Size::B128:
+        return 4;
+    }
+    throw InvalidArgument("Invalid size {}", size);
+}
 } // Anonymous namespace
 
+void TranslatorVisitor::ALD(u64 insn) {
+    union {
+        u64 raw;
+        BitField<0, 8, IR::Reg> dest_reg;
+        BitField<8, 8, IR::Reg> index_reg;
+        BitField<20, 10, u64> absolute_offset;
+        BitField<20, 11, s64> relative_offset;
+        BitField<39, 8, IR::Reg> stream_reg;
+        BitField<32, 1, u64> o;
+        BitField<31, 1, u64> patch;
+        BitField<47, 2, Size> size;
+    } const ald{insn};
+
+    if (ald.o != 0) {
+        throw NotImplementedException("O");
+    }
+    if (ald.patch != 0) {
+        throw NotImplementedException("P");
+    }
+    if (ald.index_reg != IR::Reg::RZ) {
+        throw NotImplementedException("Indexed");
+    }
+    const u64 offset{ald.absolute_offset.Value()};
+    if (offset % 4 != 0) {
+        throw NotImplementedException("Unaligned absolute offset {}", offset);
+    }
+    const int num_elements{NumElements(ald.size)};
+    for (int element = 0; element < num_elements; ++element) {
+        F(ald.dest_reg + element, ir.GetAttribute(IR::Attribute{offset / 4 + element}));
+    }
+}
+
+void TranslatorVisitor::AST(u64 insn) {
+    union {
+        u64 raw;
+        BitField<0, 8, IR::Reg> src_reg;
+        BitField<8, 8, IR::Reg> index_reg;
+        BitField<20, 10, u64> absolute_offset;
+        BitField<20, 11, s64> relative_offset;
+        BitField<31, 1, u64> patch;
+        BitField<39, 8, IR::Reg> stream_reg;
+        BitField<47, 2, Size> size;
+    } const ast{insn};
+
+    if (ast.patch != 0) {
+        throw NotImplementedException("P");
+    }
+    if (ast.stream_reg != IR::Reg::RZ) {
+        throw NotImplementedException("Stream store");
+    }
+    if (ast.index_reg != IR::Reg::RZ) {
+        throw NotImplementedException("Indexed store");
+    }
+    const u64 offset{ast.absolute_offset.Value()};
+    if (offset % 4 != 0) {
+        throw NotImplementedException("Unaligned absolute offset {}", offset);
+    }
+    const int num_elements{NumElements(ast.size)};
+    for (int element = 0; element < num_elements; ++element) {
+        ir.SetAttribute(IR::Attribute{offset / 4 + element}, F(ast.src_reg + element));
+    }
+}
+
 void TranslatorVisitor::IPA(u64 insn) {
     // IPA is the instruction used to read varyings from a fragment shader.
     // gl_FragCoord is mapped to the gl_Position attribute.
@@ -51,7 +135,7 @@ void TranslatorVisitor::IPA(u64 insn) {
     // }
     const bool is_indexed{ipa.idx != 0 && ipa.index_reg != IR::Reg::RZ};
     if (is_indexed) {
-        throw NotImplementedException("IPA.IDX");
+        throw NotImplementedException("IDX");
     }
 
     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) {
     ThrowNotImplemented(Opcode::AL2P);
 }
 
-void TranslatorVisitor::ALD(u64) {
-    ThrowNotImplemented(Opcode::ALD);
-}
-
-void TranslatorVisitor::AST(u64) {
-    ThrowNotImplemented(Opcode::AST);
-}
-
 void TranslatorVisitor::ATOM_cas(u64) {
     ThrowNotImplemented(Opcode::ATOM_cas);
 }
@@ -153,10 +145,6 @@ void TranslatorVisitor::DSETP_imm(u64) {
     ThrowNotImplemented(Opcode::DSETP_imm);
 }
 
-void TranslatorVisitor::EXIT(u64) {
-    throw LogicError("Visting EXIT instruction");
-}
-
 void TranslatorVisitor::F2F_reg(u64) {
     ThrowNotImplemented(Opcode::F2F_reg);
 }
@@ -345,8 +333,8 @@ void TranslatorVisitor::JMX(u64) {
     ThrowNotImplemented(Opcode::JMX);
 }
 
-void TranslatorVisitor::KIL(u64) {
-    ThrowNotImplemented(Opcode::KIL);
+void TranslatorVisitor::KIL() {
+    // KIL is a no-op
 }
 
 void 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) {
         BitField<36, 13, u64> cbuf_offset;
     } const tex{insn};
 
-    Impl(*this, insn, tex.aoffi != 0, tex.blod, tex.lc != 0, static_cast<u32>(tex.cbuf_offset));
+    Impl(*this, insn, tex.aoffi != 0, tex.blod, tex.lc != 0, static_cast<u32>(tex.cbuf_offset * 4));
 }
 
 void 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) {
 
 IR::Value Sample(TranslatorVisitor& v, u64 insn) {
     const Encoding texs{insn};
-    const IR::U32 handle{v.ir.Imm32(static_cast<u32>(texs.cbuf_offset))};
+    const IR::U32 handle{v.ir.Imm32(static_cast<u32>(texs.cbuf_offset * 4))};
     const IR::F32 zero{v.ir.Imm32(0.0f)};
     const IR::Reg reg_a{texs.src_reg_a};
     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) {
         return;
     }
     info.constant_buffer_mask |= 1U << index;
-    info.constant_buffer_descriptors.push_back({
-        .index{index},
-        .count{1},
-    });
+
+    auto& cbufs{info.constant_buffer_descriptors};
+    cbufs.insert(std::ranges::lower_bound(cbufs, index, {}, &ConstantBufferDescriptor::index),
+                 ConstantBufferDescriptor{
+                     .index{index},
+                     .count{1},
+                 });
+}
+
+void GetAttribute(Info& info, IR::Attribute attribute) {
+    if (IR::IsGeneric(attribute)) {
+        info.loads_generics.at(IR::GenericAttributeIndex(attribute)) = true;
+        return;
+    }
+    switch (attribute) {
+    case IR::Attribute::PositionX:
+    case IR::Attribute::PositionY:
+    case IR::Attribute::PositionZ:
+    case IR::Attribute::PositionW:
+        info.loads_position = true;
+        break;
+    default:
+        throw NotImplementedException("Get attribute {}", attribute);
+    }
+}
+
+void SetAttribute(Info& info, IR::Attribute attribute) {
+    if (IR::IsGeneric(attribute)) {
+        info.stores_generics.at(IR::GenericAttributeIndex(attribute)) = true;
+        return;
+    }
+    switch (attribute) {
+    case IR::Attribute::PositionX:
+    case IR::Attribute::PositionY:
+    case IR::Attribute::PositionZ:
+    case IR::Attribute::PositionW:
+        info.stores_position = true;
+        break;
+    default:
+        throw NotImplementedException("Set attribute {}", attribute);
+    }
 }
 
 void VisitUsages(Info& info, IR::Inst& inst) {
@@ -162,6 +199,21 @@ void VisitUsages(Info& info, IR::Inst& inst) {
         break;
     }
     switch (inst.Opcode()) {
+    case IR::Opcode::DemoteToHelperInvocation:
+        info.uses_demote_to_helper_invocation = true;
+        break;
+    case IR::Opcode::GetAttribute:
+        GetAttribute(info, inst.Arg(0).Attribute());
+        break;
+    case IR::Opcode::SetAttribute:
+        SetAttribute(info, inst.Arg(0).Attribute());
+        break;
+    case IR::Opcode::SetFragColor:
+        info.stores_frag_color[inst.Arg(0).U32()] = true;
+        break;
+    case IR::Opcode::SetFragDepth:
+        info.stores_frag_depth = true;
+        break;
     case IR::Opcode::WorkgroupId:
         info.uses_workgroup_id = true;
         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:
         const size_t num_args{phi.NumArgs()};
         for (size_t arg_index = 0; arg_index < num_args; ++arg_index) {
             const IR::Value& op{phi.Arg(arg_index)};
-            if (op == same || op == IR::Value{&phi}) {
+            if (op.Resolve() == same.Resolve() || op == IR::Value{&phi}) {
                 // Unique value or self-reference
                 continue;
             }
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 @@
+// Copyright 2018 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+#pragma once
+
+#include <array>
+#include <optional>
+
+#include "common/bit_field.h"
+#include "common/common_funcs.h"
+#include "common/common_types.h"
+
+namespace Shader {
+
+enum class OutputTopology : u32 {
+    PointList = 1,
+    LineStrip = 6,
+    TriangleStrip = 7,
+};
+
+enum class PixelImap : u8 {
+    Unused = 0,
+    Constant = 1,
+    Perspective = 2,
+    ScreenLinear = 3,
+};
+
+// Documentation in:
+// http://download.nvidia.com/open-gpu-doc/Shader-Program-Header/1/Shader-Program-Header.html
+struct ProgramHeader {
+    union {
+        BitField<0, 5, u32> sph_type;
+        BitField<5, 5, u32> version;
+        BitField<10, 4, u32> shader_type;
+        BitField<14, 1, u32> mrt_enable;
+        BitField<15, 1, u32> kills_pixels;
+        BitField<16, 1, u32> does_global_store;
+        BitField<17, 4, u32> sass_version;
+        BitField<21, 5, u32> reserved;
+        BitField<26, 1, u32> does_load_or_store;
+        BitField<27, 1, u32> does_fp64;
+        BitField<28, 4, u32> stream_out_mask;
+    } common0;
+
+    union {
+        BitField<0, 24, u32> shader_local_memory_low_size;
+        BitField<24, 8, u32> per_patch_attribute_count;
+    } common1;
+
+    union {
+        BitField<0, 24, u32> shader_local_memory_high_size;
+        BitField<24, 8, u32> threads_per_input_primitive;
+    } common2;
+
+    union {
+        BitField<0, 24, u32> shader_local_memory_crs_size;
+        BitField<24, 4, OutputTopology> output_topology;
+        BitField<28, 4, u32> reserved;
+    } common3;
+
+    union {
+        BitField<0, 12, u32> max_output_vertices;
+        BitField<12, 8, u32> store_req_start; // NOTE: not used by geometry shaders.
+        BitField<20, 4, u32> reserved;
+        BitField<24, 8, u32> store_req_end; // NOTE: not used by geometry shaders.
+    } common4;
+
+    union {
+        struct {
+            INSERT_PADDING_BYTES_NOINIT(3);  // ImapSystemValuesA
+            INSERT_PADDING_BYTES_NOINIT(1);  // ImapSystemValuesB
+            INSERT_PADDING_BYTES_NOINIT(16); // ImapGenericVector[32]
+            INSERT_PADDING_BYTES_NOINIT(2);  // ImapColor
+            union {
+                BitField<0, 8, u16> clip_distances;
+                BitField<8, 1, u16> point_sprite_s;
+                BitField<9, 1, u16> point_sprite_t;
+                BitField<10, 1, u16> fog_coordinate;
+                BitField<12, 1, u16> tessellation_eval_point_u;
+                BitField<13, 1, u16> tessellation_eval_point_v;
+                BitField<14, 1, u16> instance_id;
+                BitField<15, 1, u16> vertex_id;
+            };
+            INSERT_PADDING_BYTES_NOINIT(5);  // ImapFixedFncTexture[10]
+            INSERT_PADDING_BYTES_NOINIT(1);  // ImapReserved
+            INSERT_PADDING_BYTES_NOINIT(3);  // OmapSystemValuesA
+            INSERT_PADDING_BYTES_NOINIT(1);  // OmapSystemValuesB
+            INSERT_PADDING_BYTES_NOINIT(16); // OmapGenericVector[32]
+            INSERT_PADDING_BYTES_NOINIT(2);  // OmapColor
+            INSERT_PADDING_BYTES_NOINIT(2);  // OmapSystemValuesC
+            INSERT_PADDING_BYTES_NOINIT(5);  // OmapFixedFncTexture[10]
+            INSERT_PADDING_BYTES_NOINIT(1);  // OmapReserved
+        } vtg;
+
+        struct {
+            INSERT_PADDING_BYTES_NOINIT(3); // ImapSystemValuesA
+            INSERT_PADDING_BYTES_NOINIT(1); // ImapSystemValuesB
+
+            union {
+                BitField<0, 2, PixelImap> x;
+                BitField<2, 2, PixelImap> y;
+                BitField<4, 2, PixelImap> z;
+                BitField<6, 2, PixelImap> w;
+                u8 raw;
+            } imap_generic_vector[32];
+
+            INSERT_PADDING_BYTES_NOINIT(2);  // ImapColor
+            INSERT_PADDING_BYTES_NOINIT(2);  // ImapSystemValuesC
+            INSERT_PADDING_BYTES_NOINIT(10); // ImapFixedFncTexture[10]
+            INSERT_PADDING_BYTES_NOINIT(2);  // ImapReserved
+
+            struct {
+                u32 target;
+                union {
+                    BitField<0, 1, u32> sample_mask;
+                    BitField<1, 1, u32> depth;
+                    BitField<2, 30, u32> reserved;
+                };
+            } omap;
+
+            [[nodiscard]] std::array<bool, 4> EnabledOutputComponents(u32 rt) const noexcept {
+                const u32 bits{omap.target >> (rt * 4)};
+                return {(bits & 1) != 0, (bits & 2) != 0, (bits & 4) != 0, (bits & 8) != 0};
+            }
+
+            [[nodiscard]] std::array<PixelImap, 4> GenericInputMap(u32 attribute) const {
+                const auto& vector{imap_generic_vector[attribute]};
+                return {vector.x, vector.y, vector.z, vector.w};
+            }
+        } ps;
+
+        std::array<u32, 0xf> raw;
+    };
+
+    [[nodiscard]] u64 LocalMemorySize() const noexcept {
+        return (common1.shader_local_memory_low_size |
+                (common2.shader_local_memory_high_size << 24));
+    }
+};
+static_assert(sizeof(ProgramHeader) == 0x50, "Incorrect structure size");
+
+} // 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 @@
-// Copyright 2021 yuzu Emulator Project
-// Licensed under GPLv2 or any later version
-// Refer to the license.txt file included.
-
-#include <vector>
-
-#include "common/common_types.h"
-#include "shader_recompiler/backend/spirv/emit_spirv.h"
-#include "shader_recompiler/environment.h"
-#include "shader_recompiler/frontend/maxwell/control_flow.h"
-#include "shader_recompiler/frontend/maxwell/program.h"
-#include "shader_recompiler/object_pool.h"
-#include "shader_recompiler/recompiler.h"
-
-namespace Shader {
-
-std::pair<Info, std::vector<u32>> RecompileSPIRV(const Profile& profile, Environment& env,
-                                                 u32 start_address) {
-    ObjectPool<Maxwell::Flow::Block> flow_block_pool;
-    ObjectPool<IR::Inst> inst_pool;
-    ObjectPool<IR::Block> block_pool;
-
-    Maxwell::Flow::CFG cfg{env, flow_block_pool, start_address};
-    IR::Program program{Maxwell::TranslateProgram(inst_pool, block_pool, env, cfg)};
-    return {std::move(program.info), Backend::SPIRV::EmitSPIRV(profile, env, program)};
-}
-
-} // 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 @@
-// Copyright 2021 yuzu Emulator Project
-// Licensed under GPLv2 or any later version
-// Refer to the license.txt file included.
-
-#pragma once
-
-#include <utility>
-#include <vector>
-
-#include "common/common_types.h"
-#include "shader_recompiler/environment.h"
-#include "shader_recompiler/profile.h"
-#include "shader_recompiler/shader_info.h"
-
-namespace Shader {
-
-[[nodiscard]] std::pair<Info, std::vector<u32>> RecompileSPIRV(const Profile& profile,
-                                                               Environment& env, u32 start_address);
-
-} // 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 {
 
     bool uses_workgroup_id{};
     bool uses_local_invocation_id{};
+
+    std::array<bool, 32> loads_generics{};
+    bool loads_position{};
+
+    std::array<bool, 8> stores_frag_color{};
+    bool stores_frag_depth{};
+    std::array<bool, 32> stores_generics{};
+    bool stores_position{};
+
     bool uses_fp16{};
     bool uses_fp64{};
     bool uses_fp16_denorms_flush{};
@@ -68,6 +77,7 @@ struct Info {
     bool uses_image_1d{};
     bool uses_sampled_1d{};
     bool uses_sparse_residency{};
+    bool uses_demote_to_helper_invocation{};
 
     IR::Type used_constant_buffer_types{};
 
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 @@
+// Copyright 2021 yuzu Emulator Project
+// Licensed under GPLv2 or any later version
+// Refer to the license.txt file included.
+
+#pragma once
+
+namespace Shader {
+
+enum class Stage {
+    Compute,
+    VertexA,
+    VertexB,
+    TessellationControl,
+    TessellationEval,
+    Geometry,
+    Fragment,
+};
+
+} // namespace Shader