shader: fix crash from packed ancillary used as DPP old operand (#1769 )

shader: emit plain EXEC selects for DPP writes that reach every active lane (Street Fighter 6)
graphics: mark image heap resources as sampled reads
2026-10-23 12:26:56 +02:00 · 2026-10-23 09:51:32 +02:00 · 2026-10-23 08:08:56 +02:00 · 2026-10-23 07:57:57 +02:00 · 2026-10-23 07:42:28 +02:00 · 2026-10-23 07:22:33 +02:00
252 changed files with 16832 additions and 3592 deletions
@@ -113,7 +113,7 @@ jobs:
        shell: cmd
        run: |
          cmake --build _Build/windows --target launcher ^
-            audio_out2_port_tests pad_haptics_tests controller_settings_tests ime_dialog_tests lru_cache_tests virtual_memory_allocation_tests ^
+            audio_out2_port_tests pad_haptics_tests controller_settings_tests host_input_tests ime_dialog_tests lru_cache_tests virtual_memory_allocation_tests ^
            shader_recompiler_compute_tests ^
            resource_tracking_tests scalar_provenance_tests ^
            resource_materialization_tests shader_vertex_metadata_tests archive_file_tests avplayer_file_tests --parallel
@@ -121,7 +121,7 @@ jobs:
      - name: Test
        shell: cmd
        run: |
-          ctest --test-dir _Build/windows --output-on-failure -R "^(audio_out2_port|pad_haptics|controller_settings|ime_dialog|lru_cache|resource_tracking|scalar_provenance|resource_materialization|shader_vertex_metadata|archive_file|avplayer_file|virtual_memory_allocation)$"
+          ctest --test-dir _Build/windows --output-on-failure -R "^(audio_out2_port|pad_haptics|controller_settings|host_input|ime_dialog|lru_cache|resource_tracking|scalar_provenance|resource_materialization|shader_vertex_metadata|archive_file|avplayer_file|virtual_memory_allocation)$"

      - name: Install
        shell: cmd
@@ -202,7 +202,7 @@ jobs:
        shell: bash
        run: |
          cmake --build _Build/macos \
-            --target launcher audio_out2_port_tests pad_haptics_tests controller_settings_tests ime_dialog_tests lru_cache_tests virtual_memory_allocation_tests \
+            --target launcher audio_out2_port_tests pad_haptics_tests controller_settings_tests host_input_tests ime_dialog_tests lru_cache_tests virtual_memory_allocation_tests \
              memory_tracker_tests page_manager_tests shader_recompiler_compute_tests \
              resource_tracking_tests scalar_provenance_tests \
              resource_materialization_tests shader_vertex_metadata_tests archive_file_tests avplayer_file_tests \
@@ -212,7 +212,7 @@ jobs:
        shell: bash
        run: |
          ctest --test-dir _Build/macos --output-on-failure \
-            -R '^(audio_out2_port|pad_haptics|controller_settings|ime_dialog|lru_cache|memory_tracker|page_manager|resource_tracking|scalar_provenance|resource_materialization|shader_vertex_metadata|archive_file|avplayer_file|virtual_memory_allocation)$'
+            -R '^(audio_out2_port|pad_haptics|controller_settings|host_input|ime_dialog|lru_cache|memory_tracker|page_manager|resource_tracking|scalar_provenance|resource_materialization|shader_vertex_metadata|archive_file|avplayer_file|virtual_memory_allocation)$'

      - name: Install
        shell: bash
@@ -388,7 +388,7 @@ jobs:
        run: |
          cmake --build _Build/linux \
            --target launcher page_manager_tests memory_tracker_tests \
-              audio_out2_port_tests pad_haptics_tests controller_settings_tests ime_dialog_tests lru_cache_tests virtual_memory_allocation_tests \
+              audio_out2_port_tests pad_haptics_tests controller_settings_tests host_input_tests ime_dialog_tests lru_cache_tests virtual_memory_allocation_tests \
              resource_tracking_tests scalar_provenance_tests \
              resource_materialization_tests shader_vertex_metadata_tests archive_file_tests avplayer_file_tests \
            --parallel
@@ -397,7 +397,7 @@ jobs:
        shell: bash
        run: |
          ctest --test-dir _Build/linux --output-on-failure \
-            -R '^(audio_out2_port|pad_haptics|controller_settings|ime_dialog|lru_cache|memory_tracker|page_manager|resource_tracking|scalar_provenance|resource_materialization|shader_vertex_metadata|archive_file|avplayer_file|virtual_memory_allocation)$'
+            -R '^(audio_out2_port|pad_haptics|controller_settings|host_input|ime_dialog|lru_cache|memory_tracker|page_manager|resource_tracking|scalar_provenance|resource_materialization|shader_vertex_metadata|archive_file|avplayer_file|virtual_memory_allocation)$'

      - name: Install
        shell: bash
@@ -407,12 +407,12 @@ jobs:
      - name: Verify artifacts
        shell: bash
        run: |
-          file _Build/linux/install/launcher
+          file _Build/linux/install/kyty_launcher
          file _Build/linux/install/kyty_emulator
          test -x _Build/linux/install/kyty_updater
          file _Build/linux/install/kyty_emulator | grep -q "ELF 64-bit LSB .*x86-64"
          ldd _Build/linux/install/kyty_emulator > /dev/null
-          readelf -d _Build/linux/install/launcher | grep -q 'RPATH.*\$ORIGIN/lib'
+          readelf -d _Build/linux/install/kyty_launcher | grep -q 'RPATH.*\$ORIGIN/lib'
          while IFS= read -r -d '' binary; do
            while read -r dependency; do
              test -e "_Build/linux/install/lib/$dependency"
@@ -421,8 +421,8 @@ jobs:
                sed -n 's/.*Shared library: \[\(libQt6[^]]*\|libicu[^]]*\)\].*/\1/p'
            )
          done < <(
-            find _Build/linux/install/launcher _Build/linux/install/plugins \
-              -type f \( -name launcher -o -name '*.so' \) -print0
+            find _Build/linux/install/kyty_launcher _Build/linux/install/plugins \
+              -type f \( -name kyty_launcher -o -name '*.so' \) -print0
          )
          _Build/linux/install/kyty_emulator --help > /dev/null

@@ -497,7 +497,7 @@ jobs:
          linux_dir="artifacts/KytyPS5-Linux-x86_64"
          test -d "$linux_dir"
          chmod a+x \
-            "$linux_dir/launcher" \
+            "$linux_dir/kyty_launcher" \
            "$linux_dir/kyty_emulator" \
            "$linux_dir/kyty_updater"
          tar -C "$linux_dir" -czf \
@@ -6,3 +6,4 @@ _Build/vscode-clang/
 _Build/
 cmake-build-debug/
 _PipelineCache/
+CMakeUserPresets.json
@@ -5,7 +5,7 @@
            "name": "Debug launcher UI",
            "type": "cppvsdbg",
            "request": "launch",
-            "program": "${workspaceFolder}\\_Build\\vscode-clang\\launcher\\launcher.exe",
+            "program": "${workspaceFolder}\\_Build\\vscode-clang\\launcher\\kyty_launcher.exe",
            "args": [],
            "stopAtEntry": false,
            "cwd": "${workspaceFolder}\\_Build\\vscode-clang\\launcher",
@@ -13,7 +13,7 @@
            "externalConsole": false,
            "linux": {
                "type": "cppdbg",
-                "program": "${workspaceFolder}/_Build/vscode-clang/launcher/launcher",
+                "program": "${workspaceFolder}/_Build/vscode-clang/launcher/kyty_launcher",
                "cwd": "${workspaceFolder}/_Build/vscode-clang/launcher",
                "MIMode": "gdb",
                "miDebuggerPath": "/usr/bin/gdb"
@@ -27,7 +27,7 @@
            "args": [
                "--game", "D:/gamesps5/PPSA07429",
                "--printf-direction", "File",
-                "--graphics-debug-dump", "True"
+                "--graphics-debug-dump"
            ],
            "stopAtEntry": false,
            "cwd": "${workspaceFolder}\\_Build\\vscode-clang",
@@ -25,6 +25,23 @@ FetchContent_MakeAvailable(opus)
 set(BUILD_TESTING "${kyty_build_testing_saved}")
 install(FILES "${opus_SOURCE_DIR}/COPYING" DESTINATION "licenses/opus")

+# Headsets are driven through OpenXR. Only the static loader is built; the runtime that serves it
+# (for example SteamVR, Monado or Virtual Desktop) comes from the user's system.
+set(DYNAMIC_LOADER OFF CACHE BOOL "" FORCE)
+set(BUILD_WITH_SYSTEM_JSONCPP OFF CACHE BOOL "" FORCE)
+FetchContent_Declare(openxr
+	GIT_REPOSITORY https://github.com/KhronosGroup/OpenXR-SDK.git
+	GIT_TAG f2448a8797c85814aa892efc1ab8707900fbcc78 # release-1.1.63
+	GIT_SHALLOW FALSE)
+FetchContent_MakeAvailable(openxr)
+# Keep the loader's headers and library out of the installed package.
+set_property(DIRECTORY "${openxr_SOURCE_DIR}" PROPERTY EXCLUDE_FROM_ALL TRUE)
+# The loader uses exceptions internally.
+if(NOT KYTY_CLANG_CL)
+	target_compile_options(openxr_loader PRIVATE -fexceptions)
+endif()
+install(FILES "${openxr_SOURCE_DIR}/LICENSE" DESTINATION "licenses/openxr")
+
 # The guest instruction patcher needs a decoder/encoder and a small
 # trampoline assembler. Pin the same revisions used by shadPS4's implementation.
 set(XBYAK_BUILD_TESTS OFF CACHE BOOL "" FORCE)
@@ -265,6 +265,7 @@ set(gpu_tiler_shader_names
 	promote_d16
 	demote_d16
 	color_transform
+	color_clear
 )
 file(GLOB gpu_tiler_shader_includes CONFIGURE_DEPENDS "${gpu_tiler_shader_dir}/gpu_tiler_*.inc")
 foreach(shader_name IN LISTS gpu_tiler_shader_names)
@@ -336,7 +337,7 @@ if (KYTY_CLANG_CL)
 	set_source_files_properties(${kyty_emulator_src} PROPERTIES COMPILE_FLAGS "-Wno-pragma-pack -Wno-deprecated-declarations -D_TIMESPEC_DEFINED")
 endif()

-set(kyty_emulator_link_libraries common Vulkan::Headers spirv-tools-opt spirv-tools SDL3::SDL3-static imgui xxhash FFmpeg::ffmpeg fmt::fmt nlohmann_json::nlohmann_json LibAtrac9 kyty_git_version Zydis::Zydis xbyak::xbyak opus)
+set(kyty_emulator_link_libraries common Vulkan::Headers spirv-tools-opt spirv-tools SDL3::SDL3-static imgui xxhash FFmpeg::ffmpeg fmt::fmt nlohmann_json::nlohmann_json LibAtrac9 kyty_git_version Zydis::Zydis xbyak::xbyak opus OpenXR::openxr_loader)

 # Linux system libraries.
 if(LINUX)
@@ -350,7 +351,7 @@ set(inc_headers
 	${KYTY_THIRD_PARTY_DIR}/SPIRV-Tools/include
 	${KYTY_THIRD_PARTY_DIR}/SPIRV-Headers/include
 	${KYTY_THIRD_PARTY_DIR}/fmt/include
-	${KYTY_THIRD_PARTY_DIR}/xxhash
+	${KYTY_THIRD_PARTY_DIR}/xxHash
 	${KYTY_THIRD_PARTY_DIR}/cpuinfo/include
 	${KYTY_THIRD_PARTY_DIR}/renderdoc
 	${KYTY_THIRD_PARTY_DIR}/stb
@@ -526,6 +527,10 @@ add_executable(controller_settings_tests EXCLUDE_FROM_ALL
 target_link_libraries(controller_settings_tests common fmt::fmt SDL3::SDL3-static)
 target_include_directories(controller_settings_tests PRIVATE ${inc_headers})

+add_executable(host_input_tests EXCLUDE_FROM_ALL "${KYTY_TESTS_DIR}/HostInputTests.cpp")
+target_link_libraries(host_input_tests common fmt::fmt SDL3::SDL3-static)
+target_include_directories(host_input_tests PRIVATE ${inc_headers})
+
 add_executable(save_data_memory_tests EXCLUDE_FROM_ALL
 	"${KYTY_TESTS_DIR}/SaveDataMemoryTests.cpp"
 )
@@ -644,12 +649,17 @@ if(BUILD_TESTING)
 	add_test(NAME ngs2_sampler COMMAND $<TARGET_FILE:ngs2_sampler_tests>)
 	add_test(NAME pad_haptics COMMAND $<TARGET_FILE:pad_haptics_tests>)
 	add_test(NAME controller_settings COMMAND $<TARGET_FILE:controller_settings_tests>)
+	add_test(NAME host_input COMMAND $<TARGET_FILE:host_input_tests>)
 	add_test(NAME save_data_memory COMMAND $<TARGET_FILE:save_data_memory_tests>)
 	add_test(NAME png_enc COMMAND $<TARGET_FILE:png_enc_tests>)
 	add_test(NAME ces COMMAND $<TARGET_FILE:ces_tests>)
 	add_test(NAME save_data_dialog COMMAND $<TARGET_FILE:save_data_dialog_tests>)
 	add_test(NAME http_uri_parse COMMAND $<TARGET_FILE:http_uri_parse_tests>)
 	add_test(NAME shader_recompiler_compute COMMAND $<TARGET_FILE:shader_recompiler_compute_tests>)
+	add_test(NAME shader_recompiler_tessellation
+		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --tessellation-only)
+	add_test(NAME shader_recompiler_float_min_max
+		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --float-min-max-only)
 	add_test(NAME shader_recompiler_alignbyte
 		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --alignbyte-only)
 	add_test(NAME graphics_pipeline_rasterization
@@ -666,6 +676,12 @@ if(BUILD_TESTING)
 		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --stream-buffer-only)
 	add_test(NAME gpu_command_lane
 		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --gpu-command-lane-only)
+	add_test(NAME wait_until_safe_for_rendering
+		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --wait-until-safe-for-rendering-only)
+	add_test(NAME depth_bits
+		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --depth-bits-only)
+	add_test(NAME msaa_storage
+		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --msaa-storage-only)
 	add_test(NAME pm4_context_state
 		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --context-state-only)
 	add_test(NAME gpu_tiler
@@ -723,6 +739,7 @@ if(BUILD_TESTING)
 		ngs2_sampler_tests
 		pad_haptics_tests
 		controller_settings_tests
+		host_input_tests
 		save_data_memory_tests
 		png_enc_tests
 		ces_tests
@@ -760,6 +777,14 @@ add_custom_target(kyty_trophy_sound
 add_dependencies(kyty_emulator kyty_trophy_sound)
 install(FILES "${KYTY_TROPHY_UNLOCK_SOUND}" DESTINATION assets/sounds)

+add_custom_target(kyty_touch_cursors
+	COMMAND ${CMAKE_COMMAND} -E copy_directory
+		"${CMAKE_CURRENT_SOURCE_DIR}/assets/cursors"
+		"$<TARGET_FILE_DIR:kyty_emulator>/assets/cursors"
+	VERBATIM)
+add_dependencies(kyty_emulator kyty_touch_cursors)
+install(DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}/assets/cursors" DESTINATION assets)
+
 target_link_libraries(kyty_emulator ${kyty_emulator_link_libraries})
 if (WIN32)
 	target_link_libraries(kyty_emulator opengl32)
@@ -46,6 +46,16 @@ or graphical glitches, so please include the version you tested when reporting a
 ## Screenshots

 <table align="center">
+  <tr>
+    <td align="center">
+      <strong>Cyberpunk 2077</strong><br>
+      <img src="docs/screenshots/ps5-07.png" width="300" alt="Cyberpunk 2077 running in KytyPS5">
+    </td>
+    <td align="center">
+      <strong>MLB The Show 23</strong><br>
+      <img src="docs/screenshots/ps5-08.png" width="300" alt="MLB The Show 23 running in KytyPS5">
+    </td>
+  </tr>
  <tr>
    <td align="center">
      <strong>Astro Bot</strong><br>
@@ -275,11 +285,11 @@ From the repository root, follow the launcher configure, build and install steps
 [Building on Linux](#building-on-linux), omitting `-DCMAKE_PREFIX_PATH="$Qt6_DIR"`; CMake finds
 Fedora's packaged Qt 6 automatically.

-Run `./_Build/linux/install/launcher` on regular Fedora or inside the container. To launch a
+Run `./_Build/linux/install/kyty_launcher` on regular Fedora or inside the container. To launch a
 container build from the host, leave the container with `exit`, then run from the repository root:

 ```bash
-distrobox enter kyty-build -- ./_Build/linux/install/launcher
+distrobox enter kyty-build -- ./_Build/linux/install/kyty_launcher
 ```

 Qt libraries and plugins are bundled. You can also launch directly on a Fedora Atomic host
@@ -345,7 +355,7 @@ Use `_Build/linux` instead of `_Build/windows` for a Linux build.

 A ready-made Visual Studio Code setup is included in [`.vscode`](.vscode). It configures CMake
 Tools to build the project with Ninja and `clang-cl` and provides launch profiles for both
-`launcher.exe` and `kyty_emulator.exe`. It is Windows-only: VS Code settings cannot select a
+`kyty_launcher.exe` and `kyty_emulator.exe`. It is Windows-only: VS Code settings cannot select a
 compiler per platform, so on Linux configure from the command line as shown above.

 Before using it:
@@ -365,11 +375,11 @@ Update your graphics driver before reporting rendering problems.
 To use the graphical launcher:

 ```powershell
-.\_Build\windows\install\launcher.exe
+.\_Build\windows\install\kyty_launcher.exe
 ```

 ```bash
-./_Build/linux/install/launcher
+./_Build/linux/install/kyty_launcher
 ```

 ```bash
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "تم كسب الكأس!",
+  "user.guest": "ضيف"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Trofej získána!",
+  "user.guest": "Host"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Trofæ optjent!",
+  "user.guest": "Gæst"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Trophäe erhalten!",
+  "user.guest": "Gast"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Κερδίσατε ένα τρόπαιο!",
+  "user.guest": "Επισκ."
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Trophy earned!",
+  "user.guest": "Guest"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Trophy earned!",
+  "user.guest": "Guest"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "¡Ganaste un trofeo!",
+  "user.guest": "Invitado"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "¡Ganaste un trofeo!",
+  "user.guest": "Invitado"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Pokaali ansaittu!",
+  "user.guest": "Vieras"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Trophée obtenu !",
+  "user.guest": "Invité"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Trophée obtenu !",
+  "user.guest": "Invité"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Trófea megszerezve!",
+  "user.guest": "Vendég"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Trofi diperoleh!",
+  "user.guest": "Tamu"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Trofeo ottenuto!",
+  "user.guest": "Ospite"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "トロフィーを獲得しました！",
+  "user.guest": "ゲスト"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "트로피를 획득했습니다!",
+  "user.guest": "게스트"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Trofee behaald!",
+  "user.guest": "Gast"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Trofé opptjent!",
+  "user.guest": "Gjest"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Zdobyto trofeum!",
+  "user.guest": "Gość"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Troféu conquistado!",
+  "user.guest": "Convidado"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Troféu conquistado!",
+  "user.guest": "Convidado"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Trofeu câștigat!",
+  "user.guest": "Invitat"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Трофей получен!",
+  "user.guest": "Гость"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Trofé intjänad!",
+  "user.guest": "Gäst"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "ได้รับถ้วยรางวัลแล้ว!",
+  "user.guest": "แขก"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Kupa kazanıldı!",
+  "user.guest": "Misafir"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "Đã nhận cúp!",
+  "user.guest": "Khách"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "已获得奖杯！",
+  "user.guest": "访客"
+}
@@ -0,0 +1,4 @@
+{
+  "trophy.toast.earned": "已獲得獎盃！",
+  "user.guest": "訪客"
+}
@@ -0,0 +1,15 @@
+<svg xmlns="http://www.w3.org/2000/svg" viewBox="0 0 512 512"><defs>
+<linearGradient id="body" x1="0" y1="0" x2="0" y2="1"><stop offset="0" stop-color="#ffffff"/><stop offset="1" stop-color="#e3e9f2"/></linearGradient>
+</defs>
+<g transform="translate(256 268) scale(1.16) translate(-256 -268)" stroke-linejoin="round" stroke-linecap="round">
+<path d="M100 196 L120 104 Q124 92 134 99 L214 160 Z" fill="url(#body)"/><path d="M412 196 L392 104 Q388 92 378 99 L298 160 Z" fill="url(#body)"/>
+<path d="M128 164 L137 128 L174 156 Z" fill="#0b0d12"/><path d="M384 164 L375 128 L338 156 Z" fill="#0b0d12"/>
+<g fill="none" stroke="#0b0d12" stroke-width="16"><path d="M100 196 L120 104 Q124 92 134 99 L214 160 Z"/><path d="M412 196 L392 104 Q388 92 378 99 L298 160 Z"/></g>
+<path d="M256 154 C330 154 380 160 406 178 C432 196 447 244 459 302 C471 358 466 398 430 408 C400 416 380 398 360 368 L326 318 C316 304 306 300 290 300 L222 300 C206 300 196 304 186 318 L152 368 C132 398 112 416 82 408 C46 398 41 358 53 302 C65 244 80 196 106 178 C132 160 182 154 256 154 Z" fill="url(#body)"/>
+<path d="M256 154 C330 154 380 160 406 178 C432 196 447 244 459 302 C471 358 466 398 430 408 C400 416 380 398 360 368 L326 318 C316 304 306 300 290 300 L222 300 C206 300 196 304 186 318 L152 368 C132 398 112 416 82 408 C46 398 41 358 53 302 C65 244 80 196 106 178 C132 160 182 154 256 154 Z" fill="none" stroke="#0b0d12" stroke-width="16"/>
+<g fill="none" stroke="#0b0d12" stroke-width="9"><path d="M82 260 C100 274 118 288 128 304 C134 316 121 350 108 388"/><path d="M430 260 C412 274 394 288 384 304 C378 316 391 350 404 388"/></g>
+<path d="M64 322 L82 340 M62 340 L76 354 M448 322 L430 340 M450 340 L436 354" fill="none" stroke="#0b0d12" stroke-width="7"/>
+<path d="M168 230 Q194 202 220 230" fill="none" stroke="#0b0d12" stroke-width="15"/>
+<circle cx="322" cy="220" r="22" fill="#0b0d12"/>
+<path d="M238 252 L274 252 L256 276 Z" fill="#0b0d12" stroke="#0b0d12" stroke-width="8"/>
+</g></svg>
@@ -0,0 +1,20 @@
+# Notification localization
+
+Trophy notification subtitles use the configured console language. Trophy names still come
+from the game's metadata. The launcher trophy dialogs are unchanged.
+
+Edit `assets/localization/<locale>.json`, using flat UTF-8 string values:
+
+```json
+{
+  "trophy.toast.earned": "Trophy earned!"
+}
+```
+
+CMake converts the catalogs into C++ byte arrays in the build directory. The emulator parses
+the embedded catalogs once and falls back to `en-US.json` for missing translations. No language
+files or translation service are needed at runtime. Rebuild after editing a catalog.
+
+For local testing, rebuild `kyty_emulator`, choose a game's Console Language in the launcher,
+and unlock a trophy with notifications enabled. Installed system fonts provide non-Latin
+glyphs; the selected language needs an appropriate font on the host.
@@ -12,9 +12,23 @@ endif()

 add_library(common STATIC ${common_src} ${common_headers})

+set(localization_dir "${PROJECT_SOURCE_DIR}/assets/localization")
+file(GLOB localization_catalogs CONFIGURE_DEPENDS "${localization_dir}/*.json")
+set(localization_data "${CMAKE_CURRENT_BINARY_DIR}/localizationData.h")
+add_custom_command(
+	OUTPUT "${localization_data}"
+	COMMAND ${CMAKE_COMMAND} "-DINPUTS=${localization_catalogs}" "-DOUTPUT=${localization_data}"
+		-P "${KYTY_SOURCE_DIR}/embed_localization.cmake"
+	DEPENDS ${localization_catalogs} "${KYTY_SOURCE_DIR}/embed_localization.cmake"
+	COMMENT "Embed notification localization catalogs"
+	VERBATIM
+)
+target_sources(common PRIVATE "${localization_data}")
+
 target_include_directories(common
 	PUBLIC ${KYTY_SOURCE_DIR}
 	PRIVATE
+		${CMAKE_CURRENT_BINARY_DIR}
 		${KYTY_THIRD_PARTY_DIR}/cpuinfo/include
 )

@@ -29,6 +43,10 @@ target_link_libraries(common
 		ZArchive::zarchive
 		nlohmann_json::nlohmann_json
 )
+if(APPLE)
+	# hostStats.cpp reads GPU utilization from the IOAccelerator registry entries.
+	target_link_libraries(common PRIVATE "-framework CoreFoundation" "-framework IOKit")
+endif()

 get_property(inc_headers TARGET common PROPERTY INCLUDE_DIRECTORIES)
 set(check_headers ${CMAKE_CURRENT_SOURCE_DIR})
@@ -5,6 +5,7 @@

 #include <cstdint>
 #include <filesystem>
+#include <functional>
 #include <memory>
 #include <vector>

@@ -45,6 +46,12 @@ private:
 // Accepts either a host archive filename or a virtual path inside the archive.
 [[nodiscard]] std::shared_ptr<ArchiveReader> OpenArchive(const std::filesystem::path& path);

+// Packs a host directory into a new ZArchive outside it. progress gets 0 < done <= total and
+// returns false to cancel. A failed or cancelled pack removes the partial archive.
+[[nodiscard]] bool PackArchive(const std::filesystem::path& directory,
+                               const std::filesystem::path& archive,
+                               const std::function<bool(uint64_t done, uint64_t total)>& progress);
+
 } // namespace Common

 #endif /* KYTY_COMMON_ARCHIVE_H_ */
@@ -0,0 +1,84 @@
+// Kept out of archive.cpp so the statically linked emulator does not pull in the writer and the
+// zstd compressor.
+#include "common/archive.h"
+
+#include "common/stringUtils.h"
+
+#include <algorithm>
+#include <zarchive/zarchivewriter.h>
+
+namespace Common {
+
+bool PackArchive(const std::filesystem::path& directory, const std::filesystem::path& archive,
+                 const std::function<bool(uint64_t done, uint64_t total)>& progress) {
+	namespace fs = std::filesystem;
+	struct Output {
+		File file;
+		bool ok = true;
+	} output;
+	// Create the archive before walking, so a re-run started while a cancelled pack is still
+	// walking sees it exists instead of losing its archive to that pack's cleanup.
+	if (!output.file.Create(archive)) {
+		return false;
+	}
+	struct Entry {
+		fs::path path;
+		uint64_t size;
+		bool     is_file;
+	};
+	std::vector<Entry> entries;
+	uint64_t           total = 0;
+	std::error_code    error;
+	// Follow directory links, as the emulator does when it runs the folder.
+	fs::recursive_directory_iterator it(directory, fs::directory_options::follow_directory_symlink,
+	                                    error);
+	for (; !error && it != fs::end(it); it.increment(error)) {
+		const bool is_file = !it->is_directory(error);
+		const auto size    = is_file && !error ? it->file_size(error) : 0;
+		if (error) {
+			break;
+		}
+		entries.push_back({it->path(), size, is_file});
+		total += size;
+	}
+	output.ok = !error;
+	ZArchiveWriter writer(
+	    [](int32_t, void*) {},
+	    [](const void* data, size_t size, void* opaque) {
+		    auto&    out     = *static_cast<Output*>(opaque);
+		    uint32_t written = 0;
+		    out.file.Write(data, static_cast<uint32_t>(size), &written);
+		    out.ok = out.ok && written == size;
+	    },
+	    &output);
+	std::vector<uint8_t> buffer(1u << 20u);
+	uint64_t             done = 0;
+	for (size_t index = 0; output.ok && index < entries.size(); index++) {
+		const auto& entry = entries[index];
+		const auto  name  = PathToGenericString(entry.path.lexically_relative(directory));
+		File        input;
+		output.ok = entry.is_file ? writer.StartNewFile(name.c_str()) &&
+		                                input.Open(entry.path, File::Mode::Read)
+		                          : writer.MakeDir(name.c_str());
+		for (uint64_t left = entry.size; output.ok && left != 0;) {
+			uint32_t read = 0;
+			input.Read(buffer.data(), static_cast<uint32_t>(std::min<uint64_t>(left, buffer.size())),
+			           &read);
+			writer.AppendData(buffer.data(), read);
+			done += read;
+			left -= read;
+			output.ok = output.ok && read != 0 && progress(done, total);
+		}
+	}
+	if (output.ok) {
+		writer.Finalize();
+	}
+	output.ok = output.ok && output.file.Flush();
+	output.file.Close();
+	if (!output.ok) {
+		File::DeleteFile(archive);
+	}
+	return output.ok;
+}
+
+} // namespace Common
@@ -0,0 +1,120 @@
+#ifndef KYTY_COMMON_DEBUG_COUNTERS_H_
+#define KYTY_COMMON_DEBUG_COUNTERS_H_
+
+#include <atomic>
+#include <chrono>
+#include <cstdint>
+
+namespace Common::DebugCounters {
+
+static_assert(std::atomic<uint64_t>::is_always_lock_free && std::atomic<bool>::is_always_lock_free);
+
+// Collect events only while the panel is visible. Relaxed, lock-free atomics keep updates
+// cheap and safe in signal handlers. Rates come from comparing two snapshots.
+enum class Counter {
+	// GPU work.
+	Draws,
+	DrawInstances,
+	IndirectDraws,
+	Dispatches,
+	QueueSubmits,
+	GpuWaits, // CPU blocked on a GPU timeline semaphore.
+	GpuWaitNs,
+	GpuFaults, // Guest writes to GPU-tracked memory.
+	ShadersCompiled,
+	ShaderCompileNs,
+	PipelinesCreated,
+	PipelineCreateNs,
+	// Caches.
+	TextureUploadBytes,
+	TextureDownloadBytes,
+	TextureEvictions,
+	ColorMetadataReadbacks, // CPU blocked reading back GPU-written DCC/CMASK metadata.
+	ColorMetadataReadbackNs,
+	ColorMetadataChecks, // GPU clear-key checks of DCC/CMASK metadata.
+	BufferUploadBytes,
+	BufferDownloadBytes,
+	BufferEvictions,
+	// Presentation.
+	FlipsRequested,
+	FlipsRejected, // Guest flip queue was full.
+	LateVblanks,   // The present thread fell behind the vblank rate.
+	SwapchainRecreations,
+	DroppedPresents,
+	// Emulator.
+	EmulatedInstructions, // Host CPU lacked an instruction the guest used.
+	UnresolvedImportCalls,
+	GuestFileReadBytes,
+	GuestFileWriteBytes,
+	AudioUnderruns,
+	Count
+};
+
+enum class Peak { ShaderCompileNs, PipelineCreateNs, Count };
+
+namespace Detail {
+inline std::atomic<bool>     g_enabled {false};
+inline std::atomic<uint64_t> g_counters[static_cast<int>(Counter::Count)] {};
+inline std::atomic<uint64_t> g_peaks[static_cast<int>(Peak::Count)] {};
+} // namespace Detail
+
+[[nodiscard]] inline bool Enabled() noexcept {
+	return Detail::g_enabled.load(std::memory_order_relaxed);
+}
+
+inline void SetEnabled(bool enabled) noexcept {
+	Detail::g_enabled.store(enabled, std::memory_order_relaxed);
+}
+
+inline void Add(Counter counter, uint64_t value = 1) noexcept {
+	if (!Enabled()) {
+		return;
+	}
+	Detail::g_counters[static_cast<int>(counter)].fetch_add(value, std::memory_order_relaxed);
+}
+
+class ScopedTimer final {
+public:
+	ScopedTimer(Counter count, Counter duration, Peak peak = Peak::Count)
+	    : m_count(count), m_duration(duration), m_peak(peak),
+	      m_start(Enabled() ? std::chrono::steady_clock::now()
+	                        : std::chrono::steady_clock::time_point {}) {}
+	~ScopedTimer() {
+		if (m_start == std::chrono::steady_clock::time_point {} || !Enabled()) {
+			return;
+		}
+		const auto elapsed =
+		    static_cast<uint64_t>(std::chrono::duration_cast<std::chrono::nanoseconds>(
+		                              std::chrono::steady_clock::now() - m_start)
+		                              .count());
+		Add(m_count);
+		Add(m_duration, elapsed);
+		if (m_peak != Peak::Count) {
+			auto& peak     = Detail::g_peaks[static_cast<int>(m_peak)];
+			auto  previous = peak.load(std::memory_order_relaxed);
+			while (previous < elapsed &&
+			       !peak.compare_exchange_weak(previous, elapsed, std::memory_order_relaxed)) {
+			}
+		}
+	}
+	ScopedTimer(const ScopedTimer&)            = delete;
+	ScopedTimer& operator=(const ScopedTimer&) = delete;
+
+private:
+	Counter                               m_count;
+	Counter                               m_duration;
+	Peak                                  m_peak;
+	std::chrono::steady_clock::time_point m_start;
+};
+
+[[nodiscard]] inline uint64_t Get(Counter counter) noexcept {
+	return Detail::g_counters[static_cast<int>(counter)].load(std::memory_order_relaxed);
+}
+
+[[nodiscard]] inline uint64_t Get(Peak peak) noexcept {
+	return Detail::g_peaks[static_cast<int>(peak)].load(std::memory_order_relaxed);
+}
+
+} // namespace Common::DebugCounters
+
+#endif /* KYTY_COMMON_DEBUG_COUNTERS_H_ */
@@ -1,6 +1,7 @@
 #include "common/emulatorConfig.h"

 #include "common/assert.h"
+#include "common/debugCounters.h"

 #include <algorithm>
 #include <memory>
@@ -16,6 +17,7 @@ void Initialize() {
 }

 void Shutdown() {
+	Common::DebugCounters::SetEnabled(false);
 	g_config.reset();
 }

@@ -24,8 +26,10 @@ void Load(const ConfigOptions& cfg) {
 	EXIT_IF(cfg.user_name.empty() || cfg.user_name.size() > MAX_USER_NAME_LENGTH);
 	EXIT_IF(!IsConfiguredUserIdValid(cfg.user_id));
 	EXIT_IF(cfg.controller_speaker_volume > 100 || cfg.controller_vibration_intensity > 100);
+	EXIT_IF(cfg.keyboard_player > static_cast<uint32_t>(MAX_PLAYERS));

 	*g_config = cfg;
+	Common::DebugCounters::SetEnabled(cfg.perf_overlay_enabled);
 }

 uint32_t GetScreenWidth() {
@@ -60,6 +64,27 @@ uint32_t GetControllerVibrationIntensity() {
 	return g_config->controller_vibration_intensity;
 }

+uint32_t GetKeyboardPlayer() {
+	return g_config->keyboard_player;
+}
+
+int32_t GetPlayerUserId(int player) {
+	return g_config->user_id + player * PLAYER_USER_ID_STEP;
+}
+
+int GetPlayerFromUserId(int32_t user_id) {
+	for (int player = 0; player < MAX_PLAYERS; player++) {
+		if (user_id == GetPlayerUserId(player)) {
+			return player;
+		}
+	}
+	return -1;
+}
+
+bool IsGuestUserId(int32_t user_id) {
+	return GetPlayerFromUserId(user_id) > 0;
+}
+
 PresentMode GetPresentMode() {
 	return g_config->present_mode;
 }
@@ -160,6 +185,10 @@ bool ReadbackLinearImagesEnabled() {
 	return g_config->readback_linear_images;
 }

+bool ReadbackTiledImagesEnabled() {
+	return g_config->readback_tiled_images;
+}
+
 bool SyncRawImageBuffersEnabled() {
 	return g_config->sync_raw_image_buffers;
 }
@@ -6,6 +6,7 @@
 #include <cstddef>
 #include <array>
 #include <filesystem>
+#include <limits>
 #include <optional>
 #include <string>
 #include <vector>
@@ -35,10 +36,16 @@ constexpr uint32_t MAX_CONSOLE_LANGUAGE     = 29;
 constexpr std::size_t MAX_USER_NAME_LENGTH = 16;
 constexpr int32_t DEFAULT_USER_ID           = 1000;

+constexpr int     MAX_PLAYERS         = 4;
+constexpr int32_t PLAYER_USER_ID_STEP = 0x100;
+constexpr int32_t MAX_USER_ID =
+    std::numeric_limits<int32_t>::max() - (MAX_PLAYERS - 1) * PLAYER_USER_ID_STEP;
+
 constexpr bool IsConfiguredUserIdValid(int32_t user_id) {
 	constexpr int32_t USER_ID_EVERYONE = 0xfe;
 	constexpr int32_t USER_ID_SYSTEM   = 0xff;
-	return user_id >= 0 && user_id != USER_ID_EVERYONE && user_id != USER_ID_SYSTEM;
+	return user_id >= 0 && user_id <= MAX_USER_ID && user_id != USER_ID_EVERYONE &&
+	       user_id != USER_ID_SYSTEM;
 }

 struct ConfigOptions {
@@ -50,6 +57,8 @@ struct ConfigOptions {
 	std::optional<ControllerColor> controller_color;
 	uint32_t               controller_speaker_volume      = 50;
 	uint32_t               controller_vibration_intensity = 100;
+	// 0 shares player 1 with the first gamepad; 1..MAX_PLAYERS dedicates that player to the keyboard.
+	uint32_t               keyboard_player                = 0;
 	PresentMode            present_mode                = PresentMode::Mailbox;
 	int32_t                gpu_index                   = -1;
 	bool                   fullscreen_enabled          = false;
@@ -76,10 +85,12 @@ struct ConfigOptions {
 	bool                   gpu_assisted_validation_enabled = false;
 	bool                   renderdoc_enabled           = false;
 	bool                   readback_linear_images      = false;
+	bool                   readback_tiled_images       = false;
 	bool                   sync_raw_image_buffers      = false;
 	bool                   tessellation_enabled        = false;
 	bool                   trophy_enabled              = true;
 	uint32_t               trophy_duration_seconds     = 6;
+	bool                   perf_overlay_enabled            = false;
 	bool                   playgo_hack_enabled         = false;
 	bool                   skip_notice_screen          = false;
 #if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
@@ -100,6 +111,13 @@ const std::string& GetAudioInputDevice();
 const std::optional<ControllerColor>& GetControllerColor();
 uint32_t GetControllerSpeakerVolume();
 uint32_t GetControllerVibrationIntensity();
+uint32_t GetKeyboardPlayer();
+// Player 0 is the configured user; later players derive their IDs from it.
+int32_t  GetPlayerUserId(int player);
+// Returns -1 when the user ID does not belong to any player.
+int      GetPlayerFromUserId(int32_t user_id);
+// Players after the first are guests: they have no profile, trophies or NP account.
+bool     IsGuestUserId(int32_t user_id);
 PresentMode GetPresentMode();
 int32_t GetGpuIndex();
 bool     FullscreenEnabled();
@@ -133,6 +151,7 @@ bool GpuAssistedValidationEnabled();

 bool RenderDocEnabled();
 bool ReadbackLinearImagesEnabled();
+bool ReadbackTiledImagesEnabled();
 bool SyncRawImageBuffersEnabled();
 bool TessellationEnabled();
 bool TrophyEnabled();
@@ -0,0 +1,580 @@
+#include "common/hostStats.h"
+
+#include "common/config.h"
+
+#include <algorithm>
+#include <chrono>
+#include <thread>
+#include <unordered_map>
+
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+#ifndef NOMINMAX
+#define NOMINMAX
+#endif
+#include <windows.h> // IWYU pragma: keep
+// clang-format off
+#include <psapi.h>
+// clang-format on
+#elif defined(__APPLE__)
+#include <CoreFoundation/CoreFoundation.h>
+#include <IOKit/IOKitLib.h>
+#include <cstring>
+#include <libproc.h>
+#include <mach/mach.h>
+#include <notify.h>
+#include <sys/resource.h>
+#include <sys/sysctl.h>
+#include <unistd.h>
+#else
+#include <filesystem>
+#include <fstream>
+#include <sstream>
+#include <sys/resource.h>
+#include <unistd.h>
+#endif
+
+namespace Common {
+
+namespace {
+
+// Cumulative counters; rates come from the difference between two samples.
+struct Counters {
+	struct Thread {
+		uint64_t    cpu_ns;
+		std::string name;
+	};
+	std::optional<uint64_t>              process_cpu_ns;
+	std::optional<uint64_t>              system_busy;
+	std::optional<uint64_t>              system_total;
+	std::optional<uint64_t>              context_switches;
+	std::optional<uint64_t>              page_faults;
+	std::optional<uint64_t>              page_ins;
+	std::optional<uint64_t>              disk_read;
+	std::optional<uint64_t>              disk_write;
+	std::unordered_map<uint64_t, Thread> threads;
+	uint64_t                             wall_ns = 0;
+};
+
+uint64_t WallNanoseconds() {
+	return static_cast<uint64_t>(std::chrono::duration_cast<std::chrono::nanoseconds>(
+	                                 std::chrono::steady_clock::now().time_since_epoch())
+	                                 .count());
+}
+
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+
+uint64_t FileTimeTo100Ns(const FILETIME& time) {
+	return (static_cast<uint64_t>(time.dwHighDateTime) << 32u) | time.dwLowDateTime;
+}
+
+struct PlatformState {};
+
+void SamplePlatform(PlatformState& /*state*/, HostStats& stats, Counters& now) {
+	const auto process = GetCurrentProcess();
+	FILETIME   creation {};
+	FILETIME   exit {};
+	FILETIME   kernel {};
+	FILETIME   user {};
+	if (GetProcessTimes(process, &creation, &exit, &kernel, &user) != 0) {
+		now.process_cpu_ns = (FileTimeTo100Ns(kernel) + FileTimeTo100Ns(user)) * 100u;
+	}
+	FILETIME idle {};
+	if (GetSystemTimes(&idle, &kernel, &user) != 0) {
+		// Kernel time includes idle time.
+		now.system_total = FileTimeTo100Ns(kernel) + FileTimeTo100Ns(user);
+		now.system_busy  = *now.system_total - FileTimeTo100Ns(idle);
+	}
+
+	PROCESS_MEMORY_COUNTERS memory {};
+	if (K32GetProcessMemoryInfo(process, &memory, sizeof(memory)) != 0) {
+		stats.process_memory   = memory.WorkingSetSize;
+		stats.process_resident = memory.WorkingSetSize;
+		now.page_faults        = memory.PageFaultCount;
+	}
+	IO_COUNTERS io {};
+	if (GetProcessIoCounters(process, &io) != 0) {
+		now.disk_read  = io.ReadTransferCount;
+		now.disk_write = io.WriteTransferCount;
+	}
+
+	MEMORYSTATUSEX system {};
+	system.dwLength = sizeof(system);
+	if (GlobalMemoryStatusEx(&system) != 0) {
+		stats.system_memory_total = system.ullTotalPhys;
+		stats.system_memory_used  = system.ullTotalPhys - system.ullAvailPhys;
+		stats.system_free         = system.ullAvailPhys;
+		stats.memory_pressure     = system.dwMemoryLoad >= 90   ? MemoryPressure::Critical
+		                            : system.dwMemoryLoad >= 80 ? MemoryPressure::Warning
+		                                                        : MemoryPressure::Normal;
+	}
+}
+
+#else
+
+void SampleRusage(Counters& now) {
+	rusage usage {};
+	if (getrusage(RUSAGE_SELF, &usage) != 0) {
+		return;
+	}
+	const auto to_ns = [](const timeval& time) {
+		return static_cast<uint64_t>(time.tv_sec) * 1000000000u +
+		       static_cast<uint64_t>(time.tv_usec) * 1000u;
+	};
+	now.process_cpu_ns   = to_ns(usage.ru_utime) + to_ns(usage.ru_stime);
+	now.context_switches = static_cast<uint64_t>(usage.ru_nvcsw + usage.ru_nivcsw);
+	now.page_faults      = static_cast<uint64_t>(usage.ru_minflt + usage.ru_majflt);
+	now.page_ins         = static_cast<uint64_t>(usage.ru_majflt);
+}
+
+#if defined(__APPLE__)
+
+template <typename T>
+std::optional<T> SysctlValue(const char* name) {
+	T      value {};
+	size_t size = sizeof(value);
+	if (sysctlbyname(name, &value, &size, nullptr, 0) != 0 || size != sizeof(value)) {
+		return std::nullopt;
+	}
+	return value;
+}
+
+struct PlatformState {
+	mach_port_t             host          = mach_host_self();
+	int                     thermal_token = -1;
+	bool                    thermal_ready = false;
+	std::string             gpu_model;
+	std::optional<uint32_t> gpu_cores;
+
+	~PlatformState() {
+		if (thermal_ready) {
+			notify_cancel(thermal_token);
+		}
+		mach_port_deallocate(mach_task_self(), host);
+	}
+};
+
+std::optional<int64_t> DictionaryNumber(CFDictionaryRef dictionary, CFStringRef key) {
+	const auto* value  = CFDictionaryGetValue(dictionary, key);
+	int64_t     number = 0;
+	if (value == nullptr || CFGetTypeID(value) != CFNumberGetTypeID() ||
+	    !CFNumberGetValue(static_cast<CFNumberRef>(value), kCFNumberSInt64Type, &number) ||
+	    number < 0) {
+		return std::nullopt;
+	}
+	return number;
+}
+
+std::string RegistryString(io_object_t service, CFStringRef key) {
+	std::string result;
+	auto*       value = IORegistryEntryCreateCFProperty(service, key, kCFAllocatorDefault, 0);
+	if (value == nullptr) {
+		return result;
+	}
+	if (CFGetTypeID(value) == CFStringGetTypeID()) {
+		char buffer[128] {};
+		if (CFStringGetCString(static_cast<CFStringRef>(value), buffer, sizeof(buffer),
+		                       kCFStringEncodingUTF8)) {
+			result = buffer;
+		}
+	} else if (CFGetTypeID(value) == CFDataGetTypeID()) {
+		// Intel Macs store the model as NUL-terminated data.
+		const auto* data   = static_cast<CFDataRef>(value);
+		const auto* bytes  = reinterpret_cast<const char*>(CFDataGetBytePtr(data));
+		const auto  length = static_cast<size_t>(CFDataGetLength(data));
+		result.assign(bytes, strnlen(bytes, length));
+	}
+	CFRelease(value);
+	return result;
+}
+
+// IOAccelerator exposes whole-device statistics; there is no per-process figure.
+void SampleGpu(PlatformState& state, HostStats& stats) {
+	io_iterator_t iterator = 0;
+	if (IOServiceGetMatchingServices(MACH_PORT_NULL, IOServiceMatching("IOAccelerator"),
+	                                 &iterator) != KERN_SUCCESS) {
+		return;
+	}
+	const auto take_max = [](std::optional<double>& target, std::optional<int64_t> value) {
+		if (value) {
+			target = std::max(target.value_or(0.0), std::min(100.0, static_cast<double>(*value)));
+		}
+	};
+	const auto add = [](std::optional<uint64_t>& target, std::optional<int64_t> value) {
+		if (value) {
+			target = target.value_or(0) + static_cast<uint64_t>(*value);
+		}
+	};
+	for (io_object_t service = IOIteratorNext(iterator); service != 0;
+	     service             = IOIteratorNext(iterator)) {
+		if (state.gpu_model.empty()) {
+			state.gpu_model = RegistryString(service, CFSTR("model"));
+			auto*   cores   = IORegistryEntryCreateCFProperty(service, CFSTR("gpu-core-count"),
+			                                                  kCFAllocatorDefault, 0);
+			int64_t count   = 0;
+			if (cores != nullptr) {
+				if (CFGetTypeID(cores) == CFNumberGetTypeID() &&
+				    CFNumberGetValue(static_cast<CFNumberRef>(cores), kCFNumberSInt64Type,
+				                     &count) &&
+				    count > 0) {
+					state.gpu_cores = static_cast<uint32_t>(count);
+				}
+				CFRelease(cores);
+			}
+		}
+		auto* statistics = IORegistryEntryCreateCFProperty(service, CFSTR("PerformanceStatistics"),
+		                                                   kCFAllocatorDefault, 0);
+		if (statistics != nullptr) {
+			if (CFGetTypeID(statistics) == CFDictionaryGetTypeID()) {
+				const auto* values = static_cast<CFDictionaryRef>(statistics);
+				take_max(stats.gpu_percent,
+				         DictionaryNumber(values, CFSTR("Device Utilization %")));
+				take_max(stats.gpu_renderer_percent,
+				         DictionaryNumber(values, CFSTR("Renderer Utilization %")));
+				take_max(stats.gpu_tiler_percent,
+				         DictionaryNumber(values, CFSTR("Tiler Utilization %")));
+				add(stats.gpu_memory_in_use,
+				    DictionaryNumber(values, CFSTR("In use system memory")));
+				add(stats.gpu_memory_allocated,
+				    DictionaryNumber(values, CFSTR("Alloc system memory")));
+				add(stats.gpu_recoveries, DictionaryNumber(values, CFSTR("recoveryCount")));
+			}
+			CFRelease(statistics);
+		}
+		IOObjectRelease(service);
+	}
+	IOObjectRelease(iterator);
+	stats.gpu_model = state.gpu_model;
+	stats.gpu_cores = state.gpu_cores;
+}
+
+void SampleThreads(HostStats& stats, Counters& now) {
+	thread_act_array_t     threads = nullptr;
+	mach_msg_type_number_t count   = 0;
+	const auto             task    = mach_task_self();
+	if (task_threads(task, &threads, &count) != KERN_SUCCESS) {
+		return;
+	}
+	stats.thread_count = count;
+	for (mach_msg_type_number_t i = 0; i < count; i++) {
+		thread_identifier_info_data_t identifier {};
+		mach_msg_type_number_t        size = THREAD_IDENTIFIER_INFO_COUNT;
+		thread_basic_info_data_t      basic {};
+		mach_msg_type_number_t        basic_size = THREAD_BASIC_INFO_COUNT;
+		if (thread_info(threads[i], THREAD_IDENTIFIER_INFO,
+		                reinterpret_cast<thread_info_t>(&identifier), &size) == KERN_SUCCESS &&
+		    thread_info(threads[i], THREAD_BASIC_INFO, reinterpret_cast<thread_info_t>(&basic),
+		                &basic_size) == KERN_SUCCESS) {
+			const auto to_ns = [](const time_value_t& time) {
+				return static_cast<uint64_t>(time.seconds) * 1000000000u +
+				       static_cast<uint64_t>(time.microseconds) * 1000u;
+			};
+			auto& thread  = now.threads[identifier.thread_id];
+			thread.cpu_ns = to_ns(basic.user_time) + to_ns(basic.system_time);
+			thread_extended_info_data_t extended {};
+			mach_msg_type_number_t      extended_size = THREAD_EXTENDED_INFO_COUNT;
+			if (thread_info(threads[i], THREAD_EXTENDED_INFO,
+			                reinterpret_cast<thread_info_t>(&extended),
+			                &extended_size) == KERN_SUCCESS) {
+				thread.name = std::string(extended.pth_name,
+				                          strnlen(extended.pth_name, sizeof(extended.pth_name)));
+			}
+		}
+		mach_port_deallocate(task, threads[i]);
+	}
+	vm_deallocate(task, reinterpret_cast<vm_address_t>(threads), count * sizeof(thread_act_t));
+}
+
+void SamplePlatform(PlatformState& state, HostStats& stats, Counters& now) {
+	SampleRusage(now);
+
+	host_cpu_load_info_data_t load {};
+	mach_msg_type_number_t    load_count = HOST_CPU_LOAD_INFO_COUNT;
+	if (host_statistics(state.host, HOST_CPU_LOAD_INFO, reinterpret_cast<host_info_t>(&load),
+	                    &load_count) == KERN_SUCCESS) {
+		const uint64_t busy = static_cast<uint64_t>(load.cpu_ticks[CPU_STATE_USER]) +
+		                      load.cpu_ticks[CPU_STATE_SYSTEM] + load.cpu_ticks[CPU_STATE_NICE];
+		now.system_busy     = busy;
+		now.system_total    = busy + load.cpu_ticks[CPU_STATE_IDLE];
+	}
+
+	// Physical footprint is what Activity Monitor shows in its Memory column.
+	task_vm_info_data_t    vm_info {};
+	mach_msg_type_number_t vm_info_count = TASK_VM_INFO_COUNT;
+	if (task_info(mach_task_self(), TASK_VM_INFO, reinterpret_cast<task_info_t>(&vm_info),
+	              &vm_info_count) == KERN_SUCCESS) {
+		stats.process_memory      = vm_info.phys_footprint;
+		stats.process_resident    = vm_info.resident_size;
+		stats.process_compressed  = vm_info.compressed;
+		stats.process_file_backed = vm_info.external;
+		stats.process_virtual     = vm_info.virtual_size;
+	}
+	rusage_info_v2 usage {};
+	if (proc_pid_rusage(getpid(), RUSAGE_INFO_V2, reinterpret_cast<rusage_info_t*>(&usage)) == 0) {
+		now.disk_read  = usage.ri_diskio_bytesread;
+		now.disk_write = usage.ri_diskio_byteswritten;
+	}
+
+	stats.system_memory_total   = SysctlValue<uint64_t>("hw.memsize");
+	vm_size_t              page = 0;
+	vm_statistics64_data_t vm {};
+	mach_msg_type_number_t count = HOST_VM_INFO64_COUNT;
+	if (host_page_size(state.host, &page) == KERN_SUCCESS &&
+	    host_statistics64(state.host, HOST_VM_INFO64, reinterpret_cast<host_info64_t>(&vm),
+	                      &count) == KERN_SUCCESS) {
+		// Activity Monitor's "Memory Used": app memory + wired + compressed.
+		const uint64_t app       = vm.internal_page_count > vm.purgeable_count
+		                               ? vm.internal_page_count - vm.purgeable_count
+		                               : 0;
+		stats.system_memory_used = (app + vm.wire_count + vm.compressor_page_count) * page;
+		stats.system_wired       = static_cast<uint64_t>(vm.wire_count) * page;
+		stats.system_compressed  = static_cast<uint64_t>(vm.compressor_page_count) * page;
+		stats.system_cached =
+		    (static_cast<uint64_t>(vm.external_page_count) + vm.purgeable_count) * page;
+		stats.system_free = static_cast<uint64_t>(vm.free_count) * page;
+	}
+
+	if (const auto swap = SysctlValue<xsw_usage>("vm.swapusage")) {
+		stats.swap_used  = swap->xsu_used;
+		stats.swap_total = swap->xsu_total;
+	}
+
+	// 1 = normal, 2 = warning, 4 = critical, as reported to dispatch memory-pressure sources.
+	if (const auto level = SysctlValue<int>("kern.memorystatus_vm_pressure_level")) {
+		stats.memory_pressure = *level >= 4   ? MemoryPressure::Critical
+		                        : *level >= 2 ? MemoryPressure::Warning
+		                                      : MemoryPressure::Normal;
+	}
+
+	// 0 nominal, 1 moderate, 2 heavy, 3 trapping, 4 sleeping.
+	if (!state.thermal_ready) {
+		state.thermal_ready = notify_register_check("com.apple.system.thermalpressurelevel",
+		                                            &state.thermal_token) == NOTIFY_STATUS_OK;
+	}
+	uint64_t thermal = 0;
+	if (state.thermal_ready &&
+	    notify_get_state(state.thermal_token, &thermal) == NOTIFY_STATUS_OK) {
+		stats.thermal_pressure = thermal >= 3   ? ThermalPressure::Critical
+		                         : thermal == 2 ? ThermalPressure::Heavy
+		                         : thermal == 1 ? ThermalPressure::Moderate
+		                                        : ThermalPressure::Nominal;
+	}
+
+	SampleThreads(stats, now);
+	SampleGpu(state, stats);
+}
+
+#else
+
+using ProcValues = std::unordered_map<std::string, uint64_t>;
+
+// Reads "<key>: <number> [kB]" lines; kB values are returned in bytes.
+ProcValues ReadProcValues(const char* path) {
+	ProcValues    values;
+	std::ifstream file(path);
+	std::string   line;
+	while (std::getline(file, line)) {
+		const auto colon = line.find(':');
+		if (colon == std::string::npos) {
+			continue;
+		}
+		std::istringstream stream(line.substr(colon + 1));
+		uint64_t           number = 0;
+		std::string        unit;
+		if (stream >> number) {
+			stream >> unit;
+			values[line.substr(0, colon)] = unit == "kB" ? number * 1024u : number;
+		}
+	}
+	return values;
+}
+
+std::optional<uint64_t> Find(const ProcValues& values, const char* key) {
+	const auto it = values.find(key);
+	return it != values.end() ? std::optional(it->second) : std::nullopt;
+}
+
+struct PlatformState {};
+
+void SampleThreads(HostStats& stats, Counters& now) {
+	static const double ns_per_tick = 1e9 / static_cast<double>(std::max(sysconf(_SC_CLK_TCK), 1L));
+	std::error_code     error;
+	uint32_t            count = 0;
+	for (std::filesystem::directory_iterator it("/proc/self/task", error), end; !error && it != end;
+	     it.increment(error)) {
+		count++;
+		std::ifstream stat_file(it->path() / "stat");
+		std::string   stat;
+		uint64_t      id = 0;
+		if (!(stat_file >> id)) {
+			continue; // The thread may have exited during enumeration.
+		}
+		std::getline(stat_file, stat);
+		const auto close = stat.rfind(')');
+		if (close == std::string::npos) {
+			continue;
+		}
+		// Fields after the command name start at field 3 (state); utime and stime are 14 and 15.
+		std::istringstream fields(stat.substr(close + 2));
+		std::string        field;
+		for (int index = 3; index < 14 && fields >> field; index++) {
+		}
+		uint64_t user   = 0;
+		uint64_t system = 0;
+		if (!(fields >> user >> system)) {
+			continue;
+		}
+		auto& thread  = now.threads[id];
+		thread.cpu_ns = static_cast<uint64_t>(static_cast<double>(user + system) * ns_per_tick);
+		std::ifstream name_file(it->path() / "comm");
+		std::getline(name_file, thread.name);
+	}
+	if (!error) {
+		stats.thread_count = count;
+	}
+}
+
+void SamplePlatform(PlatformState& /*state*/, HostStats& stats, Counters& now) {
+	SampleRusage(now);
+
+	std::ifstream cpu_file("/proc/stat");
+	std::string   label;
+	if (cpu_file >> label && label == "cpu") {
+		uint64_t total = 0;
+		uint64_t idle  = 0;
+		uint64_t value = 0;
+		for (int index = 0; index < 8 && cpu_file >> value; index++) {
+			total += value;
+			if (index == 3 || index == 4) { // idle, iowait
+				idle += value;
+			}
+		}
+		now.system_total = total;
+		now.system_busy  = total - idle;
+	}
+
+	const auto status         = ReadProcValues("/proc/self/status");
+	stats.process_memory      = Find(status, "VmRSS");
+	stats.process_resident    = Find(status, "VmRSS");
+	stats.process_file_backed = Find(status, "RssFile");
+	stats.process_compressed  = Find(status, "VmSwap");
+	stats.process_virtual     = Find(status, "VmSize");
+
+	const auto io  = ReadProcValues("/proc/self/io");
+	now.disk_read  = Find(io, "read_bytes");
+	now.disk_write = Find(io, "write_bytes");
+
+	const auto memory    = ReadProcValues("/proc/meminfo");
+	const auto total     = Find(memory, "MemTotal");
+	const auto available = Find(memory, "MemAvailable");
+	if (total && available && *available <= *total) {
+		stats.system_memory_total = total;
+		stats.system_memory_used  = *total - *available;
+	}
+	stats.system_free = Find(memory, "MemFree");
+	if (const auto cached = Find(memory, "Cached")) {
+		stats.system_cached = *cached + Find(memory, "Buffers").value_or(0);
+	}
+	const auto swap_total = Find(memory, "SwapTotal");
+	const auto swap_free  = Find(memory, "SwapFree");
+	if (swap_total && swap_free && *swap_free <= *swap_total) {
+		stats.swap_total = swap_total;
+		stats.swap_used  = *swap_total - *swap_free;
+	}
+
+	// Pressure stall information: share of the last 10 s some task waited on memory.
+	std::ifstream pressure("/proc/pressure/memory");
+	std::string   kind;
+	std::string   average;
+	double        stalled = 0.0;
+	if (pressure >> kind >> average && kind == "some" && average.rfind("avg10=", 0) == 0 &&
+	    std::istringstream(average.substr(6)) >> stalled) {
+		stats.memory_pressure = stalled >= 20.0  ? MemoryPressure::Critical
+		                        : stalled >= 5.0 ? MemoryPressure::Warning
+		                                         : MemoryPressure::Normal;
+	} else if (total && available && *total != 0) {
+		const double free_share = static_cast<double>(*available) / static_cast<double>(*total);
+		stats.memory_pressure   = free_share < 0.10   ? MemoryPressure::Critical
+		                          : free_share < 0.25 ? MemoryPressure::Warning
+		                                              : MemoryPressure::Normal;
+	}
+
+	SampleThreads(stats, now);
+}
+
+#endif
+#endif
+
+constexpr size_t BUSIEST_THREADS = 6;
+
+} // namespace
+
+struct HostStatsSampler::Impl {
+	PlatformState state;
+	Counters      last;
+	uint32_t      logical_cores = std::max(std::thread::hardware_concurrency(), 1u);
+	bool          has_last      = false;
+};
+
+HostStatsSampler::HostStatsSampler(): m_impl(std::make_unique<Impl>()) {}
+
+HostStatsSampler::~HostStatsSampler() = default;
+
+HostStats HostStatsSampler::Sample() {
+	HostStats stats;
+	stats.logical_cores = m_impl->logical_cores;
+
+	Counters now;
+	now.wall_ns = WallNanoseconds();
+	SamplePlatform(m_impl->state, stats, now);
+
+	const auto& last    = m_impl->last;
+	const auto  elapsed = static_cast<double>(now.wall_ns - last.wall_ns) / 1e9;
+	if (m_impl->has_last && elapsed > 0.0) {
+		const auto rate = [&](std::optional<uint64_t> current,
+		                      std::optional<uint64_t> previous) -> std::optional<double> {
+			if (!current || !previous || *current < *previous) {
+				return std::nullopt;
+			}
+			return static_cast<double>(*current - *previous) / elapsed;
+		};
+		if (const auto cpu = rate(now.process_cpu_ns, last.process_cpu_ns)) {
+			stats.process_cpu_percent =
+			    std::clamp(*cpu / 1e9 * 100.0 / stats.logical_cores, 0.0, 100.0);
+		}
+		if (now.system_total && last.system_total && now.system_busy && last.system_busy &&
+		    *now.system_total > *last.system_total && *now.system_busy >= *last.system_busy) {
+			stats.system_cpu_percent =
+			    std::clamp(static_cast<double>(*now.system_busy - *last.system_busy) /
+			                   static_cast<double>(*now.system_total - *last.system_total) * 100.0,
+			               0.0, 100.0);
+		}
+		stats.context_switches_per_s = rate(now.context_switches, last.context_switches);
+		stats.page_faults_per_s      = rate(now.page_faults, last.page_faults);
+		stats.page_ins_per_s         = rate(now.page_ins, last.page_ins);
+		stats.disk_read_per_s        = rate(now.disk_read, last.disk_read);
+		stats.disk_write_per_s       = rate(now.disk_write, last.disk_write);
+
+		for (const auto& [id, thread]: now.threads) {
+			// A thread missing from the last sample started during this interval.
+			const auto     previous = last.threads.find(id);
+			const uint64_t before   = previous != last.threads.end() ? previous->second.cpu_ns : 0;
+			if (thread.cpu_ns < before) {
+				continue;
+			}
+			const double percent =
+			    static_cast<double>(thread.cpu_ns - before) / 1e9 / elapsed * 100.0;
+			stats.busiest_threads.push_back(
+			    {thread.name.empty() ? "thread " + std::to_string(id) : thread.name, percent});
+		}
+		std::sort(stats.busiest_threads.begin(), stats.busiest_threads.end(),
+		          [](const auto& a, const auto& b) { return a.cpu_percent > b.cpu_percent; });
+		if (stats.busiest_threads.size() > BUSIEST_THREADS) {
+			stats.busiest_threads.resize(BUSIEST_THREADS);
+		}
+	}
+	m_impl->last     = std::move(now);
+	m_impl->has_last = true;
+	return stats;
+}
+
+} // namespace Common
@@ -0,0 +1,86 @@
+#ifndef KYTY_COMMON_HOST_STATS_H_
+#define KYTY_COMMON_HOST_STATS_H_
+
+#include <cstdint>
+#include <memory>
+#include <optional>
+#include <string>
+#include <vector>
+
+namespace Common {
+
+enum class MemoryPressure { Unknown, Normal, Warning, Critical };
+enum class ThermalPressure { Unknown, Nominal, Moderate, Heavy, Critical };
+
+struct ThreadUsage {
+	std::string name;
+	double      cpu_percent = 0.0; // Of one core.
+};
+
+// Host-wide and process-wide resource usage. Missing values are std::nullopt; rates are per
+// second since the previous sample.
+struct HostStats {
+	uint32_t logical_cores = 0;
+
+	// Processing.
+	std::optional<double>    process_cpu_percent; // Share of all logical cores, 0-100.
+	std::optional<double>    system_cpu_percent;  // Whole machine, 0-100.
+	std::optional<uint32_t>  thread_count;
+	std::vector<ThreadUsage> busiest_threads; // Busiest first.
+	std::optional<double>    context_switches_per_s;
+	std::optional<double>    page_faults_per_s;
+	std::optional<double>    page_ins_per_s; // Faults that read from disk.
+	ThermalPressure          thermal_pressure = ThermalPressure::Unknown;
+
+	// Emulator process memory.
+	std::optional<uint64_t> process_memory; // Physical footprint / resident set.
+	std::optional<uint64_t> process_resident;
+	std::optional<uint64_t> process_compressed; // Compressed or swapped out.
+	std::optional<uint64_t> process_file_backed;
+	std::optional<uint64_t> process_virtual; // Reserved address space.
+
+	// System memory.
+	std::optional<uint64_t> system_memory_used;
+	std::optional<uint64_t> system_memory_total;
+	std::optional<uint64_t> system_wired;
+	std::optional<uint64_t> system_compressed;
+	std::optional<uint64_t> system_cached; // File cache, reclaimable.
+	std::optional<uint64_t> system_free;
+	std::optional<uint64_t> swap_used;
+	std::optional<uint64_t> swap_total;
+	MemoryPressure          memory_pressure = MemoryPressure::Unknown;
+
+	// Disk I/O of the emulator process.
+	std::optional<double> disk_read_per_s;
+	std::optional<double> disk_write_per_s;
+
+	// GPU, whole device (macOS IOAccelerator).
+	std::string             gpu_model;
+	std::optional<uint32_t> gpu_cores;
+	std::optional<double>   gpu_percent;
+	std::optional<double>   gpu_renderer_percent;
+	std::optional<double>   gpu_tiler_percent;
+	std::optional<uint64_t> gpu_memory_in_use; // System memory the GPU driver has in use.
+	std::optional<uint64_t> gpu_memory_allocated;
+	std::optional<uint64_t> gpu_recoveries; // GPU resets since boot.
+};
+
+// May block for a few milliseconds (thread enumeration, OS registry reads), so call it off the
+// render thread. The first sample has no rates.
+class HostStatsSampler final {
+public:
+	HostStatsSampler();
+	~HostStatsSampler();
+	HostStatsSampler(const HostStatsSampler&)            = delete;
+	HostStatsSampler& operator=(const HostStatsSampler&) = delete;
+
+	[[nodiscard]] HostStats Sample();
+
+private:
+	struct Impl;
+	std::unique_ptr<Impl> m_impl;
+};
+
+} // namespace Common
+
+#endif /* KYTY_COMMON_HOST_STATS_H_ */
@@ -0,0 +1,53 @@
+#include "common/localization.h"
+
+#include <array>
+#include <cstddef>
+#include <nlohmann/json.hpp>
+#include <string>
+
+#include "localizationData.h"
+
+namespace Common::Localization {
+namespace {
+
+constexpr std::array<std::string_view, 30> Locales = {
+    "ja-JP", "en-US", "fr-FR",   "es-ES",   "de-DE", "it-IT", "nl-NL", "pt-PT",
+    "ru-RU", "ko-KR", "zh-Hant", "zh-Hans", "fi-FI", "sv-SE", "da-DK", "no-NO",
+    "pl-PL", "pt-BR", "en-GB",   "tr-TR",   "es-419", "ar-AE", "fr-CA", "cs-CZ",
+    "hu-HU", "el-GR", "ro-RO",   "th-TH",   "vi-VN", "id-ID"};
+
+size_t LanguageIndex(int console_language) {
+	return console_language >= 0 && console_language < static_cast<int>(Locales.size())
+	           ? static_cast<size_t>(console_language)
+	           : 1;
+}
+
+} // namespace
+
+std::string_view LocaleName(int console_language) {
+	return Locales[LanguageIndex(console_language)];
+}
+
+std::string_view Get(int console_language, std::string_view key) {
+	static const auto catalogs = [] {
+		std::array<nlohmann::json, Locales.size()> result;
+		for (size_t i = 0; i < Locales.size(); ++i) {
+			const auto bytes = EmbeddedCatalog(Locales[i]);
+			result[i] = nlohmann::json::parse(bytes.begin(), bytes.end(), nullptr, false);
+		}
+		return result;
+	}();
+	for (const auto index: std::array {LanguageIndex(console_language), size_t {1}}) {
+		const auto& catalog = catalogs[index];
+		const auto  found   = catalog.find(key);
+		if (found != catalog.end() && found->is_string()) {
+			const auto& text = found->get_ref<const std::string&>();
+			if (!text.empty()) {
+				return text;
+			}
+		}
+	}
+	return key;
+}
+
+} // namespace Common::Localization
@@ -0,0 +1,13 @@
+#pragma once
+
+#include <string_view>
+
+namespace Common::Localization {
+
+// Console language indices follow the system configuration; invalid indices use English.
+[[nodiscard]] std::string_view LocaleName(int console_language);
+
+// Translations have static lifetime. Missing translations use English, then the supplied key.
+[[nodiscard]] std::string_view Get(int console_language, std::string_view key);
+
+} // namespace Common::Localization
@@ -5,6 +5,7 @@
 #include "common/emulatorConfig.h"

 #include <cstdio>
+#include <cstdlib>
 #include <filesystem>
 #include <fmt/format.h>
 #include <memory>
@@ -45,8 +46,8 @@ std::shared_ptr<spdlog::logger> MakeFileLogger(std::string                  name
 		std::filesystem::create_directories(parent);
 	}

-	auto sink =
-	    std::make_shared<spdlog::sinks::basic_file_sink_mt>(path.native(), true);
+	const bool truncate = std::getenv(Log::RelaunchedVariable) == nullptr;
+	auto       sink = std::make_shared<spdlog::sinks::basic_file_sink_mt>(path.native(), truncate);
 	return MakeLogger(std::move(name), std::move(sink));
 }

@@ -21,6 +21,10 @@ struct Lifecycle {

 enum class Direction { Silent, Console, File };

+// Set by LoadExec for the program it starts, which continues the log file of the one it replaces
+// instead of starting it over.
+inline constexpr const char* RelaunchedVariable = "KYTY_RELAUNCHED";
+
 Direction GetDirection();
 bool      IsSilent();
 void      Write(std::string_view text);
@@ -47,20 +51,26 @@ inline constexpr auto BrightWhite   = fmt::fg(fmt::terminal_color::bright_white)

 } // namespace Color

+// Out of line, so callers on small guest stacks (fibers) do not reserve the formatting frame.
+template <typename... Args>
+[[gnu::noinline]] void Printf(fmt::text_style style, const char* format, const Args&... args) {
+	Write(style, fmt::sprintf(format, args...));
+}
+
 } // namespace Log

 // NOLINTNEXTLINE(cppcoreguidelines-macro-usage)
 #define LOGF(...)                                                                                  \
 	do {                                                                                           \
 		if (!::Log::IsSilent()) {                                                                  \
-			::Log::Write(::fmt::sprintf(__VA_ARGS__));                                             \
+			::Log::Printf({}, __VA_ARGS__);                                                        \
 		}                                                                                          \
 	} while (false)
 // NOLINTNEXTLINE(cppcoreguidelines-macro-usage)
 #define LOGF_COLOR(style, ...)                                                                     \
 	do {                                                                                           \
 		if (!::Log::IsSilent()) {                                                                  \
-			::Log::Write((style), ::fmt::sprintf(__VA_ARGS__));                                    \
+			::Log::Printf((style), __VA_ARGS__);                                                   \
 		}                                                                                          \
 	} while (false)

@@ -367,7 +367,7 @@ bool SysFileDeleteFile(const std::filesystem::path& name) {

 bool SysFileFlush(sys_file_t& f) {
 	if (f.type == SYS_FILE_FILE && f.f != nullptr) {
-		return (fflush(f.f) == 0);
+		return fflush(f.f) == 0 && fsync(fileno(f.f)) == 0;
 	}

 	return false;
@@ -1,9 +1,9 @@
 #include "common/trophies.h"

 #include "common/file.h"
+#include "common/localization.h"

 #include <algorithm>
-#include <array>
 #include <charconv>
 #include <fmt/format.h>
 #include <nlohmann/json.hpp>
@@ -105,13 +105,7 @@ Package ParsePackage(std::span<const std::byte> data, int console_language) {
 	    !conf["defaultLanguage"].is_string()) {
 		return {};
 	}
-	static constexpr std::array<std::string_view, 30> locales = {
-	    "ja-JP", "en-US", "fr-FR",   "es-ES",   "de-DE",  "it-IT", "nl-NL", "pt-PT",
-	    "ru-RU", "ko-KR", "zh-Hant", "zh-Hans", "fi-FI",  "sv-SE", "da-DK", "no-NO",
-	    "pl-PL", "pt-BR", "en-GB",   "tr-TR",   "es-419", "ar-AE", "fr-CA", "cs-CZ",
-	    "hu-HU", "el-GR", "ro-RO",   "th-TH",   "vi-VN",  "id-ID"};
-	const auto locale =
-	    locales[console_language >= 0 && console_language < locales.size() ? console_language : 1];
+	const auto locale = Localization::LocaleName(console_language);
 	auto meta = ReadJson(files, fmt::format("tropmeta_{}.json", locale));
 	if (meta.is_null()) {
 		meta = ReadJson(files, fmt::format("tropmeta_{}.json",
@@ -0,0 +1,62 @@
+if(NOT DEFINED INPUTS OR NOT DEFINED OUTPUT)
+	message(FATAL_ERROR "embed_localization.cmake requires INPUTS and OUTPUT")
+endif()
+
+set(contents "// Generated from assets/localization; do not edit.\n#pragma once\n#include <span>\n#include <string_view>\n\nnamespace {\n")
+set(lookup "std::span<const unsigned char> EmbeddedCatalog(std::string_view locale) {\n")
+set(continuation "[89ab][0-9a-f]")
+string(CONCAT utf8_character
+	"[0-7][0-9a-f]|(c[2-9a-f]|d[0-9a-f])${continuation}"
+	"|e0[ab][0-9a-f]${continuation}|e[1-9abcef]${continuation}${continuation}"
+	"|ed[89][0-9a-f]${continuation}|f0[9ab][0-9a-f]${continuation}${continuation}"
+	"|f[123]${continuation}${continuation}${continuation}|f48[0-9a-f]${continuation}${continuation}")
+foreach(catalog IN LISTS INPUTS)
+	file(READ "${catalog}" bytes HEX)
+	# Validate UTF-8, excluding overlong sequences, surrogates, and values above U+10FFFF.
+	if(NOT bytes MATCHES "^(${utf8_character})*$")
+		message(FATAL_ERROR "Localization must be UTF-8: ${catalog}")
+	endif()
+	file(READ "${catalog}" json)
+	string(JSON type ERROR_VARIABLE error TYPE "${json}")
+	if(error OR NOT type STREQUAL "OBJECT")
+		message(FATAL_ERROR "Invalid localization object in ${catalog}: ${error}")
+	endif()
+	string(JSON count LENGTH "${json}")
+	if(count GREATER 0)
+		math(EXPR last "${count} - 1")
+		foreach(index RANGE ${last})
+			string(JSON key MEMBER "${json}" ${index})
+			string(JSON type TYPE "${json}" "${key}")
+			if(NOT type STREQUAL "STRING")
+				message(FATAL_ERROR "Localization value '${key}' must be a string: ${catalog}")
+			endif()
+			string(JSON text GET "${json}" "${key}")
+			if(key STREQUAL "" OR text STREQUAL "")
+				message(FATAL_ERROR "Localization keys and values must not be empty: ${catalog}")
+			endif()
+		endforeach()
+	endif()
+	get_filename_component(locale "${catalog}" NAME_WE)
+	if(locale STREQUAL "en-US")
+		string(JSON earned ERROR_VARIABLE error GET "${json}" "trophy.toast.earned")
+		if(error OR earned STREQUAL "")
+			message(FATAL_ERROR "English localization requires trophy.toast.earned: ${catalog}")
+		endif()
+		set(english_found TRUE)
+	endif()
+	string(MAKE_C_IDENTIFIER "${locale}" symbol)
+	string(REGEX REPLACE "(..)" "0x\\1," bytes "${bytes}")
+	string(APPEND contents "constexpr unsigned char catalog_${symbol}[] = {${bytes}};\n")
+	string(APPEND lookup "\tif (locale == \"${locale}\") return catalog_${symbol};\n")
+endforeach()
+if(NOT english_found)
+	message(FATAL_ERROR "Localization requires an en-US.json fallback")
+endif()
+string(APPEND contents "\n${lookup}\treturn {};\n}\n} // namespace\n")
+
+get_filename_component(output_dir "${OUTPUT}" DIRECTORY)
+file(MAKE_DIRECTORY "${output_dir}")
+file(WRITE "${OUTPUT}.tmp" "${contents}")
+execute_process(COMMAND "${CMAKE_COMMAND}" -E copy_if_different "${OUTPUT}.tmp" "${OUTPUT}"
+	COMMAND_ERROR_IS_FATAL ANY)
+file(REMOVE "${OUTPUT}.tmp")
@@ -89,42 +89,6 @@ static void MountSandboxDirs() {
 	MountOrCreateDir("_TempData/" + title_id, "/temp");
 }

-static bool ClearDirectoryContents(const std::filesystem::path& dir) {
-	bool ok = true;
-
-	for (const auto& entry: Common::File::GetDirEntries(dir)) {
-		if (entry.name == "." || entry.name == "..") {
-			continue;
-		}
-
-		auto path = dir / entry.name;
-
-		if (entry.is_file) {
-			Common::File::RemoveReadonly(path);
-			ok = Common::File::DeleteFile(path) && ok;
-		} else {
-			ok = ClearDirectoryContents(path) && ok;
-			ok = Common::File::DeleteDirectory(path) && ok;
-		}
-	}
-
-	return ok;
-}
-
-static void ClearDebugTextureFolder() {
-	const std::string debug_texture_folder = "_Textures";
-
-	if (!Common::File::IsDirectoryExisting(debug_texture_folder)) {
-		Common::File::CreateDirectories(debug_texture_folder);
-		return;
-	}
-
-	if (!ClearDirectoryContents(debug_texture_folder)) {
-		LOGF_COLOR(Log::Color::BrightYellow, "TextureDump: failed to completely clear %s\n",
-		           debug_texture_folder.c_str());
-	}
-}
-
 static void Init(const Config::ConfigOptions& cfg, const std::filesystem::path& param_json,
                 Common::Subsystems& subsystems) {
 	EXIT_IF(!Common::Thread::IsMainThread());
@@ -198,8 +162,6 @@ void Run(const RunOptions& options) {
 	Common::Subsystems subsystems(true);
 	Init(options.config, param_json, subsystems);

-	ClearDebugTextureFolder();
-
 	PrintSystemInfo();
 	std::string title_id;
 	Loader::SystemContentParamSfoGetString("TITLE_ID", &title_id);
@@ -104,7 +104,6 @@ public:
 	                       uint32_t value);
 	void PrepareCpuFlip(uint64_t request_id);
 	void SynchronizeGpu();
-	void EmitGlobalBarrier();
 	void TriggerEopEventAtEndOfPipe(uint32_t interrupt_context_id);
 	void DispatchDirect(uint32_t thread_group_x, uint32_t thread_group_y, uint32_t thread_group_z,
 	                    uint32_t mode);
@@ -155,6 +154,8 @@ private:
 	                      uint32_t interrupt_context_id);
 	void ProcessPm4(Pm4Execution& execution);
 	void SuspendPm4();
+	uint32_t            NumInstances();
+	const void*         IndexAddress(uint32_t first_index) const;
 	CommandScheduler&   GetScheduler() const { return m_renderer.GetCommandScheduler(); }
 	CommandBuffer&      CurrentBuffer() { return GetScheduler().Current(); }

@@ -170,8 +171,10 @@ private:
 	uint64_t         m_index_base_addr                  = 0;
 	uint64_t         m_draw_indirect_args_base_addr     = 0;
 	uint64_t         m_dispatch_indirect_args_base_addr = 0;
-	// Persistent draw state: indirect draws update it for subsequent draws.
-	uint32_t m_num_instances = 1;
+	// Persistent draw state. Indirect draws leave the address of their last instance count,
+	// read only when a later draw inherits it.
+	uint32_t m_num_instances      = 1;
+	uint64_t m_num_instances_addr = 0;

 	uint32_t m_de_count    = 0;
 	uint32_t m_ce_count    = 0;
@@ -73,7 +73,6 @@ constexpr auto MakeContextDispatchTable() {
 	g_hw_ctx_func[Pm4::PA_SC_BINNER_CNTL_0]                   = HwCtxSetPaScExtendedControl;
 	g_hw_ctx_func[Pm4::PA_SC_BINNER_CNTL_1]                   = HwCtxSetPaScExtendedControl;
 	g_hw_ctx_func[Pm4::PA_SC_CONSERVATIVE_RASTERIZATION_CNTL] = HwCtxSetPaScExtendedControl;
-	g_hw_ctx_func[Pm4::DB_ALPHA_TO_MASK]                      = HwCtxSetAlphaToMask;
 	g_hw_ctx_func[Pm4::VGT_DRAW_PAYLOAD_CNTL]                 = HwCtxSetDrawPayloadControl;
 	g_hw_ctx_func[Pm4::VGT_PRIMITIVEID_RESET]                 = HwCtxSetPrimitiveIdReset;
 	g_hw_ctx_func[Pm4::PA_CL_OBJPRIM_ID_CNTL]                 = HwCtxSetObjprimIdControl;
@@ -79,7 +79,6 @@ uint32_t HwCtxSetAaMask(CommandProcessor&, uint32_t, uint32_t, const uint32_t*,
 uint32_t HwCtxSetVtxControl(CommandProcessor&, uint32_t, uint32_t, const uint32_t*, uint32_t);
 uint32_t HwCtxSetPaScExtendedControl(CommandProcessor&, uint32_t, uint32_t, const uint32_t*,
                                     uint32_t);
-uint32_t HwCtxSetAlphaToMask(CommandProcessor&, uint32_t, uint32_t, const uint32_t*, uint32_t);
 uint32_t HwCtxSetDrawPayloadControl(CommandProcessor&, uint32_t, uint32_t, const uint32_t*,
                                    uint32_t);
 uint32_t HwCtxSetPrimitiveIdReset(CommandProcessor&, uint32_t, uint32_t, const uint32_t*, uint32_t);
@@ -71,12 +71,6 @@ constexpr uint32_t NormalizeRegisterOffset(uint32_t raw_offset) {
 	return raw_offset & ~RegisterSelectorMask;
 }

-bool ReleaseMemGcrNeedsBarrier(uint32_t eop_event_type, uint32_t gcr_cntl) {
-	return eop_event_type != 0x28u ||
-	       (gcr_cntl & (GcrGl2MetadataInvalidate | GcrGl0VectorInvalidate | GcrGl1Invalidate |
-	                    GcrGl2Invalidate | GcrGl2Writeback)) != 0;
-}
-
 uint32_t ReleaseMemCacheActionFromGcr(uint32_t gcr_cntl) {
 	return ((gcr_cntl & GcrGl2Writeback) != 0 ? 0x38u : 0x00u);
 }
@@ -289,12 +283,6 @@ KYTY_HW_CTX_PARSER(HwCtxSetCentroidPriority) {
 static void HwCtxIgnoreAaMaskRegister([[maybe_unused]] uint32_t cmd_offset,
                                      [[maybe_unused]] uint32_t value) {}

-static void HwCtxIgnoreAlphaToMaskRegister([[maybe_unused]] uint32_t value) {}
-
-static void HwCtxIgnoreDisabledUserClipPlane(CommandProcessor& cp, uint32_t value) {
-	EXIT_NOT_IMPLEMENTED(value != 0 || cp.GetCtx().GetClipControl().user_clip_planes != 0);
-}
-
 static void HwCtxIgnorePrimitiveIdReset([[maybe_unused]] uint32_t value) {}

 static void HwCtxIgnoreVtxControl([[maybe_unused]] uint32_t value) {}
@@ -439,16 +427,6 @@ KYTY_HW_CTX_PARSER(HwCtxSetPointState) {
 	return num_values;
 }

-KYTY_HW_CTX_PARSER(HwCtxSetAlphaToMask) {
-	auto num_values = KYTY_PM4_LEN(cmd_id) - 2u;
-
-	for (uint32_t i = 0; i < num_values; i++) {
-		HwCtxIgnoreAlphaToMaskRegister(buffer[i]);
-	}
-
-	return num_values;
-}
-
 KYTY_HW_CTX_PARSER(HwCtxSetDrawPayloadControl) {
 	cp.GetCtx().SetDrawPayloadControl(buffer[0]);
 	return KYTY_PM4_LEN(cmd_id) - 2u;
@@ -2263,30 +2241,14 @@ KYTY_CP_OP_PARSER(CpOpReleaseMem) {
 		}
 	};

-	if (data_sel == 0 || interrupt_selector == 4) {
-		if (eop_event_type != 0x28 || gcr_cntl != 0) {
-			cp.EmitGlobalBarrier();
-		}
-
+	// Cache actions only order GPU work, which the renderer's resource barriers already order.
+	if (data_sel == 0 || interrupt_selector == 4 ||
+	    (release_dst == ReleaseMemDstMemory && dst_gpu_addr == nullptr)) {
 		trigger_interrupt();

 		return 7;
 	}

-	if (release_dst == ReleaseMemDstMemory && dst_gpu_addr == nullptr) {
-		if (eop_event_type != 0x28 || gcr_cntl != 0) {
-			cp.EmitGlobalBarrier();
-		}
-
-		trigger_interrupt();
-
-		return 7;
-	}
-
-	if (ReleaseMemGcrNeedsBarrier(eop_event_type, gcr_cntl)) {
-		cp.EmitGlobalBarrier();
-	}
-
 	auto cache_action = ReleaseMemCacheActionFromGcr(gcr_cntl);
 	cache_policy      = 0;

@@ -2899,7 +2861,7 @@ void GraphicsInitJmpTablesCxIndirect() {

 	for (auto cmd_offset = Pm4::PA_CL_UCP_0_X; cmd_offset <= Pm4::PA_CL_UCP_5_W; cmd_offset++) {
 		g_hw_ctx_indirect_func[cmd_offset] = [](KYTY_HW_CTX_INDIRECT_ARGS) {
-			HwCtxIgnoreDisabledUserClipPlane(cp, value);
+			cp.GetCtx().SetUserClipPlane(cmd_offset - Pm4::PA_CL_UCP_0_X, value);
 		};
 	}

@@ -3491,7 +3453,8 @@ void GraphicsInitJmpTablesCxIndirect() {
 	g_hw_ctx_indirect_func[Pm4::PA_SC_CONSERVATIVE_RASTERIZATION_CNTL] =
 	    [](KYTY_HW_CTX_INDIRECT_ARGS) { HwCtxIgnorePaScExtendedControl(cmd_offset, value); };
 	g_hw_ctx_indirect_func[Pm4::DB_ALPHA_TO_MASK] = [](KYTY_HW_CTX_INDIRECT_ARGS) {
-		HwCtxIgnoreAlphaToMaskRegister(value);
+		cp.GetCtx().SetAlphaToMaskEnable(
+		    KYTY_PM4_GET(value, DB_ALPHA_TO_MASK, ALPHA_TO_MASK_ENABLE) != 0);
 	};
 	g_hw_ctx_indirect_func[Pm4::CB_COLOR_CONTROL] = [](KYTY_HW_CTX_INDIRECT_ARGS) {
 		cp.GetCtx().SetColorControl(DecodeColorControl(value));
@@ -1,6 +1,7 @@
 #include "graphics/guest_gpu/graphicsRun.h"

 #include "common/assert.h"
+#include "common/debugCounters.h"
 #include "common/emulatorConfig.h"
 #include "common/logging/log.h"
 #include "common/profiler.h"
@@ -29,6 +30,7 @@
 #include <mutex>
 #include <semaphore>
 #include <thread>
+#include <utility>
 #include <vector>

 namespace Libs::Graphics {
@@ -38,21 +40,6 @@ static thread_local Pm4Execution*     g_current_execution = nullptr;
 static thread_local bool              g_gpu_thread        = false;
 static thread_local GuestGpu*         g_gpu_state         = nullptr;

-struct DrawIndirectArgs {
-	uint32_t vertex_count_per_instance;
-	uint32_t instance_count;
-	uint32_t start_vertex_location;
-	uint32_t start_instance_location;
-};
-
-struct DrawIndexedIndirectArgs {
-	uint32_t index_count_per_instance;
-	uint32_t instance_count;
-	uint32_t start_index_location;
-	uint32_t base_vertex_location;
-	uint32_t start_instance_location;
-};
-
 static bool GraphicsRunDebugDumpEnabled() {
 	return Config::GraphicsDebugDumpEnabled() &&
 	       Config::GetPrintfDirection() != Config::LogDirection::Silent;
@@ -105,6 +92,12 @@ void GuestGpu::SendCommand(Common::UniqueFunction<void>&& command) {
 	m_work_available.Signal();
 }

+void GuestGpu::Wake() {
+	Common::LockGuard lock(m_queue_mutex);
+	m_wake_pending = true;
+	m_work_available.Signal();
+}
+
 void GuestGpu::ProcessCommands() {
 	EXIT_IF(!IsGpuThread());
 	while (m_pending_commands.load(std::memory_order_acquire) != 0) {
@@ -205,6 +198,7 @@ void CommandProcessor::Reset() {
 	m_index_buffer_size                = 0;
 	m_index_base_addr                  = 0;
 	m_num_instances                    = 1;
+	m_num_instances_addr               = 0;
 	m_predicate_skip                   = false;
 	m_user_data_marker                 = HW::UserSgprType::Unknown;
 	m_draw_indirect_args_base_addr     = 0;
@@ -342,12 +336,11 @@ void CommandProcessor::WriteData(uint32_t* dst, const uint32_t* src, uint32_t dw
 		return;
 	}

-	if (write_one_address) {
-		for (uint32_t i = 0; i < dw_num; i++) {
-			dst[0] = src[i];
-		}
-	} else {
-		memcpy(dst, src, static_cast<size_t>(dw_num) * sizeof(uint32_t));
+	// Completion labels can be reused immediately; memcpy may write a destination twice.
+	auto* output = static_cast<volatile uint32_t*>(dst);
+	for (uint32_t i = 0; i < dw_num; i++) {
+		*output = src[i];
+		output += !write_one_address;
 	}
 }

@@ -478,7 +471,10 @@ void GuestGpu::ThreadRun(void* data) {
 				}
 				if (selected_queue < 0) {
 					gpu->m_processing = false;
-					gpu->m_work_available.WaitFor(&gpu->m_queue_mutex, 100);
+					// A wake that came before this wait must not be lost.
+					if (!std::exchange(gpu->m_wake_pending, false)) {
+						gpu->m_work_available.WaitFor(&gpu->m_queue_mutex, 100);
+					}
 					for (auto& queue: gpu->m_queues) {
 						if (!queue.empty()) {
 							queue.front().blocked = false;
@@ -607,7 +603,6 @@ bool GuestGpu::Process(Submission& submission) {
 		case SubmissionType::SuspendPoint:
 			// Age caches at the frame boundary, not once per command list.
 			m_renderer.RunGarbageCollector();
-			cp.EmitGlobalBarrier();
 			m_renderer.GetCommandScheduler().DeferPriorityOperation(
 			    [ready = m_suspend_point_ready] { ready->release(); });
 			cp.BufferFlush();
@@ -785,7 +780,16 @@ void CommandProcessor::SetNumInstances(uint32_t num_instances) {
 		num_instances = 1;
 	}

-	m_num_instances = num_instances;
+	m_num_instances      = num_instances;
+	m_num_instances_addr = 0;
+}
+
+uint32_t CommandProcessor::NumInstances() {
+	if (m_num_instances_addr != 0) {
+		m_num_instances      = *reinterpret_cast<const uint32_t*>(m_num_instances_addr);
+		m_num_instances_addr = 0;
+	}
+	return m_num_instances;
 }

 void CommandProcessor::SetPredication(uint32_t condition, uint32_t op, uint32_t wait_op,
@@ -842,16 +846,12 @@ void CommandProcessor::SetPredication(uint32_t condition, uint32_t op, uint32_t
 void CommandProcessor::DrawIndex(DrawIndexArgs args) {
 	args.index_type_and_size = m_index_type_and_size;
 	if (args.instance_count == 0) {
-		args.instance_count = m_num_instances;
-	}
-	if (GraphicsRunDebugDumpEnabled() && (args.base_vertex != 0 || args.first_instance != 0)) {
-		LOGF("\t draw indexed offsets: base_vertex = %" PRId32 ", first_instance = %" PRIu32 "\n",
-		     args.base_vertex, args.first_instance);
+		args.instance_count = NumInstances();
 	}
 	m_renderer.GetRenderExecutor().DrawIndex(m_submit_id, CurrentBuffer(), args);
 }

-void CommandProcessor::DrawIndexOffset(uint32_t index_offset, uint32_t index_count) {
+const void* CommandProcessor::IndexAddress(uint32_t first_index) const {
 	uint64_t index_size = 0;
 	switch (m_index_type_and_size) {
 		case 0: index_size = 2; break;
@@ -859,11 +859,11 @@ void CommandProcessor::DrawIndexOffset(uint32_t index_offset, uint32_t index_cou
 		case 2: index_size = 1; break;
 		default: EXIT("unknown index_type_and_size: %u\n", m_index_type_and_size);
 	}
+	return reinterpret_cast<const void*>(m_index_base_addr + first_index * index_size);
+}

-	auto* index_addr = reinterpret_cast<const void*>(
-	    m_index_base_addr + static_cast<uint64_t>(index_offset) * index_size);
-
-	DrawIndex({.index_count = index_count, .index_addr = index_addr});
+void CommandProcessor::DrawIndexOffset(uint32_t index_offset, uint32_t index_count) {
+	DrawIndex({.index_count = index_count, .index_addr = IndexAddress(index_offset)});
 }

 static void PatchIndirectDrawOffsets(CommandProcessor& cp, IndirectDrawRegisters registers,
@@ -884,7 +884,8 @@ static void PatchIndirectDrawOffsets(CommandProcessor& cp, IndirectDrawRegisters

 void CommandProcessor::DrawIndirect(uint32_t data_offset, IndirectDrawRegisters registers,
                                    uint32_t draw_initiator, bool indexed) {
-	const auto args_size = indexed ? sizeof(DrawIndexedIndirectArgs) : sizeof(DrawIndirectArgs);
+	const auto args_size =
+	    indexed ? sizeof(vk::DrawIndexedIndirectCommand) : sizeof(vk::DrawIndirectCommand);
 	DrawIndirectMulti(data_offset, 1, nullptr, args_size, registers, draw_initiator, indexed);
 }

@@ -894,76 +895,65 @@ void CommandProcessor::DrawIndirectMulti(uint32_t data_offset, uint32_t max_coun
                                         uint32_t draw_initiator, bool indexed) {
 	EXIT_NOT_IMPLEMENTED((draw_initiator & ~0x20u) != (indexed ? 0u : 2u));
 	EXIT_NOT_IMPLEMENTED(m_draw_indirect_args_base_addr == 0);
+	// The guest records match Vulkan's indirect commands byte for byte.
+	EXIT_NOT_IMPLEMENTED((stride_in_bytes & 3u) != 0 ||
+	                     stride_in_bytes < (indexed ? sizeof(vk::DrawIndexedIndirectCommand)
+	                                                : sizeof(vk::DrawIndirectCommand)));

-	uint32_t draw_count = max_count_or_count;
-	if (count_addr != nullptr) {
-		draw_count = *count_addr;
-		if (draw_count > max_count_or_count) {
-			draw_count = max_count_or_count;
-		}
-	}
-
+	const auto draw_count = count_addr == nullptr
+	                            ? max_count_or_count
+	                            : std::min(uint32_t {*count_addr}, max_count_or_count);
 	if (draw_count == 0) {
 		return;
 	}
+	Common::DebugCounters::Add(Common::DebugCounters::Counter::IndirectDraws, draw_count);

-	const auto args_size = indexed ? sizeof(DrawIndexedIndirectArgs) : sizeof(DrawIndirectArgs);
-	EXIT_NOT_IMPLEMENTED(stride_in_bytes < args_size);
-
-	uint64_t index_size = 0;
-	if (indexed) {
-		switch (m_index_type_and_size) {
-			case 0: index_size = 2; break;
-			case 1: index_size = 4; break;
-			case 2: index_size = 1; break;
-			default: EXIT("unknown index_type_and_size: %u\n", m_index_type_and_size);
-		}
+	auto&      executor  = m_renderer.GetRenderExecutor();
+	const auto args_addr = m_draw_indirect_args_base_addr + data_offset;
+	// NUM_INSTANCES keeps the last draw's instance count.
+	m_num_instances_addr = args_addr + static_cast<uint64_t>(draw_count - 1) * stride_in_bytes +
+	                       offsetof(vk::DrawIndirectCommand, instanceCount);
+	// With no SGPRs to patch, the GPU reads the records if the shaders allow. Shaders apply
+	// the records' offsets only through those SGPRs; Vulkan applies them, which matches
+	// while they are zero.
+	const DrawIndirectArgs records {args_addr, draw_count, stride_in_bytes};
+	if (registers.vertex_offset == Pm4::SH_NOP && registers.instance_offset == Pm4::SH_NOP &&
+	    registers.index_offset == Pm4::SH_NOP && (!indexed || m_index_buffer_size != 0) &&
+	    (indexed ? executor.DrawIndex(m_submit_id, CurrentBuffer(),
+	                                  {.index_count         = m_index_buffer_size,
+	                                   .index_addr          = IndexAddress(0),
+	                                   .index_type_and_size = m_index_type_and_size,
+	                                   .indirect            = records})
+	             : executor.DrawAuto(m_submit_id, CurrentBuffer(), {.indirect = records}))) {
+		return;
 	}

+	const uint32_t max_indices = m_index_buffer_size != 0 ? m_index_buffer_size : UINT32_MAX;
 	for (uint32_t i = 0; i < draw_count; i++) {
-		const auto args_addr = m_draw_indirect_args_base_addr + data_offset +
-		                       static_cast<uint64_t>(i) * stride_in_bytes;
-
+		const auto addr = args_addr + static_cast<uint64_t>(i) * stride_in_bytes;
 		if (!indexed) {
-			auto* args = reinterpret_cast<const DrawIndirectArgs*>(args_addr);
-			PatchIndirectDrawOffsets(*this, registers, args->start_vertex_location,
-			                         args->start_instance_location, 0);
-			m_num_instances = args->instance_count;
-			DrawIndexAuto({.vertex_count   = args->vertex_count_per_instance,
-			               .instance_count = args->instance_count,
-			               .first_vertex   = args->start_vertex_location,
-			               .first_instance = args->start_instance_location,
-			               .offset_source  = DrawOffsetSource::IndirectArgs});
+			const auto& args = *reinterpret_cast<const vk::DrawIndirectCommand*>(addr);
+			PatchIndirectDrawOffsets(*this, registers, args.firstVertex, args.firstInstance, 0);
+			executor.DrawAuto(m_submit_id, CurrentBuffer(),
+			                  {.vertex_count   = args.vertexCount,
+			                   .instance_count = args.instanceCount,
+			                   .first_vertex   = args.firstVertex,
+			                   .first_instance = args.firstInstance,
+			                   .offset_source  = DrawOffsetSource::IndirectArgs});
 			continue;
 		}

-		auto* args = reinterpret_cast<const DrawIndexedIndirectArgs*>(args_addr);
-		PatchIndirectDrawOffsets(*this, registers, args->base_vertex_location,
-		                         args->start_instance_location, args->start_index_location);
-
-		auto* index_addr = reinterpret_cast<const void*>(
-		    m_index_base_addr + static_cast<uint64_t>(args->start_index_location) * index_size);
-
-		const uint32_t index_count =
-		    (m_index_buffer_size != 0
-		         ? std::min(args->index_count_per_instance, m_index_buffer_size)
-		         : args->index_count_per_instance);
-		if (GraphicsRunDebugDumpEnabled() && index_count != args->index_count_per_instance) {
-			static std::atomic<uint32_t> log_count {0};
-			if (log_count.fetch_add(1, std::memory_order_relaxed) < 64) {
-				LOGF("\t DrawIndexIndirectMulti: clamped index_count from %" PRIu32 " to %" PRIu32
-				     " using INDEX_BUFFER_SIZE\n",
-				     args->index_count_per_instance, index_count);
-			}
-		}
-
-		m_num_instances = args->instance_count;
-		DrawIndex({.index_count    = index_count,
-		           .index_addr     = index_addr,
-		           .instance_count = args->instance_count,
-		           .base_vertex    = static_cast<int32_t>(args->base_vertex_location),
-		           .first_instance = args->start_instance_location,
-		           .offset_source  = DrawOffsetSource::IndirectArgs});
+		const auto& args = *reinterpret_cast<const vk::DrawIndexedIndirectCommand*>(addr);
+		PatchIndirectDrawOffsets(*this, registers, static_cast<uint32_t>(args.vertexOffset),
+		                         args.firstInstance, args.firstIndex);
+		executor.DrawIndex(m_submit_id, CurrentBuffer(),
+		                   {.index_count         = std::min(args.indexCount, max_indices),
+		                    .index_addr          = IndexAddress(args.firstIndex),
+		                    .instance_count      = args.instanceCount,
+		                    .index_type_and_size = m_index_type_and_size,
+		                    .base_vertex         = args.vertexOffset,
+		                    .first_instance      = args.firstInstance,
+		                    .offset_source       = DrawOffsetSource::IndirectArgs});
 	}
 }

@@ -1012,16 +1002,17 @@ void CommandProcessor::DispatchIndirect(uint64_t args_addr, uint32_t mode) {

 void CommandProcessor::DrawIndexAuto(DrawAutoArgs args) {
 	if (args.instance_count == 0) {
-		args.instance_count = m_num_instances;
+		args.instance_count = NumInstances();
 	}
 	m_renderer.GetRenderExecutor().DrawAuto(m_submit_id, CurrentBuffer(), args);
 }

 void CommandProcessor::WaitFlipDone(uint32_t video_out_handle, uint32_t display_buffer_index) {
-	BufferFlush();
-
-	m_renderer.GetVideoOut().WaitFlipDone(static_cast<int>(video_out_handle),
-	                                      static_cast<int>(display_buffer_index));
+	// Stall only this command processor, not the host command lane or other queues.
+	if (!m_renderer.GetVideoOut().IsSafeForRendering(static_cast<int>(video_out_handle),
+	                                                 static_cast<int>(display_buffer_index))) {
+		SuspendPm4();
+	}
 }

 template <typename T>
@@ -1221,22 +1212,6 @@ void CommandProcessor::WriteAtEndOfPipe64(uint32_t cache_policy, uint32_t event_
 	                 interrupt_context_id);
 }

-void CommandProcessor::EmitGlobalBarrier() {
-	Common::LockGuard lock(m_renderer.GetMutex());
-
-	vk::MemoryBarrier2 barrier {};
-	barrier.srcStageMask  = vk::PipelineStageFlagBits2::eAllCommands;
-	barrier.srcAccessMask = vk::AccessFlagBits2::eMemoryWrite;
-	barrier.dstStageMask  = vk::PipelineStageFlagBits2::eAllCommands;
-	barrier.dstAccessMask = vk::AccessFlagBits2::eMemoryRead | vk::AccessFlagBits2::eMemoryWrite;
-
-	vk::DependencyInfo dependency {};
-	dependency.memoryBarrierCount = 1;
-	dependency.pMemoryBarriers    = &barrier;
-	GetScheduler().EndRendering();
-	CurrentBuffer().Handle().pipelineBarrier2(dependency);
-}
-
 void CommandProcessor::TriggerEopEventAtEndOfPipe(uint32_t interrupt_context_id) {
 	(void)CurrentBuffer().Handle();
 	GetScheduler().DeferPriorityOperation(
@@ -1256,21 +1231,17 @@ void CommandProcessor::TriggerEvent(uint32_t event_type, uint32_t event_index,
 	}

 	const auto valid_cache_event_index = event_index == 0x00000000 || event_index == 0x00000007;
+	// Partial flushes and cache writebacks only order GPU work, which the renderer's resource
+	// barriers already order.
 	switch (event_type) {
 		// CsPartialFlush, GsPartialFlush, PsPartialFlush.
 		case 0x00000007:
 		case 0x0000000f:
-		case 0x00000010: EmitGlobalBarrier(); break;
-		// CbDbDataWritebackInvalidate, CbDataWritebackInvalidate.
+		case 0x00000010: break;
+		// CbDbDataWritebackInvalidate, CbDataWritebackInvalidate, DbDataWritebackInvalidate,
+		// DbMetadataWritebackInvalidate, CbMetadataWritebackInvalidate.
 		case 0x00000016:
 		case 0x00000031:
-			if (!valid_cache_event_index) {
-				EXIT("unknown event type: 0x%08" PRIx32 ", 0x%08" PRIx32 "\n", event_type,
-				     event_index);
-			}
-			EmitGlobalBarrier();
-			break;
-		// DbDataWritebackInvalidate, DbMetadataWritebackInvalidate, CbMetadataWritebackInvalidate.
 		case 0x0000002a:
 		case 0x0000002c:
 		case 0x0000002e:
@@ -1278,7 +1249,6 @@ void CommandProcessor::TriggerEvent(uint32_t event_type, uint32_t event_index,
 				EXIT("unknown event type: 0x%08" PRIx32 ", 0x%08" PRIx32 "\n", event_type,
 				     event_index);
 			}
-			EmitGlobalBarrier();
 			break;
 		case 0x0000000d:
 		case 0x0000000e:
@@ -31,6 +31,8 @@ public:
 	[[nodiscard]] bool IsStopping();
 	void               SendCommand(Common::UniqueFunction<void>&& command);
 	void               SendCommandSync(Common::UniqueFunction<void>&& command);
+	// Retries blocked submissions now instead of after the 100 ms stall timeout.
+	void               Wake();

 	// Submitted command memory is borrowed and must remain valid until GPU execution completes.
 	void Submit(std::span<const uint32_t> draw_commands,
@@ -87,6 +89,7 @@ private:
 	bool                                           m_processing        = false;
 	bool                                           m_accepting         = true;
 	bool                                           m_stopping          = false;
+	bool                                           m_wake_pending      = false;
 	bool                                           m_shutdown_complete = false;
 	// Completion callbacks can outlive GuestGpu during renderer shutdown.
 	std::shared_ptr<std::binary_semaphore> m_suspend_point_ready =
@@ -5,6 +5,8 @@
 #include "common/common.h"
 #include "graphics/guest_gpu/gpu_defs.h"

+#include <array>
+
 namespace Libs::Graphics::HW {

 struct ColorBase {
@@ -574,6 +576,12 @@ struct ShaderRegisters {
 	uint32_t m_paScShaderControl = 0;

 	DepthShaderControl db_shader_control;
+	bool               alpha_to_mask_enable = false;
+
+	// DB_SHADER_CONTROL.ALPHA_TO_MASK_DISABLE overrides DB_ALPHA_TO_MASK per shader.
+	[[nodiscard]] bool AlphaToMask() const {
+		return alpha_to_mask_enable && !db_shader_control.alpha_to_mask_disable;
+	}

 	[[nodiscard]] uint32_t GetExportCount() const {
 		return 1u + ((m_spiVsOutConfig >> 1u) & 0x1Fu);
@@ -824,6 +832,11 @@ public:
 	void SetBlendColor(const BlendColor& color) { m_blend_color = color; }
 	[[nodiscard]] const ClipControl& GetClipControl() const { return m_clip_control; }
 	void SetClipControl(const ClipControl& control) { m_clip_control = control; }
+	// PA_CL_UCP_0_X..PA_CL_UCP_5_W as raw float bits, four per plane.
+	[[nodiscard]] const std::array<uint32_t, 24>& GetUserClipPlanes() const {
+		return m_user_clip_planes;
+	}
+	void SetUserClipPlane(uint32_t index, uint32_t value) { m_user_clip_planes[index] = value; }
 	[[nodiscard]] const RenderControl& GetRenderControl() const { return m_render_control; }
 	void SetRenderControl(const RenderControl& control) { m_render_control = control; }
 	[[nodiscard]] const DepthRenderOverride& GetDepthRenderOverride() const {
@@ -908,6 +921,7 @@ public:
 	void SetDepthShaderControl(const DepthShaderControl& value) {
 		m_sh_regs.db_shader_control = value;
 	}
+	void SetAlphaToMaskEnable(bool value) { m_sh_regs.alpha_to_mask_enable = value; }

 	void SetScShaderControl(uint32_t value) { m_sh_regs.m_paScShaderControl = value; }

@@ -924,6 +938,8 @@ private:
 	ColorControl    m_color_control;
 	ScanModeControl m_scan_mode_control;

+	std::array<uint32_t, 24> m_user_clip_planes {};
+
 	AaSampleControl m_aa_sample_control;
 	AaConfig        m_aa_config;

@@ -549,6 +549,8 @@ constexpr uint32_t VGT_SHADER_STAGES_EN                                = 0x2D5;
 constexpr uint32_t VGT_LS_HS_CONFIG                                    = 0x2D6;
 constexpr uint32_t VGT_TF_PARAM                                        = 0x2DB;
 constexpr uint32_t DB_ALPHA_TO_MASK                                    = 0x2DC;
+constexpr uint32_t DB_ALPHA_TO_MASK_ALPHA_TO_MASK_ENABLE_SHIFT         = 0;
+constexpr uint32_t DB_ALPHA_TO_MASK_ALPHA_TO_MASK_ENABLE_MASK          = 0x1;
 constexpr uint32_t PA_SU_POLY_OFFSET_DB_FMT_CNTL                       = 0x2DE;
 constexpr uint32_t PA_SU_POLY_OFFSET_DB_FMT_CNTL_NEG_NUM_DB_BITS_SHIFT = 0;
 constexpr uint32_t PA_SU_POLY_OFFSET_DB_FMT_CNTL_NEG_NUM_DB_BITS_MASK  = 0xFF;
@@ -8,6 +8,7 @@

 #include <map>
 #include <mutex>
+#include <span>
 #include <tuple>
 #include <vector>
 #include <vk_mem_alloc.h>
@@ -34,6 +35,7 @@ struct GraphicContext {
 	bool                               attachment_feedback_loop_enabled      = false;
 	bool                               provoking_vertex_last_enabled         = false;
 	bool                               maintenance8_enabled                  = false;
+	bool                               conditional_rendering_enabled         = false;
 	bool                               supports_block_texel_view              = false;
 	bool                                      mesh_shader_enabled                   = false;
 	vk::PhysicalDeviceMeshShaderPropertiesEXT mesh_shader_properties                = {};
@@ -41,6 +43,7 @@ struct GraphicContext {
 	uint32_t                           min_subgroup_size                     = 0;
 	uint32_t                           max_subgroup_size                     = 0;
 	uint32_t                           max_push_descriptors                  = 0;
+	vk::DescriptorSetLayout            image_heap_layout                     = nullptr;
 	vk::ShaderStageFlags               required_subgroup_size_stages         = {};
 	Common::Mutex                      queue_mutex;
 	uint32_t                           queue_family = static_cast<uint32_t>(-1);
@@ -99,6 +102,14 @@ struct GraphicContext {
 	[[nodiscard]] bool CanReportMemoryUsage() const noexcept { return memory_budget_ext_enabled; }
 	[[nodiscard]] uint64_t GetDeviceMemoryUsage() const;
 	[[nodiscard]] uint64_t GetTotalMemoryBudget() const;
+
+	struct HeapUsage {
+		uint64_t usage        = 0; // Whole process, as reported by VK_EXT_memory_budget.
+		uint64_t budget       = 0;
+		uint32_t allocations  = 0;
+		bool     device_local = false;
+	};
+	[[nodiscard]] uint32_t GetHeapUsage(std::span<HeapUsage, VK_MAX_MEMORY_HEAPS> heaps) const;
 	[[nodiscard]] bool     CreateImage(const vk::ImageCreateInfo& info, VulkanImage& image);
 	void                   DeleteImage(VulkanImage& image);

@@ -66,6 +66,18 @@ RegionManager* MemoryTracker::GetOrCreateRegion(uint64_t index) {
 	return ptr;
 }

+void MemoryTracker::QueueBdaSync(uint64_t vaddr, uint64_t size) {
+	const auto begin = Common::AlignDown(vaddr, TRACKER_PAGE_SIZE);
+	const auto end   = Common::AlignUp(vaddr + size, TRACKER_PAGE_SIZE);
+	std::lock_guard lock(m_bda_sync_mutex);
+	m_bda_sync.Add(begin, end - begin);
+}
+
+RangeSet MemoryTracker::TakeBdaSync() {
+	std::lock_guard lock(m_bda_sync_mutex);
+	return std::exchange(m_bda_sync, RangeSet {});
+}
+
 bool MemoryTracker::IsRegionCpuModified(uint64_t vaddr, uint64_t size) {
 	CheckNotInUploadCallback();
 	return Iterate<true>(vaddr, size, [](RegionManager* manager, uint64_t offset, uint64_t bytes) {
@@ -88,6 +100,7 @@ void MemoryTracker::MarkRegionAsCpuModified(uint64_t vaddr, uint64_t size) {
 		std::scoped_lock lock(manager->lock);
 		manager->ChangeState<DirtySource::Cpu, true>(manager->GetCpuAddr() + offset, bytes);
 	});
+	QueueBdaSync(vaddr, size);
 }

 void MemoryTracker::MarkRegionAsGpuModified(uint64_t vaddr, uint64_t size) {
@@ -128,6 +141,8 @@ void MemoryTracker::UntrackMemory(uint64_t vaddr, uint64_t size) {
 	Iterate<false>(vaddr, size, [](RegionManager* manager, uint64_t offset, uint64_t bytes) {
 		manager->ChangeState<DirtySource::Cpu, true>(manager->GetCpuAddr() + offset, bytes);
 	});
+	locks.clear();
+	QueueBdaSync(vaddr, size);
 }

 } // namespace Libs::Graphics
@@ -29,6 +29,10 @@ public:
 	void               MarkRegionAsGpuModified(uint64_t vaddr, uint64_t size);
 	void               UnmarkRegionAsGpuModified(uint64_t vaddr, uint64_t size);
 	void               UntrackMemory(uint64_t vaddr, uint64_t size);
+	// Queue after changing ownership; a concurrent drain must not consume an event
+	// before the dirty pages it describes become visible.
+	void                   QueueBdaSync(uint64_t vaddr, uint64_t size);
+	[[nodiscard]] RangeSet TakeBdaSync();
 	// Removes protection from a range and flushes GPU-owned data when required.
 	template <typename Flush>
 	void InvalidateRegion(uint64_t vaddr, uint64_t size, Flush&& on_flush) noexcept {
@@ -51,6 +55,7 @@ public:
 				on_flush();
 			}
 		});
+		QueueBdaSync(vaddr, size);
 	}
 #if KYTY_BUILD == KYTY_BUILD_DEBUG
 	void ValidateGpuDirtyPages(const RangeSet& dirty, uint64_t vaddr, uint64_t size,
@@ -149,6 +154,8 @@ private:
 	std::unique_ptr<std::atomic<RegionManager*>[]> m_regions;
 	std::vector<std::unique_ptr<RegionManager>>    m_region_storage;
 	std::mutex                                     m_region_mutex;
+	std::mutex                                     m_bda_sync_mutex;
+	RangeSet                                       m_bda_sync;
 	PageManager&                                   m_page_manager;
 };

@@ -2,6 +2,7 @@

 #include "common/alignment.h"
 #include "common/assert.h"
+#include "common/debugCounters.h"
 #include "common/logging/log.h"
 #include "common/profiler.h"
 #include "graphics/guest_gpu/graphicsRun.h"
@@ -126,6 +127,7 @@ void BufferCache::ChangeRegister(BufferId id) {
 		}
 		WriteDataBuffer(m_bda_pagetable_buffer, table_offset,
 		                addresses.data(), addresses.size() * sizeof(vk::DeviceAddress));
+		m_memory_tracker.QueueBdaSync(buffer.CpuAddress(), buffer.Size());
 	} else {
 		const auto found = m_buffers.find(buffer.CpuAddress());
 		EXIT_IF(found == m_buffers.end() || found->second != id);
@@ -172,10 +174,12 @@ bool BufferCache::DownloadBufferMemory(Buffer& buffer, uint64_t vaddr, uint64_t
 			    total_size += Common::AlignUp(end - start, 64);
 		    });
 		    m_gpu_modified_ranges.Subtract(address, bytes);
+		    ForgetColorMetadataChecks(address, bytes);
 	    });
 	if (copies.empty()) {
 		return false;
 	}
+	Common::DebugCounters::Add(Common::DebugCounters::Counter::BufferDownloadBytes, total_size);

 	auto [mapped, offset] = m_download_buffer.Map(total_size, 64);
 	std::unique_ptr<Buffer> temporary;
@@ -253,6 +257,10 @@ BufferCache::BufferCache(GraphicContext& graphics, CommandScheduler& scheduler,
 	m_gds_buffer.Flush(0, m_gds_buffer.Size());
 	SetVulkanObjectNameF(m_graphics.device, m_bda_pagetable_buffer.Handle(),
 	                     "BDA Page Table Buffer");
+	m_scheduler.BeginCommand();
+	for (auto* buffer: {&m_bda_pagetable_buffer, m_fault_manager.GetFaultBuffer()}) {
+		buffer->Fill(0, buffer->Size(), 0);
+	}
 	const auto null_id =
 	    m_slot_buffers.insert(m_graphics, m_scheduler, MemoryUsage::DeviceLocal, 0, AllFlags, 16);
 	EXIT_IF(null_id != NULL_BUFFER_ID);
@@ -466,6 +474,7 @@ vk::Buffer BufferCache::UploadCopies(Buffer& buffer, std::span<vk::BufferCopy> c
 	if (copies.empty()) {
 		return nullptr;
 	}
+	Common::DebugCounters::Add(Common::DebugCounters::Counter::BufferUploadBytes, total_size);

 	const auto source = UploadMemory(m_graphics, m_scheduler, m_staging_buffer, total_size, 4,
 	    vk::BufferUsageFlagBits::eTransferSrc, [&](uint8_t* mapped, uint64_t base_offset) {
@@ -511,6 +520,7 @@ std::pair<Buffer*, uint64_t> BufferCache::ObtainBuffer(uint64_t vaddr, uint64_t
 	(void)SynchronizeBuffer(buffer, vaddr, size, is_written, is_texel_buffer);
 	if (is_written) {
 		m_gpu_modified_ranges.Add(vaddr, size);
+		ForgetColorMetadataChecks(vaddr, size);
 	}
 	return {&buffer, buffer.Offset(vaddr)};
 }
@@ -615,6 +625,29 @@ bool BufferCache::IsRegionRegistered(uint64_t vaddr, uint64_t size) {
 	return address + m_slot_buffers[id].Size() > vaddr;
 }

+void BufferCache::MarkColorMetadataChecked(uint64_t vaddr, uint64_t size, uint32_t keys) {
+	for (uint32_t slot = 0; slot < m_checked_color_metadata.size(); slot++) {
+		if ((keys & (1u << slot)) != 0) {
+			m_checked_color_metadata[slot].Add(vaddr, size);
+		}
+	}
+}
+
+bool BufferCache::IsColorMetadataChecked(uint64_t vaddr, uint64_t size, uint32_t keys) const {
+	for (uint32_t slot = 0; slot < m_checked_color_metadata.size(); slot++) {
+		if ((keys & (1u << slot)) != 0 && !m_checked_color_metadata[slot].Contains(vaddr, size)) {
+			return false;
+		}
+	}
+	return true;
+}
+
+void BufferCache::ForgetColorMetadataChecks(uint64_t vaddr, uint64_t size) {
+	for (auto& checked: m_checked_color_metadata) {
+		checked.Subtract(vaddr, size);
+	}
+}
+
 bool BufferCache::IsRegionGpuModified(uint64_t vaddr, uint64_t size) {
 	return m_memory_tracker.IsRegionGpuModified(vaddr, size);
 }
@@ -657,6 +690,7 @@ void BufferCache::RunGarbageCollector() {
 		} else {
 			m_memory_tracker.UntrackMemory(buffer.CpuAddress(), buffer.Size());
 			DeleteBuffer(id);
+			Common::DebugCounters::Add(Common::DebugCounters::Counter::BufferEvictions);
 		}
 		return ++retire_count == limit;
 	});
@@ -685,6 +719,16 @@ void BufferCache::ProcessFaultBuffer() {
 	m_fault_manager.ProcessFaultBuffer();
 }

+void BufferCache::SynchronizeBda(const RangeSet& mapped) {
+	// Take before walking: writes queued during synchronization remain pending.
+	const auto pending = m_memory_tracker.TakeBdaSync();
+	pending.ForEach([&](uint64_t begin, uint64_t end) {
+		mapped.ForEachInRange(begin, end - begin, [this](uint64_t start, uint64_t finish) {
+			SynchronizeBuffersInRange(start, finish - start);
+		});
+	});
+}
+
 void BufferCache::SynchronizeBuffersInRange(uint64_t vaddr, uint64_t size) {
 	const auto end = vaddr + size;
 	auto       it  = m_buffers.upper_bound(vaddr);
@@ -11,6 +11,7 @@
 #include "graphics/host_gpu/renderer/cache/multiLevelPageTable.h"
 #include "graphics/host_gpu/renderer/cache/streamBuffer.h"

+#include <array>
 #include <map>
 #include <span>
 #include <utility>
@@ -79,8 +80,15 @@ public:
 	[[nodiscard]] bool HasGpuDirtyBytes(uint64_t vaddr, uint64_t size);
 	[[nodiscard]] bool IsRegionCpuModified(uint64_t vaddr, uint64_t size);
 	[[nodiscard]] bool IsRegionGpuModified(uint64_t vaddr, uint64_t size);
+	// Color metadata the GPU found free of the clear keys in `keys` (bits of code >> 5) stays so
+	// until a GPU write or download.
+	void               MarkColorMetadataChecked(uint64_t vaddr, uint64_t size, uint32_t keys);
+	[[nodiscard]] bool IsColorMetadataChecked(uint64_t vaddr, uint64_t size, uint32_t keys) const;
 	void               ProcessFaultBuffer();
-	void               SynchronizeBuffersInRange(uint64_t vaddr, uint64_t size);
+	void               SynchronizeBda(const RangeSet& mapped);
+	void               QueueBdaSync(uint64_t vaddr, uint64_t size) {
+		m_memory_tracker.QueueBdaSync(vaddr, size);
+	}
 	void               RunGarbageCollector();

 private:
@@ -114,12 +122,14 @@ private:
 	void DeleteBuffer(BufferId id);
 	[[nodiscard]] bool SynchronizeBuffer(Buffer& buffer, uint64_t vaddr, uint64_t size,
 	                                     bool is_written, bool is_texel_buffer);
+	void SynchronizeBuffersInRange(uint64_t vaddr, uint64_t size);
 	[[nodiscard]] vk::Buffer UploadCopies(Buffer& buffer, std::span<vk::BufferCopy> copies,
 	                                      uint64_t total_size);
 	[[nodiscard]] bool SynchronizeBufferFromImage(Buffer& buffer, uint64_t vaddr, uint64_t size);
 	// Synchronous downloads publish before returning; asynchronous callers wait before reuse.
 	template <bool async>
 	[[nodiscard]] bool DownloadBufferMemory(Buffer& buffer, uint64_t vaddr, uint64_t size);
+	void               ForgetColorMetadataChecks(uint64_t vaddr, uint64_t size);

 	GraphicContext&                                   m_graphics;
 	CommandScheduler&                                 m_scheduler;
@@ -131,6 +141,8 @@ private:
 	BufferMap                                         m_buffers;
 	PageTable                                         m_page_table;
 	RangeSet                                          m_gpu_modified_ranges;
+	// One per TileManager::ColorClearSlots key slot.
+	std::array<RangeSet, 8>                           m_checked_color_metadata;
 	MemoryTracker                                     m_memory_tracker;
 	StreamBuffer                                      m_staging_buffer;
 	StreamBuffer                                      m_stream_buffer;
@@ -67,11 +67,9 @@ Buffer::Buffer(GraphicContext& graphics, CommandScheduler& scheduler, MemoryUsag
 	buffer_info.usage       = flags;

 	const bool with_bda = bool(flags & vk::BufferUsageFlagBits::eShaderDeviceAddress);
-	const VmaAllocationCreateFlags bda_flag =
-	    with_bda ? VMA_ALLOCATION_CREATE_DEDICATED_MEMORY_BIT : 0;
 	VmaAllocationCreateInfo allocation_info {};
 	allocation_info.flags =
-	    VMA_ALLOCATION_CREATE_WITHIN_BUDGET_BIT | bda_flag | AllocationFlags(usage);
+	    VMA_ALLOCATION_CREATE_WITHIN_BUDGET_BIT | AllocationFlags(usage);
 	allocation_info.usage = AllocationUsage(usage);
 	allocation_info.preferredFlags = usage == MemoryUsage::DeviceLocal
 	                                     ? VkMemoryPropertyFlags {}
@@ -2,6 +2,7 @@

 #include "common/alignment.h"
 #include "common/assert.h"
+#include "common/debugCounters.h"
 #include "common/emulatorConfig.h"
 #include "common/logging/log.h"
 #include "common/profiler.h"
@@ -33,6 +34,8 @@ namespace Libs::Graphics {
 namespace {

 constexpr uint64_t NumFramesBeforeRemoval = 32;
+// Largest GPU-written tiled image (in texels) that is written back to guest memory on submit.
+constexpr uint64_t SmallTiledReadbackTexels = 64 * 64;

 [[nodiscard]] bool DecodeColorClear(const TextureCache::ImageDesc& desc, uint8_t code,
                                  vk::ClearColorValue& clear) {
@@ -121,7 +124,8 @@ TextureCache::TextureCache(GraphicContext& graphics, CommandScheduler& scheduler
      m_blit_helper(graphics, scheduler),
      m_tiler(graphics, scheduler, buffer_cache.GetUtilityBuffer(MemoryUsage::Stream)),
      m_buffer_cache(buffer_cache),
-      m_readback_linear_images(Config::ReadbackLinearImagesEnabled()) {
+      m_readback_linear_images(Config::ReadbackLinearImagesEnabled()),
+      m_readback_tiled_images(Config::ReadbackTiledImagesEnabled()) {
 	if (m_graphics.CanReportMemoryUsage()) {
 		constexpr int64_t GiB = 1024ll * 1024 * 1024;
 		const auto        budget =
@@ -225,6 +229,7 @@ void TextureCache::UnregisterImage(ImageId id) {
 	if (!image.registered) {
 		return;
 	}
+	image.HeapChanged();
 	UntrackImage(id);
 	ImagePageTable::PageRange pages {};
 	if (!ImagePageTable::TryGetPageRange(image.info.data.address, image.info.data.size, pages)) {
@@ -262,14 +267,8 @@ void TextureCache::DeleteImage(ImageId id) {
 		EXIT("TextureCache: deleting a GPU-modified image without resolving its contents\n");
 	}
 	m_download_images.erase(id);
-	if (image->info.HasMetadata()) {
-		const auto metadata = m_surface_metas.find(image->info.metadata.range.address);
-		if (metadata != m_surface_metas.end() &&
-		    image->info.metadata.kind == ImageMetadataKind::Htile &&
-		    metadata->second.type == MetaDataInfo::Type::HTile) {
-			// A later binding may have reused this address for another metadata type.
-			m_surface_metas.erase(metadata);
-		}
+	if (image->info.metadata.kind == ImageMetadataKind::Htile) {
+		m_surface_metas.erase(image->info.metadata.range.address);
 	}
 	UnregisterImage(id);
 	if (m_scheduler.Active()) {
@@ -380,6 +379,7 @@ void TextureCache::UntrackImageHead(ImageId id) {
 	const auto address = Common::AlignDown(begin + TRACKER_PAGE_SIZE, TRACKER_PAGE_SIZE);
 	const auto size    = address - begin;
 	image.track_addr   = address;
+	image.HeapChanged();
 	if (image.track_addr == image.track_addr_end) {
 		MarkAsMaybeDirty(id, image);
 	}
@@ -397,6 +397,7 @@ void TextureCache::UntrackImageTail(ImageId id) {
 	const auto address   = Common::AlignDown(end, TRACKER_PAGE_SIZE);
 	const auto size      = end - address;
 	image.track_addr_end = address;
+	image.HeapChanged();
 	if (image.track_addr == image.track_addr_end) {
 		MarkAsMaybeDirty(id, image);
 	}
@@ -406,7 +407,14 @@ void TextureCache::UntrackImageTail(ImageId id) {
 }

 void TextureCache::TrackImageDownload(ImageId id, Image& image) {
-	if (m_readback_linear_images && !image.info.IsTiled()) {
+	// Small tiled targets are written back too: games read GPU results such as reductions and
+	// simulation state from them on the CPU (PPSA04263 keeps the origin of its ocean wave grid
+	// in a 1x1 R32G32 target, and its water physics read it to sample the wave heights).
+	const auto& info   = image.info;
+	const auto  texels = static_cast<uint64_t>(info.extent.width) * info.extent.height *
+	                     info.extent.depth * info.resources.layers;
+	if (info.IsTiled() ? m_readback_tiled_images && texels <= SmallTiledReadbackTexels
+	                   : m_readback_linear_images) {
 		if (!image.IsGpuModified()) {
 			EXIT("TextureCache: cannot enroll a non-GPU-owned image for download\n");
 		}
@@ -790,7 +798,8 @@ TextureCache::OverlapResult TextureCache::ResolveOverlap(const ImageInfo& reques
 	auto&      cached       = *owner;
 	const auto current_tick = m_scheduler.CurrentTick();
 	const bool safe_to_delete =
-	    current_tick - std::min(current_tick, cached.tick_accessed_last) > NumFramesBeforeRemoval;
+	    current_tick - std::min(current_tick, cached.tick_accessed_last) > NumFramesBeforeRemoval &&
+	    cached.heap_refs.load(std::memory_order_relaxed) == 0u;

 	const uint32_t requested_block = requested.bytes_per_block * requested.samples;
 	const uint32_t cached_block    = cached.info.bytes_per_block * cached.info.samples;
@@ -954,6 +963,7 @@ void TextureCache::UploadImage(Image& image, Buffer& source, uint64_t source_off
 		const auto linear = m_tiler.Detile(source.Handle(), source_offset, info);
 		const auto copies = info.BufferCopies(linear.offset);
 		destination.Upload(copies, linear.buffer, linear.offset, linear.size);
+		Common::DebugCounters::Add(Common::DebugCounters::Counter::TextureUploadBytes, linear.size);
 		return;
 	}

@@ -999,6 +1009,7 @@ void TextureCache::UploadImage(Image& image, Buffer& source, uint64_t source_off
 		copies[0].imageExtent       = info.extent;
 	}
 	destination.Upload(copies, linear.buffer, linear.offset, linear.size);
+	Common::DebugCounters::Add(Common::DebugCounters::Counter::TextureUploadBytes, linear.size);
 }

 void TextureCache::InitializeImage(ImageId id) {
@@ -1039,7 +1050,7 @@ void TextureCache::MaterializeColorClear(ImageId id, const ImageDesc& desc,
 		std::scoped_lock lock {m_lock};
 		auto& image         = m_slot_images[id];
 		image.info.metadata = desc.info.metadata;
-		// Native color metadata must not retain a reused HTile/CMask/FMask clear flag.
+		// Native color metadata must not retain a reused HTile clear flag.
 		m_surface_metas.erase(range.address);
 		if (range.size == 0 || desc.info.resources.levels != 1) {
 			return;
@@ -1060,12 +1071,78 @@ void TextureCache::MaterializeColorClear(ImageId id, const ImageDesc& desc,
 	if (first >= layers || count > layers - first) {
 		EXIT("TextureCache: color view exceeds its native metadata slices\n");
 	}
-	// Finish native metadata writes before reading backing bytes. This can submit the scheduler,
-	// so discovery runs before final draw uploads and never holds the texture lock across it.
-	if (m_buffer_cache.IsRegionGpuModified(range.address, range.size)) {
-		m_buffer_cache.ReadMemory(range.address, range.size, false);
+	const auto       slice_size = range.size / layers;
+	const GuestRange viewed {range.address + slice_size * first, slice_size * count};
+	if (m_buffer_cache.IsRegionGpuModified(viewed.address, viewed.size)) {
+		std::array<vk::ClearAttachment, TileManager::ColorClearSlots> clears {};
+		uint32_t keys = 0;
+		for (uint32_t key = 0; key < clears.size(); key++) {
+			clears[key].aspectMask = vk::ImageAspectFlagBits::eColor;
+			if (DecodeColorClear(desc, static_cast<uint8_t>(key << 5u),
+			                     clears[key].clearValue.color)) {
+				keys |= 1u << key;
+			}
+		}
+		// No selected key to clear, or none left since the last check.
+		if (m_buffer_cache.IsColorMetadataChecked(viewed.address, viewed.size, keys)) {
+			return;
+		}
+		{
+			std::scoped_lock lock {m_lock};
+			auto&            image = m_slot_images[id];
+			// Compressed images cannot upload, so they must already hold native contents.
+			const bool native_current = (image.usage.render_target || image.IsGpuModified()) &&
+			                            !image.IsBufferModified() && !image.IsCpuDirty();
+			// 2D-array attachment views select 2D layers or, of a volume, depth slices.
+			if (m_graphics.conditional_rendering_enabled &&
+			    image.backing.image_type != vk::ImageType::e1D &&
+			    count <= m_graphics.GetPhysicalDeviceProperties().limits.maxFramebufferLayers &&
+			    (image.backing.usage & vk::ImageUsageFlagBits::eColorAttachment) &&
+			    (m_graphics.GetFormatProperties(view.format).optimalTilingFeatures &
+			     vk::FormatFeatureFlagBits::eColorAttachment) &&
+			    (desc.info.metadata.compression == VideoOutCompression::Uncompressed ||
+			     native_current)) {
+				// The GPU selects each slice's clear, so the CPU never waits for the metadata.
+				Common::DebugCounters::Add(Common::DebugCounters::Counter::ColorMetadataChecks);
+				InitializeImage(id);
+				const bool expand = desc.type != BindingType::VideoOut;
+				const auto [buffer, offset] =
+				    m_buffer_cache.ObtainBuffer(viewed.address, viewed.size, expand);
+				const auto predicates = m_tiler.CheckColorClear(
+				    {buffer->Handle(), offset, viewed.size}, count, keys, expand);
+				if (expand) {
+					m_buffer_cache.MarkColorMetadataChecked(viewed.address, viewed.size, keys);
+				}
+				const auto command = m_scheduler.Current().Handle();
+				const auto area    = BeginColorRendering(image, view.format, view.base_level,
+				                                         image_first, count, nullptr, command);
+				vk::ClearRect rect {.rect = area, .layerCount = 1};
+				for (uint32_t slice = 0; slice < count; slice++) {
+					rect.baseArrayLayer = slice;
+					for (uint32_t key = 0; key < clears.size(); key++) {
+						if ((keys & (1u << key)) != 0) {
+							const vk::ConditionalRenderingBeginInfoEXT predicate {
+							    .buffer = predicates,
+							    .offset = (slice * clears.size() + key) * sizeof(uint32_t)};
+							command.beginConditionalRenderingEXT(&predicate);
+							command.clearAttachments(1, &clears[key], 1, &rect);
+							command.endConditionalRenderingEXT();
+						}
+					}
+				}
+				command.endRendering();
+				CommitGpuWrite(image);
+				return;
+			}
+		}
+		// Finish native metadata writes before reading backing bytes. This can submit the
+		// scheduler, so discovery runs before final draw uploads and never holds the texture
+		// lock across it.
+		Common::DebugCounters::ScopedTimer timer(
+		    Common::DebugCounters::Counter::ColorMetadataReadbacks,
+		    Common::DebugCounters::Counter::ColorMetadataReadbackNs);
+		m_buffer_cache.ReadMemory(viewed.address, viewed.size, false);
 	}
-	const auto slice_size = range.size / layers;
 	for (uint32_t slice = 0; slice < count; slice++) {
 		const auto address = range.address + slice_size * (first + slice);
 		uint8_t code = 0;
@@ -1417,11 +1494,8 @@ vk::ImageView TextureCache::FindDepthTarget(ImageId id, const ImageDesc& desc) {
 	image.info.metadata = desc.info.metadata;
 	if (desc.info.HasMetadata()) {
 		const auto [metadata, inserted] = m_surface_metas.try_emplace(desc.info.metadata.range.address);
-		if (inserted) {
-			metadata->second.type = MetaDataInfo::Type::HTile;
-			if (image.info.htile_initial_clear) {
-				metadata->second.clear_mask.set();
-			}
+		if (inserted && image.info.htile_initial_clear) {
+			metadata->second.clear_mask.set();
 		}
 	}
 	RefreshImage(id);
@@ -1570,29 +1644,8 @@ void TextureCache::ClearImage(CommandBuffer& command, ImageId id, vk::Format for
 	if (format != image.backing.format || (image.info.IsVolume() && !full_image)) {
 		EXIT_NOT_IMPLEMENTED(range.aspectMask != vk::ImageAspectFlagBits::eColor ||
 		                     range.levelCount != 1);
-		ImageViewInfo view {};
-		view.format = format;
-		view.type   = range.layerCount == 1 ? vk::ImageViewType::e2D : vk::ImageViewType::e2DArray;
-		view.base_level  = range.baseMipLevel;
-		view.base_layer  = range.baseArrayLayer;
-		view.layer_count = range.layerCount;
-		view.usage       = vk::ImageUsageFlagBits::eColorAttachment;
-		image.Transit(vk::ImageLayout::eColorAttachmentOptimal,
-		              vk::AccessFlagBits2::eColorAttachmentWrite, {}, command.Handle());
-		vk::RenderingAttachmentInfo attachment {};
-		attachment.imageView   = image.FindView(view);
-		attachment.imageLayout = vk::ImageLayout::eColorAttachmentOptimal;
-		attachment.loadOp      = vk::AttachmentLoadOp::eClear;
-		attachment.storeOp     = vk::AttachmentStoreOp::eStore;
-		attachment.clearValue  = clear;
-		vk::RenderingInfo rendering {};
-		rendering.renderArea.extent = {
-		    std::max(image.info.extent.width >> range.baseMipLevel, 1u),
-		    std::max(image.info.extent.height >> range.baseMipLevel, 1u)};
-		rendering.layerCount           = range.layerCount;
-		rendering.colorAttachmentCount = 1;
-		rendering.pColorAttachments    = &attachment;
-		command.Handle().beginRendering(&rendering);
+		BeginColorRendering(image, format, range.baseMipLevel, range.baseArrayLayer,
+		                    range.layerCount, &clear, command.Handle());
 		command.Handle().endRendering();
 		CommitGpuWrite(image);
 		return;
@@ -1615,6 +1668,40 @@ void TextureCache::ClearImage(CommandBuffer& command, ImageId id, vk::Format for
 	CommitGpuWrite(image);
 }

+vk::Rect2D TextureCache::BeginColorRendering(Image& image, vk::Format format, uint32_t level,
+                                             uint32_t layer, uint32_t layers,
+                                             const vk::ClearValue* clear,
+                                             vk::CommandBuffer     command) {
+	ImageViewInfo view {};
+	view.format      = format;
+	view.type        = layers == 1 ? vk::ImageViewType::e2D : vk::ImageViewType::e2DArray;
+	view.base_level  = level;
+	view.base_layer  = layer;
+	view.layer_count = layers;
+	view.usage       = vk::ImageUsageFlagBits::eColorAttachment;
+	image.Transit(vk::ImageLayout::eColorAttachmentOptimal,
+	              vk::AccessFlagBits2::eColorAttachmentRead |
+	                  vk::AccessFlagBits2::eColorAttachmentWrite,
+	              {}, command);
+	vk::RenderingAttachmentInfo attachment {};
+	attachment.imageView   = image.FindView(view);
+	attachment.imageLayout = vk::ImageLayout::eColorAttachmentOptimal;
+	attachment.loadOp      = vk::AttachmentLoadOp::eLoad;
+	attachment.storeOp     = vk::AttachmentStoreOp::eStore;
+	if (clear != nullptr) {
+		attachment.loadOp     = vk::AttachmentLoadOp::eClear;
+		attachment.clearValue = *clear;
+	}
+	vk::RenderingInfo rendering {};
+	rendering.renderArea.extent    = {std::max(image.info.extent.width >> level, 1u),
+	                                  std::max(image.info.extent.height >> level, 1u)};
+	rendering.layerCount           = layers;
+	rendering.colorAttachmentCount = 1;
+	rendering.pColorAttachments    = &attachment;
+	command.beginRendering(&rendering);
+	return rendering.renderArea;
+}
+
 void TextureCache::InvalidateMemory(uint64_t address, uint64_t size) {
 	if (!GuestRange {address, size}.Valid()) {
 		EXIT("TextureCache: invalid memory-invalidation range\n");
@@ -1748,6 +1835,7 @@ bool BufferCache::SynchronizeBufferFromImage(Buffer& buffer, uint64_t vaddr, uin
 	     copy_size != image.info.data.size)) {
 		return false;
 	}
+	ForgetColorMetadataChecks(image.info.data.address, copy_size);
 	m_texture_cache.DownloadImage(image, buffer, buf_offset, copy_size, levels);
 	return true;
 }
@@ -1774,6 +1862,7 @@ bool TextureCache::DownloadImageMemory(ImageId id) {
 	download.Flush(offset, range.size);

 	DownloadImage(image, download, offset, range.size, image.info.resources.levels);
+	Common::DebugCounters::Add(Common::DebugCounters::Counter::TextureDownloadBytes, range.size);
 	vk::BufferMemoryBarrier barrier {};
 	barrier.srcAccessMask = vk::AccessFlagBits::eMemoryWrite | vk::AccessFlagBits::eTransferWrite |
 	                        vk::AccessFlagBits::eShaderWrite;
@@ -1930,6 +2019,11 @@ void TextureCache::RunGarbageCollector() {
 			if (owner == nullptr || !owner->registered || owner->depth_id) {
 				continue;
 			}
+			if (owner->heap_refs.load(std::memory_order_relaxed) != 0u) {
+				// Resident heaps sample it without touching it.
+				TouchImage(*owner);
+				continue;
+			}
 			if (owner->IsGpuModified()) {
 				const bool safe = owner->SafeToDownload();
 				if (safe && owner->info.IsTiled()) {
@@ -1943,6 +2037,7 @@ void TextureCache::RunGarbageCollector() {
 				}
 			}
 			FreeImage(id);
+			Common::DebugCounters::Add(Common::DebugCounters::Counter::TextureEvictions);
 			if (m_total_used_memory < m_critical_gc_memory && aggressive) {
 				deletions >>= 2;
 				aggressive = false;
@@ -1960,7 +2055,7 @@ void TextureCache::RunGarbageCollector() {
 }

 void TextureCache::ProcessDownloadImages() {
-	if (!m_readback_linear_images) {
+	if (!m_readback_linear_images && !m_readback_tiled_images) {
 		return;
 	}
 	std::scoped_lock lock {m_lock};
@@ -81,9 +81,6 @@ public:

 private:
 	struct MetaDataInfo {
-		enum class Type : uint8_t { CMask, FMask, HTile };
-
-		Type type;
 		// PS5 DB_DEPTH_VIEW encodes 13-bit base and last array-slice indices.
 		std::bitset<8192> clear_mask;
 	};
@@ -153,6 +150,12 @@ private:
 	// Caller holds m_lock. Volume layer ranges select depth slices.
 	void ClearImage(CommandBuffer& command, ImageId id, vk::Format format,
 	                const vk::ImageSubresourceRange& range, const vk::ClearValue& clear);
+	// Caller holds m_lock and ends rendering. Begins rendering `layers` layers of one color mip
+	// through a `format` view, clearing them or, when `clear` is null, loading them; returns the
+	// render area.
+	vk::Rect2D BeginColorRendering(Image& image, vk::Format format, uint32_t level,
+	                               uint32_t layer, uint32_t layers, const vk::ClearValue* clear,
+	                               vk::CommandBuffer command);
 	void PrepareImageCopy(Image& image);
 	void RefreshCopySource(ImageId id);
 	[[nodiscard]] bool CopyD16(Image& destination, Image& source);
@@ -183,10 +186,12 @@ private:
 	uint64_t         m_gc_tick                = 0;
 	mutable uint32_t m_image_query_epoch      = 0;
 	bool             m_readback_linear_images = false;
+	bool             m_readback_tiled_images  = false;

 	friend struct TextureCacheTestAccess;
 	friend class BufferCache;
 	friend class RenderExecutor;
+	friend class ImageHeap;
 	friend bool BcAlias::Matches(const ImageInfo&, const ImageInfo&);
 	friend bool BcAlias::Publish(TextureCache&, Image&);
 	friend void BcAlias::Resolve(TextureCache&, ImageId, ImageId);
@@ -72,7 +72,16 @@ void RenderExecutor::ResolveRenderColorTarget(CommandBuffer& buffer, RenderColor
 		return;
 	}
 	const auto samples = render_sample_count(rt.attrib.num_fragments);
-	if (samples == 0 || rt.attrib.num_samples != rt.attrib.num_fragments) {
+	// TODO(PPSA24156): Replace this 2x MSAA approximation with proper 4-coverage/2-fragment EQAA.
+	const bool approximate_eqaa = rt.attrib.num_samples == 2 && rt.attrib.num_fragments == 1;
+	if (approximate_eqaa) {
+		static std::atomic_bool logged = false;
+		if (!logged.exchange(true, std::memory_order_relaxed)) {
+			Log::WriteToConsoleAndLog(
+			    "Warning: approximating 4-coverage/2-fragment EQAA with 2x MSAA.\n");
+		}
+	}
+	if (samples == 0 || (rt.attrib.num_samples != rt.attrib.num_fragments && !approximate_eqaa)) {
 		EXIT("unsupported render-target sample configuration: samples=%u fragments=%u\n",
 		     rt.attrib.num_samples, rt.attrib.num_fragments);
 	}
@@ -220,8 +229,7 @@ void RenderExecutor::ResolveRenderColorTarget(CommandBuffer& buffer, RenderColor
 	}
 	const bool has_dcc        = rt.info.dcc_compression_enable && rt.dcc_addr.addr != 0;
 	const bool has_cmask      = !rt.info.dcc_compression_enable && rt.info.cmask_fast_clear_enable &&
-	                            rt.cmask.addr != 0 && samples == 1 &&
-	                            !rt.info.fmask_compression_enable &&
+	                            rt.cmask.addr != 0 &&
 	                            !rt.attrib3.write_vrs_rate_hint_to_cmask;
 	if (has_dcc || has_cmask) {
 		TileSizeAlign metadata_size {};
@@ -233,7 +241,7 @@ void RenderExecutor::ResolveRenderColorTarget(CommandBuffer& buffer, RenderColor
 		} else {
 			(void)TileGetCmaskSize(width, height, layers, levels, metadata_size);
 		}
-		// DCC owns the clear when both planes are enabled; single-sample CMASK stays expanded.
+		// DCC owns the clear when both planes are enabled.
 		desc.info.metadata.kind = has_dcc ? ImageMetadataKind::Dcc : ImageMetadataKind::Cmask;
 		desc.info.metadata.range = {has_dcc ? rt.dcc_addr.addr : rt.cmask.addr, metadata_size.size};
 		desc.info.metadata.clear_word =
@@ -1,6 +1,7 @@
 #include "graphics/host_gpu/renderer/commandScheduler.h"

 #include "common/assert.h"
+#include "common/debugCounters.h"
 #include "common/logging/log.h"
 #include "graphics/host_gpu/graphicContext.h"

@@ -326,7 +327,7 @@ void CommandScheduler::CheckActive() const {
 }

 CommandBuffer& CommandScheduler::Current() {
-	CheckActive();
+	EXIT_IF(m_command.IsInvalid());
 	return m_command;
 }

@@ -373,6 +374,7 @@ uint64_t CommandScheduler::Submit(SubmitInfo submit) {

 		result = graphics.queue.submit(1, &submit_info, nullptr);
 	}
+	Common::DebugCounters::Add(Common::DebugCounters::Counter::QueueSubmits);

 	if (result != vk::Result::eSuccess) {
 		ReportVulkanFatal("vkQueueSubmit", result, tick, m_command.m_debug_op,
@@ -104,8 +104,14 @@ void CommandBuffer::BeginRendering(const RenderState& state) const {
 	rendering.pDepthAttachment     = depth_stencil.has_depth ? &depth : nullptr;
 	rendering.pStencilAttachment   = depth_stencil.has_stencil ? &stencil : nullptr;
 	Handle().beginRendering(rendering);
+	// The open pass holds the cleared contents, so a later clear begins a new pass.
 	m_render_state = state;
-	m_rendering    = true;
+	for (auto& color: m_render_state.color_attachments) {
+		color.is_clear = false;
+	}
+	m_render_state.depth_stencil_attachment.depth_clear   = false;
+	m_render_state.depth_stencil_attachment.stencil_clear = false;
+	m_rendering                                           = true;
 }

 void CommandBuffer::EndRendering() const {
@@ -294,16 +294,6 @@ static void RtCheck(const HW::RenderTarget& rt) {
 		EXIT_NOT_IMPLEMENTED(rt.dcc.overwrite_combiner_disable != false);
 		// EXIT_NOT_IMPLEMENTED(rt.dcc.data_write_on_dcc_clear_to_reg != false);
 		EXIT_NOT_IMPLEMENTED(rt.dcc.dcc_clear_key_enable != false);
-		if (rt.cmask.addr != 0x0000000000000000 || rt.fmask.addr != 0x0000000000000000 ||
-		    rt.dcc_addr.addr != 0x0000000000000000) {
-			static bool logged = false;
-			if (!logged) {
-				LOGF("RenderTarget: temporary: ignoring PS5 metadata addresses cmask=0x%016" PRIx64
-				     " fmask=0x%016" PRIx64 " dcc=0x%016" PRIx64 "\n",
-				     rt.cmask.addr, rt.fmask.addr, rt.dcc_addr.addr);
-				logged = true;
-			}
-		}
 	}
 }

@@ -383,10 +373,9 @@ static void ClipPrint(const char* func, const HW::ClipControl& c) {
 static void ClipCheck(const HW::ClipControl& c) {
 	// dx_linear_attr_clip_enable preserves linear (noperspective) attributes at clip-generated
 	// vertices, which Vulkan provides as part of clipping and interpolation.
-	EXIT_NOT_IMPLEMENTED(c.user_clip_planes != 0 || c.user_clip_plane_mode != 0 ||
-	                     c.vertex_kill_any || c.user_clip_plane_negate_y ||
-	                     c.user_clip_plane_cull_only || c.cull_on_clipping_error_disable ||
-	                     c.force_viewport_index_from_vs_enable);
+	EXIT_NOT_IMPLEMENTED(c.user_clip_plane_mode != 0 || c.vertex_kill_any ||
+	                     c.user_clip_plane_negate_y || c.user_clip_plane_cull_only ||
+	                     c.cull_on_clipping_error_disable || c.force_viewport_index_from_vs_enable);
 }

 static void RcPrint(const char* func, const HW::RenderControl& c) {
@@ -311,7 +311,6 @@ void RenderExecutor::ResolveRenderDepthTarget(CommandBuffer& buffer, RenderDepth
 	}
 	r.desc = MakeDepthTargetDesc(buffer, z);
 	r.depth_clear_enable      = rc.depth_clear_enable;
-	r.depth_load_clear_enable = r.depth_clear_enable;
 	r.depth_clear_value       = hw.GetDepthClearValue();
 	r.depth_test_enable       = dc.z_enable;
 	r.depth_write_enable      = r.depth_test_enable && dc.z_write_enable &&
@@ -421,9 +420,9 @@ vk::ImageAspectFlags RenderDepthInfo::AttachmentWriteAspects() const {
 	}

 	const auto           available = ImageViewOps::DepthAspectMask(format);
-	vk::ImageAspectFlags writes {};
+	vk::ImageAspectFlags writes = load_clear_aspects & available;
 	if ((available & vk::ImageAspectFlagBits::eDepth) &&
-	    (depth_load_clear_enable || depth_write_enable)) {
+	    depth_write_enable) {
 		writes |= vk::ImageAspectFlagBits::eDepth;
 	}
 	if (!(available & vk::ImageAspectFlagBits::eStencil)) {
@@ -461,8 +460,7 @@ vk::ImageAspectFlags RenderDepthInfo::AttachmentWriteAspects() const {
 		       (can_pass && depth_pass && state.passOp != vk::StencilOp::eKeep) ||
 		       (can_pass && depth_fail && state.depthFailOp != vk::StencilOp::eKeep);
 	};
-	if (stencil_clear_enable ||
-	    (stencil_test_enable && (face_writes(stencil_front) || face_writes(stencil_back)))) {
+	if (stencil_test_enable && (face_writes(stencil_front) || face_writes(stencil_back))) {
 		writes |= vk::ImageAspectFlagBits::eStencil;
 	}
 	return writes;
@@ -22,7 +22,7 @@ struct RenderDepthInfo {
 	// Discovery keeps guest image information but can remap the view into a larger cache image.
 	TextureCache::ImageDesc     desc;
 	bool                        depth_clear_enable       = false;
-	bool                        depth_load_clear_enable  = false;
+	vk::ImageAspectFlags        load_clear_aspects;
 	float                       depth_clear_value        = 0.0f;
 	bool                        depth_test_enable        = false;
 	// Effective draw writes; discovery applies test, target-write and clear controls.
@@ -30,16 +30,18 @@ namespace {
 	}
 }

-[[nodiscard]] vk::ImageCreateFlags ImageCreateFlags(const GraphicContext& graphics,
-                                                   const ImageInfo& info) {
+[[nodiscard]] vk::ImageCreateFlags
+ImageCreateFlags(const GraphicContext& graphics, const ImageInfo& info, vk::ImageUsageFlags usage) {
 	vk::ImageCreateFlags flags {};
 	if (DepthAspectTransferFormat(info.pixel_format) == vk::Format::eUndefined) {
 		flags |= vk::ImageCreateFlagBits::eMutableFormat;
-		flags |= vk::ImageCreateFlagBits::eExtendedUsage;
 		if (info.IsBlock() && graphics.supports_block_texel_view) {
 			flags |= vk::ImageCreateFlagBits::eBlockTexelViewCompatible;
 		}
 	}
+	if (usage & vk::ImageUsageFlagBits::eStorage) {
+		flags |= vk::ImageCreateFlagBits::eExtendedUsage;
+	}
 	if (info.IsVolume()) {
 		flags |= vk::ImageCreateFlagBits::e2DArrayCompatible;
 	}
@@ -59,7 +61,7 @@ namespace {
 			const auto storage = usage | vk::ImageUsageFlagBits::eStorage;
 			if (graphics.GetImageFormatProperties(info.pixel_format, HostImageType(info.type),
 			                                      vk::ImageTiling::eOptimal, storage,
-			                                      ImageCreateFlags(graphics, info),
+			                                      ImageCreateFlags(graphics, info, storage),
 			                                      nullptr) == vk::Result::eSuccess) {
 				usage = storage;
 			} else {
@@ -88,7 +90,8 @@ namespace {
 	if (HasFormatFeature(properties, vk::FormatFeatureFlagBits::eColorAttachment)) {
 		usage |= vk::ImageUsageFlagBits::eColorAttachment;
 	}
-	if (info.samples == 1) {
+	if (graphics.GetPhysicalDeviceProperties().limits.storageImageSampleCounts &
+	    static_cast<vk::SampleCountFlagBits>(info.samples)) {
 		usage |= vk::ImageUsageFlagBits::eStorage;
 	}
 	return usage;
@@ -209,6 +212,13 @@ Image::Barriers Image::GetBarriers(vk::ImageLayout                      destinat
 		barriers.push_back(barrier);
 	}

+	// Resident heaps sample their images in the read state; only leaving it invalidates them.
+	if (std::ranges::any_of(barriers, [&](const vk::ImageMemoryBarrier2& barrier) {
+		    return barrier.oldLayout == ReadLayout() &&
+		           barrier.srcAccessMask == vk::AccessFlagBits2::eShaderRead;
+	    })) {
+		HeapChanged();
+	}
 	state = {destination_stage, destination_access, destination_layout};
 	return barriers;
 }
@@ -353,8 +363,12 @@ void Image::CopyImage(Image& source) {
 	std::vector<vk::ImageCopy> copies;
 	copies.reserve(levels);
 	for (uint32_t level = 0; level < levels; level++) {
-		const auto width  = std::max(source.backing.extent.width >> level, 1u);
-		const auto height = std::max(source.backing.extent.height >> level, 1u);
+		// Overlapping cache images can differ in size; copying the source extent into a smaller
+		// destination writes out of bounds and loses the device.
+		const auto width  = std::min(std::max(source.backing.extent.width >> level, 1u),
+		                             std::max(backing.extent.width >> level, 1u));
+		const auto height = std::min(std::max(source.backing.extent.height >> level, 1u),
+		                             std::max(backing.extent.height >> level, 1u));
 		const auto depth  = std::max(base_depth >> level, 1u);
 		const auto [source_layers, destination_layers] = SanitizeCopyLayers(source, *this, depth);
 		vk::ImageCopy copy {};
@@ -698,7 +712,8 @@ Image::Image(GraphicContext& graphics, CommandScheduler& scheduler, const ImageI
 	}

 	vk::ImageCreateInfo create {};
-	create.flags         = ImageCreateFlags(graphics, info);
+	create.usage         = ImageUsageFlags(graphics, info);
+	create.flags         = ImageCreateFlags(graphics, info, create.usage);
 	create.imageType     = HostImageType(info.type);
 	create.extent        = info.extent;
 	create.mipLevels     = info.resources.levels;
@@ -706,7 +721,6 @@ Image::Image(GraphicContext& graphics, CommandScheduler& scheduler, const ImageI
 	create.format        = info.pixel_format;
 	create.tiling        = vk::ImageTiling::eOptimal;
 	create.initialLayout = vk::ImageLayout::eUndefined;
-	create.usage         = ImageUsageFlags(graphics, info);
 	create.samples       = vulkan_sample_count(info.samples);

 	vk::ImageFormatProperties properties {};
@@ -7,6 +7,7 @@
 #include "graphics/host_gpu/graphicContext.h"
 #include "graphics/host_gpu/renderer/image/imageInfo.h"

+#include <atomic>
 #include <compare>
 #include <limits>
 #include <optional>
@@ -71,6 +72,7 @@ public:
 	void MarkCpuDirty() {
 		m_cpu_dirty       = true;
 		m_maybe_cpu_dirty = false;
+		HeapChanged();
 	}

 	[[nodiscard]] bool IsCpuDirty() const { return m_cpu_dirty || m_maybe_cpu_dirty; }
@@ -80,6 +82,7 @@ public:
 		if (!IsCpuDirty()) {
 			m_maybe_cpu_hash  = HashGuestEdges();
 			m_maybe_cpu_dirty = true;
+			HeapChanged();
 		}
 	}
 	[[nodiscard]] bool ResolveMaybeCpuHash(uint64_t hash) {
@@ -104,7 +107,10 @@ public:
 	void               ClearGpuModified() noexcept { m_gpu_modified = false; }

 	[[nodiscard]] bool IsBufferModified() const noexcept { return m_buffer_modified; }
-	void               MarkBufferModified() noexcept { m_buffer_modified = true; }
+	void               MarkBufferModified() noexcept {
+		m_buffer_modified = true;
+		HeapChanged();
+	}
 	void               ClearBufferModified() noexcept { m_buffer_modified = false; }

 	[[nodiscard]] bool Overlaps(uint64_t address, uint64_t size,
@@ -120,6 +126,19 @@ public:
 		return backing.image == nullptr ? 0 : Common::AlignUp(info.data.size, 1024);
 	}
 	[[nodiscard]] uint64_t HashGuestEdges() const;
+	[[nodiscard]] vk::ImageLayout ReadLayout() const noexcept {
+		return info.IsDepth() ? vk::ImageLayout::eDepthStencilReadOnlyOptimal
+		                      : vk::ImageLayout::eShaderReadOnlyOptimal;
+	}
+	// Resident image heaps sample an image without rediscovering it until one changes its
+	// contents, tracking, identity or layout.
+	void HeapChanged() const noexcept {
+		if (heap_refs.load(std::memory_order_relaxed) != 0u) {
+			heap_epoch.fetch_add(1u, std::memory_order_relaxed);
+		}
+	}
+
+	static inline std::atomic<uint64_t> heap_epoch {0};

 	ImageInfo        info;
 	VulkanImage      backing;
@@ -137,6 +156,8 @@ public:
 	ImageSubresourceRange stencil_subresources;
 	uint64_t         tick_accessed_last = 0;
 	size_t           lru_id             = 0;
+	// Resident image heap elements that select this image.
+	std::atomic<uint32_t> heap_refs {0};

 private:
 	friend struct ImageTestAccess;
@@ -101,6 +101,7 @@ IsSupportedStorageImageResource(const ShaderRecompiler::IR::ImageResource& resou
 	       (resource.dimension == ShaderRecompiler::Decoder::ImageDimension::Dim1D ||
 	        resource.dimension == ShaderRecompiler::Decoder::ImageDimension::Dim1DArray ||
 	        resource.dimension == ShaderRecompiler::Decoder::ImageDimension::Dim2D ||
+	        resource.dimension == ShaderRecompiler::Decoder::ImageDimension::Dim2DMsaa ||
 	        resource.dimension == ShaderRecompiler::Decoder::ImageDimension::Dim3D ||
 	        resource.dimension == ShaderRecompiler::Decoder::ImageDimension::Dim2DArray) &&
 	       ((resource.mip_mode == ShaderRecompiler::IR::ImageMipMode::None &&
@@ -14,6 +14,7 @@
 #include "gpu_tiler_shaders/gpu_tiler_standard64_3d_spv.h"
 #include "gpu_tiler_shaders/gpu_tiler_standard64_spv.h"
 #include "gpu_tiler_shaders/gpu_tiler_color_transform_spv.h"
+#include "gpu_tiler_shaders/gpu_tiler_color_clear_spv.h"
 #include "graphics/host_gpu/graphicContext.h"
 #include "graphics/host_gpu/renderer/cache/streamBuffer.h"
 #include "graphics/host_gpu/renderer/commandScheduler.h"
@@ -71,20 +72,11 @@ TileManager::~TileManager() {
 			m_graphics.device.destroyPipeline(pipeline, nullptr);
 		}
 	}
-	if (m_d16_to_d24 != nullptr) {
-		m_graphics.device.destroyPipeline(m_d16_to_d24, nullptr);
-	}
-	if (m_d16_to_d32 != nullptr) {
-		m_graphics.device.destroyPipeline(m_d16_to_d32, nullptr);
-	}
-	if (m_d24_to_d16 != nullptr) {
-		m_graphics.device.destroyPipeline(m_d24_to_d16, nullptr);
-	}
-	if (m_d32_to_d16 != nullptr) {
-		m_graphics.device.destroyPipeline(m_d32_to_d16, nullptr);
-	}
-	if (m_color_transform != nullptr) {
-		m_graphics.device.destroyPipeline(m_color_transform, nullptr);
+	for (auto pipeline: {m_d16_to_d24, m_d16_to_d32, m_d24_to_d16, m_d32_to_d16, m_color_transform,
+	                     m_color_clear}) {
+		if (pipeline != nullptr) {
+			m_graphics.device.destroyPipeline(pipeline, nullptr);
+		}
 	}
 	if (m_pipeline_layout != nullptr) {
 		m_graphics.device.destroyPipelineLayout(m_pipeline_layout, nullptr);
@@ -220,42 +212,37 @@ vk::Pipeline TileManager::GetPipeline(uint32_t slot) {
 	if (m_pipelines[slot] != nullptr) {
 		return m_pipelines[slot];
 	}
-	struct Shader {
-		const uint32_t* code;
-		size_t          words;
-	};
-	static constexpr std::array<Shader, FamilyCount> shaders {{
-	    {GPU_TILER_STANDARD256_SPV, std::size(GPU_TILER_STANDARD256_SPV)},
-	    {GPU_TILER_STANDARD4_SPV, std::size(GPU_TILER_STANDARD4_SPV)},
-	    {GPU_TILER_STANDARD4_3D_SPV, std::size(GPU_TILER_STANDARD4_3D_SPV)},
-	    {GPU_TILER_STANDARD64_SPV, std::size(GPU_TILER_STANDARD64_SPV)},
-	    {GPU_TILER_STANDARD64_3D_SPV, std::size(GPU_TILER_STANDARD64_3D_SPV)},
-	    {GPU_TILER_PRT_SPV, std::size(GPU_TILER_PRT_SPV)},
-	    {GPU_TILER_PRT_3D_SPV, std::size(GPU_TILER_PRT_3D_SPV)},
-	    {GPU_TILER_RENDER_TARGET_SPV, std::size(GPU_TILER_RENDER_TARGET_SPV)},
-	    {GPU_TILER_DEPTH_SPV, std::size(GPU_TILER_DEPTH_SPV)},
-	}};
-	const uint32_t                                   element_index = slot % BytesPerElementCount;
+	static constexpr std::array<std::span<const uint32_t>, FamilyCount> shaders {
+	    GPU_TILER_STANDARD256_SPV, GPU_TILER_STANDARD4_SPV,     GPU_TILER_STANDARD4_3D_SPV,
+	    GPU_TILER_STANDARD64_SPV,  GPU_TILER_STANDARD64_3D_SPV, GPU_TILER_PRT_SPV,
+	    GPU_TILER_PRT_3D_SPV,      GPU_TILER_RENDER_TARGET_SPV, GPU_TILER_DEPTH_SPV};
+	const uint32_t                   element_index   = slot % BytesPerElementCount;
 	const uint32_t                   direction_index = slot / (FamilyCount * BytesPerElementCount);
 	const uint32_t                   family_index    = (slot / BytesPerElementCount) % FamilyCount;
 	const uint32_t                   values[] {1u << element_index, direction_index};
 	const vk::SpecializationMapEntry entries[] {{0, 0, 4}, {1, 4, 4}};
 	const vk::SpecializationInfo     specialization {2, entries, sizeof(values), values};
-	const auto module =
-	    CompileSPV({shaders[family_index].code, shaders[family_index].words}, m_graphics.device);
+	m_pipelines[slot] = CreatePipeline(shaders[family_index], &specialization);
+	return m_pipelines[slot];
+}
+
+vk::Pipeline TileManager::CreatePipeline(std::span<const uint32_t>     code,
+                                         const vk::SpecializationInfo* specialization) const {
+	const auto                        module = CompileSPV(code, m_graphics.device);
 	vk::PipelineShaderStageCreateInfo stage {};
 	stage.stage               = vk::ShaderStageFlagBits::eCompute;
 	stage.module              = module;
 	stage.pName               = "main";
-	stage.pSpecializationInfo = &specialization;
+	stage.pSpecializationInfo = specialization;
 	vk::ComputePipelineCreateInfo create {};
-	create.stage  = stage;
-	create.layout = m_pipeline_layout;
-	const auto result =
-	    m_graphics.device.createComputePipelines(nullptr, 1, &create, nullptr, &m_pipelines[slot]);
+	create.stage          = stage;
+	create.layout         = m_pipeline_layout;
+	vk::Pipeline pipeline = nullptr;
+	const auto   result =
+	    m_graphics.device.createComputePipelines(nullptr, 1, &create, nullptr, &pipeline);
 	m_graphics.device.destroyShaderModule(module, nullptr);
 	RequireVulkanSuccess(result, "create TileManager pipeline");
-	return m_pipelines[slot];
+	return pipeline;
 }

 void TileManager::Record(vk::Buffer source, uint64_t source_offset, uint64_t source_capacity,
@@ -434,39 +421,21 @@ uint32_t TileManager::ConversionRows(uint64_t offset, uint64_t row_stride, uint6

 void TileManager::ConvertD16(Result source, Result target, D16Direction direction, bool d32,
                             const D16Layout& layout) {
-	vk::Pipeline* pipeline_pointer = nullptr;
+	vk::Pipeline*             pipeline_pointer = nullptr;
+	std::span<const uint32_t> code;
 	if (direction == D16Direction::Promote) {
 		pipeline_pointer = d32 ? &m_d16_to_d32 : &m_d16_to_d24;
+		code             = GPU_TILER_PROMOTE_D16_SPV;
 	} else {
 		pipeline_pointer = d32 ? &m_d32_to_d16 : &m_d24_to_d16;
+		code             = GPU_TILER_DEMOTE_D16_SPV;
 	}
 	auto& pipeline = *pipeline_pointer;
 	if (pipeline == nullptr) {
 		const uint32_t                   value = d32 ? 1u : 0u;
 		const vk::SpecializationMapEntry entry {0, 0, sizeof(value)};
 		const vk::SpecializationInfo     specialization {1, &entry, sizeof(value), &value};
-		const uint32_t*                  code  = nullptr;
-		size_t                           words = 0;
-		if (direction == D16Direction::Promote) {
-			code  = GPU_TILER_PROMOTE_D16_SPV;
-			words = std::size(GPU_TILER_PROMOTE_D16_SPV);
-		} else {
-			code  = GPU_TILER_DEMOTE_D16_SPV;
-			words = std::size(GPU_TILER_DEMOTE_D16_SPV);
-		}
-		const auto module = CompileSPV({code, words}, m_graphics.device);
-		vk::PipelineShaderStageCreateInfo stage {};
-		stage.stage               = vk::ShaderStageFlagBits::eCompute;
-		stage.module              = module;
-		stage.pName               = "main";
-		stage.pSpecializationInfo = &specialization;
-		vk::ComputePipelineCreateInfo create {};
-		create.stage  = stage;
-		create.layout = m_pipeline_layout;
-		const auto result =
-		    m_graphics.device.createComputePipelines(nullptr, 1, &create, nullptr, &pipeline);
-		m_graphics.device.destroyShaderModule(module, nullptr);
-		RequireVulkanSuccess(result, "create D16 conversion pipeline");
+		pipeline = CreatePipeline(code, &specialization);
 	}

 	const uint64_t source_element =
@@ -599,18 +568,7 @@ void TileManager::TransformColor(Result input, Result output, ColorTransform tra
 	        output.size < input.size);
 	const auto pixels = static_cast<uint32_t>(input.size / element_bytes);
 	if (m_color_transform == nullptr) {
-		const auto module = CompileSPV(GPU_TILER_COLOR_TRANSFORM_SPV, m_graphics.device);
-		vk::PipelineShaderStageCreateInfo stage {};
-		stage.stage  = vk::ShaderStageFlagBits::eCompute;
-		stage.module = module;
-		stage.pName  = "main";
-		vk::ComputePipelineCreateInfo create {};
-		create.stage  = stage;
-		create.layout = m_pipeline_layout;
-		const auto result =
-		    m_graphics.device.createComputePipelines(nullptr, 1, &create, nullptr, &m_color_transform);
-		m_graphics.device.destroyShaderModule(module, nullptr);
-		RequireVulkanSuccess(result, "create color transform pipeline");
+		m_color_transform = CreatePipeline(GPU_TILER_COLOR_TRANSFORM_SPV, nullptr);
 	}
 	const auto input_binding  = BindStorage(input, input.size);
 	const auto output_binding = BindStorage(output, input.size);
@@ -670,4 +628,60 @@ TileManager::Result TileManager::TransformColor(Result input, ColorTransform tra
 	return result;
 }

+vk::Buffer TileManager::CheckColorClear(Result metadata, uint32_t slices, uint32_t keys,
+                                        bool expand) {
+	struct Check {
+		uint32_t metadata_base;
+		uint32_t words;
+		uint32_t keys;
+		uint32_t expand;
+	};
+	static_assert(sizeof(Check) <= sizeof(Conversion));
+	const uint32_t max_slices =
+	    m_graphics.GetPhysicalDeviceProperties().limits.maxFramebufferLayers;
+	EXIT_IF(slices == 0 || slices > max_slices || metadata.size % (4ull * slices) != 0 ||
+	        metadata.size / 4 > UINT32_MAX);
+	if (m_color_clear == nullptr) {
+		m_color_clear      = CreatePipeline(GPU_TILER_COLOR_CLEAR_SPV, nullptr);
+		m_clear_predicates = std::make_unique<Buffer>(
+		    m_graphics, m_scheduler, MemoryUsage::DeviceLocal, 0,
+		    vk::BufferUsageFlagBits::eStorageBuffer |
+		        vk::BufferUsageFlagBits::eConditionalRenderingEXT,
+		    uint64_t {max_slices} * ColorClearSlots * sizeof(uint32_t));
+	}
+	const auto source = BindStorage(metadata, metadata.size);
+	const vk::DescriptorBufferInfo infos[] {source.info,
+	                                        {m_clear_predicates->Handle(), 0, VK_WHOLE_SIZE}};
+	std::array<vk::WriteDescriptorSet, 2> writes {};
+	for (uint32_t index = 0; index < writes.size(); index++) {
+		writes[index].dstBinding      = index;
+		writes[index].descriptorCount = 1;
+		writes[index].descriptorType  = vk::DescriptorType::eStorageBuffer;
+		writes[index].pBufferInfo     = &infos[index];
+	}
+	// Orders the metadata writers and the previous check's predicate reads before this check.
+	vk::MemoryBarrier barrier {
+	    .srcAccessMask = vk::AccessFlagBits::eMemoryWrite,
+	    .dstAccessMask = vk::AccessFlagBits::eShaderRead | vk::AccessFlagBits::eShaderWrite};
+	m_scheduler.EndRendering();
+	auto command = m_scheduler.Current().Handle();
+	command.pipelineBarrier(vk::PipelineStageFlagBits::eAllCommands,
+	                        vk::PipelineStageFlagBits::eComputeShader, {}, 1, &barrier, 0, nullptr,
+	                        0, nullptr);
+	command.bindPipeline(vk::PipelineBindPoint::eCompute, m_color_clear);
+	command.pushDescriptorSetKHR(vk::PipelineBindPoint::eCompute, m_pipeline_layout, 0,
+	                             static_cast<uint32_t>(writes.size()), writes.data());
+	const Check push {source.base, static_cast<uint32_t>(metadata.size / 4 / slices), keys,
+	                  expand};
+	command.pushConstants(m_pipeline_layout, vk::ShaderStageFlagBits::eCompute, 0, sizeof(push),
+	                      &push);
+	command.dispatch(slices, 1, 1);
+	barrier.srcAccessMask = vk::AccessFlagBits::eShaderWrite;
+	barrier.dstAccessMask = vk::AccessFlagBits::eMemoryRead | vk::AccessFlagBits::eMemoryWrite;
+	command.pipelineBarrier(vk::PipelineStageFlagBits::eComputeShader,
+	                        vk::PipelineStageFlagBits::eAllCommands, {}, 1, &barrier, 0, nullptr,
+	                        0, nullptr);
+	return m_clear_predicates->Handle();
+}
+
 } // namespace Libs::Graphics
@@ -54,6 +54,13 @@ public:
 	                                const D16Layout& layout);
 	[[nodiscard]] Result TransformColor(Result input, ColorTransform transform, bool to_host);
 	void TransformColor(Result input, Result output, ColorTransform transform, bool to_host);
+	// Color-clear keys are multiples of 0x20, so key >> 5 selects one of these slots.
+	static constexpr uint32_t ColorClearSlots = 8;
+	// Returns the buffer whose dword slice * ColorClearSlots + (key >> 5) is set when every byte
+	// of that metadata slice equals a key whose bit `keys` selects; a hit expands to 0xff when
+	// `expand` is set.
+	[[nodiscard]] vk::Buffer CheckColorClear(Result metadata, uint32_t slices, uint32_t keys,
+	                                         bool expand);

 private:
 	friend struct TileManagerTestAccess;
@@ -117,6 +124,8 @@ private:
 	                            vk::Buffer target, uint64_t target_offset, uint64_t target_capacity,
 	                            const Tiling& tiling, bool clear_target);
 	[[nodiscard]] vk::Pipeline GetPipeline(uint32_t slot);
+	[[nodiscard]] vk::Pipeline CreatePipeline(std::span<const uint32_t>     code,
+	                                          const vk::SpecializationInfo* specialization) const;

 	GraphicContext&                         m_graphics;
 	CommandScheduler&                       m_scheduler;
@@ -130,6 +139,8 @@ private:
 	vk::Pipeline                            m_d24_to_d16      = nullptr;
 	vk::Pipeline                            m_d32_to_d16      = nullptr;
 	vk::Pipeline                            m_color_transform = nullptr;
+	vk::Pipeline                            m_color_clear     = nullptr;
+	std::unique_ptr<Buffer>                 m_clear_predicates;
 };

 } // namespace Libs::Graphics
@@ -1,6 +1,7 @@
 #include "graphics/host_gpu/renderer/masterSemaphore.h"

 #include "common/assert.h"
+#include "common/debugCounters.h"
 #include "graphics/host_gpu/graphicContext.h"

 namespace Libs::Graphics {
@@ -49,6 +50,8 @@ void MasterSemaphore::Wait(uint64_t tick) {
 	wait_info.pSemaphores    = &m_semaphore;
 	wait_info.pValues        = &tick;

+	Common::DebugCounters::ScopedTimer timer(Common::DebugCounters::Counter::GpuWaits,
+	                                         Common::DebugCounters::Counter::GpuWaitNs);
 	const auto result = m_graphics.device.waitSemaphores(&wait_info, UINT64_MAX);
 	EXIT_NOT_IMPLEMENTED(result != vk::Result::eSuccess);
 	Refresh();
@@ -143,7 +143,8 @@ NativeStorageBuffer(RenderContext& context, const PreparedBindings::BufferSource
 	return result;
 }

-bool IsSupportedDepthTextureEncoding(const ShaderTextureResource& descriptor, bool r128) {
+// Reserved and unsupported T# fields must be zero; perf_mod only scales sampler state.
+static bool IsCanonicalTextureEncoding(const ShaderTextureResource& descriptor, bool r128) {
 	constexpr uint32_t field1_reserved_mask = 0x200fff00u;
 	constexpr uint32_t field2_reserved_mask = 0xf0003000u;
 	const uint32_t     field3_expected = descriptor.DstSelXYZW() |
@@ -154,17 +155,25 @@ bool IsSupportedDepthTextureEncoding(const ShaderTextureResource& descriptor, bo
 	const uint32_t     field4_expected = descriptor.Depth() | (descriptor.BaseArray5() << 16u);
 	const uint32_t     field5_expected = (static_cast<uint32_t>(descriptor.PerfMod5()) << 20u) |
 	                                     (static_cast<uint32_t>(descriptor.MaxMip()) << 4u);
-	const bool         common          = (descriptor.fields[1] & field1_reserved_mask) == 0 &&
-	                                     (descriptor.fields[2] & field2_reserved_mask) == 0 &&
-	                                     descriptor.fields[3] == field3_expected;
-	if (r128) {
-		return common && descriptor.fields[4] == 0 && descriptor.fields[5] == 0 &&
-		       descriptor.fields[6] == 0 && descriptor.fields[7] == 0;
+	if ((descriptor.fields[1] & field1_reserved_mask) != 0 ||
+	    (descriptor.fields[2] & field2_reserved_mask) != 0 || descriptor.fields[3] != field3_expected) {
+		return false;
 	}
-	const bool full = common && descriptor.fields[4] == field4_expected &&
-	                  descriptor.fields[5] == field5_expected;
-	if (!full ||
-	    (descriptor.MsaaDepth() && !IsMultisampledTexture(descriptor.Type()))) {
+	if (r128) {
+		return descriptor.fields[4] == 0 && descriptor.fields[5] == 0 && descriptor.fields[6] == 0 &&
+		       descriptor.fields[7] == 0;
+	}
+	return descriptor.fields[4] == field4_expected && descriptor.fields[5] == field5_expected;
+}
+
+bool IsSupportedDepthTextureEncoding(const ShaderTextureResource& descriptor, bool r128) {
+	if (!IsCanonicalTextureEncoding(descriptor, r128)) {
+		return false;
+	}
+	if (r128) {
+		return true;
+	}
+	if (descriptor.MsaaDepth() && !IsMultisampledTexture(descriptor.Type())) {
 		return false;
 	}
 	const auto metadata_control = descriptor.fields[6] & 0x00ffffffu;
@@ -245,6 +254,10 @@ static bool IsSupportedStorageTextureDescriptor(const ShaderRecompiler::IR::Imag
 	const bool is_3d = resource.dimension == ShaderRecompiler::Decoder::ImageDimension::Dim3D &&
 	                   descriptor.Type() == Prospero::ImageType::kColor3D &&
 	                   descriptor.BaseArray5() == 0;
+	const bool is_2d_msaa =
+	    resource.dimension == ShaderRecompiler::Decoder::ImageDimension::Dim2DMsaa &&
+	    descriptor.Type() == Prospero::ImageType::kColor2DMsaa && descriptor.Depth() == 0 &&
+	    descriptor.BaseArray5() == 0;
 	TileTextureBlockLayout tile_layout {};
 	bool                   supported_tile = false;
 	switch (tile) {
@@ -275,42 +288,18 @@ static bool IsSupportedStorageTextureDescriptor(const ShaderRecompiler::IR::Imag
 	const bool supported_swizzle =
 	    IsValidImageSwizzle(swizzle) &&
 	    (swizzle == DstSel(4, 5, 6, 7) || !resource.read || resource.atomic);
-	return (is_1d || is_1d_array || is_2d || is_2d_array || is_3d) && supported_tile &&
-	       descriptor.BaseLevel() <= descriptor.LastLevel() &&
+	return (is_1d || is_1d_array || is_2d || is_2d_array || is_3d || is_2d_msaa) &&
+	       supported_tile && descriptor.BaseLevel() <= descriptor.LastLevel() &&
 	       descriptor.MinLod() == 0 && supported_swizzle && descriptor.BCSwizzle() == 0 &&
 	       !descriptor.MsaaDepth();
 }

-static bool IsSupportedStorageTextureEncoding(const ShaderRecompiler::IR::ImageResource& resource,
-                                              const ShaderTextureResource& descriptor) {
-	constexpr uint32_t field1_reserved_mask = 0x200fff00u;
-	constexpr uint32_t field2_reserved_mask = 0xf0003000u;
-	constexpr uint32_t field5_expected      = 0x00700000u;
-	constexpr uint32_t field5_max_mip_mask  = 0x000000f0u;
-	const uint32_t     expected_field3 = descriptor.DstSelXYZW() |
-	                                     (static_cast<uint32_t>(descriptor.BaseLevel()) << 12u) |
-	                                     (static_cast<uint32_t>(descriptor.LastLevel()) << 16u) |
-	                                     (static_cast<uint32_t>(descriptor.TileMode()) << 20u) |
-	                                     (static_cast<uint32_t>(descriptor.Type()) << 28u);
-	const uint32_t     expected_field4 =
-	    descriptor.Depth() | (static_cast<uint32_t>(descriptor.BaseArray5()) << 16u);
-	const bool common = (descriptor.fields[1] & field1_reserved_mask) == 0 &&
-	                    (descriptor.fields[2] & field2_reserved_mask) == 0 &&
-	                    descriptor.fields[3] == expected_field3;
-	if (resource.r128) {
-		return common && descriptor.fields[4] == 0 && descriptor.fields[5] == 0 &&
-		       descriptor.fields[6] == 0 && descriptor.fields[7] == 0;
-	}
-	return common && descriptor.fields[4] == expected_field4 &&
-	       (descriptor.fields[5] & ~field5_max_mip_mask) == field5_expected;
-}
-
 void ValidateStorageTexture(const ShaderRecompiler::IR::ImageResource& resource,
                            const ShaderTextureResource& descriptor, uint64_t size) {
 	const auto format        = descriptor.Format();
 	const bool resource_ok   = IsSupportedStorageImageResource(resource);
 	const bool descriptor_ok = IsSupportedStorageTextureDescriptor(resource, descriptor);
-	const bool encoding_ok   = IsSupportedStorageTextureEncoding(resource, descriptor);
+	const bool encoding_ok   = IsCanonicalTextureEncoding(descriptor, resource.r128);
 	const bool uint_resource    = resource.numeric_class == Prospero::TextureNumericClass::Uint;
 	const bool raw_sint_storage = uint_resource &&
 	    ShaderRecompiler::IR::RawSignedStorageFormat(resource, format) != Prospero::BufferFormat::kInvalid;
@@ -371,7 +360,7 @@ static ImageViewInfo TextureViewInfo(const ShaderRecompiler::IR::ImageResource&
 	}
 	view.usage = storage ? vk::ImageUsageFlagBits::eStorage : vk::ImageUsageFlagBits::eSampled;
 	view.mapping =
-	    storage || surface_format.conversion_format != Prospero::BufferFormat::kInvalid
+	    storage || resource.conversion_format != Prospero::BufferFormat::kInvalid
 	        ? vk::ComponentMapping {}
 	        : TextureGetComponentMapping(descriptor.DstSelXYZW(), surface_format.host_to_storage);
 	switch (resource.dimension) {
@@ -530,7 +519,9 @@ TextureBinding RenderExecutor::ResolveTexture(const ShaderRecompiler::IR::ImageR
 	}
 	const auto samples = multisampled ? 1u << last_level : 1u;
 	const auto depth          = static_cast<uint32_t>(descriptor.Depth()) + 1u;
-	const auto format         = descriptor.Format();
+	const auto format         = ShaderRecompiler::IR::ReadsDepthBits(resource)
+	                                ? Prospero::BufferFormat::k16UNorm
+	                                : descriptor.Format();
 	const auto surface_format = TextureGetSurfaceFormatInfo(format);
 	const bool shader_conversion =
 	    surface_format.conversion_format != Prospero::BufferFormat::kInvalid;
@@ -594,6 +585,10 @@ TextureBinding RenderExecutor::ResolveTexture(const ShaderRecompiler::IR::ImageR
 	const auto alignment = desc.info.DataAlignment();
 	EXIT_NOT_IMPLEMENTED(desc.info.data.size == 0 || alignment == 0 ||
 	                     (address & (static_cast<uint64_t>(alignment) - 1u)) != 0);
+	// Image heaps keep the descriptors of unmapped textures, which PS5 never samples.
+	if (resource.heap && !m_context.IsMapped(address, desc.info.data.size)) {
+		return {.layout = vk::ImageLayout::eGeneral};
+	}
 	if (storage) {
 		ValidateStorageTexture(resource, descriptor, desc.info.data.size);
 	}
@@ -699,6 +694,10 @@ void RenderExecutor::PrepareBindings(const ShaderStageRuntime& runtime,
 	prepared.samplers.clear();
 	prepared.shader_data.clear();
 	for (uint32_t i = 0; i < program.info.images.size(); i++) {
+		if (program.info.images[i].heap) {
+			prepared.images[i] = {};
+			continue;
+		}
 		auto binding = ResolveTexture(program.info.images[i], snapshot.images[i]);
 		BindImage(binding.image_id, binding.desc.type == TextureCache::BindingType::Storage);
 		binding.mip_views.swap(prepared.images[i].mip_views);
@@ -740,6 +739,12 @@ void RenderExecutor::FindBuffers(std::span<PreparedBindings* const> stages) {
 			const auto descriptor = DecodeNativeDescriptor<ShaderBufferResource>(snapshot.buffers[resource]);
 			const auto address = descriptor.Base48();
 			auto size = descriptor.GetSize();
+			if (size != 0 && descriptor.Stride() == 0 && descriptor.OutOfBounds() == 2 &&
+			    !metadata.scalar) {
+				// OOB_SELECT=2 only tests NumRecords != 0; a zero-stride fetch
+				// still needs its full payload when NumRecords is smaller than it.
+				size = std::max(size, uint64_t {metadata.max_byte_extent});
+			}
 			if (address == 0 || size == 0) {
 				source = {};
 				continue;
@@ -757,14 +762,18 @@ void RenderExecutor::FindBuffers(std::span<PreparedBindings* const> stages) {
 							}
 						}
 					}
-				} else {
-					const auto& graphics = m_context.GetGraphics();
-					const auto limit = uint64_t {graphics.GetPhysicalDeviceProperties().limits.maxStorageBufferRange} -
-					                   (graphics.StorageMinAlignment() - 1);
-					size = std::min({size, uint64_t {256} * 1024 * 1024, limit});
 				}
+				//  PPSA15246
+				const auto& graphics = m_context.GetGraphics();
+				const auto  limit =
+				    uint64_t {graphics.GetPhysicalDeviceProperties().limits.maxStorageBufferRange} -
+				    (graphics.StorageMinAlignment() - 1);
+				size = std::min({size, uint64_t {256} * 1024 * 1024, limit});
+			}
+			// Finite GPU descriptors can span sparse PRT holes; clamp only oversized/sentinel ranges.
+			if (size >= (uint64_t {1} << 30u) || descriptor.NumRecords() == UINT32_MAX) {
+				size = Libs::LibKernel::Memory::ClampRangeSize(address, size);
 			}
-			size = Libs::LibKernel::Memory::ClampRangeSize(address, size);
 			source = {address, size, cache.FindBuffer(address, size)};
 		}
 	}
@@ -851,14 +860,19 @@ void RenderExecutor::RebindImages(PreparedBindings& prepared) {

 void RenderExecutor::PrepareGraphicsBindings(std::span<PreparedBindings* const> stages,
                                             std::span<RenderColorInfo> colors) {
-	bool uses_dma = false;
+	bool uses_dma        = false;
+	bool uses_image_heap = false;
 	FindBuffers(stages);
 	for (auto* stage: stages) {
 		uses_dma |= stage->runtime->program->info.uses_dma;
+		uses_image_heap |= stage->runtime->program->info.uses_image_heap;
 	}
 	if (uses_dma) {
 		m_context.PrepareBda();
 	}
+	if (uses_image_heap) {
+		m_context.GetImageHeap().Rebind(*this);
+	}
 	for (auto* stage: stages) {
 		RebindImages(*stage);
 	}
@@ -946,14 +960,21 @@ void RenderExecutor::CommitBindings(CommandBuffer&                     buffer,
 				for (const auto& resource: program.info.buffers) {
 					if (!resource.written || resource.descriptor_index == UINT32_MAX) continue;
 					const auto& written = writer->buffer_sources[resource.descriptor_index];
-					if (written.size != 0 && ImageRangeOverlaps(address, size,
-					                                          written.address, written.size)) {
+					// ADD_TID and swizzle move fixed-offset writes past write_extent.
+					const bool linear = (resource.packed_stride & ((1u << 14u) | (1u << 20u))) == 0u;
+					const auto written_size =
+					    linear ? std::min<uint64_t>(written.size, resource.write_extent) : written.size;
+					if (written_size != 0 && ImageRangeOverlaps(address, size,
+					                                          written.address, written_size)) {
 						EXIT("scalar resource reads overlap a shader buffer write\n");
 					}
 				}
 			}
 		}
 	}
+	if (pipeline.uses_image_heap) {
+		m_context.GetImageHeap().Commit(buffer, pipeline_bind_point, pipeline.pipeline_layout);
+	}
 	m_descriptor_images.resize(image_descriptor_count);
 	m_descriptor_writes.clear();
 	m_descriptor_writes.reserve(write_count);
@@ -1011,9 +1032,7 @@ void RenderExecutor::CommitBindings(CommandBuffer&                     buffer,
 				              vk::AccessFlagBits2::eShaderRead | vk::AccessFlagBits2::eShaderWrite,
 				              range, vk_buffer);
 			} else {
-				image.Transit(image.info.IsDepth() ? vk::ImageLayout::eDepthStencilReadOnlyOptimal
-				                                   : vk::ImageLayout::eShaderReadOnlyOptimal,
-				              vk::AccessFlagBits2::eShaderRead, range, vk_buffer);
+				image.Transit(image.ReadLayout(), vk::AccessFlagBits2::eShaderRead, range, vk_buffer);
 			}
 			binding.layout = image.backing.state.layout;
 		}
@@ -1084,6 +1103,7 @@ void RenderExecutor::CommitBindings(CommandBuffer&                     buffer,
 		for (uint32_t i = 0; i < program.info.images.size(); ++i) {
 			const auto& resource = program.info.images[i];
 			const auto& texture = descriptors.images[i];
+			if (resource.heap) continue;
 			EXIT_IF(texture.image_id &&
 			        resource.mip_count != (texture.mip_views.empty() ? 1u : texture.mip_views.size()));
 			const auto kind = *ShaderRecompiler::IR::DescriptorBindingForImage(resource);
@@ -0,0 +1,544 @@
+#include "graphics/host_gpu/renderer/pipeline/imageHeap.h"
+
+#include "common/alignment.h"
+#include "common/assert.h"
+#include "common/profiler.h"
+#include "graphics/host_gpu/pageManager.h"
+#include "graphics/host_gpu/regionDefinitions.h"
+#include "graphics/host_gpu/renderer/cache/streamBuffer.h"
+#include "graphics/host_gpu/renderer/render.h"
+#include "graphics/host_gpu/renderer/renderContext.h"
+#include "graphics/shader/recompiler/ir/passes/ResourceMaterialization.h"
+#include "kernel/memory.h"
+
+#include <algorithm>
+#include <xxhash.h>
+
+namespace Libs::Graphics {
+namespace {
+
+namespace IR = ShaderRecompiler::IR;
+
+// Smaller tables are cheap to materialize on each draw and are often transient.
+constexpr uint64_t MinTableBytes = 64ull * 1024ull;
+constexpr uint32_t ElementMask   = (1u << IR::ImageHeapKindShift) - 1u;
+
+template <typename Pred, typename Fn>
+void ForEachRun(size_t begin, size_t end, Pred&& pred, Fn&& fn) {
+	while (begin < end) {
+		if (!pred(begin)) {
+			++begin;
+			continue;
+		}
+		auto run = begin + 1u;
+		while (run < end && pred(run)) ++run;
+		fn(begin, run);
+		begin = run;
+	}
+}
+
+// Grows [begin, end) to cover [first, last).
+void Widen(size_t& begin, size_t& end, size_t first, size_t last) {
+	const bool empty = begin == end;
+	begin            = empty ? first : std::min(begin, first);
+	end              = empty ? last : std::max(end, last);
+}
+
+// Guest bytes of a table that its tracker pages [begin, end) cover.
+std::pair<uint64_t, uint64_t> PageRange(uint64_t base, uint64_t size, size_t begin, size_t end) {
+	const auto origin = Common::AlignDown(base, TRACKER_PAGE_SIZE);
+	return {std::max(base, origin + begin * TRACKER_PAGE_SIZE),
+	        std::min(base + size, origin + end * TRACKER_PAGE_SIZE)};
+}
+
+} // namespace
+
+size_t ImageHeap::DescriptorHash::operator()(const DescriptorValue& value) const noexcept {
+	return XXH3_64bits(value.dwords.data(), sizeof(value.dwords));
+}
+
+ImageHeap::ImageHeap(RenderContext& context, PageManager& page_manager)
+    : m_context(context), m_page_manager(page_manager) {
+	auto& graphics = context.GetGraphics();
+	vk::PhysicalDeviceVulkan12Properties limits {};
+	vk::PhysicalDeviceProperties2        properties {};
+	properties.pNext = &limits;
+	graphics.physical_device.getProperties2(&properties);
+	EXIT_NOT_IMPLEMENTED(std::min({limits.maxDescriptorSetUpdateAfterBindSampledImages,
+	                               limits.maxPerStageDescriptorUpdateAfterBindSampledImages,
+	                               limits.maxPerStageUpdateAfterBindResources}) <
+	                     IR::ImageHeapCapacity);
+
+	const vk::DescriptorBindingFlags binding_flags =
+	    vk::DescriptorBindingFlagBits::ePartiallyBound |
+	    vk::DescriptorBindingFlagBits::eUpdateAfterBind |
+	    vk::DescriptorBindingFlagBits::eUpdateUnusedWhilePending;
+	vk::DescriptorSetLayoutBindingFlagsCreateInfo flags_info {};
+	flags_info.bindingCount  = 1;
+	flags_info.pBindingFlags = &binding_flags;
+	const vk::DescriptorSetLayoutBinding binding {0, vk::DescriptorType::eSampledImage,
+	                                              IR::ImageHeapCapacity,
+	                                              vk::ShaderStageFlagBits::eAll, nullptr};
+	vk::DescriptorSetLayoutCreateInfo layout_info {};
+	layout_info.pNext        = &flags_info;
+	layout_info.flags        = vk::DescriptorSetLayoutCreateFlagBits::eUpdateAfterBindPool;
+	layout_info.bindingCount = 1;
+	layout_info.pBindings    = &binding;
+	EXIT_IF(graphics.device.createDescriptorSetLayout(&layout_info, nullptr, &m_layout) !=
+	        vk::Result::eSuccess);
+	const vk::DescriptorPoolSize pool_size {vk::DescriptorType::eSampledImage,
+	                                        IR::ImageHeapCapacity};
+	vk::DescriptorPoolCreateInfo pool_info {};
+	pool_info.flags         = vk::DescriptorPoolCreateFlagBits::eUpdateAfterBind;
+	pool_info.maxSets       = 1;
+	pool_info.poolSizeCount = 1;
+	pool_info.pPoolSizes    = &pool_size;
+	EXIT_IF(graphics.device.createDescriptorPool(&pool_info, nullptr, &m_pool) !=
+	        vk::Result::eSuccess);
+	vk::DescriptorSetAllocateInfo allocate {};
+	allocate.descriptorPool     = m_pool;
+	allocate.descriptorSetCount = 1;
+	allocate.pSetLayouts        = &m_layout;
+	EXIT_IF(graphics.device.allocateDescriptorSets(&allocate, &m_set) != vk::Result::eSuccess);
+	graphics.image_heap_layout = m_layout;
+	// Slot word 0 selects element 0, which stays a null image.
+	m_elements.push_back({.binding = {.layout = vk::ImageLayout::eShaderReadOnlyOptimal}});
+	m_writes.push_back(0u);
+	WriteDescriptors();
+}
+
+ImageHeap::~ImageHeap() {
+	for (const auto& memory: m_memories) Unwatch(*memory);
+	auto& graphics             = m_context.GetGraphics();
+	graphics.image_heap_layout = nullptr;
+	graphics.device.destroyDescriptorPool(m_pool, nullptr);
+	graphics.device.destroyDescriptorSetLayout(m_layout, nullptr);
+}
+
+bool ImageHeap::Acquire(const IR::ImageHeapRequest& request, IR::ImageHeapView& result) {
+	KYTY_PROFILER_FUNCTION();
+	const auto size = request.size & ~uint64_t {3};
+	if (size < MinTableBytes || request.count == 0u || request.step == 0u ||
+	    ((request.base | request.first | request.step) & 3u) != 0u) {
+		return false;
+	}
+	auto memory = std::ranges::find_if(m_memories, [&](const auto& candidate) {
+		return candidate->base == request.base && candidate->size == size;
+	});
+	if (memory == m_memories.end()) {
+		if (!m_context.IsMapped(request.base, size)) return false;
+		auto created   = std::make_unique<Memory>();
+		created->base  = request.base;
+		created->size  = size;
+		created->words.resize(size / 4u);
+		created->dead.resize(size / 4u);
+		created->dirty_pages.assign((Common::AlignUp(request.base + size, TRACKER_PAGE_SIZE) -
+		                             Common::AlignDown(request.base, TRACKER_PAGE_SIZE)) /
+		                                TRACKER_PAGE_SIZE,
+		                            uint8_t {1});
+		std::scoped_lock lock {m_mutex};
+		memory = m_memories.insert(m_memories.end(), std::move(created));
+	}
+	// A draw refreshes a table and a view takes its changes once, so all programs of the draw see
+	// one snapshot and one set of kinds.
+	if ((*memory)->draw != m_draw) {
+		if (!Refresh(**memory)) return false;
+		(*memory)->draw = m_draw;
+	}
+	auto&      views   = (*memory)->views;
+	const auto profile = IR::ImageHeapResource(*request.root, 0u);
+	auto       view    = std::ranges::find_if(views, [&](const auto& candidate) {
+		return candidate->step == request.step && candidate->first == request.first &&
+		       candidate->count == request.count && candidate->profile == profile;
+	});
+	if (view == views.end()) {
+		const auto epoch   = Image::heap_epoch.load(std::memory_order_relaxed);
+		auto       created = std::make_unique<View>(View {.memory        = **memory,
+		                                                  .step          = request.step,
+		                                                  .first         = request.first,
+		                                                  .count         = request.count,
+		                                                  .profile       = profile,
+		                                                  .resolve_epoch = epoch,
+		                                                  .layout_epoch  = epoch});
+		created->slots.resize(size / 4u + 1u);
+		created->buffer = std::make_unique<Buffer>(
+		    m_context.GetGraphics(), m_context.GetCommandScheduler(), MemoryUsage::DeviceLocal, 0,
+		    vk::BufferUsageFlagBits::eTransferDst | vk::BufferUsageFlagBits::eStorageBuffer |
+		        vk::BufferUsageFlagBits::eShaderDeviceAddress,
+		    size + 4u);
+		UpdateProbes(*created, 0u, size);
+		// The first upload also clears the slots no probe selects.
+		created->dirty_begin = 0;
+		created->dirty_end   = created->slots.size();
+		view = views.insert(views.end(), std::move(created));
+	} else {
+		// Empty after the first acquire in a draw, since the table refreshes once per draw.
+		UpdateProbes(**view, (*view)->changed_begin, (*view)->changed_end);
+		(*view)->changed_begin = (*view)->changed_end = 0;
+	}
+	if (std::ranges::find(m_active, view->get()) == m_active.end()) m_active.push_back(view->get());
+	const auto& selected = **view;
+	result = {.slots      = selected.buffer->BufferDeviceAddress(),
+	          .slot_count = static_cast<uint32_t>(selected.slots.size() - 1u),
+	          .kinds      = selected.kinds};
+	return true;
+}
+
+bool ImageHeap::Refresh(Memory& memory) {
+	// GPU writes leave no fault; reading the table downloads them.
+	const bool gpu_dirty = m_context.GetBufferCache().HasGpuDirtyBytes(memory.base, memory.size);
+	const auto pages     = memory.dirty_pages.size();
+	m_runs.clear();
+	{
+		std::scoped_lock lock {m_mutex};
+		// Watch before reading, so a write during the read faults the page again.
+		ForEachRun(
+		    0, pages, [&](size_t page) { return memory.dirty_pages[page] != 0u; },
+		    [&](size_t begin, size_t end) {
+			    std::fill(memory.dirty_pages.begin() + begin, memory.dirty_pages.begin() + end,
+			              uint8_t {0});
+			    Watch<true>(memory, begin, end);
+			    m_runs.emplace_back(begin, end);
+		    });
+	}
+	if (gpu_dirty) m_runs.assign(1u, {0u, pages});
+	for (size_t run = 0; run < m_runs.size(); ++run) {
+		const auto [address, end] = PageRange(memory.base, memory.size, m_runs[run].first,
+		                                      m_runs[run].second);
+		m_scratch.resize((end - address) / 4u);
+		if (!LibKernel::Memory::TryReadBufferBacking(address, m_scratch.data(), end - address)) {
+			std::scoped_lock lock {m_mutex};
+			for (; run < m_runs.size(); ++run) {
+				MarkDirty(memory, m_runs[run].first, m_runs[run].second);
+			}
+			return false;
+		}
+		const auto words = std::span(memory.words).subspan((address - memory.base) / 4u,
+		                                                   m_scratch.size());
+		size_t     first = 0;
+		size_t     last  = words.size();
+		while (first < last && words[first] == m_scratch[first]) ++first;
+		while (first < last && words[last - 1u] == m_scratch[last - 1u]) --last;
+		if (first == last) continue;
+		std::copy(m_scratch.begin() + first, m_scratch.begin() + last, words.begin() + first);
+		const auto offset = address - memory.base;
+		for (auto& view: memory.views) {
+			Widen(view->changed_begin, view->changed_end, offset + first * 4u, offset + last * 4u);
+		}
+	}
+	return true;
+}
+
+void ImageHeap::MarkDirty(Memory& memory, size_t begin, size_t end) {
+	ForEachRun(
+	    begin, end, [&](size_t page) { return memory.dirty_pages[page] == 0u; },
+	    [&](size_t first, size_t last) {
+		    std::fill(memory.dirty_pages.begin() + first, memory.dirty_pages.begin() + last,
+		              uint8_t {1});
+		    Watch<false>(memory, first, last);
+	    });
+}
+
+template <bool track>
+void ImageHeap::Watch(const Memory& memory, size_t begin, size_t end) {
+	const auto [address, limit] = PageRange(memory.base, memory.size, begin, end);
+	m_page_manager.UpdatePageWatchers<track>(address, limit - address);
+}
+
+void ImageHeap::Unwatch(const Memory& memory) {
+	ForEachRun(
+	    0, memory.dirty_pages.size(), [&](size_t page) { return memory.dirty_pages[page] == 0u; },
+	    [&](size_t begin, size_t end) { Watch<false>(memory, begin, end); });
+}
+
+void ImageHeap::InvalidateMemory(uint64_t address, uint64_t size) {
+	std::scoped_lock lock {m_mutex};
+	for (auto& memory: m_memories) {
+		const auto begin = std::max(address, memory->base);
+		const auto end   = std::min(address + size, memory->base + memory->size);
+		if (begin >= end) continue;
+		const auto origin = Common::AlignDown(memory->base, TRACKER_PAGE_SIZE);
+		MarkDirty(*memory, (begin - origin) / TRACKER_PAGE_SIZE,
+		          (end - 1u - origin) / TRACKER_PAGE_SIZE + 1u);
+	}
+}
+
+void ImageHeap::UpdateProbes(View& view, uint64_t begin, uint64_t end) {
+	// A probe reads its descriptor from first + probe * step.
+	const uint64_t bytes = view.profile.r128 ? 16u : 32u;
+	auto probe = begin < view.first + bytes ? 0u : (begin - view.first - bytes) / view.step + 1u;
+	for (; probe < view.count && view.first + probe * view.step < end; ++probe) {
+		UpdateProbe(view, probe);
+	}
+}
+
+void ImageHeap::UpdateProbe(View& view, uint64_t probe) {
+	const auto&     words  = view.memory.words;
+	const auto      offset = (view.first + probe * view.step) / 4u;
+	DescriptorValue value {.dword_count = 8u};
+	std::copy_n(words.begin() + offset,
+	            std::min<size_t>(view.profile.r128 ? 4u : 8u, words.size() - offset),
+	            value.dwords.begin());
+	// A buried word stays null in every view until its descriptor changes.
+	auto& dead = view.memory.dead[offset];
+	if (dead != 0u && dead != DescriptorHash {}(value)) dead = 0u;
+	const auto kind = dead != 0u ? 0u : IR::ImageHeapKind(view.profile, value);
+	Assign(view, offset, kind == 0u ? 0u : Intern(view, kind, value));
+}
+
+void ImageHeap::Assign(View& view, size_t slot, uint32_t entry) {
+	const auto previous = std::exchange(view.slots[slot], entry);
+	Release(view, previous);
+	if (previous != entry) Widen(view.dirty_begin, view.dirty_end, slot, slot + 1u);
+}
+
+void ImageHeap::Bury() {
+	if (m_unmapped.empty()) return;
+	for (const auto& memory: m_memories) {
+		for (const auto& view: memory->views) {
+			bool buried = false;
+			for (const auto& [value, index]: view->elements) {
+				const auto& binding = m_elements[index].binding;
+				if (binding.image_id && std::ranges::any_of(m_unmapped, [&](GuestRange range) {
+					    return ImageRangeOverlaps(binding.desc.info.data, range);
+				    })) {
+					m_elements[index].stale = buried = true;
+				}
+			}
+			if (!buried) continue;
+			for (uint64_t probe = 0; probe < view->count; ++probe) {
+				const auto  slot    = (view->first + probe * view->step) / 4u;
+				const auto& element = m_elements[view->slots[slot] & ElementMask];
+				if (!element.stale) continue;
+				memory->dead[slot] = DescriptorHash {}(element.value);
+				Assign(*view, slot, 0u);
+			}
+		}
+	}
+	m_unmapped.clear();
+}
+
+uint32_t ImageHeap::Intern(View& view, uint32_t kind, const DescriptorValue& value) {
+	const auto arm = static_cast<uint32_t>(std::ranges::find(view.kind_keys, kind) -
+	                                       view.kind_keys.begin());
+	if (arm == view.kind_keys.size()) {
+		view.kind_keys.push_back(kind);
+		view.kinds.push_back(value);
+	}
+	auto [found, inserted] = view.elements.try_emplace(value, 0u);
+	if (inserted) {
+		found->second    = Allocate();
+		auto& element    = m_elements[found->second];
+		element.value    = value;
+		element.kind     = kind;
+		view.pending.push_back(found->second);
+	}
+	++m_elements[found->second].refs;
+	return arm << IR::ImageHeapKindShift | found->second;
+}
+
+void ImageHeap::Release(View& view, uint32_t entry) {
+	const auto index = entry & ElementMask;
+	if (index == 0u || --m_elements[index].refs != 0u) return;
+	auto& element = m_elements[index];
+	if (const auto found = view.elements.find(element.value);
+	    found != view.elements.end() && found->second == index) {
+		view.elements.erase(found);
+	}
+	Unreference(element);
+	view.released.push_back(index);
+}
+
+void ImageHeap::Unreference(Element& element) {
+	if (auto* image = m_context.GetTextureCache().m_slot_images.try_get(element.binding.image_id)) {
+		image->heap_refs.fetch_sub(1u, std::memory_order_relaxed);
+		image->tick_accessed_last = m_context.GetCommandScheduler().CurrentTick();
+	}
+}
+
+uint32_t ImageHeap::Allocate() {
+	auto& scheduler = m_context.GetCommandScheduler();
+	while (!m_retired.empty() && scheduler.IsFree(m_retired.front().first)) {
+		m_free.push_back(m_retired.front().second);
+		m_retired.pop_front();
+	}
+	if (m_free.empty()) {
+		if (m_elements.size() == IR::ImageHeapCapacity) {
+			EXIT("image heaps exceed %u resident images\n", IR::ImageHeapCapacity);
+		}
+		m_elements.emplace_back();
+		return static_cast<uint32_t>(m_elements.size() - 1u);
+	}
+	const auto index = m_free.back();
+	m_free.pop_back();
+	m_elements[index] = {};
+	return index;
+}
+
+void ImageHeap::Retire(View& view) {
+	const auto tick = m_context.GetCommandScheduler().CurrentTick();
+	for (const auto index: view.released) m_retired.emplace_back(tick, index);
+	view.released.clear();
+}
+
+void ImageHeap::Rebind(RenderExecutor& executor) {
+	KYTY_PROFILER_FUNCTION();
+	auto&      textures = m_context.GetTextureCache();
+	const auto epoch    = Image::heap_epoch.load(std::memory_order_relaxed);
+	for (auto* view: m_active) {
+		bool       stale = false;
+		const auto mark  = [&](uint32_t index) { m_elements[index].stale = stale = true; };
+		if (view->resolve_epoch != epoch) {
+			view->resolve_epoch = epoch;
+			for (const auto& [value, index]: view->elements) {
+				const auto& binding = m_elements[index].binding;
+				if (!binding.image_id) continue;
+				const auto* image = textures.m_slot_images.try_get(binding.image_id);
+				if (image == nullptr || !image->registered || image->depth_id ||
+				    image->binding.needs_rebind) {
+					mark(index);
+				} else {
+					(void)textures.FindTexture(binding.image_id, binding.desc);
+				}
+			}
+		}
+		// Discovery applies the color metadata clears written since the last draw.
+		std::erase_if(view->dcc, [&](uint32_t index) {
+			const auto& element = m_elements[index];
+			if (element.refs == 0u || element.stale) return true;
+			if (executor.ResolveTexture(IR::ImageHeapResource(view->profile, element.kind),
+			                            element.value).image_id == element.binding.image_id) {
+				return false;
+			}
+			mark(index);
+			return true;
+		});
+		if (stale) {
+			// A changed image needs a new element; pending draws still read the old one. The new
+			// element keeps the old value, since a view takes table changes only when acquired.
+			std::erase_if(view->elements,
+			              [&](const auto& item) { return m_elements[item.second].stale; });
+			for (uint64_t probe = 0; probe < view->count; ++probe) {
+				const auto  slot    = (view->first + probe * view->step) / 4u;
+				const auto& element = m_elements[view->slots[slot] & ElementMask];
+				if (!element.stale) continue;
+				const auto value = element.value; // Intern may grow m_elements.
+				Assign(*view, slot, Intern(*view, element.kind, value));
+			}
+		}
+		for (const auto index: view->pending) {
+			if (m_elements[index].refs == 0u) continue;
+			auto binding = executor.ResolveTexture(
+			    IR::ImageHeapResource(view->profile, m_elements[index].kind),
+			    m_elements[index].value);
+			if (binding.image_id) {
+				auto& image = textures.GetImage(binding.image_id);
+				image.heap_refs.fetch_add(1u, std::memory_order_relaxed);
+				binding.layout     = image.ReadLayout();
+				binding.image_view = textures.FindTexture(binding.image_id, binding.desc);
+				if (binding.desc.info.metadata.kind == ImageMetadataKind::Dcc) {
+					view->dcc.push_back(index);
+				}
+				view->away.push_back(index);
+			}
+			m_elements[index].binding = std::move(binding);
+			m_writes.push_back(index);
+		}
+		view->pending.clear();
+	}
+	WriteDescriptors();
+}
+
+void ImageHeap::WriteDescriptors() {
+	if (m_writes.empty()) return;
+	std::vector<vk::DescriptorImageInfo> infos;
+	std::vector<vk::WriteDescriptorSet>  writes(m_writes.size());
+	infos.reserve(m_writes.size());
+	for (size_t write = 0; write < m_writes.size(); ++write) {
+		const auto& binding = m_elements[m_writes[write]].binding;
+		infos.push_back({nullptr, binding.image_view, binding.layout});
+		writes[write].dstSet          = m_set;
+		writes[write].dstArrayElement = m_writes[write];
+		writes[write].descriptorCount = 1;
+		writes[write].descriptorType  = vk::DescriptorType::eSampledImage;
+		writes[write].pImageInfo      = &infos.back();
+	}
+	m_context.GetGraphics().device.updateDescriptorSets(static_cast<uint32_t>(writes.size()),
+	                                                    writes.data(), 0, nullptr);
+	m_writes.clear();
+}
+
+void ImageHeap::Commit(CommandBuffer& buffer, vk::PipelineBindPoint bind_point,
+                       vk::PipelineLayout layout) {
+	KYTY_PROFILER_FUNCTION();
+	auto&      textures = m_context.GetTextureCache();
+	const auto epoch    = Image::heap_epoch.load(std::memory_order_relaxed);
+	for (auto* view: m_active) {
+		if (view->dirty_begin != view->dirty_end) {
+			const auto slots = std::span(view->slots).subspan(view->dirty_begin,
+			                                                  view->dirty_end - view->dirty_begin);
+			auto& stream = m_context.GetBufferCache().GetUtilityBuffer(MemoryUsage::Stream);
+			const auto offset = stream.Copy(slots.data(), slots.size_bytes(), sizeof(uint32_t));
+			view->buffer->CopyFrom(buffer, stream, offset, view->dirty_begin * sizeof(uint32_t),
+			                       slots.size_bytes());
+			view->dirty_begin = view->dirty_end = 0;
+		}
+		Retire(*view);
+		if (view->layout_epoch != epoch) {
+			view->layout_epoch = epoch;
+			view->away.clear();
+			for (const auto& [value, index]: view->elements) view->away.push_back(index);
+		}
+		// This draw's own bindings own the layouts of its attachments and storage images. An image
+		// this draw's discovery replaced lives until the tick ends, and its slots still select it.
+		std::erase_if(view->away, [&](uint32_t index) {
+			const auto& element = m_elements[index];
+			auto*       image   = textures.m_slot_images.try_get(element.binding.image_id);
+			if (element.refs == 0u || image == nullptr) return true;
+			if (image->binding.is_target || image->binding.shader_write ||
+			    image->binding.force_general) {
+				return false;
+			}
+			const auto& view_info = element.binding.desc.view_info;
+			image->Transit(image->ReadLayout(), vk::AccessFlagBits2::eShaderRead,
+			               ImageSubresourceRange {view_info.base_level, view_info.level_count,
+			                                      view_info.base_layer, view_info.layer_count},
+			               buffer.Handle());
+			return true;
+		});
+	}
+	m_active.clear();
+	buffer.Handle().bindDescriptorSets(bind_point, layout, IR::ImageHeapSet, 1, &m_set, 0,
+	                                   nullptr);
+}
+
+void ImageHeap::UnmapMemory(uint64_t address, uint64_t size) {
+	// Descriptors of textures in the range stay null even once it is mapped again.
+	if (!m_unmapped.empty() && m_unmapped.back().End() == address) {
+		m_unmapped.back().size += size;
+	} else {
+		m_unmapped.push_back({address, size});
+	}
+	auto&           scheduler = m_context.GetCommandScheduler();
+	std::scoped_lock lock {m_mutex};
+	std::erase_if(m_memories, [&](const std::unique_ptr<Memory>& memory) {
+		if (address >= memory->base + memory->size || memory->base >= address + size) {
+			return false;
+		}
+		Unwatch(*memory);
+		for (auto& view: memory->views) {
+			std::erase(m_active, view.get());
+			for (const auto& [value, index]: view->elements) {
+				Unreference(m_elements[index]);
+				view->released.push_back(index);
+			}
+			Retire(*view);
+			if (scheduler.Active()) {
+				scheduler.DeferOperation([buffer = std::move(view->buffer)] {});
+			}
+		}
+		return true;
+	});
+}
+
+} // namespace Libs::Graphics
@@ -0,0 +1,138 @@
+#ifndef EMULATOR_SRC_GRAPHICS_HOST_GPU_RENDERER_IMAGEHEAP_H_
+#define EMULATOR_SRC_GRAPHICS_HOST_GPU_RENDERER_IMAGEHEAP_H_
+
+#include "common/common.h"
+#include "graphics/host_gpu/renderer/pipeline/descriptors.h"
+#include "graphics/shader/recompiler/ir/passes/SrtWalker.h"
+
+#include <deque>
+#include <memory>
+#include <mutex>
+#include <unordered_map>
+#include <utility>
+#include <vector>
+
+namespace Libs::Graphics {
+
+class Buffer;
+class CommandBuffer;
+class PageManager;
+class RenderContext;
+class RenderExecutor;
+struct ImageHeapTestAccess;
+
+// Guest image tables that shaders index by GPU-computed offsets, kept resident in one bindless
+// sampled-image array. Each view of a table owns device slot words, one per table dword, that
+// select an array element. CPU writes to a table fault in its snapshot, which each draw refreshes
+// once; a view takes the changes when a draw first acquires it. Changes to the images behind
+// elements bump Image::heap_epoch.
+class ImageHeap final: public ShaderRecompiler::IR::ImageHeapSource {
+public:
+	ImageHeap(RenderContext& context, PageManager& page_manager);
+	~ImageHeap();
+	KYTY_CLASS_NO_COPY(ImageHeap);
+
+	// Starts a draw or dispatch, before its programs acquire views.
+	void BeginDraw() {
+		++m_draw;
+		Bury();
+	}
+	bool Acquire(const ShaderRecompiler::IR::ImageHeapRequest& request,
+	             ShaderRecompiler::IR::ImageHeapView&          result) override;
+	// Image discovery: resolves new elements and rediscovers changed images.
+	void Rebind(RenderExecutor& executor);
+	// Recording: uploads slots, moves resident images to their read layout and binds the array.
+	void Commit(CommandBuffer& buffer, vk::PipelineBindPoint bind_point, vk::PipelineLayout layout);
+	// CPU write fault, on any thread.
+	void InvalidateMemory(uint64_t address, uint64_t size);
+	void UnmapMemory(uint64_t address, uint64_t size);
+
+private:
+	using DescriptorValue = ShaderRecompiler::IR::DescriptorValue;
+
+	struct DescriptorHash {
+		size_t operator()(const DescriptorValue& value) const noexcept;
+	};
+
+	struct Element {
+		DescriptorValue value;
+		TextureBinding  binding;
+		uint32_t        kind  = 0;
+		uint32_t        refs  = 0; // Slots of the owning view that select it.
+		bool            stale = false;
+	};
+
+	struct Memory;
+
+	struct View {
+		Memory&                             memory;
+		uint64_t                            step  = 0;
+		uint64_t                            first = 0;
+		uint64_t                            count = 0;
+		ShaderRecompiler::IR::ImageResource profile;
+		std::vector<uint32_t>               slots; // Plus a trailing null slot.
+		std::vector<uint32_t>               kind_keys;
+		std::vector<DescriptorValue>        kinds; // One representative per kind_keys entry.
+		std::unordered_map<DescriptorValue, uint32_t, DescriptorHash> elements;
+		std::vector<uint32_t>   pending;  // Not yet resolved.
+		std::vector<uint32_t>   released; // No slot selects them after the next upload.
+		std::vector<uint32_t>   away;     // Not yet moved to the read layout.
+		std::vector<uint32_t>   dcc;      // Color metadata clears apply on each discovery.
+		std::unique_ptr<Buffer> buffer;
+		size_t                  dirty_begin   = 0; // Slots to upload.
+		size_t                  dirty_end     = 0;
+		size_t                  changed_begin = 0; // Table bytes changed since it was probed.
+		size_t                  changed_end   = 0;
+		uint64_t                resolve_epoch = 0;
+		uint64_t                layout_epoch  = 0;
+	};
+
+	struct Memory {
+		uint64_t                           base = 0;
+		uint64_t                           size = 0;
+		std::vector<uint32_t>              words;
+		std::vector<uint8_t>               dirty_pages; // Unwatched, so unread since a write.
+		std::vector<uint64_t>              dead;        // Per word, a buried descriptor's hash.
+		uint64_t                           draw = 0;    // Last draw that refreshed it.
+		std::vector<std::unique_ptr<View>> views;
+	};
+
+	[[nodiscard]] bool Refresh(Memory& memory);
+	void               MarkDirty(Memory& memory, size_t begin, size_t end);
+	template <bool track>
+	void     Watch(const Memory& memory, size_t begin, size_t end);
+	void     Unwatch(const Memory& memory);
+	void     UpdateProbes(View& view, uint64_t begin, uint64_t end);
+	void     UpdateProbe(View& view, uint64_t probe);
+	void     Assign(View& view, size_t slot, uint32_t entry);
+	void     Bury();
+	uint32_t Intern(View& view, uint32_t kind, const DescriptorValue& value);
+	void     Release(View& view, uint32_t entry);
+	void     Unreference(Element& element);
+	uint32_t Allocate();
+	void     Retire(View& view);
+	void     WriteDescriptors();
+
+	RenderContext&                       m_context;
+	PageManager&                         m_page_manager;
+	std::mutex                           m_mutex; // Fault-visible Memory state and m_memories.
+	std::vector<std::unique_ptr<Memory>> m_memories;
+	std::vector<View*>                   m_active; // Acquired since the last commit.
+	uint64_t                             m_draw = 1;
+	std::vector<Element>                 m_elements;
+	std::vector<uint32_t>                m_free;
+	std::deque<std::pair<uint64_t, uint32_t>> m_retired; // Tick, element.
+	std::vector<std::pair<size_t, size_t>>    m_runs;
+	std::vector<uint32_t>                     m_scratch;
+	std::vector<uint32_t>                     m_writes; // Elements whose descriptors changed.
+	std::vector<GuestRange>                   m_unmapped; // Since the last Bury.
+	vk::DescriptorSetLayout                   m_layout = nullptr;
+	vk::DescriptorPool                        m_pool   = nullptr;
+	vk::DescriptorSet                         m_set    = nullptr;
+
+	friend struct ImageHeapTestAccess;
+};
+
+} // namespace Libs::Graphics
+
+#endif // EMULATOR_SRC_GRAPHICS_HOST_GPU_RENDERER_IMAGEHEAP_H_
@@ -1,6 +1,7 @@
 #include "graphics/host_gpu/renderer/pipeline/pipelineCache.h"

 #include "common/assert.h"
+#include "common/debugCounters.h"
 #include "common/emulatorConfig.h"
 #include "common/file.h"
 #include "common/logging/log.h"
@@ -18,13 +19,13 @@
 #include "graphics/shader/recompiler/ShaderRecompiler.h"
 #include "graphics/shader/shaderCompiler.h"
 #include "kernel/memory.h"
-#include "kytyGitVersion.h"
 #include "loader/systemContent.h"

 #include <algorithm>
 #include <array>
 #include <atomic>
 #include <cctype>
+#include <chrono>
 #include <cstdio>
 #include <cstring>
 #include <fmt/format.h>
@@ -81,7 +82,8 @@ std::string DriverCacheSignature(const vk::PhysicalDeviceProperties& properties)
 		uuid[i * 2]     = hex[properties.pipelineCacheUUID[i] >> 4u];
 		uuid[i * 2 + 1] = hex[properties.pipelineCacheUUID[i] & 0xfu];
 	}
-	return fmt::format("KytyPC1:{}:{:08x}:{:08x}:{:08x}:{}\n", KYTY_GIT_REVISION,
+	// Driver cache entries already distinguish shader code and pipeline state across builds.
+	return fmt::format("KytyPC2:{:08x}:{:08x}:{:08x}:{}\n",
 	                   properties.vendorID, properties.deviceID, properties.driverVersion, uuid);
 }

@@ -273,6 +275,7 @@ struct PipelineCache::ProgramCache {
 		ShaderRecompiler::IR::ResourceSnapshot       resources;
 		ShaderRecompiler::IR::ResourceSpecialization specialization;
 		std::vector<Permutation>                    permutations;
+		bool                                       skip_dispatch = false;
 	};

 	struct ProgramKeyHash {
@@ -299,10 +302,13 @@ struct PipelineCache::ProgramCache {
 	                               ShaderRecompiler::IR::ResourceSpecialization specialization,
 	                               uint32_t push_data_start_dword,
 	                               std::span<const uint32_t> function_code) {
+		Common::DebugCounters::ScopedTimer timer(Common::DebugCounters::Counter::ShadersCompiled,
+		                                         Common::DebugCounters::Counter::ShaderCompileNs,
+		                                         Common::DebugCounters::Peak::ShaderCompileNs);
 		auto result = ShaderRecompiler::CompileProgram(std::move(translated), options,
 		                                               specialization, push_data_start_dword);
 		if (!ValidateShaderSpirv(options.dump_label, options.shader_hash, result.spirv,
-		                         options.maintenance8_enabled)) {
+		                         options.profile->maintenance8)) {
 			DumpShaderSpirv(stage_name, options.shader_hash, result.spirv);
 			EXIT("%s failed hash=0x%016" PRIx64 ": SPIR-V validation failed\n", options.dump_label,
 			     options.shader_hash);
@@ -347,6 +353,7 @@ struct PipelineCache::ProgramCache {
 		    .user_data                  = user_data,
 		    .shader_base                = params.Base(),
 		    .read_specialization_memory = ReadShaderGuestMemory,
+		    .image_heap                 = image_heap,
 		};
 		if constexpr (std::is_same_v<InputInfo, ShaderComputeInputInfo>) {
 			runtime.workgroup_counts = input_info.workgroup_counts;
@@ -355,6 +362,7 @@ struct PipelineCache::ProgramCache {
 		bool source_changed = false;
 		if (entry != programs.end()) {
 			auto& cached = entry->second;
+			if (cached.skip_dispatch) return {};
 			Permutation* permutation = nullptr;
 			if (cached.call_source) {
 				source_changed = cached.RefreshCallSource(runtime);
@@ -391,7 +399,7 @@ struct PipelineCache::ProgramCache {
 		}
 		ShaderRecompiler::CompileOptions options;
 		options.stage       = stage;
-		options.maintenance8_enabled = graphics.maintenance8_enabled;
+		options.profile     = &profile;
 		options.shader_hash = params.hash;
 		options.user_data   = user_data;
 		options.back_code      = params.back_code;
@@ -414,10 +422,18 @@ struct PipelineCache::ProgramCache {
 		DumpShaderOriginal(stage_name, options.shader_hash, params.code);
 		ShaderRecompiler::ShaderSource source;
 		if (entry == programs.end()) source = ShaderRecompiler::PrepareShaderSource(params.code, options);
+		if (source.unsupported_call_pc != UINT32_MAX) {
+			PipelineCacheLog("!!! TEMPORARY COMPUTE DISPATCH SKIP !!! shader=0x{:016x} pc=0x{:08x} "
+			                 "reason=unsupported GPU-selected callable target; rendering is incomplete",
+			                 params.hash, source.unsupported_call_pc);
+			programs.try_emplace(lookup_key, ShaderRecompiler::IR::ResourcePlan{}, nullptr)
+			    .first->second.skip_dispatch = true;
+			return {};
+		}
 		auto* current = entry == programs.end() ? &source : entry->second.call_source.get();
 		ShaderRecompiler::TranslateResult translated;
 		if (current != nullptr) {
-			if (current->call) {
+			if (!current->calls.empty()) {
 				if (entry == programs.end()) ShaderRecompiler::RefreshShaderSource(*current, runtime);
 				translated = ShaderRecompiler::TranslateProgram(current->linked->decoded, options);
 				translated.program.source_reads = current->reads;
@@ -428,7 +444,7 @@ struct PipelineCache::ProgramCache {
 			translated = ShaderRecompiler::TranslateProgram(params.code, options);
 		}
 		if (entry == programs.end()) {
-			auto retained = source.call ? std::make_unique<ShaderRecompiler::ShaderSource>(std::move(source)) : nullptr;
+			auto retained = !source.calls.empty() ? std::make_unique<ShaderRecompiler::ShaderSource>(std::move(source)) : nullptr;
 			entry = programs.try_emplace(lookup_key,
 			    ShaderRecompiler::IR::ExtractResourcePlan(translated.program), std::move(retained)).first;
 			EXIT_IF(!ShaderRecompiler::IR::MaterializeResources(
@@ -465,7 +481,13 @@ struct PipelineCache::ProgramCache {
 		return permutation.handle;
 	}

-	explicit ProgramCache(GraphicContext& graphics): graphics(graphics) {
+	ProgramCache(GraphicContext& graphics, ShaderRecompiler::IR::ImageHeapSource* image_heap)
+	    : graphics(graphics), image_heap(image_heap),
+	      // AMD GPUs (vendor 0x1002) run NMin/NMax as v_min/v_max, which give the guest's NaN and
+	      // signed-zero results; GLSL.std.450 alone does not order -0 below +0.
+	      profile {.maintenance8 = graphics.maintenance8_enabled,
+		           .native_float_min_max =
+		               graphics.physical_device_properties.vendorID == 0x1002u} {
 		lookup_key.static_state.reserve(MaxStaticKeyWords);
 	}
 	~ProgramCache() {
@@ -480,11 +502,14 @@ struct PipelineCache::ProgramCache {
 	std::unordered_map<ProgramKey, SourceEntry, ProgramKeyHash> programs;
 	ProgramKey                                                  lookup_key;
 	GraphicContext&                                             graphics;
+	ShaderRecompiler::IR::ImageHeapSource*                      image_heap;
+	const ShaderRecompiler::Spirv::Profile                      profile;
 	uint64_t                                                    next_shader_id = 0;
 };

-PipelineCache::PipelineCache(GraphicContext& graphics)
-    : m_graphics(graphics), m_program_cache(std::make_unique<ProgramCache>(graphics)),
+PipelineCache::PipelineCache(GraphicContext&                         graphics,
+                             ShaderRecompiler::IR::ImageHeapSource* image_heap)
+    : m_graphics(graphics), m_program_cache(std::make_unique<ProgramCache>(graphics, image_heap)),
      m_next_driver_cache_save(std::chrono::steady_clock::now() + DriverCacheSaveInterval) {
 	EXIT_NOT_IMPLEMENTED(!Common::Thread::IsMainThread());
 	InitializeDriverCache();
@@ -516,17 +541,6 @@ void PipelineCache::InitializeDriverCache() {
 		PipelineCacheLog("Vulkan pipeline cache: disabled (non-Release build)");
 		return;
 	}
-	const std::string_view git_hash     = KYTY_GIT_HASH;
-	const std::string_view git_revision = KYTY_GIT_REVISION;
-	if (git_hash == "unknown" || git_revision == "unknown") {
-		PipelineCacheLog("Vulkan pipeline cache: disabled (unknown git revision)");
-		return;
-	}
-	if (git_hash.ends_with("-dirty")) {
-		PipelineCacheLog("Vulkan pipeline cache: disabled (dirty build)");
-		return;
-	}
-
 	m_driver_cache_path     = std::filesystem::path("_PipelineCache") / (title_id + ".bin");
 	const auto path         = Common::PathToString(m_driver_cache_path);
 	const bool cache_exists = Common::File::IsFileExisting(m_driver_cache_path);
@@ -559,7 +573,7 @@ void PipelineCache::InitializeDriverCache() {
 			    XXH3_64bits(initial_data.data(), initial_data.size()) != payload_hash) {
 				initial_data.clear();
 				PipelineCacheLog(
-				    "Vulkan pipeline cache: invalidating {} (driver, emulator, or data mismatch)",
+				    "Vulkan pipeline cache: invalidating {} (driver or data mismatch)",
 				    path);
 			}
 		} else {
@@ -586,8 +600,8 @@ void PipelineCache::InitializeDriverCache() {
 		return;
 	}
 	if (!initial_data.empty()) {
-		PipelineCacheLog("Vulkan pipeline cache: loaded {} bytes from {}", initial_data.size(),
-		                 path);
+		PipelineCacheLog("Vulkan pipeline cache: loaded {:.2f} MB from {}",
+		                 static_cast<double>(initial_data.size()) / 1'000'000.0, path);
 	} else {
 		PipelineCacheLog("Vulkan pipeline cache: initialized empty");
 	}
@@ -627,8 +641,8 @@ void PipelineCache::Save() {
 	}
 	if (result != vk::Result::eSuccess || size == 0 ||
 	    size > std::numeric_limits<uint32_t>::max()) {
-		PipelineCacheLog("Vulkan pipeline cache: save failed ({}, {} bytes)",
-		                 vk::to_string(result), size);
+		PipelineCacheLog("Vulkan pipeline cache: save failed ({}, {:.2f} MB)",
+		                 vk::to_string(result), static_cast<double>(size) / 1'000'000.0);
 		return;
 	}
 	payload.resize(size);
@@ -660,7 +674,8 @@ void PipelineCache::Save() {
 		                 Common::PathToString(m_driver_cache_path));
 		return;
 	}
-	PipelineCacheLog("Vulkan pipeline cache: saved {} bytes to {}", payload.size(),
+	PipelineCacheLog("Vulkan pipeline cache: saved {:.2f} MB to {}",
+	                 static_cast<double>(payload.size()) / 1'000'000.0,
 	                 Common::PathToString(m_driver_cache_path));
 	m_driver_cache_dirty = false;
 }
@@ -775,10 +790,12 @@ PipelineCache::GraphicsPrograms PipelineCache::GetGraphicsPrograms(
 			}
 		}
 	}
-	if (context.GetClipControl().clip_disable) {
+	auto&       last_stage   = vertex_info[tess_active ? 2u : 0u];
+	const auto& clip_control = context.GetClipControl();
+	if (clip_control.clip_disable) {
 		const auto& viewport = context.GetScreenViewport().viewports[0];
 		const auto& limits   = m_graphics.GetPhysicalDeviceProperties().limits;
-		auto&       clip     = vertex_info[tess_active ? 2u : 0u].clip_space;
+		auto&       clip     = last_stage.clip_space;
 		clip.scale[0]        = viewport.xscale;
 		clip.scale[1]        = viewport.yscale;
 		clip.offset[0]       = viewport.xoffset;
@@ -788,9 +805,16 @@ PipelineCache::GraphicsPrograms PipelineCache::GetGraphicsPrograms(
 		clip.half_extent[1] =
 		    static_cast<float>(std::min(limits.maxViewportDimensions[1], 16384u)) * 0.5f;
 		clip.enabled = true;
+	} else if (clip_control.user_clip_planes != 0) {
+		// CLIP_DISABLE also disables user planes. How they combine with shader clip and cull
+		// distances is unverified; draw push constants only reach VS/mesh and PS, and the
+		// generated RECTLIST stages do not forward clip distances.
+		EXIT_NOT_IMPLEMENTED((last_stage.pa_cl_vs_out_cntl & 0xffffu) != 0 || tess_active ||
+		                     Prospero::IsRectList(user_config.GetPrimType()));
+		last_stage.user_clip_planes = clip_control.user_clip_planes;
 	}
 	uint32_t push_data_cursor = ShaderRecompiler::IR::PushData::GraphicsDrawDwords(
-	    mesh_active, object_primitive_id);
+	    mesh_active, object_primitive_id, last_stage.user_clip_planes);
 	GraphicsPrograms  result;
 	if (pixel_active) {
 		result.pixel = m_program_cache->Get(pixel_params, pixel_info, push_data_cursor);
@@ -957,6 +981,7 @@ PipelineCache::Pipeline& PipelineCache::GetGraphicsPipeline(
 	if (static_params.sample_shading_enable && !m_graphics.sample_rate_shading_enabled) {
 		EXIT("Pipeline: sample-rate shading is required but unsupported by the host\n");
 	}
+	static_params.alpha_to_coverage_enable = ps_active && ctx.GetShaderRegisters().AlphaToMask();
 	const bool rect_list = Prospero::IsRectList(command.GetUserConfig().GetPrimType());
 	static_params.cull_back  = !rect_list && mc.cull_back;
 	static_params.cull_front = !rect_list && mc.cull_front;
@@ -1004,8 +1029,13 @@ PipelineCache::Pipeline& PipelineCache::GetGraphicsPipeline(

 	auto cached = std::make_unique<Pipeline>();
 	LogPipelineTrace("CreatePipelineInternal begin", vs_id, ps_id);
-	CreatePipelineInternal(m_graphics, *cached, rendering, key.vertex_input, vertex_info,
-	                       ps_input_info, programs, static_params, m_driver_cache);
+	{
+		Common::DebugCounters::ScopedTimer timer(Common::DebugCounters::Counter::PipelinesCreated,
+		                                         Common::DebugCounters::Counter::PipelineCreateNs,
+		                                         Common::DebugCounters::Peak::PipelineCreateNs);
+		CreatePipelineInternal(m_graphics, *cached, rendering, key.vertex_input, vertex_info,
+		                       ps_input_info, programs, static_params, m_driver_cache);
+	}
 	LogPipelineTrace("CreatePipelineInternal done", vs_id, ps_id);

 	EXIT_NOT_IMPLEMENTED(cached->pipeline == nullptr);
@@ -1035,7 +1065,13 @@ PipelineCache::GetComputePipeline(const ShaderComputeInputInfo& input_info,
 	}

 	auto cached = std::make_unique<Pipeline>();
-	CreatePipelineInternal(m_graphics, *cached, input_info, compute_program.module, m_driver_cache);
+	{
+		Common::DebugCounters::ScopedTimer timer(Common::DebugCounters::Counter::PipelinesCreated,
+		                                         Common::DebugCounters::Counter::PipelineCreateNs,
+		                                         Common::DebugCounters::Peak::PipelineCreateNs);
+		CreatePipelineInternal(m_graphics, *cached, input_info, compute_program.module,
+		                       m_driver_cache);
+	}

 	EXIT_NOT_IMPLEMENTED(cached->pipeline == nullptr);
 	EXIT_NOT_IMPLEMENTED(cached->pipeline_layout == nullptr);
@@ -30,6 +30,10 @@ class UserConfig;
 struct ComputeShaderInfo;
 } // namespace HW

+namespace ShaderRecompiler::IR {
+class ImageHeapSource;
+} // namespace ShaderRecompiler::IR
+
 #pragma pack(push, 1)

 struct PipelineStaticParameters {
@@ -39,6 +43,7 @@ struct PipelineStaticParameters {
 	bool                       primitive_restart_enable = false;
 	uint32_t                   samples                  = 1;
 	bool                       sample_shading_enable    = false;
+	bool                       alpha_to_coverage_enable = false;
 	uint32_t                   color_mask[RENDER_COLOR_ATTACHMENTS_MAX]           = {};
 	bool                       cull_front                                         = false;
 	bool                       cull_back                                          = false;
@@ -62,7 +67,7 @@ struct PipelineStaticParameters {
 static_assert(std::is_trivially_copyable_v<PipelineStaticParameters>);
 static_assert(std::is_standard_layout_v<PipelineStaticParameters>);
 static_assert(alignof(PipelineStaticParameters) == 1);
-static_assert(sizeof(PipelineStaticParameters) == 116);
+static_assert(sizeof(PipelineStaticParameters) == 117);

 struct PipelineRenderingState {
 	std::array<vk::Format, RENDER_COLOR_ATTACHMENTS_MAX> color_formats {};
@@ -103,7 +108,7 @@ struct ShaderProgram {
 // The owning renderer serializes access, including saves while the GPU is running.
 class PipelineCache {
 public:
-	explicit PipelineCache(GraphicContext& graphics);
+	PipelineCache(GraphicContext& graphics, ShaderRecompiler::IR::ImageHeapSource* image_heap);
 	~PipelineCache();
 	KYTY_CLASS_NO_COPY(PipelineCache);
 	void Save();
@@ -113,6 +118,7 @@ public:
 		vk::Pipeline            pipeline              = nullptr;
 		vk::DescriptorSetLayout descriptor_set_layout = nullptr;
 		bool                    uses_push_descriptors = false;
+		bool                    uses_image_heap       = false;
 	};

 	struct GraphicsPrograms {
@@ -52,6 +52,7 @@ vk::MemoryBarrier MakeShaderWriteDependency() {
 	barrier.dstAccessMask = vk::AccessFlagBits::eShaderRead | vk::AccessFlagBits::eShaderWrite |
 	                        vk::AccessFlagBits::eVertexAttributeRead |
 	                        vk::AccessFlagBits::eIndexRead | vk::AccessFlagBits::eUniformRead |
+	                        vk::AccessFlagBits::eIndirectCommandRead |
 	                        vk::AccessFlagBits::eTransferRead | vk::AccessFlagBits::eTransferWrite |
 	                        vk::AccessFlagBits::eColorAttachmentRead |
 	                        vk::AccessFlagBits::eColorAttachmentWrite;
@@ -92,7 +93,8 @@ bool HasShaderBufferWrites(const ShaderStageRuntime& runtime) {
 	const auto& program   = *runtime.program;
 	const auto& resources = *runtime.resources;
 	EXIT_IF(resources.buffers.size() != program.info.buffers.size());
-	bool has_writes = false;
+	// Address writes reach memory outside the bound buffer descriptors.
+	bool has_writes = program.has_address_writes;
 	for (uint32_t i = 0; i < program.info.buffers.size(); i++) {
 		if (!program.info.buffers[i].written) {
 			continue;
@@ -196,21 +196,27 @@ static void CreateDescriptorLayout(GraphicContext& graphics, PipelineCache::Pipe
 	const auto& limits = graphics.GetPhysicalDeviceProperties().limits;
 	uint64_t descriptor_count = 0;
 	uint64_t storage_buffers = 0;
-	std::array<uint64_t, 32> stage_buffers {}, stage_resources {};
+	uint64_t samplers = 0;
+	std::array<uint64_t, 32> stage_buffers {}, stage_resources {}, stage_samplers {};
 	for (const auto& binding: bindings) {
 		descriptor_count += binding.descriptorCount;
 		const bool storage = binding.descriptorType == vk::DescriptorType::eStorageBuffer;
+		const bool sampler = binding.descriptorType == vk::DescriptorType::eSampler;
 		if (storage) storage_buffers += binding.descriptorCount;
+		if (sampler) samplers += binding.descriptorCount;
 		for (auto stages = static_cast<uint32_t>(binding.stageFlags); stages != 0; stages &= stages - 1u) {
 			const auto stage = std::countr_zero(stages);
 			if (storage) stage_buffers[stage] += binding.descriptorCount;
-			if (binding.descriptorType != vk::DescriptorType::eSampler)
+			if (sampler) stage_samplers[stage] += binding.descriptorCount;
+			else
 				stage_resources[stage] += binding.descriptorCount;
 			if (stage_buffers[stage] > limits.maxPerStageDescriptorStorageBuffers ||
-			    stage_resources[stage] > limits.maxPerStageResources) {
+			    stage_resources[stage] > limits.maxPerStageResources ||
+			    stage_samplers[stage] > limits.maxPerStageDescriptorSamplers) {
 				EXIT("native descriptor stage %u exceeds device limits: buffers=%" PRIu64
-				     "/%u resources=%" PRIu64 "/%u\n", stage, stage_buffers[stage],
-				     limits.maxPerStageDescriptorStorageBuffers, stage_resources[stage], limits.maxPerStageResources);
+				     "/%u resources=%" PRIu64 "/%u samplers=%" PRIu64 "/%u\n", stage, stage_buffers[stage],
+				     limits.maxPerStageDescriptorStorageBuffers, stage_resources[stage], limits.maxPerStageResources,
+				     stage_samplers[stage], limits.maxPerStageDescriptorSamplers);
 			}
 		}
 	}
@@ -218,6 +224,10 @@ static void CreateDescriptorLayout(GraphicContext& graphics, PipelineCache::Pipe
 		EXIT("native storage-buffer descriptors exceed device limit: %" PRIu64 "/%u\n",
 		     storage_buffers, limits.maxDescriptorSetStorageBuffers);
 	}
+	if (samplers > limits.maxDescriptorSetSamplers) {
+		EXIT("native sampler descriptors exceed device limit: %" PRIu64 "/%u\n",
+		     samplers, limits.maxDescriptorSetSamplers);
+	}
 	pipeline.uses_push_descriptors = descriptor_count <= graphics.max_push_descriptors;

 	vk::DescriptorSetLayoutCreateInfo create {};
@@ -412,6 +422,8 @@ void CreatePipelineInternal(GraphicContext& graphics, PipelineCache::Pipeline& p
 	multisampling.sampleShadingEnable  = static_params.sample_shading_enable ? VK_TRUE : VK_FALSE;
 	multisampling.rasterizationSamples = vulkan_sample_count(static_params.samples);
 	multisampling.minSampleShading     = 1.0f;
+	multisampling.alphaToCoverageEnable =
+	    static_params.alpha_to_coverage_enable ? VK_TRUE : VK_FALSE;

 	vk::PipelineColorBlendAttachmentState color_blend_attachment[RENDER_COLOR_ATTACHMENTS_MAX] = {};
 	for (uint32_t i = 0; i < rendering.color_count; i++) {
@@ -461,20 +473,23 @@ void CreatePipelineInternal(GraphicContext& graphics, PipelineCache::Pipeline& p
 	for (const auto& stage: vertex_info) {
 		const auto native_stage = NativeShaderStage(stage.logical_stage);
 		AddLayoutBindings(descriptor_bindings, *stage.stage.program, native_stage);
+		pipeline.uses_image_heap |= stage.stage.program->info.uses_image_heap;
 		graphics_stages |= native_stage;
 	}
 	if (ps_active) {
 		EXIT_IF(!ps_input_info->stage);
 		AddLayoutBindings(descriptor_bindings, *ps_input_info->stage.program,
 		                  vk::ShaderStageFlagBits::eFragment);
+		pipeline.uses_image_heap |= ps_input_info->stage.program->info.uses_image_heap;
 	}
 	CreateDescriptorLayout(graphics, pipeline, descriptor_bindings);
 	const vk::PushConstantRange push_constants {graphics_stages, 0,
 	                                            ShaderRecompiler::IR::NativePushConstantSize};
+	const std::array set_layouts {pipeline.descriptor_set_layout, graphics.image_heap_layout};

 	vk::PipelineLayoutCreateInfo pipeline_layout_info {};
-	pipeline_layout_info.setLayoutCount         = 1;
-	pipeline_layout_info.pSetLayouts            = &pipeline.descriptor_set_layout;
+	pipeline_layout_info.setLayoutCount         = pipeline.uses_image_heap ? 2u : 1u;
+	pipeline_layout_info.pSetLayouts            = set_layouts.data();
 	pipeline_layout_info.pushConstantRangeCount = 1;
 	pipeline_layout_info.pPushConstantRanges    = &push_constants;

@@ -603,13 +618,15 @@ void CreatePipelineInternal(GraphicContext& graphics, PipelineCache::Pipeline& p
 	std::vector<vk::DescriptorSetLayoutBinding> descriptor_bindings;
 	AddLayoutBindings(descriptor_bindings, *input_info.stage.program,
 	                  vk::ShaderStageFlagBits::eCompute);
+	pipeline.uses_image_heap = input_info.stage.program->info.uses_image_heap;
 	CreateDescriptorLayout(graphics, pipeline, descriptor_bindings);
 	const vk::PushConstantRange push_constants {vk::ShaderStageFlagBits::eCompute, 0,
 	                                            ShaderRecompiler::IR::NativePushConstantSize};
+	const std::array set_layouts {pipeline.descriptor_set_layout, graphics.image_heap_layout};

 	vk::PipelineLayoutCreateInfo pipeline_layout_info {};
-	pipeline_layout_info.setLayoutCount         = 1;
-	pipeline_layout_info.pSetLayouts            = &pipeline.descriptor_set_layout;
+	pipeline_layout_info.setLayoutCount         = pipeline.uses_image_heap ? 2u : 1u;
+	pipeline_layout_info.pSetLayouts            = set_layouts.data();
 	pipeline_layout_info.pushConstantRangeCount = 1;
 	pipeline_layout_info.pPushConstantRanges    = &push_constants;

@@ -626,6 +643,9 @@ void CreatePipelineInternal(GraphicContext& graphics, PipelineCache::Pipeline& p
 	EXIT_NOT_IMPLEMENTED(pipeline.pipeline_layout == nullptr);

 	vk::ComputePipelineCreateInfo info {};
+	if (input_info.stage.program->info.uses_gds) {
+		info.flags = vk::PipelineCreateFlagBits::eDispatchBase;
+	}
 	info.stage             = comp_shader_stage_info;
 	info.layout            = pipeline.pipeline_layout;
 	info.basePipelineIndex = -1;
@@ -48,6 +48,13 @@ enum class DrawOffsetSource : uint8_t {
 	IndirectArgs,
 };

+// Nonzero address: the GPU reads count Vulkan indirect draw records, stride bytes apart.
+struct DrawIndirectArgs {
+	uint64_t address = 0;
+	uint32_t count   = 0;
+	uint32_t stride  = 0;
+};
+
 struct DrawIndexArgs {
 	uint32_t         index_count                = 0;
 	const void*      index_addr                 = nullptr;
@@ -57,6 +64,8 @@ struct DrawIndexArgs {
 	uint32_t         first_instance             = 0;
 	DrawOffsetSource offset_source              = DrawOffsetSource::DrawState;
 	uint32_t         render_target_slice_offset = 0;
+	// With indirect records, index_count is INDEX_BUFFER_SIZE.
+	DrawIndirectArgs indirect;
 };

 struct DrawAutoArgs {
@@ -66,6 +75,7 @@ struct DrawAutoArgs {
 	uint32_t         first_instance             = 0;
 	DrawOffsetSource offset_source              = DrawOffsetSource::DrawState;
 	uint32_t         render_target_slice_offset = 0;
+	DrawIndirectArgs indirect;
 };

 struct SubmitInfo {
@@ -161,18 +171,19 @@ public:
 	void CommitBindings(CommandBuffer& buffer, vk::PipelineBindPoint pipeline_bind_point,
 	                    const PipelineCache::Pipeline&     pipeline,
 	                    std::span<PreparedBindings* const> bindings);
+	[[nodiscard]] TextureBinding ResolveTexture(const ShaderRecompiler::IR::ImageResource&   resource,
+	                                            const ShaderRecompiler::IR::DescriptorValue& value);

 private:
-	void DrawIndex(uint64_t submit_id, CommandBuffer& buffer, const DrawIndexArgs& args);
-	void DrawAuto(uint64_t submit_id, CommandBuffer& buffer, const DrawAutoArgs& args);
+	// False, with nothing drawn, when the draw's indirect records need CPU counts.
+	bool DrawIndex(uint64_t submit_id, CommandBuffer& buffer, const DrawIndexArgs& args);
+	bool DrawAuto(uint64_t submit_id, CommandBuffer& buffer, const DrawAutoArgs& args);

 	struct GraphicsBindings {
 		std::array<PreparedBindings, 3> vertex;
 		std::optional<PreparedBindings> pixel;
 	};

-	[[nodiscard]] TextureBinding ResolveTexture(const ShaderRecompiler::IR::ImageResource& resource,
-	                                            const ShaderRecompiler::IR::DescriptorValue& value);
 	void PrepareGraphicsBindings(std::span<PreparedBindings* const> stages,
 	                             std::span<RenderColorInfo> colors);
 	void ResolveRenderColorTarget(CommandBuffer& buffer, RenderColorInfo& target,
@@ -201,11 +212,11 @@ private:
 	void                      BindImage(ImageId id, bool storage);
 	void                      BindRenderTarget(ImageId id);
 	void                      ResetBindings();
-	void                     UpdateComputeMetaClears(const ShaderComputeInputInfo& input,
-	                                                 const CommandBuffer&          buffer);
-	[[nodiscard]] bool TryConsumeComputeImageClear(const ShaderComputeInputInfo& input,
-	                                              CommandBuffer& command, uint32_t group_x,
-	                                              uint32_t group_y, uint32_t group_z, uint32_t mode);
+	[[nodiscard]] bool        TryConsumeComputeMetaClear(const ShaderComputeInputInfo& input,
+	                                                     const CommandBuffer&          buffer);
+	[[nodiscard]] bool TryConsumeComputeFill(const ShaderComputeInputInfo& input,
+	                                        CommandBuffer& command, uint32_t group_x,
+	                                        uint32_t group_y, uint32_t group_z, uint32_t mode);

 	RenderContext&                        m_context;
 	GraphicsBindings                     m_graphics_bindings;
@@ -1,5 +1,6 @@
 #include "common/assert.h"
 #include "common/common.h"
+#include "common/debugCounters.h"
 #include "common/emulatorConfig.h"
 #include "common/file.h"
 #include "common/logging/log.h"
@@ -50,8 +51,8 @@ static bool FillSourcesDisjoint(std::span<const ShaderRecompiler::IR::Descriptor
 	return true;
 }

-void RenderExecutor::UpdateComputeMetaClears(const ShaderComputeInputInfo& input,
-                                            const CommandBuffer&          buffer) {
+bool RenderExecutor::TryConsumeComputeMetaClear(const ShaderComputeInputInfo& input,
+                                                const CommandBuffer&          buffer) {
 	const auto& program   = *input.stage.program;
 	const auto& resources = *input.stage.resources;
 	if (resources.buffers.size() != program.info.buffers.size()) {
@@ -61,9 +62,10 @@ void RenderExecutor::UpdateComputeMetaClears(const ShaderComputeInputInfo& input
 	for (uint32_t i = 0; i < program.info.buffers.size(); i++) {
 		const auto& resource   = program.info.buffers[i];
 		const auto  descriptor = DecodeNativeDescriptor<ShaderBufferResource>(resources.buffers[i]);
-		// Reads and transforms of metadata do not establish a clear state.
+		// A metadata resource that is also read is not proven to be a full overwrite. Execute it
+		// conservatively instead of replacing the dispatch with a coarse full-surface clear.
 		if ((!resource.written || resource.read) && cache.IsMeta(descriptor.Base48())) {
-			return;
+			return false;
 		}
 	}

@@ -73,10 +75,13 @@ void RenderExecutor::UpdateComputeMetaClears(const ShaderComputeInputInfo& input
 			if (resource.written) {
 				const auto descriptor =
 				    DecodeNativeDescriptor<ShaderBufferResource>(resources.buffers[i]);
-				(void)cache.ClearMeta(descriptor.Base48());
+				if (cache.ClearMeta(descriptor.Base48())) {
+					return true;
+				}
 			}
 		}
 	}
+	return false;
 }

 bool ResolveComputeBufferFill(const ShaderComputeInputInfo& input, uint32_t group_x,
@@ -94,9 +99,10 @@ bool ResolveComputeBufferFill(const ShaderComputeInputInfo& input, uint32_t grou
 	constexpr std::array formats {
 	    Prospero::BufferFormat::k32UInt, Prospero::BufferFormat::k32_32UInt,
 	    Prospero::BufferFormat::k32_32_32UInt, Prospero::BufferFormat::k32_32_32_32UInt};
+	// Offset-or-swizzle bounds do not describe the element range proved by the shader.
 	if (descriptor.Stride() != element_size || descriptor.Format() != formats[fill.words - 1] ||
 	    descriptor.SwizzleEnabled() || descriptor.IndexStride() != 0 || descriptor.AddTid() ||
-	    descriptor.Base48() == 0) {
+	    descriptor.OutOfBounds() == 3 || descriptor.Base48() == 0) {
 		return false;
 	}
 	if (input.threads_num[0] == 0 || input.threads_num[0] != fill.group_stride[0] ||
@@ -108,7 +114,11 @@ bool ResolveComputeBufferFill(const ShaderComputeInputInfo& input, uint32_t grou
 	                                 ? group_x
 	                                 : static_cast<uint64_t>(group_x) * input.threads_num[0];
 	const auto     size        = descriptor.GetSize();
-	if (invocations != descriptor.NumRecords() || size == 0 || size > UINT32_MAX ||
+	const auto elements = descriptor.NumRecords();
+	const bool covered = fill.bounded_elements != 0
+	                         ? fill.bounded_elements == elements && invocations >= elements
+	                         : invocations == elements;
+	if (!covered || invocations > uint64_t {UINT32_MAX} + 1 || size == 0 || size > UINT32_MAX ||
 	    (input.dispatch_thread_dimensions &&
 	     (group_x % input.threads_num[0] != 0 || input.dispatch_threads_num[0] != group_x ||
 	      input.dispatch_threads_num[1] != 1 || input.dispatch_threads_num[2] != 1))) {
@@ -122,9 +132,9 @@ bool ResolveComputeBufferFill(const ShaderComputeInputInfo& input, uint32_t grou
 	return true;
 }

-bool RenderExecutor::TryConsumeComputeImageClear(const ShaderComputeInputInfo& input,
-                                                CommandBuffer& command, uint32_t group_x,
-                                                uint32_t group_y, uint32_t group_z, uint32_t mode) {
+bool RenderExecutor::TryConsumeComputeFill(const ShaderComputeInputInfo& input,
+                                          CommandBuffer& command, uint32_t group_x,
+                                          uint32_t group_y, uint32_t group_z, uint32_t mode) {
 	const auto& program   = *input.stage.program;
 	const auto& resources = *input.stage.resources;
 	const auto& fill      = resources.uniform_fill;
@@ -180,14 +190,26 @@ bool RenderExecutor::TryConsumeComputeImageClear(const ShaderComputeInputInfo& i
 	                              size)) {
 		return false;
 	}
-	if (!cache.ClearImageFromBuffer(command, descriptor.Base48(), size, packed_clear)) {
-		return false;
+	const bool image_clear =
+	    cache.ClearImageFromBuffer(command, descriptor.Base48(), size, packed_clear);
+	if (!image_clear) {
+		// CPU fills invalidate overlapping cached images. Preserve the native dispatch for
+		// partial or ambiguous image backings, including pooled metadata aliases.
+		if ((descriptor.Base48() & 3u) != 0 ||
+		    !m_context.IsMapped(descriptor.Base48(), size) ||
+		    cache.IsRegionRegistered(descriptor.Base48(), size)) {
+			return false;
+		}
+		// Metadata need not have an image backing yet. Keep proven fills in guest memory
+		// when CPU-owned, so color-clear discovery can read them without a GPU download.
+		m_context.GetBufferCache().FillBuffer(descriptor.Base48(), size, packed_clear, false);
 	}
 	static std::atomic<uint32_t> logged_clears {0};
 	if (logged_clears.fetch_add(1, std::memory_order_relaxed) < 32) {
-		LOGF("GraphicsRenderDispatchDirect: compute image clear shader=0x%016" PRIx64
+		LOGF("GraphicsRenderDispatchDirect: compute %s shader=0x%016" PRIx64
 		     " addr=0x%016" PRIx64 " size=0x%016" PRIx64 " value=0x%08" PRIx32 "\n",
-		     input.stage.program->shader_hash, descriptor.Base48(), size, packed_clear);
+		     image_clear ? "image clear" : "buffer fill", input.stage.program->shader_hash,
+		     descriptor.Base48(), size, packed_clear);
 	}
 	return true;
 }
@@ -252,10 +274,6 @@ void RenderExecutor::DispatchDirect(uint64_t submit_id, CommandBuffer& buffer,
 		return;
 	}

-	if (sh_ctx.GetCs().cs_regs.data_addr == 0) {
-		return;
-	}
-
 	constexpr uint32_t DISPATCH_INITIATOR_USE_THREAD_DIMENSIONS = 1u << 5u;
 	constexpr uint32_t DISPATCH_INITIATOR_BASE_BITS             = 0x41u;
 	constexpr uint32_t DISPATCH_INITIATOR_MODIFIER_BITS         = 0xa038u;
@@ -293,8 +311,10 @@ void RenderExecutor::DispatchDirect(uint64_t submit_id, CommandBuffer& buffer,
 			input_info.workgroup_counts[axis] = threads / size + (threads % size != 0u);
 		}
 	}
+	m_context.GetImageHeap().BeginDraw();
 	const auto compute_program =
 	    m_context.GetPipelineCache().GetComputeProgram(cs_regs, sh_regs, input_info);
+	if (!compute_program) return;
 	if (use_thread_dimensions) {
 		input_info.dispatch_threads_num[0]    = thread_group_x;
 		input_info.dispatch_threads_num[1]    = thread_group_y;
@@ -303,14 +323,12 @@ void RenderExecutor::DispatchDirect(uint64_t submit_id, CommandBuffer& buffer,

 	const auto& program   = *input_info.stage.program;
 	const auto& resources = *input_info.stage.resources;
-	if (resources.specialization_reads.empty()) {
-		// Cached metadata addresses can be reused as ordinary buffers; preserve their guest writes.
-		UpdateComputeMetaClears(input_info, buffer);
-		if (TryConsumeComputeImageClear(input_info, buffer, thread_group_x, thread_group_y,
-		                               thread_group_z, mode)) {
-			ResetBindings();
-			return;
-		}
+	if (resources.specialization_reads.empty() &&
+	    (TryConsumeComputeMetaClear(input_info, buffer) ||
+	     TryConsumeComputeFill(input_info, buffer, thread_group_x, thread_group_y,
+	                           thread_group_z, mode))) {
+		ResetBindings();
+		return;
 	}

 	if (use_thread_dimensions) {
@@ -346,6 +364,9 @@ void RenderExecutor::DispatchDirect(uint64_t submit_id, CommandBuffer& buffer,
 	if (program.info.uses_dma) {
 		m_context.PrepareBda();
 	}
+	if (program.info.uses_image_heap) {
+		m_context.GetImageHeap().Rebind(*this);
+	}
 	RebindImages(bindings);
 	BindSharedMemory(m_context, input_info, bindings);
 	RebindBuffers(bindings);
@@ -430,7 +451,28 @@ void RenderExecutor::DispatchDirect(uint64_t submit_id, CommandBuffer& buffer,
 	}

 	vk_buffer.bindPipeline(vk::PipelineBindPoint::eCompute, pipeline.pipeline);
-	vk_buffer.dispatch(thread_group_x, thread_group_y, thread_group_z);
+	constexpr uint32_t DISPATCH_INITIATOR_ORDERED_APPEND_ENABLE = 1u << 3u;
+	const uint64_t     group_threads = uint64_t {std::max(input_info.threads_num[0], 1u)} *
+	                                   std::max(input_info.threads_num[1], 1u) *
+	                                   std::max(input_info.threads_num[2], 1u);
+	if ((mode & DISPATCH_INITIATOR_ORDERED_APPEND_ENABLE) != 0 && program.info.uses_gds &&
+	    group_threads <= input_info.wave_size && !input_info.lds_storage) {
+		// A single guest wave per group can run in launch order without a device-wide
+		// spinlock. Keep the guest workgroup IDs when splitting the dispatch.
+		for (uint32_t z = 0; z < thread_group_z; ++z) {
+			for (uint32_t y = 0; y < thread_group_y; ++y) {
+				for (uint32_t x = 0; x < thread_group_x; ++x) {
+					if ((x | y | z) != 0u) {
+						ShaderAccessBarrier(vk_buffer, vk::PipelineStageFlagBits::eComputeShader);
+					}
+					vk_buffer.dispatchBase(x, y, z, 1, 1, 1);
+				}
+			}
+		}
+	} else {
+		vk_buffer.dispatch(thread_group_x, thread_group_y, thread_group_z);
+	}
+	Common::DebugCounters::Add(Common::DebugCounters::Counter::Dispatches);

 	// The removed host fence also ordered read-only dispatches before later writers.
 	ShaderAccessBarrier(vk_buffer, vk::PipelineStageFlagBits::eComputeShader);
@@ -459,8 +501,10 @@ void RenderExecutor::DispatchIndirect(uint64_t submit_id, CommandBuffer& buffer,
 	ShaderComputeInputInfo input_info {};
 	input_info.async_compute = async_compute;
 	input_info.indirect_args_address = args_addr;
+	m_context.GetImageHeap().BeginDraw();
 	const auto compute_program = m_context.GetPipelineCache().GetComputeProgram(
 	    cs_regs, buffer.GetRegisters().GetShaderRegisters(), input_info);
+	if (!compute_program) return;
 	buffer.EndRendering();
 	auto& pipeline = m_context.GetPipelineCache().GetComputePipeline(input_info, compute_program);
 	auto& bindings = m_compute_bindings;
@@ -472,6 +516,9 @@ void RenderExecutor::DispatchIndirect(uint64_t submit_id, CommandBuffer& buffer,
 		m_context.PrepareBda();
 	}
 	BindSharedMemory(m_context, input_info, bindings, args_addr);
+	if (program.info.uses_image_heap) {
+		m_context.GetImageHeap().Rebind(*this);
+	}
 	RebindImages(bindings);
 	// Acquiring arguments can merge cache buffers; finalize shader bindings afterward.
 	const auto [args_buffer, args_offset] = m_context.GetBufferCache().ObtainBuffer(
@@ -500,6 +547,7 @@ void RenderExecutor::DispatchIndirect(uint64_t submit_id, CommandBuffer& buffer,
 	                          1, &barrier, 0, nullptr, 0, nullptr);
 	vk_buffer.bindPipeline(vk::PipelineBindPoint::eCompute, pipeline.pipeline);
 	vk_buffer.dispatchIndirect(args_buffer->Handle(), args_offset);
+	Common::DebugCounters::Add(Common::DebugCounters::Counter::Dispatches);
 	ShaderAccessBarrier(vk_buffer, vk::PipelineStageFlagBits::eComputeShader);
 	ResetBindings();
 }
@@ -13,9 +13,10 @@ namespace Libs::Graphics {
 RenderContext::RenderContext(GraphicContext& graphics)
    : m_graphics(graphics), m_render_executor(*this), m_command_scheduler(*this, graphics),
      m_descriptor_heap(graphics, m_command_scheduler.GetMasterSemaphore()),
-      m_pipeline_cache(graphics), m_sampler_cache(graphics),
+      m_pipeline_cache(graphics, &m_image_heap), m_sampler_cache(graphics),
      m_buffer_cache(graphics, m_command_scheduler, m_page_manager, m_texture_cache),
-      m_texture_cache(graphics, m_command_scheduler, m_page_manager, m_buffer_cache) {
+      m_texture_cache(graphics, m_command_scheduler, m_page_manager, m_buffer_cache),
+      m_image_heap(*this, m_page_manager) {
 	EXIT_NOT_IMPLEMENTED(!Common::Thread::IsMainThread());
 }

@@ -64,6 +65,7 @@ bool RenderContext::HandleFault(PageFaultAccess access, uint64_t fault_vaddr) no
 	if (access == PageFaultAccess::Write) {
 		m_buffer_cache.InvalidateMemory(fault_vaddr, fault_size);
 		m_texture_cache.InvalidateMemory(fault_vaddr, fault_size);
+		m_image_heap.InvalidateMemory(fault_vaddr, fault_size);
 	} else {
 		m_buffer_cache.ReadMemory(fault_vaddr, fault_size);
 	}
@@ -76,6 +78,7 @@ bool RenderContext::InvalidateMemory(uint64_t vaddr, uint64_t size) {
 	}
 	m_buffer_cache.InvalidateMemory(vaddr, size);
 	m_texture_cache.InvalidateMemory(vaddr, size);
+	m_image_heap.InvalidateMemory(vaddr, size);
 	return true;
 }

@@ -90,6 +93,7 @@ bool RenderContext::IsMapped(uint64_t vaddr, uint64_t size) const noexcept {
 void RenderContext::MapMemory(uint64_t vaddr, uint64_t size) {
 	std::lock_guard lock(m_mapped_ranges_mutex);
 	m_mapped_ranges.Add(vaddr, size);
+	m_buffer_cache.QueueBdaSync(vaddr, size);
 }

 void RenderContext::UnmapMemory(uint64_t vaddr, uint64_t size) {
@@ -111,6 +115,7 @@ void RenderContext::UnmapMemory(uint64_t vaddr, uint64_t size) {
 		}
 		m_buffer_cache.InvalidateMemory(vaddr, size);
 		m_texture_cache.UnmapMemory(vaddr, size);
+		m_image_heap.UnmapMemory(vaddr, size);
 		std::lock_guard lock(m_mapped_ranges_mutex);
 		m_mapped_ranges.Subtract(vaddr, size);
 	};
@@ -129,9 +134,7 @@ void RenderContext::PrepareBda() {
 		m_bda_logged = true;
 	}
 	std::shared_lock lock(m_mapped_ranges_mutex);
-	m_mapped_ranges.ForEach([this](uint64_t start, uint64_t end) {
-		m_buffer_cache.SynchronizeBuffersInRange(start, end - start);
-	});
+	m_buffer_cache.SynchronizeBda(m_mapped_ranges);
 	m_fault_process_pending = true;
 }

@@ -12,6 +12,7 @@
 #include "graphics/host_gpu/renderer/cache/textureCache.h"
 #include "graphics/host_gpu/renderer/commandScheduler.h"
 #include "graphics/host_gpu/renderer/pipeline/descriptorHeap.h"
+#include "graphics/host_gpu/renderer/pipeline/imageHeap.h"
 #include "graphics/host_gpu/renderer/pipeline/pipelineCache.h"
 #include "kernel/eventQueue.h"

@@ -43,6 +44,7 @@ public:
 	CommandScheduler&   GetCommandScheduler() { return m_command_scheduler; }
 	PipelineCache&      GetPipelineCache() { return m_pipeline_cache; }
 	DescriptorHeap&     GetDescriptorHeap() { return m_descriptor_heap; }
+	ImageHeap&          GetImageHeap() { return m_image_heap; }
 	SamplerCache&       GetSamplerCache() { return m_sampler_cache; }
 	BufferCache&        GetBufferCache() { return m_buffer_cache; }
 	TextureCache&       GetTextureCache() { return m_texture_cache; }
@@ -76,6 +78,7 @@ private:
 	PageManager               m_page_manager;
 	BufferCache               m_buffer_cache;
 	TextureCache              m_texture_cache;
+	ImageHeap                 m_image_heap;
 	mutable std::shared_mutex m_mapped_ranges_mutex;
 	RangeSet                  m_mapped_ranges;
 	std::unique_ptr<GuestGpu> m_gpu;
@@ -2,6 +2,7 @@

 #include "common/assert.h"
 #include "common/common.h"
+#include "common/debugCounters.h"
 #include "common/emulatorConfig.h"
 #include "common/file.h"
 #include "common/logging/log.h"
@@ -435,7 +436,7 @@ static bool DrawHasValidVertexShader(const HW::Shader& sh_ctx) {
 static bool PixelShaderHasDepthOrCoverageSideEffects(const HW::ShaderRegisters& sh_regs) {
 	const auto& db = sh_regs.db_shader_control;
 	return db.shader_kill_enable || db.shader_z_export_enable || db.shader_mask_export_enable ||
-	       db.shader_dual_export_enable || db.shader_execute_on_noop;
+	       db.shader_dual_export_enable || db.shader_execute_on_noop || sh_regs.AlphaToMask();
 }

 struct DrawRenderState {
@@ -507,14 +508,15 @@ RenderState RenderExecutor::AcquireRenderTargets(CommandBuffer& buffer, RenderCo
 		}
 		const auto  image_view = cache.FindDepthTarget(depth.image_id, depth.desc);
 		const auto& metadata   = depth.desc.info.metadata;
-		if (metadata.kind == ImageMetadataKind::Htile && depth.depth_clear_enable &&
-		    !cache.ClearMeta(metadata.range.address)) {
-			EXIT("failed to acquire HTile metadata for a depth clear\n");
-		}
 		const bool meta_clear =
 		    metadata.kind == ImageMetadataKind::Htile &&
 		    cache.IsMetaCleared(metadata.range.address, depth.desc.view_info.base_layer);
-		depth.depth_load_clear_enable = depth.depth_clear_enable || meta_clear;
+		depth.load_clear_aspects = {};
+		if (depth.depth_clear_enable || meta_clear)
+			depth.load_clear_aspects |= vk::ImageAspectFlagBits::eDepth;
+		// Apply pending clears to compressed stencil as well as depth.
+		if (depth.stencil_clear_enable || (meta_clear && metadata.stencil_compressed))
+			depth.load_clear_aspects |= vk::ImageAspectFlagBits::eStencil;
 		if (meta_clear &&
 		    !cache.TouchMeta(metadata.range.address, depth.desc.view_info.base_layer, false)) {
 			EXIT("failed to consume HTile clear state\n");
@@ -536,7 +538,10 @@ RenderState RenderExecutor::AcquireRenderTargets(CommandBuffer& buffer, RenderCo
 			}
 		}
 		if (feedback_aspects && !m_context.GetGraphics().attachment_feedback_loop_enabled) {
-			EXIT("depth attachment feedback loop is not supported by the host\n");
+			static std::atomic_bool logged = false;
+			if (!logged.exchange(true, std::memory_order_relaxed)) {
+				std::printf("depth attachment feedback loop is not supported by the host\n");
+			}
 		}
 		auto layout = depth_attachment_layout(depth);
 		if (sampled_aspects & ~DepthReadableAspects(layout)) {
@@ -564,9 +569,11 @@ RenderState RenderExecutor::AcquireRenderTargets(CommandBuffer& buffer, RenderCo
 		attachment.clear_value[0] = std::bit_cast<uint32_t>(depth.depth_clear_value);
 		attachment.clear_value[1] = depth.stencil_clear_value;
 		attachment.has_depth      = static_cast<bool>(aspects & vk::ImageAspectFlagBits::eDepth);
-		attachment.depth_clear    = depth.depth_load_clear_enable;
+		attachment.depth_clear =
+		    static_cast<bool>(depth.load_clear_aspects & vk::ImageAspectFlagBits::eDepth);
 		attachment.has_stencil    = static_cast<bool>(aspects & vk::ImageAspectFlagBits::eStencil);
-		attachment.stencil_clear  = depth.stencil_clear_enable;
+		attachment.stencil_clear =
+		    static_cast<bool>(depth.load_clear_aspects & vk::ImageAspectFlagBits::eStencil);
 	}
 	if (color_count == 0 && !depth.image_id) {
 		const auto& limits = buffer.GetGraphics().GetPhysicalDeviceProperties().limits;
@@ -624,10 +631,11 @@ bool RenderExecutor::ConsumeMetadataColorOperation(CommandBuffer& buffer,
 	       mode == static_cast<uint8_t>(CbColorMode::DccDecompress);
 }

+// vertex_offset is the first vertex of non-indexed draws.
 struct DrawEmitInfo {
-	int32_t  vertex_offset = 0;
-	uint32_t first_vertex  = 0;
-	uint32_t first_instance = 0;
+	int32_t          vertex_offset  = 0;
+	uint32_t         first_instance = 0;
+	DrawIndirectArgs indirect;
 };

 struct DrawIndexBufferSource {
@@ -885,6 +893,7 @@ static void RefreshShaders(CommandBuffer& buffer, const DrawCallInfo& draw,
 		}
 	}
 	auto& pipeline_cache = buffer.GetContext().GetPipelineCache();
+	buffer.GetContext().GetImageHeap().BeginDraw();
 	if (draw.IsIndexed()) {
 		LogDrawPhase(draw.Name(), "GetGraphicsPrograms");
 	}
@@ -928,6 +937,29 @@ bool RenderExecutor::PrepareDrawRenderState(CommandBuffer& buffer, const DrawCal
 		LogDrawPhase(draw.Name(), "ResolveRenderDepthTarget");
 	}
 	ResolveRenderDepthTarget(buffer, state.depth_info);
+	// TODO(PPSA09243): Vulkan passes take one sample count. Where PS5 attachments mix counts,
+	// keep those matching the raster count instead of resolving multisampled shadows.
+	const auto colors       = std::span {state.color_info, state.color_count};
+	const auto pass_samples = state.depth_info.image_id ? state.depth_info.desc.info.samples
+	                                                    : state.color_info[0].desc.info.samples;
+	if (std::ranges::any_of(colors, [&](const RenderColorInfo& color) {
+		    return color.desc.info.samples != pass_samples;
+	    })) {
+		const auto raster =
+		    render_sample_count(buffer.GetRegisters().GetAaConfig().msaa_num_samples);
+		const auto dropped = std::ranges::remove_if(colors, [&](const RenderColorInfo& color) {
+			return color.desc.info.samples != raster;
+		});
+		state.color_count -= static_cast<uint32_t>(dropped.size());
+		if (state.depth_info.image_id && state.depth_info.desc.info.samples != raster) {
+			state.depth_info = {};
+		}
+		static std::atomic_bool logged = false;
+		if (!logged.exchange(true, std::memory_order_relaxed)) {
+			Log::WriteToConsoleAndLog(
+			    "Warning: dropping attachments whose sample count differs from the raster.\n");
+		}
+	}

 	if (state.color_count == 0 && !state.depth_info.image_id && !state.ps_active) {
 		LogFramebufferSkip(draw.Name(), state.color_info[0], state.depth_info, buffer,
@@ -1013,8 +1045,8 @@ static void EmitDrawPrimitives(const HW::UserConfig& ucfg, vk::CommandBuffer vk_
 				vk_buffer.drawIndexed(draw.index_count, draw.instance_count, 0, emit.vertex_offset,
 				                      emit.first_instance);
 			} else {
-				vk_buffer.draw(draw.index_count, draw.instance_count, emit.first_vertex,
-				               emit.first_instance);
+				vk_buffer.draw(draw.index_count, draw.instance_count,
+				               static_cast<uint32_t>(emit.vertex_offset), emit.first_instance);
 			}
 			break;
 		case Prospero::PrimitiveType::kQuadListLegacy:
@@ -1024,7 +1056,8 @@ static void EmitDrawPrimitives(const HW::UserConfig& ucfg, vk::CommandBuffer vk_
 					vk_buffer.drawIndexed(4, draw.instance_count, i, emit.vertex_offset,
 					                      emit.first_instance);
 				} else {
-					vk_buffer.draw(4, draw.instance_count, i + emit.first_vertex,
+					vk_buffer.draw(4, draw.instance_count,
+					               i + static_cast<uint32_t>(emit.vertex_offset),
 					               emit.first_instance);
 				}
 			}
@@ -1045,7 +1078,7 @@ bool RenderExecutor::CopyDepthToColor(CommandBuffer& buffer, const DrawCallInfo&
 	EXIT_NOT_IMPLEMENTED(!m_context.GetGraphics().maintenance8_enabled ||
 	                     !Prospero::IsRectList(ucfg.GetPrimType()) ||
 	                     draw.IsIndexed() || draw.index_count != 3 || draw.instance_count != 1 ||
-	                     emit.first_vertex != 0 || emit.first_instance != 0 ||
+	                     emit.vertex_offset != 0 || emit.first_instance != 0 ||
 	                     state.color_count != 1 || color.target_slot != 0 || !depth.image_id ||
 	                     color.desc.info.samples != 1 ||
 	                     depth.desc.info.extent.width != color.Extent().width ||
@@ -1057,7 +1090,7 @@ bool RenderExecutor::CopyDepthToColor(CommandBuffer& buffer, const DrawCallInfo&
 	                     state.ps_input_info.ps_sample_mask_export_enable);
 	vk::ImageAspectFlags feedback;
 	const auto rendering = AcquireRenderTargets(buffer, &color, 1, depth, feedback);
-	if (depth.depth_load_clear_enable) {
+	if (depth.load_clear_aspects) {
 		m_context.GetCommandScheduler().BeginRendering(rendering);
 	}
 	auto& cache = m_context.GetTextureCache();
@@ -1159,6 +1192,14 @@ void RenderExecutor::ExecutePreparedDraw(uint64_t submit_id, CommandBuffer& buff
 		vertex_bindings = AcquireVertexBuffers(buffer, state.vertex_info[0]);
 		index_binding   = PrepareIndexBuffer(buffer, index_source);
 	}
+	std::pair<Buffer*, uint64_t> args_buffer {};
+	if (emit.indirect.address != 0) {
+		const auto size = static_cast<uint64_t>(emit.indirect.count - 1) * emit.indirect.stride +
+		                  (draw.IsIndexed() ? sizeof(vk::DrawIndexedIndirectCommand)
+		                                    : sizeof(vk::DrawIndirectCommand));
+		args_buffer = m_context.GetBufferCache().ObtainBuffer(emit.indirect.address, size, false);
+		EXIT_IF((args_buffer.second & 3u) != 0);
+	}
 	if (draw.IsIndexed()) {
 		LogDrawPhase(draw.Name(), "CreatePipeline");
 	}
@@ -1186,26 +1227,33 @@ void RenderExecutor::ExecutePreparedDraw(uint64_t submit_id, CommandBuffer& buff
 	if (!mesh_active) {
 		CommitIndexBuffer(vk_buffer, index_binding);
 	}
-	const auto draw_dwords = ShaderRecompiler::IR::PushData::GraphicsDrawDwords(
-	    mesh_active, state.ps_input_info.object_primitive_id);
+	using ShaderRecompiler::IR::PushData;
+	const auto draw_stages =
+	    vk::ShaderStageFlagBits::eFragment |
+	    (mesh_active ? vk::ShaderStageFlagBits::eMeshEXT : vk::ShaderStageFlagBits::eVertex);
+	const auto draw_dwords =
+	    PushData::GraphicsDrawDwords(mesh_active, state.ps_input_info.object_primitive_id);
 	if (draw_dwords != 0) {
-		std::array<uint32_t, ShaderRecompiler::IR::PushData::MeshDrawDwordCount> draw_data {};
-		const auto draw_stages = vk::ShaderStageFlagBits::eFragment |
-		                         (mesh_active ? vk::ShaderStageFlagBits::eMeshEXT
-		                                      : vk::ShaderStageFlagBits::eVertex);
+		std::array<uint32_t, PushData::MeshDrawDwordCount> draw_data {};
 		if (mesh_active) {
-			draw_data = {
-			    draw.index_count,
-			    draw.IsIndexed() ? static_cast<uint32_t>(emit.vertex_offset) : emit.first_vertex,
-			    emit.first_instance, index_source.guest_element_size,
-			    static_cast<uint32_t>(index_source.address),
-			    static_cast<uint32_t>(index_source.address >> 32u)};
+			draw_data = {draw.index_count,
+			             static_cast<uint32_t>(emit.vertex_offset),
+			             emit.first_instance,
+			             index_source.guest_element_size,
+			             static_cast<uint32_t>(index_source.address),
+			             static_cast<uint32_t>(index_source.address >> 32u)};
 		} else {
 			draw_data[0] = buffer.GetUserConfig().GetObjectId();
 		}
 		vk_buffer.pushConstants(pipeline.pipeline_layout, draw_stages, 0,
 		                        draw_dwords * sizeof(uint32_t), draw_data.data());
 	}
+	if (const auto planes = vertex_stages.back().user_clip_planes; planes != 0) {
+		vk_buffer.pushConstants(pipeline.pipeline_layout, draw_stages,
+		                        PushData::UserClipPlaneDword * sizeof(uint32_t),
+		                        std::bit_width(planes) * 4u * sizeof(uint32_t),
+		                        buffer.GetRegisters().GetUserClipPlanes().data());
+	}

 	auto& dynamic_state = m_context.GetCommandScheduler().GetDynamicState();
 	UpdateGraphicsDynamicState(buffer, dynamic_state, vertex_stages.back(), state.depth_info,
@@ -1213,6 +1261,20 @@ void RenderExecutor::ExecutePreparedDraw(uint64_t submit_id, CommandBuffer& buff
 	dynamic_state.SetAttachmentFeedbackLoopAspects(feedback_aspects);
 	dynamic_state.Commit(m_context.GetGraphics(), vk_buffer);

+	vk::PipelineStageFlags shader_write_stages = {};
+	for (const auto& stage: vertex_stages) {
+		if (HasShaderBufferWrites(stage.stage)) {
+			shader_write_stages |= ShaderPipelineStages(NativeShaderStage(stage.logical_stage));
+		}
+	}
+	if (state.ps_active && HasShaderBufferWrites(state.ps_input_info.stage)) {
+		shader_write_stages |= vk::PipelineStageFlagBits::eFragmentShader;
+	}
+	if (shader_write_stages) {
+		// As for dispatches, buffer writes wait for earlier accesses to the memory.
+		m_context.GetCommandScheduler().EndRendering();
+		ShaderWriteHazardBarrier(vk_buffer, shader_write_stages);
+	}
 	LogDrawPhase(draw.Name(), "BeginRendering");
 	if (!draw.IsIndexed()) {
 		SetDrawDebugPhase(buffer, submit_id, draw, 0x400u);
@@ -1224,22 +1286,21 @@ void RenderExecutor::ExecutePreparedDraw(uint64_t submit_id, CommandBuffer& buff
 	}
 	if (mesh_active) {
 		vk_buffer.drawMeshTasksEXT(mesh_groups, draw.instance_count, 1);
-	} else {
+	} else if (args_buffer.first == nullptr) {
 		EmitDrawPrimitives(ucfg, vk_buffer, draw, emit);
+	} else if (draw.IsIndexed()) {
+		vk_buffer.drawIndexedIndirect(args_buffer.first->Handle(), args_buffer.second,
+		                              emit.indirect.count, emit.indirect.stride);
+	} else {
+		vk_buffer.drawIndirect(args_buffer.first->Handle(), args_buffer.second,
+		                       emit.indirect.count, emit.indirect.stride);
 	}
+	Common::DebugCounters::Add(Common::DebugCounters::Counter::Draws);
+	Common::DebugCounters::Add(Common::DebugCounters::Counter::DrawInstances, draw.instance_count);

 	if (!draw.IsIndexed()) {
 		SetDrawDebugPhase(buffer, submit_id, draw, 0x600u);
 	}
-	vk::PipelineStageFlags shader_write_stages = {};
-	for (const auto& stage: vertex_stages) {
-		if (HasShaderBufferWrites(stage.stage)) {
-			shader_write_stages |= ShaderPipelineStages(NativeShaderStage(stage.logical_stage));
-		}
-	}
-	if (state.ps_active && HasShaderBufferWrites(state.ps_input_info.stage)) {
-		shader_write_stages |= vk::PipelineStageFlagBits::eFragmentShader;
-	}
 	if (shader_write_stages) {
 		m_context.GetCommandScheduler().EndRendering();
 		ShaderWriteBarrier(vk_buffer, shader_write_stages);
@@ -1250,7 +1311,30 @@ void RenderExecutor::ExecutePreparedDraw(uint64_t submit_id, CommandBuffer& buff
 	}
 }

-void RenderExecutor::DrawIndex(uint64_t submit_id, CommandBuffer& buffer,
+// Vulkan offsets for a prepared draw, or nullopt when its indirect records need CPU counts, as
+// mesh and quad-list draws do. Indirect records carry their own offsets.
+static std::optional<DrawEmitInfo> ResolveDrawEmit(const HW::UserConfig&        ucfg,
+                                                   const ShaderVertexInputInfo& vs,
+                                                   DrawOffsetSource             source,
+                                                   int32_t                      first_vertex,
+                                                   uint32_t                     first_instance,
+                                                   const DrawIndirectArgs&      indirect) {
+	if (indirect.address != 0) {
+		if (vs.stage.program->stage == ShaderType::Mesh ||
+		    ucfg.GetPrimType() == Prospero::PrimitiveType::kQuadListLegacy) {
+			return std::nullopt;
+		}
+		return DrawEmitInfo {.indirect = indirect};
+	}
+	if (source == DrawOffsetSource::IndirectArgs) {
+		// Native fetches consume the patched SGPRs; rewritten fetches need Vulkan offsets.
+		return vs.fetch_embedded ? DrawEmitInfo {first_vertex, first_instance} : DrawEmitInfo {};
+	}
+	const auto [vertex_offset, instance_offset] = ResolveDrawOffsets(ucfg.GetIndexOffset(), vs);
+	return DrawEmitInfo {vertex_offset + first_vertex, instance_offset};
+}
+
+bool RenderExecutor::DrawIndex(uint64_t submit_id, CommandBuffer& buffer,
                               const DrawIndexArgs& args) {
 	KYTY_PROFILER_FUNCTION();

@@ -1265,19 +1349,19 @@ void RenderExecutor::DrawIndex(uint64_t submit_id, CommandBuffer& buffer,
 	                    reinterpret_cast<uint64_t>(args.index_addr));

 	Common::LockGuard lock(m_context.GetMutex());
-	if (args.index_count == 0 || args.instance_count == 0) {
-		return;
+	if (args.index_count == 0 || (args.instance_count == 0 && args.indirect.address == 0)) {
+		return true;
 	}

 	if (ConsumeMetadataColorOperation(buffer, args.render_target_slice_offset) ||
 	    DepthStencilCopy(buffer) ||
 	    ResolveColorTargets(buffer, args.render_target_slice_offset)) {
 		ResetBindings();
-		return;
+		return true;
 	}

 	if (!DrawHasValidVertexShader(sh_ctx)) {
-		return;
+		return true;
 	}

 	if (graphics_debug_dump_enabled()) {
@@ -1303,7 +1387,7 @@ void RenderExecutor::DrawIndex(uint64_t submit_id, CommandBuffer& buffer,

 	vk::PrimitiveTopology topology = vk::PrimitiveTopology::ePointList;
 	if (!GetDrawTopology(ucfg, topology)) {
-		return;
+		return true;
 	}

 	DrawIndexBufferSource index_source {};
@@ -1342,30 +1426,23 @@ void RenderExecutor::DrawIndex(uint64_t submit_id, CommandBuffer& buffer,
 	DrawRenderState state {};
 	if (!PrepareDrawRenderState(buffer, draw, args.render_target_slice_offset, state)) {
 		ResetBindings();
-		return;
+		return true;
 	}

 	LogDrawStateIfNeeded(buffer, draw, state, args.index_type_and_size,
 	                     args.index_addr);
-
-	const bool indirect = args.offset_source == DrawOffsetSource::IndirectArgs;
-	const auto [vertex_offset, instance_offset] =
-	    indirect ? std::pair<int32_t, uint32_t> {0, args.first_instance}
-	             : ResolveDrawOffsets(ucfg.GetIndexOffset(), state.vertex_info[0]);
-
-	DrawEmitInfo emit {};
-	// Native fetches consume the patched SGPRs; rewritten fetches need Vulkan offsets.
-	const bool native_indirect = indirect && !state.vertex_info[0].fetch_embedded;
-	emit.vertex_offset  = native_indirect ? 0 : vertex_offset + args.base_vertex;
-	emit.first_instance = native_indirect ? 0 : instance_offset;
-
-	ExecutePreparedDraw(submit_id, buffer, draw, state, topology, emit, index_source,
-	                    primitive_restart);
+	const auto emit = ResolveDrawEmit(ucfg, state.vertex_info[0], args.offset_source,
+	                                  args.base_vertex, args.first_instance, args.indirect);
+	if (emit) {
+		ExecutePreparedDraw(submit_id, buffer, draw, state, topology, *emit, index_source,
+		                    primitive_restart);
+	}
 	ResetBindings();
+	return emit.has_value();
 }

 // NOLINTNEXTLINE(readability-function-cognitive-complexity)
-void RenderExecutor::DrawAuto(uint64_t submit_id, CommandBuffer& buffer, const DrawAutoArgs& args) {
+bool RenderExecutor::DrawAuto(uint64_t submit_id, CommandBuffer& buffer, const DrawAutoArgs& args) {
 	KYTY_PROFILER_FUNCTION();

 	EXIT_IF(buffer.IsInvalid());
@@ -1379,19 +1456,19 @@ void RenderExecutor::DrawAuto(uint64_t submit_id, CommandBuffer& buffer, const D
 	                    args.first_instance);

 	Common::LockGuard lock(m_context.GetMutex());
-	if (args.vertex_count == 0 || args.instance_count == 0) {
-		return;
+	if (args.indirect.address == 0 && (args.vertex_count == 0 || args.instance_count == 0)) {
+		return true;
 	}

 	if (ConsumeMetadataColorOperation(buffer, args.render_target_slice_offset) ||
 	    DepthStencilCopy(buffer) ||
 	    ResolveColorTargets(buffer, args.render_target_slice_offset)) {
 		ResetBindings();
-		return;
+		return true;
 	}

 	if (!DrawHasValidVertexShader(sh_ctx)) {
-		return;
+		return true;
 	}

 	if (graphics_debug_dump_enabled()) {
@@ -1417,12 +1494,12 @@ void RenderExecutor::DrawAuto(uint64_t submit_id, CommandBuffer& buffer, const D
 	vk::PrimitiveTopology topology = vk::PrimitiveTopology::ePointList;
 	if (!GetDrawTopology(ucfg, topology)) {
 		ResetBindings();
-		return;
+		return true;
 	}
 	DrawRenderState state {};
 	if (!PrepareDrawRenderState(buffer, draw, args.render_target_slice_offset, state)) {
 		ResetBindings();
-		return;
+		return true;
 	}

 	const bool rect_list = Prospero::IsRectList(ucfg.GetPrimType());
@@ -1436,24 +1513,18 @@ void RenderExecutor::DrawAuto(uint64_t submit_id, CommandBuffer& buffer, const D
 			     sh_ctx.GetVs().es_regs.data_addr, sh_ctx.GetVs().gs_regs.data_addr);
 		}
 		ResetBindings();
-		return;
+		return true;
 	}

 	LogDrawStateIfNeeded(buffer, draw, state, 0, nullptr);
-
-	const bool indirect = args.offset_source == DrawOffsetSource::IndirectArgs;
-	const auto [vertex_offset, instance_offset] =
-	    indirect ? std::pair<int32_t, uint32_t> {0, args.first_instance}
-	             : ResolveDrawOffsets(ucfg.GetIndexOffset(), state.vertex_info[0]);
-	DrawEmitInfo emit {};
-	const bool native_indirect = indirect && !state.vertex_info[0].fetch_embedded;
-	emit.first_vertex = native_indirect ? 0 :
-	    static_cast<uint32_t>(vertex_offset + static_cast<int32_t>(args.first_vertex));
-	emit.first_instance = native_indirect ? 0 : instance_offset;
-
-	DrawIndexBufferSource index_source {};
-	ExecutePreparedDraw(submit_id, buffer, draw, state, topology, emit, index_source, false);
+	const auto emit = ResolveDrawEmit(ucfg, state.vertex_info[0], args.offset_source,
+	                                  static_cast<int32_t>(args.first_vertex), args.first_instance,
+	                                  args.indirect);
+	if (emit) {
+		ExecutePreparedDraw(submit_id, buffer, draw, state, topology, *emit, {}, false);
+	}
 	ResetBindings();
+	return emit.has_value();
 }

 bool RenderExecutor::ResolveColorTargets(CommandBuffer& buffer, uint32_t render_target_slice_offset) {
@@ -0,0 +1,50 @@
+#version 450
+
+// One workgroup per metadata slice; see TileManager::CheckColorClear.
+layout(local_size_x = 256) in;
+
+layout(set = 0, binding = 0, std430) buffer Metadata { uint data[]; } metadata;
+layout(set = 0, binding = 1, std430) writeonly buffer Predicates { uint data[]; } predicates;
+
+layout(push_constant, std430) uniform Check {
+	uint metadata_base;
+	uint words;
+	uint keys;
+	uint expand;
+} params;
+
+const uint Slots = 8u; // TileManager::ColorClearSlots
+
+shared uint key;
+shared uint mismatch;
+
+void main() {
+	const uint slice = gl_WorkGroupID.x;
+	const uint base  = (params.metadata_base >> 2u) + slice * params.words;
+	// Each invocation reads, and may rewrite, only its own words.
+	if (gl_LocalInvocationIndex == 0u) {
+		key      = metadata.data[base] & 0xffu;
+		mismatch = 0u;
+	}
+	barrier();
+	const uint pattern = key * 0x01010101u;
+	bool       same    = true;
+	for (uint i = gl_LocalInvocationIndex; same && i < params.words; i += gl_WorkGroupSize.x) {
+		same = metadata.data[base + i] == pattern;
+	}
+	if (!same) {
+		atomicOr(mismatch, 1u);
+	}
+	barrier();
+	const uint slot = key >> 5u;
+	const bool hit  = mismatch == 0u && (key & 0x1fu) == 0u && (params.keys & (1u << slot)) != 0u;
+	if (gl_LocalInvocationIndex < Slots) {
+		predicates.data[slice * Slots + gl_LocalInvocationIndex] =
+		    uint(hit && slot == gl_LocalInvocationIndex);
+	}
+	if (hit && params.expand != 0u) {
+		for (uint i = gl_LocalInvocationIndex; i < params.words; i += gl_WorkGroupSize.x) {
+			metadata.data[base + i] = 0xffffffffu;
+		}
+	}
+}
@@ -97,6 +97,24 @@ uint64_t GraphicContext::GetDeviceMemoryUsage() const {
 	return usage;
 }

+uint32_t GraphicContext::GetHeapUsage(std::span<HeapUsage, VK_MAX_MEMORY_HEAPS> heaps) const {
+	if (allocator == nullptr) {
+		return 0;
+	}
+	VmaBudget budgets[VK_MAX_MEMORY_HEAPS] {};
+	vmaGetHeapBudgets(allocator, budgets);
+	const auto count = physical_device_memory_properties.memoryHeapCount;
+	for (uint32_t heap = 0; heap < count; heap++) {
+		const auto& properties = physical_device_memory_properties.memoryHeaps[heap];
+		heaps[heap] = {.usage        = budgets[heap].usage,
+		               .budget       = budgets[heap].budget,
+		               .allocations  = budgets[heap].statistics.allocationCount,
+		               .device_local = static_cast<bool>(properties.flags &
+		                                                 vk::MemoryHeapFlagBits::eDeviceLocal)};
+	}
+	return count;
+}
+
 uint64_t GraphicContext::GetTotalMemoryBudget() const {
 	if (allocator == nullptr) {
 		return 0;
--- a/Show More
+++ b/Show More