diff --git a/CMakeLists.txt b/CMakeLists.txt
index 09fddb3d7..d8fe5f68b 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -653,6 +653,7 @@ set(COMMON src/common/logging/backend.cpp
            src/common/arch.h
            src/common/assert.cpp
            src/common/assert.h
+           src/common/bit_array.h
            src/common/bit_field.h
            src/common/bounded_threadsafe_queue.h
            src/common/concepts.h
@@ -913,9 +914,10 @@ set(VIDEO_CORE src/video_core/amdgpu/liverpool.cpp
                src/video_core/buffer_cache/buffer.h
                src/video_core/buffer_cache/buffer_cache.cpp
                src/video_core/buffer_cache/buffer_cache.h
-               src/video_core/buffer_cache/memory_tracker_base.h
+               src/video_core/buffer_cache/memory_tracker.h
                src/video_core/buffer_cache/range_set.h
-               src/video_core/buffer_cache/word_manager.h
+               src/video_core/buffer_cache/region_definitions.h
+               src/video_core/buffer_cache/region_manager.h
                src/video_core/renderer_vulkan/liverpool_to_vk.cpp
                src/video_core/renderer_vulkan/liverpool_to_vk.h
                src/video_core/renderer_vulkan/vk_common.cpp
diff --git a/src/common/bit_array.h b/src/common/bit_array.h
new file mode 100644
index 000000000..f211bbf95
--- /dev/null
+++ b/src/common/bit_array.h
@@ -0,0 +1,411 @@
+// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
+// SPDX-License-Identifier: GPL-2.0-or-later
+
+#pragma once
+
+#include <array>
+#include <cstddef>
+#include "common/types.h"
+
+#ifdef __AVX2__
+#define BIT_ARRAY_USE_AVX
+#include <immintrin.h>
+#endif
+
+namespace Common {
+
+template <size_t N>
+class BitArray {
+    static_assert(N % 64 == 0, "BitArray size must be a multiple of 64 bits.");
+
+    static constexpr size_t BITS_PER_WORD = 64;
+    static constexpr size_t WORD_COUNT = N / BITS_PER_WORD;
+    static constexpr size_t WORDS_PER_AVX = 4;
+    static constexpr size_t AVX_WORD_COUNT = WORD_COUNT / WORDS_PER_AVX;
+
+public:
+    using Range = std::pair<size_t, size_t>;
+
+    class Iterator {
+    public:
+        explicit Iterator(const BitArray& bit_array_, u64 start) : bit_array(bit_array_) {
+            range = bit_array.FirstRangeFrom(start);
+        }
+
+        Iterator& operator++() {
+            range = bit_array.FirstRangeFrom(range.second);
+            return *this;
+        }
+
+        bool operator==(const Iterator& other) const {
+            return range == other.range;
+        }
+
+        bool operator!=(const Iterator& other) const {
+            return !(*this == other);
+        }
+
+        const Range& operator*() const {
+            return range;
+        }
+
+        const Range* operator->() const {
+            return &range;
+        }
+
+    private:
+        const BitArray& bit_array;
+        Range range;
+    };
+
+    using const_iterator = Iterator;
+    using iterator_category = std::forward_iterator_tag;
+    using value_type = Range;
+    using difference_type = std::ptrdiff_t;
+    using pointer = const Range*;
+    using reference = const Range&;
+
+    BitArray() = default;
+    BitArray(const BitArray& other) = default;
+    BitArray& operator=(const BitArray& other) = default;
+    BitArray(BitArray&& other) noexcept = default;
+    BitArray& operator=(BitArray&& other) noexcept = default;
+    ~BitArray() = default;
+
+    BitArray(const BitArray& other, size_t start, size_t end) {
+        if (start >= end || end > N) {
+            return;
+        }
+        const size_t first_word = start / BITS_PER_WORD;
+        const size_t last_word = (end - 1) / BITS_PER_WORD;
+        const size_t start_bit = start % BITS_PER_WORD;
+        const size_t end_bit = (end - 1) % BITS_PER_WORD;
+        const u64 start_mask = ~((1ULL << start_bit) - 1);
+        const u64 end_mask = end_bit == BITS_PER_WORD - 1 ? ~0ULL : (1ULL << (end_bit + 1)) - 1;
+        if (first_word == last_word) {
+            data[first_word] = other.data[first_word] & (start_mask & end_mask);
+        } else {
+            data[first_word] = other.data[first_word] & start_mask;
+            size_t i = first_word + 1;
+#ifdef BIT_ARRAY_USE_AVX
+            for (; i + WORDS_PER_AVX <= last_word; i += WORDS_PER_AVX) {
+                const __m256i current =
+                    _mm256_loadu_si256(reinterpret_cast<const __m256i*>(&other.data[i]));
+                _mm256_storeu_si256(reinterpret_cast<__m256i*>(&data[i]), current);
+            }
+#endif
+            for (; i < last_word; ++i) {
+                data[i] = other.data[i];
+            }
+            data[last_word] = other.data[last_word] & end_mask;
+        }
+    }
+
+    BitArray(const BitArray& other, const Range& range)
+        : BitArray(other, range.first, range.second) {}
+
+    const_iterator begin() const {
+        return Iterator(*this, 0);
+    }
+    const_iterator end() const {
+        return Iterator(*this, N);
+    }
+
+    inline constexpr void Set(size_t idx) {
+        data[idx / BITS_PER_WORD] |= (1ULL << (idx % BITS_PER_WORD));
+    }
+
+    inline constexpr void Unset(size_t idx) {
+        data[idx / BITS_PER_WORD] &= ~(1ULL << (idx % BITS_PER_WORD));
+    }
+
+    inline constexpr bool Get(size_t idx) const {
+        return (data[idx / BITS_PER_WORD] & (1ULL << (idx % BITS_PER_WORD))) != 0;
+    }
+
+    inline void SetRange(size_t start, size_t end) {
+        if (start >= end || end > N) {
+            return;
+        }
+        const size_t first_word = start / BITS_PER_WORD;
+        const size_t last_word = (end - 1) / BITS_PER_WORD;
+        const size_t start_bit = start % BITS_PER_WORD;
+        const size_t end_bit = (end - 1) % BITS_PER_WORD;
+        const u64 start_mask = ~((1ULL << start_bit) - 1);
+        const u64 end_mask = end_bit == BITS_PER_WORD - 1 ? ~0ULL : (1ULL << (end_bit + 1)) - 1;
+        if (first_word == last_word) {
+            data[first_word] |= start_mask & end_mask;
+        } else {
+            data[first_word] |= start_mask;
+            size_t i = first_word + 1;
+#ifdef BIT_ARRAY_USE_AVX
+            const __m256i value = _mm256_set1_epi64x(-1);
+            for (; i + WORDS_PER_AVX <= last_word; i += WORDS_PER_AVX) {
+                _mm256_storeu_si256(reinterpret_cast<__m256i*>(&data[i]), value);
+            }
+#endif
+            for (; i < last_word; ++i) {
+                data[i] = ~0ULL;
+            }
+            data[last_word] |= end_mask;
+        }
+    }
+
+    inline void UnsetRange(size_t start, size_t end) {
+        if (start >= end || end > N) {
+            return;
+        }
+        size_t first_word = start / BITS_PER_WORD;
+        const size_t last_word = (end - 1) / BITS_PER_WORD;
+        const size_t start_bit = start % BITS_PER_WORD;
+        const size_t end_bit = (end - 1) % BITS_PER_WORD;
+        const u64 start_mask = (1ULL << start_bit) - 1;
+        const u64 end_mask = end_bit == BITS_PER_WORD - 1 ? 0ULL : ~((1ULL << (end_bit + 1)) - 1);
+        if (first_word == last_word) {
+            data[first_word] &= start_mask | end_mask;
+        } else {
+            data[first_word] &= start_mask;
+            size_t i = first_word + 1;
+#ifdef BIT_ARRAY_USE_AVX
+            const __m256i value = _mm256_setzero_si256();
+            for (; i + WORDS_PER_AVX <= last_word; i += WORDS_PER_AVX) {
+                _mm256_storeu_si256(reinterpret_cast<__m256i*>(&data[i]), value);
+            }
+#endif
+            for (; i < last_word; ++i) {
+                data[i] = 0ULL;
+            }
+            data[last_word] &= end_mask;
+        }
+    }
+
+    inline constexpr void SetRange(const Range& range) {
+        SetRange(range.first, range.second);
+    }
+
+    inline constexpr void UnsetRange(const Range& range) {
+        UnsetRange(range.first, range.second);
+    }
+
+    inline constexpr void Clear() {
+        data.fill(0);
+    }
+
+    inline constexpr void Fill() {
+        data.fill(~0ULL);
+    }
+
+    inline constexpr bool None() const {
+        u64 result = 0;
+        for (const auto& word : data) {
+            result |= word;
+        }
+        return result == 0;
+    }
+
+    inline constexpr bool Any() const {
+        return !None();
+    }
+
+    Range FirstRangeFrom(size_t start) const {
+        if (start >= N) {
+            return {N, N};
+        }
+        const auto find_end_bit = [&](size_t word) {
+#ifdef BIT_ARRAY_USE_AVX
+            const __m256i all_one = _mm256_set1_epi64x(-1);
+            for (; word + WORDS_PER_AVX <= WORD_COUNT; word += WORDS_PER_AVX) {
+                const __m256i current =
+                    _mm256_loadu_si256(reinterpret_cast<const __m256i*>(&data[word]));
+                const __m256i cmp = _mm256_cmpeq_epi64(current, all_one);
+                if (_mm256_movemask_epi8(cmp) != 0xFFFFFFFF) {
+                    break;
+                }
+            }
+#endif
+            for (; word < WORD_COUNT; ++word) {
+                if (data[word] != ~0ULL) {
+                    return (word * BITS_PER_WORD) + std::countr_one(data[word]);
+                }
+            }
+            return N;
+        };
+
+        const auto word_bits = [&](size_t index, u64 word) {
+            const int empty_bits = std::countr_zero(word);
+            const int ones_count = std::countr_one(word >> empty_bits);
+            const size_t start_bit = index * BITS_PER_WORD + empty_bits;
+            if (ones_count + empty_bits < BITS_PER_WORD) {
+                return Range{start_bit, start_bit + ones_count};
+            }
+            return Range{start_bit, find_end_bit(index + 1)};
+        };
+
+        const size_t start_word = start / BITS_PER_WORD;
+        const size_t start_bit = start % BITS_PER_WORD;
+        const u64 masked_first = data[start_word] & (~((1ULL << start_bit) - 1));
+        if (masked_first) {
+            return word_bits(start_word, masked_first);
+        }
+
+        size_t word = start_word + 1;
+#ifdef BIT_ARRAY_USE_AVX
+        for (; word + WORDS_PER_AVX <= WORD_COUNT; word += WORDS_PER_AVX) {
+            const __m256i current =
+                _mm256_loadu_si256(reinterpret_cast<const __m256i*>(&data[word]));
+            if (!_mm256_testz_si256(current, current)) {
+                break;
+            }
+        }
+#endif
+        for (; word < WORD_COUNT; ++word) {
+            if (data[word] != 0) {
+                return word_bits(word, data[word]);
+            }
+        }
+        return {N, N};
+    }
+
+    inline constexpr Range FirstRange() const {
+        return FirstRangeFrom(0);
+    }
+
+    Range LastRangeFrom(size_t end) const {
+        if (end == 0) {
+            return {0, 0};
+        }
+        if (end > N) {
+            end = N;
+        }
+        const auto find_start_bit = [&](size_t word) {
+#ifdef BIT_ARRAY_USE_AVX
+            const __m256i all_zero = _mm256_setzero_si256();
+            for (; word >= WORDS_PER_AVX; word -= WORDS_PER_AVX) {
+                const __m256i current = _mm256_loadu_si256(
+                    reinterpret_cast<const __m256i*>(&data[word - WORDS_PER_AVX]));
+                const __m256i cmp = _mm256_cmpeq_epi64(current, all_zero);
+                if (_mm256_movemask_epi8(cmp) != 0xFFFFFFFF) {
+                    break;
+                }
+            }
+#endif
+            for (; word > 0; --word) {
+                if (data[word - 1] != ~0ULL) {
+                    return word * BITS_PER_WORD - std::countl_one(data[word - 1]);
+                }
+            }
+            return size_t(0);
+        };
+        const auto word_bits = [&](size_t index, u64 word) {
+            const int empty_bits = std::countl_zero(word);
+            const int ones_count = std::countl_one(word << empty_bits);
+            const size_t end_bit = index * BITS_PER_WORD - empty_bits;
+            if (empty_bits + ones_count < BITS_PER_WORD) {
+                return Range{end_bit - ones_count, end_bit};
+            }
+            return Range{find_start_bit(index - 1), end_bit};
+        };
+        const size_t end_word = ((end - 1) / BITS_PER_WORD) + 1;
+        const size_t end_bit = (end - 1) % BITS_PER_WORD;
+        u64 masked_last = data[end_word - 1];
+        if (end_bit < BITS_PER_WORD - 1) {
+            masked_last &= (1ULL << (end_bit + 1)) - 1;
+        }
+        if (masked_last) {
+            return word_bits(end_word, masked_last);
+        }
+        size_t word = end_word - 1;
+#ifdef BIT_ARRAY_USE_AVX
+        for (; word >= WORDS_PER_AVX; word -= WORDS_PER_AVX) {
+            const __m256i current =
+                _mm256_loadu_si256(reinterpret_cast<const __m256i*>(&data[word - WORDS_PER_AVX]));
+            if (!_mm256_testz_si256(current, current)) {
+                break;
+            }
+        }
+#endif
+        for (; word > 0; --word) {
+            if (data[word - 1] != 0) {
+                return word_bits(word, data[word - 1]);
+            }
+        }
+        return {0, 0};
+    }
+
+    inline constexpr Range LastRange() const {
+        return LastRangeFrom(N);
+    }
+
+    inline constexpr size_t Size() const {
+        return N;
+    }
+
+    inline constexpr BitArray& operator|=(const BitArray& other) {
+        for (size_t i = 0; i < WORD_COUNT; ++i) {
+            data[i] |= other.data[i];
+        }
+        return *this;
+    }
+
+    inline constexpr BitArray& operator&=(const BitArray& other) {
+        for (size_t i = 0; i < WORD_COUNT; ++i) {
+            data[i] &= other.data[i];
+        }
+        return *this;
+    }
+
+    inline constexpr BitArray& operator^=(const BitArray& other) {
+        for (size_t i = 0; i < WORD_COUNT; ++i) {
+            data[i] ^= other.data[i];
+        }
+        return *this;
+    }
+
+    inline constexpr BitArray& operator~() {
+        for (size_t i = 0; i < WORD_COUNT; ++i) {
+            data[i] = ~data[i];
+        }
+        return *this;
+    }
+
+    inline constexpr BitArray operator|(const BitArray& other) const {
+        BitArray result = *this;
+        result |= other;
+        return result;
+    }
+
+    inline constexpr BitArray operator&(const BitArray& other) const {
+        BitArray result = *this;
+        result &= other;
+        return result;
+    }
+
+    inline constexpr BitArray operator^(const BitArray& other) const {
+        BitArray result = *this;
+        result ^= other;
+        return result;
+    }
+
+    inline constexpr BitArray operator~() const {
+        BitArray result = *this;
+        result = ~result;
+        return result;
+    }
+
+    inline constexpr bool operator==(const BitArray& other) const {
+        u64 result = 0;
+        for (size_t i = 0; i < WORD_COUNT; ++i) {
+            result |= data[i] ^ other.data[i];
+        }
+        return result == 0;
+    }
+
+    inline constexpr bool operator!=(const BitArray& other) const {
+        return !(*this == other);
+    }
+
+private:
+    std::array<u64, WORD_COUNT> data{};
+};
+
+} // namespace Common
\ No newline at end of file
diff --git a/src/common/config.cpp b/src/common/config.cpp
index d8f46a17d..9c316949a 100644
--- a/src/common/config.cpp
+++ b/src/common/config.cpp
@@ -42,7 +42,6 @@ static std::string logFilter;
 static std::string logType = "sync";
 static std::string userName = "shadPS4";
 static std::string chooseHomeTab;
-static std::string backButtonBehavior = "left";
 static bool useSpecialPad = false;
 static int specialPadClass = 1;
 static bool isMotionControlsEnabled = true;
@@ -81,10 +80,6 @@ static std::vector<GameInstallDir> settings_install_dirs = {};
 std::vector<bool> install_dirs_enabled = {};
 std::filesystem::path settings_addon_install_dir = {};
 std::filesystem::path save_data_path = {};
-u32 mw_themes = 0;
-std::vector<std::string> m_elf_viewer;
-std::vector<std::string> m_recent_files;
-std::string emulator_language = "en_US";
 static bool isFullscreen = false;
 static std::string fullscreenMode = "Windowed";
 static bool isHDRAllowed = false;
@@ -209,10 +204,6 @@ std::string getChooseHomeTab() {
     return chooseHomeTab;
 }
 
-std::string getBackButtonBehavior() {
-    return backButtonBehavior;
-}
-
 bool getUseSpecialPad() {
     return useSpecialPad;
 }
@@ -428,10 +419,6 @@ void setChooseHomeTab(const std::string& type) {
     chooseHomeTab = type;
 }
 
-void setBackButtonBehavior(const std::string& type) {
-    backButtonBehavior = type;
-}
-
 void setUseSpecialPad(bool use) {
     useSpecialPad = use;
 }
@@ -484,24 +471,6 @@ void setAddonInstallDir(const std::filesystem::path& dir) {
     settings_addon_install_dir = dir;
 }
 
-void setMainWindowTheme(u32 theme) {
-    mw_themes = theme;
-}
-
-void setElfViewer(const std::vector<std::string>& elfList) {
-    m_elf_viewer.resize(elfList.size());
-    m_elf_viewer = elfList;
-}
-
-void setRecentFiles(const std::vector<std::string>& recentFiles) {
-    m_recent_files.resize(recentFiles.size());
-    m_recent_files = recentFiles;
-}
-
-void setEmulatorLanguage(std::string language) {
-    emulator_language = language;
-}
-
 void setGameInstallDirs(const std::vector<std::filesystem::path>& dirs_config) {
     settings_install_dirs.clear();
     for (const auto& dir : dirs_config) {
@@ -543,22 +512,6 @@ std::filesystem::path getAddonInstallDir() {
     return settings_addon_install_dir;
 }
 
-u32 getMainWindowTheme() {
-    return mw_themes;
-}
-
-std::vector<std::string> getElfViewer() {
-    return m_elf_viewer;
-}
-
-std::vector<std::string> getRecentFiles() {
-    return m_recent_files;
-}
-
-std::string getEmulatorLanguage() {
-    return emulator_language;
-}
-
 u32 GetLanguage() {
     return m_language;
 }
@@ -620,7 +573,6 @@ void load(const std::filesystem::path& path) {
 
         cursorState = toml::find_or<int>(input, "cursorState", HideCursorState::Idle);
         cursorHideTimeout = toml::find_or<int>(input, "cursorHideTimeout", 5);
-        backButtonBehavior = toml::find_or<std::string>(input, "backButtonBehavior", "left");
         useSpecialPad = toml::find_or<bool>(input, "useSpecialPad", false);
         specialPadClass = toml::find_or<int>(input, "specialPadClass", 1);
         isMotionControlsEnabled = toml::find_or<bool>(input, "isMotionControlsEnabled", true);
@@ -668,7 +620,6 @@ void load(const std::filesystem::path& path) {
         const toml::value& gui = data.at("GUI");
 
         load_game_size = toml::find_or<bool>(gui, "loadGameSizeEnabled", true);
-        mw_themes = toml::find_or<int>(gui, "theme", 0);
 
         const auto install_dir_array =
             toml::find_or<std::vector<std::u8string>>(gui, "installDirs", {});
@@ -693,9 +644,6 @@ void load(const std::filesystem::path& path) {
         save_data_path = toml::find_fs_path_or(gui, "saveDataPath", {});
 
         settings_addon_install_dir = toml::find_fs_path_or(gui, "addonInstallDir", {});
-        m_elf_viewer = toml::find_or<std::vector<std::string>>(gui, "elfDirs", {});
-        m_recent_files = toml::find_or<std::vector<std::string>>(gui, "recentFiles", {});
-        emulator_language = toml::find_or<std::string>(gui, "emulatorLanguage", "en_US");
     }
 
     if (data.contains("Settings")) {
@@ -708,19 +656,6 @@ void load(const std::filesystem::path& path) {
         const toml::value& keys = data.at("Keys");
         trophyKey = toml::find_or<std::string>(keys, "TrophyKey", "");
     }
-
-    // Check if the loaded language is in the allowed list
-    const std::vector<std::string> allowed_languages = {
-        "ar_SA", "da_DK", "de_DE", "el_GR", "en_US", "es_ES", "fa_IR", "fi_FI",
-        "fr_FR", "hu_HU", "id_ID", "it_IT", "ja_JP", "ko_KR", "lt_LT", "nb_NO",
-        "nl_NL", "pl_PL", "pt_BR", "pt_PT", "ro_RO", "ru_RU", "sq_AL", "sv_SE",
-        "tr_TR", "uk_UA", "vi_VN", "zh_CN", "zh_TW", "ca_ES", "sr_CS"};
-
-    if (std::find(allowed_languages.begin(), allowed_languages.end(), emulator_language) ==
-        allowed_languages.end()) {
-        emulator_language = "en_US"; // Default to en_US if not in the list
-        save(path);
-    }
 }
 
 void sortTomlSections(toml::ordered_value& data) {
@@ -792,7 +727,6 @@ void save(const std::filesystem::path& path) {
     data["General"]["checkCompatibilityOnStartup"] = checkCompatibilityOnStartup;
     data["Input"]["cursorState"] = cursorState;
     data["Input"]["cursorHideTimeout"] = cursorHideTimeout;
-    data["Input"]["backButtonBehavior"] = backButtonBehavior;
     data["Input"]["useSpecialPad"] = useSpecialPad;
     data["Input"]["specialPadClass"] = specialPadClass;
     data["Input"]["isMotionControlsEnabled"] = isMotionControlsEnabled;
@@ -855,7 +789,6 @@ void save(const std::filesystem::path& path) {
 
     data["GUI"]["addonInstallDir"] =
         std::string{fmt::UTF(settings_addon_install_dir.u8string()).data};
-    data["GUI"]["emulatorLanguage"] = emulator_language;
     data["Settings"]["consoleLanguage"] = m_language;
 
     // Sorting of TOML sections
@@ -864,42 +797,6 @@ void save(const std::filesystem::path& path) {
     std::ofstream file(path, std::ios::binary);
     file << data;
     file.close();
-
-    saveMainWindow(path);
-}
-
-void saveMainWindow(const std::filesystem::path& path) {
-    toml::ordered_value data;
-
-    std::error_code error;
-    if (std::filesystem::exists(path, error)) {
-        try {
-            std::ifstream ifs;
-            ifs.exceptions(std::ifstream::failbit | std::ifstream::badbit);
-            ifs.open(path, std::ios_base::binary);
-            data = toml::parse<toml::ordered_type_config>(
-                ifs, std::string{fmt::UTF(path.filename().u8string()).data});
-        } catch (const std::exception& ex) {
-            fmt::print("Exception trying to parse config file. Exception: {}\n", ex.what());
-            return;
-        }
-    } else {
-        if (error) {
-            fmt::print("Filesystem error: {}\n", error.message());
-        }
-        fmt::print("Saving new configuration file {}\n", fmt::UTF(path.u8string()));
-    }
-
-    data["GUI"]["theme"] = mw_themes;
-    data["GUI"]["elfDirs"] = m_elf_viewer;
-    data["GUI"]["recentFiles"] = m_recent_files;
-
-    // Sorting of TOML sections
-    sortTomlSections(data);
-
-    std::ofstream file(path, std::ios::binary);
-    file << data;
-    file.close();
 }
 
 void setDefaultValues() {
@@ -920,7 +817,6 @@ void setDefaultValues() {
     cursorState = HideCursorState::Idle;
     cursorHideTimeout = 5;
     trophyNotificationDuration = 6.0;
-    backButtonBehavior = "left";
     useSpecialPad = false;
     specialPadClass = 1;
     isDebugDump = false;
@@ -937,7 +833,6 @@ void setDefaultValues() {
     vkHostMarkers = false;
     vkGuestMarkers = false;
     rdocEnable = false;
-    emulator_language = "en_US";
     m_language = 1;
     gpuId = -1;
     compatibilityData = false;
@@ -967,7 +862,7 @@ l3 = x
 r3 = m
 
 options = enter
-touchpad = space
+touchpad_center = space
 
 pad_up = up
 pad_down = down
@@ -999,7 +894,7 @@ r2 = r2
 r3 = r3
 
 options = options
-touchpad = back
+touchpad_center = back
 
 pad_up = pad_up
 pad_down = pad_down
diff --git a/src/common/config.h b/src/common/config.h
index 414bc122e..38114983f 100644
--- a/src/common/config.h
+++ b/src/common/config.h
@@ -18,77 +18,96 @@ enum HideCursorState : int { Never, Idle, Always };
 
 void load(const std::filesystem::path& path);
 void save(const std::filesystem::path& path);
-void saveMainWindow(const std::filesystem::path& path);
 
 std::string getTrophyKey();
 void setTrophyKey(std::string key);
+bool getIsFullscreen();
+void setIsFullscreen(bool enable);
+std::string getFullscreenMode();
+void setFullscreenMode(std::string mode);
+u32 getScreenWidth();
+u32 getScreenHeight();
+void setScreenWidth(u32 width);
+void setScreenHeight(u32 height);
+bool debugDump();
+void setDebugDump(bool enable);
+s32 getGpuId();
+void setGpuId(s32 selectedGpuId);
+bool allowHDR();
+void setAllowHDR(bool enable);
+bool collectShadersForDebug();
+void setCollectShaderForDebug(bool enable);
+bool showSplash();
+void setShowSplash(bool enable);
+std::string sideTrophy();
+void setSideTrophy(std::string side);
+bool nullGpu();
+void setNullGpu(bool enable);
+bool copyGPUCmdBuffers();
+void setCopyGPUCmdBuffers(bool enable);
+bool dumpShaders();
+void setDumpShaders(bool enable);
+u32 vblankDiv();
+void setVblankDiv(u32 value);
+bool getisTrophyPopupDisabled();
+void setisTrophyPopupDisabled(bool disable);
+s16 getCursorState();
+void setCursorState(s16 cursorState);
+bool vkValidationEnabled();
+void setVkValidation(bool enable);
+bool vkValidationSyncEnabled();
+void setVkSyncValidation(bool enable);
+bool getVkCrashDiagnosticEnabled();
+void setVkCrashDiagnosticEnabled(bool enable);
+bool getVkHostMarkersEnabled();
+void setVkHostMarkersEnabled(bool enable);
+bool getVkGuestMarkersEnabled();
+void setVkGuestMarkersEnabled(bool enable);
+bool getEnableDiscordRPC();
+void setEnableDiscordRPC(bool enable);
+bool isRdocEnabled();
+void setRdocEnabled(bool enable);
+std::string getLogType();
+void setLogType(const std::string& type);
+std::string getLogFilter();
+void setLogFilter(const std::string& type);
+double getTrophyNotificationDuration();
+void setTrophyNotificationDuration(double newTrophyNotificationDuration);
+int getCursorHideTimeout();
+void setCursorHideTimeout(int newcursorHideTimeout);
+void setSeparateLogFilesEnabled(bool enabled);
+bool getSeparateLogFilesEnabled();
+u32 GetLanguage();
+void setLanguage(u32 language);
+void setUseSpecialPad(bool use);
+bool getUseSpecialPad();
+void setSpecialPadClass(int type);
+int getSpecialPadClass();
+bool getPSNSignedIn();
+void setPSNSignedIn(bool sign); // no ui setting
+bool patchShaders();            // no set
+bool fpsColor();                // no set
+bool isNeoModeConsole();
+void setNeoMode(bool enable);  // no ui setting
+bool isDevKitConsole();        // no set
+bool vkValidationGpuEnabled(); // no set
+bool getIsMotionControlsEnabled();
+void setIsMotionControlsEnabled(bool use);
+
+// TODO
 bool GetLoadGameSizeEnabled();
 std::filesystem::path GetSaveDataPath();
 void setLoadGameSizeEnabled(bool enable);
-bool getIsFullscreen();
-std::string getFullscreenMode();
-bool isNeoModeConsole();
-bool isDevKitConsole();
-bool getisTrophyPopupDisabled();
-bool getEnableDiscordRPC();
 bool getCompatibilityEnabled();
 bool getCheckCompatibilityOnStartup();
-bool getPSNSignedIn();
-
-std::string getLogFilter();
-std::string getLogType();
 std::string getUserName();
 std::string getChooseHomeTab();
-
-s16 getCursorState();
-int getCursorHideTimeout();
-double getTrophyNotificationDuration();
-std::string getBackButtonBehavior();
-bool getUseSpecialPad();
-int getSpecialPadClass();
-bool getIsMotionControlsEnabled();
 bool GetUseUnifiedInputConfig();
 void SetUseUnifiedInputConfig(bool use);
 bool GetOverrideControllerColor();
 void SetOverrideControllerColor(bool enable);
 int* GetControllerCustomColor();
 void SetControllerCustomColor(int r, int b, int g);
-
-u32 getScreenWidth();
-u32 getScreenHeight();
-s32 getGpuId();
-bool allowHDR();
-
-bool debugDump();
-bool collectShadersForDebug();
-bool showSplash();
-std::string sideTrophy();
-bool nullGpu();
-bool copyGPUCmdBuffers();
-bool dumpShaders();
-bool patchShaders();
-bool isRdocEnabled();
-bool fpsColor();
-u32 vblankDiv();
-
-void setDebugDump(bool enable);
-void setCollectShaderForDebug(bool enable);
-void setShowSplash(bool enable);
-void setSideTrophy(std::string side);
-void setNullGpu(bool enable);
-void setAllowHDR(bool enable);
-void setCopyGPUCmdBuffers(bool enable);
-void setDumpShaders(bool enable);
-void setVblankDiv(u32 value);
-void setGpuId(s32 selectedGpuId);
-void setScreenWidth(u32 width);
-void setScreenHeight(u32 height);
-void setIsFullscreen(bool enable);
-void setFullscreenMode(std::string mode);
-void setisTrophyPopupDisabled(bool disable);
-void setEnableDiscordRPC(bool enable);
-void setLanguage(u32 language);
-void setNeoMode(bool enable);
 void setUserName(const std::string& type);
 void setChooseHomeTab(const std::string& type);
 void setGameInstallDirs(const std::vector<std::filesystem::path>& dirs_config);
@@ -96,57 +115,19 @@ void setAllGameInstallDirs(const std::vector<GameInstallDir>& dirs_config);
 void setSaveDataPath(const std::filesystem::path& path);
 void setCompatibilityEnabled(bool use);
 void setCheckCompatibilityOnStartup(bool use);
-void setPSNSignedIn(bool sign);
-
-void setCursorState(s16 cursorState);
-void setCursorHideTimeout(int newcursorHideTimeout);
-void setTrophyNotificationDuration(double newTrophyNotificationDuration);
-void setBackButtonBehavior(const std::string& type);
-void setUseSpecialPad(bool use);
-void setSpecialPadClass(int type);
-void setIsMotionControlsEnabled(bool use);
-
-void setLogType(const std::string& type);
-void setLogFilter(const std::string& type);
-void setSeparateLogFilesEnabled(bool enabled);
-bool getSeparateLogFilesEnabled();
-void setVkValidation(bool enable);
-void setVkSyncValidation(bool enable);
-void setRdocEnabled(bool enable);
-
-bool vkValidationEnabled();
-bool vkValidationSyncEnabled();
-bool vkValidationGpuEnabled();
-bool getVkCrashDiagnosticEnabled();
-bool getVkHostMarkersEnabled();
-bool getVkGuestMarkersEnabled();
-void setVkCrashDiagnosticEnabled(bool enable);
-void setVkHostMarkersEnabled(bool enable);
-void setVkGuestMarkersEnabled(bool enable);
-
 // Gui
 bool addGameInstallDir(const std::filesystem::path& dir, bool enabled = true);
 void removeGameInstallDir(const std::filesystem::path& dir);
 void setGameInstallDirEnabled(const std::filesystem::path& dir, bool enabled);
 void setAddonInstallDir(const std::filesystem::path& dir);
-void setMainWindowTheme(u32 theme);
-void setElfViewer(const std::vector<std::string>& elfList);
-void setRecentFiles(const std::vector<std::string>& recentFiles);
-void setEmulatorLanguage(std::string language);
 
 const std::vector<std::filesystem::path> getGameInstallDirs();
 const std::vector<bool> getGameInstallDirsEnabled();
 std::filesystem::path getAddonInstallDir();
-u32 getMainWindowTheme();
-std::vector<std::string> getElfViewer();
-std::vector<std::string> getRecentFiles();
-std::string getEmulatorLanguage();
 
 void setDefaultValues();
 
 // todo: name and function location pending
 std::filesystem::path GetFoolproofKbmConfigFile(const std::string& game_id = "");
 
-// settings
-u32 GetLanguage();
 }; // namespace Config
diff --git a/src/core/libraries/pad/pad.cpp b/src/core/libraries/pad/pad.cpp
index 42582783b..59964fa58 100644
--- a/src/core/libraries/pad/pad.cpp
+++ b/src/core/libraries/pad/pad.cpp
@@ -447,21 +447,18 @@ int PS4_SYSV_ABI scePadReadState(s32 handle, OrbisPadData* pData) {
 
     // Only do this on handle 1 for now
     if (engine && handle == 1) {
-        const auto gyro_poll_rate = engine->GetAccelPollRate();
-        if (gyro_poll_rate != 0.0f) {
-            auto now = std::chrono::steady_clock::now();
-            float deltaTime = std::chrono::duration_cast<std::chrono::microseconds>(
-                                  now - controller->GetLastUpdate())
-                                  .count() /
-                              1000000.0f;
-            controller->SetLastUpdate(now);
-            Libraries::Pad::OrbisFQuaternion lastOrientation = controller->GetLastOrientation();
-            Libraries::Pad::OrbisFQuaternion outputOrientation = {0.0f, 0.0f, 0.0f, 1.0f};
-            GameController::CalculateOrientation(pData->acceleration, pData->angularVelocity,
-                                                 deltaTime, lastOrientation, outputOrientation);
-            pData->orientation = outputOrientation;
-            controller->SetLastOrientation(outputOrientation);
-        }
+        auto now = std::chrono::steady_clock::now();
+        float deltaTime =
+            std::chrono::duration_cast<std::chrono::microseconds>(now - controller->GetLastUpdate())
+                .count() /
+            1000000.0f;
+        controller->SetLastUpdate(now);
+        Libraries::Pad::OrbisFQuaternion lastOrientation = controller->GetLastOrientation();
+        Libraries::Pad::OrbisFQuaternion outputOrientation = {0.0f, 0.0f, 0.0f, 1.0f};
+        GameController::CalculateOrientation(pData->acceleration, pData->angularVelocity, deltaTime,
+                                             lastOrientation, outputOrientation);
+        pData->orientation = outputOrientation;
+        controller->SetLastOrientation(outputOrientation);
     }
     pData->touchData.touchNum =
         (state.touchpad[0].state ? 1 : 0) + (state.touchpad[1].state ? 1 : 0);
diff --git a/src/core/libraries/save_data/save_instance.cpp b/src/core/libraries/save_data/save_instance.cpp
index a7ce3d35f..05253eb23 100644
--- a/src/core/libraries/save_data/save_instance.cpp
+++ b/src/core/libraries/save_data/save_instance.cpp
@@ -22,25 +22,25 @@ static Core::FileSys::MntPoints* g_mnt = Common::Singleton<Core::FileSys::MntPoi
 namespace fs = std::filesystem;
 
 // clang-format off
-static const std::unordered_map<std::string, std::string> default_title = {
-    {"ja_JP", "セーブデータ"},
-    {"en_US", "Saved Data"},
-    {"fr_FR", "Données sauvegardées"},
-    {"es_ES", "Datos guardados"},
-    {"de_DE", "Gespeicherte Daten"},
-    {"it_IT", "Dati salvati"},
-    {"nl_NL", "Opgeslagen data"},
-    {"pt_PT", "Dados guardados"},
-    {"ru_RU", "Сохраненные данные"},
-    {"ko_KR", "저장 데이터"},
-    {"zh_CN", "保存数据"},
-    {"fi_FI", "Tallennetut tiedot"},
-    {"sv_SE", "Sparade data"},
-    {"da_DK", "Gemte data"},
-    {"no_NO", "Lagrede data"},
-    {"pl_PL", "Zapisane dane"},
-    {"pt_BR", "Dados salvos"},
-    {"tr_TR", "Kayıtlı Veriler"},
+static const std::unordered_map<int, std::string> default_title = {
+    {0/*"ja_JP"*/, "セーブデータ"},
+    {1/*"en_US"*/, "Saved Data"},
+    {2/*"fr_FR"*/, "Données sauvegardées"},
+    {3/*"es_ES"*/, "Datos guardados"},
+    {4/*"de_DE"*/, "Gespeicherte Daten"},
+    {5/*"it_IT"*/, "Dati salvati"},
+    {6/*"nl_NL"*/, "Opgeslagen data"},
+    {7/*"pt_PT"*/, "Dados guardados"},
+    {8/*"ru_RU"*/, "Сохраненные данные"},
+    {9/*"ko_KR"*/, "저장 데이터"},
+    {10/*"zh_CN"*/, "保存数据"},
+    {12/*"fi_FI"*/, "Tallennetut tiedot"},
+    {13/*"sv_SE"*/, "Sparade data"},
+    {14/*"da_DK"*/, "Gemte data"},
+    {15/*"no_NO"*/, "Lagrede data"},
+    {16/*"pl_PL"*/, "Zapisane dane"},
+    {17/*"pt_BR"*/, "Dados salvos"},
+    {19/*"tr_TR"*/, "Kayıtlı Veriler"},
 };
 // clang-format on
 
@@ -71,9 +71,9 @@ fs::path SaveInstance::GetParamSFOPath(const fs::path& dir_path) {
 
 void SaveInstance::SetupDefaultParamSFO(PSF& param_sfo, std::string dir_name,
                                         std::string game_serial) {
-    std::string locale = Config::getEmulatorLanguage();
+    int locale = Config::GetLanguage();
     if (!default_title.contains(locale)) {
-        locale = "en_US";
+        locale = 1; // default to en_US if not found
     }
 
 #define P(type, key, ...) param_sfo.Add##type(std::string{key}, __VA_ARGS__)
diff --git a/src/emulator.cpp b/src/emulator.cpp
index f50147818..99fd50af5 100644
--- a/src/emulator.cpp
+++ b/src/emulator.cpp
@@ -58,10 +58,7 @@ Emulator::Emulator() {
 #endif
 }
 
-Emulator::~Emulator() {
-    const auto config_dir = Common::FS::GetUserPath(Common::FS::PathType::UserDir);
-    Config::saveMainWindow(config_dir / "config.toml");
-}
+Emulator::~Emulator() {}
 
 void Emulator::Run(std::filesystem::path file, const std::vector<std::string> args) {
     if (std::filesystem::is_directory(file)) {
diff --git a/src/input/input_handler.cpp b/src/input/input_handler.cpp
index ab0dbe21c..7c4e19103 100644
--- a/src/input/input_handler.cpp
+++ b/src/input/input_handler.cpp
@@ -66,22 +66,25 @@ auto output_array = std::array{
     ControllerOutput(LEFTJOYSTICK_HALFMODE),
     ControllerOutput(RIGHTJOYSTICK_HALFMODE),
     ControllerOutput(KEY_TOGGLE),
+    ControllerOutput(MOUSE_GYRO_ROLL_MODE),
 
     // Button mappings
-    ControllerOutput(SDL_GAMEPAD_BUTTON_NORTH),          // Triangle
-    ControllerOutput(SDL_GAMEPAD_BUTTON_EAST),           // Circle
-    ControllerOutput(SDL_GAMEPAD_BUTTON_SOUTH),          // Cross
-    ControllerOutput(SDL_GAMEPAD_BUTTON_WEST),           // Square
-    ControllerOutput(SDL_GAMEPAD_BUTTON_LEFT_SHOULDER),  // L1
-    ControllerOutput(SDL_GAMEPAD_BUTTON_LEFT_STICK),     // L3
-    ControllerOutput(SDL_GAMEPAD_BUTTON_RIGHT_SHOULDER), // R1
-    ControllerOutput(SDL_GAMEPAD_BUTTON_RIGHT_STICK),    // R3
-    ControllerOutput(SDL_GAMEPAD_BUTTON_START),          // Options
-    ControllerOutput(SDL_GAMEPAD_BUTTON_TOUCHPAD),       // TouchPad
-    ControllerOutput(SDL_GAMEPAD_BUTTON_DPAD_UP),        // Up
-    ControllerOutput(SDL_GAMEPAD_BUTTON_DPAD_DOWN),      // Down
-    ControllerOutput(SDL_GAMEPAD_BUTTON_DPAD_LEFT),      // Left
-    ControllerOutput(SDL_GAMEPAD_BUTTON_DPAD_RIGHT),     // Right
+    ControllerOutput(SDL_GAMEPAD_BUTTON_NORTH),           // Triangle
+    ControllerOutput(SDL_GAMEPAD_BUTTON_EAST),            // Circle
+    ControllerOutput(SDL_GAMEPAD_BUTTON_SOUTH),           // Cross
+    ControllerOutput(SDL_GAMEPAD_BUTTON_WEST),            // Square
+    ControllerOutput(SDL_GAMEPAD_BUTTON_LEFT_SHOULDER),   // L1
+    ControllerOutput(SDL_GAMEPAD_BUTTON_LEFT_STICK),      // L3
+    ControllerOutput(SDL_GAMEPAD_BUTTON_RIGHT_SHOULDER),  // R1
+    ControllerOutput(SDL_GAMEPAD_BUTTON_RIGHT_STICK),     // R3
+    ControllerOutput(SDL_GAMEPAD_BUTTON_START),           // Options
+    ControllerOutput(SDL_GAMEPAD_BUTTON_TOUCHPAD_LEFT),   // TouchPad
+    ControllerOutput(SDL_GAMEPAD_BUTTON_TOUCHPAD_CENTER), // TouchPad
+    ControllerOutput(SDL_GAMEPAD_BUTTON_TOUCHPAD_RIGHT),  // TouchPad
+    ControllerOutput(SDL_GAMEPAD_BUTTON_DPAD_UP),         // Up
+    ControllerOutput(SDL_GAMEPAD_BUTTON_DPAD_DOWN),       // Down
+    ControllerOutput(SDL_GAMEPAD_BUTTON_DPAD_LEFT),       // Left
+    ControllerOutput(SDL_GAMEPAD_BUTTON_DPAD_RIGHT),      // Right
 
     // Axis mappings
     // ControllerOutput(SDL_GAMEPAD_BUTTON_INVALID, SDL_GAMEPAD_AXIS_LEFTX, false),
@@ -130,6 +133,12 @@ static OrbisPadButtonDataOffset SDLGamepadToOrbisButton(u8 button) {
         return OPBDO::Options;
     case SDL_GAMEPAD_BUTTON_TOUCHPAD:
         return OPBDO::TouchPad;
+    case SDL_GAMEPAD_BUTTON_TOUCHPAD_LEFT:
+        return OPBDO::TouchPad;
+    case SDL_GAMEPAD_BUTTON_TOUCHPAD_CENTER:
+        return OPBDO::TouchPad;
+    case SDL_GAMEPAD_BUTTON_TOUCHPAD_RIGHT:
+        return OPBDO::TouchPad;
     case SDL_GAMEPAD_BUTTON_BACK:
         return OPBDO::TouchPad;
     case SDL_GAMEPAD_BUTTON_LEFT_SHOULDER:
@@ -502,14 +511,21 @@ void ControllerOutput::FinalizeUpdate() {
     }
     old_button_state = new_button_state;
     old_param = *new_param;
-    float touchpad_x = 0;
     if (button != SDL_GAMEPAD_BUTTON_INVALID) {
         switch (button) {
-        case SDL_GAMEPAD_BUTTON_TOUCHPAD:
-            touchpad_x = Config::getBackButtonBehavior() == "left"    ? 0.25f
-                         : Config::getBackButtonBehavior() == "right" ? 0.75f
-                                                                      : 0.5f;
-            controller->SetTouchpadState(0, new_button_state, touchpad_x, 0.5f);
+        case SDL_GAMEPAD_BUTTON_TOUCHPAD_LEFT:
+            LOG_INFO(Input, "Topuchpad left");
+            controller->SetTouchpadState(0, new_button_state, 0.25f, 0.5f);
+            controller->CheckButton(0, SDLGamepadToOrbisButton(button), new_button_state);
+            break;
+        case SDL_GAMEPAD_BUTTON_TOUCHPAD_CENTER:
+            LOG_INFO(Input, "Topuchpad center");
+            controller->SetTouchpadState(0, new_button_state, 0.50f, 0.5f);
+            controller->CheckButton(0, SDLGamepadToOrbisButton(button), new_button_state);
+            break;
+        case SDL_GAMEPAD_BUTTON_TOUCHPAD_RIGHT:
+            LOG_INFO(Input, "Topuchpad right");
+            controller->SetTouchpadState(0, new_button_state, 0.75f, 0.5f);
             controller->CheckButton(0, SDLGamepadToOrbisButton(button), new_button_state);
             break;
         case LEFTJOYSTICK_HALFMODE:
@@ -522,6 +538,9 @@ void ControllerOutput::FinalizeUpdate() {
         // to do it, and it would be inconvenient to force it here, when AddUpdate does the job just
         // fine, and a toggle doesn't have to checked against every input that's bound to it, it's
         // enough that one is pressed
+        case MOUSE_GYRO_ROLL_MODE:
+            SetMouseGyroRollMode(new_button_state);
+            break;
         default: // is a normal key (hopefully)
             controller->CheckButton(0, SDLGamepadToOrbisButton(button), new_button_state);
             break;
diff --git a/src/input/input_handler.h b/src/input/input_handler.h
index 4aca133b1..745906620 100644
--- a/src/input/input_handler.h
+++ b/src/input/input_handler.h
@@ -23,6 +23,10 @@
 #define SDL_MOUSE_WHEEL_LEFT SDL_EVENT_MOUSE_WHEEL + 5
 #define SDL_MOUSE_WHEEL_RIGHT SDL_EVENT_MOUSE_WHEEL + 7
 
+#define SDL_GAMEPAD_BUTTON_TOUCHPAD_LEFT SDL_GAMEPAD_BUTTON_COUNT + 1
+#define SDL_GAMEPAD_BUTTON_TOUCHPAD_CENTER SDL_GAMEPAD_BUTTON_COUNT + 2
+#define SDL_GAMEPAD_BUTTON_TOUCHPAD_RIGHT SDL_GAMEPAD_BUTTON_COUNT + 3
+
 // idk who already used what where so I just chose a big number
 #define SDL_EVENT_MOUSE_WHEEL_OFF SDL_EVENT_USER + 10
 
@@ -31,6 +35,7 @@
 #define BACK_BUTTON 0x00040000
 
 #define KEY_TOGGLE 0x00200000
+#define MOUSE_GYRO_ROLL_MODE 0x00400000
 
 #define SDL_UNMAPPED UINT32_MAX - 1
 
@@ -98,7 +103,9 @@ const std::map<std::string, u32> string_to_cbutton_map = {
     {"options", SDL_GAMEPAD_BUTTON_START},
 
     // these are outputs only (touchpad can only be bound to itself)
-    {"touchpad", SDL_GAMEPAD_BUTTON_TOUCHPAD},
+    {"touchpad_left", SDL_GAMEPAD_BUTTON_TOUCHPAD_LEFT},
+    {"touchpad_center", SDL_GAMEPAD_BUTTON_TOUCHPAD_CENTER},
+    {"touchpad_right", SDL_GAMEPAD_BUTTON_TOUCHPAD_RIGHT},
     {"leftjoystick_halfmode", LEFTJOYSTICK_HALFMODE},
     {"rightjoystick_halfmode", RIGHTJOYSTICK_HALFMODE},
 
@@ -108,6 +115,7 @@ const std::map<std::string, u32> string_to_cbutton_map = {
     {"lpaddle_low", SDL_GAMEPAD_BUTTON_LEFT_PADDLE2},
     {"rpaddle_high", SDL_GAMEPAD_BUTTON_RIGHT_PADDLE1},
     {"rpaddle_low", SDL_GAMEPAD_BUTTON_RIGHT_PADDLE2},
+    {"mouse_gyro_roll_mode", MOUSE_GYRO_ROLL_MODE},
 };
 
 const std::map<std::string, AxisMapping> string_to_axis_map = {
diff --git a/src/input/input_mouse.cpp b/src/input/input_mouse.cpp
index 5eb0aab3e..3c718dbd5 100644
--- a/src/input/input_mouse.cpp
+++ b/src/input/input_mouse.cpp
@@ -3,6 +3,7 @@
 
 #include <cmath>
 
+#include "common/assert.h"
 #include "common/types.h"
 #include "input/controller.h"
 #include "input_mouse.h"
@@ -13,12 +14,19 @@ namespace Input {
 
 int mouse_joystick_binding = 0;
 float mouse_deadzone_offset = 0.5, mouse_speed = 1, mouse_speed_offset = 0.1250;
+bool mouse_gyro_roll_mode = false;
 Uint32 mouse_polling_id = 0;
-bool mouse_enabled = false;
+MouseMode mouse_mode = MouseMode::Off;
 
-// We had to go through 3 files of indirection just to update a flag
-void ToggleMouseEnabled() {
-    mouse_enabled = !mouse_enabled;
+// Switches mouse to a set mode or turns mouse emulation off if it was already in that mode.
+// Returns whether the mode is turned on.
+bool ToggleMouseModeTo(MouseMode m) {
+    if (mouse_mode == m) {
+        mouse_mode = MouseMode::Off;
+    } else {
+        mouse_mode = m;
+    }
+    return mouse_mode == m;
 }
 
 void SetMouseToJoystick(int joystick) {
@@ -31,10 +39,11 @@ void SetMouseParams(float mdo, float ms, float mso) {
     mouse_speed_offset = mso;
 }
 
-Uint32 MousePolling(void* param, Uint32 id, Uint32 interval) {
-    auto* controller = (GameController*)param;
-    if (!mouse_enabled)
-        return interval;
+void SetMouseGyroRollMode(bool mode) {
+    mouse_gyro_roll_mode = mode;
+}
+
+void EmulateJoystick(GameController* controller, u32 interval) {
 
     Axis axis_x, axis_y;
     switch (mouse_joystick_binding) {
@@ -47,7 +56,7 @@ Uint32 MousePolling(void* param, Uint32 id, Uint32 interval) {
         axis_y = Axis::RightY;
         break;
     default:
-        return interval; // no update needed
+        return; // no update needed
     }
 
     float d_x = 0, d_y = 0;
@@ -67,7 +76,35 @@ Uint32 MousePolling(void* param, Uint32 id, Uint32 interval) {
         controller->Axis(0, axis_x, GetAxis(-0x80, 0x7f, 0));
         controller->Axis(0, axis_y, GetAxis(-0x80, 0x7f, 0));
     }
+}
 
+constexpr float constant_down_accel[3] = {0.0f, 10.0f, 0.0f};
+void EmulateGyro(GameController* controller, u32 interval) {
+    // LOG_INFO(Input, "todo gyro");
+    float d_x = 0, d_y = 0;
+    SDL_GetRelativeMouseState(&d_x, &d_y);
+    controller->Acceleration(1, constant_down_accel);
+    float gyro_from_mouse[3] = {-d_y / 100, -d_x / 100, 0.0f};
+    if (mouse_gyro_roll_mode) {
+        gyro_from_mouse[1] = 0.0f;
+        gyro_from_mouse[2] = -d_x / 100;
+    }
+    controller->Gyro(1, gyro_from_mouse);
+}
+
+Uint32 MousePolling(void* param, Uint32 id, Uint32 interval) {
+    auto* controller = (GameController*)param;
+    switch (mouse_mode) {
+    case MouseMode::Joystick:
+        EmulateJoystick(controller, interval);
+        break;
+    case MouseMode::Gyro:
+        EmulateGyro(controller, interval);
+        break;
+
+    default:
+        break;
+    }
     return interval;
 }
 
diff --git a/src/input/input_mouse.h b/src/input/input_mouse.h
index da18ee04e..a56ef2d8f 100644
--- a/src/input/input_mouse.h
+++ b/src/input/input_mouse.h
@@ -8,11 +8,21 @@
 
 namespace Input {
 
-void ToggleMouseEnabled();
+enum MouseMode {
+    Off = 0,
+    Joystick,
+    Gyro,
+};
+
+bool ToggleMouseModeTo(MouseMode m);
 void SetMouseToJoystick(int joystick);
 void SetMouseParams(float mouse_deadzone_offset, float mouse_speed, float mouse_speed_offset);
+void SetMouseGyroRollMode(bool mode);
 
-// Polls the mouse for changes, and simulates joystick movement from it.
+void EmulateJoystick(GameController* controller, u32 interval);
+void EmulateGyro(GameController* controller, u32 interval);
+
+// Polls the mouse for changes
 Uint32 MousePolling(void* param, Uint32 id, Uint32 interval);
 
 } // namespace Input
diff --git a/src/qt_gui/about_dialog.cpp b/src/qt_gui/about_dialog.cpp
index 90fb14236..627a0c052 100644
--- a/src/qt_gui/about_dialog.cpp
+++ b/src/qt_gui/about_dialog.cpp
@@ -12,7 +12,8 @@
 #include "main_window_themes.h"
 #include "ui_about_dialog.h"
 
-AboutDialog::AboutDialog(QWidget* parent) : QDialog(parent), ui(new Ui::AboutDialog) {
+AboutDialog::AboutDialog(std::shared_ptr<gui_settings> gui_settings, QWidget* parent)
+    : QDialog(parent), ui(new Ui::AboutDialog), m_gui_settings(std::move(gui_settings)) {
     ui->setupUi(this);
     preloadImages();
 
@@ -57,7 +58,7 @@ void AboutDialog::preloadImages() {
 }
 
 void AboutDialog::updateImagesForCurrentTheme() {
-    Theme currentTheme = static_cast<Theme>(Config::getMainWindowTheme());
+    Theme currentTheme = static_cast<Theme>(m_gui_settings->GetValue(gui::gen_theme).toInt());
     bool isDarkTheme = (currentTheme == Theme::Dark || currentTheme == Theme::Green ||
                         currentTheme == Theme::Blue || currentTheme == Theme::Violet);
     if (isDarkTheme) {
@@ -188,7 +189,7 @@ void AboutDialog::removeHoverEffect(QLabel* label) {
 }
 
 bool AboutDialog::isDarkTheme() const {
-    Theme currentTheme = static_cast<Theme>(Config::getMainWindowTheme());
+    Theme currentTheme = static_cast<Theme>(m_gui_settings->GetValue(gui::gen_theme).toInt());
     return currentTheme == Theme::Dark || currentTheme == Theme::Green ||
            currentTheme == Theme::Blue || currentTheme == Theme::Violet;
 }
diff --git a/src/qt_gui/about_dialog.h b/src/qt_gui/about_dialog.h
index 42e8d557a..b74cdfd1a 100644
--- a/src/qt_gui/about_dialog.h
+++ b/src/qt_gui/about_dialog.h
@@ -8,6 +8,7 @@
 #include <QLabel>
 #include <QPixmap>
 #include <QUrl>
+#include "gui_settings.h"
 
 namespace Ui {
 class AboutDialog;
@@ -17,7 +18,7 @@ class AboutDialog : public QDialog {
     Q_OBJECT
 
 public:
-    explicit AboutDialog(QWidget* parent = nullptr);
+    explicit AboutDialog(std::shared_ptr<gui_settings> gui_settings, QWidget* parent = nullptr);
     ~AboutDialog();
     bool eventFilter(QObject* obj, QEvent* event);
 
@@ -33,4 +34,5 @@ private:
 
     QPixmap originalImages[5];
     QPixmap invertedImages[5];
+    std::shared_ptr<gui_settings> m_gui_settings;
 };
diff --git a/src/qt_gui/control_settings.h b/src/qt_gui/control_settings.h
index e686f044d..b1fff1dad 100644
--- a/src/qt_gui/control_settings.h
+++ b/src/qt_gui/control_settings.h
@@ -39,14 +39,28 @@ private:
         "pad_left",     "pad_right", "axis_left_x", "axis_left_y", "axis_right_x",
         "axis_right_y", "back"};
 
-    const QStringList ButtonOutputs = {"cross",    "circle",    "square",   "triangle", "l1",
-                                       "r1",       "l2",        "r2",       "l3",
+    const QStringList ButtonOutputs = {"cross",
+                                       "circle",
+                                       "square",
+                                       "triangle",
+                                       "l1",
+                                       "r1",
+                                       "l2",
+                                       "r2",
+                                       "l3",
 
-                                       "r3",       "options",   "pad_up",
+                                       "r3",
+                                       "options",
+                                       "pad_up",
 
                                        "pad_down",
 
-                                       "pad_left", "pad_right", "touchpad", "unmapped"};
+                                       "pad_left",
+                                       "pad_right",
+                                       "touchpad_left",
+                                       "touchpad_center",
+                                       "touchpad_right",
+                                       "unmapped"};
 
     const QStringList StickOutputs = {"axis_left_x", "axis_left_y", "axis_right_x", "axis_right_y",
                                       "unmapped"};
diff --git a/src/qt_gui/elf_viewer.cpp b/src/qt_gui/elf_viewer.cpp
index e80fa25c1..8d472755b 100644
--- a/src/qt_gui/elf_viewer.cpp
+++ b/src/qt_gui/elf_viewer.cpp
@@ -3,10 +3,12 @@
 
 #include "elf_viewer.h"
 
-ElfViewer::ElfViewer(QWidget* parent) : QTableWidget(parent) {
-    dir_list_std = Config::getElfViewer();
-    for (const auto& str : dir_list_std) {
-        dir_list.append(QString::fromStdString(str));
+ElfViewer::ElfViewer(std::shared_ptr<gui_settings> gui_settings, QWidget* parent)
+    : QTableWidget(parent), m_gui_settings(std::move(gui_settings)) {
+
+    list = gui_settings::Var2List(m_gui_settings->GetValue(gui::gen_elfDirs));
+    for (const auto& str : list) {
+        dir_list.append(str);
     }
 
     CheckElfFolders();
@@ -55,11 +57,11 @@ void ElfViewer::OpenElfFolder() {
         }
         std::ranges::sort(m_elf_list);
         OpenElfFiles();
-        dir_list_std.clear();
+        list.clear();
         for (auto dir : dir_list) {
-            dir_list_std.push_back(dir.toStdString());
+            list.push_back(dir);
         }
-        Config::setElfViewer(dir_list_std);
+        m_gui_settings->SetValue(gui::gen_elfDirs, gui_settings::List2Var(list));
     } else {
         // qDebug() << "Folder selection canceled.";
     }
diff --git a/src/qt_gui/elf_viewer.h b/src/qt_gui/elf_viewer.h
index 1a65d70de..6256abf31 100644
--- a/src/qt_gui/elf_viewer.h
+++ b/src/qt_gui/elf_viewer.h
@@ -11,7 +11,7 @@
 class ElfViewer : public QTableWidget {
     Q_OBJECT
 public:
-    explicit ElfViewer(QWidget* parent = nullptr);
+    explicit ElfViewer(std::shared_ptr<gui_settings> gui_settings, QWidget* parent = nullptr);
     QStringList m_elf_list;
 
 private:
@@ -21,7 +21,8 @@ private:
     Core::Loader::Elf m_elf_file;
     QStringList dir_list;
     QStringList elf_headers_list;
-    std::vector<std::string> dir_list_std;
+    QList<QString> list;
+    std::shared_ptr<gui_settings> m_gui_settings;
 
     void SetTableItem(QTableWidget* game_list, int row, int column, QString itemStr) {
         QTableWidgetItem* item = new QTableWidgetItem();
diff --git a/src/qt_gui/game_grid_frame.cpp b/src/qt_gui/game_grid_frame.cpp
index 66679dc71..8a5219da1 100644
--- a/src/qt_gui/game_grid_frame.cpp
+++ b/src/qt_gui/game_grid_frame.cpp
@@ -34,7 +34,8 @@ GameGridFrame::GameGridFrame(std::shared_ptr<gui_settings> gui_settings,
     connect(this->horizontalScrollBar(), &QScrollBar::valueChanged, this,
             &GameGridFrame::RefreshGridBackgroundImage);
     connect(this, &QTableWidget::customContextMenuRequested, this, [=, this](const QPoint& pos) {
-        m_gui_context_menus.RequestGameMenu(pos, m_game_info->m_games, m_compat_info, this, false);
+        m_gui_context_menus.RequestGameMenu(pos, m_game_info->m_games, m_compat_info,
+                                            m_gui_settings, this, false);
     });
 }
 
diff --git a/src/qt_gui/game_list_frame.cpp b/src/qt_gui/game_list_frame.cpp
index dd10e0f8b..45a9a4810 100644
--- a/src/qt_gui/game_list_frame.cpp
+++ b/src/qt_gui/game_list_frame.cpp
@@ -75,7 +75,8 @@ GameListFrame::GameListFrame(std::shared_ptr<gui_settings> gui_settings,
         });
 
     connect(this, &QTableWidget::customContextMenuRequested, this, [=, this](const QPoint& pos) {
-        m_gui_context_menus.RequestGameMenu(pos, m_game_info->m_games, m_compat_info, this, true);
+        m_gui_context_menus.RequestGameMenu(pos, m_game_info->m_games, m_compat_info,
+                                            m_gui_settings, this, true);
     });
 
     connect(this, &QTableWidget::cellClicked, this, [=, this](int row, int column) {
diff --git a/src/qt_gui/gui_context_menus.h b/src/qt_gui/gui_context_menus.h
index 46a40c5cd..ba82da261 100644
--- a/src/qt_gui/gui_context_menus.h
+++ b/src/qt_gui/gui_context_menus.h
@@ -32,8 +32,10 @@ class GuiContextMenus : public QObject {
 public:
     void RequestGameMenu(const QPoint& pos, QVector<GameInfo>& m_games,
                          std::shared_ptr<CompatibilityInfoClass> m_compat_info,
-                         QTableWidget* widget, bool isList) {
+                         std::shared_ptr<gui_settings> settings, QTableWidget* widget,
+                         bool isList) {
         QPoint global_pos = widget->viewport()->mapToGlobal(pos);
+        std::shared_ptr<gui_settings> m_gui_settings = std::move(settings);
         int itemID = 0;
         if (isList) {
             itemID = widget->currentRow();
@@ -357,7 +359,7 @@ public:
 
             QString gameName = QString::fromStdString(m_games[itemID].name);
             TrophyViewer* trophyViewer =
-                new TrophyViewer(trophyPath, gameTrpPath, gameName, allTrophyGames);
+                new TrophyViewer(m_gui_settings, trophyPath, gameTrpPath, gameName, allTrophyGames);
             trophyViewer->show();
             connect(widget->parent(), &QWidget::destroyed, trophyViewer,
                     [trophyViewer]() { trophyViewer->deleteLater(); });
diff --git a/src/qt_gui/gui_settings.h b/src/qt_gui/gui_settings.h
index da5542956..0fa807d70 100644
--- a/src/qt_gui/gui_settings.h
+++ b/src/qt_gui/gui_settings.h
@@ -17,6 +17,12 @@ const QString game_grid = "game_grid";
 const gui_value gen_checkForUpdates = gui_value(general_settings, "checkForUpdates", false);
 const gui_value gen_showChangeLog = gui_value(general_settings, "showChangeLog", false);
 const gui_value gen_updateChannel = gui_value(general_settings, "updateChannel", "Release");
+const gui_value gen_recentFiles =
+    gui_value(main_window, "recentFiles", QVariant::fromValue(QList<QString>()));
+const gui_value gen_guiLanguage = gui_value(general_settings, "guiLanguage", "en_US");
+const gui_value gen_elfDirs =
+    gui_value(main_window, "elfDirs", QVariant::fromValue(QList<QString>()));
+const gui_value gen_theme = gui_value(general_settings, "theme", 0);
 
 // main window settings
 const gui_value mw_geometry = gui_value(main_window, "geometry", QByteArray());
diff --git a/src/qt_gui/kbm_gui.cpp b/src/qt_gui/kbm_gui.cpp
index df8efe9c0..1f7743412 100644
--- a/src/qt_gui/kbm_gui.cpp
+++ b/src/qt_gui/kbm_gui.cpp
@@ -32,14 +32,34 @@ KBMSettings::KBMSettings(std::shared_ptr<GameInfoClass> game_info_get, QWidget*
         ui->ProfileComboBox->addItem(QString::fromStdString(m_game_info->m_games[i].serial));
     }
 
-    ButtonsList = {
-        ui->CrossButton,    ui->CircleButton,     ui->TriangleButton,   ui->SquareButton,
-        ui->L1Button,       ui->R1Button,         ui->L2Button,         ui->R2Button,
-        ui->L3Button,       ui->R3Button,         ui->OptionsButton,    ui->TouchpadButton,
-        ui->DpadUpButton,   ui->DpadDownButton,   ui->DpadLeftButton,   ui->DpadRightButton,
-        ui->LStickUpButton, ui->LStickDownButton, ui->LStickLeftButton, ui->LStickRightButton,
-        ui->RStickUpButton, ui->RStickDownButton, ui->RStickLeftButton, ui->RStickRightButton,
-        ui->LHalfButton,    ui->RHalfButton};
+    ButtonsList = {ui->CrossButton,
+                   ui->CircleButton,
+                   ui->TriangleButton,
+                   ui->SquareButton,
+                   ui->L1Button,
+                   ui->R1Button,
+                   ui->L2Button,
+                   ui->R2Button,
+                   ui->L3Button,
+                   ui->R3Button,
+                   ui->OptionsButton,
+                   ui->TouchpadLeftButton,
+                   ui->TouchpadCenterButton,
+                   ui->TouchpadRightButton,
+                   ui->DpadUpButton,
+                   ui->DpadDownButton,
+                   ui->DpadLeftButton,
+                   ui->DpadRightButton,
+                   ui->LStickUpButton,
+                   ui->LStickDownButton,
+                   ui->LStickLeftButton,
+                   ui->LStickRightButton,
+                   ui->RStickUpButton,
+                   ui->RStickDownButton,
+                   ui->RStickLeftButton,
+                   ui->RStickRightButton,
+                   ui->LHalfButton,
+                   ui->RHalfButton};
 
     ButtonConnects();
     SetUIValuestoMappings("default");
@@ -187,8 +207,10 @@ void KBMSettings::SaveKBMConfig(bool close_on_save) {
 
     lines.push_back("");
 
+    add_mapping(ui->TouchpadLeftButton->text(), "touchpad_left");
+    add_mapping(ui->TouchpadCenterButton->text(), "touchpad_center");
+    add_mapping(ui->TouchpadRightButton->text(), "touchpad_right");
     add_mapping(ui->OptionsButton->text(), "options");
-    add_mapping(ui->TouchpadButton->text(), "touchpad");
 
     lines.push_back("");
 
@@ -317,7 +339,9 @@ void KBMSettings::SetDefault() {
     ui->R2Button->setText("o");
     ui->R3Button->setText("m");
 
-    ui->TouchpadButton->setText("space");
+    ui->TouchpadLeftButton->setText("space");
+    ui->TouchpadCenterButton->setText("unmapped");
+    ui->TouchpadRightButton->setText("unmapped");
     ui->OptionsButton->setText("enter");
 
     ui->DpadUpButton->setText("up");
@@ -396,8 +420,12 @@ void KBMSettings::SetUIValuestoMappings(std::string config_id) {
                 ui->DpadRightButton->setText(QString::fromStdString(input_string));
             } else if (output_string == "options") {
                 ui->OptionsButton->setText(QString::fromStdString(input_string));
-            } else if (output_string == "touchpad") {
-                ui->TouchpadButton->setText(QString::fromStdString(input_string));
+            } else if (output_string == "touchpad_left") {
+                ui->TouchpadLeftButton->setText(QString::fromStdString(input_string));
+            } else if (output_string == "touchpad_center") {
+                ui->TouchpadCenterButton->setText(QString::fromStdString(input_string));
+            } else if (output_string == "touchpad_right") {
+                ui->TouchpadRightButton->setText(QString::fromStdString(input_string));
             } else if (output_string == "axis_left_x_minus") {
                 ui->LStickLeftButton->setText(QString::fromStdString(input_string));
             } else if (output_string == "axis_left_x_plus") {
diff --git a/src/qt_gui/kbm_gui.ui b/src/qt_gui/kbm_gui.ui
index 109423aa8..eb393254d 100644
--- a/src/qt_gui/kbm_gui.ui
+++ b/src/qt_gui/kbm_gui.ui
@@ -11,8 +11,8 @@
    <rect>
     <x>0</x>
     <y>0</y>
-    <width>1234</width>
-    <height>796</height>
+    <width>1235</width>
+    <height>842</height>
    </rect>
   </property>
   <property name="sizePolicy">
@@ -44,8 +44,8 @@
        <rect>
         <x>0</x>
         <y>0</y>
-        <width>1214</width>
-        <height>746</height>
+        <width>1215</width>
+        <height>792</height>
        </rect>
       </property>
       <widget class="QWidget" name="layoutWidget">
@@ -54,7 +54,7 @@
          <x>0</x>
          <y>0</y>
          <width>1211</width>
-         <height>741</height>
+         <height>791</height>
         </rect>
        </property>
        <layout class="QHBoxLayout" name="RemapLayout">
@@ -793,7 +793,7 @@
              </layout>
             </item>
             <item>
-             <layout class="QVBoxLayout" name="layout_system_buttons" stretch="0">
+             <layout class="QVBoxLayout" name="layout_system_buttons" stretch="0,0">
               <item>
                <widget class="QGroupBox" name="groupBox_4">
                 <property name="sizePolicy">
@@ -825,8 +825,11 @@
                      <verstretch>0</verstretch>
                     </sizepolicy>
                    </property>
-                   <property name="minimumHeight">
-                    <number>48</number>
+                   <property name="minimumSize">
+                    <size>
+                     <width>0</width>
+                     <height>24</height>
+                    </size>
                    </property>
                    <property name="focusPolicy">
                     <enum>Qt::FocusPolicy::NoFocus</enum>
@@ -844,8 +847,11 @@
                      <verstretch>0</verstretch>
                     </sizepolicy>
                    </property>
-                   <property name="minimumHeight">
-                    <number>48</number>
+                   <property name="minimumSize">
+                    <size>
+                     <width>0</width>
+                     <height>24</height>
+                    </size>
                    </property>
                    <property name="focusPolicy">
                     <enum>Qt::FocusPolicy::NoFocus</enum>
@@ -858,6 +864,55 @@
                 </layout>
                </widget>
               </item>
+              <item>
+               <widget class="QGroupBox" name="gb_options">
+                <property name="sizePolicy">
+                 <sizepolicy hsizetype="Expanding" vsizetype="Preferred">
+                  <horstretch>0</horstretch>
+                  <verstretch>0</verstretch>
+                 </sizepolicy>
+                </property>
+                <property name="minimumSize">
+                 <size>
+                  <width>160</width>
+                  <height>0</height>
+                 </size>
+                </property>
+                <property name="title">
+                 <string>Options</string>
+                </property>
+                <layout class="QVBoxLayout" name="gb_start_layout">
+                 <property name="leftMargin">
+                  <number>5</number>
+                 </property>
+                 <property name="topMargin">
+                  <number>5</number>
+                 </property>
+                 <property name="rightMargin">
+                  <number>5</number>
+                 </property>
+                 <property name="bottomMargin">
+                  <number>5</number>
+                 </property>
+                 <item>
+                  <widget class="QPushButton" name="OptionsButton">
+                   <property name="sizePolicy">
+                    <sizepolicy hsizetype="Preferred" vsizetype="Fixed">
+                     <horstretch>0</horstretch>
+                     <verstretch>0</verstretch>
+                    </sizepolicy>
+                   </property>
+                   <property name="focusPolicy">
+                    <enum>Qt::FocusPolicy::NoFocus</enum>
+                   </property>
+                   <property name="text">
+                    <string>unmapped</string>
+                   </property>
+                  </widget>
+                 </item>
+                </layout>
+               </widget>
+              </item>
              </layout>
             </item>
             <item>
@@ -1067,34 +1122,13 @@
                </widget>
               </item>
               <item>
-               <widget class="QGroupBox" name="gb_touchpad">
-                <property name="sizePolicy">
-                 <sizepolicy hsizetype="Expanding" vsizetype="Preferred">
-                  <horstretch>0</horstretch>
-                  <verstretch>0</verstretch>
-                 </sizepolicy>
-                </property>
-                <property name="minimumSize">
-                 <size>
-                  <width>160</width>
-                  <height>0</height>
-                 </size>
-                </property>
+               <widget class="QGroupBox" name="gb_touchpadleft">
                 <property name="title">
-                 <string>Touchpad Click</string>
+                 <string>Touchpad Left</string>
                 </property>
-                <layout class="QVBoxLayout" name="verticalLayout_9">
+                <layout class="QVBoxLayout" name="verticalLayout_17">
                  <item>
-                  <widget class="QPushButton" name="TouchpadButton">
-                   <property name="sizePolicy">
-                    <sizepolicy hsizetype="Preferred" vsizetype="Fixed">
-                     <horstretch>0</horstretch>
-                     <verstretch>0</verstretch>
-                    </sizepolicy>
-                   </property>
-                   <property name="focusPolicy">
-                    <enum>Qt::FocusPolicy::NoFocus</enum>
-                   </property>
+                  <widget class="QPushButton" name="TouchpadLeftButton">
                    <property name="text">
                     <string>unmapped</string>
                    </property>
@@ -1150,6 +1184,22 @@
                 </layout>
                </widget>
               </item>
+              <item>
+               <widget class="QGroupBox" name="gb_touchpadcenter">
+                <property name="title">
+                 <string>Touchpad Center</string>
+                </property>
+                <layout class="QVBoxLayout" name="verticalLayout_11">
+                 <item>
+                  <widget class="QPushButton" name="TouchpadCenterButton">
+                   <property name="text">
+                    <string>unmapped</string>
+                   </property>
+                  </widget>
+                 </item>
+                </layout>
+               </widget>
+              </item>
              </layout>
             </item>
             <item>
@@ -1204,7 +1254,7 @@
                </widget>
               </item>
               <item>
-               <widget class="QGroupBox" name="gb_options">
+               <widget class="QGroupBox" name="gb_touchpadright">
                 <property name="sizePolicy">
                  <sizepolicy hsizetype="Expanding" vsizetype="Preferred">
                   <horstretch>0</horstretch>
@@ -1218,23 +1268,11 @@
                  </size>
                 </property>
                 <property name="title">
-                 <string>Options</string>
+                 <string>Touchpad Right</string>
                 </property>
-                <layout class="QVBoxLayout" name="gb_start_layout">
-                 <property name="leftMargin">
-                  <number>5</number>
-                 </property>
-                 <property name="topMargin">
-                  <number>5</number>
-                 </property>
-                 <property name="rightMargin">
-                  <number>5</number>
-                 </property>
-                 <property name="bottomMargin">
-                  <number>5</number>
-                 </property>
+                <layout class="QVBoxLayout" name="verticalLayout_9">
                  <item>
-                  <widget class="QPushButton" name="OptionsButton">
+                  <widget class="QPushButton" name="TouchpadRightButton">
                    <property name="sizePolicy">
                     <sizepolicy hsizetype="Preferred" vsizetype="Fixed">
                      <horstretch>0</horstretch>
diff --git a/src/qt_gui/main_window.cpp b/src/qt_gui/main_window.cpp
index c6da49182..9379519c2 100644
--- a/src/qt_gui/main_window.cpp
+++ b/src/qt_gui/main_window.cpp
@@ -39,8 +39,6 @@ MainWindow::MainWindow(QWidget* parent) : QMainWindow(parent), ui(new Ui::MainWi
 
 MainWindow::~MainWindow() {
     SaveWindowState();
-    const auto config_dir = Common::FS::GetUserPath(Common::FS::PathType::UserDir);
-    Config::saveMainWindow(config_dir / "config.toml");
 }
 
 bool MainWindow::Init() {
@@ -297,7 +295,7 @@ void MainWindow::CreateDockWindows() {
     m_game_list_frame->setObjectName("gamelist");
     m_game_grid_frame.reset(new GameGridFrame(m_gui_settings, m_game_info, m_compat_info, this));
     m_game_grid_frame->setObjectName("gamegridlist");
-    m_elf_viewer.reset(new ElfViewer(this));
+    m_elf_viewer.reset(new ElfViewer(m_gui_settings, this));
     m_elf_viewer->setObjectName("elflist");
 
     int table_mode = m_gui_settings->GetValue(gui::gl_mode).toInt();
@@ -492,7 +490,7 @@ void MainWindow::CreateConnects() {
 #endif
 
     connect(ui->aboutAct, &QAction::triggered, this, [this]() {
-        auto aboutDialog = new AboutDialog(this);
+        auto aboutDialog = new AboutDialog(m_gui_settings, this);
         aboutDialog->exec();
     });
 
@@ -771,14 +769,14 @@ void MainWindow::CreateConnects() {
 
         QString gameName = QString::fromStdString(firstGame.name);
         TrophyViewer* trophyViewer =
-            new TrophyViewer(trophyPath, gameTrpPath, gameName, allTrophyGames);
+            new TrophyViewer(m_gui_settings, trophyPath, gameTrpPath, gameName, allTrophyGames);
         trophyViewer->show();
     });
 
     // Themes
     connect(ui->setThemeDark, &QAction::triggered, &m_window_themes, [this]() {
         m_window_themes.SetWindowTheme(Theme::Dark, ui->mw_searchbar);
-        Config::setMainWindowTheme(static_cast<int>(Theme::Dark));
+        m_gui_settings->SetValue(gui::gen_theme, static_cast<int>(Theme::Dark));
         if (isIconBlack) {
             SetUiIcons(false);
             isIconBlack = false;
@@ -786,7 +784,7 @@ void MainWindow::CreateConnects() {
     });
     connect(ui->setThemeLight, &QAction::triggered, &m_window_themes, [this]() {
         m_window_themes.SetWindowTheme(Theme::Light, ui->mw_searchbar);
-        Config::setMainWindowTheme(static_cast<int>(Theme::Light));
+        m_gui_settings->SetValue(gui::gen_theme, static_cast<int>(Theme::Light));
         if (!isIconBlack) {
             SetUiIcons(true);
             isIconBlack = true;
@@ -794,7 +792,7 @@ void MainWindow::CreateConnects() {
     });
     connect(ui->setThemeGreen, &QAction::triggered, &m_window_themes, [this]() {
         m_window_themes.SetWindowTheme(Theme::Green, ui->mw_searchbar);
-        Config::setMainWindowTheme(static_cast<int>(Theme::Green));
+        m_gui_settings->SetValue(gui::gen_theme, static_cast<int>(Theme::Green));
         if (isIconBlack) {
             SetUiIcons(false);
             isIconBlack = false;
@@ -802,7 +800,7 @@ void MainWindow::CreateConnects() {
     });
     connect(ui->setThemeBlue, &QAction::triggered, &m_window_themes, [this]() {
         m_window_themes.SetWindowTheme(Theme::Blue, ui->mw_searchbar);
-        Config::setMainWindowTheme(static_cast<int>(Theme::Blue));
+        m_gui_settings->SetValue(gui::gen_theme, static_cast<int>(Theme::Blue));
         if (isIconBlack) {
             SetUiIcons(false);
             isIconBlack = false;
@@ -810,7 +808,7 @@ void MainWindow::CreateConnects() {
     });
     connect(ui->setThemeViolet, &QAction::triggered, &m_window_themes, [this]() {
         m_window_themes.SetWindowTheme(Theme::Violet, ui->mw_searchbar);
-        Config::setMainWindowTheme(static_cast<int>(Theme::Violet));
+        m_gui_settings->SetValue(gui::gen_theme, static_cast<int>(Theme::Violet));
         if (isIconBlack) {
             SetUiIcons(false);
             isIconBlack = false;
@@ -818,7 +816,7 @@ void MainWindow::CreateConnects() {
     });
     connect(ui->setThemeGruvbox, &QAction::triggered, &m_window_themes, [this]() {
         m_window_themes.SetWindowTheme(Theme::Gruvbox, ui->mw_searchbar);
-        Config::setMainWindowTheme(static_cast<int>(Theme::Gruvbox));
+        m_gui_settings->SetValue(gui::gen_theme, static_cast<int>(Theme::Gruvbox));
         if (isIconBlack) {
             SetUiIcons(false);
             isIconBlack = false;
@@ -826,7 +824,7 @@ void MainWindow::CreateConnects() {
     });
     connect(ui->setThemeTokyoNight, &QAction::triggered, &m_window_themes, [this]() {
         m_window_themes.SetWindowTheme(Theme::TokyoNight, ui->mw_searchbar);
-        Config::setMainWindowTheme(static_cast<int>(Theme::TokyoNight));
+        m_gui_settings->SetValue(gui::gen_theme, static_cast<int>(Theme::TokyoNight));
         if (isIconBlack) {
             SetUiIcons(false);
             isIconBlack = false;
@@ -834,7 +832,7 @@ void MainWindow::CreateConnects() {
     });
     connect(ui->setThemeOled, &QAction::triggered, &m_window_themes, [this]() {
         m_window_themes.SetWindowTheme(Theme::Oled, ui->mw_searchbar);
-        Config::setMainWindowTheme(static_cast<int>(Theme::Oled));
+        m_gui_settings->SetValue(gui::gen_theme, static_cast<int>(Theme::Oled));
         if (isIconBlack) {
             SetUiIcons(false);
             isIconBlack = false;
@@ -981,7 +979,7 @@ void MainWindow::InstallDirectory() {
 }
 
 void MainWindow::SetLastUsedTheme() {
-    Theme lastTheme = static_cast<Theme>(Config::getMainWindowTheme());
+    Theme lastTheme = static_cast<Theme>(m_gui_settings->GetValue(gui::gen_theme).toInt());
     m_window_themes.SetWindowTheme(lastTheme, ui->mw_searchbar);
 
     switch (lastTheme) {
@@ -1122,33 +1120,32 @@ void MainWindow::HandleResize(QResizeEvent* event) {
 }
 
 void MainWindow::AddRecentFiles(QString filePath) {
-    std::vector<std::string> vec = Config::getRecentFiles();
-    if (!vec.empty()) {
-        if (filePath.toStdString() == vec.at(0)) {
+    QList<QString> list = gui_settings::Var2List(m_gui_settings->GetValue(gui::gen_recentFiles));
+    if (!list.empty()) {
+        if (filePath == list.at(0)) {
             return;
         }
-        auto it = std::find(vec.begin(), vec.end(), filePath.toStdString());
-        if (it != vec.end()) {
-            vec.erase(it);
+        auto it = std::find(list.begin(), list.end(), filePath);
+        if (it != list.end()) {
+            list.erase(it);
         }
     }
-    vec.insert(vec.begin(), filePath.toStdString());
-    if (vec.size() > 6) {
-        vec.pop_back();
+    list.insert(list.begin(), filePath);
+    if (list.size() > 6) {
+        list.pop_back();
     }
-    Config::setRecentFiles(vec);
-    const auto config_dir = Common::FS::GetUserPath(Common::FS::PathType::UserDir);
-    Config::saveMainWindow(config_dir / "config.toml");
+    m_gui_settings->SetValue(gui::gen_recentFiles, gui_settings::List2Var(list));
     CreateRecentGameActions(); // Refresh the QActions.
 }
 
 void MainWindow::CreateRecentGameActions() {
     m_recent_files_group = new QActionGroup(this);
     ui->menuRecent->clear();
-    std::vector<std::string> vec = Config::getRecentFiles();
-    for (int i = 0; i < vec.size(); i++) {
+    QList<QString> list = gui_settings::Var2List(m_gui_settings->GetValue(gui::gen_recentFiles));
+
+    for (int i = 0; i < list.size(); i++) {
         QAction* recentFileAct = new QAction(this);
-        recentFileAct->setText(QString::fromStdString(vec.at(i)));
+        recentFileAct->setText(list.at(i));
         ui->menuRecent->addAction(recentFileAct);
         m_recent_files_group->addAction(recentFileAct);
     }
@@ -1165,7 +1162,7 @@ void MainWindow::CreateRecentGameActions() {
 }
 
 void MainWindow::LoadTranslation() {
-    auto language = QString::fromStdString(Config::getEmulatorLanguage());
+    auto language = m_gui_settings->GetValue(gui::gen_guiLanguage).toString();
 
     const QString base_dir = QStringLiteral(":/translations");
     QString base_path = QStringLiteral("%1/%2.qm").arg(base_dir).arg(language);
@@ -1190,8 +1187,8 @@ void MainWindow::LoadTranslation() {
     }
 }
 
-void MainWindow::OnLanguageChanged(const std::string& locale) {
-    Config::setEmulatorLanguage(locale);
+void MainWindow::OnLanguageChanged(const QString& locale) {
+    m_gui_settings->SetValue(gui::gen_guiLanguage, locale);
 
     LoadTranslation();
 }
diff --git a/src/qt_gui/main_window.h b/src/qt_gui/main_window.h
index 7f11f7310..eec1a65de 100644
--- a/src/qt_gui/main_window.h
+++ b/src/qt_gui/main_window.h
@@ -47,7 +47,7 @@ private Q_SLOTS:
     void ShowGameList();
     void RefreshGameTable();
     void HandleResize(QResizeEvent* event);
-    void OnLanguageChanged(const std::string& locale);
+    void OnLanguageChanged(const QString& locale);
     void toggleLabelsUnderIcons();
 
 private:
diff --git a/src/qt_gui/settings.cpp b/src/qt_gui/settings.cpp
index 44133dac5..4a9c1d375 100644
--- a/src/qt_gui/settings.cpp
+++ b/src/qt_gui/settings.cpp
@@ -75,3 +75,17 @@ void settings::SetValue(const QString& key, const QString& name, const QVariant&
         }
     }
 }
+QVariant settings::List2Var(const QList<QString>& list) {
+    QByteArray ba;
+    QDataStream stream(&ba, QIODevice::WriteOnly);
+    stream << list;
+    return QVariant(ba);
+}
+
+QList<QString> settings::Var2List(const QVariant& var) {
+    QList<QString> list;
+    QByteArray ba = var.toByteArray();
+    QDataStream stream(&ba, QIODevice::ReadOnly);
+    stream >> list;
+    return list;
+}
\ No newline at end of file
diff --git a/src/qt_gui/settings.h b/src/qt_gui/settings.h
index da71fe01a..837804d00 100644
--- a/src/qt_gui/settings.h
+++ b/src/qt_gui/settings.h
@@ -35,6 +35,8 @@ public:
 
     QVariant GetValue(const QString& key, const QString& name, const QVariant& def) const;
     QVariant GetValue(const gui_value& entry) const;
+    static QVariant List2Var(const QList<QString>& list);
+    static QList<QString> Var2List(const QVariant& var);
 
 public Q_SLOTS:
     /** Remove entry */
diff --git a/src/qt_gui/settings_dialog.cpp b/src/qt_gui/settings_dialog.cpp
index da2b0dde3..c9d264587 100644
--- a/src/qt_gui/settings_dialog.cpp
+++ b/src/qt_gui/settings_dialog.cpp
@@ -123,11 +123,6 @@ SettingsDialog::SettingsDialog(std::shared_ptr<gui_settings> gui_settings,
     ui->hideCursorComboBox->addItem(tr("Idle"));
     ui->hideCursorComboBox->addItem(tr("Always"));
 
-    ui->backButtonBehaviorComboBox->addItem(tr("Touchpad Left"), "left");
-    ui->backButtonBehaviorComboBox->addItem(tr("Touchpad Center"), "center");
-    ui->backButtonBehaviorComboBox->addItem(tr("Touchpad Right"), "right");
-    ui->backButtonBehaviorComboBox->addItem(tr("None"), "none");
-
     InitializeEmulatorLanguages();
     LoadValuesFromConfig();
 
@@ -366,7 +361,6 @@ SettingsDialog::SettingsDialog(std::shared_ptr<gui_settings> gui_settings,
         // Input
         ui->hideCursorGroupBox->installEventFilter(this);
         ui->idleTimeoutGroupBox->installEventFilter(this);
-        ui->backButtonBehaviorGroupBox->installEventFilter(this);
 
         // Graphics
         ui->graphicsAdapterGroupBox->installEventFilter(this);
@@ -534,10 +528,6 @@ void SettingsDialog::LoadValuesFromConfig() {
         indexTab = 0;
     ui->tabWidgetSettings->setCurrentIndex(indexTab);
 
-    QString backButtonBehavior = QString::fromStdString(
-        toml::find_or<std::string>(data, "Input", "backButtonBehavior", "left"));
-    int index = ui->backButtonBehaviorComboBox->findData(backButtonBehavior);
-    ui->backButtonBehaviorComboBox->setCurrentIndex(index != -1 ? index : 0);
     ui->motionControlsCheckBox->setChecked(
         toml::find_or<bool>(data, "Input", "isMotionControlsEnabled", true));
 
@@ -594,7 +584,7 @@ void SettingsDialog::OnLanguageChanged(int index) {
 
     ui->retranslateUi(this);
 
-    emit LanguageChanged(ui->emulatorLanguageComboBox->itemData(index).toString().toStdString());
+    emit LanguageChanged(ui->emulatorLanguageComboBox->itemData(index).toString());
 }
 
 void SettingsDialog::OnCursorStateChanged(s16 index) {
@@ -666,8 +656,6 @@ void SettingsDialog::updateNoteTextEdit(const QString& elementName) {
         text = tr("Hide Cursor:\\nChoose when the cursor will disappear:\\nNever: You will always see the mouse.\\nidle: Set a time for it to disappear after being idle.\\nAlways: you will never see the mouse.");
     } else if (elementName == "idleTimeoutGroupBox") {
         text = tr("Hide Idle Cursor Timeout:\\nThe duration (seconds) after which the cursor that has been idle hides itself.");
-    } else if (elementName == "backButtonBehaviorGroupBox") {
-        text = tr("Back Button Behavior:\\nSets the controller's back button to emulate tapping the specified position on the PS4 touchpad.");
     }
 
     // Graphics
@@ -745,8 +733,6 @@ bool SettingsDialog::eventFilter(QObject* obj, QEvent* event) {
 
 void SettingsDialog::UpdateSettings() {
 
-    const QVector<std::string> TouchPadIndex = {"left", "center", "right", "none"};
-    Config::setBackButtonBehavior(TouchPadIndex[ui->backButtonBehaviorComboBox->currentIndex()]);
     Config::setIsFullscreen(screenModeMap.value(ui->displayModeComboBox->currentText()) !=
                             "Windowed");
     Config::setFullscreenMode(
@@ -886,4 +872,5 @@ void SettingsDialog::setDefaultValues() {
     } else {
         m_gui_settings->SetValue(gui::gen_updateChannel, "Nightly");
     }
+    m_gui_settings->SetValue(gui::gen_guiLanguage, "en_US");
 }
\ No newline at end of file
diff --git a/src/qt_gui/settings_dialog.h b/src/qt_gui/settings_dialog.h
index db1bcf772..d9fbcb214 100644
--- a/src/qt_gui/settings_dialog.h
+++ b/src/qt_gui/settings_dialog.h
@@ -32,7 +32,7 @@ public:
     int exec() override;
 
 signals:
-    void LanguageChanged(const std::string& locale);
+    void LanguageChanged(const QString& locale);
     void CompatibilityChanged();
     void BackgroundOpacityChanged(int opacity);
 
diff --git a/src/qt_gui/settings_dialog.ui b/src/qt_gui/settings_dialog.ui
index 20e26775d..8d239b58c 100644
--- a/src/qt_gui/settings_dialog.ui
+++ b/src/qt_gui/settings_dialog.ui
@@ -1613,36 +1613,6 @@
                <property name="bottomMargin">
                 <number>11</number>
                </property>
-               <item>
-                <widget class="QGroupBox" name="backButtonBehaviorGroupBox">
-                 <property name="enabled">
-                  <bool>true</bool>
-                 </property>
-                 <property name="sizePolicy">
-                  <sizepolicy hsizetype="Preferred" vsizetype="Fixed">
-                   <horstretch>0</horstretch>
-                   <verstretch>0</verstretch>
-                  </sizepolicy>
-                 </property>
-                 <property name="minimumSize">
-                  <size>
-                   <width>0</width>
-                   <height>0</height>
-                  </size>
-                 </property>
-                 <property name="title">
-                  <string>Back Button Behavior</string>
-                 </property>
-                 <layout class="QVBoxLayout" name="BackButtonLayout">
-                  <property name="leftMargin">
-                   <number>11</number>
-                  </property>
-                  <item>
-                   <widget class="QComboBox" name="backButtonBehaviorComboBox"/>
-                  </item>
-                 </layout>
-                </widget>
-               </item>
                <item>
                 <widget class="QCheckBox" name="motionControlsCheckBox">
                  <property name="text">
diff --git a/src/qt_gui/trophy_viewer.cpp b/src/qt_gui/trophy_viewer.cpp
index bed487605..675dba799 100644
--- a/src/qt_gui/trophy_viewer.cpp
+++ b/src/qt_gui/trophy_viewer.cpp
@@ -104,14 +104,16 @@ void TrophyViewer::updateTableFilters() {
     }
 }
 
-TrophyViewer::TrophyViewer(QString trophyPath, QString gameTrpPath, QString gameName,
+TrophyViewer::TrophyViewer(std::shared_ptr<gui_settings> gui_settings, QString trophyPath,
+                           QString gameTrpPath, QString gameName,
                            const QVector<TrophyGameInfo>& allTrophyGames)
-    : QMainWindow(), allTrophyGames_(allTrophyGames), currentGameName_(gameName) {
+    : QMainWindow(), allTrophyGames_(allTrophyGames), currentGameName_(gameName),
+      m_gui_settings(std::move(gui_settings)) {
     this->setWindowTitle(tr("Trophy Viewer") + " - " + currentGameName_);
     this->setAttribute(Qt::WA_DeleteOnClose);
     tabWidget = new QTabWidget(this);
 
-    auto lan = Config::getEmulatorLanguage();
+    auto lan = m_gui_settings->GetValue(gui::gen_guiLanguage).toString();
     if (lan == "en_US" || lan == "zh_CN" || lan == "zh_TW" || lan == "ja_JP" || lan == "ko_KR" ||
         lan == "lt_LT" || lan == "nb_NO" || lan == "nl_NL") {
         useEuropeanDateFormat = false;
@@ -463,7 +465,7 @@ void TrophyViewer::SetTableItem(QTableWidget* parent, int row, int column, QStri
         item->setTextAlignment(Qt::AlignCenter);
     item->setFont(QFont("Arial", 12, QFont::Bold));
 
-    Theme theme = static_cast<Theme>(Config::getMainWindowTheme());
+    Theme theme = static_cast<Theme>(m_gui_settings->GetValue(gui::gen_theme).toInt());
 
     if (theme == Theme::Light) {
         item->setForeground(QBrush(Qt::black));
diff --git a/src/qt_gui/trophy_viewer.h b/src/qt_gui/trophy_viewer.h
index c63171774..60ffe8420 100644
--- a/src/qt_gui/trophy_viewer.h
+++ b/src/qt_gui/trophy_viewer.h
@@ -23,6 +23,7 @@
 
 #include "common/types.h"
 #include "core/file_format/trp.h"
+#include "gui_settings.h"
 
 struct TrophyGameInfo {
     QString name;
@@ -34,7 +35,8 @@ class TrophyViewer : public QMainWindow {
     Q_OBJECT
 public:
     explicit TrophyViewer(
-        QString trophyPath, QString gameTrpPath, QString gameName = "",
+        std::shared_ptr<gui_settings> gui_settings, QString trophyPath, QString gameTrpPath,
+        QString gameName = "",
         const QVector<TrophyGameInfo>& allTrophyGames = QVector<TrophyGameInfo>());
 
     void updateTrophyInfo();
@@ -77,4 +79,5 @@ private:
         }
         return "Unknown";
     }
+    std::shared_ptr<gui_settings> m_gui_settings;
 };
diff --git a/src/sdl_window.cpp b/src/sdl_window.cpp
index e369240c6..735f14639 100644
--- a/src/sdl_window.cpp
+++ b/src/sdl_window.cpp
@@ -474,11 +474,16 @@ void WindowSDL::OnKeyboardMouseInput(const SDL_Event* event) {
             Input::ParseInputConfig(std::string(Common::ElfInfo::Instance().GameSerial()));
             return;
         }
-        // Toggle mouse capture and movement input
+        // Toggle mouse capture and joystick input emulation
         else if (input_id == SDLK_F7) {
-            Input::ToggleMouseEnabled();
             SDL_SetWindowRelativeMouseMode(this->GetSDLWindow(),
-                                           !SDL_GetWindowRelativeMouseMode(this->GetSDLWindow()));
+                                           Input::ToggleMouseModeTo(Input::MouseMode::Joystick));
+            return;
+        }
+        // Toggle mouse capture and gyro input emulation
+        else if (input_id == SDLK_F6) {
+            SDL_SetWindowRelativeMouseMode(this->GetSDLWindow(),
+                                           Input::ToggleMouseModeTo(Input::MouseMode::Gyro));
             return;
         }
         // Toggle fullscreen
diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp
index 93fb81df4..02f290140 100644
--- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp
@@ -271,7 +271,8 @@ void SetupCapabilities(const Info& info, const Profile& profile, EmitContext& ct
     if (info.has_image_query) {
         ctx.AddCapability(spv::Capability::ImageQuery);
     }
-    if (info.uses_atomic_float_min_max && profile.supports_image_fp32_atomic_min_max) {
+    if ((info.uses_image_atomic_float_min_max && profile.supports_image_fp32_atomic_min_max) ||
+        (info.uses_buffer_atomic_float_min_max && profile.supports_buffer_fp32_atomic_min_max)) {
         ctx.AddExtension("SPV_EXT_shader_atomic_float_min_max");
         ctx.AddCapability(spv::Capability::AtomicFloat32MinMaxEXT);
     }
diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_atomic.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_atomic.cpp
index 79f47a6a0..97e455ff8 100644
--- a/src/shader_recompiler/backend/spirv/emit_spirv_atomic.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_atomic.cpp
@@ -50,9 +50,17 @@ Id SharedAtomicU64(EmitContext& ctx, Id offset, Id value,
     });
 }
 
+template <bool is_float = false>
 Id BufferAtomicU32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id value,
                    Id (Sirit::Module::*atomic_func)(Id, Id, Id, Id, Id)) {
     const auto& buffer = ctx.buffers[handle];
+    const auto type = [&] {
+        if constexpr (is_float) {
+            return ctx.F32[1];
+        } else {
+            return ctx.U32[1];
+        }
+    }();
     if (Sirit::ValidId(buffer.offset)) {
         address = ctx.OpIAdd(ctx.U32[1], address, buffer.offset);
     }
@@ -60,8 +68,8 @@ Id BufferAtomicU32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id
     const auto [id, pointer_type] = buffer[EmitContext::PointerType::U32];
     const Id ptr = ctx.OpAccessChain(pointer_type, id, ctx.u32_zero_value, index);
     const auto [scope, semantics]{AtomicArgs(ctx)};
-    return AccessBoundsCheck<32>(ctx, index, buffer.size_dwords, [&] {
-        return (ctx.*atomic_func)(ctx.U32[1], ptr, scope, semantics, value);
+    return AccessBoundsCheck<32, 1, is_float>(ctx, index, buffer.size_dwords, [&] {
+        return (ctx.*atomic_func)(type, ptr, scope, semantics, value);
     });
 }
 
@@ -196,6 +204,24 @@ Id EmitBufferAtomicUMin32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id addre
     return BufferAtomicU32(ctx, inst, handle, address, value, &Sirit::Module::OpAtomicUMin);
 }
 
+Id EmitBufferAtomicFMin32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id value) {
+    if (ctx.profile.supports_buffer_fp32_atomic_min_max) {
+        return BufferAtomicU32<true>(ctx, inst, handle, address, value,
+                                     &Sirit::Module::OpAtomicFMin);
+    }
+
+    const auto u32_value = ctx.OpBitcast(ctx.U32[1], value);
+    const auto sign_bit_set =
+        ctx.OpBitFieldUExtract(ctx.U32[1], u32_value, ctx.ConstU32(31u), ctx.ConstU32(1u));
+
+    const auto result = ctx.OpSelect(
+        ctx.F32[1], sign_bit_set,
+        EmitBitCastF32U32(ctx, EmitBufferAtomicUMax32(ctx, inst, handle, address, u32_value)),
+        EmitBitCastF32U32(ctx, EmitBufferAtomicSMin32(ctx, inst, handle, address, u32_value)));
+
+    return result;
+}
+
 Id EmitBufferAtomicSMax32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id value) {
     return BufferAtomicU32(ctx, inst, handle, address, value, &Sirit::Module::OpAtomicSMax);
 }
@@ -204,6 +230,24 @@ Id EmitBufferAtomicUMax32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id addre
     return BufferAtomicU32(ctx, inst, handle, address, value, &Sirit::Module::OpAtomicUMax);
 }
 
+Id EmitBufferAtomicFMax32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id value) {
+    if (ctx.profile.supports_buffer_fp32_atomic_min_max) {
+        return BufferAtomicU32<true>(ctx, inst, handle, address, value,
+                                     &Sirit::Module::OpAtomicFMax);
+    }
+
+    const auto u32_value = ctx.OpBitcast(ctx.U32[1], value);
+    const auto sign_bit_set =
+        ctx.OpBitFieldUExtract(ctx.U32[1], u32_value, ctx.ConstU32(31u), ctx.ConstU32(1u));
+
+    const auto result = ctx.OpSelect(
+        ctx.F32[1], sign_bit_set,
+        EmitBitCastF32U32(ctx, EmitBufferAtomicUMin32(ctx, inst, handle, address, u32_value)),
+        EmitBitCastF32U32(ctx, EmitBufferAtomicSMax32(ctx, inst, handle, address, u32_value)));
+
+    return result;
+}
+
 Id EmitBufferAtomicInc32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address) {
     return BufferAtomicU32IncDec(ctx, inst, handle, address, &Sirit::Module::OpAtomicIIncrement);
 }
diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_instructions.h b/src/shader_recompiler/backend/spirv/emit_spirv_instructions.h
index daf1b973e..12d4fa671 100644
--- a/src/shader_recompiler/backend/spirv/emit_spirv_instructions.h
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_instructions.h
@@ -92,8 +92,10 @@ Id EmitBufferAtomicIAdd64(EmitContext& ctx, IR::Inst* inst, u32 handle, Id addre
 Id EmitBufferAtomicISub32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id value);
 Id EmitBufferAtomicSMin32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id value);
 Id EmitBufferAtomicUMin32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id value);
+Id EmitBufferAtomicFMin32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id value);
 Id EmitBufferAtomicSMax32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id value);
 Id EmitBufferAtomicUMax32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id value);
+Id EmitBufferAtomicFMax32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id value);
 Id EmitBufferAtomicInc32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address);
 Id EmitBufferAtomicDec32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address);
 Id EmitBufferAtomicAnd32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id value);
diff --git a/src/shader_recompiler/frontend/translate/vector_memory.cpp b/src/shader_recompiler/frontend/translate/vector_memory.cpp
index 3451358b6..a102ebf99 100644
--- a/src/shader_recompiler/frontend/translate/vector_memory.cpp
+++ b/src/shader_recompiler/frontend/translate/vector_memory.cpp
@@ -90,6 +90,10 @@ void Translator::EmitVectorMemory(const GcnInst& inst) {
         return BUFFER_ATOMIC(AtomicOp::Inc, inst);
     case Opcode::BUFFER_ATOMIC_DEC:
         return BUFFER_ATOMIC(AtomicOp::Dec, inst);
+    case Opcode::BUFFER_ATOMIC_FMIN:
+        return BUFFER_ATOMIC(AtomicOp::Fmin, inst);
+    case Opcode::BUFFER_ATOMIC_FMAX:
+        return BUFFER_ATOMIC(AtomicOp::Fmax, inst);
 
         // MIMG
         // Image load operations
@@ -357,6 +361,10 @@ void Translator::BUFFER_ATOMIC(AtomicOp op, const GcnInst& inst) {
             return ir.BufferAtomicInc(handle, address, buffer_info);
         case AtomicOp::Dec:
             return ir.BufferAtomicDec(handle, address, buffer_info);
+        case AtomicOp::Fmin:
+            return ir.BufferAtomicFMin(handle, address, vdata_val, buffer_info);
+        case AtomicOp::Fmax:
+            return ir.BufferAtomicFMax(handle, address, vdata_val, buffer_info);
         default:
             UNREACHABLE();
         }
diff --git a/src/shader_recompiler/info.h b/src/shader_recompiler/info.h
index f9b932c1d..6c931da31 100644
--- a/src/shader_recompiler/info.h
+++ b/src/shader_recompiler/info.h
@@ -215,7 +215,8 @@ struct Info {
     bool has_image_query{};
     bool has_perspective_interp{};
     bool has_linear_interp{};
-    bool uses_atomic_float_min_max{};
+    bool uses_buffer_atomic_float_min_max{};
+    bool uses_image_atomic_float_min_max{};
     bool uses_lane_id{};
     bool uses_group_quad{};
     bool uses_group_ballot{};
diff --git a/src/shader_recompiler/ir/ir_emitter.cpp b/src/shader_recompiler/ir/ir_emitter.cpp
index 82712c441..ab6535af2 100644
--- a/src/shader_recompiler/ir/ir_emitter.cpp
+++ b/src/shader_recompiler/ir/ir_emitter.cpp
@@ -504,12 +504,22 @@ Value IREmitter::BufferAtomicIMin(const Value& handle, const Value& address, con
                      : Inst(Opcode::BufferAtomicUMin32, Flags{info}, handle, address, value);
 }
 
+Value IREmitter::BufferAtomicFMin(const Value& handle, const Value& address, const Value& value,
+                                  BufferInstInfo info) {
+    return Inst(Opcode::BufferAtomicFMin32, Flags{info}, handle, address, value);
+}
+
 Value IREmitter::BufferAtomicIMax(const Value& handle, const Value& address, const Value& value,
                                   bool is_signed, BufferInstInfo info) {
     return is_signed ? Inst(Opcode::BufferAtomicSMax32, Flags{info}, handle, address, value)
                      : Inst(Opcode::BufferAtomicUMax32, Flags{info}, handle, address, value);
 }
 
+Value IREmitter::BufferAtomicFMax(const Value& handle, const Value& address, const Value& value,
+                                  BufferInstInfo info) {
+    return Inst(Opcode::BufferAtomicFMax32, Flags{info}, handle, address, value);
+}
+
 Value IREmitter::BufferAtomicInc(const Value& handle, const Value& address, BufferInstInfo info) {
     return Inst(Opcode::BufferAtomicInc32, Flags{info}, handle, address);
 }
diff --git a/src/shader_recompiler/ir/ir_emitter.h b/src/shader_recompiler/ir/ir_emitter.h
index 982c2dee4..9e2f79978 100644
--- a/src/shader_recompiler/ir/ir_emitter.h
+++ b/src/shader_recompiler/ir/ir_emitter.h
@@ -140,8 +140,12 @@ public:
                                          const Value& value, BufferInstInfo info);
     [[nodiscard]] Value BufferAtomicIMin(const Value& handle, const Value& address,
                                          const Value& value, bool is_signed, BufferInstInfo info);
+    [[nodiscard]] Value BufferAtomicFMin(const Value& handle, const Value& address,
+                                         const Value& value, BufferInstInfo info);
     [[nodiscard]] Value BufferAtomicIMax(const Value& handle, const Value& address,
                                          const Value& value, bool is_signed, BufferInstInfo info);
+    [[nodiscard]] Value BufferAtomicFMax(const Value& handle, const Value& address,
+                                         const Value& value, BufferInstInfo info);
     [[nodiscard]] Value BufferAtomicInc(const Value& handle, const Value& address,
                                         BufferInstInfo info);
     [[nodiscard]] Value BufferAtomicDec(const Value& handle, const Value& address,
diff --git a/src/shader_recompiler/ir/microinstruction.cpp b/src/shader_recompiler/ir/microinstruction.cpp
index c2311afea..1ea5c0967 100644
--- a/src/shader_recompiler/ir/microinstruction.cpp
+++ b/src/shader_recompiler/ir/microinstruction.cpp
@@ -71,8 +71,10 @@ bool Inst::MayHaveSideEffects() const noexcept {
     case Opcode::BufferAtomicISub32:
     case Opcode::BufferAtomicSMin32:
     case Opcode::BufferAtomicUMin32:
+    case Opcode::BufferAtomicFMin32:
     case Opcode::BufferAtomicSMax32:
     case Opcode::BufferAtomicUMax32:
+    case Opcode::BufferAtomicFMax32:
     case Opcode::BufferAtomicInc32:
     case Opcode::BufferAtomicDec32:
     case Opcode::BufferAtomicAnd32:
diff --git a/src/shader_recompiler/ir/opcodes.inc b/src/shader_recompiler/ir/opcodes.inc
index 0380cb0e6..179a01945 100644
--- a/src/shader_recompiler/ir/opcodes.inc
+++ b/src/shader_recompiler/ir/opcodes.inc
@@ -125,8 +125,10 @@ OPCODE(BufferAtomicIAdd64,                                  U64,            Opaq
 OPCODE(BufferAtomicISub32,                                  U32,            Opaque,         Opaque,         U32                                             )
 OPCODE(BufferAtomicSMin32,                                  U32,            Opaque,         Opaque,         U32                                             )
 OPCODE(BufferAtomicUMin32,                                  U32,            Opaque,         Opaque,         U32                                             )
+OPCODE(BufferAtomicFMin32,                                  U32,            Opaque,         Opaque,         F32                                             )
 OPCODE(BufferAtomicSMax32,                                  U32,            Opaque,         Opaque,         U32                                             )
 OPCODE(BufferAtomicUMax32,                                  U32,            Opaque,         Opaque,         U32                                             )
+OPCODE(BufferAtomicFMax32,                                  U32,            Opaque,         Opaque,         F32                                             )
 OPCODE(BufferAtomicInc32,                                   U32,            Opaque,         Opaque,                                                         )
 OPCODE(BufferAtomicDec32,                                   U32,            Opaque,         Opaque,                                                         )
 OPCODE(BufferAtomicAnd32,                                   U32,            Opaque,         Opaque,         U32,                                            )
diff --git a/src/shader_recompiler/ir/passes/resource_tracking_pass.cpp b/src/shader_recompiler/ir/passes/resource_tracking_pass.cpp
index a209f7126..e278d10f8 100644
--- a/src/shader_recompiler/ir/passes/resource_tracking_pass.cpp
+++ b/src/shader_recompiler/ir/passes/resource_tracking_pass.cpp
@@ -21,8 +21,10 @@ bool IsBufferAtomic(const IR::Inst& inst) {
     case IR::Opcode::BufferAtomicISub32:
     case IR::Opcode::BufferAtomicSMin32:
     case IR::Opcode::BufferAtomicUMin32:
+    case IR::Opcode::BufferAtomicFMin32:
     case IR::Opcode::BufferAtomicSMax32:
     case IR::Opcode::BufferAtomicUMax32:
+    case IR::Opcode::BufferAtomicFMax32:
     case IR::Opcode::BufferAtomicInc32:
     case IR::Opcode::BufferAtomicDec32:
     case IR::Opcode::BufferAtomicAnd32:
diff --git a/src/shader_recompiler/ir/passes/shader_info_collection_pass.cpp b/src/shader_recompiler/ir/passes/shader_info_collection_pass.cpp
index 4cd16d18f..b3b4ac36a 100644
--- a/src/shader_recompiler/ir/passes/shader_info_collection_pass.cpp
+++ b/src/shader_recompiler/ir/passes/shader_info_collection_pass.cpp
@@ -92,7 +92,11 @@ void Visit(Info& info, const IR::Inst& inst) {
         break;
     case IR::Opcode::ImageAtomicFMax32:
     case IR::Opcode::ImageAtomicFMin32:
-        info.uses_atomic_float_min_max = true;
+        info.uses_image_atomic_float_min_max = true;
+        break;
+    case IR::Opcode::BufferAtomicFMax32:
+    case IR::Opcode::BufferAtomicFMin32:
+        info.uses_buffer_atomic_float_min_max = true;
         break;
     case IR::Opcode::LaneId:
         info.uses_lane_id = true;
diff --git a/src/shader_recompiler/profile.h b/src/shader_recompiler/profile.h
index 7d313180f..bcdf86962 100644
--- a/src/shader_recompiler/profile.h
+++ b/src/shader_recompiler/profile.h
@@ -28,6 +28,7 @@ struct Profile {
     bool supports_native_cube_calc{};
     bool supports_trinary_minmax{};
     bool supports_robust_buffer_access{};
+    bool supports_buffer_fp32_atomic_min_max{};
     bool supports_image_fp32_atomic_min_max{};
     bool supports_workgroup_explicit_memory_layout{};
     bool has_broken_spirv_clamp{};
diff --git a/src/video_core/buffer_cache/buffer_cache.h b/src/video_core/buffer_cache/buffer_cache.h
index d7d753213..651ba84dc 100644
--- a/src/video_core/buffer_cache/buffer_cache.h
+++ b/src/video_core/buffer_cache/buffer_cache.h
@@ -9,7 +9,7 @@
 #include "common/slot_vector.h"
 #include "common/types.h"
 #include "video_core/buffer_cache/buffer.h"
-#include "video_core/buffer_cache/memory_tracker_base.h"
+#include "video_core/buffer_cache/memory_tracker.h"
 #include "video_core/buffer_cache/range_set.h"
 #include "video_core/multi_level_page_table.h"
 
diff --git a/src/video_core/buffer_cache/memory_tracker_base.h b/src/video_core/buffer_cache/memory_tracker.h
similarity index 99%
rename from src/video_core/buffer_cache/memory_tracker_base.h
rename to src/video_core/buffer_cache/memory_tracker.h
index c60aa9c80..37fafa2d6 100644
--- a/src/video_core/buffer_cache/memory_tracker_base.h
+++ b/src/video_core/buffer_cache/memory_tracker.h
@@ -9,7 +9,7 @@
 #include <vector>
 #include "common/debug.h"
 #include "common/types.h"
-#include "video_core/buffer_cache/word_manager.h"
+#include "video_core/buffer_cache/region_manager.h"
 
 namespace VideoCore {
 
diff --git a/src/video_core/buffer_cache/region_definitions.h b/src/video_core/buffer_cache/region_definitions.h
new file mode 100644
index 000000000..80c6afdc6
--- /dev/null
+++ b/src/video_core/buffer_cache/region_definitions.h
@@ -0,0 +1,28 @@
+// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
+// SPDX-License-Identifier: GPL-2.0-or-later
+
+#pragma once
+
+#include <array>
+#include "common/bit_array.h"
+#include "common/types.h"
+
+namespace VideoCore {
+
+constexpr u64 PAGES_PER_WORD = 64;
+constexpr u64 BYTES_PER_PAGE = 4_KB;
+
+constexpr u64 HIGHER_PAGE_BITS = 22;
+constexpr u64 HIGHER_PAGE_SIZE = 1ULL << HIGHER_PAGE_BITS;
+constexpr u64 HIGHER_PAGE_MASK = HIGHER_PAGE_SIZE - 1ULL;
+constexpr u64 NUM_REGION_PAGES = HIGHER_PAGE_SIZE / BYTES_PER_PAGE;
+
+enum class Type {
+    CPU,
+    GPU,
+    Writeable,
+};
+
+using RegionBits = Common::BitArray<NUM_REGION_PAGES>;
+
+} // namespace VideoCore
\ No newline at end of file
diff --git a/src/video_core/buffer_cache/region_manager.h b/src/video_core/buffer_cache/region_manager.h
new file mode 100644
index 000000000..07ffee36b
--- /dev/null
+++ b/src/video_core/buffer_cache/region_manager.h
@@ -0,0 +1,208 @@
+// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
+// SPDX-License-Identifier: GPL-2.0-or-later
+
+#pragma once
+
+#include <mutex>
+#include <utility>
+#include "common/div_ceil.h"
+
+#ifdef __linux__
+#include "common/adaptive_mutex.h"
+#else
+#include "common/spin_lock.h"
+#endif
+#include "common/debug.h"
+#include "common/types.h"
+#include "video_core/buffer_cache/region_definitions.h"
+#include "video_core/page_manager.h"
+
+namespace VideoCore {
+
+/**
+ * Allows tracking CPU and GPU modification of pages in a contigious 4MB virtual address region.
+ * Information is stored in bitsets for spacial locality and fast update of single pages.
+ */
+class RegionManager {
+public:
+    explicit RegionManager(PageManager* tracker_, VAddr cpu_addr_)
+        : tracker{tracker_}, cpu_addr{cpu_addr_} {
+        cpu.Fill();
+        gpu.Clear();
+        writeable.Fill();
+    }
+    explicit RegionManager() = default;
+
+    void SetCpuAddress(VAddr new_cpu_addr) {
+        cpu_addr = new_cpu_addr;
+    }
+
+    VAddr GetCpuAddr() const {
+        return cpu_addr;
+    }
+
+    static constexpr size_t SanitizeAddress(size_t address) {
+        return static_cast<size_t>(std::max<s64>(static_cast<s64>(address), 0LL));
+    }
+
+    template <Type type>
+    RegionBits& GetRegionBits() noexcept {
+        static_assert(type != Type::Writeable);
+        if constexpr (type == Type::CPU) {
+            return cpu;
+        } else if constexpr (type == Type::GPU) {
+            return gpu;
+        } else if constexpr (type == Type::Writeable) {
+            return writeable;
+        } else {
+            static_assert(false, "Invalid type");
+        }
+    }
+
+    template <Type type>
+    const RegionBits& GetRegionBits() const noexcept {
+        static_assert(type != Type::Writeable);
+        if constexpr (type == Type::CPU) {
+            return cpu;
+        } else if constexpr (type == Type::GPU) {
+            return gpu;
+        } else if constexpr (type == Type::Writeable) {
+            return writeable;
+        } else {
+            static_assert(false, "Invalid type");
+        }
+    }
+
+    /**
+     * Change the state of a range of pages
+     *
+     * @param dirty_addr    Base address to mark or unmark as modified
+     * @param size          Size in bytes to mark or unmark as modified
+     */
+    template <Type type, bool enable>
+    void ChangeRegionState(u64 dirty_addr, u64 size) noexcept(type == Type::GPU) {
+        RENDERER_TRACE;
+        const size_t offset = dirty_addr - cpu_addr;
+        const size_t start_page = SanitizeAddress(offset) / BYTES_PER_PAGE;
+        const size_t end_page = Common::DivCeil(SanitizeAddress(offset + size), BYTES_PER_PAGE);
+        if (start_page >= NUM_REGION_PAGES || end_page <= start_page) {
+            return;
+        }
+        std::scoped_lock lk{lock};
+        static_assert(type != Type::Writeable);
+
+        RegionBits& bits = GetRegionBits<type>();
+        if constexpr (enable) {
+            bits.SetRange(start_page, end_page);
+        } else {
+            bits.UnsetRange(start_page, end_page);
+        }
+        if constexpr (type == Type::CPU) {
+            UpdateProtection<!enable>();
+        }
+    }
+
+    /**
+     * Loop over each page in the given range, turn off those bits and notify the tracker if
+     * needed. Call the given function on each turned off range.
+     *
+     * @param query_cpu_range Base CPU address to loop over
+     * @param size            Size in bytes of the CPU range to loop over
+     * @param func            Function to call for each turned off region
+     */
+    template <Type type, bool clear>
+    void ForEachModifiedRange(VAddr query_cpu_range, s64 size, auto&& func) {
+        RENDERER_TRACE;
+        const size_t offset = query_cpu_range - cpu_addr;
+        const size_t start_page = SanitizeAddress(offset) / BYTES_PER_PAGE;
+        const size_t end_page = Common::DivCeil(SanitizeAddress(offset + size), BYTES_PER_PAGE);
+        if (start_page >= NUM_REGION_PAGES || end_page <= start_page) {
+            return;
+        }
+        std::scoped_lock lk{lock};
+        static_assert(type != Type::Writeable);
+
+        RegionBits& bits = GetRegionBits<type>();
+        RegionBits mask(bits, start_page, end_page);
+
+        // TODO: this will not be needed once we handle readbacks
+        if constexpr (type == Type::GPU) {
+            mask &= ~writeable;
+        }
+
+        for (const auto& [start, end] : mask) {
+            func(cpu_addr + start * BYTES_PER_PAGE, (end - start) * BYTES_PER_PAGE);
+        }
+
+        if constexpr (clear) {
+            bits.UnsetRange(start_page, end_page);
+            if constexpr (type == Type::CPU) {
+                UpdateProtection<true>();
+            }
+        }
+    }
+
+    /**
+     * Returns true when a region has been modified
+     *
+     * @param offset Offset in bytes from the start of the buffer
+     * @param size   Size in bytes of the region to query for modifications
+     */
+    template <Type type>
+    [[nodiscard]] bool IsRegionModified(u64 offset, u64 size) const noexcept {
+        RENDERER_TRACE;
+        const size_t start_page = SanitizeAddress(offset) / BYTES_PER_PAGE;
+        const size_t end_page = Common::DivCeil(SanitizeAddress(offset + size), BYTES_PER_PAGE);
+        if (start_page >= NUM_REGION_PAGES || end_page <= start_page) {
+            return false;
+        }
+        // std::scoped_lock lk{lock}; // Is this needed?
+        static_assert(type != Type::Writeable);
+
+        const RegionBits& bits = GetRegionBits<type>();
+        RegionBits test(bits, start_page, end_page);
+
+        // TODO: this will not be needed once we handle readbacks
+        if constexpr (type == Type::GPU) {
+            test &= ~writeable;
+        }
+
+        return test.Any();
+    }
+
+private:
+    /**
+     * Notify tracker about changes in the CPU tracking state of a word in the buffer
+     *
+     * @param word_index   Index to the word to notify to the tracker
+     * @param current_bits Current state of the word
+     * @param new_bits     New state of the word
+     *
+     * @tparam add_to_tracker True when the tracker should start tracking the new pages
+     */
+    template <bool add_to_tracker>
+    void UpdateProtection() {
+        RENDERER_TRACE;
+        RegionBits mask = cpu ^ writeable;
+
+        if (mask.None()) {
+            return; // No changes to the CPU tracking state
+        }
+
+        writeable = cpu;
+        tracker->UpdatePageWatchersForRegion<add_to_tracker>(cpu_addr, mask);
+    }
+
+#ifdef PTHREAD_ADAPTIVE_MUTEX_INITIALIZER_NP
+    Common::AdaptiveMutex lock;
+#else
+    Common::SpinLock lock;
+#endif
+    PageManager* tracker;
+    VAddr cpu_addr = 0;
+    RegionBits cpu;
+    RegionBits gpu;
+    RegionBits writeable;
+};
+
+} // namespace VideoCore
diff --git a/src/video_core/buffer_cache/word_manager.h b/src/video_core/buffer_cache/word_manager.h
deleted file mode 100644
index 51a912c62..000000000
--- a/src/video_core/buffer_cache/word_manager.h
+++ /dev/null
@@ -1,296 +0,0 @@
-// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
-// SPDX-License-Identifier: GPL-2.0-or-later
-
-#pragma once
-
-#include <array>
-#include <mutex>
-#include <span>
-#include <utility>
-
-#ifdef __linux__
-#include "common/adaptive_mutex.h"
-#else
-#include "common/spin_lock.h"
-#endif
-#include "common/debug.h"
-#include "common/types.h"
-#include "video_core/page_manager.h"
-
-namespace VideoCore {
-
-constexpr u64 PAGES_PER_WORD = 64;
-constexpr u64 BYTES_PER_PAGE = 4_KB;
-constexpr u64 BYTES_PER_WORD = PAGES_PER_WORD * BYTES_PER_PAGE;
-
-constexpr u64 HIGHER_PAGE_BITS = 22;
-constexpr u64 HIGHER_PAGE_SIZE = 1ULL << HIGHER_PAGE_BITS;
-constexpr u64 HIGHER_PAGE_MASK = HIGHER_PAGE_SIZE - 1ULL;
-constexpr u64 NUM_REGION_WORDS = HIGHER_PAGE_SIZE / BYTES_PER_WORD;
-
-enum class Type {
-    CPU,
-    GPU,
-    Untracked,
-};
-
-using WordsArray = std::array<u64, NUM_REGION_WORDS>;
-
-/**
- * Allows tracking CPU and GPU modification of pages in a contigious 4MB virtual address region.
- * Information is stored in bitsets for spacial locality and fast update of single pages.
- */
-class RegionManager {
-public:
-    explicit RegionManager(PageManager* tracker_, VAddr cpu_addr_)
-        : tracker{tracker_}, cpu_addr{cpu_addr_} {
-        cpu.fill(~u64{0});
-        gpu.fill(0);
-        untracked.fill(~u64{0});
-    }
-    explicit RegionManager() = default;
-
-    void SetCpuAddress(VAddr new_cpu_addr) {
-        cpu_addr = new_cpu_addr;
-    }
-
-    VAddr GetCpuAddr() const {
-        return cpu_addr;
-    }
-
-    static constexpr u64 ExtractBits(u64 word, size_t page_start, size_t page_end) {
-        constexpr size_t number_bits = sizeof(u64) * 8;
-        const size_t limit_page_end = number_bits - std::min(page_end, number_bits);
-        u64 bits = (word >> page_start) << page_start;
-        bits = (bits << limit_page_end) >> limit_page_end;
-        return bits;
-    }
-
-    static constexpr std::pair<size_t, size_t> GetWordPage(VAddr address) {
-        const size_t converted_address = static_cast<size_t>(address);
-        const size_t word_number = converted_address / BYTES_PER_WORD;
-        const size_t amount_pages = converted_address % BYTES_PER_WORD;
-        return std::make_pair(word_number, amount_pages / BYTES_PER_PAGE);
-    }
-
-    template <typename Func>
-    void IterateWords(size_t offset, size_t size, Func&& func) const {
-        RENDERER_TRACE;
-        using FuncReturn = std::invoke_result_t<Func, std::size_t, u64>;
-        static constexpr bool BOOL_BREAK = std::is_same_v<FuncReturn, bool>;
-        const size_t start = static_cast<size_t>(std::max<s64>(static_cast<s64>(offset), 0LL));
-        const size_t end = static_cast<size_t>(std::max<s64>(static_cast<s64>(offset + size), 0LL));
-        if (start >= HIGHER_PAGE_SIZE || end <= start) {
-            return;
-        }
-        auto [start_word, start_page] = GetWordPage(start);
-        auto [end_word, end_page] = GetWordPage(end + BYTES_PER_PAGE - 1ULL);
-        constexpr size_t num_words = NUM_REGION_WORDS;
-        start_word = std::min(start_word, num_words);
-        end_word = std::min(end_word, num_words);
-        const size_t diff = end_word - start_word;
-        end_word += (end_page + PAGES_PER_WORD - 1ULL) / PAGES_PER_WORD;
-        end_word = std::min(end_word, num_words);
-        end_page += diff * PAGES_PER_WORD;
-        constexpr u64 base_mask{~0ULL};
-        for (size_t word_index = start_word; word_index < end_word; word_index++) {
-            const u64 mask = ExtractBits(base_mask, start_page, end_page);
-            start_page = 0;
-            end_page -= PAGES_PER_WORD;
-            if constexpr (BOOL_BREAK) {
-                if (func(word_index, mask)) {
-                    return;
-                }
-            } else {
-                func(word_index, mask);
-            }
-        }
-    }
-
-    void IteratePages(u64 mask, auto&& func) const {
-        RENDERER_TRACE;
-        size_t offset = 0;
-        while (mask != 0) {
-            const size_t empty_bits = std::countr_zero(mask);
-            offset += empty_bits;
-            mask >>= empty_bits;
-
-            const size_t continuous_bits = std::countr_one(mask);
-            func(offset, continuous_bits);
-            mask = continuous_bits < PAGES_PER_WORD ? (mask >> continuous_bits) : 0;
-            offset += continuous_bits;
-        }
-    }
-
-    /**
-     * Change the state of a range of pages
-     *
-     * @param dirty_addr    Base address to mark or unmark as modified
-     * @param size          Size in bytes to mark or unmark as modified
-     */
-    template <Type type, bool enable>
-    void ChangeRegionState(u64 dirty_addr, u64 size) noexcept(type == Type::GPU) {
-        std::scoped_lock lk{lock};
-        std::span<u64> state_words = Span<type>();
-        IterateWords(dirty_addr - cpu_addr, size, [&](size_t index, u64 mask) {
-            if constexpr (type == Type::CPU) {
-                UpdateProtection<!enable>(index, untracked[index], mask);
-            }
-            if constexpr (enable) {
-                state_words[index] |= mask;
-                if constexpr (type == Type::CPU) {
-                    untracked[index] |= mask;
-                }
-            } else {
-                state_words[index] &= ~mask;
-                if constexpr (type == Type::CPU) {
-                    untracked[index] &= ~mask;
-                }
-            }
-        });
-    }
-
-    /**
-     * Loop over each page in the given range, turn off those bits and notify the tracker if
-     * needed. Call the given function on each turned off range.
-     *
-     * @param query_cpu_range Base CPU address to loop over
-     * @param size            Size in bytes of the CPU range to loop over
-     * @param func            Function to call for each turned off region
-     */
-    template <Type type, bool clear>
-    void ForEachModifiedRange(VAddr query_cpu_range, s64 size, auto&& func) {
-        RENDERER_TRACE;
-        std::scoped_lock lk{lock};
-        static_assert(type != Type::Untracked);
-
-        std::span<u64> state_words = Span<type>();
-        const size_t offset = query_cpu_range - cpu_addr;
-        bool pending = false;
-        size_t pending_offset{};
-        size_t pending_pointer{};
-        const auto release = [&]() {
-            func(cpu_addr + pending_offset * BYTES_PER_PAGE,
-                 (pending_pointer - pending_offset) * BYTES_PER_PAGE);
-        };
-        IterateWords(offset, size, [&](size_t index, u64 mask) {
-            RENDERER_TRACE;
-            if constexpr (type == Type::GPU) {
-                mask &= ~untracked[index];
-            }
-            const u64 word = state_words[index] & mask;
-            if constexpr (clear) {
-                if constexpr (type == Type::CPU) {
-                    UpdateProtection<true>(index, untracked[index], mask);
-                    untracked[index] &= ~mask;
-                }
-                state_words[index] &= ~mask;
-            }
-            const size_t base_offset = index * PAGES_PER_WORD;
-            IteratePages(word, [&](size_t pages_offset, size_t pages_size) {
-                RENDERER_TRACE;
-                const auto reset = [&]() {
-                    pending_offset = base_offset + pages_offset;
-                    pending_pointer = base_offset + pages_offset + pages_size;
-                };
-                if (!pending) {
-                    reset();
-                    pending = true;
-                    return;
-                }
-                if (pending_pointer == base_offset + pages_offset) {
-                    pending_pointer += pages_size;
-                    return;
-                }
-                release();
-                reset();
-            });
-        });
-        if (pending) {
-            release();
-        }
-    }
-
-    /**
-     * Returns true when a region has been modified
-     *
-     * @param offset Offset in bytes from the start of the buffer
-     * @param size   Size in bytes of the region to query for modifications
-     */
-    template <Type type>
-    [[nodiscard]] bool IsRegionModified(u64 offset, u64 size) const noexcept {
-        static_assert(type != Type::Untracked);
-
-        const std::span<const u64> state_words = Span<type>();
-        bool result = false;
-        IterateWords(offset, size, [&](size_t index, u64 mask) {
-            if constexpr (type == Type::GPU) {
-                mask &= ~untracked[index];
-            }
-            const u64 word = state_words[index] & mask;
-            if (word != 0) {
-                result = true;
-                return true;
-            }
-            return false;
-        });
-        return result;
-    }
-
-private:
-    /**
-     * Notify tracker about changes in the CPU tracking state of a word in the buffer
-     *
-     * @param word_index   Index to the word to notify to the tracker
-     * @param current_bits Current state of the word
-     * @param new_bits     New state of the word
-     *
-     * @tparam add_to_tracker True when the tracker should start tracking the new pages
-     */
-    template <bool add_to_tracker>
-    void UpdateProtection(u64 word_index, u64 current_bits, u64 new_bits) const {
-        RENDERER_TRACE;
-        constexpr s32 delta = add_to_tracker ? 1 : -1;
-        u64 changed_bits = (add_to_tracker ? current_bits : ~current_bits) & new_bits;
-        VAddr addr = cpu_addr + word_index * BYTES_PER_WORD;
-        IteratePages(changed_bits, [&](size_t offset, size_t size) {
-            tracker->UpdatePageWatchers<delta>(addr + offset * BYTES_PER_PAGE,
-                                               size * BYTES_PER_PAGE);
-        });
-    }
-
-    template <Type type>
-    std::span<u64> Span() noexcept {
-        if constexpr (type == Type::CPU) {
-            return cpu;
-        } else if constexpr (type == Type::GPU) {
-            return gpu;
-        } else if constexpr (type == Type::Untracked) {
-            return untracked;
-        }
-    }
-
-    template <Type type>
-    std::span<const u64> Span() const noexcept {
-        if constexpr (type == Type::CPU) {
-            return cpu;
-        } else if constexpr (type == Type::GPU) {
-            return gpu;
-        } else if constexpr (type == Type::Untracked) {
-            return untracked;
-        }
-    }
-
-#ifdef PTHREAD_ADAPTIVE_MUTEX_INITIALIZER_NP
-    Common::AdaptiveMutex lock;
-#else
-    Common::SpinLock lock;
-#endif
-    PageManager* tracker;
-    VAddr cpu_addr = 0;
-    WordsArray cpu;
-    WordsArray gpu;
-    WordsArray untracked;
-};
-
-} // namespace VideoCore
diff --git a/src/video_core/page_manager.cpp b/src/video_core/page_manager.cpp
index 39c03e7da..145779070 100644
--- a/src/video_core/page_manager.cpp
+++ b/src/video_core/page_manager.cpp
@@ -48,19 +48,15 @@ struct PageManager::Impl {
         u8 AddDelta() {
             if constexpr (delta == 1) {
                 return ++num_watchers;
-            } else {
+            } else if constexpr (delta == -1) {
                 ASSERT_MSG(num_watchers > 0, "Not enough watchers");
                 return --num_watchers;
+            } else {
+                return num_watchers;
             }
         }
     };
 
-    struct UpdateProtectRange {
-        VAddr addr;
-        u64 size;
-        Core::MemoryPermission perms;
-    };
-
     static constexpr size_t ADDRESS_BITS = 40;
     static constexpr size_t NUM_ADDRESS_PAGES = 1ULL << (40 - PAGE_BITS);
     inline static Vulkan::Rasterizer* rasterizer;
@@ -190,66 +186,122 @@ struct PageManager::Impl {
     }
 
 #endif
-    template <s32 delta>
+    template <bool track>
     void UpdatePageWatchers(VAddr addr, u64 size) {
         RENDERER_TRACE;
-        boost::container::small_vector<UpdateProtectRange, 16> update_ranges;
-        {
-            std::scoped_lock lk(lock);
 
-            size_t page = addr >> PAGE_BITS;
-            auto perms = cached_pages[page].Perm();
-            u64 range_begin = 0;
-            u64 range_bytes = 0;
+        size_t page = addr >> PAGE_BITS;
+        auto perms = cached_pages[page].Perm();
+        u64 range_begin = 0;
+        u64 range_bytes = 0;
 
-            const auto release_pending = [&] {
-                if (range_bytes > 0) {
-                    RENDERER_TRACE;
-                    // Add pending (un)protect action
-                    update_ranges.push_back({range_begin << PAGE_BITS, range_bytes, perms});
-                    range_bytes = 0;
-                }
-            };
+        const auto release_pending = [&] {
+            if (range_bytes > 0) {
+                RENDERER_TRACE;
+                // Perform pending (un)protect action
+                Protect(range_begin << PAGE_BITS, range_bytes, perms);
+                range_bytes = 0;
+            }
+        };
 
-            // Iterate requested pages
-            const u64 page_end = Common::DivCeil(addr + size, PAGE_SIZE);
-            const u64 aligned_addr = page << PAGE_BITS;
-            const u64 aligned_end = page_end << PAGE_BITS;
-            ASSERT_MSG(rasterizer->IsMapped(aligned_addr, aligned_end - aligned_addr),
-                       "Attempted to track non-GPU memory at address {:#x}, size {:#x}.",
-                       aligned_addr, aligned_end - aligned_addr);
+        std::scoped_lock lk(lock);
 
-            for (; page != page_end; ++page) {
-                PageState& state = cached_pages[page];
+        // Iterate requested pages
+        const u64 page_end = Common::DivCeil(addr + size, PAGE_SIZE);
+        const u64 aligned_addr = page << PAGE_BITS;
+        const u64 aligned_end = page_end << PAGE_BITS;
+        ASSERT_MSG(rasterizer->IsMapped(aligned_addr, aligned_end - aligned_addr),
+                   "Attempted to track non-GPU memory at address {:#x}, size {:#x}.", aligned_addr,
+                   aligned_end - aligned_addr);
 
-                // Apply the change to the page state
-                const u8 new_count = state.AddDelta<delta>();
+        for (; page != page_end; ++page) {
+            PageState& state = cached_pages[page];
 
+            // Apply the change to the page state
+            const u8 new_count = state.AddDelta<track ? 1 : -1>();
+
+            if (auto new_perms = state.Perm(); new_perms != perms) [[unlikely]] {
                 // If the protection changed add pending (un)protect action
-                if (auto new_perms = state.Perm(); new_perms != perms) [[unlikely]] {
-                    release_pending();
-                    perms = new_perms;
-                }
-
-                // If the page must be (un)protected, add it to the pending range
-                if ((new_count == 0 && delta < 0) || (new_count == 1 && delta > 0)) {
-                    if (range_bytes == 0) {
-                        range_begin = page;
-                    }
-                    range_bytes += PAGE_SIZE;
-                } else {
-                    release_pending();
-                }
+                release_pending();
+                perms = new_perms;
+            } else if (range_bytes != 0) {
+                // If the protection did not change, extend the current range
+                range_bytes += PAGE_SIZE;
             }
 
-            // Add pending (un)protect action
-            release_pending();
+            // Only start a new range if the page must be (un)protected
+            if (range_bytes == 0 && ((new_count == 0 && !track) || (new_count == 1 && track))) {
+                range_begin = page;
+                range_bytes = PAGE_SIZE;
+            }
         }
 
-        // Flush deferred protects
-        for (const auto& range : update_ranges) {
-            Protect(range.addr, range.size, range.perms);
+        // Add pending (un)protect action
+        release_pending();
+    }
+
+    template <bool track>
+    void UpdatePageWatchersForRegion(VAddr base_addr, RegionBits& mask) {
+        RENDERER_TRACE;
+        auto start_range = mask.FirstRange();
+        auto end_range = mask.LastRange();
+
+        if (start_range.second == end_range.second) {
+            // Optimization: if all pages are contiguous, use the regular UpdatePageWatchers
+            const VAddr start_addr = base_addr + (start_range.first << PAGE_BITS);
+            const u64 size = (start_range.second - start_range.first) << PAGE_BITS;
+
+            UpdatePageWatchers<track>(start_addr, size);
+            return;
         }
+
+        size_t base_page = (base_addr >> PAGE_BITS);
+        auto perms = cached_pages[base_page + start_range.first].Perm();
+        u64 range_begin = 0;
+        u64 range_bytes = 0;
+
+        const auto release_pending = [&] {
+            if (range_bytes > 0) {
+                RENDERER_TRACE;
+                // Perform pending (un)protect action
+                Protect((range_begin << PAGE_BITS), range_bytes, perms);
+                range_bytes = 0;
+            }
+        };
+
+        std::scoped_lock lk(lock);
+
+        // Iterate pages
+        for (size_t page = start_range.first; page < end_range.second; ++page) {
+            PageState& state = cached_pages[base_page + page];
+            const bool update = mask.Get(page);
+
+            // Apply the change to the page state
+            const u8 new_count = update ? state.AddDelta<track ? 1 : -1>() : state.AddDelta<0>();
+
+            if (auto new_perms = state.Perm(); new_perms != perms) [[unlikely]] {
+                // If the protection changed add pending (un)protect action
+                release_pending();
+                perms = new_perms;
+            } else if (range_bytes != 0) {
+                // If the protection did not change, extend the current range
+                range_bytes += PAGE_SIZE;
+            }
+
+            // If the page is not being updated, skip it
+            if (!update) {
+                continue;
+            }
+
+            // Only start a new range if the page must be (un)protected
+            if (range_bytes == 0 && ((new_count == 0 && !track) || (new_count == 1 && track))) {
+                range_begin = base_page + page;
+                range_bytes = PAGE_SIZE;
+            }
+        }
+
+        // Add pending (un)protect action
+        release_pending();
     }
 
     std::array<PageState, NUM_ADDRESS_PAGES> cached_pages{};
@@ -273,12 +325,21 @@ void PageManager::OnGpuUnmap(VAddr address, size_t size) {
     impl->OnUnmap(address, size);
 }
 
-template <s32 delta>
+template <bool track>
 void PageManager::UpdatePageWatchers(VAddr addr, u64 size) const {
-    impl->UpdatePageWatchers<delta>(addr, size);
+    impl->UpdatePageWatchers<track>(addr, size);
 }
 
-template void PageManager::UpdatePageWatchers<1>(VAddr addr, u64 size) const;
-template void PageManager::UpdatePageWatchers<-1>(VAddr addr, u64 size) const;
+template <bool track>
+void PageManager::UpdatePageWatchersForRegion(VAddr base_addr, RegionBits& mask) const {
+    impl->UpdatePageWatchersForRegion<track>(base_addr, mask);
+}
+
+template void PageManager::UpdatePageWatchers<true>(VAddr addr, u64 size) const;
+template void PageManager::UpdatePageWatchers<false>(VAddr addr, u64 size) const;
+template void PageManager::UpdatePageWatchersForRegion<true>(VAddr base_addr,
+                                                             RegionBits& mask) const;
+template void PageManager::UpdatePageWatchersForRegion<false>(VAddr base_addr,
+                                                              RegionBits& mask) const;
 
 } // namespace VideoCore
diff --git a/src/video_core/page_manager.h b/src/video_core/page_manager.h
index 98dd099af..157b34984 100644
--- a/src/video_core/page_manager.h
+++ b/src/video_core/page_manager.h
@@ -6,6 +6,7 @@
 #include <memory>
 #include "common/alignment.h"
 #include "common/types.h"
+#include "video_core/buffer_cache//region_definitions.h"
 
 namespace Vulkan {
 class Rasterizer;
@@ -28,9 +29,14 @@ public:
     void OnGpuUnmap(VAddr address, size_t size);
 
     /// Updates watches in the pages touching the specified region.
-    template <s32 delta>
+    template <bool track>
     void UpdatePageWatchers(VAddr addr, u64 size) const;
 
+    /// Updates watches in the pages touching the specified region
+    /// using a mask.
+    template <bool track>
+    void UpdatePageWatchersForRegion(VAddr base_addr, RegionBits& mask) const;
+
     /// Returns page aligned address.
     static constexpr VAddr GetPageAddr(VAddr addr) {
         return Common::AlignDown(addr, PAGE_SIZE);
diff --git a/src/video_core/renderer_vulkan/vk_instance.cpp b/src/video_core/renderer_vulkan/vk_instance.cpp
index fb489ec78..61ddd3f05 100644
--- a/src/video_core/renderer_vulkan/vk_instance.cpp
+++ b/src/video_core/renderer_vulkan/vk_instance.cpp
@@ -281,6 +281,8 @@ bool Instance::CreateDevice() {
     if (shader_atomic_float2) {
         shader_atomic_float2_features =
             feature_chain.get<vk::PhysicalDeviceShaderAtomicFloat2FeaturesEXT>();
+        LOG_INFO(Render_Vulkan, "- shaderBufferFloat32AtomicMinMax: {}",
+                 shader_atomic_float2_features.shaderBufferFloat32AtomicMinMax);
         LOG_INFO(Render_Vulkan, "- shaderImageFloat32AtomicMinMax: {}",
                  shader_atomic_float2_features.shaderImageFloat32AtomicMinMax);
     }
@@ -433,6 +435,8 @@ bool Instance::CreateDevice() {
             .legacyVertexAttributes = true,
         },
         vk::PhysicalDeviceShaderAtomicFloat2FeaturesEXT{
+            .shaderBufferFloat32AtomicMinMax =
+                shader_atomic_float2_features.shaderBufferFloat32AtomicMinMax,
             .shaderImageFloat32AtomicMinMax =
                 shader_atomic_float2_features.shaderImageFloat32AtomicMinMax,
         },
diff --git a/src/video_core/renderer_vulkan/vk_instance.h b/src/video_core/renderer_vulkan/vk_instance.h
index c687e6f67..991bfb031 100644
--- a/src/video_core/renderer_vulkan/vk_instance.h
+++ b/src/video_core/renderer_vulkan/vk_instance.h
@@ -165,6 +165,13 @@ public:
         return amd_shader_trinary_minmax;
     }
 
+    /// Returns true when the shaderBufferFloat32AtomicMinMax feature of
+    /// VK_EXT_shader_atomic_float2 is supported.
+    bool IsShaderAtomicFloatBuffer32MinMaxSupported() const {
+        return shader_atomic_float2 &&
+               shader_atomic_float2_features.shaderBufferFloat32AtomicMinMax;
+    }
+
     /// Returns true when the shaderImageFloat32AtomicMinMax feature of
     /// VK_EXT_shader_atomic_float2 is supported.
     bool IsShaderAtomicFloatImage32MinMaxSupported() const {
@@ -324,6 +331,9 @@ public:
         return driver_id != vk::DriverId::eMoltenvk;
     }
 
+    /// Determines if a format is supported for a set of feature flags.
+    [[nodiscard]] bool IsFormatSupported(vk::Format format, vk::FormatFeatureFlags2 flags) const;
+
 private:
     /// Creates the logical device opportunistically enabling extensions
     bool CreateDevice();
@@ -338,9 +348,6 @@ private:
     /// Gets the supported feature flags for a format.
     [[nodiscard]] vk::FormatFeatureFlags2 GetFormatFeatureFlags(vk::Format format) const;
 
-    /// Determines if a format is supported for a set of feature flags.
-    [[nodiscard]] bool IsFormatSupported(vk::Format format, vk::FormatFeatureFlags2 flags) const;
-
 private:
     vk::UniqueInstance instance;
     vk::PhysicalDevice physical_device;
diff --git a/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp b/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp
index 2c3f4ba2f..1d8ac4823 100644
--- a/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp
+++ b/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp
@@ -216,6 +216,8 @@ PipelineCache::PipelineCache(const Instance& instance_, Scheduler& scheduler_,
         .supports_trinary_minmax = instance_.IsAmdShaderTrinaryMinMaxSupported(),
         // TODO: Emitted bounds checks cause problems with phi control flow; needs to be fixed.
         .supports_robust_buffer_access = true, // instance_.IsRobustBufferAccess2Supported(),
+        .supports_buffer_fp32_atomic_min_max =
+            instance_.IsShaderAtomicFloatBuffer32MinMaxSupported(),
         .supports_image_fp32_atomic_min_max = instance_.IsShaderAtomicFloatImage32MinMaxSupported(),
         .supports_workgroup_explicit_memory_layout =
             instance_.IsWorkgroupMemoryExplicitLayoutSupported(),
@@ -346,8 +348,15 @@ bool PipelineCache::RefreshGraphicsKey() {
                                         col_buf.GetDataFmt() == AmdGpu::DataFormat::Format8_8 ||
                                         col_buf.GetDataFmt() == AmdGpu::DataFormat::Format8_8_8_8);
 
-        key.color_formats[remapped_cb] =
+        const auto format =
             LiverpoolToVK::SurfaceFormat(col_buf.GetDataFmt(), col_buf.GetNumberFmt());
+        key.color_formats[remapped_cb] = format;
+        if (!instance.IsFormatSupported(format, vk::FormatFeatureFlagBits2::eColorAttachment)) {
+            LOG_WARNING(Render_Vulkan,
+                        "color buffer format {} does not support COLOR_ATTACHMENT_BIT",
+                        vk::to_string(format));
+        }
+
         key.color_buffers[remapped_cb] = Shader::PsColorBuffer{
             .num_format = col_buf.GetNumberFmt(),
             .num_conversion = col_buf.GetNumberConversion(),
diff --git a/src/video_core/texture_cache/image.cpp b/src/video_core/texture_cache/image.cpp
index ab9111e6b..7b8ff4403 100644
--- a/src/video_core/texture_cache/image.cpp
+++ b/src/video_core/texture_cache/image.cpp
@@ -130,11 +130,24 @@ Image::Image(const Vulkan::Instance& instance_, Vulkan::Scheduler& scheduler_,
 
     constexpr auto tiling = vk::ImageTiling::eOptimal;
     const auto supported_format = instance->GetSupportedFormat(info.pixel_format, format_features);
-    const auto properties = instance->GetPhysicalDevice().getImageFormatProperties(
-        supported_format, info.type, tiling, usage_flags, flags);
-    const auto supported_samples = properties.result == vk::Result::eSuccess
-                                       ? properties.value.sampleCounts
-                                       : vk::SampleCountFlagBits::e1;
+    const vk::PhysicalDeviceImageFormatInfo2 format_info{
+        .format = supported_format,
+        .type = info.type,
+        .tiling = tiling,
+        .usage = usage_flags,
+        .flags = flags,
+    };
+    const auto image_format_properties =
+        instance->GetPhysicalDevice().getImageFormatProperties2(format_info);
+    if (image_format_properties.result == vk::Result::eErrorFormatNotSupported) {
+        LOG_ERROR(Render_Vulkan, "image format {} type {} is not supported (flags {}, usage {})",
+                  vk::to_string(supported_format), vk::to_string(info.type),
+                  vk::to_string(format_info.flags), vk::to_string(format_info.usage));
+    }
+    const auto supported_samples =
+        image_format_properties.result == vk::Result::eSuccess
+            ? image_format_properties.value.imageFormatProperties.sampleCounts
+            : vk::SampleCountFlagBits::e1;
 
     const vk::ImageCreateInfo image_ci = {
         .flags = flags,
diff --git a/src/video_core/texture_cache/image_view.cpp b/src/video_core/texture_cache/image_view.cpp
index 7befb5259..2e162ce83 100644
--- a/src/video_core/texture_cache/image_view.cpp
+++ b/src/video_core/texture_cache/image_view.cpp
@@ -29,6 +29,24 @@ vk::ImageViewType ConvertImageViewType(AmdGpu::ImageType type) {
     }
 }
 
+bool IsViewTypeCompatible(vk::ImageViewType view_type, vk::ImageType image_type) {
+    switch (view_type) {
+    case vk::ImageViewType::e1D:
+    case vk::ImageViewType::e1DArray:
+        return image_type == vk::ImageType::e1D;
+    case vk::ImageViewType::e2D:
+    case vk::ImageViewType::e2DArray:
+        return image_type == vk::ImageType::e2D || image_type == vk::ImageType::e3D;
+    case vk::ImageViewType::eCube:
+    case vk::ImageViewType::eCubeArray:
+        return image_type == vk::ImageType::e2D;
+    case vk::ImageViewType::e3D:
+        return image_type == vk::ImageType::e3D;
+    default:
+        UNREACHABLE();
+    }
+}
+
 ImageViewInfo::ImageViewInfo(const AmdGpu::Image& image, const Shader::ImageResource& desc) noexcept
     : is_storage{desc.is_written} {
     const auto dfmt = image.GetDataFmt();
@@ -106,6 +124,11 @@ ImageView::ImageView(const Vulkan::Instance& instance, const ImageViewInfo& info
             .layerCount = info.range.extent.layers,
         },
     };
+    if (!IsViewTypeCompatible(image_view_ci.viewType, image.info.type)) {
+        LOG_ERROR(Render_Vulkan, "image view type {} is incompatible with image type {}",
+                  vk::to_string(image_view_ci.viewType), vk::to_string(image.info.type));
+    }
+
     auto [view_result, view] = instance.GetDevice().createImageViewUnique(image_view_ci);
     ASSERT_MSG(view_result == vk::Result::eSuccess, "Failed to create image view: {}",
                vk::to_string(view_result));
diff --git a/src/video_core/texture_cache/texture_cache.cpp b/src/video_core/texture_cache/texture_cache.cpp
index a1ff5db8a..a50601af6 100644
--- a/src/video_core/texture_cache/texture_cache.cpp
+++ b/src/video_core/texture_cache/texture_cache.cpp
@@ -761,7 +761,7 @@ void TextureCache::UntrackImage(ImageId image_id) {
     image.track_addr = 0;
     image.track_addr_end = 0;
     if (size != 0) {
-        tracker.UpdatePageWatchers<-1>(addr, size);
+        tracker.UpdatePageWatchers<false>(addr, size);
     }
 }
 
@@ -780,7 +780,7 @@ void TextureCache::UntrackImageHead(ImageId image_id) {
         // Cehck its hash later.
         MarkAsMaybeDirty(image_id, image);
     }
-    tracker.UpdatePageWatchers<-1>(image_begin, size);
+    tracker.UpdatePageWatchers<false>(image_begin, size);
 }
 
 void TextureCache::UntrackImageTail(ImageId image_id) {
@@ -799,7 +799,7 @@ void TextureCache::UntrackImageTail(ImageId image_id) {
         // Cehck its hash later.
         MarkAsMaybeDirty(image_id, image);
     }
-    tracker.UpdatePageWatchers<-1>(addr, size);
+    tracker.UpdatePageWatchers<false>(addr, size);
 }
 
 void TextureCache::DeleteImage(ImageId image_id) {