[shader_recompiler, spir-v] Coalesce storage emission

user doesn't need it and just wastes resources
Signed-off-by: lizzie <lizzie@eden-emu.dev>

Reviewed-on: https://git.eden-emu.dev/eden-emu/eden/pulls/3293
Reviewed-by: Maufeat <sahyno1996@gmail.com>
Reviewed-by: DraVee <dravee@eden-emu.dev>
Co-authored-by: lizzie <lizzie@eden-emu.dev>
Co-committed-by: lizzie <lizzie@eden-emu.dev>

.patch/mbedtls/0002-arm64-aes-fix.patch

@@ -13,7 +13,7 @@ index 50f2a29..c60d9dc 100644
  #define MBEDTLS_HAVE_NEON_INTRINSICS
 -#elif defined(MBEDTLS_PLATFORM_IS_WINDOWS_ON_ARM64)
 -#include <arm64_neon.h>
-+#elif defined(__ANDROID__) || defined(__ARM_NEON)
++#elif (defined(__ANDROID__) && defined(__ARM_FP)) || defined(__ARM_NEON)
 +#include <arm_neon.h>
  #define MBEDTLS_HAVE_NEON_INTRINSICS
  #endif

README.md

@@ -15,8 +15,8 @@
   <br>
 </h1>
 
-<h4 align="center"><b>Eden</b> is an open-source Nintendo Switch emulator, forked from the Yuzu emulator — started by former Citron developer Camille LaVey and the Eden team.
-It is written in C++ with portability in mind, and we actively maintain builds for Windows, Linux and Android.
+<h4 align="center"><b>Eden</b> is a free and opensource (FOSS) Switch 1 emulator, derived from Yuzu and Sudachi - started by developer Camille LaVey.
+It's written in C++ with portability in mind, with builds for Windows, Linux, macOS, Android, FreeBSD and more.
 </h4>
 
 <p align="center">

dist/dev.eden_emu.eden.desktop

@@ -9,7 +9,7 @@ Version=1.0
 Type=Application
 Name=Eden
 GenericName=Switch Emulator
-Comment=Nintendo Switch video game console emulator
+Comment=Multiplatform FOSS Switch 1 emulator written in C++, derived from Yuzu and Sudachi
 Icon=dev.eden_emu.eden
 TryExec=eden
 Exec=eden %f

dist/dev.eden_emu.eden.metainfo.xml

@@ -16,7 +16,7 @@ SPDX-License-Identifier: CC0-1.0
   <name>eden</name>
   <summary>Nintendo Switch emulator</summary>
   <description>
-    <p>Nintendo Switch video game console emulator</p>
+    <p>Multiplatform FOSS Switch 1 emulator written in C++, derived from Yuzu and Sudachi</p>
   </description>
   <categories>
     <category>Game</category>

docs/build/Android.md

@@ -76,12 +76,12 @@ The APK and AAB will be output into "artifacts".
 
 Examples:
 
-* Build legacy release with update checker for ChromeOS:
-  * `.ci/android/build.sh -c -r -t legacy`
+* Build legacy release with update checker:
+  * `.ci/android/build.sh -r -t legacy`
 * Build standard release with debug info without update checker for phones:
   * `.ci/android/build.sh -b RelWithDebInfo`
-* Build optimized release with update checker:
-  * `.ci/android/build.sh -r -t optimized`
+* Build optimized release with update checker for ChromeOS:
+  * `.ci/android/build.sh -c -r -t optimized`
 
 ### Additional Resources
 

externals/cpmfile.json

@@ -71,9 +71,9 @@
         "hash": "9697e80a7d5d9bcb3ce51051a9a24962fb90ca79d215f1f03ae6b58da8ba13a63b5dda1b4dde3d26ac6445029696b8ef2883f4e5a777b342bba01283ed293856"
     },
     "libadrenotools": {
-        "repo": "bylaws/libadrenotools",
-        "sha": "8fae8ce254",
-        "hash": "db4a74ce15559c75e01d1868a90701519b655d77f2a343bbee283a42f8332dc9046960fb022dc969f205e457348a3f99cb8be6e1cd91264d2ae1235294b9f9b2",
+        "repo": "eden-emulator/libadrenotools",
+        "sha": "8ba23b42d7",
+        "hash": "f6526620cb752876edc5ed4c0925d57b873a8218ee09ad10859ee476e9333259784f61c1dcc55a2bcba597352d18aff22cd2e4c1925ec2ae94074e09d7da2265",
         "patches": [
             "0001-linkerns-cpm.patch"
         ]
@@ -180,8 +180,8 @@
     "discord-rpc": {
         "package": "DiscordRPC",
         "repo": "eden-emulator/discord-rpc",
-        "sha": "1cf7772bb6",
-        "hash": "9a6c35887dcacceb4ba1bf3141edb73b05b2abc719a8d81dad9cb9dd5b039ce203946787335d9d738af669c10cf2534638b645635a22096fc28dcae2475e0cbe",
+        "sha": "0d8b2d6a37",
+        "hash": "8213c43dcb0f7d479f5861091d111ed12fbdec1e62e6d729d65a4bc181d82f48a35d5fd3cd5c291f2393ac7c9681eabc6b76609755f55376284c8a8d67e148f3",
         "find_args": "MODULE"
     },
     "simpleini": {

src/CMakeLists.txt

@@ -7,8 +7,7 @@
 include_directories(.)
 
 # Dynarmic
-if ((ARCHITECTURE_x86_64 OR ARCHITECTURE_arm64))
-    set(DYNARMIC_IGNORE_ASSERTS ON)
+if (ARCHITECTURE_x86_64 OR ARCHITECTURE_arm64)
     add_subdirectory(dynarmic)
     add_library(dynarmic::dynarmic ALIAS dynarmic)
 endif()

src/android/app/src/main/java/org/yuzu/yuzu_emu/YuzuApplication.kt

@@ -61,6 +61,11 @@ class YuzuApplication : Application() {
         application = this
         documentsTree = DocumentsTree()
         DirectoryInitialization.start()
+
+        // Initialize Freedreno config BEFORE loading native library
+        // This ensures GPU driver environment variables are set before adrenotools initializes
+        GpuDriverHelper.initializeFreedrenoConfigEarly()
+
         NativeLibrary.playTimeManagerInit()
         GpuDriverHelper.initializeDriverParameters()
         NativeInput.reloadInputDevices()

src/android/app/src/main/java/org/yuzu/yuzu_emu/adapters/FreedrenoPresetAdapter.kt

@@ -0,0 +1,58 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
+package org.yuzu.yuzu_emu.adapters
+
+import android.content.Context
+import android.view.LayoutInflater
+import android.view.ViewGroup
+import androidx.recyclerview.widget.DiffUtil
+import androidx.recyclerview.widget.ListAdapter
+import androidx.recyclerview.widget.RecyclerView
+import org.yuzu.yuzu_emu.utils.FreedrenoPreset
+import org.yuzu.yuzu_emu.databinding.ListItemFreedrenoPresetBinding
+
+/**
+ * Adapter for displaying Freedreno preset configurations in a horizontal list.
+ */
+class FreedrenoPresetAdapter(
+    private val onPresetClicked: (FreedrenoPreset) -> Unit
+) : ListAdapter<FreedrenoPreset, FreedrenoPresetAdapter.PresetViewHolder>(DiffCallback) {
+
+    override fun onCreateViewHolder(parent: ViewGroup, viewType: Int): PresetViewHolder {
+        val binding = ListItemFreedrenoPresetBinding.inflate(
+            LayoutInflater.from(parent.context),
+            parent,
+            false
+        )
+        return PresetViewHolder(binding)
+    }
+
+    override fun onBindViewHolder(holder: PresetViewHolder, position: Int) {
+        holder.bind(getItem(position))
+    }
+
+    inner class PresetViewHolder(private val binding: ListItemFreedrenoPresetBinding) :
+        RecyclerView.ViewHolder(binding.root) {
+
+        fun bind(preset: FreedrenoPreset) {
+            binding.presetButton.apply {
+                text = preset.name
+                setOnClickListener {
+                    onPresetClicked(preset)
+                }
+                contentDescription = "${preset.name}: ${preset.description}"
+            }
+        }
+    }
+
+    companion object {
+        private val DiffCallback = object : DiffUtil.ItemCallback<FreedrenoPreset>() {
+            override fun areItemsTheSame(oldItem: FreedrenoPreset, newItem: FreedrenoPreset): Boolean =
+                oldItem.name == newItem.name
+
+            override fun areContentsTheSame(oldItem: FreedrenoPreset, newItem: FreedrenoPreset): Boolean =
+                oldItem == newItem
+        }
+    }
+}

src/android/app/src/main/java/org/yuzu/yuzu_emu/adapters/FreedrenoVariableAdapter.kt

@@ -0,0 +1,61 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
+package org.yuzu.yuzu_emu.adapters
+
+import android.content.Context
+import android.view.LayoutInflater
+import android.view.ViewGroup
+import androidx.recyclerview.widget.DiffUtil
+import androidx.recyclerview.widget.ListAdapter
+import androidx.recyclerview.widget.RecyclerView
+import org.yuzu.yuzu_emu.databinding.ListItemFreedrenoVariableBinding
+import org.yuzu.yuzu_emu.fragments.FreedrenoVariable
+import org.yuzu.yuzu_emu.utils.NativeFreedrenoConfig
+
+/**
+ * Adapter for displaying currently set Freedreno environment variables in a list.
+ */
+class FreedrenoVariableAdapter(
+    private val context: Context,
+    private val onItemClicked: (FreedrenoVariable, () -> Unit) -> Unit
+) : ListAdapter<FreedrenoVariable, FreedrenoVariableAdapter.VariableViewHolder>(DiffCallback) {
+
+    override fun onCreateViewHolder(parent: ViewGroup, viewType: Int): VariableViewHolder {
+        val binding = ListItemFreedrenoVariableBinding.inflate(
+            LayoutInflater.from(parent.context),
+            parent,
+            false
+        )
+        return VariableViewHolder(binding)
+    }
+
+    override fun onBindViewHolder(holder: VariableViewHolder, position: Int) {
+        holder.bind(getItem(position))
+    }
+
+    inner class VariableViewHolder(private val binding: ListItemFreedrenoVariableBinding) :
+        RecyclerView.ViewHolder(binding.root) {
+
+        fun bind(variable: FreedrenoVariable) {
+            binding.variableName.text = variable.name
+            binding.variableValue.text = variable.value
+
+            binding.buttonDelete.setOnClickListener {
+                onItemClicked(variable) {
+                    NativeFreedrenoConfig.clearFreedrenoEnv(variable.name)
+                }
+            }
+        }
+    }
+
+    companion object {
+        private val DiffCallback = object : DiffUtil.ItemCallback<FreedrenoVariable>() {
+            override fun areItemsTheSame(oldItem: FreedrenoVariable, newItem: FreedrenoVariable): Boolean =
+                oldItem.name == newItem.name
+
+            override fun areContentsTheSame(oldItem: FreedrenoVariable, newItem: FreedrenoVariable): Boolean =
+                oldItem == newItem
+        }
+    }
+}

src/android/app/src/main/java/org/yuzu/yuzu_emu/features/settings/model/Settings.kt

@@ -27,6 +27,7 @@ object Settings {
         SECTION_APP_SETTINGS(R.string.app_settings),
         SECTION_CUSTOM_PATHS(R.string.preferences_custom_paths),
         SECTION_DEBUG(R.string.preferences_debug),
+        SECTION_FREEDRENO(R.string.gpu_driver_settings),
         SECTION_APPLETS(R.string.applets_menu);
     }
 

src/android/app/src/main/java/org/yuzu/yuzu_emu/features/settings/ui/SettingsAdapter.kt

@@ -29,6 +29,7 @@ import org.yuzu.yuzu_emu.databinding.ListItemSettingsHeaderBinding
 import org.yuzu.yuzu_emu.features.input.NativeInput
 import org.yuzu.yuzu_emu.features.input.model.AnalogDirection
 import org.yuzu.yuzu_emu.features.settings.model.AbstractIntSetting
+import org.yuzu.yuzu_emu.features.settings.model.Settings
 import org.yuzu.yuzu_emu.features.settings.model.view.*
 import org.yuzu.yuzu_emu.features.settings.ui.viewholder.*
 import org.yuzu.yuzu_emu.utils.ParamPackage
@@ -212,8 +213,15 @@ class SettingsAdapter(
     }
 
     fun onSubmenuClick(item: SubmenuSetting) {
-        val action = SettingsNavigationDirections.actionGlobalSettingsFragment(item.menuKey, null)
-        fragment.view?.findNavController()?.navigate(action)
+        // Check if this is the Freedreno Settings submenu
+        if (item.menuKey == Settings.MenuTag.SECTION_FREEDRENO) {
+            fragment.view?.findNavController()?.navigate(
+                R.id.action_settingsFragment_to_freedrenoSettingsFragment
+            )
+        } else {
+            val action = SettingsNavigationDirections.actionGlobalSettingsFragment(item.menuKey, null)
+            fragment.view?.findNavController()?.navigate(action)
+        }
     }
 
     fun onLaunchableClick(item: LaunchableSetting) {

src/android/app/src/main/java/org/yuzu/yuzu_emu/features/settings/ui/SettingsFragmentPresenter.kt

@@ -27,6 +27,7 @@ import org.yuzu.yuzu_emu.features.settings.model.Settings.MenuTag
 import org.yuzu.yuzu_emu.features.settings.model.ShortSetting
 import org.yuzu.yuzu_emu.features.settings.model.StringSetting
 import org.yuzu.yuzu_emu.features.settings.model.view.*
+import org.yuzu.yuzu_emu.utils.GpuDriverHelper
 import org.yuzu.yuzu_emu.utils.InputHandler
 import org.yuzu.yuzu_emu.utils.NativeConfig
 import org.yuzu.yuzu_emu.utils.DirectoryInitialization
@@ -109,6 +110,7 @@ class SettingsFragmentPresenter(
             MenuTag.SECTION_INPUT_PLAYER_EIGHT -> addInputPlayer(sl, 7)
             MenuTag.SECTION_APP_SETTINGS -> addThemeSettings(sl)
             MenuTag.SECTION_DEBUG -> addDebugSettings(sl)
+            MenuTag.SECTION_FREEDRENO -> addFreedrenoSettings(sl)
             MenuTag.SECTION_APPLETS -> addAppletSettings(sl)
             MenuTag.SECTION_CUSTOM_PATHS -> addCustomPathsSettings(sl)
         }
@@ -181,6 +183,16 @@ class SettingsFragmentPresenter(
                     menuKey = MenuTag.SECTION_DEBUG
                 )
             )
+            if (GpuDriverHelper.isAdrenoGpu() && !NativeConfig.isPerGameConfigLoaded()) {
+                add(
+                    SubmenuSetting(
+                        titleId = R.string.gpu_driver_settings,
+                        descriptionId = R.string.freedreno_settings_title,
+                        iconId = R.drawable.ic_graphics,
+                        menuKey = MenuTag.SECTION_FREEDRENO
+                    )
+                )
+            }
             add(
                 SubmenuSetting(
                     titleId = R.string.applets_menu,
@@ -498,6 +510,11 @@ class SettingsFragmentPresenter(
         }
     }
 
+    private fun addFreedrenoSettings(sl: ArrayList<SettingsItem>) {
+        // No additional settings needed here - the SubmenuSetting handles navigation
+        // This method is kept for consistency with other menu sections
+    }
+
     private fun addAppletSettings(sl: ArrayList<SettingsItem>) {
         sl.apply {
             add(IntSetting.SWKBD_APPLET.key)

src/android/app/src/main/java/org/yuzu/yuzu_emu/fragments/DriverFetcherFragment.kt

@@ -45,7 +45,7 @@ class DriverFetcherFragment : Fragment() {
     private val client = OkHttpClient()
 
     private val gpuModel: String?
-        get() = GpuDriverHelper.getGpuModel()
+        get() = GpuDriverHelper.hookLibPath?.let { GpuDriverHelper.getGpuModel(hookLibPath = it) }
 
     private val adrenoModel: Int
         get() = parseAdrenoModel()

src/android/app/src/main/java/org/yuzu/yuzu_emu/fragments/EmulationFragment.kt

@@ -87,6 +87,7 @@ import org.yuzu.yuzu_emu.utils.GameIconUtils
 import org.yuzu.yuzu_emu.utils.GpuDriverHelper
 import org.yuzu.yuzu_emu.utils.Log
 import org.yuzu.yuzu_emu.utils.NativeConfig
+import org.yuzu.yuzu_emu.utils.NativeFreedrenoConfig
 import org.yuzu.yuzu_emu.utils.ViewUtils
 import org.yuzu.yuzu_emu.utils.ViewUtils.setVisible
 import org.yuzu.yuzu_emu.utils.collect
@@ -303,6 +304,18 @@ class EmulationFragment : Fragment(), SurfaceHolder.Callback {
                 throw fallbackException
             }
         }
+        try {
+            if (GpuDriverHelper.isAdrenoGpu()) {
+                val programIdHex = game!!.programIdHex
+                if (NativeFreedrenoConfig.loadPerGameConfigWithGlobalFallback(programIdHex)) {
+                    Log.info("[EmulationFragment] Loaded per-game Freedreno config for $programIdHex")
+                } else {
+                    Log.info("[EmulationFragment] Using global Freedreno config for $programIdHex")
+                }
+            }
+        } catch (e: Exception) {
+            Log.warning("[EmulationFragment] Failed to load Freedreno config: ${e.message}")
+        }
 
         emulationState = EmulationState(game!!.path) {
             return@EmulationState driverViewModel.isInteractionAllowed.value
@@ -616,7 +629,7 @@ class EmulationFragment : Fragment(), SurfaceHolder.Callback {
         }
         Log.info("[EmulationFragment] Starting view setup for game: ${game?.title}")
 
-        gpuModel = GpuDriverHelper.getGpuModel().toString()
+        gpuModel = GpuDriverHelper.hookLibPath?.let { GpuDriverHelper.getGpuModel(hookLibPath = it).toString() } ?: "Unknown"
         fwVersion = NativeLibrary.firmwareVersion()
 
         updateQuickOverlayMenuEntry(BooleanSetting.SHOW_INPUT_OVERLAY.getBoolean())

src/android/app/src/main/java/org/yuzu/yuzu_emu/fragments/FreedrenoSettingsFragment.kt

@@ -0,0 +1,204 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
+package org.yuzu.yuzu_emu.fragments
+
+import android.os.Bundle
+import android.view.LayoutInflater
+import android.view.View
+import android.view.ViewGroup
+import androidx.core.view.ViewCompat
+import androidx.core.view.WindowInsetsCompat
+import androidx.core.view.updatePadding
+import androidx.fragment.app.Fragment
+import androidx.navigation.fragment.navArgs
+import androidx.recyclerview.widget.LinearLayoutManager
+import com.google.android.material.snackbar.Snackbar
+import com.google.android.material.transition.MaterialSharedAxis
+import org.yuzu.yuzu_emu.R
+import org.yuzu.yuzu_emu.adapters.FreedrenoPresetAdapter
+import org.yuzu.yuzu_emu.adapters.FreedrenoVariableAdapter
+import org.yuzu.yuzu_emu.databinding.FragmentFreedrenoSettingsBinding
+import org.yuzu.yuzu_emu.model.Game
+import org.yuzu.yuzu_emu.utils.NativeFreedrenoConfig
+import org.yuzu.yuzu_emu.utils.FreedrenoPresets
+
+
+class FreedrenoSettingsFragment : Fragment() {
+    private var _binding: FragmentFreedrenoSettingsBinding? = null
+    private val binding get() = _binding!!
+    private val args by navArgs<FreedrenoSettingsFragmentArgs>()
+    private val game: Game? get() = args.game
+    private val isPerGameConfig: Boolean get() = game != null
+
+    private lateinit var presetAdapter: FreedrenoPresetAdapter
+    private lateinit var settingsAdapter: FreedrenoVariableAdapter
+
+    override fun onCreate(savedInstanceState: Bundle?) {
+        super.onCreate(savedInstanceState)
+        enterTransition = MaterialSharedAxis(MaterialSharedAxis.X, true)
+        returnTransition = MaterialSharedAxis(MaterialSharedAxis.X, false)
+        reenterTransition = MaterialSharedAxis(MaterialSharedAxis.X, false)
+        exitTransition = MaterialSharedAxis(MaterialSharedAxis.X, true)
+    }
+
+    override fun onCreateView(
+        inflater: LayoutInflater,
+        container: ViewGroup?,
+        savedInstanceState: Bundle?
+    ): View {
+        _binding = FragmentFreedrenoSettingsBinding.inflate(layoutInflater, container, false)
+        return binding.root
+    }
+
+    override fun onViewCreated(view: View, savedInstanceState: Bundle?) {
+        super.onViewCreated(view, savedInstanceState)
+
+        NativeFreedrenoConfig.setFreedrenoBasePath(requireContext().cacheDir.absolutePath)
+        NativeFreedrenoConfig.initializeFreedrenoConfig()
+
+        if (isPerGameConfig) {
+            NativeFreedrenoConfig.loadPerGameConfig(game!!.programIdHex)
+        } else {
+            NativeFreedrenoConfig.reloadFreedrenoConfig()
+        }
+
+        setupToolbar()
+        setupAdapters()
+        loadCurrentSettings()
+        setupButtonListeners()
+        setupWindowInsets()
+    }
+
+    private fun setupToolbar() {
+        binding.toolbarFreedreno.setNavigationOnClickListener {
+            requireActivity().onBackPressedDispatcher.onBackPressed()
+        }
+        if (isPerGameConfig) {
+            binding.toolbarFreedreno.title = getString(R.string.freedreno_per_game_title)
+            binding.toolbarFreedreno.subtitle = game!!.title
+        }
+    }
+
+    private fun setupAdapters() {
+        // Setup presets adapter (horizontal list)
+        presetAdapter = FreedrenoPresetAdapter { preset ->
+            applyPreset(preset)
+        }
+        binding.listFreedrenoPresets.apply {
+            adapter = presetAdapter
+            layoutManager = LinearLayoutManager(requireContext(), LinearLayoutManager.HORIZONTAL, false)
+        }
+        presetAdapter.submitList(FreedrenoPresets.ALL_PRESETS)
+
+        // Setup current settings adapter (vertical list)
+        settingsAdapter = FreedrenoVariableAdapter(requireContext()) { variable, onDelete ->
+            onDelete()
+            loadCurrentSettings() // Refresh list after deletion
+        }
+        binding.listFreedrenoSettings.apply {
+            adapter = settingsAdapter
+            layoutManager = LinearLayoutManager(requireContext())
+        }
+    }
+
+    private fun loadCurrentSettings() {
+        // Load all currently set environment variables
+        val variables = mutableListOf<FreedrenoVariable>()
+
+        // Common variables to check
+        val commonVars = listOf(
+            "TU_DEBUG", "FD_MESA_DEBUG", "IR3_SHADER_DEBUG",
+            "FD_RD_DUMP", "FD_RD_DUMP_FRAMES", "FD_RD_DUMP_TESTNAME",
+            "TU_BREADCRUMBS"
+        )
+
+        for (varName in commonVars) {
+            if (NativeFreedrenoConfig.isFreedrenoEnvSet(varName)) {
+                val value = NativeFreedrenoConfig.getFreedrenoEnv(varName)
+                variables.add(FreedrenoVariable(varName, value))
+            }
+        }
+
+        settingsAdapter.submitList(variables)
+    }
+
+    private fun setupButtonListeners() {
+        binding.buttonAddVariable.setOnClickListener {
+            val varName = binding.variableNameInput.text.toString().trim()
+            val varValue = binding.variableValueInput.text.toString().trim()
+
+            if (varName.isEmpty()) {
+                showSnackbar(getString(R.string.freedreno_error_empty_name))
+                return@setOnClickListener
+            }
+
+            if (NativeFreedrenoConfig.setFreedrenoEnv(varName, varValue)) {
+                showSnackbar(getString(R.string.freedreno_variable_added, varName))
+                binding.variableNameInput.text?.clear()
+                binding.variableValueInput.text?.clear()
+                loadCurrentSettings()
+            } else {
+                showSnackbar(getString(R.string.freedreno_error_setting_variable))
+            }
+        }
+
+        binding.buttonClearAll.setOnClickListener {
+            NativeFreedrenoConfig.clearAllFreedrenoEnv()
+            showSnackbar(getString(R.string.freedreno_cleared_all))
+            loadCurrentSettings()
+        }
+
+        binding.buttonSave.setOnClickListener {
+            if (isPerGameConfig) {
+                NativeFreedrenoConfig.savePerGameConfig(game!!.programIdHex)
+                showSnackbar(getString(R.string.freedreno_per_game_saved))
+            } else {
+                NativeFreedrenoConfig.saveFreedrenoConfig()
+                showSnackbar(getString(R.string.freedreno_saved))
+            }
+        }
+    }
+
+    private fun applyPreset(preset: org.yuzu.yuzu_emu.utils.FreedrenoPreset) {
+        // Clear all first for consistency
+        NativeFreedrenoConfig.clearAllFreedrenoEnv()
+
+        // Apply all variables in the preset
+        for ((varName, varValue) in preset.variables) {
+            NativeFreedrenoConfig.setFreedrenoEnv(varName, varValue)
+        }
+
+        showSnackbar(getString(R.string.freedreno_preset_applied, preset.name))
+        loadCurrentSettings()
+    }
+
+    private fun setupWindowInsets() {
+        ViewCompat.setOnApplyWindowInsetsListener(binding.root) { _, insets ->
+            val systemInsets = insets.getInsets(WindowInsetsCompat.Type.systemBars())
+            binding.root.updatePadding(
+                left = systemInsets.left,
+                right = systemInsets.right,
+                bottom = systemInsets.bottom
+            )
+            insets
+        }
+    }
+
+    private fun showSnackbar(message: String) {
+        Snackbar.make(binding.root, message, Snackbar.LENGTH_SHORT).show()
+    }
+
+    override fun onDestroyView() {
+        super.onDestroyView()
+        _binding = null
+    }
+}
+
+/**
+ * Data class representing a Freedreno environment variable.
+ */
+data class FreedrenoVariable(
+    val name: String,
+    val value: String
+)

src/android/app/src/main/java/org/yuzu/yuzu_emu/fragments/GamePropertiesFragment.kt

@@ -325,6 +325,20 @@ class GamePropertiesFragment : Fragment() {
                     )
                 )
             }
+            if (GpuDriverHelper.isAdrenoGpu()) {
+                add(
+                    SubmenuProperty(
+                        R.string.freedreno_per_game_title,
+                        R.string.freedreno_per_game_description,
+                        R.drawable.ic_graphics,
+                        action = {
+                            val action = GamePropertiesFragmentDirections
+                                .actionPerGamePropertiesFragmentToFreedrenoSettingsFragment(args.game)
+                            binding.root.findNavController().navigate(action)
+                        }
+                    )
+                )
+            }
 
             if (!args.game.isHomebrew) {
                 add(

src/android/app/src/main/java/org/yuzu/yuzu_emu/utils/GpuDriverHelper.kt

@@ -23,10 +23,16 @@ object GpuDriverHelper {
     private const val META_JSON_FILENAME = "meta.json"
     private var fileRedirectionPath: String? = null
     var driverInstallationPath: String? = null
-    private var hookLibPath: String? = null
+    internal var hookLibPath: String? = null
 
     val driverStoragePath get() = DirectoryInitialization.userDirectory!! + "/gpu_drivers/"
 
+    fun initializeFreedrenoConfigEarly() {
+        NativeFreedrenoConfig.setFreedrenoBasePath(YuzuApplication.appContext.cacheDir.absolutePath)
+        NativeFreedrenoConfig.initializeFreedrenoConfig()
+        NativeFreedrenoConfig.reloadFreedrenoConfig()
+    }
+
     fun initializeDriverParameters() {
         try {
             // Initialize the file redirection directory.
@@ -40,11 +46,9 @@ object GpuDriverHelper {
             throw RuntimeException(e)
         }
 
-        // Initialize directories.
         initializeDirectories()
-
-        // Initialize hook libraries directory.
         hookLibPath = YuzuApplication.appContext.applicationInfo.nativeLibraryDir + "/"
+        NativeFreedrenoConfig.reloadFreedrenoConfig()
 
         // Initialize GPU driver.
         NativeLibrary.initializeGpuDriver(
@@ -211,9 +215,17 @@ object GpuDriverHelper {
 
     external fun getGpuModel(
         surface: Surface = Surface(SurfaceTexture(true)),
-        hookLibPath: String = GpuDriverHelper.hookLibPath!!
+        hookLibPath: String
     ): String?
 
+    fun isAdrenoGpu(): Boolean {
+        return try {
+            supportsCustomDriverLoading()
+        } catch (e: Exception) {
+            false
+        }
+    }
+
     // Parse the custom driver metadata to retrieve the name.
     val installedCustomDriverData: GpuDriverMetadata
         get() = GpuDriverMetadata(File(driverInstallationPath + META_JSON_FILENAME))

src/android/app/src/main/java/org/yuzu/yuzu_emu/utils/NativeFreedrenoConfig.kt

@@ -0,0 +1,164 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
+package org.yuzu.yuzu_emu.utils
+
+/**
+ * Provides access to Freedreno/Turnip driver configuration through JNI bindings.
+ *
+ * This class allows Java/Kotlin code to configure Freedreno environment variables
+ * for the GPU driver (Turnip/Freedreno) that runs in the emulator on Android.
+ *
+ * Variables must be set BEFORE starting emulation for them to take effect.
+ *
+ * See https://docs.mesa3d.org/drivers/freedreno.html for documentation.
+ */
+object NativeFreedrenoConfig {
+
+    @Synchronized
+    external fun setFreedrenoBasePath(basePath: String)
+
+    @Synchronized
+    external fun initializeFreedrenoConfig()
+
+    @Synchronized
+    external fun saveFreedrenoConfig()
+
+    @Synchronized
+    external fun reloadFreedrenoConfig()
+
+    @Synchronized
+    external fun setFreedrenoEnv(varName: String, value: String): Boolean
+
+    @Synchronized
+    external fun getFreedrenoEnv(varName: String): String
+
+    @Synchronized
+    external fun isFreedrenoEnvSet(varName: String): Boolean
+
+    @Synchronized
+    external fun clearFreedrenoEnv(varName: String): Boolean
+
+    @Synchronized
+    external fun clearAllFreedrenoEnv()
+
+    @Synchronized
+    external fun getFreedrenoEnvSummary(): String
+
+    @Synchronized
+    external fun setCurrentProgramId(programId: String)
+
+    @Synchronized
+    external fun loadPerGameConfig(programId: String): Boolean
+
+    @Synchronized
+    external fun loadPerGameConfigWithGlobalFallback(programId: String): Boolean
+
+    @Synchronized
+    external fun savePerGameConfig(programId: String): Boolean
+
+    @Synchronized
+    external fun hasPerGameConfig(programId: String): Boolean
+
+    @Synchronized
+    external fun deletePerGameConfig(programId: String): Boolean
+}
+
+/**
+ * Data class representing a Freedreno preset configuration.
+ * Presets are commonly used debugging/profiling configurations.
+ */
+data class FreedrenoPreset(
+    val name: String,           // Display name (e.g., "Debug - CPU Memory")
+    val description: String,    // Description of what this preset does
+    val icon: String,           // Icon identifier
+    val variables: Map<String, String> // Map of env vars to set
+)
+
+/**
+ * Predefined Freedreno presets for quick configuration.
+ */
+object FreedrenoPresets {
+
+    val DEBUG_CPU_MEMORY = FreedrenoPreset(
+        name = "Debug - CPU Memory",
+        description = "Use CPU memory (slower but more stable)",
+        icon = "ic_debug_cpu",
+        variables = mapOf(
+            "TU_DEBUG" to "sysmem"
+        )
+    )
+
+    val DEBUG_UBWC_DISABLED = FreedrenoPreset(
+        name = "Debug - No UBWC",
+        description = "Disable UBWC compression for debugging",
+        icon = "ic_debug_ubwc",
+        variables = mapOf(
+            "TU_DEBUG" to "noubwc"
+        )
+    )
+
+    val DEBUG_NO_BINNING = FreedrenoPreset(
+        name = "Debug - No Binning",
+        description = "Disable binning optimization",
+        icon = "ic_debug_bin",
+        variables = mapOf(
+            "TU_DEBUG" to "nobin"
+        )
+    )
+
+    val CAPTURE_RENDERPASS = FreedrenoPreset(
+        name = "Capture - Renderpass",
+        description = "Capture command stream data for debugging",
+        icon = "ic_capture",
+        variables = mapOf(
+            "FD_RD_DUMP" to "enable"
+        )
+    )
+
+    val CAPTURE_FRAMES = FreedrenoPreset(
+        name = "Capture - First 100 Frames",
+        description = "Capture command stream for first 100 frames only",
+        icon = "ic_capture",
+        variables = mapOf(
+            "FD_RD_DUMP" to "enable",
+            "FD_RD_DUMP_FRAMES" to "0-100"
+        )
+    )
+
+    val SHADER_DEBUG = FreedrenoPreset(
+        name = "Shader Debug",
+        description = "Enable IR3 shader compiler debugging",
+        icon = "ic_shader",
+        variables = mapOf(
+            "IR3_SHADER_DEBUG" to "nouboopt,spillall"
+        )
+    )
+
+    val GPU_HANG_TRACE = FreedrenoPreset(
+        name = "GPU Hang Trace",
+        description = "Trace GPU progress for debugging hangs",
+        icon = "ic_hang_trace",
+        variables = mapOf(
+            "TU_BREADCRUMBS" to "1"
+        )
+    )
+
+    val PERFORMANCE_DEFAULT = FreedrenoPreset(
+        name = "Performance - Default",
+        description = "Clear all debug options for performance",
+        icon = "ic_performance",
+        variables = emptyMap()  // Clears all when applied
+    )
+
+    val ALL_PRESETS = listOf(
+        DEBUG_CPU_MEMORY,
+        DEBUG_UBWC_DISABLED,
+        DEBUG_NO_BINNING,
+        CAPTURE_RENDERPASS,
+        CAPTURE_FRAMES,
+        SHADER_DEBUG,
+        GPU_HANG_TRACE,
+        PERFORMANCE_DEFAULT
+    )
+}

src/android/app/src/main/jni/CMakeLists.txt

@@ -10,6 +10,7 @@ add_library(yuzu-android SHARED
     native.cpp
     native.h
     native_config.cpp
+    native_freedreno.cpp
     android_settings.cpp
     game_metadata.cpp
     native_log.cpp

src/android/app/src/main/jni/native.cpp

@@ -730,11 +730,28 @@ jboolean Java_org_yuzu_yuzu_1emu_NativeLibrary_doesUpdateMatchProgram(JNIEnv* en
     return false;
 }
 
-void JNICALL Java_org_yuzu_yuzu_1emu_NativeLibrary_initializeGpuDriver(JNIEnv* env, jclass clazz,
+void JNICALL Java_org_yuzu_yuzu_1emu_NativeLibrary_initializeGpuDriver(JNIEnv* env,
+                                                                       [[maybe_unused]] jclass clazz,
                                                                        jstring hook_lib_dir,
                                                                        jstring custom_driver_dir,
                                                                        jstring custom_driver_name,
                                                                        jstring file_redirect_dir) {
+    // Log active Freedreno environment variables
+    const char* tu_debug = getenv("TU_DEBUG");
+    const char* fd_debug = getenv("FD_MESA_DEBUG");
+    const char* ir3_debug = getenv("IR3_SHADER_DEBUG");
+    const char* fd_rd_dump = getenv("FD_RD_DUMP");
+    const char* tu_breadcrumbs = getenv("TU_BREADCRUMBS");
+
+    if (tu_debug || fd_debug || ir3_debug || fd_rd_dump || tu_breadcrumbs) {
+        LOG_INFO(Frontend, "[Freedreno] Initializing GPU driver with configuration:");
+        if (tu_debug) LOG_INFO(Frontend, "[Freedreno]   TU_DEBUG={}", tu_debug);
+        if (fd_debug) LOG_INFO(Frontend, "[Freedreno]   FD_MESA_DEBUG={}", fd_debug);
+        if (ir3_debug) LOG_INFO(Frontend, "[Freedreno]   IR3_SHADER_DEBUG={}", ir3_debug);
+        if (fd_rd_dump) LOG_INFO(Frontend, "[Freedreno]   FD_RD_DUMP={}", fd_rd_dump);
+        if (tu_breadcrumbs) LOG_INFO(Frontend, "[Freedreno]   TU_BREADCRUMBS={}", tu_breadcrumbs);
+    }
+
     EmulationSession::GetInstance().InitializeGpuDriver(
         Common::Android::GetJString(env, hook_lib_dir),
         Common::Android::GetJString(env, custom_driver_dir),

src/android/app/src/main/jni/native_freedreno.cpp

@@ -0,0 +1,477 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
+/**
+ * @file native_freedreno.cpp
+ * @brief JNI bindings for Freedreno/Turnip GPU driver configuration.
+ *
+ * Provides runtime configuration of Mesa Freedreno environment variables
+ * for the Turnip Vulkan driver on Adreno GPUs.
+ *
+ * @see https://docs.mesa3d.org/drivers/freedreno.html
+ */
+
+#include <algorithm>
+#include <cerrno>
+#include <cstdio>
+#include <cstring>
+#include <map>
+#include <memory>
+#include <string>
+#include <sys/stat.h>
+
+#include <jni.h>
+
+#include "common/android/android_common.h"
+#include "common/logging/log.h"
+#include "native.h"
+
+namespace {
+
+struct FreedrenoConfig {
+    std::map<std::string, std::string> env_vars;
+    std::string config_file_path;
+};
+
+std::unique_ptr<FreedrenoConfig> g_config;
+std::string g_base_path;
+std::string g_current_program_id;
+
+constexpr const char* kConfigFileName = ".freedreno.conf";
+constexpr const char* kPerGameConfigDir = "freedreno_games";
+
+void LogActiveVariables() {
+    if (!g_config || g_config->env_vars.empty()) {
+        return;
+    }
+    for (const auto& [key, value] : g_config->env_vars) {
+        LOG_INFO(Frontend, "[Freedreno] {}={}", key, value);
+    }
+}
+
+bool ApplyEnvironmentVariable(const std::string& key, const std::string& value) {
+    if (setenv(key.c_str(), value.c_str(), 1) != 0) {
+        LOG_ERROR(Frontend, "[Freedreno] Failed to set {}={} (errno: {})", key, value, errno);
+        return false;
+    }
+    return true;
+}
+
+void ClearAllEnvironmentVariables() {
+    if (!g_config) return;
+    for (const auto& [key, value] : g_config->env_vars) {
+        unsetenv(key.c_str());
+    }
+    g_config->env_vars.clear();
+}
+
+std::string GetConfigPath() {
+    return g_base_path + "/" + kConfigFileName;
+}
+
+std::string GetPerGameConfigPath(const std::string& program_id) {
+    return g_base_path + "/" + kPerGameConfigDir + "/" + program_id + ".conf";
+}
+
+void EnsurePerGameConfigDir() {
+    std::string dir_path = g_base_path + "/" + kPerGameConfigDir;
+    mkdir(dir_path.c_str(), 0755);
+}
+
+bool LoadConfigFromFile(const std::string& config_path) {
+    if (!g_config) return false;
+
+    FILE* file = fopen(config_path.c_str(), "r");
+    if (!file) {
+        return false;
+    }
+
+    char line[512];
+    int count = 0;
+    while (fgets(line, sizeof(line), file)) {
+        size_t len = strlen(line);
+        if (len > 0 && line[len - 1] == '\n') {
+            line[len - 1] = '\0';
+            len--;
+        }
+
+        if (len == 0 || line[0] == '#') {
+            continue;
+        }
+
+        const char* eq = strchr(line, '=');
+        if (!eq) {
+            continue;
+        }
+
+        std::string key(line, eq - line);
+        std::string value(eq + 1);
+
+        g_config->env_vars[key] = value;
+        ApplyEnvironmentVariable(key, value);
+        count++;
+    }
+
+    fclose(file);
+    return count > 0;
+}
+
+bool SaveConfigToFile(const std::string& config_path) {
+    if (!g_config) return false;
+
+    FILE* file = fopen(config_path.c_str(), "w");
+    if (!file) {
+        LOG_ERROR(Frontend, "[Freedreno] Failed to open {} for writing", config_path);
+        return false;
+    }
+
+    fprintf(file, "# Freedreno/Turnip Configuration\n");
+    fprintf(file, "# Auto-generated by Eden Emulator\n\n");
+
+    for (const auto& [key, value] : g_config->env_vars) {
+        fprintf(file, "%s=%s\n", key.c_str(), value.c_str());
+    }
+
+    fclose(file);
+    return true;
+}
+
+} // anonymous namespace
+
+extern "C" {
+
+JNIEXPORT void JNICALL
+Java_org_yuzu_yuzu_1emu_utils_NativeFreedrenoConfig_setFreedrenoBasePath(
+    JNIEnv* env, [[maybe_unused]] jobject obj, jstring jbasePath) {
+    g_base_path = Common::Android::GetJString(env, jbasePath);
+}
+
+JNIEXPORT void JNICALL
+Java_org_yuzu_yuzu_1emu_utils_NativeFreedrenoConfig_initializeFreedrenoConfig(
+    [[maybe_unused]] JNIEnv* env, [[maybe_unused]] jobject obj) {
+    if (!g_config) {
+        g_config = std::make_unique<FreedrenoConfig>();
+        LOG_INFO(Frontend, "[Freedreno] Configuration system initialized");
+    }
+}
+
+JNIEXPORT void JNICALL
+Java_org_yuzu_yuzu_1emu_utils_NativeFreedrenoConfig_saveFreedrenoConfig(
+    [[maybe_unused]] JNIEnv* env, [[maybe_unused]] jobject obj) {
+    if (!g_config) {
+        LOG_WARNING(Frontend, "[Freedreno] Cannot save: not initialized");
+        return;
+    }
+
+    const std::string config_path = GetConfigPath();
+    FILE* file = fopen(config_path.c_str(), "w");
+    if (!file) {
+        LOG_ERROR(Frontend, "[Freedreno] Failed to open {} for writing", config_path);
+        return;
+    }
+
+    fprintf(file, "# Freedreno/Turnip Configuration\n");
+    fprintf(file, "# Auto-generated by Eden Emulator\n\n");
+
+    for (const auto& [key, value] : g_config->env_vars) {
+        fprintf(file, "%s=%s\n", key.c_str(), value.c_str());
+    }
+
+    fclose(file);
+    g_config->config_file_path = config_path;
+
+    LOG_INFO(Frontend, "[Freedreno] Saved {} variables to {}",
+             g_config->env_vars.size(), config_path);
+}
+
+JNIEXPORT void JNICALL
+Java_org_yuzu_yuzu_1emu_utils_NativeFreedrenoConfig_reloadFreedrenoConfig(
+    [[maybe_unused]] JNIEnv* env, [[maybe_unused]] jobject obj) {
+    if (!g_config) {
+        LOG_WARNING(Frontend, "[Freedreno] Cannot reload: not initialized");
+        return;
+    }
+
+    const std::string config_path = GetConfigPath();
+    g_config->env_vars.clear();
+
+    FILE* file = fopen(config_path.c_str(), "r");
+    if (!file) {
+        LOG_DEBUG(Frontend, "[Freedreno] No config file found at {}", config_path);
+        return;
+    }
+
+    char line[512];
+    while (fgets(line, sizeof(line), file)) {
+        // Remove trailing newline
+        size_t len = strlen(line);
+        if (len > 0 && line[len - 1] == '\n') {
+            line[len - 1] = '\0';
+            len--;
+        }
+
+        // Skip empty lines and comments
+        if (len == 0 || line[0] == '#') {
+            continue;
+        }
+
+        // Parse key=value
+        const char* eq = strchr(line, '=');
+        if (!eq) {
+            continue;
+        }
+
+        std::string key(line, eq - line);
+        std::string value(eq + 1);
+
+        g_config->env_vars[key] = value;
+        ApplyEnvironmentVariable(key, value);
+    }
+
+    fclose(file);
+    g_config->config_file_path = config_path;
+
+    if (!g_config->env_vars.empty()) {
+        LOG_INFO(Frontend, "[Freedreno] Loaded {} variables:", g_config->env_vars.size());
+        LogActiveVariables();
+    }
+}
+
+JNIEXPORT jboolean JNICALL
+Java_org_yuzu_yuzu_1emu_utils_NativeFreedrenoConfig_setFreedrenoEnv(
+    JNIEnv* env, [[maybe_unused]] jobject obj, jstring jvarName, jstring jvalue) {
+    if (!g_config) {
+        return JNI_FALSE;
+    }
+
+    auto var_name = Common::Android::GetJString(env, jvarName);
+    auto value = Common::Android::GetJString(env, jvalue);
+
+    if (var_name.empty()) {
+        return JNI_FALSE;
+    }
+
+    g_config->env_vars[var_name] = value;
+
+    if (!ApplyEnvironmentVariable(var_name, value)) {
+        return JNI_FALSE;
+    }
+
+    LOG_INFO(Frontend, "[Freedreno] Set {}={}", var_name, value);
+    return JNI_TRUE;
+}
+
+JNIEXPORT jstring JNICALL
+Java_org_yuzu_yuzu_1emu_utils_NativeFreedrenoConfig_getFreedrenoEnv(
+    JNIEnv* env, [[maybe_unused]] jobject obj, jstring jvarName) {
+    if (!g_config) {
+        return env->NewStringUTF("");
+    }
+
+    auto var_name = Common::Android::GetJString(env, jvarName);
+    auto it = g_config->env_vars.find(var_name);
+
+    if (it != g_config->env_vars.end()) {
+        return env->NewStringUTF(it->second.c_str());
+    }
+
+    return env->NewStringUTF("");
+}
+
+JNIEXPORT jboolean JNICALL
+Java_org_yuzu_yuzu_1emu_utils_NativeFreedrenoConfig_isFreedrenoEnvSet(
+    JNIEnv* env, [[maybe_unused]] jobject obj, jstring jvarName) {
+    if (!g_config) {
+        return JNI_FALSE;
+    }
+
+    auto var_name = Common::Android::GetJString(env, jvarName);
+    auto it = g_config->env_vars.find(var_name);
+
+    return (it != g_config->env_vars.end() && !it->second.empty()) ? JNI_TRUE : JNI_FALSE;
+}
+
+JNIEXPORT jboolean JNICALL
+Java_org_yuzu_yuzu_1emu_utils_NativeFreedrenoConfig_clearFreedrenoEnv(
+    JNIEnv* env, [[maybe_unused]] jobject obj, jstring jvarName) {
+    if (!g_config) {
+        return JNI_FALSE;
+    }
+
+    auto var_name = Common::Android::GetJString(env, jvarName);
+    auto it = g_config->env_vars.find(var_name);
+
+    if (it != g_config->env_vars.end()) {
+        g_config->env_vars.erase(it);
+        unsetenv(var_name.c_str());
+        LOG_INFO(Frontend, "[Freedreno] Cleared {}", var_name);
+        return JNI_TRUE;
+    }
+
+    return JNI_FALSE;
+}
+
+JNIEXPORT void JNICALL
+Java_org_yuzu_yuzu_1emu_utils_NativeFreedrenoConfig_clearAllFreedrenoEnv(
+    [[maybe_unused]] JNIEnv* env, [[maybe_unused]] jobject obj) {
+    if (!g_config) {
+        return;
+    }
+
+    for (const auto& [key, value] : g_config->env_vars) {
+        unsetenv(key.c_str());
+    }
+
+    size_t count = g_config->env_vars.size();
+    g_config->env_vars.clear();
+
+    if (count > 0) {
+        LOG_INFO(Frontend, "[Freedreno] Cleared all {} variables", count);
+    }
+}
+
+JNIEXPORT jstring JNICALL
+Java_org_yuzu_yuzu_1emu_utils_NativeFreedrenoConfig_getFreedrenoEnvSummary(
+    JNIEnv* env, [[maybe_unused]] jobject obj) {
+    if (!g_config || g_config->env_vars.empty()) {
+        return env->NewStringUTF("");
+    }
+
+    std::string summary;
+    for (const auto& [key, value] : g_config->env_vars) {
+        if (!summary.empty()) {
+            summary += ",";
+        }
+        summary += key + "=" + value;
+    }
+
+    return env->NewStringUTF(summary.c_str());
+}
+
+JNIEXPORT void JNICALL
+Java_org_yuzu_yuzu_1emu_utils_NativeFreedrenoConfig_setCurrentProgramId(
+    JNIEnv* env, [[maybe_unused]] jobject obj, jstring jprogramId) {
+    g_current_program_id = Common::Android::GetJString(env, jprogramId);
+}
+
+JNIEXPORT jboolean JNICALL
+Java_org_yuzu_yuzu_1emu_utils_NativeFreedrenoConfig_loadPerGameConfig(
+    JNIEnv* env, [[maybe_unused]] jobject obj, jstring jprogramId) {
+    if (!g_config) {
+        return JNI_FALSE;
+    }
+
+    auto program_id = Common::Android::GetJString(env, jprogramId);
+    if (program_id.empty()) {
+        return JNI_FALSE;
+    }
+
+    // Clear current environment variables first
+    ClearAllEnvironmentVariables();
+    g_current_program_id = program_id;
+
+    // Try to load per-game config - do NOT fall back to global
+    // Per-game config should start empty if no config exists yet
+    std::string per_game_path = GetPerGameConfigPath(program_id);
+    if (LoadConfigFromFile(per_game_path)) {
+        LOG_INFO(Frontend, "[Freedreno] Loaded per-game config for {}", program_id);
+        LogActiveVariables();
+        return JNI_TRUE;
+    }
+
+    // No per-game config exists - start with empty config
+    LOG_INFO(Frontend, "[Freedreno] No per-game config for {}, starting empty", program_id);
+    return JNI_FALSE;
+}
+
+JNIEXPORT jboolean JNICALL
+Java_org_yuzu_yuzu_1emu_utils_NativeFreedrenoConfig_loadPerGameConfigWithGlobalFallback(
+    JNIEnv* env, [[maybe_unused]] jobject obj, jstring jprogramId) {
+    if (!g_config) {
+        return JNI_FALSE;
+    }
+
+    auto program_id = Common::Android::GetJString(env, jprogramId);
+    if (program_id.empty()) {
+        return JNI_FALSE;
+    }
+
+    // Clear current environment variables first
+    ClearAllEnvironmentVariables();
+    g_current_program_id = program_id;
+
+    // Try to load per-game config first
+    std::string per_game_path = GetPerGameConfigPath(program_id);
+    if (LoadConfigFromFile(per_game_path)) {
+        LOG_INFO(Frontend, "[Freedreno] Loaded per-game config for {}", program_id);
+        LogActiveVariables();
+        return JNI_TRUE;
+    }
+
+    // Fall back to global config for emulation
+    std::string global_path = GetConfigPath();
+    if (LoadConfigFromFile(global_path)) {
+        LOG_INFO(Frontend, "[Freedreno] No per-game config for {}, using global for emulation", program_id);
+        LogActiveVariables();
+    }
+
+    return JNI_FALSE;
+}
+
+JNIEXPORT jboolean JNICALL
+Java_org_yuzu_yuzu_1emu_utils_NativeFreedrenoConfig_savePerGameConfig(
+    JNIEnv* env, [[maybe_unused]] jobject obj, jstring jprogramId) {
+    if (!g_config) {
+        return JNI_FALSE;
+    }
+
+    auto program_id = Common::Android::GetJString(env, jprogramId);
+    if (program_id.empty()) {
+        return JNI_FALSE;
+    }
+
+    EnsurePerGameConfigDir();
+    std::string config_path = GetPerGameConfigPath(program_id);
+
+    if (SaveConfigToFile(config_path)) {
+        LOG_INFO(Frontend, "[Freedreno] Saved per-game config for {}", program_id);
+        return JNI_TRUE;
+    }
+
+    return JNI_FALSE;
+}
+
+JNIEXPORT jboolean JNICALL
+Java_org_yuzu_yuzu_1emu_utils_NativeFreedrenoConfig_hasPerGameConfig(
+    JNIEnv* env, [[maybe_unused]] jobject obj, jstring jprogramId) {
+    auto program_id = Common::Android::GetJString(env, jprogramId);
+    if (program_id.empty()) {
+        return JNI_FALSE;
+    }
+
+    std::string config_path = GetPerGameConfigPath(program_id);
+    FILE* file = fopen(config_path.c_str(), "r");
+    if (file) {
+        fclose(file);
+        return JNI_TRUE;
+    }
+    return JNI_FALSE;
+}
+
+JNIEXPORT jboolean JNICALL
+Java_org_yuzu_yuzu_1emu_utils_NativeFreedrenoConfig_deletePerGameConfig(
+    JNIEnv* env, [[maybe_unused]] jobject obj, jstring jprogramId) {
+    auto program_id = Common::Android::GetJString(env, jprogramId);
+    if (program_id.empty()) {
+        return JNI_FALSE;
+    }
+
+    std::string config_path = GetPerGameConfigPath(program_id);
+    if (remove(config_path.c_str()) == 0) {
+        LOG_INFO(Frontend, "[Freedreno] Deleted per-game config for {}", program_id);
+        return JNI_TRUE;
+    }
+    return JNI_FALSE;
+}
+
+} // extern "C"

src/android/app/src/main/res/layout/fragment_freedreno_settings.xml

@@ -0,0 +1,196 @@
+<?xml version="1.0" encoding="utf-8"?>
+<androidx.coordinatorlayout.widget.CoordinatorLayout xmlns:android="http://schemas.android.com/apk/res/android"
+    xmlns:app="http://schemas.android.com/apk/res-auto"
+    android:id="@+id/coordinator_main"
+    android:layout_width="match_parent"
+    android:layout_height="match_parent"
+    android:background="?attr/colorSurface">
+
+    <com.google.android.material.appbar.AppBarLayout
+        android:id="@+id/appbar_freedreno"
+        android:layout_width="match_parent"
+        android:layout_height="wrap_content"
+        android:fitsSystemWindows="true"
+        android:touchscreenBlocksFocus="false"
+        app:elevation="0dp">
+
+        <com.google.android.material.appbar.CollapsingToolbarLayout
+            android:id="@+id/toolbar_freedreno_layout"
+            style="?attr/collapsingToolbarLayoutMediumStyle"
+            android:layout_width="match_parent"
+            android:layout_height="?attr/collapsingToolbarLayoutMediumSize"
+            app:layout_scrollFlags="scroll|exitUntilCollapsed|snap"
+            app:contentScrim="?attr/colorSurface"
+            app:scrimVisibleHeightTrigger="100dp">
+
+            <com.google.android.material.appbar.MaterialToolbar
+                android:id="@+id/toolbar_freedreno"
+                android:layout_width="match_parent"
+                android:layout_height="?attr/actionBarSize"
+                android:touchscreenBlocksFocus="false"
+                app:layout_collapseMode="pin"
+                app:navigationIcon="@drawable/ic_back"
+                app:title="@string/gpu_driver_settings" />
+
+        </com.google.android.material.appbar.CollapsingToolbarLayout>
+
+    </com.google.android.material.appbar.AppBarLayout>
+
+    <androidx.core.widget.NestedScrollView
+        android:layout_width="match_parent"
+        android:layout_height="match_parent"
+        app:layout_behavior="@string/appbar_scrolling_view_behavior">
+
+        <LinearLayout
+            android:layout_width="match_parent"
+            android:layout_height="wrap_content"
+            android:orientation="vertical"
+            android:paddingStart="16dp"
+            android:paddingEnd="16dp"
+            android:paddingTop="8dp"
+            android:paddingBottom="16dp">
+
+            <!-- Presets Section -->
+            <com.google.android.material.textview.MaterialTextView
+                android:layout_width="match_parent"
+                android:layout_height="wrap_content"
+                android:layout_marginTop="16dp"
+                android:layout_marginBottom="8dp"
+                android:text="@string/freedreno_presets"
+                android:textAppearance="?attr/textAppearanceTitleMedium" />
+
+            <androidx.recyclerview.widget.RecyclerView
+                android:id="@+id/list_freedreno_presets"
+                android:layout_width="match_parent"
+                android:layout_height="wrap_content"
+                android:orientation="horizontal"
+                android:scrollbars="horizontal"
+                app:layoutManager="androidx.recyclerview.widget.LinearLayoutManager" />
+
+            <!-- Current Settings Section -->
+            <com.google.android.material.textview.MaterialTextView
+                android:layout_width="match_parent"
+                android:layout_height="wrap_content"
+                android:layout_marginTop="24dp"
+                android:layout_marginBottom="8dp"
+                android:text="@string/freedreno_current_settings"
+                android:textAppearance="?attr/textAppearanceTitleMedium" />
+
+            <!-- Settings List -->
+            <androidx.recyclerview.widget.RecyclerView
+                android:id="@+id/list_freedreno_settings"
+                android:layout_width="match_parent"
+                android:layout_height="wrap_content"
+                android:orientation="vertical"
+                app:layoutManager="androidx.recyclerview.widget.LinearLayoutManager" />
+
+            <!-- Debug Section -->
+            <com.google.android.material.textview.MaterialTextView
+                android:layout_width="match_parent"
+                android:layout_height="wrap_content"
+                android:layout_marginTop="24dp"
+                android:layout_marginBottom="8dp"
+                android:text="@string/freedreno_debug"
+                android:textAppearance="?attr/textAppearanceTitleMedium" />
+
+            <!-- Manual Variable Input -->
+            <com.google.android.material.textfield.TextInputLayout
+                android:id="@+id/variable_name_input_layout"
+                android:layout_width="match_parent"
+                android:layout_height="wrap_content"
+                android:layout_marginBottom="8dp"
+                android:hint="@string/freedreno_var_name">
+
+                <com.google.android.material.textfield.TextInputEditText
+                    android:id="@+id/variable_name_input"
+                    android:layout_width="match_parent"
+                    android:layout_height="wrap_content"
+                    android:inputType="text" />
+
+            </com.google.android.material.textfield.TextInputLayout>
+
+            <com.google.android.material.textfield.TextInputLayout
+                android:id="@+id/variable_value_input_layout"
+                android:layout_width="match_parent"
+                android:layout_height="wrap_content"
+                android:layout_marginBottom="8dp"
+                android:hint="@string/freedreno_var_value">
+
+                <com.google.android.material.textfield.TextInputEditText
+                    android:id="@+id/variable_value_input"
+                    android:layout_width="match_parent"
+                    android:layout_height="wrap_content"
+                    android:inputType="text" />
+
+            </com.google.android.material.textfield.TextInputLayout>
+
+            <Button
+                android:id="@+id/button_add_variable"
+                android:layout_width="match_parent"
+                android:layout_height="wrap_content"
+                android:layout_marginBottom="8dp"
+                android:text="@string/freedreno_add_variable" />
+
+            <!-- Action Buttons -->
+            <LinearLayout
+                android:layout_width="match_parent"
+                android:layout_height="wrap_content"
+                android:orientation="horizontal"
+                android:layout_marginTop="16dp"
+                android:spacing="8dp">
+
+                <Button
+                    android:id="@+id/button_clear_all"
+                    android:layout_width="0dp"
+                    android:layout_height="wrap_content"
+                    android:layout_weight="1"
+                    android:text="@string/freedreno_clear_all"
+                    style="?attr/materialButtonOutlinedStyle" />
+
+                <Button
+                    android:id="@+id/button_save"
+                    android:layout_width="0dp"
+                    android:layout_height="wrap_content"
+                    android:layout_weight="1"
+                    android:text="@string/save" />
+
+            </LinearLayout>
+
+            <!-- Info Section -->
+            <com.google.android.material.card.MaterialCardView
+                android:layout_width="match_parent"
+                android:layout_height="wrap_content"
+                android:layout_marginTop="24dp"
+                app:cardElevation="0dp"
+                app:strokeWidth="1dp"
+                app:strokeColor="?attr/colorOutline">
+
+                <LinearLayout
+                    android:layout_width="match_parent"
+                    android:layout_height="wrap_content"
+                    android:orientation="vertical"
+                    android:padding="12dp">
+
+                    <com.google.android.material.textview.MaterialTextView
+                        android:layout_width="match_parent"
+                        android:layout_height="wrap_content"
+                        android:text="@string/freedreno_info_title"
+                        android:textAppearance="?attr/textAppearanceTitleSmall" />
+
+                    <com.google.android.material.textview.MaterialTextView
+                        android:layout_width="match_parent"
+                        android:layout_height="wrap_content"
+                        android:layout_marginTop="8dp"
+                        android:text="@string/freedreno_info_description"
+                        android:textAppearance="?attr/textAppearanceBodySmall"
+                        android:textColor="?attr/colorOnSurfaceVariant" />
+
+                </LinearLayout>
+
+            </com.google.android.material.card.MaterialCardView>
+
+        </LinearLayout>
+
+    </androidx.core.widget.NestedScrollView>
+
+</androidx.coordinatorlayout.widget.CoordinatorLayout>

src/android/app/src/main/res/layout/list_item_freedreno_preset.xml

@@ -0,0 +1,18 @@
+<?xml version="1.0" encoding="utf-8"?>
+<LinearLayout xmlns:android="http://schemas.android.com/apk/res/android"
+    xmlns:app="http://schemas.android.com/apk/res-auto"
+    android:layout_width="wrap_content"
+    android:layout_height="wrap_content"
+    android:orientation="vertical"
+    android:paddingStart="4dp"
+    android:paddingEnd="4dp">
+
+    <com.google.android.material.button.MaterialButton
+        android:id="@+id/preset_button"
+        android:layout_width="wrap_content"
+        android:layout_height="wrap_content"
+        android:minWidth="100dp"
+        android:text="Preset"
+        style="?attr/materialButtonOutlinedStyle" />
+
+</LinearLayout>

src/android/app/src/main/res/layout/list_item_freedreno_variable.xml

@@ -0,0 +1,62 @@
+<?xml version="1.0" encoding="utf-8"?>
+<com.google.android.material.card.MaterialCardView xmlns:android="http://schemas.android.com/apk/res/android"
+    xmlns:app="http://schemas.android.com/apk/res-auto"
+    android:layout_width="match_parent"
+    android:layout_height="wrap_content"
+    android:layout_marginStart="0dp"
+    android:layout_marginEnd="0dp"
+    android:layout_marginTop="4dp"
+    android:layout_marginBottom="4dp"
+    app:cardElevation="0dp"
+    app:strokeWidth="1dp"
+    app:strokeColor="?attr/colorOutline">
+
+    <LinearLayout
+        android:layout_width="match_parent"
+        android:layout_height="wrap_content"
+        android:orientation="horizontal"
+        android:padding="12dp"
+        android:gravity="center_vertical">
+
+        <LinearLayout
+            android:layout_width="0dp"
+            android:layout_height="wrap_content"
+            android:layout_weight="1"
+            android:orientation="vertical">
+
+            <com.google.android.material.textview.MaterialTextView
+                android:id="@+id/variable_name"
+                android:layout_width="match_parent"
+                android:layout_height="wrap_content"
+                android:text="VARIABLE_NAME"
+                android:textAppearance="?attr/textAppearanceTitleSmall"
+                android:textColor="?attr/colorOnSurface" />
+
+            <com.google.android.material.textview.MaterialTextView
+                android:id="@+id/variable_value"
+                android:layout_width="match_parent"
+                android:layout_height="wrap_content"
+                android:layout_marginTop="4dp"
+                android:text="variable_value"
+                android:textAppearance="?attr/textAppearanceBodySmall"
+                android:textColor="?attr/colorOnSurfaceVariant"
+                android:maxLines="1"
+                android:ellipsize="end" />
+
+        </LinearLayout>
+
+        <com.google.android.material.button.MaterialButton
+            android:id="@+id/button_delete"
+            android:layout_width="wrap_content"
+            android:layout_height="wrap_content"
+            android:layout_marginStart="8dp"
+            android:text="@string/delete"
+            android:minWidth="0dp"
+            android:minHeight="36dp"
+            android:paddingStart="8dp"
+            android:paddingEnd="8dp"
+            style="?attr/materialButtonOutlinedStyle" />
+
+    </LinearLayout>
+
+</com.google.android.material.card.MaterialCardView>

src/android/app/src/main/res/navigation/home_navigation.xml

@@ -148,6 +148,9 @@
         <action
             android:id="@+id/action_perGamePropertiesFragment_to_driverManagerFragment"
             app:destination="@id/driverManagerFragment" />
+        <action
+            android:id="@+id/action_perGamePropertiesFragment_to_freedrenoSettingsFragment"
+            app:destination="@id/freedrenoSettingsFragment" />
     </fragment>
     <action
         android:id="@+id/action_global_perGamePropertiesFragment"
@@ -173,5 +176,15 @@
         android:name="org.yuzu.yuzu_emu.fragments.DriverFetcherFragment"
         android:label="fragment_driver_fetcher"
         tools:layout="@layout/fragment_driver_fetcher" />
+    <fragment
+        android:id="@+id/freedrenoSettingsFragment"
+        android:name="org.yuzu.yuzu_emu.fragments.FreedrenoSettingsFragment"
+        android:label="@string/freedreno_settings_title">
+        <argument
+            android:name="game"
+            app:argType="org.yuzu.yuzu_emu.model.Game"
+            app:nullable="true"
+            android:defaultValue="@null" />
+    </fragment>
 
 </navigation>

src/android/app/src/main/res/navigation/settings_navigation.xml

@@ -29,4 +29,19 @@
         android:name="org.yuzu.yuzu_emu.features.settings.ui.SettingsSearchFragment"
         android:label="SettingsSearchFragment" />
 
+    <fragment
+        android:id="@+id/freedrenoSettingsFragment"
+        android:name="org.yuzu.yuzu_emu.fragments.FreedrenoSettingsFragment"
+        android:label="@string/freedreno_settings_title">
+        <argument
+            android:name="game"
+            app:argType="org.yuzu.yuzu_emu.model.Game"
+            app:nullable="true"
+            android:defaultValue="@null" />
+    </fragment>
+
+    <action
+        android:id="@+id/action_settingsFragment_to_freedrenoSettingsFragment"
+        app:destination="@id/freedrenoSettingsFragment" />
+
 </navigation>

src/android/app/src/main/res/values/strings.xml

@@ -1051,6 +1051,28 @@
     <string name="cpu_accuracy_paranoid">Paranoid</string>
     <string name="cpu_accuracy_debugging">Debugging</string>
 
+    <!-- Freedreno Settings -->
+    <string name="freedreno_settings_title">Freedreno Settings</string>
+    <string name="gpu_driver_settings">GPU Driver Settings</string>
+    <string name="freedreno_presets">Quick Presets</string>
+    <string name="freedreno_current_settings">Current Settings</string>
+    <string name="freedreno_debug">Advanced Settings</string>
+    <string name="freedreno_var_name">Variable Name (e.g., TU_DEBUG)</string>
+    <string name="freedreno_var_value">Variable Value</string>
+    <string name="freedreno_add_variable">Add Variable</string>
+    <string name="freedreno_clear_all">Clear All</string>
+    <string name="freedreno_saved">Freedreno configuration saved</string>
+    <string name="freedreno_cleared_all">All Freedreno variables cleared</string>
+    <string name="freedreno_variable_added">Variable %1$s added</string>
+    <string name="freedreno_preset_applied">Preset \'%1$s\' applied</string>
+    <string name="freedreno_error_empty_name">Variable name cannot be empty</string>
+    <string name="freedreno_error_setting_variable">Failed to set variable</string>
+    <string name="freedreno_info_title">About Freedreno Configuration</string>
+    <string name="freedreno_info_description">Configure Freedreno/Turnip GPU driver options for debugging, profiling, and performance optimization. Changes are saved automatically. See https://docs.mesa3d.org/drivers/freedreno.html for detailed documentation.</string>
+    <string name="freedreno_per_game_title">Freedreno Settings</string>
+    <string name="freedreno_per_game_description">Configure GPU driver settings for this game</string>
+    <string name="freedreno_per_game_saved">Freedreno configuration saved</string>
+
     <!-- Gamepad Buttons -->
     <string name="gamepad_d_pad">D-pad</string>
     <string name="gamepad_left_stick">Left stick</string>

src/common/settings.h

@@ -536,7 +536,13 @@ struct Values {
                                                 Category::RendererExtensions,
                                                 Specialization::Scalar};
 
-    SwitchableSetting<bool> vertex_input_dynamic_state{linkage, true, "vertex_input_dynamic_state", Category::RendererExtensions};
+    SwitchableSetting<bool> vertex_input_dynamic_state{linkage,
+#if defined (ANDROID)
+                                                       false,
+#else
+                                                       true,
+#endif
+                                                       "vertex_input_dynamic_state", Category::RendererExtensions};
     SwitchableSetting<bool> provoking_vertex{linkage, false, "provoking_vertex", Category::RendererExtensions};
     SwitchableSetting<bool> descriptor_indexing{linkage, false, "descriptor_indexing", Category::RendererExtensions};
 

src/core/hle/kernel/kernel.cpp

@@ -50,13 +50,38 @@
 
 namespace Kernel {
 
+// Can only be used by a single implementation PER THREAD
+struct ThreadLocalData {
+    std::optional<KThread> raw_thread;
+    KThread* current_thread = nullptr;
+    KThread* thread = nullptr;
+    u8 host_thread_id = UINT8_MAX;
+    bool is_phantom_mode_for_singlecore = false;
+    bool lock = false;
+};
+
 struct KernelCore::Impl {
     static constexpr size_t ApplicationMemoryBlockSlabHeapSize = 20000;
     static constexpr size_t SystemMemoryBlockSlabHeapSize = 10000;
     static constexpr size_t BlockInfoSlabHeapSize = 4000;
     static constexpr size_t ReservedDynamicPageCount = 64;
+    // Be very careful when handling TLS data
+    // We do not want to concern ourselves with the appropriate way to manage them
+    // across **all** threads, we just need these for a few spare threads (+host/guest threads)
+    //
+    // Do not just read straight from here, use a reference beforehand, the cost of reading
+    // from TLS is greater than the cost of reading normal variables.
+    // But account that this Impl() is instanced once per program, and shared across threads
+    // so we can't use a reference for now.
+    //
+    // And we have the guarantee that the data won't move out of the way so we can safely
+    // take a reference to it. This isn't always universally true but this is "global" data
+    // so it will be statically given a TLS slot anyways.
+    static inline thread_local ThreadLocalData tls_data = {};
 
-    explicit Impl(Core::System& system_, KernelCore& kernel_) : system{system_} {}
+    explicit Impl(Core::System& system_, KernelCore& kernel_) : system{system_} {
+        tls_data.lock = true;
+    }
 
     void SetMulticore(bool is_multi) {
         is_multicore = is_multi;
@@ -69,8 +94,6 @@ struct KernelCore::Impl {
         global_object_list_container = std::make_unique<KAutoObjectWithListContainer>(kernel);
         global_scheduler_context = std::make_unique<Kernel::GlobalSchedulerContext>(kernel);
 
-        is_phantom_mode_for_singlecore = false;
-
         // Derive the initial memory layout from the emulated board
         Init::InitializeSlabResourceCounts(kernel);
         DeriveInitialMemoryLayout();
@@ -88,9 +111,7 @@ struct KernelCore::Impl {
         {
             const auto& pt_heap_region = memory_layout->GetPageTableHeapRegion();
             ASSERT(pt_heap_region.GetEndAddress() != 0);
-
-            InitializeResourceManagers(kernel, pt_heap_region.GetAddress(),
-                                       pt_heap_region.GetSize());
+            InitializeResourceManagers(kernel, pt_heap_region.GetAddress(), pt_heap_region.GetSize());
         }
 
         InitializeHackSharedMemory(kernel);
@@ -222,17 +243,11 @@ struct KernelCore::Impl {
         const auto kernel_size{sizes.second};
 
         // If setting the default system values fails, then something seriously wrong has occurred.
-        ASSERT(
-            system_resource_limit->SetLimitValue(LimitableResource::PhysicalMemoryMax, total_size)
-                .IsSuccess());
-        ASSERT(system_resource_limit->SetLimitValue(LimitableResource::ThreadCountMax, 800)
-                   .IsSuccess());
-        ASSERT(system_resource_limit->SetLimitValue(LimitableResource::EventCountMax, 900)
-                   .IsSuccess());
-        ASSERT(system_resource_limit->SetLimitValue(LimitableResource::TransferMemoryCountMax, 200)
-                   .IsSuccess());
-        ASSERT(system_resource_limit->SetLimitValue(LimitableResource::SessionCountMax, 1133)
-                   .IsSuccess());
+        ASSERT(system_resource_limit->SetLimitValue(LimitableResource::PhysicalMemoryMax, total_size).IsSuccess());
+        ASSERT(system_resource_limit->SetLimitValue(LimitableResource::ThreadCountMax, 800).IsSuccess());
+        ASSERT(system_resource_limit->SetLimitValue(LimitableResource::EventCountMax, 900).IsSuccess());
+        ASSERT(system_resource_limit->SetLimitValue(LimitableResource::TransferMemoryCountMax, 200).IsSuccess());
+        ASSERT(system_resource_limit->SetLimitValue(LimitableResource::SessionCountMax, 1133).IsSuccess());
         system_resource_limit->Reserve(LimitableResource::PhysicalMemoryMax, kernel_size);
 
         // Reserve secure applet memory, introduced in firmware 5.0.0
@@ -242,16 +257,13 @@ struct KernelCore::Impl {
     }
 
     void InitializePreemption(KernelCore& kernel) {
-        preemption_event = Core::Timing::CreateEvent(
-            "PreemptionCallback",
-            [this, &kernel](s64 time,
-                            std::chrono::nanoseconds) -> std::optional<std::chrono::nanoseconds> {
-                {
-                    KScopedSchedulerLock lock(kernel);
-                    global_scheduler_context->PreemptThreads();
-                }
-                return std::nullopt;
-            });
+        preemption_event = Core::Timing::CreateEvent("PreemptionCallback", [this, &kernel](s64 time, std::chrono::nanoseconds) -> std::optional<std::chrono::nanoseconds> {
+            {
+                KScopedSchedulerLock lock(kernel);
+                global_scheduler_context->PreemptThreads();
+            }
+            return std::nullopt;
+        });
 
         const auto time_interval = std::chrono::nanoseconds{std::chrono::milliseconds(10)};
         system.CoreTiming().ScheduleLoopingEvent(time_interval, time_interval, preemption_event);
@@ -263,15 +275,13 @@ struct KernelCore::Impl {
         ASSERT(Common::IsAligned(size, PageSize));
 
         // Ensure that we have space for our reference counts.
-        const size_t rc_size =
-            Common::AlignUp(KPageTableSlabHeap::CalculateReferenceCountSize(size), PageSize);
+        const size_t rc_size = Common::AlignUp(KPageTableSlabHeap::CalculateReferenceCountSize(size), PageSize);
         ASSERT(rc_size < size);
         size -= rc_size;
 
         // Initialize the resource managers' shared page manager.
         resource_manager_page_manager = std::make_unique<KDynamicPageManager>();
-        resource_manager_page_manager->Initialize(
-            address, size, std::max<size_t>(PageSize, KPageBufferSlabHeap::BufferSize));
+        resource_manager_page_manager->Initialize(address, size, std::max<size_t>(PageSize, KPageBufferSlabHeap::BufferSize));
 
         // Initialize the KPageBuffer slab heap.
         page_buffer_slab_heap.Initialize(system);
@@ -280,16 +290,12 @@ struct KernelCore::Impl {
         app_memory_block_heap = std::make_unique<KMemoryBlockSlabHeap>();
         sys_memory_block_heap = std::make_unique<KMemoryBlockSlabHeap>();
         block_info_heap = std::make_unique<KBlockInfoSlabHeap>();
-        app_memory_block_heap->Initialize(resource_manager_page_manager.get(),
-                                          ApplicationMemoryBlockSlabHeapSize);
-        sys_memory_block_heap->Initialize(resource_manager_page_manager.get(),
-                                          SystemMemoryBlockSlabHeapSize);
+        app_memory_block_heap->Initialize(resource_manager_page_manager.get(), ApplicationMemoryBlockSlabHeapSize);
+        sys_memory_block_heap->Initialize(resource_manager_page_manager.get(), SystemMemoryBlockSlabHeapSize);
         block_info_heap->Initialize(resource_manager_page_manager.get(), BlockInfoSlabHeapSize);
 
         // Reserve all but a fixed number of remaining pages for the page table heap.
-        const size_t num_pt_pages = resource_manager_page_manager->GetCount() -
-                                    resource_manager_page_manager->GetUsed() -
-                                    ReservedDynamicPageCount;
+        const size_t num_pt_pages = resource_manager_page_manager->GetCount() - resource_manager_page_manager->GetUsed() - ReservedDynamicPageCount;
         page_table_heap = std::make_unique<KPageTableSlabHeap>();
 
         // TODO(bunnei): Pass in address once we support kernel virtual memory allocations.
@@ -301,8 +307,8 @@ struct KernelCore::Impl {
         KDynamicPageManager* const app_dynamic_page_manager = nullptr;
         KDynamicPageManager* const sys_dynamic_page_manager =
             /*KTargetSystem::IsDynamicResourceLimitsEnabled()*/ true
-                ? resource_manager_page_manager.get()
-                : nullptr;
+            ? resource_manager_page_manager.get()
+            : nullptr;
         app_memory_block_manager = std::make_unique<KMemoryBlockSlabManager>();
         sys_memory_block_manager = std::make_unique<KMemoryBlockSlabManager>();
         app_block_info_manager = std::make_unique<KBlockInfoManager>();
@@ -320,9 +326,7 @@ struct KernelCore::Impl {
         sys_page_table_manager->Initialize(sys_dynamic_page_manager, page_table_heap.get());
 
         // Check that we have the correct number of dynamic pages available.
-        ASSERT(resource_manager_page_manager->GetCount() -
-                   resource_manager_page_manager->GetUsed() ==
-               ReservedDynamicPageCount);
+        ASSERT(resource_manager_page_manager->GetCount() - resource_manager_page_manager->GetUsed() == ReservedDynamicPageCount);
 
         // Create the system page table managers.
         app_system_resource = std::make_unique<KSystemResource>(kernel);
@@ -331,18 +335,15 @@ struct KernelCore::Impl {
         KAutoObject::Create(std::addressof(*sys_system_resource));
 
         // Set the managers for the system resources.
-        app_system_resource->SetManagers(*app_memory_block_manager, *app_block_info_manager,
-                                         *app_page_table_manager);
-        sys_system_resource->SetManagers(*sys_memory_block_manager, *sys_block_info_manager,
-                                         *sys_page_table_manager);
+        app_system_resource->SetManagers(*app_memory_block_manager, *app_block_info_manager, *app_page_table_manager);
+        sys_system_resource->SetManagers(*sys_memory_block_manager, *sys_block_info_manager, *sys_page_table_manager);
     }
 
     void InitializeShutdownThreads() {
         for (u32 core_id = 0; core_id < Core::Hardware::NUM_CPU_CORES; core_id++) {
             shutdown_threads[core_id] = KThread::Create(system.Kernel());
-            ASSERT(KThread::InitializeHighPriorityThread(system, shutdown_threads[core_id], {}, {},
-                                                         core_id)
-                       .IsSuccess());
+            ASSERT(KThread::InitializeHighPriorityThread(system, shutdown_threads[core_id], {}, {}, core_id)
+                .IsSuccess());
             KThread::Register(system.Kernel(), shutdown_threads[core_id]);
         }
     }
@@ -356,83 +357,77 @@ struct KernelCore::Impl {
         application_process->Open();
     }
 
-    static inline thread_local u8 host_thread_id = UINT8_MAX;
-
     /// Sets the host thread ID for the caller.
     u32 SetHostThreadId(std::size_t core_id) {
         // This should only be called during core init.
-        ASSERT(host_thread_id == UINT8_MAX);
+        ASSERT(tls_data.host_thread_id == UINT8_MAX);
 
         // The first four slots are reserved for CPU core threads
         ASSERT(core_id < Core::Hardware::NUM_CPU_CORES);
-        host_thread_id = static_cast<u8>(core_id);
-        return host_thread_id;
+        tls_data.host_thread_id = u8(core_id);
+        return tls_data.host_thread_id;
     }
 
     /// Gets the host thread ID for the caller
     u32 GetHostThreadId() const {
-        return host_thread_id;
+        return tls_data.host_thread_id;
     }
 
     // Gets the dummy KThread for the caller, allocating a new one if this is the first time
     KThread* GetHostDummyThread(KThread* existing_thread) {
-        const auto initialize{[](KThread* thread) {
-            ASSERT(KThread::InitializeDummyThread(thread, nullptr).IsSuccess());
-            return thread;
-        }};
-        thread_local KThread raw_thread{system.Kernel()};
-        thread_local KThread* thread = existing_thread ? existing_thread : initialize(&raw_thread);
-        return thread;
+        if (tls_data.thread == nullptr) {
+            auto const initialize{[](KThread* thread) {
+                ASSERT(KThread::InitializeDummyThread(thread, nullptr).IsSuccess());
+                return thread;
+            }};
+            tls_data.raw_thread.emplace(system.Kernel());
+            tls_data.thread = existing_thread ? existing_thread : initialize(&*tls_data.raw_thread);
+            ASSERT(tls_data.thread != nullptr);
+        }
+        return tls_data.thread;
     }
 
     /// Registers a CPU core thread by allocating a host thread ID for it
     void RegisterCoreThread(std::size_t core_id) {
         ASSERT(core_id < Core::Hardware::NUM_CPU_CORES);
         const auto this_id = SetHostThreadId(core_id);
-        if (!is_multicore) {
+        if (!is_multicore)
             single_core_thread_id = this_id;
-        }
     }
 
     /// Registers a new host thread by allocating a host thread ID for it
     void RegisterHostThread(KThread* existing_thread) {
-        [[maybe_unused]] const auto dummy_thread = GetHostDummyThread(existing_thread);
+        (void)GetHostDummyThread(existing_thread);
     }
 
     [[nodiscard]] u32 GetCurrentHostThreadID() {
-        const auto this_id = GetHostThreadId();
-        if (!is_multicore && single_core_thread_id == this_id) {
-            return static_cast<u32>(system.GetCpuManager().CurrentCore());
-        }
+        auto const this_id = GetHostThreadId();
+        if (!is_multicore && single_core_thread_id == this_id)
+            return u32(system.GetCpuManager().CurrentCore());
         return this_id;
     }
 
-    static inline thread_local bool is_phantom_mode_for_singlecore{false};
-
+    // Forces singlecore
     bool IsPhantomModeForSingleCore() const {
-        return is_phantom_mode_for_singlecore;
+        return tls_data.is_phantom_mode_for_singlecore;
     }
-
     void SetIsPhantomModeForSingleCore(bool value) {
         ASSERT(!is_multicore);
-        is_phantom_mode_for_singlecore = value;
+        tls_data.is_phantom_mode_for_singlecore = value;
     }
 
     bool IsShuttingDown() const {
         return is_shutting_down.load(std::memory_order_relaxed);
     }
 
-    static inline thread_local KThread* current_thread{nullptr};
-
     KThread* GetCurrentEmuThread() {
-        if (!current_thread) {
-            current_thread = GetHostDummyThread(nullptr);
-        }
-        return current_thread;
+        if (!tls_data.current_thread)
+            tls_data.current_thread = GetHostDummyThread(nullptr);
+        return tls_data.current_thread;
     }
 
     void SetCurrentEmuThread(KThread* thread) {
-        current_thread = thread;
+        tls_data.current_thread = thread;
     }
 
     void DeriveInitialMemoryLayout() {

src/core/hle/service/acc/acc.cpp

@@ -100,6 +100,7 @@ public:
             {133, nullptr, "GetNintendoAccountVerificationUrlCache"}, // 9.0.0+
             {134, nullptr, "RefreshNintendoAccountVerificationUrlCache"}, // 9.0.0+
             {135, nullptr, "RefreshNintendoAccountVerificationUrlCacheAsyncIfSecondsElapsed"}, // 9.0.0+
+            {136, nullptr, "GetNintendoAccountUserResourceCache"}, // 19.0.0+
             {140, nullptr, "GetNetworkServiceLicenseCache"}, // 5.0.0+
             {141, nullptr, "RefreshNetworkServiceLicenseCacheAsync"}, // 5.0.0+
             {142, nullptr, "RefreshNetworkServiceLicenseCacheAsyncIfSecondsElapsed"}, // 5.0.0+

src/core/hle/service/am/service/application_functions.cpp

@@ -90,10 +90,10 @@ IApplicationFunctions::IApplicationFunctions(Core::System& system_, std::shared_
         {200, nullptr, "GetLastApplicationExitReason"},
         {210, D<&IApplicationFunctions::GetUnknownEvent210>, "Unknown210"},
         {220, nullptr, "Unknown220"}, // [20.0.0+]
-        {300, nullptr, "Unknown300"}, // [20.0.0+]
+        {300, nullptr, "CreateMovieWriter"}, // [19.0.0+]
         {310, nullptr, "Unknown310"}, // [20.0.0+]
         {320, nullptr, "Unknown320"}, // [20.0.0+]
-        {330, nullptr, "Unknown330"}, // [20.0.0+]
+        {330, D<&IApplicationFunctions::Unknown330>, "Unknown330"}, // [20.0.0+]
         {500, nullptr, "StartContinuousRecordingFlushForDebug"},
         {1000, nullptr, "CreateMovieMaker"},
         {1001, D<&IApplicationFunctions::PrepareForJit>, "PrepareForJit"},
@@ -503,6 +503,12 @@ Result IApplicationFunctions::GetUnknownEvent210(
     R_SUCCEED();
 }
 
+Result IApplicationFunctions::Unknown330(Out<u8> out) {
+    LOG_DEBUG(Service_AM, "called");
+    *out = 0;
+    R_SUCCEED();
+}
+
 Result IApplicationFunctions::PrepareForJit() {
     LOG_WARNING(Service_AM, "(STUBBED) called");
 

src/core/hle/service/am/service/application_functions.h

@@ -80,6 +80,7 @@ private:
     Result GetNotificationStorageChannelEvent(OutCopyHandle<Kernel::KReadableEvent> out_event);
     Result GetHealthWarningDisappearedSystemEvent(OutCopyHandle<Kernel::KReadableEvent> out_event);
     Result GetUnknownEvent210(OutCopyHandle<Kernel::KReadableEvent> out_event);
+    Result Unknown330(Out<u8> out);
     Result PrepareForJit();
 
     const std::shared_ptr<Applet> m_applet;

src/core/hle/service/am/service/common_state_getter.cpp

@@ -36,19 +36,20 @@ ICommonStateGetter::ICommonStateGetter(Core::System& system_, std::shared_ptr<Ap
         {11, D<&ICommonStateGetter::ReleaseSleepLock>, "ReleaseSleepLock"},
         {12, D<&ICommonStateGetter::ReleaseSleepLockTransiently>, "ReleaseSleepLockTransiently"},
         {13, D<&ICommonStateGetter::GetAcquiredSleepLockEvent>, "GetAcquiredSleepLockEvent"},
-        {14, nullptr, "GetWakeupCount"},
+        {14, nullptr, "GetWakeupCount"}, //11.0.0+
+        {15, nullptr, "Unknown15"}, //19.0.0+
         {20, D<&ICommonStateGetter::PushToGeneralChannel>, "PushToGeneralChannel"},
         {30, nullptr, "GetHomeButtonReaderLockAccessor"},
-        {31, D<&ICommonStateGetter::GetReaderLockAccessorEx>, "GetReaderLockAccessorEx"},
-        {32, D<&ICommonStateGetter::GetWriterLockAccessorEx>, "GetWriterLockAccessorEx"},
-        {40, nullptr, "GetCradleFwVersion"},
-        {50, D<&ICommonStateGetter::IsVrModeEnabled>, "IsVrModeEnabled"},
-        {51, D<&ICommonStateGetter::SetVrModeEnabled>, "SetVrModeEnabled"},
-        {52, D<&ICommonStateGetter::SetLcdBacklighOffEnabled>, "SetLcdBacklighOffEnabled"},
-        {53, D<&ICommonStateGetter::BeginVrModeEx>, "BeginVrModeEx"},
-        {54, D<&ICommonStateGetter::EndVrModeEx>, "EndVrModeEx"},
-        {55, D<&ICommonStateGetter::IsInControllerFirmwareUpdateSection>, "IsInControllerFirmwareUpdateSection"},
-        {59, nullptr, "SetVrPositionForDebug"},
+        {31, D<&ICommonStateGetter::GetReaderLockAccessorEx>, "GetReaderLockAccessorEx"}, //2.0.0+
+        {32, D<&ICommonStateGetter::GetWriterLockAccessorEx>, "GetWriterLockAccessorEx"}, //7.0.0+
+        {40, nullptr, "GetCradleFwVersion"}, //2.0.0+
+        {50, D<&ICommonStateGetter::IsVrModeEnabled>, "IsVrModeEnabled"}, //3.0.0+
+        {51, D<&ICommonStateGetter::SetVrModeEnabled>, "SetVrModeEnabled"}, //3.0.0+
+        {52, D<&ICommonStateGetter::SetLcdBacklighOffEnabled>, "SetLcdBacklighOffEnabled"}, //4.0.0+
+        {53, D<&ICommonStateGetter::BeginVrModeEx>, "BeginVrModeEx"}, //7.0.0+
+        {54, D<&ICommonStateGetter::EndVrModeEx>, "EndVrModeEx"}, //7.0.0+
+        {55, D<&ICommonStateGetter::IsInControllerFirmwareUpdateSection>, "IsInControllerFirmwareUpdateSection"}, //3.0.0+
+        {59, nullptr, "SetVrPositionForDebug"}, //1.0.0+
         {60, D<&ICommonStateGetter::GetDefaultDisplayResolution>, "GetDefaultDisplayResolution"},
         {61, D<&ICommonStateGetter::GetDefaultDisplayResolutionChangeEvent>, "GetDefaultDisplayResolutionChangeEvent"},
         {62, D<&ICommonStateGetter::GetHdcpAuthenticationState>, "GetHdcpAuthenticationState"},
@@ -63,7 +64,8 @@ ICommonStateGetter::ICommonStateGetter(Core::System& system_, std::shared_ptr<Ap
         {91, nullptr, "GetCurrentPerformanceConfiguration"},
         {100, D<&ICommonStateGetter::SetHandlingHomeButtonShortPressedEnabled>, "SetHandlingHomeButtonShortPressedEnabled"},
         {110, nullptr, "OpenMyGpuErrorHandler"},
-        {120, D<&ICommonStateGetter::GetAppletLaunchedHistory>, "GetAppletLaunchedHistory"},
+        {120, D<&ICommonStateGetter::GetAppletLaunchedHistory>, "GetAppletLaunchedHistory"}, //13.0.0+
+        {130, nullptr, "Unknown130"}, //21.0.0+
         {200, D<&ICommonStateGetter::GetOperationModeSystemInfo>, "GetOperationModeSystemInfo"},
         {300, D<&ICommonStateGetter::GetSettingsPlatformRegion>, "GetSettingsPlatformRegion"},
         {400, nullptr, "ActivateMigrationService"},
@@ -72,7 +74,13 @@ ICommonStateGetter::ICommonStateGetter(Core::System& system_, std::shared_ptr<Ap
         {501, nullptr, "SuppressDisablingSleepTemporarily"},
         {502, nullptr, "IsSleepEnabled"},
         {503, nullptr, "IsDisablingSleepSuppressed"},
-        {900, D<&ICommonStateGetter::SetRequestExitToLibraryAppletAtExecuteNextProgramEnabled>, "SetRequestExitToLibraryAppletAtExecuteNextProgramEnabled"},
+        {600, nullptr, "Unknown600"}, //20.0.0+
+        {610, nullptr, "Unknown610"}, //21.0.0+
+        {900, D<&ICommonStateGetter::SetRequestExitToLibraryAppletAtExecuteNextProgramEnabled>, "SetRequestExitToLibraryAppletAtExecuteNextProgramEnabled"}, //11.0.0+
+        {910, nullptr, "GetLaunchRequiredTick"}, //17.0.0+
+        {1000, nullptr, "BeginVrMode3d"}, //19.0.0+
+        {1001, nullptr, "EndVrMode3d"}, //19.0.0+
+        {1002, nullptr, "IsVrModeEnabled3d"}, //19.0.0+
     };
     // clang-format on
 

src/core/hle/service/am/service/library_applet_accessor.cpp

@@ -27,8 +27,10 @@ ILibraryAppletAccessor::ILibraryAppletAccessor(Core::System& system_,
         {25, D<&ILibraryAppletAccessor::Terminate>, "Terminate"},
         {30, D<&ILibraryAppletAccessor::GetResult>, "GetResult"},
         {50, nullptr, "SetOutOfFocusApplicationSuspendingEnabled"},
-        {60, D<&ILibraryAppletAccessor::PresetLibraryAppletGpuTimeSliceZero>, "PresetLibraryAppletGpuTimeSliceZero"},
-        {90, D<&ILibraryAppletAccessor::Unknown90>, "Unknown90"},
+        {60, D<&ILibraryAppletAccessor::PresetLibraryAppletGpuTimeSliceZero>, "PresetLibraryAppletGpuTimeSliceZero"}, //10.0.0+
+        {80, nullptr, "RequestForLibraryAppletToGetForeground"}, //19.0.0+
+        {81, nullptr, "GetCurrentChildLibraryApplet"}, //19.0.0+
+        {90, D<&ILibraryAppletAccessor::Unknown90>, "Unknown90"}, //20.0.0+
         {100, D<&ILibraryAppletAccessor::PushInData>, "PushInData"},
         {101, D<&ILibraryAppletAccessor::PopOutData>, "PopOutData"},
         {102, nullptr, "PushExtraStorage"},
@@ -39,7 +41,7 @@ ILibraryAppletAccessor::ILibraryAppletAccessor(Core::System& system_,
         {110, nullptr, "NeedsToExitProcess"},
         {120, nullptr, "GetLibraryAppletInfo"},
         {150, nullptr, "RequestForAppletToGetForeground"},
-        {160, D<&ILibraryAppletAccessor::GetIndirectLayerConsumerHandle>, "GetIndirectLayerConsumerHandle"},
+        {160, D<&ILibraryAppletAccessor::GetIndirectLayerConsumerHandle>, "GetIndirectLayerConsumerHandle"}, //2.0.0+
     };
     // clang-format on
 

src/core/hle/service/audio/audio_renderer.cpp

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2024 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
@@ -28,10 +31,10 @@ IAudioRenderer::IAudioRenderer(Core::System& system_, Manager& manager_,
         {7, D<&IAudioRenderer::QuerySystemEvent>, "QuerySystemEvent"},
         {8, D<&IAudioRenderer::SetRenderingTimeLimit>, "SetRenderingTimeLimit"},
         {9, D<&IAudioRenderer::GetRenderingTimeLimit>, "GetRenderingTimeLimit"},
-        {10, D<&IAudioRenderer::RequestUpdateAuto>, "RequestUpdateAuto"},
-        {11, nullptr, "ExecuteAudioRendererRendering"},
-        {12, D<&IAudioRenderer::SetVoiceDropParameter>, "SetVoiceDropParameter"},
-        {13, D<&IAudioRenderer::GetVoiceDropParameter>, "GetVoiceDropParameter"},
+        {10, D<&IAudioRenderer::RequestUpdateAuto>, "RequestUpdateAuto"}, //3.0.0+
+        {11, nullptr, "ExecuteAudioRendererRendering"}, //3.0.0+
+        {12, D<&IAudioRenderer::SetVoiceDropParameter>, "SetVoiceDropParameter"}, //15.0.0+
+        {13, D<&IAudioRenderer::GetVoiceDropParameter>, "GetVoiceDropParameter"}, //15.0.0+
     };
     // clang-format on
     RegisterHandlers(functions);

src/core/hle/service/service.cpp

@@ -19,16 +19,13 @@
 namespace Service {
 
 /// @brief Creates a function string for logging, complete with the name (or header code, depending
-/// on what's passed in) the port name, and all the cmd_buff arguments.
-[[maybe_unused]] static std::string MakeFunctionString(std::string_view name, std::string_view port_name, const u32* cmd_buff) {
+/// on what's passed in) the port name, and all the cmd_buf arguments.
+[[maybe_unused]] static std::string MakeFunctionString(std::string_view name, std::string_view port_name, const u32* cmd_buf) {
     // Number of params == bits 0-5 + bits 6-11
-    int num_params = (cmd_buff[0] & 0x3F) + ((cmd_buff[0] >> 6) & 0x3F);
-
-    std::string function_string = fmt::format("function '{}': port={}", name, port_name);
-
+    int num_params = (cmd_buf[0] & 0x3F) + ((cmd_buf[0] >> 6) & 0x3F);
+    std::string function_string = fmt::format("fn '{}': port={}", name, port_name);
     for (int i = 1; i <= num_params; ++i)
-        function_string += fmt::format(", cmd_buff[{}]={:#X}", i, cmd_buff[i]);
-
+        function_string += fmt::format(", cmd_buf[{}]={:#X}", i, cmd_buf[i]);
     return function_string;
 }
 

src/core/hle/service/sm/sm_controller.cpp

@@ -119,7 +119,7 @@ Controller::Controller(Core::System& system_) : ServiceFramework{system_, "IpcCo
         {2, &Controller::CloneCurrentObject, "CloneCurrentObject"},
         {3, &Controller::QueryPointerBufferSize, "QueryPointerBufferSize"},
         {4, &Controller::CloneCurrentObjectEx, "CloneCurrentObjectEx"},
-        {5, &Controller::SetPointerBufferSize, "SetPointerBufferSize"},
+        {5, &Controller::SetPointerBufferSize, "SetPointerBufferSize"}, //TODO: where does this come from
     };
     RegisterHandlers(functions);
 }

src/core/hle/service/sockets/bsd.cpp

@@ -1015,11 +1015,23 @@ BSD::BSD(Core::System& system_, const char* name)
         {26, &BSD::Close, "Close"},
         {27, &BSD::DuplicateSocket, "DuplicateSocket"},
         {28, nullptr, "GetResourceStatistics"},
-        {29, nullptr, "RecvMMsg"},
-        {30, nullptr, "SendMMsg"},
-        {31, &BSD::EventFd, "EventFd"},
-        {32, nullptr, "RegisterResourceStatisticsName"},
-        {33, nullptr, "Initialize2"},
+        {29, nullptr, "RecvMMsg"}, //3.0.0+
+        {30, nullptr, "SendMMsg"}, //3.0.0+
+        {31, &BSD::EventFd, "EventFd"}, //7.0.0+
+        {32, nullptr, "RegisterResourceStatisticsName"}, //7.0.0+
+        {33, nullptr, "RegisterClientShared"}, //10.0.0+
+        {34, nullptr, "GetSocketStatistics"}, //15.0.0+
+        {35, nullptr, "NifIoctl"}, //17.0.0+
+        {36, nullptr, "Unknown36"}, //18.0.0+
+        {37, nullptr, "Unknown37"}, //18.0.0+
+        {38, nullptr, "Unknown38"}, //18.0.0+
+        {39, nullptr, "Unknown39"}, //20.0.0+
+        {40, nullptr, "Unknown40"}, //20.0.0+
+        {41, nullptr, "Unknown41"}, //21.0.0+
+        {42, nullptr, "Unknown42"}, //21.0.0+
+        {43, nullptr, "Unknown43"}, //21.0.0+
+        {200, nullptr, "SetThreadCoreMask"}, //15.0.0+
+        {201, nullptr, "GetThreadCoreMask"}, //15.0.0+
     };
     // clang-format on
 

src/core/hle/service/sockets/sockets.h

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2018 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
@@ -82,6 +85,10 @@ enum class OptName : u32 {
     RCVTIMEO = 0x1006,
     ERROR_ = 0x1007,   // avoid name collision with Windows macro
     NOSIGPIPE = 0x800, // at least according to libnx
+    ACCEPTFILTER = 0x1000,
+    BINTIME = 0x2000,
+    NO_OFFLOAD = 0x4000,
+    NO_DDP = 0x8000,
 };
 
 enum class ShutdownHow : s32 {

src/dynarmic/CMakeLists.txt

@@ -9,6 +9,8 @@ set(dynarmic_VERSION_PATCH 0)
 
 set(dynarmic_VERSION ${dynarmic_VERSION_MAJOR}.${dynarmic_VERSION_MINOR}.${dynarmic_VERSION_PATCH})
 
+project(dynarmic LANGUAGES C CXX ASM VERSION ${dynarmic_VERSION})
+
 # Determine if we're built as a subproject (using add_subdirectory)
 # or if this is the master project.
 set(MASTER_PROJECT OFF)
@@ -30,7 +32,7 @@ else()
 endif()
 option(DYNARMIC_ENABLE_NO_EXECUTE_SUPPORT "Enables support for systems that require W^X" ${REQUIRE_WX})
 
-option(DYNARMIC_IGNORE_ASSERTS "Ignore asserts" OFF)
+option(DYNARMIC_IGNORE_ASSERTS "Ignore asserts" ON)
 option(DYNARMIC_TESTS_USE_UNICORN "Enable fuzzing tests against unicorn" OFF)
 CMAKE_DEPENDENT_OPTION(DYNARMIC_USE_LLVM "Support disassembly of jitted x86_64 code using LLVM" OFF "NOT YUZU_DISABLE_LLVM" OFF)
 

src/dynarmic/src/dynarmic/backend/x64/a32_emit_x64.cpp

@@ -102,9 +102,6 @@ A32EmitX64::BlockDescriptor A32EmitX64::Emit(IR::Block& block) {
     }
 
     code.EnableWriting();
-    SCOPE_EXIT {
-        code.DisableWriting();
-    };
 
     const boost::container::static_vector<HostLoc, 28> gpr_order = [this] {
         boost::container::static_vector<HostLoc, 28> gprs{any_gpr};
@@ -126,37 +123,31 @@ A32EmitX64::BlockDescriptor A32EmitX64::Emit(IR::Block& block) {
 
     EmitCondPrelude(ctx);
 
-    auto const loop_all_inst = [this, &block, &ctx](auto const func) {
-        for (auto iter = block.begin(); iter != block.end(); ++iter) [[likely]] {
-            auto* inst = &*iter;
-            // Call the relevant Emit* member function.
-            switch (inst->GetOpcode()) {
+    for (auto iter = block.begin(); iter != block.end(); ++iter) [[likely]] {
+        auto* inst = &*iter;
+        // Call the relevant Emit* member function.
+        switch (inst->GetOpcode()) {
 #define OPCODE(name, type, ...)                     \
-            case IR::Opcode::name:                  \
-                A32EmitX64::Emit##name(ctx, inst);  \
-                break;
+        case IR::Opcode::name:                  \
+            A32EmitX64::Emit##name(ctx, inst);  \
+            break;
 #define A32OPC(name, type, ...)                     \
-            case IR::Opcode::A32##name:             \
-                A32EmitX64::EmitA32##name(ctx, inst);\
-                break;
+        case IR::Opcode::A32##name:             \
+            A32EmitX64::EmitA32##name(ctx, inst);\
+            break;
 #define A64OPC(...)
 #include "dynarmic/ir/opcodes.inc"
 #undef OPCODE
 #undef A32OPC
 #undef A64OPC
-            default:
-                UNREACHABLE();
-            }
-            reg_alloc.EndOfAllocScope();
-            func(reg_alloc);
+        default:
+            UNREACHABLE();
         }
-    };
-    if (!conf.very_verbose_debugging_output) [[likely]] {
-        loop_all_inst([](auto&) { /*noop*/ });
-    } else [[unlikely]] {
-        loop_all_inst([this](auto& reg_alloc) {
+        reg_alloc.EndOfAllocScope();
+#ifndef NDEBUG
+        if (conf.very_verbose_debugging_output)
             EmitVerboseDebuggingOutput(reg_alloc);
-        });
+#endif
     }
 
     reg_alloc.AssertNoMoreUses();
@@ -172,7 +163,7 @@ A32EmitX64::BlockDescriptor A32EmitX64::Emit(IR::Block& block) {
     }
     code.int3();
 
-    const size_t size = static_cast<size_t>(code.getCurr() - entrypoint);
+    const size_t size = size_t(code.getCurr() - entrypoint);
 
     const A32::LocationDescriptor descriptor{block.Location()};
     const A32::LocationDescriptor end_location{block.EndLocation()};
@@ -180,7 +171,9 @@ A32EmitX64::BlockDescriptor A32EmitX64::Emit(IR::Block& block) {
     const auto range = boost::icl::discrete_interval<u32>::closed(descriptor.PC(), end_location.PC() - 1);
     block_ranges.AddRange(range, descriptor);
 
-    return RegisterBlock(descriptor, entrypoint, size);
+    auto const bdesc = RegisterBlock(descriptor, entrypoint, size);
+    code.DisableWriting();
+    return bdesc;
 }
 
 void A32EmitX64::ClearCache() {

src/dynarmic/src/dynarmic/backend/x64/a64_emit_x64.cpp

@@ -76,10 +76,6 @@ A64EmitX64::BlockDescriptor A64EmitX64::Emit(IR::Block& block) noexcept {
     }
 
     code.EnableWriting();
-    SCOPE_EXIT {
-        code.DisableWriting();
-    };
-
     const boost::container::static_vector<HostLoc, 28> gpr_order = [this] {
         boost::container::static_vector<HostLoc, 28> gprs{any_gpr};
         if (conf.fastmem_pointer) {
@@ -141,9 +137,10 @@ a64_branch:
         (this->*a64_handlers[size_t(opcode) - std::size(opcode_handlers)])(ctx, &inst);
 finish_this_inst:
         ctx.reg_alloc.EndOfAllocScope();
-        if (conf.very_verbose_debugging_output) [[unlikely]] {
+#ifndef NDEBUG
+        if (conf.very_verbose_debugging_output)
             EmitVerboseDebuggingOutput(reg_alloc);
-        }
+#endif
     }
 
     reg_alloc.AssertNoMoreUses();
@@ -167,7 +164,9 @@ finish_this_inst:
     const auto range = boost::icl::discrete_interval<u64>::closed(descriptor.PC(), end_location.PC() - 1);
     block_ranges.AddRange(range, descriptor);
 
-    return RegisterBlock(descriptor, entrypoint, size);
+    auto bdesc = RegisterBlock(descriptor, entrypoint, size);
+    code.DisableWriting();
+    return bdesc;
 }
 
 void A64EmitX64::ClearCache() {

src/dynarmic/src/dynarmic/backend/x64/emit_x64.cpp

@@ -105,6 +105,7 @@ void EmitX64::PushRSBHelper(Xbyak::Reg64 loc_desc_reg, Xbyak::Reg64 index_reg, I
     code.mov(dword[code.ABI_JIT_PTR + code.GetJitStateInfo().offsetof_rsb_ptr], index_reg.cvt32());
 }
 
+#ifndef NDEBUG
 void EmitX64::EmitVerboseDebuggingOutput(RegAlloc& reg_alloc) {
     code.lea(rsp, ptr[rsp - sizeof(RegisterData)]);
     code.stmxcsr(dword[rsp + offsetof(RegisterData, mxcsr)]);
@@ -134,6 +135,7 @@ void EmitX64::EmitVerboseDebuggingOutput(RegAlloc& reg_alloc) {
     code.ldmxcsr(dword[rsp + offsetof(RegisterData, mxcsr)]);
     code.add(rsp, sizeof(RegisterData));
 }
+#endif
 
 void EmitX64::EmitPushRSB(EmitContext& ctx, IR::Inst* inst) {
     auto args = ctx.reg_alloc.GetArgumentInfo(inst);

src/dynarmic/src/dynarmic/backend/x64/emit_x64.h

@@ -113,7 +113,9 @@ public:
     BlockDescriptor RegisterBlock(const IR::LocationDescriptor& location_descriptor, CodePtr entrypoint, size_t size);
     void PushRSBHelper(Xbyak::Reg64 loc_desc_reg, Xbyak::Reg64 index_reg, IR::LocationDescriptor target);
 
+#ifndef NDEBUG
     void EmitVerboseDebuggingOutput(RegAlloc& reg_alloc);
+#endif
     virtual void EmitTerminal(IR::Terminal terminal, IR::LocationDescriptor initial_location, bool is_single_step) noexcept = 0;
 
     // Patching

src/dynarmic/src/dynarmic/backend/x64/hostloc.h

@@ -58,7 +58,7 @@ enum class HostLoc : std::uint8_t {
     FirstSpill,
 };
 
-constexpr size_t NonSpillHostLocCount = static_cast<size_t>(HostLoc::FirstSpill);
+constexpr size_t NonSpillHostLocCount = size_t(HostLoc::FirstSpill);
 
 constexpr bool HostLocIsGPR(HostLoc reg) {
     return reg >= HostLoc::RAX && reg <= HostLoc::R15;

src/dynarmic/src/dynarmic/backend/x64/reg_alloc.cpp

@@ -56,26 +56,23 @@ static inline bool IsValuelessType(const IR::Type type) noexcept {
 }
 
 void HostLocInfo::ReleaseOne() noexcept {
-    is_being_used_count--;
+    ASSERT(is_being_used_count > 0);
+    --is_being_used_count;
     is_scratch = false;
-
-    if (current_references == 0)
-        return;
-
-    ASSERT(size_t(accumulated_uses) + 1 < (std::numeric_limits<uint16_t>::max)());
-    accumulated_uses++;
-    current_references--;
-
-    if (current_references == 0)
-        ReleaseAll();
+    if (current_references > 0) {
+        ASSERT(size_t(accumulated_uses) + 1 < (std::numeric_limits<decltype(accumulated_uses)>::max)());
+        ++accumulated_uses;
+        --current_references;
+        if (current_references == 0)
+            ReleaseAll();
+    }
 }
 
 void HostLocInfo::ReleaseAll() noexcept {
+    ASSERT(size_t(accumulated_uses) + current_references < (std::numeric_limits<decltype(accumulated_uses)>::max)());
     accumulated_uses += current_references;
     current_references = 0;
-
     is_set_last_use = false;
-
     if (total_uses == accumulated_uses) {
         values.clear();
         accumulated_uses = 0;
@@ -87,17 +84,19 @@ void HostLocInfo::ReleaseAll() noexcept {
     is_scratch = false;
 }
 
-void HostLocInfo::AddValue(IR::Inst* inst) noexcept {
+void HostLocInfo::AddValue(HostLoc loc, IR::Inst* inst) noexcept {
     if (is_set_last_use) {
         is_set_last_use = false;
         values.clear();
     }
     values.push_back(inst);
-    ASSERT(size_t(total_uses) + inst->UseCount() < (std::numeric_limits<uint16_t>::max)());
+
+    ASSERT(size_t(total_uses) + inst->UseCount() < (std::numeric_limits<decltype(total_uses)>::max)());
     total_uses += inst->UseCount();
     max_bit_width = std::max<uint8_t>(max_bit_width, std::countr_zero(GetBitWidth(inst->GetType())));
 }
 
+#ifndef NDEBUG
 void HostLocInfo::EmitVerboseDebuggingOutput(BlockOfCode& code, size_t host_loc_index) const noexcept {
     using namespace Xbyak::util;
     for (auto const value : values) {
@@ -108,6 +107,7 @@ void HostLocInfo::EmitVerboseDebuggingOutput(BlockOfCode& code, size_t host_loc_
         code.CallFunction(PrintVerboseDebuggingOutputLine);
     }
 }
+#endif
 
 bool Argument::FitsInImmediateU32() const noexcept {
     if (!IsImmediate())
@@ -197,12 +197,13 @@ RegAlloc::ArgumentInfo RegAlloc::GetArgumentInfo(const IR::Inst* inst) noexcept
         Argument{},
         Argument{}
     };
-    for (size_t i = 0; i < inst->NumArgs(); i++) {
+    for (size_t i = 0; i < inst->NumArgs() && i < 4; i++) {
         const auto arg = inst->GetArg(i);
         ret[i].value = arg;
         if (!arg.IsImmediate() && !IsValuelessType(arg.GetType())) {
-            ASSERT(ValueLocation(arg.GetInst()) && "argument must already been defined");
-            LocInfo(*ValueLocation(arg.GetInst())).AddArgReference();
+            auto const loc = ValueLocation(arg.GetInst());
+            ASSERT(loc && "argument must already been defined");
+            LocInfo(*loc).AddArgReference();
         }
     }
     return ret;
@@ -211,9 +212,9 @@ RegAlloc::ArgumentInfo RegAlloc::GetArgumentInfo(const IR::Inst* inst) noexcept
 void RegAlloc::RegisterPseudoOperation(const IR::Inst* inst) noexcept {
     ASSERT(IsValueLive(inst) || !inst->HasUses());
     for (size_t i = 0; i < inst->NumArgs(); i++) {
-        const auto arg = inst->GetArg(i);
+        auto const arg = inst->GetArg(i);
         if (!arg.IsImmediate() && !IsValuelessType(arg.GetType())) {
-            if (const auto loc = ValueLocation(arg.GetInst())) {
+            if (auto const loc = ValueLocation(arg.GetInst())) {
                 // May not necessarily have a value (e.g. CMP variant of Sub32).
                 LocInfo(*loc).AddArgReference();
             }
@@ -262,9 +263,8 @@ HostLoc RegAlloc::UseImpl(BlockOfCode& code, IR::Value use_value, const boost::c
         return LoadImmediate(code, use_value, ScratchImpl(code, desired_locations));
     }
 
-    const auto* use_inst = use_value.GetInst();
-    const HostLoc current_location = *ValueLocation(use_inst);
-    const size_t max_bit_width = LocInfo(current_location).GetMaxBitWidth();
+    auto const* use_inst = use_value.GetInst();
+    HostLoc const current_location = *ValueLocation(use_inst);
 
     const bool can_use_current_location = std::find(desired_locations.begin(), desired_locations.end(), current_location) != desired_locations.end();
     if (can_use_current_location) {
@@ -276,7 +276,8 @@ HostLoc RegAlloc::UseImpl(BlockOfCode& code, IR::Value use_value, const boost::c
         return UseScratchImpl(code, use_value, desired_locations);
     }
 
-    const HostLoc destination_location = SelectARegister(desired_locations);
+    size_t const max_bit_width = LocInfo(current_location).GetMaxBitWidth();
+    HostLoc const destination_location = SelectARegister(desired_locations);
     if (max_bit_width > HostLocBitWidth(destination_location)) {
         return UseScratchImpl(code, use_value, desired_locations);
     } else if (CanExchange(destination_location, current_location)) {
@@ -300,10 +301,10 @@ HostLoc RegAlloc::UseScratchImpl(BlockOfCode& code, IR::Value use_value, const b
 
     const bool can_use_current_location = std::find(desired_locations.begin(), desired_locations.end(), current_location) != desired_locations.end();
     if (can_use_current_location && !LocInfo(current_location).IsLocked()) {
-        if (!LocInfo(current_location).IsLastUse()) {
-            MoveOutOfTheWay(code, current_location);
+        if (LocInfo(current_location).IsLastUse()) {
+            LocInfo(current_location).is_set_last_use = true;
         } else {
-            LocInfo(current_location).SetLastUse();
+            MoveOutOfTheWay(code, current_location);
         }
         LocInfo(current_location).WriteLock();
         return current_location;
@@ -455,29 +456,31 @@ HostLoc RegAlloc::SelectARegister(const boost::container::static_vector<HostLoc,
 
 std::optional<HostLoc> RegAlloc::ValueLocation(const IR::Inst* value) const noexcept {
     for (size_t i = 0; i < hostloc_info.size(); i++)
-        if (hostloc_info[i].ContainsValue(value))
+        if (hostloc_info[i].ContainsValue(value)) {
+            //for (size_t j = 0; j < hostloc_info.size(); ++j)
+            //    ASSERT((i == j || !hostloc_info[j].ContainsValue(value)) && "duplicate defs");
             return HostLoc(i);
+        }
     return std::nullopt;
 }
 
 void RegAlloc::DefineValueImpl(BlockOfCode& code, IR::Inst* def_inst, HostLoc host_loc) noexcept {
     ASSERT(!ValueLocation(def_inst) && "def_inst has already been defined");
-    LocInfo(host_loc).AddValue(def_inst);
+    LocInfo(host_loc).AddValue(host_loc, def_inst);
+    ASSERT(*ValueLocation(def_inst) == host_loc);
 }
 
 void RegAlloc::DefineValueImpl(BlockOfCode& code, IR::Inst* def_inst, const IR::Value& use_inst) noexcept {
     ASSERT(!ValueLocation(def_inst) && "def_inst has already been defined");
-
     if (use_inst.IsImmediate()) {
         const HostLoc location = ScratchImpl(code, gpr_order);
         DefineValueImpl(code, def_inst, location);
         LoadImmediate(code, use_inst, location);
-        return;
+    } else {
+        ASSERT(ValueLocation(use_inst.GetInst()) && "use_inst must already be defined");
+        const HostLoc location = *ValueLocation(use_inst.GetInst());
+        DefineValueImpl(code, def_inst, location);
     }
-
-    ASSERT(ValueLocation(use_inst.GetInst()) && "use_inst must already be defined");
-    const HostLoc location = *ValueLocation(use_inst.GetInst());
-    DefineValueImpl(code, def_inst, location);
 }
 
 void RegAlloc::Move(BlockOfCode& code, HostLoc to, HostLoc from) noexcept {

src/dynarmic/src/dynarmic/backend/x64/reg_alloc.h

@@ -46,61 +46,55 @@ public:
         return is_being_used_count == 0 && values.empty();
     }
     inline bool IsLastUse() const {
-        return is_being_used_count == 0 && current_references == 1 && accumulated_uses + 1 == total_uses;
-    }
-    inline void SetLastUse() noexcept {
-        ASSERT(IsLastUse());
-        is_set_last_use = true;
+        return is_being_used_count == 0 && current_references == 1 && size_t(accumulated_uses) + 1 == size_t(total_uses);
     }
     inline void ReadLock() noexcept {
-        ASSERT(size_t(is_being_used_count) + 1 < (std::numeric_limits<uint16_t>::max)());
+        ASSERT(size_t(is_being_used_count) + 1 < (std::numeric_limits<decltype(is_being_used_count)>::max)());
         ASSERT(!is_scratch);
         is_being_used_count++;
     }
     inline void WriteLock() noexcept {
-        ASSERT(size_t(is_being_used_count) + 1 < (std::numeric_limits<uint16_t>::max)());
         ASSERT(is_being_used_count == 0);
         is_being_used_count++;
         is_scratch = true;
     }
     inline void AddArgReference() noexcept {
-        ASSERT(size_t(current_references) + 1 < (std::numeric_limits<uint16_t>::max)());
-        current_references++;
-        ASSERT(accumulated_uses + current_references <= total_uses);
+        ASSERT(size_t(current_references) + 1 < (std::numeric_limits<decltype(current_references)>::max)());
+        ++current_references;
+        ASSERT(size_t(accumulated_uses) + current_references <= size_t(total_uses));
     }
     void ReleaseOne() noexcept;
     void ReleaseAll() noexcept;
-
+    constexpr size_t GetMaxBitWidth() const noexcept { return 1 << max_bit_width; }
+    void AddValue(HostLoc loc, IR::Inst* inst) noexcept;
     /// Checks if the given instruction is in our values set
     /// SAFETY: Const is casted away, irrelevant since this is only used for checking
-    inline bool ContainsValue(const IR::Inst* inst) const noexcept {
-        //return values.contains(const_cast<IR::Inst*>(inst));
+    [[nodiscard]] bool ContainsValue(const IR::Inst* inst) const noexcept {
         return std::find(values.begin(), values.end(), inst) != values.end();
     }
-    inline size_t GetMaxBitWidth() const noexcept {
-        return 1 << max_bit_width;
-    }
-    void AddValue(IR::Inst* inst) noexcept;
+#ifndef NDEBUG
     void EmitVerboseDebuggingOutput(BlockOfCode& code, size_t host_loc_index) const noexcept;
+#endif
 private:
-//non trivial
     boost::container::small_vector<IR::Inst*, 3> values; //24
-    // Block state
-    uint16_t total_uses = 0; //8
-    //sometimes zeroed
-    uint16_t accumulated_uses = 0; //8
+//non trivial
+    // Block state, the total amount of uses for this particular arg
+    uint16_t total_uses = 0; //2
+    // Sometimes zeroed, accumulated (non referenced) uses
+    uint16_t accumulated_uses = 0; //2
 //always zeroed
     // Current instruction state
-    uint16_t is_being_used_count = 0; //8
-    uint16_t current_references = 0; //8
-    // Value state
+    uint8_t current_references = 0; //1
+    uint8_t is_being_used_count = 0; //1
+    // Value state, count for LRU selection in registers
     uint8_t lru_counter : 2 = 0; //1
-    uint8_t max_bit_width : 4 = 0; //Valid values: log2(1,2,4,8,16,32,128) = (0, 1, 2, 3, 4, 5, 6)
+    // Log 2 of bit width, valid values: log2(1,2,4,8,16,32,128) = (0, 1, 2, 3, 4, 5, 6)
+    uint8_t max_bit_width : 4 = 0;
     bool is_scratch : 1 = false; //1
     bool is_set_last_use : 1 = false; //1
     friend class RegAlloc;
 };
-static_assert(sizeof(HostLocInfo) == 64);
+//static_assert(sizeof(HostLocInfo) == 64);
 
 struct Argument {
 public:
@@ -213,10 +207,12 @@ public:
     inline void AssertNoMoreUses() noexcept {
         ASSERT(std::all_of(hostloc_info.begin(), hostloc_info.end(), [](const auto& i) noexcept { return i.IsEmpty(); }));
     }
+#ifndef NDEBUG
     inline void EmitVerboseDebuggingOutput(BlockOfCode& code) noexcept {
         for (size_t i = 0; i < hostloc_info.size(); i++)
             hostloc_info[i].EmitVerboseDebuggingOutput(code, i);
     }
+#endif
 private:
     friend struct Argument;
 
@@ -238,11 +234,11 @@ private:
     HostLoc FindFreeSpill(bool is_xmm) const noexcept;
 
     inline HostLocInfo& LocInfo(const HostLoc loc) noexcept {
-        ASSERT(loc != HostLoc::RSP && loc != ABI_JIT_PTR);
+        DEBUG_ASSERT(loc != HostLoc::RSP && loc != ABI_JIT_PTR);
         return hostloc_info[size_t(loc)];
     }
     inline const HostLocInfo& LocInfo(const HostLoc loc) const noexcept {
-        ASSERT(loc != HostLoc::RSP && loc != ABI_JIT_PTR);
+        DEBUG_ASSERT(loc != HostLoc::RSP && loc != ABI_JIT_PTR);
         return hostloc_info[size_t(loc)];
     }
 
@@ -256,6 +252,6 @@ private:
     size_t reserved_stack_space = 0;
 };
 // Ensure a cache line (or less) is used, this is primordial
-static_assert(sizeof(boost::container::static_vector<HostLoc, 28>) == 40);
+static_assert(sizeof(boost::container::static_vector<HostLoc, 28>) < 64);
 
 }  // namespace Dynarmic::Backend::X64

src/dynarmic/src/dynarmic/ir/microinstruction.h

@@ -74,7 +74,6 @@ public:
     void SetName(unsigned value) { name = value; }
     unsigned GetName() const { return name; }
 
-private:
     void Use(const Value& value);
     void UndoUse(const Value& value);
 
@@ -87,6 +86,6 @@ private:
     unsigned name = 0; //4 (4)
     alignas(64) std::array<Value, max_arg_count> args; //16 * 4 = 64 (1 cache line)
 };
-static_assert(sizeof(Inst) == 128);
+//static_assert(sizeof(Inst) == 128);
 
 }  // namespace Dynarmic::IR

src/qt_common/discord/discord_impl.cpp

@@ -84,6 +84,7 @@ void DiscordImpl::UpdateGameStatus(bool use_default) {
     presence.smallImageKey = DEFAULT_DISCORD_IMAGE;
     presence.smallImageText = DEFAULT_DISCORD_TEXT;
     presence.state = game_title.c_str();
+    presence.status_display_type = DiscordStatusDisplayType_State;
     presence.details = "Currently in game";
     presence.startTimestamp = start_time;
     Discord_UpdatePresence(&presence);
@@ -123,6 +124,7 @@ void DiscordImpl::Update() {
     DiscordRichPresence presence{};
     presence.largeImageKey = DEFAULT_DISCORD_IMAGE;
     presence.largeImageText = DEFAULT_DISCORD_TEXT;
+    presence.status_display_type = DiscordStatusDisplayType_Name;
     presence.details = "Currently not in game";
     presence.startTimestamp = start_time;
     Discord_UpdatePresence(&presence);

src/shader_recompiler/backend/spirv/emit_spirv.cpp

@@ -4,6 +4,7 @@
 // SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
+#include <array>
 #include <span>
 #include <tuple>
 #include <type_traits>
@@ -24,6 +25,120 @@ template <class Func>
 struct FuncTraits {};
     thread_local std::unique_ptr<spvtools::Optimizer> thread_optimizer;
 
+static bool TryEmitCoalescedStorage(EmitContext& ctx, IR::Block::iterator& it,
+                                    IR::Block::iterator end) {
+    IR::Inst& inst = *it;
+    const auto opcode = inst.GetOpcode();
+
+    const bool is_u8 = opcode == IR::Opcode::WriteStorageU8 || opcode == IR::Opcode::WriteStorageS8;
+    const bool is_u16 = opcode == IR::Opcode::WriteStorageU16 || opcode == IR::Opcode::WriteStorageS16;
+
+    if ((is_u8 && ctx.profile.support_int8) || (is_u16 && ctx.profile.support_int16)) {
+        return false;
+    }
+
+    if (!is_u8 && !is_u16) {
+        return false;
+    }
+
+    if (!inst.Arg(0).IsImmediate() || !inst.Arg(1).IsImmediate()) {
+        return false;
+    }
+
+    const u32 binding = inst.Arg(0).U32();
+    const u32 base_offset = inst.Arg(1).U32();
+    const u32 base_word = base_offset / 4;
+    const u32 base_byte = base_offset % 4;
+
+    const u32 step_bytes = is_u8 ? 1u : 2u;
+    const u32 max_bytes = is_u8 ? 4u : 4u;
+
+    if (is_u16 && (base_offset % 2u) != 0) {
+        return false;
+    }
+
+    std::array<IR::Inst*, 4> grouped{};
+    grouped[0] = &inst;
+    u32 grouped_bytes = step_bytes;
+    auto look_ahead = it;
+
+    while (grouped_bytes < max_bytes) {
+        auto next = std::next(look_ahead);
+        if (next == end) {
+            break;
+        }
+        const auto next_opcode = next->GetOpcode();
+        const bool next_is_u8 = next_opcode == IR::Opcode::WriteStorageU8 || next_opcode == IR::Opcode::WriteStorageS8;
+        const bool next_is_u16 = next_opcode == IR::Opcode::WriteStorageU16 || next_opcode == IR::Opcode::WriteStorageS16;
+        if (next_is_u8 != is_u8 || next_is_u16 != is_u16) {
+            break;
+        }
+        if (!next->Arg(0).IsImmediate() || !next->Arg(1).IsImmediate()) {
+            break;
+        }
+        const u32 next_binding = next->Arg(0).U32();
+        const u32 next_offset = next->Arg(1).U32();
+        if (next_binding != binding) {
+            break;
+        }
+        if (next_offset / 4 != base_word) {
+            break;
+        }
+
+        const u32 expected_offset = base_offset + grouped_bytes;
+        if (next_offset != expected_offset) {
+            break;
+        }
+
+        grouped[grouped_bytes / step_bytes] = &*next;
+        grouped_bytes += step_bytes;
+        look_ahead = next;
+    }
+
+    const u32 count = grouped_bytes / step_bytes;
+    if (count <= 1) {
+        return false;
+    }
+
+    Id combined = ctx.u32_zero_value;
+    for (u32 i = 0; i < count; ++i) {
+        IR::Inst* current = grouped[i];
+        Id raw = ctx.Def(current->Arg(2));
+        if (is_u8) {
+            raw = ctx.OpBitwiseAnd(ctx.U32[1], raw, ctx.Const(0xFFu));
+        } else {
+            raw = ctx.OpBitwiseAnd(ctx.U32[1], raw, ctx.Const(0xFFFFu));
+        }
+        const u32 shift_bits = (is_u8 ? i * 8u : i * 16u);
+        Id shifted = raw;
+        if (shift_bits != 0) {
+            shifted = ctx.OpShiftLeftLogical(ctx.U32[1], raw, ctx.Const(shift_bits));
+        }
+        if (i == 0) {
+            combined = shifted;
+        } else {
+            combined = ctx.OpBitwiseOr(ctx.U32[1], combined, shifted);
+        }
+    }
+
+    const u32 bit_offset_val = is_u8 ? base_byte * 8u : ((base_offset / 2u) % 2u) * 16u;
+    const u32 bit_count_val = grouped_bytes * 8u;
+
+    const Id ssbo = ctx.ssbos[binding].U32;
+    const Id index = ctx.Const(base_word);
+    const Id pointer = ctx.OpAccessChain(ctx.storage_types.U32.element, ssbo, ctx.u32_zero_value, index);
+
+    if (bit_count_val == 32 && bit_offset_val == 0) {
+        ctx.OpStore(pointer, combined);
+    } else {
+        ctx.OpFunctionCall(ctx.TypeVoid(), ctx.write_storage_cas_loop_func, pointer, combined,
+                           ctx.Const(bit_offset_val), ctx.Const(bit_count_val));
+    }
+
+    it = look_ahead;
+    return true;
+}
+
     spvtools::Optimizer& GetThreadOptimizer() {
         if (!thread_optimizer) {
             thread_optimizer = std::make_unique<spvtools::Optimizer>(SPV_ENV_VULKAN_1_3);
@@ -140,8 +255,11 @@ void Traverse(EmitContext& ctx, IR::Program& program) {
             }
             current_block = node.data.block;
             ctx.AddLabel(label);
-            for (IR::Inst& inst : node.data.block->Instructions()) {
-                EmitInst(ctx, &inst);
+            for (auto it = node.data.block->begin(); it != node.data.block->end(); ++it) {
+                if (TryEmitCoalescedStorage(ctx, it, node.data.block->end())) {
+                    continue;
+                }
+                EmitInst(ctx, &*it);
             }
             break;
         }

src/shader_recompiler/backend/spirv/spirv_emit_context.cpp

@@ -315,9 +315,6 @@ void DefineSsbos(EmitContext& ctx, StorageTypeDefinition& type_def,
         ctx.Decorate(id, spv::Decoration::Binding, binding);
         ctx.Decorate(id, spv::Decoration::DescriptorSet, 0U);
         ctx.Name(id, fmt::format("ssbo{}", index));
-        if (!desc.is_written) {
-            ctx.Decorate(id, spv::Decoration::NonWritable);
-        }
         if (ctx.profile.supported_spirv >= 0x00010400) {
             ctx.interfaces.push_back(id);
         }
@@ -909,6 +906,9 @@ void EmitContext::DefineWriteStorageCasLoopFunction(const Info& info) {
     const Id continue_label{OpLabel()};
     const Id endloop_label{OpLabel()};
     const Id beginloop_label{OpLabel()};
+    const Id max_iterations{Const(16u)};
+    const Id iteration_counter{OpVariable(TypePointer(spv::StorageClass::Function, U32[1]),
+                                          spv::StorageClass::Function, u32_zero_value)};
     OpBranch(beginloop_label);
 
     AddLabel(beginloop_label);
@@ -916,11 +916,20 @@ void EmitContext::DefineWriteStorageCasLoopFunction(const Info& info) {
     OpBranch(body_label);
 
     AddLabel(body_label);
+        const Id current_iteration{OpLoad(U32[1], iteration_counter)};
+        const Id iteration_exceeded{OpUGreaterThanEqual(U1, current_iteration, max_iterations)};
+        const Id bailout_label{OpLabel()};
+        OpSelectionMerge(bailout_label, spv::SelectionControlMask::MaskNone);
+        OpBranchConditional(iteration_exceeded, bailout_label, bailout_label);
+        AddLabel(bailout_label);
+
     const Id expected_value{OpLoad(U32[1], pointer)};
     const Id desired_value{OpBitFieldInsert(U32[1], expected_value, value, bit_offset, bit_count)};
     const Id actual_value{OpAtomicCompareExchange(U32[1], pointer, scope_device, ordering_relaxed,
                                                   ordering_relaxed, desired_value, expected_value)};
     const Id store_successful{OpIEqual(U1, expected_value, actual_value)};
+    const Id next_iteration{OpIAdd(U32[1], current_iteration, Const(1u))};
+    OpStore(iteration_counter, next_iteration);
     OpBranchConditional(store_successful, endloop_label, continue_label);
 
     AddLabel(endloop_label);

src/video_core/buffer_cache/buffer_cache.h

@@ -407,6 +407,12 @@ void BufferCache<P>::SetComputeUniformBufferState(u32 mask,
 
 template <class P>
 void BufferCache<P>::UnbindGraphicsStorageBuffers(size_t stage) {
+    if constexpr (requires { runtime.ShouldLimitDynamicStorageBuffers(); }) {
+        if (runtime.ShouldLimitDynamicStorageBuffers()) {
+            channel_state->total_graphics_storage_buffers -=
+                static_cast<u32>(std::popcount(channel_state->enabled_storage_buffers[stage]));
+        }
+    }
     channel_state->enabled_storage_buffers[stage] = 0;
     channel_state->written_storage_buffers[stage] = 0;
 }
@@ -414,8 +420,26 @@ void BufferCache<P>::UnbindGraphicsStorageBuffers(size_t stage) {
 template <class P>
 bool BufferCache<P>::BindGraphicsStorageBuffer(size_t stage, size_t ssbo_index, u32 cbuf_index,
                                                u32 cbuf_offset, bool is_written) {
+    const bool already_enabled =
+        ((channel_state->enabled_storage_buffers[stage] >> ssbo_index) & 1U) != 0;
+    if constexpr (requires { runtime.ShouldLimitDynamicStorageBuffers(); }) {
+        if (runtime.ShouldLimitDynamicStorageBuffers() && !already_enabled) {
+            const u32 max_bindings = runtime.GetMaxDynamicStorageBuffers();
+            if (channel_state->total_graphics_storage_buffers >= max_bindings) {
+                LOG_WARNING(HW_GPU,
+                            "Skipping graphics storage buffer {} due to driver limit {}",
+                            ssbo_index, max_bindings);
+                return false;
+            }
+        }
+    }
     channel_state->enabled_storage_buffers[stage] |= 1U << ssbo_index;
     channel_state->written_storage_buffers[stage] |= (is_written ? 1U : 0U) << ssbo_index;
+    if constexpr (requires { runtime.ShouldLimitDynamicStorageBuffers(); }) {
+        if (runtime.ShouldLimitDynamicStorageBuffers() && !already_enabled) {
+            ++channel_state->total_graphics_storage_buffers;
+        }
+    }
 
     const auto& cbufs = maxwell3d->state.shader_stages[stage];
     const GPUVAddr ssbo_addr = cbufs.const_buffers[cbuf_index].address + cbuf_offset;
@@ -446,6 +470,12 @@ void BufferCache<P>::BindGraphicsTextureBuffer(size_t stage, size_t tbo_index, G
 
 template <class P>
 void BufferCache<P>::UnbindComputeStorageBuffers() {
+    if constexpr (requires { runtime.ShouldLimitDynamicStorageBuffers(); }) {
+        if (runtime.ShouldLimitDynamicStorageBuffers()) {
+            channel_state->total_compute_storage_buffers -=
+                static_cast<u32>(std::popcount(channel_state->enabled_compute_storage_buffers));
+        }
+    }
     channel_state->enabled_compute_storage_buffers = 0;
     channel_state->written_compute_storage_buffers = 0;
     channel_state->image_compute_texture_buffers = 0;
@@ -459,8 +489,26 @@ void BufferCache<P>::BindComputeStorageBuffer(size_t ssbo_index, u32 cbuf_index,
                   ssbo_index);
         return;
     }
+    const bool already_enabled =
+        ((channel_state->enabled_compute_storage_buffers >> ssbo_index) & 1U) != 0;
+    if constexpr (requires { runtime.ShouldLimitDynamicStorageBuffers(); }) {
+        if (runtime.ShouldLimitDynamicStorageBuffers() && !already_enabled) {
+            const u32 max_bindings = runtime.GetMaxDynamicStorageBuffers();
+            if (channel_state->total_compute_storage_buffers >= max_bindings) {
+                LOG_WARNING(HW_GPU,
+                            "Skipping compute storage buffer {} due to driver limit {}",
+                            ssbo_index, max_bindings);
+                return;
+            }
+        }
+    }
     channel_state->enabled_compute_storage_buffers |= 1U << ssbo_index;
     channel_state->written_compute_storage_buffers |= (is_written ? 1U : 0U) << ssbo_index;
+    if constexpr (requires { runtime.ShouldLimitDynamicStorageBuffers(); }) {
+        if (runtime.ShouldLimitDynamicStorageBuffers() && !already_enabled) {
+            ++channel_state->total_compute_storage_buffers;
+        }
+    }
 
     const auto& launch_desc = kepler_compute->launch_description;
     if (((launch_desc.const_buffer_enable_mask >> cbuf_index) & 1) == 0) {
@@ -793,9 +841,23 @@ void BufferCache<P>::BindHostGraphicsUniformBuffer(size_t stage, u32 index, u32
     const u32 size = (std::min)(binding.size, (*channel_state->uniform_buffer_sizes)[stage][index]);
     Buffer& buffer = slot_buffers[binding.buffer_id];
     TouchBuffer(buffer, binding.buffer_id);
-    const bool use_fast_buffer = binding.buffer_id != NULL_BUFFER_ID &&
-                                 size <= channel_state->uniform_buffer_skip_cache_size &&
-                                 !memory_tracker.IsRegionGpuModified(device_addr, size);
+    const bool has_host_buffer = binding.buffer_id != NULL_BUFFER_ID;
+    const u32 offset = has_host_buffer ? buffer.Offset(device_addr) : 0;
+    const bool needs_alignment_stream = [&]() {
+        if constexpr (IS_OPENGL) {
+            return false;
+        } else {
+            if (!has_host_buffer) {
+                return false;
+            }
+            const u32 alignment = runtime.GetUniformBufferAlignment();
+            return alignment > 1 && (offset % alignment) != 0;
+        }
+    }();
+    const bool use_fast_buffer = needs_alignment_stream ||
+                                 (has_host_buffer &&
+                                  size <= channel_state->uniform_buffer_skip_cache_size &&
+                                  !memory_tracker.IsRegionGpuModified(device_addr, size));
     if (use_fast_buffer) {
         if constexpr (IS_OPENGL) {
             if (runtime.HasFastBufferSubData()) {
@@ -834,7 +896,6 @@ void BufferCache<P>::BindHostGraphicsUniformBuffer(size_t stage, u32 index, u32
     if (!needs_bind) {
         return;
     }
-    const u32 offset = buffer.Offset(device_addr);
     if constexpr (IS_OPENGL) {
         // Mark the index as dirty if offset doesn't match
         const bool is_copy_bind = offset != 0 && !runtime.SupportsNonZeroUniformOffset();
@@ -951,9 +1012,30 @@ void BufferCache<P>::BindHostComputeUniformBuffers() {
         TouchBuffer(buffer, binding.buffer_id);
         const u32 size =
             (std::min)(binding.size, (*channel_state->compute_uniform_buffer_sizes)[index]);
+        const bool has_host_buffer = binding.buffer_id != NULL_BUFFER_ID;
+        const u32 offset = has_host_buffer ? buffer.Offset(binding.device_addr) : 0;
+        const bool needs_alignment_stream = [&]() {
+            if constexpr (IS_OPENGL) {
+                return false;
+            } else {
+                if (!has_host_buffer) {
+                    return false;
+                }
+                const u32 alignment = runtime.GetUniformBufferAlignment();
+                return alignment > 1 && (offset % alignment) != 0;
+            }
+        }();
+        if constexpr (!IS_OPENGL) {
+            if (needs_alignment_stream) {
+                const std::span<u8> span =
+                    runtime.BindMappedUniformBuffer(0, binding_index, size);
+                device_memory.ReadBlockUnsafe(binding.device_addr, span.data(), size);
+                return;
+            }
+        }
+
         SynchronizeBuffer(buffer, binding.device_addr, size);
 
-        const u32 offset = buffer.Offset(binding.device_addr);
         buffer.MarkUsage(offset, size);
         if constexpr (NEEDS_BIND_UNIFORM_INDEX) {
             runtime.BindComputeUniformBuffer(binding_index, buffer, offset, size);

src/video_core/buffer_cache/buffer_cache_base.h

@@ -8,6 +8,7 @@
 
 #include <algorithm>
 #include <array>
+#include <bit>
 #include <functional>
 #include <memory>
 #include <mutex>
@@ -132,6 +133,9 @@ public:
     u32 enabled_compute_storage_buffers = 0;
     u32 written_compute_storage_buffers = 0;
 
+    u32 total_graphics_storage_buffers = 0;
+    u32 total_compute_storage_buffers = 0;
+
     std::array<u32, NUM_STAGES> enabled_texture_buffers{};
     std::array<u32, NUM_STAGES> written_texture_buffers{};
     std::array<u32, NUM_STAGES> image_texture_buffers{};

src/video_core/host1x/vic.h

@@ -605,20 +605,20 @@ public:
         SetOutputSurfaceChromaUnusedOffset = offsetof(VicRegisters, output_surface.chroma_v)
     };
 
-    explicit Vic(Host1x& host1x, s32 id, u32 syncpt, FrameQueue& frame_queue);
-    ~Vic();
+    explicit Vic(Host1x& host1x, s32 id, u32 syncpt, FrameQueue& frame_queue) noexcept;
+    ~Vic() noexcept;
 
     /// Write to the device state.
-    void ProcessMethod(u32 method, u32 arg) override;
+    void ProcessMethod(u32 method, u32 arg) noexcept override;
 
 private:
-    void Execute();
-    void Blend(const ConfigStruct& config, const SlotStruct& slot);
-    void ReadProgressiveY8__V8U8_N420(const SlotStruct& slot, std::span<const PlaneOffsets> offsets, std::shared_ptr<const FFmpeg::Frame> frame, bool planar, bool interlaced);
-    void ReadInterlacedY8__V8U8_N420(const SlotStruct& slot, std::span<const PlaneOffsets> offsets, std::shared_ptr<const FFmpeg::Frame> frame, bool planar, bool top_field);
-    void ReadY8__V8U8_N420(const SlotStruct& slot, std::span<const PlaneOffsets> offsets, std::shared_ptr<const FFmpeg::Frame> frame, bool planar);
-    void WriteY8__V8U8_N420(const OutputSurfaceConfig& output_surface_config);
-    void WriteABGR(const OutputSurfaceConfig& output_surface_config, VideoPixelFormat format);
+    void Execute() noexcept;
+    void Blend(const ConfigStruct& config, const SlotStruct& slot, VideoPixelFormat format) noexcept;
+    void ReadProgressiveY8__V8U8_N420(const SlotStruct& slot, std::span<const PlaneOffsets> offsets, std::shared_ptr<const FFmpeg::Frame> frame, bool planar, bool interlaced) noexcept;
+    void ReadInterlacedY8__V8U8_N420(const SlotStruct& slot, std::span<const PlaneOffsets> offsets, std::shared_ptr<const FFmpeg::Frame> frame, bool planar, bool top_field) noexcept;
+    void ReadY8__V8U8_N420(const SlotStruct& slot, std::span<const PlaneOffsets> offsets, std::shared_ptr<const FFmpeg::Frame> frame, bool planar) noexcept;
+    void WriteY8__V8U8_N420(const OutputSurfaceConfig& output_surface_config) noexcept;
+    void WriteABGR(const OutputSurfaceConfig& output_surface_config, VideoPixelFormat format) noexcept;
 
     s32 id;
     s32 nvdec_id{-1};
@@ -627,11 +627,11 @@ private:
     VicRegisters regs{};
     FrameQueue& frame_queue;
 
+    Common::ScratchBuffer<u8> swizzle_scratch;
     Common::ScratchBuffer<Pixel> output_surface;
     Common::ScratchBuffer<Pixel> slot_surface;
     Common::ScratchBuffer<u8> luma_scratch;
     Common::ScratchBuffer<u8> chroma_scratch;
-    Common::ScratchBuffer<u8> swizzle_scratch;
 };
 
 } // namespace Tegra::Host1x

src/video_core/renderer_opengl/gl_buffer_cache.h

@@ -198,6 +198,10 @@ public:
         return device.CanReportMemoryUsage();
     }
 
+    u32 GetUniformBufferAlignment() const {
+        return static_cast<u32>(device.GetUniformBufferAlignment());
+    }
+
     u32 GetStorageBufferAlignment() const {
         return static_cast<u32>(device.GetShaderStorageBufferAlignment());
     }

src/video_core/renderer_vulkan/vk_buffer_cache.cpp

@@ -333,16 +333,18 @@ BufferCacheRuntime::BufferCacheRuntime(const Device& device_, MemoryAllocator& m
       staging_pool{staging_pool_}, guest_descriptor_queue{guest_descriptor_queue_},
       quad_index_pass(device, scheduler, descriptor_pool, staging_pool,
                       compute_pass_descriptor_queue) {
+    const VkDriverIdKHR driver_id = device.GetDriverID();
+    limit_dynamic_storage_buffers = driver_id == VK_DRIVER_ID_QUALCOMM_PROPRIETARY ||
+                                    driver_id == VK_DRIVER_ID_MESA_TURNIP ||
+                                    driver_id == VK_DRIVER_ID_ARM_PROPRIETARY;
+    if (limit_dynamic_storage_buffers) {
+        max_dynamic_storage_buffers = device.GetMaxDescriptorSetStorageBuffersDynamic();
+    }
     if (device.GetDriverID() != VK_DRIVER_ID_QUALCOMM_PROPRIETARY) {
         // TODO: FixMe: Uint8Pass compute shader does not build on some Qualcomm drivers.
         uint8_pass = std::make_unique<Uint8Pass>(device, scheduler, descriptor_pool, staging_pool,
                                                  compute_pass_descriptor_queue);
     }
-    const u32 ubo_align = static_cast<u32>(
-            device.GetUniformBufferAlignment() //check if the device has it
-    );
-    // add the ability to change the size in settings in future
-    uniform_ring.Init(memory_allocator, 8 * 1024 * 1024 /* 8 MiB */, ubo_align ? ubo_align : 256);
     quad_array_index_buffer = std::make_shared<QuadArrayIndexBuffer>(device_, memory_allocator_,
                                                                      scheduler_, staging_pool_);
     quad_strip_index_buffer = std::make_shared<QuadStripIndexBuffer>(device_, memory_allocator_,
@@ -361,41 +363,6 @@ void BufferCacheRuntime::FreeDeferredStagingBuffer(StagingBufferRef& ref) {
     staging_pool.FreeDeferred(ref);
 }
 
-void BufferCacheRuntime::UniformRing::Init(MemoryAllocator& alloc, u64 bytes, u32 alignment)
-{
-    for (size_t i = 0; i < NUM_FRAMES; ++i) {
-        VkBufferCreateInfo ci{
-            .sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO,
-            .pNext = nullptr,
-            .flags = 0,
-            .size  = bytes,
-            .usage = VK_BUFFER_USAGE_UNIFORM_BUFFER_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT,
-            .sharingMode = VK_SHARING_MODE_EXCLUSIVE,
-            .queueFamilyIndexCount = 0,
-            .pQueueFamilyIndices = nullptr,
-        };
-        buffers[i] = alloc.CreateBuffer(ci, MemoryUsage::Upload);
-        mapped[i] = buffers[i].Mapped().data();
-    }
-    size   = bytes;
-    align  = alignment ? alignment : 256;
-    head   = 0;
-    current_frame = 0;
-}
-
-std::span<u8> BufferCacheRuntime::UniformRing::Alloc(u32 bytes, u32& out_offset) {
-    const u64 aligned = Common::AlignUp(head, static_cast<u64>(align));
-    u64 end = aligned + bytes;
-
-    if (end > size) {
-       return {}; // Fallback to staging pool
-    }
-
-    out_offset = static_cast<u32>(aligned);
-    head = end;
-    return {mapped[current_frame] + out_offset, bytes};
-}
-
 u64 BufferCacheRuntime::GetDeviceLocalMemory() const {
     return device.GetDeviceLocalMemory();
 }
@@ -408,6 +375,10 @@ bool BufferCacheRuntime::CanReportMemoryUsage() const {
     return device.CanReportMemoryUsage();
 }
 
+u32 BufferCacheRuntime::GetUniformBufferAlignment() const {
+    return static_cast<u32>(device.GetUniformBufferAlignment());
+}
+
 u32 BufferCacheRuntime::GetStorageBufferAlignment() const {
     return static_cast<u32>(device.GetStorageBufferAlignment());
 }
@@ -416,7 +387,6 @@ void BufferCacheRuntime::TickFrame(Common::SlotVector<Buffer>& slot_buffers) noe
     for (auto it = slot_buffers.begin(); it != slot_buffers.end(); it++) {
         it->ResetUsageTracking();
     }
-    uniform_ring.BeginFrame();
 }
 
 void BufferCacheRuntime::Finish() {

src/video_core/renderer_vulkan/vk_buffer_cache.h

@@ -6,6 +6,8 @@
 
 #pragma once
 
+#include <limits>
+
 #include "video_core/buffer_cache/buffer_cache_base.h"
 #include "video_core/buffer_cache/memory_tracker_base.h"
 #include "video_core/buffer_cache/usage_tracker.h"
@@ -94,6 +96,8 @@ public:
 
     bool CanReportMemoryUsage() const;
 
+    u32 GetUniformBufferAlignment() const;
+
     u32 GetStorageBufferAlignment() const;
 
     [[nodiscard]] StagingBufferRef UploadStagingBuffer(size_t size);
@@ -127,15 +131,9 @@ public:
 
     void BindTransformFeedbackBuffers(VideoCommon::HostBindings<Buffer>& bindings);
 
-    std::span<u8> BindMappedUniformBuffer([[maybe_unused]] size_t /*stage*/,
-                                          [[maybe_unused]] u32 /*binding_index*/,
+    std::span<u8> BindMappedUniformBuffer([[maybe_unused]] size_t stage,
+                                          [[maybe_unused]] u32 binding_index,
                                           u32 size) {
-        u32 offset = 0;
-        if (auto span = uniform_ring.Alloc(size, offset); !span.empty()) {
-            BindBuffer(*uniform_ring.buffers[uniform_ring.current_frame], offset, size);
-            return span;
-        }
-        // Fallback for giant requests
         const StagingBufferRef ref = staging_pool.Request(size, MemoryUsage::Upload);
         BindBuffer(ref.buffer, static_cast<u32>(ref.offset), size);
         return ref.mapped_span;
@@ -155,6 +153,14 @@ public:
         guest_descriptor_queue.AddTexelBuffer(buffer.View(offset, size, format));
     }
 
+    bool ShouldLimitDynamicStorageBuffers() const {
+        return limit_dynamic_storage_buffers;
+    }
+
+    u32 GetMaxDynamicStorageBuffers() const {
+        return max_dynamic_storage_buffers;
+    }
+
 private:
     void BindBuffer(VkBuffer buffer, u32 offset, u32 size) {
         guest_descriptor_queue.AddBuffer(buffer, offset, size);
@@ -163,24 +169,6 @@ private:
     void ReserveNullBuffer();
     vk::Buffer CreateNullBuffer();
 
-    struct UniformRing {
-        static constexpr size_t NUM_FRAMES = 3;
-        std::array<vk::Buffer, NUM_FRAMES> buffers{};
-        std::array<u8*, NUM_FRAMES> mapped{};
-        u64 size = 0;
-        u64 head = 0;
-        u32 align = 256;
-        size_t current_frame = 0;
-
-        void Init(MemoryAllocator& alloc, u64 bytes, u32 alignment);
-        void BeginFrame() {
-            current_frame = (current_frame + 1) % NUM_FRAMES;
-            head = 0;
-        }
-        std::span<u8> Alloc(u32 bytes, u32& out_offset);
-    };
-    UniformRing uniform_ring;
-
     const Device& device;
     MemoryAllocator& memory_allocator;
     Scheduler& scheduler;
@@ -194,6 +182,9 @@ private:
 
     std::unique_ptr<Uint8Pass> uint8_pass;
     QuadIndexedPass quad_index_pass;
+
+    bool limit_dynamic_storage_buffers = false;
+    u32 max_dynamic_storage_buffers = std::numeric_limits<u32>::max();
 };
 
 struct BufferCacheParams {

src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp

@@ -816,23 +816,12 @@ void GraphicsPipeline::MakePipeline(VkRenderPass render_pass) {
         .blendConstants = {}
     };
     static_vector<VkDynamicState, 34> dynamic_states{
-        VK_DYNAMIC_STATE_VIEWPORT,
-        VK_DYNAMIC_STATE_SCISSOR,
-        VK_DYNAMIC_STATE_DEPTH_BIAS,
+        VK_DYNAMIC_STATE_VIEWPORT,           VK_DYNAMIC_STATE_SCISSOR,
+        VK_DYNAMIC_STATE_DEPTH_BIAS,         VK_DYNAMIC_STATE_BLEND_CONSTANTS,
+        VK_DYNAMIC_STATE_DEPTH_BOUNDS,       VK_DYNAMIC_STATE_STENCIL_COMPARE_MASK,
+        VK_DYNAMIC_STATE_STENCIL_WRITE_MASK, VK_DYNAMIC_STATE_STENCIL_REFERENCE,
         VK_DYNAMIC_STATE_LINE_WIDTH,
     };
-
-    if (device.UsesAdvancedCoreDynamicState()) {
-        static constexpr std::array core_dynamic_states{
-            VK_DYNAMIC_STATE_BLEND_CONSTANTS,
-            VK_DYNAMIC_STATE_DEPTH_BOUNDS,
-            VK_DYNAMIC_STATE_STENCIL_COMPARE_MASK,
-            VK_DYNAMIC_STATE_STENCIL_WRITE_MASK,
-            VK_DYNAMIC_STATE_STENCIL_REFERENCE,
-        };
-        dynamic_states.insert(dynamic_states.end(), core_dynamic_states.begin(),
-                              core_dynamic_states.end());
-    }
     if (key.state.extended_dynamic_state) {
         static constexpr std::array extended{
             VK_DYNAMIC_STATE_CULL_MODE_EXT,

src/video_core/renderer_vulkan/vk_pipeline_cache.cpp

@@ -278,8 +278,9 @@ size_t GetTotalPipelineWorkers() {
     const size_t max_core_threads =
         std::max<size_t>(static_cast<size_t>(std::thread::hardware_concurrency()), 2ULL) - 1ULL;
 #ifdef ANDROID
-    // Leave at least a few cores free in android
-    constexpr size_t free_cores = 3ULL;
+    // Leave at least one core free on Android. Previously we reserved two, but
+    // shipping builds benefit from one extra compilation worker.
+    constexpr size_t free_cores = 1ULL;
     if (max_core_threads <= free_cores) {
         return 1ULL;
     }
@@ -797,6 +798,19 @@ std::unique_ptr<ComputePipeline> PipelineCache::CreateComputePipeline(
     }
 
     auto program{TranslateProgram(pools.inst, pools.block, env, cfg, host_info)};
+    const VkDriverIdKHR driver_id = device.GetDriverID();
+    const bool needs_shared_mem_clamp =
+        driver_id == VK_DRIVER_ID_QUALCOMM_PROPRIETARY ||
+        driver_id == VK_DRIVER_ID_ARM_PROPRIETARY;
+    const u32 max_shared_memory = device.GetMaxComputeSharedMemorySize();
+    if (needs_shared_mem_clamp && program.shared_memory_size > max_shared_memory) {
+        LOG_WARNING(Render_Vulkan,
+                    "Compute shader 0x{:016x} requests {}KB shared memory but device max is {}KB - clamping",
+                    key.unique_hash,
+                    program.shared_memory_size / 1024,
+                    max_shared_memory / 1024);
+        program.shared_memory_size = max_shared_memory;
+    }
     const std::vector<u32> code{EmitSPIRV(profile, program, this->optimize_spirv_output)};
     device.SaveShader(code);
     vk::ShaderModule spv_module{BuildShader(device, code)};

src/video_core/renderer_vulkan/vk_query_cache.cpp

@@ -156,19 +156,6 @@ public:
 
         ReserveHostQuery();
 
-        // Ensure outside render pass
-        scheduler.RequestOutsideRenderPassOperationContext();
-
-        // Reset query pool outside render pass
-        scheduler.Record([query_pool = current_query_pool,
-                                 query_index = current_bank_slot](vk::CommandBuffer cmdbuf) {
-            cmdbuf.ResetQueryPool(query_pool, static_cast<u32>(query_index), 1);
-        });
-
-        // Manually restart the render pass (required for vkCmdClearAttachments, etc.)
-        scheduler.RequestRenderpass(texture_cache.GetFramebuffer());
-
-        // Begin query inside the newly started render pass
         scheduler.Record([query_pool = current_query_pool,
                                  query_index = current_bank_slot](vk::CommandBuffer cmdbuf) {
             const bool use_precise = Settings::IsGPULevelHigh();

src/video_core/renderer_vulkan/vk_rasterizer.cpp

@@ -1233,9 +1233,6 @@ void RasterizerVulkan::UpdateBlendConstants(Tegra::Engines::Maxwell3D::Regs& reg
     if (!state_tracker.TouchBlendConstants()) {
         return;
     }
-    if (!device.UsesAdvancedCoreDynamicState()) {
-        return;
-    }
     const std::array blend_color = {regs.blend_color.r, regs.blend_color.g, regs.blend_color.b,
                                     regs.blend_color.a};
     scheduler.Record(
@@ -1246,9 +1243,6 @@ void RasterizerVulkan::UpdateDepthBounds(Tegra::Engines::Maxwell3D::Regs& regs)
     if (!state_tracker.TouchDepthBounds()) {
         return;
     }
-    if (!device.UsesAdvancedCoreDynamicState() || !device.IsDepthBoundsSupported()) {
-        return;
-    }
     scheduler.Record([min = regs.depth_bounds[0], max = regs.depth_bounds[1]](
                          vk::CommandBuffer cmdbuf) { cmdbuf.SetDepthBounds(min, max); });
 }
@@ -1257,10 +1251,6 @@ void RasterizerVulkan::UpdateStencilFaces(Tegra::Engines::Maxwell3D::Regs& regs)
     if (!state_tracker.TouchStencilProperties()) {
         return;
     }
-    if (!device.UsesAdvancedCoreDynamicState()) {
-        state_tracker.ClearStencilReset();
-        return;
-    }
     bool update_references = state_tracker.TouchStencilReference();
     bool update_write_mask = state_tracker.TouchStencilWriteMask();
     bool update_compare_masks = state_tracker.TouchStencilCompare();

src/video_core/renderer_vulkan/vk_scheduler.cpp

@@ -302,8 +302,6 @@ void Scheduler::EndRenderPass()
                        images = renderpass_images,
                        ranges = renderpass_image_ranges](vk::CommandBuffer cmdbuf) {
             std::array<VkImageMemoryBarrier, 9> barriers;
-            VkPipelineStageFlags src_stages = 0;
-
             for (size_t i = 0; i < num_images; ++i) {
                 const VkImageSubresourceRange& range = ranges[i];
                 const bool is_color = (range.aspectMask & VK_IMAGE_ASPECT_COLOR_BIT) != 0;
@@ -312,20 +310,14 @@ void Scheduler::EndRenderPass()
                                                  | VK_IMAGE_ASPECT_STENCIL_BIT)) !=0;
 
                 VkAccessFlags src_access = 0;
-                VkPipelineStageFlags this_stage = 0;
 
-                if (is_color) {
+                if (is_color)
                     src_access |= VK_ACCESS_COLOR_ATTACHMENT_WRITE_BIT;
-                    this_stage |= VK_PIPELINE_STAGE_COLOR_ATTACHMENT_OUTPUT_BIT;
-                }
-
-                if (is_depth_stencil) {
+                else if (is_depth_stencil)
                     src_access |= VK_ACCESS_DEPTH_STENCIL_ATTACHMENT_WRITE_BIT;
-                    this_stage |= VK_PIPELINE_STAGE_EARLY_FRAGMENT_TESTS_BIT
-                                  | VK_PIPELINE_STAGE_LATE_FRAGMENT_TESTS_BIT;
-                }
-
-                src_stages |= this_stage;
+                else
+                    src_access |= VK_ACCESS_COLOR_ATTACHMENT_WRITE_BIT
+                                  | VK_ACCESS_DEPTH_STENCIL_ATTACHMENT_WRITE_BIT;
 
                 barriers[i] = VkImageMemoryBarrier{
                         .sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER,
@@ -344,15 +336,10 @@ void Scheduler::EndRenderPass()
                         .subresourceRange = range,
                 };
             }
-
-            // Graft: ensure explicit fragment tests + color output stages are always synchronized (AMD/Windows fix)
-            src_stages |= VK_PIPELINE_STAGE_EARLY_FRAGMENT_TESTS_BIT |
-                    VK_PIPELINE_STAGE_LATE_FRAGMENT_TESTS_BIT |
-                    VK_PIPELINE_STAGE_COLOR_ATTACHMENT_OUTPUT_BIT;
-
             cmdbuf.EndRenderPass();
-
-            cmdbuf.PipelineBarrier(src_stages,
+            cmdbuf.PipelineBarrier(VK_PIPELINE_STAGE_EARLY_FRAGMENT_TESTS_BIT |
+                                   VK_PIPELINE_STAGE_LATE_FRAGMENT_TESTS_BIT |
+                                   VK_PIPELINE_STAGE_COLOR_ATTACHMENT_OUTPUT_BIT,
                                    VK_PIPELINE_STAGE_ALL_COMMANDS_BIT,
                                    0,
                                    nullptr,

src/video_core/renderer_vulkan/vk_texture_cache.cpp

@@ -1527,6 +1527,10 @@ bool TextureCacheRuntime::CanReportMemoryUsage() const {
     return device.CanReportMemoryUsage();
 }
 
+std::optional<size_t> TextureCacheRuntime::GetSamplerHeapBudget() const {
+    return device.GetSamplerHeapBudget();
+}
+
 void TextureCacheRuntime::TickFrame() {}
 
 Image::Image(TextureCacheRuntime& runtime_, const ImageInfo& info_, GPUVAddr gpu_addr_,

src/video_core/renderer_vulkan/vk_texture_cache.h

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2019 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-3.0-or-later
 
@@ -62,6 +65,8 @@ public:
 
     bool CanReportMemoryUsage() const;
 
+    std::optional<size_t> GetSamplerHeapBudget() const;
+
     void BlitImage(Framebuffer* dst_framebuffer, ImageView& dst, ImageView& src,
                    const Region2D& dst_region, const Region2D& src_region,
                    Tegra::Engines::Fermi2D::Filter filter,

src/video_core/texture_cache/texture_cache.h

@@ -6,6 +6,8 @@
 
 #pragma once
 
+#include <limits>
+#include <optional>
 #include <unordered_set>
 #include <boost/container/small_vector.hpp>
 
@@ -1736,11 +1738,89 @@ SamplerId TextureCache<P>::FindSampler(const TSCEntry& config) {
     }
     const auto [pair, is_new] = channel_state->samplers.try_emplace(config);
     if (is_new) {
+        EnforceSamplerBudget();
         pair->second = slot_samplers.insert(runtime, config);
     }
     return pair->second;
 }
 
+template <class P>
+std::optional<size_t> TextureCache<P>::QuerySamplerBudget() const {
+    if constexpr (requires { runtime.GetSamplerHeapBudget(); }) {
+        return runtime.GetSamplerHeapBudget();
+    } else {
+        return std::nullopt;
+    }
+}
+
+template <class P>
+void TextureCache<P>::EnforceSamplerBudget() {
+    const auto budget = QuerySamplerBudget();
+    if (!budget) {
+        return;
+    }
+    if (slot_samplers.size() < *budget) {
+        return;
+    }
+    if (!channel_state) {
+        return;
+    }
+    if (last_sampler_gc_frame == frame_tick) {
+        return;
+    }
+    last_sampler_gc_frame = frame_tick;
+    TrimInactiveSamplers(*budget);
+}
+
+template <class P>
+void TextureCache<P>::TrimInactiveSamplers(size_t budget) {
+    if (channel_state->samplers.empty()) {
+        return;
+    }
+    static constexpr size_t SAMPLER_GC_SLACK = 1024;
+    auto mark_active = [](auto& set, SamplerId id) {
+        if (!id || id == CORRUPT_ID || id == NULL_SAMPLER_ID) {
+            return;
+        }
+        set.insert(id);
+    };
+    std::unordered_set<SamplerId> active;
+    active.reserve(channel_state->graphics_sampler_ids.size() +
+                   channel_state->compute_sampler_ids.size());
+    for (const SamplerId id : channel_state->graphics_sampler_ids) {
+        mark_active(active, id);
+    }
+    for (const SamplerId id : channel_state->compute_sampler_ids) {
+        mark_active(active, id);
+    }
+
+    size_t removed = 0;
+    auto& sampler_map = channel_state->samplers;
+    for (auto it = sampler_map.begin(); it != sampler_map.end();) {
+        const SamplerId sampler_id = it->second;
+        if (!sampler_id || sampler_id == CORRUPT_ID) {
+            it = sampler_map.erase(it);
+            continue;
+        }
+        if (active.find(sampler_id) != active.end()) {
+            ++it;
+            continue;
+        }
+        slot_samplers.erase(sampler_id);
+        it = sampler_map.erase(it);
+        ++removed;
+        if (slot_samplers.size() + SAMPLER_GC_SLACK <= budget) {
+            break;
+        }
+    }
+
+    if (removed != 0) {
+        LOG_WARNING(HW_GPU,
+                    "Sampler cache exceeded {} entries on this driver; reclaimed {} inactive samplers",
+                    budget, removed);
+    }
+}
+
 template <class P>
 ImageViewId TextureCache<P>::FindColorBuffer(size_t index) {
     const auto& regs = maxwell3d->regs;

src/video_core/texture_cache/texture_cache_base.h

@@ -429,6 +429,9 @@ private:
 
     void QueueAsyncDecode(Image& image, ImageId image_id);
     void TickAsyncDecode();
+    void EnforceSamplerBudget();
+    void TrimInactiveSamplers(size_t budget);
+    std::optional<size_t> QuerySamplerBudget() const;
 
     Runtime& runtime;
 
@@ -500,6 +503,7 @@ private:
 
     u64 modification_tick = 0;
     u64 frame_tick = 0;
+    u64 last_sampler_gc_frame = (std::numeric_limits<u64>::max)();
 
     Common::ThreadWorker texture_decode_worker{1, "TextureDecoder"};
     std::vector<std::unique_ptr<AsyncDecodeContext>> async_decodes;

src/video_core/textures/decoders.cpp

@@ -18,8 +18,7 @@
 
 namespace Tegra::Texture {
 namespace {
-template <u32 mask>
-constexpr u32 pdep(u32 value) {
+constexpr u32 pdep(u32 mask, u32 value) {
     u32 result = 0;
     u32 m = mask;
     for (u32 bit = 1; m; bit += bit) {
@@ -30,15 +29,12 @@ constexpr u32 pdep(u32 value) {
     return result;
 }
 
-template <u32 mask, u32 incr_amount>
-void incrpdep(u32& value) {
-    static constexpr u32 swizzled_incr = pdep<mask>(incr_amount);
+void incrpdep(u32 mask, u32 incr_amount, u32& value) {
+    u32 swizzled_incr = pdep(mask, incr_amount);
     value = ((value | ~mask) + swizzled_incr) & mask;
 }
 
-template <bool TO_LINEAR, u32 BYTES_PER_PIXEL>
-void SwizzleImpl(std::span<u8> output, std::span<const u8> input, u32 width, u32 height, u32 depth,
-                 u32 block_height, u32 block_depth, u32 stride) {
+void SwizzleImpl(bool to_linear, u32 bytes_per_pixel, std::span<u8> output, std::span<const u8> input, u32 width, u32 height, u32 depth, u32 block_height, u32 block_depth, u32 stride) {
     // The origin of the transformation can be configured here, leave it as zero as the current API
     // doesn't expose it.
     static constexpr u32 origin_x = 0;
@@ -46,13 +42,12 @@ void SwizzleImpl(std::span<u8> output, std::span<const u8> input, u32 width, u32
     static constexpr u32 origin_z = 0;
 
     // We can configure here a custom pitch
-    // As it's not exposed 'width * BYTES_PER_PIXEL' will be the expected pitch.
-    const u32 pitch = width * BYTES_PER_PIXEL;
+    // As it's not exposed 'width * bytes_per_pixel' will be the expected pitch.
+    const u32 pitch = width * bytes_per_pixel;
 
     const u32 gobs_in_x = Common::DivCeilLog2(stride, GOB_SIZE_X_SHIFT);
     const u32 block_size = gobs_in_x << (GOB_SIZE_SHIFT + block_height + block_depth);
-    const u32 slice_size =
-        Common::DivCeilLog2(height, block_height + GOB_SIZE_Y_SHIFT) * block_size;
+    const u32 slice_size = Common::DivCeilLog2(height, block_height + GOB_SIZE_Y_SHIFT) * block_size;
 
     const u32 block_height_mask = (1U << block_height) - 1;
     const u32 block_depth_mask = (1U << block_depth) - 1;
@@ -64,50 +59,42 @@ void SwizzleImpl(std::span<u8> output, std::span<const u8> input, u32 width, u32
                              ((z & block_depth_mask) << (GOB_SIZE_SHIFT + block_height));
         for (u32 line = 0; line < height; ++line) {
             const u32 y = line + origin_y;
-            const u32 swizzled_y = pdep<SWIZZLE_Y_BITS>(y);
+            const u32 swizzled_y = pdep(SWIZZLE_Y_BITS, y);
 
             const u32 block_y = y >> GOB_SIZE_Y_SHIFT;
-            const u32 offset_y = (block_y >> block_height) * block_size +
-                                 ((block_y & block_height_mask) << GOB_SIZE_SHIFT);
+            const u32 offset_y = (block_y >> block_height) * block_size + ((block_y & block_height_mask) << GOB_SIZE_SHIFT);
 
-            u32 swizzled_x = pdep<SWIZZLE_X_BITS>(origin_x * BYTES_PER_PIXEL);
-            for (u32 column = 0; column < width;
-                 ++column, incrpdep<SWIZZLE_X_BITS, BYTES_PER_PIXEL>(swizzled_x)) {
-                const u32 x = (column + origin_x) * BYTES_PER_PIXEL;
+            u32 swizzled_x = pdep(SWIZZLE_X_BITS, origin_x * bytes_per_pixel);
+            for (u32 column = 0; column < width; ++column, incrpdep(SWIZZLE_X_BITS, bytes_per_pixel, swizzled_x)) {
+                const u32 x = (column + origin_x) * bytes_per_pixel;
                 const u32 offset_x = (x >> GOB_SIZE_X_SHIFT) << x_shift;
 
                 const u32 base_swizzled_offset = offset_z + offset_y + offset_x;
                 const u32 swizzled_offset = base_swizzled_offset + (swizzled_x | swizzled_y);
 
-                const u32 unswizzled_offset =
-                    slice * pitch * height + line * pitch + column * BYTES_PER_PIXEL;
+                const u32 unswizzled_offset = slice * pitch * height + line * pitch + column * bytes_per_pixel;
+                u8* const dst = &output[to_linear ? swizzled_offset : unswizzled_offset];
+                const u8* const src = &input[to_linear ? unswizzled_offset : swizzled_offset];
 
-                u8* const dst = &output[TO_LINEAR ? swizzled_offset : unswizzled_offset];
-                const u8* const src = &input[TO_LINEAR ? unswizzled_offset : swizzled_offset];
-
-                std::memcpy(dst, src, BYTES_PER_PIXEL);
+                std::memcpy(dst, src, bytes_per_pixel);
             }
         }
     }
 }
 
-template <bool TO_LINEAR, u32 BYTES_PER_PIXEL>
-void SwizzleSubrectImpl(std::span<u8> output, std::span<const u8> input, u32 width, u32 height,
-                        u32 depth, u32 origin_x, u32 origin_y, u32 extent_x, u32 num_lines,
-                        u32 block_height, u32 block_depth, u32 pitch_linear) {
+void SwizzleSubrectImpl(bool to_linear, u32 bytes_per_pixel, std::span<u8> output, std::span<const u8> input, u32 width, u32 height, u32 depth, u32 origin_x, u32 origin_y, u32 extent_x, u32 num_lines, u32 block_height, u32 block_depth, u32 pitch_linear) {
     // The origin of the transformation can be configured here, leave it as zero as the current API
     // doesn't expose it.
     static constexpr u32 origin_z = 0;
 
     // We can configure here a custom pitch
-    // As it's not exposed 'width * BYTES_PER_PIXEL' will be the expected pitch.
+    // As it's not exposed 'width * bytes_per_pixel' will be the expected pitch.
     const u32 pitch = pitch_linear;
-    const u32 stride = Common::AlignUpLog2(width * BYTES_PER_PIXEL, GOB_SIZE_X_SHIFT);
+    const u32 stride = Common::AlignUpLog2(width * bytes_per_pixel, GOB_SIZE_X_SHIFT);
 
     const u32 gobs_in_x = Common::DivCeilLog2(stride, GOB_SIZE_X_SHIFT);
     const u32 block_size = gobs_in_x << (GOB_SIZE_SHIFT + block_height + block_depth);
-    const u32 slice_size =
-        Common::DivCeilLog2(height, block_height + GOB_SIZE_Y_SHIFT) * block_size;
+    const u32 slice_size = Common::DivCeilLog2(height, block_height + GOB_SIZE_Y_SHIFT) * block_size;
 
     const u32 block_height_mask = (1U << block_height) - 1;
     const u32 block_depth_mask = (1U << block_depth) - 1;
@@ -123,28 +110,25 @@ void SwizzleSubrectImpl(std::span<u8> output, std::span<const u8> input, u32 wid
         const u32 lines_in_y = (std::min)(unprocessed_lines, extent_y);
         for (u32 line = 0; line < lines_in_y; ++line) {
             const u32 y = line + origin_y;
-            const u32 swizzled_y = pdep<SWIZZLE_Y_BITS>(y);
+            const u32 swizzled_y = pdep(SWIZZLE_Y_BITS, y);
 
             const u32 block_y = y >> GOB_SIZE_Y_SHIFT;
-            const u32 offset_y = (block_y >> block_height) * block_size +
-                                 ((block_y & block_height_mask) << GOB_SIZE_SHIFT);
+            const u32 offset_y = (block_y >> block_height) * block_size + ((block_y & block_height_mask) << GOB_SIZE_SHIFT);
 
-            u32 swizzled_x = pdep<SWIZZLE_X_BITS>(origin_x * BYTES_PER_PIXEL);
-            for (u32 column = 0; column < extent_x;
-                 ++column, incrpdep<SWIZZLE_X_BITS, BYTES_PER_PIXEL>(swizzled_x)) {
-                const u32 x = (column + origin_x) * BYTES_PER_PIXEL;
+            u32 swizzled_x = pdep(SWIZZLE_X_BITS, origin_x * bytes_per_pixel);
+            for (u32 column = 0; column < extent_x; ++column, incrpdep(SWIZZLE_X_BITS, bytes_per_pixel, swizzled_x)) {
+                const u32 x = (column + origin_x) * bytes_per_pixel;
                 const u32 offset_x = (x >> GOB_SIZE_X_SHIFT) << x_shift;
 
                 const u32 base_swizzled_offset = offset_z + offset_y + offset_x;
                 const u32 swizzled_offset = base_swizzled_offset + (swizzled_x | swizzled_y);
 
-                const u32 unswizzled_offset =
-                    slice * pitch * height + line * pitch + column * BYTES_PER_PIXEL;
+                const u32 unswizzled_offset = slice * pitch * height + line * pitch + column * bytes_per_pixel;
 
-                u8* const dst = &output[TO_LINEAR ? swizzled_offset : unswizzled_offset];
-                const u8* const src = &input[TO_LINEAR ? unswizzled_offset : swizzled_offset];
+                u8* const dst = &output[to_linear ? swizzled_offset : unswizzled_offset];
+                const u8* const src = &input[to_linear ? unswizzled_offset : swizzled_offset];
 
-                std::memcpy(dst, src, BYTES_PER_PIXEL);
+                std::memcpy(dst, src, bytes_per_pixel);
             }
         }
         unprocessed_lines -= lines_in_y;
@@ -154,23 +138,17 @@ void SwizzleSubrectImpl(std::span<u8> output, std::span<const u8> input, u32 wid
     }
 }
 
-template <bool TO_LINEAR>
-void Swizzle(std::span<u8> output, std::span<const u8> input, u32 bytes_per_pixel, u32 width,
-             u32 height, u32 depth, u32 block_height, u32 block_depth, u32 stride_alignment) {
+void Swizzle(bool to_linear, std::span<u8> output, std::span<const u8> input, u32 bytes_per_pixel, u32 width, u32 height, u32 depth, u32 block_height, u32 block_depth, u32 stride_alignment) {
     switch (bytes_per_pixel) {
-#define BPP_CASE(x)                                                                                \
-    case x:                                                                                        \
-        return SwizzleImpl<TO_LINEAR, x>(output, input, width, height, depth, block_height,        \
-                                         block_depth, stride_alignment);
-        BPP_CASE(1)
-        BPP_CASE(2)
-        BPP_CASE(3)
-        BPP_CASE(4)
-        BPP_CASE(6)
-        BPP_CASE(8)
-        BPP_CASE(12)
-        BPP_CASE(16)
-#undef BPP_CASE
+    case 1:
+    case 2:
+    case 3:
+    case 4:
+    case 6:
+    case 8:
+    case 12:
+    case 16:
+        return SwizzleImpl(to_linear, bytes_per_pixel, output, input, width, height, depth, block_height, block_depth, stride_alignment);
     default:
         ASSERT_MSG(false, "Invalid bytes_per_pixel={}", bytes_per_pixel);
         break;
@@ -179,78 +157,57 @@ void Swizzle(std::span<u8> output, std::span<const u8> input, u32 bytes_per_pixe
 
 } // Anonymous namespace
 
-void UnswizzleTexture(std::span<u8> output, std::span<const u8> input, u32 bytes_per_pixel,
-                      u32 width, u32 height, u32 depth, u32 block_height, u32 block_depth,
-                      u32 stride_alignment) {
+void UnswizzleTexture(std::span<u8> output, std::span<const u8> input, u32 bytes_per_pixel, u32 width, u32 height, u32 depth, u32 block_height, u32 block_depth, u32 stride_alignment) {
     const u32 stride = Common::AlignUpLog2(width, stride_alignment) * bytes_per_pixel;
-    const u32 new_bpp = (std::min)(4U, static_cast<u32>(std::countr_zero(width * bytes_per_pixel)));
+    const u32 new_bpp = (std::min)(4U, u32(std::countr_zero(width * bytes_per_pixel)));
     width = (width * bytes_per_pixel) >> new_bpp;
     bytes_per_pixel = 1U << new_bpp;
-    Swizzle<false>(output, input, bytes_per_pixel, width, height, depth, block_height, block_depth,
-                   stride);
+    Swizzle(false, output, input, bytes_per_pixel, width, height, depth, block_height, block_depth, stride);
 }
 
-void SwizzleTexture(std::span<u8> output, std::span<const u8> input, u32 bytes_per_pixel, u32 width,
-                    u32 height, u32 depth, u32 block_height, u32 block_depth,
-                    u32 stride_alignment) {
+void SwizzleTexture(std::span<u8> output, std::span<const u8> input, u32 bytes_per_pixel, u32 width, u32 height, u32 depth, u32 block_height, u32 block_depth, u32 stride_alignment) {
     const u32 stride = Common::AlignUpLog2(width, stride_alignment) * bytes_per_pixel;
-    const u32 new_bpp = (std::min)(4U, static_cast<u32>(std::countr_zero(width * bytes_per_pixel)));
+    const u32 new_bpp = (std::min)(4U, u32(std::countr_zero(width * bytes_per_pixel)));
     width = (width * bytes_per_pixel) >> new_bpp;
     bytes_per_pixel = 1U << new_bpp;
-    Swizzle<true>(output, input, bytes_per_pixel, width, height, depth, block_height, block_depth,
-                  stride);
+    Swizzle(true, output, input, bytes_per_pixel, width, height, depth, block_height, block_depth, stride);
 }
 
-void SwizzleSubrect(std::span<u8> output, std::span<const u8> input, u32 bytes_per_pixel, u32 width,
-                    u32 height, u32 depth, u32 origin_x, u32 origin_y, u32 extent_x, u32 extent_y,
-                    u32 block_height, u32 block_depth, u32 pitch_linear) {
+void SwizzleSubrect(std::span<u8> output, std::span<const u8> input, u32 bytes_per_pixel, u32 width, u32 height, u32 depth, u32 origin_x, u32 origin_y, u32 extent_x, u32 extent_y, u32 block_height, u32 block_depth, u32 pitch_linear) {
     switch (bytes_per_pixel) {
-#define BPP_CASE(x)                                                                                \
-    case x:                                                                                        \
-        return SwizzleSubrectImpl<true, x>(output, input, width, height, depth, origin_x,          \
-                                           origin_y, extent_x, extent_y, block_height,             \
-                                           block_depth, pitch_linear);
-        BPP_CASE(1)
-        BPP_CASE(2)
-        BPP_CASE(3)
-        BPP_CASE(4)
-        BPP_CASE(6)
-        BPP_CASE(8)
-        BPP_CASE(12)
-        BPP_CASE(16)
-#undef BPP_CASE
+    case 1:
+    case 2:
+    case 3:
+    case 4:
+    case 6:
+    case 8:
+    case 12:
+    case 16:
+        return SwizzleSubrectImpl(true, bytes_per_pixel, output, input, width, height, depth, origin_x, origin_y, extent_x, extent_y, block_height, block_depth, pitch_linear);
     default:
         ASSERT_MSG(false, "Invalid bytes_per_pixel={}", bytes_per_pixel);
         break;
     }
 }
 
-void UnswizzleSubrect(std::span<u8> output, std::span<const u8> input, u32 bytes_per_pixel,
-                      u32 width, u32 height, u32 depth, u32 origin_x, u32 origin_y, u32 extent_x,
-                      u32 extent_y, u32 block_height, u32 block_depth, u32 pitch_linear) {
+void UnswizzleSubrect(std::span<u8> output, std::span<const u8> input, u32 bytes_per_pixel, u32 width, u32 height, u32 depth, u32 origin_x, u32 origin_y, u32 extent_x, u32 extent_y, u32 block_height, u32 block_depth, u32 pitch_linear) {
     switch (bytes_per_pixel) {
-#define BPP_CASE(x)                                                                                \
-    case x:                                                                                        \
-        return SwizzleSubrectImpl<false, x>(output, input, width, height, depth, origin_x,         \
-                                            origin_y, extent_x, extent_y, block_height,            \
-                                            block_depth, pitch_linear);
-        BPP_CASE(1)
-        BPP_CASE(2)
-        BPP_CASE(3)
-        BPP_CASE(4)
-        BPP_CASE(6)
-        BPP_CASE(8)
-        BPP_CASE(12)
-        BPP_CASE(16)
-#undef BPP_CASE
+    case 1:
+    case 2:
+    case 3:
+    case 4:
+    case 6:
+    case 8:
+    case 12:
+    case 16:
+        return SwizzleSubrectImpl(false, bytes_per_pixel, output, input, width, height, depth, origin_x, origin_y, extent_x, extent_y, block_height, block_depth, pitch_linear);
     default:
         ASSERT_MSG(false, "Invalid bytes_per_pixel={}", bytes_per_pixel);
         break;
     }
 }
 
-std::size_t CalculateSize(bool tiled, u32 bytes_per_pixel, u32 width, u32 height, u32 depth,
-                          u32 block_height, u32 block_depth) {
+std::size_t CalculateSize(bool tiled, u32 bytes_per_pixel, u32 width, u32 height, u32 depth, u32 block_height, u32 block_depth) {
     if (tiled) {
         const u32 aligned_width = Common::AlignUpLog2(width * bytes_per_pixel, GOB_SIZE_X_SHIFT);
         const u32 aligned_height = Common::AlignUpLog2(height, GOB_SIZE_Y_SHIFT + block_height);
@@ -261,8 +218,7 @@ std::size_t CalculateSize(bool tiled, u32 bytes_per_pixel, u32 width, u32 height
     }
 }
 
-u64 GetGOBOffset(u32 width, u32 height, u32 dst_x, u32 dst_y, u32 block_height,
-                 u32 bytes_per_pixel) {
+u64 GetGOBOffset(u32 width, u32 height, u32 dst_x, u32 dst_y, u32 block_height, u32 bytes_per_pixel) {
     auto div_ceil = [](const u32 x, const u32 y) { return ((x + y - 1) / y); };
     const u32 gobs_in_block = 1 << block_height;
     const u32 y_blocks = GOB_SIZE_Y << block_height;

src/video_core/vulkan_common/vulkan_device.cpp

@@ -605,6 +605,18 @@ Device::Device(VkInstance instance_, vk::PhysicalDevice physical_, VkSurfaceKHR
         if (version < VK_MAKE_API_VERSION(0, 255, 615, 512)) {
             has_broken_parallel_compiling = true;
         }
+        const size_t sampler_limit = properties.properties.limits.maxSamplerAllocationCount;
+        if (sampler_limit > 0) {
+            constexpr size_t MIN_SAMPLER_BUDGET = 1024U;
+            const size_t reserved = sampler_limit / 4U;
+            const size_t derived_budget =
+                (std::max)(MIN_SAMPLER_BUDGET, sampler_limit - reserved);
+            sampler_heap_budget = derived_budget;
+            LOG_WARNING(Render_Vulkan,
+                        "Qualcomm driver reports max {} samplers; reserving {} (25%) and "
+                        "allowing Eden to use {} (75%) to avoid heap exhaustion",
+                        sampler_limit, reserved, sampler_heap_budget);
+        }
     }
 
     if (extensions.sampler_filter_minmax && is_amd) {
@@ -1529,6 +1541,13 @@ void Device::SetupFamilies(VkSurfaceKHR surface) {
     }
 }
 
+std::optional<size_t> Device::GetSamplerHeapBudget() const {
+    if (sampler_heap_budget == 0) {
+        return std::nullopt;
+    }
+    return sampler_heap_budget;
+}
+
 u64 Device::GetDeviceMemoryUsage() const {
     VkPhysicalDeviceMemoryBudgetPropertiesEXT budget;
     budget.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_MEMORY_BUDGET_PROPERTIES_EXT;

src/video_core/vulkan_common/vulkan_device.h

@@ -6,6 +6,7 @@
 
 #pragma once
 
+#include <optional>
 #include <set>
 #include <span>
 #include <string>
@@ -326,6 +327,16 @@ public:
         return properties.properties.limits.maxComputeSharedMemorySize;
     }
 
+    /// Returns the maximum number of dynamic storage buffer descriptors per set.
+    u32 GetMaxDescriptorSetStorageBuffersDynamic() const {
+        return properties.properties.limits.maxDescriptorSetStorageBuffersDynamic;
+    }
+
+    /// Returns the maximum number of dynamic uniform buffer descriptors per set.
+    u32 GetMaxDescriptorSetUniformBuffersDynamic() const {
+        return properties.properties.limits.maxDescriptorSetUniformBuffersDynamic;
+    }
+
     /// Returns float control properties of the device.
     const VkPhysicalDeviceFloatControlsPropertiesKHR& FloatControlProperties() const {
         return properties.float_controls;
@@ -682,11 +693,6 @@ public:
         return dynamic_state3_alpha_to_one;
     }
 
-    /// Returns true when the user enabled extended core dynamic states (level > 0).
-    bool UsesAdvancedCoreDynamicState() const {
-        return u32(Settings::values.dyna_state.GetValue()) > 0;
-    }
-
     /// Returns true if the device supports VK_EXT_vertex_input_dynamic_state.
     bool IsExtVertexInputDynamicStateSupported() const {
         return extensions.vertex_input_dynamic_state;
@@ -749,6 +755,8 @@ public:
         return has_broken_parallel_compiling;
     }
 
+    std::optional<size_t> GetSamplerHeapBudget() const;
+
     /// Returns the vendor name reported from Vulkan.
     std::string_view GetVendorName() const {
         return properties.driver.driverName;
@@ -1045,6 +1053,7 @@ private:
     bool dynamic_state3_alpha_to_coverage{};
     bool dynamic_state3_alpha_to_one{};
     bool supports_conditional_barriers{};      ///< Allows barriers in conditional control flow.
+    size_t sampler_heap_budget{};              ///< Sampler budget for buggy drivers (0 = unlimited).
     u64 device_access_memory{};                ///< Total size of device local memory in bytes.
     u32 sets_per_pool{};                       ///< Sets per Description Pool
     NvidiaArchitecture nvidia_arch{NvidiaArchitecture::Arch_AmpereOrNewer};

src/video_core/vulkan_common/vulkan_wrapper.cpp

@@ -121,7 +121,6 @@ void Load(VkDevice device, DeviceDispatch& dld) noexcept {
     X(vkCmdEndConditionalRenderingEXT);
     X(vkCmdEndQuery);
     X(vkCmdEndRenderPass);
-    X(vkCmdResetQueryPool);
     X(vkCmdEndTransformFeedbackEXT);
     X(vkCmdEndDebugUtilsLabelEXT);
     X(vkCmdFillBuffer);

src/video_core/vulkan_common/vulkan_wrapper.h

@@ -221,7 +221,6 @@ struct DeviceDispatch : InstanceDispatch {
     PFN_vkCmdEndConditionalRenderingEXT vkCmdEndConditionalRenderingEXT{};
     PFN_vkCmdEndDebugUtilsLabelEXT vkCmdEndDebugUtilsLabelEXT{};
     PFN_vkCmdEndQuery vkCmdEndQuery{};
-    PFN_vkCmdResetQueryPool vkCmdResetQueryPool{};
     PFN_vkCmdEndRenderPass vkCmdEndRenderPass{};
     PFN_vkCmdEndTransformFeedbackEXT vkCmdEndTransformFeedbackEXT{};
     PFN_vkCmdFillBuffer vkCmdFillBuffer{};
@@ -1144,9 +1143,6 @@ public:
     VkCommandBuffer operator*() const noexcept {
         return handle;
     }
-    void ResetQueryPool(VkQueryPool query_pool, uint32_t first, uint32_t count) const noexcept {
-        dld->vkCmdResetQueryPool(handle, query_pool, first, count);
-    }
     void Begin(const VkCommandBufferBeginInfo& begin_info) const {
         Check(dld->vkBeginCommandBuffer(handle, &begin_info));
     }