[vulkan] Changed info.color_output_type gate

Fix for current crash on master.
Just reverted only the necessary stuff so that PresentManager can hold a reference to khr and resist death upon application hold/restore.
@Lizzie shall judge.

Reviewed-on: https://git.eden-emu.dev/eden-emu/eden/pulls/3685
Co-authored-by: xbzk <xbzk@eden-emu.dev>
Co-committed-by: xbzk <xbzk@eden-emu.dev>

src/android/app/src/main/java/org/yuzu/yuzu_emu/features/fetcher/SpacingItemDecoration.kt

@@ -1,10 +1,11 @@
-// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
 // SPDX-License-Identifier: GPL-3.0-or-later
 
 package org.yuzu.yuzu_emu.features.fetcher
 
 import android.graphics.Rect
 import android.view.View
+import androidx.recyclerview.widget.GridLayoutManager
 import androidx.recyclerview.widget.RecyclerView
 
 class SpacingItemDecoration(private val spacing: Int) : RecyclerView.ItemDecoration() {
@@ -15,8 +16,20 @@ class SpacingItemDecoration(private val spacing: Int) : RecyclerView.ItemDecorat
         state: RecyclerView.State
     ) {
         outRect.bottom = spacing
-        if (parent.getChildAdapterPosition(view) == 0) {
+
+        val position = parent.getChildAdapterPosition(view)
+        if (position == RecyclerView.NO_POSITION) return
+
+        if (position == 0) {
             outRect.top = spacing
+            return
+        }
+
+        // If the item is in the first row, but NOT in first column add top spacing as well
+        val layoutManager = parent.layoutManager
+        if (layoutManager is GridLayoutManager && layoutManager.spanSizeLookup.getSpanGroupIndex(position, layoutManager.spanCount) == 0) {
+            outRect.top = spacing
+            return
         }
     }
 }

src/core/hle/service/prepo/prepo.cpp

@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
 // SPDX-License-Identifier: GPL-3.0-or-later
 
 // SPDX-FileCopyrightText: Copyright 2018 yuzu Emulator Project
@@ -28,8 +28,10 @@ public:
             {10101, &PlayReport::SaveReportWithUser<Core::Reporter::PlayReportType::Old>, "SaveReportWithUserOld"},
             {10102, &PlayReport::SaveReport<Core::Reporter::PlayReportType::Old2>, "SaveReportOld2"},
             {10103, &PlayReport::SaveReportWithUser<Core::Reporter::PlayReportType::Old2>, "SaveReportWithUserOld2"},
-            {10104, &PlayReport::SaveReport<Core::Reporter::PlayReportType::New>, "SaveReport"},
-            {10105, &PlayReport::SaveReportWithUser<Core::Reporter::PlayReportType::New>, "SaveReportWithUser"},
+            {10104, &PlayReport::SaveReport<Core::Reporter::PlayReportType::Old3>, "SaveReportOld3"},
+            {10105, &PlayReport::SaveReportWithUser<Core::Reporter::PlayReportType::Old3>, "SaveReportWithUserOld3"},
+            {10106, &PlayReport::SaveReport<Core::Reporter::PlayReportType::New>, "SaveReport"},
+            {10107, &PlayReport::SaveReportWithUser<Core::Reporter::PlayReportType::New>, "SaveReportWithUser"},
             {10200, &PlayReport::RequestImmediateTransmission, "RequestImmediateTransmission"},
             {10300, &PlayReport::GetTransmissionStatus, "GetTransmissionStatus"},
             {10400, &PlayReport::GetSystemSessionId, "GetSystemSessionId"},

src/core/reporter.h

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2019 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
@@ -53,6 +56,7 @@ public:
     enum class PlayReportType {
         Old,
         Old2,
+        Old3,
         New,
         System,
     };

src/qt_common/config/shared_translation.cpp

@@ -414,6 +414,9 @@ std::unique_ptr<TranslationMap> InitializeTranslations(QObject* parent)
               "their resolution, details and supported controllers and depending on this setting.\n"
               "Setting to Handheld can help improve performance for low end systems."));
     INSERT(Settings, current_user, QString(), QString());
+    INSERT(Settings, serial_unit, tr("Unit Serial"), QString());
+    INSERT(Settings, serial_battery, tr("Battery Serial"), QString());
+    INSERT(Settings, debug_knobs, tr("Debug knobs"), QString());
 
     // Controls
 

src/shader_recompiler/backend/spirv/emit_spirv.cpp

@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
 // SPDX-License-Identifier: GPL-3.0-or-later
 
 // SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
@@ -11,15 +11,159 @@
 #include <vector>
 #include <spirv-tools/optimizer.hpp>
 
+#include "common/logging/log.h"
 #include "common/settings.h"
 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
 #include "shader_recompiler/backend/spirv/spirv_emit_context.h"
 #include "shader_recompiler/frontend/ir/basic_block.h"
+#include "shader_recompiler/frontend/ir/modifiers.h"
 #include "shader_recompiler/frontend/ir/program.h"
 
 namespace Shader::Backend::SPIRV {
 namespace {
+[[nodiscard]] constexpr std::string_view StageName(Stage stage) noexcept {
+    switch (stage) {
+    case Stage::VertexA:
+        return "VertexA";
+    case Stage::VertexB:
+        return "VertexB";
+    case Stage::TessellationControl:
+        return "TessellationControl";
+    case Stage::TessellationEval:
+        return "TessellationEval";
+    case Stage::Geometry:
+        return "Geometry";
+    case Stage::Fragment:
+        return "Fragment";
+    case Stage::Compute:
+        return "Compute";
+    }
+    return "Unknown";
+}
+
+[[nodiscard]] constexpr std::string_view DenormModeName(bool flush, bool preserve) noexcept {
+    if (flush && preserve) {
+        return "Flush+Preserve";
+    }
+    if (flush) {
+        return "Flush";
+    }
+    if (preserve) {
+        return "Preserve";
+    }
+    return "None";
+}
+
+[[nodiscard]] constexpr bool IsFp32RoundingRelevantOpcode(IR::Opcode opcode) noexcept {
+    switch (opcode) {
+    case IR::Opcode::FPAdd32:
+    case IR::Opcode::FPFma32:
+    case IR::Opcode::FPMul32:
+    case IR::Opcode::FPRoundEven32:
+    case IR::Opcode::FPFloor32:
+    case IR::Opcode::FPCeil32:
+    case IR::Opcode::FPTrunc32:
+    case IR::Opcode::FPOrdEqual32:
+    case IR::Opcode::FPUnordEqual32:
+    case IR::Opcode::FPOrdNotEqual32:
+    case IR::Opcode::FPUnordNotEqual32:
+    case IR::Opcode::FPOrdLessThan32:
+    case IR::Opcode::FPUnordLessThan32:
+    case IR::Opcode::FPOrdGreaterThan32:
+    case IR::Opcode::FPUnordGreaterThan32:
+    case IR::Opcode::FPOrdLessThanEqual32:
+    case IR::Opcode::FPUnordLessThanEqual32:
+    case IR::Opcode::FPOrdGreaterThanEqual32:
+    case IR::Opcode::FPUnordGreaterThanEqual32:
+    case IR::Opcode::ConvertF16F32:
+    case IR::Opcode::ConvertF64F32:
+        return true;
+    default:
+        return false;
+    }
+}
+
+struct Fp32RoundingUsage {
+    u32 rz_count{};
+    bool has_conflicting_rounding{};
+};
+
+Fp32RoundingUsage CollectFp32RoundingUsage(const IR::Program& program) {
+    Fp32RoundingUsage usage{};
+    for (const IR::Block* const block : program.post_order_blocks) {
+        for (const IR::Inst& inst : block->Instructions()) {
+            if (!IsFp32RoundingRelevantOpcode(inst.GetOpcode())) {
+                continue;
+            }
+            switch (inst.Flags<IR::FpControl>().rounding) {
+            case IR::FpRounding::RZ:
+                ++usage.rz_count;
+                break;
+            case IR::FpRounding::RN:
+            case IR::FpRounding::RM:
+            case IR::FpRounding::RP:
+                usage.has_conflicting_rounding = true;
+                break;
+            case IR::FpRounding::DontCare:
+                break;
+            }
+        }
+    }
+    return usage;
+}
+
+void LogRzBackendSummary(const Profile& profile, const IR::Program& program, bool optimize) {
+    if (!Settings::values.renderer_debug) {
+        return;
+    }
+    const Fp32RoundingUsage usage{CollectFp32RoundingUsage(program)};
+    if (usage.rz_count == 0) {
+        return;
+    }
+
+    LOG_INFO(Shader_SPIRV,
+             "SPV_RZ {} start={:#010x} optimize={} support_float_controls={} separate_denorm_behavior={} separate_rounding_mode={} support_fp32_rounding_rtz={} broken_fp16_float_controls={} fp16_denorm={} fp32_denorm={} signed_nan16={} signed_nan32={} signed_nan64={} rz_inst_count={} mixed_fp32_rounding={}",
+             StageName(program.stage), program.start_address, optimize,
+             profile.support_float_controls, profile.support_separate_denorm_behavior,
+             profile.support_separate_rounding_mode, profile.support_fp32_rounding_rtz,
+             profile.has_broken_fp16_float_controls,
+             DenormModeName(program.info.uses_fp16_denorms_flush,
+                            program.info.uses_fp16_denorms_preserve),
+             DenormModeName(program.info.uses_fp32_denorms_flush,
+                            program.info.uses_fp32_denorms_preserve),
+             profile.support_fp16_signed_zero_nan_preserve,
+             profile.support_fp32_signed_zero_nan_preserve,
+             profile.support_fp64_signed_zero_nan_preserve, usage.rz_count,
+             usage.has_conflicting_rounding);
+}
+
+void SetupRoundingControl(const Profile& profile, const IR::Program& program, EmitContext& ctx,
+                         Id main_func) {
+    const Fp32RoundingUsage usage{CollectFp32RoundingUsage(program)};
+    if (usage.rz_count == 0) {
+        return;
+    }
+    if (usage.has_conflicting_rounding) {
+        if (Settings::values.renderer_debug) {
+            LOG_INFO(Shader_SPIRV,
+                     "SPV_RZ {} start={:#010x} skipping_fp32_rtz_execution_mode reason=mixed_rounding",
+                     StageName(program.stage), program.start_address);
+        }
+        return;
+    }
+    if (!profile.support_fp32_rounding_rtz) {
+        if (Settings::values.renderer_debug) {
+            LOG_INFO(Shader_SPIRV,
+                     "SPV_RZ {} start={:#010x} skipping_fp32_rtz_execution_mode reason=unsupported_fp32_rtz",
+                     StageName(program.stage), program.start_address);
+        }
+        return;
+    }
+    ctx.AddCapability(spv::Capability::RoundingModeRTZ);
+    ctx.AddExecutionMode(main_func, spv::ExecutionMode::RoundingModeRTZ, 32U);
+}
+
 template <class Func>
 struct FuncTraits {};
     thread_local std::unique_ptr<spvtools::Optimizer> thread_optimizer;
@@ -503,12 +647,14 @@ void PatchPhiNodes(IR::Program& program, EmitContext& ctx) {
 
 std::vector<u32> EmitSPIRV(const Profile& profile, const RuntimeInfo& runtime_info,
                            IR::Program& program, Bindings& bindings, bool optimize) {
+    LogRzBackendSummary(profile, program, optimize);
     EmitContext ctx{profile, runtime_info, program, bindings};
     const Id main{DefineMain(ctx, program)};
     DefineEntryPoint(program, ctx, main);
     if (profile.support_float_controls) {
         ctx.AddExtension("SPV_KHR_float_controls");
         SetupDenormControl(profile, program, ctx, main);
+        SetupRoundingControl(profile, program, ctx, main);
         SetupSignedNanCapabilities(profile, program, ctx, main);
     }
     SetupCapabilities(profile, program.info, ctx);
@@ -516,6 +662,12 @@ std::vector<u32> EmitSPIRV(const Profile& profile, const RuntimeInfo& runtime_in
     PatchPhiNodes(program, ctx);
 
     if (!optimize) {
+        if (Settings::values.renderer_debug && ctx.log_rz_fp_controls) {
+            const std::vector<u32> spirv{ctx.Assemble()};
+            LOG_INFO(Shader_SPIRV, "SPV_RZ {} start={:#010x} assembled_words={} optimized_words={} validator_run=false",
+                     StageName(program.stage), program.start_address, spirv.size(), spirv.size());
+            return spirv;
+        }
         return ctx.Assemble();
     } else {
         std::vector<u32> spirv = ctx.Assemble();
@@ -535,6 +687,11 @@ std::vector<u32> EmitSPIRV(const Profile& profile, const RuntimeInfo& runtime_in
                       "Failed to optimize SPIRV shader output, continuing without optimization");
             result = std::move(spirv);
         }
+        if (Settings::values.renderer_debug && ctx.log_rz_fp_controls) {
+            LOG_INFO(Shader_SPIRV,
+                     "SPV_RZ {} start={:#010x} assembled_words={} optimized_words={} validator_run=false",
+                     StageName(program.stage), program.start_address, spirv.size(), result.size());
+        }
         return result;
     }
 }

src/shader_recompiler/backend/spirv/emit_spirv_floating_point.cpp

@@ -1,16 +1,60 @@
-// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
 // SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
+#include "common/logging/log.h"
+#include "common/settings.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
 #include "shader_recompiler/backend/spirv/spirv_emit_context.h"
 #include "shader_recompiler/frontend/ir/modifiers.h"
 
 namespace Shader::Backend::SPIRV {
 namespace {
+[[nodiscard]] constexpr std::string_view StageName(Stage stage) noexcept {
+    switch (stage) {
+    case Stage::VertexA:
+        return "VertexA";
+    case Stage::VertexB:
+        return "VertexB";
+    case Stage::TessellationControl:
+        return "TessellationControl";
+    case Stage::TessellationEval:
+        return "TessellationEval";
+    case Stage::Geometry:
+        return "Geometry";
+    case Stage::Fragment:
+        return "Fragment";
+    case Stage::Compute:
+        return "Compute";
+    }
+    return "Unknown";
+}
+
+[[nodiscard]] constexpr std::string_view FmzName(IR::FmzMode fmz_mode) noexcept {
+    switch (fmz_mode) {
+    case IR::FmzMode::DontCare:
+        return "DontCare";
+    case IR::FmzMode::FTZ:
+        return "FTZ";
+    case IR::FmzMode::FMZ:
+        return "FMZ";
+    case IR::FmzMode::None:
+        return "None";
+    }
+    return "Unknown";
+}
+
 Id Decorate(EmitContext& ctx, IR::Inst* inst, Id op) {
     const auto flags{inst->Flags<IR::FpControl>()};
+    if (Settings::values.renderer_debug && ctx.log_rz_fp_controls &&
+        flags.rounding == IR::FpRounding::RZ) {
+        LOG_INFO(Shader_SPIRV,
+                 "SPV_RZ_EMIT {} start={:#010x} ir_opcode={} spirv_op=OpFMul no_contraction={} fmz={} float_controls_ext={}",
+                 StageName(ctx.stage), ctx.start_address, inst->GetOpcode(),
+                 flags.no_contraction, FmzName(flags.fmz_mode), ctx.profile.support_float_controls);
+    }
     if (flags.no_contraction) {
         ctx.Decorate(op, spv::Decoration::NoContraction);
     }

src/shader_recompiler/backend/spirv/spirv_emit_context.cpp

@@ -17,6 +17,7 @@
 #include "common/div_ceil.h"
 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/spirv_emit_context.h"
+#include "shader_recompiler/frontend/ir/modifiers.h"
 
 namespace Shader::Backend::SPIRV {
 namespace {
@@ -473,7 +474,44 @@ void VectorTypes::Define(Sirit::Module& sirit_ctx, Id base_type, std::string_vie
 EmitContext::EmitContext(const Profile& profile_, const RuntimeInfo& runtime_info_,
                          IR::Program& program, Bindings& bindings)
     : Sirit::Module(profile_.supported_spirv), profile{profile_}, runtime_info{runtime_info_},
-      stage{program.stage}, texture_rescaling_index{bindings.texture_scaling_index},
+      stage{program.stage}, start_address{program.start_address},
+      log_rz_fp_controls{std::ranges::any_of(program.post_order_blocks, [](const IR::Block* block) {
+          return std::ranges::any_of(block->Instructions(), [](const IR::Inst& inst) {
+              switch (inst.GetOpcode()) {
+              case IR::Opcode::FPAdd16:
+              case IR::Opcode::FPFma16:
+              case IR::Opcode::FPMul16:
+              case IR::Opcode::FPRoundEven16:
+              case IR::Opcode::FPFloor16:
+              case IR::Opcode::FPCeil16:
+              case IR::Opcode::FPTrunc16:
+              case IR::Opcode::FPAdd32:
+              case IR::Opcode::FPFma32:
+              case IR::Opcode::FPMul32:
+              case IR::Opcode::FPRoundEven32:
+              case IR::Opcode::FPFloor32:
+              case IR::Opcode::FPCeil32:
+              case IR::Opcode::FPTrunc32:
+              case IR::Opcode::FPOrdEqual32:
+              case IR::Opcode::FPUnordEqual32:
+              case IR::Opcode::FPOrdNotEqual32:
+              case IR::Opcode::FPUnordNotEqual32:
+              case IR::Opcode::FPOrdLessThan32:
+              case IR::Opcode::FPUnordLessThan32:
+              case IR::Opcode::FPOrdGreaterThan32:
+              case IR::Opcode::FPUnordGreaterThan32:
+              case IR::Opcode::FPOrdLessThanEqual32:
+              case IR::Opcode::FPUnordLessThanEqual32:
+              case IR::Opcode::FPOrdGreaterThanEqual32:
+              case IR::Opcode::FPUnordGreaterThanEqual32:
+              case IR::Opcode::ConvertF16F32:
+              case IR::Opcode::ConvertF64F32:
+                  return inst.Flags<IR::FpControl>().rounding == IR::FpRounding::RZ;
+              default:
+                  return false;
+              }
+          });
+      })}, texture_rescaling_index{bindings.texture_scaling_index},
       image_rescaling_index{bindings.image_scaling_index} {
     const bool is_unified{profile.unified_descriptor_binding};
     u32& uniform_binding{is_unified ? bindings.unified : bindings.uniform_buffer};

src/shader_recompiler/backend/spirv/spirv_emit_context.h

@@ -216,6 +216,8 @@ public:
     const Profile& profile;
     const RuntimeInfo& runtime_info;
     Stage stage{};
+    u32 start_address{};
+    bool log_rz_fp_controls{};
 
     Id void_id{};
     Id U1{};

src/shader_recompiler/frontend/ir/program.h

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
@@ -20,6 +23,7 @@ struct Program {
     BlockList post_order_blocks;
     Info info;
     Stage stage{};
+    u32 start_address{};
     std::array<u32, 3> workgroup_size{};
     OutputTopology output_topology{};
     u32 output_vertices{};

src/shader_recompiler/frontend/maxwell/translate_program.cpp

@@ -458,6 +458,7 @@ IR::Program TranslateProgram(ObjectPool<IR::Inst>& inst_pool, ObjectPool<IR::Blo
     program.blocks = GenerateBlocks(program.syntax_list);
     program.post_order_blocks = PostOrder(program.syntax_list.front());
     program.stage = env.ShaderStage();
+    program.start_address = env.StartAddress();
     program.local_memory_size = env.LocalMemorySize();
     switch (program.stage) {
     case Stage::TessellationControl: {
@@ -554,6 +555,7 @@ IR::Program MergeDualVertexPrograms(IR::Program& vertex_a, IR::Program& vertex_b
         result.post_order_blocks.push_back(block);
     }
     result.stage = Stage::VertexB;
+    result.start_address = env_vertex_b.StartAddress();
     result.info = vertex_a.info;
     result.local_memory_size = (std::max)(vertex_a.local_memory_size, vertex_b.local_memory_size);
     result.info.loads.mask |= vertex_b.info.loads.mask;

src/shader_recompiler/profile.h

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
@@ -18,6 +21,7 @@ struct Profile {
     bool support_float_controls{};
     bool support_separate_denorm_behavior{};
     bool support_separate_rounding_mode{};
+    bool support_fp32_rounding_rtz{};
     bool support_fp16_denorm_preserve{};
     bool support_fp32_denorm_preserve{};
     bool support_fp16_denorm_flush{};

src/video_core/buffer_cache/buffer_cache_base.h

@@ -14,9 +14,12 @@
 #include <mutex>
 #include <numeric>
 #include <span>
-#include <ankerl/unordered_dense.h>
 #include <vector>
 
+#include <ankerl/unordered_dense.h>
+#include <boost/container/static_vector.hpp>
+#include <boost/container/small_vector.hpp>
+
 #include "common/common_types.h"
 #include "common/div_ceil.h"
 #include "common/literals.h"
@@ -94,10 +97,10 @@ static constexpr Binding NULL_BINDING{
 
 template <typename Buffer>
 struct HostBindings {
-    boost::container::small_vector<Buffer*, NUM_VERTEX_BUFFERS> buffers;
-    boost::container::small_vector<u64, NUM_VERTEX_BUFFERS> offsets;
-    boost::container::small_vector<u64, NUM_VERTEX_BUFFERS> sizes;
-    boost::container::small_vector<u64, NUM_VERTEX_BUFFERS> strides;
+    boost::container::static_vector<Buffer*, NUM_VERTEX_BUFFERS> buffers;
+    boost::container::static_vector<u64, NUM_VERTEX_BUFFERS> offsets;
+    boost::container::static_vector<u64, NUM_VERTEX_BUFFERS> sizes;
+    boost::container::static_vector<u64, NUM_VERTEX_BUFFERS> strides;
     u32 min_index{NUM_VERTEX_BUFFERS};
     u32 max_index{0};
 };

src/video_core/control/channel_state.cpp

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: 2022 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-3.0-or-later
 
@@ -19,12 +22,12 @@ ChannelState::ChannelState(s32 bind_id_) : bind_id{bind_id_}, initialized{} {}
 void ChannelState::Init(Core::System& system, GPU& gpu, u64 program_id_) {
     ASSERT(memory_manager);
     program_id = program_id_;
-    dma_pusher = std::make_unique<Tegra::DmaPusher>(system, gpu, *memory_manager, *this);
-    maxwell_3d = std::make_unique<Engines::Maxwell3D>(system, *memory_manager);
-    fermi_2d = std::make_unique<Engines::Fermi2D>(*memory_manager);
-    kepler_compute = std::make_unique<Engines::KeplerCompute>(system, *memory_manager);
-    maxwell_dma = std::make_unique<Engines::MaxwellDMA>(system, *memory_manager);
-    kepler_memory = std::make_unique<Engines::KeplerMemory>(system, *memory_manager);
+    dma_pusher.emplace(system, gpu, *memory_manager, *this);
+    maxwell_3d.emplace(system, *memory_manager);
+    fermi_2d.emplace(*memory_manager);
+    kepler_compute.emplace(system, *memory_manager);
+    maxwell_dma.emplace(system, *memory_manager);
+    kepler_memory.emplace(system, *memory_manager);
     initialized = true;
 }
 

src/video_core/control/channel_state.h

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: 2022 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-3.0-or-later
 
@@ -6,6 +9,12 @@
 #include <memory>
 
 #include "common/common_types.h"
+#include "video_core/engines/fermi_2d.h"
+#include "video_core/engines/kepler_memory.h"
+#include "video_core/engines/kepler_compute.h"
+#include "video_core/engines/maxwell_3d.h"
+#include "video_core/engines/maxwell_dma.h"
+#include "video_core/dma_pusher.h"
 
 namespace Core {
 class System;
@@ -18,49 +27,34 @@ class RasterizerInterface;
 namespace Tegra {
 
 class GPU;
-
-namespace Engines {
-class Puller;
-class Fermi2D;
-class Maxwell3D;
-class MaxwellDMA;
-class KeplerCompute;
-class KeplerMemory;
-} // namespace Engines
-
 class MemoryManager;
-class DmaPusher;
 
 namespace Control {
 
 struct ChannelState {
     explicit ChannelState(s32 bind_id);
-    ChannelState(const ChannelState& state) = delete;
-    ChannelState& operator=(const ChannelState&) = delete;
-    ChannelState(ChannelState&& other) noexcept = default;
-    ChannelState& operator=(ChannelState&& other) noexcept = default;
 
     void Init(Core::System& system, GPU& gpu, u64 program_id);
 
     void BindRasterizer(VideoCore::RasterizerInterface* rasterizer);
 
-    s32 bind_id = -1;
-    u64 program_id = 0;
     /// 3D engine
-    std::unique_ptr<Engines::Maxwell3D> maxwell_3d;
+    std::optional<Engines::Maxwell3D> maxwell_3d;
     /// 2D engine
-    std::unique_ptr<Engines::Fermi2D> fermi_2d;
+    std::optional<Engines::Fermi2D> fermi_2d;
     /// Compute engine
-    std::unique_ptr<Engines::KeplerCompute> kepler_compute;
+    std::optional<Engines::KeplerCompute> kepler_compute;
     /// DMA engine
-    std::unique_ptr<Engines::MaxwellDMA> maxwell_dma;
+    std::optional<Engines::MaxwellDMA> maxwell_dma;
     /// Inline memory engine
-    std::unique_ptr<Engines::KeplerMemory> kepler_memory;
-
+    std::optional<Engines::KeplerMemory> kepler_memory;
+    /// NV01 Timer
+    std::optional<Engines::KeplerMemory> nv01_timer;
+    std::optional<DmaPusher> dma_pusher;
     std::shared_ptr<MemoryManager> memory_manager;
 
-    std::unique_ptr<DmaPusher> dma_pusher;
-
+    s32 bind_id = -1;
+    u64 program_id = 0;
     bool initialized{};
 };
 

src/video_core/engines/engine_interface.h

@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
 // SPDX-License-Identifier: GPL-3.0-or-later
 
 // SPDX-FileCopyrightText: Copyright 2020 yuzu Emulator Project
@@ -15,6 +15,7 @@
 namespace Tegra::Engines {
 
 enum class EngineTypes : u32 {
+    Nv01Timer,
     KeplerCompute,
     Maxwell3D,
     Fermi2D,

src/video_core/engines/maxwell_3d.cpp

@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
 // SPDX-License-Identifier: GPL-3.0-or-later
 
 // SPDX-FileCopyrightText: Copyright 2018 yuzu Emulator Project
@@ -26,8 +26,15 @@ namespace Tegra::Engines {
 constexpr u32 MacroRegistersStart = 0xE00;
 
 Maxwell3D::Maxwell3D(Core::System& system_, MemoryManager& memory_manager_)
-    : draw_manager{std::make_unique<DrawManager>(this)}, system{system_},
-      memory_manager{memory_manager_}, macro_engine{GetMacroEngine(*this)}, upload_state{memory_manager, regs.upload} {
+    : draw_manager{std::make_unique<DrawManager>(this)}, system{system_}
+    , memory_manager{memory_manager_}
+#ifdef ARCHITECTURE_x86_64
+    , macro_engine(bool(Settings::values.disable_macro_jit))
+#else
+    , macro_engine(true)
+#endif
+    , upload_state{memory_manager, regs.upload}
+{
     dirty.flags.flip();
     InitializeRegisterDefaults();
     execution_mask.reset();
@@ -328,9 +335,9 @@ void Maxwell3D::ProcessMethodCall(u32 method, u32 argument, u32 nonshadow_argume
         shadow_state.shadow_ram_control = static_cast<Regs::ShadowRamControl>(nonshadow_argument);
         return;
     case MAXWELL3D_REG_INDEX(load_mme.instruction_ptr):
-        return macro_engine->ClearCode(regs.load_mme.instruction_ptr);
+        return macro_engine.ClearCode(regs.load_mme.instruction_ptr);
     case MAXWELL3D_REG_INDEX(load_mme.instruction):
-        return macro_engine->AddCode(regs.load_mme.instruction_ptr, argument);
+        return macro_engine.AddCode(regs.load_mme.instruction_ptr, argument);
     case MAXWELL3D_REG_INDEX(load_mme.start_address):
         return ProcessMacroBind(argument);
     case MAXWELL3D_REG_INDEX(falcon[4]):
@@ -398,7 +405,7 @@ void Maxwell3D::CallMacroMethod(u32 method, const std::vector<u32>& parameters)
         ((method - MacroRegistersStart) >> 1) % static_cast<u32>(macro_positions.size());
 
     // Execute the current macro.
-    macro_engine->Execute(macro_positions[entry], parameters);
+    macro_engine.Execute(*this, macro_positions[entry], parameters);
 
     draw_manager->DrawDeferred();
 }
@@ -464,7 +471,7 @@ void Maxwell3D::CallMultiMethod(u32 method, const u32* base_start, u32 amount,
 }
 
 void Maxwell3D::ProcessMacroUpload(u32 data) {
-    macro_engine->AddCode(regs.load_mme.instruction_ptr++, data);
+    macro_engine.AddCode(regs.load_mme.instruction_ptr++, data);
 }
 
 void Maxwell3D::ProcessMacroBind(u32 data) {

src/video_core/engines/maxwell_3d.h

@@ -2258,7 +2258,7 @@ public:
             /// Returns whether the vertex array specified by index is supposed to be
             /// accessed per instance or not.
             bool IsInstancingEnabled(std::size_t index) const {
-                return is_instanced[index];
+                return bool(is_instanced[index]); //FUCK YOU MSVC
             }
         };
 
@@ -3203,7 +3203,7 @@ private:
     std::vector<u32> macro_params;
 
     /// Interpreter for the macro codes uploaded to the GPU.
-    std::optional<MacroEngine> macro_engine;
+    MacroEngine macro_engine;
 
     Upload::State upload_state;
 

src/video_core/engines/nv01_timer.h

@@ -0,0 +1,52 @@
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
+// SPDX-FileCopyrightText: Copyright 2018 yuzu Emulator Project
+// SPDX-License-Identifier: GPL-2.0-or-later
+
+#pragma once
+
+#include <array>
+#include <cstddef>
+#include "common/bit_field.h"
+#include "common/common_funcs.h"
+#include "common/common_types.h"
+#include "common/logging/log.h"
+#include "video_core/engines/engine_interface.h"
+#include "video_core/engines/engine_upload.h"
+
+namespace Core {
+class System;
+}
+
+namespace Tegra {
+class MemoryManager;
+}
+
+namespace Tegra::Engines {
+class Nv01Timer final : public EngineInterface {
+public:
+    explicit Nv01Timer(Core::System& system_, MemoryManager& memory_manager)
+        : system{system_}
+    {}
+    ~Nv01Timer() override;
+
+    /// Write the value to the register identified by method.
+    void CallMethod(u32 method, u32 method_argument, bool is_last_call) override {
+        LOG_DEBUG(HW_GPU, "method={}, argument={}, is_last_call={}", method, method_argument, is_last_call);
+    }
+
+    /// Write multiple values to the register identified by method.
+    void CallMultiMethod(u32 method, const u32* base_start, u32 amount, u32 methods_pending) override {
+        LOG_DEBUG(HW_GPU, "method={}, base_start={}, amount={}, pending={}", method, fmt::ptr(base_start), amount, methods_pending);
+    }
+
+    struct Regs {
+        // No fucking idea
+        INSERT_PADDING_BYTES_NOINIT(0x48);
+    } regs{};
+private:
+    void ConsumeSinkImpl() override {}
+    Core::System& system;
+};
+}

src/video_core/engines/puller.cpp

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: 2022 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-3.0-or-later
 
@@ -34,24 +37,22 @@ void Puller::ProcessBindMethod(const MethodCall& method_call) {
     bound_engines[method_call.subchannel] = engine_id;
     switch (engine_id) {
     case EngineID::FERMI_TWOD_A:
-        dma_pusher.BindSubchannel(channel_state.fermi_2d.get(), method_call.subchannel,
-                                  EngineTypes::Fermi2D);
+        dma_pusher.BindSubchannel(&*channel_state.fermi_2d, method_call.subchannel, EngineTypes::Fermi2D);
         break;
     case EngineID::MAXWELL_B:
-        dma_pusher.BindSubchannel(channel_state.maxwell_3d.get(), method_call.subchannel,
-                                  EngineTypes::Maxwell3D);
+        dma_pusher.BindSubchannel(&*channel_state.maxwell_3d, method_call.subchannel, EngineTypes::Maxwell3D);
         break;
     case EngineID::KEPLER_COMPUTE_B:
-        dma_pusher.BindSubchannel(channel_state.kepler_compute.get(), method_call.subchannel,
-                                  EngineTypes::KeplerCompute);
+        dma_pusher.BindSubchannel(&*channel_state.kepler_compute, method_call.subchannel, EngineTypes::KeplerCompute);
         break;
     case EngineID::MAXWELL_DMA_COPY_A:
-        dma_pusher.BindSubchannel(channel_state.maxwell_dma.get(), method_call.subchannel,
-                                  EngineTypes::MaxwellDMA);
+        dma_pusher.BindSubchannel(&*channel_state.maxwell_dma, method_call.subchannel, EngineTypes::MaxwellDMA);
         break;
     case EngineID::KEPLER_INLINE_TO_MEMORY_B:
-        dma_pusher.BindSubchannel(channel_state.kepler_memory.get(), method_call.subchannel,
-                                  EngineTypes::KeplerMemory);
+        dma_pusher.BindSubchannel(&*channel_state.kepler_memory, method_call.subchannel, EngineTypes::KeplerMemory);
+        break;
+    case EngineID::NV01_TIMER:
+        dma_pusher.BindSubchannel(&*channel_state.nv01_timer, method_call.subchannel, EngineTypes::Nv01Timer);
         break;
     default:
         UNIMPLEMENTED_MSG("Unimplemented engine {:04X}", engine_id);
@@ -209,24 +210,22 @@ void Puller::CallEngineMethod(const MethodCall& method_call) {
 
     switch (engine) {
     case EngineID::FERMI_TWOD_A:
-        channel_state.fermi_2d->CallMethod(method_call.method, method_call.argument,
-                                           method_call.IsLastCall());
+        channel_state.fermi_2d->CallMethod(method_call.method, method_call.argument, method_call.IsLastCall());
         break;
     case EngineID::MAXWELL_B:
-        channel_state.maxwell_3d->CallMethod(method_call.method, method_call.argument,
-                                             method_call.IsLastCall());
+        channel_state.maxwell_3d->CallMethod(method_call.method, method_call.argument, method_call.IsLastCall());
         break;
     case EngineID::KEPLER_COMPUTE_B:
-        channel_state.kepler_compute->CallMethod(method_call.method, method_call.argument,
-                                                 method_call.IsLastCall());
+        channel_state.kepler_compute->CallMethod(method_call.method, method_call.argument, method_call.IsLastCall());
         break;
     case EngineID::MAXWELL_DMA_COPY_A:
-        channel_state.maxwell_dma->CallMethod(method_call.method, method_call.argument,
-                                              method_call.IsLastCall());
+        channel_state.maxwell_dma->CallMethod(method_call.method, method_call.argument, method_call.IsLastCall());
         break;
     case EngineID::KEPLER_INLINE_TO_MEMORY_B:
-        channel_state.kepler_memory->CallMethod(method_call.method, method_call.argument,
-                                                method_call.IsLastCall());
+        channel_state.kepler_memory->CallMethod(method_call.method, method_call.argument, method_call.IsLastCall());
+        break;
+    case EngineID::NV01_TIMER:
+        channel_state.nv01_timer->CallMethod(method_call.method, method_call.argument, method_call.IsLastCall());
         break;
     default:
         UNIMPLEMENTED_MSG("Unimplemented engine");
@@ -255,6 +254,9 @@ void Puller::CallEngineMultiMethod(u32 method, u32 subchannel, const u32* base_s
     case EngineID::KEPLER_INLINE_TO_MEMORY_B:
         channel_state.kepler_memory->CallMultiMethod(method, base_start, amount, methods_pending);
         break;
+    case EngineID::NV01_TIMER:
+        channel_state.nv01_timer->CallMultiMethod(method, base_start, amount, methods_pending);
+        break;
     default:
         UNIMPLEMENTED_MSG("Unimplemented engine");
         break;

src/video_core/engines/puller.h

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: 2022 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-3.0-or-later
 
@@ -20,6 +23,7 @@ class MemoryManager;
 class DmaPusher;
 
 enum class EngineID {
+    NV01_TIMER = 0x0004,
     FERMI_TWOD_A = 0x902D, // 2D Engine
     MAXWELL_B = 0xB197,    // 3D Engine
     KEPLER_COMPUTE_B = 0xB1C0,

src/video_core/macro.cpp

@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
 // SPDX-License-Identifier: GPL-3.0-or-later
 
 // SPDX-FileCopyrightText: Copyright 2020 yuzu Emulator Project
@@ -10,6 +10,7 @@
 #include <span>
 
 #include <fstream>
+#include <variant>
 #ifdef ARCHITECTURE_x86_64
 // xbyak hates human beings
 #ifdef __GNUC__
@@ -73,26 +74,12 @@ bool IsTopologySafe(Maxwell3D::Regs::PrimitiveTopology topology) {
     }
 }
 
-class HLEMacroImpl : public CachedMacro {
-public:
-    explicit HLEMacroImpl(Maxwell3D& maxwell3d_)
-        : CachedMacro(maxwell3d_)
-    {}
-};
+} // Anonymous namespace
 
-/// @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 {
+void HLE_DrawArraysIndirect::Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method) {
     auto topology = static_cast<Maxwell3D::Regs::PrimitiveTopology>(parameters[0]);
     if (!maxwell3d.AnyParametersDirty() || !IsTopologySafe(topology)) {
-            Fallback(parameters);
+        Fallback(maxwell3d, parameters);
         return;
     }
 
@@ -117,10 +104,8 @@ public:
         maxwell3d.engine_state = Maxwell3D::EngineHint::None;
         maxwell3d.replace_table.clear();
     }
-    }
-
-private:
-    void Fallback(const std::vector<u32>& parameters) {
+}
+void HLE_DrawArraysIndirect::Fallback(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters) {
     SCOPE_EXIT {
         if (extended) {
             maxwell3d.engine_state = Maxwell3D::EngineHint::None;
@@ -129,52 +114,35 @@ private:
     };
     maxwell3d.RefreshParameters();
     const u32 instance_count = (maxwell3d.GetRegisterValue(0xD1B) & parameters[2]);
-
-        auto topology = static_cast<Maxwell3D::Regs::PrimitiveTopology>(parameters[0]);
+    auto topology = Maxwell3D::Regs::PrimitiveTopology(parameters[0]);
     const u32 vertex_first = parameters[3];
     const u32 vertex_count = parameters[1];
-
     if (!IsTopologySafe(topology) && size_t(maxwell3d.GetMaxCurrentVertices()) < size_t(vertex_first) + size_t(vertex_count)) {
         ASSERT(false && "Faulty draw!");
         return;
     }
-
     const u32 base_instance = parameters[4];
     if (extended) {
         maxwell3d.regs.global_base_instance_index = base_instance;
         maxwell3d.engine_state = Maxwell3D::EngineHint::OnHLEMacro;
-            maxwell3d.SetHLEReplacementAttributeType(
-                0, 0x640, Maxwell3D::HLEReplacementAttributeType::BaseInstance);
+        maxwell3d.SetHLEReplacementAttributeType(0, 0x640, Maxwell3D::HLEReplacementAttributeType::BaseInstance);
     }
-
-        maxwell3d.draw_manager->DrawArray(topology, vertex_first, vertex_count, base_instance,
-                                          instance_count);
-
+    maxwell3d.draw_manager->DrawArray(topology, vertex_first, vertex_count, base_instance, instance_count);
     if (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 {
+void HLE_DrawIndexedIndirect::Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method) {
     auto topology = static_cast<Maxwell3D::Regs::PrimitiveTopology>(parameters[0]);
     if (!maxwell3d.AnyParametersDirty() || !IsTopologySafe(topology)) {
-            Fallback(parameters);
+        Fallback(maxwell3d, parameters);
         return;
     }
 
-        const u32 estimate = static_cast<u32>(maxwell3d.EstimateIndexBufferSize());
+    const u32 estimate = u32(maxwell3d.EstimateIndexBufferSize());
     const u32 element_base = parameters[4];
     const u32 base_instance = parameters[5];
     maxwell3d.regs.vertex_id_base = element_base;
@@ -204,10 +172,8 @@ public:
         maxwell3d.engine_state = Maxwell3D::EngineHint::None;
         maxwell3d.replace_table.clear();
     }
-    }
-
-private:
-    void Fallback(const std::vector<u32>& parameters) {
+}
+void HLE_DrawIndexedIndirect::Fallback(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters) {
     maxwell3d.RefreshParameters();
     const u32 instance_count = (maxwell3d.GetRegisterValue(0xD1B) & parameters[2]);
     const u32 element_base = parameters[4];
@@ -221,9 +187,7 @@ private:
         maxwell3d.SetHLEReplacementAttributeType(0, 0x640, Maxwell3D::HLEReplacementAttributeType::BaseVertex);
         maxwell3d.SetHLEReplacementAttributeType(0, 0x644, Maxwell3D::HLEReplacementAttributeType::BaseInstance);
     }
-
     maxwell3d.draw_manager->DrawIndex(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;
@@ -231,14 +195,8 @@ private:
         maxwell3d.engine_state = Maxwell3D::EngineHint::None;
         maxwell3d.replace_table.clear();
     }
-    }
-};
-
-class HLE_MultiLayerClear final : public HLEMacroImpl {
-public:
-    explicit HLE_MultiLayerClear(Maxwell3D& maxwell3d_) : HLEMacroImpl(maxwell3d_) {}
-
-    void Execute(const std::vector<u32>& parameters, [[maybe_unused]] u32 method) override {
+}
+void HLE_MultiLayerClear::Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method) {
     maxwell3d.RefreshParameters();
     ASSERT(parameters.size() == 1);
 
@@ -249,17 +207,11 @@ public:
 
     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 {
+}
+void HLE_MultiDrawIndexedIndirectCount::Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method) {
     const auto topology = Maxwell3D::Regs::PrimitiveTopology(parameters[2]);
     if (!IsTopologySafe(topology)) {
-            Fallback(parameters);
+        Fallback(maxwell3d, parameters);
         return;
     }
 
@@ -289,19 +241,14 @@ public:
     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.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) {
+}
+void HLE_MultiDrawIndexedIndirectCount::Fallback(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters) {
     SCOPE_EXIT {
         // Clean everything.
         maxwell3d.regs.vertex_id_base = 0x0;
@@ -318,41 +265,29 @@ private:
     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.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]);
+        maxwell3d.draw_manager->DrawIndex(topology, parameters[base + 2], parameters[base], base_vertex, base_instance, parameters[base + 1]);
     }
-    }
-};
-
-class HLE_DrawIndirectByteCount final : public HLEMacroImpl {
-public:
-    explicit HLE_DrawIndirectByteCount(Maxwell3D& maxwell3d_) : HLEMacroImpl(maxwell3d_) {}
-
-    void Execute(const std::vector<u32>& parameters, [[maybe_unused]] u32 method) override {
+}
+void HLE_DrawIndirectByteCount::Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method) {
     const bool force = maxwell3d.Rasterizer().HasDrawTransformFeedback();
-
-        auto topology = static_cast<Maxwell3D::Regs::PrimitiveTopology>(parameters[0] & 0xFFFFU);
+    auto topology = Maxwell3D::Regs::PrimitiveTopology(parameters[0] & 0xFFFFU);
     if (!force && (!maxwell3d.AnyParametersDirty() || !IsTopologySafe(topology))) {
-            Fallback(parameters);
+        Fallback(maxwell3d, parameters);
         return;
     }
     auto& params = maxwell3d.draw_manager->GetIndirectParams();
@@ -367,12 +302,9 @@ public:
     maxwell3d.regs.draw.begin = parameters[0];
     maxwell3d.regs.draw_auto_stride = parameters[1];
     maxwell3d.regs.draw_auto_byte_count = parameters[2];
-
     maxwell3d.draw_manager->DrawArrayIndirect(topology);
-    }
-
-private:
-    void Fallback(const std::vector<u32>& parameters) {
+}
+void HLE_DrawIndirectByteCount::Fallback(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters) {
     maxwell3d.RefreshParameters();
 
     maxwell3d.regs.draw.begin = parameters[0];
@@ -382,14 +314,8 @@ private:
     maxwell3d.draw_manager->DrawArray(
         maxwell3d.regs.draw.topology, 0,
         maxwell3d.regs.draw_auto_byte_count / maxwell3d.regs.draw_auto_stride, 0, 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 {
+}
+void HLE_C713C83D8F63CCF3::Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method) {
     maxwell3d.RefreshParameters();
     const u32 offset = (parameters[0] & 0x3FFFFFFF) << 2;
     const u32 address = maxwell3d.regs.shadow_scratch[24];
@@ -398,14 +324,8 @@ public:
     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 {
+}
+void HLE_D7333D26E0A93EDE::Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method) {
     maxwell3d.RefreshParameters();
     const size_t index = parameters[0];
     const u32 address = maxwell3d.regs.shadow_scratch[42 + index];
@@ -414,14 +334,8 @@ public:
     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 {
+}
+void HLE_BindShader::Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method) {
     maxwell3d.RefreshParameters();
     auto& regs = maxwell3d.regs;
     const u32 index = parameters[0];
@@ -444,14 +358,8 @@ public:
     auto& bind_group = regs.bind_groups[bind_group_id];
     bind_group.raw_config = 0x11;
     maxwell3d.ProcessCBBind(bind_group_id);
-    }
-};
-
-class HLE_SetRasterBoundingBox final : public HLEMacroImpl {
-public:
-    explicit HLE_SetRasterBoundingBox(Maxwell3D& maxwell3d_) : HLEMacroImpl(maxwell3d_) {}
-
-    void Execute(const std::vector<u32>& parameters, [[maybe_unused]] u32 method) override {
+}
+void HLE_SetRasterBoundingBox::Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method) {
     maxwell3d.RefreshParameters();
     const u32 raster_mode = parameters[0];
     auto& regs = maxwell3d.regs;
@@ -459,33 +367,19 @@ public:
     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 {
+}
+void HLE_ClearConstBuffer::Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method) {
+    static constexpr std::array<u32, 0x7000> zeroes{}; //must be bigger than either 7000 or 5F00
     maxwell3d.RefreshParameters();
-        static constexpr std::array<u32, base_size> zeroes{};
     auto& regs = maxwell3d.regs;
     regs.const_buffer.size = 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 {
+}
+void HLE_ClearMemory::Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method) {
     maxwell3d.RefreshParameters();
-
     const u32 needed_memory = parameters[2] / sizeof(u32);
     if (needed_memory > zero_memory.size()) {
         zero_memory.resize(needed_memory, 0);
@@ -497,177 +391,94 @@ public:
     regs.upload.dest.address_low = parameters[1];
     maxwell3d.CallMethod(size_t(MAXWELL3D_REG_INDEX(launch_dma)), 0x1011, true);
     maxwell3d.CallMultiMethod(size_t(MAXWELL3D_REG_INDEX(inline_data)), zero_memory.data(), needed_memory, needed_memory);
-    }
-
-private:
-    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 {
+}
+void HLE_TransformFeedbackSetup::Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method) {
     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(size_t(MAXWELL3D_REG_INDEX(launch_dma)), 0x1011, true);
     maxwell3d.CallMethod(size_t(MAXWELL3D_REG_INDEX(inline_data)), regs.transform_feedback.controls[0].stride, true);
-
     maxwell3d.Rasterizer().RegisterTransformFeedback(regs.upload.dest.Address());
-    }
-};
+}
 
-} // Anonymous namespace
+#define HLE_MACRO_LIST \
+    HLE_MACRO_ELEM(0x0D61FC9FAAC9FCADULL, HLE_DrawArraysIndirect, (false)) \
+    HLE_MACRO_ELEM(0x8A4D173EB99A8603ULL, HLE_DrawArraysIndirect, (true)) \
+    HLE_MACRO_ELEM(0x771BB18C62444DA0ULL, HLE_DrawIndexedIndirect, (false)) \
+    HLE_MACRO_ELEM(0x0217920100488FF7ULL, HLE_DrawIndexedIndirect, (true)) \
+    HLE_MACRO_ELEM(0x3F5E74B9C9A50164ULL, HLE_MultiDrawIndexedIndirectCount, ()) \
+    HLE_MACRO_ELEM(0xEAD26C3E2109B06BULL, HLE_MultiLayerClear, ()) \
+    HLE_MACRO_ELEM(0xC713C83D8F63CCF3ULL, HLE_C713C83D8F63CCF3, ()) \
+    HLE_MACRO_ELEM(0xD7333D26E0A93EDEULL, HLE_D7333D26E0A93EDE, ()) \
+    HLE_MACRO_ELEM(0xEB29B2A09AA06D38ULL, HLE_BindShader, ()) \
+    HLE_MACRO_ELEM(0xDB1341DBEB4C8AF7ULL, HLE_SetRasterBoundingBox, ()) \
+    HLE_MACRO_ELEM(0x6C97861D891EDf7EULL, HLE_ClearConstBuffer, (0x5F00)) \
+    HLE_MACRO_ELEM(0xD246FDDF3A6173D7ULL, HLE_ClearConstBuffer, (0x7000)) \
+    HLE_MACRO_ELEM(0xEE4D0004BEC8ECF4ULL, HLE_ClearMemory, ()) \
+    HLE_MACRO_ELEM(0xFC0CF27F5FFAA661ULL, HLE_TransformFeedbackSetup, ()) \
+    HLE_MACRO_ELEM(0xB5F74EDB717278ECULL, HLE_DrawIndirectByteCount, ()) \
 
-HLEMacro::HLEMacro(Maxwell3D& maxwell3d_) : maxwell3d{maxwell3d_} {}
-
-HLEMacro::~HLEMacro() = default;
-
-std::unique_ptr<CachedMacro> HLEMacro::GetHLEProgram(u64 hash) const {
+// Allocates and returns a cached macro if the hash matches a known function.
+[[nodiscard]] inline AnyCachedMacro GetHLEProgram(u64 hash) noexcept {
     // Compiler will make you a GREAT job at making an ad-hoc hash table :)
     switch (hash) {
-    case 0x0D61FC9FAAC9FCADULL: return std::make_unique<HLE_DrawArraysIndirect<false>>(maxwell3d);
-    case 0x8A4D173EB99A8603ULL: return std::make_unique<HLE_DrawArraysIndirect<true>>(maxwell3d);
-    case 0x771BB18C62444DA0ULL: return std::make_unique<HLE_DrawIndexedIndirect<false>>(maxwell3d);
-    case 0x0217920100488FF7ULL: return std::make_unique<HLE_DrawIndexedIndirect<true>>(maxwell3d);
-    case 0x3F5E74B9C9A50164ULL: return std::make_unique<HLE_MultiDrawIndexedIndirectCount>(maxwell3d);
-    case 0xEAD26C3E2109B06BULL: return std::make_unique<HLE_MultiLayerClear>(maxwell3d);
-    case 0xC713C83D8F63CCF3ULL: return std::make_unique<HLE_C713C83D8F63CCF3>(maxwell3d);
-    case 0xD7333D26E0A93EDEULL: return std::make_unique<HLE_D7333D26E0A93EDE>(maxwell3d);
-    case 0xEB29B2A09AA06D38ULL: return std::make_unique<HLE_BindShader>(maxwell3d);
-    case 0xDB1341DBEB4C8AF7ULL: return std::make_unique<HLE_SetRasterBoundingBox>(maxwell3d);
-    case 0x6C97861D891EDf7EULL: return std::make_unique<HLE_ClearConstBuffer<0x5F00>>(maxwell3d);
-    case 0xD246FDDF3A6173D7ULL: return std::make_unique<HLE_ClearConstBuffer<0x7000>>(maxwell3d);
-    case 0xEE4D0004BEC8ECF4ULL: return std::make_unique<HLE_ClearMemory>(maxwell3d);
-    case 0xFC0CF27F5FFAA661ULL: return std::make_unique<HLE_TransformFeedbackSetup>(maxwell3d);
-    case 0xB5F74EDB717278ECULL: return std::make_unique<HLE_DrawIndirectByteCount>(maxwell3d);
-    default:
-        return nullptr;
+#define HLE_MACRO_ELEM(HASH, TY, VAL) case HASH: return TY VAL;
+    HLE_MACRO_LIST
+#undef HLE_MACRO_ELEM
+    default: return std::monostate{};
+    }
+}
+[[nodiscard]] inline bool CanBeHLEProgram(u64 hash) noexcept {
+    switch (hash) {
+#define HLE_MACRO_ELEM(HASH, TY, VAL) case HASH: return true;
+    HLE_MACRO_LIST
+#undef HLE_MACRO_ELEM
+    default: return false;
     }
 }
 
-namespace {
-class MacroInterpreterImpl final : public CachedMacro {
-public:
-    explicit MacroInterpreterImpl(Engines::Maxwell3D& maxwell3d_, const std::vector<u32>& code_)
-        : CachedMacro(maxwell3d_)
-        , code{code_}
-    {}
-
-    void Execute(const std::vector<u32>& params, u32 method) override;
-
-private:
-    /// Resets the execution engine state, zeroing registers, etc.
-    void Reset();
-
-    /**
-     * Executes a single macro instruction located at the current program counter. Returns whether
-     * the interpreter should keep running.
-     *
-     * @param is_delay_slot Whether the current step is being executed due to a delay slot in a
-     *                      previous instruction.
-     */
-    bool Step(bool is_delay_slot);
-
-    /// Calculates the result of an ALU operation. src_a OP src_b;
-    u32 GetALUResult(Macro::ALUOperation operation, u32 src_a, u32 src_b);
-
-    /// Performs the result operation on the input result and stores it in the specified register
-    /// (if necessary).
-    void ProcessResult(Macro::ResultOperation operation, u32 reg, u32 result);
-
-    /// Evaluates the branch condition and returns whether the branch should be taken or not.
-    bool EvaluateBranchCondition(Macro::BranchCondition cond, u32 value) const;
-
-    /// Reads an opcode at the current program counter location.
-    Macro::Opcode GetOpcode() const;
-
-    /// Returns the specified register's value. Register 0 is hardcoded to always return 0.
-    u32 GetRegister(u32 register_id) const;
-
-    /// Sets the register to the input value.
-    void SetRegister(u32 register_id, u32 value);
-
-    /// Sets the method address to use for the next Send instruction.
-    void SetMethodAddress(u32 address);
-
-    /// Calls a GPU Engine method with the input parameter.
-    void Send(u32 value);
-
-    /// Reads a GPU register located at the method address.
-    u32 Read(u32 method) const;
-
-    /// Returns the next parameter in the parameter queue.
-    u32 FetchParameter();
-
-    /// Current program counter
-    u32 pc{};
-    /// Program counter to execute at after the delay slot is executed.
-    std::optional<u32> delayed_pc;
-
-    /// General purpose macro registers.
-    std::array<u32, Macro::NUM_MACRO_REGISTERS> registers = {};
-
-    /// Method address to use for the next Send instruction.
-    Macro::MethodAddress method_address = {};
-
-    /// Input parameters of the current macro.
-    std::unique_ptr<u32[]> parameters;
-    std::size_t num_parameters = 0;
-    std::size_t parameters_capacity = 0;
-    /// Index of the next parameter that will be fetched by the 'parm' instruction.
-    u32 next_parameter_index = 0;
-
-    bool carry_flag = false;
-    const std::vector<u32>& code;
-};
-
-void MacroInterpreterImpl::Execute(const std::vector<u32>& params, u32 method) {
+void MacroInterpreterImpl::Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> params, u32 method) {
     Reset();
 
     registers[1] = params[0];
-    num_parameters = params.size();
-
-    if (num_parameters > parameters_capacity) {
-        parameters_capacity = num_parameters;
-        parameters = std::make_unique<u32[]>(num_parameters);
-    }
-    std::memcpy(parameters.get(), params.data(), num_parameters * sizeof(u32));
+    parameters.resize(params.size());
+    std::memcpy(parameters.data(), params.data(), params.size() * sizeof(u32));
 
     // Execute the code until we hit an exit condition.
     bool keep_executing = true;
     while (keep_executing) {
-        keep_executing = Step(false);
+        keep_executing = Step(maxwell3d, false);
     }
 
     // Assert the the macro used all the input parameters
-    ASSERT(next_parameter_index == num_parameters);
+    ASSERT(next_parameter_index == parameters.size());
 }
 
+/// Resets the execution engine state, zeroing registers, etc.
 void MacroInterpreterImpl::Reset() {
     registers = {};
     pc = 0;
     delayed_pc = {};
     method_address.raw = 0;
-    num_parameters = 0;
+    // Vector must hold its last indices otherwise wonky shit will happen
     // The next parameter index starts at 1, because $r1 already has the value of the first
     // parameter.
     next_parameter_index = 1;
     carry_flag = false;
 }
 
-bool MacroInterpreterImpl::Step(bool is_delay_slot) {
+/// @brief Executes a single macro instruction located at the current program counter. Returns whether
+/// the interpreter should keep running.
+/// @param is_delay_slot Whether the current step is being executed due to a delay slot in a previous instruction.
+bool MacroInterpreterImpl::Step(Engines::Maxwell3D& maxwell3d, bool is_delay_slot) {
     u32 base_address = pc;
 
     Macro::Opcode opcode = GetOpcode();
@@ -682,14 +493,12 @@ bool MacroInterpreterImpl::Step(bool is_delay_slot) {
 
     switch (opcode.operation) {
     case Macro::Operation::ALU: {
-        u32 result = GetALUResult(opcode.alu_operation, GetRegister(opcode.src_a),
-                                  GetRegister(opcode.src_b));
-        ProcessResult(opcode.result_operation, opcode.dst, result);
+        u32 result = GetALUResult(opcode.alu_operation, GetRegister(opcode.src_a), GetRegister(opcode.src_b));
+        ProcessResult(maxwell3d, opcode.result_operation, opcode.dst, result);
         break;
     }
     case Macro::Operation::AddImmediate: {
-        ProcessResult(opcode.result_operation, opcode.dst,
-                      GetRegister(opcode.src_a) + opcode.immediate);
+        ProcessResult(maxwell3d, opcode.result_operation, opcode.dst, GetRegister(opcode.src_a) + opcode.immediate);
         break;
     }
     case Macro::Operation::ExtractInsert: {
@@ -699,7 +508,7 @@ bool MacroInterpreterImpl::Step(bool is_delay_slot) {
         src = (src >> opcode.bf_src_bit) & opcode.GetBitfieldMask();
         dst &= ~(opcode.GetBitfieldMask() << opcode.bf_dst_bit);
         dst |= src << opcode.bf_dst_bit;
-        ProcessResult(opcode.result_operation, opcode.dst, dst);
+        ProcessResult(maxwell3d, opcode.result_operation, opcode.dst, dst);
         break;
     }
     case Macro::Operation::ExtractShiftLeftImmediate: {
@@ -708,7 +517,7 @@ bool MacroInterpreterImpl::Step(bool is_delay_slot) {
 
         u32 result = ((src >> dst) & opcode.GetBitfieldMask()) << opcode.bf_dst_bit;
 
-        ProcessResult(opcode.result_operation, opcode.dst, result);
+        ProcessResult(maxwell3d, opcode.result_operation, opcode.dst, result);
         break;
     }
     case Macro::Operation::ExtractShiftLeftRegister: {
@@ -717,12 +526,12 @@ bool MacroInterpreterImpl::Step(bool is_delay_slot) {
 
         u32 result = ((src >> opcode.bf_src_bit) & opcode.GetBitfieldMask()) << dst;
 
-        ProcessResult(opcode.result_operation, opcode.dst, result);
+        ProcessResult(maxwell3d, opcode.result_operation, opcode.dst, result);
         break;
     }
     case Macro::Operation::Read: {
-        u32 result = Read(GetRegister(opcode.src_a) + opcode.immediate);
-        ProcessResult(opcode.result_operation, opcode.dst, result);
+        u32 result = Read(maxwell3d, GetRegister(opcode.src_a) + opcode.immediate);
+        ProcessResult(maxwell3d, opcode.result_operation, opcode.dst, result);
         break;
     }
     case Macro::Operation::Branch: {
@@ -738,7 +547,7 @@ bool MacroInterpreterImpl::Step(bool is_delay_slot) {
 
             delayed_pc = base_address + opcode.GetBranchTarget();
             // Execute one more instruction due to the delay slot.
-            return Step(true);
+            return Step(maxwell3d, true);
         }
         break;
     }
@@ -751,13 +560,13 @@ bool MacroInterpreterImpl::Step(bool is_delay_slot) {
     // cause an exit if it's executed inside a delay slot.
     if (opcode.is_exit && !is_delay_slot) {
         // Exit has a delay slot, execute the next instruction
-        Step(true);
+        Step(maxwell3d, true);
         return false;
     }
-
     return true;
 }
 
+/// Calculates the result of an ALU operation. src_a OP src_b;
 u32 MacroInterpreterImpl::GetALUResult(Macro::ALUOperation operation, u32 src_a, u32 src_b) {
     switch (operation) {
     case Macro::ALUOperation::Add: {
@@ -797,7 +606,8 @@ u32 MacroInterpreterImpl::GetALUResult(Macro::ALUOperation operation, u32 src_a,
     }
 }
 
-void MacroInterpreterImpl::ProcessResult(Macro::ResultOperation operation, u32 reg, u32 result) {
+/// Performs the result operation on the input result and stores it in the specified register (if necessary).
+void MacroInterpreterImpl::ProcessResult(Engines::Maxwell3D& maxwell3d, Macro::ResultOperation operation, u32 reg, u32 result) {
     switch (operation) {
     case Macro::ResultOperation::IgnoreAndFetch:
         // Fetch parameter and ignore result.
@@ -815,12 +625,12 @@ void MacroInterpreterImpl::ProcessResult(Macro::ResultOperation operation, u32 r
     case Macro::ResultOperation::FetchAndSend:
         // Fetch parameter and send result.
         SetRegister(reg, FetchParameter());
-        Send(result);
+        Send(maxwell3d, result);
         break;
     case Macro::ResultOperation::MoveAndSend:
         // Move and send result.
         SetRegister(reg, result);
-        Send(result);
+        Send(maxwell3d, result);
         break;
     case Macro::ResultOperation::FetchAndSetMethod:
         // Fetch parameter and use result as Method Address.
@@ -831,13 +641,13 @@ void MacroInterpreterImpl::ProcessResult(Macro::ResultOperation operation, u32 r
         // Move result and use as Method Address, then fetch and send parameter.
         SetRegister(reg, result);
         SetMethodAddress(result);
-        Send(FetchParameter());
+        Send(maxwell3d, FetchParameter());
         break;
     case Macro::ResultOperation::MoveAndSetMethodSend:
         // Move result and use as Method Address, then send bits 12:17 of result.
         SetRegister(reg, result);
         SetMethodAddress(result);
-        Send((result >> 12) & 0b111111);
+        Send(maxwell3d, (result >> 12) & 0b111111);
         break;
     default:
         UNIMPLEMENTED_MSG("Unimplemented result operation {}", operation);
@@ -845,6 +655,7 @@ void MacroInterpreterImpl::ProcessResult(Macro::ResultOperation operation, u32 r
     }
 }
 
+/// Evaluates the branch condition and returns whether the branch should be taken or not.
 bool MacroInterpreterImpl::EvaluateBranchCondition(Macro::BranchCondition cond, u32 value) const {
     switch (cond) {
     case Macro::BranchCondition::Zero:
@@ -855,46 +666,44 @@ bool MacroInterpreterImpl::EvaluateBranchCondition(Macro::BranchCondition cond,
     UNREACHABLE();
 }
 
+/// Reads an opcode at the current program counter location.
 Macro::Opcode MacroInterpreterImpl::GetOpcode() const {
     ASSERT((pc % sizeof(u32)) == 0);
     ASSERT(pc < code.size() * sizeof(u32));
     return {code[pc / sizeof(u32)]};
 }
 
+/// Returns the specified register's value. Register 0 is hardcoded to always return 0.
 u32 MacroInterpreterImpl::GetRegister(u32 register_id) const {
-    return registers.at(register_id);
+    return registers[register_id];
 }
 
+/// Sets the register to the input value.
 void MacroInterpreterImpl::SetRegister(u32 register_id, u32 value) {
     // Register 0 is hardwired as the zero register.
     // Ensure no writes to it actually occur.
-    if (register_id == 0) {
+    if (register_id == 0)
         return;
-    }
-
-    registers.at(register_id) = value;
+    registers[register_id] = value;
 }
 
-void MacroInterpreterImpl::SetMethodAddress(u32 address) {
-    method_address.raw = address;
-}
-
-void MacroInterpreterImpl::Send(u32 value) {
+/// Calls a GPU Engine method with the input parameter.
+void MacroInterpreterImpl::Send(Engines::Maxwell3D& maxwell3d, u32 value) {
     maxwell3d.CallMethod(method_address.address, value, true);
     // Increment the method address by the method increment.
-    method_address.address.Assign(method_address.address.Value() +
-                                  method_address.increment.Value());
+    method_address.address.Assign(method_address.address.Value() + method_address.increment.Value());
 }
 
-u32 MacroInterpreterImpl::Read(u32 method) const {
+/// Reads a GPU register located at the method address.
+u32 MacroInterpreterImpl::Read(Engines::Maxwell3D& maxwell3d, u32 method) const {
     return maxwell3d.GetRegisterValue(method);
 }
 
+/// Returns the next parameter in the parameter queue.
 u32 MacroInterpreterImpl::FetchParameter() {
-    ASSERT(next_parameter_index < num_parameters);
+    ASSERT(next_parameter_index < parameters.size());
     return parameters[next_parameter_index++];
 }
-} // Anonymous namespace
 
 #ifdef ARCHITECTURE_x86_64
 namespace {
@@ -930,17 +739,15 @@ static const auto default_cg_mode = Xbyak::DontSetProtectRWE;
 static const auto default_cg_mode = nullptr; //Allow RWE
 #endif
 
-class MacroJITx64Impl final : public Xbyak::CodeGenerator, public CachedMacro {
-public:
-    explicit MacroJITx64Impl(Engines::Maxwell3D& maxwell3d_, const std::vector<u32>& code_)
+struct MacroJITx64Impl final : public Xbyak::CodeGenerator, public DynamicCachedMacro {
+    explicit MacroJITx64Impl(std::span<const u32> code_)
         : Xbyak::CodeGenerator(MAX_CODE_SIZE, default_cg_mode)
-        , CachedMacro(maxwell3d_)
         , code{code_}
     {
         Compile();
     }
 
-    void Execute(const std::vector<u32>& parameters, u32 method) override;
+    void Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, u32 method) override;
 
     void Compile_ALU(Macro::Opcode opcode);
     void Compile_AddImmediate(Macro::Opcode opcode);
@@ -950,18 +757,13 @@ public:
     void Compile_Read(Macro::Opcode opcode);
     void Compile_Branch(Macro::Opcode opcode);
 
-private:
     void Optimizer_ScanFlags();
-
     void Compile();
     bool Compile_NextInstruction();
-
     Xbyak::Reg32 Compile_FetchParameter();
     Xbyak::Reg32 Compile_GetRegister(u32 index, Xbyak::Reg32 dst);
-
     void Compile_ProcessResult(Macro::ResultOperation operation, u32 reg);
     void Compile_Send(Xbyak::Reg32 value);
-
     Macro::Opcode GetOpCode() const;
 
     struct JITState {
@@ -981,21 +783,17 @@ private:
         bool enable_asserts{};
     };
     OptimizerState optimizer{};
-
     std::optional<Macro::Opcode> next_opcode{};
     ProgramType program{nullptr};
-
     std::array<Xbyak::Label, MAX_CODE_SIZE> labels;
     std::array<Xbyak::Label, MAX_CODE_SIZE> delay_skip;
     Xbyak::Label end_of_code{};
-
     bool is_delay_slot{};
     u32 pc{};
-
-    const std::vector<u32>& code;
+    std::span<const u32> code;
 };
 
-void MacroJITx64Impl::Execute(const std::vector<u32>& parameters, u32 method) {
+void MacroJITx64Impl::Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, u32 method) {
     ASSERT_OR_EXECUTE(program != nullptr, { return; });
     JITState state{};
     state.maxwell3d = &maxwell3d;
@@ -1231,7 +1029,7 @@ void MacroJITx64Impl::Compile_Read(Macro::Opcode opcode) {
     Compile_ProcessResult(opcode.result_operation, opcode.dst);
 }
 
-void Send(Engines::Maxwell3D* maxwell3d, Macro::MethodAddress method_address, u32 value) {
+static void MacroJIT_SendThunk(Engines::Maxwell3D* maxwell3d, Macro::MethodAddress method_address, u32 value) {
     maxwell3d->CallMethod(method_address.address, value, true);
 }
 
@@ -1240,7 +1038,7 @@ void MacroJITx64Impl::Compile_Send(Xbyak::Reg32 value) {
     mov(Common::X64::ABI_PARAM1, qword[STATE]);
     mov(Common::X64::ABI_PARAM2.cvt32(), METHOD_ADDRESS);
     mov(Common::X64::ABI_PARAM3.cvt32(), value);
-    Common::X64::CallFarFunction(*this, &Send);
+    Common::X64::CallFarFunction(*this, &MacroJIT_SendThunk);
     Common::X64::ABI_PopRegistersAndAdjustStack(*this, PersistentCallerSavedRegs(), 0);
 
     Xbyak::Label dont_process{};
@@ -1452,10 +1250,8 @@ bool MacroJITx64Impl::Compile_NextInstruction() {
     return true;
 }
 
-static void WarnInvalidParameter(uintptr_t parameter, uintptr_t max_parameter) {
-    LOG_CRITICAL(HW_GPU,
-                 "Macro JIT: invalid parameter access 0x{:x} (0x{:x} is the last parameter)",
-                 parameter, max_parameter - sizeof(u32));
+static void MacroJIT_ErrorThunk(uintptr_t parameter, uintptr_t max_parameter) {
+    LOG_CRITICAL(HW_GPU, "Macro JIT: invalid parameter access 0x{:x} (0x{:x} is the last parameter)", parameter, max_parameter - sizeof(u32));
 }
 
 Xbyak::Reg32 MacroJITx64Impl::Compile_FetchParameter() {
@@ -1465,7 +1261,7 @@ Xbyak::Reg32 MacroJITx64Impl::Compile_FetchParameter() {
     Common::X64::ABI_PushRegistersAndAdjustStack(*this, PersistentCallerSavedRegs(), 0);
     mov(Common::X64::ABI_PARAM1, PARAMETERS);
     mov(Common::X64::ABI_PARAM2, MAX_PARAMETER);
-    Common::X64::CallFarFunction(*this, &WarnInvalidParameter);
+    Common::X64::CallFarFunction(*this, &MacroJIT_ErrorThunk);
     Common::X64::ABI_PopRegistersAndAdjustStack(*this, PersistentCallerSavedRegs(), 0);
     L(parameter_ok);
     mov(eax, dword[PARAMETERS]);
@@ -1574,33 +1370,42 @@ static void Dump(u64 hash, std::span<const u32> code, bool decompiled = false) {
     macro_file.write(reinterpret_cast<const char*>(code.data()), code.size_bytes());
 }
 
-MacroEngine::MacroEngine(Engines::Maxwell3D& maxwell3d_, bool is_interpreted_)
-    : hle_macros{std::make_optional<Tegra::HLEMacro>(maxwell3d_)}
-    , maxwell3d{maxwell3d_}
-    , is_interpreted{is_interpreted_}
-{}
-
-MacroEngine::~MacroEngine() = default;
-
-void MacroEngine::AddCode(u32 method, u32 data) {
-    uploaded_macro_code[method].push_back(data);
-}
-
-void MacroEngine::ClearCode(u32 method) {
-    macro_cache.erase(method);
-    uploaded_macro_code.erase(method);
-}
-
-void MacroEngine::Execute(u32 method, const std::vector<u32>& parameters) {
-    auto compiled_macro = macro_cache.find(method);
-    if (compiled_macro != macro_cache.end()) {
-        const auto& cache_info = compiled_macro->second;
-        if (cache_info.has_hle_program) {
-            cache_info.hle_program->Execute(parameters, method);
-        } else {
-            maxwell3d.RefreshParameters();
-            cache_info.lle_program->Execute(parameters, method);
-        }
+void MacroEngine::Execute(Engines::Maxwell3D& maxwell3d, u32 method, std::span<const u32> parameters) {
+    auto const execute_variant = [&maxwell3d, &parameters, method](AnyCachedMacro& acm) {
+        if (auto a = std::get_if<HLE_DrawArraysIndirect>(&acm))
+            a->Execute(maxwell3d, parameters, method);
+        if (auto a = std::get_if<HLE_DrawIndexedIndirect>(&acm))
+            a->Execute(maxwell3d, parameters, method);
+        if (auto a = std::get_if<HLE_MultiDrawIndexedIndirectCount>(&acm))
+            a->Execute(maxwell3d, parameters, method);
+        if (auto a = std::get_if<HLE_MultiLayerClear>(&acm))
+            a->Execute(maxwell3d, parameters, method);
+        if (auto a = std::get_if<HLE_C713C83D8F63CCF3>(&acm))
+            a->Execute(maxwell3d, parameters, method);
+        if (auto a = std::get_if<HLE_D7333D26E0A93EDE>(&acm))
+            a->Execute(maxwell3d, parameters, method);
+        if (auto a = std::get_if<HLE_BindShader>(&acm))
+            a->Execute(maxwell3d, parameters, method);
+        if (auto a = std::get_if<HLE_SetRasterBoundingBox>(&acm))
+            a->Execute(maxwell3d, parameters, method);
+        if (auto a = std::get_if<HLE_ClearConstBuffer>(&acm))
+            a->Execute(maxwell3d, parameters, method);
+        if (auto a = std::get_if<HLE_ClearMemory>(&acm))
+            a->Execute(maxwell3d, parameters, method);
+        if (auto a = std::get_if<HLE_TransformFeedbackSetup>(&acm))
+            a->Execute(maxwell3d, parameters, method);
+        if (auto a = std::get_if<HLE_DrawIndirectByteCount>(&acm))
+            a->Execute(maxwell3d, parameters, method);
+        if (auto a = std::get_if<MacroInterpreterImpl>(&acm))
+            a->Execute(maxwell3d, parameters, method);
+        if (auto a = std::get_if<std::unique_ptr<DynamicCachedMacro>>(&acm))
+            a->get()->Execute(maxwell3d, parameters, method);
+    };
+    if (auto const it = macro_cache.find(method); it != macro_cache.end()) {
+        auto& ci = it->second;
+        if (!CanBeHLEProgram(ci.hash) || Settings::values.disable_macro_hle)
+            maxwell3d.RefreshParameters(); //LLE must reload parameters
+        execute_variant(ci.program);
     } else {
         // Macro not compiled, check if it's uploaded and if so, compile it
         std::optional<u32> mid_method;
@@ -1617,51 +1422,37 @@ void MacroEngine::Execute(u32 method, const std::vector<u32>& parameters) {
                 return;
             }
         }
-        auto& cache_info = macro_cache[method];
-
-        if (!mid_method.has_value()) {
-            cache_info.lle_program = Compile(macro_code->second);
-            cache_info.hash = Common::HashValue(macro_code->second);
-        } else {
+        auto& ci = macro_cache[method];
+        if (mid_method) {
             const auto& macro_cached = uploaded_macro_code[mid_method.value()];
             const auto rebased_method = method - mid_method.value();
             auto& code = uploaded_macro_code[method];
             code.resize(macro_cached.size() - rebased_method);
             std::memcpy(code.data(), macro_cached.data() + rebased_method, code.size() * sizeof(u32));
-            cache_info.hash = Common::HashValue(code);
-            cache_info.lle_program = Compile(code);
-        }
-
-        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);
+            ci.hash = Common::HashValue(code);
+            ci.program = Compile(maxwell3d, code);
         } else {
-            cache_info.has_hle_program = true;
-            cache_info.hle_program = std::move(hle_program);
-            cache_info.hle_program->Execute(parameters, method);
+            ci.program = Compile(maxwell3d, macro_code->second);
+            ci.hash = Common::HashValue(macro_code->second);
         }
-
+        if (CanBeHLEProgram(ci.hash) && !Settings::values.disable_macro_hle) {
+            ci.program = GetHLEProgram(ci.hash);
+        } else {
+            maxwell3d.RefreshParameters();
+        }
+        execute_variant(ci.program);
         if (Settings::values.dump_macros) {
-            Dump(cache_info.hash, macro_code->second, cache_info.has_hle_program);
+            Dump(ci.hash, macro_code->second, !std::holds_alternative<std::monostate>(ci.program));
         }
     }
 }
 
-std::unique_ptr<CachedMacro> MacroEngine::Compile(const std::vector<u32>& code) {
+AnyCachedMacro MacroEngine::Compile(Engines::Maxwell3D& maxwell3d, std::span<const u32> code) {
 #ifdef ARCHITECTURE_x86_64
     if (!is_interpreted)
-        return std::make_unique<MacroJITx64Impl>(maxwell3d, code);
-#endif
-    return std::make_unique<MacroInterpreterImpl>(maxwell3d, code);
-}
-
-std::optional<MacroEngine> GetMacroEngine(Engines::Maxwell3D& maxwell3d) {
-#ifdef ARCHITECTURE_x86_64
-    return std::make_optional<MacroEngine>(maxwell3d, bool(Settings::values.disable_macro_jit));
-#else
-    return std::make_optional<MacroEngine>(maxwell3d, true);
+        return std::make_unique<MacroJITx64Impl>(code);
 #endif
+    return MacroInterpreterImpl(code);
 }
 
 } // namespace Tegra

src/video_core/macro.h

@@ -7,8 +7,10 @@
 #pragma once
 
 #include <memory>
-#include <ankerl/unordered_dense.h>
+#include <span>
+#include <variant>
 #include <vector>
+#include <ankerl/unordered_dense.h>
 #include "common/bit_field.h"
 #include "common/common_types.h"
 
@@ -98,62 +100,142 @@ union MethodAddress {
 
 } // namespace Macro
 
-class CachedMacro {
-public:
-    CachedMacro(Engines::Maxwell3D& maxwell3d_)
-        : maxwell3d{maxwell3d_}
-    {}
-    virtual ~CachedMacro() = default;
+struct HLEMacro {
+};
+/// @note: these macros have two versions, a normal and extended version, with the extended version
+/// also assigning the base vertex/instance.
+struct HLE_DrawArraysIndirect final {
+    HLE_DrawArraysIndirect(bool extended_) noexcept : extended{extended_} {}
+    void Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method);
+    void Fallback(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters);
+    bool extended;
+};
+/// @note: these macros have two versions, a normal and extended version, with the extended version
+/// also assigning the base vertex/instance.
+struct HLE_DrawIndexedIndirect final {
+    explicit HLE_DrawIndexedIndirect(bool extended_) noexcept : extended{extended_} {}
+    void Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method);
+    void Fallback(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters);
+    bool extended;
+};
+struct HLE_MultiLayerClear final {
+    void Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method);
+};
+struct HLE_MultiDrawIndexedIndirectCount final {
+    void Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method);
+    void Fallback(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters);
+};
+struct HLE_DrawIndirectByteCount final {
+    void Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method);
+    void Fallback(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters);
+};
+struct HLE_C713C83D8F63CCF3 final {
+    void Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method);
+};
+struct HLE_D7333D26E0A93EDE final {
+    void Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method);
+};
+struct HLE_BindShader final {
+    void Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method);
+};
+struct HLE_SetRasterBoundingBox final {
+    void Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method);
+};
+struct HLE_ClearConstBuffer final {
+    HLE_ClearConstBuffer(size_t base_size_) noexcept : base_size{base_size_} {}
+    void Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method);
+    size_t base_size;
+};
+struct HLE_ClearMemory final {
+    void Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method);
+    std::vector<u32> zero_memory;
+};
+struct HLE_TransformFeedbackSetup final {
+    void Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, [[maybe_unused]] u32 method);
+};
+struct MacroInterpreterImpl final {
+    MacroInterpreterImpl() {}
+    MacroInterpreterImpl(std::span<const u32> code_) : code{code_} {}
+    void Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> params, u32 method);
+    void Reset();
+    bool Step(Engines::Maxwell3D& maxwell3d, bool is_delay_slot);
+    u32 GetALUResult(Macro::ALUOperation operation, u32 src_a, u32 src_b);
+    void ProcessResult(Engines::Maxwell3D& maxwell3d, Macro::ResultOperation operation, u32 reg, u32 result);
+    bool EvaluateBranchCondition(Macro::BranchCondition cond, u32 value) const;
+    Macro::Opcode GetOpcode() const;
+    u32 GetRegister(u32 register_id) const;
+    void SetRegister(u32 register_id, u32 value);
+    /// Sets the method address to use for the next Send instruction.
+    [[nodiscard]] inline void SetMethodAddress(u32 address) noexcept {
+        method_address.raw = address;
+    }
+    void Send(Engines::Maxwell3D& maxwell3d, u32 value);
+    u32 Read(Engines::Maxwell3D& maxwell3d, u32 method) const;
+    u32 FetchParameter();
+    /// General purpose macro registers.
+    std::array<u32, Macro::NUM_MACRO_REGISTERS> registers = {};
+    /// Input parameters of the current macro.
+    std::vector<u32> parameters;
+    std::span<const u32> code;
+    /// Program counter to execute at after the delay slot is executed.
+    std::optional<u32> delayed_pc;
+    /// Method address to use for the next Send instruction.
+    Macro::MethodAddress method_address = {};
+    /// Current program counter
+    u32 pc{};
+    /// Index of the next parameter that will be fetched by the 'parm' instruction.
+    u32 next_parameter_index = 0;
+    bool carry_flag = false;
+};
+struct DynamicCachedMacro {
+    virtual ~DynamicCachedMacro() = default;
     /// Executes the macro code with the specified input parameters.
     /// @param parameters The parameters of the macro
     /// @param method     The method to execute
-    virtual void Execute(const std::vector<u32>& parameters, u32 method) = 0;
-    Engines::Maxwell3D& maxwell3d;
+    virtual void Execute(Engines::Maxwell3D& maxwell3d, std::span<const u32> parameters, u32 method) = 0;
 };
 
-class HLEMacro {
-public:
-    explicit HLEMacro(Engines::Maxwell3D& maxwell3d_);
-    ~HLEMacro();
-    // Allocates and returns a cached macro if the hash matches a known function.
-    // Returns nullptr otherwise.
-    [[nodiscard]] std::unique_ptr<CachedMacro> GetHLEProgram(u64 hash) const;
-private:
-    Engines::Maxwell3D& maxwell3d;
-};
-
-class MacroEngine {
-public:
-    explicit MacroEngine(Engines::Maxwell3D& maxwell3d, bool is_interpreted);
-    ~MacroEngine();
+using AnyCachedMacro = std::variant<
+    std::monostate,
+    HLEMacro,
+    HLE_DrawArraysIndirect,
+    HLE_DrawIndexedIndirect,
+    HLE_MultiDrawIndexedIndirectCount,
+    HLE_MultiLayerClear,
+    HLE_C713C83D8F63CCF3,
+    HLE_D7333D26E0A93EDE,
+    HLE_BindShader,
+    HLE_SetRasterBoundingBox,
+    HLE_ClearConstBuffer,
+    HLE_ClearMemory,
+    HLE_TransformFeedbackSetup,
+    HLE_DrawIndirectByteCount,
+    MacroInterpreterImpl,
+    // Used for JIT x86 macro
+    std::unique_ptr<DynamicCachedMacro>
+>;
 
+struct MacroEngine {
+    MacroEngine(bool is_interpreted_) noexcept : is_interpreted{is_interpreted_} {}
     // Store the uploaded macro code to compile them when they're called.
-    void AddCode(u32 method, u32 data);
-
+    inline void AddCode(u32 method, u32 data) noexcept {
+        uploaded_macro_code[method].push_back(data);
+    }
     // Clear the code associated with a method.
-    void ClearCode(u32 method);
-
+    inline void ClearCode(u32 method) noexcept {
+        macro_cache.erase(method);
+        uploaded_macro_code.erase(method);
+    }
     // Compiles the macro if its not in the cache, and executes the compiled macro
-    void Execute(u32 method, const std::vector<u32>& parameters);
-
-protected:
-    std::unique_ptr<CachedMacro> Compile(const std::vector<u32>& code);
-
-private:
+    void Execute(Engines::Maxwell3D& maxwell3d, u32 method, std::span<const u32> parameters);
+    AnyCachedMacro Compile(Engines::Maxwell3D& maxwell3d, std::span<const u32> code);
     struct CacheInfo {
-        std::unique_ptr<CachedMacro> lle_program{};
-        std::unique_ptr<CachedMacro> hle_program{};
+        AnyCachedMacro program;
         u64 hash{};
-        bool has_hle_program{};
     };
-
     ankerl::unordered_dense::map<u32, CacheInfo> macro_cache;
     ankerl::unordered_dense::map<u32, std::vector<u32>> uploaded_macro_code;
-    std::optional<HLEMacro> hle_macros;
-    Engines::Maxwell3D& maxwell3d;
     bool is_interpreted;
 };
 
-std::optional<MacroEngine> GetMacroEngine(Engines::Maxwell3D& maxwell3d);
-
 } // namespace Tegra

src/video_core/renderer_opengl/gl_shader_cache.cpp

@@ -194,6 +194,7 @@ ShaderCache::ShaderCache(Tegra::MaxwellDeviceMemoryManager& device_memory_,
           .support_float_controls = false,
           .support_separate_denorm_behavior = false,
           .support_separate_rounding_mode = false,
+          .support_fp32_rounding_rtz = false,
           .support_fp16_denorm_preserve = false,
           .support_fp32_denorm_preserve = false,
           .support_fp16_denorm_flush = false,

src/video_core/renderer_vulkan/blit_image.cpp

@@ -1045,7 +1045,7 @@ void BlitImageHelper::ConvertDepthToColorPipeline(vk::Pipeline& pipeline, VkRend
     VkShaderModule frag_shader = *convert_float_to_depth_frag;
     const std::array stages = MakeStages(*full_screen_vert, frag_shader);
     const VkPipelineInputAssemblyStateCreateInfo input_assembly_ci = GetPipelineInputAssemblyStateCreateInfo(device);
-    pipeline = device.GetLogical().CreateGraphicsPipeline({
+    pipeline = device.GetLogical().CreateGraphicsPipeline(VkGraphicsPipelineCreateInfo{
         .sType = VK_STRUCTURE_TYPE_GRAPHICS_PIPELINE_CREATE_INFO,
         .pNext = nullptr,
         .flags = 0,
@@ -1075,7 +1075,7 @@ void BlitImageHelper::ConvertColorToDepthPipeline(vk::Pipeline& pipeline, VkRend
     VkShaderModule frag_shader = *convert_depth_to_float_frag;
     const std::array stages = MakeStages(*full_screen_vert, frag_shader);
     const VkPipelineInputAssemblyStateCreateInfo input_assembly_ci = GetPipelineInputAssemblyStateCreateInfo(device);
-    pipeline = device.GetLogical().CreateGraphicsPipeline({
+    pipeline = device.GetLogical().CreateGraphicsPipeline(VkGraphicsPipelineCreateInfo{
         .sType = VK_STRUCTURE_TYPE_GRAPHICS_PIPELINE_CREATE_INFO,
         .pNext = nullptr,
         .flags = 0,
@@ -1106,7 +1106,7 @@ void BlitImageHelper::ConvertPipelineEx(vk::Pipeline& pipeline, VkRenderPass ren
     }
     const std::array stages = MakeStages(*full_screen_vert, *module);
     const VkPipelineInputAssemblyStateCreateInfo input_assembly_ci = GetPipelineInputAssemblyStateCreateInfo(device);
-    pipeline = device.GetLogical().CreateGraphicsPipeline({
+    pipeline = device.GetLogical().CreateGraphicsPipeline(VkGraphicsPipelineCreateInfo{
         .sType = VK_STRUCTURE_TYPE_GRAPHICS_PIPELINE_CREATE_INFO,
         .pNext = nullptr,
         .flags = 0,
@@ -1148,7 +1148,7 @@ void BlitImageHelper::ConvertPipeline(vk::Pipeline& pipeline, VkRenderPass rende
         is_target_depth ? *convert_float_to_depth_frag : *convert_depth_to_float_frag;
     const std::array stages = MakeStages(*full_screen_vert, frag_shader);
     const VkPipelineInputAssemblyStateCreateInfo input_assembly_ci = GetPipelineInputAssemblyStateCreateInfo(device);
-    pipeline = device.GetLogical().CreateGraphicsPipeline({
+    pipeline = device.GetLogical().CreateGraphicsPipeline(VkGraphicsPipelineCreateInfo{
         .sType = VK_STRUCTURE_TYPE_GRAPHICS_PIPELINE_CREATE_INFO,
         .pNext = nullptr,
         .flags = 0,

src/video_core/renderer_vulkan/pipeline_helper.h

@@ -15,6 +15,7 @@
 #include "shader_recompiler/shader_info.h"
 #include "video_core/renderer_vulkan/vk_texture_cache.h"
 #include "video_core/renderer_vulkan/vk_update_descriptor.h"
+#include "video_core/surface.h"
 #include "video_core/texture_cache/types.h"
 #include "video_core/vulkan_common/vulkan_device.h"
 
@@ -232,10 +233,16 @@ inline void PushImageDescriptors(TextureCache& texture_cache,
             ImageView& image_view{texture_cache.GetImageView(image_view_id)};
             const VkImageView vk_image_view{image_view.Handle(desc.type)};
             const Sampler& sampler{texture_cache.GetSampler(sampler_id)};
+            const auto surface_type{VideoCore::Surface::GetFormatType(image_view.format)};
+            const bool allow_depth_compare =
+                desc.is_depth && (surface_type == VideoCore::Surface::SurfaceType::Depth ||
+                                  surface_type == VideoCore::Surface::SurfaceType::DepthStencil);
             const bool use_fallback_sampler{sampler.HasAddedAnisotropy() &&
                                             !image_view.SupportsAnisotropy()};
-            const VkSampler vk_sampler{use_fallback_sampler ? sampler.HandleWithDefaultAnisotropy()
-                                                            : sampler.Handle()};
+            const VkSampler vk_sampler{use_fallback_sampler
+                                           ? sampler.HandleWithDefaultAnisotropy(
+                                                 allow_depth_compare)
+                                           : sampler.Handle(allow_depth_compare)};
             guest_descriptor_queue.AddSampledImage(vk_image_view, vk_sampler);
             rescaling.PushTexture(texture_cache.IsRescaling(image_view));
         }

src/video_core/renderer_vulkan/renderer_vulkan.cpp

@@ -137,14 +137,8 @@ try
                       memory_allocator,
                       scheduler,
                       swapchain,
-#ifdef ANDROID
                       surface)
-    ,
-#else
-                      *surface)
-    ,
-#endif
-    blit_swapchain(device_memory,
+    , blit_swapchain(device_memory,
                    device,
                    memory_allocator,
                    present_manager,

src/video_core/renderer_vulkan/vk_buffer_cache.cpp

@@ -10,6 +10,7 @@
 #include <span>
 #include <vector>
 
+#include "video_core/buffer_cache/buffer_cache_base.h"
 #include "video_core/renderer_vulkan/vk_buffer_cache.h"
 
 #include "video_core/renderer_vulkan/maxwell_to_vk.h"
@@ -597,18 +598,18 @@ void BufferCacheRuntime::BindVertexBuffer(u32 index, VkBuffer buffer, u32 offset
 }
 
 void BufferCacheRuntime::BindVertexBuffers(VideoCommon::HostBindings<Buffer>& bindings) {
-    boost::container::small_vector<VkBuffer, 32> buffer_handles;
-    for (u32 index = 0; index < bindings.buffers.size(); ++index) {
-        auto handle = bindings.buffers[index]->Handle();
+    boost::container::static_vector<VkBuffer, VideoCommon::NUM_VERTEX_BUFFERS> buffer_handles(bindings.buffers.size());
+    for (u32 i = 0; i < bindings.buffers.size(); ++i) {
+        auto handle = bindings.buffers[i]->Handle();
         if (handle == VK_NULL_HANDLE) {
-            bindings.offsets[index] = 0;
-            bindings.sizes[index] = VK_WHOLE_SIZE;
+            bindings.offsets[i] = 0;
+            bindings.sizes[i] = VK_WHOLE_SIZE;
             if (!device.HasNullDescriptor()) {
                 ReserveNullBuffer();
                 handle = *null_buffer;
             }
         }
-        buffer_handles.push_back(handle);
+        buffer_handles[i] = handle;
     }
     const u32 device_max = device.GetMaxVertexInputBindings();
     const u32 min_binding = (std::min)(bindings.min_index, device_max);
@@ -651,15 +652,21 @@ void BufferCacheRuntime::BindTransformFeedbackBuffers(VideoCommon::HostBindings<
         // Already logged in the rasterizer
         return;
     }
-    boost::container::small_vector<VkBuffer, 4> buffer_handles;
-    for (u32 index = 0; index < bindings.buffers.size(); ++index) {
-        buffer_handles.push_back(bindings.buffers[index]->Handle());
+    boost::container::static_vector<VkBuffer, VideoCommon::NUM_VERTEX_BUFFERS> buffer_handles(bindings.buffers.size());
+    for (u32 i = 0; i < bindings.buffers.size(); ++i) {
+        auto handle = bindings.buffers[i]->Handle();
+        if (handle == VK_NULL_HANDLE) {
+            bindings.offsets[i] = 0;
+            bindings.sizes[i] = VK_WHOLE_SIZE;
+            if (!device.HasNullDescriptor()) {
+                ReserveNullBuffer();
+                handle = *null_buffer;
             }
-    scheduler.Record([bindings_ = std::move(bindings),
-                      buffer_handles_ = std::move(buffer_handles)](vk::CommandBuffer cmdbuf) {
-        cmdbuf.BindTransformFeedbackBuffersEXT(0, static_cast<u32>(buffer_handles_.size()),
-                                               buffer_handles_.data(), bindings_.offsets.data(),
-                                               bindings_.sizes.data());
+        }
+        buffer_handles[i] = handle;
+    }
+    scheduler.Record([bindings_ = std::move(bindings), buffer_handles_ = std::move(buffer_handles)](vk::CommandBuffer cmdbuf) {
+        cmdbuf.BindTransformFeedbackBuffersEXT(0, u32(buffer_handles_.size()), buffer_handles_.data(), bindings_.offsets.data(), bindings_.sizes.data());
     });
 }
 

src/video_core/renderer_vulkan/vk_compute_pass.cpp

@@ -315,7 +315,7 @@ ComputePass::ComputePass(const Device& device_, DescriptorPool& descriptor_pool,
         .requiredSubgroupSize = optional_subgroup_size ? *optional_subgroup_size : 32U,
     };
     bool use_setup_size = device.IsExtSubgroupSizeControlSupported() && optional_subgroup_size;
-    pipeline = device.GetLogical().CreateComputePipeline({
+    pipeline = device.GetLogical().CreateComputePipeline(VkComputePipelineCreateInfo{
         .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
         .pNext = nullptr,
         .flags = 0,
@@ -329,7 +329,7 @@ ComputePass::ComputePass(const Device& device_, DescriptorPool& descriptor_pool,
             .pSpecializationInfo = nullptr,
         },
         .layout = *layout,
-        .basePipelineHandle = nullptr,
+        .basePipelineHandle = {},
         .basePipelineIndex = 0,
     });
 }
@@ -974,7 +974,7 @@ MSAACopyPass::MSAACopyPass(const Device& device_, Scheduler& scheduler_,
             .codeSize = static_cast<u32>(code.size_bytes()),
             .pCode = code.data(),
         });
-        pipelines[i] = device.GetLogical().CreateComputePipeline({
+        pipelines[i] = device.GetLogical().CreateComputePipeline(VkComputePipelineCreateInfo{
             .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
             .pNext = nullptr,
             .flags = 0,
@@ -988,7 +988,7 @@ MSAACopyPass::MSAACopyPass(const Device& device_, Scheduler& scheduler_,
                 .pSpecializationInfo = nullptr,
             },
             .layout = *layout,
-            .basePipelineHandle = nullptr,
+            .basePipelineHandle = {},
             .basePipelineIndex = 0,
         });
     };

src/video_core/renderer_vulkan/vk_compute_pipeline.cpp

@@ -67,8 +67,7 @@ ComputePipeline::ComputePipeline(const Device& device_, vk::PipelineCache& pipel
         if (device.IsKhrPipelineExecutablePropertiesEnabled() && Settings::values.renderer_debug.GetValue()) {
             flags |= VK_PIPELINE_CREATE_CAPTURE_STATISTICS_BIT_KHR;
         }
-        pipeline = device.GetLogical().CreateComputePipeline(
-            {
+        pipeline = device.GetLogical().CreateComputePipeline(VkComputePipelineCreateInfo{
             .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
             .pNext = nullptr,
             .flags = flags,
@@ -85,8 +84,7 @@ ComputePipeline::ComputePipeline(const Device& device_, vk::PipelineCache& pipel
             .layout = *pipeline_layout,
             .basePipelineHandle = 0,
             .basePipelineIndex = 0,
-            },
-            *pipeline_cache);
+        }, *pipeline_cache);
 
         // Log compute pipeline creation
         if (Settings::values.gpu_logging_enabled.GetValue()) {

src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp

@@ -972,8 +972,7 @@ void GraphicsPipeline::MakePipeline(VkRenderPass render_pass) {
         flags |= VK_PIPELINE_CREATE_CAPTURE_STATISTICS_BIT_KHR;
     }
 
-    pipeline = device.GetLogical().CreateGraphicsPipeline(
-        {
+    pipeline = device.GetLogical().CreateGraphicsPipeline({
         .sType = VK_STRUCTURE_TYPE_GRAPHICS_PIPELINE_CREATE_INFO,
         .pNext = nullptr,
         .flags = flags,
@@ -993,8 +992,7 @@ void GraphicsPipeline::MakePipeline(VkRenderPass render_pass) {
         .subpass = 0,
         .basePipelineHandle = nullptr,
         .basePipelineIndex = 0,
-        },
-        *pipeline_cache);
+    }, *pipeline_cache);
 
     // Log graphics pipeline creation
     if (Settings::values.gpu_logging_enabled.GetValue()) {

src/video_core/renderer_vulkan/vk_pipeline_cache.cpp

@@ -260,7 +260,6 @@ Shader::RuntimeInfo MakeRuntimeInfo(std::span<const Shader::IR::Program> program
             info.active_color_outputs[0] = true;
         }
 
-        if (device.IsMoltenVK()) {
         for (size_t i = 0; i < 8; ++i) {
             const auto format = static_cast<Tegra::RenderTargetFormat>(key.state.color_formats[i]);
             const auto pixel_format = VideoCore::Surface::PixelFormatFromRenderTargetFormat(format);
@@ -274,7 +273,6 @@ Shader::RuntimeInfo MakeRuntimeInfo(std::span<const Shader::IR::Program> program
                 info.color_output_types[i] = Shader::AttributeType::Float;
             }
         }
-        }
         break;
     }
     default:
@@ -379,6 +377,7 @@ PipelineCache::PipelineCache(Tegra::MaxwellDeviceMemoryManager& device_memory_,
             float_control.denormBehaviorIndependence == VK_SHADER_FLOAT_CONTROLS_INDEPENDENCE_ALL,
         .support_separate_rounding_mode =
             float_control.roundingModeIndependence == VK_SHADER_FLOAT_CONTROLS_INDEPENDENCE_ALL,
+        .support_fp32_rounding_rtz = float_control.shaderRoundingModeRTZFloat32 != VK_FALSE,
         .support_fp16_denorm_preserve = float_control.shaderDenormPreserveFloat16 != VK_FALSE,
         .support_fp32_denorm_preserve = float_control.shaderDenormPreserveFloat32 != VK_FALSE,
         .support_fp16_denorm_flush = float_control.shaderDenormFlushToZeroFloat16 != VK_FALSE,

src/video_core/renderer_vulkan/vk_present_manager.cpp

@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
 // SPDX-License-Identifier: GPL-3.0-or-later
 
 // SPDX-FileCopyrightText: Copyright 2023 yuzu Emulator Project
@@ -101,22 +101,14 @@ PresentManager::PresentManager(const vk::Instance& instance_,
                                MemoryAllocator& memory_allocator_,
                                Scheduler& scheduler_,
                                Swapchain& swapchain_,
-#ifdef ANDROID
                                vk::SurfaceKHR& surface_)
-#else
-                               VkSurfaceKHR_T* surface_handle_)
-#endif
     : instance{instance_}
     , render_window{render_window_}
     , device{device_}
     , memory_allocator{memory_allocator_}
     , scheduler{scheduler_}
     , swapchain{swapchain_}
-#ifdef ANDROID
     , surface{surface_}
-#else
-    , surface_handle{surface_handle_}
-#endif
     , blit_supported{CanBlitToSwapchain(device.GetPhysical(), swapchain.GetImageViewFormat())}
     , use_present_thread{Settings::values.async_presentation.GetValue()}
 {
@@ -299,11 +291,7 @@ void PresentManager::PresentThread(std::stop_token token) {
 }
 
 void PresentManager::RecreateSwapchain(Frame* frame) {
-#ifndef ANDROID
-    swapchain.Create(surface_handle, frame->width, frame->height); // Pass raw pointer
-#else
     swapchain.Create(*surface, frame->width, frame->height); // Pass raw pointer
-#endif
     SetImageCount();
 }
 

src/video_core/renderer_vulkan/vk_present_manager.h

@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
 // SPDX-License-Identifier: GPL-3.0-or-later
 
 // SPDX-FileCopyrightText: Copyright 2023 yuzu Emulator Project
@@ -15,8 +15,6 @@
 #include "video_core/vulkan_common/vulkan_memory_allocator.h"
 #include "video_core/vulkan_common/vulkan_wrapper.h"
 
-struct VkSurfaceKHR_T;
-
 namespace Core::Frontend {
 class EmuWindow;
 } // namespace Core::Frontend
@@ -46,11 +44,7 @@ public:
                    MemoryAllocator& memory_allocator,
                    Scheduler& scheduler,
                    Swapchain& swapchain,
-#ifdef ANDROID
                    vk::SurfaceKHR& surface);
-#else
-                   VkSurfaceKHR_T* surface_handle);
-#endif
     ~PresentManager();
 
     /// Returns the last used presentation frame
@@ -84,11 +78,7 @@ private:
     MemoryAllocator& memory_allocator;
     Scheduler& scheduler;
     Swapchain& swapchain;
-#ifdef ANDROID
     vk::SurfaceKHR& surface;
-#else
-    VkSurfaceKHR_T* surface_handle;
-#endif
     vk::CommandPool cmdpool;
     std::vector<Frame> frames;
     boost::container::deque<Frame*> present_queue;

src/video_core/renderer_vulkan/vk_query_cache.cpp

@@ -1288,7 +1288,7 @@ void QueryCacheRuntime::EndHostConditionalRendering() {
     PauseHostConditionalRendering();
     impl->hcr_is_set = false;
     impl->is_hcr_running = false;
-    impl->hcr_buffer = nullptr;
+    impl->hcr_buffer = VkBuffer{};
     impl->hcr_offset = 0;
 }
 

src/video_core/renderer_vulkan/vk_query_cache.h

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2023 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-3.0-or-later
 
@@ -35,7 +38,7 @@ public:
     ~QueryCacheRuntime();
 
     template <typename SyncValuesType>
-    void SyncValues(std::span<SyncValuesType> values, VkBuffer base_src_buffer = nullptr);
+    void SyncValues(std::span<SyncValuesType> values, VkBuffer base_src_buffer = VkBuffer{});
 
     void Barriers(bool is_prebarrier);
 

src/video_core/renderer_vulkan/vk_scheduler.cpp

@@ -383,7 +383,7 @@ void Scheduler::EndRenderPass()
                                    vk::Span(barriers.data(), num_images));
         });
 
-        state.renderpass = nullptr;
+        state.renderpass = VkRenderPass{};
         num_renderpass_images = 0;
     }
 

src/video_core/renderer_vulkan/vk_scheduler.h

@@ -44,10 +44,10 @@ public:
     ~Scheduler();
 
     /// Sends the current execution context to the GPU.
-    u64 Flush(VkSemaphore signal_semaphore = nullptr, VkSemaphore wait_semaphore = nullptr);
+    u64 Flush(VkSemaphore signal_semaphore = {}, VkSemaphore wait_semaphore = {});
 
     /// Sends the current execution context to the GPU and waits for it to complete.
-    void Finish(VkSemaphore signal_semaphore = nullptr, VkSemaphore wait_semaphore = nullptr);
+    void Finish(VkSemaphore signal_semaphore = {}, VkSemaphore wait_semaphore = {});
 
     /// Waits for the worker thread to finish executing everything. After this function returns it's
     /// safe to touch worker resources.
@@ -237,8 +237,8 @@ private:
     };
 
     struct State {
-        VkRenderPass renderpass = nullptr;
-        VkFramebuffer framebuffer = nullptr;
+        VkRenderPass renderpass{};
+        VkFramebuffer framebuffer{};
         VkExtent2D render_area = {0, 0};
         GraphicsPipeline* graphics_pipeline = nullptr;
         bool is_rescaling = false;

src/video_core/renderer_vulkan/vk_swapchain.cpp

@@ -109,38 +109,22 @@ VkCompositeAlphaFlagBitsKHR ChooseAlphaFlags(const VkSurfaceCapabilitiesKHR& cap
 } // Anonymous namespace
 
 Swapchain::Swapchain(
-#ifdef ANDROID
-    VkSurfaceKHR surface_,
-#else
-    VkSurfaceKHR_T* surface_handle_,
-#endif
+    VkSurfaceKHR_T* surface_,
     const Device& device_,
     Scheduler& scheduler_,
     u32 width_,
     u32 height_)
-#ifdef ANDROID
     : surface(surface_)
-#else
-    : surface_handle{surface_handle_}
-#endif
     , device{device_}
     , scheduler{scheduler_}
 {
-#ifdef ANDROID
     Create(surface, width_, height_);
-#else
-    Create(surface_handle, width_, height_);
-#endif
 }
 
 Swapchain::~Swapchain() = default;
 
 void Swapchain::Create(
-#ifdef ANDROID
-    VkSurfaceKHR surface_,
-#else
-    VkSurfaceKHR_T* surface_handle_,
-#endif
+    VkSurfaceKHR_T* surface_,
     u32 width_,
     u32 height_)
 {
@@ -148,18 +132,10 @@ void Swapchain::Create(
     is_suboptimal = false;
     width = width_;
     height = height_;
-#ifdef ANDROID
     surface = surface_;
-#else
-    surface_handle = surface_handle_;
-#endif
 
     const auto physical_device = device.GetPhysical();
-#ifdef ANDROID
-    const auto capabilities{physical_device.GetSurfaceCapabilitiesKHR(surface)};
-#else
-    const auto capabilities{physical_device.GetSurfaceCapabilitiesKHR(surface_handle)};
-#endif
+    const auto capabilities{physical_device.GetSurfaceCapabilitiesKHR(VkSurfaceKHR(surface))};
     if (capabilities.maxImageExtent.width == 0 || capabilities.maxImageExtent.height == 0) {
         return;
     }
@@ -254,14 +230,8 @@ void Swapchain::Present(VkSemaphore render_semaphore) {
 
 void Swapchain::CreateSwapchain(const VkSurfaceCapabilitiesKHR& capabilities) {
     const auto physical_device{device.GetPhysical()};
-
-#ifdef ANDROID
-    const auto formats{physical_device.GetSurfaceFormatsKHR(surface)};
-    const auto present_modes = physical_device.GetSurfacePresentModesKHR(surface);
-#else
-    const auto formats{physical_device.GetSurfaceFormatsKHR(surface_handle)};
-    const auto present_modes = physical_device.GetSurfacePresentModesKHR(surface_handle);
-#endif
+    const auto formats{physical_device.GetSurfaceFormatsKHR(VkSurfaceKHR(surface))};
+    const auto present_modes = physical_device.GetSurfacePresentModesKHR(VkSurfaceKHR(surface));
 
     has_mailbox = std::find(present_modes.begin(), present_modes.end(), VK_PRESENT_MODE_MAILBOX_KHR)
                   != present_modes.end();
@@ -290,11 +260,7 @@ void Swapchain::CreateSwapchain(const VkSurfaceCapabilitiesKHR& capabilities) {
         .sType = VK_STRUCTURE_TYPE_SWAPCHAIN_CREATE_INFO_KHR,
         .pNext = nullptr,
         .flags = 0,
-#ifdef ANDROID
-        .surface = surface,
-#else
-        .surface = surface_handle,
-#endif
+        .surface = VkSurfaceKHR(surface),
         .minImageCount = requested_image_count,
         .imageFormat = surface_format.format,
         .imageColorSpace = surface_format.colorSpace,
@@ -313,7 +279,7 @@ void Swapchain::CreateSwapchain(const VkSurfaceCapabilitiesKHR& capabilities) {
         .compositeAlpha = alpha_flags,
         .presentMode = present_mode,
         .clipped = VK_FALSE,
-        .oldSwapchain = nullptr,
+        .oldSwapchain = VkSwapchainKHR{},
     };
     const u32 graphics_family{device.GetGraphicsFamily()};
     const u32 present_family{device.GetPresentFamily()};
@@ -345,11 +311,7 @@ void Swapchain::CreateSwapchain(const VkSurfaceCapabilitiesKHR& capabilities) {
         swapchain_ci.flags |= VK_SWAPCHAIN_CREATE_MUTABLE_FORMAT_BIT_KHR;
     }
     // Request the size again to reduce the possibility of a TOCTOU race condition.
-#ifdef ANDROID
-    const auto updated_capabilities = physical_device.GetSurfaceCapabilitiesKHR(surface);
-#else
-    const auto updated_capabilities = physical_device.GetSurfaceCapabilitiesKHR(surface_handle);
-#endif
+    const auto updated_capabilities = physical_device.GetSurfaceCapabilitiesKHR(VkSurfaceKHR(surface));
     swapchain_ci.imageExtent = ChooseSwapExtent(updated_capabilities, width, height);
     // Don't add code within this and the swapchain creation.
     swapchain = device.GetLogical().CreateSwapchainKHR(swapchain_ci);

src/video_core/renderer_vulkan/vk_swapchain.h

@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
 // SPDX-License-Identifier: GPL-3.0-or-later
 
 // SPDX-FileCopyrightText: Copyright 2019 yuzu Emulator Project
@@ -11,8 +11,6 @@
 #include "common/common_types.h"
 #include "video_core/vulkan_common/vulkan_wrapper.h"
 
-struct VkSurfaceKHR_T;
-
 namespace Layout {
 struct FramebufferLayout;
 }
@@ -25,11 +23,7 @@ class Scheduler;
 class Swapchain {
 public:
     explicit Swapchain(
-#ifdef ANDROID
-        VkSurfaceKHR surface,
-#else
-        VkSurfaceKHR_T* surface_handle,
-#endif
+        VkSurfaceKHR_T* surface,
         const Device& device,
         Scheduler& scheduler,
         u32 width,
@@ -38,11 +32,7 @@ public:
 
     /// Creates (or recreates) the swapchain with a given size.
     void Create(
-#ifdef ANDROID
-        VkSurfaceKHR surface,
-#else
-        VkSurfaceKHR_T* surface_handle,
-#endif
+        VkSurfaceKHR_T* surface,
         u32 width,
         u32 height);
 
@@ -128,11 +118,7 @@ private:
 
     bool NeedsPresentModeUpdate() const;
 
-#ifdef ANDROID
-    VkSurfaceKHR surface;
-#else
-    VkSurfaceKHR_T* surface_handle;
-#endif
+    VkSurfaceKHR_T* surface;
 
     const Device& device;
     Scheduler& scheduler;

src/video_core/renderer_vulkan/vk_texture_cache.cpp

@@ -2314,6 +2314,7 @@ vk::ImageView ImageView::MakeView(VkFormat vk_format, VkImageAspectFlags aspect_
 
 Sampler::Sampler(TextureCacheRuntime& runtime, const Tegra::Texture::TSCEntry& tsc) {
     const auto& device = runtime.device;
+    has_depth_compare = tsc.depth_compare_enabled != 0;
     // Check if custom border colors are supported
     const bool has_custom_border_colors = runtime.device.IsCustomBorderColorsSupported();
     const bool has_format_undefined = runtime.device.IsCustomBorderColorWithoutFormatSupported();
@@ -2354,7 +2355,7 @@ Sampler::Sampler(TextureCacheRuntime& runtime, const Tegra::Texture::TSCEntry& t
     // Some games have samplers with garbage. Sanitize them here.
     const f32 max_anisotropy = std::clamp(tsc.MaxAnisotropy(), 1.0f, 16.0f);
 
-    const auto create_sampler = [&](const f32 anisotropy) {
+    const auto create_sampler = [&](const f32 anisotropy, bool enable_depth_compare) {
         return device.GetLogical().CreateSampler(VkSamplerCreateInfo{
             .sType = VK_STRUCTURE_TYPE_SAMPLER_CREATE_INFO,
             .pNext = pnext,
@@ -2368,7 +2369,7 @@ Sampler::Sampler(TextureCacheRuntime& runtime, const Tegra::Texture::TSCEntry& t
             .mipLodBias = tsc.LodBias(),
             .anisotropyEnable = static_cast<VkBool32>(anisotropy > 1.0f ? VK_TRUE : VK_FALSE),
             .maxAnisotropy = anisotropy,
-            .compareEnable = tsc.depth_compare_enabled,
+            .compareEnable = enable_depth_compare,
             .compareOp = MaxwellToVK::Sampler::DepthCompareFunction(tsc.depth_compare_func),
             .minLod = tsc.mipmap_filter == TextureMipmapFilter::None ? 0.0f : tsc.MinLod(),
             .maxLod = tsc.mipmap_filter == TextureMipmapFilter::None ? 0.25f : tsc.MaxLod(),
@@ -2378,11 +2379,18 @@ Sampler::Sampler(TextureCacheRuntime& runtime, const Tegra::Texture::TSCEntry& t
         });
     };
 
-    sampler = create_sampler(max_anisotropy);
+    sampler = create_sampler(max_anisotropy, has_depth_compare);
+    if (has_depth_compare) {
+        sampler_no_compare = create_sampler(max_anisotropy, false);
+    }
 
     const f32 max_anisotropy_default = static_cast<f32>(1U << tsc.max_anisotropy);
     if (max_anisotropy > max_anisotropy_default) {
-        sampler_default_anisotropy = create_sampler(max_anisotropy_default);
+        sampler_default_anisotropy = create_sampler(max_anisotropy_default, has_depth_compare);
+        if (has_depth_compare) {
+            sampler_default_anisotropy_no_compare =
+                create_sampler(max_anisotropy_default, false);
+        }
     }
 }
 

src/video_core/renderer_vulkan/vk_texture_cache.h

@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
 // SPDX-License-Identifier: GPL-3.0-or-later
 
 // SPDX-FileCopyrightText: Copyright 2019 yuzu Emulator Project
@@ -396,11 +396,18 @@ class Sampler {
 public:
     explicit Sampler(TextureCacheRuntime&, const Tegra::Texture::TSCEntry&);
 
-    [[nodiscard]] VkSampler Handle() const noexcept {
+    [[nodiscard]] VkSampler Handle(bool enable_depth_compare = true) const noexcept {
+        if (!enable_depth_compare && sampler_no_compare) {
+            return *sampler_no_compare;
+        }
         return *sampler;
     }
 
-    [[nodiscard]] VkSampler HandleWithDefaultAnisotropy() const noexcept {
+    [[nodiscard]] VkSampler HandleWithDefaultAnisotropy(
+        bool enable_depth_compare = true) const noexcept {
+        if (!enable_depth_compare && sampler_default_anisotropy_no_compare) {
+            return *sampler_default_anisotropy_no_compare;
+        }
         return *sampler_default_anisotropy;
     }
 
@@ -408,9 +415,16 @@ public:
         return static_cast<bool>(sampler_default_anisotropy);
     }
 
+    [[nodiscard]] bool HasDepthCompareEnabled() const noexcept {
+        return has_depth_compare;
+    }
+
 private:
     vk::Sampler sampler;
+    vk::Sampler sampler_no_compare;
     vk::Sampler sampler_default_anisotropy;
+    vk::Sampler sampler_default_anisotropy_no_compare;
+    bool has_depth_compare = false;
 };
 
 struct TextureCacheParams {

src/video_core/vulkan_common/vulkan.h

@@ -37,3 +37,6 @@
 #undef False
 #undef None
 #undef True
+
+// "Catch-all" handle for both Android and.. the rest of platforms
+struct VkSurfaceKHR_T;

src/video_core/vulkan_common/vulkan_device.cpp

@@ -419,7 +419,7 @@ Device::Device(VkInstance instance_, vk::PhysicalDevice physical_, VkSurfaceKHR
     : instance{instance_}, dld{dld_}, physical{physical_},
     format_properties(GetFormatProperties(physical)) {
     // Get suitability and device properties.
-    const bool is_suitable = GetSuitability(surface != nullptr);
+    const bool is_suitable = GetSuitability(surface != VkSurfaceKHR{});
 
     const VkDriverId driver_id = properties.driver.driverID;
 

src/video_core/vulkan_common/vulkan_surface.cpp

@@ -15,7 +15,7 @@ vk::SurfaceKHR CreateSurface(
     const vk::Instance& instance,
     [[maybe_unused]] const Core::Frontend::EmuWindow::WindowSystemInfo& window_info) {
     [[maybe_unused]] const vk::InstanceDispatch& dld = instance.Dispatch();
-    VkSurfaceKHR unsafe_surface = nullptr;
+    VkSurfaceKHR unsafe_surface = VkSurfaceKHR{};
 
 #ifdef _WIN32
     if (window_info.type == Core::Frontend::WindowSystemType::Windows) {

src/video_core/vulkan_common/vulkan_wrapper.h

@@ -395,13 +395,13 @@ public:
 
     /// Construct a handle transferring the ownership from another handle.
     Handle(Handle&& rhs) noexcept
-        : handle{std::exchange(rhs.handle, nullptr)}, owner{rhs.owner}, dld{rhs.dld} {}
+        : handle{std::exchange(rhs.handle, Type{})}, owner{rhs.owner}, dld{rhs.dld} {}
 
     /// Assign the current handle transferring the ownership from another handle.
     /// Destroys any previously held object.
     Handle& operator=(Handle&& rhs) noexcept {
         Release();
-        handle = std::exchange(rhs.handle, nullptr);
+        handle = std::exchange(rhs.handle, Type{});
         owner = rhs.owner;
         dld = rhs.dld;
         return *this;
@@ -415,7 +415,7 @@ public:
     /// Destroys any held object.
     void reset() noexcept {
         Release();
-        handle = nullptr;
+        handle = Type{};
     }
 
     /// Returns the address of the held object.
@@ -431,7 +431,7 @@ public:
 
     /// Returns true when there's a held object.
     explicit operator bool() const noexcept {
-        return handle != nullptr;
+        return handle != Type{};
     }
 
 #ifndef ANDROID
@@ -446,7 +446,7 @@ public:
 #endif
 
 protected:
-    Type handle = nullptr;
+    Type handle{};
     OwnerType owner = nullptr;
     const Dispatch* dld = nullptr;
 
@@ -454,7 +454,7 @@ private:
     /// Destroys the held object if it exists.
     void Release() noexcept {
         if (handle) {
-            Destroy(owner, handle, *dld);
+            Destroy(OwnerType(owner), Type(handle), *dld);
         }
     }
 };
@@ -497,7 +497,7 @@ public:
     /// Destroys any held object.
     void reset() noexcept {
         Release();
-        handle = nullptr;
+        handle = {};
     }
 
     /// Returns the address of the held object.
@@ -513,7 +513,7 @@ public:
 
     /// Returns true when there's a held object.
     explicit operator bool() const noexcept {
-        return handle != nullptr;
+        return handle != Type{};
     }
 
 #ifndef ANDROID
@@ -528,7 +528,7 @@ public:
 #endif
 
 protected:
-    Type handle = nullptr;
+    Type handle{};
     const Dispatch* dld = nullptr;
 
 private:
@@ -598,7 +598,7 @@ private:
     std::unique_ptr<AllocationType[]> allocations;
     std::size_t num = 0;
     VkDevice device = nullptr;
-    PoolType pool = nullptr;
+    PoolType pool{};
     const DeviceDispatch* dld = nullptr;
 };
 
@@ -660,12 +660,12 @@ public:
     Image& operator=(const Image&) = delete;
 
     Image(Image&& rhs) noexcept
-        : handle{std::exchange(rhs.handle, nullptr)}, usage{rhs.usage}, owner{rhs.owner},
+        : handle{std::exchange(rhs.handle, VkImage{})}, usage{rhs.usage}, owner{rhs.owner},
           allocator{rhs.allocator}, allocation{rhs.allocation}, dld{rhs.dld} {}
 
     Image& operator=(Image&& rhs) noexcept {
         Release();
-        handle = std::exchange(rhs.handle, nullptr);
+        handle = std::exchange(rhs.handle, VkImage{});
         usage = rhs.usage;
         owner = rhs.owner;
         allocator = rhs.allocator;
@@ -684,11 +684,11 @@ public:
 
     void reset() noexcept {
         Release();
-        handle = nullptr;
+        handle = VkImage{};
     }
 
     explicit operator bool() const noexcept {
-        return handle != nullptr;
+        return handle != VkImage{};
     }
 
     void SetObjectNameEXT(const char* name) const;
@@ -700,7 +700,7 @@ public:
 private:
     void Release() const noexcept;
 
-    VkImage handle = nullptr;
+    VkImage handle{};
     VkImageUsageFlags usage{};
     VkDevice owner = nullptr;
     VmaAllocator allocator = nullptr;
@@ -721,13 +721,13 @@ public:
     Buffer& operator=(const Buffer&) = delete;
 
     Buffer(Buffer&& rhs) noexcept
-        : handle{std::exchange(rhs.handle, nullptr)}, owner{rhs.owner}, allocator{rhs.allocator},
+        : handle{std::exchange(rhs.handle, VkBuffer{})}, owner{rhs.owner}, allocator{rhs.allocator},
           allocation{rhs.allocation}, mapped{rhs.mapped},
           is_coherent{rhs.is_coherent}, dld{rhs.dld} {}
 
     Buffer& operator=(Buffer&& rhs) noexcept {
         Release();
-        handle = std::exchange(rhs.handle, nullptr);
+        handle = std::exchange(rhs.handle, VkBuffer{});
         owner = rhs.owner;
         allocator = rhs.allocator;
         allocation = rhs.allocation;
@@ -747,11 +747,11 @@ public:
 
     void reset() noexcept {
         Release();
-        handle = nullptr;
+        handle = VkBuffer{};
     }
 
     explicit operator bool() const noexcept {
-        return handle != nullptr;
+        return handle != VkBuffer{};
     }
 
     /// Returns the host mapped memory, an empty span otherwise.
@@ -777,7 +777,7 @@ public:
 private:
     void Release() const noexcept;
 
-    VkBuffer handle = nullptr;
+    VkBuffer handle{};
     VkDevice owner = nullptr;
     VmaAllocator allocator = nullptr;
     VmaAllocation allocation = nullptr;
@@ -1011,10 +1011,10 @@ public:
     [[nodiscard]] PipelineLayout CreatePipelineLayout(const VkPipelineLayoutCreateInfo& ci) const;
 
     [[nodiscard]] Pipeline CreateGraphicsPipeline(const VkGraphicsPipelineCreateInfo& ci,
-                                                  VkPipelineCache cache = nullptr) const;
+                                                  VkPipelineCache cache = {}) const;
 
     [[nodiscard]] Pipeline CreateComputePipeline(const VkComputePipelineCreateInfo& ci,
-                                                 VkPipelineCache cache = nullptr) const;
+                                                 VkPipelineCache cache = {}) const;
 
     [[nodiscard]] Sampler CreateSampler(const VkSamplerCreateInfo& ci) const;
 

src/yuzu/configuration/configure_per_game_addons.cpp

@@ -10,13 +10,14 @@
 
 #include <fmt/format.h>
 
+#include <QDesktopServices>
 #include <QHeaderView>
 #include <QMenu>
 #include <QStandardItemModel>
+#include <QStandardPaths>
 #include <QString>
 #include <QTimer>
 #include <QTreeView>
-#include <QStandardPaths>
 
 #include "common/common_types.h"
 #include "common/fs/fs.h"
@@ -42,7 +43,7 @@ ConfigurePerGameAddons::ConfigurePerGameAddons(Core::System& system_, QWidget* p
     item_model = new QStandardItemModel(tree_view);
     tree_view->setModel(item_model);
     tree_view->setAlternatingRowColors(true);
-    tree_view->setSelectionMode(QHeaderView::MultiSelection);
+    tree_view->setSelectionMode(QHeaderView::ExtendedSelection);
     tree_view->setSelectionBehavior(QHeaderView::SelectRows);
     tree_view->setVerticalScrollMode(QHeaderView::ScrollPerPixel);
     tree_view->setHorizontalScrollMode(QHeaderView::ScrollPerPixel);
@@ -248,8 +249,11 @@ void ConfigurePerGameAddons::AddonDeleteRequested(QList<QModelIndex> selected) {
 
 void ConfigurePerGameAddons::showContextMenu(const QPoint& pos) {
     const QModelIndex index = tree_view->indexAt(pos);
-    auto selected = tree_view->selectionModel()->selectedIndexes();
-    if (index.isValid() && selected.empty()) selected = {index};
+    auto selected = tree_view->selectionModel()->selectedRows();
+    if (index.isValid() && selected.empty()) {
+        QModelIndex idx = item_model->index(index.row(), 0);
+        if (idx.isValid()) selected << idx;
+    }
 
     if (selected.empty()) return;
 
@@ -260,6 +264,15 @@ void ConfigurePerGameAddons::showContextMenu(const QPoint& pos) {
         AddonDeleteRequested(selected);
     });
 
+    if (selected.length() == 1) {
+        auto loc = selected.at(0).data(PATCH_LOCATION).toString();
+        if (QFileInfo::exists(loc)) {
+            QAction* open = menu.addAction(tr("&Open in File Manager"));
+            connect(open, &QAction::triggered, this,
+                    [selected, loc]() { QDesktopServices::openUrl(QUrl::fromLocalFile(loc)); });
+        }
+    }
+
     menu.exec(tree_view->viewport()->mapToGlobal(pos));
 }