Merge pull request #26 from Djihads80/nx2-experimental

Nx2 experimental

65 files changed, 3123 insertions, 218 deletions

diff --git a/src/common/container/unordered_map.h b/src/common/container/unordered_map.h
index 7d9e49fd80..c4d4a72663 100644
--- a/src/common/container/unordered_map.h
+++ b/src/common/container/unordered_map.h
@@ -9,7 +9,7 @@
 
 namespace Common {
 
-template <class Key, class T, class Hash = std::hash<Key>, class Pred = std::equal_to<Key>,
+template <class Key, class T, class Hash = FastHash<Key>, class Pred = std::equal_to<Key>,
           class Allocator = std::allocator<std::pair<const Key, T>>>
 using unordered_map = boost::unordered::unordered_flat_map<Key, T, Hash, Pred, Allocator>;
 
diff --git a/src/common/container/unordered_set.h b/src/common/container/unordered_set.h
index 54d71406a1..690c33286f 100644
--- a/src/common/container/unordered_set.h
+++ b/src/common/container/unordered_set.h
@@ -9,7 +9,7 @@
 
 namespace Common {
 
-template <class Key, class Hash = std::hash<Key>, class Pred = std::equal_to<Key>,
+template <class Key, class Hash = FastHash<Key>, class Pred = std::equal_to<Key>,
           class Allocator = std::allocator<Key>>
 using unordered_set = boost::unordered::unordered_flat_set<Key, Hash, Pred, Allocator>;
 
diff --git a/src/common/container_hash.h b/src/common/container_hash.h
index d7477f7e4f..f06f69ad8b 100644
--- a/src/common/container_hash.h
+++ b/src/common/container_hash.h
@@ -19,6 +19,28 @@
 
 namespace Common {
 
+/// The default hash of Common::unordered_map/set. std::hash of an integer or a pointer is a
+/// byte-wise FNV-1a on MSVC, and boost's flat containers mix a hash that does not say it
+/// avalanches anyway: pass those keys through (the GPU thread's page maps looked keys up
+/// tens of thousands of times a second with it). Other keys keep std::hash.
+template <class Key>
+struct FastHash : std::hash<Key> {};
+
+template <class Key>
+    requires std::is_integral_v<Key>
+struct FastHash<Key> {
+    std::size_t operator()(Key key) const noexcept {
+        return static_cast<std::size_t>(key);
+    }
+};
+
+template <class T>
+struct FastHash<T*> {
+    std::size_t operator()(T* pointer) const noexcept {
+        return reinterpret_cast<std::size_t>(pointer);
+    }
+};
+
 namespace detail {
 
 template <typename T>
diff --git a/src/common/cpu_features.cpp b/src/common/cpu_features.cpp
index 52fd003c44..93eecb1dc5 100644
--- a/src/common/cpu_features.cpp
+++ b/src/common/cpu_features.cpp
@@ -278,8 +278,16 @@ std::chrono::milliseconds WallClock::GetTimeMS() const {
 
 s64 WallClock::GetCNTPCT() const {
     if (!invariant)
-        return GetUptime() * NsToCNTPCTRatio::num / NsToCNTPCTRatio::den;
-    return MultiplyHigh(GetUptime(), cntpct_rdtsc_factor);
+        return MultiplyAndDivide64(GetUptime(), GetGuestCNTFRQ(), NsRatio::den);
+    return MultiplyHigh(GetUptime(), cntpct_rdtsc_factor.load(std::memory_order_relaxed));
+}
+
+void WallClock::SetGuestCNTFRQ(u64 frequency) {
+    guest_cntfrq.store(frequency, std::memory_order_relaxed);
+    if (invariant) {
+        cntpct_rdtsc_factor.store(GetFixedPoint64Factor(frequency, rdtsc_frequency),
+                                  std::memory_order_relaxed);
+    }
 }
 
 s64 WallClock::GetGPUTick() const {
@@ -336,7 +344,7 @@ namespace {
 } // namespace
 
 WallClock::WallClock(bool invariant_, u64 rdtsc_frequency_) noexcept {
-    const u64 host_cntfrq = std::max<u64>(GetHostCNTFRQ(), 1);
+    host_cntfrq = std::max<u64>(GetHostCNTFRQ(), 1);
     ns_cntfrq_factor = GetFixedPointFactor(NsRatio::den, host_cntfrq);
     us_cntfrq_factor = GetFixedPointFactor(UsRatio::den, host_cntfrq);
     ms_cntfrq_factor = GetFixedPointFactor(MsRatio::den, host_cntfrq);
@@ -361,6 +369,11 @@ s64 WallClock::GetCNTPCT() const {
     return MultiplyHigh(GetUptime(), guest_cntfrq_factor);
 }
 
+void WallClock::SetGuestCNTFRQ(u64 frequency) {
+    guest_cntfrq.store(frequency, std::memory_order_relaxed);
+    guest_cntfrq_factor = GetFixedPointFactor(frequency, host_cntfrq);
+}
+
 s64 WallClock::GetGPUTick() const {
     return MultiplyHigh(GetUptime(), gputick_cntfrq_factor);
 }
@@ -399,7 +412,11 @@ std::chrono::milliseconds WallClock::GetTimeMS() const {
 }
 
 s64 WallClock::GetCNTPCT() const {
-    return GetUptime() * NsToCNTPCTRatio::num / NsToCNTPCTRatio::den;
+    return MultiplyAndDivide64(GetUptime(), GetGuestCNTFRQ(), NsRatio::den);
+}
+
+void WallClock::SetGuestCNTFRQ(u64 frequency) {
+    guest_cntfrq.store(frequency, std::memory_order_relaxed);
 }
 
 s64 WallClock::GetGPUTick() const {
@@ -419,9 +436,13 @@ u64 WallClock::NsToTicks(std::chrono::nanoseconds ns) const {
 }
 #endif
 
+u64 WallClock::CPUTickToGuestCNTPCT(u64 cpu_tick) const {
+    return MultiplyAndDivide64(cpu_tick, GetGuestCNTFRQ(), CPUTickFreq);
+}
+
 // Wall clock MUST be initialized AFTER g_cpu_caps
 // C++ only guarantees ctor init in the order they appear in TU
-const WallClock g_wall_clock = [] {
+WallClock g_wall_clock = [] {
 #if defined(ARCHITECTURE_x86_64)
     auto const& caps = Common::g_cpu_caps;
     return WallClock(caps.invariant_tsc && caps.tsc_frequency > std::nano::den, caps.tsc_frequency);
diff --git a/src/common/cpu_features.h b/src/common/cpu_features.h
index 61aa2767ab..31be2dc956 100644
--- a/src/common/cpu_features.h
+++ b/src/common/cpu_features.h
@@ -7,6 +7,7 @@
 
 #pragma once
 
+#include <atomic>
 #include <optional>
 #include <string_view>
 #include <chrono>
@@ -20,11 +21,24 @@ namespace Common {
 class WallClock {
 public:
     static constexpr u64 CNTFRQ = 19'200'000;         // CNTPCT_EL0 Frequency = 19.2 MHz
+    static constexpr u64 Switch2CNTFRQ = 31'250'000;  // Switch 2 CNTPCT_EL0 Frequency = 31.25 MHz
     static constexpr u64 GPUTickFreq = 614'400'000;   // GM20B GPU Tick Frequency = 614.4 MHz
     static constexpr u64 CPUTickFreq = 1'020'000'000; // T210/4 A57 CPU Tick Frequency = 1020.0 MHz
 
     explicit WallClock(bool invariant, u64 rdtsc_frequency_) noexcept;
 
+    /// Sets the guest counter's frequency: CNTFRQ, or Switch2CNTFRQ for Switch 2 titles (their SDK
+    /// converts ticks with a constant 31.25 MHz). Called at boot, before guest code runs.
+    void SetGuestCNTFRQ(u64 frequency);
+
+    /// @returns The guest counter's frequency (CNTFRQ_EL0).
+    [[nodiscard]] u64 GetGuestCNTFRQ() const {
+        return guest_cntfrq.load(std::memory_order_relaxed);
+    }
+
+    /// @returns Emulated CPU ticks (single core timing) as guest CNTPCT ticks.
+    [[nodiscard]] u64 CPUTickToGuestCNTPCT(u64 cpu_tick) const;
+
     /// @returns The time in nanoseconds since the construction of this clock.
     std::chrono::nanoseconds GetTimeNS() const;
 
@@ -34,7 +48,7 @@ public:
     /// @returns The time in milliseconds since the construction of this clock.
     std::chrono::milliseconds GetTimeMS() const;
 
-    /// @returns The guest CNTPCT ticks since the construction of this clock.
+    /// @returns The guest CNTPCT ticks (at GetGuestCNTFRQ()) since the construction of this clock.
     s64 GetCNTPCT() const;
 
     /// @returns The guest GPU ticks since the construction of this clock.
@@ -91,6 +105,7 @@ public:
     using CPUTickToCNTPCTRatio = std::ratio<CNTFRQ, CPUTickFreq>;
     using CPUTickToGPUTickRatio = std::ratio<GPUTickFreq, CPUTickFreq>;
 
+    std::atomic<u64> guest_cntfrq{CNTFRQ};
 #if defined(ARCHITECTURE_x86_64)
     u64 rdtsc_frequency;
     u64 ns_rdtsc_factor;
@@ -98,7 +113,7 @@ public:
     u64 ms_rdtsc_factor;
     u64 rdtsc_ns_integer;
     u64 rdtsc_ns_factor;
-    u64 cntpct_rdtsc_factor;
+    std::atomic<u64> cntpct_rdtsc_factor;
     u64 gputick_rdtsc_factor;
     bool invariant;
 #elif defined(HAS_NCE)
@@ -106,6 +121,7 @@ public:
     [[nodiscard]] inline FactorType GetGuestCNTFRQFactor() const {
         return guest_cntfrq_factor;
     }
+    u64 host_cntfrq;
     FactorType ns_cntfrq_factor;
     FactorType us_cntfrq_factor;
     FactorType ms_cntfrq_factor;
@@ -183,6 +199,6 @@ std::optional<int> GetProcessorCount();
 /// @brief Global cpu caps
 extern const CPUCaps g_cpu_caps;
 /// @brief Global wall clock
-extern const WallClock g_wall_clock;
+extern WallClock g_wall_clock;
 
 } // namespace Common
diff --git a/src/common/range_sets.h b/src/common/range_sets.h
index f8fcee4837..e5cb53b058 100644
--- a/src/common/range_sets.h
+++ b/src/common/range_sets.h
@@ -26,6 +26,9 @@ public:
     void Clear();
     bool Empty() const;
 
+    /// Whether any range overlaps [base_address, base_address + size): one lookup.
+    bool Intersects(AddressType base_address, size_t size) const;
+
     template <typename Func>
     void ForEach(Func&& func) const;
 
diff --git a/src/common/range_sets.inc b/src/common/range_sets.inc
index 3edd8c8a43..aeda7bc351 100644
--- a/src/common/range_sets.inc
+++ b/src/common/range_sets.inc
@@ -116,28 +116,25 @@ struct OverlapRangeSet<AddressType>::OverlapRangeSetImpl {
         }
         AddressType end_address = base_address + static_cast<AddressType>(size);
         IntervalType interval{base_address, end_address};
-        bool any_removals = false;
         m_split_ranges_set += std::make_pair(interval, -amount);
-        do {
-            any_removals = false;
-            auto it = m_split_ranges_set.lower_bound(interval);
-            if (it == m_split_ranges_set.end()) {
-                return;
+        // One pass, erasing the segments that dropped to zero or below in address order (the
+        // search restarted after every erase: quadratic in the segments of the range, and the
+        // GPU fencing thread does this for each downloaded buffer range while holding the buffer
+        // cache's lock). Erasing a node leaves the other iterators valid.
+        auto it = m_split_ranges_set.lower_bound(interval);
+        const auto end_it = m_split_ranges_set.upper_bound(interval);
+        while (it != end_it) {
+            if (it->second > 0) {
+                ++it;
+                continue;
             }
-            auto end_it = m_split_ranges_set.upper_bound(interval);
-            for (; it != end_it; it++) {
-                if (it->second <= 0) {
-                    if constexpr (has_on_delete) {
-                        if (it->second == 0) {
-                            on_delete(it->first.lower(), it->first.upper());
-                        }
-                    }
-                    any_removals = true;
-                    m_split_ranges_set.erase(it);
-                    break;
+            if constexpr (has_on_delete) {
+                if (it->second == 0) {
+                    on_delete(it->first.lower(), it->first.upper());
                 }
             }
-        } while (any_removals);
+            m_split_ranges_set.erase(it++);
+        }
     }
 
     template <typename Func>
@@ -224,6 +221,18 @@ bool RangeSet<AddressType>::Empty() const {
 }
 
 template <typename AddressType>
+bool RangeSet<AddressType>::Intersects(AddressType base_address, size_t size) const {
+    const auto& set = m_impl->m_ranges_set;
+    if (set.empty() || size == 0) {
+        return false;
+    }
+    const AddressType end_address = base_address + static_cast<AddressType>(size);
+    // The first range not entirely below the query; it overlaps when it starts before its end.
+    const auto it = set.lower_bound(typename RangeSetImpl::IntervalType{base_address, end_address});
+    return it != set.end() && it->lower() < end_address;
+}
+
+template <typename AddressType>
 template <typename Func>
 void RangeSet<AddressType>::ForEach(Func&& func) const {
     m_impl->ForEach(std::move(func));
diff --git a/src/core/arm/dynarmic/arm_dynarmic_64.cpp b/src/core/arm/dynarmic/arm_dynarmic_64.cpp
index c2bb05afba..e65e38dd9e 100644
--- a/src/core/arm/dynarmic/arm_dynarmic_64.cpp
+++ b/src/core/arm/dynarmic/arm_dynarmic_64.cpp
@@ -5,6 +5,10 @@
 // SPDX-License-Identifier: GPL-2.0-or-later
 
 #include <atomic>
+#include <cstdlib>
+#if defined(_M_X64) || defined(__x86_64__)
+#include <immintrin.h>
+#endif
 #include "common/settings.h"
 #include "core/arm/dynarmic/arm_dynarmic.h"
 #include "core/arm/dynarmic/arm_dynarmic_64.h"
@@ -166,7 +170,26 @@ u64 DynarmicCallbacks64::GetTicksRemaining() {
 }
 
 u64 DynarmicCallbacks64::GetCNTPCT() {
-    return m_parent.m_system.CoreTiming().GetClockTicks();
+    const u64 ticks = m_parent.m_system.CoreTiming().GetClockTicks();
+    // Guest code that waits by polling the counter (title 1's main thread reads it ~17 million
+    // times a second while waiting for the GPU) keeps a host core busy at full speed, and the
+    // GPU thread is often on the other half of that SMT core. Two reads within ~1.7 us of each
+    // other are such a loop: pause, which hands the core's resources to the sibling thread for a
+    // moment. The guest sees the same counter. EDEN_NX2_CNTPCT_PAUSE=0 turns it off.
+    static const bool pause_spins = [] {
+        const char* value = std::getenv("EDEN_NX2_CNTPCT_PAUSE");
+        return value == nullptr || value[0] != '0';
+    }();
+    if (pause_spins && ticks - m_last_cntpct < m_spin_ticks) {
+#if defined(_M_X64) || defined(__x86_64__)
+        _mm_pause();
+        _mm_pause();
+#elif defined(__aarch64__)
+        asm volatile("yield");
+#endif
+    }
+    m_last_cntpct = ticks;
+    return ticks;
 }
 
 bool DynarmicCallbacks64::CheckMemoryAccess(u64 addr, u64 size, Kernel::DebugWatchpointType type) {
@@ -246,7 +269,9 @@ void ArmDynarmic64::MakeJit(Common::PageTable* page_table, std::size_t address_s
     config.tpidr_el0 = &m_cb->m_tpidr_el0;
     config.dczid_el0 = 4;
     config.ctr_el0 = 0x8444c004;
-    config.cntfrq_el0 = Hardware::CNTFREQ;
+    // 31.25 MHz for Switch 2 titles (Common::WallClock::SetGuestCNTFRQ, set at boot).
+    config.cntfrq_el0 = static_cast<u32>(Common::g_wall_clock.GetGuestCNTFRQ());
+    m_cb->m_spin_ticks = config.cntfrq_el0 / 600'000;
 
     // Unpredictable instructions
     config.define_unpredictable_behaviour = true;
diff --git a/src/core/arm/dynarmic/arm_dynarmic_64.h b/src/core/arm/dynarmic/arm_dynarmic_64.h
index 5ea5dde8a6..66b07474c1 100644
--- a/src/core/arm/dynarmic/arm_dynarmic_64.h
+++ b/src/core/arm/dynarmic/arm_dynarmic_64.h
@@ -65,6 +65,8 @@ public:
     Core::Memory::Memory& m_memory;
     u64 m_tpidrro_el0{};
     u64 m_tpidr_el0{};
+    u64 m_last_cntpct{}; ///< The previous counter read (GetCNTPCT spin detection)
+    u64 m_spin_ticks{32}; ///< Reads closer than this (~1.7 us) count as a polling loop
     Kernel::KProcess* m_process{};
     const bool m_debugger_enabled{};
     const bool m_check_memory_access{};
diff --git a/src/core/arm/nce/patcher.cpp b/src/core/arm/nce/patcher.cpp
index f21d7cacdc..72c3497f03 100644
--- a/src/core/arm/nce/patcher.cpp
+++ b/src/core/arm/nce/patcher.cpp
@@ -627,7 +627,7 @@ void Patcher::WriteMsrHandler(ModuleDestLabel module_dest, oaknut::XReg src_reg,
 }
 
 void Patcher::WriteCntfrqHandler(ModuleDestLabel module_dest, oaknut::XReg dest_reg, oaknut::VectorCodeGenerator& cg) {
-    cg.MOV(dest_reg, Common::WallClock::CNTFRQ);
+    cg.MOV(dest_reg, Common::g_wall_clock.GetGuestCNTFRQ());
 
     // Jump back to the instruction after the emulated MRS.
     if (&cg == &c_pre)
@@ -642,6 +642,7 @@ void Patcher::WriteCntpctHandler(ModuleDestLabel module_dest, oaknut::XReg dest_
 #else
     static Common::WallClock clock(true, 1);
 #endif
+    clock.SetGuestCNTFRQ(Common::g_wall_clock.GetGuestCNTFRQ());
     const auto factor = clock.GetGuestCNTFRQFactor();
     const auto raw_factor = std::bit_cast<std::array<u64, 2>>(factor);
 
diff --git a/src/core/core.cpp b/src/core/core.cpp
index f3d224b1ce..3728ccf8f7 100644
--- a/src/core/core.cpp
+++ b/src/core/core.cpp
@@ -2,11 +2,13 @@
 // SPDX-License-Identifier: GPL-3.0-or-later
 
 #include <array>
+#include <cstdlib>
 #include <atomic>
 #include <memory>
 #include <utility>
 
 #include "game_settings.h"
+#include "common/cpu_features.h"
 #include "audio_core/audio_core.h"
 #include "common/fs/fs.h"
 #include "common/logging.h"
@@ -313,16 +315,33 @@ struct System::Impl {
             fs_controller.InitTempStorage();
         }
 
+        const bool switch2_program =
+            virtual_filesystem != nullptr &&
+            IsSwitch2Program(GetGameFileFromPath(virtual_filesystem, filepath));
+
         // Switch 2 titles size their heap from the total memory and need the 12 GiB layout. Set it
         // as a per-game value, so RestoreGlobalState brings back the configured layout afterwards.
-        if (virtual_filesystem != nullptr &&
-            Settings::values.memory_layout_mode.GetValue() != Settings::MemoryLayout::Memory_12Gb &&
-            IsSwitch2Program(GetGameFileFromPath(virtual_filesystem, filepath))) {
+        if (switch2_program &&
+            Settings::values.memory_layout_mode.GetValue() != Settings::MemoryLayout::Memory_12Gb) {
             LOG_INFO(Core, "Switch 2 program: using the 12 GiB memory layout");
             Settings::values.memory_layout_mode.SetGlobal(false);
             Settings::values.memory_layout_mode.SetValue(Settings::MemoryLayout::Memory_12Gb);
         }
 
+        // The Switch 2 counter runs at 31.25 MHz and its SDK converts ticks with that constant
+        // (nn::os::GetSystemTickFrequency, TickManager::ConvertToTimeSpan/ConvertToTick). At 19.2 MHz
+        // every tick-based time in the game was off by 1.63x: elapsed times read short and a frame
+        // deadline in ticks waited ~27 ms. EDEN_NX2_CNTFRQ=<Hz> overrides it (A/B tests).
+        static const u64 forced_cntfrq = [] {
+            const char* value = std::getenv("EDEN_NX2_CNTFRQ");
+            return value != nullptr ? std::strtoull(value, nullptr, 10) : 0ULL;
+        }();
+        const u64 cntfrq = forced_cntfrq != 0 ? forced_cntfrq
+                           : switch2_program  ? Common::WallClock::Switch2CNTFRQ
+                                              : Common::WallClock::CNTFRQ;
+        Common::g_wall_clock.SetGuestCNTFRQ(cntfrq);
+        LOG_INFO(Core, "Guest counter frequency: {} Hz", cntfrq);
+
         InitializeKernel(system);
 
         if (params.applet_type == Service::AM::AppletType::Application) {
diff --git a/src/core/core_timing.cpp b/src/core/core_timing.cpp
index b85bd30a58..dfc3f6f2a6 100644
--- a/src/core/core_timing.cpp
+++ b/src/core/core_timing.cpp
@@ -224,7 +224,7 @@ void CoreTiming::ResetTicks() {
 }
 
 u64 CoreTiming::GetClockTicks() const {
-    u64 fres = is_multicore ? Common::g_wall_clock.GetCNTPCT() : Common::WallClock::CPUTickToCNTPCT(cpu_ticks);
+    u64 fres = is_multicore ? Common::g_wall_clock.GetCNTPCT() : Common::g_wall_clock.CPUTickToGuestCNTPCT(cpu_ticks);
     if (const u32 target = CpuClockTargetMhz(Settings::values.cpu_clock.GetValue());
         target != CPU_CLOCK_BASE_MHZ) {
         fres = fres * target / CPU_CLOCK_BASE_MHZ;
diff --git a/src/core/device_memory_manager.h b/src/core/device_memory_manager.h
index 9ed0b6f5e5..46cda14e78 100644
--- a/src/core/device_memory_manager.h
+++ b/src/core/device_memory_manager.h
@@ -132,6 +132,27 @@ private:
         u8* host_ptr{};
     };
 
+    /// A thread's last four page translations of one manager. Several threads read through
+    /// them (the GPU thread, the fencing thread, service threads): one cache shared by all was
+    /// cleared by Map on a service thread between the GPU thread's null check of an entry and
+    /// its use, and the GPU thread copied from address 0 (title 1 crashed every few minutes).
+    struct TranslationCache {
+        const void* owner{};
+        u64 generation{};
+        std::array<TranslationEntry, 4> slots{};
+        u32 cursor{};
+    };
+    static inline thread_local TranslationCache t_cache{};
+
+    /// This thread's translations, emptied when Map or Unmap ran since they were made.
+    TranslationCache& Translations() {
+        const u64 current = translation_generation.load(std::memory_order_acquire);
+        if (t_cache.owner != this || t_cache.generation != current) {
+            t_cache = TranslationCache{.owner = this, .generation = current};
+        }
+        return t_cache;
+    }
+
     // Internal helper that performs the update assuming the caller already holds the necessary lock.
     void UpdatePagesCachedCountNoLock(DAddr addr, size_t size, s32 delta);
 
@@ -204,8 +225,7 @@ private:
         tracked_entries.GetUnchecked(page_index).cpu_backing_address = address | (asid.id << asid_start_bit);
     }
 
-    std::array<TranslationEntry, 4> t_slot{};
-    u32 cache_cursor = 0;
+    std::atomic<u64> translation_generation{1}; ///< Bumped by Map and Unmap
     using CounterType = u8;
     using CounterAtomicType = std::atomic_uint8_t;
     static constexpr size_t subentries = 8 / sizeof(CounterType);
diff --git a/src/core/device_memory_manager.inc b/src/core/device_memory_manager.inc
index 97c2706abb..d70444ef8c 100644
--- a/src/core/device_memory_manager.inc
+++ b/src/core/device_memory_manager.inc
@@ -234,7 +234,7 @@ void DeviceMemoryManager<Traits>::Map(DAddr address, VAddr virtual_address, size
         }
         impl->multi_dev_address.Register(new_dev, start_id);
     }
-    t_slot = {};
+    translation_generation.fetch_add(1, std::memory_order_release);
     if (track) {
         TrackContinuityImpl(address, virtual_address, size, asid);
     }
@@ -268,7 +268,7 @@ void DeviceMemoryManager<Traits>::Unmap(DAddr address, size_t size) {
             base_dev = new_start | MULTI_FLAG;
         }
     }
-    t_slot = {};
+    translation_generation.fetch_add(1, std::memory_order_release);
 }
 template <typename Traits>
 void DeviceMemoryManager<Traits>::TrackContinuityImpl(DAddr address, VAddr virtual_address,
@@ -404,9 +404,10 @@ void DeviceMemoryManager<Traits>::ReadBlock(DAddr address, void* dest_pointer, s
     const std::size_t page_offset = address & Memory::YUZU_PAGEMASK;
     if (size <= Memory::YUZU_PAGESIZE - page_offset) {
         const DAddr guest_page = address & ~static_cast<DAddr>(Memory::YUZU_PAGEMASK);
-        for (size_t i = 0; i < 4; ++i) {
-            if (t_slot[i].guest_page == guest_page && t_slot[i].host_ptr != nullptr) {
-                std::memcpy(dest_pointer, t_slot[i].host_ptr + page_offset, size);
+        TranslationCache& cache = Translations();
+        for (const TranslationEntry& entry : cache.slots) {
+            if (entry.guest_page == guest_page && entry.host_ptr != nullptr) {
+                std::memcpy(dest_pointer, entry.host_ptr + page_offset, size);
                 return;
             }
         }
@@ -415,8 +416,8 @@ void DeviceMemoryManager<Traits>::ReadBlock(DAddr address, void* dest_pointer, s
         const auto phys_addr = tracked_entries[page_index].compressed_physical_ptr;
         if (phys_addr != 0) {
             auto* const mem_ptr = GetPointerFromRaw<u8>((PAddr(phys_addr - 1) << Memory::YUZU_PAGEBITS));
-            t_slot[cache_cursor % t_slot.size()] = TranslationEntry{.guest_page = guest_page, .host_ptr = mem_ptr};
-            cache_cursor = (cache_cursor + 1) & 3U;
+            cache.slots[cache.cursor] = TranslationEntry{.guest_page = guest_page, .host_ptr = mem_ptr};
+            cache.cursor = (cache.cursor + 1) & 3U;
             std::memcpy(dest_pointer, mem_ptr + page_offset, size);
             return;
         }
@@ -462,9 +463,10 @@ void DeviceMemoryManager<Traits>::ReadBlockUnsafe(DAddr address, void* dest_poin
     const std::size_t page_offset = address & Memory::YUZU_PAGEMASK;
     if (size <= Memory::YUZU_PAGESIZE - page_offset) {
         const DAddr guest_page = address & ~static_cast<DAddr>(Memory::YUZU_PAGEMASK);
-        for (size_t i = 0; i < 4; ++i) {
-            if (t_slot[i].guest_page == guest_page && t_slot[i].host_ptr != nullptr) {
-                std::memcpy(dest_pointer, t_slot[i].host_ptr + page_offset, size);
+        TranslationCache& cache = Translations();
+        for (const TranslationEntry& entry : cache.slots) {
+            if (entry.guest_page == guest_page && entry.host_ptr != nullptr) {
+                std::memcpy(dest_pointer, entry.host_ptr + page_offset, size);
                 return;
             }
         }
@@ -473,8 +475,8 @@ void DeviceMemoryManager<Traits>::ReadBlockUnsafe(DAddr address, void* dest_poin
         const auto phys_addr = tracked_entries[page_index].compressed_physical_ptr;
         if (phys_addr != 0) {
             auto* const mem_ptr = GetPointerFromRaw<u8>((PAddr(phys_addr - 1) << Memory::YUZU_PAGEBITS));
-            t_slot[cache_cursor % t_slot.size()] = TranslationEntry{.guest_page = guest_page, .host_ptr = mem_ptr};
-            cache_cursor = (cache_cursor + 1) & 3U;
+            cache.slots[cache.cursor] = TranslationEntry{.guest_page = guest_page, .host_ptr = mem_ptr};
+            cache.cursor = (cache.cursor + 1) & 3U;
             std::memcpy(dest_pointer, mem_ptr + page_offset, size);
             return;
         }
diff --git a/src/core/hle/service/nx2_gsv.h b/src/core/hle/service/nx2_gsv.h
index 5eab9d420f..3f1f25b6bf 100644
--- a/src/core/hle/service/nx2_gsv.h
+++ b/src/core/hle/service/nx2_gsv.h
@@ -210,7 +210,7 @@ private:
         const u32 syncpoint = ~cmd_buf[3];
         const u64 value = u64{cmd_buf[4]} | (u64{cmd_buf[5]} << 32);
         const u32 mode = cmd_buf[6];
-        LOG_INFO(Service, "nx2:syncpt TIPC 201 slot={:#x} syncpoint={} value={:#x} mode={}", slot,
+        LOG_DEBUG(Service, "nx2:syncpt TIPC 201 slot={:#x} syncpoint={} value={:#x} mode={}", slot,
                   syncpoint, value, mode);
         auto*& event = m_events[slot];
         if (event == nullptr) {
@@ -249,7 +249,7 @@ private:
 
     void CancelSlotEvent(HLERequestContext& ctx) {
         const u32 slot = ctx.CommandBuffer()[2];
-        LOG_INFO(Service, "nx2:syncpt TIPC 202 slot={:#x}", slot);
+        LOG_DEBUG(Service, "nx2:syncpt TIPC 202 slot={:#x}", slot);
         if (auto it = m_requests.find(slot); it != m_requests.end()) {
             Cancel(it->second);
         }
@@ -402,7 +402,7 @@ private:
         const auto* b = ctx.CommandBuffer();
         const u64 start = u64{b[2]} | (u64{b[3]} << 32);
         const u64 last = u64{b[4]} | (u64{b[5]} << 32);
-        LOG_INFO(Service, "nx2:gpu sub 101 start={:#x} last={:#x} [6]={:#x} [7]={:#x}", start,
+        LOG_DEBUG(Service, "nx2:gpu sub 101 start={:#x} last={:#x} [6]={:#x} [7]={:#x}", start,
                     last, b[6], b[7]);
         if (last > start) {
             Nvidia::Devices::Nx2GpuMemoryRegistry::Get().Register(start, last - start + 1);
@@ -410,7 +410,7 @@ private:
             // read the first (see video_core/nx2.h).
             if (last - start + 1 == 0x2000) {
                 Tegra::Nx2AddSemaphoreRegion(start);
-                LOG_INFO(Service, "nx2:gpu semaphore page {:#x} mirrored to {:#x}", start + 0x1000,
+                LOG_DEBUG(Service, "nx2:gpu semaphore page {:#x} mirrored to {:#x}", start + 0x1000,
                          start);
             }
         }
@@ -425,7 +425,7 @@ private:
     template <size_t Slot>
     void SubSlotJob(HLERequestContext& ctx) {
         const auto* w = ctx.CommandBuffer();
-        LOG_INFO(Service, "(STUBBED) nx2:gpu {} words {:08x} slot={}", ctx.GetCommand(),
+        LOG_DEBUG(Service, "(STUBBED) nx2:gpu {} words {:08x} slot={}", ctx.GetCommand(),
                  fmt::join(std::span<const u32>{w + 2, 18}, " "), w[Slot]);
         EventFor(w[Slot])->Signal(system.Kernel());
         IPC::ResponseBuilder rb{ctx, 2};
@@ -443,7 +443,7 @@ private:
         const u64 c = u64{cmd_buf[8]} | (u64{cmd_buf[9]} << 32);
         // Words 3-4: first GPU address, 5-6: last, 7: memory handle, 8-9: offset into it.
         const u32 handle = cmd_buf[7];
-        LOG_INFO(Service, "nx2:gpu 211 map [{:#x}, {:#x}] handle={} offset={:#x} slot={}", a, b,
+        LOG_DEBUG(Service, "nx2:gpu 211 map [{:#x}, {:#x}] handle={} offset={:#x} slot={}", a, b,
                  handle, c, slot);
         if (const auto it = m_memory.find(handle); it != m_memory.end() && b > a) {
             // Also at the memory's own address: an earlier 213 may have unmapped the mapping 101
@@ -528,7 +528,7 @@ private:
         const u64 start = u64{w[3]} | (u64{w[4]} << 32);
         const u64 last = u64{w[5]} | (u64{w[6]} << 32);
         const u32 slot = w[11];
-        LOG_INFO(Service, "nx2:gpu 214 words {:#x} {:#x} {:#x} {:#x} {:#x} {:#x} {:#x} {:#x} {:#x} {:#x} {:#x} {:#x}",
+        LOG_DEBUG(Service, "nx2:gpu 214 words {:#x} {:#x} {:#x} {:#x} {:#x} {:#x} {:#x} {:#x} {:#x} {:#x} {:#x} {:#x}",
                  w[1], w[2], w[3], w[4], w[5], w[6], w[7], w[8], w[9], w[10], w[11], w[12]);
         if (last > start && ((last + 1 - start) % Core::Memory::YUZU_PAGESIZE) == 0) {
             const u64 size = last + 1 - start;
@@ -554,7 +554,7 @@ private:
         const auto* w = ctx.CommandBuffer();
         const u64 first = u64{w[3]} | (u64{w[4]} << 32);
         const u64 last = u64{w[5]} | (u64{w[6]} << 32);
-        LOG_INFO(Service, "nx2:gpu 215 destroy sparse region [{:#x}, {:#x}] words {:#x} {:#x} {:#x} {:#x}",
+        LOG_DEBUG(Service, "nx2:gpu 215 destroy sparse region [{:#x}, {:#x}] words {:#x} {:#x} {:#x} {:#x}",
                  first, last, w[1], w[2], w[7], w[8]);
         FreeReleasedRegions();
         if (last > first) {
@@ -879,7 +879,10 @@ private:
         // game is blocked here while Eden's GPU thread catches up (fence) and composites.
         {
             using us = std::chrono::microseconds;
-            const auto to_us = [](auto d) { return std::chrono::duration_cast<us>(d).count(); };
+            // long long like the stats: microseconds::rep is long on Linux (std::max needs one type).
+            const auto to_us = [](auto d) -> long long {
+                return std::chrono::duration_cast<us>(d).count();
+            };
             m_stats.fence_us += to_us(waited - entered);
             m_stats.composite_us += to_us(composited - waited);
             m_stats.pace_us += to_us(paced - composited);
@@ -1014,10 +1017,18 @@ private:
         }
         u32 crop_width = layout.width;
         u32 crop_height = layout.height;
-        if (system.GetApplicationProcessProgramID() == 0x0400C3F00006E000ULL &&
-            layout.width == 3840 && layout.height == 2160) {
-            crop_width = 2560;
-            crop_height = 1440;
+        // Mario Kart World draws into the top-left corner of its present images: 2560x1440 docked
+        // (in 3840x2160 images), 1920x1080 handheld. It keeps its docked images after a switch to
+        // handheld, and showing all of them put the frame in a corner with stale docked frames
+        // around it.
+        if (system.GetApplicationProcessProgramID() == 0x0400C3F00006E000ULL) {
+            if (!Settings::IsDockedMode()) {
+                crop_width = (std::min)(crop_width, 1920U);
+                crop_height = (std::min)(crop_height, 1080U);
+            } else if (layout.width == 3840 && layout.height == 2160) {
+                crop_width = 2560;
+                crop_height = 1440;
+            }
         }
         static const u32 env_crop_w = [] {
             const char* v = std::getenv("EDEN_NX2_CROP_WIDTH");
@@ -1031,7 +1042,10 @@ private:
             crop_width = env_crop_w;
             crop_height = env_crop_h;
         }
-        if (m_presents <= 4) {
+        const u64 shown = (u64{layout.width} << 48) | (u64{layout.height} << 32) |
+                          (u64{crop_width} << 16) | crop_height;
+        if (shown != m_shown_layout) {
+            m_shown_layout = shown;
             LOG_INFO(Service, "nx2:vi2 composite image {} layout {}x{}, crop {}x{}",
                      image_id, layout.width, layout.height, crop_width, crop_height);
         }
@@ -1214,6 +1228,7 @@ private:
     std::shared_ptr<VsyncState> m_vsync = std::make_shared<VsyncState>();
     std::shared_ptr<Core::Timing::EventType> m_vsync_timer;
     u64 m_presents{};
+    u64 m_shown_layout{}; // image and crop size last logged
     std::shared_ptr<std::atomic<bool>> m_capture_pending = std::make_shared<std::atomic<bool>>(false);
 };
 
diff --git a/src/core/hle/service/psc/time/common.h b/src/core/hle/service/psc/time/common.h
index 2183fa72b8..6f7d2cb88b 100644
--- a/src/core/hle/service/psc/time/common.h
+++ b/src/core/hle/service/psc/time/common.h
@@ -134,12 +134,13 @@ struct OperationEvent : public Common::IntrusiveListBaseNode<OperationEvent> {
     Kernel::KEvent* m_event{};
 };
 
-constexpr inline std::chrono::nanoseconds ConvertToTimeSpan(s64 ticks) {
+inline std::chrono::nanoseconds ConvertToTimeSpan(s64 ticks) {
     constexpr auto one_second_ns{
         std::chrono::duration_cast<std::chrono::nanoseconds>(std::chrono::seconds(1)).count()};
 
-    constexpr s64 max{Common::WallClock::CNTFRQ *
-                      ((std::numeric_limits<s64>::max)() / one_second_ns)};
+    // The guest counter's frequency: 31.25 MHz for Switch 2 titles.
+    const s64 frequency{static_cast<s64>(Common::g_wall_clock.GetGuestCNTFRQ())};
+    const s64 max{frequency * ((std::numeric_limits<s64>::max)() / one_second_ns)};
 
     if (ticks > max) {
         return std::chrono::nanoseconds((std::numeric_limits<s64>::max)());
@@ -147,8 +148,8 @@ constexpr inline std::chrono::nanoseconds ConvertToTimeSpan(s64 ticks) {
         return std::chrono::nanoseconds((std::numeric_limits<s64>::min)());
     }
 
-    auto a{ticks / Common::WallClock::CNTFRQ * one_second_ns};
-    auto b{((ticks % Common::WallClock::CNTFRQ) * one_second_ns) / Common::WallClock::CNTFRQ};
+    auto a{ticks / frequency * one_second_ns};
+    auto b{((ticks % frequency) * one_second_ns) / frequency};
 
     return std::chrono::nanoseconds(a + b);
 }
diff --git a/src/core/hle/service/sm/sm.cpp b/src/core/hle/service/sm/sm.cpp
index c75c5437df..60f0dd89cb 100644
--- a/src/core/hle/service/sm/sm.cpp
+++ b/src/core/hle/service/sm/sm.cpp
@@ -11,11 +11,19 @@
 #include "core/hle/kernel/k_client_port.h"
 #include "core/hle/kernel/k_client_session.h"
 #include "core/hle/kernel/k_port.h"
+#include "core/hle/kernel/k_process.h"
+#include "core/hle/kernel/k_resource_limit.h"
 #include "core/hle/kernel/k_scoped_resource_reservation.h"
 #include "core/hle/kernel/k_server_port.h"
+#include "core/hle/kernel/k_server_session.h"
+#include "core/hle/kernel/k_session.h"
+#include "core/hle/kernel/k_thread.h"
+#include "core/hle/kernel/kernel.h"
 #include "core/hle/result.h"
 #include "core/hle/service/cmif_types.h"
+#include "core/hle/service/hle_ipc.h"
 #include "core/hle/service/ipc_helpers.h"
+#include "core/hle/service/nx2_stub_interface.h"
 #include "core/hle/service/server_manager.h"
 #include "core/hle/service/sm/sm.h"
 #include "core/hle/service/sm/sm_controller.h"
@@ -170,6 +178,30 @@ static std::string PopServiceName(IPC::RequestParser& rp) {
     return result;
 }
 
+namespace {
+
+// A session served by a Switch 2 stub interface named after the missing service (see
+// nx2_stub_interface.h), set up like a session handed out by a service.
+Kernel::KClientSession* CreateStubSession(Core::System& system, HLERequestContext& ctx,
+                                          const std::string& name) {
+    auto& kernel = system.Kernel();
+    ASSERT(Kernel::GetCurrentProcess(kernel).GetResourceLimit()->Reserve(
+        kernel, Kernel::LimitableResource::SessionCountMax, 1));
+
+    auto* session = Kernel::KSession::Create(kernel);
+    session->Initialize(kernel, nullptr, 0);
+    Kernel::KSession::Register(kernel, session);
+
+    auto manager =
+        std::make_shared<SessionRequestManager>(kernel, ctx.GetManager()->GetServerManager());
+    manager->SetSessionHandler(std::make_shared<INx2StubInterface>(
+        system, INx2StubInterface::InternName(fmt::format("{}{}", Nx2StubPrefix, name))));
+    ctx.GetManager()->GetServerManager().RegisterSession(&session->GetServerSession(), manager);
+    return &session->GetClientSession();
+}
+
+} // namespace
+
 Result SM::GetServiceImpl(Kernel::KClientSession** out_client_session, HLERequestContext& ctx) {
     if (!ctx.GetManager()->GetIsInitializedForSm()) {
         return Service::SM::ResultInvalidClient;
@@ -187,6 +219,17 @@ Result SM::GetServiceImpl(Kernel::KClientSession** out_client_session, HLEReques
     }
 
     if (port_result != ResultSuccess) {
+        // Switch 2 programs also ask for services Eden does not have at all (newer Mario Kart
+        // World updates open chat:u, GameChat, while booting). Nothing ever registers those, so
+        // waiting would hang them for good: hand out a placeholder whose commands are all
+        // auto-stubbed, like their unknown commands on existing services.
+        const auto* const caller = ctx.GetThread().GetOwnerProcess();
+        if (caller != nullptr && caller->IsSwitch2()) {
+            LOG_WARNING(Service_SM, "(STUBBED) Switch 2 service {} does not exist, returning a stub",
+                        name);
+            *out_client_session = CreateStubSession(system, ctx, name);
+            return ResultSuccess;
+        }
         LOG_INFO(Service_SM, "Waiting for service {} to become available", name);
         ctx.SetIsDeferred();
         return Service::SM::ResultNotRegistered;
diff --git a/src/qt_common/game_list/worker.cpp b/src/qt_common/game_list/worker.cpp
index 6b6f93447d..70577cb270 100644
--- a/src/qt_common/game_list/worker.cpp
+++ b/src/qt_common/game_list/worker.cpp
@@ -167,6 +167,19 @@ QString FormatGameName(const std::string& physical_name) {
     return physical_name_as_qstring;
 }
 
+// Loose ExeFS dumps (<game>/exefs/main, updates in <game>/updates/exefs/main) have no control
+// data, so they are named after their game folder.
+QString LooseExeFSName(const std::string& physical_name) {
+    QDir dir = QFileInfo(QString::fromStdString(physical_name)).dir();
+    if (dir.dirName().compare(QStringLiteral("exefs"), Qt::CaseInsensitive) != 0 || !dir.cdUp()) {
+        return {};
+    }
+    if (dir.dirName().compare(QStringLiteral("updates"), Qt::CaseInsensitive) == 0 && dir.cdUp()) {
+        return QObject::tr("%1 (update)").arg(dir.dirName());
+    }
+    return dir.dirName();
+}
+
 QString FormatPatchNameVersions(const FileSys::PatchManager& patch_manager,
                                 Loader::AppLoader& loader, bool updatable = true) {
     QString out;
@@ -392,7 +405,13 @@ void GameListWorker::ScanFileSystem(ScanTarget target, const std::string& dir_pa
                     [[maybe_unused]] const auto res1 = app_loader->ReadIcon(icon);
 
                     std::string name = " ";
-                    [[maybe_unused]] const auto res3 = app_loader->ReadTitle(name);
+                    const auto res3 = app_loader->ReadTitle(name);
+                    if ((res3 != Loader::ResultStatus::Success || name.empty() || name == " ") &&
+                        IsExtractedNCAMain(physical_name)) {
+                        if (const QString loose = LooseExeFSName(physical_name); !loose.isEmpty()) {
+                            name = loose.toStdString();
+                        }
+                    }
 
                     const FileSys::PatchManager patch{id, system.GetFileSystemController(),
                                                       system.GetContentProvider()};
diff --git a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h
index 9e5dd0313f..0bed72eaec 100644
--- a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h
+++ b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h
@@ -758,20 +758,35 @@ inline LoweredProgram LowerToMaxwell(std::span<const Instruction> code) {
                  (u64{insn.Bits(64, 8)} << 20) | (u64{maxwell} << 28));
             break;
         }
-        case 0xb6c: { // TXD ARRAY_2D, four channels and consecutive split destinations.
+        case 0xb6c: { // TXD 2D or ARRAY_2D, four channels; the two destination pairs anywhere.
             const u32 binding = insn.Bits(54, 5), offset = insn.Bits(40, 14);
             const u32 coordinates = insn.Bits(24, 8), derivatives = insn.Bits(32, 8);
+            // Ampere dimensions: 1 = 2D, 5 = ARRAY_2D; Maxwell TXD's are 2 and 3. The coordinates
+            // (x, y[, layer]) and the four gradients sit in the same registers for both.
+            const u32 dimension = insn.Bits(61, 3);
+            const u32 high = insn.Bits(64, 8); // destination of the third and fourth channels
             if ((insn.words[2] & ~0xffU) != 0x001e0f00 || insn.Bits(59, 2) ||
-                insn.Bits(61, 3) != 5 || binding >= 18 || offset > 0x1fff ||
-                (dest & 1) || dest >= 248 || insn.Bits(64, 8) != dest + 2 ||
+                (dimension != 1 && dimension != 5) || binding >= 18 || offset > 0x1fff ||
+                (dest & 1) || dest >= 248 || (high & 1) || (high >= 248 && high != 255) ||
                 coordinates >= 248 || derivatives >= 247)
-                return fail(index, "native array gradient texture mode");
+                return fail(index, "native gradient texture mode");
             if (result.required_texture_buffer != ~0U && result.required_texture_buffer != binding)
                 return fail(index, "native texture descriptor buffer conflict");
             result.required_texture_buffer = binding;
-            emit(base(0xde00000000000000ULL) | (u64{coordinates} << 8) |
-                 (u64{derivatives} << 20) | (u64{3} << 28) | (u64{15} << 31) |
+            // Maxwell writes four consecutive registers: when the second pair is elsewhere (Splatoon
+            // Raiders: R12/R18, R18/R44), write them to scratch and move each pair into place.
+            const bool consecutive = high == dest + 2;
+            const u64 target = consecutive ? dest : 248;
+            emit(0xde00000000000000ULL | predicate | target | (u64{coordinates} << 8) |
+                 (u64{derivatives} << 20) | (u64{dimension == 5 ? 3U : 2U} << 28) | (u64{15} << 31) |
                  (u64{offset} << 36) | (u64{insn.Bits(90, 1)} << 49) | (u64{7} << 51));
+            if (!consecutive) {
+                for (u32 i = 0; i < 2; ++i) {
+                    emit(0x5c98000000000000ULL | predicate | (dest + i) | (u64{15} << 39) | (u64{248 + i} << 20));
+                    if (high != 255)
+                        emit(0x5c98000000000000ULL | predicate | (high + i) | (u64{15} << 39) | (u64{250 + i} << 20));
+                }
+            }
             break;
         }
         case 0x989: { // SHFL, immediate lane and register clamp/segmentation mask, any mode.
@@ -1706,7 +1721,7 @@ inline LoweredProgram LowerToMaxwell(std::span<const Instruction> code) {
         case 0x220: case 0x820: case 0xa20: case 0xc20: { // FMUL
             const u32 scale = insn.Bits(84, 3);
             if (insn.Bits(76, 1) || scale == 0 || insn.Bits(81, 3) ||
-                insn.Bits(87, 4) || insn.Bits(91, 1) != (opcode == 0xc20) ||
+                insn.Bits(87, 4) || insn.Bits(91, 1) != static_cast<u32>(opcode == 0xc20) ||
                 (scale != 4 && (!insn.Bits(80, 1) || insn.Bits(78, 2))))
                 return fail(index, "native FMUL scale/mode");
             const u32 a = operand(0), b = operand(1);
@@ -1715,7 +1730,8 @@ inline LoweredProgram LowerToMaxwell(std::span<const Instruction> code) {
             break;
         }
         case 0x221: case 0x421: case 0x621: case 0x821: case 0xa21: case 0xe21: { // FADD
-            if (insn.Bits(81, 10) || insn.Bits(91, 1) != (opcode == 0xe21)) return fail(index, "native FADD mode");
+            if (insn.Bits(81, 10) || insn.Bits(91, 1) != static_cast<u32>(opcode == 0xe21))
+                return fail(index, "native FADD mode");
             const u32 a = operand(0), b = operand(insn.Bits(9, 3) == 2 || insn.Bits(9, 3) == 3 || insn.Bits(9, 3) == 7 ? 2 : 1);
             emit(base(0x5c58000000000000ULL) | (u64{a} << 8) | (u64{b} << 20) |
                 (u64{insn.Bits(78, 2)} << 39) | (u64{insn.Bits(80, 1)} << 44) | (u64{insn.Bits(77, 1)} << 50));
diff --git a/src/shader_recompiler/ir_opt/constant_propagation_pass.cpp b/src/shader_recompiler/ir_opt/constant_propagation_pass.cpp
index 6285fef16f..0c0d4fb1a2 100644
--- a/src/shader_recompiler/ir_opt/constant_propagation_pass.cpp
+++ b/src/shader_recompiler/ir_opt/constant_propagation_pass.cpp
@@ -953,6 +953,36 @@ void FoldDriverConstBuffer(Environment& env, IR::Block& block, IR::Inst& inst, u
     }
 }
 
+void FoldSMax32(IR::Block& block, IR::Inst& inst) {
+    for (size_t index = 0; index < 2; ++index) {
+        const IR::Value value{inst.Arg(index).Resolve()};
+        const IR::Inst* const negation{inst.Arg(1 - index).TryInstRecursive()};
+        if (!negation || negation->GetOpcode() != IR::Opcode::INeg32 ||
+            negation->Arg(0).Resolve() != value) {
+            continue;
+        }
+        const IR::Inst* const select{value.TryInstRecursive()};
+        if (!select || select->GetOpcode() != IR::Opcode::SelectU32) {
+            continue;
+        }
+        const IR::Value true_value{select->Arg(1).Resolve()};
+        const IR::Value false_value{select->Arg(2).Resolve()};
+        const auto is_mask = [](const IR::Value& candidate) {
+            return candidate.IsImmediate() &&
+                   (candidate.U32() == 0 || candidate.U32() == 1 ||
+                    candidate.U32() == 0xffffffffU);
+        };
+        if (!is_mask(true_value) || !is_mask(false_value)) {
+            continue;
+        }
+        // max(x, -x) is 0 or 1 for boolean and predicate-mask constants.
+        IR::IREmitter ir{block, IR::Block::InstructionList::s_iterator_to(inst)};
+        inst.ReplaceUsesWith(ir.Select(IR::U1{select->Arg(0)},
+                                      ir.Imm32(true_value.U32() != 0 ? 1U : 0U),
+                                      ir.Imm32(false_value.U32() != 0 ? 1U : 0U)));
+        return;
+    }
+}
 void ConstantPropagation(Environment& env, IR::Block& block, IR::Inst& inst) {
     switch (inst.GetOpcode()) {
     case IR::Opcode::GetRegister:
@@ -963,6 +993,8 @@ void ConstantPropagation(Environment& env, IR::Block& block, IR::Inst& inst) {
         return FoldAdd<u32>(block, inst);
     case IR::Opcode::ISub32:
         return FoldISub32(inst);
+    case IR::Opcode::SMax32:
+        return FoldSMax32(block, inst);
     case IR::Opcode::IMul32:
         FoldWhenAllImmediates(inst, [](u32 a, u32 b) { return a * b; });
         return;
diff --git a/src/shader_recompiler/ir_opt/global_memory_to_storage_buffer_pass.cpp b/src/shader_recompiler/ir_opt/global_memory_to_storage_buffer_pass.cpp
index bdb4494a15..6aee241dc2 100644
--- a/src/shader_recompiler/ir_opt/global_memory_to_storage_buffer_pass.cpp
+++ b/src/shader_recompiler/ir_opt/global_memory_to_storage_buffer_pass.cpp
@@ -414,7 +414,7 @@ void CollectStorageBuffers(IR::Block& block, IR::Inst& inst, StorageInfo& info)
             LOG_WARNING(Shader, "Storage buffer failed to track, using global memory fallbacks");
             return;
         }
-        LOG_WARNING(Shader, "Storage buffer tracked without bias, index {} offset {}",
+        LOG_DEBUG(Shader, "Storage buffer tracked without bias, index {} offset {}",
                     storage_buffer->index, storage_buffer->offset);
     }
     // Collect storage buffer and the instruction
diff --git a/src/video_core/CMakeLists.txt b/src/video_core/CMakeLists.txt
index 163caa32d0..7f83e0d3c2 100644
--- a/src/video_core/CMakeLists.txt
+++ b/src/video_core/CMakeLists.txt
@@ -97,6 +97,8 @@ add_library(video_core STATIC
     fence_manager.h
     gpu.cpp
     gpu.h
+    gpu_split.cpp
+    gpu_split.h
     gpu_thread.cpp
     gpu_thread.h
     guest_memory.h
diff --git a/src/video_core/buffer_cache/buffer_cache.h b/src/video_core/buffer_cache/buffer_cache.h
index ad2e200438..42d7ba066d 100644
--- a/src/video_core/buffer_cache/buffer_cache.h
+++ b/src/video_core/buffer_cache/buffer_cache.h
@@ -17,6 +17,7 @@
 #include "video_core/buffer_cache/buffer_cache_base.h"
 #include "video_core/guest_memory.h"
 #include "video_core/host1x/gpu_device_memory_manager.h"
+#include "video_core/nx2_profile.h"
 #include "video_core/texture_cache/util.h"
 
 namespace VideoCommon {
@@ -57,21 +58,39 @@ template <class P>
 BufferCache<P>::~BufferCache() = default;
 
 template <class P>
-void BufferCache<P>::RunGarbageCollector() {
-    const bool aggressive_gc = total_used_memory >= critical_memory;
+void BufferCache<P>::RunGarbageCollector(bool aggressive_gc) {
     const u64 ticks_to_destroy = aggressive_gc ? 60 : 120;
     int num_iterations = aggressive_gc ? 64 : 32;
-    const auto clean_up = [this, &num_iterations](BufferId buffer_id) {
+    // Their GPU-written data is downloaded first, all with one wait for the GPU: one wait per
+    // buffer stalled the GPU thread every frame on title 1 (an 8 GB GPU near its budget, no
+    // per-fence downloads keeping the buffers clean), ~2.5% of its time.
+    boost::container::small_vector<BufferId, 64> evicted;
+    lru_cache.ForEachItemBelow(frame_tick - ticks_to_destroy, [&](BufferId buffer_id) {
         if (num_iterations == 0) {
             return true;
         }
         --num_iterations;
-        auto& buffer = slot_buffers[buffer_id];
-        DownloadBufferMemory(buffer);
-        DeleteBuffer(buffer_id);
+        // Below the critical threshold, buffers holding GPU-written data stay: evicting one
+        // means downloading it first. The thresholds compare the whole device's use (mostly
+        // textures) with fixed headroom, so on an 8 GB GPU the collector ran every frame once a
+        // title 1 course was loaded, and the downloads were ~2.7% of the GPU thread.
+        const Buffer& buffer = slot_buffers[buffer_id];
+        if (!aggressive_gc && gpu_modified_ranges.Intersects(buffer.CpuAddr(), buffer.SizeBytes())) {
+            return false;
+        }
+        evicted.push_back(buffer_id);
         return false;
-    };
-    lru_cache.ForEachItemBelow(frame_tick - ticks_to_destroy, clean_up);
+    });
+    if constexpr (USE_MEMORY_MAPS) {
+        DownloadBuffersMemory(evicted);
+    } else {
+        for (const BufferId buffer_id : evicted) {
+            DownloadBufferMemory(slot_buffers[buffer_id]);
+        }
+    }
+    for (const BufferId buffer_id : evicted) {
+        DeleteBuffer(buffer_id);
+    }
 }
 
 template <class P>
@@ -117,8 +136,17 @@ void BufferCache<P>::TickFrame() {
     if (runtime.CanReportMemoryUsage()) {
         total_used_memory = runtime.GetDeviceMemoryUsage();
     }
-    if (total_used_memory >= minimum_memory) {
-        RunGarbageCollector();
+    if (gpu_memory != nullptr && gpu_memory->UsesNativeTiling()) {
+        // Switch 2 titles fill the device with textures (their collector watches the budget);
+        // title 1's buffers were ~200 MiB of an 8 GB GPU's ~6.8 GB in use, yet this collector
+        // ran aggressively every frame on the device-wide figure, evicting up to 64 buffers (and
+        // downloading the GPU-written ones, ~5% of the GPU thread) that were recreated later.
+        // Collect when the buffers themselves are large.
+        if (buffer_bytes >= 1_GiB) {
+            RunGarbageCollector(buffer_bytes >= 2_GiB);
+        }
+    } else if (total_used_memory >= minimum_memory) {
+        RunGarbageCollector(total_used_memory >= critical_memory);
     }
     ++frame_tick;
     delayed_destruction_ring.Tick();
@@ -251,7 +279,9 @@ bool BufferCache<P>::DMACopy(GPUVAddr src_address, GPUVAddr dest_address, u64 am
         const DAddr diff = base_address - *cpu_src_address;
         const DAddr new_base_address = *cpu_dest_address + diff;
         tmp_intervals.push_back({new_base_address, size});
-        uncommitted_gpu_modified_ranges.Add(new_base_address, size);
+        if (AsyncBufferDownloads()) {
+            uncommitted_gpu_modified_ranges.Add(new_base_address, size);
+        }
     };
     gpu_modified_ranges.ForEachInRange(*cpu_src_address, amount, mirror);
     // This subtraction in this order is important for overlapping copies.
@@ -666,6 +696,8 @@ void BufferCache<P>::CommitAsyncFlushesHigh() {
         async_buffers.emplace_back(std::optional<Async_Buffer>{});
         return;
     }
+    Nx2Profile::Add(Nx2Profile::Counter::DownloadBytes, total_size_bytes);
+    Nx2Profile::Add(Nx2Profile::Counter::DownloadCopies, downloads.size());
     auto download_staging = runtime.DownloadStagingBuffer(total_size_bytes, true);
     boost::container::small_vector<BufferCopy, 4> normalized_copies;
     runtime.PreCopyBarrier();
@@ -676,6 +708,19 @@ void BufferCache<P>::CommitAsyncFlushesHigh() {
         Buffer& buffer = slot_buffers[buffer_id];
         second_copy.src_offset = static_cast<size_t>(buffer.CpuAddr()) + copy.src_offset;
         const DAddr orig_device_addr = static_cast<DAddr>(second_copy.src_offset);
+        if (Nx2Profile::Enabled()) {
+            u64 on_dma_pages = 0;
+            const auto& pages = Nx2Profile::DmaReadPages();
+            for (DAddr page = orig_device_addr >> 12; page <= (orig_device_addr + copy.size - 1) >> 12;
+                 ++page) {
+                if (pages.contains(page)) {
+                    const DAddr begin = (std::max)(orig_device_addr, page << 12);
+                    const DAddr end = (std::min)(orig_device_addr + copy.size, (page + 1) << 12);
+                    on_dma_pages += end - begin;
+                }
+            }
+            Nx2Profile::Add(Nx2Profile::Counter::DownloadBytesOnDmaPages, on_dma_pages);
+        }
         async_downloads.Add(orig_device_addr, copy.size);
         buffer.MarkUsage(copy.src_offset, copy.size);
         runtime.CopyBuffer(download_staging.buffer, buffer, copies, false);
@@ -728,9 +773,9 @@ void BufferCache<P>::PopAsyncBuffers() {
 
 template <class P>
 bool BufferCache<P>::IsRegionGpuModified(DAddr addr, size_t size) {
-    bool is_dirty = false;
-    gpu_modified_ranges.ForEachInRange(addr, size, [&](DAddr, DAddr) { is_dirty = true; });
-    return is_dirty;
+    // One lookup instead of visiting every range inside (~4% of the GPU thread for title 1's
+    // inline uploads once GPU-written ranges are no longer cleared by downloads).
+    return gpu_modified_ranges.Intersects(addr, size);
 }
 
 template <class P>
@@ -1558,7 +1603,9 @@ void BufferCache<P>::MarkWrittenBuffer(BufferId buffer_id, DAddr device_addr, u3
     }
     memory_tracker.MarkRegionAsGpuModified(device_addr, size);
     gpu_modified_ranges.Add(device_addr, size);
-    uncommitted_gpu_modified_ranges.Add(device_addr, size);
+    if (AsyncBufferDownloads()) {
+        uncommitted_gpu_modified_ranges.Add(device_addr, size);
+    }
 }
 
 template <class P>
@@ -1739,9 +1786,11 @@ void BufferCache<P>::ChangeRegister(BufferId buffer_id) {
     Buffer& buffer = slot_buffers[buffer_id];
     const auto size = buffer.SizeBytes();
     if (insert) {
+        buffer_bytes += Common::AlignUp(size, 1024);
         total_used_memory += Common::AlignUp(size, 1024);
         buffer.setLRUID(lru_cache.Insert(buffer_id, frame_tick));
     } else {
+        buffer_bytes -= Common::AlignUp(size, 1024);
         total_used_memory -= Common::AlignUp(size, 1024);
         lru_cache.Free(buffer.getLRUID());
     }
@@ -1932,6 +1981,52 @@ void BufferCache<P>::DownloadBufferMemory(Buffer& buffer) {
 }
 
 template <class P>
+void BufferCache<P>::DownloadBuffersMemory(std::span<const BufferId> buffer_ids) {
+    struct Download {
+        BufferId buffer_id;
+        BufferCopy copy;
+    };
+    std::vector<Download> downloads;
+    u64 total_size_bytes = 0;
+    for (const BufferId buffer_id : buffer_ids) {
+        Buffer& buffer = slot_buffers[buffer_id];
+        const DAddr buffer_addr = buffer.CpuAddr();
+        memory_tracker.ForEachDownloadRangeAndClear(
+            buffer_addr, buffer.SizeBytes(), [&](u64 device_addr_out, u64 range_size) {
+                gpu_modified_ranges.ForEachInRange(device_addr_out, range_size,
+                                                   [&](DAddr start, DAddr end) {
+                    downloads.push_back({buffer_id, BufferCopy{
+                                                        .src_offset = start - buffer_addr,
+                                                        .dst_offset = total_size_bytes,
+                                                        .size = end - start,
+                                                    }});
+                    total_size_bytes += Common::AlignUp(end - start, 64);
+                });
+                ClearDownload(device_addr_out, range_size);
+                gpu_modified_ranges.Subtract(device_addr_out, range_size);
+            });
+    }
+    if (total_size_bytes == 0) {
+        return;
+    }
+    auto download_staging = runtime.DownloadStagingBuffer(total_size_bytes);
+    for (Download& download : downloads) {
+        download.copy.dst_offset += download_staging.offset;
+        Buffer& buffer = slot_buffers[download.buffer_id];
+        buffer.MarkUsage(download.copy.src_offset, download.copy.size);
+        runtime.CopyBuffer(download_staging.buffer, buffer, std::span{&download.copy, 1}, true);
+    }
+    runtime.Finish();
+    const u8* const mapped_memory = download_staging.mapped_span.data();
+    for (const Download& download : downloads) {
+        const DAddr device_addr = slot_buffers[download.buffer_id].CpuAddr() + download.copy.src_offset;
+        device_memory.WriteBlockUnsafe(
+            device_addr, mapped_memory + (download.copy.dst_offset - download_staging.offset),
+            download.copy.size);
+    }
+}
+
+template <class P>
 void BufferCache<P>::DownloadBufferMemory(Buffer& buffer, DAddr device_addr, u64 size) {
     boost::container::small_vector<BufferCopy, 1> copies;
     u64 total_size_bytes = 0;
diff --git a/src/video_core/buffer_cache/buffer_cache_base.h b/src/video_core/buffer_cache/buffer_cache_base.h
index 57011752cc..45a1836341 100644
--- a/src/video_core/buffer_cache/buffer_cache_base.h
+++ b/src/video_core/buffer_cache/buffer_cache_base.h
@@ -7,6 +7,7 @@
 #pragma once
 
 #include <algorithm>
+#include <cstdlib>
 #include <array>
 #include <bit>
 #include <functional>
@@ -377,7 +378,7 @@ private:
                ((device_addr + size) & ~Core::DEVICE_PAGEMASK);
     }
 
-    void RunGarbageCollector();
+    void RunGarbageCollector(bool aggressive_gc);
 
     void BindHostIndexBuffer();
 
@@ -467,6 +468,9 @@ private:
 
     void DownloadBufferMemory(Buffer& buffer_id, DAddr device_addr, u64 size);
 
+    /// Downloads the GPU-written data of several buffers with one wait for the GPU.
+    void DownloadBuffersMemory(std::span<const BufferId> buffer_ids);
+
     void DeleteBuffer(BufferId buffer_id, bool do_not_mark = false);
 
     [[nodiscard]] Binding StorageBufferBinding(GPUVAddr ssbo_addr, u32 cbuf_index,
@@ -514,6 +518,24 @@ private:
 
     MemoryTracker memory_tracker;
     Common::RangeSet<DAddr> uncommitted_gpu_modified_ranges;
+
+    /// Whether GPU-written buffer ranges are copied back to guest memory at the next fence;
+    /// otherwise only when something flushes them (as upstream yuzu below High GPU accuracy).
+    /// Not for Switch 2 titles: title 1 had 250-330 MB/s copied back (a fence per compute
+    /// dispatch, the fencing thread holding the buffer cache while it wrote them), while its CPU
+    /// read GPU-written memory at most 64 times in 200 s (counted with reactive flushing) and
+    /// the native MME's DMA reads never touched those pages. Without them a fence costs ~1 us
+    /// instead of ~5 and draw setup ~7% less. EDEN_NX2_BUFFER_DOWNLOADS=0/1 forces it.
+    bool AsyncBufferDownloads() const {
+        static const int forced = [] {
+            const char* value = std::getenv("EDEN_NX2_BUFFER_DOWNLOADS");
+            return value == nullptr ? -1 : (value[0] == '0' ? 0 : 1);
+        }();
+        if (forced >= 0) {
+            return forced == 1;
+        }
+        return gpu_memory == nullptr || !gpu_memory->UsesNativeTiling();
+    }
     Common::RangeSet<DAddr> gpu_modified_ranges;
     std::deque<Common::RangeSet<DAddr>> committed_gpu_modified_ranges;
 
@@ -538,6 +560,7 @@ private:
     std::vector<MultiRangeSegment> compute_segments;
     u64 frame_tick = 0;
     u64 total_used_memory = 0;
+    u64 buffer_bytes = 0; ///< The buffers' own size (total_used_memory may be the device's)
     u64 minimum_memory = 0;
     u64 critical_memory = 0;
     BufferId inline_buffer_id;
diff --git a/src/video_core/control/scheduler.cpp b/src/video_core/control/scheduler.cpp
index 24f1db1f7c..dc42906752 100644
--- a/src/video_core/control/scheduler.cpp
+++ b/src/video_core/control/scheduler.cpp
@@ -9,12 +9,28 @@
 #include <vector>
 
 #include "common/assert.h"
+#include "core/core.h"
 #include "video_core/control/channel_state.h"
 #include "video_core/control/scheduler.h"
 #include "video_core/gpu.h"
+#include "video_core/gpu_split.h"
 
 namespace Tegra::Control {
 
+void Scheduler::BindChannel(GPU& gpu, s32 bind_id) {
+    // With the GPU thread split the rasterizer's channel switches with the methods around it, on
+    // the back thread (deferred outside the lock: the back takes it to bind).
+    if (auto& split = VideoCommon::GpuSplit::Get(); split.Active()) {
+        split.Defer([this, gpu_ptr = &gpu, bind_id] {
+            std::scoped_lock lk{scheduling_guard};
+            gpu_ptr->BindChannel(bind_id);
+        });
+        return;
+    }
+    std::scoped_lock lk{scheduling_guard};
+    gpu.BindChannel(bind_id);
+}
+
 void Scheduler::Push(GPU& gpu, s32 channel, CommandList&& entries) {
     std::shared_ptr<ChannelState> channel_state;
     {
@@ -22,8 +38,8 @@ void Scheduler::Push(GPU& gpu, s32 channel, CommandList&& entries) {
         auto it = channels.find(channel);
         ASSERT(it != channels.end());
         channel_state = it->second;
-        gpu.BindChannel(channel_state->bind_id);
     }
+    BindChannel(gpu, channel_state->bind_id);
     // Process commands outside the lock to reduce contention.
     // Multiple channels can prepare their commands in parallel.
     channel_state->payload->dma_pusher.Push(std::move(entries));
@@ -44,13 +60,16 @@ void Scheduler::ResumePending(GPU& gpu) {
         progress = false;
         for (const auto& state : pending_channels) {
             auto& dma_pusher = state->payload->dma_pusher;
+            // Powered off (shutting down), DispatchCalls returns at once and the acquire stays
+            // pending: resuming it again counted as progress forever and the GPU thread never
+            // reached its stop request (closing the game hung).
+            if (!dma_pusher.system.IsPoweredOn()) {
+                return;
+            }
             if (!dma_pusher.HasPendingAcquire() || !dma_pusher.CanResumeAcquire()) {
                 continue;
             }
-            {
-                std::scoped_lock lk{scheduling_guard};
-                gpu.BindChannel(state->bind_id);
-            }
+            BindChannel(gpu, state->bind_id);
             dma_pusher.DispatchCalls();
             progress = true;
         }
diff --git a/src/video_core/control/scheduler.h b/src/video_core/control/scheduler.h
index 3a3dcc58cf..9a0f1b919f 100644
--- a/src/video_core/control/scheduler.h
+++ b/src/video_core/control/scheduler.h
@@ -34,6 +34,8 @@ public:
     void DeclareChannel(std::shared_ptr<ChannelState> new_channel);
 
 private:
+    void BindChannel(GPU& gpu, s32 bind_id);
+
     ::Common::unordered_map<s32, std::shared_ptr<ChannelState>> channels;
     std::mutex scheduling_guard;
     std::vector<std::shared_ptr<ChannelState>> pending_channels; ///< Suspended on an acquire
diff --git a/src/video_core/dma_pusher.cpp b/src/video_core/dma_pusher.cpp
index 3cc5afe967..62f7e6bfdb 100644
--- a/src/video_core/dma_pusher.cpp
+++ b/src/video_core/dma_pusher.cpp
@@ -12,10 +12,13 @@
 #include "common/settings.h"
 #include "core/core.h"
 #include "core/memory.h"
+#include "video_core/nx2_profile.h"
 #include "video_core/dma_pusher.h"
 #include "video_core/control/channel_state.h"
+#include "video_core/engines/kepler_compute.h"
 #include "video_core/engines/maxwell_3d.h"
 #include "video_core/gpu.h"
+#include "video_core/gpu_split.h"
 #include "video_core/guest_memory.h"
 #include "video_core/memory_manager.h"
 #include "video_core/nx2.h"
@@ -44,6 +47,14 @@ void DmaPusher::DispatchCalls() {
             break;
         }
     }
+    if (auto& split = VideoCommon::GpuSplit::Get(); split.Active()) {
+        split.Defer([gpu = &system.GPU()] {
+            gpu->FlushCommands();
+            gpu->OnCommandListEnd();
+        });
+        split.Publish();
+        return;
+    }
     system.GPU().FlushCommands();
     system.GPU().OnCommandListEnd();
 }
@@ -97,7 +108,16 @@ bool DmaPusher::Step() {
             const bool macro_payload = engine == Engines::EngineTypes::Maxwell3D && dma_state.method >= MacroRegistersStart;
             if (kepler_payload || macro_payload) {
                 const size_t words = std::min<size_t>(dma_state.method_count, header.size);
-                subchannels[dma_state.subchannel]->current_dirty = memory_manager.IsMemoryDirty(dma_state.dma_get, words * sizeof(u32));
+                auto* const engine_state = subchannels[dma_state.subchannel];
+                if (auto& split = VideoCommon::GpuSplit::Get(); split.Active()) {
+                    // The caches behind IsMemoryDirty belong to the split's back thread.
+                    split.Defer([engine_state, memory = &memory_manager, address = GPUVAddr{dma_state.dma_get},
+                                 bytes = words * sizeof(u32)] {
+                        engine_state->current_dirty = memory->IsMemoryDirty(address, bytes);
+                    });
+                } else {
+                    engine_state->current_dirty = memory_manager.IsMemoryDirty(dma_state.dma_get, words * sizeof(u32));
+                }
             }
         }
         if (command_resume_offset == 0) {
@@ -115,9 +135,9 @@ bool DmaPusher::Step() {
                              u64{header.allow_flush}, split_method_end, parked_methods.size());
                     std::array<u32, 8> next_words{};
                     std::array<u32, 4> prev_words{};
-                    memory_manager.ReadBlock(header.addr, next_words.data(),
+                    memory_manager.ReadBlockUnsafe(header.addr, next_words.data(),
                                              std::min<size_t>(header.size, 8) * sizeof(u32));
-                    memory_manager.ReadBlock(split_method_end - sizeof(prev_words), prev_words.data(),
+                    memory_manager.ReadBlockUnsafe(split_method_end - sizeof(prev_words), prev_words.data(),
                                              sizeof(prev_words));
                     LOG_DEBUG(HW_GPU, "  previous entry ends {:08x}; next entry starts {:08x}",
                               fmt::join(prev_words, " "), fmt::join(next_words, " "));
@@ -175,6 +195,7 @@ bool DmaPusher::Step() {
     }
 
     if (signal_sync) {
+        VideoCommon::GpuSplit::Get().Drain();
         rasterizer->SignalFence([this]() {
             std::scoped_lock lk(sync_mutex);
             synced = true;
@@ -315,10 +336,18 @@ void DmaPusher::WatchArgument(u32 method, u32 argument) const {
     }
 }
 
+bool DmaPusher::WantsSegment() const {
+    // Only macro parameters and compute uploads read the engine's current_dma_segment.
+    const auto engine = subchannel_type[dma_state.subchannel];
+    return engine == Engines::EngineTypes::KeplerCompute ||
+           (engine == Engines::EngineTypes::Maxwell3D && dma_state.method >= MacroRegistersStart);
+}
+
 void DmaPusher::CallMethod(u32 argument) {
     if (g_semaphore_watch != 0) [[unlikely]] {
         WatchArgument(dma_state.method, argument);
     }
+    auto& split = VideoCommon::GpuSplit::Get();
     if (dma_state.method < non_puller_methods) {
         puller.CallPullerMethod(*this, Engines::Puller::MethodCall{
             dma_state.method,
@@ -332,9 +361,34 @@ void DmaPusher::CallMethod(u32 argument) {
             ReportUnboundSubchannel();
             return;
         }
+        if (split.Active()) {
+            const GPUVAddr segment = dma_state.dma_get + dma_state.dma_word_offset;
+            if (split.FrontMacros()) {
+                const auto type = subchannel_type[dma_state.subchannel];
+                if (type == Engines::EngineTypes::Maxwell3D) {
+                    auto* const maxwell3d = static_cast<Engines::Maxwell3D*>(subchannel);
+                    if (maxwell3d->GetNativeMacroEngine().native_turing) {
+                        maxwell3d->FrontCallMethod(system, split, dma_state.method, argument,
+                                                   dma_state.is_last_call);
+                        return;
+                    }
+                } else if (type == Engines::EngineTypes::KeplerCompute) {
+                    auto* const compute = static_cast<Engines::KeplerCompute*>(subchannel);
+                    if (compute->FrontRunsMacros()) {
+                        compute->FrontCallMethod(system, split, dma_state.method, argument,
+                                                 dma_state.is_last_call, &segment);
+                        return;
+                    }
+                }
+            }
+            split.Method(subchannel, dma_state.method, argument, dma_state.is_last_call,
+                         WantsSegment() ? &segment : nullptr);
+            return;
+        }
         if (!subchannel->execution_mask[dma_state.method]) {
             subchannel->method_sink.emplace_back(dma_state.method, argument);
         } else {
+            VideoCommon::Nx2Profile::Scope method_scope{VideoCommon::Nx2Profile::Zone::Method};
             subchannel->ConsumeSink(system);
             subchannel->current_dma_segment = dma_state.dma_get + dma_state.dma_word_offset;
             subchannel->CallMethod(system, dma_state.method, argument, dma_state.is_last_call);
@@ -348,6 +402,7 @@ void DmaPusher::CallMultiMethod(const u32* base_start, u32 num_methods) {
             WatchArgument(dma_state.method + (dma_state.non_incrementing ? 0 : i), base_start[i]);
         }
     }
+    auto& split = VideoCommon::GpuSplit::Get();
     if (dma_state.method < non_puller_methods) {
         puller.CallMultiMethod(*this, dma_state.method, dma_state.subchannel, base_start, num_methods, dma_state.method_count);
     } else {
@@ -356,6 +411,31 @@ void DmaPusher::CallMultiMethod(const u32* base_start, u32 num_methods) {
             ReportUnboundSubchannel();
             return;
         }
+        if (split.Active()) {
+            const GPUVAddr segment = dma_state.dma_get + dma_state.dma_word_offset;
+            if (split.FrontMacros()) {
+                const auto type = subchannel_type[dma_state.subchannel];
+                if (type == Engines::EngineTypes::Maxwell3D) {
+                    auto* const maxwell3d = static_cast<Engines::Maxwell3D*>(subchannel);
+                    if (maxwell3d->GetNativeMacroEngine().native_turing) {
+                        maxwell3d->FrontCallMultiMethod(system, split, dma_state.method, base_start,
+                                                        num_methods, dma_state.method_count);
+                        return;
+                    }
+                } else if (type == Engines::EngineTypes::KeplerCompute) {
+                    auto* const compute = static_cast<Engines::KeplerCompute*>(subchannel);
+                    if (compute->FrontRunsMacros()) {
+                        compute->FrontCallMultiMethod(system, split, dma_state.method, base_start,
+                                                      num_methods, dma_state.method_count, &segment);
+                        return;
+                    }
+                }
+            }
+            split.MultiMethod(subchannel, dma_state.method, base_start, num_methods,
+                              dma_state.method_count, WantsSegment() ? &segment : nullptr);
+            return;
+        }
+        VideoCommon::Nx2Profile::Scope method_scope{VideoCommon::Nx2Profile::Zone::Method};
         subchannel->ConsumeSink(system);
         subchannel->current_dma_segment = dma_state.dma_get + dma_state.dma_word_offset;
         subchannel->CallMultiMethod(system, dma_state.method, base_start, num_methods, dma_state.method_count);
diff --git a/src/video_core/dma_pusher.h b/src/video_core/dma_pusher.h
index bbcc345c84..82d29edb4b 100644
--- a/src/video_core/dma_pusher.h
+++ b/src/video_core/dma_pusher.h
@@ -172,6 +172,7 @@ private:
 
     void CallMethod(u32 argument);
     void CallMultiMethod(const u32* base_start, u32 num_methods);
+    [[nodiscard]] bool WantsSegment() const;
     void ReportUnboundSubchannel();
     void WatchArgument(u32 method, u32 argument) const;
 
diff --git a/src/video_core/engines/kepler_compute.cpp b/src/video_core/engines/kepler_compute.cpp
index 7792bc9b5d..ab2ba10162 100644
--- a/src/video_core/engines/kepler_compute.cpp
+++ b/src/video_core/engines/kepler_compute.cpp
@@ -17,12 +17,17 @@
 #include "common/logging.h"
 #include "shader_recompiler/exception.h"
 #include "core/core.h"
+#include "video_core/gpu_split.h"
+#include "video_core/nx2_profile.h"
 #include "video_core/nx2.h"
 #include "video_core/engines/fermi_2d.h"
 #include "video_core/engines/kepler_compute.h"
 #include "video_core/engines/maxwell_3d.h"
 #include "video_core/memory_manager.h"
 #include "video_core/native_query_trace.h"
+#ifdef ARCHITECTURE_x86_64
+#include "video_core/macro_turing_jit.h"
+#endif
 #include "video_core/rasterizer_interface.h"
 #include "video_core/textures/decoders.h"
 
@@ -43,6 +48,20 @@ KeplerCompute::KeplerCompute(MemoryManager& memory_manager_)
     execution_mask[KEPLER_COMPUTE_REG_INDEX(exec_upload)] = true;
     execution_mask[KEPLER_COMPUTE_REG_INDEX(data_upload)] = true;
     execution_mask[KEPLER_COMPUTE_REG_INDEX(launch)] = true;
+    // Everything ProcessNativeMethod acts on besides storing the value; the QMD and 0x5a
+    // trace range too while native draw diagnostics trace it.
+    native_plain_methods.set();
+    for (const u32 method : {0x45U, 0x46U, 0x47U, 0x48U, 0xaeU, 0xafU, 0xb0U, 0x107U, 0x6c3U,
+                             static_cast<u32>(KEPLER_COMPUTE_REG_INDEX(exec_upload)),
+                             static_cast<u32>(KEPLER_COMPUTE_REG_INDEX(data_upload))}) {
+        native_plain_methods[method] = false;
+    }
+    if (std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS") != nullptr) {
+        native_plain_methods[0x5a] = false;
+        for (u32 method = 0xc6; method <= 0x108; ++method) {
+            native_plain_methods[method] = false;
+        }
+    }
 }
 
 KeplerCompute::~KeplerCompute() = default;
@@ -84,8 +103,8 @@ void KeplerCompute::ProcessNativeReportSemaphore(u32 argument) {
         rasterizer->SignalFence([this, address, payload, reduction] {
             const u32 old_value = memory_manager.Read<u32>(address);
             const u32 value = reduction == 0   ? old_value + payload
-                              : reduction == 1 ? std::min(old_value, payload)
-                              : reduction == 2 ? std::max(old_value, payload)
+                              : reduction == 1 ? (std::min)(old_value, payload)
+                              : reduction == 2 ? (std::max)(old_value, payload)
                               : reduction == 3 ? (old_value >= payload ? 0 : old_value + 1)
                               : reduction == 4 ? (old_value == 0 || old_value > payload ? payload : old_value - 1)
                               : reduction == 5 ? (old_value & payload)
@@ -102,7 +121,7 @@ void KeplerCompute::ProcessNativeReportSemaphore(u32 argument) {
 }
 
 bool KeplerCompute::ProcessNativeMethod(Core::System& system, u32 method, u32 argument, bool is_last_call) {
-    if (native_context) {
+    if (native_context && !front_macros) {
         if (method == 0x45 || method == 0x47) native_compute_uploaded = true;
         native_macro = native_compute_uploaded ? &native_context->compute_engine : &native_context->engine;
         native_macro_positions = native_compute_uploaded ? std::span<u32>{native_context->compute_positions} : std::span<u32>{native_context->positions};
@@ -154,9 +173,36 @@ bool KeplerCompute::ProcessNativeMethod(Core::System& system, u32 method, u32 ar
             if (watched && log_emits) LOG_INFO(HW_GPU, "Semaphore watch: compute macro emit {:#x} = {:#x}", address, value);
             if (diagnostic_writes++ < 4096) LOG_DEBUG(HW_GPU, "Native compute MME emit method={:#x}, value={:#x}", address, value);
             if (address >= 0xe00) return false;
+            VideoCommon::Nx2Profile::Scope method_scope{VideoCommon::Nx2Profile::Zone::Method};
             CallMethod(system, address, value, true);
             return true;
         };
+#ifdef ARCHITECTURE_x86_64
+        // Compiled macro code, as for 3D macros (macro_turing_jit.h); EDEN_NX2_MME_JIT=0 off.
+        static const bool use_jit = [] {
+            const char* value = std::getenv("EDEN_NX2_MME_JIT");
+            return value == nullptr || value[0] != '0';
+        }();
+        if (use_jit && !watched) {
+            if (!native_macro->native_jit) {
+                native_macro->native_jit = std::make_shared<TuringMacroJit>();
+            }
+            native_macro->native_jit->Prepare(native_macro->native_code,
+                                              native_macro->native_code_generation);
+            const auto result = native_macro->native_jit->Execute(
+                start, native_parameters,
+                std::span<u32>{*native_macro->native_interpreter.active_dram},
+                [&](u32 address) { return address < native_registers.size() ? native_registers[address] : 0U; },
+                emit);
+            if (!result.success) {
+                native_macro->native_execution_failed = true;
+                LOG_ERROR(HW_GPU, "Native MME (JIT) failed: compute start={:#x}, consumed={}/{}, error={}",
+                          start, result.consumed, native_parameters.size(), result.error);
+            }
+            native_parameters.clear();
+            return true;
+        }
+#endif
         // Repeated calls replay their writes, as for 3D macros (macro_memo.h).
         static const bool use_memo = [] {
             const char* value = std::getenv("EDEN_NX2_MME_MEMO");
@@ -263,6 +309,9 @@ bool KeplerCompute::ProcessNativeMethod(Core::System& system, u32 method, u32 ar
         }
         return true;
     }
+    if (front_macros && method >= 0x45 && method <= 0x48) {
+        return true; // Loaded and bound by the split GPU thread's front (FrontApply).
+    }
     switch (method) {
     case 0x45: native_macro->ClearCode(argument); return true;
     case 0x46: native_macro->AddCode(0, argument); return true;
@@ -334,7 +383,7 @@ bool KeplerCompute::ProcessNativeMethod(Core::System& system, u32 method, u32 ar
                     // releases also advance: a plain write would move it backwards and leave
                     // GEQ acquires waiting forever.
                     rasterizer->SignalFence([this, address, payload] {
-                        const u32 value = std::max(memory_manager.Read<u32>(address), payload);
+                        const u32 value = (std::max)(memory_manager.Read<u32>(address), payload);
                         memory_manager.Write<u32>(address, value);
                         if (const auto device = memory_manager.GpuToCpuAddress(address)) {
                             Tegra::Nx2SemaphoreWaiters::Get().OnWrite(*device, value);
@@ -390,7 +439,19 @@ bool KeplerCompute::ProcessNativeMethod(Core::System& system, u32 method, u32 ar
 }
 
 void KeplerCompute::CallMethod(Core::System& system, u32 method, u32 method_argument, bool is_last_call) {
-    if (native_mode && ProcessNativeMethod(system, method, method_argument, is_last_call)) return;
+    if (native_mode) {
+        // Most of the compute class's register writes (QMD fields replayed by its macros for
+        // each of thousands of dispatches a second) only store the value: ~4% of title 1's GPU
+        // thread went through ProcessNativeMethod's checks for them.
+        if (method < native_plain_methods.size() && native_plain_methods[method]) {
+            native_registers[method] = method_argument;
+            if (method < Regs::NUM_REGS) {
+                regs.reg_array[method] = method_argument;
+            }
+            return;
+        }
+        if (ProcessNativeMethod(system, method, method_argument, is_last_call)) return;
+    }
     if (method >= Regs::NUM_REGS) {
         LOG_ERROR(HW_GPU, "Compute register out of range: {:#x}", method);
         return;
@@ -697,4 +758,162 @@ Texture::TSCEntry KeplerCompute::GetTSCEntry(u32 tsc_index) const {
     return tsc_entry;
 }
 
+// Split GPU thread with macros on its front half (gpu_split.h, EDEN_NX2_SPLIT=2): the front
+// runs compute macros against its own copy of the native registers (which hold the last value
+// written to each, on the back too) and loads and binds the compute macro code itself.
+struct KeplerCompute::FrontState {
+    std::array<u32, 0x1000> registers{};
+    std::optional<u32> executing_macro;
+    std::vector<u32> parameters;
+    bool compute_uploaded{};
+    u32 bind_pointer{};
+    MacroEngine* macro{}; ///< Without a native context: the one macro engine
+};
+
+KeplerCompute::FrontState& KeplerCompute::Front(VideoCommon::GpuSplit& split) {
+    if (!front) [[unlikely]] {
+        // Copy the back's state and take the macros over while it is idle.
+        split.Drain();
+        front = std::make_unique<FrontState>();
+        front->registers = native_registers;
+        front->executing_macro = native_executing_macro;
+        front->parameters = std::move(native_parameters);
+        native_parameters.clear();
+        native_executing_macro.reset();
+        front->compute_uploaded = native_compute_uploaded;
+        front->bind_pointer = native_bind_pointer;
+        front->macro = native_macro;
+        front_macros = true;
+    }
+    return *front;
+}
+
+void KeplerCompute::FrontApply(FrontState& fs, u32 method, u32 argument) {
+    // As ProcessNativeMethod does with the macro code and bindings.
+    if (method == 0x45 || method == 0x47) {
+        fs.compute_uploaded = true;
+    }
+    fs.registers[method] = argument;
+    switch (method) {
+    case 0x45:
+        FrontMacroEngine(fs).ClearCode(argument);
+        break;
+    case 0x46:
+        FrontMacroEngine(fs).AddCode(0, argument);
+        break;
+    case 0x47:
+        fs.bind_pointer = argument;
+        break;
+    case 0x48: {
+        const auto positions = FrontMacroPositions(fs);
+        if (fs.bind_pointer < positions.size()) {
+            positions[fs.bind_pointer++] = argument;
+        } else {
+            LOG_ERROR(HW_GPU, "Native compute macro bind out of bounds");
+        }
+        break;
+    }
+    default:
+        break;
+    }
+}
+
+MacroEngine& KeplerCompute::FrontMacroEngine(FrontState& fs) {
+    if (native_context) {
+        return fs.compute_uploaded ? native_context->compute_engine : native_context->engine;
+    }
+    return *fs.macro;
+}
+
+std::span<u32> KeplerCompute::FrontMacroPositions(FrontState& fs) {
+    if (native_context) {
+        return fs.compute_uploaded ? std::span<u32>{native_context->compute_positions}
+                                      : std::span<u32>{native_context->positions};
+    }
+    return native_macro_positions;
+}
+
+void KeplerCompute::FrontCallMethod(Core::System& system, VideoCommon::GpuSplit& split,
+                                    u32 method, u32 argument, bool is_last_call,
+                                    const GPUVAddr* segment) {
+    auto& fs = Front(split);
+    if (method >= fs.registers.size()) {
+        split.Method(this, method, argument, is_last_call, segment);
+        return;
+    }
+    if (method >= 0xe00) {
+        if (!fs.executing_macro) {
+            if (method & 1) {
+                LOG_ERROR(HW_GPU, "Native compute macro starts on argument method");
+                return;
+            }
+            fs.executing_macro = method;
+        }
+        fs.parameters.push_back(argument);
+        if (is_last_call) {
+            FrontRunMacro(system, split, fs);
+        }
+        return;
+    }
+    FrontApply(fs, method, argument);
+    split.Method(this, method, argument, is_last_call, segment);
+}
+
+void KeplerCompute::FrontCallMultiMethod(Core::System& system, VideoCommon::GpuSplit& split,
+                                         u32 method, const u32* base_start, u32 amount,
+                                         u32 methods_pending, const GPUVAddr* segment) {
+    auto& fs = Front(split);
+    if (method >= 0xe00 && method < fs.registers.size()) {
+        for (u32 i = 0; i < amount; ++i) {
+            FrontCallMethod(system, split, method, base_start[i], methods_pending - i <= 1, segment);
+        }
+        return;
+    }
+    if (method < fs.registers.size()) {
+        for (u32 i = 0; i < amount; ++i) {
+            FrontApply(fs, method, base_start[i]);
+        }
+    }
+    split.MultiMethod(this, method, base_start, amount, methods_pending, segment);
+}
+
+void KeplerCompute::FrontRunMacro(Core::System& system, VideoCommon::GpuSplit& split,
+                                  FrontState& fs) {
+#ifdef ARCHITECTURE_x86_64
+    const auto positions = FrontMacroPositions(fs);
+    const u32 start = positions[((*fs.executing_macro - 0xe00) / 2) % positions.size()];
+    fs.executing_macro.reset();
+    auto& macro = FrontMacroEngine(fs);
+    if (native_data_ram) {
+        macro.native_interpreter.active_dram = native_data_ram;
+    }
+    VideoCommon::Nx2Profile::Scope profile_scope{VideoCommon::Nx2Profile::Zone::Macro};
+    if (!macro.native_jit) {
+        macro.native_jit = std::make_shared<TuringMacroJit>();
+    }
+    macro.native_jit->Prepare(macro.native_code, macro.native_code_generation);
+    const auto result = macro.native_jit->Execute(
+        start, fs.parameters, std::span<u32>{*macro.native_interpreter.active_dram},
+        [&fs](u32 address) {
+            return address < fs.registers.size() ? fs.registers[address] : 0U;
+        },
+        [&](u32 address, u32 value) {
+            if (address >= 0xe00) {
+                return false;
+            }
+            FrontApply(fs, address, value);
+            split.MacroWrite(this, address, value);
+            return true;
+        });
+    if (!result.success) {
+        macro.native_execution_failed = true;
+        LOG_ERROR(HW_GPU, "Native MME (JIT, split front) failed: compute start={:#x}, consumed={}, error={}",
+                  start, result.consumed, result.error);
+    }
+#else
+    fs.executing_macro.reset();
+#endif
+    fs.parameters.clear();
+}
+
 } // namespace Tegra::Engines
diff --git a/src/video_core/engines/kepler_compute.h b/src/video_core/engines/kepler_compute.h
index 5ebf8edf38..b525636dde 100644
--- a/src/video_core/engines/kepler_compute.h
+++ b/src/video_core/engines/kepler_compute.h
@@ -9,7 +9,9 @@
 #include <array>
 #include <atomic>
 #include <cstddef>
+#include <memory>
 #include <optional>
+#include <span>
 #include <vector>
 #include "common/bit_field.h"
 #include "common/common_funcs.h"
@@ -24,6 +26,10 @@ namespace Core {
 class System;
 }
 
+namespace VideoCommon {
+class GpuSplit;
+}
+
 namespace Tegra {
 class MemoryManager;
 }
@@ -237,7 +243,26 @@ public:
 
     bool IsNativeMode() const { return native_mode; }
 
+    /// GPU thread split with macros on its front half (gpu_split.h, EDEN_NX2_SPLIT=2), called by
+    /// the front instead of CallMethod for native compute engines: runs their macros and macro
+    /// uploads against a copy of the registers and sends the register writes to the back.
+    bool FrontRunsMacros() const { return native_mode; }
+    void FrontCallMethod(Core::System& system, VideoCommon::GpuSplit& split, u32 method,
+                         u32 argument, bool is_last_call, const GPUVAddr* segment);
+    void FrontCallMultiMethod(Core::System& system, VideoCommon::GpuSplit& split, u32 method,
+                              const u32* base_start, u32 amount, u32 methods_pending,
+                              const GPUVAddr* segment);
+
 private:
+    struct FrontState;
+    FrontState& Front(VideoCommon::GpuSplit& split);
+    void FrontApply(FrontState& state, u32 method, u32 argument);
+    MacroEngine& FrontMacroEngine(FrontState& state);
+    std::span<u32> FrontMacroPositions(FrontState& state);
+    void FrontRunMacro(Core::System& system, VideoCommon::GpuSplit& split, FrontState& state);
+    std::unique_ptr<FrontState> front;
+    bool front_macros{}; ///< The front owns the macro code and bindings (the back skips them)
+
     bool DispatchNativeCompute(const NativeQmdScheduler::Descriptor& descriptor);
     void ProcessNativeReportSemaphore(u32 argument);
     void DescribeNativeResources(const NativeQmdScheduler::Descriptor& qmd);
@@ -252,6 +277,9 @@ private:
     bool native_compute_uploaded{};
     MacroEngine* native_macro{};
     std::array<u32, 0x1000> native_registers{};
+    /// Native methods that only store their value (no macro, QMD, upload, semaphore or code
+    /// side effect): CallMethod writes those directly.
+    std::bitset<0xe00> native_plain_methods;
     std::span<u32> native_macro_positions;
     std::vector<u32> native_parameters;
     std::vector<u32> native_memo_parameters; ///< The current call's parameters, for the memo
diff --git a/src/video_core/engines/maxwell_3d.cpp b/src/video_core/engines/maxwell_3d.cpp
index 60efef9b8e..6adfcf89b8 100644
--- a/src/video_core/engines/maxwell_3d.cpp
+++ b/src/video_core/engines/maxwell_3d.cpp
@@ -4,6 +4,8 @@
 // SPDX-FileCopyrightText: Copyright 2018 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
+#include <algorithm>
+#include <array>
 #include <cstring>
 #include <limits>
 #include <optional>
@@ -17,9 +19,14 @@
 #include "video_core/engines/maxwell_3d.h"
 #include "video_core/engines/maxwell_3d.h"
 #include "video_core/gpu.h"
+#include "video_core/gpu_split.h"
 #include "video_core/memory_manager.h"
+#include "video_core/nx2_profile.h"
 #include "video_core/rasterizer_interface.h"
 #include "video_core/textures/texture.h"
+#ifdef ARCHITECTURE_x86_64
+#include "video_core/macro_turing_jit.h"
+#endif
 
 namespace Tegra::Engines {
 
@@ -208,11 +215,29 @@ void Maxwell3D::ProcessMacro(Core::System& system, u32 method, const u32* base_s
         executing_macro = method;
     }
 
-    macro_params.insert(macro_params.end(), base_start, base_start + amount);
     macro_segments.emplace_back(current_dma_segment, amount);
     current_macro_dirty |= current_dirty;
     current_dirty = false;
 
+    // Nearly every call arrives whole in one method run (title 1 makes ~1,000,000 a second):
+    // copy it to the stack instead of growing the parameter vector. Still a copy: the command
+    // list is often guest memory read in place, which the guest may reuse while the macro runs
+    // (and the macro memo stores the parameters after it ran). Only accurate GPU mode rewrites
+    // the parameters in place (RefreshParameters), so it keeps the vector.
+    // EDEN_NX2_MACRO_COPY=1: always use the vector (for comparisons).
+    static const bool always_copy = std::getenv("EDEN_NX2_MACRO_COPY") != nullptr;
+    if (is_last_call && macro_params.empty() && amount <= 64 && !always_copy &&
+        !Settings::IsGPULevelHigh()) {
+        std::array<u32, 64> parameters;
+        std::copy_n(base_start, amount, parameters.begin());
+        ConsumeSink(system);
+        CallMacroMethod(system, executing_macro, std::span<const u32>{parameters.data(), amount});
+        macro_segments.clear();
+        current_macro_dirty = false;
+        return;
+    }
+    macro_params.insert(macro_params.end(), base_start, base_start + amount);
+
     // Call the macro when there are no more parameters in the command buffer
     if (is_last_call) {
         ConsumeSink(system);
@@ -257,7 +282,7 @@ u32 Maxwell3D::GetMaxCurrentVertices() {
         const u64 stream_size = GetNativeMacroEngine().native_turing
                                     ? regs.NativeVertexStreamSize(index)
                                     : gpu_addr_end - gpu_addr_begin;
-        const u32 address_size = static_cast<u32>(std::min<u64>(stream_size, std::numeric_limits<u32>::max()));
+        const u32 address_size = static_cast<u32>((std::min<u64>)(stream_size, (std::numeric_limits<u32>::max)()));
         num_vertices = (std::max)(
             num_vertices, address_size / (std::max)(attribute.SizeInBytes(), array.stride.Value()));
         break;
@@ -287,17 +312,25 @@ u32 Maxwell3D::ProcessShadowRam(u32 method, u32 argument) {
 
 void Maxwell3D::ConsumeSinkImpl(Core::System& system) {
     const auto control = shadow_state.shadow_ram_control;
+    // Sunk methods are state writes only (not in execution_mask): skip unchanged values, as
+    // CallMethod does.
+    const bool skip_same = SkipSameValueWrites();
+    const auto process = [&](u32 method, u32 value) {
+        if (!skip_same || regs.reg_array[method] != value) {
+            ProcessDirtyRegisters(method, value);
+        }
+    };
     if (control == Regs::ShadowRamControl::Track || control == Regs::ShadowRamControl::TrackWithFilter) {
         for (auto [method, value] : method_sink) {
             shadow_state.reg_array[method] = value;
-            ProcessDirtyRegisters(method, value);
+            process(method, value);
         }
     } else if (control == Regs::ShadowRamControl::Replay) {
         for (auto [method, value] : method_sink)
-            ProcessDirtyRegisters(method, shadow_state.reg_array[method]);
+            process(method, shadow_state.reg_array[method]);
     } else {
         for (auto [method, value] : method_sink)
-            ProcessDirtyRegisters(method, value);
+            process(method, value);
     }
     method_sink.clear();
 }
@@ -318,11 +351,18 @@ void Maxwell3D::ProcessMethodCall(u32 method, u32 argument, u32 nonshadow_argume
     case MAXWELL3D_REG_INDEX(shadow_ram_control):
         shadow_state.shadow_ram_control = static_cast<Regs::ShadowRamControl>(nonshadow_argument);
         return;
+    // With macros on the split GPU thread's front (FrontWrite), the front loads and binds them.
     case MAXWELL3D_REG_INDEX(load_mme.instruction_ptr):
+        if (front_macros) return;
         return GetNativeMacroEngine().ClearCode(regs.load_mme.instruction_ptr);
     case MAXWELL3D_REG_INDEX(load_mme.instruction):
+        if (front_macros) return;
         return GetNativeMacroEngine().AddCode(regs.load_mme.instruction_ptr, argument);
     case MAXWELL3D_REG_INDEX(load_mme.start_address):
+        if (front_macros) {
+            ++regs.load_mme.start_address_ptr;
+            return;
+        }
         return ProcessMacroBind(argument);
     case MAXWELL3D_REG_INDEX(falcon[4]):
         return ProcessFirmwareCall4();
@@ -379,7 +419,7 @@ void Maxwell3D::ProcessMethodCall(u32 method, u32 argument, u32 nonshadow_argume
     }
 }
 
-void Maxwell3D::CallMacroMethod(Core::System& system, u32 method, const std::vector<u32>& parameters) {
+void Maxwell3D::CallMacroMethod(Core::System& system, u32 method, std::span<const u32> parameters) {
     // Reset the current macro.
     executing_macro = 0;
 
@@ -425,7 +465,29 @@ void Maxwell3D::CallMethod(Core::System& system, u32 method, u32 method_argument
     }
 
     ASSERT(method < Regs::NUM_REGS && "Invalid Maxwell3D register, increase the size of the Regs structure");
+    if (VideoCommon::Nx2Profile::Enabled()) [[unlikely]] {
+        VideoCommon::Nx2Profile::Add(VideoCommon::Nx2Profile::Counter::MethodWrites, 1);
+        if (regs.reg_array[method] == method_argument) {
+            VideoCommon::Nx2Profile::Add(VideoCommon::Nx2Profile::Counter::SameValueWrites, 1);
+        }
+    }
+    if (method == MAXWELL3D_REG_INDEX(shadow_ram_control) && front_macros) [[unlikely]] {
+        back_control_writes[back_control_count++ % back_control_writes.size()] = {
+            method_argument, static_cast<u32>(shadow_state.shadow_ram_control), is_last_call ? 1U : 0U};
+    }
     const u32 argument = ProcessShadowRam(method, method_argument);
+    // A state register written with the value it holds changes nothing: no dirty flags (they
+    // would make the next draw re-check that state), no method processing. Title 1 makes ~6
+    // million 3D register writes a second, ~60% of them same-value, most from macro replays.
+    // Methods with side effects (execution_mask: draws, clears, constant buffer data and
+    // binds, semaphores, ...) always run. EDEN_NX2_SKIP_SAME=0 turns it off.
+    if (!execution_mask[method]) {
+        // ProcessMethodCall (and the draw manager behind it) acts only on executable methods.
+        if (!SkipSameValueWrites() || regs.reg_array[method] != argument) {
+            ProcessDirtyRegisters(method, argument);
+        }
+        return;
+    }
     ProcessDirtyRegisters(method, argument);
     ProcessMethodCall(method, argument, method_argument, is_last_call);
 }
@@ -720,4 +782,416 @@ void Maxwell3D::SetHLEReplacementAttributeType(u32 bank, u32 offset,
     replace_table.emplace(key, name);
 }
 
+// Split GPU thread with macros on its front half (gpu_split.h, EDEN_NX2_SPLIT=2). The front keeps
+// what macros read: the registers as the back thread will hold them when it reaches the same
+// point of the method stream. That is every write the front sends, with the shadow RAM applied,
+// plus the few changes the back makes on its own (constant buffer offset, shadow_scratch[0] on
+// falcon[4], the macro bind pointer); the front loads and binds the macro code itself.
+struct Maxwell3D::FrontState {
+    std::array<ControlWrite, 16> control_writes{};
+    u32 control_count{};
+    std::array<u32, Regs::NUM_REGS> regs{};
+    std::array<u32, Regs::NUM_REGS> shadow{};
+    Regs::ShadowRamControl control{};
+    u32 executing_macro{};
+    std::vector<u32> parameters;
+    std::vector<u32> fifo;
+    u64 calls{};
+    /// A method the draw manager handles was sent since the back last ran DrawDeferred: only
+    /// then can it have a deferred instanced draw to issue at the end of a macro.
+    bool draw_pending{true};
+    bool skip_same{}; ///< SkipSameValueWrites() && FrontSkipSame(), read once
+};
+
+namespace {
+
+// EDEN_NX2_SPLIT_VERIFY=1: every 4096 front macro calls, wait for the back and compare the
+// front's registers with the back's (logs and resynchronizes differences).
+bool FrontVerifyEnabled() {
+    static const bool verify = std::getenv("EDEN_NX2_SPLIT_VERIFY") != nullptr;
+    return verify;
+}
+
+// Executable methods ProcessMethodCall handles itself; the rest go to the draw manager.
+bool FrontHandledOutsideDrawManager(u32 method) {
+    constexpr u32 cb_data = MAXWELL3D_REG_INDEX(const_buffer.buffer);
+    if (method >= cb_data && method < cb_data + 16) {
+        return true;
+    }
+    switch (method) {
+    case MAXWELL3D_REG_INDEX(wait_for_idle):
+    case MAXWELL3D_REG_INDEX(shadow_ram_control):
+    case MAXWELL3D_REG_INDEX(load_mme.instruction_ptr):
+    case MAXWELL3D_REG_INDEX(load_mme.instruction):
+    case MAXWELL3D_REG_INDEX(load_mme.start_address):
+    case MAXWELL3D_REG_INDEX(falcon[4]):
+    case MAXWELL3D_REG_INDEX(bind_groups[0].raw_config):
+    case MAXWELL3D_REG_INDEX(bind_groups[1].raw_config):
+    case MAXWELL3D_REG_INDEX(bind_groups[2].raw_config):
+    case MAXWELL3D_REG_INDEX(bind_groups[3].raw_config):
+    case MAXWELL3D_REG_INDEX(bind_groups[4].raw_config):
+    case MAXWELL3D_REG_INDEX(report_semaphore.query):
+    case MAXWELL3D_REG_INDEX(render_enable.mode):
+    case MAXWELL3D_REG_INDEX(clear_report_value):
+    case MAXWELL3D_REG_INDEX(sync_info):
+    case MAXWELL3D_REG_INDEX(launch_dma):
+    case MAXWELL3D_REG_INDEX(inline_data):
+    case MAXWELL3D_REG_INDEX(fragment_barrier):
+    case MAXWELL3D_REG_INDEX(invalidate_texture_data_cache):
+    case MAXWELL3D_REG_INDEX(tiled_cache_barrier):
+        return true;
+    default:
+        return false;
+    }
+}
+
+// EDEN_NX2_SPLIT_SKIP=0: send state writes that keep the register's value too.
+bool FrontSkipSame() {
+    static const bool skip = [] {
+        const char* value = std::getenv("EDEN_NX2_SPLIT_SKIP");
+        return value == nullptr || value[0] != '0';
+    }();
+    return skip;
+}
+
+} // Anonymous namespace
+
+Maxwell3D::FrontState& Maxwell3D::Front(Core::System& system, VideoCommon::GpuSplit& split) {
+    if (!front) [[unlikely]] {
+        // Copy the back's state and take the macros over while it is idle (state writes it
+        // still holds in its method sink first).
+        split.Drain();
+        ConsumeSink(system);
+        front = std::make_unique<FrontState>();
+        std::memcpy(front->regs.data(), regs.reg_array.data(), sizeof(front->regs));
+        std::memcpy(front->shadow.data(), shadow_state.reg_array.data(), sizeof(front->shadow));
+        front->control = shadow_state.shadow_ram_control;
+        front->skip_same = SkipSameValueWrites() && FrontSkipSame();
+        front->executing_macro = executing_macro;
+        front->parameters = std::move(macro_params);
+        macro_params.clear();
+        executing_macro = 0;
+        front_macros = true;
+    }
+    return *front;
+}
+
+void Maxwell3D::FrontWrite(VideoCommon::GpuSplit& split, FrontState& fs, u32 method,
+                           u32 argument, bool is_last_call, bool from_macro) {
+    namespace Profile = VideoCommon::Nx2Profile;
+    // As ProcessShadowRam does on the back.
+    u32 value = argument;
+    bool shadow_kept = true;
+    if (fs.control == Regs::ShadowRamControl::Track ||
+        fs.control == Regs::ShadowRamControl::TrackWithFilter) {
+        shadow_kept = fs.shadow[method] == argument;
+        fs.shadow[method] = argument;
+    } else if (fs.control == Regs::ShadowRamControl::Replay) {
+        value = fs.shadow[method];
+    }
+    if (Profile::Enabled()) [[unlikely]] {
+        Profile::Add(Profile::Counter::MethodWrites, 1);
+    }
+    if (method == MAXWELL3D_REG_INDEX(shadow_ram_control) && FrontVerifyEnabled()) [[unlikely]] {
+        fs.control_writes[fs.control_count++ % fs.control_writes.size()] = {
+            argument, static_cast<u32>(fs.control), from_macro ? 2U : 1U};
+    }
+    if (!execution_mask[method]) {
+        // Nothing would change on the back: CallMethod and ConsumeSinkImpl skip same-value state
+        // writes, and the shadow RAM (tracked by default) holds the value already.
+        if (fs.regs[method] == value && shadow_kept && fs.skip_same) {
+            if (Profile::Enabled()) [[unlikely]] {
+                Profile::Add(Profile::Counter::SameValueWrites, 1);
+            }
+            return;
+        }
+        fs.regs[method] = value;
+    } else {
+        fs.regs[method] = value;
+        constexpr u32 cb_data = MAXWELL3D_REG_INDEX(const_buffer.buffer);
+        switch (method) {
+        case MAXWELL3D_REG_INDEX(shadow_ram_control):
+            // The back's mode is shadow_state.shadow_ram_control: the shadow copy of this very
+            // register, which a later Replay write reads back.
+            fs.control = static_cast<Regs::ShadowRamControl>(argument);
+            fs.shadow[method] = argument;
+            break;
+        case MAXWELL3D_REG_INDEX(load_mme.instruction_ptr):
+            GetNativeMacroEngine().ClearCode(fs.regs[MAXWELL3D_REG_INDEX(load_mme.instruction_ptr)]);
+            break;
+        case MAXWELL3D_REG_INDEX(load_mme.instruction):
+            GetNativeMacroEngine().AddCode(fs.regs[MAXWELL3D_REG_INDEX(load_mme.instruction_ptr)], value);
+            break;
+        case MAXWELL3D_REG_INDEX(load_mme.start_address): {
+            const auto positions = GetNativeMacroPositions();
+            u32& pointer = fs.regs[MAXWELL3D_REG_INDEX(load_mme.start_address_ptr)];
+            if (pointer < positions.size()) {
+                positions[pointer] = value;
+            }
+            ++pointer;
+            break;
+        }
+        case MAXWELL3D_REG_INDEX(falcon[4]):
+            fs.regs[MAXWELL3D_REG_INDEX(shadow_scratch)] = 1;
+            break;
+        default:
+            if (method >= cb_data && method < cb_data + 16) {
+                fs.regs[MAXWELL3D_REG_INDEX(const_buffer.offset)] += sizeof(u32);
+            } else if (!FrontHandledOutsideDrawManager(method)) {
+                fs.draw_pending = true;
+            }
+            break;
+        }
+    }
+    if (from_macro) {
+        split.MacroWrite(this, method, argument);
+    } else {
+        split.Method(this, method, argument, is_last_call, nullptr);
+    }
+}
+
+void Maxwell3D::FrontCallMethod(Core::System& system, VideoCommon::GpuSplit& split, u32 method,
+                                u32 argument, bool is_last_call) {
+    auto& fs = Front(system, split);
+    if (method >= MacroRegistersStart) {
+        if (fs.executing_macro == 0) {
+            fs.executing_macro = method;
+        }
+        fs.parameters.push_back(argument);
+        if (is_last_call) {
+            FrontRunMacro(system, split, fs);
+        }
+        return;
+    }
+    FrontWrite(split, fs, method, argument, is_last_call, false);
+}
+
+void Maxwell3D::FrontCallMultiMethod(Core::System& system, VideoCommon::GpuSplit& split,
+                                     u32 method, const u32* base_start, u32 amount,
+                                     u32 methods_pending) {
+    auto& fs = Front(system, split);
+    if (method >= MacroRegistersStart) {
+        if (fs.executing_macro == 0) {
+            fs.executing_macro = method;
+        }
+        fs.parameters.insert(fs.parameters.end(), base_start, base_start + amount);
+        if (amount == methods_pending) {
+            FrontRunMacro(system, split, fs);
+        }
+        return;
+    }
+    // What CallMultiMethod does to the registers on the back, then one entry for all the words.
+    constexpr u32 cb_data = MAXWELL3D_REG_INDEX(const_buffer.buffer);
+    const auto tracked = [&fs] {
+        return fs.control == Regs::ShadowRamControl::Track ||
+               fs.control == Regs::ShadowRamControl::TrackWithFilter;
+    };
+    if (method >= cb_data && method < cb_data + 16) {
+        fs.regs[MAXWELL3D_REG_INDEX(const_buffer.offset)] += amount * sizeof(u32);
+    } else if (method == MAXWELL3D_REG_INDEX(inline_data)) {
+        // Upload data only.
+    } else if ((method == MAXWELL3D_REG_INDEX(draw_inline_index) ||
+                method == MAXWELL3D_REG_INDEX(inline_index_2x16.even) ||
+                method == MAXWELL3D_REG_INDEX(inline_index_4x8.index0)) &&
+               fs.control != Regs::ShadowRamControl::Replay) {
+        if (amount != 0) {
+            const u32 last = base_start[amount - 1];
+            if (tracked()) {
+                fs.shadow[method] = last;
+            }
+            fs.regs[method] = last;
+        }
+        fs.draw_pending = true;
+    } else {
+        // CallMethod for each word, macro uploads among them (load_mme.instruction). A word may
+        // change the shadow RAM mode for the next (a run of SET_SHADOW_RAM_CONTROL writes).
+        for (u32 i = 0; i < amount; ++i) {
+            const u32 argument = base_start[i];
+            u32 value = argument;
+            if (method == MAXWELL3D_REG_INDEX(shadow_ram_control) && FrontVerifyEnabled()) [[unlikely]] {
+                fs.control_writes[fs.control_count++ % fs.control_writes.size()] = {
+                    argument, static_cast<u32>(fs.control), 3U};
+            }
+            if (tracked()) {
+                fs.shadow[method] = argument;
+            } else if (fs.control == Regs::ShadowRamControl::Replay) {
+                value = fs.shadow[method];
+            }
+            fs.regs[method] = value;
+            if (!execution_mask[method]) {
+                continue;
+            }
+            switch (method) {
+            case MAXWELL3D_REG_INDEX(shadow_ram_control):
+                fs.control = static_cast<Regs::ShadowRamControl>(argument);
+                fs.shadow[method] = argument; // The same word on the back (see FrontWrite)
+                break;
+            case MAXWELL3D_REG_INDEX(load_mme.instruction_ptr):
+                GetNativeMacroEngine().ClearCode(value);
+                break;
+            case MAXWELL3D_REG_INDEX(load_mme.instruction):
+                GetNativeMacroEngine().AddCode(fs.regs[MAXWELL3D_REG_INDEX(load_mme.instruction_ptr)], value);
+                break;
+            case MAXWELL3D_REG_INDEX(load_mme.start_address): {
+                const auto positions = GetNativeMacroPositions();
+                u32& pointer = fs.regs[MAXWELL3D_REG_INDEX(load_mme.start_address_ptr)];
+                if (pointer < positions.size()) {
+                    positions[pointer] = value;
+                }
+                ++pointer;
+                break;
+            }
+            case MAXWELL3D_REG_INDEX(falcon[4]):
+                fs.regs[MAXWELL3D_REG_INDEX(shadow_scratch)] = 1;
+                break;
+            default:
+                if (!FrontHandledOutsideDrawManager(method)) {
+                    fs.draw_pending = true;
+                }
+                break;
+            }
+        }
+    }
+    split.MultiMethod(this, method, base_start, amount, methods_pending, nullptr);
+}
+
+void Maxwell3D::FrontRunMacro(Core::System& system, VideoCommon::GpuSplit& split,
+                              FrontState& fs) {
+#ifdef ARCHITECTURE_x86_64
+    const u32 method = fs.executing_macro;
+    fs.executing_macro = 0;
+    auto& engine = GetNativeMacroEngine();
+    const auto positions = GetNativeMacroPositions();
+    const u32 start = positions[((method - MacroRegistersStart) >> 1) % positions.size()];
+    {
+        VideoCommon::Nx2Profile::Scope profile_scope{VideoCommon::Nx2Profile::Zone::Macro};
+        engine.native_interpreter.active_dram = &native_data_ram;
+        if (!engine.native_jit) {
+            engine.native_jit = std::make_shared<TuringMacroJit>();
+        }
+        engine.native_jit->Prepare(engine.native_code, engine.native_code_generation);
+        fs.fifo.assign(fs.parameters.begin(), fs.parameters.end());
+        const auto result = engine.native_jit->Execute(
+            start, fs.fifo, std::span<u32>{native_data_ram},
+            [&fs](u32 address) { return address < Regs::NUM_REGS ? fs.regs[address] : 0U; },
+            [&](u32 address, u32 value) {
+                if (address >= MacroRegistersStart) {
+                    return false;
+                }
+                FrontWrite(split, fs, address, value, true, true);
+                // MME_DMA_READ (into data RAM) and MME_DMA_READ_FIFOED (into the load FIFO).
+                if (address == 0x157 || address == 0x158) {
+                    FrontDmaRead(split, fs, address, value);
+                }
+                return true;
+            });
+        if (!result.success) {
+            engine.native_execution_failed = true;
+            LOG_ERROR(HW_GPU, "Native MME (JIT, split front) failed: start={:#x}, consumed={}/{}, error={}",
+                      start, result.consumed, fs.parameters.size(), result.error);
+        }
+    }
+    fs.parameters.clear();
+    if (fs.draw_pending) {
+        fs.draw_pending = false;
+        split.Defer([this, system_ptr = &system] { FinishFrontMacro(*system_ptr); });
+    }
+    if (FrontVerifyEnabled() && (++fs.calls & 4095) == 0) [[unlikely]] {
+        FrontVerify(split, fs);
+    }
+#else
+    fs.executing_macro = 0;
+    fs.parameters.clear();
+#endif
+}
+
+void Maxwell3D::FrontDmaRead(VideoCommon::GpuSplit& split, FrontState& fs, u32 method,
+                             u32 count) {
+    // The words may be written by methods still queued for the back (indirect draw arguments
+    // from a culling dispatch), and the caches that may hold newer ones belong to it.
+    split.Drain();
+    const GPUVAddr source = (GPUVAddr{fs.regs[0x154] & 0x1ffffff} << 32) | fs.regs[0x155];
+    const u32 ram = fs.regs[0x156];
+    const u32 words = std::min<u32>(count, 0x1000);
+    const size_t read_bytes = size_t{words} * sizeof(u32);
+    if (VideoCommon::Nx2Profile::Enabled()) {
+        VideoCommon::Nx2Profile::Add(VideoCommon::Nx2Profile::Counter::DmaReadBytes, read_bytes);
+    }
+    if (memory_manager.IsMemoryDirty(source, read_bytes, VideoCommon::CacheType::BufferCache)) {
+        memory_manager.FlushRegion(source, read_bytes, VideoCommon::CacheType::BufferCache);
+    }
+    for (u32 i = 0; i < words; ++i) {
+        const u32 word = memory_manager.Read<u32>(source + u64{i} * sizeof(u32));
+        if (method == 0x158) {
+            fs.fifo.push_back(word);
+        } else if (ram + i < native_data_ram.size()) {
+            native_data_ram[ram + i] = word;
+        }
+    }
+}
+
+void Maxwell3D::FrontVerify(VideoCommon::GpuSplit& split, FrontState& fs) {
+    split.Drain();
+    static u64 checks{};
+    static u64 differences{};
+    ++checks;
+    // State writes the back still holds in its method sink count as done.
+    std::array<u32, Regs::NUM_REGS> expected;
+    std::memcpy(expected.data(), regs.reg_array.data(), sizeof(expected));
+    for (const auto& [method, value] : method_sink) {
+        expected[method] = shadow_state.shadow_ram_control == Regs::ShadowRamControl::Replay
+                               ? shadow_state.reg_array[method]
+                               : value;
+    }
+    static u64 control_differences{};
+    if ((fs.control != shadow_state.shadow_ram_control ||
+         fs.shadow[MAXWELL3D_REG_INDEX(shadow_ram_control)] != shadow_state.reg_array[MAXWELL3D_REG_INDEX(shadow_ram_control)] ||
+         fs.regs[MAXWELL3D_REG_INDEX(shadow_ram_control)] != expected[MAXWELL3D_REG_INDEX(shadow_ram_control)]) &&
+        control_differences < 4) {
+        LOG_WARNING(HW_GPU, "Split shadow RAM control writes, front ({} total) / back ({} total):",
+                    fs.control_count, back_control_count);
+        for (u32 i = 0; i < 16; ++i) {
+            const auto& f = fs.control_writes[(fs.control_count + i) % 16];
+            const auto& b = back_control_writes[(back_control_count + i) % 16];
+            LOG_WARNING(HW_GPU, "  front {:#x} (mode {}, path {}) | back {:#x} (mode {}, last {})",
+                        f.argument, f.before, f.path, b.argument, b.before, b.path);
+        }
+    }
+    for (u32 i = 0; i < Regs::NUM_REGS; ++i) {
+        if (fs.regs[i] != expected[i]) {
+            if (differences++ < 64) {
+                LOG_WARNING(HW_GPU, "Split front register {:#x}: front {:#x}, back {:#x}", i,
+                            fs.regs[i], expected[i]);
+            }
+            fs.regs[i] = expected[i];
+        }
+    }
+    if (fs.control != shadow_state.shadow_ram_control) {
+        if (control_differences++ < 16) {
+            LOG_WARNING(HW_GPU, "Split front shadow RAM control {} vs back {}",
+                        static_cast<u32>(fs.control),
+                        static_cast<u32>(shadow_state.shadow_ram_control));
+        }
+        fs.control = shadow_state.shadow_ram_control;
+    }
+    for (u32 i = 0; i < Regs::NUM_REGS; ++i) {
+        if (fs.shadow[i] != shadow_state.reg_array[i]) {
+            if (control_differences++ < 16) {
+                LOG_WARNING(HW_GPU, "Split front shadow register {:#x}: front {:#x}, back {:#x}", i,
+                            fs.shadow[i], shadow_state.reg_array[i]);
+            }
+            fs.shadow[i] = shadow_state.reg_array[i];
+        }
+    }
+    if ((checks & 63) == 1) {
+        LOG_INFO(HW_GPU, "Split front register checks: {} done, {} differences", checks, differences);
+    }
+}
+
+void Maxwell3D::FinishFrontMacro(Core::System& system) {
+    // What ProcessMacro and CallMacroMethod do around the macro.
+    ConsumeSink(system);
+    draw_manager.DrawDeferred(*this);
+}
+
 } // namespace Tegra::Engines
diff --git a/src/video_core/engines/maxwell_3d.h b/src/video_core/engines/maxwell_3d.h
index 82e716b822..0121f1a01c 100644
--- a/src/video_core/engines/maxwell_3d.h
+++ b/src/video_core/engines/maxwell_3d.h
@@ -9,8 +9,10 @@
 #include <array>
 #include <bitset>
 #include <cmath>
+#include <cstdlib>
 #include <limits>
 #include <optional>
+#include <span>
 #include <type_traits>
 #include <vector>
 
@@ -35,6 +37,10 @@ namespace Tegra {
 class MemoryManager;
 }
 
+namespace VideoCommon {
+class GpuSplit;
+}
+
 namespace VideoCore {
 class RasterizerInterface;
 }
@@ -3161,6 +3167,14 @@ public:
     /// Write multiple values to the register identified by method.
     void CallMultiMethod(Core::System& system, u32 method, const u32* base_start, u32 amount, u32 methods_pending) override;
 
+    /// GPU thread split with macros on its front half (gpu_split.h, EDEN_NX2_SPLIT=2), called by
+    /// the front instead of CallMethod: runs macros and macro uploads against a copy of the
+    /// registers and sends the register writes to the back thread.
+    void FrontCallMethod(Core::System& system, VideoCommon::GpuSplit& split, u32 method,
+                         u32 argument, bool is_last_call);
+    void FrontCallMultiMethod(Core::System& system, VideoCommon::GpuSplit& split, u32 method,
+                              const u32* base_start, u32 amount, u32 methods_pending);
+
     bool ShouldExecute() const {
         return execute_on;
     }
@@ -3247,7 +3261,16 @@ private:
      * @param method Method to call
      * @param parameters Arguments to the method call
      */
-    void CallMacroMethod(Core::System& system, u32 method, const std::vector<u32>& parameters);
+    void CallMacroMethod(Core::System& system, u32 method, std::span<const u32> parameters);
+
+    /// Whether state register writes that keep the register's value are dropped (CallMethod).
+    static bool SkipSameValueWrites() {
+        static const bool skip = [] {
+            const char* value = std::getenv("EDEN_NX2_SKIP_SAME");
+            return value == nullptr || value[0] != '0';
+        }();
+        return skip;
+    }
 
     /// Handles writes to the macro uploading register.
     void ProcessMacroUpload(u32 data);
@@ -3313,6 +3336,28 @@ private:
 
     std::vector<std::pair<GPUVAddr, size_t>> macro_segments;
     bool current_macro_dirty{};
+
+    // Split GPU thread, macros on the front (FrontCallMethod): the front's state, and whether
+    // the front owns the macro code and bindings (the back thread then leaves them alone).
+    struct FrontState;
+    FrontState& Front(Core::System& system, VideoCommon::GpuSplit& split);
+    void FrontWrite(VideoCommon::GpuSplit& split, FrontState& state, u32 method, u32 argument,
+                    bool is_last_call, bool from_macro);
+    void FrontRunMacro(Core::System& system, VideoCommon::GpuSplit& split, FrontState& state);
+    void FrontDmaRead(VideoCommon::GpuSplit& split, FrontState& state, u32 method, u32 count);
+    void FrontVerify(VideoCommon::GpuSplit& split, FrontState& state);
+    void FinishFrontMacro(Core::System& system);
+    std::unique_ptr<FrontState> front;
+    bool front_macros{};
+    // EDEN_NX2_SPLIT_VERIFY: the last SET_SHADOW_RAM_CONTROL writes each side saw (argument,
+    // mode before, path), printed when the sides disagree.
+    struct ControlWrite {
+        u32 argument;
+        u32 before;
+        u32 path;
+    };
+    std::array<ControlWrite, 16> back_control_writes{};
+    u32 back_control_count{};
 };
 
 #define ASSERT_REG_POSITION(field_name, position)                                                  \
diff --git a/src/video_core/engines/puller.cpp b/src/video_core/engines/puller.cpp
index b607f20161..0bb1809aa3 100644
--- a/src/video_core/engines/puller.cpp
+++ b/src/video_core/engines/puller.cpp
@@ -11,6 +11,7 @@
 #include "common/logging.h"
 #include "common/settings.h"
 #include "core/core.h"
+#include "video_core/nx2_profile.h"
 #include "video_core/control/channel_state.h"
 #include "video_core/dma_pusher.h"
 #include "video_core/engines/fermi_2d.h"
@@ -20,6 +21,7 @@
 #include "video_core/engines/maxwell_dma.h"
 #include "video_core/engines/puller.h"
 #include "video_core/gpu.h"
+#include "video_core/gpu_split.h"
 #include "video_core/memory_manager.h"
 #include "video_core/nx2.h"
 #include "video_core/rasterizer_interface.h"
@@ -42,6 +44,21 @@ u32 ReadSemaphoreWord(DmaPusher& dma_pusher, GPUVAddr address) {
     return dma_pusher.memory_manager.Read<u32>(address);
 }
 
+// With the GPU thread split (gpu_split.h) the rasterizer belongs to the back thread: work that
+// only feeds it runs there in order; reads of what it wrote first wait for it (DrainBack).
+template <typename Func>
+void OnBack(const Func& func) {
+    if (auto& split = VideoCommon::GpuSplit::Get(); split.Active()) {
+        split.Defer(func);
+    } else {
+        func();
+    }
+}
+
+void DrainBack() {
+    VideoCommon::GpuSplit::Get().Drain();
+}
+
 } // Anonymous namespace
 
 bool Puller::PendingAcquire::Matches(u32 current) const {
@@ -61,6 +78,9 @@ bool Puller::CanResumeAcquire(DmaPusher& dma_pusher) {
     if (!pending_acquire) {
         return true;
     }
+    // The release may still be queued for the split GPU thread's back half (and the caches
+    // ReadSemaphoreWord asks are its state).
+    DrainBack();
     const u32 current = ReadSemaphoreWord(dma_pusher, pending_acquire->address);
     if (Nx2SemaphoreWaiters::Get().Latched(this)) {
         if (static std::atomic<u32> rescued{};
@@ -160,10 +180,12 @@ void Puller::ProcessFenceActionMethod(DmaPusher& dma_pusher) {
     case Puller::FenceOperation::Acquire:
         // UNIMPLEMENTED_MSG("Channel Scheduling pending.");
         // WaitFence(regs.fence_action.syncpoint_id, regs.fence_value);
-        dma_pusher.rasterizer->ReleaseFences();
+        OnBack([rasterizer = dma_pusher.rasterizer] { rasterizer->ReleaseFences(); });
         break;
     case Puller::FenceOperation::Increment:
-        dma_pusher.rasterizer->SignalSyncPoint(regs.fence_action.syncpoint_id);
+        OnBack([rasterizer = dma_pusher.rasterizer, id = static_cast<u32>(regs.fence_action.syncpoint_id.Value())] {
+            rasterizer->SignalSyncPoint(id);
+        });
         break;
     default:
         UNIMPLEMENTED_MSG("Unimplemented operation {}", regs.fence_action.op.Value());
@@ -195,12 +217,16 @@ void Puller::ProcessSemaphoreTriggerMethod(DmaPusher& dma_pusher) {
         const auto properties = g_nx2_ampere && (regs.semaphore_trigger & (1U << 24)) != 0
                                     ? VideoCommon::QueryPropertiesFlags::IsAFence
                                     : VideoCommon::QueryPropertiesFlags::HasTimeout;
-        dma_pusher.rasterizer->Query(sequence_address, VideoCommon::QueryType::Payload, properties, payload, 0);
-        if (const GPUVAddr mirror = g_nx2_ampere ? Nx2SemaphoreMirror(sequence_address) : 0;
-            mirror != 0) {
+        const GPUVAddr mirror = g_nx2_ampere ? Nx2SemaphoreMirror(sequence_address) : 0;
+        if (mirror != 0) {
             LOG_DEBUG(HW_GPU, "Switch 2 semaphore {:#x}={:#x} mirrored to {:#x}", sequence_address, payload, mirror);
-            dma_pusher.rasterizer->Query(mirror, VideoCommon::QueryType::Payload, properties, payload, 0);
         }
+        OnBack([rasterizer = dma_pusher.rasterizer, sequence_address, mirror, properties, payload] {
+            rasterizer->Query(sequence_address, VideoCommon::QueryType::Payload, properties, payload, 0);
+            if (mirror != 0) {
+                rasterizer->Query(mirror, VideoCommon::QueryType::Payload, properties, payload, 0);
+            }
+        });
     } else if (g_nx2_ampere) {
         // Switch 2 waits on semaphores the CPU (or another channel) releases, e.g. once it has
         // written the commands this channel runs next: suspend the channel until the value is
@@ -217,6 +243,8 @@ void Puller::ProcessSemaphoreTriggerMethod(DmaPusher& dma_pusher) {
             LOG_ERROR(HW_GPU, "Invalid semaphore operation {}", static_cast<u32>(op));
             return;
         }
+        VideoCommon::Nx2Profile::Add(VideoCommon::Nx2Profile::Counter::AcquireChecks, 1);
+        DrainBack();
         if (acquire.Matches(ReadSemaphoreWord(dma_pusher, acquire.address))) {
             return;
         }
@@ -229,10 +257,12 @@ void Puller::ProcessSemaphoreTriggerMethod(DmaPusher& dma_pusher) {
         dma_pusher.rasterizer->ReleaseFences();
         if (!waiters.Latched(this) && !acquire.Matches(ReadSemaphoreWord(dma_pusher, acquire.address))) {
             pending_acquire = acquire;
+            VideoCommon::Nx2Profile::Add(VideoCommon::Nx2Profile::Counter::AcquireWaits, 1);
         } else {
             waiters.Remove(this);
         }
     } else {
+        DrainBack();
         do {
             const u32 word = dma_pusher.memory_manager.Read<u32>(regs.semaphore_address.SemaphoreAddress());
             regs.acquire_source = true;
@@ -266,10 +296,14 @@ void Puller::ProcessSemaphoreTriggerMethod(DmaPusher& dma_pusher) {
 void Puller::ProcessSemaphoreRelease(DmaPusher& dma_pusher) {
     const GPUVAddr sequence_address{regs.semaphore_address.SemaphoreAddress()};
     const u32 payload = regs.semaphore_release;
-    dma_pusher.rasterizer->Query(sequence_address, VideoCommon::QueryType::Payload, VideoCommon::QueryPropertiesFlags::IsAFence, payload, 0);
+    OnBack([rasterizer = dma_pusher.rasterizer, sequence_address, payload] {
+        rasterizer->Query(sequence_address, VideoCommon::QueryType::Payload,
+                          VideoCommon::QueryPropertiesFlags::IsAFence, payload, 0);
+    });
 }
 
 void Puller::ProcessSemaphoreAcquire(DmaPusher& dma_pusher) {
+    DrainBack();
     u32 word = dma_pusher.memory_manager.Read<u32>(regs.semaphore_address.SemaphoreAddress());
     const auto value = regs.semaphore_acquire;
     while (word != value) {
@@ -291,11 +325,14 @@ void Puller::ProcessSemaphoreExecute(DmaPusher& dma_pusher, u32 argument) {
     const u32 payload = regs.reg_array[0x19];
     const u32 operation = argument & 7;
     if (operation == 1) { // RELEASE
-        dma_pusher.rasterizer->Query(address, VideoCommon::QueryType::Payload,
-                                     VideoCommon::QueryPropertiesFlags::IsAFence, payload, 0);
+        OnBack([rasterizer = dma_pusher.rasterizer, address, payload] {
+            rasterizer->Query(address, VideoCommon::QueryType::Payload,
+                              VideoCommon::QueryPropertiesFlags::IsAFence, payload, 0);
+        });
         return;
     }
     // Acquires: check once after flushing pending fences; never block the GPU thread on them.
+    DrainBack();
     dma_pusher.rasterizer->ReleaseFences();
     const u32 word = dma_pusher.memory_manager.Read<u32>(address);
     const bool ready = operation == 0   ? word == payload
@@ -336,13 +373,13 @@ void Puller::CallPullerMethod(DmaPusher& dma_pusher, const MethodCall& method_ca
     case BufferMethods::WrcacheFlush:
         break;
     case BufferMethods::RefCnt:
-        dma_pusher.rasterizer->SignalReference();
+        OnBack([rasterizer = dma_pusher.rasterizer] { rasterizer->SignalReference(); });
         break;
     case BufferMethods::SyncpointOperation:
         ProcessFenceActionMethod(dma_pusher);
         break;
     case BufferMethods::WaitForIdle:
-        dma_pusher.rasterizer->WaitForIdle();
+        OnBack([rasterizer = dma_pusher.rasterizer] { rasterizer->WaitForIdle(); });
         break;
     case BufferMethods::SemaphoreOperation: {
         ProcessSemaphoreTriggerMethod(dma_pusher);
@@ -358,7 +395,7 @@ void Puller::CallPullerMethod(DmaPusher& dma_pusher, const MethodCall& method_ca
     }
     case BufferMethods::MemOpB: {
         // Implement this better.
-        dma_pusher.rasterizer->InvalidateGPUCache();
+        OnBack([rasterizer = dma_pusher.rasterizer] { rasterizer->InvalidateGPUCache(); });
         break;
     }
     case BufferMethods::MemOpC:
diff --git a/src/video_core/fence_manager.h b/src/video_core/fence_manager.h
index eb72529da0..61c9844cf2 100644
--- a/src/video_core/fence_manager.h
+++ b/src/video_core/fence_manager.h
@@ -8,11 +8,13 @@
 
 #include <algorithm>
 #include <condition_variable>
+#include <cstdlib>
 #include <cstring>
 #include <deque>
 #include <functional>
 #include <mutex>
 #include <thread>
+#include <vector>
 #include <queue>
 
 #include "common/common_types.h"
@@ -22,6 +24,7 @@
 #include "video_core/gpu.h"
 #include "video_core/host1x/host1x.h"
 #include "video_core/host1x/syncpoint_manager.h"
+#include "video_core/nx2_profile.h"
 #include "video_core/rasterizer_interface.h"
 
 namespace VideoCommon {
@@ -76,12 +79,21 @@ public:
     }
 
     void SignalFence(std::function<void()>&& func) {
+        Nx2Profile::Scope profile_scope{Nx2Profile::Zone::Fence};
         const bool delay_fence = Settings::IsGPUFenceBehaviorDefault() ? Settings::IsGPULevelHigh() : Settings::IsGPUFenceBehaviorBalanced() || Settings::IsGPUFenceBehaviorAccurate();
-        const bool should_flush = ShouldFlush();
         if constexpr (!can_async_check) {
             TryReleasePendingFences<false>();
         }
-        const u8 committed_flushes = CommitAsyncFlushes();
+        auto [should_flush, committed_flushes] = CommitCacheFlushes();
+        if (!delay_fence && committed_flushes == 0 && SkipEmptyFence()) {
+            Nx2Profile::Scope skipped_scope{Nx2Profile::Zone::FenceSkipped};
+            func();
+            rasterizer.InvalidateGPUCache();
+            return;
+        }
+        skipped_fences = 0;
+        should_flush |= query_cache.HasUncommittedFlushes();
+        query_cache.CommitAsyncFlushes();
         TFence new_fence = CreateFence(!should_flush);
         if constexpr (can_async_check) {
             guard.lock();
@@ -207,9 +219,29 @@ private:
         Common::SetCurrentThreadPriority(Common::ThreadPriority::High);
         Common::SetCurrentThreadToPerformanceCores();
 
-        TFence current_fence;
-        std::deque<std::function<void()>> current_operations;
-        u8 current_flushes{};
+        // Every fence the host GPU has already passed is released together: Switch 2 titles
+        // signal thousands a second and the host GPU completes them in bursts, so each cache's
+        // lock is taken once per burst instead of once per fence (the GPU thread, which holds
+        // them for nearly every command, spun, slept and woke this thread for each one).
+        // EDEN_NX2_FENCE_BATCH=0: one fence at a time (for comparisons).
+        static const size_t MaxBatch = [] {
+            const char* value = std::getenv("EDEN_NX2_FENCE_BATCH");
+            return value != nullptr && value[0] == '0' ? size_t{1} : size_t{64};
+        }();
+        struct Released {
+            TFence fence;
+            std::deque<std::function<void()>> operations;
+            u8 flushes;
+        };
+        std::vector<Released> batch;
+        batch.reserve(MaxBatch);
+        const auto take_front = [&] {
+            batch.push_back({std::move(fences.front()), std::move(pending_operations.front()),
+                             pending_flushes.front()});
+            fences.pop();
+            pending_operations.pop_front();
+            pending_flushes.pop_front();
+        };
         while (!stop_token.stop_requested()) {
             {
                 std::unique_lock lock(guard);
@@ -217,24 +249,50 @@ private:
                 if (stop_token.stop_requested()) [[unlikely]] {
                     return;
                 }
-                current_fence = std::move(fences.front());
-                current_operations = std::move(pending_operations.front());
-                current_flushes = pending_flushes.front();
-                fences.pop();
-                pending_operations.pop_front();
-                pending_flushes.pop_front();
+                take_front();
             }
-            if (!current_fence->IsStubbed()) {
-                WaitFence(current_fence);
+            if (!batch.front().fence->IsStubbed()) {
+                WaitFence(batch.front().fence);
             }
-            PopAsyncFlushes(current_flushes);
-            for (auto& operation : current_operations) {
-                operation();
+            {
+                std::unique_lock lock(guard);
+                while (batch.size() < MaxBatch && !fences.empty() &&
+                       (fences.front()->IsStubbed() || IsFenceSignaled(fences.front()))) {
+                    take_front();
+                }
+            }
+            size_t texture_pops = 0;
+            size_t buffer_pops = 0;
+            for (const Released& released : batch) {
+                texture_pops += (released.flushes & TextureFlushes) != 0 ? 1 : 0;
+                buffer_pops += (released.flushes & BufferFlushes) != 0 ? 1 : 0;
+            }
+            if (texture_pops != 0) {
+                std::scoped_lock lock{texture_cache.mutex};
+                for (size_t i = 0; i < texture_pops; ++i) {
+                    texture_cache.PopAsyncFlushes();
+                }
+            }
+            if (buffer_pops != 0) {
+                std::scoped_lock lock{buffer_cache.mutex};
+                for (size_t i = 0; i < buffer_pops; ++i) {
+                    buffer_cache.PopAsyncFlushes();
+                }
+            }
+            // Queries and operations keep their per-fence order.
+            for (Released& released : batch) {
+                query_cache.PopAsyncFlushes();
+                for (auto& operation : released.operations) {
+                    operation();
+                }
             }
             {
                 std::unique_lock lock(ring_guard);
-                delayed_destruction_ring.Push(std::move(current_fence));
+                for (Released& released : batch) {
+                    delayed_destruction_ring.Push(std::move(released.fence));
+                }
             }
+            batch.clear();
         }
     }
 
@@ -244,12 +302,6 @@ private:
                query_cache.ShouldWaitAsyncFlushes();
     }
 
-    bool ShouldFlush() const {
-        std::scoped_lock lock{buffer_cache.mutex, texture_cache.mutex};
-        return texture_cache.HasUncommittedFlushes() || buffer_cache.HasUncommittedFlushes() ||
-               query_cache.HasUncommittedFlushes();
-    }
-
     static constexpr u8 TextureFlushes = 1;
     static constexpr u8 BufferFlushes = 2;
 
@@ -266,35 +318,66 @@ private:
         query_cache.PopAsyncFlushes();
     }
 
-    /// Returns which caches committed downloads for this fence (TextureFlushes, BufferFlushes).
-    /// Switch 2 titles signal a fence per compute dispatch (thousands a second), mostly with
-    /// nothing or only buffer data to download: the fencing thread then takes only the lock of
-    /// the cache with work, instead of holding both while it copies (the GPU thread's draws and
-    /// uploads waited for it).
-    u8 CommitAsyncFlushes() {
-        u8 committed = 0;
+    struct CacheFlushes {
+        bool should_flush;  ///< Downloads to wait for: the fence must reach the host GPU
+        u8 committed;       ///< TextureFlushes, BufferFlushes: the caches that committed downloads
+    };
+
+    /// Commits the texture and buffer caches' downloads for a new fence (the query cache's are
+    /// committed once a fence is made). Switch 2 titles signal a fence per compute dispatch
+    /// (thousands a second), mostly with nothing or only buffer data to download: each cache's
+    /// lock is taken once, alone, and the fencing thread later takes only the lock of the cache
+    /// with work, instead of holding both while it copies (the GPU thread's draws and uploads
+    /// waited for it).
+    CacheFlushes CommitCacheFlushes() {
+        CacheFlushes flushes{};
         {
             std::scoped_lock lock{texture_cache.mutex};
+            flushes.should_flush = texture_cache.HasUncommittedFlushes();
             if (texture_cache.HasAsyncFlushWork()) {
                 texture_cache.CommitAsyncFlushes();
-                committed |= TextureFlushes;
+                flushes.committed |= TextureFlushes;
             }
         }
         {
             std::scoped_lock lock{buffer_cache.mutex};
             if (buffer_cache.HasUncommittedFlushes()) {
                 buffer_cache.CommitAsyncFlushes();
-                committed |= BufferFlushes;
+                flushes.should_flush = true;
+                flushes.committed |= BufferFlushes;
             }
         }
-        query_cache.CommitAsyncFlushes();
-        return committed;
+        return flushes;
+    }
+
+    /// Whether a fence with no downloads would do nothing, so need not be made: with fast GPU
+    /// fences its callback runs at once anyway, and most of the per-dispatch fences of Switch 2
+    /// titles have no queries or queued operations either. Making, queueing and retiring them
+    /// cost the GPU thread ~5% and woke the fencing thread for each. One in 64 is still made, so
+    /// the operations fences carry along (query slot cleanup) keep running.
+    bool SkipEmptyFence() {
+        if constexpr (requires { query_cache.HasPendingFenceWork(); }) {
+            if (skipped_fences >= 64 || query_cache.HasPendingFenceWork()) {
+                return false;
+            }
+            {
+                std::scoped_lock lock{uncommitted_guard};
+                if (!uncommitted_operations.empty()) {
+                    return false;
+                }
+            }
+            ++skipped_fences;
+            return true;
+        } else {
+            return false;
+        }
     }
 
     std::queue<TFence> fences;
     std::deque<std::function<void()>> uncommitted_operations;
     std::deque<std::deque<std::function<void()>>> pending_operations;
     std::deque<u8> pending_flushes; ///< Per fence: the caches that committed downloads
+    u32 skipped_fences{}; ///< Empty fences skipped since the last one made (GPU thread only)
 
     std::mutex guard;
     std::mutex ring_guard;
diff --git a/src/video_core/gpu.cpp b/src/video_core/gpu.cpp
index 6441b9511a..53f8f063f7 100644
--- a/src/video_core/gpu.cpp
+++ b/src/video_core/gpu.cpp
@@ -221,6 +221,13 @@ struct GPU::Impl {
             return raster_area;
         }
         raster_area.preemtive = true;
+        // A guest CPU read of GPU-written memory (reactive flushing): count them, they decide
+        // whether buffer data must be copied back at every fence (EDEN_NX2_BUFFER_DOWNLOADS).
+        if (static std::atomic<u64> reads{}; (reads.fetch_add(1, std::memory_order_relaxed) & 63) == 0) {
+            LOG_INFO(HW_GPU, "CPU read of GPU-written memory [{:#x}, {:#x}) ({} so far)",
+                     raster_area.start_address, raster_area.end_address,
+                     reads.load(std::memory_order_relaxed));
+        }
         const u64 fence = RequestSyncOperation([this, &raster_area]() {
             renderer->ReadRasterizer()->FlushRegion(raster_area.start_address, raster_area.end_address - raster_area.start_address);
         });
diff --git a/src/video_core/gpu_split.cpp b/src/video_core/gpu_split.cpp
new file mode 100644
index 0000000000..f913ca82ce
--- /dev/null
+++ b/src/video_core/gpu_split.cpp
@@ -0,0 +1,350 @@
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
+#include <algorithm>
+#include <cstdlib>
+#include <thread>
+
+#ifdef _MSC_VER
+#include <intrin.h>
+#endif
+#ifdef _WIN32
+#ifndef NOMINMAX
+#define NOMINMAX
+#endif
+#include <windows.h>
+#endif
+#include <cstdio>
+
+#include "common/logging.h"
+#include "common/settings.h"
+#include "common/thread.h"
+#include "core/core.h"
+#include "video_core/engines/engine_interface.h"
+#include "video_core/gpu_split.h"
+#include "video_core/nx2.h"
+#include "video_core/nx2_profile.h"
+
+namespace VideoCommon {
+
+namespace {
+
+void Pause() {
+#if defined(_MSC_VER) && (defined(_M_X64) || defined(_M_IX86))
+    _mm_pause();
+#elif defined(__x86_64__) || defined(__i386__)
+    __builtin_ia32_pause();
+#else
+    std::this_thread::yield();
+#endif
+}
+
+std::atomic<u32> g_back_thread_id{};
+std::atomic<bool> g_back_busy{};
+
+} // Anonymous namespace
+
+void GpuSplit::CheckOwner(const char* where) {
+#ifdef _WIN32
+    if (!g_back_busy.load(std::memory_order_acquire)) {
+        return;
+    }
+    const DWORD thread = GetCurrentThreadId();
+    if (thread == g_back_thread_id.load(std::memory_order_relaxed)) {
+        return;
+    }
+    static std::atomic<u32> reports{};
+    if (reports.fetch_add(1) >= 12) {
+        return;
+    }
+    std::fprintf(stderr, "GPU split owner check: %s on thread %lu while the back runs\n", where,
+                 thread);
+    void* frames[40];
+    const USHORT count = RtlCaptureStackBackTrace(1, 40, frames, nullptr);
+    for (USHORT i = 0; i < count; ++i) {
+        HMODULE module{};
+        char name[MAX_PATH] = "?";
+        if (GetModuleHandleExA(GET_MODULE_HANDLE_EX_FLAG_FROM_ADDRESS |
+                                   GET_MODULE_HANDLE_EX_FLAG_UNCHANGED_REFCOUNT,
+                               static_cast<LPCSTR>(frames[i]), &module)) {
+            GetModuleFileNameA(module, name, sizeof(name));
+        }
+        const char* base_name = std::strrchr(name, '\\');
+        std::fprintf(stderr, "  frame %s+%#llx\n", base_name ? base_name + 1 : name,
+                     static_cast<unsigned long long>(reinterpret_cast<uintptr_t>(frames[i]) -
+                                                     reinterpret_cast<uintptr_t>(module)));
+    }
+    std::fflush(stderr);
+#else
+    (void)where;
+#endif
+}
+
+GpuSplit GpuSplit::instance;
+
+void GpuSplit::MaybeStart(Core::System& system_) {
+    static const char mode = [] {
+        const char* value = std::getenv("EDEN_NX2_SPLIT");
+        return value != nullptr ? value[0] : '2';
+    }();
+    const bool wanted = mode == '1' || mode == '2';
+    if (active || !wanted || !Tegra::g_nx2_ampere.load(std::memory_order_relaxed)) {
+        return;
+    }
+    // The front reads command lists straight from guest memory; a read that flushes the caches
+    // first (safe DMA, accurate GPU) would touch back-owned state.
+    const bool safe_reads = Settings::IsDMALevelDefault() ? Settings::IsGPULevelHigh()
+                                                          : Settings::IsDMALevelSafe();
+    if (safe_reads || Settings::IsGPULevelHigh() ||
+        Settings::values.renderer_backend.GetValue() != Settings::RendererBackend::Vulkan) {
+        return;
+    }
+    system = &system_;
+    check_owner = std::getenv("EDEN_NX2_SPLIT_CHECK") != nullptr;
+    ring = std::make_unique<u64[]>(RingWords);
+    back_thread = std::jthread([this](std::stop_token stop_token) { BackLoop(stop_token); });
+    active = true;
+#ifdef ARCHITECTURE_x86_64
+    // Front macros run compiled (macro_turing_jit.h) only.
+    const char* jit = std::getenv("EDEN_NX2_MME_JIT");
+    front_macros = mode == '2' && (jit == nullptr || jit[0] != '0');
+#endif
+    LOG_INFO(HW_GPU, "GPU thread split: engine methods run on a second thread{}",
+             front_macros ? ", macros on the first" : "");
+}
+
+void GpuSplit::Stop() {
+    if (!back_thread.joinable()) {
+        return;
+    }
+    back_thread.request_stop();
+    wake.fetch_add(1, std::memory_order_seq_cst);
+    wake.notify_one();
+    back_thread.join();
+    active = false;
+    // The split is static and outlives the emulation session. The next session's back thread
+    // starts reading at 0: with the old positions kept it skipped through the new, empty ring a
+    // wrap at a time and never reached the old write position (the next game's first GPU work
+    // never ran, it waited on its syncpoint forever), and the engines it held were freed.
+    ring.reset();
+    write_position = 0;
+    published_local = 0;
+    consumed_cache = 0;
+    publish_every = PublishEvery;
+    front_engine = nullptr;
+    back_engine = nullptr;
+    front_macros = false;
+    published.store(0, std::memory_order_relaxed);
+    consumed.store(0, std::memory_order_relaxed);
+    back_waiting.store(false, std::memory_order_relaxed);
+}
+
+u64* GpuSplit::ReserveSlow(size_t words) {
+    size_t index = write_position & RingMask;
+    if (index + words > RingWords) {
+        // Entries never wrap: skip to the start of the ring.
+        const size_t rest = RingWords - index;
+        WaitForSpace(rest);
+        ring[index] = static_cast<u64>(Kind::Wrap);
+        write_position += rest;
+        index = 0;
+    }
+    WaitForSpace(words);
+    return &ring[index];
+}
+
+void GpuSplit::WaitForSpace(size_t words) {
+    if (write_position + words - consumed_cache <= RingWords) {
+        return;
+    }
+    Publish();
+    Nx2Profile::Scope scope{Nx2Profile::Zone::SplitDrain};
+    for (u32 spins = 0;; ++spins) {
+        consumed_cache = consumed.load(std::memory_order_acquire);
+        if (write_position + words - consumed_cache <= RingWords) {
+            return;
+        }
+        if (spins < 4096) {
+            Pause();
+        } else {
+            std::this_thread::yield();
+        }
+    }
+}
+
+void GpuSplit::Publish() {
+    if (published_local == write_position) {
+        return;
+    }
+    published_local = write_position;
+    publish_every = std::min<size_t>(publish_every * 2, PublishEvery);
+    published.store(write_position, std::memory_order_seq_cst);
+    if (back_waiting.load(std::memory_order_seq_cst)) {
+        wake.fetch_add(1, std::memory_order_seq_cst);
+        wake.notify_one();
+    }
+}
+
+void GpuSplit::Drain() {
+    consumed_cache = consumed.load(std::memory_order_acquire);
+    if (consumed_cache == write_position) {
+        return;
+    }
+    Publish();
+    Nx2Profile::Scope scope{Nx2Profile::Zone::SplitDrain};
+    for (u32 spins = 0;; ++spins) {
+        consumed_cache = consumed.load(std::memory_order_acquire);
+        if (consumed_cache == write_position) {
+            publish_every = 8;
+            return;
+        }
+        if (spins < 16384) {
+            Pause();
+        } else {
+            std::this_thread::yield();
+        }
+    }
+}
+
+void GpuSplit::SwitchEngine(Tegra::Engines::EngineInterface* engine) {
+    front_engine = engine;
+    u64* const entry = Reserve(2);
+    entry[0] = static_cast<u64>(Kind::Engine);
+    entry[1] = reinterpret_cast<u64>(engine);
+    Commit(2);
+}
+
+void GpuSplit::MultiMethod(Tegra::Engines::EngineInterface* engine, u32 method, const u32* words,
+                           u32 amount, u32 methods_pending, const GPUVAddr* segment) {
+    SelectEngine(engine);
+    // Long runs (constant buffer and inline uploads) go in pieces, as when a method's data
+    // continues in the next GPFIFO entry.
+    while (amount > 0) {
+        const u32 piece = std::min(amount, MaxMultiWords);
+        const size_t head = segment != nullptr ? 3 : 2;
+        const size_t size = head + (piece + 1) / 2;
+        u64* const entry = Reserve(size);
+        entry[0] = static_cast<u64>(Kind::Multi) | (segment != nullptr ? HasSegment : 0) |
+                   (u64{method & 0xffff} << 8) | (u64{piece} << 32);
+        entry[1] = methods_pending;
+        if (segment != nullptr) {
+            entry[2] = *segment;
+        }
+        entry[size - 1] = 0;
+        std::memcpy(entry + head, words, size_t{piece} * sizeof(u32));
+        Commit(size);
+        words += piece;
+        amount -= piece;
+        methods_pending -= piece;
+    }
+}
+
+void GpuSplit::DeferFunction(std::function<void()>&& func) {
+    auto* const heap = new std::function<void()>(std::move(func));
+    Defer([heap] {
+        (*heap)();
+        delete heap;
+    });
+}
+
+size_t GpuSplit::Execute(const u64* entry) {
+    const u64 head = entry[0];
+    switch (static_cast<Kind>(head & 0xf)) {
+    case Kind::Wrap:
+        return 0;
+    case Kind::Engine:
+        back_engine = reinterpret_cast<Tegra::Engines::EngineInterface*>(entry[1]);
+        return 2;
+    case Kind::Method: {
+        const u32 method = static_cast<u32>((head >> 8) & 0xffff);
+        const u32 argument = static_cast<u32>(head >> 32);
+        auto* const engine = back_engine;
+        if ((head & Direct) != 0) {
+            Nx2Profile::Scope method_scope{Nx2Profile::Zone::Method};
+            engine->ConsumeSink(*system);
+            engine->CallMethod(*system, method, argument, true);
+        } else if (!engine->execution_mask[method]) {
+            engine->method_sink.emplace_back(method, argument);
+        } else {
+            Nx2Profile::Scope method_scope{Nx2Profile::Zone::Method};
+            engine->ConsumeSink(*system);
+            if ((head & HasSegment) != 0) {
+                engine->current_dma_segment = entry[1];
+            }
+            engine->CallMethod(*system, method, argument, (head & IsLast) != 0);
+        }
+        return (head & HasSegment) != 0 ? 2 : 1;
+    }
+    case Kind::Multi: {
+        const u32 method = static_cast<u32>((head >> 8) & 0xffff);
+        const u32 amount = static_cast<u32>(head >> 32);
+        const size_t data = (head & HasSegment) != 0 ? 3 : 2;
+        auto* const engine = back_engine;
+        Nx2Profile::Scope method_scope{Nx2Profile::Zone::Method};
+        engine->ConsumeSink(*system);
+        if ((head & HasSegment) != 0) {
+            engine->current_dma_segment = entry[2];
+        }
+        engine->CallMultiMethod(*system, method, reinterpret_cast<const u32*>(entry + data),
+                                amount, static_cast<u32>(entry[1]));
+        return data + (amount + 1) / 2;
+    }
+    case Kind::Call: {
+        const size_t payload = static_cast<size_t>((head >> 8) & 0xffffff);
+        reinterpret_cast<void (*)(const u64*)>(entry[1])(entry + 2);
+        return 2 + payload;
+    }
+    }
+    return 1;
+}
+
+void GpuSplit::BackLoop(std::stop_token stop_token) {
+    Common::SetCurrentThreadName("GPUBack");
+    Common::SetCurrentThreadPriority(Common::ThreadPriority::Critical);
+    Common::SetCurrentThreadToPerformanceCores();
+    system->RegisterHostThread();
+#ifdef _WIN32
+    g_back_thread_id.store(GetCurrentThreadId(), std::memory_order_relaxed);
+#endif
+    u64 read_position = 0;
+    while (!stop_token.stop_requested()) {
+        Nx2Profile::Report("GPU back thread");
+        u64 available = published.load(std::memory_order_acquire);
+        if (available == read_position) {
+            Nx2Profile::Scope scope{Nx2Profile::Zone::SplitWait};
+            // The front usually sends more within microseconds: spin briefly before sleeping.
+            for (u32 spins = 0; spins < 8192 && available == read_position; ++spins) {
+                Pause();
+                available = published.load(std::memory_order_acquire);
+            }
+            if (available == read_position) {
+                // Publish bumps `wake` when it sees back_waiting: read it before the last look
+                // at `published`, so a publish in between ends the wait at once.
+                back_waiting.store(true, std::memory_order_seq_cst);
+                const u32 seen = wake.load(std::memory_order_seq_cst);
+                available = published.load(std::memory_order_seq_cst);
+                if (available == read_position && !stop_token.stop_requested()) {
+                    wake.wait(seen, std::memory_order_seq_cst);
+                }
+                back_waiting.store(false, std::memory_order_relaxed);
+                continue;
+            }
+        }
+        u64 last_store = read_position;
+        g_back_busy.store(true, std::memory_order_relaxed);
+        while (read_position != available) {
+            const size_t index = read_position & RingMask;
+            const size_t size = Execute(&ring[index]);
+            read_position += size != 0 ? size : RingWords - index;
+            // Hand ring space back now and then while working through a long batch.
+            if (read_position - last_store >= 4096 && read_position != available) {
+                consumed.store(read_position, std::memory_order_release);
+                last_store = read_position;
+            }
+        }
+        g_back_busy.store(false, std::memory_order_release);
+        consumed.store(read_position, std::memory_order_release);
+    }
+}
+
+} // namespace VideoCommon
diff --git a/src/video_core/gpu_split.h b/src/video_core/gpu_split.h
new file mode 100644
index 0000000000..d6f91e4141
--- /dev/null
+++ b/src/video_core/gpu_split.h
@@ -0,0 +1,177 @@
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
+#pragma once
+
+#include <atomic>
+#include <cstring>
+#include <functional>
+#include <memory>
+#include <thread>
+#include <type_traits>
+
+#include "common/common_types.h"
+
+namespace Core {
+class System;
+}
+
+namespace Tegra::Engines {
+class EngineInterface;
+}
+
+namespace VideoCommon {
+
+// Splits the GPU thread in two (NX2_GPU_THREAD_SPLIT.md). The GPU thread (the front) takes the
+// CPU's commands, parses command lists and runs the puller; engine method calls and everything
+// behind them (register state, macros, draws, the rasterizer and its caches) run in order on a
+// second thread (the back), fed through a single-producer single-consumer ring. Work the front
+// would do on back-owned state is deferred into the same ring; the front waits for the back to
+// catch up (Drain) only where it needs a result from that state, e.g. a semaphore it acquires.
+// On for Switch 2 titles with the 3D and compute engines' macros on the front, against copies of
+// the registers they read (Maxwell3D and KeplerCompute Front*), which sends the writes they make.
+// EDEN_NX2_SPLIT=1 keeps the macros on the back; EDEN_NX2_SPLIT=0 runs everything on one thread.
+class GpuSplit {
+public:
+    static GpuSplit& Get() {
+        return instance;
+    }
+
+    /// Whether the front sends work to the back (front thread only).
+    [[nodiscard]] bool Active() const {
+        return active;
+    }
+    /// Whether the front runs the engines' macros (EDEN_NX2_SPLIT=2).
+    [[nodiscard]] bool FrontMacros() const {
+        return front_macros;
+    }
+
+    /// Starts the back thread once the split is wanted (EDEN_NX2_SPLIT=1, a Switch 2 title,
+    /// commands read without cache flushes). Called by the front between commands.
+    void MaybeStart(Core::System& system);
+    void Stop();
+
+    /// An engine method write, executed as DmaPusher would (method sink, ConsumeSink, call).
+    /// `segment`: the engine's current_dma_segment for the call, when the engine reads it.
+    void Method(Tegra::Engines::EngineInterface* engine, u32 method, u32 argument, bool is_last,
+                const GPUVAddr* segment) {
+        SelectEngine(engine);
+        const size_t words = segment != nullptr ? 2 : 1;
+        u64* const entry = Reserve(words);
+        entry[0] = static_cast<u64>(Kind::Method) | (is_last ? IsLast : 0) |
+                   (segment != nullptr ? HasSegment : 0) | (u64{method & 0xffff} << 8) |
+                   (u64{argument} << 32);
+        if (segment != nullptr) {
+            entry[1] = *segment;
+        }
+        Commit(words);
+    }
+    void MultiMethod(Tegra::Engines::EngineInterface* engine, u32 method, const u32* words,
+                     u32 amount, u32 methods_pending, const GPUVAddr* segment);
+    /// A write made by a macro: CallMethod without the method sink, as the engines' macros do.
+    void MacroWrite(Tegra::Engines::EngineInterface* engine, u32 method, u32 argument) {
+        SelectEngine(engine);
+        u64* const entry = Reserve(1);
+        entry[0] = static_cast<u64>(Kind::Method) | IsLast | Direct | (u64{method & 0xffff} << 8) |
+                   (u64{argument} << 32);
+        Commit(1);
+    }
+
+    /// Runs `func` on the back after everything sent before it. Must be trivially copyable.
+    template <typename Func>
+    void Defer(const Func& func) {
+        static_assert(std::is_trivially_copyable_v<Func> && alignof(Func) <= alignof(u64));
+        constexpr size_t payload = (sizeof(Func) + sizeof(u64) - 1) / sizeof(u64);
+        u64* const entry = Reserve(2 + payload);
+        entry[0] = static_cast<u64>(Kind::Call) | (u64{payload} << 8);
+        entry[1] = reinterpret_cast<u64>(&CallThunk<Func>);
+        std::memcpy(entry + 2, &func, sizeof(Func));
+        Commit(2 + payload);
+    }
+    void DeferFunction(std::function<void()>&& func);
+
+    /// EDEN_NX2_SPLIT_CHECK=1 (debug): reports, with its stack, a thread other than the back
+    /// using back-owned state (the Vulkan scheduler) while the back runs entries.
+    static inline bool check_owner{};
+    static void CheckOwner(const char* where);
+
+    /// Makes what was sent visible to the back.
+    void Publish();
+    /// Waits until the back executed everything sent so far.
+    void Drain();
+
+private:
+    enum class Kind : u64 {
+        Wrap,   ///< The rest of the ring is unused: continue at its start
+        Engine, ///< The engine the next methods go to: word 1
+        Method, ///< Bits 8-23 method, 32-63 argument; flags IsLast, HasSegment (word 1)
+        Multi,  ///< Bits 8-23 method, 32-63 amount; word 1 methods pending, [segment], words
+        Call,   ///< Bits 8-31 payload size in words; word 1 thunk, then the payload
+    };
+    static constexpr u64 IsLast = 1 << 4;
+    static constexpr u64 HasSegment = 1 << 5;
+    static constexpr u64 Direct = 1 << 6; ///< Method: call it even if not executable (macros)
+
+    static constexpr size_t RingWords = size_t{1} << 20; // 8 MiB
+    static constexpr size_t RingMask = RingWords - 1;
+    static constexpr u32 MaxMultiWords = 8192;
+    static constexpr size_t PublishEvery = 256;
+
+    template <typename Func>
+    static void CallThunk(const u64* payload) {
+        (*reinterpret_cast<const Func*>(payload))();
+    }
+
+    /// Room for an entry of `words` at the write position (never across the ring's end).
+    u64* Reserve(size_t words) {
+        const size_t index = write_position & RingMask;
+        if (index + words <= RingWords && write_position + words - consumed_cache <= RingWords)
+            [[likely]] {
+            return &ring[index];
+        }
+        return ReserveSlow(words);
+    }
+    u64* ReserveSlow(size_t words);
+    void Commit(size_t words) {
+        write_position += words;
+        if (write_position - published_local >= publish_every) {
+            Publish();
+        }
+    }
+    void WaitForSpace(size_t words);
+    void SelectEngine(Tegra::Engines::EngineInterface* engine) {
+        if (engine != front_engine) [[unlikely]] {
+            SwitchEngine(engine);
+        }
+    }
+    void SwitchEngine(Tegra::Engines::EngineInterface* engine);
+
+    void BackLoop(std::stop_token stop_token);
+    size_t Execute(const u64* entry);
+
+    Core::System* system{};
+    bool active{};
+    bool front_macros{};
+    std::unique_ptr<u64[]> ring;
+
+    // Front side.
+    u64 write_position{};
+    u64 published_local{};
+    u64 consumed_cache{};
+    /// Words between publishes: small right after a drain, while the back has nothing to do.
+    size_t publish_every{PublishEvery};
+    Tegra::Engines::EngineInterface* front_engine{};
+
+    // Back side.
+    Tegra::Engines::EngineInterface* back_engine{};
+
+    static GpuSplit instance;
+
+    alignas(64) std::atomic<u64> published{};
+    alignas(64) std::atomic<u64> consumed{};
+    alignas(64) std::atomic<bool> back_waiting{};
+    std::atomic<u32> wake{};
+    std::jthread back_thread;
+};
+
+} // namespace VideoCommon
diff --git a/src/video_core/gpu_thread.cpp b/src/video_core/gpu_thread.cpp
index 72da1904e1..55200b0647 100644
--- a/src/video_core/gpu_thread.cpp
+++ b/src/video_core/gpu_thread.cpp
@@ -17,12 +17,28 @@
 #include "video_core/dma_pusher.h"
 #include "video_core/gpu.h"
 #include "video_core/gpu_thread.h"
+#include "video_core/gpu_split.h"
 #include "video_core/nx2_profile.h"
 #include "video_core/host1x/host1x.h"
 #include "video_core/renderer_base.h"
 
 namespace VideoCommon::GPUThread {
 
+namespace {
+
+// Tells a CPU thread waiting in PushCommand that its command ran.
+void SignalCommand(SynchState& state, u64 fence, bool block) {
+    state.signaled_fence.store(fence);
+    if (block) {
+        // We have to lock the write_lock to ensure that the condition_variable wait not get a
+        // race between the check and the lock itself.
+        std::scoped_lock lk{state.write_lock};
+        state.cv.notify_all();
+    }
+}
+
+} // Anonymous namespace
+
 ThreadManager::ThreadManager(Core::System& system_)
     : system{system_}
 {}
@@ -42,8 +58,10 @@ void ThreadManager::StartThread(VideoCore::RendererBase& renderer, Core::Fronten
         // When a pending channel last moved or a command arrived (see the polling below).
         auto last_progress = std::chrono::steady_clock::now();
         namespace Profile = VideoCommon::Nx2Profile;
+        auto& split = VideoCommon::GpuSplit::Get();
         while (!stop_token.stop_requested()) {
             Profile::Report();
+            split.MaybeStart(system);
             if (scheduler.HasPendingChannels()) {
                 // Suspended channels wait on semaphores the CPU releases: poll them meanwhile.
                 {
@@ -71,6 +89,34 @@ void ThreadManager::StartThread(VideoCore::RendererBase& renderer, Core::Fronten
                 break;
             }
             last_progress = std::chrono::steady_clock::now();
+            if (split.Active()) {
+                // Only command lists are parsed here; everything else runs on the back thread
+                // in order, and so does the signal that the command ran.
+                auto* const rasterizer_ptr = renderer.ReadRasterizer();
+                if (auto* submit_list = std::get_if<SubmitListCommand>(&next.data)) {
+                    Profile::Scope scope{Profile::Zone::Push};
+                    scheduler.Push(system.GPU(), submit_list->channel, std::move(submit_list->entries));
+                } else if (std::holds_alternative<GPUTickCommand>(next.data)) {
+                    split.Defer([gpu = &system.GPU()] { gpu->TickWork(); });
+                } else if (auto* run = std::get_if<RunCommand>(&next.data)) {
+                    split.DeferFunction(std::move(run->func));
+                } else if (const auto* flush = std::get_if<FlushRegionCommand>(&next.data)) {
+                    split.Defer([rasterizer_ptr, addr = flush->addr, size = flush->size] {
+                        rasterizer_ptr->FlushRegion(addr, size);
+                    });
+                } else if (const auto* invalidate = std::get_if<InvalidateRegionCommand>(&next.data)) {
+                    split.Defer([rasterizer_ptr, addr = invalidate->addr, size = invalidate->size] {
+                        rasterizer_ptr->OnCacheInvalidation(addr, size);
+                    });
+                } else {
+                    ASSERT(false);
+                }
+                split.Defer([state_ptr = &state, fence = next.fence, block = next.block] {
+                    SignalCommand(*state_ptr, fence, block);
+                });
+                split.Publish();
+                continue;
+            }
             if (auto* submit_list = std::get_if<SubmitListCommand>(&next.data)) {
                 Profile::Scope scope{Profile::Zone::Push};
                 scheduler.Push(system.GPU(), submit_list->channel, std::move(submit_list->entries));
@@ -89,13 +135,7 @@ void ThreadManager::StartThread(VideoCore::RendererBase& renderer, Core::Fronten
             } else {
                 ASSERT(false);
             }
-            state.signaled_fence.store(next.fence);
-            if (next.block) {
-                // We have to lock the write_lock to ensure that the condition_variable wait not get a
-                // race between the check and the lock itself.
-                std::scoped_lock lk{state.write_lock};
-                state.cv.notify_all();
-            }
+            SignalCommand(state, next.fence, next.block);
         }
     });
 }
@@ -161,6 +201,7 @@ void ThreadManager::NotifyShutdown() {
         thread.request_stop();
         thread.join();
     }
+    VideoCommon::GpuSplit::Get().Stop();
 }
 
 } // namespace VideoCommon::GPUThread
diff --git a/src/video_core/macro.cpp b/src/video_core/macro.cpp
index 3422cb16fa..485b8eaf10 100644
--- a/src/video_core/macro.cpp
+++ b/src/video_core/macro.cpp
@@ -17,6 +17,7 @@
 
 #ifdef ARCHITECTURE_x86_64
 #include "common/x64/xbyak.h"
+#include "video_core/macro_turing_jit.h"
 #endif
 
 #include "common/assert.h"
@@ -1404,7 +1405,10 @@ void MacroEngine::Execute(Core::System& system, Engines::Maxwell3D& maxwell3d, u
                                  modes[0], modes[1], modes[2], modes[3]);
                     }
                 }
-                maxwell3d.CallMethod(system, address, value, true);
+                {
+                    VideoCommon::Nx2Profile::Scope method_scope{VideoCommon::Nx2Profile::Zone::Method};
+                    maxwell3d.CallMethod(system, address, value, true);
+                }
                 // Turing MME memory reads (methods 0x550 SET_MME_MEM_ADDRESS_A/B, 0x558 data RAM
                 // address, 0x55c MME_DMA_READ into data RAM, 0x560 MME_DMA_READ_FIFOED into
                 // the load FIFO, which the macro then reads after a load barrier).
@@ -1420,6 +1424,14 @@ void MacroEngine::Execute(Core::System& system, Engines::Maxwell3D& maxwell3d, u
                     // from a culling dispatch) and still be only in host buffers: flush them first,
                     // or the macro draws with an older frame's arguments (objects pop in and out).
                     const size_t read_bytes = size_t{std::min<u32>(value, 0x1000)} * 4;
+                    if (VideoCommon::Nx2Profile::Enabled()) {
+                        VideoCommon::Nx2Profile::Add(VideoCommon::Nx2Profile::Counter::DmaReadBytes, read_bytes);
+                        for (GPUVAddr page = source & ~GPUVAddr{0xfff}; page < source + read_bytes; page += 0x1000) {
+                            if (const auto device = memory.GpuToCpuAddress(page)) {
+                                VideoCommon::Nx2Profile::DmaReadPages().insert(*device >> 12);
+                            }
+                        }
+                    }
                     if (memory.IsMemoryDirty(source, read_bytes, VideoCommon::CacheType::BufferCache)) {
                         memory.FlushRegion(source, read_bytes, VideoCommon::CacheType::BufferCache);
                         static u64 flushes{};
@@ -1451,6 +1463,33 @@ void MacroEngine::Execute(Core::System& system, Engines::Maxwell3D& maxwell3d, u
                 }
                 return true;
             };
+#ifdef ARCHITECTURE_x86_64
+        // Macro code runs compiled to x86-64 (macro_turing_jit.h) instead of through the memo
+        // and the interpreter; EDEN_NX2_MME_JIT=0 goes back to those.
+        static const bool use_jit = [] {
+            const char* value = std::getenv("EDEN_NX2_MME_JIT");
+            return value == nullptr || value[0] != '0';
+        }();
+        if (use_jit) {
+            if (!native_jit) {
+                native_jit = std::make_shared<TuringMacroJit>();
+            }
+            native_jit->Prepare(native_code, native_code_generation);
+            fifo.assign(parameters.begin(), parameters.end());
+            const auto result = native_jit->Execute(
+                method, fifo, std::span<u32>{*native_interpreter.active_dram},
+                [&](u32 address) {
+                    return address < Maxwell3D::Regs::NUM_REGS ? maxwell3d.GetRegisterValue(address) : 0U;
+                },
+                emit);
+            if (!result.success) {
+                native_execution_failed = true;
+                LOG_ERROR(HW_GPU, "Native MME (JIT) failed: start={:#x}, consumed={}/{}, error={}",
+                          method, result.consumed, parameters.size(), result.error);
+            }
+            return;
+        }
+#endif
         // EDEN_NX2_MME_MEMO=0: interpret every call (see macro_memo.h).
         static const bool use_memo = [] {
             const char* value = std::getenv("EDEN_NX2_MME_MEMO");
@@ -1462,10 +1501,11 @@ void MacroEngine::Execute(Core::System& system, Engines::Maxwell3D& maxwell3d, u
                 memo_stats && (++native_calls & 0xffffff) == 0) {
                 const auto& stats = native_memo.GetStatistics();
                 LOG_INFO(HW_GPU,
-                         "Macro memo: {} hits, {} misses, {} guard failures, {} stores ({} in "
-                         "place), {} clears; {} entries, {} KiB",
-                         stats.hits, stats.misses, stats.guard_failures, stats.stores,
-                         stats.overwrites, stats.clears, native_memo.Size(),
+                         "Macro memo: {} hits ({} on the last call's slot), {} misses, {} guard "
+                         "failures, {} stores ({} in place), {} compactions, {} clears; {} entries, "
+                         "{} KiB",
+                         stats.hits, stats.slot_hits, stats.misses, stats.guard_failures, stats.stores,
+                         stats.overwrites, stats.compactions, stats.clears, native_memo.Size(),
                          native_memo.ArenaWords() * 4 / 1024);
             }
             native_memo.Synchronize(native_code_generation);
diff --git a/src/video_core/macro.h b/src/video_core/macro.h
index a6ca4c4bd2..0d574636e9 100644
--- a/src/video_core/macro.h
+++ b/src/video_core/macro.h
@@ -9,6 +9,10 @@
 #include <memory>
 #include "video_core/macro_memo.h"
 #include "video_core/macro_turing.h"
+
+namespace Tegra {
+class TuringMacroJit;
+}
 #include <span>
 #include <variant>
 #include <vector>
@@ -257,6 +261,8 @@ struct MacroEngine {
     // Calls replayed instead of interpreted (macro_memo.h), and what the current call wrote.
     u64 native_code_generation{};
     BasicMacroMemo<::Common::unordered_map<u64, u32, MacroMemoHash>> native_memo;
+    /// The macro code compiled to x86-64 (EDEN_NX2_MME_JIT=0 turns it off), created on first use.
+    std::shared_ptr<TuringMacroJit> native_jit;
     MacroMemoPairs native_writes;
     u64 native_calls{};
 
diff --git a/src/video_core/macro_memo.h b/src/video_core/macro_memo.h
index 0e994fc78a..7fe914b4f2 100644
--- a/src/video_core/macro_memo.h
+++ b/src/video_core/macro_memo.h
@@ -41,7 +41,8 @@ public:
     };
 
     struct Statistics {
-        std::uint64_t hits{}, misses{}, guard_failures{}, stores{}, overwrites{}, clears{};
+        std::uint64_t hits{}, slot_hits{}, misses{}, guard_failures{}, stores{}, overwrites{},
+            compactions{}, clears{};
     };
 
     BasicMacroMemo() {
@@ -57,6 +58,7 @@ public:
         if (last.first == start && last.second != NoRecord &&
             Match(last.second, start, parameters, data_ram, effects)) {
             ++statistics.hits;
+            ++statistics.slot_hits;
             return true;
         }
         const auto it = entries.find(Hash(start, parameters));
@@ -127,8 +129,7 @@ public:
             }
         }
         if (arena.size() >= MaxArenaWords) {
-            Clear();
-            ++statistics.clears;
+            Compact();
         }
         const auto offset = static_cast<std::uint32_t>(arena.size());
         arena.push_back(start);
@@ -186,8 +187,31 @@ private:
         last_by_start.fill({0, NoRecord});
     }
 
+    /// Copies the records still in the map to a new arena. A call stored again with another
+    /// shape leaves its old record behind: title 1 filled the arena with them every ~2 s, and
+    /// clearing it made the next ~20,000 calls interpret again. Clears only when the live
+    /// records alone would fill most of it.
+    void Compact() {
+        compacted.clear();
+        for (auto& entry : entries) {
+            const std::uint32_t* record = arena.data() + entry.second;
+            const std::size_t words =
+                HeaderWords + record[1] + std::size_t{2} * (record[2] + record[3] + record[4]);
+            entry.second = static_cast<std::uint32_t>(compacted.size());
+            compacted.insert(compacted.end(), record, record + words);
+        }
+        arena.swap(compacted);
+        last_by_start.fill({0, NoRecord});
+        ++statistics.compactions;
+        if (arena.size() >= MaxArenaWords / 4 * 3) {
+            Clear();
+            ++statistics.clears;
+        }
+    }
+
     Map entries;
     std::vector<std::uint32_t> arena;
+    std::vector<std::uint32_t> compacted; ///< Compact's new arena (kept for its capacity)
     /// The last record found or stored per macro start (direct mapped).
     std::array<std::pair<std::uint32_t, std::uint32_t>, 256> last_by_start{};
     std::uint64_t generation{};
diff --git a/src/video_core/macro_turing_jit.h b/src/video_core/macro_turing_jit.h
new file mode 100644
index 0000000000..d0f24080a3
--- /dev/null
+++ b/src/video_core/macro_turing_jit.h
@@ -0,0 +1,662 @@
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+#pragma once
+
+#include <array>
+#include <cstddef>
+#include <cstdint>
+#include <memory>
+#include <span>
+#include <string_view>
+#include <vector>
+
+// Eden's Xbyak configuration (its hash maps): every translation unit must see the same one.
+#include "common/x64/xbyak.h"
+
+#include "video_core/macro_turing.h"
+
+namespace Tegra {
+
+// x86-64 code for the TU104/GA10x MME: every instruction of the uploaded macro memory compiled
+// once (each a block that falls through to the next one, branches and loops through a jump
+// table), with the semantics of TuringMacroInterpreter::Execute. Title 1 calls macros ~1.3
+// million times a second; the interpreter and the macro memo's lookups were ~10% of its GPU
+// thread. Emits (method writes) and register reads call back into C++.
+class TuringMacroJit {
+public:
+    using Word = std::uint32_t;
+    using Result = TuringMacroInterpreter::Result;
+
+    struct State {
+        std::array<Word, 24> regs;
+        Word carry;
+        Word method;
+        Word increment;
+        Word loop_count;
+        Word loop_start;
+        Word loop_end;
+        Word next_pc;
+        Word steps;
+        Word start;
+        Word consumed;
+        Word params_size;
+        Word fail_pc;
+        std::uint8_t end_next;
+        std::uint8_t stopped;
+        std::uint8_t cancel_end;
+        std::uint8_t pad;
+        Word scratch_x;
+        Word scratch_y;
+        const Word* params;
+        Word* dram;
+        void* user;
+        bool (*emit)(State*, Word method, Word value);
+        Word (*read)(State*, Word address);
+        const std::vector<Word>* fifo;
+    };
+
+    enum Error : Word {
+        None = 0,
+        FifoUnderflow,
+        InvalidSource,
+        InvalidDestination,
+        UnsupportedExtended,
+        InvalidMerge,
+        InvalidLoop,
+        DataRamReadOutOfBounds,
+        DataRamWriteOutOfBounds,
+        InvalidOutput,
+        UnsupportedEmit,
+        OutsideCode,
+        StepLimit,
+    };
+
+    static std::string_view ErrorText(Word error) {
+        constexpr std::array<std::string_view, 13> texts{
+            "",
+            "parameter FIFO underflow",
+            "invalid source register",
+            "invalid destination register",
+            "unsupported extended operation",
+            "invalid merge destination",
+            "invalid loop",
+            "data RAM read out of bounds",
+            "data RAM write out of bounds",
+            "invalid output selector",
+            "unsupported emitted method",
+            "instruction outside uploaded RAM",
+            "instruction limit exceeded",
+        };
+        return error < texts.size() ? texts[error] : "unknown";
+    }
+
+    /// Compiles `code` (3 words per instruction) when it differs from what was compiled.
+    void Prepare(std::span<const Word> code, std::uint64_t generation) {
+        if (program && generation == compiled_generation) {
+            return;
+        }
+        program = std::make_unique<Program>(code);
+        compiled_generation = generation;
+    }
+
+    template <typename Read, typename Emit>
+    Result Execute(Word start, std::vector<Word>& parameters, std::span<Word> data_ram, Read read,
+                   Emit emit) {
+        State state;
+        state.regs.fill(0);
+        state.carry = state.method = state.increment = 0;
+        state.loop_count = state.loop_start = state.loop_end = 0;
+        state.consumed = 0;
+        state.fail_pc = 0;
+        state.end_next = state.stopped = state.cancel_end = 0;
+        state.start = start;
+        state.steps = 1000000;
+        state.params = parameters.data();
+        state.params_size = static_cast<Word>(parameters.size());
+        state.dram = data_ram.data();
+        state.fifo = &parameters;
+        Thunks<Read, Emit> thunks{read, emit};
+        state.user = &thunks;
+        state.emit = &Thunks<Read, Emit>::CallEmit;
+        state.read = &Thunks<Read, Emit>::CallRead;
+        const Word error = program->entry(&state);
+        if (error == None) {
+            return {true, state.fail_pc, state.consumed, {}};
+        }
+        return {false, state.fail_pc, state.consumed, ErrorText(error)};
+    }
+
+private:
+    template <typename Read, typename Emit>
+    struct Thunks {
+        Read& read;
+        Emit& emit;
+        static bool CallEmit(State* state, Word method, Word value) {
+            const bool ok = static_cast<Thunks*>(state->user)->emit(method, value);
+            // MME_DMA_READ_FIFOED appends to the FIFO from inside the emit.
+            state->params = state->fifo->data();
+            state->params_size = static_cast<Word>(state->fifo->size());
+            return ok;
+        }
+        static Word CallRead(State* state, Word address) {
+            return static_cast<Thunks*>(state->user)->read(address);
+        }
+    };
+
+    struct Program : Xbyak::CodeGenerator {
+        explicit Program(std::span<const Word> code)
+            : Xbyak::CodeGenerator(1 << 20, Xbyak::AutoGrow) {
+            Generate(code);
+            ready();
+            entry = getCode<Word (*)(State*)>();
+        }
+
+        Word (*entry)(State*) = nullptr;
+
+    private:
+        using Decoded = TuringMacroInterpreter::Decoded;
+#ifdef _WIN32
+        const Xbyak::Reg64 arg0 = rcx, arg1 = rdx, arg2 = r8;
+#else
+        const Xbyak::Reg64 arg0 = rdi, arg1 = rsi, arg2 = rdx;
+#endif
+        // rbx: State*; r12d/r13d: the instruction's FIFO loads; r14d/r15d: its ALU results.
+        const Xbyak::Reg64 st = rbx;
+
+        static constexpr std::size_t RegOffset(Word reg) {
+            return offsetof(State, regs) + reg * sizeof(Word);
+        }
+
+        void Generate(std::span<const Word> code) {
+            const Word count = static_cast<Word>(code.size() / 3);
+            labels.resize(count);
+            std::vector<Decoded> decoded(count);
+            for (Word pc = 0; pc < count; ++pc) {
+                decoded[pc] = TuringMacroInterpreter::Decode(
+                    {code[pc * 3 + 2], code[pc * 3 + 1], code[pc * 3]});
+            }
+            Analyze(decoded);
+            push(rbx); push(rbp); push(r12); push(r13); push(r14); push(r15);
+            sub(rsp, 40);
+            mov(st, arg0);
+            mov(eax, dword[st + offsetof(State, start)]);
+            cmp(eax, count);
+            jae(fail_outside, T_NEAR);
+            lea(rcx, ptr[rip + table]);
+            jmp(ptr[rcx + rax * 8]);
+
+            for (Word pc = 0; pc < count; ++pc) {
+                L(labels[pc]);
+                Instruction(pc, count, decoded[pc]);
+            }
+
+            // Dispatch on eax (the next pc). Every backward edge comes through here, so the
+            // instruction limit is counted here.
+            L(dispatch);
+            sub(dword[st + offsetof(State, steps)], 1);
+            jz(fail_steps, T_NEAR);
+            cmp(eax, count);
+            jae(fail_outside, T_NEAR);
+            lea(rcx, ptr[rip + table]);
+            jmp(ptr[rcx + rax * 8]);
+
+            L(success);
+            xor_(eax, eax);
+            L(leave);
+            add(rsp, 40);
+            pop(r15); pop(r14); pop(r13); pop(r12); pop(rbp); pop(rbx);
+            ret();
+            const auto fail_with = [&](Xbyak::Label& label, Word error) {
+                L(label);
+                mov(eax, error);
+                jmp(leave, T_NEAR);
+            };
+            fail_with(fail_underflow, FifoUnderflow);
+            fail_with(fail_source, InvalidSource);
+            fail_with(fail_destination, InvalidDestination);
+            fail_with(fail_extended, UnsupportedExtended);
+            fail_with(fail_merge, InvalidMerge);
+            fail_with(fail_loop, InvalidLoop);
+            fail_with(fail_dread, DataRamReadOutOfBounds);
+            fail_with(fail_dwrite, DataRamWriteOutOfBounds);
+            fail_with(fail_output, InvalidOutput);
+            fail_with(fail_emit, UnsupportedEmit);
+            fail_with(fail_outside, OutsideCode);
+            fail_with(fail_steps, StepLimit);
+
+            align(8);
+            L(table);
+            for (Word pc = 0; pc < count; ++pc) {
+                putL(labels[pc]);
+            }
+        }
+
+        // Loads MME register `reg` (as an operand of `lane`) into `dst`.
+        void Get(const Xbyak::Reg32& dst, const Decoded& d, Word reg, Word lane) {
+            if (reg < 24) {
+                mov(dst, dword[st + RegOffset(reg)]);
+                return;
+            }
+            switch (reg) {
+            case 24: xor_(dst, dst); return;
+            case 25: mov(dst, Word(std::int32_t(std::int16_t(d.imm[lane])))); return;
+            case 26: mov(dst, Word(std::int32_t(std::int16_t(d.imm[1 - lane])))); return;
+            case 27: mov(dst, (d.imm[0] << 16) | d.imm[1]); return;
+            case 28: mov(dst, r12d); return;
+            case 29: mov(dst, r13d); return;
+            default: jmp(fail_source, T_NEAR); return;
+            }
+        }
+
+        static bool IsFlow(Word op) {
+            return op >= 23 && op <= 29;
+        }
+
+        // Where a loop can end (its back edge is checked only there) and which instructions can
+        // run right after an END (only those check for the end of the macro).
+        void Analyze(const std::vector<Decoded>& decoded) {
+            const Word count = static_cast<Word>(decoded.size());
+            loop_end.assign(count, false);
+            after_end.assign(count, false);
+            std::vector<Word> loop_starts_of_end(count, ~0U);
+            const auto mark = [&](std::vector<bool>& set, std::int64_t pc) {
+                if (pc >= 0 && pc < count) set[static_cast<std::size_t>(pc)] = true;
+            };
+            for (Word pc = 0; pc < count; ++pc) {
+                for (Word lane = 0; lane < 2; ++lane) {
+                    const Decoded& d = decoded[pc];
+                    if (d.op[lane] == 23 && d.imm[lane] >= 2 && pc + d.imm[lane] - 1 < count) {
+                        loop_end[pc + d.imm[lane] - 1] = true;
+                    }
+                }
+            }
+            for (Word pc = 0; pc < count; ++pc) {
+                const Decoded& d = decoded[pc];
+                if (!d.end) continue;
+                mark(after_end, std::int64_t{pc} + 1);
+                for (Word lane = 0; lane < 2; ++lane) {
+                    const Word imm = d.imm[lane];
+                    switch (d.op[lane]) {
+                    case 23: mark(after_end, std::int64_t{pc} + imm); break;
+                    case 24: mark(after_end, std::int64_t{pc} + (imm & 0x7fff)); break;
+                    case 25: case 26: case 27: case 28: case 29:
+                        mark(after_end, std::int64_t{pc} + (std::int32_t(imm << 19) >> 19));
+                        break;
+                    default: break;
+                    }
+                }
+            }
+            // An END at a loop's last instruction can be followed by the loop's second one.
+            for (Word pc = 0; pc < count; ++pc) {
+                for (Word lane = 0; lane < 2; ++lane) {
+                    const Decoded& d = decoded[pc];
+                    if (d.op[lane] == 23 && d.imm[lane] >= 2) {
+                        const Word end = pc + d.imm[lane] - 1;
+                        if (end < count && decoded[end].end) mark(after_end, std::int64_t{pc} + 1);
+                    }
+                }
+            }
+        }
+
+        void Instruction(Word pc, Word count, const Decoded& d) {
+            // FIFO loads, in lane order.
+            for (Word i = 0; i < 2; ++i) {
+                const Xbyak::Reg32 load = i == 0 ? r12d : r13d;
+                if (!d.load[i]) {
+                    xor_(load, load);
+                    continue;
+                }
+                mov(eax, dword[st + offsetof(State, consumed)]);
+                cmp(eax, dword[st + offsetof(State, params_size)]);
+                jae(fail_underflow, T_NEAR);
+                mov(rcx, qword[st + offsetof(State, params)]);
+                mov(load, dword[rcx + rax * 4]);
+                inc(eax);
+                mov(dword[st + offsetof(State, consumed)], eax);
+            }
+            if (!d.pred_mode) {
+                Body(pc, count, d, 15);
+                return;
+            }
+            Xbyak::Label when_false;
+            Get(eax, d, d.pred_reg, 0);
+            test(eax, eax);
+            jz(when_false, T_NEAR);
+            Body(pc, count, d, d.mask_true);
+            L(when_false);
+            Body(pc, count, d, d.mask_false);
+        }
+
+        void Body(Word pc, Word count, const Decoded& d, Word mask) {
+            // An inactive lane's result reads as 0 when an output selects it.
+            for (Word lane = 0; lane < 2; ++lane) {
+                if (mask & (1U << lane)) continue;
+                bool selected = false;
+                for (Word out = 0; out < 2; ++out) {
+                    if (!(mask & (4U << out))) continue;
+                    selected |= d.out_method[out] == lane + 1 || d.out_emit[out] == lane + 1;
+                }
+                if (selected) {
+                    xor_(lane == 0 ? r14d : r15d, lane == 0 ? r14d : r15d);
+                }
+            }
+            const bool flow = ((mask & 1) && IsFlow(d.op[0])) || ((mask & 2) && IsFlow(d.op[1]));
+            if (flow) {
+                mov(dword[st + offsetof(State, next_pc)], pc + 1);
+                mov(byte[st + offsetof(State, cancel_end)], 0);
+            }
+            // Lane 1's operands are read before lane 0 commits its destination (or calls out).
+            const bool lane0_active = (mask & 1) != 0;
+            const bool early_operands =
+                (mask & 2) && lane0_active &&
+                (d.op[0] == 22 || (d.dst[0] < 24 && (d.dst[0] == d.src_a[1] || d.dst[0] == d.src_b[1])));
+            if (early_operands) {
+                Get(eax, d, d.src_a[1], 1);
+                mov(dword[st + offsetof(State, scratch_x)], eax);
+                Get(eax, d, d.src_b[1], 1);
+                mov(dword[st + offsetof(State, scratch_y)], eax);
+            }
+            for (Word lane = 0; lane < 2; ++lane) {
+                if (!(mask & (1U << lane))) {
+                    continue;
+                }
+                if (lane == 1 && early_operands) {
+                    mov(ecx, dword[st + offsetof(State, scratch_x)]);
+                    mov(edx, dword[st + offsetof(State, scratch_y)]);
+                } else {
+                    Get(ecx, d, d.src_a[lane], lane);
+                    Get(edx, d, d.src_b[lane], lane);
+                }
+                Lane(pc, d, lane);
+                // eax holds the lane's value.
+                mov(lane == 0 ? r14d : r15d, eax);
+                if (d.dst[lane] < 24) {
+                    mov(dword[st + RegOffset(d.dst[lane])], eax);
+                } else if (d.dst[lane] != 24) {
+                    jmp(fail_destination, T_NEAR);
+                }
+            }
+            for (Word lane = 0; lane < 2; ++lane) {
+                if (!(mask & (4U << lane))) {
+                    continue;
+                }
+                if (d.out_method[lane]) {
+                    Output(eax, d, d.out_method[lane]);
+                    mov(ecx, eax);
+                    and_(ecx, 0xfff);
+                    mov(dword[st + offsetof(State, method)], ecx);
+                    shr(eax, 12);
+                    and_(eax, 15);
+                    mov(dword[st + offsetof(State, increment)], eax);
+                }
+                if (d.out_emit[lane]) {
+                    Output(eax, d, d.out_emit[lane]);
+                    mov(arg2.cvt32(), eax);
+                    mov(arg1.cvt32(), dword[st + offsetof(State, method)]);
+                    mov(arg0, st);
+                    call(ptr[st + offsetof(State, emit)]);
+                    test(al, al);
+                    jz(fail_emit, T_NEAR);
+                    mov(eax, dword[st + offsetof(State, increment)]);
+                    add(dword[st + offsetof(State, method)], eax);
+                }
+            }
+            // An END in the previous instruction ends the macro after this one.
+            if (after_end[pc]) {
+                cmp(byte[st + offsetof(State, end_next)], 0);
+                jne(success, T_NEAR);
+            }
+            if (flow) {
+                cmp(byte[st + offsetof(State, stopped)], 0);
+                jne(success, T_NEAR);
+            }
+            if (d.end) {
+                if (flow) {
+                    mov(al, byte[st + offsetof(State, cancel_end)]);
+                    xor_(al, 1);
+                    mov(byte[st + offsetof(State, end_next)], al);
+                } else {
+                    mov(byte[st + offsetof(State, end_next)], 1);
+                }
+            }
+            if (loop_end[pc]) {
+                Xbyak::Label no_loop;
+                mov(eax, dword[st + offsetof(State, loop_count)]);
+                test(eax, eax);
+                jz(no_loop);
+                cmp(dword[st + offsetof(State, loop_end)], pc);
+                jne(no_loop);
+                dec(dword[st + offsetof(State, loop_count)]);
+                mov(eax, dword[st + offsetof(State, loop_start)]);
+                inc(eax);
+                jmp(dispatch, T_NEAR);
+                L(no_loop);
+            }
+            if (flow) {
+                mov(eax, dword[st + offsetof(State, next_pc)]);
+                if (pc + 1 < count) {
+                    cmp(eax, pc + 1);
+                    jne(dispatch, T_NEAR);
+                    jmp(labels[pc + 1], T_NEAR);
+                } else {
+                    jmp(dispatch, T_NEAR);
+                }
+            } else if (pc + 1 < count) {
+                jmp(labels[pc + 1], T_NEAR);
+            } else {
+                jmp(fail_outside, T_NEAR);
+            }
+        }
+
+        void Output(const Xbyak::Reg32& dst, const Decoded& d, Word selector) {
+            switch (selector) {
+            case 1: mov(dst, r14d); return;
+            case 2: mov(dst, r15d); return;
+            case 3: mov(dst, r12d); return;
+            case 4: mov(dst, r13d); return;
+            case 5: mov(dst, d.imm[0]); return;
+            case 6: mov(dst, d.imm[1]); return;
+            case 8: mov(dst, d.imm[0] >> 12); return;
+            case 9: mov(dst, d.imm[1] >> 12); return;
+            case 10: mov(dst, (d.imm[0] << 16) | d.imm[1]); return;
+            default: jmp(fail_output, T_NEAR); return;
+            }
+        }
+
+        // ecx = x, edx = y; leaves the value in eax.
+        void Compare(Word op) {
+            cmp(ecx, edx);
+            switch (op) {
+            case 17: case 25: setl(al); break;
+            case 18: case 26: setb(al); break;
+            case 19: case 27: setle(al); break;
+            case 20: case 28: setbe(al); break;
+            default: sete(al); break;
+            }
+            movzx(eax, al);
+        }
+
+        void Lane(Word pc, const Decoded& d, Word lane) {
+            const Word imm = d.imm[lane];
+            const Word op = d.op[lane];
+            switch (op) {
+            case 0:
+                mov(eax, ecx);
+                add(eax, edx);
+                setc(r8b);
+                movzx(r8d, r8b);
+                mov(dword[st + offsetof(State, carry)], r8d);
+                return;
+            case 1:
+                lea(eax, ptr[rcx + rdx]);
+                add(eax, dword[st + offsetof(State, carry)]);
+                return;
+            case 2:
+                mov(eax, ecx);
+                sub(eax, edx);
+                setc(r8b);
+                movzx(r8d, r8b);
+                mov(dword[st + offsetof(State, carry)], r8d);
+                return;
+            case 3:
+                mov(eax, ecx);
+                sub(eax, edx);
+                sub(eax, dword[st + offsetof(State, carry)]);
+                return;
+            case 4:
+                movsxd(rax, ecx);
+                movsxd(rdx, edx);
+                imul(rax, rdx);
+                mov(rcx, rax);
+                shr(rcx, 32);
+                mov(dword[st + offsetof(State, carry)], ecx);
+                return;
+            case 6:
+                mov(eax, ecx);
+                mov(edx, edx);
+                imul(rax, rdx);
+                mov(rcx, rax);
+                shr(rcx, 32);
+                mov(dword[st + offsetof(State, carry)], ecx);
+                return;
+            case 5:
+                mov(eax, dword[st + offsetof(State, carry)]);
+                return;
+            case 7:
+                cmp(ecx, 0x1000);
+                jne(fail_extended, T_NEAR);
+                xor_(eax, eax);
+                return;
+            case 8: {
+                Xbyak::Label zero, done;
+                test(ecx, ecx);
+                jz(zero);
+                bsr(eax, ecx);
+                xor_(eax, 31);
+                jmp(done);
+                L(zero);
+                mov(eax, 32);
+                L(done);
+                return;
+            }
+            case 9:
+                mov(eax, ecx);
+                mov(ecx, edx);
+                and_(ecx, 31);
+                shl(eax, cl);
+                return;
+            case 10:
+                mov(eax, ecx);
+                mov(ecx, edx);
+                and_(ecx, 31);
+                shr(eax, cl);
+                return;
+            case 11:
+                mov(eax, ecx);
+                mov(ecx, edx);
+                and_(ecx, 31);
+                sar(eax, cl);
+                return;
+            case 12: mov(eax, ecx); and_(eax, edx); return;
+            case 13: mov(eax, ecx); and_(eax, edx); not_(eax); return;
+            case 14: mov(eax, ecx); or_(eax, edx); return;
+            case 15: mov(eax, ecx); xor_(eax, edx); return;
+            case 16: {
+                const Word dest = (imm >> 10) & 63, bit_count = (imm >> 5) & 31;
+                const Word source = imm & 31;
+                if (dest >= 32) {
+                    jmp(fail_merge, T_NEAR);
+                    return;
+                }
+                const Word bits_mask = (1U << bit_count) - 1;
+                mov(eax, ecx);
+                and_(eax, ~(bits_mask << dest));
+                shr(edx, source);
+                and_(edx, bits_mask);
+                shl(edx, dest);
+                or_(eax, edx);
+                return;
+            }
+            case 17: case 18: case 19: case 20: case 21:
+                Compare(op);
+                neg(eax);
+                return;
+            case 22:
+                lea(eax, ptr[rcx + rdx]);
+                and_(eax, 0xffff);
+                mov(arg1.cvt32(), eax);
+                mov(arg0, st);
+                call(ptr[st + offsetof(State, read)]);
+                return;
+            case 23: {
+                if (imm < 2) {
+                    jmp(fail_loop, T_NEAR);
+                    return;
+                }
+                Xbyak::Label skip, done;
+                cmp(dword[st + offsetof(State, loop_count)], 0);
+                jne(fail_loop, T_NEAR);
+                test(ecx, ecx);
+                jz(skip);
+                dec(ecx);
+                mov(dword[st + offsetof(State, loop_count)], ecx);
+                mov(dword[st + offsetof(State, loop_start)], pc);
+                mov(dword[st + offsetof(State, loop_end)], pc + imm - 1);
+                jmp(done);
+                L(skip);
+                mov(dword[st + offsetof(State, next_pc)], pc + imm);
+                L(done);
+                xor_(eax, eax);
+                return;
+            }
+            case 24:
+                mov(dword[st + offsetof(State, next_pc)], pc + (imm & 0x7fff));
+                xor_(eax, eax);
+                return;
+            case 25: case 26: case 27: case 28: case 29: {
+                Xbyak::Label not_taken;
+                Compare(op);
+                cmp(eax, (imm & 0x8000) ? 1 : 0);
+                jne(not_taken);
+                const auto offset = std::int32_t(imm << 19) >> 19;
+                if (offset == -4096) {
+                    mov(byte[st + offsetof(State, stopped)], 1);
+                } else {
+                    mov(dword[st + offsetof(State, next_pc)], Word(std::int32_t(pc) + offset));
+                    mov(byte[st + offsetof(State, cancel_end)], (imm & 0x4000) ? 1 : 0);
+                }
+                L(not_taken);
+                xor_(eax, eax);
+                return;
+            }
+            case 30:
+                cmp(ecx, 0xc00);
+                jae(fail_dread, T_NEAR);
+                mov(rax, qword[st + offsetof(State, dram)]);
+                mov(eax, dword[rax + rcx * 4]);
+                return;
+            case 31:
+                cmp(ecx, 0xc00);
+                jae(fail_dwrite, T_NEAR);
+                mov(rax, qword[st + offsetof(State, dram)]);
+                mov(dword[rax + rcx * 4], edx);
+                xor_(eax, eax);
+                return;
+            default:
+                xor_(eax, eax);
+                return;
+            }
+        }
+
+        std::vector<Xbyak::Label> labels;
+        std::vector<bool> loop_end, after_end;
+        Xbyak::Label table, dispatch, success, leave;
+        Xbyak::Label fail_underflow, fail_source, fail_destination, fail_extended, fail_merge,
+            fail_loop, fail_dread, fail_dwrite, fail_output, fail_emit, fail_outside, fail_steps;
+    };
+
+    std::unique_ptr<Program> program;
+    std::uint64_t compiled_generation{~std::uint64_t{0}};
+};
+
+} // namespace Tegra
diff --git a/src/video_core/memory_manager.h b/src/video_core/memory_manager.h
index 105cb14e23..bd5025e9fd 100644
--- a/src/video_core/memory_manager.h
+++ b/src/video_core/memory_manager.h
@@ -165,6 +165,12 @@ public:
 
     void FlushCaching();
 
+    /// Leaves the caches' invalidation of a range already written for the next FlushCaching
+    /// (before every draw, dispatch, clear and copy), merged with neighbouring ranges.
+    void AccumulateInvalidation(GPUVAddr gpu_addr, std::size_t size) {
+        accumulator.Add(gpu_addr, size);
+    }
+
     const u8* GetSpan(const GPUVAddr src_addr, const std::size_t size) const;
     u8* GetSpan(const GPUVAddr src_addr, const std::size_t size);
 
diff --git a/src/video_core/nx2_profile.h b/src/video_core/nx2_profile.h
index a162ce16a3..a921b72082 100644
--- a/src/video_core/nx2_profile.h
+++ b/src/video_core/nx2_profile.h
@@ -7,6 +7,7 @@
 #include <chrono>
 #include <cstdlib>
 #include <string>
+#include <unordered_set>
 
 #include "common/common_types.h"
 #include "common/logging.h"
@@ -45,6 +46,11 @@ enum class Zone : u32 {
     DynamicState,
     DrawRecord,
     TextureUpload,
+    Fence,        ///< Fences made (SignalFence)
+    FenceSkipped, ///< Empty fences not made
+    Method,       ///< Engine method calls (what a split GPU thread's back half would run)
+    SplitDrain,   ///< Split GPU thread: the front waits for the back (gpu_split.h)
+    SplitWait,    ///< Split GPU thread: the back waits for work
     Count,
 };
 
@@ -54,12 +60,15 @@ inline constexpr std::array<const char*, static_cast<size_t>(Zone::Count)> zone_
     "clear",         "compute",       "flush work",     "pipeline lookup", "configure",
     "sync descriptors", "read handles", "fill image views", "texture buffers", "update buffers",
     "descriptor acquire", "bind stage buffers", "push descriptors", "render targets",
-    "configure draw", "dynamic state", "draw record", "texture upload",
+    "configure draw", "dynamic state", "draw record", "texture upload", "fence",
+    "fence skipped", "method", "split drain", "split wait",
 };
 
+// Read once at startup: Enabled() runs several times per method write (no guard to check).
+inline const bool g_enabled = std::getenv("EDEN_NX2_GPU_PROFILE") != nullptr;
+
 inline bool Enabled() {
-    static const bool enabled = std::getenv("EDEN_NX2_GPU_PROFILE") != nullptr;
-    return enabled;
+    return g_enabled;
 }
 
 // Nanoseconds; from the host's own counter where the wall clock reads it directly (the zones
@@ -71,6 +80,7 @@ inline u64 Now() {
 struct ThreadData {
     std::array<u64, static_cast<size_t>(Zone::Count)> exclusive_ns{};
     std::array<u64, static_cast<size_t>(Zone::Count)> count{};
+    std::array<u64, 8> counters{}; ///< See Counter
     u64 child_ns{};    ///< Time of the zones nested in the current one
     u64 window_start{};
 };
@@ -110,8 +120,34 @@ private:
     u64 parent_child_ns{};
 };
 
-/// Called by the GPU thread between commands: logs and resets the totals every 10 s.
-inline void Report() {
+/// Amounts counted on the GPU thread, logged per second with the zones.
+enum class Counter : u32 {
+    DownloadBytes,  ///< Buffer data copied back from the host GPU at fences
+    DownloadCopies, ///< Buffer ranges copied back
+    DmaReadBytes,   ///< Guest memory the native MME read with DMA
+    DownloadBytesOnDmaPages, ///< Of DownloadBytes: on pages the MME has DMA-read
+    MethodWrites,   ///< 3D register writes (Maxwell3D::CallMethod below the macro range)
+    SameValueWrites, ///< Of MethodWrites: the register already held the value
+    AcquireChecks,  ///< Semaphore acquires evaluated
+    AcquireWaits,   ///< Of AcquireChecks: not met yet (the channel waits)
+};
+
+/// Device pages (4 KiB) the native MME has DMA-read since start (GPU thread only; filled only
+/// while profiling): which buffer downloads anything on the GPU thread reads back.
+inline std::unordered_set<u64>& DmaReadPages() {
+    static std::unordered_set<u64> pages;
+    return pages;
+}
+
+inline void Add(Counter counter, u64 amount) {
+    if (Enabled()) {
+        t_data.counters[static_cast<size_t>(counter)] += amount;
+    }
+}
+
+/// Called by the GPU thread (and the split's back thread) between commands: logs and resets
+/// the totals every 10 s.
+inline void Report(const char* thread_name = "GPU thread") {
     if (!Enabled()) {
         return;
     }
@@ -142,7 +178,19 @@ inline void Report() {
         const double calls = static_cast<double>(data.count[i]) * 1e9 / window;
         line += fmt::format(" | {} {:.1f} ms ({:.0f}/s)", zone_names[i], ms, calls);
     }
-    LOG_INFO(HW_GPU, "GPU thread per second:{}", line);
+    line += fmt::format(" | buffer downloads {:.0f} KiB in {:.0f} copies ({:.0f} KiB on pages "
+                        "MME DMA reads) | MME DMA reads {:.0f} KiB",
+                        static_cast<double>(data.counters[0]) * 1e9 / 1024.0 / window,
+                        static_cast<double>(data.counters[1]) * 1e9 / window,
+                        static_cast<double>(data.counters[3]) * 1e9 / 1024.0 / window,
+                        static_cast<double>(data.counters[2]) * 1e9 / 1024.0 / window);
+    line += fmt::format(" | 3D register writes {:.0f} ({:.0f} same value)",
+                        static_cast<double>(data.counters[4]) * 1e9 / window,
+                        static_cast<double>(data.counters[5]) * 1e9 / window);
+    line += fmt::format(" | acquires {:.0f} ({:.0f} waited)",
+                        static_cast<double>(data.counters[6]) * 1e9 / window,
+                        static_cast<double>(data.counters[7]) * 1e9 / window);
+    LOG_INFO(HW_GPU, "{} per second:{}", thread_name, line);
     data = {};
     data.window_start = now;
 }
diff --git a/src/video_core/query_cache/query_cache.h b/src/video_core/query_cache/query_cache.h
index 245ab1996c..60ca50eb1f 100644
--- a/src/video_core/query_cache/query_cache.h
+++ b/src/video_core/query_cache/query_cache.h
@@ -520,6 +520,16 @@ bool QueryCacheBase<Traits>::HasUncommittedFlushes() const {
 }
 
 template <typename Traits>
+bool QueryCacheBase<Traits>::HasPendingFenceWork() const {
+    bool result = false;
+    impl->ForEachStreamer([&result](StreamerInterface* streamer) {
+        result |= streamer->HasUnsyncedQueries() || streamer->HasPendingSync();
+        return result;
+    });
+    return result;
+}
+
+template <typename Traits>
 bool QueryCacheBase<Traits>::ShouldWaitAsyncFlushes() {
     std::scoped_lock lk(impl->flush_guard);
     return !impl->flushes_pending.empty() && impl->flushes_pending.front() != 0ULL;
diff --git a/src/video_core/query_cache/query_cache_base.h b/src/video_core/query_cache/query_cache_base.h
index d73e9039f7..5a60e42986 100644
--- a/src/video_core/query_cache/query_cache_base.h
+++ b/src/video_core/query_cache/query_cache_base.h
@@ -109,6 +109,9 @@ public:
 
     bool HasUncommittedFlushes() const;
 
+    /// Whether a new fence has queries to flush or host writes to sync for this cache.
+    bool HasPendingFenceWork() const;
+
     bool ShouldWaitAsyncFlushes();
 
     void PopAsyncFlushes();
diff --git a/src/video_core/renderer_opengl/gl_graphics_pipeline.cpp b/src/video_core/renderer_opengl/gl_graphics_pipeline.cpp
index 83545463ac..ecd6544b10 100644
--- a/src/video_core/renderer_opengl/gl_graphics_pipeline.cpp
+++ b/src/video_core/renderer_opengl/gl_graphics_pipeline.cpp
@@ -20,7 +20,9 @@
 #include "video_core/shader_notify.h"
 #include "video_core/texture_cache/texture_cache.h"
 
-#if defined(_MSC_VER) && defined(NDEBUG)
+#if defined(__clang__) && defined(NDEBUG)
+#define LAMBDA_FORCEINLINE __attribute__((always_inline))
+#elif defined(_MSC_VER) && defined(NDEBUG)
 #define LAMBDA_FORCEINLINE [[msvc::forceinline]]
 #else
 #define LAMBDA_FORCEINLINE
diff --git a/src/video_core/renderer_vulkan/fixed_pipeline_state.cpp b/src/video_core/renderer_vulkan/fixed_pipeline_state.cpp
index 0ced890929..9dac98c830 100644
--- a/src/video_core/renderer_vulkan/fixed_pipeline_state.cpp
+++ b/src/video_core/renderer_vulkan/fixed_pipeline_state.cpp
@@ -179,7 +179,14 @@ void FixedPipelineState::Refresh(Tegra::Engines::Maxwell3D& maxwell3d, DynamicFe
     for (size_t i = 0; i < regs.rt.size(); ++i) {
         color_formats[i] = static_cast<u8>(regs.rt[i].format);
     }
-    alpha_test_ref = std::bit_cast<u32>(regs.alpha_test_ref);
+    // The reference is part of the pipeline only while the test compares with it: title 1
+    // animates the register with the test off (text fading in and out), which made a new
+    // pipeline, translated on the GPU thread, for every value (~84% of its runtime pipelines).
+    const u32 packed_alpha_test = alpha_test_func.Value();
+    alpha_test_ref = packed_alpha_test != PackComparisonOp(Maxwell::ComparisonOp::Never_GL) &&
+                             packed_alpha_test != PackComparisonOp(Maxwell::ComparisonOp::Always_GL)
+                         ? std::bit_cast<u32>(regs.alpha_test_ref)
+                         : 0;
     point_size = std::bit_cast<u32>(regs.point_size);
 
     if (maxwell3d.dirty.flags[Dirty::VertexInput]) {
diff --git a/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp b/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp
index f737ad9f70..94bace432c 100644
--- a/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp
+++ b/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp
@@ -31,7 +31,9 @@
 #include "common/settings.h"
 #include "video_core/nx2_profile.h"
 
-#if defined(_MSC_VER) && defined(NDEBUG)
+#if defined(__clang__) && defined(NDEBUG)
+#define LAMBDA_FORCEINLINE __attribute__((always_inline))
+#elif defined(_MSC_VER) && defined(NDEBUG)
 #define LAMBDA_FORCEINLINE [[msvc::forceinline]]
 #else
 #define LAMBDA_FORCEINLINE
diff --git a/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp b/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp
index 60bb040984..17f57983e3 100644
--- a/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp
+++ b/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp
@@ -768,10 +768,66 @@ void PipelineCache::QueueVulkanPipelineCacheFlush() {
     });
 }
 
+namespace {
+// EDEN_NX2_PIPELINE_DIFF=1: for a new graphics pipeline whose shaders an earlier one uses too,
+// log which parts of the fixed state differ from that one (finds state that splits the same
+// draws into many pipelines, each translated on the GPU thread).
+void ReportNewPipelineKey(const GraphicsPipelineCacheKey& key) {
+    static const bool enabled = std::getenv("EDEN_NX2_PIPELINE_DIFF") != nullptr;
+    if (!enabled) {
+        return;
+    }
+    static std::unordered_map<u64, GraphicsPipelineCacheKey> last_by_shaders;
+    static u32 logged{};
+    const u64 shaders = Common::CityHash64(reinterpret_cast<const char*>(key.unique_hashes.data()),
+                                           sizeof(key.unique_hashes));
+    const auto [it, first] = last_by_shaders.try_emplace(shaders, key);
+    if (first) {
+        return;
+    }
+    const FixedPipelineState& now = key.state;
+    const FixedPipelineState& before = it->second.state;
+    std::string diffs;
+    if (now.raw1 != before.raw1) diffs += fmt::format(" raw1^{:#x}", now.raw1 ^ before.raw1);
+    if (now.raw2 != before.raw2) diffs += fmt::format(" raw2^{:#x}", now.raw2 ^ before.raw2);
+    if (now.color_formats != before.color_formats) {
+        diffs += fmt::format(" color_formats {:02x} -> {:02x}", fmt::join(before.color_formats, ""),
+                             fmt::join(now.color_formats, ""));
+    }
+    if (now.alpha_test_ref != before.alpha_test_ref) {
+        diffs += fmt::format(" alpha_test_ref {:#x} -> {:#x} (func {})", before.alpha_test_ref,
+                             now.alpha_test_ref, now.alpha_test_func.Value());
+    }
+    if (now.point_size != before.point_size) {
+        diffs += fmt::format(" point_size {:#x} -> {:#x}", before.point_size, now.point_size);
+    }
+    if (now.viewport_swizzles != before.viewport_swizzles) diffs += " viewport_swizzles";
+    if (now.attribute_types != before.attribute_types) {
+        diffs += fmt::format(" attribute_types^{:#x}", now.attribute_types ^ before.attribute_types);
+    }
+    if (std::memcmp(&now.dynamic_state, &before.dynamic_state, sizeof(now.dynamic_state)) != 0) {
+        diffs += " dynamic_state";
+    }
+    if (now.driver_id != before.driver_id || now.driver_version != before.driver_version) {
+        diffs += " driver";
+    }
+    const size_t size = now.Size();
+    if (diffs.empty() && size > offsetof(FixedPipelineState, attachments)) {
+        diffs = fmt::format(" past attachments (key size {})", size);
+    }
+    it->second = key;
+    if (logged++ < 400) {
+        LOG_INFO(Render_Vulkan, "New pipeline for known shaders {:016x} differs in:{}", shaders,
+                 diffs.empty() ? std::string{" nothing in the key prefix"} : diffs);
+    }
+}
+} // Anonymous namespace
+
 GraphicsPipeline* PipelineCache::CurrentGraphicsPipelineSlowPath() {
     const auto [pair, is_new]{graphics_cache.try_emplace(graphics_key)};
     auto& pipeline{pair->second};
     if (is_new) {
+        ReportNewPipelineKey(graphics_key);
         pipeline = CreateGraphicsPipeline();
     }
     if (!pipeline) {
diff --git a/src/video_core/renderer_vulkan/vk_present_manager.cpp b/src/video_core/renderer_vulkan/vk_present_manager.cpp
index 9298764546..c1724e1ab4 100644
--- a/src/video_core/renderer_vulkan/vk_present_manager.cpp
+++ b/src/video_core/renderer_vulkan/vk_present_manager.cpp
@@ -434,7 +434,7 @@ void PresentManager::CopyToSwapchainImpl(Frame* frame) {
         RecreateSwapchain(frame);
     }
 
-    while (swapchain.AcquireNextImage()) {
+    while (swapchain.AcquireNextImage(!use_present_thread)) {
         RecreateSwapchain(frame);
     }
 
diff --git a/src/video_core/renderer_vulkan/vk_rasterizer.cpp b/src/video_core/renderer_vulkan/vk_rasterizer.cpp
index 429b2d63ef..0826820714 100644
--- a/src/video_core/renderer_vulkan/vk_rasterizer.cpp
+++ b/src/video_core/renderer_vulkan/vk_rasterizer.cpp
@@ -1250,20 +1250,37 @@ void RasterizerVulkan::AccelerateInlineToMemory(GPUVAddr address, size_t copy_si
         return;
     }
     gpu_memory->WriteBlockUnsafe(address, memory.data(), copy_size);
+    // EDEN_NX2_INLINE_DEFER=0: invalidate every cache at once (for comparisons).
+    static const bool defer = [] {
+        const char* value = std::getenv("EDEN_NX2_INLINE_DEFER");
+        return value == nullptr || value[0] != '0';
+    }();
     // One cache at a time: holding both while waiting for either (std::lock) waited longer, the
     // fencing thread and CPU write traps holding one of them most of the time.
+    bool uploaded;
     {
         std::scoped_lock lock{buffer_cache.mutex};
-        if (!buffer_cache.InlineMemory(*cpu_addr, copy_size, memory)) {
+        uploaded = buffer_cache.InlineMemory(*cpu_addr, copy_size, memory);
+        if (!uploaded && !defer) {
             buffer_cache.WriteMemory(*cpu_addr, copy_size);
         }
     }
+    // Pending query results must not land over the new words later: drop them now.
+    query_cache.InvalidateRegion(*cpu_addr, copy_size);
+    if (!uploaded && defer) {
+        // Title 1 makes thousands of small inline writes between draws (~10% of the GPU
+        // thread invalidating every cache for each). Guest memory is written; the buffer,
+        // texture and shader caches see the merged ranges at the next FlushCaching, which runs
+        // before anything of theirs is used. Writes into GPU-written buffer data were uploaded
+        // in place above, as before.
+        gpu_memory->AccumulateInvalidation(address, copy_size);
+        return;
+    }
     {
         std::scoped_lock lock{texture_cache.mutex};
         texture_cache.WriteMemory(*cpu_addr, copy_size);
     }
     pipeline_cache.InvalidateRegion(*cpu_addr, copy_size);
-    query_cache.InvalidateRegion(*cpu_addr, copy_size);
 }
 
 std::optional<FramebufferTextureInfo> RasterizerVulkan::AccelerateDisplay(
@@ -1309,8 +1326,10 @@ void RasterizerVulkan::FlushWork() {
     static constexpr u32 DRAWS_TO_DISPATCH = 512;
     static constexpr u32 CHECK_MASK = 3;
 #else
+    // Every 32 draws (was 8): each dispatch woke the worker thread, and Switch 2 titles make
+    // ~55,000 draws a second. A full chunk is dispatched on its own (Scheduler::Record).
     static constexpr u32 DRAWS_TO_DISPATCH = 4096;
-    static constexpr u32 CHECK_MASK = 7;
+    static constexpr u32 CHECK_MASK = 31;
 #endif // __ANDROID__
 
     static_assert(DRAWS_TO_DISPATCH % (CHECK_MASK + 1) == 0);
diff --git a/src/video_core/renderer_vulkan/vk_scheduler.cpp b/src/video_core/renderer_vulkan/vk_scheduler.cpp
index 4a7a7d5e78..e1e989c79f 100644
--- a/src/video_core/renderer_vulkan/vk_scheduler.cpp
+++ b/src/video_core/renderer_vulkan/vk_scheduler.cpp
@@ -55,6 +55,9 @@ Scheduler::Scheduler(const Device& device_, StateTracker& state_tracker_)
 Scheduler::~Scheduler() = default;
 
 u64 Scheduler::Flush(VkSemaphore signal_semaphore, VkSemaphore wait_semaphore) {
+    if (VideoCommon::GpuSplit::check_owner) [[unlikely]] {
+        VideoCommon::GpuSplit::CheckOwner("Scheduler::Flush");
+    }
     // When flushing, we only send data to the worker thread; no waiting is necessary.
     const u64 signal_value = SubmitExecution(signal_semaphore, wait_semaphore);
     AllocateNewContext();
@@ -62,6 +65,9 @@ u64 Scheduler::Flush(VkSemaphore signal_semaphore, VkSemaphore wait_semaphore) {
 }
 
 void Scheduler::Finish(VkSemaphore signal_semaphore, VkSemaphore wait_semaphore) {
+    if (VideoCommon::GpuSplit::check_owner) [[unlikely]] {
+        VideoCommon::GpuSplit::CheckOwner("Scheduler::Finish");
+    }
     // When finishing, we need to wait for the submission to have executed on the device.
     const u64 presubmit_tick = CurrentTick();
     SubmitExecution(signal_semaphore, wait_semaphore);
@@ -83,12 +89,21 @@ void Scheduler::WaitWorker() {
 }
 
 void Scheduler::DispatchWork() {
+    if (VideoCommon::GpuSplit::check_owner) [[unlikely]] {
+        VideoCommon::GpuSplit::CheckOwner("Scheduler::DispatchWork");
+    }
     if (chunk && !chunk->Empty()) {
+        // Waking the worker is a system call (~2% of the GPU thread at one dispatch per 8
+        // draws): skip it while the worker is still busy, it pops the queue before waiting.
+        bool wake;
         {
             std::scoped_lock ql{queue_mutex};
             work_queue.push(std::move(chunk));
+            wake = worker_waiting;
+        }
+        if (wake) {
+            event_cv.notify_all();
         }
-        event_cv.notify_all();
         AcquireNewChunk();
     }
 }
@@ -290,7 +305,11 @@ void Scheduler::WorkerThread(std::stop_token stop_token) {
             std::unique_lock lk{queue_mutex};
 
             // Wait for work.
-            event_cv.wait(lk, stop_token, [&] { return TryPopQueue(work); });
+            event_cv.wait(lk, stop_token, [&] {
+                worker_waiting = !TryPopQueue(work);
+                return !worker_waiting;
+            });
+            worker_waiting = false;
 
             // If we've been asked to stop, we're done.
             if (stop_token.stop_requested()) {
diff --git a/src/video_core/renderer_vulkan/vk_scheduler.h b/src/video_core/renderer_vulkan/vk_scheduler.h
index d4bc8d0ec5..568063ddbd 100644
--- a/src/video_core/renderer_vulkan/vk_scheduler.h
+++ b/src/video_core/renderer_vulkan/vk_scheduler.h
@@ -18,6 +18,7 @@
 #include "common/common_types.h"
 #include "common/polyfill_thread.h"
 #include "common/settings.h"
+#include "video_core/gpu_split.h"
 #include "video_core/renderer_vulkan/vk_master_semaphore.h"
 #include "video_core/vulkan_common/vulkan_wrapper.h"
 
@@ -103,6 +104,9 @@ public:
     template <typename T>
         requires std::is_invocable_v<T, vk::CommandBuffer, vk::CommandBuffer>
     void RecordWithUploadBuffer(T&& command) {
+        if (VideoCommon::GpuSplit::check_owner) [[unlikely]] {
+            VideoCommon::GpuSplit::CheckOwner("Scheduler::Record");
+        }
         if (chunk->Record(command)) {
             return;
         }
@@ -119,6 +123,18 @@ public:
             });
     }
 
+    /// Whether the recording thread submitted `tick` within 20 ms.
+    [[nodiscard]] bool WaitForSubmission(u64 tick) const {
+        const auto deadline = std::chrono::steady_clock::now() + std::chrono::milliseconds{20};
+        while (tick >= master_semaphore->CurrentTick()) {
+            if (std::chrono::steady_clock::now() >= deadline) {
+                return false;
+            }
+            std::this_thread::sleep_for(std::chrono::microseconds{100});
+        }
+        return true;
+    }
+
     /// Returns the current command buffer tick.
     [[nodiscard]] u64 CurrentTick() const noexcept {
         return master_semaphore->CurrentTick();
@@ -129,9 +145,13 @@ public:
         return master_semaphore->IsFree(tick);
     }
 
-    /// Waits for the given GPU tick, optionally pacing frames.
-    void Wait(u64 tick, double target_fps = 0.0) {
-        if (tick > 0) {
+    /// Waits for the given GPU tick, optionally pacing frames. Threads other than the one that
+    /// records (the async present thread) pass may_flush = false: flushing from there records
+    /// into the current chunk while that thread does (corrupt command buffers, crashes). They
+    /// wait for the recording thread to submit the tick instead, and give up after 20 ms (it may
+    /// be waiting for them).
+    void Wait(u64 tick, double target_fps = 0.0, bool may_flush = true) {
+        if (tick > 0 && (may_flush || WaitForSubmission(tick))) {
             if (tick >= master_semaphore->CurrentTick()) {
                 Flush();
             }
@@ -323,6 +343,7 @@ private:
     std::mutex reserve_mutex;
     std::mutex queue_mutex;
     std::condition_variable_any event_cv;
+    bool worker_waiting{}; ///< The worker is waiting for work (guarded by queue_mutex)
     std::jthread worker_thread;
 
     std::chrono::steady_clock::duration frame_interval{};
diff --git a/src/video_core/renderer_vulkan/vk_swapchain.cpp b/src/video_core/renderer_vulkan/vk_swapchain.cpp
index fea5398557..a0ed5a4493 100644
--- a/src/video_core/renderer_vulkan/vk_swapchain.cpp
+++ b/src/video_core/renderer_vulkan/vk_swapchain.cpp
@@ -158,7 +158,7 @@ void Swapchain::Create(
     resource_ticks.resize(image_count);
 }
 
-bool Swapchain::AcquireNextImage() {
+bool Swapchain::AcquireNextImage(bool may_flush) {
     const VkResult result = device.GetLogical().AcquireNextImageKHR(
         *swapchain, (std::numeric_limits<u64>::max)(), *present_semaphores[frame_index],
         VK_NULL_HANDLE, &image_index);
@@ -179,29 +179,29 @@ bool Swapchain::AcquireNextImage() {
         break;
     }
 
-    const auto wait_with_frame_pacing = [this] {
+    const auto wait_with_frame_pacing = [this, may_flush] {
     switch (Settings::values.frame_pacing_mode.GetValue()) {
     case Settings::FramePacingMode::Target_Auto:
-        scheduler.Wait(resource_ticks[image_index]);
+        scheduler.Wait(resource_ticks[image_index], 0.0, may_flush);
         break;
     case Settings::FramePacingMode::Target_30:
-        scheduler.Wait(resource_ticks[image_index], 30.0);
+        scheduler.Wait(resource_ticks[image_index], 30.0, may_flush);
         break;
     case Settings::FramePacingMode::Target_60:
-        scheduler.Wait(resource_ticks[image_index], 60.0);
+        scheduler.Wait(resource_ticks[image_index], 60.0, may_flush);
         break;
     case Settings::FramePacingMode::Target_90:
-        scheduler.Wait(resource_ticks[image_index], 90.0);
+        scheduler.Wait(resource_ticks[image_index], 90.0, may_flush);
         break;
     case Settings::FramePacingMode::Target_120:
-        scheduler.Wait(resource_ticks[image_index], 120.0);
+        scheduler.Wait(resource_ticks[image_index], 120.0, may_flush);
         break;
     }
     };
 
 #ifdef __ANDROID__
     if (android_get_device_api_level() >= 30) {
-        scheduler.Wait(resource_ticks[image_index]);
+        scheduler.Wait(resource_ticks[image_index], 0.0, may_flush);
     } else {
         wait_with_frame_pacing();
     }
diff --git a/src/video_core/renderer_vulkan/vk_swapchain.h b/src/video_core/renderer_vulkan/vk_swapchain.h
index d926cc118a..415160046c 100644
--- a/src/video_core/renderer_vulkan/vk_swapchain.h
+++ b/src/video_core/renderer_vulkan/vk_swapchain.h
@@ -37,7 +37,8 @@ public:
         u32 height);
 
     /// Acquires the next image in the swapchain, waits as needed.
-    bool AcquireNextImage();
+    /// `may_flush`: false on the async present thread (see Scheduler::Wait).
+    bool AcquireNextImage(bool may_flush = true);
 
     /// Presents the rendered image to the swapchain.
     void Present(VkSemaphore render_semaphore);
diff --git a/src/video_core/shader_environment.cpp b/src/video_core/shader_environment.cpp
index 96306d9a27..9d189ad4f0 100644
--- a/src/video_core/shader_environment.cpp
+++ b/src/video_core/shader_environment.cpp
@@ -458,7 +458,7 @@ void GenericEnvironment::InitializeNativeCode(u32 source_header_size) {
     Shader::Ampere::AppendNativeAddressMap(code, lowered.source_instruction_offsets);
     cached_lowest = 0;
     cached_highest = static_cast<u32>((code.size() - 1) * sizeof(u64));
-    LOG_INFO(Shader, "Lowered native shader at {:#x}: {} source bytes, {} instructions",
+    LOG_DEBUG(Shader, "Lowered native shader at {:#x}: {} source bytes, {} instructions",
              program_base, native_source_size, lowered.instructions.size());
 }
 
diff --git a/src/video_core/texture_cache/image_base.cpp b/src/video_core/texture_cache/image_base.cpp
index 9b0f33220d..3185dd8428 100644
--- a/src/video_core/texture_cache/image_base.cpp
+++ b/src/video_core/texture_cache/image_base.cpp
@@ -109,7 +109,9 @@ std::optional<SubresourceBase> ImageBase::TryFindBase(GPUVAddr other_addr) const
 }
 
 ImageViewId ImageBase::FindView(const ImageViewInfo& view_info) const noexcept {
-    const auto it = std::ranges::find(image_view_infos, view_info);
+    const auto it = std::ranges::find_if(image_view_infos, [&view_info](const ImageViewInfo& candidate) {
+        return candidate == view_info;
+    });
     if (it == image_view_infos.end()) {
         return ImageViewId{};
     }
diff --git a/src/video_core/texture_cache/texture_cache.h b/src/video_core/texture_cache/texture_cache.h
index fb8adef161..62cd523073 100644
--- a/src/video_core/texture_cache/texture_cache.h
+++ b/src/video_core/texture_cache/texture_cache.h
@@ -129,6 +129,12 @@ void TextureCache<P>::RunGarbageCollector() {
         high_priority_mode = total_used_memory >= expected_memory;
         aggressive_mode = allow_aggressive && total_used_memory >= critical_memory;
         ticks_to_destroy = aggressive_mode ? 10ULL : high_priority_mode ? 25ULL : 50ULL;
+        // EDEN_NX2_TC_MIN_AGE=<frames>: never evict images used more recently (thrash tests).
+        static const u64 min_age = [] {
+            const char* value = std::getenv("EDEN_NX2_TC_MIN_AGE");
+            return value != nullptr ? std::strtoull(value, nullptr, 10) : 0ULL;
+        }();
+        ticks_to_destroy = (std::max)(ticks_to_destroy, min_age);
         num_iterations = aggressive_mode ? 40 : (high_priority_mode ? 20 : 10);
     };
     const auto Cleanup = [this, &num_iterations, &high_priority_mode, &aggressive_mode](ImageId image_id) {
@@ -150,6 +156,7 @@ void TextureCache<P>::RunGarbageCollector() {
             image.DownloadMemory(map, copies);
             runtime.Finish();
             SwizzleImage(*gpu_memory, image.gpu_addr, image.info, copies, map.mapped_span, swizzle_data_buffer);
+            ++gc_download_deletes;
         }
         if (True(image.flags & ImageFlagBits::Tracked)) {
             UntrackImage(image, image_id);
@@ -253,7 +260,7 @@ void TextureCache<P>::TickFrame() {
                           largest.end(), [](const auto& a, const auto& b) { return a.first > b.first; });
         for (size_t i = 0; i < std::min<size_t>(12, largest.size()); ++i) {
             const Image& image = *largest[i].second;
-            LOG_INFO(HW_GPU, "  {} MiB: {:#x} {}x{}x{} layers {} levels {} format {} type {} flags {:#x}",
+            LOG_DEBUG(HW_GPU, "  {} MiB: {:#x} {}x{}x{} layers {} levels {} format {} type {} flags {:#x}",
                      largest[i].first >> 20, image.gpu_addr, image.info.size.width, image.info.size.height,
                      image.info.size.depth, image.info.resources.layers, image.info.resources.levels,
                      static_cast<u32>(image.info.format), static_cast<u32>(image.info.type),
@@ -263,14 +270,15 @@ void TextureCache<P>::TickFrame() {
         for (const auto& [format, total] : by_format) formats.emplace_back(total.first, format);
         std::sort(formats.rbegin(), formats.rend());
         for (size_t i = 0; i < std::min<size_t>(8, formats.size()); ++i) {
-            LOG_INFO(HW_GPU, "  format {}: {} MiB in {} images", formats[i].second, formats[i].first >> 20,
+            LOG_DEBUG(HW_GPU, "  format {}: {} MiB in {} images", formats[i].second, formats[i].first >> 20,
                      by_format[formats[i].second].second);
         }
         if constexpr (requires { runtime.LogMemoryStatistics(); }) {
             runtime.LogMemoryStatistics();
         }
-        LOG_INFO(HW_GPU, "Texture cache last 600 frames: created {} images, deleted {} (collector {}, {} costly), {} async decodes, ASTC {} opaque / {} alpha",
+        LOG_INFO(HW_GPU, "Texture cache last 600 frames: created {} images, deleted {} (collector {}, {} costly, {} downloaded), {} async decodes, ASTC {} opaque / {} alpha",
                  images_created, images_deleted, gc_deletes + gc_costly_deletes, gc_costly_deletes,
+                 gc_download_deletes,
                  async_decodes_queued, astc_opaque_images, astc_alpha_images);
         astc_opaque_images = astc_alpha_images = 0;
         std::vector<std::pair<u32, u64>> shapes;
@@ -278,13 +286,13 @@ void TextureCache<P>::TickFrame() {
         std::sort(shapes.rbegin(), shapes.rend());
         for (size_t i = 0; i < std::min<size_t>(10, shapes.size()); ++i) {
             const u64 key = shapes[i].second;
-            LOG_INFO(HW_GPU, "  created {}x: format {} type {} levels {} layers {} {}x{} (last at {:#x})",
+            LOG_DEBUG(HW_GPU, "  created {}x: format {} type {} levels {} layers {} {}x{} (last at {:#x})",
                      shapes[i].first, key >> 56, (key >> 52) & 0xF, (key >> 48) & 0xF,
                      (key >> 36) & 0xFFF, (key >> 18) & 0x3FFFF, key & 0x3FFFF,
                      created_shapes[key].second);
         }
         created_shapes.clear();
-        gc_deletes = gc_costly_deletes = async_decodes_queued = 0;
+        gc_deletes = gc_costly_deletes = gc_download_deletes = async_decodes_queued = 0;
         images_created = images_deleted = 0;
     }
     if (total_used_memory > minimum_memory) {
@@ -1755,7 +1763,7 @@ ImageId TextureCache<P>::InsertImage(const ImageInfo& info, GPUVAddr gpu_addr,
     storage_info.native_tiling = gpu_memory->UsesNativeTiling();
     if (storage_info.native_tiling && info.type == ImageType::e3D) {
         // Switch 2 3D images (lighting/fog volumes) are rare: log each one's shape.
-        LOG_INFO(HW_GPU, "3D image at {:#x}: {}x{}x{} levels {} format {} block {}x{}x{}",
+        LOG_DEBUG(HW_GPU, "3D image at {:#x}: {}x{}x{} levels {} format {} block {}x{}x{}",
                  gpu_addr, info.size.width, info.size.height, info.size.depth, info.resources.levels,
                  static_cast<u32>(info.format), info.block.width, info.block.height, info.block.depth);
     }
diff --git a/src/video_core/texture_cache/texture_cache_base.h b/src/video_core/texture_cache/texture_cache_base.h
index 11ac8a2421..2e375595de 100644
--- a/src/video_core/texture_cache/texture_cache_base.h
+++ b/src/video_core/texture_cache/texture_cache_base.h
@@ -125,6 +125,7 @@ class TextureCache : public VideoCommon::ChannelSetupCaches<TextureCacheChannelI
     // Periodic statistics (TickFrame).
     u64 gc_deletes{};
     u64 gc_costly_deletes{};
+    u64 gc_download_deletes{}; ///< Evictions that downloaded the image first (a full GPU wait)
     u64 async_decodes_queued{};
     u64 images_created{};
     bool native_thresholds_applied{};
diff --git a/src/video_core/vulkan_common/vulkan_memory_allocator.cpp b/src/video_core/vulkan_common/vulkan_memory_allocator.cpp
index 27f8fa6b4c..8bf692b00e 100644
--- a/src/video_core/vulkan_common/vulkan_memory_allocator.cpp
+++ b/src/video_core/vulkan_common/vulkan_memory_allocator.cpp
@@ -465,7 +465,7 @@ void LogMemoryStatistics(const Device& device) {
         if (s.blockCount == 0) {
             continue;
         }
-        LOG_INFO(Render_Vulkan, "VMA type {} (heap {}, flags {:#x}): {} blocks {} MiB, {} allocations {} MiB",
+        LOG_DEBUG(Render_Vulkan, "VMA type {} (heap {}, flags {:#x}): {} blocks {} MiB, {} allocations {} MiB",
                  type, properties->memoryTypes[type].heapIndex,
                  static_cast<u32>(properties->memoryTypes[type].propertyFlags), s.blockCount,
                  s.blockBytes >> 20, s.allocationCount, s.allocationBytes >> 20);
diff --git a/src/yuzu/CMakeLists.txt b/src/yuzu/CMakeLists.txt
index 62c2a2bd45..e6906a4ef0 100644
--- a/src/yuzu/CMakeLists.txt
+++ b/src/yuzu/CMakeLists.txt
@@ -432,8 +432,21 @@ if(UNIX AND NOT APPLE)
 endif()
 
 if (WIN32 AND NOT YUZU_USE_BUNDLED_QT)
-    set(YUZU_EXE_DIR "$<TARGET_FILE_DIR:yuzu>")
-    add_custom_command(TARGET yuzu POST_BUILD COMMAND ${WINDEPLOYQT_EXECUTABLE} "${YUZU_EXE_DIR}/eglantine.exe" --dir "${YUZU_EXE_DIR}" --libdir "${YUZU_EXE_DIR}" --plugindir "${YUZU_EXE_DIR}/plugins" --no-compiler-runtime --no-opengl-sw --no-system-d3d-compiler --no-translations --verbose 0)
+    # Copy the Qt DLLs and plugins next to the executable. WINDEPLOYQT_EXECUTABLE was never set, so
+    # this step ran the emulator itself with windeployqt's arguments and failed every build that
+    # uses an installed Qt (the Windows release workflow).
+    if (TARGET Qt6::windeployqt)
+        set(WINDEPLOYQT_EXECUTABLE "$<TARGET_FILE:Qt6::windeployqt>")
+    else()
+        find_program(WINDEPLOYQT_EXECUTABLE NAMES windeployqt6 windeployqt
+                     HINTS "${QT_HOST_PATH}/bin" "${QT_TARGET_PATH}/bin")
+    endif()
+    if (WINDEPLOYQT_EXECUTABLE)
+        set(YUZU_EXE_DIR "$<TARGET_FILE_DIR:yuzu>")
+        add_custom_command(TARGET yuzu POST_BUILD COMMAND ${WINDEPLOYQT_EXECUTABLE} "$<TARGET_FILE:yuzu>" --dir "${YUZU_EXE_DIR}" --libdir "${YUZU_EXE_DIR}" --plugindir "${YUZU_EXE_DIR}/plugins" --no-compiler-runtime --no-opengl-sw --no-system-d3d-compiler --no-translations --verbose 0)
+    else()
+        message(WARNING "windeployqt not found: copy the Qt DLLs next to the executable yourself")
+    endif()
 endif()
 
 if (ARCHITECTURE_x86_64 OR ARCHITECTURE_arm64)
diff --git a/src/yuzu_cmd/CMakeLists.txt b/src/yuzu_cmd/CMakeLists.txt
index ca8db523ee..b44cc586f5 100644
--- a/src/yuzu_cmd/CMakeLists.txt
+++ b/src/yuzu_cmd/CMakeLists.txt
@@ -58,8 +58,14 @@ endif()
 if(WIN32)
     # compile as a win32 gui application instead of a console application
     if(MSVC)
-        # UNICODE is defined, so SDL's entry point is wmain, not main
-        set_target_properties(yuzu-cmd PROPERTIES LINK_FLAGS_RELEASE "/SUBSYSTEM:WINDOWS /ENTRY:wmainCRTStartup")
+        # SDL's main implementation provides wmain when UNICODE is defined, which src/CMakeLists.txt
+        # does for cl.exe but not for clang-cl; forcing mainCRTStartup failed to link with cl.exe.
+        if(CXX_CLANG)
+            set(YUZU_CMD_ENTRY mainCRTStartup)
+        else()
+            set(YUZU_CMD_ENTRY wmainCRTStartup)
+        endif()
+        set_target_properties(yuzu-cmd PROPERTIES LINK_FLAGS_RELEASE "/SUBSYSTEM:WINDOWS /ENTRY:${YUZU_CMD_ENTRY}")
     elseif(MINGW)
         set_target_properties(yuzu-cmd PROPERTIES LINK_FLAGS_RELEASE "-Wl,--subsystem,windows")
     endif()
diff --git a/src/yuzu_cmd/nx2_sampler.cpp b/src/yuzu_cmd/nx2_sampler.cpp
index c01b812650..1488a3a078 100644
--- a/src/yuzu_cmd/nx2_sampler.cpp
+++ b/src/yuzu_cmd/nx2_sampler.cpp
@@ -240,9 +240,9 @@ void Report(const std::vector<Target>& targets, const std::vector<Sample>& sampl
             }
             std::sort(sorted.rbegin(), sorted.rend());
             std::fprintf(file, "-- %s\n", title);
-            for (size_t i = 0; i < std::min(limit, sorted.size()); ++i) {
+            for (size_t i = 0; i < (std::min)(limit, sorted.size()); ++i) {
                 std::fprintf(file, "%6u %5.1f%% %s\n", sorted[i].first,
-                             100.0 * sorted[i].first / std::max(total, 1U),
+                             100.0 * sorted[i].first / (std::max)(total, 1U),
                              sorted[i].second.c_str());
             }
         };
diff --git a/src/yuzu_cmd/yuzu.cpp b/src/yuzu_cmd/yuzu.cpp
index ffa1876c4f..3d9f50d4e5 100644
--- a/src/yuzu_cmd/yuzu.cpp
+++ b/src/yuzu_cmd/yuzu.cpp
@@ -434,7 +434,7 @@ static void RunKeyScript(InputCommon::InputSubsystem& input_subsystem) {
             }
             size_t start = 0;
             while (start < list.size()) {
-                const size_t end = std::min(list.find(',', start), list.size());
+                const size_t end = (std::min)(list.find(',', start), list.size());
                 const std::string item = list.substr(start, end - start);
                 if (const size_t colon = item.find(':'); colon != std::string::npos) {
                     const std::string when = item.substr(0, colon);