Merge pull request #9501 from FernandoS27/yfc-rel-2

Yuzu Fried Chicken Part 1.5: MacroHLE Rework and Dynamic State

79 files changed, 3008 insertions, 573 deletions

diff --git a/src/common/CMakeLists.txt b/src/common/CMakeLists.txt
index eb05e46a8..45332cf95 100644
--- a/src/common/CMakeLists.txt
+++ b/src/common/CMakeLists.txt
@@ -97,6 +97,7 @@ add_library(common STATIC
     point.h
     precompiled_headers.h
     quaternion.h
+    range_map.h
     reader_writer_queue.h
     ring_buffer.h
     ${CMAKE_CURRENT_BINARY_DIR}/scm_rev.cpp
diff --git a/src/common/range_map.h b/src/common/range_map.h
new file mode 100644
index 000000000..79c7ef547
--- /dev/null
+++ b/src/common/range_map.h
@@ -0,0 +1,139 @@
+// SPDX-FileCopyrightText: Copyright 2022 yuzu Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
+#pragma once
+
+#include <map>
+#include <type_traits>
+
+#include "common/common_types.h"
+
+namespace Common {
+
+template <typename KeyTBase, typename ValueT>
+class RangeMap {
+private:
+    using KeyT =
+        std::conditional_t<std::is_signed_v<KeyTBase>, KeyTBase, std::make_signed_t<KeyTBase>>;
+
+public:
+    explicit RangeMap(ValueT null_value_) : null_value{null_value_} {
+        container.emplace(std::numeric_limits<KeyT>::min(), null_value);
+    };
+    ~RangeMap() = default;
+
+    void Map(KeyTBase address, KeyTBase address_end, ValueT value) {
+        KeyT new_address = static_cast<KeyT>(address);
+        KeyT new_address_end = static_cast<KeyT>(address_end);
+        if (new_address < 0) {
+            new_address = 0;
+        }
+        if (new_address_end < 0) {
+            new_address_end = 0;
+        }
+        InternalMap(new_address, new_address_end, value);
+    }
+
+    void Unmap(KeyTBase address, KeyTBase address_end) {
+        Map(address, address_end, null_value);
+    }
+
+    [[nodiscard]] size_t GetContinousSizeFrom(KeyTBase address) const {
+        const KeyT new_address = static_cast<KeyT>(address);
+        if (new_address < 0) {
+            return 0;
+        }
+        return ContinousSizeInternal(new_address);
+    }
+
+    [[nodiscard]] ValueT GetValueAt(KeyT address) const {
+        const KeyT new_address = static_cast<KeyT>(address);
+        if (new_address < 0) {
+            return null_value;
+        }
+        return GetValueInternal(new_address);
+    }
+
+private:
+    using MapType = std::map<KeyT, ValueT>;
+    using IteratorType = typename MapType::iterator;
+    using ConstIteratorType = typename MapType::const_iterator;
+
+    size_t ContinousSizeInternal(KeyT address) const {
+        const auto it = GetFirstElementBeforeOrOn(address);
+        if (it == container.end() || it->second == null_value) {
+            return 0;
+        }
+        const auto it_end = std::next(it);
+        if (it_end == container.end()) {
+            return std::numeric_limits<KeyT>::max() - address;
+        }
+        return it_end->first - address;
+    }
+
+    ValueT GetValueInternal(KeyT address) const {
+        const auto it = GetFirstElementBeforeOrOn(address);
+        if (it == container.end()) {
+            return null_value;
+        }
+        return it->second;
+    }
+
+    ConstIteratorType GetFirstElementBeforeOrOn(KeyT address) const {
+        auto it = container.lower_bound(address);
+        if (it == container.begin()) {
+            return it;
+        }
+        if (it != container.end() && (it->first == address)) {
+            return it;
+        }
+        --it;
+        return it;
+    }
+
+    ValueT GetFirstValueWithin(KeyT address) {
+        auto it = container.lower_bound(address);
+        if (it == container.begin()) {
+            return it->second;
+        }
+        if (it == container.end()) [[unlikely]] { // this would be a bug
+            return null_value;
+        }
+        --it;
+        return it->second;
+    }
+
+    ValueT GetLastValueWithin(KeyT address) {
+        auto it = container.upper_bound(address);
+        if (it == container.end()) {
+            return null_value;
+        }
+        if (it == container.begin()) [[unlikely]] { // this would be a bug
+            return it->second;
+        }
+        --it;
+        return it->second;
+    }
+
+    void InternalMap(KeyT address, KeyT address_end, ValueT value) {
+        const bool must_add_start = GetFirstValueWithin(address) != value;
+        const ValueT last_value = GetLastValueWithin(address_end);
+        const bool must_add_end = last_value != value;
+        auto it = container.lower_bound(address);
+        const auto it_end = container.upper_bound(address_end);
+        while (it != it_end) {
+            it = container.erase(it);
+        }
+        if (must_add_start) {
+            container.emplace(address, value);
+        }
+        if (must_add_end) {
+            container.emplace(address_end, last_value);
+        }
+    }
+
+    ValueT null_value;
+    MapType container;
+};
+
+} // namespace Common
diff --git a/src/common/settings.h b/src/common/settings.h
index 6b199af93..5017951c5 100644
--- a/src/common/settings.h
+++ b/src/common/settings.h
@@ -531,6 +531,7 @@ struct Values {
     Setting<bool> reporting_services{false, "reporting_services"};
     Setting<bool> quest_flag{false, "quest_flag"};
     Setting<bool> disable_macro_jit{false, "disable_macro_jit"};
+    Setting<bool> disable_macro_hle{false, "disable_macro_hle"};
     Setting<bool> extended_logging{false, "extended_logging"};
     Setting<bool> use_debug_asserts{false, "use_debug_asserts"};
     Setting<bool> use_auto_stub{false, "use_auto_stub"};
diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_context_get_set.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_context_get_set.cpp
index f0bd84ab2..c7d7d5fef 100644
--- a/src/shader_recompiler/backend/glasm/emit_glasm_context_get_set.cpp
+++ b/src/shader_recompiler/backend/glasm/emit_glasm_context_get_set.cpp
@@ -137,6 +137,15 @@ void EmitGetAttribute(EmitContext& ctx, IR::Inst& inst, IR::Attribute attr, Scal
     case IR::Attribute::VertexId:
         ctx.Add("MOV.F {}.x,{}.id;", inst, ctx.attrib_name);
         break;
+    case IR::Attribute::BaseInstance:
+        ctx.Add("MOV.F {}.x,{}.baseInstance;", inst, ctx.attrib_name);
+        break;
+    case IR::Attribute::BaseVertex:
+        ctx.Add("MOV.F {}.x,{}.baseVertex;", inst, ctx.attrib_name);
+        break;
+    case IR::Attribute::DrawID:
+        ctx.Add("MOV.F {}.x,{}.draw.id;", inst, ctx.attrib_name);
+        break;
     case IR::Attribute::FrontFace:
         ctx.Add("CMP.F {}.x,{}.facing.x,0,-1;", inst, ctx.attrib_name);
         break;
@@ -156,6 +165,15 @@ void EmitGetAttributeU32(EmitContext& ctx, IR::Inst& inst, IR::Attribute attr, S
     case IR::Attribute::VertexId:
         ctx.Add("MOV.S {}.x,{}.id;", inst, ctx.attrib_name);
         break;
+    case IR::Attribute::BaseInstance:
+        ctx.Add("MOV.S {}.x,{}.baseInstance;", inst, ctx.attrib_name);
+        break;
+    case IR::Attribute::BaseVertex:
+        ctx.Add("MOV.S {}.x,{}.baseVertex;", inst, ctx.attrib_name);
+        break;
+    case IR::Attribute::DrawID:
+        ctx.Add("MOV.S {}.x,{}.draw.id;", inst, ctx.attrib_name);
+        break;
     default:
         throw NotImplementedException("Get U32 attribute {}", attr);
     }
diff --git a/src/shader_recompiler/backend/glsl/emit_glsl.cpp b/src/shader_recompiler/backend/glsl/emit_glsl.cpp
index e8a4390f6..d91e04446 100644
--- a/src/shader_recompiler/backend/glsl/emit_glsl.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl.cpp
@@ -219,7 +219,7 @@ std::string EmitGLSL(const Profile& profile, const RuntimeInfo& runtime_info, IR
     EmitContext ctx{program, bindings, profile, runtime_info};
     Precolor(program);
     EmitCode(ctx, program);
-    const std::string version{fmt::format("#version 450{}\n", GlslVersionSpecifier(ctx))};
+    const std::string version{fmt::format("#version 460{}\n", GlslVersionSpecifier(ctx))};
     ctx.header.insert(0, version);
     if (program.shared_memory_size > 0) {
         const auto requested_size{program.shared_memory_size};
diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_context_get_set.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_context_get_set.cpp
index 39579cf5d..2e369ed72 100644
--- a/src/shader_recompiler/backend/glsl/emit_glsl_context_get_set.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_glsl_context_get_set.cpp
@@ -234,6 +234,15 @@ void EmitGetAttribute(EmitContext& ctx, IR::Inst& inst, IR::Attribute attr,
     case IR::Attribute::FrontFace:
         ctx.AddF32("{}=itof(gl_FrontFacing?-1:0);", inst);
         break;
+    case IR::Attribute::BaseInstance:
+        ctx.AddF32("{}=itof(gl_BaseInstance);", inst);
+        break;
+    case IR::Attribute::BaseVertex:
+        ctx.AddF32("{}=itof(gl_BaseVertex);", inst);
+        break;
+    case IR::Attribute::DrawID:
+        ctx.AddF32("{}=itof(gl_DrawID);", inst);
+        break;
     default:
         throw NotImplementedException("Get attribute {}", attr);
     }
@@ -250,6 +259,15 @@ void EmitGetAttributeU32(EmitContext& ctx, IR::Inst& inst, IR::Attribute attr, s
     case IR::Attribute::VertexId:
         ctx.AddU32("{}=uint(gl_VertexID);", inst);
         break;
+    case IR::Attribute::BaseInstance:
+        ctx.AddU32("{}=uint(gl_BaseInstance);", inst);
+        break;
+    case IR::Attribute::BaseVertex:
+        ctx.AddU32("{}=uint(gl_BaseVertex);", inst);
+        break;
+    case IR::Attribute::DrawID:
+        ctx.AddU32("{}=uint(gl_DrawID);", inst);
+        break;
     default:
         throw NotImplementedException("Get U32 attribute {}", attr);
     }
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 73b67f0af..db9c94ce8 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
@@ -339,6 +339,12 @@ Id EmitGetAttribute(EmitContext& ctx, IR::Attribute attr, Id vertex) {
             const Id base{ctx.OpLoad(ctx.U32[1], ctx.base_vertex)};
             return ctx.OpBitcast(ctx.F32[1], ctx.OpISub(ctx.U32[1], index, base));
         }
+    case IR::Attribute::BaseInstance:
+        return ctx.OpBitcast(ctx.F32[1], ctx.OpLoad(ctx.U32[1], ctx.base_instance));
+    case IR::Attribute::BaseVertex:
+        return ctx.OpBitcast(ctx.F32[1], ctx.OpLoad(ctx.U32[1], ctx.base_vertex));
+    case IR::Attribute::DrawID:
+        return ctx.OpBitcast(ctx.F32[1], ctx.OpLoad(ctx.U32[1], ctx.draw_index));
     case IR::Attribute::FrontFace:
         return ctx.OpSelect(ctx.F32[1], ctx.OpLoad(ctx.U1, ctx.front_face),
                             ctx.OpBitcast(ctx.F32[1], ctx.Const(std::numeric_limits<u32>::max())),
@@ -380,6 +386,12 @@ Id EmitGetAttributeU32(EmitContext& ctx, IR::Attribute attr, Id) {
             const Id base{ctx.OpLoad(ctx.U32[1], ctx.base_vertex)};
             return ctx.OpISub(ctx.U32[1], index, base);
         }
+    case IR::Attribute::BaseInstance:
+        return ctx.OpLoad(ctx.U32[1], ctx.base_instance);
+    case IR::Attribute::BaseVertex:
+        return ctx.OpLoad(ctx.U32[1], ctx.base_vertex);
+    case IR::Attribute::DrawID:
+        return ctx.OpLoad(ctx.U32[1], ctx.draw_index);
     default:
         throw NotImplementedException("Read U32 attribute {}", attr);
     }
diff --git a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp
index 41dc6d031..ecb2db494 100644
--- a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp
+++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp
@@ -1379,18 +1379,31 @@ void EmitContext::DefineInputs(const IR::Program& program) {
     if (loads[IR::Attribute::InstanceId]) {
         if (profile.support_vertex_instance_id) {
             instance_id = DefineInput(*this, U32[1], true, spv::BuiltIn::InstanceId);
+            if (loads[IR::Attribute::BaseInstance]) {
+                base_instance = DefineInput(*this, U32[1], true, spv::BuiltIn::BaseVertex);
+            }
         } else {
             instance_index = DefineInput(*this, U32[1], true, spv::BuiltIn::InstanceIndex);
             base_instance = DefineInput(*this, U32[1], true, spv::BuiltIn::BaseInstance);
         }
+    } else if (loads[IR::Attribute::BaseInstance]) {
+        base_instance = DefineInput(*this, U32[1], true, spv::BuiltIn::BaseInstance);
     }
     if (loads[IR::Attribute::VertexId]) {
         if (profile.support_vertex_instance_id) {
             vertex_id = DefineInput(*this, U32[1], true, spv::BuiltIn::VertexId);
+            if (loads[IR::Attribute::BaseVertex]) {
+                base_vertex = DefineInput(*this, U32[1], true, spv::BuiltIn::BaseVertex);
+            }
         } else {
             vertex_index = DefineInput(*this, U32[1], true, spv::BuiltIn::VertexIndex);
             base_vertex = DefineInput(*this, U32[1], true, spv::BuiltIn::BaseVertex);
         }
+    } else if (loads[IR::Attribute::BaseVertex]) {
+        base_vertex = DefineInput(*this, U32[1], true, spv::BuiltIn::BaseVertex);
+    }
+    if (loads[IR::Attribute::DrawID]) {
+        draw_index = DefineInput(*this, U32[1], true, spv::BuiltIn::DrawIndex);
     }
     if (loads[IR::Attribute::FrontFace]) {
         front_face = DefineInput(*this, U1, true, spv::BuiltIn::FrontFacing);
diff --git a/src/shader_recompiler/backend/spirv/spirv_emit_context.h b/src/shader_recompiler/backend/spirv/spirv_emit_context.h
index dde45b4bc..4414a5169 100644
--- a/src/shader_recompiler/backend/spirv/spirv_emit_context.h
+++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.h
@@ -218,6 +218,7 @@ public:
     Id base_instance{};
     Id vertex_id{};
     Id vertex_index{};
+    Id draw_index{};
     Id base_vertex{};
     Id front_face{};
     Id point_coord{};
diff --git a/src/shader_recompiler/environment.h b/src/shader_recompiler/environment.h
index 402f2664f..26e8307c1 100644
--- a/src/shader_recompiler/environment.h
+++ b/src/shader_recompiler/environment.h
@@ -34,6 +34,11 @@ public:
 
     [[nodiscard]] virtual std::array<u32, 3> WorkgroupSize() const = 0;
 
+    [[nodiscard]] virtual bool HasHLEMacroState() const = 0;
+
+    [[nodiscard]] virtual std::optional<ReplaceConstant> GetReplaceConstBuffer(u32 bank,
+                                                                               u32 offset) = 0;
+
     virtual void Dump(u64 hash) = 0;
 
     [[nodiscard]] const ProgramHeader& SPH() const noexcept {
@@ -52,11 +57,16 @@ public:
         return start_address;
     }
 
+    [[nodiscard]] bool IsPropietaryDriver() const noexcept {
+        return is_propietary_driver;
+    }
+
 protected:
     ProgramHeader sph{};
     std::array<u32, 8> gp_passthrough_mask{};
     Stage stage{};
     u32 start_address{};
+    bool is_propietary_driver{};
 };
 
 } // namespace Shader
diff --git a/src/shader_recompiler/frontend/ir/attribute.cpp b/src/shader_recompiler/frontend/ir/attribute.cpp
index 7d3d882e4..1bf9db935 100644
--- a/src/shader_recompiler/frontend/ir/attribute.cpp
+++ b/src/shader_recompiler/frontend/ir/attribute.cpp
@@ -446,6 +446,12 @@ std::string NameOf(Attribute attribute) {
         return "ViewportMask";
     case Attribute::FrontFace:
         return "FrontFace";
+    case Attribute::BaseInstance:
+        return "BaseInstance";
+    case Attribute::BaseVertex:
+        return "BaseVertex";
+    case Attribute::DrawID:
+        return "DrawID";
     }
     return fmt::format("<reserved attribute {}>", static_cast<int>(attribute));
 }
diff --git a/src/shader_recompiler/frontend/ir/attribute.h b/src/shader_recompiler/frontend/ir/attribute.h
index 6ee3947b1..5f039b6f6 100644
--- a/src/shader_recompiler/frontend/ir/attribute.h
+++ b/src/shader_recompiler/frontend/ir/attribute.h
@@ -219,6 +219,11 @@ enum class Attribute : u64 {
     FixedFncTexture9Q = 231,
     ViewportMask = 232,
     FrontFace = 255,
+
+    // Implementation attributes
+    BaseInstance = 256,
+    BaseVertex = 257,
+    DrawID = 258,
 };
 
 constexpr size_t NUM_GENERICS = 32;
diff --git a/src/shader_recompiler/frontend/ir/ir_emitter.cpp b/src/shader_recompiler/frontend/ir/ir_emitter.cpp
index 0cdac0eff..eb2e49a68 100644
--- a/src/shader_recompiler/frontend/ir/ir_emitter.cpp
+++ b/src/shader_recompiler/frontend/ir/ir_emitter.cpp
@@ -294,6 +294,14 @@ F32 IREmitter::GetAttribute(IR::Attribute attribute, const U32& vertex) {
     return Inst<F32>(Opcode::GetAttribute, attribute, vertex);
 }
 
+U32 IREmitter::GetAttributeU32(IR::Attribute attribute) {
+    return GetAttributeU32(attribute, Imm32(0));
+}
+
+U32 IREmitter::GetAttributeU32(IR::Attribute attribute, const U32& vertex) {
+    return Inst<U32>(Opcode::GetAttributeU32, attribute, vertex);
+}
+
 void IREmitter::SetAttribute(IR::Attribute attribute, const F32& value, const U32& vertex) {
     Inst(Opcode::SetAttribute, attribute, value, vertex);
 }
diff --git a/src/shader_recompiler/frontend/ir/ir_emitter.h b/src/shader_recompiler/frontend/ir/ir_emitter.h
index 2df992feb..7aaaa4ab0 100644
--- a/src/shader_recompiler/frontend/ir/ir_emitter.h
+++ b/src/shader_recompiler/frontend/ir/ir_emitter.h
@@ -74,6 +74,8 @@ public:
 
     [[nodiscard]] F32 GetAttribute(IR::Attribute attribute);
     [[nodiscard]] F32 GetAttribute(IR::Attribute attribute, const U32& vertex);
+    [[nodiscard]] U32 GetAttributeU32(IR::Attribute attribute);
+    [[nodiscard]] U32 GetAttributeU32(IR::Attribute attribute, const U32& vertex);
     void SetAttribute(IR::Attribute attribute, const F32& value, const U32& vertex);
 
     [[nodiscard]] F32 GetAttributeIndexed(const U32& phys_address);
diff --git a/src/shader_recompiler/frontend/maxwell/translate_program.cpp b/src/shader_recompiler/frontend/maxwell/translate_program.cpp
index 3adbd2b16..ac159d24b 100644
--- a/src/shader_recompiler/frontend/maxwell/translate_program.cpp
+++ b/src/shader_recompiler/frontend/maxwell/translate_program.cpp
@@ -219,7 +219,7 @@ IR::Program TranslateProgram(ObjectPool<IR::Inst>& inst_pool, ObjectPool<IR::Blo
     }
     Optimization::SsaRewritePass(program);
 
-    Optimization::ConstantPropagationPass(program);
+    Optimization::ConstantPropagationPass(env, program);
 
     Optimization::PositionPass(env, program);
 
diff --git a/src/shader_recompiler/ir_opt/constant_propagation_pass.cpp b/src/shader_recompiler/ir_opt/constant_propagation_pass.cpp
index 826f9a54a..4d81e9336 100644
--- a/src/shader_recompiler/ir_opt/constant_propagation_pass.cpp
+++ b/src/shader_recompiler/ir_opt/constant_propagation_pass.cpp
@@ -7,6 +7,7 @@
 #include <type_traits>
 
 #include "common/bit_cast.h"
+#include "shader_recompiler/environment.h"
 #include "shader_recompiler/exception.h"
 #include "shader_recompiler/frontend/ir/ir_emitter.h"
 #include "shader_recompiler/frontend/ir/value.h"
@@ -515,6 +516,9 @@ void FoldBitCast(IR::Inst& inst, IR::Opcode reverse) {
             case IR::Attribute::PrimitiveId:
             case IR::Attribute::InstanceId:
             case IR::Attribute::VertexId:
+            case IR::Attribute::BaseVertex:
+            case IR::Attribute::BaseInstance:
+            case IR::Attribute::DrawID:
                 break;
             default:
                 return;
@@ -644,7 +648,63 @@ void FoldFSwizzleAdd(IR::Block& block, IR::Inst& inst) {
     }
 }
 
-void ConstantPropagation(IR::Block& block, IR::Inst& inst) {
+void FoldConstBuffer(Environment& env, IR::Block& block, IR::Inst& inst) {
+    const IR::Value bank{inst.Arg(0)};
+    const IR::Value offset{inst.Arg(1)};
+    if (!bank.IsImmediate() || !offset.IsImmediate()) {
+        return;
+    }
+    const auto bank_value = bank.U32();
+    const auto offset_value = offset.U32();
+    auto replacement = env.GetReplaceConstBuffer(bank_value, offset_value);
+    if (!replacement) {
+        return;
+    }
+    const auto new_attribute = [replacement]() {
+        switch (*replacement) {
+        case ReplaceConstant::BaseInstance:
+            return IR::Attribute::BaseInstance;
+        case ReplaceConstant::BaseVertex:
+            return IR::Attribute::BaseVertex;
+        case ReplaceConstant::DrawID:
+            return IR::Attribute::DrawID;
+        default:
+            throw NotImplementedException("Not implemented replacement variable {}", *replacement);
+        }
+    }();
+    IR::IREmitter ir{block, IR::Block::InstructionList::s_iterator_to(inst)};
+    if (inst.GetOpcode() == IR::Opcode::GetCbufU32) {
+        inst.ReplaceUsesWith(ir.GetAttributeU32(new_attribute));
+    } else {
+        inst.ReplaceUsesWith(ir.GetAttribute(new_attribute));
+    }
+}
+
+void FoldDriverConstBuffer(Environment& env, IR::Block& block, IR::Inst& inst, u32 which_bank,
+                           u32 offset_start = 0, u32 offset_end = std::numeric_limits<u16>::max()) {
+    const IR::Value bank{inst.Arg(0)};
+    const IR::Value offset{inst.Arg(1)};
+    if (!bank.IsImmediate() || !offset.IsImmediate()) {
+        return;
+    }
+    const auto bank_value = bank.U32();
+    if (bank_value != which_bank) {
+        return;
+    }
+    const auto offset_value = offset.U32();
+    if (offset_value < offset_start || offset_value >= offset_end) {
+        return;
+    }
+    IR::IREmitter ir{block, IR::Block::InstructionList::s_iterator_to(inst)};
+    if (inst.GetOpcode() == IR::Opcode::GetCbufU32) {
+        inst.ReplaceUsesWith(IR::Value{env.ReadCbufValue(bank_value, offset_value)});
+    } else {
+        inst.ReplaceUsesWith(
+            IR::Value{Common::BitCast<f32>(env.ReadCbufValue(bank_value, offset_value))});
+    }
+}
+
+void ConstantPropagation(Environment& env, IR::Block& block, IR::Inst& inst) {
     switch (inst.GetOpcode()) {
     case IR::Opcode::GetRegister:
         return FoldGetRegister(inst);
@@ -789,18 +849,28 @@ void ConstantPropagation(IR::Block& block, IR::Inst& inst) {
                                     IR::Opcode::CompositeInsertF16x4);
     case IR::Opcode::FSwizzleAdd:
         return FoldFSwizzleAdd(block, inst);
+    case IR::Opcode::GetCbufF32:
+    case IR::Opcode::GetCbufU32:
+        if (env.HasHLEMacroState()) {
+            FoldConstBuffer(env, block, inst);
+        }
+        if (env.IsPropietaryDriver()) {
+            FoldDriverConstBuffer(env, block, inst, 1);
+        }
+        break;
     default:
         break;
     }
 }
+
 } // Anonymous namespace
 
-void ConstantPropagationPass(IR::Program& program) {
+void ConstantPropagationPass(Environment& env, IR::Program& program) {
     const auto end{program.post_order_blocks.rend()};
     for (auto it = program.post_order_blocks.rbegin(); it != end; ++it) {
         IR::Block* const block{*it};
         for (IR::Inst& inst : block->Instructions()) {
-            ConstantPropagation(*block, inst);
+            ConstantPropagation(env, *block, inst);
         }
     }
 }
diff --git a/src/shader_recompiler/ir_opt/passes.h b/src/shader_recompiler/ir_opt/passes.h
index 11bfe801a..1f8f2ba95 100644
--- a/src/shader_recompiler/ir_opt/passes.h
+++ b/src/shader_recompiler/ir_opt/passes.h
@@ -13,7 +13,7 @@ struct HostTranslateInfo;
 namespace Shader::Optimization {
 
 void CollectShaderInfoPass(Environment& env, IR::Program& program);
-void ConstantPropagationPass(IR::Program& program);
+void ConstantPropagationPass(Environment& env, IR::Program& program);
 void DeadCodeEliminationPass(IR::Program& program);
 void GlobalMemoryToStorageBufferPass(IR::Program& program);
 void IdentityRemovalPass(IR::Program& program);
diff --git a/src/shader_recompiler/shader_info.h b/src/shader_recompiler/shader_info.h
index d9c6e92db..44236b6b1 100644
--- a/src/shader_recompiler/shader_info.h
+++ b/src/shader_recompiler/shader_info.h
@@ -16,6 +16,12 @@
 
 namespace Shader {
 
+enum class ReplaceConstant : u32 {
+    BaseInstance,
+    BaseVertex,
+    DrawID,
+};
+
 enum class TextureType : u32 {
     Color1D,
     ColorArray1D,
diff --git a/src/shader_recompiler/varying_state.h b/src/shader_recompiler/varying_state.h
index 7b28a285f..18a9aaf50 100644
--- a/src/shader_recompiler/varying_state.h
+++ b/src/shader_recompiler/varying_state.h
@@ -11,7 +11,7 @@
 namespace Shader {
 
 struct VaryingState {
-    std::bitset<256> mask{};
+    std::bitset<512> mask{};
 
     void Set(IR::Attribute attribute, bool state = true) {
         mask[static_cast<size_t>(attribute)] = state;
diff --git a/src/tests/CMakeLists.txt b/src/tests/CMakeLists.txt
index 6a4022e45..9b65e79cb 100644
--- a/src/tests/CMakeLists.txt
+++ b/src/tests/CMakeLists.txt
@@ -7,6 +7,7 @@ add_executable(tests
     common/fibers.cpp
     common/host_memory.cpp
     common/param_package.cpp
+    common/range_map.cpp
     common/ring_buffer.cpp
     common/scratch_buffer.cpp
     common/unique_function.cpp
diff --git a/src/tests/common/range_map.cpp b/src/tests/common/range_map.cpp
new file mode 100644
index 000000000..5a4630a38
--- /dev/null
+++ b/src/tests/common/range_map.cpp
@@ -0,0 +1,70 @@
+// SPDX-FileCopyrightText: Copyright 2022 yuzu Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
+#include <stdexcept>
+
+#include <catch2/catch.hpp>
+
+#include "common/range_map.h"
+
+enum class MappedEnum : u32 {
+    Invalid = 0,
+    Valid_1 = 1,
+    Valid_2 = 2,
+    Valid_3 = 3,
+};
+
+TEST_CASE("Range Map: Setup", "[video_core]") {
+    Common::RangeMap<u64, MappedEnum> my_map(MappedEnum::Invalid);
+    my_map.Map(3000, 3500, MappedEnum::Valid_1);
+    my_map.Unmap(3200, 3600);
+    my_map.Map(4000, 4500, MappedEnum::Valid_2);
+    my_map.Map(4200, 4400, MappedEnum::Valid_2);
+    my_map.Map(4200, 4400, MappedEnum::Valid_1);
+    REQUIRE(my_map.GetContinousSizeFrom(4200) == 200);
+    REQUIRE(my_map.GetContinousSizeFrom(3000) == 200);
+    REQUIRE(my_map.GetContinousSizeFrom(2900) == 0);
+
+    REQUIRE(my_map.GetValueAt(2900) == MappedEnum::Invalid);
+    REQUIRE(my_map.GetValueAt(3100) == MappedEnum::Valid_1);
+    REQUIRE(my_map.GetValueAt(3000) == MappedEnum::Valid_1);
+    REQUIRE(my_map.GetValueAt(3200) == MappedEnum::Invalid);
+
+    REQUIRE(my_map.GetValueAt(4199) == MappedEnum::Valid_2);
+    REQUIRE(my_map.GetValueAt(4200) == MappedEnum::Valid_1);
+    REQUIRE(my_map.GetValueAt(4400) == MappedEnum::Valid_2);
+    REQUIRE(my_map.GetValueAt(4500) == MappedEnum::Invalid);
+    REQUIRE(my_map.GetValueAt(4600) == MappedEnum::Invalid);
+
+    my_map.Unmap(0, 6000);
+    for (u64 address = 0; address < 10000; address += 1000) {
+        REQUIRE(my_map.GetContinousSizeFrom(address) == 0);
+    }
+
+    my_map.Map(1000, 3000, MappedEnum::Valid_1);
+    my_map.Map(4000, 5000, MappedEnum::Valid_1);
+    my_map.Map(2500, 4100, MappedEnum::Valid_1);
+    REQUIRE(my_map.GetContinousSizeFrom(1000) == 4000);
+
+    my_map.Map(1000, 3000, MappedEnum::Valid_1);
+    my_map.Map(4000, 5000, MappedEnum::Valid_2);
+    my_map.Map(2500, 4100, MappedEnum::Valid_3);
+    REQUIRE(my_map.GetContinousSizeFrom(1000) == 1500);
+    REQUIRE(my_map.GetContinousSizeFrom(2500) == 1600);
+    REQUIRE(my_map.GetContinousSizeFrom(4100) == 900);
+    REQUIRE(my_map.GetValueAt(900) == MappedEnum::Invalid);
+    REQUIRE(my_map.GetValueAt(1000) == MappedEnum::Valid_1);
+    REQUIRE(my_map.GetValueAt(2500) == MappedEnum::Valid_3);
+    REQUIRE(my_map.GetValueAt(4100) == MappedEnum::Valid_2);
+    REQUIRE(my_map.GetValueAt(5000) == MappedEnum::Invalid);
+
+    my_map.Map(2000, 6000, MappedEnum::Valid_3);
+    REQUIRE(my_map.GetContinousSizeFrom(1000) == 1000);
+    REQUIRE(my_map.GetContinousSizeFrom(3000) == 3000);
+    REQUIRE(my_map.GetValueAt(1000) == MappedEnum::Valid_1);
+    REQUIRE(my_map.GetValueAt(1999) == MappedEnum::Valid_1);
+    REQUIRE(my_map.GetValueAt(1500) == MappedEnum::Valid_1);
+    REQUIRE(my_map.GetValueAt(2001) == MappedEnum::Valid_3);
+    REQUIRE(my_map.GetValueAt(5999) == MappedEnum::Valid_3);
+    REQUIRE(my_map.GetValueAt(6000) == MappedEnum::Invalid);
+}
diff --git a/src/video_core/CMakeLists.txt b/src/video_core/CMakeLists.txt
index fd71bf186..aa271a377 100644
--- a/src/video_core/CMakeLists.txt
+++ b/src/video_core/CMakeLists.txt
@@ -13,6 +13,7 @@ add_library(video_core STATIC
     buffer_cache/buffer_base.h
     buffer_cache/buffer_cache.cpp
     buffer_cache/buffer_cache.h
+    cache_types.h
     cdma_pusher.cpp
     cdma_pusher.h
     compatible_formats.cpp
diff --git a/src/video_core/buffer_cache/buffer_cache.h b/src/video_core/buffer_cache/buffer_cache.h
index f1c60d1f3..06fd40851 100644
--- a/src/video_core/buffer_cache/buffer_cache.h
+++ b/src/video_core/buffer_cache/buffer_cache.h
@@ -200,7 +200,16 @@ public:
     /// Return true when a CPU region is modified from the CPU
     [[nodiscard]] bool IsRegionCpuModified(VAddr addr, size_t size);
 
-    std::mutex mutex;
+    void SetDrawIndirect(
+        const Tegra::Engines::DrawManager::IndirectParams* current_draw_indirect_) {
+        current_draw_indirect = current_draw_indirect_;
+    }
+
+    [[nodiscard]] std::pair<Buffer*, u32> GetDrawIndirectCount();
+
+    [[nodiscard]] std::pair<Buffer*, u32> GetDrawIndirectBuffer();
+
+    std::recursive_mutex mutex;
     Runtime& runtime;
 
 private:
@@ -272,6 +281,8 @@ private:
 
     void BindHostVertexBuffers();
 
+    void BindHostDrawIndirectBuffers();
+
     void BindHostGraphicsUniformBuffers(size_t stage);
 
     void BindHostGraphicsUniformBuffer(size_t stage, u32 index, u32 binding_index, bool needs_bind);
@@ -298,6 +309,8 @@ private:
 
     void UpdateVertexBuffer(u32 index);
 
+    void UpdateDrawIndirect();
+
     void UpdateUniformBuffers(size_t stage);
 
     void UpdateStorageBuffers(size_t stage);
@@ -372,6 +385,8 @@ private:
     SlotVector<Buffer> slot_buffers;
     DelayedDestructionRing<Buffer, 8> delayed_destruction_ring;
 
+    const Tegra::Engines::DrawManager::IndirectParams* current_draw_indirect{};
+
     u32 last_index_count = 0;
 
     Binding index_buffer;
@@ -380,6 +395,8 @@ private:
     std::array<std::array<Binding, NUM_STORAGE_BUFFERS>, NUM_STAGES> storage_buffers;
     std::array<std::array<TextureBufferBinding, NUM_TEXTURE_BUFFERS>, NUM_STAGES> texture_buffers;
     std::array<Binding, NUM_TRANSFORM_FEEDBACK_BUFFERS> transform_feedback_buffers;
+    Binding count_buffer_binding;
+    Binding indirect_buffer_binding;
 
     std::array<Binding, NUM_COMPUTE_UNIFORM_BUFFERS> compute_uniform_buffers;
     std::array<Binding, NUM_STORAGE_BUFFERS> compute_storage_buffers;
@@ -674,6 +691,9 @@ void BufferCache<P>::BindHostGeometryBuffers(bool is_indexed) {
     }
     BindHostVertexBuffers();
     BindHostTransformFeedbackBuffers();
+    if (current_draw_indirect) {
+        BindHostDrawIndirectBuffers();
+    }
 }
 
 template <class P>
@@ -823,6 +843,7 @@ bool BufferCache<P>::ShouldWaitAsyncFlushes() const noexcept {
 template <class P>
 void BufferCache<P>::CommitAsyncFlushesHigh() {
     AccumulateFlushes();
+
     if (committed_ranges.empty()) {
         return;
     }
@@ -869,7 +890,7 @@ void BufferCache<P>::CommitAsyncFlushesHigh() {
                                 buffer_id,
                             });
                             // Align up to avoid cache conflicts
-                            constexpr u64 align = 256ULL;
+                            constexpr u64 align = 8ULL;
                             constexpr u64 mask = ~(align - 1ULL);
                             total_size_bytes += (new_size + align - 1) & mask;
                             largest_copy = std::max(largest_copy, new_size);
@@ -1042,6 +1063,19 @@ void BufferCache<P>::BindHostVertexBuffers() {
 }
 
 template <class P>
+void BufferCache<P>::BindHostDrawIndirectBuffers() {
+    const auto bind_buffer = [this](const Binding& binding) {
+        Buffer& buffer = slot_buffers[binding.buffer_id];
+        TouchBuffer(buffer, binding.buffer_id);
+        SynchronizeBuffer(buffer, binding.cpu_addr, binding.size);
+    };
+    if (current_draw_indirect->include_count) {
+        bind_buffer(count_buffer_binding);
+    }
+    bind_buffer(indirect_buffer_binding);
+}
+
+template <class P>
 void BufferCache<P>::BindHostGraphicsUniformBuffers(size_t stage) {
     u32 dirty = ~0U;
     if constexpr (HAS_PERSISTENT_UNIFORM_BUFFER_BINDINGS) {
@@ -1272,6 +1306,9 @@ void BufferCache<P>::DoUpdateGraphicsBuffers(bool is_indexed) {
             UpdateStorageBuffers(stage);
             UpdateTextureBuffers(stage);
         }
+        if (current_draw_indirect) {
+            UpdateDrawIndirect();
+        }
     } while (has_deleted_buffers);
 }
 
@@ -1289,7 +1326,7 @@ void BufferCache<P>::UpdateIndexBuffer() {
     const auto& draw_state = maxwell3d->draw_manager->GetDrawState();
     const auto& index_array = draw_state.index_buffer;
     auto& flags = maxwell3d->dirty.flags;
-    if (!flags[Dirty::IndexBuffer] && last_index_count == index_array.count) {
+    if (!flags[Dirty::IndexBuffer]) {
         return;
     }
     flags[Dirty::IndexBuffer] = false;
@@ -1362,6 +1399,27 @@ void BufferCache<P>::UpdateVertexBuffer(u32 index) {
 }
 
 template <class P>
+void BufferCache<P>::UpdateDrawIndirect() {
+    const auto update = [this](GPUVAddr gpu_addr, size_t size, Binding& binding) {
+        const std::optional<VAddr> cpu_addr = gpu_memory->GpuToCpuAddress(gpu_addr);
+        if (!cpu_addr) {
+            binding = NULL_BINDING;
+            return;
+        }
+        binding = Binding{
+            .cpu_addr = *cpu_addr,
+            .size = static_cast<u32>(size),
+            .buffer_id = FindBuffer(*cpu_addr, static_cast<u32>(size)),
+        };
+    };
+    if (current_draw_indirect->include_count) {
+        update(current_draw_indirect->count_start_address, sizeof(u32), count_buffer_binding);
+    }
+    update(current_draw_indirect->indirect_start_address, current_draw_indirect->buffer_size,
+           indirect_buffer_binding);
+}
+
+template <class P>
 void BufferCache<P>::UpdateUniformBuffers(size_t stage) {
     ForEachEnabledBit(enabled_uniform_buffer_masks[stage], [&](u32 index) {
         Binding& binding = uniform_buffers[stage][index];
@@ -1941,4 +1999,16 @@ bool BufferCache<P>::HasFastUniformBufferBound(size_t stage, u32 binding_index)
     }
 }
 
+template <class P>
+std::pair<typename BufferCache<P>::Buffer*, u32> BufferCache<P>::GetDrawIndirectCount() {
+    auto& buffer = slot_buffers[count_buffer_binding.buffer_id];
+    return std::make_pair(&buffer, buffer.Offset(count_buffer_binding.cpu_addr));
+}
+
+template <class P>
+std::pair<typename BufferCache<P>::Buffer*, u32> BufferCache<P>::GetDrawIndirectBuffer() {
+    auto& buffer = slot_buffers[indirect_buffer_binding.buffer_id];
+    return std::make_pair(&buffer, buffer.Offset(indirect_buffer_binding.cpu_addr));
+}
+
 } // namespace VideoCommon
diff --git a/src/video_core/cache_types.h b/src/video_core/cache_types.h
new file mode 100644
index 000000000..1a5db3c55
--- /dev/null
+++ b/src/video_core/cache_types.h
@@ -0,0 +1,24 @@
+// SPDX-FileCopyrightText: Copyright 2020 yuzu Emulator Project
+// SPDX-License-Identifier: GPL-2.0-or-later
+
+#pragma once
+
+#include "common/common_funcs.h"
+#include "common/common_types.h"
+
+namespace VideoCommon {
+
+enum class CacheType : u32 {
+    None = 0,
+    TextureCache = 1 << 0,
+    QueryCache = 1 << 1,
+    BufferCache = 1 << 2,
+    ShaderCache = 1 << 3,
+    NoTextureCache = QueryCache | BufferCache | ShaderCache,
+    NoBufferCache = TextureCache | QueryCache | ShaderCache,
+    NoQueryCache = TextureCache | BufferCache | ShaderCache,
+    All = TextureCache | QueryCache | BufferCache | ShaderCache,
+};
+DECLARE_ENUM_FLAG_OPERATORS(CacheType)
+
+} // namespace VideoCommon
diff --git a/src/video_core/dma_pusher.cpp b/src/video_core/dma_pusher.cpp
index 322de2606..551929824 100644
--- a/src/video_core/dma_pusher.cpp
+++ b/src/video_core/dma_pusher.cpp
@@ -61,7 +61,7 @@ bool DmaPusher::Step() {
     } else {
         const CommandListHeader command_list_header{
             command_list.command_lists[dma_pushbuffer_subindex++]};
-        const GPUVAddr dma_get = command_list_header.addr;
+        dma_state.dma_get = command_list_header.addr;
 
         if (dma_pushbuffer_subindex >= command_list.command_lists.size()) {
             // We've gone through the current list, remove it from the queue
@@ -75,12 +75,22 @@ bool DmaPusher::Step() {
 
         // Push buffer non-empty, read a word
         command_headers.resize_destructive(command_list_header.size);
-        if (Settings::IsGPULevelHigh()) {
-            memory_manager.ReadBlock(dma_get, command_headers.data(),
-                                     command_list_header.size * sizeof(u32));
+        constexpr u32 MacroRegistersStart = 0xE00;
+        if (dma_state.method < MacroRegistersStart) {
+            if (Settings::IsGPULevelHigh()) {
+                memory_manager.ReadBlock(dma_state.dma_get, command_headers.data(),
+                                         command_list_header.size * sizeof(u32));
+            } else {
+                memory_manager.ReadBlockUnsafe(dma_state.dma_get, command_headers.data(),
+                                               command_list_header.size * sizeof(u32));
+            }
         } else {
-            memory_manager.ReadBlockUnsafe(dma_get, command_headers.data(),
-                                           command_list_header.size * sizeof(u32));
+            const size_t copy_size = command_list_header.size * sizeof(u32);
+            if (subchannels[dma_state.subchannel]) {
+                subchannels[dma_state.subchannel]->current_dirty =
+                    memory_manager.IsMemoryDirty(dma_state.dma_get, copy_size);
+            }
+            memory_manager.ReadBlockUnsafe(dma_state.dma_get, command_headers.data(), copy_size);
         }
         ProcessCommands(command_headers);
     }
@@ -94,6 +104,7 @@ void DmaPusher::ProcessCommands(std::span<const CommandHeader> commands) {
 
         if (dma_state.method_count) {
             // Data word of methods command
+            dma_state.dma_word_offset = static_cast<u32>(index * sizeof(u32));
             if (dma_state.non_incrementing) {
                 const u32 max_write = static_cast<u32>(
                     std::min<std::size_t>(index + dma_state.method_count, commands.size()) - index);
@@ -132,6 +143,8 @@ void DmaPusher::ProcessCommands(std::span<const CommandHeader> commands) {
             case SubmissionMode::Inline:
                 dma_state.method = command_header.method;
                 dma_state.subchannel = command_header.subchannel;
+                dma_state.dma_word_offset = static_cast<u64>(
+                    -static_cast<s64>(dma_state.dma_get)); // negate to set address as 0
                 CallMethod(command_header.arg_count);
                 dma_state.non_incrementing = true;
                 dma_increment_once = false;
@@ -164,8 +177,14 @@ void DmaPusher::CallMethod(u32 argument) const {
             dma_state.method_count,
         });
     } else {
-        subchannels[dma_state.subchannel]->CallMethod(dma_state.method, argument,
-                                                      dma_state.is_last_call);
+        auto subchannel = subchannels[dma_state.subchannel];
+        if (!subchannel->execution_mask[dma_state.method]) [[likely]] {
+            subchannel->method_sink.emplace_back(dma_state.method, argument);
+            return;
+        }
+        subchannel->ConsumeSink();
+        subchannel->current_dma_segment = dma_state.dma_get + dma_state.dma_word_offset;
+        subchannel->CallMethod(dma_state.method, argument, dma_state.is_last_call);
     }
 }
 
@@ -174,8 +193,11 @@ void DmaPusher::CallMultiMethod(const u32* base_start, u32 num_methods) const {
         puller.CallMultiMethod(dma_state.method, dma_state.subchannel, base_start, num_methods,
                                dma_state.method_count);
     } else {
-        subchannels[dma_state.subchannel]->CallMultiMethod(dma_state.method, base_start,
-                                                           num_methods, dma_state.method_count);
+        auto subchannel = subchannels[dma_state.subchannel];
+        subchannel->ConsumeSink();
+        subchannel->current_dma_segment = dma_state.dma_get + dma_state.dma_word_offset;
+        subchannel->CallMultiMethod(dma_state.method, base_start, num_methods,
+                                    dma_state.method_count);
     }
 }
 
diff --git a/src/video_core/dma_pusher.h b/src/video_core/dma_pusher.h
index 6f00de937..1cdb690ed 100644
--- a/src/video_core/dma_pusher.h
+++ b/src/video_core/dma_pusher.h
@@ -156,6 +156,8 @@ private:
         u32 subchannel;        ///< Current subchannel
         u32 method_count;      ///< Current method count
         u32 length_pending;    ///< Large NI command length pending
+        GPUVAddr dma_get;      ///< Currently read segment
+        u64 dma_word_offset;   ///< Current word ofset from address
         bool non_incrementing; ///< Current command's NI flag
         bool is_last_call;
     };
diff --git a/src/video_core/engines/draw_manager.cpp b/src/video_core/engines/draw_manager.cpp
index 3a78421f6..2437121ce 100644
--- a/src/video_core/engines/draw_manager.cpp
+++ b/src/video_core/engines/draw_manager.cpp
@@ -91,6 +91,23 @@ void DrawManager::DrawIndex(PrimitiveTopology topology, u32 index_first, u32 ind
     ProcessDraw(true, num_instances);
 }
 
+void DrawManager::DrawArrayIndirect(PrimitiveTopology topology) {
+    draw_state.topology = topology;
+
+    ProcessDrawIndirect();
+}
+
+void DrawManager::DrawIndexedIndirect(PrimitiveTopology topology, u32 index_first,
+                                      u32 index_count) {
+    const auto& regs{maxwell3d->regs};
+    draw_state.topology = topology;
+    draw_state.index_buffer = regs.index_buffer;
+    draw_state.index_buffer.first = index_first;
+    draw_state.index_buffer.count = index_count;
+
+    ProcessDrawIndirect();
+}
+
 void DrawManager::SetInlineIndexBuffer(u32 index) {
     draw_state.inline_index_draw_indexes.push_back(static_cast<u8>(index & 0x000000ff));
     draw_state.inline_index_draw_indexes.push_back(static_cast<u8>((index & 0x0000ff00) >> 8));
@@ -198,4 +215,18 @@ void DrawManager::ProcessDraw(bool draw_indexed, u32 instance_count) {
         maxwell3d->rasterizer->Draw(draw_indexed, instance_count);
     }
 }
+
+void DrawManager::ProcessDrawIndirect() {
+    LOG_TRACE(
+        HW_GPU,
+        "called, topology={}, is_indexed={}, includes_count={}, buffer_size={}, max_draw_count={}",
+        draw_state.topology, indirect_state.is_indexed, indirect_state.include_count,
+        indirect_state.buffer_size, indirect_state.max_draw_counts);
+
+    UpdateTopology();
+
+    if (maxwell3d->ShouldExecute()) {
+        maxwell3d->rasterizer->DrawIndirect();
+    }
+}
 } // namespace Tegra::Engines
diff --git a/src/video_core/engines/draw_manager.h b/src/video_core/engines/draw_manager.h
index 0e6930a9c..58d1b2d59 100644
--- a/src/video_core/engines/draw_manager.h
+++ b/src/video_core/engines/draw_manager.h
@@ -32,6 +32,16 @@ public:
         std::vector<u8> inline_index_draw_indexes;
     };
 
+    struct IndirectParams {
+        bool is_indexed;
+        bool include_count;
+        GPUVAddr count_start_address;
+        GPUVAddr indirect_start_address;
+        size_t buffer_size;
+        size_t max_draw_counts;
+        size_t stride;
+    };
+
     explicit DrawManager(Maxwell3D* maxwell_3d);
 
     void ProcessMethodCall(u32 method, u32 argument);
@@ -46,10 +56,22 @@ public:
     void DrawIndex(PrimitiveTopology topology, u32 index_first, u32 index_count, u32 base_index,
                    u32 base_instance, u32 num_instances);
 
+    void DrawArrayIndirect(PrimitiveTopology topology);
+
+    void DrawIndexedIndirect(PrimitiveTopology topology, u32 index_first, u32 index_count);
+
     const State& GetDrawState() const {
         return draw_state;
     }
 
+    IndirectParams& GetIndirectParams() {
+        return indirect_state;
+    }
+
+    const IndirectParams& GetIndirectParams() const {
+        return indirect_state;
+    }
+
 private:
     void SetInlineIndexBuffer(u32 index);
 
@@ -63,7 +85,10 @@ private:
 
     void ProcessDraw(bool draw_indexed, u32 instance_count);
 
+    void ProcessDrawIndirect();
+
     Maxwell3D* maxwell3d{};
     State draw_state{};
+    IndirectParams indirect_state{};
 };
 } // namespace Tegra::Engines
diff --git a/src/video_core/engines/engine_interface.h b/src/video_core/engines/engine_interface.h
index 26cde8584..392322358 100644
--- a/src/video_core/engines/engine_interface.h
+++ b/src/video_core/engines/engine_interface.h
@@ -3,6 +3,10 @@
 
 #pragma once
 
+#include <bitset>
+#include <limits>
+#include <vector>
+
 #include "common/common_types.h"
 
 namespace Tegra::Engines {
@@ -17,6 +21,26 @@ public:
     /// Write multiple values to the register identified by method.
     virtual void CallMultiMethod(u32 method, const u32* base_start, u32 amount,
                                  u32 methods_pending) = 0;
+
+    void ConsumeSink() {
+        if (method_sink.empty()) {
+            return;
+        }
+        ConsumeSinkImpl();
+    }
+
+    std::bitset<std::numeric_limits<u16>::max()> execution_mask{};
+    std::vector<std::pair<u32, u32>> method_sink{};
+    bool current_dirty{};
+    GPUVAddr current_dma_segment;
+
+protected:
+    virtual void ConsumeSinkImpl() {
+        for (auto [method, value] : method_sink) {
+            CallMethod(method, value, true);
+        }
+        method_sink.clear();
+    }
 };
 
 } // namespace Tegra::Engines
diff --git a/src/video_core/engines/fermi_2d.cpp b/src/video_core/engines/fermi_2d.cpp
index c6478ae85..e655e7254 100644
--- a/src/video_core/engines/fermi_2d.cpp
+++ b/src/video_core/engines/fermi_2d.cpp
@@ -25,6 +25,9 @@ Fermi2D::Fermi2D(MemoryManager& memory_manager_) {
     // Nvidia's OpenGL driver seems to assume these values
     regs.src.depth = 1;
     regs.dst.depth = 1;
+
+    execution_mask.reset();
+    execution_mask[FERMI2D_REG_INDEX(pixels_from_memory.src_y0) + 1] = true;
 }
 
 Fermi2D::~Fermi2D() = default;
@@ -49,6 +52,13 @@ void Fermi2D::CallMultiMethod(u32 method, const u32* base_start, u32 amount, u32
     }
 }
 
+void Fermi2D::ConsumeSinkImpl() {
+    for (auto [method, value] : method_sink) {
+        regs.reg_array[method] = value;
+    }
+    method_sink.clear();
+}
+
 void Fermi2D::Blit() {
     MICROPROFILE_SCOPE(GPU_BlitEngine);
     LOG_DEBUG(HW_GPU, "called. source address=0x{:x}, destination address=0x{:x}",
diff --git a/src/video_core/engines/fermi_2d.h b/src/video_core/engines/fermi_2d.h
index 100b21bac..523fbdec2 100644
--- a/src/video_core/engines/fermi_2d.h
+++ b/src/video_core/engines/fermi_2d.h
@@ -309,6 +309,8 @@ private:
     /// Performs the copy from the source surface to the destination surface as configured in the
     /// registers.
     void Blit();
+
+    void ConsumeSinkImpl() override;
 };
 
 #define ASSERT_REG_POSITION(field_name, position)                                                  \
diff --git a/src/video_core/engines/kepler_compute.cpp b/src/video_core/engines/kepler_compute.cpp
index e5c622155..601095f03 100644
--- a/src/video_core/engines/kepler_compute.cpp
+++ b/src/video_core/engines/kepler_compute.cpp
@@ -14,7 +14,12 @@
 namespace Tegra::Engines {
 
 KeplerCompute::KeplerCompute(Core::System& system_, MemoryManager& memory_manager_)
-    : system{system_}, memory_manager{memory_manager_}, upload_state{memory_manager, regs.upload} {}
+    : system{system_}, memory_manager{memory_manager_}, upload_state{memory_manager, regs.upload} {
+    execution_mask.reset();
+    execution_mask[KEPLER_COMPUTE_REG_INDEX(exec_upload)] = true;
+    execution_mask[KEPLER_COMPUTE_REG_INDEX(data_upload)] = true;
+    execution_mask[KEPLER_COMPUTE_REG_INDEX(launch)] = true;
+}
 
 KeplerCompute::~KeplerCompute() = default;
 
@@ -23,6 +28,13 @@ void KeplerCompute::BindRasterizer(VideoCore::RasterizerInterface* rasterizer_)
     upload_state.BindRasterizer(rasterizer);
 }
 
+void KeplerCompute::ConsumeSinkImpl() {
+    for (auto [method, value] : method_sink) {
+        regs.reg_array[method] = value;
+    }
+    method_sink.clear();
+}
+
 void KeplerCompute::CallMethod(u32 method, u32 method_argument, bool is_last_call) {
     ASSERT_MSG(method < Regs::NUM_REGS,
                "Invalid KeplerCompute register, increase the size of the Regs structure");
diff --git a/src/video_core/engines/kepler_compute.h b/src/video_core/engines/kepler_compute.h
index e154e3f06..2092e685f 100644
--- a/src/video_core/engines/kepler_compute.h
+++ b/src/video_core/engines/kepler_compute.h
@@ -204,6 +204,8 @@ public:
 private:
     void ProcessLaunch();
 
+    void ConsumeSinkImpl() override;
+
     /// Retrieves information about a specific TIC entry from the TIC buffer.
     Texture::TICEntry GetTICEntry(u32 tic_index) const;
 
diff --git a/src/video_core/engines/kepler_memory.cpp b/src/video_core/engines/kepler_memory.cpp
index 08045d1cf..c026801a3 100644
--- a/src/video_core/engines/kepler_memory.cpp
+++ b/src/video_core/engines/kepler_memory.cpp
@@ -18,6 +18,17 @@ KeplerMemory::~KeplerMemory() = default;
 
 void KeplerMemory::BindRasterizer(VideoCore::RasterizerInterface* rasterizer_) {
     upload_state.BindRasterizer(rasterizer_);
+
+    execution_mask.reset();
+    execution_mask[KEPLERMEMORY_REG_INDEX(exec)] = true;
+    execution_mask[KEPLERMEMORY_REG_INDEX(data)] = true;
+}
+
+void KeplerMemory::ConsumeSinkImpl() {
+    for (auto [method, value] : method_sink) {
+        regs.reg_array[method] = value;
+    }
+    method_sink.clear();
 }
 
 void KeplerMemory::CallMethod(u32 method, u32 method_argument, bool is_last_call) {
diff --git a/src/video_core/engines/kepler_memory.h b/src/video_core/engines/kepler_memory.h
index 5fe7489f0..fb1eecbba 100644
--- a/src/video_core/engines/kepler_memory.h
+++ b/src/video_core/engines/kepler_memory.h
@@ -73,6 +73,8 @@ public:
     } regs{};
 
 private:
+    void ConsumeSinkImpl() override;
+
     Core::System& system;
     Upload::State upload_state;
 };
diff --git a/src/video_core/engines/maxwell_3d.cpp b/src/video_core/engines/maxwell_3d.cpp
index 9b182b653..fbfd1ddd2 100644
--- a/src/video_core/engines/maxwell_3d.cpp
+++ b/src/video_core/engines/maxwell_3d.cpp
@@ -4,6 +4,8 @@
 #include <cstring>
 #include <optional>
 #include "common/assert.h"
+#include "common/scope_exit.h"
+#include "common/settings.h"
 #include "core/core.h"
 #include "core/core_timing.h"
 #include "video_core/dirty_flags.h"
@@ -28,6 +30,10 @@ Maxwell3D::Maxwell3D(Core::System& system_, MemoryManager& memory_manager_)
                                                                                 regs.upload} {
     dirty.flags.flip();
     InitializeRegisterDefaults();
+    execution_mask.reset();
+    for (size_t i = 0; i < execution_mask.size(); i++) {
+        execution_mask[i] = IsMethodExecutable(static_cast<u32>(i));
+    }
 }
 
 Maxwell3D::~Maxwell3D() = default;
@@ -121,6 +127,71 @@ void Maxwell3D::InitializeRegisterDefaults() {
     shadow_state = regs;
 }
 
+bool Maxwell3D::IsMethodExecutable(u32 method) {
+    if (method >= MacroRegistersStart) {
+        return true;
+    }
+    switch (method) {
+    case MAXWELL3D_REG_INDEX(draw.end):
+    case MAXWELL3D_REG_INDEX(draw.begin):
+    case MAXWELL3D_REG_INDEX(vertex_buffer.first):
+    case MAXWELL3D_REG_INDEX(vertex_buffer.count):
+    case MAXWELL3D_REG_INDEX(index_buffer.first):
+    case MAXWELL3D_REG_INDEX(index_buffer.count):
+    case MAXWELL3D_REG_INDEX(draw_inline_index):
+    case MAXWELL3D_REG_INDEX(index_buffer32_subsequent):
+    case MAXWELL3D_REG_INDEX(index_buffer16_subsequent):
+    case MAXWELL3D_REG_INDEX(index_buffer8_subsequent):
+    case MAXWELL3D_REG_INDEX(index_buffer32_first):
+    case MAXWELL3D_REG_INDEX(index_buffer16_first):
+    case MAXWELL3D_REG_INDEX(index_buffer8_first):
+    case MAXWELL3D_REG_INDEX(inline_index_2x16.even):
+    case MAXWELL3D_REG_INDEX(inline_index_4x8.index0):
+    case MAXWELL3D_REG_INDEX(vertex_array_instance_first):
+    case MAXWELL3D_REG_INDEX(vertex_array_instance_subsequent):
+    case MAXWELL3D_REG_INDEX(wait_for_idle):
+    case MAXWELL3D_REG_INDEX(shadow_ram_control):
+    case MAXWELL3D_REG_INDEX(load_mme.instruction_ptr):
+    case MAXWELL3D_REG_INDEX(load_mme.instruction):
+    case MAXWELL3D_REG_INDEX(load_mme.start_address):
+    case MAXWELL3D_REG_INDEX(falcon[4]):
+    case MAXWELL3D_REG_INDEX(const_buffer.buffer):
+    case MAXWELL3D_REG_INDEX(const_buffer.buffer) + 1:
+    case MAXWELL3D_REG_INDEX(const_buffer.buffer) + 2:
+    case MAXWELL3D_REG_INDEX(const_buffer.buffer) + 3:
+    case MAXWELL3D_REG_INDEX(const_buffer.buffer) + 4:
+    case MAXWELL3D_REG_INDEX(const_buffer.buffer) + 5:
+    case MAXWELL3D_REG_INDEX(const_buffer.buffer) + 6:
+    case MAXWELL3D_REG_INDEX(const_buffer.buffer) + 7:
+    case MAXWELL3D_REG_INDEX(const_buffer.buffer) + 8:
+    case MAXWELL3D_REG_INDEX(const_buffer.buffer) + 9:
+    case MAXWELL3D_REG_INDEX(const_buffer.buffer) + 10:
+    case MAXWELL3D_REG_INDEX(const_buffer.buffer) + 11:
+    case MAXWELL3D_REG_INDEX(const_buffer.buffer) + 12:
+    case MAXWELL3D_REG_INDEX(const_buffer.buffer) + 13:
+    case MAXWELL3D_REG_INDEX(const_buffer.buffer) + 14:
+    case MAXWELL3D_REG_INDEX(const_buffer.buffer) + 15:
+    case MAXWELL3D_REG_INDEX(bind_groups[0].raw_config):
+    case MAXWELL3D_REG_INDEX(bind_groups[1].raw_config):
+    case MAXWELL3D_REG_INDEX(bind_groups[2].raw_config):
+    case MAXWELL3D_REG_INDEX(bind_groups[3].raw_config):
+    case MAXWELL3D_REG_INDEX(bind_groups[4].raw_config):
+    case MAXWELL3D_REG_INDEX(topology_override):
+    case MAXWELL3D_REG_INDEX(clear_surface):
+    case MAXWELL3D_REG_INDEX(report_semaphore.query):
+    case MAXWELL3D_REG_INDEX(render_enable.mode):
+    case MAXWELL3D_REG_INDEX(clear_report_value):
+    case MAXWELL3D_REG_INDEX(sync_info):
+    case MAXWELL3D_REG_INDEX(launch_dma):
+    case MAXWELL3D_REG_INDEX(inline_data):
+    case MAXWELL3D_REG_INDEX(fragment_barrier):
+    case MAXWELL3D_REG_INDEX(tiled_cache_barrier):
+        return true;
+    default:
+        return false;
+    }
+}
+
 void Maxwell3D::ProcessMacro(u32 method, const u32* base_start, u32 amount, bool is_last_call) {
     if (executing_macro == 0) {
         // A macro call must begin by writing the macro method's register, not its argument.
@@ -130,14 +201,72 @@ void Maxwell3D::ProcessMacro(u32 method, const u32* base_start, u32 amount, bool
     }
 
     macro_params.insert(macro_params.end(), base_start, base_start + amount);
+    for (size_t i = 0; i < amount; i++) {
+        macro_addresses.push_back(current_dma_segment + i * sizeof(u32));
+    }
+    macro_segments.emplace_back(current_dma_segment, amount);
+    current_macro_dirty |= current_dirty;
+    current_dirty = false;
 
     // Call the macro when there are no more parameters in the command buffer
     if (is_last_call) {
+        ConsumeSink();
         CallMacroMethod(executing_macro, macro_params);
         macro_params.clear();
+        macro_addresses.clear();
+        macro_segments.clear();
+        current_macro_dirty = false;
     }
 }
 
+void Maxwell3D::RefreshParametersImpl() {
+    size_t current_index = 0;
+    for (auto& segment : macro_segments) {
+        if (segment.first == 0) {
+            current_index += segment.second;
+            continue;
+        }
+        memory_manager.ReadBlock(segment.first, &macro_params[current_index],
+                                 sizeof(u32) * segment.second);
+        current_index += segment.second;
+    }
+}
+
+u32 Maxwell3D::GetMaxCurrentVertices() {
+    u32 num_vertices = 0;
+    for (size_t index = 0; index < Regs::NumVertexArrays; ++index) {
+        const auto& array = regs.vertex_streams[index];
+        if (array.enable == 0) {
+            continue;
+        }
+        const auto& attribute = regs.vertex_attrib_format[index];
+        if (attribute.constant) {
+            num_vertices = std::max(num_vertices, 1U);
+            continue;
+        }
+        const auto& limit = regs.vertex_stream_limits[index];
+        const GPUVAddr gpu_addr_begin = array.Address();
+        const GPUVAddr gpu_addr_end = limit.Address() + 1;
+        const u32 address_size = static_cast<u32>(gpu_addr_end - gpu_addr_begin);
+        num_vertices = std::max(
+            num_vertices, address_size / std::max(attribute.SizeInBytes(), array.stride.Value()));
+    }
+    return num_vertices;
+}
+
+size_t Maxwell3D::EstimateIndexBufferSize() {
+    GPUVAddr start_address = regs.index_buffer.StartAddress();
+    GPUVAddr end_address = regs.index_buffer.EndAddress();
+    constexpr std::array<size_t, 4> max_sizes = {
+        std::numeric_limits<u8>::max(), std::numeric_limits<u16>::max(),
+        std::numeric_limits<u32>::max(), std::numeric_limits<u32>::max()};
+    const size_t byte_size = regs.index_buffer.FormatSizeInBytes();
+    return std::min<size_t>(
+        memory_manager.GetMemoryLayoutSize(start_address, byte_size * max_sizes[byte_size]) /
+            byte_size,
+        static_cast<size_t>(end_address - start_address));
+}
+
 u32 Maxwell3D::ProcessShadowRam(u32 method, u32 argument) {
     // Keep track of the register value in shadow_state when requested.
     const auto control = shadow_state.shadow_ram_control;
@@ -152,6 +281,29 @@ u32 Maxwell3D::ProcessShadowRam(u32 method, u32 argument) {
     return argument;
 }
 
+void Maxwell3D::ConsumeSinkImpl() {
+    SCOPE_EXIT({ method_sink.clear(); });
+    const auto control = shadow_state.shadow_ram_control;
+    if (control == Regs::ShadowRamControl::Track ||
+        control == Regs::ShadowRamControl::TrackWithFilter) {
+
+        for (auto [method, value] : method_sink) {
+            shadow_state.reg_array[method] = value;
+            ProcessDirtyRegisters(method, value);
+        }
+        return;
+    }
+    if (control == Regs::ShadowRamControl::Replay) {
+        for (auto [method, value] : method_sink) {
+            ProcessDirtyRegisters(method, shadow_state.reg_array[method]);
+        }
+        return;
+    }
+    for (auto [method, value] : method_sink) {
+        ProcessDirtyRegisters(method, value);
+    }
+}
+
 void Maxwell3D::ProcessDirtyRegisters(u32 method, u32 argument) {
     if (regs.reg_array[method] == argument) {
         return;
@@ -263,7 +415,6 @@ void Maxwell3D::CallMethod(u32 method, u32 method_argument, bool is_last_call) {
 
     const u32 argument = ProcessShadowRam(method, method_argument);
     ProcessDirtyRegisters(method, argument);
-
     ProcessMethodCall(method, argument, method_argument, is_last_call);
 }
 
@@ -294,9 +445,11 @@ void Maxwell3D::CallMultiMethod(u32 method, const u32* base_start, u32 amount,
     case MAXWELL3D_REG_INDEX(const_buffer.buffer) + 15:
         ProcessCBMultiData(base_start, amount);
         break;
-    case MAXWELL3D_REG_INDEX(inline_data):
+    case MAXWELL3D_REG_INDEX(inline_data): {
+        ASSERT(methods_pending == amount);
         upload_state.ProcessData(base_start, amount);
         return;
+    }
     default:
         for (u32 i = 0; i < amount; i++) {
             CallMethod(method, base_start[i], methods_pending - i <= 1);
@@ -389,7 +542,11 @@ void Maxwell3D::ProcessQueryCondition() {
     case Regs::RenderEnable::Override::NeverRender:
         execute_on = false;
         break;
-    case Regs::RenderEnable::Override::UseRenderEnable:
+    case Regs::RenderEnable::Override::UseRenderEnable: {
+        if (rasterizer->AccelerateConditionalRendering()) {
+            execute_on = true;
+            return;
+        }
         switch (regs.render_enable.mode) {
         case Regs::RenderEnable::Mode::True: {
             execute_on = true;
@@ -427,6 +584,7 @@ void Maxwell3D::ProcessQueryCondition() {
         }
         break;
     }
+    }
 }
 
 void Maxwell3D::ProcessCounterReset() {
@@ -463,7 +621,8 @@ std::optional<u64> Maxwell3D::GetQueryResult() {
 }
 
 void Maxwell3D::ProcessCBBind(size_t stage_index) {
-    // Bind the buffer currently in CB_ADDRESS to the specified index in the desired shader stage.
+    // Bind the buffer currently in CB_ADDRESS to the specified index in the desired shader
+    // stage.
     const auto& bind_data = regs.bind_groups[stage_index];
     auto& buffer = state.shader_stages[stage_index].const_buffers[bind_data.shader_slot];
     buffer.enabled = bind_data.valid.Value() != 0;
@@ -524,4 +683,10 @@ u32 Maxwell3D::GetRegisterValue(u32 method) const {
     return regs.reg_array[method];
 }
 
+void Maxwell3D::SetHLEReplacementAttributeType(u32 bank, u32 offset,
+                                               HLEReplacementAttributeType name) {
+    const u64 key = (static_cast<u64>(bank) << 32) | offset;
+    replace_table.emplace(key, name);
+}
+
 } // namespace Tegra::Engines
diff --git a/src/video_core/engines/maxwell_3d.h b/src/video_core/engines/maxwell_3d.h
index 22b904319..0b2fd2928 100644
--- a/src/video_core/engines/maxwell_3d.h
+++ b/src/video_core/engines/maxwell_3d.h
@@ -272,6 +272,7 @@ public:
             };
 
             union {
+                u32 raw;
                 BitField<0, 1, Mode> mode;
                 BitField<4, 8, u32> pad;
             };
@@ -1217,10 +1218,12 @@ public:
 
         struct Window {
             union {
+                u32 raw_x;
                 BitField<0, 16, u32> x_min;
                 BitField<16, 16, u32> x_max;
             };
             union {
+                u32 raw_y;
                 BitField<0, 16, u32> y_min;
                 BitField<16, 16, u32> y_max;
             };
@@ -2708,7 +2711,7 @@ public:
                 u32 post_z_pixel_imask;                                                ///< 0x0F1C
                 INSERT_PADDING_BYTES_NOINIT(0x20);
                 ConstantColorRendering const_color_rendering;                          ///< 0x0F40
-                s32 stencil_back_ref;                                                  ///< 0x0F54
+                u32 stencil_back_ref;                                                  ///< 0x0F54
                 u32 stencil_back_mask;                                                 ///< 0x0F58
                 u32 stencil_back_func_mask;                                            ///< 0x0F5C
                 INSERT_PADDING_BYTES_NOINIT(0x14);
@@ -2832,9 +2835,9 @@ public:
                 Blend blend;                                                           ///< 0x133C
                 u32 stencil_enable;                                                    ///< 0x1380
                 StencilOp stencil_front_op;                                            ///< 0x1384
-                s32 stencil_front_ref;                                                 ///< 0x1394
-                s32 stencil_front_func_mask;                                           ///< 0x1398
-                s32 stencil_front_mask;                                                ///< 0x139C
+                u32 stencil_front_ref;                                                 ///< 0x1394
+                u32 stencil_front_func_mask;                                           ///< 0x1398
+                u32 stencil_front_mask;                                                ///< 0x139C
                 INSERT_PADDING_BYTES_NOINIT(0x4);
                 u32 draw_auto_start_byte_count;                                        ///< 0x13A4
                 PsSaturate frag_color_clamp;                                           ///< 0x13A8
@@ -3020,6 +3023,24 @@ public:
     /// Store temporary hw register values, used by some calls to restore state after a operation
     Regs shadow_state;
 
+    // None Engine
+    enum class EngineHint : u32 {
+        None = 0x0,
+        OnHLEMacro = 0x1,
+    };
+
+    EngineHint engine_state{EngineHint::None};
+
+    enum class HLEReplacementAttributeType : u32 {
+        BaseVertex = 0x0,
+        BaseInstance = 0x1,
+        DrawID = 0x2,
+    };
+
+    void SetHLEReplacementAttributeType(u32 bank, u32 offset, HLEReplacementAttributeType name);
+
+    std::unordered_map<u64, HLEReplacementAttributeType> replace_table;
+
     static_assert(sizeof(Regs) == Regs::NUM_REGS * sizeof(u32), "Maxwell3D Regs has wrong size");
     static_assert(std::is_trivially_copyable_v<Regs>, "Maxwell3D Regs must be trivially copyable");
 
@@ -3067,6 +3088,35 @@ public:
     std::unique_ptr<DrawManager> draw_manager;
     friend class DrawManager;
 
+    GPUVAddr GetMacroAddress(size_t index) const {
+        return macro_addresses[index];
+    }
+
+    void RefreshParameters() {
+        if (!current_macro_dirty) {
+            return;
+        }
+        RefreshParametersImpl();
+    }
+
+    bool AnyParametersDirty() const {
+        return current_macro_dirty;
+    }
+
+    u32 GetMaxCurrentVertices();
+
+    size_t EstimateIndexBufferSize();
+
+    /// Handles a write to the CLEAR_BUFFERS register.
+    void ProcessClearBuffers(u32 layer_count);
+
+    /// Handles a write to the CB_BIND register.
+    void ProcessCBBind(size_t stage_index);
+
+    /// Handles a write to the CB_DATA[i] register.
+    void ProcessCBData(u32 value);
+    void ProcessCBMultiData(const u32* start_base, u32 amount);
+
 private:
     void InitializeRegisterDefaults();
 
@@ -3076,6 +3126,8 @@ private:
 
     void ProcessDirtyRegisters(u32 method, u32 argument);
 
+    void ConsumeSinkImpl() override;
+
     void ProcessMethodCall(u32 method, u32 argument, u32 nonshadow_argument, bool is_last_call);
 
     /// Retrieves information about a specific TIC entry from the TIC buffer.
@@ -3116,16 +3168,13 @@ private:
     /// Handles writes to syncing register.
     void ProcessSyncPoint();
 
-    /// Handles a write to the CB_DATA[i] register.
-    void ProcessCBData(u32 value);
-    void ProcessCBMultiData(const u32* start_base, u32 amount);
-
-    /// Handles a write to the CB_BIND register.
-    void ProcessCBBind(size_t stage_index);
-
     /// Returns a query's value or an empty object if the value will be deferred through a cache.
     std::optional<u64> GetQueryResult();
 
+    void RefreshParametersImpl();
+
+    bool IsMethodExecutable(u32 method);
+
     Core::System& system;
     MemoryManager& memory_manager;
 
@@ -3145,6 +3194,10 @@ private:
     Upload::State upload_state;
 
     bool execute_on{true};
+
+    std::vector<std::pair<GPUVAddr, size_t>> macro_segments;
+    std::vector<GPUVAddr> macro_addresses;
+    bool current_macro_dirty{};
 };
 
 #define ASSERT_REG_POSITION(field_name, position)                                                  \
diff --git a/src/video_core/engines/maxwell_dma.cpp b/src/video_core/engines/maxwell_dma.cpp
index f73d7bf0f..01f70ea9e 100644
--- a/src/video_core/engines/maxwell_dma.cpp
+++ b/src/video_core/engines/maxwell_dma.cpp
@@ -21,7 +21,10 @@ namespace Tegra::Engines {
 using namespace Texture;
 
 MaxwellDMA::MaxwellDMA(Core::System& system_, MemoryManager& memory_manager_)
-    : system{system_}, memory_manager{memory_manager_} {}
+    : system{system_}, memory_manager{memory_manager_} {
+    execution_mask.reset();
+    execution_mask[offsetof(Regs, launch_dma) / sizeof(u32)] = true;
+}
 
 MaxwellDMA::~MaxwellDMA() = default;
 
@@ -29,6 +32,13 @@ void MaxwellDMA::BindRasterizer(VideoCore::RasterizerInterface* rasterizer_) {
     rasterizer = rasterizer_;
 }
 
+void MaxwellDMA::ConsumeSinkImpl() {
+    for (auto [method, value] : method_sink) {
+        regs.reg_array[method] = value;
+    }
+    method_sink.clear();
+}
+
 void MaxwellDMA::CallMethod(u32 method, u32 method_argument, bool is_last_call) {
     ASSERT_MSG(method < NUM_REGS, "Invalid MaxwellDMA register");
 
diff --git a/src/video_core/engines/maxwell_dma.h b/src/video_core/engines/maxwell_dma.h
index c88191a61..0e594fa74 100644
--- a/src/video_core/engines/maxwell_dma.h
+++ b/src/video_core/engines/maxwell_dma.h
@@ -231,6 +231,8 @@ private:
 
     void ReleaseSemaphore();
 
+    void ConsumeSinkImpl() override;
+
     Core::System& system;
 
     MemoryManager& memory_manager;
diff --git a/src/video_core/macro/macro.cpp b/src/video_core/macro/macro.cpp
index 505d81c1e..82ad0477d 100644
--- a/src/video_core/macro/macro.cpp
+++ b/src/video_core/macro/macro.cpp
@@ -12,7 +12,9 @@
 #include "common/assert.h"
 #include "common/fs/fs.h"
 #include "common/fs/path_util.h"
+#include "common/microprofile.h"
 #include "common/settings.h"
+#include "video_core/engines/maxwell_3d.h"
 #include "video_core/macro/macro.h"
 #include "video_core/macro/macro_hle.h"
 #include "video_core/macro/macro_interpreter.h"
@@ -21,6 +23,8 @@
 #include "video_core/macro/macro_jit_x64.h"
 #endif
 
+MICROPROFILE_DEFINE(MacroHLE, "GPU", "Execute macro HLE", MP_RGB(128, 192, 192));
+
 namespace Tegra {
 
 static void Dump(u64 hash, std::span<const u32> code) {
@@ -40,8 +44,8 @@ static void Dump(u64 hash, std::span<const u32> code) {
     macro_file.write(reinterpret_cast<const char*>(code.data()), code.size_bytes());
 }
 
-MacroEngine::MacroEngine(Engines::Maxwell3D& maxwell3d)
-    : hle_macros{std::make_unique<Tegra::HLEMacro>(maxwell3d)} {}
+MacroEngine::MacroEngine(Engines::Maxwell3D& maxwell3d_)
+    : hle_macros{std::make_unique<Tegra::HLEMacro>(maxwell3d_)}, maxwell3d{maxwell3d_} {}
 
 MacroEngine::~MacroEngine() = default;
 
@@ -59,8 +63,10 @@ void MacroEngine::Execute(u32 method, const std::vector<u32>& parameters) {
     if (compiled_macro != macro_cache.end()) {
         const auto& cache_info = compiled_macro->second;
         if (cache_info.has_hle_program) {
+            MICROPROFILE_SCOPE(MacroHLE);
             cache_info.hle_program->Execute(parameters, method);
         } else {
+            maxwell3d.RefreshParameters();
             cache_info.lle_program->Execute(parameters, method);
         }
     } else {
@@ -101,12 +107,15 @@ void MacroEngine::Execute(u32 method, const std::vector<u32>& parameters) {
             }
         }
 
-        if (auto hle_program = hle_macros->GetHLEProgram(cache_info.hash)) {
+        auto hle_program = hle_macros->GetHLEProgram(cache_info.hash);
+        if (!hle_program || Settings::values.disable_macro_hle) {
+            maxwell3d.RefreshParameters();
+            cache_info.lle_program->Execute(parameters, method);
+        } else {
             cache_info.has_hle_program = true;
             cache_info.hle_program = std::move(hle_program);
+            MICROPROFILE_SCOPE(MacroHLE);
             cache_info.hle_program->Execute(parameters, method);
-        } else {
-            cache_info.lle_program->Execute(parameters, method);
         }
     }
 }
diff --git a/src/video_core/macro/macro.h b/src/video_core/macro/macro.h
index 07d97ba39..737ced9a4 100644
--- a/src/video_core/macro/macro.h
+++ b/src/video_core/macro/macro.h
@@ -137,6 +137,7 @@ private:
     std::unordered_map<u32, CacheInfo> macro_cache;
     std::unordered_map<u32, std::vector<u32>> uploaded_macro_code;
     std::unique_ptr<HLEMacro> hle_macros;
+    Engines::Maxwell3D& maxwell3d;
 };
 
 std::unique_ptr<MacroEngine> GetMacroEngine(Engines::Maxwell3D& maxwell3d);
diff --git a/src/video_core/macro/macro_hle.cpp b/src/video_core/macro/macro_hle.cpp
index 8549db2e4..a5476e795 100644
--- a/src/video_core/macro/macro_hle.cpp
+++ b/src/video_core/macro/macro_hle.cpp
@@ -1,143 +1,593 @@
-// SPDX-FileCopyrightText: Copyright 2020 yuzu Emulator Project
-// SPDX-License-Identifier: GPL-2.0-or-later
+// SPDX-FileCopyrightText: Copyright 2023 yuzu Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
 
 #include <array>
 #include <vector>
+#include "common/assert.h"
 #include "common/scope_exit.h"
 #include "video_core/dirty_flags.h"
 #include "video_core/engines/draw_manager.h"
 #include "video_core/engines/maxwell_3d.h"
 #include "video_core/macro/macro.h"
 #include "video_core/macro/macro_hle.h"
+#include "video_core/memory_manager.h"
 #include "video_core/rasterizer_interface.h"
 
 namespace Tegra {
-namespace {
 
-using HLEFunction = void (*)(Engines::Maxwell3D& maxwell3d, const std::vector<u32>& parameters);
+using Maxwell3D = Engines::Maxwell3D;
 
-// HLE'd functions
-void HLE_771BB18C62444DA0(Engines::Maxwell3D& maxwell3d, const std::vector<u32>& parameters) {
-    const u32 instance_count = parameters[2] & maxwell3d.GetRegisterValue(0xD1B);
-    maxwell3d.draw_manager->DrawIndex(
-        static_cast<Tegra::Engines::Maxwell3D::Regs::PrimitiveTopology>(parameters[0] & 0x3ffffff),
-        parameters[4], parameters[1], parameters[3], parameters[5], instance_count);
-}
+namespace {
 
-void HLE_0D61FC9FAAC9FCAD(Engines::Maxwell3D& maxwell3d, const std::vector<u32>& parameters) {
-    const u32 instance_count = (maxwell3d.GetRegisterValue(0xD1B) & parameters[2]);
-    maxwell3d.draw_manager->DrawArray(
-        static_cast<Tegra::Engines::Maxwell3D::Regs::PrimitiveTopology>(parameters[0]),
-        parameters[3], parameters[1], parameters[4], instance_count);
+bool IsTopologySafe(Maxwell3D::Regs::PrimitiveTopology topology) {
+    switch (topology) {
+    case Maxwell3D::Regs::PrimitiveTopology::Points:
+    case Maxwell3D::Regs::PrimitiveTopology::Lines:
+    case Maxwell3D::Regs::PrimitiveTopology::LineLoop:
+    case Maxwell3D::Regs::PrimitiveTopology::LineStrip:
+    case Maxwell3D::Regs::PrimitiveTopology::Triangles:
+    case Maxwell3D::Regs::PrimitiveTopology::TriangleStrip:
+    case Maxwell3D::Regs::PrimitiveTopology::TriangleFan:
+    case Maxwell3D::Regs::PrimitiveTopology::LinesAdjacency:
+    case Maxwell3D::Regs::PrimitiveTopology::LineStripAdjacency:
+    case Maxwell3D::Regs::PrimitiveTopology::TrianglesAdjacency:
+    case Maxwell3D::Regs::PrimitiveTopology::TriangleStripAdjacency:
+    case Maxwell3D::Regs::PrimitiveTopology::Patches:
+        return true;
+    case Maxwell3D::Regs::PrimitiveTopology::Quads:
+    case Maxwell3D::Regs::PrimitiveTopology::QuadStrip:
+    case Maxwell3D::Regs::PrimitiveTopology::Polygon:
+    default:
+        return false;
+    }
 }
 
-void HLE_0217920100488FF7(Engines::Maxwell3D& maxwell3d, const std::vector<u32>& parameters) {
-    const u32 instance_count = (maxwell3d.GetRegisterValue(0xD1B) & parameters[2]);
-    const u32 element_base = parameters[4];
-    const u32 base_instance = parameters[5];
-    maxwell3d.regs.vertex_id_base = element_base;
-    maxwell3d.dirty.flags[VideoCommon::Dirty::IndexBuffer] = true;
-    maxwell3d.CallMethod(0x8e3, 0x640, true);
-    maxwell3d.CallMethod(0x8e4, element_base, true);
-    maxwell3d.CallMethod(0x8e5, base_instance, true);
-
-    maxwell3d.draw_manager->DrawIndex(
-        static_cast<Tegra::Engines::Maxwell3D::Regs::PrimitiveTopology>(parameters[0]),
-        parameters[3], parameters[1], element_base, base_instance, instance_count);
-
-    maxwell3d.regs.vertex_id_base = 0x0;
-    maxwell3d.CallMethod(0x8e3, 0x640, true);
-    maxwell3d.CallMethod(0x8e4, 0x0, true);
-    maxwell3d.CallMethod(0x8e5, 0x0, true);
-}
+class HLEMacroImpl : public CachedMacro {
+public:
+    explicit HLEMacroImpl(Maxwell3D& maxwell3d_) : maxwell3d{maxwell3d_} {}
 
-// Multidraw Indirect
-void HLE_3F5E74B9C9A50164(Engines::Maxwell3D& maxwell3d, const std::vector<u32>& parameters) {
-    SCOPE_EXIT({
-        // Clean everything.
-        maxwell3d.regs.vertex_id_base = 0x0;
-        maxwell3d.CallMethod(0x8e3, 0x640, true);
-        maxwell3d.CallMethod(0x8e4, 0x0, true);
-        maxwell3d.CallMethod(0x8e5, 0x0, true);
+protected:
+    Maxwell3D& maxwell3d;
+};
+
+class HLE_DrawArrays final : public HLEMacroImpl {
+public:
+    explicit HLE_DrawArrays(Maxwell3D& maxwell3d_) : HLEMacroImpl(maxwell3d_) {}
+
+    void Execute(const std::vector<u32>& parameters, [[maybe_unused]] u32 method) override {
+        maxwell3d.RefreshParameters();
+
+        auto topology = static_cast<Maxwell3D::Regs::PrimitiveTopology>(parameters[0]);
+        maxwell3d.draw_manager->DrawArray(topology, parameters[1], parameters[2],
+                                          maxwell3d.regs.global_base_instance_index, 1);
+    }
+};
+
+class HLE_DrawIndexed final : public HLEMacroImpl {
+public:
+    explicit HLE_DrawIndexed(Maxwell3D& maxwell3d_) : HLEMacroImpl(maxwell3d_) {}
+
+    void Execute(const std::vector<u32>& parameters, [[maybe_unused]] u32 method) override {
+        maxwell3d.RefreshParameters();
+        maxwell3d.regs.index_buffer.start_addr_high = parameters[1];
+        maxwell3d.regs.index_buffer.start_addr_low = parameters[2];
+        maxwell3d.regs.index_buffer.format =
+            static_cast<Engines::Maxwell3D::Regs::IndexFormat>(parameters[3]);
         maxwell3d.dirty.flags[VideoCommon::Dirty::IndexBuffer] = true;
-    });
-    const u32 start_indirect = parameters[0];
-    const u32 end_indirect = parameters[1];
-    if (start_indirect >= end_indirect) {
-        // Nothing to do.
-        return;
-    }
-    const u32 padding = parameters[3];
-    const std::size_t max_draws = parameters[4];
-
-    const u32 indirect_words = 5 + padding;
-    const std::size_t first_draw = start_indirect;
-    const std::size_t effective_draws = end_indirect - start_indirect;
-    const std::size_t last_draw = start_indirect + std::min(effective_draws, max_draws);
-
-    for (std::size_t index = first_draw; index < last_draw; index++) {
-        const std::size_t base = index * indirect_words + 5;
-        const u32 base_vertex = parameters[base + 3];
-        const u32 base_instance = parameters[base + 4];
-        maxwell3d.regs.vertex_id_base = base_vertex;
-        maxwell3d.CallMethod(0x8e3, 0x640, true);
-        maxwell3d.CallMethod(0x8e4, base_vertex, true);
-        maxwell3d.CallMethod(0x8e5, base_instance, true);
+
+        auto topology = static_cast<Maxwell3D::Regs::PrimitiveTopology>(parameters[0]);
+        maxwell3d.draw_manager->DrawIndex(topology, 0, parameters[4],
+                                          maxwell3d.regs.global_base_vertex_index,
+                                          maxwell3d.regs.global_base_instance_index, 1);
+    }
+};
+
+/*
+ * @note: these macros have two versions, a normal and extended version, with the extended version
+ * also assigning the base vertex/instance.
+ */
+template <bool extended>
+class HLE_DrawArraysIndirect final : public HLEMacroImpl {
+public:
+    explicit HLE_DrawArraysIndirect(Maxwell3D& maxwell3d_) : HLEMacroImpl(maxwell3d_) {}
+
+    void Execute(const std::vector<u32>& parameters, [[maybe_unused]] u32 method) override {
+        auto topology = static_cast<Maxwell3D::Regs::PrimitiveTopology>(parameters[0]);
+        if (!maxwell3d.AnyParametersDirty() || !IsTopologySafe(topology)) {
+            Fallback(parameters);
+            return;
+        }
+
+        auto& params = maxwell3d.draw_manager->GetIndirectParams();
+        params.is_indexed = false;
+        params.include_count = false;
+        params.count_start_address = 0;
+        params.indirect_start_address = maxwell3d.GetMacroAddress(1);
+        params.buffer_size = 4 * sizeof(u32);
+        params.max_draw_counts = 1;
+        params.stride = 0;
+
+        if constexpr (extended) {
+            maxwell3d.engine_state = Maxwell3D::EngineHint::OnHLEMacro;
+            maxwell3d.SetHLEReplacementAttributeType(
+                0, 0x640, Maxwell3D::HLEReplacementAttributeType::BaseInstance);
+        }
+
+        maxwell3d.draw_manager->DrawArrayIndirect(topology);
+
+        if constexpr (extended) {
+            maxwell3d.engine_state = Maxwell3D::EngineHint::None;
+            maxwell3d.replace_table.clear();
+        }
+    }
+
+private:
+    void Fallback(const std::vector<u32>& parameters) {
+        SCOPE_EXIT({
+            if (extended) {
+                maxwell3d.engine_state = Maxwell3D::EngineHint::None;
+                maxwell3d.replace_table.clear();
+            }
+        });
+        maxwell3d.RefreshParameters();
+        const u32 instance_count = (maxwell3d.GetRegisterValue(0xD1B) & parameters[2]);
+
+        auto topology = static_cast<Maxwell3D::Regs::PrimitiveTopology>(parameters[0]);
+        const u32 vertex_first = parameters[3];
+        const u32 vertex_count = parameters[1];
+
+        if (!IsTopologySafe(topology) &&
+            static_cast<size_t>(maxwell3d.GetMaxCurrentVertices()) <
+                static_cast<size_t>(vertex_first) + static_cast<size_t>(vertex_count)) {
+            ASSERT_MSG(false, "Faulty draw!");
+            return;
+        }
+
+        const u32 base_instance = parameters[4];
+        if constexpr (extended) {
+            maxwell3d.regs.global_base_instance_index = base_instance;
+            maxwell3d.engine_state = Maxwell3D::EngineHint::OnHLEMacro;
+            maxwell3d.SetHLEReplacementAttributeType(
+                0, 0x640, Maxwell3D::HLEReplacementAttributeType::BaseInstance);
+        }
+
+        maxwell3d.draw_manager->DrawArray(topology, vertex_first, vertex_count, base_instance,
+                                          instance_count);
+
+        if constexpr (extended) {
+            maxwell3d.regs.global_base_instance_index = 0;
+            maxwell3d.engine_state = Maxwell3D::EngineHint::None;
+            maxwell3d.replace_table.clear();
+        }
+    }
+};
+
+/*
+ * @note: these macros have two versions, a normal and extended version, with the extended version
+ * also assigning the base vertex/instance.
+ */
+template <bool extended>
+class HLE_DrawIndexedIndirect final : public HLEMacroImpl {
+public:
+    explicit HLE_DrawIndexedIndirect(Maxwell3D& maxwell3d_) : HLEMacroImpl(maxwell3d_) {}
+
+    void Execute(const std::vector<u32>& parameters, [[maybe_unused]] u32 method) override {
+        auto topology = static_cast<Maxwell3D::Regs::PrimitiveTopology>(parameters[0]);
+        if (!maxwell3d.AnyParametersDirty() || !IsTopologySafe(topology)) {
+            Fallback(parameters);
+            return;
+        }
+
+        const u32 estimate = static_cast<u32>(maxwell3d.EstimateIndexBufferSize());
+        const u32 element_base = parameters[4];
+        const u32 base_instance = parameters[5];
+        maxwell3d.regs.vertex_id_base = element_base;
+        maxwell3d.regs.global_base_vertex_index = element_base;
+        maxwell3d.regs.global_base_instance_index = base_instance;
+        maxwell3d.dirty.flags[VideoCommon::Dirty::IndexBuffer] = true;
+        if constexpr (extended) {
+            maxwell3d.engine_state = Maxwell3D::EngineHint::OnHLEMacro;
+            maxwell3d.SetHLEReplacementAttributeType(
+                0, 0x640, Maxwell3D::HLEReplacementAttributeType::BaseVertex);
+            maxwell3d.SetHLEReplacementAttributeType(
+                0, 0x644, Maxwell3D::HLEReplacementAttributeType::BaseInstance);
+        }
+        auto& params = maxwell3d.draw_manager->GetIndirectParams();
+        params.is_indexed = true;
+        params.include_count = false;
+        params.count_start_address = 0;
+        params.indirect_start_address = maxwell3d.GetMacroAddress(1);
+        params.buffer_size = 5 * sizeof(u32);
+        params.max_draw_counts = 1;
+        params.stride = 0;
         maxwell3d.dirty.flags[VideoCommon::Dirty::IndexBuffer] = true;
+        maxwell3d.draw_manager->DrawIndexedIndirect(topology, 0, estimate);
+        maxwell3d.regs.vertex_id_base = 0x0;
+        maxwell3d.regs.global_base_vertex_index = 0x0;
+        maxwell3d.regs.global_base_instance_index = 0x0;
+        if constexpr (extended) {
+            maxwell3d.engine_state = Maxwell3D::EngineHint::None;
+            maxwell3d.replace_table.clear();
+        }
+    }
+
+private:
+    void Fallback(const std::vector<u32>& parameters) {
+        maxwell3d.RefreshParameters();
+        const u32 instance_count = (maxwell3d.GetRegisterValue(0xD1B) & parameters[2]);
+        const u32 element_base = parameters[4];
+        const u32 base_instance = parameters[5];
+        maxwell3d.regs.vertex_id_base = element_base;
+        maxwell3d.regs.global_base_vertex_index = element_base;
+        maxwell3d.regs.global_base_instance_index = base_instance;
+        maxwell3d.dirty.flags[VideoCommon::Dirty::IndexBuffer] = true;
+        if constexpr (extended) {
+            maxwell3d.engine_state = Maxwell3D::EngineHint::OnHLEMacro;
+            maxwell3d.SetHLEReplacementAttributeType(
+                0, 0x640, Maxwell3D::HLEReplacementAttributeType::BaseVertex);
+            maxwell3d.SetHLEReplacementAttributeType(
+                0, 0x644, Maxwell3D::HLEReplacementAttributeType::BaseInstance);
+        }
+
         maxwell3d.draw_manager->DrawIndex(
-            static_cast<Tegra::Engines::Maxwell3D::Regs::PrimitiveTopology>(parameters[2]),
-            parameters[base + 2], parameters[base], base_vertex, base_instance,
-            parameters[base + 1]);
+            static_cast<Tegra::Maxwell3D::Regs::PrimitiveTopology>(parameters[0]), parameters[3],
+            parameters[1], element_base, base_instance, instance_count);
+
+        maxwell3d.regs.vertex_id_base = 0x0;
+        maxwell3d.regs.global_base_vertex_index = 0x0;
+        maxwell3d.regs.global_base_instance_index = 0x0;
+        if constexpr (extended) {
+            maxwell3d.engine_state = Maxwell3D::EngineHint::None;
+            maxwell3d.replace_table.clear();
+        }
     }
-}
+};
 
-// Multi-layer Clear
-void HLE_EAD26C3E2109B06B(Engines::Maxwell3D& maxwell3d, const std::vector<u32>& parameters) {
-    ASSERT(parameters.size() == 1);
+class HLE_MultiLayerClear final : public HLEMacroImpl {
+public:
+    explicit HLE_MultiLayerClear(Maxwell3D& maxwell3d_) : HLEMacroImpl(maxwell3d_) {}
 
-    const Engines::Maxwell3D::Regs::ClearSurface clear_params{parameters[0]};
-    const u32 rt_index = clear_params.RT;
-    const u32 num_layers = maxwell3d.regs.rt[rt_index].depth;
-    ASSERT(clear_params.layer == 0);
+    void Execute(const std::vector<u32>& parameters, [[maybe_unused]] u32 method) override {
+        maxwell3d.RefreshParameters();
+        ASSERT(parameters.size() == 1);
 
-    maxwell3d.regs.clear_surface.raw = clear_params.raw;
-    maxwell3d.draw_manager->Clear(num_layers);
-}
+        const Maxwell3D::Regs::ClearSurface clear_params{parameters[0]};
+        const u32 rt_index = clear_params.RT;
+        const u32 num_layers = maxwell3d.regs.rt[rt_index].depth;
+        ASSERT(clear_params.layer == 0);
 
-constexpr std::array<std::pair<u64, HLEFunction>, 5> hle_funcs{{
-    {0x771BB18C62444DA0, &HLE_771BB18C62444DA0},
-    {0x0D61FC9FAAC9FCAD, &HLE_0D61FC9FAAC9FCAD},
-    {0x0217920100488FF7, &HLE_0217920100488FF7},
-    {0x3F5E74B9C9A50164, &HLE_3F5E74B9C9A50164},
-    {0xEAD26C3E2109B06B, &HLE_EAD26C3E2109B06B},
-}};
+        maxwell3d.regs.clear_surface.raw = clear_params.raw;
+        maxwell3d.draw_manager->Clear(num_layers);
+    }
+};
+
+class HLE_MultiDrawIndexedIndirectCount final : public HLEMacroImpl {
+public:
+    explicit HLE_MultiDrawIndexedIndirectCount(Maxwell3D& maxwell3d_) : HLEMacroImpl(maxwell3d_) {}
+
+    void Execute(const std::vector<u32>& parameters, [[maybe_unused]] u32 method) override {
+        const auto topology = static_cast<Maxwell3D::Regs::PrimitiveTopology>(parameters[2]);
+        if (!IsTopologySafe(topology)) {
+            Fallback(parameters);
+            return;
+        }
+
+        const u32 start_indirect = parameters[0];
+        const u32 end_indirect = parameters[1];
+        if (start_indirect >= end_indirect) {
+            // Nothing to do.
+            return;
+        }
+
+        const u32 padding = parameters[3]; // padding is in words
+
+        // size of each indirect segment
+        const u32 indirect_words = 5 + padding;
+        const u32 stride = indirect_words * sizeof(u32);
+        const std::size_t draw_count = end_indirect - start_indirect;
+        const u32 estimate = static_cast<u32>(maxwell3d.EstimateIndexBufferSize());
+        maxwell3d.dirty.flags[VideoCommon::Dirty::IndexBuffer] = true;
+        auto& params = maxwell3d.draw_manager->GetIndirectParams();
+        params.is_indexed = true;
+        params.include_count = true;
+        params.count_start_address = maxwell3d.GetMacroAddress(4);
+        params.indirect_start_address = maxwell3d.GetMacroAddress(5);
+        params.buffer_size = stride * draw_count;
+        params.max_draw_counts = draw_count;
+        params.stride = stride;
+        maxwell3d.dirty.flags[VideoCommon::Dirty::IndexBuffer] = true;
+        maxwell3d.engine_state = Maxwell3D::EngineHint::OnHLEMacro;
+        maxwell3d.SetHLEReplacementAttributeType(
+            0, 0x640, Maxwell3D::HLEReplacementAttributeType::BaseVertex);
+        maxwell3d.SetHLEReplacementAttributeType(
+            0, 0x644, Maxwell3D::HLEReplacementAttributeType::BaseInstance);
+        maxwell3d.SetHLEReplacementAttributeType(0, 0x648,
+                                                 Maxwell3D::HLEReplacementAttributeType::DrawID);
+        maxwell3d.draw_manager->DrawIndexedIndirect(topology, 0, estimate);
+        maxwell3d.engine_state = Maxwell3D::EngineHint::None;
+        maxwell3d.replace_table.clear();
+    }
+
+private:
+    void Fallback(const std::vector<u32>& parameters) {
+        SCOPE_EXIT({
+            // Clean everything.
+            maxwell3d.regs.vertex_id_base = 0x0;
+            maxwell3d.engine_state = Maxwell3D::EngineHint::None;
+            maxwell3d.replace_table.clear();
+        });
+        maxwell3d.RefreshParameters();
+        const u32 start_indirect = parameters[0];
+        const u32 end_indirect = parameters[1];
+        if (start_indirect >= end_indirect) {
+            // Nothing to do.
+            return;
+        }
+        const auto topology = static_cast<Maxwell3D::Regs::PrimitiveTopology>(parameters[2]);
+        const u32 padding = parameters[3];
+        const std::size_t max_draws = parameters[4];
+
+        const u32 indirect_words = 5 + padding;
+        const std::size_t first_draw = start_indirect;
+        const std::size_t effective_draws = end_indirect - start_indirect;
+        const std::size_t last_draw = start_indirect + std::min(effective_draws, max_draws);
+
+        for (std::size_t index = first_draw; index < last_draw; index++) {
+            const std::size_t base = index * indirect_words + 5;
+            const u32 base_vertex = parameters[base + 3];
+            const u32 base_instance = parameters[base + 4];
+            maxwell3d.regs.vertex_id_base = base_vertex;
+            maxwell3d.engine_state = Maxwell3D::EngineHint::OnHLEMacro;
+            maxwell3d.SetHLEReplacementAttributeType(
+                0, 0x640, Maxwell3D::HLEReplacementAttributeType::BaseVertex);
+            maxwell3d.SetHLEReplacementAttributeType(
+                0, 0x644, Maxwell3D::HLEReplacementAttributeType::BaseInstance);
+            maxwell3d.CallMethod(0x8e3, 0x648, true);
+            maxwell3d.CallMethod(0x8e4, static_cast<u32>(index), true);
+            maxwell3d.dirty.flags[VideoCommon::Dirty::IndexBuffer] = true;
+            maxwell3d.draw_manager->DrawIndex(topology, parameters[base + 2], parameters[base],
+                                              base_vertex, base_instance, parameters[base + 1]);
+        }
+    }
+};
+
+class HLE_C713C83D8F63CCF3 final : public HLEMacroImpl {
+public:
+    explicit HLE_C713C83D8F63CCF3(Maxwell3D& maxwell3d_) : HLEMacroImpl(maxwell3d_) {}
+
+    void Execute(const std::vector<u32>& parameters, [[maybe_unused]] u32 method) override {
+        maxwell3d.RefreshParameters();
+        const u32 offset = (parameters[0] & 0x3FFFFFFF) << 2;
+        const u32 address = maxwell3d.regs.shadow_scratch[24];
+        auto& const_buffer = maxwell3d.regs.const_buffer;
+        const_buffer.size = 0x7000;
+        const_buffer.address_high = (address >> 24) & 0xFF;
+        const_buffer.address_low = address << 8;
+        const_buffer.offset = offset;
+    }
+};
+
+class HLE_D7333D26E0A93EDE final : public HLEMacroImpl {
+public:
+    explicit HLE_D7333D26E0A93EDE(Maxwell3D& maxwell3d_) : HLEMacroImpl(maxwell3d_) {}
+
+    void Execute(const std::vector<u32>& parameters, [[maybe_unused]] u32 method) override {
+        maxwell3d.RefreshParameters();
+        const size_t index = parameters[0];
+        const u32 address = maxwell3d.regs.shadow_scratch[42 + index];
+        const u32 size = maxwell3d.regs.shadow_scratch[47 + index];
+        auto& const_buffer = maxwell3d.regs.const_buffer;
+        const_buffer.size = size;
+        const_buffer.address_high = (address >> 24) & 0xFF;
+        const_buffer.address_low = address << 8;
+    }
+};
+
+class HLE_BindShader final : public HLEMacroImpl {
+public:
+    explicit HLE_BindShader(Maxwell3D& maxwell3d_) : HLEMacroImpl(maxwell3d_) {}
+
+    void Execute(const std::vector<u32>& parameters, [[maybe_unused]] u32 method) override {
+        maxwell3d.RefreshParameters();
+        auto& regs = maxwell3d.regs;
+        const u32 index = parameters[0];
+        if ((parameters[1] - regs.shadow_scratch[28 + index]) == 0) {
+            return;
+        }
+
+        regs.pipelines[index & 0xF].offset = parameters[2];
+        maxwell3d.dirty.flags[VideoCommon::Dirty::Shaders] = true;
+        regs.shadow_scratch[28 + index] = parameters[1];
+        regs.shadow_scratch[34 + index] = parameters[2];
+
+        const u32 address = parameters[4];
+        auto& const_buffer = regs.const_buffer;
+        const_buffer.size = 0x10000;
+        const_buffer.address_high = (address >> 24) & 0xFF;
+        const_buffer.address_low = address << 8;
+
+        const size_t bind_group_id = parameters[3] & 0x7F;
+        auto& bind_group = regs.bind_groups[bind_group_id];
+        bind_group.raw_config = 0x11;
+        maxwell3d.ProcessCBBind(bind_group_id);
+    }
+};
 
-class HLEMacroImpl final : public CachedMacro {
+class HLE_SetRasterBoundingBox final : public HLEMacroImpl {
 public:
-    explicit HLEMacroImpl(Engines::Maxwell3D& maxwell3d_, HLEFunction func_)
-        : maxwell3d{maxwell3d_}, func{func_} {}
+    explicit HLE_SetRasterBoundingBox(Maxwell3D& maxwell3d_) : HLEMacroImpl(maxwell3d_) {}
 
-    void Execute(const std::vector<u32>& parameters, u32 method) override {
-        func(maxwell3d, parameters);
+    void Execute(const std::vector<u32>& parameters, [[maybe_unused]] u32 method) override {
+        maxwell3d.RefreshParameters();
+        const u32 raster_mode = parameters[0];
+        auto& regs = maxwell3d.regs;
+        const u32 raster_enabled = maxwell3d.regs.conservative_raster_enable;
+        const u32 scratch_data = maxwell3d.regs.shadow_scratch[52];
+        regs.raster_bounding_box.raw = raster_mode & 0xFFFFF00F;
+        regs.raster_bounding_box.pad.Assign(scratch_data & raster_enabled);
+    }
+};
+
+template <size_t base_size>
+class HLE_ClearConstBuffer final : public HLEMacroImpl {
+public:
+    explicit HLE_ClearConstBuffer(Maxwell3D& maxwell3d_) : HLEMacroImpl(maxwell3d_) {}
+
+    void Execute(const std::vector<u32>& parameters, [[maybe_unused]] u32 method) override {
+        maxwell3d.RefreshParameters();
+        static constexpr std::array<u32, base_size> zeroes{};
+        auto& regs = maxwell3d.regs;
+        regs.const_buffer.size = static_cast<u32>(base_size);
+        regs.const_buffer.address_high = parameters[0];
+        regs.const_buffer.address_low = parameters[1];
+        regs.const_buffer.offset = 0;
+        maxwell3d.ProcessCBMultiData(zeroes.data(), parameters[2] * 4);
+    }
+};
+
+class HLE_ClearMemory final : public HLEMacroImpl {
+public:
+    explicit HLE_ClearMemory(Maxwell3D& maxwell3d_) : HLEMacroImpl(maxwell3d_) {}
+
+    void Execute(const std::vector<u32>& parameters, [[maybe_unused]] u32 method) override {
+        maxwell3d.RefreshParameters();
+
+        const u32 needed_memory = parameters[2] / sizeof(u32);
+        if (needed_memory > zero_memory.size()) {
+            zero_memory.resize(needed_memory, 0);
+        }
+        auto& regs = maxwell3d.regs;
+        regs.upload.line_length_in = parameters[2];
+        regs.upload.line_count = 1;
+        regs.upload.dest.address_high = parameters[0];
+        regs.upload.dest.address_low = parameters[1];
+        maxwell3d.CallMethod(static_cast<size_t>(MAXWELL3D_REG_INDEX(launch_dma)), 0x1011, true);
+        maxwell3d.CallMultiMethod(static_cast<size_t>(MAXWELL3D_REG_INDEX(inline_data)),
+                                  zero_memory.data(), needed_memory, needed_memory);
     }
 
 private:
-    Engines::Maxwell3D& maxwell3d;
-    HLEFunction func;
+    std::vector<u32> zero_memory;
+};
+
+class HLE_TransformFeedbackSetup final : public HLEMacroImpl {
+public:
+    explicit HLE_TransformFeedbackSetup(Maxwell3D& maxwell3d_) : HLEMacroImpl(maxwell3d_) {}
+
+    void Execute(const std::vector<u32>& parameters, [[maybe_unused]] u32 method) override {
+        maxwell3d.RefreshParameters();
+
+        auto& regs = maxwell3d.regs;
+        regs.transform_feedback_enabled = 1;
+        regs.transform_feedback.buffers[0].start_offset = 0;
+        regs.transform_feedback.buffers[1].start_offset = 0;
+        regs.transform_feedback.buffers[2].start_offset = 0;
+        regs.transform_feedback.buffers[3].start_offset = 0;
+
+        regs.upload.line_length_in = 4;
+        regs.upload.line_count = 1;
+        regs.upload.dest.address_high = parameters[0];
+        regs.upload.dest.address_low = parameters[1];
+        maxwell3d.CallMethod(static_cast<size_t>(MAXWELL3D_REG_INDEX(launch_dma)), 0x1011, true);
+        maxwell3d.CallMethod(static_cast<size_t>(MAXWELL3D_REG_INDEX(inline_data)),
+                             regs.transform_feedback.controls[0].stride, true);
+    }
 };
 
 } // Anonymous namespace
 
-HLEMacro::HLEMacro(Engines::Maxwell3D& maxwell3d_) : maxwell3d{maxwell3d_} {}
+HLEMacro::HLEMacro(Maxwell3D& maxwell3d_) : maxwell3d{maxwell3d_} {
+    builders.emplace(0xDD6A7FA92A7D2674ULL,
+                     std::function<std::unique_ptr<CachedMacro>(Maxwell3D&)>(
+                         [](Maxwell3D& maxwell3d__) -> std::unique_ptr<CachedMacro> {
+                             return std::make_unique<HLE_DrawArrays>(maxwell3d__);
+                         }));
+    builders.emplace(0x0D61FC9FAAC9FCADULL,
+                     std::function<std::unique_ptr<CachedMacro>(Maxwell3D&)>(
+                         [](Maxwell3D& maxwell3d__) -> std::unique_ptr<CachedMacro> {
+                             return std::make_unique<HLE_DrawArraysIndirect<false>>(maxwell3d__);
+                         }));
+    builders.emplace(0x8A4D173EB99A8603ULL,
+                     std::function<std::unique_ptr<CachedMacro>(Maxwell3D&)>(
+                         [](Maxwell3D& maxwell3d__) -> std::unique_ptr<CachedMacro> {
+                             return std::make_unique<HLE_DrawArraysIndirect<true>>(maxwell3d__);
+                         }));
+    builders.emplace(0x2DB33AADB741839CULL,
+                     std::function<std::unique_ptr<CachedMacro>(Maxwell3D&)>(
+                         [](Maxwell3D& maxwell3d__) -> std::unique_ptr<CachedMacro> {
+                             return std::make_unique<HLE_DrawIndexed>(maxwell3d__);
+                         }));
+    builders.emplace(0x771BB18C62444DA0ULL,
+                     std::function<std::unique_ptr<CachedMacro>(Maxwell3D&)>(
+                         [](Maxwell3D& maxwell3d__) -> std::unique_ptr<CachedMacro> {
+                             return std::make_unique<HLE_DrawIndexedIndirect<false>>(maxwell3d__);
+                         }));
+    builders.emplace(0x0217920100488FF7ULL,
+                     std::function<std::unique_ptr<CachedMacro>(Maxwell3D&)>(
+                         [](Maxwell3D& maxwell3d__) -> std::unique_ptr<CachedMacro> {
+                             return std::make_unique<HLE_DrawIndexedIndirect<true>>(maxwell3d__);
+                         }));
+    builders.emplace(0x3F5E74B9C9A50164ULL,
+                     std::function<std::unique_ptr<CachedMacro>(Maxwell3D&)>(
+                         [](Maxwell3D& maxwell3d__) -> std::unique_ptr<CachedMacro> {
+                             return std::make_unique<HLE_MultiDrawIndexedIndirectCount>(
+                                 maxwell3d__);
+                         }));
+    builders.emplace(0xEAD26C3E2109B06BULL,
+                     std::function<std::unique_ptr<CachedMacro>(Maxwell3D&)>(
+                         [](Maxwell3D& maxwell3d__) -> std::unique_ptr<CachedMacro> {
+                             return std::make_unique<HLE_MultiLayerClear>(maxwell3d__);
+                         }));
+    builders.emplace(0xC713C83D8F63CCF3ULL,
+                     std::function<std::unique_ptr<CachedMacro>(Maxwell3D&)>(
+                         [](Maxwell3D& maxwell3d__) -> std::unique_ptr<CachedMacro> {
+                             return std::make_unique<HLE_C713C83D8F63CCF3>(maxwell3d__);
+                         }));
+    builders.emplace(0xD7333D26E0A93EDEULL,
+                     std::function<std::unique_ptr<CachedMacro>(Maxwell3D&)>(
+                         [](Maxwell3D& maxwell3d__) -> std::unique_ptr<CachedMacro> {
+                             return std::make_unique<HLE_D7333D26E0A93EDE>(maxwell3d__);
+                         }));
+    builders.emplace(0xEB29B2A09AA06D38ULL,
+                     std::function<std::unique_ptr<CachedMacro>(Maxwell3D&)>(
+                         [](Maxwell3D& maxwell3d__) -> std::unique_ptr<CachedMacro> {
+                             return std::make_unique<HLE_BindShader>(maxwell3d__);
+                         }));
+    builders.emplace(0xDB1341DBEB4C8AF7ULL,
+                     std::function<std::unique_ptr<CachedMacro>(Maxwell3D&)>(
+                         [](Maxwell3D& maxwell3d__) -> std::unique_ptr<CachedMacro> {
+                             return std::make_unique<HLE_SetRasterBoundingBox>(maxwell3d__);
+                         }));
+    builders.emplace(0x6C97861D891EDf7EULL,
+                     std::function<std::unique_ptr<CachedMacro>(Maxwell3D&)>(
+                         [](Maxwell3D& maxwell3d__) -> std::unique_ptr<CachedMacro> {
+                             return std::make_unique<HLE_ClearConstBuffer<0x5F00>>(maxwell3d__);
+                         }));
+    builders.emplace(0xD246FDDF3A6173D7ULL,
+                     std::function<std::unique_ptr<CachedMacro>(Maxwell3D&)>(
+                         [](Maxwell3D& maxwell3d__) -> std::unique_ptr<CachedMacro> {
+                             return std::make_unique<HLE_ClearConstBuffer<0x7000>>(maxwell3d__);
+                         }));
+    builders.emplace(0xEE4D0004BEC8ECF4ULL,
+                     std::function<std::unique_ptr<CachedMacro>(Maxwell3D&)>(
+                         [](Maxwell3D& maxwell3d__) -> std::unique_ptr<CachedMacro> {
+                             return std::make_unique<HLE_ClearMemory>(maxwell3d__);
+                         }));
+    builders.emplace(0xFC0CF27F5FFAA661ULL,
+                     std::function<std::unique_ptr<CachedMacro>(Maxwell3D&)>(
+                         [](Maxwell3D& maxwell3d__) -> std::unique_ptr<CachedMacro> {
+                             return std::make_unique<HLE_TransformFeedbackSetup>(maxwell3d__);
+                         }));
+}
+
 HLEMacro::~HLEMacro() = default;
 
 std::unique_ptr<CachedMacro> HLEMacro::GetHLEProgram(u64 hash) const {
-    const auto it = std::find_if(hle_funcs.cbegin(), hle_funcs.cend(),
-                                 [hash](const auto& pair) { return pair.first == hash; });
-    if (it == hle_funcs.end()) {
+    const auto it = builders.find(hash);
+    if (it == builders.end()) {
         return nullptr;
     }
-    return std::make_unique<HLEMacroImpl>(maxwell3d, it->second);
+    return it->second(maxwell3d);
 }
 
 } // namespace Tegra
diff --git a/src/video_core/macro/macro_hle.h b/src/video_core/macro/macro_hle.h
index 625332c9d..33f92fab1 100644
--- a/src/video_core/macro/macro_hle.h
+++ b/src/video_core/macro/macro_hle.h
@@ -3,7 +3,10 @@
 
 #pragma once
 
+#include <functional>
 #include <memory>
+#include <unordered_map>
+
 #include "common/common_types.h"
 
 namespace Tegra {
@@ -23,6 +26,8 @@ public:
 
 private:
     Engines::Maxwell3D& maxwell3d;
+    std::unordered_map<u64, std::function<std::unique_ptr<CachedMacro>(Engines::Maxwell3D&)>>
+        builders;
 };
 
 } // namespace Tegra
diff --git a/src/video_core/memory_manager.cpp b/src/video_core/memory_manager.cpp
index 8c8dfcca6..3a5cdeb39 100644
--- a/src/video_core/memory_manager.cpp
+++ b/src/video_core/memory_manager.cpp
@@ -25,7 +25,8 @@ MemoryManager::MemoryManager(Core::System& system_, u64 address_space_bits_, u64
       address_space_bits{address_space_bits_}, page_bits{page_bits_}, big_page_bits{big_page_bits_},
       entries{}, big_entries{}, page_table{address_space_bits, address_space_bits + page_bits - 38,
                                            page_bits != big_page_bits ? page_bits : 0},
-      unique_identifier{unique_identifier_generator.fetch_add(1, std::memory_order_acq_rel)} {
+      kind_map{PTEKind::INVALID}, unique_identifier{unique_identifier_generator.fetch_add(
+                                      1, std::memory_order_acq_rel)} {
     address_space_size = 1ULL << address_space_bits;
     page_size = 1ULL << page_bits;
     page_mask = page_size - 1ULL;
@@ -41,11 +42,7 @@ MemoryManager::MemoryManager(Core::System& system_, u64 address_space_bits_, u64
     big_entries.resize(big_page_table_size / 32, 0);
     big_page_table_cpu.resize(big_page_table_size);
     big_page_continous.resize(big_page_table_size / continous_bits, 0);
-    std::array<PTEKind, 32> kind_valus;
-    kind_valus.fill(PTEKind::INVALID);
-    big_kinds.resize(big_page_table_size / 32, kind_valus);
     entries.resize(page_table_size / 32, 0);
-    kinds.resize(page_table_size / 32, kind_valus);
 }
 
 MemoryManager::~MemoryManager() = default;
@@ -83,38 +80,7 @@ void MemoryManager::SetEntry(size_t position, MemoryManager::EntryType entry) {
 }
 
 PTEKind MemoryManager::GetPageKind(GPUVAddr gpu_addr) const {
-    auto entry = GetEntry<true>(gpu_addr);
-    if (entry == EntryType::Mapped || entry == EntryType::Reserved) [[likely]] {
-        return GetKind<true>(gpu_addr);
-    } else {
-        return GetKind<false>(gpu_addr);
-    }
-}
-
-template <bool is_big_page>
-PTEKind MemoryManager::GetKind(size_t position) const {
-    if constexpr (is_big_page) {
-        position = position >> big_page_bits;
-        const size_t sub_index = position % 32;
-        return big_kinds[position / 32][sub_index];
-    } else {
-        position = position >> page_bits;
-        const size_t sub_index = position % 32;
-        return kinds[position / 32][sub_index];
-    }
-}
-
-template <bool is_big_page>
-void MemoryManager::SetKind(size_t position, PTEKind kind) {
-    if constexpr (is_big_page) {
-        position = position >> big_page_bits;
-        const size_t sub_index = position % 32;
-        big_kinds[position / 32][sub_index] = kind;
-    } else {
-        position = position >> page_bits;
-        const size_t sub_index = position % 32;
-        kinds[position / 32][sub_index] = kind;
-    }
+    return kind_map.GetValueAt(gpu_addr);
 }
 
 inline bool MemoryManager::IsBigPageContinous(size_t big_page_index) const {
@@ -141,7 +107,6 @@ GPUVAddr MemoryManager::PageTableOp(GPUVAddr gpu_addr, [[maybe_unused]] VAddr cp
         const GPUVAddr current_gpu_addr = gpu_addr + offset;
         [[maybe_unused]] const auto current_entry_type = GetEntry<false>(current_gpu_addr);
         SetEntry<false>(current_gpu_addr, entry_type);
-        SetKind<false>(current_gpu_addr, kind);
         if (current_entry_type != entry_type) {
             rasterizer->ModifyGPUMemory(unique_identifier, gpu_addr, page_size);
         }
@@ -153,6 +118,7 @@ GPUVAddr MemoryManager::PageTableOp(GPUVAddr gpu_addr, [[maybe_unused]] VAddr cp
         }
         remaining_size -= page_size;
     }
+    kind_map.Map(gpu_addr, gpu_addr + size, kind);
     return gpu_addr;
 }
 
@@ -164,7 +130,6 @@ GPUVAddr MemoryManager::BigPageTableOp(GPUVAddr gpu_addr, [[maybe_unused]] VAddr
         const GPUVAddr current_gpu_addr = gpu_addr + offset;
         [[maybe_unused]] const auto current_entry_type = GetEntry<true>(current_gpu_addr);
         SetEntry<true>(current_gpu_addr, entry_type);
-        SetKind<true>(current_gpu_addr, kind);
         if (current_entry_type != entry_type) {
             rasterizer->ModifyGPUMemory(unique_identifier, gpu_addr, big_page_size);
         }
@@ -193,6 +158,7 @@ GPUVAddr MemoryManager::BigPageTableOp(GPUVAddr gpu_addr, [[maybe_unused]] VAddr
         }
         remaining_size -= big_page_size;
     }
+    kind_map.Map(gpu_addr, gpu_addr + size, kind);
     return gpu_addr;
 }
 
@@ -325,9 +291,15 @@ template <bool is_big_pages, typename FuncMapped, typename FuncReserved, typenam
 inline void MemoryManager::MemoryOperation(GPUVAddr gpu_src_addr, std::size_t size,
                                            FuncMapped&& func_mapped, FuncReserved&& func_reserved,
                                            FuncUnmapped&& func_unmapped) const {
-    static constexpr bool BOOL_BREAK_MAPPED = std::is_same_v<FuncMapped, bool>;
-    static constexpr bool BOOL_BREAK_RESERVED = std::is_same_v<FuncReserved, bool>;
-    static constexpr bool BOOL_BREAK_UNMAPPED = std::is_same_v<FuncUnmapped, bool>;
+    using FuncMappedReturn =
+        typename std::invoke_result<FuncMapped, std::size_t, std::size_t, std::size_t>::type;
+    using FuncReservedReturn =
+        typename std::invoke_result<FuncReserved, std::size_t, std::size_t, std::size_t>::type;
+    using FuncUnmappedReturn =
+        typename std::invoke_result<FuncUnmapped, std::size_t, std::size_t, std::size_t>::type;
+    static constexpr bool BOOL_BREAK_MAPPED = std::is_same_v<FuncMappedReturn, bool>;
+    static constexpr bool BOOL_BREAK_RESERVED = std::is_same_v<FuncReservedReturn, bool>;
+    static constexpr bool BOOL_BREAK_UNMAPPED = std::is_same_v<FuncUnmappedReturn, bool>;
     u64 used_page_size;
     u64 used_page_mask;
     u64 used_page_bits;
@@ -384,8 +356,8 @@ inline void MemoryManager::MemoryOperation(GPUVAddr gpu_src_addr, std::size_t si
 }
 
 template <bool is_safe>
-void MemoryManager::ReadBlockImpl(GPUVAddr gpu_src_addr, void* dest_buffer,
-                                  std::size_t size) const {
+void MemoryManager::ReadBlockImpl(GPUVAddr gpu_src_addr, void* dest_buffer, std::size_t size,
+                                  [[maybe_unused]] VideoCommon::CacheType which) const {
     auto set_to_zero = [&]([[maybe_unused]] std::size_t page_index,
                            [[maybe_unused]] std::size_t offset, std::size_t copy_amount) {
         std::memset(dest_buffer, 0, copy_amount);
@@ -395,7 +367,7 @@ void MemoryManager::ReadBlockImpl(GPUVAddr gpu_src_addr, void* dest_buffer,
         const VAddr cpu_addr_base =
             (static_cast<VAddr>(page_table[page_index]) << cpu_page_bits) + offset;
         if constexpr (is_safe) {
-            rasterizer->FlushRegion(cpu_addr_base, copy_amount);
+            rasterizer->FlushRegion(cpu_addr_base, copy_amount, which);
         }
         u8* physical = memory.GetPointer(cpu_addr_base);
         std::memcpy(dest_buffer, physical, copy_amount);
@@ -405,7 +377,7 @@ void MemoryManager::ReadBlockImpl(GPUVAddr gpu_src_addr, void* dest_buffer,
         const VAddr cpu_addr_base =
             (static_cast<VAddr>(big_page_table_cpu[page_index]) << cpu_page_bits) + offset;
         if constexpr (is_safe) {
-            rasterizer->FlushRegion(cpu_addr_base, copy_amount);
+            rasterizer->FlushRegion(cpu_addr_base, copy_amount, which);
         }
         if (!IsBigPageContinous(page_index)) [[unlikely]] {
             memory.ReadBlockUnsafe(cpu_addr_base, dest_buffer, copy_amount);
@@ -423,18 +395,19 @@ void MemoryManager::ReadBlockImpl(GPUVAddr gpu_src_addr, void* dest_buffer,
     MemoryOperation<true>(gpu_src_addr, size, mapped_big, set_to_zero, read_short_pages);
 }
 
-void MemoryManager::ReadBlock(GPUVAddr gpu_src_addr, void* dest_buffer, std::size_t size) const {
-    ReadBlockImpl<true>(gpu_src_addr, dest_buffer, size);
+void MemoryManager::ReadBlock(GPUVAddr gpu_src_addr, void* dest_buffer, std::size_t size,
+                              VideoCommon::CacheType which) const {
+    ReadBlockImpl<true>(gpu_src_addr, dest_buffer, size, which);
 }
 
 void MemoryManager::ReadBlockUnsafe(GPUVAddr gpu_src_addr, void* dest_buffer,
                                     const std::size_t size) const {
-    ReadBlockImpl<false>(gpu_src_addr, dest_buffer, size);
+    ReadBlockImpl<false>(gpu_src_addr, dest_buffer, size, VideoCommon::CacheType::None);
 }
 
 template <bool is_safe>
-void MemoryManager::WriteBlockImpl(GPUVAddr gpu_dest_addr, const void* src_buffer,
-                                   std::size_t size) {
+void MemoryManager::WriteBlockImpl(GPUVAddr gpu_dest_addr, const void* src_buffer, std::size_t size,
+                                   [[maybe_unused]] VideoCommon::CacheType which) {
     auto just_advance = [&]([[maybe_unused]] std::size_t page_index,
                             [[maybe_unused]] std::size_t offset, std::size_t copy_amount) {
         src_buffer = static_cast<const u8*>(src_buffer) + copy_amount;
@@ -443,7 +416,7 @@ void MemoryManager::WriteBlockImpl(GPUVAddr gpu_dest_addr, const void* src_buffe
         const VAddr cpu_addr_base =
             (static_cast<VAddr>(page_table[page_index]) << cpu_page_bits) + offset;
         if constexpr (is_safe) {
-            rasterizer->InvalidateRegion(cpu_addr_base, copy_amount);
+            rasterizer->InvalidateRegion(cpu_addr_base, copy_amount, which);
         }
         u8* physical = memory.GetPointer(cpu_addr_base);
         std::memcpy(physical, src_buffer, copy_amount);
@@ -453,7 +426,7 @@ void MemoryManager::WriteBlockImpl(GPUVAddr gpu_dest_addr, const void* src_buffe
         const VAddr cpu_addr_base =
             (static_cast<VAddr>(big_page_table_cpu[page_index]) << cpu_page_bits) + offset;
         if constexpr (is_safe) {
-            rasterizer->InvalidateRegion(cpu_addr_base, copy_amount);
+            rasterizer->InvalidateRegion(cpu_addr_base, copy_amount, which);
         }
         if (!IsBigPageContinous(page_index)) [[unlikely]] {
             memory.WriteBlockUnsafe(cpu_addr_base, src_buffer, copy_amount);
@@ -471,16 +444,18 @@ void MemoryManager::WriteBlockImpl(GPUVAddr gpu_dest_addr, const void* src_buffe
     MemoryOperation<true>(gpu_dest_addr, size, mapped_big, just_advance, write_short_pages);
 }
 
-void MemoryManager::WriteBlock(GPUVAddr gpu_dest_addr, const void* src_buffer, std::size_t size) {
-    WriteBlockImpl<true>(gpu_dest_addr, src_buffer, size);
+void MemoryManager::WriteBlock(GPUVAddr gpu_dest_addr, const void* src_buffer, std::size_t size,
+                               VideoCommon::CacheType which) {
+    WriteBlockImpl<true>(gpu_dest_addr, src_buffer, size, which);
 }
 
 void MemoryManager::WriteBlockUnsafe(GPUVAddr gpu_dest_addr, const void* src_buffer,
                                      std::size_t size) {
-    WriteBlockImpl<false>(gpu_dest_addr, src_buffer, size);
+    WriteBlockImpl<false>(gpu_dest_addr, src_buffer, size, VideoCommon::CacheType::None);
 }
 
-void MemoryManager::FlushRegion(GPUVAddr gpu_addr, size_t size) const {
+void MemoryManager::FlushRegion(GPUVAddr gpu_addr, size_t size,
+                                VideoCommon::CacheType which) const {
     auto do_nothing = [&]([[maybe_unused]] std::size_t page_index,
                           [[maybe_unused]] std::size_t offset,
                           [[maybe_unused]] std::size_t copy_amount) {};
@@ -488,12 +463,12 @@ void MemoryManager::FlushRegion(GPUVAddr gpu_addr, size_t size) const {
     auto mapped_normal = [&](std::size_t page_index, std::size_t offset, std::size_t copy_amount) {
         const VAddr cpu_addr_base =
             (static_cast<VAddr>(page_table[page_index]) << cpu_page_bits) + offset;
-        rasterizer->FlushRegion(cpu_addr_base, copy_amount);
+        rasterizer->FlushRegion(cpu_addr_base, copy_amount, which);
     };
     auto mapped_big = [&](std::size_t page_index, std::size_t offset, std::size_t copy_amount) {
         const VAddr cpu_addr_base =
             (static_cast<VAddr>(big_page_table_cpu[page_index]) << cpu_page_bits) + offset;
-        rasterizer->FlushRegion(cpu_addr_base, copy_amount);
+        rasterizer->FlushRegion(cpu_addr_base, copy_amount, which);
     };
     auto flush_short_pages = [&](std::size_t page_index, std::size_t offset,
                                  std::size_t copy_amount) {
@@ -503,7 +478,8 @@ void MemoryManager::FlushRegion(GPUVAddr gpu_addr, size_t size) const {
     MemoryOperation<true>(gpu_addr, size, mapped_big, do_nothing, flush_short_pages);
 }
 
-bool MemoryManager::IsMemoryDirty(GPUVAddr gpu_addr, size_t size) const {
+bool MemoryManager::IsMemoryDirty(GPUVAddr gpu_addr, size_t size,
+                                  VideoCommon::CacheType which) const {
     bool result = false;
     auto do_nothing = [&]([[maybe_unused]] std::size_t page_index,
                           [[maybe_unused]] std::size_t offset,
@@ -512,13 +488,13 @@ bool MemoryManager::IsMemoryDirty(GPUVAddr gpu_addr, size_t size) const {
     auto mapped_normal = [&](std::size_t page_index, std::size_t offset, std::size_t copy_amount) {
         const VAddr cpu_addr_base =
             (static_cast<VAddr>(page_table[page_index]) << cpu_page_bits) + offset;
-        result |= rasterizer->MustFlushRegion(cpu_addr_base, copy_amount);
+        result |= rasterizer->MustFlushRegion(cpu_addr_base, copy_amount, which);
         return result;
     };
     auto mapped_big = [&](std::size_t page_index, std::size_t offset, std::size_t copy_amount) {
         const VAddr cpu_addr_base =
             (static_cast<VAddr>(big_page_table_cpu[page_index]) << cpu_page_bits) + offset;
-        result |= rasterizer->MustFlushRegion(cpu_addr_base, copy_amount);
+        result |= rasterizer->MustFlushRegion(cpu_addr_base, copy_amount, which);
         return result;
     };
     auto check_short_pages = [&](std::size_t page_index, std::size_t offset,
@@ -571,7 +547,12 @@ size_t MemoryManager::MaxContinousRange(GPUVAddr gpu_addr, size_t size) const {
     return range_so_far;
 }
 
-void MemoryManager::InvalidateRegion(GPUVAddr gpu_addr, size_t size) const {
+size_t MemoryManager::GetMemoryLayoutSize(GPUVAddr gpu_addr, size_t max_size) const {
+    return kind_map.GetContinousSizeFrom(gpu_addr);
+}
+
+void MemoryManager::InvalidateRegion(GPUVAddr gpu_addr, size_t size,
+                                     VideoCommon::CacheType which) const {
     auto do_nothing = [&]([[maybe_unused]] std::size_t page_index,
                           [[maybe_unused]] std::size_t offset,
                           [[maybe_unused]] std::size_t copy_amount) {};
@@ -579,12 +560,12 @@ void MemoryManager::InvalidateRegion(GPUVAddr gpu_addr, size_t size) const {
     auto mapped_normal = [&](std::size_t page_index, std::size_t offset, std::size_t copy_amount) {
         const VAddr cpu_addr_base =
             (static_cast<VAddr>(page_table[page_index]) << cpu_page_bits) + offset;
-        rasterizer->InvalidateRegion(cpu_addr_base, copy_amount);
+        rasterizer->InvalidateRegion(cpu_addr_base, copy_amount, which);
     };
     auto mapped_big = [&](std::size_t page_index, std::size_t offset, std::size_t copy_amount) {
         const VAddr cpu_addr_base =
             (static_cast<VAddr>(big_page_table_cpu[page_index]) << cpu_page_bits) + offset;
-        rasterizer->InvalidateRegion(cpu_addr_base, copy_amount);
+        rasterizer->InvalidateRegion(cpu_addr_base, copy_amount, which);
     };
     auto invalidate_short_pages = [&](std::size_t page_index, std::size_t offset,
                                       std::size_t copy_amount) {
@@ -594,14 +575,15 @@ void MemoryManager::InvalidateRegion(GPUVAddr gpu_addr, size_t size) const {
     MemoryOperation<true>(gpu_addr, size, mapped_big, do_nothing, invalidate_short_pages);
 }
 
-void MemoryManager::CopyBlock(GPUVAddr gpu_dest_addr, GPUVAddr gpu_src_addr, std::size_t size) {
+void MemoryManager::CopyBlock(GPUVAddr gpu_dest_addr, GPUVAddr gpu_src_addr, std::size_t size,
+                              VideoCommon::CacheType which) {
     std::vector<u8> tmp_buffer(size);
-    ReadBlock(gpu_src_addr, tmp_buffer.data(), size);
+    ReadBlock(gpu_src_addr, tmp_buffer.data(), size, which);
 
     // The output block must be flushed in case it has data modified from the GPU.
     // Fixes NPC geometry in Zombie Panic in Wonderland DX
-    FlushRegion(gpu_dest_addr, size);
-    WriteBlock(gpu_dest_addr, tmp_buffer.data(), size);
+    FlushRegion(gpu_dest_addr, size, which);
+    WriteBlock(gpu_dest_addr, tmp_buffer.data(), size, which);
 }
 
 bool MemoryManager::IsGranularRange(GPUVAddr gpu_addr, std::size_t size) const {
diff --git a/src/video_core/memory_manager.h b/src/video_core/memory_manager.h
index ab4bc9ec6..828e13439 100644
--- a/src/video_core/memory_manager.h
+++ b/src/video_core/memory_manager.h
@@ -10,7 +10,9 @@
 
 #include "common/common_types.h"
 #include "common/multi_level_page_table.h"
+#include "common/range_map.h"
 #include "common/virtual_buffer.h"
+#include "video_core/cache_types.h"
 #include "video_core/pte_kind.h"
 
 namespace VideoCore {
@@ -59,9 +61,12 @@ public:
      * in the Host Memory counterpart. Note: This functions cause Host GPU Memory
      * Flushes and Invalidations, respectively to each operation.
      */
-    void ReadBlock(GPUVAddr gpu_src_addr, void* dest_buffer, std::size_t size) const;
-    void WriteBlock(GPUVAddr gpu_dest_addr, const void* src_buffer, std::size_t size);
-    void CopyBlock(GPUVAddr gpu_dest_addr, GPUVAddr gpu_src_addr, std::size_t size);
+    void ReadBlock(GPUVAddr gpu_src_addr, void* dest_buffer, std::size_t size,
+                   VideoCommon::CacheType which = VideoCommon::CacheType::All) const;
+    void WriteBlock(GPUVAddr gpu_dest_addr, const void* src_buffer, std::size_t size,
+                    VideoCommon::CacheType which = VideoCommon::CacheType::All);
+    void CopyBlock(GPUVAddr gpu_dest_addr, GPUVAddr gpu_src_addr, std::size_t size,
+                   VideoCommon::CacheType which = VideoCommon::CacheType::All);
 
     /**
      * ReadBlockUnsafe and WriteBlockUnsafe are special versions of ReadBlock and
@@ -104,11 +109,14 @@ public:
     GPUVAddr MapSparse(GPUVAddr gpu_addr, std::size_t size, bool is_big_pages = true);
     void Unmap(GPUVAddr gpu_addr, std::size_t size);
 
-    void FlushRegion(GPUVAddr gpu_addr, size_t size) const;
+    void FlushRegion(GPUVAddr gpu_addr, size_t size,
+                     VideoCommon::CacheType which = VideoCommon::CacheType::All) const;
 
-    void InvalidateRegion(GPUVAddr gpu_addr, size_t size) const;
+    void InvalidateRegion(GPUVAddr gpu_addr, size_t size,
+                          VideoCommon::CacheType which = VideoCommon::CacheType::All) const;
 
-    bool IsMemoryDirty(GPUVAddr gpu_addr, size_t size) const;
+    bool IsMemoryDirty(GPUVAddr gpu_addr, size_t size,
+                       VideoCommon::CacheType which = VideoCommon::CacheType::All) const;
 
     size_t MaxContinousRange(GPUVAddr gpu_addr, size_t size) const;
 
@@ -118,16 +126,21 @@ public:
 
     PTEKind GetPageKind(GPUVAddr gpu_addr) const;
 
+    size_t GetMemoryLayoutSize(GPUVAddr gpu_addr,
+                               size_t max_size = std::numeric_limits<size_t>::max()) const;
+
 private:
     template <bool is_big_pages, typename FuncMapped, typename FuncReserved, typename FuncUnmapped>
     inline void MemoryOperation(GPUVAddr gpu_src_addr, std::size_t size, FuncMapped&& func_mapped,
                                 FuncReserved&& func_reserved, FuncUnmapped&& func_unmapped) const;
 
     template <bool is_safe>
-    void ReadBlockImpl(GPUVAddr gpu_src_addr, void* dest_buffer, std::size_t size) const;
+    void ReadBlockImpl(GPUVAddr gpu_src_addr, void* dest_buffer, std::size_t size,
+                       VideoCommon::CacheType which) const;
 
     template <bool is_safe>
-    void WriteBlockImpl(GPUVAddr gpu_dest_addr, const void* src_buffer, std::size_t size);
+    void WriteBlockImpl(GPUVAddr gpu_dest_addr, const void* src_buffer, std::size_t size,
+                        VideoCommon::CacheType which);
 
     template <bool is_big_page>
     [[nodiscard]] std::size_t PageEntryIndex(GPUVAddr gpu_addr) const {
@@ -183,16 +196,8 @@ private:
     template <bool is_big_page>
     inline void SetEntry(size_t position, EntryType entry);
 
-    std::vector<std::array<PTEKind, 32>> kinds;
-    std::vector<std::array<PTEKind, 32>> big_kinds;
-
-    template <bool is_big_page>
-    inline PTEKind GetKind(size_t position) const;
-
-    template <bool is_big_page>
-    inline void SetKind(size_t position, PTEKind kind);
-
     Common::MultiLevelPageTable<u32> page_table;
+    Common::RangeMap<GPUVAddr, PTEKind> kind_map;
     Common::VirtualBuffer<u32> big_page_table_cpu;
 
     std::vector<u64> big_page_continous;
diff --git a/src/video_core/rasterizer_interface.h b/src/video_core/rasterizer_interface.h
index b6907463c..f44c7df50 100644
--- a/src/video_core/rasterizer_interface.h
+++ b/src/video_core/rasterizer_interface.h
@@ -8,6 +8,7 @@
 #include <span>
 #include "common/common_types.h"
 #include "common/polyfill_thread.h"
+#include "video_core/cache_types.h"
 #include "video_core/engines/fermi_2d.h"
 #include "video_core/gpu.h"
 
@@ -42,6 +43,9 @@ public:
     /// Dispatches a draw invocation
     virtual void Draw(bool is_indexed, u32 instance_count) = 0;
 
+    /// Dispatches an indirect draw invocation
+    virtual void DrawIndirect() {}
+
     /// Clear the current framebuffer
     virtual void Clear(u32 layer_count) = 0;
 
@@ -80,13 +84,16 @@ public:
     virtual void FlushAll() = 0;
 
     /// Notify rasterizer that any caches of the specified region should be flushed to Switch memory
-    virtual void FlushRegion(VAddr addr, u64 size) = 0;
+    virtual void FlushRegion(VAddr addr, u64 size,
+                             VideoCommon::CacheType which = VideoCommon::CacheType::All) = 0;
 
     /// Check if the the specified memory area requires flushing to CPU Memory.
-    virtual bool MustFlushRegion(VAddr addr, u64 size) = 0;
+    virtual bool MustFlushRegion(VAddr addr, u64 size,
+                                 VideoCommon::CacheType which = VideoCommon::CacheType::All) = 0;
 
     /// Notify rasterizer that any caches of the specified region should be invalidated
-    virtual void InvalidateRegion(VAddr addr, u64 size) = 0;
+    virtual void InvalidateRegion(VAddr addr, u64 size,
+                                  VideoCommon::CacheType which = VideoCommon::CacheType::All) = 0;
 
     /// Notify rasterizer that any caches of the specified region are desync with guest
     virtual void OnCPUWrite(VAddr addr, u64 size) = 0;
@@ -102,7 +109,8 @@ public:
 
     /// Notify rasterizer that any caches of the specified region should be flushed to Switch memory
     /// and invalidated
-    virtual void FlushAndInvalidateRegion(VAddr addr, u64 size) = 0;
+    virtual void FlushAndInvalidateRegion(
+        VAddr addr, u64 size, VideoCommon::CacheType which = VideoCommon::CacheType::All) = 0;
 
     /// Notify the host renderer to wait for previous primitive and compute operations.
     virtual void WaitForIdle() = 0;
@@ -119,6 +127,10 @@ public:
     /// Notify rasterizer that a frame is about to finish
     virtual void TickFrame() = 0;
 
+    virtual bool AccelerateConditionalRendering() {
+        return false;
+    }
+
     /// Attempt to use a faster method to perform a surface copy
     [[nodiscard]] virtual bool AccelerateSurfaceCopy(
         const Tegra::Engines::Fermi2D::Surface& src, const Tegra::Engines::Fermi2D::Surface& dst,
diff --git a/src/video_core/renderer_null/null_rasterizer.cpp b/src/video_core/renderer_null/null_rasterizer.cpp
index 9734d84bc..2c11345d7 100644
--- a/src/video_core/renderer_null/null_rasterizer.cpp
+++ b/src/video_core/renderer_null/null_rasterizer.cpp
@@ -39,11 +39,11 @@ void RasterizerNull::BindGraphicsUniformBuffer(size_t stage, u32 index, GPUVAddr
                                                u32 size) {}
 void RasterizerNull::DisableGraphicsUniformBuffer(size_t stage, u32 index) {}
 void RasterizerNull::FlushAll() {}
-void RasterizerNull::FlushRegion(VAddr addr, u64 size) {}
-bool RasterizerNull::MustFlushRegion(VAddr addr, u64 size) {
+void RasterizerNull::FlushRegion(VAddr addr, u64 size, VideoCommon::CacheType) {}
+bool RasterizerNull::MustFlushRegion(VAddr addr, u64 size, VideoCommon::CacheType) {
     return false;
 }
-void RasterizerNull::InvalidateRegion(VAddr addr, u64 size) {}
+void RasterizerNull::InvalidateRegion(VAddr addr, u64 size, VideoCommon::CacheType) {}
 void RasterizerNull::OnCPUWrite(VAddr addr, u64 size) {}
 void RasterizerNull::InvalidateGPUCache() {}
 void RasterizerNull::UnmapMemory(VAddr addr, u64 size) {}
@@ -61,7 +61,7 @@ void RasterizerNull::SignalSyncPoint(u32 value) {
 }
 void RasterizerNull::SignalReference() {}
 void RasterizerNull::ReleaseFences() {}
-void RasterizerNull::FlushAndInvalidateRegion(VAddr addr, u64 size) {}
+void RasterizerNull::FlushAndInvalidateRegion(VAddr addr, u64 size, VideoCommon::CacheType) {}
 void RasterizerNull::WaitForIdle() {}
 void RasterizerNull::FragmentBarrier() {}
 void RasterizerNull::TiledCacheBarrier() {}
diff --git a/src/video_core/renderer_null/null_rasterizer.h b/src/video_core/renderer_null/null_rasterizer.h
index ecf77ba42..2112aa70e 100644
--- a/src/video_core/renderer_null/null_rasterizer.h
+++ b/src/video_core/renderer_null/null_rasterizer.h
@@ -38,9 +38,12 @@ public:
     void BindGraphicsUniformBuffer(size_t stage, u32 index, GPUVAddr gpu_addr, u32 size) override;
     void DisableGraphicsUniformBuffer(size_t stage, u32 index) override;
     void FlushAll() override;
-    void FlushRegion(VAddr addr, u64 size) override;
-    bool MustFlushRegion(VAddr addr, u64 size) override;
-    void InvalidateRegion(VAddr addr, u64 size) override;
+    void FlushRegion(VAddr addr, u64 size,
+                     VideoCommon::CacheType which = VideoCommon::CacheType::All) override;
+    bool MustFlushRegion(VAddr addr, u64 size,
+                         VideoCommon::CacheType which = VideoCommon::CacheType::All) override;
+    void InvalidateRegion(VAddr addr, u64 size,
+                          VideoCommon::CacheType which = VideoCommon::CacheType::All) override;
     void OnCPUWrite(VAddr addr, u64 size) override;
     void InvalidateGPUCache() override;
     void UnmapMemory(VAddr addr, u64 size) override;
@@ -50,7 +53,8 @@ public:
     void SignalSyncPoint(u32 value) override;
     void SignalReference() override;
     void ReleaseFences() override;
-    void FlushAndInvalidateRegion(VAddr addr, u64 size) override;
+    void FlushAndInvalidateRegion(
+        VAddr addr, u64 size, VideoCommon::CacheType which = VideoCommon::CacheType::All) override;
     void WaitForIdle() override;
     void FragmentBarrier() override;
     void TiledCacheBarrier() override;
diff --git a/src/video_core/renderer_opengl/gl_graphics_pipeline.h b/src/video_core/renderer_opengl/gl_graphics_pipeline.h
index ea53ddb46..1c06b3655 100644
--- a/src/video_core/renderer_opengl/gl_graphics_pipeline.h
+++ b/src/video_core/renderer_opengl/gl_graphics_pipeline.h
@@ -40,6 +40,7 @@ struct GraphicsPipelineKey {
         BitField<6, 2, Maxwell::Tessellation::DomainType> tessellation_primitive;
         BitField<8, 2, Maxwell::Tessellation::Spacing> tessellation_spacing;
         BitField<10, 1, u32> tessellation_clockwise;
+        BitField<11, 3, Tegra::Engines::Maxwell3D::EngineHint> app_stage;
     };
     std::array<u32, 3> padding;
     VideoCommon::TransformFeedbackState xfb_state;
diff --git a/src/video_core/renderer_opengl/gl_rasterizer.cpp b/src/video_core/renderer_opengl/gl_rasterizer.cpp
index a44b8c454..7d48af8e1 100644
--- a/src/video_core/renderer_opengl/gl_rasterizer.cpp
+++ b/src/video_core/renderer_opengl/gl_rasterizer.cpp
@@ -202,7 +202,8 @@ void RasterizerOpenGL::Clear(u32 layer_count) {
     ++num_queued_commands;
 }
 
-void RasterizerOpenGL::Draw(bool is_indexed, u32 instance_count) {
+template <typename Func>
+void RasterizerOpenGL::PrepareDraw(bool is_indexed, Func&& draw_func) {
     MICROPROFILE_SCOPE(OpenGL_Drawing);
 
     SCOPE_EXIT({ gpu.TickWork(); });
@@ -226,48 +227,97 @@ void RasterizerOpenGL::Draw(bool is_indexed, u32 instance_count) {
     const GLenum primitive_mode = MaxwellToGL::PrimitiveTopology(draw_state.topology);
     BeginTransformFeedback(pipeline, primitive_mode);
 
-    const GLuint base_instance = static_cast<GLuint>(draw_state.base_instance);
-    const GLsizei num_instances = static_cast<GLsizei>(instance_count);
-    if (is_indexed) {
-        const GLint base_vertex = static_cast<GLint>(draw_state.base_index);
-        const GLsizei num_vertices = static_cast<GLsizei>(draw_state.index_buffer.count);
-        const GLvoid* const offset = buffer_cache_runtime.IndexOffset();
-        const GLenum format = MaxwellToGL::IndexFormat(draw_state.index_buffer.format);
-        if (num_instances == 1 && base_instance == 0 && base_vertex == 0) {
-            glDrawElements(primitive_mode, num_vertices, format, offset);
-        } else if (num_instances == 1 && base_instance == 0) {
-            glDrawElementsBaseVertex(primitive_mode, num_vertices, format, offset, base_vertex);
-        } else if (base_vertex == 0 && base_instance == 0) {
-            glDrawElementsInstanced(primitive_mode, num_vertices, format, offset, num_instances);
-        } else if (base_vertex == 0) {
-            glDrawElementsInstancedBaseInstance(primitive_mode, num_vertices, format, offset,
-                                                num_instances, base_instance);
-        } else if (base_instance == 0) {
-            glDrawElementsInstancedBaseVertex(primitive_mode, num_vertices, format, offset,
-                                              num_instances, base_vertex);
-        } else {
-            glDrawElementsInstancedBaseVertexBaseInstance(primitive_mode, num_vertices, format,
-                                                          offset, num_instances, base_vertex,
-                                                          base_instance);
-        }
-    } else {
-        const GLint base_vertex = static_cast<GLint>(draw_state.vertex_buffer.first);
-        const GLsizei num_vertices = static_cast<GLsizei>(draw_state.vertex_buffer.count);
-        if (num_instances == 1 && base_instance == 0) {
-            glDrawArrays(primitive_mode, base_vertex, num_vertices);
-        } else if (base_instance == 0) {
-            glDrawArraysInstanced(primitive_mode, base_vertex, num_vertices, num_instances);
-        } else {
-            glDrawArraysInstancedBaseInstance(primitive_mode, base_vertex, num_vertices,
-                                              num_instances, base_instance);
-        }
-    }
+    draw_func(primitive_mode);
+
     EndTransformFeedback();
 
     ++num_queued_commands;
     has_written_global_memory |= pipeline->WritesGlobalMemory();
 }
 
+void RasterizerOpenGL::Draw(bool is_indexed, u32 instance_count) {
+    PrepareDraw(is_indexed, [this, is_indexed, instance_count](GLenum primitive_mode) {
+        const auto& draw_state = maxwell3d->draw_manager->GetDrawState();
+        const GLuint base_instance = static_cast<GLuint>(draw_state.base_instance);
+        const GLsizei num_instances = static_cast<GLsizei>(instance_count);
+        if (is_indexed) {
+            const GLint base_vertex = static_cast<GLint>(draw_state.base_index);
+            const GLsizei num_vertices = static_cast<GLsizei>(draw_state.index_buffer.count);
+            const GLvoid* const offset = buffer_cache_runtime.IndexOffset();
+            const GLenum format = MaxwellToGL::IndexFormat(draw_state.index_buffer.format);
+            if (num_instances == 1 && base_instance == 0 && base_vertex == 0) {
+                glDrawElements(primitive_mode, num_vertices, format, offset);
+            } else if (num_instances == 1 && base_instance == 0) {
+                glDrawElementsBaseVertex(primitive_mode, num_vertices, format, offset, base_vertex);
+            } else if (base_vertex == 0 && base_instance == 0) {
+                glDrawElementsInstanced(primitive_mode, num_vertices, format, offset,
+                                        num_instances);
+            } else if (base_vertex == 0) {
+                glDrawElementsInstancedBaseInstance(primitive_mode, num_vertices, format, offset,
+                                                    num_instances, base_instance);
+            } else if (base_instance == 0) {
+                glDrawElementsInstancedBaseVertex(primitive_mode, num_vertices, format, offset,
+                                                  num_instances, base_vertex);
+            } else {
+                glDrawElementsInstancedBaseVertexBaseInstance(primitive_mode, num_vertices, format,
+                                                              offset, num_instances, base_vertex,
+                                                              base_instance);
+            }
+        } else {
+            const GLint base_vertex = static_cast<GLint>(draw_state.vertex_buffer.first);
+            const GLsizei num_vertices = static_cast<GLsizei>(draw_state.vertex_buffer.count);
+            if (num_instances == 1 && base_instance == 0) {
+                glDrawArrays(primitive_mode, base_vertex, num_vertices);
+            } else if (base_instance == 0) {
+                glDrawArraysInstanced(primitive_mode, base_vertex, num_vertices, num_instances);
+            } else {
+                glDrawArraysInstancedBaseInstance(primitive_mode, base_vertex, num_vertices,
+                                                  num_instances, base_instance);
+            }
+        }
+    });
+}
+
+void RasterizerOpenGL::DrawIndirect() {
+    const auto& params = maxwell3d->draw_manager->GetIndirectParams();
+    buffer_cache.SetDrawIndirect(&params);
+    PrepareDraw(params.is_indexed, [this, &params](GLenum primitive_mode) {
+        const auto [buffer, offset] = buffer_cache.GetDrawIndirectBuffer();
+        const GLvoid* const gl_offset =
+            reinterpret_cast<const GLvoid*>(static_cast<uintptr_t>(offset));
+        glBindBuffer(GL_DRAW_INDIRECT_BUFFER, buffer->Handle());
+        if (params.include_count) {
+            const auto [draw_buffer, offset_base] = buffer_cache.GetDrawIndirectCount();
+            glBindBuffer(GL_PARAMETER_BUFFER, draw_buffer->Handle());
+
+            if (params.is_indexed) {
+                const GLenum format = MaxwellToGL::IndexFormat(maxwell3d->regs.index_buffer.format);
+                glMultiDrawElementsIndirectCount(primitive_mode, format, gl_offset,
+                                                 static_cast<GLintptr>(offset_base),
+                                                 static_cast<GLsizei>(params.max_draw_counts),
+                                                 static_cast<GLsizei>(params.stride));
+            } else {
+                glMultiDrawArraysIndirectCount(primitive_mode, gl_offset,
+                                               static_cast<GLintptr>(offset_base),
+                                               static_cast<GLsizei>(params.max_draw_counts),
+                                               static_cast<GLsizei>(params.stride));
+            }
+            return;
+        }
+        if (params.is_indexed) {
+            const GLenum format = MaxwellToGL::IndexFormat(maxwell3d->regs.index_buffer.format);
+            glMultiDrawElementsIndirect(primitive_mode, format, gl_offset,
+                                        static_cast<GLsizei>(params.max_draw_counts),
+                                        static_cast<GLsizei>(params.stride));
+        } else {
+            glMultiDrawArraysIndirect(primitive_mode, gl_offset,
+                                      static_cast<GLsizei>(params.max_draw_counts),
+                                      static_cast<GLsizei>(params.stride));
+        }
+    });
+    buffer_cache.SetDrawIndirect(nullptr);
+}
+
 void RasterizerOpenGL::DispatchCompute() {
     ComputePipeline* const pipeline{shader_cache.CurrentComputePipeline()};
     if (!pipeline) {
@@ -302,46 +352,60 @@ void RasterizerOpenGL::DisableGraphicsUniformBuffer(size_t stage, u32 index) {
 
 void RasterizerOpenGL::FlushAll() {}
 
-void RasterizerOpenGL::FlushRegion(VAddr addr, u64 size) {
+void RasterizerOpenGL::FlushRegion(VAddr addr, u64 size, VideoCommon::CacheType which) {
     MICROPROFILE_SCOPE(OpenGL_CacheManagement);
     if (addr == 0 || size == 0) {
         return;
     }
-    {
+    if (True(which & VideoCommon::CacheType::TextureCache)) {
         std::scoped_lock lock{texture_cache.mutex};
         texture_cache.DownloadMemory(addr, size);
     }
-    {
+    if ((True(which & VideoCommon::CacheType::BufferCache))) {
         std::scoped_lock lock{buffer_cache.mutex};
         buffer_cache.DownloadMemory(addr, size);
     }
-    query_cache.FlushRegion(addr, size);
+    if ((True(which & VideoCommon::CacheType::QueryCache))) {
+        query_cache.FlushRegion(addr, size);
+    }
 }
 
-bool RasterizerOpenGL::MustFlushRegion(VAddr addr, u64 size) {
-    std::scoped_lock lock{buffer_cache.mutex, texture_cache.mutex};
+bool RasterizerOpenGL::MustFlushRegion(VAddr addr, u64 size, VideoCommon::CacheType which) {
+    if ((True(which & VideoCommon::CacheType::BufferCache))) {
+        std::scoped_lock lock{buffer_cache.mutex};
+        if (buffer_cache.IsRegionGpuModified(addr, size)) {
+            return true;
+        }
+    }
     if (!Settings::IsGPULevelHigh()) {
-        return buffer_cache.IsRegionGpuModified(addr, size);
+        return false;
     }
-    return texture_cache.IsRegionGpuModified(addr, size) ||
-           buffer_cache.IsRegionGpuModified(addr, size);
+    if (True(which & VideoCommon::CacheType::TextureCache)) {
+        std::scoped_lock lock{texture_cache.mutex};
+        return texture_cache.IsRegionGpuModified(addr, size);
+    }
+    return false;
 }
 
-void RasterizerOpenGL::InvalidateRegion(VAddr addr, u64 size) {
+void RasterizerOpenGL::InvalidateRegion(VAddr addr, u64 size, VideoCommon::CacheType which) {
     MICROPROFILE_SCOPE(OpenGL_CacheManagement);
     if (addr == 0 || size == 0) {
         return;
     }
-    {
+    if (True(which & VideoCommon::CacheType::TextureCache)) {
         std::scoped_lock lock{texture_cache.mutex};
         texture_cache.WriteMemory(addr, size);
     }
-    {
+    if (True(which & VideoCommon::CacheType::BufferCache)) {
         std::scoped_lock lock{buffer_cache.mutex};
         buffer_cache.WriteMemory(addr, size);
     }
-    shader_cache.InvalidateRegion(addr, size);
-    query_cache.InvalidateRegion(addr, size);
+    if (True(which & VideoCommon::CacheType::ShaderCache)) {
+        shader_cache.InvalidateRegion(addr, size);
+    }
+    if (True(which & VideoCommon::CacheType::QueryCache)) {
+        query_cache.InvalidateRegion(addr, size);
+    }
 }
 
 void RasterizerOpenGL::OnCPUWrite(VAddr addr, u64 size) {
@@ -408,11 +472,12 @@ void RasterizerOpenGL::ReleaseFences() {
     fence_manager.WaitPendingFences();
 }
 
-void RasterizerOpenGL::FlushAndInvalidateRegion(VAddr addr, u64 size) {
+void RasterizerOpenGL::FlushAndInvalidateRegion(VAddr addr, u64 size,
+                                                VideoCommon::CacheType which) {
     if (Settings::IsGPULevelExtreme()) {
-        FlushRegion(addr, size);
+        FlushRegion(addr, size, which);
     }
-    InvalidateRegion(addr, size);
+    InvalidateRegion(addr, size, which);
 }
 
 void RasterizerOpenGL::WaitForIdle() {
@@ -460,6 +525,21 @@ void RasterizerOpenGL::TickFrame() {
     }
 }
 
+bool RasterizerOpenGL::AccelerateConditionalRendering() {
+    if (Settings::IsGPULevelHigh()) {
+        // Reimplement Host conditional rendering.
+        return false;
+    }
+    // Medium / Low Hack: stub any checks on queries writen into the buffer cache.
+    const GPUVAddr condition_address{maxwell3d->regs.render_enable.Address()};
+    Maxwell::ReportSemaphore::Compare cmp;
+    if (gpu_memory->IsMemoryDirty(condition_address, sizeof(cmp),
+                                  VideoCommon::CacheType::BufferCache)) {
+        return true;
+    }
+    return false;
+}
+
 bool RasterizerOpenGL::AccelerateSurfaceCopy(const Tegra::Engines::Fermi2D::Surface& src,
                                              const Tegra::Engines::Fermi2D::Surface& dst,
                                              const Tegra::Engines::Fermi2D::Config& copy_config) {
@@ -481,7 +561,7 @@ void RasterizerOpenGL::AccelerateInlineToMemory(GPUVAddr address, size_t copy_si
     }
     gpu_memory->WriteBlockUnsafe(address, memory.data(), copy_size);
     {
-        std::unique_lock<std::mutex> lock{buffer_cache.mutex};
+        std::unique_lock<std::recursive_mutex> lock{buffer_cache.mutex};
         if (!buffer_cache.InlineMemory(*cpu_addr, copy_size, memory)) {
             buffer_cache.WriteMemory(*cpu_addr, copy_size);
         }
diff --git a/src/video_core/renderer_opengl/gl_rasterizer.h b/src/video_core/renderer_opengl/gl_rasterizer.h
index fc183c3ca..be4f76c18 100644
--- a/src/video_core/renderer_opengl/gl_rasterizer.h
+++ b/src/video_core/renderer_opengl/gl_rasterizer.h
@@ -69,6 +69,7 @@ public:
     ~RasterizerOpenGL() override;
 
     void Draw(bool is_indexed, u32 instance_count) override;
+    void DrawIndirect() override;
     void Clear(u32 layer_count) override;
     void DispatchCompute() override;
     void ResetCounter(VideoCore::QueryType type) override;
@@ -76,9 +77,12 @@ public:
     void BindGraphicsUniformBuffer(size_t stage, u32 index, GPUVAddr gpu_addr, u32 size) override;
     void DisableGraphicsUniformBuffer(size_t stage, u32 index) override;
     void FlushAll() override;
-    void FlushRegion(VAddr addr, u64 size) override;
-    bool MustFlushRegion(VAddr addr, u64 size) override;
-    void InvalidateRegion(VAddr addr, u64 size) override;
+    void FlushRegion(VAddr addr, u64 size,
+                     VideoCommon::CacheType which = VideoCommon::CacheType::All) override;
+    bool MustFlushRegion(VAddr addr, u64 size,
+                         VideoCommon::CacheType which = VideoCommon::CacheType::All) override;
+    void InvalidateRegion(VAddr addr, u64 size,
+                          VideoCommon::CacheType which = VideoCommon::CacheType::All) override;
     void OnCPUWrite(VAddr addr, u64 size) override;
     void InvalidateGPUCache() override;
     void UnmapMemory(VAddr addr, u64 size) override;
@@ -88,12 +92,14 @@ public:
     void SignalSyncPoint(u32 value) override;
     void SignalReference() override;
     void ReleaseFences() override;
-    void FlushAndInvalidateRegion(VAddr addr, u64 size) override;
+    void FlushAndInvalidateRegion(
+        VAddr addr, u64 size, VideoCommon::CacheType which = VideoCommon::CacheType::All) override;
     void WaitForIdle() override;
     void FragmentBarrier() override;
     void TiledCacheBarrier() override;
     void FlushCommands() override;
     void TickFrame() override;
+    bool AccelerateConditionalRendering() override;
     bool AccelerateSurfaceCopy(const Tegra::Engines::Fermi2D::Surface& src,
                                const Tegra::Engines::Fermi2D::Surface& dst,
                                const Tegra::Engines::Fermi2D::Config& copy_config) override;
@@ -121,6 +127,9 @@ private:
     static constexpr size_t MAX_IMAGES = 48;
     static constexpr size_t MAX_IMAGE_VIEWS = MAX_TEXTURES + MAX_IMAGES;
 
+    template <typename Func>
+    void PrepareDraw(bool is_indexed, Func&&);
+
     /// Syncs state to match guest's
     void SyncState();
 
diff --git a/src/video_core/renderer_opengl/gl_shader_cache.cpp b/src/video_core/renderer_opengl/gl_shader_cache.cpp
index f8868a012..03b6314ff 100644
--- a/src/video_core/renderer_opengl/gl_shader_cache.cpp
+++ b/src/video_core/renderer_opengl/gl_shader_cache.cpp
@@ -51,7 +51,7 @@ using VideoCommon::LoadPipelines;
 using VideoCommon::SerializePipeline;
 using Context = ShaderContext::Context;
 
-constexpr u32 CACHE_VERSION = 7;
+constexpr u32 CACHE_VERSION = 9;
 
 template <typename Container>
 auto MakeSpan(Container& container) {
@@ -350,6 +350,7 @@ GraphicsPipeline* ShaderCache::CurrentGraphicsPipeline() {
         regs.tessellation.params.output_primitives.Value() ==
         Maxwell::Tessellation::OutputPrimitives::Triangles_CW);
     graphics_key.xfb_enabled.Assign(regs.transform_feedback_enabled != 0 ? 1 : 0);
+    graphics_key.app_stage.Assign(maxwell3d->engine_state);
     if (graphics_key.xfb_enabled) {
         SetXfbState(graphics_key.xfb_state, regs);
     }
diff --git a/src/video_core/renderer_opengl/gl_texture_cache.h b/src/video_core/renderer_opengl/gl_texture_cache.h
index 113528e9b..5d9d370f2 100644
--- a/src/video_core/renderer_opengl/gl_texture_cache.h
+++ b/src/video_core/renderer_opengl/gl_texture_cache.h
@@ -354,6 +354,7 @@ struct TextureCacheParams {
     static constexpr bool FRAMEBUFFER_BLITS = true;
     static constexpr bool HAS_EMULATED_COPIES = true;
     static constexpr bool HAS_DEVICE_MEMORY_INFO = true;
+    static constexpr bool IMPLEMENTS_ASYNC_DOWNLOADS = false;
 
     using Runtime = OpenGL::TextureCacheRuntime;
     using Image = OpenGL::Image;
@@ -361,6 +362,7 @@ struct TextureCacheParams {
     using ImageView = OpenGL::ImageView;
     using Sampler = OpenGL::Sampler;
     using Framebuffer = OpenGL::Framebuffer;
+    using AsyncBuffer = u32;
 };
 
 using TextureCache = VideoCommon::TextureCache<TextureCacheParams>;
diff --git a/src/video_core/renderer_vulkan/fixed_pipeline_state.cpp b/src/video_core/renderer_vulkan/fixed_pipeline_state.cpp
index e62b36822..3d328a250 100644
--- a/src/video_core/renderer_vulkan/fixed_pipeline_state.cpp
+++ b/src/video_core/renderer_vulkan/fixed_pipeline_state.cpp
@@ -48,43 +48,30 @@ void RefreshXfbState(VideoCommon::TransformFeedbackState& state, const Maxwell&
 }
 } // Anonymous namespace
 
-void FixedPipelineState::Refresh(Tegra::Engines::Maxwell3D& maxwell3d,
-                                 bool has_extended_dynamic_state, bool has_dynamic_vertex_input) {
+void FixedPipelineState::Refresh(Tegra::Engines::Maxwell3D& maxwell3d, DynamicFeatures& features) {
     const Maxwell& regs = maxwell3d.regs;
     const auto topology_ = maxwell3d.draw_manager->GetDrawState().topology;
-    const std::array enabled_lut{
-        regs.polygon_offset_point_enable,
-        regs.polygon_offset_line_enable,
-        regs.polygon_offset_fill_enable,
-    };
-    const u32 topology_index = static_cast<u32>(topology_);
 
     raw1 = 0;
-    extended_dynamic_state.Assign(has_extended_dynamic_state ? 1 : 0);
-    dynamic_vertex_input.Assign(has_dynamic_vertex_input ? 1 : 0);
+    extended_dynamic_state.Assign(features.has_extended_dynamic_state ? 1 : 0);
+    extended_dynamic_state_2.Assign(features.has_extended_dynamic_state_2 ? 1 : 0);
+    extended_dynamic_state_2_extra.Assign(features.has_extended_dynamic_state_2_extra ? 1 : 0);
+    extended_dynamic_state_3_blend.Assign(features.has_extended_dynamic_state_3_blend ? 1 : 0);
+    extended_dynamic_state_3_enables.Assign(features.has_extended_dynamic_state_3_enables ? 1 : 0);
+    dynamic_vertex_input.Assign(features.has_dynamic_vertex_input ? 1 : 0);
     xfb_enabled.Assign(regs.transform_feedback_enabled != 0);
-    primitive_restart_enable.Assign(regs.primitive_restart.enabled != 0 ? 1 : 0);
-    depth_bias_enable.Assign(enabled_lut[POLYGON_OFFSET_ENABLE_LUT[topology_index]] != 0 ? 1 : 0);
-    depth_clamp_disabled.Assign(regs.viewport_clip_control.geometry_clip ==
-                                    Maxwell::ViewportClipControl::GeometryClip::Passthrough ||
-                                regs.viewport_clip_control.geometry_clip ==
-                                    Maxwell::ViewportClipControl::GeometryClip::FrustumXYZ ||
-                                regs.viewport_clip_control.geometry_clip ==
-                                    Maxwell::ViewportClipControl::GeometryClip::FrustumZ);
     ndc_minus_one_to_one.Assign(regs.depth_mode == Maxwell::DepthMode::MinusOneToOne ? 1 : 0);
     polygon_mode.Assign(PackPolygonMode(regs.polygon_mode_front));
-    patch_control_points_minus_one.Assign(regs.patch_vertices - 1);
     tessellation_primitive.Assign(static_cast<u32>(regs.tessellation.params.domain_type.Value()));
     tessellation_spacing.Assign(static_cast<u32>(regs.tessellation.params.spacing.Value()));
     tessellation_clockwise.Assign(regs.tessellation.params.output_primitives.Value() ==
                                   Maxwell::Tessellation::OutputPrimitives::Triangles_CW);
-    logic_op_enable.Assign(regs.logic_op.enable != 0 ? 1 : 0);
-    logic_op.Assign(PackLogicOp(regs.logic_op.op));
+    patch_control_points_minus_one.Assign(regs.patch_vertices - 1);
     topology.Assign(topology_);
     msaa_mode.Assign(regs.anti_alias_samples_mode);
 
     raw2 = 0;
-    rasterize_enable.Assign(regs.rasterize_enable != 0 ? 1 : 0);
+
     const auto test_func =
         regs.alpha_test_enabled != 0 ? regs.alpha_test_func : Maxwell::ComparisonOp::Always_GL;
     alpha_test_func.Assign(PackComparisonOp(test_func));
@@ -97,6 +84,7 @@ void FixedPipelineState::Refresh(Tegra::Engines::Maxwell3D& maxwell3d,
     smooth_lines.Assign(regs.line_anti_alias_enable != 0 ? 1 : 0);
     alpha_to_coverage_enabled.Assign(regs.anti_alias_alpha_control.alpha_to_coverage != 0 ? 1 : 0);
     alpha_to_one_enabled.Assign(regs.anti_alias_alpha_control.alpha_to_one != 0 ? 1 : 0);
+    app_stage.Assign(maxwell3d.engine_state);
 
     for (size_t i = 0; i < regs.rt.size(); ++i) {
         color_formats[i] = static_cast<u8>(regs.rt[i].format);
@@ -105,7 +93,7 @@ void FixedPipelineState::Refresh(Tegra::Engines::Maxwell3D& maxwell3d,
     point_size = Common::BitCast<u32>(regs.point_size);
 
     if (maxwell3d.dirty.flags[Dirty::VertexInput]) {
-        if (has_dynamic_vertex_input) {
+        if (features.has_dynamic_vertex_input) {
             // Dirty flag will be reset by the command buffer update
             static constexpr std::array LUT{
                 0u, // Invalid
@@ -144,12 +132,6 @@ void FixedPipelineState::Refresh(Tegra::Engines::Maxwell3D& maxwell3d,
             }
         }
     }
-    if (maxwell3d.dirty.flags[Dirty::Blending]) {
-        maxwell3d.dirty.flags[Dirty::Blending] = false;
-        for (size_t index = 0; index < attachments.size(); ++index) {
-            attachments[index].Refresh(regs, index);
-        }
-    }
     if (maxwell3d.dirty.flags[Dirty::ViewportSwizzles]) {
         maxwell3d.dirty.flags[Dirty::ViewportSwizzles] = false;
         const auto& transform = regs.viewport_transform;
@@ -157,8 +139,27 @@ void FixedPipelineState::Refresh(Tegra::Engines::Maxwell3D& maxwell3d,
             return static_cast<u16>(viewport.swizzle.raw);
         });
     }
+    dynamic_state.raw1 = 0;
+    dynamic_state.raw2 = 0;
     if (!extended_dynamic_state) {
         dynamic_state.Refresh(regs);
+        std::ranges::transform(regs.vertex_streams, vertex_strides.begin(), [](const auto& array) {
+            return static_cast<u16>(array.stride.Value());
+        });
+    }
+    if (!extended_dynamic_state_2_extra) {
+        dynamic_state.Refresh2(regs, topology, extended_dynamic_state_2);
+    }
+    if (!extended_dynamic_state_3_blend) {
+        if (maxwell3d.dirty.flags[Dirty::Blending]) {
+            maxwell3d.dirty.flags[Dirty::Blending] = false;
+            for (size_t index = 0; index < attachments.size(); ++index) {
+                attachments[index].Refresh(regs, index);
+            }
+        }
+    }
+    if (!extended_dynamic_state_3_enables) {
+        dynamic_state.Refresh3(regs);
     }
     if (xfb_enabled) {
         RefreshXfbState(xfb_state, regs);
@@ -175,12 +176,11 @@ void FixedPipelineState::BlendingAttachment::Refresh(const Maxwell& regs, size_t
     mask_a.Assign(mask.A);
 
     // TODO: C++20 Use templated lambda to deduplicate code
+    if (!regs.blend.enable[index]) {
+        return;
+    }
 
-    if (!regs.blend_per_target_enabled) {
-        if (!regs.blend.enable[index]) {
-            return;
-        }
-        const auto& src = regs.blend;
+    const auto setup_blend = [&]<typename T>(const T& src) {
         equation_rgb.Assign(PackBlendEquation(src.color_op));
         equation_a.Assign(PackBlendEquation(src.alpha_op));
         factor_source_rgb.Assign(PackBlendFactor(src.color_source));
@@ -188,20 +188,13 @@ void FixedPipelineState::BlendingAttachment::Refresh(const Maxwell& regs, size_t
         factor_source_a.Assign(PackBlendFactor(src.alpha_source));
         factor_dest_a.Assign(PackBlendFactor(src.alpha_dest));
         enable.Assign(1);
-        return;
-    }
+    };
 
-    if (!regs.blend.enable[index]) {
+    if (!regs.blend_per_target_enabled) {
+        setup_blend(regs.blend);
         return;
     }
-    const auto& src = regs.blend_per_target[index];
-    equation_rgb.Assign(PackBlendEquation(src.color_op));
-    equation_a.Assign(PackBlendEquation(src.alpha_op));
-    factor_source_rgb.Assign(PackBlendFactor(src.color_source));
-    factor_dest_rgb.Assign(PackBlendFactor(src.color_dest));
-    factor_source_a.Assign(PackBlendFactor(src.alpha_source));
-    factor_dest_a.Assign(PackBlendFactor(src.alpha_dest));
-    enable.Assign(1);
+    setup_blend(regs.blend_per_target[index]);
 }
 
 void FixedPipelineState::DynamicState::Refresh(const Maxwell& regs) {
@@ -211,8 +204,6 @@ void FixedPipelineState::DynamicState::Refresh(const Maxwell& regs) {
         packed_front_face = 1 - packed_front_face;
     }
 
-    raw1 = 0;
-    raw2 = 0;
     front.action_stencil_fail.Assign(PackStencilOp(regs.stencil_front_op.fail));
     front.action_depth_fail.Assign(PackStencilOp(regs.stencil_front_op.zfail));
     front.action_depth_pass.Assign(PackStencilOp(regs.stencil_front_op.zpass));
@@ -236,9 +227,37 @@ void FixedPipelineState::DynamicState::Refresh(const Maxwell& regs) {
     depth_test_func.Assign(PackComparisonOp(regs.depth_test_func));
     cull_face.Assign(PackCullFace(regs.gl_cull_face));
     cull_enable.Assign(regs.gl_cull_test_enabled != 0 ? 1 : 0);
-    std::ranges::transform(regs.vertex_streams, vertex_strides.begin(), [](const auto& array) {
-        return static_cast<u16>(array.stride.Value());
-    });
+}
+
+void FixedPipelineState::DynamicState::Refresh2(const Maxwell& regs,
+                                                Maxwell::PrimitiveTopology topology_,
+                                                bool base_feautures_supported) {
+    logic_op.Assign(PackLogicOp(regs.logic_op.op));
+
+    if (base_feautures_supported) {
+        return;
+    }
+
+    const std::array enabled_lut{
+        regs.polygon_offset_point_enable,
+        regs.polygon_offset_line_enable,
+        regs.polygon_offset_fill_enable,
+    };
+    const u32 topology_index = static_cast<u32>(topology_);
+
+    rasterize_enable.Assign(regs.rasterize_enable != 0 ? 1 : 0);
+    primitive_restart_enable.Assign(regs.primitive_restart.enabled != 0 ? 1 : 0);
+    depth_bias_enable.Assign(enabled_lut[POLYGON_OFFSET_ENABLE_LUT[topology_index]] != 0 ? 1 : 0);
+}
+
+void FixedPipelineState::DynamicState::Refresh3(const Maxwell& regs) {
+    logic_op_enable.Assign(regs.logic_op.enable != 0 ? 1 : 0);
+    depth_clamp_disabled.Assign(regs.viewport_clip_control.geometry_clip ==
+                                    Maxwell::ViewportClipControl::GeometryClip::Passthrough ||
+                                regs.viewport_clip_control.geometry_clip ==
+                                    Maxwell::ViewportClipControl::GeometryClip::FrustumXYZ ||
+                                regs.viewport_clip_control.geometry_clip ==
+                                    Maxwell::ViewportClipControl::GeometryClip::FrustumZ);
 }
 
 size_t FixedPipelineState::Hash() const noexcept {
diff --git a/src/video_core/renderer_vulkan/fixed_pipeline_state.h b/src/video_core/renderer_vulkan/fixed_pipeline_state.h
index ab79fb8f3..98ea20b42 100644
--- a/src/video_core/renderer_vulkan/fixed_pipeline_state.h
+++ b/src/video_core/renderer_vulkan/fixed_pipeline_state.h
@@ -17,6 +17,15 @@ namespace Vulkan {
 
 using Maxwell = Tegra::Engines::Maxwell3D::Regs;
 
+struct DynamicFeatures {
+    bool has_extended_dynamic_state;
+    bool has_extended_dynamic_state_2;
+    bool has_extended_dynamic_state_2_extra;
+    bool has_extended_dynamic_state_3_blend;
+    bool has_extended_dynamic_state_3_enables;
+    bool has_dynamic_vertex_input;
+};
+
 struct FixedPipelineState {
     static u32 PackComparisonOp(Maxwell::ComparisonOp op) noexcept;
     static Maxwell::ComparisonOp UnpackComparisonOp(u32 packed) noexcept;
@@ -133,6 +142,17 @@ struct FixedPipelineState {
     struct DynamicState {
         union {
             u32 raw1;
+            BitField<0, 2, u32> cull_face;
+            BitField<2, 1, u32> cull_enable;
+            BitField<3, 1, u32> primitive_restart_enable;
+            BitField<4, 1, u32> depth_bias_enable;
+            BitField<5, 1, u32> rasterize_enable;
+            BitField<6, 4, u32> logic_op;
+            BitField<10, 1, u32> logic_op_enable;
+            BitField<11, 1, u32> depth_clamp_disabled;
+        };
+        union {
+            u32 raw2;
             StencilFace<0> front;
             StencilFace<12> back;
             BitField<24, 1, u32> stencil_enable;
@@ -142,15 +162,11 @@ struct FixedPipelineState {
             BitField<28, 1, u32> front_face;
             BitField<29, 3, u32> depth_test_func;
         };
-        union {
-            u32 raw2;
-            BitField<0, 2, u32> cull_face;
-            BitField<2, 1, u32> cull_enable;
-        };
-        // Vertex stride is a 12 bits value, we have 4 bits to spare per element
-        std::array<u16, Maxwell::NumVertexArrays> vertex_strides;
 
         void Refresh(const Maxwell& regs);
+        void Refresh2(const Maxwell& regs, Maxwell::PrimitiveTopology topology,
+                      bool base_feautures_supported);
+        void Refresh3(const Maxwell& regs);
 
         Maxwell::ComparisonOp DepthTestFunc() const noexcept {
             return UnpackComparisonOp(depth_test_func);
@@ -168,25 +184,24 @@ struct FixedPipelineState {
     union {
         u32 raw1;
         BitField<0, 1, u32> extended_dynamic_state;
-        BitField<1, 1, u32> dynamic_vertex_input;
-        BitField<2, 1, u32> xfb_enabled;
-        BitField<3, 1, u32> primitive_restart_enable;
-        BitField<4, 1, u32> depth_bias_enable;
-        BitField<5, 1, u32> depth_clamp_disabled;
-        BitField<6, 1, u32> ndc_minus_one_to_one;
-        BitField<7, 2, u32> polygon_mode;
-        BitField<9, 5, u32> patch_control_points_minus_one;
-        BitField<14, 2, u32> tessellation_primitive;
-        BitField<16, 2, u32> tessellation_spacing;
-        BitField<18, 1, u32> tessellation_clockwise;
-        BitField<19, 1, u32> logic_op_enable;
-        BitField<20, 4, u32> logic_op;
+        BitField<1, 1, u32> extended_dynamic_state_2;
+        BitField<2, 1, u32> extended_dynamic_state_2_extra;
+        BitField<3, 1, u32> extended_dynamic_state_3_blend;
+        BitField<4, 1, u32> extended_dynamic_state_3_enables;
+        BitField<5, 1, u32> dynamic_vertex_input;
+        BitField<6, 1, u32> xfb_enabled;
+        BitField<7, 1, u32> ndc_minus_one_to_one;
+        BitField<8, 2, u32> polygon_mode;
+        BitField<10, 2, u32> tessellation_primitive;
+        BitField<12, 2, u32> tessellation_spacing;
+        BitField<14, 1, u32> tessellation_clockwise;
+        BitField<15, 5, u32> patch_control_points_minus_one;
+
         BitField<24, 4, Maxwell::PrimitiveTopology> topology;
         BitField<28, 4, Tegra::Texture::MsaaMode> msaa_mode;
     };
     union {
         u32 raw2;
-        BitField<0, 1, u32> rasterize_enable;
         BitField<1, 3, u32> alpha_test_func;
         BitField<4, 1, u32> early_z;
         BitField<5, 1, u32> depth_enabled;
@@ -197,25 +212,28 @@ struct FixedPipelineState {
         BitField<14, 1, u32> smooth_lines;
         BitField<15, 1, u32> alpha_to_coverage_enabled;
         BitField<16, 1, u32> alpha_to_one_enabled;
+        BitField<17, 3, Tegra::Engines::Maxwell3D::EngineHint> app_stage;
     };
     std::array<u8, Maxwell::NumRenderTargets> color_formats;
 
     u32 alpha_test_ref;
     u32 point_size;
-    std::array<BlendingAttachment, Maxwell::NumRenderTargets> attachments;
     std::array<u16, Maxwell::NumViewports> viewport_swizzles;
     union {
         u64 attribute_types; // Used with VK_EXT_vertex_input_dynamic_state
         u64 enabled_divisors;
     };
+
+    DynamicState dynamic_state;
+    std::array<BlendingAttachment, Maxwell::NumRenderTargets> attachments;
     std::array<VertexAttribute, Maxwell::NumVertexAttributes> attributes;
     std::array<u32, Maxwell::NumVertexArrays> binding_divisors;
+    // Vertex stride is a 12 bits value, we have 4 bits to spare per element
+    std::array<u16, Maxwell::NumVertexArrays> vertex_strides;
 
-    DynamicState dynamic_state;
     VideoCommon::TransformFeedbackState xfb_state;
 
-    void Refresh(Tegra::Engines::Maxwell3D& maxwell3d, bool has_extended_dynamic_state,
-                 bool has_dynamic_vertex_input);
+    void Refresh(Tegra::Engines::Maxwell3D& maxwell3d, DynamicFeatures& features);
 
     size_t Hash() const noexcept;
 
@@ -230,13 +248,17 @@ struct FixedPipelineState {
             // When transform feedback is enabled, use the whole struct
             return sizeof(*this);
         }
-        if (dynamic_vertex_input) {
+        if (dynamic_vertex_input && extended_dynamic_state_3_blend) {
             // Exclude dynamic state and attributes
+            return offsetof(FixedPipelineState, dynamic_state);
+        }
+        if (dynamic_vertex_input) {
+            // Exclude dynamic state
             return offsetof(FixedPipelineState, attributes);
         }
         if (extended_dynamic_state) {
             // Exclude dynamic state
-            return offsetof(FixedPipelineState, dynamic_state);
+            return offsetof(FixedPipelineState, vertex_strides);
         }
         // Default
         return offsetof(FixedPipelineState, xfb_state);
diff --git a/src/video_core/renderer_vulkan/vk_buffer_cache.cpp b/src/video_core/renderer_vulkan/vk_buffer_cache.cpp
index 6b54d7111..487d8b416 100644
--- a/src/video_core/renderer_vulkan/vk_buffer_cache.cpp
+++ b/src/video_core/renderer_vulkan/vk_buffer_cache.cpp
@@ -56,7 +56,8 @@ vk::Buffer CreateBuffer(const Device& device, u64 size) {
         VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT |
         VK_BUFFER_USAGE_UNIFORM_TEXEL_BUFFER_BIT | VK_BUFFER_USAGE_STORAGE_TEXEL_BUFFER_BIT |
         VK_BUFFER_USAGE_UNIFORM_BUFFER_BIT | VK_BUFFER_USAGE_STORAGE_BUFFER_BIT |
-        VK_BUFFER_USAGE_INDEX_BUFFER_BIT | VK_BUFFER_USAGE_VERTEX_BUFFER_BIT;
+        VK_BUFFER_USAGE_INDEX_BUFFER_BIT | VK_BUFFER_USAGE_VERTEX_BUFFER_BIT |
+        VK_BUFFER_USAGE_INDIRECT_BUFFER_BIT;
     if (device.IsExtTransformFeedbackSupported()) {
         flags |= VK_BUFFER_USAGE_TRANSFORM_FEEDBACK_BUFFER_BIT_EXT;
     }
@@ -516,6 +517,7 @@ void BufferCacheRuntime::ReserveNullBuffer() {
     if (device.IsExtTransformFeedbackSupported()) {
         create_info.usage |= VK_BUFFER_USAGE_TRANSFORM_FEEDBACK_BUFFER_BIT_EXT;
     }
+    create_info.usage |= VK_BUFFER_USAGE_INDIRECT_BUFFER_BIT;
     null_buffer = device.GetLogical().CreateBuffer(create_info);
     if (device.HasDebuggingToolAttached()) {
         null_buffer.SetObjectNameEXT("Null buffer");
diff --git a/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp b/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp
index 515d8d869..d11383bf1 100644
--- a/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp
+++ b/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp
@@ -201,6 +201,22 @@ struct SimpleVertexSpec {
     static constexpr bool has_images = false;
 };
 
+struct SimpleStorageSpec {
+    static constexpr std::array<bool, 5> enabled_stages{true, false, false, false, true};
+    static constexpr bool has_storage_buffers = true;
+    static constexpr bool has_texture_buffers = false;
+    static constexpr bool has_image_buffers = false;
+    static constexpr bool has_images = false;
+};
+
+struct SimpleImageSpec {
+    static constexpr std::array<bool, 5> enabled_stages{true, false, false, false, true};
+    static constexpr bool has_storage_buffers = false;
+    static constexpr bool has_texture_buffers = false;
+    static constexpr bool has_image_buffers = false;
+    static constexpr bool has_images = true;
+};
+
 struct DefaultSpec {
     static constexpr std::array<bool, 5> enabled_stages{true, true, true, true, true};
     static constexpr bool has_storage_buffers = true;
@@ -211,7 +227,8 @@ struct DefaultSpec {
 
 ConfigureFuncPtr ConfigureFunc(const std::array<vk::ShaderModule, NUM_STAGES>& modules,
                                const std::array<Shader::Info, NUM_STAGES>& infos) {
-    return FindSpec<SimpleVertexSpec, SimpleVertexFragmentSpec, DefaultSpec>(modules, infos);
+    return FindSpec<SimpleVertexSpec, SimpleVertexFragmentSpec, SimpleStorageSpec, SimpleImageSpec,
+                    DefaultSpec>(modules, infos);
 }
 } // Anonymous namespace
 
@@ -524,6 +541,8 @@ void GraphicsPipeline::MakePipeline(VkRenderPass render_pass) {
     FixedPipelineState::DynamicState dynamic{};
     if (!key.state.extended_dynamic_state) {
         dynamic = key.state.dynamic_state;
+    } else {
+        dynamic.raw1 = key.state.dynamic_state.raw1;
     }
     static_vector<VkVertexInputBindingDescription, 32> vertex_bindings;
     static_vector<VkVertexInputBindingDivisorDescriptionEXT, 32> vertex_binding_divisors;
@@ -561,7 +580,7 @@ void GraphicsPipeline::MakePipeline(VkRenderPass render_pass) {
                 instanced ? VK_VERTEX_INPUT_RATE_INSTANCE : VK_VERTEX_INPUT_RATE_VERTEX;
             vertex_bindings.push_back({
                 .binding = static_cast<u32>(index),
-                .stride = dynamic.vertex_strides[index],
+                .stride = key.state.vertex_strides[index],
                 .inputRate = rate,
             });
             if (instanced) {
@@ -625,7 +644,7 @@ void GraphicsPipeline::MakePipeline(VkRenderPass render_pass) {
         .pNext = nullptr,
         .flags = 0,
         .topology = input_assembly_topology,
-        .primitiveRestartEnable = key.state.primitive_restart_enable != 0 &&
+        .primitiveRestartEnable = dynamic.primitive_restart_enable != 0 &&
                                   ((input_assembly_topology != VK_PRIMITIVE_TOPOLOGY_PATCH_LIST &&
                                     device.IsTopologyListPrimitiveRestartSupported()) ||
                                    SupportsPrimitiveRestart(input_assembly_topology) ||
@@ -672,15 +691,15 @@ void GraphicsPipeline::MakePipeline(VkRenderPass render_pass) {
         .pNext = nullptr,
         .flags = 0,
         .depthClampEnable =
-            static_cast<VkBool32>(key.state.depth_clamp_disabled == 0 ? VK_TRUE : VK_FALSE),
+            static_cast<VkBool32>(dynamic.depth_clamp_disabled == 0 ? VK_TRUE : VK_FALSE),
         .rasterizerDiscardEnable =
-            static_cast<VkBool32>(key.state.rasterize_enable == 0 ? VK_TRUE : VK_FALSE),
+            static_cast<VkBool32>(dynamic.rasterize_enable == 0 ? VK_TRUE : VK_FALSE),
         .polygonMode =
             MaxwellToVK::PolygonMode(FixedPipelineState::UnpackPolygonMode(key.state.polygon_mode)),
         .cullMode = static_cast<VkCullModeFlags>(
             dynamic.cull_enable ? MaxwellToVK::CullFace(dynamic.CullFace()) : VK_CULL_MODE_NONE),
         .frontFace = MaxwellToVK::FrontFace(dynamic.FrontFace()),
-        .depthBiasEnable = key.state.depth_bias_enable,
+        .depthBiasEnable = (dynamic.depth_bias_enable == 0 ? VK_TRUE : VK_FALSE),
         .depthBiasConstantFactor = 0.0f,
         .depthBiasClamp = 0.0f,
         .depthBiasSlopeFactor = 0.0f,
@@ -782,13 +801,13 @@ void GraphicsPipeline::MakePipeline(VkRenderPass render_pass) {
         .sType = VK_STRUCTURE_TYPE_PIPELINE_COLOR_BLEND_STATE_CREATE_INFO,
         .pNext = nullptr,
         .flags = 0,
-        .logicOpEnable = key.state.logic_op_enable != 0,
-        .logicOp = static_cast<VkLogicOp>(key.state.logic_op.Value()),
+        .logicOpEnable = dynamic.logic_op_enable != 0,
+        .logicOp = static_cast<VkLogicOp>(dynamic.logic_op.Value()),
         .attachmentCount = static_cast<u32>(cb_attachments.size()),
         .pAttachments = cb_attachments.data(),
         .blendConstants = {},
     };
-    static_vector<VkDynamicState, 19> dynamic_states{
+    static_vector<VkDynamicState, 28> dynamic_states{
         VK_DYNAMIC_STATE_VIEWPORT,           VK_DYNAMIC_STATE_SCISSOR,
         VK_DYNAMIC_STATE_DEPTH_BIAS,         VK_DYNAMIC_STATE_BLEND_CONSTANTS,
         VK_DYNAMIC_STATE_DEPTH_BOUNDS,       VK_DYNAMIC_STATE_STENCIL_COMPARE_MASK,
@@ -811,6 +830,32 @@ void GraphicsPipeline::MakePipeline(VkRenderPass render_pass) {
             dynamic_states.push_back(VK_DYNAMIC_STATE_VERTEX_INPUT_EXT);
         }
         dynamic_states.insert(dynamic_states.end(), extended.begin(), extended.end());
+        if (key.state.extended_dynamic_state_2) {
+            static constexpr std::array extended2{
+                VK_DYNAMIC_STATE_DEPTH_BIAS_ENABLE_EXT,
+                VK_DYNAMIC_STATE_PRIMITIVE_RESTART_ENABLE_EXT,
+                VK_DYNAMIC_STATE_RASTERIZER_DISCARD_ENABLE_EXT,
+            };
+            dynamic_states.insert(dynamic_states.end(), extended2.begin(), extended2.end());
+        }
+        if (key.state.extended_dynamic_state_2_extra) {
+            dynamic_states.push_back(VK_DYNAMIC_STATE_LOGIC_OP_EXT);
+        }
+        if (key.state.extended_dynamic_state_3_blend) {
+            static constexpr std::array extended3{
+                VK_DYNAMIC_STATE_COLOR_BLEND_ENABLE_EXT,
+                VK_DYNAMIC_STATE_COLOR_BLEND_EQUATION_EXT,
+                VK_DYNAMIC_STATE_COLOR_WRITE_MASK_EXT,
+            };
+            dynamic_states.insert(dynamic_states.end(), extended3.begin(), extended3.end());
+        }
+        if (key.state.extended_dynamic_state_3_enables) {
+            static constexpr std::array extended3{
+                VK_DYNAMIC_STATE_DEPTH_CLAMP_ENABLE_EXT,
+                VK_DYNAMIC_STATE_LOGIC_OP_ENABLE_EXT,
+            };
+            dynamic_states.insert(dynamic_states.end(), extended3.begin(), extended3.end());
+        }
     }
     const VkPipelineDynamicStateCreateInfo dynamic_state_ci{
         .sType = VK_STRUCTURE_TYPE_PIPELINE_DYNAMIC_STATE_CREATE_INFO,
diff --git a/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp b/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp
index e7262420c..3046b72ab 100644
--- a/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp
+++ b/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp
@@ -54,7 +54,7 @@ using VideoCommon::FileEnvironment;
 using VideoCommon::GenericEnvironment;
 using VideoCommon::GraphicsEnvironment;
 
-constexpr u32 CACHE_VERSION = 8;
+constexpr u32 CACHE_VERSION = 10;
 
 template <typename Container>
 auto MakeSpan(Container& container) {
@@ -351,6 +351,15 @@ PipelineCache::PipelineCache(RasterizerVulkan& rasterizer_, const Device& device
         LOG_WARNING(Render_Vulkan, "maxVertexInputBindings is too low: {} < {}",
                     device.GetMaxVertexInputBindings(), Maxwell::NumVertexArrays);
     }
+
+    dynamic_features = DynamicFeatures{
+        .has_extended_dynamic_state = device.IsExtExtendedDynamicStateSupported(),
+        .has_extended_dynamic_state_2 = device.IsExtExtendedDynamicState2Supported(),
+        .has_extended_dynamic_state_2_extra = device.IsExtExtendedDynamicState2ExtrasSupported(),
+        .has_extended_dynamic_state_3_blend = device.IsExtExtendedDynamicState3BlendingSupported(),
+        .has_extended_dynamic_state_3_enables = device.IsExtExtendedDynamicState3EnablesSupported(),
+        .has_dynamic_vertex_input = device.IsExtVertexInputDynamicStateSupported(),
+    };
 }
 
 PipelineCache::~PipelineCache() = default;
@@ -362,8 +371,7 @@ GraphicsPipeline* PipelineCache::CurrentGraphicsPipeline() {
         current_pipeline = nullptr;
         return nullptr;
     }
-    graphics_key.state.Refresh(*maxwell3d, device.IsExtExtendedDynamicStateSupported(),
-                               device.IsExtVertexInputDynamicStateSupported());
+    graphics_key.state.Refresh(*maxwell3d, dynamic_features);
 
     if (current_pipeline) {
         GraphicsPipeline* const next{current_pipeline->Next(graphics_key)};
@@ -439,14 +447,21 @@ void PipelineCache::LoadDiskResources(u64 title_id, std::stop_token stop_loading
         });
         ++state.total;
     }};
-    const bool extended_dynamic_state = device.IsExtExtendedDynamicStateSupported();
-    const bool dynamic_vertex_input = device.IsExtVertexInputDynamicStateSupported();
     const auto load_graphics{[&](std::ifstream& file, std::vector<FileEnvironment> envs) {
         GraphicsPipelineCacheKey key;
         file.read(reinterpret_cast<char*>(&key), sizeof(key));
 
-        if ((key.state.extended_dynamic_state != 0) != extended_dynamic_state ||
-            (key.state.dynamic_vertex_input != 0) != dynamic_vertex_input) {
+        if ((key.state.extended_dynamic_state != 0) !=
+                dynamic_features.has_extended_dynamic_state ||
+            (key.state.extended_dynamic_state_2 != 0) !=
+                dynamic_features.has_extended_dynamic_state_2 ||
+            (key.state.extended_dynamic_state_2_extra != 0) !=
+                dynamic_features.has_extended_dynamic_state_2_extra ||
+            (key.state.extended_dynamic_state_3_blend != 0) !=
+                dynamic_features.has_extended_dynamic_state_3_blend ||
+            (key.state.extended_dynamic_state_3_enables != 0) !=
+                dynamic_features.has_extended_dynamic_state_3_enables ||
+            (key.state.dynamic_vertex_input != 0) != dynamic_features.has_dynamic_vertex_input) {
             return;
         }
         workers.QueueWork([this, key, envs = std::move(envs), &state, &callback]() mutable {
diff --git a/src/video_core/renderer_vulkan/vk_pipeline_cache.h b/src/video_core/renderer_vulkan/vk_pipeline_cache.h
index 61f9e9366..b4f593ef5 100644
--- a/src/video_core/renderer_vulkan/vk_pipeline_cache.h
+++ b/src/video_core/renderer_vulkan/vk_pipeline_cache.h
@@ -160,6 +160,7 @@ private:
 
     Common::ThreadWorker workers;
     Common::ThreadWorker serialization_thread;
+    DynamicFeatures dynamic_features;
 };
 
 } // namespace Vulkan
diff --git a/src/video_core/renderer_vulkan/vk_rasterizer.cpp b/src/video_core/renderer_vulkan/vk_rasterizer.cpp
index ac1eb9895..242bf9602 100644
--- a/src/video_core/renderer_vulkan/vk_rasterizer.cpp
+++ b/src/video_core/renderer_vulkan/vk_rasterizer.cpp
@@ -180,7 +180,8 @@ RasterizerVulkan::RasterizerVulkan(Core::Frontend::EmuWindow& emu_window_, Tegra
 
 RasterizerVulkan::~RasterizerVulkan() = default;
 
-void RasterizerVulkan::Draw(bool is_indexed, u32 instance_count) {
+template <typename Func>
+void RasterizerVulkan::PrepareDraw(bool is_indexed, Func&& draw_func) {
     MICROPROFILE_SCOPE(Vulkan_Drawing);
 
     SCOPE_EXIT({ gpu.TickWork(); });
@@ -201,20 +202,67 @@ void RasterizerVulkan::Draw(bool is_indexed, u32 instance_count) {
 
     UpdateDynamicStates();
 
-    const auto& draw_state = maxwell3d->draw_manager->GetDrawState();
-    const u32 num_instances{instance_count};
-    const DrawParams draw_params{MakeDrawParams(draw_state, num_instances, is_indexed)};
-    scheduler.Record([draw_params](vk::CommandBuffer cmdbuf) {
-        if (draw_params.is_indexed) {
-            cmdbuf.DrawIndexed(draw_params.num_vertices, draw_params.num_instances,
-                               draw_params.first_index, draw_params.base_vertex,
-                               draw_params.base_instance);
-        } else {
-            cmdbuf.Draw(draw_params.num_vertices, draw_params.num_instances,
-                        draw_params.base_vertex, draw_params.base_instance);
+    draw_func();
+
+    EndTransformFeedback();
+}
+
+void RasterizerVulkan::Draw(bool is_indexed, u32 instance_count) {
+    PrepareDraw(is_indexed, [this, is_indexed, instance_count] {
+        const auto& draw_state = maxwell3d->draw_manager->GetDrawState();
+        const u32 num_instances{instance_count};
+        const DrawParams draw_params{MakeDrawParams(draw_state, num_instances, is_indexed)};
+        scheduler.Record([draw_params](vk::CommandBuffer cmdbuf) {
+            if (draw_params.is_indexed) {
+                cmdbuf.DrawIndexed(draw_params.num_vertices, draw_params.num_instances,
+                                   draw_params.first_index, draw_params.base_vertex,
+                                   draw_params.base_instance);
+            } else {
+                cmdbuf.Draw(draw_params.num_vertices, draw_params.num_instances,
+                            draw_params.base_vertex, draw_params.base_instance);
+            }
+        });
+    });
+}
+
+void RasterizerVulkan::DrawIndirect() {
+    const auto& params = maxwell3d->draw_manager->GetIndirectParams();
+    buffer_cache.SetDrawIndirect(&params);
+    PrepareDraw(params.is_indexed, [this, &params] {
+        const auto indirect_buffer = buffer_cache.GetDrawIndirectBuffer();
+        const auto& buffer = indirect_buffer.first;
+        const auto& offset = indirect_buffer.second;
+        if (params.include_count) {
+            const auto count = buffer_cache.GetDrawIndirectCount();
+            const auto& draw_buffer = count.first;
+            const auto& offset_base = count.second;
+            scheduler.Record([draw_buffer_obj = draw_buffer->Handle(),
+                              buffer_obj = buffer->Handle(), offset_base, offset,
+                              params](vk::CommandBuffer cmdbuf) {
+                if (params.is_indexed) {
+                    cmdbuf.DrawIndexedIndirectCount(
+                        buffer_obj, offset, draw_buffer_obj, offset_base,
+                        static_cast<u32>(params.max_draw_counts), static_cast<u32>(params.stride));
+                } else {
+                    cmdbuf.DrawIndirectCount(buffer_obj, offset, draw_buffer_obj, offset_base,
+                                             static_cast<u32>(params.max_draw_counts),
+                                             static_cast<u32>(params.stride));
+                }
+            });
+            return;
         }
+        scheduler.Record([buffer_obj = buffer->Handle(), offset, params](vk::CommandBuffer cmdbuf) {
+            if (params.is_indexed) {
+                cmdbuf.DrawIndexedIndirect(buffer_obj, offset,
+                                           static_cast<u32>(params.max_draw_counts),
+                                           static_cast<u32>(params.stride));
+            } else {
+                cmdbuf.DrawIndirect(buffer_obj, offset, static_cast<u32>(params.max_draw_counts),
+                                    static_cast<u32>(params.stride));
+            }
+        });
     });
-    EndTransformFeedback();
+    buffer_cache.SetDrawIndirect(nullptr);
 }
 
 void RasterizerVulkan::Clear(u32 layer_count) {
@@ -379,44 +427,58 @@ void Vulkan::RasterizerVulkan::DisableGraphicsUniformBuffer(size_t stage, u32 in
 
 void RasterizerVulkan::FlushAll() {}
 
-void RasterizerVulkan::FlushRegion(VAddr addr, u64 size) {
+void RasterizerVulkan::FlushRegion(VAddr addr, u64 size, VideoCommon::CacheType which) {
     if (addr == 0 || size == 0) {
         return;
     }
-    {
+    if (True(which & VideoCommon::CacheType::TextureCache)) {
         std::scoped_lock lock{texture_cache.mutex};
         texture_cache.DownloadMemory(addr, size);
     }
-    {
+    if ((True(which & VideoCommon::CacheType::BufferCache))) {
         std::scoped_lock lock{buffer_cache.mutex};
         buffer_cache.DownloadMemory(addr, size);
     }
-    query_cache.FlushRegion(addr, size);
+    if ((True(which & VideoCommon::CacheType::QueryCache))) {
+        query_cache.FlushRegion(addr, size);
+    }
 }
 
-bool RasterizerVulkan::MustFlushRegion(VAddr addr, u64 size) {
-    std::scoped_lock lock{texture_cache.mutex, buffer_cache.mutex};
+bool RasterizerVulkan::MustFlushRegion(VAddr addr, u64 size, VideoCommon::CacheType which) {
+    if ((True(which & VideoCommon::CacheType::BufferCache))) {
+        std::scoped_lock lock{buffer_cache.mutex};
+        if (buffer_cache.IsRegionGpuModified(addr, size)) {
+            return true;
+        }
+    }
     if (!Settings::IsGPULevelHigh()) {
-        return buffer_cache.IsRegionGpuModified(addr, size);
+        return false;
+    }
+    if (True(which & VideoCommon::CacheType::TextureCache)) {
+        std::scoped_lock lock{texture_cache.mutex};
+        return texture_cache.IsRegionGpuModified(addr, size);
     }
-    return texture_cache.IsRegionGpuModified(addr, size) ||
-           buffer_cache.IsRegionGpuModified(addr, size);
+    return false;
 }
 
-void RasterizerVulkan::InvalidateRegion(VAddr addr, u64 size) {
+void RasterizerVulkan::InvalidateRegion(VAddr addr, u64 size, VideoCommon::CacheType which) {
     if (addr == 0 || size == 0) {
         return;
     }
-    {
+    if (True(which & VideoCommon::CacheType::TextureCache)) {
         std::scoped_lock lock{texture_cache.mutex};
         texture_cache.WriteMemory(addr, size);
     }
-    {
+    if ((True(which & VideoCommon::CacheType::BufferCache))) {
         std::scoped_lock lock{buffer_cache.mutex};
         buffer_cache.WriteMemory(addr, size);
     }
-    pipeline_cache.InvalidateRegion(addr, size);
-    query_cache.InvalidateRegion(addr, size);
+    if ((True(which & VideoCommon::CacheType::QueryCache))) {
+        query_cache.InvalidateRegion(addr, size);
+    }
+    if ((True(which & VideoCommon::CacheType::ShaderCache))) {
+        pipeline_cache.InvalidateRegion(addr, size);
+    }
 }
 
 void RasterizerVulkan::OnCPUWrite(VAddr addr, u64 size) {
@@ -481,11 +543,12 @@ void RasterizerVulkan::ReleaseFences() {
     fence_manager.WaitPendingFences();
 }
 
-void RasterizerVulkan::FlushAndInvalidateRegion(VAddr addr, u64 size) {
+void RasterizerVulkan::FlushAndInvalidateRegion(VAddr addr, u64 size,
+                                                VideoCommon::CacheType which) {
     if (Settings::IsGPULevelExtreme()) {
-        FlushRegion(addr, size);
+        FlushRegion(addr, size, which);
     }
-    InvalidateRegion(addr, size);
+    InvalidateRegion(addr, size, which);
 }
 
 void RasterizerVulkan::WaitForIdle() {
@@ -541,6 +604,21 @@ void RasterizerVulkan::TickFrame() {
     }
 }
 
+bool RasterizerVulkan::AccelerateConditionalRendering() {
+    if (Settings::IsGPULevelHigh()) {
+        // TODO(Blinkhawk): Reimplement Host conditional rendering.
+        return false;
+    }
+    // Medium / Low Hack: stub any checks on queries writen into the buffer cache.
+    const GPUVAddr condition_address{maxwell3d->regs.render_enable.Address()};
+    Maxwell::ReportSemaphore::Compare cmp;
+    if (gpu_memory->IsMemoryDirty(condition_address, sizeof(cmp),
+                                  VideoCommon::CacheType::BufferCache)) {
+        return true;
+    }
+    return false;
+}
+
 bool RasterizerVulkan::AccelerateSurfaceCopy(const Tegra::Engines::Fermi2D::Surface& src,
                                              const Tegra::Engines::Fermi2D::Surface& dst,
                                              const Tegra::Engines::Fermi2D::Config& copy_config) {
@@ -561,7 +639,7 @@ void RasterizerVulkan::AccelerateInlineToMemory(GPUVAddr address, size_t copy_si
     }
     gpu_memory->WriteBlockUnsafe(address, memory.data(), copy_size);
     {
-        std::unique_lock<std::mutex> lock{buffer_cache.mutex};
+        std::unique_lock<std::recursive_mutex> lock{buffer_cache.mutex};
         if (!buffer_cache.InlineMemory(*cpu_addr, copy_size, memory)) {
             buffer_cache.WriteMemory(*cpu_addr, copy_size);
         }
@@ -639,16 +717,35 @@ void RasterizerVulkan::UpdateDynamicStates() {
     UpdateLineWidth(regs);
     if (device.IsExtExtendedDynamicStateSupported()) {
         UpdateCullMode(regs);
-        UpdateDepthBoundsTestEnable(regs);
-        UpdateDepthTestEnable(regs);
-        UpdateDepthWriteEnable(regs);
         UpdateDepthCompareOp(regs);
         UpdateFrontFace(regs);
         UpdateStencilOp(regs);
-        UpdateStencilTestEnable(regs);
+
         if (device.IsExtVertexInputDynamicStateSupported()) {
             UpdateVertexInput(regs);
         }
+
+        if (state_tracker.TouchStateEnable()) {
+            UpdateDepthBoundsTestEnable(regs);
+            UpdateDepthTestEnable(regs);
+            UpdateDepthWriteEnable(regs);
+            UpdateStencilTestEnable(regs);
+            if (device.IsExtExtendedDynamicState2Supported()) {
+                UpdatePrimitiveRestartEnable(regs);
+                UpdateRasterizerDiscardEnable(regs);
+                UpdateDepthBiasEnable(regs);
+            }
+            if (device.IsExtExtendedDynamicState3EnablesSupported()) {
+                UpdateLogicOpEnable(regs);
+                UpdateDepthClampEnable(regs);
+            }
+        }
+        if (device.IsExtExtendedDynamicState2ExtrasSupported()) {
+            UpdateLogicOp(regs);
+        }
+        if (device.IsExtExtendedDynamicState3Supported()) {
+            UpdateBlending(regs);
+        }
     }
 }
 
@@ -789,32 +886,92 @@ void RasterizerVulkan::UpdateStencilFaces(Tegra::Engines::Maxwell3D::Regs& regs)
     if (!state_tracker.TouchStencilProperties()) {
         return;
     }
-    if (regs.stencil_two_side_enable) {
-        // Separate values per face
-        scheduler.Record(
-            [front_ref = regs.stencil_front_ref, front_write_mask = regs.stencil_front_mask,
-             front_test_mask = regs.stencil_front_func_mask, back_ref = regs.stencil_back_ref,
-             back_write_mask = regs.stencil_back_mask,
-             back_test_mask = regs.stencil_back_func_mask](vk::CommandBuffer cmdbuf) {
+    bool update_references = state_tracker.TouchStencilReference();
+    bool update_write_mask = state_tracker.TouchStencilWriteMask();
+    bool update_compare_masks = state_tracker.TouchStencilCompare();
+    if (state_tracker.TouchStencilSide(regs.stencil_two_side_enable != 0)) {
+        update_references = true;
+        update_write_mask = true;
+        update_compare_masks = true;
+    }
+    if (update_references) {
+        [&]() {
+            if (regs.stencil_two_side_enable) {
+                if (!state_tracker.CheckStencilReferenceFront(regs.stencil_front_ref) &&
+                    !state_tracker.CheckStencilReferenceBack(regs.stencil_back_ref)) {
+                    return;
+                }
+            } else {
+                if (!state_tracker.CheckStencilReferenceFront(regs.stencil_front_ref)) {
+                    return;
+                }
+            }
+            scheduler.Record([front_ref = regs.stencil_front_ref, back_ref = regs.stencil_back_ref,
+                              two_sided = regs.stencil_two_side_enable](vk::CommandBuffer cmdbuf) {
+                const bool set_back = two_sided && front_ref != back_ref;
                 // Front face
-                cmdbuf.SetStencilReference(VK_STENCIL_FACE_FRONT_BIT, front_ref);
-                cmdbuf.SetStencilWriteMask(VK_STENCIL_FACE_FRONT_BIT, front_write_mask);
-                cmdbuf.SetStencilCompareMask(VK_STENCIL_FACE_FRONT_BIT, front_test_mask);
-
-                // Back face
-                cmdbuf.SetStencilReference(VK_STENCIL_FACE_BACK_BIT, back_ref);
-                cmdbuf.SetStencilWriteMask(VK_STENCIL_FACE_BACK_BIT, back_write_mask);
-                cmdbuf.SetStencilCompareMask(VK_STENCIL_FACE_BACK_BIT, back_test_mask);
+                cmdbuf.SetStencilReference(set_back ? VK_STENCIL_FACE_FRONT_BIT
+                                                    : VK_STENCIL_FACE_FRONT_AND_BACK,
+                                           front_ref);
+                if (set_back) {
+                    cmdbuf.SetStencilReference(VK_STENCIL_FACE_BACK_BIT, back_ref);
+                }
             });
-    } else {
-        // Front face defines both faces
-        scheduler.Record([ref = regs.stencil_front_ref, write_mask = regs.stencil_front_mask,
-                          test_mask = regs.stencil_front_func_mask](vk::CommandBuffer cmdbuf) {
-            cmdbuf.SetStencilReference(VK_STENCIL_FACE_FRONT_AND_BACK, ref);
-            cmdbuf.SetStencilWriteMask(VK_STENCIL_FACE_FRONT_AND_BACK, write_mask);
-            cmdbuf.SetStencilCompareMask(VK_STENCIL_FACE_FRONT_AND_BACK, test_mask);
-        });
+        }();
+    }
+    if (update_write_mask) {
+        [&]() {
+            if (regs.stencil_two_side_enable) {
+                if (!state_tracker.CheckStencilWriteMaskFront(regs.stencil_front_mask) &&
+                    !state_tracker.CheckStencilWriteMaskBack(regs.stencil_back_mask)) {
+                    return;
+                }
+            } else {
+                if (!state_tracker.CheckStencilWriteMaskFront(regs.stencil_front_mask)) {
+                    return;
+                }
+            }
+            scheduler.Record([front_write_mask = regs.stencil_front_mask,
+                              back_write_mask = regs.stencil_back_mask,
+                              two_sided = regs.stencil_two_side_enable](vk::CommandBuffer cmdbuf) {
+                const bool set_back = two_sided && front_write_mask != back_write_mask;
+                // Front face
+                cmdbuf.SetStencilWriteMask(set_back ? VK_STENCIL_FACE_FRONT_BIT
+                                                    : VK_STENCIL_FACE_FRONT_AND_BACK,
+                                           front_write_mask);
+                if (set_back) {
+                    cmdbuf.SetStencilWriteMask(VK_STENCIL_FACE_BACK_BIT, back_write_mask);
+                }
+            });
+        }();
+    }
+    if (update_compare_masks) {
+        [&]() {
+            if (regs.stencil_two_side_enable) {
+                if (!state_tracker.CheckStencilCompareMaskFront(regs.stencil_front_func_mask) &&
+                    !state_tracker.CheckStencilCompareMaskBack(regs.stencil_back_func_mask)) {
+                    return;
+                }
+            } else {
+                if (!state_tracker.CheckStencilCompareMaskFront(regs.stencil_front_func_mask)) {
+                    return;
+                }
+            }
+            scheduler.Record([front_test_mask = regs.stencil_front_func_mask,
+                              back_test_mask = regs.stencil_back_func_mask,
+                              two_sided = regs.stencil_two_side_enable](vk::CommandBuffer cmdbuf) {
+                const bool set_back = two_sided && front_test_mask != back_test_mask;
+                // Front face
+                cmdbuf.SetStencilCompareMask(set_back ? VK_STENCIL_FACE_FRONT_BIT
+                                                      : VK_STENCIL_FACE_FRONT_AND_BACK,
+                                             front_test_mask);
+                if (set_back) {
+                    cmdbuf.SetStencilCompareMask(VK_STENCIL_FACE_BACK_BIT, back_test_mask);
+                }
+            });
+        }();
     }
+    state_tracker.ClearStencilReset();
 }
 
 void RasterizerVulkan::UpdateLineWidth(Tegra::Engines::Maxwell3D::Regs& regs) {
@@ -868,6 +1025,82 @@ void RasterizerVulkan::UpdateDepthWriteEnable(Tegra::Engines::Maxwell3D::Regs& r
     });
 }
 
+void RasterizerVulkan::UpdatePrimitiveRestartEnable(Tegra::Engines::Maxwell3D::Regs& regs) {
+    if (!state_tracker.TouchPrimitiveRestartEnable()) {
+        return;
+    }
+    scheduler.Record([enable = regs.primitive_restart.enabled](vk::CommandBuffer cmdbuf) {
+        cmdbuf.SetPrimitiveRestartEnableEXT(enable);
+    });
+}
+
+void RasterizerVulkan::UpdateRasterizerDiscardEnable(Tegra::Engines::Maxwell3D::Regs& regs) {
+    if (!state_tracker.TouchRasterizerDiscardEnable()) {
+        return;
+    }
+    scheduler.Record([disable = regs.rasterize_enable](vk::CommandBuffer cmdbuf) {
+        cmdbuf.SetRasterizerDiscardEnableEXT(disable == 0);
+    });
+}
+
+void RasterizerVulkan::UpdateDepthBiasEnable(Tegra::Engines::Maxwell3D::Regs& regs) {
+    if (!state_tracker.TouchDepthBiasEnable()) {
+        return;
+    }
+    constexpr size_t POINT = 0;
+    constexpr size_t LINE = 1;
+    constexpr size_t POLYGON = 2;
+    static constexpr std::array POLYGON_OFFSET_ENABLE_LUT = {
+        POINT,   // Points
+        LINE,    // Lines
+        LINE,    // LineLoop
+        LINE,    // LineStrip
+        POLYGON, // Triangles
+        POLYGON, // TriangleStrip
+        POLYGON, // TriangleFan
+        POLYGON, // Quads
+        POLYGON, // QuadStrip
+        POLYGON, // Polygon
+        LINE,    // LinesAdjacency
+        LINE,    // LineStripAdjacency
+        POLYGON, // TrianglesAdjacency
+        POLYGON, // TriangleStripAdjacency
+        POLYGON, // Patches
+    };
+    const std::array enabled_lut{
+        regs.polygon_offset_point_enable,
+        regs.polygon_offset_line_enable,
+        regs.polygon_offset_fill_enable,
+    };
+    const u32 topology_index = static_cast<u32>(maxwell3d->draw_manager->GetDrawState().topology);
+    const u32 enable = enabled_lut[POLYGON_OFFSET_ENABLE_LUT[topology_index]];
+    scheduler.Record(
+        [enable](vk::CommandBuffer cmdbuf) { cmdbuf.SetDepthBiasEnableEXT(enable != 0); });
+}
+
+void RasterizerVulkan::UpdateLogicOpEnable(Tegra::Engines::Maxwell3D::Regs& regs) {
+    if (!state_tracker.TouchLogicOpEnable()) {
+        return;
+    }
+    scheduler.Record([enable = regs.logic_op.enable](vk::CommandBuffer cmdbuf) {
+        cmdbuf.SetLogicOpEnableEXT(enable != 0);
+    });
+}
+
+void RasterizerVulkan::UpdateDepthClampEnable(Tegra::Engines::Maxwell3D::Regs& regs) {
+    if (!state_tracker.TouchDepthClampEnable()) {
+        return;
+    }
+    bool is_enabled = !(regs.viewport_clip_control.geometry_clip ==
+                            Maxwell::ViewportClipControl::GeometryClip::Passthrough ||
+                        regs.viewport_clip_control.geometry_clip ==
+                            Maxwell::ViewportClipControl::GeometryClip::FrustumXYZ ||
+                        regs.viewport_clip_control.geometry_clip ==
+                            Maxwell::ViewportClipControl::GeometryClip::FrustumZ);
+    scheduler.Record(
+        [is_enabled](vk::CommandBuffer cmdbuf) { cmdbuf.SetDepthClampEnableEXT(is_enabled); });
+}
+
 void RasterizerVulkan::UpdateDepthCompareOp(Tegra::Engines::Maxwell3D::Regs& regs) {
     if (!state_tracker.TouchDepthCompareOp()) {
         return;
@@ -925,6 +1158,78 @@ void RasterizerVulkan::UpdateStencilOp(Tegra::Engines::Maxwell3D::Regs& regs) {
     }
 }
 
+void RasterizerVulkan::UpdateLogicOp(Tegra::Engines::Maxwell3D::Regs& regs) {
+    if (!state_tracker.TouchLogicOp()) {
+        return;
+    }
+    const auto op_value = static_cast<u32>(regs.logic_op.op);
+    auto op = op_value >= 0x1500 && op_value < 0x1510 ? static_cast<VkLogicOp>(op_value - 0x1500)
+                                                      : VK_LOGIC_OP_NO_OP;
+    scheduler.Record([op](vk::CommandBuffer cmdbuf) { cmdbuf.SetLogicOpEXT(op); });
+}
+
+void RasterizerVulkan::UpdateBlending(Tegra::Engines::Maxwell3D::Regs& regs) {
+    if (!state_tracker.TouchBlending()) {
+        return;
+    }
+
+    if (state_tracker.TouchColorMask()) {
+        std::array<VkColorComponentFlags, Maxwell::NumRenderTargets> setup_masks{};
+        for (size_t index = 0; index < Maxwell::NumRenderTargets; index++) {
+            const auto& mask = regs.color_mask[regs.color_mask_common ? 0 : index];
+            auto& current = setup_masks[index];
+            if (mask.R) {
+                current |= VK_COLOR_COMPONENT_R_BIT;
+            }
+            if (mask.G) {
+                current |= VK_COLOR_COMPONENT_G_BIT;
+            }
+            if (mask.B) {
+                current |= VK_COLOR_COMPONENT_B_BIT;
+            }
+            if (mask.A) {
+                current |= VK_COLOR_COMPONENT_A_BIT;
+            }
+        }
+        scheduler.Record([setup_masks](vk::CommandBuffer cmdbuf) {
+            cmdbuf.SetColorWriteMaskEXT(0, setup_masks);
+        });
+    }
+
+    if (state_tracker.TouchBlendEnable()) {
+        std::array<VkBool32, Maxwell::NumRenderTargets> setup_enables{};
+        std::ranges::transform(
+            regs.blend.enable, setup_enables.begin(),
+            [&](const auto& is_enabled) { return is_enabled != 0 ? VK_TRUE : VK_FALSE; });
+        scheduler.Record([setup_enables](vk::CommandBuffer cmdbuf) {
+            cmdbuf.SetColorBlendEnableEXT(0, setup_enables);
+        });
+    }
+
+    if (state_tracker.TouchBlendEquations()) {
+        std::array<VkColorBlendEquationEXT, Maxwell::NumRenderTargets> setup_blends{};
+        for (size_t index = 0; index < Maxwell::NumRenderTargets; index++) {
+            const auto blend_setup = [&]<typename T>(const T& guest_blend) {
+                auto& host_blend = setup_blends[index];
+                host_blend.srcColorBlendFactor = MaxwellToVK::BlendFactor(guest_blend.color_source);
+                host_blend.dstColorBlendFactor = MaxwellToVK::BlendFactor(guest_blend.color_dest);
+                host_blend.colorBlendOp = MaxwellToVK::BlendEquation(guest_blend.color_op);
+                host_blend.srcAlphaBlendFactor = MaxwellToVK::BlendFactor(guest_blend.alpha_source);
+                host_blend.dstAlphaBlendFactor = MaxwellToVK::BlendFactor(guest_blend.alpha_dest);
+                host_blend.alphaBlendOp = MaxwellToVK::BlendEquation(guest_blend.alpha_op);
+            };
+            if (!regs.blend_per_target_enabled) {
+                blend_setup(regs.blend);
+                continue;
+            }
+            blend_setup(regs.blend_per_target[index]);
+        }
+        scheduler.Record([setup_blends](vk::CommandBuffer cmdbuf) {
+            cmdbuf.SetColorBlendEquationEXT(0, setup_blends);
+        });
+    }
+}
+
 void RasterizerVulkan::UpdateStencilTestEnable(Tegra::Engines::Maxwell3D::Regs& regs) {
     if (!state_tracker.TouchStencilTestEnable()) {
         return;
diff --git a/src/video_core/renderer_vulkan/vk_rasterizer.h b/src/video_core/renderer_vulkan/vk_rasterizer.h
index ee483cfd9..c661e5b19 100644
--- a/src/video_core/renderer_vulkan/vk_rasterizer.h
+++ b/src/video_core/renderer_vulkan/vk_rasterizer.h
@@ -65,6 +65,7 @@ public:
     ~RasterizerVulkan() override;
 
     void Draw(bool is_indexed, u32 instance_count) override;
+    void DrawIndirect() override;
     void Clear(u32 layer_count) override;
     void DispatchCompute() override;
     void ResetCounter(VideoCore::QueryType type) override;
@@ -72,9 +73,12 @@ public:
     void BindGraphicsUniformBuffer(size_t stage, u32 index, GPUVAddr gpu_addr, u32 size) override;
     void DisableGraphicsUniformBuffer(size_t stage, u32 index) override;
     void FlushAll() override;
-    void FlushRegion(VAddr addr, u64 size) override;
-    bool MustFlushRegion(VAddr addr, u64 size) override;
-    void InvalidateRegion(VAddr addr, u64 size) override;
+    void FlushRegion(VAddr addr, u64 size,
+                     VideoCommon::CacheType which = VideoCommon::CacheType::All) override;
+    bool MustFlushRegion(VAddr addr, u64 size,
+                         VideoCommon::CacheType which = VideoCommon::CacheType::All) override;
+    void InvalidateRegion(VAddr addr, u64 size,
+                          VideoCommon::CacheType which = VideoCommon::CacheType::All) override;
     void OnCPUWrite(VAddr addr, u64 size) override;
     void InvalidateGPUCache() override;
     void UnmapMemory(VAddr addr, u64 size) override;
@@ -84,12 +88,14 @@ public:
     void SignalSyncPoint(u32 value) override;
     void SignalReference() override;
     void ReleaseFences() override;
-    void FlushAndInvalidateRegion(VAddr addr, u64 size) override;
+    void FlushAndInvalidateRegion(
+        VAddr addr, u64 size, VideoCommon::CacheType which = VideoCommon::CacheType::All) override;
     void WaitForIdle() override;
     void FragmentBarrier() override;
     void TiledCacheBarrier() override;
     void FlushCommands() override;
     void TickFrame() override;
+    bool AccelerateConditionalRendering() override;
     bool AccelerateSurfaceCopy(const Tegra::Engines::Fermi2D::Surface& src,
                                const Tegra::Engines::Fermi2D::Surface& dst,
                                const Tegra::Engines::Fermi2D::Config& copy_config) override;
@@ -114,6 +120,9 @@ private:
 
     static constexpr VkDeviceSize DEFAULT_BUFFER_SIZE = 4 * sizeof(float);
 
+    template <typename Func>
+    void PrepareDraw(bool is_indexed, Func&&);
+
     void FlushWork();
 
     void UpdateDynamicStates();
@@ -135,9 +144,16 @@ private:
     void UpdateDepthTestEnable(Tegra::Engines::Maxwell3D::Regs& regs);
     void UpdateDepthWriteEnable(Tegra::Engines::Maxwell3D::Regs& regs);
     void UpdateDepthCompareOp(Tegra::Engines::Maxwell3D::Regs& regs);
+    void UpdatePrimitiveRestartEnable(Tegra::Engines::Maxwell3D::Regs& regs);
+    void UpdateRasterizerDiscardEnable(Tegra::Engines::Maxwell3D::Regs& regs);
+    void UpdateDepthBiasEnable(Tegra::Engines::Maxwell3D::Regs& regs);
+    void UpdateLogicOpEnable(Tegra::Engines::Maxwell3D::Regs& regs);
+    void UpdateDepthClampEnable(Tegra::Engines::Maxwell3D::Regs& regs);
     void UpdateFrontFace(Tegra::Engines::Maxwell3D::Regs& regs);
     void UpdateStencilOp(Tegra::Engines::Maxwell3D::Regs& regs);
     void UpdateStencilTestEnable(Tegra::Engines::Maxwell3D::Regs& regs);
+    void UpdateLogicOp(Tegra::Engines::Maxwell3D::Regs& regs);
+    void UpdateBlending(Tegra::Engines::Maxwell3D::Regs& regs);
 
     void UpdateVertexInput(Tegra::Engines::Maxwell3D::Regs& regs);
 
diff --git a/src/video_core/renderer_vulkan/vk_staging_buffer_pool.cpp b/src/video_core/renderer_vulkan/vk_staging_buffer_pool.cpp
index 06f68d09a..74ca77216 100644
--- a/src/video_core/renderer_vulkan/vk_staging_buffer_pool.cpp
+++ b/src/video_core/renderer_vulkan/vk_staging_buffer_pool.cpp
@@ -1,5 +1,5 @@
-// SPDX-FileCopyrightText: Copyright 2019 yuzu Emulator Project
-// SPDX-License-Identifier: GPL-2.0-or-later
+// SPDX-FileCopyrightText: Copyright 2022 yuzu Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
 
 #include <algorithm>
 #include <utility>
@@ -94,7 +94,8 @@ StagingBufferPool::StagingBufferPool(const Device& device_, MemoryAllocator& mem
         .flags = 0,
         .size = STREAM_BUFFER_SIZE,
         .usage = VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_UNIFORM_BUFFER_BIT |
-                 VK_BUFFER_USAGE_INDEX_BUFFER_BIT | VK_BUFFER_USAGE_STORAGE_BUFFER_BIT,
+                 VK_BUFFER_USAGE_INDEX_BUFFER_BIT | VK_BUFFER_USAGE_STORAGE_BUFFER_BIT |
+                 VK_BUFFER_USAGE_TRANSFORM_FEEDBACK_BUFFER_BIT_EXT,
         .sharingMode = VK_SHARING_MODE_EXCLUSIVE,
         .queueFamilyIndexCount = 0,
         .pQueueFamilyIndices = nullptr,
@@ -142,11 +143,23 @@ StagingBufferPool::StagingBufferPool(const Device& device_, MemoryAllocator& mem
 
 StagingBufferPool::~StagingBufferPool() = default;
 
-StagingBufferRef StagingBufferPool::Request(size_t size, MemoryUsage usage) {
-    if (usage == MemoryUsage::Upload && size <= MAX_STREAM_BUFFER_REQUEST_SIZE) {
+StagingBufferRef StagingBufferPool::Request(size_t size, MemoryUsage usage, bool deferred) {
+    if (!deferred && usage == MemoryUsage::Upload && size <= MAX_STREAM_BUFFER_REQUEST_SIZE) {
         return GetStreamBuffer(size);
     }
-    return GetStagingBuffer(size, usage);
+    return GetStagingBuffer(size, usage, deferred);
+}
+
+void StagingBufferPool::FreeDeferred(StagingBufferRef& ref) {
+    auto& entries = GetCache(ref.usage)[ref.log2_level].entries;
+    const auto is_this_one = [&ref](const StagingBuffer& entry) {
+        return entry.index == ref.index;
+    };
+    auto it = std::find_if(entries.begin(), entries.end(), is_this_one);
+    ASSERT(it != entries.end());
+    ASSERT(it->deferred);
+    it->tick = scheduler.CurrentTick();
+    it->deferred = false;
 }
 
 void StagingBufferPool::TickFrame() {
@@ -187,6 +200,9 @@ StagingBufferRef StagingBufferPool::GetStreamBuffer(size_t size) {
         .buffer = *stream_buffer,
         .offset = static_cast<VkDeviceSize>(offset),
         .mapped_span = std::span<u8>(stream_pointer + offset, size),
+        .usage{},
+        .log2_level{},
+        .index{},
     };
 }
 
@@ -196,19 +212,21 @@ bool StagingBufferPool::AreRegionsActive(size_t region_begin, size_t region_end)
                        [gpu_tick](u64 sync_tick) { return gpu_tick < sync_tick; });
 };
 
-StagingBufferRef StagingBufferPool::GetStagingBuffer(size_t size, MemoryUsage usage) {
-    if (const std::optional<StagingBufferRef> ref = TryGetReservedBuffer(size, usage)) {
+StagingBufferRef StagingBufferPool::GetStagingBuffer(size_t size, MemoryUsage usage,
+                                                     bool deferred) {
+    if (const std::optional<StagingBufferRef> ref = TryGetReservedBuffer(size, usage, deferred)) {
         return *ref;
     }
-    return CreateStagingBuffer(size, usage);
+    return CreateStagingBuffer(size, usage, deferred);
 }
 
 std::optional<StagingBufferRef> StagingBufferPool::TryGetReservedBuffer(size_t size,
-                                                                        MemoryUsage usage) {
+                                                                        MemoryUsage usage,
+                                                                        bool deferred) {
     StagingBuffers& cache_level = GetCache(usage)[Common::Log2Ceil64(size)];
 
     const auto is_free = [this](const StagingBuffer& entry) {
-        return scheduler.IsFree(entry.tick);
+        return !entry.deferred && scheduler.IsFree(entry.tick);
     };
     auto& entries = cache_level.entries;
     const auto hint_it = entries.begin() + cache_level.iterate_index;
@@ -220,11 +238,14 @@ std::optional<StagingBufferRef> StagingBufferPool::TryGetReservedBuffer(size_t s
         }
     }
     cache_level.iterate_index = std::distance(entries.begin(), it) + 1;
-    it->tick = scheduler.CurrentTick();
+    it->tick = deferred ? std::numeric_limits<u64>::max() : scheduler.CurrentTick();
+    ASSERT(!it->deferred);
+    it->deferred = deferred;
     return it->Ref();
 }
 
-StagingBufferRef StagingBufferPool::CreateStagingBuffer(size_t size, MemoryUsage usage) {
+StagingBufferRef StagingBufferPool::CreateStagingBuffer(size_t size, MemoryUsage usage,
+                                                        bool deferred) {
     const u32 log2 = Common::Log2Ceil64(size);
     vk::Buffer buffer = device.GetLogical().CreateBuffer({
         .sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO,
@@ -233,7 +254,8 @@ StagingBufferRef StagingBufferPool::CreateStagingBuffer(size_t size, MemoryUsage
         .size = 1ULL << log2,
         .usage = VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT |
                  VK_BUFFER_USAGE_UNIFORM_BUFFER_BIT | VK_BUFFER_USAGE_STORAGE_BUFFER_BIT |
-                 VK_BUFFER_USAGE_INDEX_BUFFER_BIT | VK_BUFFER_USAGE_VERTEX_BUFFER_BIT,
+                 VK_BUFFER_USAGE_INDEX_BUFFER_BIT | VK_BUFFER_USAGE_VERTEX_BUFFER_BIT |
+                 VK_BUFFER_USAGE_TRANSFORM_FEEDBACK_BUFFER_BIT_EXT,
         .sharingMode = VK_SHARING_MODE_EXCLUSIVE,
         .queueFamilyIndexCount = 0,
         .pQueueFamilyIndices = nullptr,
@@ -249,7 +271,11 @@ StagingBufferRef StagingBufferPool::CreateStagingBuffer(size_t size, MemoryUsage
         .buffer = std::move(buffer),
         .commit = std::move(commit),
         .mapped_span = mapped_span,
-        .tick = scheduler.CurrentTick(),
+        .usage = usage,
+        .log2_level = log2,
+        .index = unique_ids++,
+        .tick = deferred ? std::numeric_limits<u64>::max() : scheduler.CurrentTick(),
+        .deferred = deferred,
     });
     return entry.Ref();
 }
diff --git a/src/video_core/renderer_vulkan/vk_staging_buffer_pool.h b/src/video_core/renderer_vulkan/vk_staging_buffer_pool.h
index 91dc84da8..4fd15f11a 100644
--- a/src/video_core/renderer_vulkan/vk_staging_buffer_pool.h
+++ b/src/video_core/renderer_vulkan/vk_staging_buffer_pool.h
@@ -1,5 +1,5 @@
-// SPDX-FileCopyrightText: Copyright 2019 yuzu Emulator Project
-// SPDX-License-Identifier: GPL-2.0-or-later
+// SPDX-FileCopyrightText: Copyright 2022 yuzu Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
 
 #pragma once
 
@@ -20,6 +20,9 @@ struct StagingBufferRef {
     VkBuffer buffer;
     VkDeviceSize offset;
     std::span<u8> mapped_span;
+    MemoryUsage usage;
+    u32 log2_level;
+    u64 index;
 };
 
 class StagingBufferPool {
@@ -30,7 +33,8 @@ public:
                                Scheduler& scheduler);
     ~StagingBufferPool();
 
-    StagingBufferRef Request(size_t size, MemoryUsage usage);
+    StagingBufferRef Request(size_t size, MemoryUsage usage, bool deferred = false);
+    void FreeDeferred(StagingBufferRef& ref);
 
     void TickFrame();
 
@@ -44,13 +48,20 @@ private:
         vk::Buffer buffer;
         MemoryCommit commit;
         std::span<u8> mapped_span;
+        MemoryUsage usage;
+        u32 log2_level;
+        u64 index;
         u64 tick = 0;
+        bool deferred{};
 
         StagingBufferRef Ref() const noexcept {
             return {
                 .buffer = *buffer,
                 .offset = 0,
                 .mapped_span = mapped_span,
+                .usage = usage,
+                .log2_level = log2_level,
+                .index = index,
             };
         }
     };
@@ -68,11 +79,12 @@ private:
 
     bool AreRegionsActive(size_t region_begin, size_t region_end) const;
 
-    StagingBufferRef GetStagingBuffer(size_t size, MemoryUsage usage);
+    StagingBufferRef GetStagingBuffer(size_t size, MemoryUsage usage, bool deferred = false);
 
-    std::optional<StagingBufferRef> TryGetReservedBuffer(size_t size, MemoryUsage usage);
+    std::optional<StagingBufferRef> TryGetReservedBuffer(size_t size, MemoryUsage usage,
+                                                         bool deferred);
 
-    StagingBufferRef CreateStagingBuffer(size_t size, MemoryUsage usage);
+    StagingBufferRef CreateStagingBuffer(size_t size, MemoryUsage usage, bool deferred);
 
     StagingBuffersCache& GetCache(MemoryUsage usage);
 
@@ -99,6 +111,7 @@ private:
 
     size_t current_delete_level = 0;
     u64 buffer_index = 0;
+    u64 unique_ids{};
 };
 
 } // namespace Vulkan
diff --git a/src/video_core/renderer_vulkan/vk_state_tracker.cpp b/src/video_core/renderer_vulkan/vk_state_tracker.cpp
index edb41b171..d56558a83 100644
--- a/src/video_core/renderer_vulkan/vk_state_tracker.cpp
+++ b/src/video_core/renderer_vulkan/vk_state_tracker.cpp
@@ -27,10 +27,37 @@ using Flags = Maxwell3D::DirtyState::Flags;
 
 Flags MakeInvalidationFlags() {
     static constexpr int INVALIDATION_FLAGS[]{
-        Viewports,         Scissors,       DepthBias, BlendConstants,    DepthBounds,
-        StencilProperties, LineWidth,      CullMode,  DepthBoundsEnable, DepthTestEnable,
-        DepthWriteEnable,  DepthCompareOp, FrontFace, StencilOp,         StencilTestEnable,
-        VertexBuffers,     VertexInput,
+        Viewports,
+        Scissors,
+        DepthBias,
+        BlendConstants,
+        DepthBounds,
+        StencilProperties,
+        StencilReference,
+        StencilWriteMask,
+        StencilCompare,
+        LineWidth,
+        CullMode,
+        DepthBoundsEnable,
+        DepthTestEnable,
+        DepthWriteEnable,
+        DepthCompareOp,
+        FrontFace,
+        StencilOp,
+        StencilTestEnable,
+        VertexBuffers,
+        VertexInput,
+        StateEnable,
+        PrimitiveRestartEnable,
+        RasterizerDiscardEnable,
+        DepthBiasEnable,
+        LogicOpEnable,
+        DepthClampEnable,
+        LogicOp,
+        Blending,
+        ColorMask,
+        BlendEquations,
+        BlendEnable,
     };
     Flags flags{};
     for (const int flag : INVALIDATION_FLAGS) {
@@ -75,14 +102,17 @@ void SetupDirtyDepthBounds(Tables& tables) {
 }
 
 void SetupDirtyStencilProperties(Tables& tables) {
-    auto& table = tables[0];
-    table[OFF(stencil_two_side_enable)] = StencilProperties;
-    table[OFF(stencil_front_ref)] = StencilProperties;
-    table[OFF(stencil_front_mask)] = StencilProperties;
-    table[OFF(stencil_front_func_mask)] = StencilProperties;
-    table[OFF(stencil_back_ref)] = StencilProperties;
-    table[OFF(stencil_back_mask)] = StencilProperties;
-    table[OFF(stencil_back_func_mask)] = StencilProperties;
+    const auto setup = [&](size_t position, u8 flag) {
+        tables[0][position] = flag;
+        tables[1][position] = StencilProperties;
+    };
+    tables[0][OFF(stencil_two_side_enable)] = StencilProperties;
+    setup(OFF(stencil_front_ref), StencilReference);
+    setup(OFF(stencil_front_mask), StencilWriteMask);
+    setup(OFF(stencil_front_func_mask), StencilCompare);
+    setup(OFF(stencil_back_ref), StencilReference);
+    setup(OFF(stencil_back_mask), StencilWriteMask);
+    setup(OFF(stencil_back_func_mask), StencilCompare);
 }
 
 void SetupDirtyLineWidth(Tables& tables) {
@@ -96,16 +126,22 @@ void SetupDirtyCullMode(Tables& tables) {
     table[OFF(gl_cull_test_enabled)] = CullMode;
 }
 
-void SetupDirtyDepthBoundsEnable(Tables& tables) {
-    tables[0][OFF(depth_bounds_enable)] = DepthBoundsEnable;
-}
-
-void SetupDirtyDepthTestEnable(Tables& tables) {
-    tables[0][OFF(depth_test_enable)] = DepthTestEnable;
-}
-
-void SetupDirtyDepthWriteEnable(Tables& tables) {
-    tables[0][OFF(depth_write_enabled)] = DepthWriteEnable;
+void SetupDirtyStateEnable(Tables& tables) {
+    const auto setup = [&](size_t position, u8 flag) {
+        tables[0][position] = flag;
+        tables[1][position] = StateEnable;
+    };
+    setup(OFF(depth_bounds_enable), DepthBoundsEnable);
+    setup(OFF(depth_test_enable), DepthTestEnable);
+    setup(OFF(depth_write_enabled), DepthWriteEnable);
+    setup(OFF(stencil_enable), StencilTestEnable);
+    setup(OFF(primitive_restart.enabled), PrimitiveRestartEnable);
+    setup(OFF(rasterize_enable), RasterizerDiscardEnable);
+    setup(OFF(polygon_offset_point_enable), DepthBiasEnable);
+    setup(OFF(polygon_offset_line_enable), DepthBiasEnable);
+    setup(OFF(polygon_offset_fill_enable), DepthBiasEnable);
+    setup(OFF(logic_op.enable), LogicOpEnable);
+    setup(OFF(viewport_clip_control.geometry_clip), DepthClampEnable);
 }
 
 void SetupDirtyDepthCompareOp(Tables& tables) {
@@ -133,16 +169,22 @@ void SetupDirtyStencilOp(Tables& tables) {
     tables[1][OFF(stencil_two_side_enable)] = StencilOp;
 }
 
-void SetupDirtyStencilTestEnable(Tables& tables) {
-    tables[0][OFF(stencil_enable)] = StencilTestEnable;
-}
-
 void SetupDirtyBlending(Tables& tables) {
     tables[0][OFF(color_mask_common)] = Blending;
+    tables[1][OFF(color_mask_common)] = ColorMask;
     tables[0][OFF(blend_per_target_enabled)] = Blending;
+    tables[1][OFF(blend_per_target_enabled)] = BlendEquations;
     FillBlock(tables[0], OFF(color_mask), NUM(color_mask), Blending);
+    FillBlock(tables[1], OFF(color_mask), NUM(color_mask), ColorMask);
     FillBlock(tables[0], OFF(blend), NUM(blend), Blending);
+    FillBlock(tables[1], OFF(blend), NUM(blend), BlendEquations);
+    FillBlock(tables[1], OFF(blend.enable), NUM(blend.enable), BlendEnable);
     FillBlock(tables[0], OFF(blend_per_target), NUM(blend_per_target), Blending);
+    FillBlock(tables[1], OFF(blend_per_target), NUM(blend_per_target), BlendEquations);
+}
+
+void SetupDirtySpecialOps(Tables& tables) {
+    tables[0][OFF(logic_op.op)] = LogicOp;
 }
 
 void SetupDirtyViewportSwizzles(Tables& tables) {
@@ -185,17 +227,15 @@ void StateTracker::SetupTables(Tegra::Control::ChannelState& channel_state) {
     SetupDirtyStencilProperties(tables);
     SetupDirtyLineWidth(tables);
     SetupDirtyCullMode(tables);
-    SetupDirtyDepthBoundsEnable(tables);
-    SetupDirtyDepthTestEnable(tables);
-    SetupDirtyDepthWriteEnable(tables);
+    SetupDirtyStateEnable(tables);
     SetupDirtyDepthCompareOp(tables);
     SetupDirtyFrontFace(tables);
     SetupDirtyStencilOp(tables);
-    SetupDirtyStencilTestEnable(tables);
     SetupDirtyBlending(tables);
     SetupDirtyViewportSwizzles(tables);
     SetupDirtyVertexAttributes(tables);
     SetupDirtyVertexBindings(tables);
+    SetupDirtySpecialOps(tables);
 }
 
 void StateTracker::ChangeChannel(Tegra::Control::ChannelState& channel_state) {
@@ -204,6 +244,8 @@ void StateTracker::ChangeChannel(Tegra::Control::ChannelState& channel_state) {
 
 void StateTracker::InvalidateState() {
     flags->set();
+    current_topology = INVALID_TOPOLOGY;
+    stencil_reset = true;
 }
 
 StateTracker::StateTracker()
diff --git a/src/video_core/renderer_vulkan/vk_state_tracker.h b/src/video_core/renderer_vulkan/vk_state_tracker.h
index 2296dea60..8010ad26c 100644
--- a/src/video_core/renderer_vulkan/vk_state_tracker.h
+++ b/src/video_core/renderer_vulkan/vk_state_tracker.h
@@ -35,6 +35,9 @@ enum : u8 {
     BlendConstants,
     DepthBounds,
     StencilProperties,
+    StencilReference,
+    StencilWriteMask,
+    StencilCompare,
     LineWidth,
 
     CullMode,
@@ -45,8 +48,18 @@ enum : u8 {
     FrontFace,
     StencilOp,
     StencilTestEnable,
+    PrimitiveRestartEnable,
+    RasterizerDiscardEnable,
+    DepthBiasEnable,
+    StateEnable,
+    LogicOp,
+    LogicOpEnable,
+    DepthClampEnable,
 
     Blending,
+    BlendEnable,
+    BlendEquations,
+    ColorMask,
     ViewportSwizzles,
 
     Last,
@@ -64,6 +77,7 @@ public:
     void InvalidateCommandBufferState() {
         (*flags) |= invalidation_flags;
         current_topology = INVALID_TOPOLOGY;
+        stencil_reset = true;
     }
 
     void InvalidateViewports() {
@@ -103,6 +117,57 @@ public:
         return Exchange(Dirty::StencilProperties, false);
     }
 
+    bool TouchStencilReference() {
+        return Exchange(Dirty::StencilReference, false);
+    }
+
+    bool TouchStencilWriteMask() {
+        return Exchange(Dirty::StencilWriteMask, false);
+    }
+
+    bool TouchStencilCompare() {
+        return Exchange(Dirty::StencilCompare, false);
+    }
+
+    template <typename T>
+    bool ExchangeCheck(T& old_value, T new_value) {
+        bool result = old_value != new_value;
+        old_value = new_value;
+        return result;
+    }
+
+    bool TouchStencilSide(bool two_sided_stencil_new) {
+        return ExchangeCheck(two_sided_stencil, two_sided_stencil_new) || stencil_reset;
+    }
+
+    bool CheckStencilReferenceFront(u32 new_value) {
+        return ExchangeCheck(front.ref, new_value) || stencil_reset;
+    }
+
+    bool CheckStencilReferenceBack(u32 new_value) {
+        return ExchangeCheck(back.ref, new_value) || stencil_reset;
+    }
+
+    bool CheckStencilWriteMaskFront(u32 new_value) {
+        return ExchangeCheck(front.write_mask, new_value) || stencil_reset;
+    }
+
+    bool CheckStencilWriteMaskBack(u32 new_value) {
+        return ExchangeCheck(back.write_mask, new_value) || stencil_reset;
+    }
+
+    bool CheckStencilCompareMaskFront(u32 new_value) {
+        return ExchangeCheck(front.compare_mask, new_value) || stencil_reset;
+    }
+
+    bool CheckStencilCompareMaskBack(u32 new_value) {
+        return ExchangeCheck(back.compare_mask, new_value) || stencil_reset;
+    }
+
+    void ClearStencilReset() {
+        stencil_reset = false;
+    }
+
     bool TouchLineWidth() const {
         return Exchange(Dirty::LineWidth, false);
     }
@@ -111,6 +176,10 @@ public:
         return Exchange(Dirty::CullMode, false);
     }
 
+    bool TouchStateEnable() {
+        return Exchange(Dirty::StateEnable, false);
+    }
+
     bool TouchDepthBoundsTestEnable() {
         return Exchange(Dirty::DepthBoundsEnable, false);
     }
@@ -123,6 +192,26 @@ public:
         return Exchange(Dirty::DepthWriteEnable, false);
     }
 
+    bool TouchPrimitiveRestartEnable() {
+        return Exchange(Dirty::PrimitiveRestartEnable, false);
+    }
+
+    bool TouchRasterizerDiscardEnable() {
+        return Exchange(Dirty::RasterizerDiscardEnable, false);
+    }
+
+    bool TouchDepthBiasEnable() {
+        return Exchange(Dirty::DepthBiasEnable, false);
+    }
+
+    bool TouchLogicOpEnable() {
+        return Exchange(Dirty::LogicOpEnable, false);
+    }
+
+    bool TouchDepthClampEnable() {
+        return Exchange(Dirty::DepthClampEnable, false);
+    }
+
     bool TouchDepthCompareOp() {
         return Exchange(Dirty::DepthCompareOp, false);
     }
@@ -135,10 +224,30 @@ public:
         return Exchange(Dirty::StencilOp, false);
     }
 
+    bool TouchBlending() {
+        return Exchange(Dirty::Blending, false);
+    }
+
+    bool TouchBlendEnable() {
+        return Exchange(Dirty::BlendEnable, false);
+    }
+
+    bool TouchBlendEquations() {
+        return Exchange(Dirty::BlendEquations, false);
+    }
+
+    bool TouchColorMask() {
+        return Exchange(Dirty::ColorMask, false);
+    }
+
     bool TouchStencilTestEnable() {
         return Exchange(Dirty::StencilTestEnable, false);
     }
 
+    bool TouchLogicOp() {
+        return Exchange(Dirty::LogicOp, false);
+    }
+
     bool ChangePrimitiveTopology(Maxwell::PrimitiveTopology new_topology) {
         const bool has_changed = current_topology != new_topology;
         current_topology = new_topology;
@@ -160,10 +269,20 @@ private:
         return is_dirty;
     }
 
+    struct StencilProperties {
+        u32 ref = 0;
+        u32 write_mask = 0;
+        u32 compare_mask = 0;
+    };
+
     Tegra::Engines::Maxwell3D::DirtyState::Flags* flags;
     Tegra::Engines::Maxwell3D::DirtyState::Flags default_flags;
     Tegra::Engines::Maxwell3D::DirtyState::Flags invalidation_flags;
     Maxwell::PrimitiveTopology current_topology = INVALID_TOPOLOGY;
+    bool two_sided_stencil = false;
+    StencilProperties front{};
+    StencilProperties back{};
+    bool stencil_reset = false;
 };
 
 } // namespace Vulkan
diff --git a/src/video_core/renderer_vulkan/vk_texture_cache.cpp b/src/video_core/renderer_vulkan/vk_texture_cache.cpp
index a65bbeb1c..d39372ec4 100644
--- a/src/video_core/renderer_vulkan/vk_texture_cache.cpp
+++ b/src/video_core/renderer_vulkan/vk_texture_cache.cpp
@@ -812,8 +812,12 @@ StagingBufferRef TextureCacheRuntime::UploadStagingBuffer(size_t size) {
     return staging_buffer_pool.Request(size, MemoryUsage::Upload);
 }
 
-StagingBufferRef TextureCacheRuntime::DownloadStagingBuffer(size_t size) {
-    return staging_buffer_pool.Request(size, MemoryUsage::Download);
+StagingBufferRef TextureCacheRuntime::DownloadStagingBuffer(size_t size, bool deferred) {
+    return staging_buffer_pool.Request(size, MemoryUsage::Download, deferred);
+}
+
+void TextureCacheRuntime::FreeDeferredStagingBuffer(StagingBufferRef& ref) {
+    staging_buffer_pool.FreeDeferred(ref);
 }
 
 bool TextureCacheRuntime::ShouldReinterpret(Image& dst, Image& src) {
diff --git a/src/video_core/renderer_vulkan/vk_texture_cache.h b/src/video_core/renderer_vulkan/vk_texture_cache.h
index 7ec0df134..1f27a3589 100644
--- a/src/video_core/renderer_vulkan/vk_texture_cache.h
+++ b/src/video_core/renderer_vulkan/vk_texture_cache.h
@@ -51,7 +51,9 @@ public:
 
     StagingBufferRef UploadStagingBuffer(size_t size);
 
-    StagingBufferRef DownloadStagingBuffer(size_t size);
+    StagingBufferRef DownloadStagingBuffer(size_t size, bool deferred = false);
+
+    void FreeDeferredStagingBuffer(StagingBufferRef& ref);
 
     void TickFrame();
 
@@ -347,6 +349,7 @@ struct TextureCacheParams {
     static constexpr bool FRAMEBUFFER_BLITS = false;
     static constexpr bool HAS_EMULATED_COPIES = false;
     static constexpr bool HAS_DEVICE_MEMORY_INFO = true;
+    static constexpr bool IMPLEMENTS_ASYNC_DOWNLOADS = true;
 
     using Runtime = Vulkan::TextureCacheRuntime;
     using Image = Vulkan::Image;
@@ -354,6 +357,7 @@ struct TextureCacheParams {
     using ImageView = Vulkan::ImageView;
     using Sampler = Vulkan::Sampler;
     using Framebuffer = Vulkan::Framebuffer;
+    using AsyncBuffer = Vulkan::StagingBufferRef;
 };
 
 using TextureCache = VideoCommon::TextureCache<TextureCacheParams>;
diff --git a/src/video_core/shader_environment.cpp b/src/video_core/shader_environment.cpp
index 958810747..574760f80 100644
--- a/src/video_core/shader_environment.cpp
+++ b/src/video_core/shader_environment.cpp
@@ -202,12 +202,15 @@ void GenericEnvironment::Serialize(std::ofstream& file) const {
     const u64 num_texture_types{static_cast<u64>(texture_types.size())};
     const u64 num_texture_pixel_formats{static_cast<u64>(texture_pixel_formats.size())};
     const u64 num_cbuf_values{static_cast<u64>(cbuf_values.size())};
+    const u64 num_cbuf_replacement_values{static_cast<u64>(cbuf_replacements.size())};
 
     file.write(reinterpret_cast<const char*>(&code_size), sizeof(code_size))
         .write(reinterpret_cast<const char*>(&num_texture_types), sizeof(num_texture_types))
         .write(reinterpret_cast<const char*>(&num_texture_pixel_formats),
                sizeof(num_texture_pixel_formats))
         .write(reinterpret_cast<const char*>(&num_cbuf_values), sizeof(num_cbuf_values))
+        .write(reinterpret_cast<const char*>(&num_cbuf_replacement_values),
+               sizeof(num_cbuf_replacement_values))
         .write(reinterpret_cast<const char*>(&local_memory_size), sizeof(local_memory_size))
         .write(reinterpret_cast<const char*>(&texture_bound), sizeof(texture_bound))
         .write(reinterpret_cast<const char*>(&start_address), sizeof(start_address))
@@ -229,6 +232,10 @@ void GenericEnvironment::Serialize(std::ofstream& file) const {
         file.write(reinterpret_cast<const char*>(&key), sizeof(key))
             .write(reinterpret_cast<const char*>(&type), sizeof(type));
     }
+    for (const auto& [key, type] : cbuf_replacements) {
+        file.write(reinterpret_cast<const char*>(&key), sizeof(key))
+            .write(reinterpret_cast<const char*>(&type), sizeof(type));
+    }
     if (stage == Shader::Stage::Compute) {
         file.write(reinterpret_cast<const char*>(&workgroup_size), sizeof(workgroup_size))
             .write(reinterpret_cast<const char*>(&shared_memory_size), sizeof(shared_memory_size));
@@ -318,6 +325,9 @@ GraphicsEnvironment::GraphicsEnvironment(Tegra::Engines::Maxwell3D& maxwell3d_,
     ASSERT(local_size <= std::numeric_limits<u32>::max());
     local_memory_size = static_cast<u32>(local_size) + sph.common3.shader_local_memory_crs_size;
     texture_bound = maxwell3d->regs.bindless_texture_const_buffer_slot;
+    is_propietary_driver = texture_bound == 2;
+    has_hle_engine_state =
+        maxwell3d->engine_state == Tegra::Engines::Maxwell3D::EngineHint::OnHLEMacro;
 }
 
 u32 GraphicsEnvironment::ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) {
@@ -331,6 +341,32 @@ u32 GraphicsEnvironment::ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) {
     return value;
 }
 
+std::optional<Shader::ReplaceConstant> GraphicsEnvironment::GetReplaceConstBuffer(u32 bank,
+                                                                                  u32 offset) {
+    if (!has_hle_engine_state) {
+        return std::nullopt;
+    }
+    const u64 key = (static_cast<u64>(bank) << 32) | static_cast<u64>(offset);
+    auto it = maxwell3d->replace_table.find(key);
+    if (it == maxwell3d->replace_table.end()) {
+        return std::nullopt;
+    }
+    const auto converted_value = [](Tegra::Engines::Maxwell3D::HLEReplacementAttributeType name) {
+        switch (name) {
+        case Tegra::Engines::Maxwell3D::HLEReplacementAttributeType::BaseVertex:
+            return Shader::ReplaceConstant::BaseVertex;
+        case Tegra::Engines::Maxwell3D::HLEReplacementAttributeType::BaseInstance:
+            return Shader::ReplaceConstant::BaseInstance;
+        case Tegra::Engines::Maxwell3D::HLEReplacementAttributeType::DrawID:
+            return Shader::ReplaceConstant::DrawID;
+        default:
+            UNREACHABLE();
+        }
+    }(it->second);
+    cbuf_replacements.emplace(key, converted_value);
+    return converted_value;
+}
+
 Shader::TextureType GraphicsEnvironment::ReadTextureType(u32 handle) {
     const auto& regs{maxwell3d->regs};
     const bool via_header_index{regs.sampler_binding == Maxwell::SamplerBinding::ViaHeaderBinding};
@@ -366,6 +402,7 @@ ComputeEnvironment::ComputeEnvironment(Tegra::Engines::KeplerCompute& kepler_com
     stage = Shader::Stage::Compute;
     local_memory_size = qmd.local_pos_alloc + qmd.local_crs_alloc;
     texture_bound = kepler_compute->regs.tex_cb_index;
+    is_propietary_driver = texture_bound == 2;
     shared_memory_size = qmd.shared_alloc;
     workgroup_size = {qmd.block_dim_x, qmd.block_dim_y, qmd.block_dim_z};
 }
@@ -409,11 +446,14 @@ void FileEnvironment::Deserialize(std::ifstream& file) {
     u64 num_texture_types{};
     u64 num_texture_pixel_formats{};
     u64 num_cbuf_values{};
+    u64 num_cbuf_replacement_values{};
     file.read(reinterpret_cast<char*>(&code_size), sizeof(code_size))
         .read(reinterpret_cast<char*>(&num_texture_types), sizeof(num_texture_types))
         .read(reinterpret_cast<char*>(&num_texture_pixel_formats),
               sizeof(num_texture_pixel_formats))
         .read(reinterpret_cast<char*>(&num_cbuf_values), sizeof(num_cbuf_values))
+        .read(reinterpret_cast<char*>(&num_cbuf_replacement_values),
+              sizeof(num_cbuf_replacement_values))
         .read(reinterpret_cast<char*>(&local_memory_size), sizeof(local_memory_size))
         .read(reinterpret_cast<char*>(&texture_bound), sizeof(texture_bound))
         .read(reinterpret_cast<char*>(&start_address), sizeof(start_address))
@@ -444,6 +484,13 @@ void FileEnvironment::Deserialize(std::ifstream& file) {
             .read(reinterpret_cast<char*>(&value), sizeof(value));
         cbuf_values.emplace(key, value);
     }
+    for (size_t i = 0; i < num_cbuf_replacement_values; ++i) {
+        u64 key;
+        Shader::ReplaceConstant value;
+        file.read(reinterpret_cast<char*>(&key), sizeof(key))
+            .read(reinterpret_cast<char*>(&value), sizeof(value));
+        cbuf_replacements.emplace(key, value);
+    }
     if (stage == Shader::Stage::Compute) {
         file.read(reinterpret_cast<char*>(&workgroup_size), sizeof(workgroup_size))
             .read(reinterpret_cast<char*>(&shared_memory_size), sizeof(shared_memory_size));
@@ -455,6 +502,7 @@ void FileEnvironment::Deserialize(std::ifstream& file) {
             file.read(reinterpret_cast<char*>(&gp_passthrough_mask), sizeof(gp_passthrough_mask));
         }
     }
+    is_propietary_driver = texture_bound == 2;
 }
 
 void FileEnvironment::Dump(u64 hash) {
@@ -512,6 +560,16 @@ std::array<u32, 3> FileEnvironment::WorkgroupSize() const {
     return workgroup_size;
 }
 
+std::optional<Shader::ReplaceConstant> FileEnvironment::GetReplaceConstBuffer(u32 bank,
+                                                                              u32 offset) {
+    const u64 key = (static_cast<u64>(bank) << 32) | static_cast<u64>(offset);
+    auto it = cbuf_replacements.find(key);
+    if (it == cbuf_replacements.end()) {
+        return std::nullopt;
+    }
+    return it->second;
+}
+
 void SerializePipeline(std::span<const char> key, std::span<const GenericEnvironment* const> envs,
                        const std::filesystem::path& filename, u32 cache_version) try {
     std::ofstream file(filename, std::ios::binary | std::ios::ate | std::ios::app);
diff --git a/src/video_core/shader_environment.h b/src/video_core/shader_environment.h
index 1342fab1e..d75987a52 100644
--- a/src/video_core/shader_environment.h
+++ b/src/video_core/shader_environment.h
@@ -60,6 +60,10 @@ public:
 
     void Serialize(std::ofstream& file) const;
 
+    bool HasHLEMacroState() const override {
+        return has_hle_engine_state;
+    }
+
 protected:
     std::optional<u64> TryFindSize();
 
@@ -73,6 +77,7 @@ protected:
     std::unordered_map<u32, Shader::TextureType> texture_types;
     std::unordered_map<u32, Shader::TexturePixelFormat> texture_pixel_formats;
     std::unordered_map<u64, u32> cbuf_values;
+    std::unordered_map<u64, Shader::ReplaceConstant> cbuf_replacements;
 
     u32 local_memory_size{};
     u32 texture_bound{};
@@ -89,6 +94,7 @@ protected:
     u32 viewport_transform_state = 1;
 
     bool has_unbound_instructions = false;
+    bool has_hle_engine_state = false;
 };
 
 class GraphicsEnvironment final : public GenericEnvironment {
@@ -109,6 +115,8 @@ public:
 
     u32 ReadViewportTransformState() override;
 
+    std::optional<Shader::ReplaceConstant> GetReplaceConstBuffer(u32 bank, u32 offset) override;
+
 private:
     Tegra::Engines::Maxwell3D* maxwell3d{};
     size_t stage_index{};
@@ -131,6 +139,11 @@ public:
 
     u32 ReadViewportTransformState() override;
 
+    std::optional<Shader::ReplaceConstant> GetReplaceConstBuffer(
+        [[maybe_unused]] u32 bank, [[maybe_unused]] u32 offset) override {
+        return std::nullopt;
+    }
+
 private:
     Tegra::Engines::KeplerCompute* kepler_compute{};
 };
@@ -166,6 +179,13 @@ public:
 
     [[nodiscard]] std::array<u32, 3> WorkgroupSize() const override;
 
+    [[nodiscard]] std::optional<Shader::ReplaceConstant> GetReplaceConstBuffer(u32 bank,
+                                                                               u32 offset) override;
+
+    [[nodiscard]] bool HasHLEMacroState() const override {
+        return cbuf_replacements.size() != 0;
+    }
+
     void Dump(u64 hash) override;
 
 private:
@@ -173,6 +193,7 @@ private:
     std::unordered_map<u32, Shader::TextureType> texture_types;
     std::unordered_map<u32, Shader::TexturePixelFormat> texture_pixel_formats;
     std::unordered_map<u64, u32> cbuf_values;
+    std::unordered_map<u64, Shader::ReplaceConstant> cbuf_replacements;
     std::array<u32, 3> workgroup_size{};
     u32 local_memory_size{};
     u32 shared_memory_size{};
diff --git a/src/video_core/texture_cache/texture_cache.h b/src/video_core/texture_cache/texture_cache.h
index 27c82cd20..87152c8e9 100644
--- a/src/video_core/texture_cache/texture_cache.h
+++ b/src/video_core/texture_cache/texture_cache.h
@@ -646,7 +646,28 @@ bool TextureCache<P>::ShouldWaitAsyncFlushes() const noexcept {
 template <class P>
 void TextureCache<P>::CommitAsyncFlushes() {
     // This is intentionally passing the value by copy
-    committed_downloads.push(uncommitted_downloads);
+    if constexpr (IMPLEMENTS_ASYNC_DOWNLOADS) {
+        const std::span<const ImageId> download_ids = uncommitted_downloads;
+        if (download_ids.empty()) {
+            committed_downloads.emplace_back(std::move(uncommitted_downloads));
+            uncommitted_downloads.clear();
+            async_buffers.emplace_back(std::optional<AsyncBuffer>{});
+            return;
+        }
+        size_t total_size_bytes = 0;
+        for (const ImageId image_id : download_ids) {
+            total_size_bytes += slot_images[image_id].unswizzled_size_bytes;
+        }
+        auto download_map = runtime.DownloadStagingBuffer(total_size_bytes, true);
+        for (const ImageId image_id : download_ids) {
+            Image& image = slot_images[image_id];
+            const auto copies = FullDownloadCopies(image.info);
+            image.DownloadMemory(download_map, copies);
+            download_map.offset += Common::AlignUp(image.unswizzled_size_bytes, 64);
+        }
+        async_buffers.emplace_back(download_map);
+    }
+    committed_downloads.emplace_back(std::move(uncommitted_downloads));
     uncommitted_downloads.clear();
 }
 
@@ -655,37 +676,58 @@ void TextureCache<P>::PopAsyncFlushes() {
     if (committed_downloads.empty()) {
         return;
     }
-    const std::span<const ImageId> download_ids = committed_downloads.front();
-    if (download_ids.empty()) {
-        committed_downloads.pop();
-        return;
-    }
-    size_t total_size_bytes = 0;
-    for (const ImageId image_id : download_ids) {
-        total_size_bytes += slot_images[image_id].unswizzled_size_bytes;
-    }
-    auto download_map = runtime.DownloadStagingBuffer(total_size_bytes);
-    const size_t original_offset = download_map.offset;
-    for (const ImageId image_id : download_ids) {
-        Image& image = slot_images[image_id];
-        const auto copies = FullDownloadCopies(image.info);
-        image.DownloadMemory(download_map, copies);
-        download_map.offset += image.unswizzled_size_bytes;
-    }
-    // Wait for downloads to finish
-    runtime.Finish();
-
-    download_map.offset = original_offset;
-    std::span<u8> download_span = download_map.mapped_span;
-    for (const ImageId image_id : download_ids) {
-        const ImageBase& image = slot_images[image_id];
-        const auto copies = FullDownloadCopies(image.info);
-        SwizzleImage(*gpu_memory, image.gpu_addr, image.info, copies, download_span,
-                     swizzle_data_buffer);
-        download_map.offset += image.unswizzled_size_bytes;
-        download_span = download_span.subspan(image.unswizzled_size_bytes);
+    if constexpr (IMPLEMENTS_ASYNC_DOWNLOADS) {
+        const std::span<const ImageId> download_ids = committed_downloads.front();
+        if (download_ids.empty()) {
+            committed_downloads.pop_front();
+            async_buffers.pop_front();
+            return;
+        }
+        auto download_map = *async_buffers.front();
+        std::span<u8> download_span = download_map.mapped_span;
+        for (size_t i = download_ids.size(); i > 0; i--) {
+            const ImageBase& image = slot_images[download_ids[i - 1]];
+            const auto copies = FullDownloadCopies(image.info);
+            download_map.offset -= Common::AlignUp(image.unswizzled_size_bytes, 64);
+            std::span<u8> download_span_alt = download_span.subspan(download_map.offset);
+            SwizzleImage(*gpu_memory, image.gpu_addr, image.info, copies, download_span_alt,
+                         swizzle_data_buffer);
+        }
+        runtime.FreeDeferredStagingBuffer(download_map);
+        committed_downloads.pop_front();
+        async_buffers.pop_front();
+    } else {
+        const std::span<const ImageId> download_ids = committed_downloads.front();
+        if (download_ids.empty()) {
+            committed_downloads.pop_front();
+            return;
+        }
+        size_t total_size_bytes = 0;
+        for (const ImageId image_id : download_ids) {
+            total_size_bytes += slot_images[image_id].unswizzled_size_bytes;
+        }
+        auto download_map = runtime.DownloadStagingBuffer(total_size_bytes);
+        const size_t original_offset = download_map.offset;
+        for (const ImageId image_id : download_ids) {
+            Image& image = slot_images[image_id];
+            const auto copies = FullDownloadCopies(image.info);
+            image.DownloadMemory(download_map, copies);
+            download_map.offset += image.unswizzled_size_bytes;
+        }
+        // Wait for downloads to finish
+        runtime.Finish();
+        download_map.offset = original_offset;
+        std::span<u8> download_span = download_map.mapped_span;
+        for (const ImageId image_id : download_ids) {
+            const ImageBase& image = slot_images[image_id];
+            const auto copies = FullDownloadCopies(image.info);
+            SwizzleImage(*gpu_memory, image.gpu_addr, image.info, copies, download_span,
+                         swizzle_data_buffer);
+            download_map.offset += image.unswizzled_size_bytes;
+            download_span = download_span.subspan(image.unswizzled_size_bytes);
+        }
+        committed_downloads.pop_front();
     }
-    committed_downloads.pop();
 }
 
 template <class P>
@@ -740,7 +782,8 @@ void TextureCache<P>::UploadImageContents(Image& image, StagingBuffer& staging)
     const GPUVAddr gpu_addr = image.gpu_addr;
 
     if (True(image.flags & ImageFlagBits::AcceleratedUpload)) {
-        gpu_memory->ReadBlockUnsafe(gpu_addr, mapped_span.data(), mapped_span.size_bytes());
+        gpu_memory->ReadBlock(gpu_addr, mapped_span.data(), mapped_span.size_bytes(),
+                              VideoCommon::CacheType::NoTextureCache);
         const auto uploads = FullUploadSwizzles(image.info);
         runtime.AccelerateImageUpload(image, staging, uploads);
         return;
diff --git a/src/video_core/texture_cache/texture_cache_base.h b/src/video_core/texture_cache/texture_cache_base.h
index 4fd677a80..4eea1f609 100644
--- a/src/video_core/texture_cache/texture_cache_base.h
+++ b/src/video_core/texture_cache/texture_cache_base.h
@@ -92,6 +92,8 @@ class TextureCache : public VideoCommon::ChannelSetupCaches<TextureCacheChannelI
     static constexpr bool HAS_EMULATED_COPIES = P::HAS_EMULATED_COPIES;
     /// True when the API can provide info about the memory of the device.
     static constexpr bool HAS_DEVICE_MEMORY_INFO = P::HAS_DEVICE_MEMORY_INFO;
+    /// True when the API can do asynchronous texture downloads.
+    static constexpr bool IMPLEMENTS_ASYNC_DOWNLOADS = P::IMPLEMENTS_ASYNC_DOWNLOADS;
 
     static constexpr size_t UNSET_CHANNEL{std::numeric_limits<size_t>::max()};
 
@@ -106,6 +108,7 @@ class TextureCache : public VideoCommon::ChannelSetupCaches<TextureCacheChannelI
     using ImageView = typename P::ImageView;
     using Sampler = typename P::Sampler;
     using Framebuffer = typename P::Framebuffer;
+    using AsyncBuffer = typename P::AsyncBuffer;
 
     struct BlitImages {
         ImageId dst_id;
@@ -203,7 +206,7 @@ public:
     /// Create channel state.
     void CreateChannel(Tegra::Control::ChannelState& channel) final override;
 
-    std::mutex mutex;
+    std::recursive_mutex mutex;
 
 private:
     /// Iterate over all page indices in a range
@@ -403,7 +406,8 @@ private:
 
     // TODO: This data structure is not optimal and it should be reworked
     std::vector<ImageId> uncommitted_downloads;
-    std::queue<std::vector<ImageId>> committed_downloads;
+    std::deque<std::vector<ImageId>> committed_downloads;
+    std::deque<std::optional<AsyncBuffer>> async_buffers;
 
     struct LRUItemParams {
         using ObjectType = ImageId;
diff --git a/src/video_core/vulkan_common/vulkan_device.cpp b/src/video_core/vulkan_common/vulkan_device.cpp
index c4d31681a..5c5bfa18d 100644
--- a/src/video_core/vulkan_common/vulkan_device.cpp
+++ b/src/video_core/vulkan_common/vulkan_device.cpp
@@ -350,8 +350,8 @@ Device::Device(VkInstance instance_, vk::PhysicalDevice physical_, VkSurfaceKHR
             .sampleRateShading = true,
             .dualSrcBlend = true,
             .logicOp = true,
-            .multiDrawIndirect = false,
-            .drawIndirectFirstInstance = false,
+            .multiDrawIndirect = true,
+            .drawIndirectFirstInstance = true,
             .depthClamp = true,
             .depthBiasClamp = true,
             .fillModeNonSolid = true,
@@ -569,6 +569,67 @@ Device::Device(VkInstance instance_, vk::PhysicalDevice physical_, VkSurfaceKHR
         LOG_INFO(Render_Vulkan, "Device doesn't support extended dynamic state");
     }
 
+    VkPhysicalDeviceExtendedDynamicState2FeaturesEXT dynamic_state_2;
+    if (ext_extended_dynamic_state_2) {
+        dynamic_state_2 = {
+            .sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_EXTENDED_DYNAMIC_STATE_2_FEATURES_EXT,
+            .pNext = nullptr,
+            .extendedDynamicState2 = VK_TRUE,
+            .extendedDynamicState2LogicOp = ext_extended_dynamic_state_2_extra ? VK_TRUE : VK_FALSE,
+            .extendedDynamicState2PatchControlPoints = VK_FALSE,
+        };
+        SetNext(next, dynamic_state_2);
+    } else {
+        LOG_INFO(Render_Vulkan, "Device doesn't support extended dynamic state 2");
+    }
+
+    VkPhysicalDeviceExtendedDynamicState3FeaturesEXT dynamic_state_3;
+    if (ext_extended_dynamic_state_3) {
+        dynamic_state_3 = {
+            .sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_EXTENDED_DYNAMIC_STATE_3_FEATURES_EXT,
+            .pNext = nullptr,
+            .extendedDynamicState3TessellationDomainOrigin = VK_FALSE,
+            .extendedDynamicState3DepthClampEnable =
+                ext_extended_dynamic_state_3_enables ? VK_TRUE : VK_FALSE,
+            .extendedDynamicState3PolygonMode = VK_FALSE,
+            .extendedDynamicState3RasterizationSamples = VK_FALSE,
+            .extendedDynamicState3SampleMask = VK_FALSE,
+            .extendedDynamicState3AlphaToCoverageEnable = VK_FALSE,
+            .extendedDynamicState3AlphaToOneEnable = VK_FALSE,
+            .extendedDynamicState3LogicOpEnable =
+                ext_extended_dynamic_state_3_enables ? VK_TRUE : VK_FALSE,
+            .extendedDynamicState3ColorBlendEnable =
+                ext_extended_dynamic_state_3_blend ? VK_TRUE : VK_FALSE,
+            .extendedDynamicState3ColorBlendEquation =
+                ext_extended_dynamic_state_3_blend ? VK_TRUE : VK_FALSE,
+            .extendedDynamicState3ColorWriteMask =
+                ext_extended_dynamic_state_3_blend ? VK_TRUE : VK_FALSE,
+            .extendedDynamicState3RasterizationStream = VK_FALSE,
+            .extendedDynamicState3ConservativeRasterizationMode = VK_FALSE,
+            .extendedDynamicState3ExtraPrimitiveOverestimationSize = VK_FALSE,
+            .extendedDynamicState3DepthClipEnable = VK_FALSE,
+            .extendedDynamicState3SampleLocationsEnable = VK_FALSE,
+            .extendedDynamicState3ColorBlendAdvanced = VK_FALSE,
+            .extendedDynamicState3ProvokingVertexMode = VK_FALSE,
+            .extendedDynamicState3LineRasterizationMode = VK_FALSE,
+            .extendedDynamicState3LineStippleEnable = VK_FALSE,
+            .extendedDynamicState3DepthClipNegativeOneToOne = VK_FALSE,
+            .extendedDynamicState3ViewportWScalingEnable = VK_FALSE,
+            .extendedDynamicState3ViewportSwizzle = VK_FALSE,
+            .extendedDynamicState3CoverageToColorEnable = VK_FALSE,
+            .extendedDynamicState3CoverageToColorLocation = VK_FALSE,
+            .extendedDynamicState3CoverageModulationMode = VK_FALSE,
+            .extendedDynamicState3CoverageModulationTableEnable = VK_FALSE,
+            .extendedDynamicState3CoverageModulationTable = VK_FALSE,
+            .extendedDynamicState3CoverageReductionMode = VK_FALSE,
+            .extendedDynamicState3RepresentativeFragmentTestEnable = VK_FALSE,
+            .extendedDynamicState3ShadingRateImageEnable = VK_FALSE,
+        };
+        SetNext(next, dynamic_state_3);
+    } else {
+        LOG_INFO(Render_Vulkan, "Device doesn't support extended dynamic state 3");
+    }
+
     VkPhysicalDeviceLineRasterizationFeaturesEXT line_raster;
     if (ext_line_rasterization) {
         line_raster = {
@@ -695,6 +756,8 @@ Device::Device(VkInstance instance_, vk::PhysicalDevice physical_, VkSurfaceKHR
     CollectToolingInfo();
 
     if (driver_id == VK_DRIVER_ID_NVIDIA_PROPRIETARY_KHR) {
+        const u32 nv_major_version = (properties.driverVersion >> 22) & 0x3ff;
+
         const auto arch = GetNvidiaArchitecture(physical, supported_extensions);
         switch (arch) {
         case NvidiaArchitecture::AmpereOrNewer:
@@ -704,11 +767,13 @@ Device::Device(VkInstance instance_, vk::PhysicalDevice physical_, VkSurfaceKHR
         case NvidiaArchitecture::Turing:
             break;
         case NvidiaArchitecture::VoltaOrOlder:
-            LOG_WARNING(Render_Vulkan, "Blacklisting Volta and older from VK_KHR_push_descriptor");
-            khr_push_descriptor = false;
+            if (nv_major_version < 527) {
+                LOG_WARNING(Render_Vulkan,
+                            "Blacklisting Volta and older from VK_KHR_push_descriptor");
+                khr_push_descriptor = false;
+            }
             break;
         }
-        const u32 nv_major_version = (properties.driverVersion >> 22) & 0x3ff;
         if (nv_major_version >= 510) {
             LOG_WARNING(Render_Vulkan, "NVIDIA Drivers >= 510 do not support MSAA image blits");
             cant_blit_msaa = true;
@@ -735,6 +800,16 @@ Device::Device(VkInstance instance_, vk::PhysicalDevice physical_, VkSurfaceKHR
             ext_vertex_input_dynamic_state = false;
         }
     }
+    if (ext_extended_dynamic_state_2 && is_radv) {
+        const u32 version = (properties.driverVersion << 3) >> 3;
+        if (version < VK_MAKE_API_VERSION(0, 22, 3, 1)) {
+            LOG_WARNING(
+                Render_Vulkan,
+                "RADV versions older than 22.3.1 have broken VK_EXT_extended_dynamic_state2");
+            ext_extended_dynamic_state_2 = false;
+            ext_extended_dynamic_state_2_extra = false;
+        }
+    }
     sets_per_pool = 64;
 
     const bool is_amd =
@@ -763,8 +838,11 @@ Device::Device(VkInstance instance_, vk::PhysicalDevice physical_, VkSurfaceKHR
     const bool is_intel_windows = driver_id == VK_DRIVER_ID_INTEL_PROPRIETARY_WINDOWS;
     const bool is_intel_anv = driver_id == VK_DRIVER_ID_INTEL_OPEN_SOURCE_MESA;
     if (ext_vertex_input_dynamic_state && is_intel_windows) {
-        LOG_WARNING(Render_Vulkan, "Blacklisting Intel for VK_EXT_vertex_input_dynamic_state");
-        ext_vertex_input_dynamic_state = false;
+        const u32 version = (properties.driverVersion << 3) >> 3;
+        if (version < VK_MAKE_API_VERSION(27, 20, 100, 0)) {
+            LOG_WARNING(Render_Vulkan, "Blacklisting Intel for VK_EXT_vertex_input_dynamic_state");
+            ext_vertex_input_dynamic_state = false;
+        }
     }
     if (is_float16_supported && is_intel_windows) {
         // Intel's compiler crashes when using fp16 on Astral Chain, disable it for the time being.
@@ -1024,6 +1102,8 @@ void Device::CheckSuitability(bool requires_swapchain) const {
         std::make_pair(features.vertexPipelineStoresAndAtomics, "vertexPipelineStoresAndAtomics"),
         std::make_pair(features.imageCubeArray, "imageCubeArray"),
         std::make_pair(features.independentBlend, "independentBlend"),
+        std::make_pair(features.multiDrawIndirect, "multiDrawIndirect"),
+        std::make_pair(features.drawIndirectFirstInstance, "drawIndirectFirstInstance"),
         std::make_pair(features.depthClamp, "depthClamp"),
         std::make_pair(features.samplerAnisotropy, "samplerAnisotropy"),
         std::make_pair(features.largePoints, "largePoints"),
@@ -1089,6 +1169,8 @@ std::vector<const char*> Device::LoadExtensions(bool requires_surface) {
     bool has_ext_transform_feedback{};
     bool has_ext_custom_border_color{};
     bool has_ext_extended_dynamic_state{};
+    bool has_ext_extended_dynamic_state_2{};
+    bool has_ext_extended_dynamic_state_3{};
     bool has_ext_shader_atomic_int64{};
     bool has_ext_provoking_vertex{};
     bool has_ext_vertex_input_dynamic_state{};
@@ -1117,6 +1199,7 @@ std::vector<const char*> Device::LoadExtensions(bool requires_surface) {
         test(khr_spirv_1_4, VK_KHR_SPIRV_1_4_EXTENSION_NAME, true);
         test(khr_push_descriptor, VK_KHR_PUSH_DESCRIPTOR_EXTENSION_NAME, true);
         test(has_khr_shader_float16_int8, VK_KHR_SHADER_FLOAT16_INT8_EXTENSION_NAME, false);
+        test(khr_draw_indirect_count, VK_KHR_DRAW_INDIRECT_COUNT_EXTENSION_NAME, true);
         test(ext_depth_range_unrestricted, VK_EXT_DEPTH_RANGE_UNRESTRICTED_EXTENSION_NAME, true);
         test(ext_index_type_uint8, VK_EXT_INDEX_TYPE_UINT8_EXTENSION_NAME, true);
         test(has_ext_primitive_topology_list_restart,
@@ -1132,6 +1215,10 @@ std::vector<const char*> Device::LoadExtensions(bool requires_surface) {
         test(has_ext_transform_feedback, VK_EXT_TRANSFORM_FEEDBACK_EXTENSION_NAME, false);
         test(has_ext_custom_border_color, VK_EXT_CUSTOM_BORDER_COLOR_EXTENSION_NAME, false);
         test(has_ext_extended_dynamic_state, VK_EXT_EXTENDED_DYNAMIC_STATE_EXTENSION_NAME, false);
+        test(has_ext_extended_dynamic_state_2, VK_EXT_EXTENDED_DYNAMIC_STATE_2_EXTENSION_NAME,
+             false);
+        test(has_ext_extended_dynamic_state_3, VK_EXT_EXTENDED_DYNAMIC_STATE_3_EXTENSION_NAME,
+             false);
         test(has_ext_subgroup_size_control, VK_EXT_SUBGROUP_SIZE_CONTROL_EXTENSION_NAME, true);
         test(has_ext_provoking_vertex, VK_EXT_PROVOKING_VERTEX_EXTENSION_NAME, false);
         test(has_ext_vertex_input_dynamic_state, VK_EXT_VERTEX_INPUT_DYNAMIC_STATE_EXTENSION_NAME,
@@ -1281,6 +1368,44 @@ std::vector<const char*> Device::LoadExtensions(bool requires_surface) {
             ext_extended_dynamic_state = true;
         }
     }
+    if (has_ext_extended_dynamic_state_2) {
+        VkPhysicalDeviceExtendedDynamicState2FeaturesEXT extended_dynamic_state_2;
+        extended_dynamic_state_2.sType =
+            VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_EXTENDED_DYNAMIC_STATE_2_FEATURES_EXT;
+        extended_dynamic_state_2.pNext = nullptr;
+        features.pNext = &extended_dynamic_state_2;
+        physical.GetFeatures2(features);
+
+        if (extended_dynamic_state_2.extendedDynamicState2) {
+            extensions.push_back(VK_EXT_EXTENDED_DYNAMIC_STATE_2_EXTENSION_NAME);
+            ext_extended_dynamic_state_2 = true;
+            ext_extended_dynamic_state_2_extra =
+                extended_dynamic_state_2.extendedDynamicState2LogicOp;
+        }
+    }
+    if (has_ext_extended_dynamic_state_3) {
+        VkPhysicalDeviceExtendedDynamicState3FeaturesEXT extended_dynamic_state_3;
+        extended_dynamic_state_3.sType =
+            VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_EXTENDED_DYNAMIC_STATE_3_FEATURES_EXT;
+        extended_dynamic_state_3.pNext = nullptr;
+        features.pNext = &extended_dynamic_state_3;
+        physical.GetFeatures2(features);
+
+        ext_extended_dynamic_state_3_blend =
+            extended_dynamic_state_3.extendedDynamicState3ColorBlendEnable &&
+            extended_dynamic_state_3.extendedDynamicState3ColorBlendEquation &&
+            extended_dynamic_state_3.extendedDynamicState3ColorWriteMask;
+
+        ext_extended_dynamic_state_3_enables =
+            extended_dynamic_state_3.extendedDynamicState3DepthClampEnable &&
+            extended_dynamic_state_3.extendedDynamicState3LogicOpEnable;
+
+        ext_extended_dynamic_state_3 =
+            ext_extended_dynamic_state_3_blend || ext_extended_dynamic_state_3_enables;
+        if (ext_extended_dynamic_state_3) {
+            extensions.push_back(VK_EXT_EXTENDED_DYNAMIC_STATE_3_EXTENSION_NAME);
+        }
+    }
     if (has_ext_line_rasterization) {
         VkPhysicalDeviceLineRasterizationFeaturesEXT line_raster;
         line_raster.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_LINE_RASTERIZATION_FEATURES_EXT;
diff --git a/src/video_core/vulkan_common/vulkan_device.h b/src/video_core/vulkan_common/vulkan_device.h
index 6a26c4e6e..920a8f4e3 100644
--- a/src/video_core/vulkan_common/vulkan_device.h
+++ b/src/video_core/vulkan_common/vulkan_device.h
@@ -286,6 +286,30 @@ public:
         return ext_extended_dynamic_state;
     }
 
+    /// Returns true if the device supports VK_EXT_extended_dynamic_state2.
+    bool IsExtExtendedDynamicState2Supported() const {
+        return ext_extended_dynamic_state_2;
+    }
+
+    bool IsExtExtendedDynamicState2ExtrasSupported() const {
+        return ext_extended_dynamic_state_2_extra;
+    }
+
+    /// Returns true if the device supports VK_EXT_extended_dynamic_state3.
+    bool IsExtExtendedDynamicState3Supported() const {
+        return ext_extended_dynamic_state_3;
+    }
+
+    /// Returns true if the device supports VK_EXT_extended_dynamic_state3.
+    bool IsExtExtendedDynamicState3BlendingSupported() const {
+        return ext_extended_dynamic_state_3_blend;
+    }
+
+    /// Returns true if the device supports VK_EXT_extended_dynamic_state3.
+    bool IsExtExtendedDynamicState3EnablesSupported() const {
+        return ext_extended_dynamic_state_3_enables;
+    }
+
     /// Returns true if the device supports VK_EXT_line_rasterization.
     bool IsExtLineRasterizationSupported() const {
         return ext_line_rasterization;
@@ -451,6 +475,7 @@ private:
     bool nv_viewport_swizzle{};                 ///< Support for VK_NV_viewport_swizzle.
     bool nv_viewport_array2{};                  ///< Support for VK_NV_viewport_array2.
     bool nv_geometry_shader_passthrough{};      ///< Support for VK_NV_geometry_shader_passthrough.
+    bool khr_draw_indirect_count{};             ///< Support for VK_KHR_draw_indirect_count.
     bool khr_uniform_buffer_standard_layout{};  ///< Support for scalar uniform buffer layouts.
     bool khr_spirv_1_4{};                       ///< Support for VK_KHR_spirv_1_4.
     bool khr_workgroup_memory_explicit_layout{}; ///< Support for explicit workgroup layouts.
@@ -461,28 +486,33 @@ private:
     bool ext_sampler_filter_minmax{};            ///< Support for VK_EXT_sampler_filter_minmax.
     bool ext_depth_clip_control{};               ///< Support for VK_EXT_depth_clip_control
     bool ext_depth_range_unrestricted{};         ///< Support for VK_EXT_depth_range_unrestricted.
-    bool ext_shader_viewport_index_layer{}; ///< Support for VK_EXT_shader_viewport_index_layer.
-    bool ext_tooling_info{};                ///< Support for VK_EXT_tooling_info.
-    bool ext_subgroup_size_control{};       ///< Support for VK_EXT_subgroup_size_control.
-    bool ext_transform_feedback{};          ///< Support for VK_EXT_transform_feedback.
-    bool ext_custom_border_color{};         ///< Support for VK_EXT_custom_border_color.
-    bool ext_extended_dynamic_state{};      ///< Support for VK_EXT_extended_dynamic_state.
-    bool ext_line_rasterization{};          ///< Support for VK_EXT_line_rasterization.
-    bool ext_vertex_input_dynamic_state{};  ///< Support for VK_EXT_vertex_input_dynamic_state.
-    bool ext_shader_stencil_export{};       ///< Support for VK_EXT_shader_stencil_export.
-    bool ext_shader_atomic_int64{};         ///< Support for VK_KHR_shader_atomic_int64.
-    bool ext_conservative_rasterization{};  ///< Support for VK_EXT_conservative_rasterization.
-    bool ext_provoking_vertex{};            ///< Support for VK_EXT_provoking_vertex.
-    bool ext_memory_budget{};               ///< Support for VK_EXT_memory_budget.
-    bool nv_device_diagnostics_config{};    ///< Support for VK_NV_device_diagnostics_config.
-    bool has_broken_cube_compatibility{};   ///< Has broken cube compatiblity bit
-    bool has_renderdoc{};                   ///< Has RenderDoc attached
-    bool has_nsight_graphics{};             ///< Has Nsight Graphics attached
-    bool supports_d24_depth{};              ///< Supports D24 depth buffers.
-    bool cant_blit_msaa{};                  ///< Does not support MSAA<->MSAA blitting.
-    bool must_emulate_bgr565{};             ///< Emulates BGR565 by swizzling RGB565 format.
-    u32 max_vertex_input_attributes{};      ///< Max vertex input attributes in pipeline
-    u32 max_vertex_input_bindings{};        ///< Max vertex input buffers in pipeline
+    bool ext_shader_viewport_index_layer{};    ///< Support for VK_EXT_shader_viewport_index_layer.
+    bool ext_tooling_info{};                   ///< Support for VK_EXT_tooling_info.
+    bool ext_subgroup_size_control{};          ///< Support for VK_EXT_subgroup_size_control.
+    bool ext_transform_feedback{};             ///< Support for VK_EXT_transform_feedback.
+    bool ext_custom_border_color{};            ///< Support for VK_EXT_custom_border_color.
+    bool ext_extended_dynamic_state{};         ///< Support for VK_EXT_extended_dynamic_state.
+    bool ext_extended_dynamic_state_2{};       ///< Support for VK_EXT_extended_dynamic_state2.
+    bool ext_extended_dynamic_state_2_extra{}; ///< Support for VK_EXT_extended_dynamic_state2.
+    bool ext_extended_dynamic_state_3{};       ///< Support for VK_EXT_extended_dynamic_state3.
+    bool ext_extended_dynamic_state_3_blend{}; ///< Support for VK_EXT_extended_dynamic_state3.
+    bool ext_extended_dynamic_state_3_enables{}; ///< Support for VK_EXT_extended_dynamic_state3.
+    bool ext_line_rasterization{};               ///< Support for VK_EXT_line_rasterization.
+    bool ext_vertex_input_dynamic_state{};       ///< Support for VK_EXT_vertex_input_dynamic_state.
+    bool ext_shader_stencil_export{};            ///< Support for VK_EXT_shader_stencil_export.
+    bool ext_shader_atomic_int64{};              ///< Support for VK_KHR_shader_atomic_int64.
+    bool ext_conservative_rasterization{};       ///< Support for VK_EXT_conservative_rasterization.
+    bool ext_provoking_vertex{};                 ///< Support for VK_EXT_provoking_vertex.
+    bool ext_memory_budget{};                    ///< Support for VK_EXT_memory_budget.
+    bool nv_device_diagnostics_config{};         ///< Support for VK_NV_device_diagnostics_config.
+    bool has_broken_cube_compatibility{};        ///< Has broken cube compatiblity bit
+    bool has_renderdoc{};                        ///< Has RenderDoc attached
+    bool has_nsight_graphics{};                  ///< Has Nsight Graphics attached
+    bool supports_d24_depth{};                   ///< Supports D24 depth buffers.
+    bool cant_blit_msaa{};                       ///< Does not support MSAA<->MSAA blitting.
+    bool must_emulate_bgr565{};                  ///< Emulates BGR565 by swizzling RGB565 format.
+    u32 max_vertex_input_attributes{};           ///< Max vertex input attributes in pipeline
+    u32 max_vertex_input_bindings{};             ///< Max vertex input buffers in pipeline
 
     // Telemetry parameters
     std::string vendor_name;                       ///< Device's driver name.
diff --git a/src/video_core/vulkan_common/vulkan_wrapper.cpp b/src/video_core/vulkan_common/vulkan_wrapper.cpp
index 7dca7341c..861767c13 100644
--- a/src/video_core/vulkan_common/vulkan_wrapper.cpp
+++ b/src/video_core/vulkan_common/vulkan_wrapper.cpp
@@ -94,6 +94,10 @@ void Load(VkDevice device, DeviceDispatch& dld) noexcept {
     X(vkCmdDispatch);
     X(vkCmdDraw);
     X(vkCmdDrawIndexed);
+    X(vkCmdDrawIndirect);
+    X(vkCmdDrawIndexedIndirect);
+    X(vkCmdDrawIndirectCountKHR);
+    X(vkCmdDrawIndexedIndirectCountKHR);
     X(vkCmdEndQuery);
     X(vkCmdEndRenderPass);
     X(vkCmdEndTransformFeedbackEXT);
@@ -118,12 +122,22 @@ void Load(VkDevice device, DeviceDispatch& dld) noexcept {
     X(vkCmdSetDepthCompareOpEXT);
     X(vkCmdSetDepthTestEnableEXT);
     X(vkCmdSetDepthWriteEnableEXT);
+    X(vkCmdSetPrimitiveRestartEnableEXT);
+    X(vkCmdSetRasterizerDiscardEnableEXT);
+    X(vkCmdSetDepthBiasEnableEXT);
+    X(vkCmdSetLogicOpEnableEXT);
+    X(vkCmdSetDepthClampEnableEXT);
     X(vkCmdSetFrontFaceEXT);
+    X(vkCmdSetLogicOpEXT);
+    X(vkCmdSetPatchControlPointsEXT);
     X(vkCmdSetLineWidth);
     X(vkCmdSetPrimitiveTopologyEXT);
     X(vkCmdSetStencilOpEXT);
     X(vkCmdSetStencilTestEnableEXT);
     X(vkCmdSetVertexInputEXT);
+    X(vkCmdSetColorWriteMaskEXT);
+    X(vkCmdSetColorBlendEnableEXT);
+    X(vkCmdSetColorBlendEquationEXT);
     X(vkCmdResolveImage);
     X(vkCreateBuffer);
     X(vkCreateBufferView);
diff --git a/src/video_core/vulkan_common/vulkan_wrapper.h b/src/video_core/vulkan_common/vulkan_wrapper.h
index 8bd4fd4d9..accfad8c1 100644
--- a/src/video_core/vulkan_common/vulkan_wrapper.h
+++ b/src/video_core/vulkan_common/vulkan_wrapper.h
@@ -213,6 +213,10 @@ struct DeviceDispatch : InstanceDispatch {
     PFN_vkCmdDispatch vkCmdDispatch{};
     PFN_vkCmdDraw vkCmdDraw{};
     PFN_vkCmdDrawIndexed vkCmdDrawIndexed{};
+    PFN_vkCmdDrawIndirect vkCmdDrawIndirect{};
+    PFN_vkCmdDrawIndexedIndirect vkCmdDrawIndexedIndirect{};
+    PFN_vkCmdDrawIndirectCountKHR vkCmdDrawIndirectCountKHR{};
+    PFN_vkCmdDrawIndexedIndirectCountKHR vkCmdDrawIndexedIndirectCountKHR{};
     PFN_vkCmdEndDebugUtilsLabelEXT vkCmdEndDebugUtilsLabelEXT{};
     PFN_vkCmdEndQuery vkCmdEndQuery{};
     PFN_vkCmdEndRenderPass vkCmdEndRenderPass{};
@@ -230,8 +234,15 @@ struct DeviceDispatch : InstanceDispatch {
     PFN_vkCmdSetDepthCompareOpEXT vkCmdSetDepthCompareOpEXT{};
     PFN_vkCmdSetDepthTestEnableEXT vkCmdSetDepthTestEnableEXT{};
     PFN_vkCmdSetDepthWriteEnableEXT vkCmdSetDepthWriteEnableEXT{};
+    PFN_vkCmdSetPrimitiveRestartEnableEXT vkCmdSetPrimitiveRestartEnableEXT{};
+    PFN_vkCmdSetRasterizerDiscardEnableEXT vkCmdSetRasterizerDiscardEnableEXT{};
+    PFN_vkCmdSetDepthBiasEnableEXT vkCmdSetDepthBiasEnableEXT{};
+    PFN_vkCmdSetLogicOpEnableEXT vkCmdSetLogicOpEnableEXT{};
+    PFN_vkCmdSetDepthClampEnableEXT vkCmdSetDepthClampEnableEXT{};
     PFN_vkCmdSetEvent vkCmdSetEvent{};
     PFN_vkCmdSetFrontFaceEXT vkCmdSetFrontFaceEXT{};
+    PFN_vkCmdSetPatchControlPointsEXT vkCmdSetPatchControlPointsEXT{};
+    PFN_vkCmdSetLogicOpEXT vkCmdSetLogicOpEXT{};
     PFN_vkCmdSetLineWidth vkCmdSetLineWidth{};
     PFN_vkCmdSetPrimitiveTopologyEXT vkCmdSetPrimitiveTopologyEXT{};
     PFN_vkCmdSetScissor vkCmdSetScissor{};
@@ -242,6 +253,9 @@ struct DeviceDispatch : InstanceDispatch {
     PFN_vkCmdSetStencilWriteMask vkCmdSetStencilWriteMask{};
     PFN_vkCmdSetVertexInputEXT vkCmdSetVertexInputEXT{};
     PFN_vkCmdSetViewport vkCmdSetViewport{};
+    PFN_vkCmdSetColorWriteMaskEXT vkCmdSetColorWriteMaskEXT{};
+    PFN_vkCmdSetColorBlendEnableEXT vkCmdSetColorBlendEnableEXT{};
+    PFN_vkCmdSetColorBlendEquationEXT vkCmdSetColorBlendEquationEXT{};
     PFN_vkCmdWaitEvents vkCmdWaitEvents{};
     PFN_vkCreateBuffer vkCreateBuffer{};
     PFN_vkCreateBufferView vkCreateBufferView{};
@@ -1019,6 +1033,29 @@ public:
                               first_instance);
     }
 
+    void DrawIndirect(VkBuffer src_buffer, VkDeviceSize src_offset, u32 draw_count,
+                      u32 stride) const noexcept {
+        dld->vkCmdDrawIndirect(handle, src_buffer, src_offset, draw_count, stride);
+    }
+
+    void DrawIndexedIndirect(VkBuffer src_buffer, VkDeviceSize src_offset, u32 draw_count,
+                             u32 stride) const noexcept {
+        dld->vkCmdDrawIndexedIndirect(handle, src_buffer, src_offset, draw_count, stride);
+    }
+
+    void DrawIndirectCount(VkBuffer src_buffer, VkDeviceSize src_offset, VkBuffer count_buffer,
+                           VkDeviceSize count_offset, u32 draw_count, u32 stride) const noexcept {
+        dld->vkCmdDrawIndirectCountKHR(handle, src_buffer, src_offset, count_buffer, count_offset,
+                                       draw_count, stride);
+    }
+
+    void DrawIndexedIndirectCount(VkBuffer src_buffer, VkDeviceSize src_offset,
+                                  VkBuffer count_buffer, VkDeviceSize count_offset, u32 draw_count,
+                                  u32 stride) const noexcept {
+        dld->vkCmdDrawIndexedIndirectCountKHR(handle, src_buffer, src_offset, count_buffer,
+                                              count_offset, draw_count, stride);
+    }
+
     void ClearAttachments(Span<VkClearAttachment> attachments,
                           Span<VkClearRect> rects) const noexcept {
         dld->vkCmdClearAttachments(handle, attachments.size(), attachments.data(), rects.size(),
@@ -1192,10 +1229,51 @@ public:
         dld->vkCmdSetDepthWriteEnableEXT(handle, enable ? VK_TRUE : VK_FALSE);
     }
 
+    void SetPrimitiveRestartEnableEXT(bool enable) const noexcept {
+        dld->vkCmdSetPrimitiveRestartEnableEXT(handle, enable ? VK_TRUE : VK_FALSE);
+    }
+
+    void SetRasterizerDiscardEnableEXT(bool enable) const noexcept {
+        dld->vkCmdSetRasterizerDiscardEnableEXT(handle, enable ? VK_TRUE : VK_FALSE);
+    }
+
+    void SetDepthBiasEnableEXT(bool enable) const noexcept {
+        dld->vkCmdSetDepthBiasEnableEXT(handle, enable ? VK_TRUE : VK_FALSE);
+    }
+
+    void SetLogicOpEnableEXT(bool enable) const noexcept {
+        dld->vkCmdSetLogicOpEnableEXT(handle, enable ? VK_TRUE : VK_FALSE);
+    }
+
+    void SetDepthClampEnableEXT(bool enable) const noexcept {
+        dld->vkCmdSetDepthClampEnableEXT(handle, enable ? VK_TRUE : VK_FALSE);
+    }
+
     void SetFrontFaceEXT(VkFrontFace front_face) const noexcept {
         dld->vkCmdSetFrontFaceEXT(handle, front_face);
     }
 
+    void SetLogicOpEXT(VkLogicOp logic_op) const noexcept {
+        dld->vkCmdSetLogicOpEXT(handle, logic_op);
+    }
+
+    void SetPatchControlPointsEXT(uint32_t patch_control_points) const noexcept {
+        dld->vkCmdSetPatchControlPointsEXT(handle, patch_control_points);
+    }
+
+    void SetColorWriteMaskEXT(u32 first, Span<VkColorComponentFlags> masks) const noexcept {
+        dld->vkCmdSetColorWriteMaskEXT(handle, first, masks.size(), masks.data());
+    }
+
+    void SetColorBlendEnableEXT(u32 first, Span<VkBool32> enables) const noexcept {
+        dld->vkCmdSetColorBlendEnableEXT(handle, first, enables.size(), enables.data());
+    }
+
+    void SetColorBlendEquationEXT(u32 first,
+                                  Span<VkColorBlendEquationEXT> equations) const noexcept {
+        dld->vkCmdSetColorBlendEquationEXT(handle, first, equations.size(), equations.data());
+    }
+
     void SetLineWidth(float line_width) const noexcept {
         dld->vkCmdSetLineWidth(handle, line_width);
     }
diff --git a/src/yuzu/configuration/config.cpp b/src/yuzu/configuration/config.cpp
index 3e51426c8..e9425b5bd 100644
--- a/src/yuzu/configuration/config.cpp
+++ b/src/yuzu/configuration/config.cpp
@@ -562,6 +562,7 @@ void Config::ReadDebuggingValues() {
     ReadBasicSetting(Settings::values.reporting_services);
     ReadBasicSetting(Settings::values.quest_flag);
     ReadBasicSetting(Settings::values.disable_macro_jit);
+    ReadBasicSetting(Settings::values.disable_macro_hle);
     ReadBasicSetting(Settings::values.extended_logging);
     ReadBasicSetting(Settings::values.use_debug_asserts);
     ReadBasicSetting(Settings::values.use_auto_stub);
@@ -1198,6 +1199,7 @@ void Config::SaveDebuggingValues() {
     WriteBasicSetting(Settings::values.quest_flag);
     WriteBasicSetting(Settings::values.use_debug_asserts);
     WriteBasicSetting(Settings::values.disable_macro_jit);
+    WriteBasicSetting(Settings::values.disable_macro_hle);
     WriteBasicSetting(Settings::values.enable_all_controllers);
     WriteBasicSetting(Settings::values.create_crash_dumps);
     WriteBasicSetting(Settings::values.perform_vulkan_check);
diff --git a/src/yuzu/configuration/configure_debug.cpp b/src/yuzu/configuration/configure_debug.cpp
index dacc75a20..cbeb8f168 100644
--- a/src/yuzu/configuration/configure_debug.cpp
+++ b/src/yuzu/configuration/configure_debug.cpp
@@ -73,6 +73,8 @@ void ConfigureDebug::SetConfiguration() {
     ui->dump_macros->setChecked(Settings::values.dump_macros.GetValue());
     ui->disable_macro_jit->setEnabled(runtime_lock);
     ui->disable_macro_jit->setChecked(Settings::values.disable_macro_jit.GetValue());
+    ui->disable_macro_hle->setEnabled(runtime_lock);
+    ui->disable_macro_hle->setChecked(Settings::values.disable_macro_hle.GetValue());
     ui->disable_loop_safety_checks->setEnabled(runtime_lock);
     ui->disable_loop_safety_checks->setChecked(
         Settings::values.disable_shader_loop_safety_checks.GetValue());
@@ -117,6 +119,7 @@ void ConfigureDebug::ApplyConfiguration() {
     Settings::values.disable_shader_loop_safety_checks =
         ui->disable_loop_safety_checks->isChecked();
     Settings::values.disable_macro_jit = ui->disable_macro_jit->isChecked();
+    Settings::values.disable_macro_hle = ui->disable_macro_hle->isChecked();
     Settings::values.extended_logging = ui->extended_logging->isChecked();
     Settings::values.perform_vulkan_check = ui->perform_vulkan_check->isChecked();
     UISettings::values.disable_web_applet = ui->disable_web_applet->isChecked();
diff --git a/src/yuzu/configuration/configure_debug.ui b/src/yuzu/configuration/configure_debug.ui
index 102c8c66c..15acefe33 100644
--- a/src/yuzu/configuration/configure_debug.ui
+++ b/src/yuzu/configuration/configure_debug.ui
@@ -176,7 +176,7 @@
         </property>
        </widget>
       </item>
-      <item row="0" column="2">
+      <item row="1" column="2">
        <widget class="QCheckBox" name="dump_macros">
         <property name="enabled">
          <bool>true</bool>
@@ -202,6 +202,19 @@
         </property>
        </widget>
       </item>
+      <item row="0" column="2">
+       <widget class="QCheckBox" name="disable_macro_hle">
+        <property name="enabled">
+         <bool>true</bool>
+        </property>
+        <property name="toolTip">
+         <string>When checked, it disables the macro HLE functions. Enabling this makes games run slower</string>
+        </property>
+        <property name="text">
+         <string>Disable Macro HLE</string>
+        </property>
+       </widget>
+      </item>
       <item row="1" column="0">
        <widget class="QCheckBox" name="enable_shader_feedback">
         <property name="toolTip">
diff --git a/src/yuzu_cmd/config.cpp b/src/yuzu_cmd/config.cpp
index de9b220da..1e45e57bc 100644
--- a/src/yuzu_cmd/config.cpp
+++ b/src/yuzu_cmd/config.cpp
@@ -348,6 +348,7 @@ void Config::ReadValues() {
     ReadSetting("Debugging", Settings::values.use_debug_asserts);
     ReadSetting("Debugging", Settings::values.use_auto_stub);
     ReadSetting("Debugging", Settings::values.disable_macro_jit);
+    ReadSetting("Debugging", Settings::values.disable_macro_hle);
     ReadSetting("Debugging", Settings::values.use_gdbstub);
     ReadSetting("Debugging", Settings::values.gdbstub_port);