usbd: return an empty device list from sceUsbdGetDeviceList (Gran Turismo 7) (#1831 )

graphics: draw non-indexed and native-fetch indirect draws from GPU-read records
graphics: draw indexed indirect multi-draws from GPU-read arguments
2026-10-22 20:06:54 +02:00 · 2026-10-22 16:18:28 +02:00 · 2026-10-22 16:12:20 +02:00 · 2026-10-22 16:12:20 +02:00 · 2026-10-22 16:12:20 +02:00 · 2026-10-22 16:12:20 +02:00
273 changed files with 33768 additions and 6211 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
@@ -432,6 +433,14 @@ add_executable(avplayer_file_tests EXCLUDE_FROM_ALL
 target_link_libraries(avplayer_file_tests common FFmpeg::ffmpeg ZArchive::zarchive)
 target_include_directories(avplayer_file_tests PRIVATE ${inc_headers})

+add_executable(vdecsw_tests EXCLUDE_FROM_ALL
+	"${KYTY_TESTS_DIR}/VdecswTests.cpp"
+	"${KYTY_SOURCE_DIR}/libs/videoDec2Decoder.cpp"
+	"${KYTY_SOURCE_DIR}/loader/timer.cpp"
+)
+target_link_libraries(vdecsw_tests common FFmpeg::ffmpeg)
+target_include_directories(vdecsw_tests PRIVATE ${inc_headers})
+
 add_executable(bit_array_tests EXCLUDE_FROM_ALL
 	"${KYTY_TESTS_DIR}/BitArrayTests.cpp"
 )
@@ -518,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"
 )
@@ -622,6 +635,7 @@ if(BUILD_TESTING)
 	add_test(NAME archive_file COMMAND $<TARGET_FILE:archive_file_tests>)
 	add_test(NAME trophy_system COMMAND $<TARGET_FILE:trophy_system_tests>)
 	add_test(NAME avplayer_file COMMAND $<TARGET_FILE:avplayer_file_tests>)
+	add_test(NAME vdecsw COMMAND $<TARGET_FILE:vdecsw_tests>)
 	add_test(NAME bit_array COMMAND $<TARGET_FILE:bit_array_tests>)
 	add_test(NAME lru_cache COMMAND $<TARGET_FILE:lru_cache_tests>)
 	add_test(NAME shader_vertex_metadata COMMAND $<TARGET_FILE:shader_vertex_metadata_tests>)
@@ -635,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
@@ -698,6 +717,7 @@ if(BUILD_TESTING)
 		archive_file_tests
 		trophy_system_tests
 		avplayer_file_tests
+		vdecsw_tests
 		bit_array_tests
 		lru_cache_tests
 		shader_vertex_metadata_tests
@@ -713,6 +733,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
@@ -750,6 +771,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)
@@ -275,11 +275,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 +345,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 +365,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})
@@ -0,0 +1,112 @@
+#include "common/additionalContent.h"
+
+#include "common/stringUtils.h"
+
+#include <algorithm>
+#include <fstream>
+#include <map>
+#include <nlohmann/json.hpp>
+#include <regex>
+#include <system_error>
+
+namespace Common {
+
+namespace {
+
+std::string GetString(const nlohmann::json& object, const std::string& key) {
+	const auto it = object.find(key);
+	return it != object.end() && it->is_string() ? it->get<std::string>() : std::string {};
+}
+
+std::string GetTitle(const nlohmann::json& param) {
+	const auto localized = param.find("localizedParameters");
+	if (localized != param.end() && localized->is_object()) {
+		const auto preferred = localized->find(GetString(*localized, "defaultLanguage"));
+		if (preferred != localized->end()) {
+			if (auto title = GetString(*preferred, "titleName"); !title.empty()) {
+				return title;
+			}
+		}
+		for (const auto& language: *localized) {
+			if (auto title = GetString(language, "titleName"); !title.empty()) {
+				return title;
+			}
+		}
+	}
+	return GetString(param, "titleName");
+}
+
+} // namespace
+
+std::vector<AddContentEntry> FindAddContent(const std::filesystem::path& dir,
+                                            const std::string&           title_id) {
+	if (dir.empty() || title_id.empty()) {
+		return {};
+	}
+
+	std::error_code ec;
+	const auto      root = std::filesystem::absolute(dir, ec);
+	if (ec) {
+		return {};
+	}
+
+	// Scan directories separately so an unreadable child does not discard its siblings.
+	std::vector<std::filesystem::path> pending {root};
+	std::vector<std::filesystem::path> params;
+	while (!pending.empty()) {
+		const auto current = std::move(pending.back());
+		pending.pop_back();
+		std::filesystem::directory_iterator it(
+		    current, std::filesystem::directory_options::skip_permission_denied, ec);
+		const std::filesystem::directory_iterator end;
+		for (; !ec && it != end; it.increment(ec)) {
+			std::error_code entry_ec;
+			if (it->is_directory(entry_ec) && !it->is_symlink(entry_ec) && !entry_ec) {
+				pending.push_back(it->path());
+			} else if (it->is_regular_file(entry_ec) && !entry_ec &&
+			           EqualNoCase(PathToString(it->path().filename()), "param.json") &&
+			           EqualNoCase(PathToString(it->path().parent_path().filename()), "sce_sys")) {
+				params.push_back(it->path());
+			}
+		}
+	}
+	std::sort(params.begin(), params.end());
+
+	static const std::regex content_id_pattern(
+	    "^[A-Z]{2}[0-9]{4}-[A-Z]{4}[0-9]{5}_[0-9]{2}-[A-Z0-9]{16}$");
+	std::map<std::string, AddContentEntry> by_label;
+	for (const auto& file: params) {
+		const auto size = std::filesystem::file_size(file, ec);
+		if (ec || size > 1024 * 1024) {
+			continue;
+		}
+		std::ifstream stream(file, std::ios::binary);
+		if (!stream) {
+			continue;
+		}
+		const auto param = nlohmann::json::parse(stream, nullptr, false);
+		// applicationCategoryType is required for applications, not additional content.
+		if (!param.is_object() || param.contains("applicationCategoryType")) {
+			continue;
+		}
+		const auto content_id = GetString(param, "contentId");
+		if (!std::regex_match(content_id, content_id_pattern) ||
+		    content_id.compare(7, 9, title_id) != 0) {
+			continue;
+		}
+
+		const auto mount_dir = file.parent_path().parent_path();
+		const auto label     = content_id.substr(20);
+		by_label.try_emplace(label,
+		                     AddContentEntry {content_id, label, mount_dir, GetTitle(param)});
+	}
+
+	std::vector<AddContentEntry> entries;
+	entries.reserve(by_label.size());
+	for (auto& [label, entry]: by_label) {
+		entries.push_back(std::move(entry));
+	}
+	return entries;
+}
+
+} // namespace Common
@@ -0,0 +1,23 @@
+#ifndef KYTY_COMMON_ADDITIONAL_CONTENT_H_
+#define KYTY_COMMON_ADDITIONAL_CONTENT_H_
+
+#include <filesystem>
+#include <string>
+#include <vector>
+
+namespace Common {
+
+struct AddContentEntry {
+	std::string           content_id;
+	std::string           entitlement_label;
+	std::filesystem::path mount_dir;
+	std::string           title;
+};
+
+// Finds packages for this title, deduplicated by entitlement label.
+std::vector<AddContentEntry> FindAddContent(const std::filesystem::path& dir,
+                                            const std::string&           title_id);
+
+} // namespace Common
+
+#endif /* KYTY_COMMON_ADDITIONAL_CONTENT_H_ */
@@ -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;
 }
@@ -172,6 +201,10 @@ bool TrophyEnabled() {
 	return g_config->trophy_enabled;
 }

+uint32_t GetTrophyDurationSeconds() {
+	return std::clamp(g_config->trophy_duration_seconds, 5u, 10u);
+}
+
 bool PlayGoHackEnabled() {
 	return g_config->playgo_hack_enabled;
 }
@@ -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,9 +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
@@ -99,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();
@@ -132,9 +151,11 @@ bool GpuAssistedValidationEnabled();

 bool RenderDocEnabled();
 bool ReadbackLinearImagesEnabled();
+bool ReadbackTiledImagesEnabled();
 bool SyncRawImageBuffersEnabled();
 bool TessellationEnabled();
 bool TrophyEnabled();
+uint32_t GetTrophyDurationSeconds();
 bool PlayGoHackEnabled();
 bool SkipNoticeScreen();
 #if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
@@ -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)

@@ -152,42 +152,29 @@ void SysFileWrite(const void* data, uint32_t size, sys_file_t& f, uint32_t* byte
 	}
 }

-sys_file_t* SysFileCreate(const std::filesystem::path& file_name) {
+static sys_file_t* OpenHostFile(const std::filesystem::path& file_name, const char* mode,
+                               sys_file_cache_type_t cache_type) {
 	auto* ret = new sys_file_t;
-
-	auto real_name     = get_internal_name(file_name);
-	auto real_name_str = real_name.string();
-
-	ret->f = fopen(real_name_str.c_str(), "w+");
-
-	if (ret->f == nullptr) {
-		printf("can't create file: %s\n", real_name_str.c_str());
+	ret->f = fopen(get_internal_name(file_name).c_str(), mode);
+	ret->type = ret->f == nullptr ? SYS_FILE_ERROR : SYS_FILE_FILE;
+	if (ret->f != nullptr) {
+		// Independent handles must observe writes and truncation immediately.
+		setvbuf(ret->f, nullptr, _IONBF, 0);
+		apply_cache_hint(ret->f, cache_type);
 	}
+	return ret;
+}

-	ret->type = SYS_FILE_FILE;
-
+sys_file_t* SysFileCreate(const std::filesystem::path& file_name) {
+	auto* ret = OpenHostFile(file_name, "w+", SYS_FILE_CACHE_AUTO);
+	if (ret->f == nullptr) {
+		printf("can't create file: %s\n", file_name.c_str());
+	}
 	return ret;
 }

 sys_file_t* SysFileOpenR(const std::filesystem::path& file_name, sys_file_cache_type_t cache_type) {
-	auto* ret = new sys_file_t;
-
-	ret->type = SYS_FILE_FILE;
-
-	auto internal_name     = get_internal_name(file_name);
-	auto internal_name_str = internal_name.string();
-
-	FILE* f = fopen(internal_name_str.c_str(), "r");
-
-	if (f == nullptr) {
-		ret->type = SYS_FILE_ERROR;
-	}
-
-	apply_cache_hint(f, cache_type);
-
-	ret->f = f;
-
-	return ret;
+	return OpenHostFile(file_name, "r", cache_type);
 }

 sys_file_t* SysFileOpen(uint8_t* buf, uint32_t buf_size) {
@@ -215,46 +202,12 @@ sys_file_t* SysFileCreate() {
 }

 sys_file_t* SysFileOpenW(const std::filesystem::path& file_name, sys_file_cache_type_t cache_type) {
-	auto* ret = new sys_file_t;
-
-	auto real_name     = get_internal_name(file_name);
-	auto real_name_str = real_name.string();
-
-	FILE* f = fopen(real_name_str.c_str(), "r+");
-
-	if (f == nullptr) {
-		ret->type = SYS_FILE_ERROR;
-	} else {
-		ret->type = SYS_FILE_FILE;
-	}
-
-	apply_cache_hint(f, cache_type);
-
-	ret->f = f;
-
-	return ret;
+	return OpenHostFile(file_name, "r+", cache_type);
 }

 sys_file_t* SysFileOpenRw(const std::filesystem::path& file_name,
                          sys_file_cache_type_t        cache_type) {
-	auto* ret = new sys_file_t;
-
-	auto real_name     = get_internal_name(file_name);
-	auto real_name_str = real_name.string();
-
-	FILE* f = fopen(real_name_str.c_str(), "r+");
-
-	if (f == nullptr) {
-		ret->type = SYS_FILE_ERROR;
-	} else {
-		ret->type = SYS_FILE_FILE;
-	}
-
-	apply_cache_hint(f, cache_type);
-
-	ret->f = f;
-
-	return ret;
+	return OpenHostFile(file_name, "r+", cache_type);
 }

 void SysFileClose(sys_file_t* f) {
@@ -414,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;
@@ -8,10 +8,6 @@
 #include <condition_variable> // IWYU pragma: keep
 #include <mutex>

-#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS && KYTY_COMPILER == KYTY_COMPILER_CLANG
-#define KYTY_WIN_CS
-#endif
-
 // macOS has no clock_nanosleep.
 #if KYTY_PLATFORM != KYTY_PLATFORM_WINDOWS && !defined(__APPLE__)
 #define KYTY_POSIX_HIGH_RES_SLEEP
@@ -22,7 +18,7 @@
 #include <string>
 #include <thread>

-#ifdef KYTY_WIN_CS
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
 #include <windows.h> // IWYU pragma: keep
 // IWYU pragma: no_include <winbase.h>
 constexpr DWORD KYTY_CS_SPIN_COUNT = 4000;
@@ -133,7 +129,7 @@ namespace Common {
 using thread_id_t = std::thread::id;

 struct MutexPrivate {
-#ifdef KYTY_WIN_CS
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
 	MutexPrivate() { InitializeCriticalSectionAndSpinCount(&m_cs, KYTY_CS_SPIN_COUNT); }
 	~MutexPrivate() { DeleteCriticalSection(&m_cs); }
 	KYTY_CLASS_NO_COPY(MutexPrivate);
@@ -144,7 +140,7 @@ struct MutexPrivate {
 };

 struct CondVarPrivate {
-#ifdef KYTY_WIN_CS
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
 	CondVarPrivate() {
 		static auto func = ResolveInitializeConditionVariable();
 		EXIT_NOT_IMPLEMENTED(func == nullptr);
@@ -216,7 +212,7 @@ void Thread::Detach() {
 }

 void Thread::SleepMicro(uint32_t micros) {
-#ifdef KYTY_WIN_CS
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
 	SleepHighResolution100ns(static_cast<uint64_t>(micros) * 10);
 #elif defined(KYTY_POSIX_HIGH_RES_SLEEP)
 	SleepHighResolutionNanos(static_cast<uint64_t>(micros) * 1000);
@@ -226,7 +222,7 @@ void Thread::SleepMicro(uint32_t micros) {
 }

 void Thread::SleepNano(uint64_t nanos) {
-#ifdef KYTY_WIN_CS
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
 	SleepHighResolution100ns((nanos + 99) / 100);
 #elif defined(KYTY_POSIX_HIGH_RES_SLEEP)
 	SleepHighResolutionNanos(nanos);
@@ -262,7 +258,7 @@ Mutex::~Mutex() {
 }

 void Mutex::Lock() {
-#ifdef KYTY_WIN_CS
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
 	EnterCriticalSection(&m_mutex->m_cs);
 #else
 	m_mutex->m_mutex.lock();
@@ -270,7 +266,7 @@ void Mutex::Lock() {
 }

 void Mutex::Unlock() {
-#ifdef KYTY_WIN_CS
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
 	LeaveCriticalSection(&m_mutex->m_cs);
 #else
 	m_mutex->m_mutex.unlock();
@@ -278,7 +274,7 @@ void Mutex::Unlock() {
 }

 bool Mutex::TryLock() {
-#ifdef KYTY_WIN_CS
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
 	return (TryEnterCriticalSection(&m_mutex->m_cs) != 0);
 #else
 	return m_mutex->m_mutex.try_lock();
@@ -292,7 +288,7 @@ CondVar::~CondVar() {
 }

 void CondVar::Wait(Mutex* mutex) {
-#ifndef KYTY_WIN_CS
+#if KYTY_PLATFORM != KYTY_PLATFORM_WINDOWS
 	std::unique_lock<std::recursive_mutex> cpp_lock(mutex->m_mutex->m_mutex, std::adopt_lock_t());
 #endif
 	auto poll_callback = [&] {
@@ -300,7 +296,7 @@ void CondVar::Wait(Mutex* mutex) {
 		if (callback == nullptr) {
 			return;
 		}
-#if defined(KYTY_WIN_CS)
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
 		LeaveCriticalSection(&mutex->m_mutex->m_cs);
 		callback();
 		EnterCriticalSection(&mutex->m_mutex->m_cs);
@@ -310,7 +306,7 @@ void CondVar::Wait(Mutex* mutex) {
 		cpp_lock.lock();
 #endif
 	};
-#ifdef KYTY_WIN_CS
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
 	static auto func = ResolveSleepConditionVariableCS();
 	EXIT_NOT_IMPLEMENTED(func == nullptr);
 	if (g_cond_wait_poll_callback == nullptr) {
@@ -340,10 +336,10 @@ void CondVar::SetWaitPollCallback(wait_poll_func_t callback) {

 bool CondVar::WaitFor(Mutex* mutex, uint32_t micros) {
 	bool ok = false;
-#ifndef KYTY_WIN_CS
+#if KYTY_PLATFORM != KYTY_PLATFORM_WINDOWS
 	std::unique_lock<std::recursive_mutex> cpp_lock(mutex->m_mutex->m_mutex, std::adopt_lock_t());
 #endif
-#ifdef KYTY_WIN_CS
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
 	static auto func = ResolveSleepConditionVariableCS();
 	EXIT_NOT_IMPLEMENTED(func == nullptr);
 	ok = !(func(&m_cond_var->m_cv, &mutex->m_mutex->m_cs, (micros < 1000 ? 1 : micros / 1000)) ==
@@ -358,7 +354,7 @@ bool CondVar::WaitFor(Mutex* mutex, uint32_t micros) {
 }

 void CondVar::Signal() {
-#ifdef KYTY_WIN_CS
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
 	static auto func = ResolveWakeConditionVariable();
 	EXIT_NOT_IMPLEMENTED(func == nullptr);
 	func(&m_cond_var->m_cv);
@@ -368,7 +364,7 @@ void CondVar::Signal() {
 }

 void CondVar::SignalAll() {
-#ifdef KYTY_WIN_CS
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
 	static auto func = ResolveWakeAllConditionVariable();
 	EXIT_NOT_IMPLEMENTED(func == nullptr);
 	func(&m_cond_var->m_cv);
@@ -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")
@@ -22,6 +22,7 @@
 #include "libs/controller.h"
 #include "libs/libs.h"
 #include "libs/network.h"
+#include "loader/additionalContent.h"
 #include "loader/runtimeLinker.h"
 #include "loader/systemContent.h"
 #include "loader/timer.h"
@@ -88,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());
@@ -197,11 +162,17 @@ void Run(const RunOptions& options) {
 	Common::Subsystems subsystems(true);
 	Init(options.config, param_json, subsystems);

-	ClearDebugTextureFolder();
-
 	PrintSystemInfo();
 	std::string title_id;
-	if (Loader::SystemContentParamSfoGetString("TITLE_ID", &title_id) && !title_id.empty()) {
+	Loader::SystemContentParamSfoGetString("TITLE_ID", &title_id);
+	if (title_id.empty()) {
+		std::string content_id;
+		if (Loader::SystemContentParamSfoGetString("CONTENT_ID", &content_id) &&
+		    content_id.size() == 36 && content_id[6] == '-' && content_id[16] == '_') {
+			title_id = content_id.substr(7, 9);
+		}
+	}
+	if (!title_id.empty()) {
 		Log::WriteToConsoleAndLog(fmt::format("Title ID: {}\n", title_id));
 	}

@@ -217,6 +188,7 @@ void Run(const RunOptions& options) {

 	MountSandboxDirs();

+	Loader::AddContentLoad(options.addcont_dir, title_id);
 	auto* rt = Common::Singleton<Loader::RuntimeLinker>::Instance();
 	Libs::InitAll(rt->Symbols());

@@ -13,6 +13,7 @@ struct RunOptions {
 	std::filesystem::path app0_dir;
 	std::filesystem::path elf;
 	std::filesystem::path game_patch;
+	std::filesystem::path addcont_dir;
 };

 void Run(const RunOptions& options);
@@ -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;
@@ -231,7 +230,6 @@ constexpr auto MakeCustomOpcodeDispatchTable() {
 	g_cp_op_custom_func[Pm4::R_WAIT_FLIP_DONE] = CpOpWaitFlipDone;
 	g_cp_op_custom_func[Pm4::R_PUSH_MARKER]    = CpOpPushMarker;
 	g_cp_op_custom_func[Pm4::R_POP_MARKER]     = CpOpPopMarker;
-	g_cp_op_custom_func[Pm4::R_ACQUIRE_MEM]    = CpOpAcquireMem;
 	g_cp_op_custom_func[Pm4::R_FLIP]           = CpOpFlip;
 	g_cp_op_custom_func[Pm4::R_RELEASE_MEM]    = CpOpReleaseMem;
 	g_cp_op_custom_func[Pm4::R_CONTEXT_STATE]  = CpOpContextState;
@@ -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,16 +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 HwCtxIgnoreDrawPayloadControl([[maybe_unused]] uint32_t value) {}
-
-static void HwCtxIgnoreObjprimIdControl([[maybe_unused]] uint32_t value) {}
-
 static void HwCtxIgnorePrimitiveIdReset([[maybe_unused]] uint32_t value) {}

 static void HwCtxIgnoreVtxControl([[maybe_unused]] uint32_t value) {}
@@ -443,34 +427,14 @@ 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) {
-	auto num_values = KYTY_PM4_LEN(cmd_id) - 2u;
-
-	for (uint32_t i = 0; i < num_values; i++) {
-		HwCtxIgnoreDrawPayloadControl(buffer[i]);
-	}
-
-	return num_values;
+	cp.GetCtx().SetDrawPayloadControl(buffer[0]);
+	return KYTY_PM4_LEN(cmd_id) - 2u;
 }

 KYTY_HW_CTX_PARSER(HwCtxSetObjprimIdControl) {
-	auto num_values = KYTY_PM4_LEN(cmd_id) - 2u;
-
-	for (uint32_t i = 0; i < num_values; i++) {
-		HwCtxIgnoreObjprimIdControl(buffer[i]);
-	}
-
-	return num_values;
+	cp.GetCtx().SetObjprimIdControl(buffer[0]);
+	return KYTY_PM4_LEN(cmd_id) - 2u;
 }

 KYTY_HW_CTX_PARSER(HwCtxSetPrimitiveIdReset) {
@@ -1304,8 +1268,8 @@ KYTY_HW_UC_PARSER(HwUcSetGdsOaRegisters) {
 KYTY_CP_OP_PARSER(CpOpAcquireMem) {
 	KYTY_PROFILER_FUNCTION();

-	EXIT_NOT_IMPLEMENTED(cmd_id != 0xC0055800 && cmd_id != 0xc0061050);
-	return (cmd_id == 0xc0061050 ? 7 : 6);
+	EXIT_NOT_IMPLEMENTED(cmd_id != 0xc0055800 && cmd_id != 0xc0065800);
+	return KYTY_PM4_LEN(cmd_id) - 1;
 }

 KYTY_CP_OP_PARSER(CpOpDispatchDirect) {
@@ -2142,27 +2106,23 @@ KYTY_CP_OP_PARSER(CpOpIndirectUcRegs) {
 KYTY_CP_OP_PARSER(CpOpMarker) {
 	KYTY_PROFILER_FUNCTION();

-	// EXIT_NOT_IMPLEMENTED(cmd_id != 0xC0001000);
-
-	uint32_t id     = buffer[0] & 0xfffu;
 	uint32_t len_dw = ((cmd_id >> 16u) & 0x3fffu);

-	switch (id) {
-		case 0x0: break;
-		case 0x4: cp.SetUserDataMarker(HW::UserSgprType::Vsharp); break;
-		case 0xd: cp.SetUserDataMarker(HW::UserSgprType::Region); break;
-		case 0x777: {
+	switch (buffer[0]) {
+		case 0x68750004: cp.SetUserDataMarker(HW::UserSgprType::Vsharp); break;
+		case 0x6875000d: cp.SetUserDataMarker(HW::UserSgprType::Region); break;
+		case 0x68750777: {
 			cp.Flip();
 			break;
 		}
-		case 0x778: {
+		case 0x68750778: {
 			auto* addr =
 			    reinterpret_cast<void*>(buffer[1] | (static_cast<uint64_t>(buffer[2]) << 32u));
 			uint32_t value = buffer[3];
 			cp.Flip(addr, value);
 			break;
 		}
-		case 0x781: {
+		case 0x68750781: {
 			auto* addr =
 			    reinterpret_cast<void*>(buffer[1] | (static_cast<uint64_t>(buffer[2]) << 32u));
 			uint32_t value          = buffer[3];
@@ -2171,7 +2131,8 @@ KYTY_CP_OP_PARSER(CpOpMarker) {
 			cp.FlipWithInterrupt(eop_event_type, cache_action, addr, value);
 			break;
 		}
-		default: EXIT("unknown marker at %05" PRIx32 ": 0x%" PRIx32 "\n", num_dw - dw, id);
+		// NOP payloads are caller-owned data; only complete known marker tags execute.
+		default: break;
 	}

 	return len_dw + 1;
@@ -2280,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;

@@ -2916,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);
 		};
 	}

@@ -3109,7 +3054,7 @@ void GraphicsInitJmpTablesCxIndirect() {
 		cp.GetCtx().SetTfParam(value);
 	};
 	g_hw_ctx_indirect_func[Pm4::VGT_DRAW_PAYLOAD_CNTL] = [](KYTY_HW_CTX_INDIRECT_ARGS) {
-		HwCtxIgnoreDrawPayloadControl(value);
+		cp.GetCtx().SetDrawPayloadControl(value);
 	};
 	g_hw_ctx_indirect_func[Pm4::VGT_PRIMITIVEID_RESET] = [](KYTY_HW_CTX_INDIRECT_ARGS) {
 		HwCtxIgnorePrimitiveIdReset(value);
@@ -3118,7 +3063,7 @@ void GraphicsInitJmpTablesCxIndirect() {
 		cp.GetCtx().SetShaderZFormat(value);
 	};
 	g_hw_ctx_indirect_func[Pm4::PA_CL_OBJPRIM_ID_CNTL] = [](KYTY_HW_CTX_INDIRECT_ARGS) {
-		HwCtxIgnoreObjprimIdControl(value);
+		cp.GetCtx().SetObjprimIdControl(value);
 	};

 	for (auto cmd_offset = Pm4::PA_SC_CLIPRECT_RULE; cmd_offset < Pm4::PA_SC_CLIPRECT_0_TL + 8u;
@@ -3508,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));
@@ -23,6 +23,7 @@ constexpr FormatInfo kFormatInfo[] = {
 	{BufferFormat::k8UScaled, 1, 0, 0, true, false},
 	{BufferFormat::k8SNorm, 1, 0, 1, true, false},
 	{BufferFormat::k8UInt, 1, 0, 1, true, true},
+	{BufferFormat::k8SInt, 1, 0, 1, true, false, true},
 	{BufferFormat::k16UNorm, 2, 0, 2, true, false},
 	{BufferFormat::k16SNorm, 2, 0, 2, true, false},
 	{BufferFormat::k16UInt, 2, 0, 2, true, true},
@@ -41,6 +42,7 @@ constexpr FormatInfo kFormatInfo[] = {
 	{BufferFormat::k16_16UInt, 4, 0, 4, true, true},
 	{BufferFormat::k16_16SInt, 4, 0, 4, true, false, true},
 	{BufferFormat::k16_16Float, 4, 0, 4, true, false},
+	{BufferFormat::k11_11_10UNorm, 4, 0, 4, true, false},
 	{BufferFormat::k11_11_10UInt, 4, 0, 4, true, true},
 	{BufferFormat::k11_11_10Float, 4, 0, 4, true, false},
 	{BufferFormat::k10_11_11Float, 4, 0, 4, true, false},
@@ -237,6 +239,7 @@ BufferFormat RemapTextureFormat(BufferFormat format) {
 	switch (format) {
 		case BufferFormat::k8UScaled: return BufferFormat::k8UNorm;
 		case BufferFormat::k8_8UScaled: return BufferFormat::k8_8UNorm;
+		case BufferFormat::k11_11_10UNorm:
 		case BufferFormat::k11_11_10UInt: return BufferFormat::k32UInt;
 		default: return format;
 	}
@@ -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"
@@ -38,21 +39,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;
@@ -205,6 +191,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 +329,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;
 	}
 }

@@ -607,7 +593,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 +770,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 +836,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 +849,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 +874,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 +885,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,7 +992,7 @@ 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);
 }
@@ -1221,22 +1201,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 +1220,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 +1238,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:
@@ -5,6 +5,8 @@
 #include "common/common.h"
 #include "graphics/guest_gpu/gpu_defs.h"

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

 struct ColorBase {
@@ -556,6 +558,8 @@ struct ShaderRegisters {
 	uint32_t m_vgtGsMaxVertOut        = 0;
 	uint32_t m_vgtGsOutPrimType       = 0;
 	uint32_t m_vgtPrimitiveIdEn       = 0;
+	uint32_t m_vgtDrawPayloadCntl     = 0;
+	uint32_t m_paClObjprimIdCntl      = 0;
 	uint32_t m_vgtReuseOff            = 0;
 	uint32_t m_vgtTessDistribution    = 0;
 	uint32_t m_vgtLsHsConfig          = 0;
@@ -572,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);
@@ -679,6 +689,10 @@ public:

 	void                   SetRenderTargetMask(uint32_t mask) { m_render_target_mask = mask; }
 	[[nodiscard]] uint32_t GetRenderTargetMask() const { return m_render_target_mask; }
+	[[nodiscard]] bool HasObjectPrimitiveIdOutput() const {
+		return (m_sh_regs.m_vgtDrawPayloadCntl & 1u) != 0 &&
+		       (m_render_target_mask & 0xf0000000u) != 0;
+	}

 	void                   SetShaderStages(uint32_t flags) { m_shader_stages = flags; }
 	[[nodiscard]] uint32_t GetShaderStages() const { return m_shader_stages; }
@@ -818,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 {
@@ -877,6 +896,8 @@ public:
 	void SetGsMaxVertOut(uint32_t value) { m_sh_regs.m_vgtGsMaxVertOut = value; }
 	void SetGsOutPrimType(uint32_t value) { m_sh_regs.m_vgtGsOutPrimType = value; }
 	void SetPrimitiveIdEn(uint32_t value) { m_sh_regs.m_vgtPrimitiveIdEn = value; }
+	void SetDrawPayloadControl(uint32_t value) { m_sh_regs.m_vgtDrawPayloadCntl = value; }
+	void SetObjprimIdControl(uint32_t value) { m_sh_regs.m_paClObjprimIdCntl = value; }
 	void SetReuseOff(uint32_t value) { m_sh_regs.m_vgtReuseOff = value; }
 	void SetTessDistribution(uint32_t value) { m_sh_regs.m_vgtTessDistribution = value; }
 	void SetLsHsConfig(uint32_t value) { m_sh_regs.m_vgtLsHsConfig = value; }
@@ -900,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; }

@@ -916,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;

@@ -23,7 +23,6 @@ constexpr auto MakeRegisterNames() {
 	names[R_DISPATCH_RESET] = "R_DISPATCH_RESET";
 	names[R_PUSH_MARKER]    = "R_PUSH_MARKER";
 	names[R_POP_MARKER]     = "R_POP_MARKER";
-	names[R_ACQUIRE_MEM]    = "R_ACQUIRE_MEM";
 	names[R_WRITE_DATA]     = "R_WRITE_DATA";
 	names[R_FLIP]           = "R_FLIP";
 	names[R_RELEASE_MEM]    = "R_RELEASE_MEM";
@@ -86,7 +86,6 @@ constexpr uint32_t R_WAIT_FLIP_DONE = 0x06;
 constexpr uint32_t R_DISPATCH_RESET = 0x09;
 constexpr uint32_t R_PUSH_MARKER    = 0x0B;
 constexpr uint32_t R_POP_MARKER     = 0x0C;
-constexpr uint32_t R_ACQUIRE_MEM    = 0x14;
 constexpr uint32_t R_WRITE_DATA     = 0x15;
 constexpr uint32_t R_FLIP           = 0x17;
 constexpr uint32_t R_RELEASE_MEM    = 0x18;
@@ -550,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>
@@ -33,7 +34,8 @@ struct GraphicContext {
 	// bool fp64_denorm_preserve = false; // Temporarily disabled.
 	bool                               attachment_feedback_loop_enabled      = false;
 	bool                               provoking_vertex_last_enabled         = false;
-	bool                               depth_color_copy_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                = {};
@@ -99,6 +101,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"
@@ -27,6 +28,26 @@ namespace {
 constexpr uint64_t MiB           = 1024 * 1024;
 constexpr uint64_t GdsBufferSize = 64 * 1024;

+std::pair<Buffer*, uint64_t> UploadMemory(GraphicContext& graphics, CommandScheduler& scheduler,
+                                        StreamBuffer& staging, uint64_t size, uint64_t alignment,
+                                        vk::BufferUsageFlags usage, const auto& write) {
+	const auto [mapped, offset] = staging.Map(size, alignment);
+	if (mapped != nullptr) {
+		write(mapped, offset);
+		staging.Commit();
+		return {&staging, offset};
+	}
+
+	// Large transfers use the same one-time host allocation as ordinary buffer uploads.
+	auto temporary = std::make_unique<Buffer>(graphics, scheduler, MemoryUsage::Upload, 0,
+	                                         usage, size);
+	write(temporary->Mapped().data(), 0);
+	temporary->Flush(0, size);
+	auto* buffer = temporary.get();
+	scheduler.DeferOperation([owner = std::move(temporary)]() mutable { owner.reset(); });
+	return {buffer, 0};
+}
+
 } // namespace

 void BufferCache::WriteDataBuffer(Buffer& buffer, uint64_t address, const void* source,
@@ -106,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);
@@ -152,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;
@@ -233,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);
@@ -446,29 +474,19 @@ 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);

-	auto [mapped, base_offset] = m_staging_buffer.Map(total_size, 4);
-	if (mapped != nullptr) {
+	const auto source = UploadMemory(m_graphics, m_scheduler, m_staging_buffer, total_size, 4,
+	    vk::BufferUsageFlagBits::eTransferSrc, [&](uint8_t* mapped, uint64_t base_offset) {
 		for (auto& copy: copies) {
 			const auto address = buffer.CpuAddress() + copy.dstOffset;
-			std::memcpy(mapped + copy.srcOffset, reinterpret_cast<const void*>(address), copy.size);
+			if (!LibKernel::Memory::CopySparseMemory(address, mapped + copy.srcOffset, copy.size)) {
+				EXIT("BufferCache: failed to read sparse guest buffer memory\n");
+			}
 			copy.srcOffset += base_offset;
 		}
-		m_staging_buffer.Commit();
-		return m_staging_buffer.Handle();
-	}
-
-	auto temporary = std::make_unique<Buffer>(m_graphics, m_scheduler, MemoryUsage::Upload, 0,
-	                                         vk::BufferUsageFlagBits::eTransferSrc, total_size);
-	for (const auto& copy: copies) {
-		const auto address = buffer.CpuAddress() + copy.dstOffset;
-		std::memcpy(temporary->Mapped().data() + copy.srcOffset,
-		            reinterpret_cast<const void*>(address), copy.size);
-	}
-	temporary->Flush(0, total_size);
-	const auto handle = temporary->Handle();
-	m_scheduler.DeferOperation([owner = std::move(temporary)]() mutable { owner.reset(); });
-	return handle;
+	});
+	return source.first->Handle();
 }

 std::pair<Buffer*, uint64_t> BufferCache::ObtainBuffer(uint64_t vaddr, uint64_t size,
@@ -486,7 +504,9 @@ std::pair<Buffer*, uint64_t> BufferCache::ObtainBuffer(uint64_t vaddr, uint64_t
 		    m_graphics.physical_device_properties.limits.minUniformBufferOffsetAlignment, 1);
 		auto [mapped, offset] = m_stream_buffer.Map(size, alignment, false);
 		if (mapped != nullptr) {
-			std::memcpy(mapped, reinterpret_cast<const void*>(vaddr), size);
+			if (!LibKernel::Memory::CopySparseMemory(vaddr, mapped, size)) {
+				EXIT("BufferCache: failed to read sparse guest buffer memory\n");
+			}
 			m_stream_buffer.Commit();
 			return {&m_stream_buffer, offset};
 		}
@@ -500,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)};
 }
@@ -521,12 +542,13 @@ std::pair<Buffer*, uint64_t> BufferCache::ObtainBufferForImage(uint64_t vaddr, u
 		return ObtainBuffer(vaddr, size, false, false);
 	}

-	auto [staging, stage_offset] = m_staging_buffer.Map(size, 16);
-	if (staging == nullptr || !Libs::LibKernel::Memory::TryReadSparseBacking(vaddr, staging, size)) {
-		EXIT("BufferCache: failed to read mapped guest image backing\n");
-	}
-	m_staging_buffer.Commit();
-	return {&m_staging_buffer, stage_offset};
+	return UploadMemory(m_graphics, m_scheduler, m_staging_buffer, size, 16,
+	    vk::BufferUsageFlagBits::eTransferSrc | vk::BufferUsageFlagBits::eStorageBuffer,
+	    [&](uint8_t* mapped, uint64_t) {
+		if (!Libs::LibKernel::Memory::TryReadSparseBacking(vaddr, mapped, size)) {
+			EXIT("BufferCache: failed to read mapped guest image backing\n");
+		}
+	});
 }

 void BufferCache::FillBuffer(uint64_t vaddr, uint64_t size, uint32_t value, bool is_gds) {
@@ -603,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);
 }
@@ -645,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;
 	});
@@ -673,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"
@@ -25,6 +26,7 @@
 #include <mutex>
 #include <span>
 #include <tuple>
+#include <utility>
 #include <vulkan/vulkan_format_traits.hpp>

 namespace Libs::Graphics {
@@ -32,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) {
@@ -120,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 =
@@ -198,9 +203,7 @@ TextureCache::BindingType TextureCache::UploadBinding(const Image& image) {

 ImageId TextureCache::InsertImage(const ImageInfo& info) {
 	const auto id = m_slot_images.insert(m_graphics, m_scheduler, info);
-	if (!info.data.Empty()) {
-		RegisterImage(id);
-	}
+	RegisterImage(id);
 	return id;
 }

@@ -251,14 +254,11 @@ void TextureCache::DeleteImage(ImageId id) {
 	if (image == nullptr || !image->registered) {
 		return;
 	}
-	if (!image->depth_id) {
-		std::vector<ImageId> associations;
-		m_slot_images.ForEach([&](ImageId candidate, const Image& associated) {
-			if (associated.depth_id == id) {
-				associations.push_back(candidate);
-			}
-		});
-		for (const auto association: associations) {
+	if (image->depth_id) {
+		std::erase(m_slot_images[image->depth_id].stencil_ids, id);
+	} else {
+		// Detach first: deleting each association also unlinks it from its depth owner.
+		for (const auto association: std::exchange(image->stencil_ids, {})) {
 			FreeImage(association);
 		}
 	}
@@ -266,14 +266,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()) {
@@ -299,9 +293,6 @@ void TextureCache::TouchImage(Image& image) {

 void TextureCache::MarkAsMaybeDirty(ImageId id, Image& image) {
 	image.MarkMaybeCpuDirty();
-	if (image.NeedsMaybeCpuHash()) {
-		image.SetMaybeCpuHash(image.HashGuestEdges());
-	}
 	UntrackImage(id);
 }

@@ -413,7 +404,14 @@ void TextureCache::UntrackImageTail(ImageId id) {
 }

 void TextureCache::TrackImageDownload(ImageId id, Image& image) {
-	if (m_readback_linear_images && !image.info.IsTiled() && !image.info.data.Empty()) {
+	// 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");
 		}
@@ -457,27 +455,6 @@ TextureCache::ImageIds TextureCache::FindImagesInRegion(uint64_t address, uint64
 	return result;
 }

-ImageId TextureCache::GetNullImage(const ImageDesc& desc) {
-	const auto format = desc.info.pixel_format;
-	if (const auto found = m_null_images.find(format); found != m_null_images.end()) {
-		return found->second;
-	}
-	ImageInfo info {};
-	info.pixel_format    = desc.info.pixel_format;
-	info.guest_format    = desc.info.guest_format;
-	info.type            = Prospero::ImageType::kColor2D;
-	info.extent          = {1, 1, 1};
-	info.resources       = {1, 1};
-	info.pitch           = 1;
-	info.bytes_per_block = std::max(desc.info.bytes_per_block, 1u);
-	info.samples         = 1;
-	info.tile_mode       = Prospero::TileMode::kLinear;
-	info.mip_layout[0]   = {0, info.bytes_per_block, 1, 1};
-	const auto id        = InsertImage(info);
-	m_null_images.emplace(format, id);
-	return id;
-}
-
 void TextureCache::ValidateImageDesc(const ImageDesc& desc) const {
 	ImageOps::Validate(desc.info);
 	if (desc.view_info.format == vk::Format::eUndefined || desc.view_info.level_count == 0 ||
@@ -720,7 +697,7 @@ ImageId TextureCache::ResolveDepthOverlap(const ImageInfo& requested, BindingTyp
 		info.linear_slice_stride = cached.info.linear_slice_stride;
 		info.first_tail_level    = cached.info.first_tail_level;
 	}
-	info.htile_clear_mask     = 0;
+	info.htile_initial_clear  = false;
 	const auto replacement_id = InsertImage(info);
 	auto&      replacement    = m_slot_images[replacement_id];
 	replacement.usage         = cached.usage;
@@ -982,6 +959,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;
 	}

@@ -1027,13 +1005,11 @@ 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) {
 	auto& image = m_slot_images[id];
-	if (image.info.data.Empty()) {
-		return;
-	}
 	TrackImage(id);
 	if (image.info.metadata.compression != VideoOutCompression::Uncompressed) {
 		if (image.IsCpuDirty()) {
@@ -1070,7 +1046,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;
@@ -1091,12 +1067,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;
@@ -1146,12 +1188,7 @@ void TextureCache::RefreshImage(ImageId id) {
 	}
 	TrackImage(id);
 	if (image.IsMaybeCpuDirty()) {
-		const auto hash = image.HashGuestEdges();
-		if (image.NeedsMaybeCpuHash()) {
-			image.SetMaybeCpuHash(hash);
-			return;
-		}
-		(void)image.ResolveMaybeCpuHash(hash);
+		(void)image.ResolveMaybeCpuHash(image.HashGuestEdges());
 	}
 	bool cpu_dirty = image.IsBufferModified() || image.IsDefinitelyCpuDirty();
 	if (image.info.metadata.compression != VideoOutCompression::Uncompressed) {
@@ -1205,6 +1242,12 @@ ImageId TextureCache::AssociateStencil(ImageId depth_id, GuestRange stencil) {
 		    plane.IsTiled() ? TileGetDepthPitch(plane.extent.width, 1, 0) : depth.info.pitch;
 		plane.bytes_per_block = 1;
 		plane.UpdateSize();
+		if (record.depth_id != depth_id) {
+			if (record.depth_id) {
+				std::erase(m_slot_images[record.depth_id].stencil_ids, association);
+			}
+			depth.stencil_ids.push_back(association);
+		}
 		record.info                 = plane;
 		record.depth_id             = depth_id;
 		record.stencil_subresources = depth.stencil_subresources;
@@ -1219,10 +1262,7 @@ ImageId TextureCache::FindImage(ImageDesc& desc, bool exact_format) {
 		EXIT("TextureCache: image lookup requires a valid command buffer\n");
 	}
 	ValidateImageDesc(desc);
-	if (desc.info.data.Empty()) {
-		std::scoped_lock lock {m_lock};
-		return GetNullImage(desc);
-	}
+	EXIT_IF(desc.info.data.Empty());
 	const auto metadata_base_layer = desc.view_info.base_layer;

 	ImageId result {};
@@ -1374,36 +1414,30 @@ vk::ImageView TextureCache::FindTexture(ImageId id, const ImageDesc& desc) {
 	std::scoped_lock lock {m_lock};
 	auto&            image = m_slot_images[id];
 	TouchImage(image);
-	if (!image.info.data.Empty()) {
-		if (!image.registered || image.depth_id || image.binding.needs_rebind) {
-			EXIT("TextureCache: texture requires rediscovery before final acquisition\n");
-		}
+	if (!image.registered || image.depth_id || image.binding.needs_rebind) {
+		EXIT("TextureCache: texture requires rediscovery before final acquisition\n");
 	}
 	if (desc.type == BindingType::Storage) {
 		image.MarkGpuModified();
 	}
-	if (!image.info.data.Empty()) {
-		RefreshImage(id);
-		if (image.info.HasStencil() &&
-		    desc.info.data.address >= image.info.stencil.address &&
-		    desc.info.data.End() <= image.info.stencil.End()) {
-			for (const auto stencil_id:
-			     FindImagesInRegion(image.info.stencil.address, image.info.stencil.size, false)) {
-				const auto* stencil = m_slot_images.try_get(stencil_id);
-				if (stencil != nullptr && stencil->depth_id == id &&
-				    stencil->info.data == image.info.stencil) {
-					RefreshImage(stencil_id);
-					break;
-				}
+	RefreshImage(id);
+	if (image.info.HasStencil() &&
+	    desc.info.data.address >= image.info.stencil.address &&
+	    desc.info.data.End() <= image.info.stencil.End()) {
+		for (const auto stencil_id:
+		     FindImagesInRegion(image.info.stencil.address, image.info.stencil.size, false)) {
+			const auto* stencil = m_slot_images.try_get(stencil_id);
+			if (stencil != nullptr && stencil->depth_id == id &&
+			    stencil->info.data == image.info.stencil) {
+				RefreshImage(stencil_id);
+				break;
 			}
 		}
 	}
 	switch (desc.type) {
 		case BindingType::Texture: break;
 		case BindingType::Storage:
-			if (!image.info.data.Empty()) {
-				CommitGpuWrite(image);
-			}
+			CommitGpuWrite(image);
 			TrackImageDownload(id, image);
 			break;
 		default: EXIT("TextureCache: invalid texture binding\n");
@@ -1440,7 +1474,6 @@ vk::ImageView TextureCache::FindDepthTarget(ImageId id, const ImageDesc& desc) {
 	}
 	TouchImage(image);
 	image.MarkGpuModified();
-	image.usage.depth_target = true;
 	if (desc.info.HasStencil()) {
 		const auto layers = image.info.resources.layers;
 		EXIT_IF(layers == 0 || image.info.data.size % layers != 0 ||
@@ -1456,9 +1489,10 @@ vk::ImageView TextureCache::FindDepthTarget(ImageId id, const ImageDesc& desc) {
 	image.info.stencil = desc.info.stencil;
 	image.info.metadata = desc.info.metadata;
 	if (desc.info.HasMetadata()) {
-		m_surface_metas.emplace(desc.info.metadata.range.address,
-		                        MetaDataInfo {.type       = MetaDataInfo::Type::HTile,
-		                                      .clear_mask = image.info.htile_clear_mask});
+		const auto [metadata, inserted] = m_surface_metas.try_emplace(desc.info.metadata.range.address);
+		if (inserted && image.info.htile_initial_clear) {
+			metadata->second.clear_mask.set();
+		}
 	}
 	RefreshImage(id);
 	CommitGpuWrite(image);
@@ -1606,29 +1640,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;
@@ -1651,6 +1664,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");
@@ -1784,6 +1831,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;
 }
@@ -1810,6 +1858,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;
@@ -1856,7 +1905,7 @@ void TextureCache::InvalidateCpuAliases(uint64_t address, uint64_t size) {
 			continue;
 		}
 		if (owner->Overlaps(address, size)) {
-			owner->InvalidateCpuWrite(address, size);
+			owner->MarkCpuDirty();
 			UntrackImage(id);
 			continue;
 		}
@@ -1881,10 +1930,10 @@ bool TextureCache::IsMeta(uint64_t address) {
 bool TextureCache::IsMetaCleared(uint64_t address, uint32_t slice) {
 	std::scoped_lock lock {m_lock};
 	const auto       found = m_surface_metas.find(address);
-	if (found == m_surface_metas.end() || slice >= 32) {
+	if (found == m_surface_metas.end() || slice >= found->second.clear_mask.size()) {
 		return false;
 	}
-	return (found->second.clear_mask & (1u << slice)) != 0;
+	return found->second.clear_mask.test(slice);
 }

 bool TextureCache::ClearMeta(uint64_t address) {
@@ -1893,21 +1942,17 @@ bool TextureCache::ClearMeta(uint64_t address) {
 	if (found == m_surface_metas.end()) {
 		return false;
 	}
-	found->second.clear_mask = UINT32_MAX;
+	found->second.clear_mask.set();
 	return true;
 }

 bool TextureCache::TouchMeta(uint64_t address, uint32_t slice, bool is_clear) {
 	std::scoped_lock lock {m_lock};
 	const auto       found = m_surface_metas.find(address);
-	if (found == m_surface_metas.end() || slice >= 32) {
+	if (found == m_surface_metas.end() || slice >= found->second.clear_mask.size()) {
 		return false;
 	}
-	if (is_clear) {
-		found->second.clear_mask |= 1u << slice;
-	} else {
-		found->second.clear_mask &= ~(1u << slice);
-	}
+	found->second.clear_mask.set(slice, is_clear);
 	return true;
 }

@@ -1983,6 +2028,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;
@@ -2000,7 +2046,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};
@@ -12,9 +12,9 @@
 #include "graphics/host_gpu/renderer/image/image.h"
 #include "graphics/host_gpu/renderer/image/tiler.h"

+#include <bitset>
 #include <map>
 #include <type_traits>
-#include <unordered_map>
 #include <unordered_set>
 #include <vector>

@@ -81,10 +81,8 @@ public:

 private:
 	struct MetaDataInfo {
-		enum class Type : uint8_t { CMask, FMask, HTile };
-
-		Type     type;
-		uint32_t clear_mask = UINT32_MAX;
+		// PS5 DB_DEPTH_VIEW encodes 13-bit base and last array-slice indices.
+		std::bitset<8192> clear_mask;
 	};

 	struct OverlapResult {
@@ -114,7 +112,6 @@ private:
 	}

 	[[nodiscard]] ImageId     InsertImage(const ImageInfo& info);
-	[[nodiscard]] ImageId     GetNullImage(const ImageDesc& desc);
 	void                      RegisterImage(ImageId id);
 	void                      UnregisterImage(ImageId id);
 	void                      DeleteImage(ImageId id);
@@ -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);
@@ -173,7 +176,6 @@ private:
 	BufferCache&                                      m_buffer_cache;
 	Common::SlotVector<Image>                         m_slot_images;
 	ImagePageTable                                    m_image_page_table;
-	std::unordered_map<vk::Format, ImageId>           m_null_images;
 	Common::LeastRecentlyUsedCache<ImageId, uint64_t> m_lru_cache;
 	std::unordered_set<ImageId>                       m_download_images;
 	std::map<uint64_t, MetaDataInfo>                  m_surface_metas;
@@ -184,6 +186,7 @@ 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;
@@ -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,10 +241,11 @@ 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           = rt.clear_word0.word0;
+		desc.info.metadata.clear_word =
+		    rt.clear_word0.word0 | (uint64_t {rt.clear_word1.word1} << 32u);
 		desc.info.metadata.clear_register_valid = true;
 	}
 	desc.view_info.format = target_format.format;
@@ -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) {
@@ -874,14 +863,33 @@ void hw_check(const CommandBuffer& buffer) {
 	log_phase("done");

 	if (graphics_debug_dump_enabled() && RenderTargetMaskHasBoundMrt(buffer)) {
-		LOGF("MRT render target mask: 0x%08" PRIx32 "\n", hw.GetRenderTargetMask());
+		const auto& sh = hw.GetShaderRegisters();
+		LOGF("MRT render target mask: 0x%08" PRIx32
+		     " export_modes={%u,%u,%u,%u,%u,%u,%u,%u}\n",
+		     hw.GetRenderTargetMask(), sh.target_output_mode[0], sh.target_output_mode[1],
+		     sh.target_output_mode[2], sh.target_output_mode[3], sh.target_output_mode[4],
+		     sh.target_output_mode[5], sh.target_output_mode[6], sh.target_output_mode[7]);
 		for (uint32_t i = 0; i < 8; i++) {
 			const auto& mrt = hw.GetRenderTarget(i);
+			const auto& bc  = hw.GetBlendControl(i);
 			LOGF("\t RT%u addr=0x%010" PRIx64 " mask=0x%x fmt=0x%08" PRIx32
-			     " width=%u height=%u tile=%u\n",
+			     " type=%u order=%u shader_mask=0x%x"
+			     " width=%u height=%u tile=%u blend=%u bypass=%u clamp=%u round_mode=%u"
+			     " samples=%u fragments=%u dcc=%u dcc_addr=0x%010" PRIx64
+			     " cmask_clear=%u cmask_addr=0x%010" PRIx64
+			     " dcc_clear_key=%u clear={%08" PRIx32 ",%08" PRIx32 "}\n",
 			     i, mrt.base.addr, (hw.GetRenderTargetMask() >> (i * 4u)) & 0x0fu,
-			     static_cast<uint32_t>(mrt.info.format), mrt.attrib2.width + 1,
-			     mrt.attrib2.height + 1, static_cast<uint32_t>(mrt.attrib3.tile_mode));
+			     static_cast<uint32_t>(mrt.info.format), static_cast<uint32_t>(mrt.info.channel_type),
+			     static_cast<uint32_t>(mrt.info.channel_order), (sh.m_cbShaderMask >> (i * 4u)) & 0x0fu,
+			     mrt.attrib2.width + 1, mrt.attrib2.height + 1,
+			     static_cast<uint32_t>(mrt.attrib3.tile_mode), static_cast<uint32_t>(bc.enable),
+			     static_cast<uint32_t>(mrt.info.blend_bypass), static_cast<uint32_t>(mrt.info.blend_clamp),
+			     static_cast<uint32_t>(mrt.info.round_mode), mrt.attrib.num_samples,
+			     mrt.attrib.num_fragments,
+			     static_cast<uint32_t>(mrt.info.dcc_compression_enable), mrt.dcc_addr.addr,
+			     static_cast<uint32_t>(mrt.info.cmask_fast_clear_enable), mrt.cmask.addr,
+			     static_cast<uint32_t>(mrt.dcc.dcc_clear_key_enable), mrt.clear_word0.word0,
+			     mrt.clear_word1.word1);
 		}
 	}
 	if (rc.depth_clear_enable && hw.GetDepthClearValue() != 0.0f &&
@@ -177,9 +177,6 @@ static TextureCache::ImageDesc MakeDepthTargetDesc(const CommandBuffer& buffer,
 		if (z.htile_data_base_addr == 0 || (z.htile_data_base_addr & 0x7fffu) != 0) {
 			DepthFatal("invalid HTile metadata address");
 		}
-		if (z.depth_view.slice_max >= 32) {
-			DepthFatal("HTile clear tracking supports at most 32 slices");
-		}
 	}
 	if (!z.size.valid) {
 		DepthFatal("missing depth extent");
@@ -314,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 &&
@@ -424,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)) {
@@ -464,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 {
@@ -353,8 +355,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 +704,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 +713,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 {};
@@ -29,10 +29,8 @@ struct CachedImageView {
 };

 struct ImageUsage {
-	bool texture       = false;
 	bool storage       = false;
 	bool render_target = false;
-	bool depth_target  = false;
 	bool video_out     = false;
 };

@@ -70,40 +68,25 @@ public:
 	void CopyImageWithBuffer(Image& source, Buffer& buffer, TileManager& tiler);
 	void CopyMip(Image& source, uint32_t mip, uint32_t layer);

-	void InvalidateCpuWrite(uint64_t vaddr, uint64_t size) {
-		if (ImageRangeOverlaps(info.data.address, info.data.size, vaddr, size)) {
-			m_cpu_dirty        = true;
-			m_maybe_cpu_dirty  = false;
-			m_maybe_hash_valid = false;
-		} else if (ImagePageRangesOverlap(info.data.address, info.data.size, vaddr, size)) {
-			m_maybe_cpu_dirty = true;
-		}
+	void MarkCpuDirty() {
+		m_cpu_dirty       = true;
+		m_maybe_cpu_dirty = false;
 	}

 	[[nodiscard]] bool IsCpuDirty() const { return m_cpu_dirty || m_maybe_cpu_dirty; }
 	[[nodiscard]] bool IsDefinitelyCpuDirty() const { return m_cpu_dirty; }
 	[[nodiscard]] bool IsMaybeCpuDirty() const { return m_maybe_cpu_dirty; }
 	void               MarkMaybeCpuDirty() {
-		if (!m_cpu_dirty) {
+		if (!IsCpuDirty()) {
+			m_maybe_cpu_hash  = HashGuestEdges();
 			m_maybe_cpu_dirty = true;
 		}
 	}
-	[[nodiscard]] bool NeedsMaybeCpuHash() const {
-		return m_maybe_cpu_dirty && !m_maybe_hash_valid;
-	}
-	void SetMaybeCpuHash(uint64_t hash) {
-		if (!NeedsMaybeCpuHash()) {
-			EXIT("image cannot initialize maybe-dirty hash\n");
-		}
-		m_maybe_cpu_hash   = hash;
-		m_maybe_hash_valid = true;
-	}
 	[[nodiscard]] bool ResolveMaybeCpuHash(uint64_t hash) {
-		if (!m_maybe_cpu_dirty || !m_maybe_hash_valid || m_cpu_dirty) {
+		if (!m_maybe_cpu_dirty || m_cpu_dirty) {
 			EXIT("image cannot resolve maybe-dirty hash\n");
 		}
-		m_maybe_cpu_dirty  = false;
-		m_maybe_hash_valid = false;
+		m_maybe_cpu_dirty = false;
 		m_cpu_dirty |= hash != m_maybe_cpu_hash;
 		return m_cpu_dirty;
 	}
@@ -112,9 +95,8 @@ public:
 		if (!IsCpuDirty()) {
 			EXIT("clean image cannot complete a refresh\n");
 		}
-		m_cpu_dirty        = false;
-		m_maybe_cpu_dirty  = false;
-		m_maybe_hash_valid = false;
+		m_cpu_dirty       = false;
+		m_maybe_cpu_dirty = false;
 	}

 	[[nodiscard]] bool IsGpuModified() const noexcept { return m_gpu_modified; }
@@ -149,6 +131,8 @@ public:
 	uint64_t         track_addr     = 0;
 	uint64_t         track_addr_end = 0;
 	ImageId          depth_id {};
+	// Stencil associations that refer to this depth image.
+	std::vector<ImageId> stencil_ids;
 	// The current stencil plane's mapping into the depth image, also retained by its association.
 	ImageSubresourceRange stencil_subresources;
 	uint64_t         tick_accessed_last = 0;
@@ -165,12 +149,11 @@ private:

 	GraphicContext&   m_graphics;
 	CommandScheduler& m_scheduler;
-	uint64_t          m_maybe_cpu_hash   = 0;
-	bool              m_cpu_dirty        = false;
-	bool              m_maybe_cpu_dirty  = false;
-	bool              m_maybe_hash_valid = false;
-	bool              m_gpu_modified     = false;
-	bool              m_buffer_modified  = false;
+	uint64_t          m_maybe_cpu_hash  = 0;
+	bool              m_cpu_dirty       = false;
+	bool              m_maybe_cpu_dirty = false;
+	bool              m_gpu_modified    = false;
+	bool              m_buffer_modified = false;
 };

 namespace ImageOps {
@@ -21,6 +21,53 @@ bool IsDepthPlane(const ImageInfo& info) {
 	        info.tile_mode == Prospero::TileMode::kLinear);
 }

+GuestRange TiledMipRange(const ImageInfo& info, uint32_t level) {
+	const auto& block = info.tiling.block;
+	if (info.tile_mode != Prospero::TileMode::kStandard4KB ||
+	    block.family != TileBlockFamily::Standard4KB || info.IsVolume() || info.samples != 1 ||
+	    info.HasMetadata() || level >= info.resources.levels || level >= info.mip_layout.size() ||
+	    info.tiling.texel_width == 0 || info.tiling.texel_height == 0) {
+		return {};
+	}
+	const auto extent = info.MipExtent(level);
+	const auto width  = (extent.width + info.tiling.texel_width - 1u) / info.tiling.texel_width;
+	const auto height = (extent.height + info.tiling.texel_height - 1u) / info.tiling.texel_height;
+	const auto& mip   = info.mip_layout[level];
+	uint32_t last = 0, origin = 0;
+	if (!std::has_single_bit(width) || !std::has_single_bit(height) ||
+	    width > block.block_width || height > block.block_height ||
+	    mip.tail_x > block.block_width - width || mip.tail_y > block.block_height - height ||
+	    mip.tail_x % width != 0 || mip.tail_y % height != 0 ||
+	    !TileGetBlockOffset(block, width - 1u, height - 1u, 0, last) ||
+	    !TileGetBlockOffset(block, mip.tail_x, mip.tail_y, 0, origin)) {
+		return {};
+	}
+	// Standard4KB interleaves monotonically ordered X/Y bits, so the last
+	// corner bounds every active texel. Tiling preserves any holes in this span.
+	const uint64_t size = uint64_t {last} + block.bytes_per_element;
+	if ((origin & (std::bit_ceil(size) - 1u)) != 0 || size > block.block_size - origin ||
+	    mip.offset > UINT64_MAX - origin) {
+		return {};
+	}
+	return {mip.offset + origin, size};
+}
+
+uint32_t CompactMipSize(const ImageInfo& info) {
+	if (info.resources.levels != 1 || info.resources.layers != 1) {
+		return 0;
+	}
+	const auto range = TiledMipRange(info, 0);
+	if (range.size == 0 || range.size % 4 != 0 || range.address != 0) {
+		return 0;
+	}
+	const auto alignment = std::max<uint64_t>(256, std::bit_ceil(range.size));
+	const auto in_block  = info.data.address & (info.tiling.block.block_size - 1u);
+	return (info.data.address & (alignment - 1u)) == 0 &&
+	               range.size <= info.tiling.block.block_size - in_block
+	           ? static_cast<uint32_t>(range.size)
+	           : 0;
+}
+
 } // namespace

 vk::Extent3D ImageInfo::MipExtent(uint32_t level) const {
@@ -34,6 +81,7 @@ void ImageInfo::UpdateSize() {
 	        extent.height == 0 || extent.depth == 0 || resources.levels == 0 ||
 	        resources.levels > mip_layout.size() || resources.layers == 0 || samples == 0 ||
 	        !std::has_single_bit(samples));
+	const bool auto_size = data.size == 0;
 	TileTextureElementLayout element {};
 	EXIT_NOT_IMPLEMENTED(!TileGetTextureElementLayout(guest_format, element));
 	if (bytes_per_block == 0) {
@@ -164,9 +212,86 @@ void ImageInfo::UpdateSize() {
 			linear_offset     = Add(mip.linear_offset, Multiply(mip.linear_size, active.depth));
 			mip_pitch         = std::max(mip_pitch >> 1u, 1u);
 		}
+		const auto compact_size = CompactMipSize(*this);
+		if (compact_size != 0 && (auto_size || data.size == compact_size)) {
+			data.size = mip_layout[0].size = compact_size;
+		}
 	}
 }

+uint32_t ImageInfo::DataAlignment() const {
+	if (!IsTiled()) {
+		return 256;
+	}
+	const auto compact_size = CompactMipSize(*this);
+	if (compact_size != 0 && data.size == compact_size && mip_layout[0].size == compact_size) {
+		return std::max(256u, std::bit_ceil(compact_size));
+	}
+	return tiling.block.block_size;
+}
+
+int32_t ImageInfo::SliceOf(const ImageInfo& container, int32_t mip) const noexcept {
+	if (!IsCompatible(container) || tile_mode != container.tile_mode || type != container.type ||
+	    IsVolume() || resources.levels != 1 || resources.layers == 0 ||
+	    container.resources.layers == 0 || mip < 0 ||
+	    static_cast<uint32_t>(mip) >= container.resources.levels ||
+	    container.resources.levels > container.mip_layout.size() ||
+	    !data.Valid() || !container.data.Valid() || data.address < container.data.address ||
+	    data.End() > container.data.End()) {
+		return -1;
+	}
+	if (HasStencil() != container.HasStencil() ||
+	    (HasStencil() && (stencil.address < container.stencil.address ||
+	                      stencil.End() > container.stencil.End()))) {
+		return -1;
+	}
+
+	const auto  level  = static_cast<uint32_t>(mip);
+	const auto& layout = container.mip_layout[level];
+	if (extent.width != std::max(container.extent.width >> level, 1u) ||
+	    extent.height != std::max(container.extent.height >> level, 1u) ||
+	    extent.depth != container.extent.depth || mip_layout[0].pitch != layout.pitch ||
+	    mip_layout[0].height != layout.height || layout.size == 0 ||
+	    layout.size % container.resources.layers != 0 ||
+	    container.data.size % container.resources.layers != 0 ||
+	    data.size % resources.layers != 0) {
+		return -1;
+	}
+
+	// PS5 block slices contain the entire reversed mip chain. A rebased tail
+	// descriptor covers its addressed span within the shared tail block.
+	const auto slice_stride = container.data.size / container.resources.layers;
+	const auto child_stride = data.size / resources.layers;
+	auto mip_offset = layout.offset;
+	auto mip_size   = layout.size / container.resources.layers;
+	const auto compact_size = CompactMipSize(*this);
+	if (compact_size != 0 && child_stride == compact_size && mip_layout[0].size == compact_size) {
+		const auto range = TiledMipRange(container, level);
+		if (range.size != compact_size) {
+			return -1;
+		}
+		mip_offset = range.address;
+		mip_size   = range.size;
+	} else if (mip_layout[0].tail_x != layout.tail_x || mip_layout[0].tail_y != layout.tail_y) {
+		return -1;
+	}
+	if (child_stride != mip_size || mip_offset >= slice_stride ||
+	    child_stride > slice_stride - mip_offset ||
+	    (resources.layers > 1 && child_stride != slice_stride)) {
+		return -1;
+	}
+	const auto address_delta = data.address - container.data.address;
+	if (address_delta < mip_offset || (address_delta - mip_offset) % slice_stride != 0) {
+		return -1;
+	}
+	const auto layer = (address_delta - mip_offset) / slice_stride;
+	if (layer > INT32_MAX || layer >= container.resources.layers ||
+	    resources.layers > container.resources.layers - layer) {
+		return -1;
+	}
+	return static_cast<int32_t>(layer);
+}
+
 uint64_t ImageInfo::LinearSize(uint32_t levels) const {
 	levels = levels == 0 ? resources.levels : levels;
 	EXIT_IF(levels == 0 || levels > resources.levels || levels > mip_layout.size());
@@ -27,7 +27,7 @@ struct ImageMetadataInfo {
 	GuestRange          range;
 	ImageMetadataKind   kind                 = ImageMetadataKind::None;
 	uint32_t            control              = 0;
-	uint32_t            clear_word           = 0;
+	uint64_t            clear_word           = 0;
 	VideoOutCompression compression          = VideoOutCompression::Uncompressed;
 	bool                stencil_compressed   = false;
 	bool                clear_register_valid = false;
@@ -69,7 +69,7 @@ struct ImageInfo {
 	GuestRange                   data;
 	GuestRange                   stencil;
 	ImageMetadataInfo            metadata;
-	uint32_t                     htile_clear_mask = UINT32_MAX;
+	bool                         htile_initial_clear = true;
 	vk::Format                   pixel_format     = vk::Format::eUndefined;
 	Prospero::BufferFormat       guest_format     = Prospero::BufferFormat::kInvalid;
 	Prospero::ImageType          type             = Prospero::ImageType::kColor2D;
@@ -88,6 +88,7 @@ struct ImageInfo {
 	uint32_t first_tail_level    = 16;

 	void                                           UpdateSize();
+	[[nodiscard]] uint32_t                         DataAlignment() const;
 	[[nodiscard]] uint64_t                         LinearSize(uint32_t levels = 0) const;
 	[[nodiscard]] vk::Extent3D                     MipExtent(uint32_t level) const;
 	[[nodiscard]] std::vector<vk::BufferImageCopy> BufferCopies(uint64_t buffer_offset = 0,
@@ -136,53 +137,7 @@ struct ImageInfo {
 		}
 		return -1;
 	}
-	[[nodiscard]] int32_t SliceOf(const ImageInfo& container, int32_t mip) const noexcept {
-		if (!IsCompatible(container) || tile_mode != container.tile_mode || type != container.type ||
-		    IsVolume() || resources.levels != 1 || resources.layers == 0 ||
-		    container.resources.layers == 0 || mip < 0 ||
-		    static_cast<uint32_t>(mip) >= container.resources.levels ||
-		    container.resources.levels > container.mip_layout.size() ||
-		    !data.Valid() || !container.data.Valid() || data.address < container.data.address ||
-		    data.End() > container.data.End()) {
-			return -1;
-		}
-		if (HasStencil() != container.HasStencil() ||
-		    (HasStencil() && (stencil.address < container.stencil.address ||
-		                      stencil.End() > container.stencil.End()))) {
-			return -1;
-		}
-
-		const auto  level  = static_cast<uint32_t>(mip);
-		const auto& layout = container.mip_layout[level];
-		if (extent.width != std::max(container.extent.width >> level, 1u) ||
-		    extent.height != std::max(container.extent.height >> level, 1u) ||
-		    extent.depth != container.extent.depth || mip_layout[0].pitch != layout.pitch ||
-		    mip_layout[0].height != layout.height || layout.size == 0 ||
-		    layout.size % container.resources.layers != 0 ||
-		    container.data.size % container.resources.layers != 0 ||
-		    data.size % resources.layers != 0) {
-			return -1;
-		}
-
-		// PS5 block slices contain the entire reversed mip chain.
-		const auto slice_stride = container.data.size / container.resources.layers;
-		const auto child_stride = data.size / resources.layers;
-		if (child_stride != layout.size / container.resources.layers ||
-		    layout.offset >= slice_stride || child_stride > slice_stride - layout.offset ||
-		    (resources.layers > 1 && child_stride != slice_stride)) {
-			return -1;
-		}
-		const auto address_delta = data.address - container.data.address;
-		if (address_delta < layout.offset || (address_delta - layout.offset) % slice_stride != 0) {
-			return -1;
-		}
-		const auto layer = (address_delta - layout.offset) / slice_stride;
-		if (layer > INT32_MAX || layer >= container.resources.layers ||
-		    resources.layers > container.resources.layers - layer) {
-			return -1;
-		}
-		return static_cast<int32_t>(layer);
-	}
+	[[nodiscard]] int32_t SliceOf(const ImageInfo& container, int32_t mip) const noexcept;
 };

 struct ImageViewInfo {
@@ -414,9 +369,10 @@ inline constexpr std::array<VideoOutFormatPolicy, 8> VIDEO_OUT_FORMAT_POLICIES {
 	return false;
 }

-[[nodiscard]] inline bool DecodePackedColorClear(vk::Format format, uint32_t packed,
+[[nodiscard]] inline bool DecodePackedColorClear(vk::Format format, uint64_t clear_words,
                                                 vk::ClearColorValue& clear) {
 	vk::ClearColorValue next {};
+	const auto packed = static_cast<uint32_t>(clear_words);
 	const auto unorm8 = [](uint32_t value) { return static_cast<float>(value & 0xffu) / 255.0f; };
 	const auto srgb8  = [](uint32_t value) {
 		const auto encoded = static_cast<float>(value & 0xffu) / 255.0f;
@@ -434,6 +390,15 @@ inline constexpr std::array<VideoOutFormatPolicy, 8> VIDEO_OUT_FORMAT_POLICIES {
 		                            (mantissa << 13u));
 	};
 	switch (format) {
+		case vk::Format::eB10G11R11UfloatPack32:
+			next.float32[0] = float16((packed & 0x7ffu) << 4u);
+			next.float32[1] = float16(((packed >> 11u) & 0x7ffu) << 4u);
+			next.float32[2] = float16((packed >> 22u) << 5u);
+			break;
+		case vk::Format::eR16G16B16A16Sfloat:
+			next.float32[2] = float16(static_cast<uint32_t>(clear_words >> 32u));
+			next.float32[3] = float16(static_cast<uint32_t>(clear_words >> 48u));
+			[[fallthrough]];
 		case vk::Format::eR16G16Sfloat:
 			next.float32[1] = float16(packed >> 16u);
 			[[fallthrough]];
@@ -444,9 +409,10 @@ inline constexpr std::array<VideoOutFormatPolicy, 8> VIDEO_OUT_FORMAT_POLICIES {
 		case vk::Format::eR16Unorm:
 			next.float32[0] = static_cast<float>(packed & 0xffffu) / 65535.0f;
 			break;
-		// A single-plane float target carries its clear as raw float bits, the same encoding the
-		// depth decoder below uses. Without this the clear is discarded and the target keeps stale
-		// contents.
+		// Each clear register stores one raw 32-bit float channel.
+		case vk::Format::eR32G32Sfloat:
+			next.float32[1] = std::bit_cast<float>(static_cast<uint32_t>(clear_words >> 32u));
+			[[fallthrough]];
 		case vk::Format::eR32Sfloat: next.float32[0] = std::bit_cast<float>(packed); break;
 		case vk::Format::eR32Uint: next.uint32[0] = packed; break;
 		case vk::Format::eR32Sint: next.int32[0] = static_cast<int32_t>(packed); break;
@@ -507,7 +473,7 @@ inline constexpr std::array<VideoOutFormatPolicy, 8> VIDEO_OUT_FORMAT_POLICIES {
 		clear = vk::ClearColorValue {};
 		return true;
 	}
-	return DecodePackedColorClear(format, packed, clear);
+	return DecodePackedColorClear(format, uint64_t {packed} | (uint64_t {packed} << 32u), clear);
 }

 [[nodiscard]] inline bool DecodePackedStencilClear(uint32_t packed, uint8_t& clear) {
@@ -31,18 +31,6 @@ namespace ImageViewOps {
 	       ImageViewOps::IsFormatDepthCompatible(view_format);
 }

-[[nodiscard]] inline bool IsValidImageSwizzle(uint32_t swizzle,
-                                             uint32_t component_count = 4) noexcept {
-	if ((swizzle & ~0xfffu) != 0) {
-		return false;
-	}
-	for (uint32_t channel = 0; channel < 4; channel++) {
-		const auto selector = GetDstSel(swizzle, channel);
-		if (selector > 1 && (selector < 4 || selector >= 4 + component_count)) return false;
-	}
-	return true;
-}
-
 [[noreturn]] inline void UnsupportedColorView(const char* usage, vk::Format image_format,
                                              vk::Format view_format, uint32_t swizzle) noexcept {
 	EXIT("unsupported %s color image view: image_format=%d view_format=%d swizzle=0x%03x\n", usage,
@@ -64,7 +64,10 @@ RenderTargetFormatInfo TextureGetRenderTargetFormat(Prospero::ChannelLayout layo
                                                    Prospero::ChannelOrder  order) {
 	const auto encoding = Prospero::ResolveRenderTargetFormat(layout, type);
 	if (encoding.IsValid() && encoding.SupportsOrder(order)) {
-		const auto host_format = ResolveHostFormat(encoding.buffer_format, order);
+		const auto backing_format = encoding.buffer_format == Prospero::BufferFormat::k11_11_10UNorm
+		                                ? Prospero::RemapTextureFormat(encoding.buffer_format)
+		                                : encoding.buffer_format;
+		const auto host_format = ResolveHostFormat(backing_format, order);
 		const auto bytes       = Prospero::RenderTargetBytesPerElement(encoding.buffer_format);
 		if (host_format.format != vk::Format::eUndefined && bytes != 0) {
 			const auto order_mapping =
@@ -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();
@@ -76,6 +76,10 @@ vk::DescriptorType NativeDescriptorType(BindingKind kind) {
 }

 vk::DescriptorImageInfo MakeImageInfo(const TextureBinding& texture, uint32_t element) {
+	if (!texture.image_id) {
+		EXIT_IF(texture.image_view != nullptr || !texture.mip_views.empty());
+		return {nullptr, nullptr, vk::ImageLayout::eGeneral};
+	}
 	vk::ImageView view = nullptr;
 	if (texture.mip_views.empty()) {
 		if (element == 0u) {
@@ -84,7 +88,7 @@ vk::DescriptorImageInfo MakeImageInfo(const TextureBinding& texture, uint32_t el
 	} else if (element < texture.mip_views.size()) {
 		view = texture.mip_views[element];
 	}
-	EXIT_IF(!texture.image_id || view == nullptr || texture.layout == vk::ImageLayout::eUndefined);
+	EXIT_IF(view == nullptr || texture.layout == vk::ImageLayout::eUndefined);
 	return {nullptr, view, texture.layout};
 }

@@ -139,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() |
@@ -150,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;
@@ -277,40 +290,17 @@ static bool IsSupportedStorageTextureDescriptor(const ShaderRecompiler::IR::Imag
 	       !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 = format == Prospero::BufferFormat::k32SInt && uint_resource &&
-	                              resource.written && !resource.read && !resource.atomic;
-	const auto numeric_class = Prospero::SampledTextureNumericClass(format);
+	const bool raw_sint_storage = uint_resource &&
+	    ShaderRecompiler::IR::RawSignedStorageFormat(resource, format) != Prospero::BufferFormat::kInvalid;
+	const auto numeric_class =
+	    Prospero::SampledTextureNumericClass(Prospero::RemapTextureFormat(format));
 	const bool raw_float_atomic = format == Prospero::BufferFormat::k32Float && uint_resource &&
 	                              resource.atomic && !resource.atomic64;
 	const bool format_ok =
@@ -347,39 +337,6 @@ void ValidateStorageTexture(const ShaderRecompiler::IR::ImageResource& resource,
 	     descriptor.fields[5], descriptor.fields[6], descriptor.fields[7]);
 }

-static TextureCache::ImageDesc NullTextureDesc(const ShaderRecompiler::IR::ImageResource& resource,
-                                               TextureCache::BindingType                  binding) {
-	TextureCache::ImageDesc desc {};
-	switch (resource.numeric_class) {
-		case Prospero::TextureNumericClass::Float:
-			desc.info.guest_format = Prospero::BufferFormat::k32Float;
-			break;
-		case Prospero::TextureNumericClass::Uint:
-			desc.info.guest_format = resource.atomic64 ? Prospero::BufferFormat::k32_32UInt
-			                                         : Prospero::BufferFormat::k32UInt;
-			break;
-		case Prospero::TextureNumericClass::Sint:
-			desc.info.guest_format = Prospero::BufferFormat::k32SInt;
-			break;
-		default: EXIT("null image has unsupported numeric class\n");
-	}
-	desc.info.pixel_format    = VulkanFormat(desc.info.guest_format);
-	desc.info.type            = Prospero::ImageType::kColor2D;
-	desc.info.extent          = {1, 1, 1};
-	desc.info.resources       = {1, 1};
-	desc.info.bytes_per_block = Prospero::NumBytesPerElement(desc.info.guest_format);
-	desc.info.samples         = 1;
-	desc.info.mip_layout[0]   = {0, 0, 1, 1};
-	desc.view_info.format     = resource.atomic64 ? vk::Format::eR64Uint : desc.info.pixel_format;
-	desc.view_info.type       = vk::ImageViewType::e2D;
-	desc.view_info.aspect     = vk::ImageAspectFlagBits::eColor;
-	desc.view_info.usage      = binding == TextureCache::BindingType::Storage
-	                                ? vk::ImageUsageFlagBits::eStorage
-	                                : vk::ImageUsageFlagBits::eSampled;
-	desc.type                 = binding;
-	return desc;
-}
-
 static ImageViewInfo TextureViewInfo(const ShaderRecompiler::IR::ImageResource& resource,
                                     const ShaderTextureResource& descriptor, vk::Format format,
                                     const SurfaceFormatInfo& surface_format, bool storage,
@@ -515,10 +472,7 @@ TextureBinding RenderExecutor::ResolveTexture(const ShaderRecompiler::IR::ImageR

 	auto& texture_cache = m_context.GetTextureCache();
 	if (descriptor.IsNull()) {
-		auto       desc = NullTextureDesc(resource, storage ? TextureCache::BindingType::Storage
-		                                                    : TextureCache::BindingType::Texture);
-		const auto id   = texture_cache.FindImage(desc);
-		return {id, nullptr, std::move(desc)};
+		return {.layout = vk::ImageLayout::eGeneral};
 	}

 	const auto address         = descriptor.Base40();
@@ -566,7 +520,8 @@ TextureBinding RenderExecutor::ResolveTexture(const ShaderRecompiler::IR::ImageR
 	const bool shader_conversion =
 	    surface_format.conversion_format != Prospero::BufferFormat::kInvalid;
 	const bool sampled_numeric_class =
-	    storage || resource.numeric_class == Prospero::SampledTextureNumericClass(format);
+	    storage || resource.numeric_class ==
+	                   Prospero::SampledTextureNumericClass(Prospero::RemapTextureFormat(format));
 	if (!storage && resource.resource_class == ShaderRecompiler::IR::ImageResourceClass::Sampled &&
 	    !sampled_numeric_class) {
 		EXIT("sampled image numeric class mismatch: numeric=%u format=%u addr=0x%016" PRIx64 "\n",
@@ -585,10 +540,11 @@ TextureBinding RenderExecutor::ResolveTexture(const ShaderRecompiler::IR::ImageR
 			pixel_format = depth_format->depth_attachment_format;
 		}
 	}
+	const auto raw_sint_format = ShaderRecompiler::IR::RawSignedStorageFormat(resource, format);
 	const auto storage_view_format = resource.atomic64 ? vk::Format::eR64Uint
-	                                 : storage && (resource.atomic ||
-	                                               format == Prospero::BufferFormat::k32SInt)
-	                                     ? vk::Format::eR32Uint
+	                                 : storage && resource.atomic ? vk::Format::eR32Uint
+	                                 : raw_sint_format != Prospero::BufferFormat::kInvalid
+	                                     ? VulkanFormat(raw_sint_format)
 	                                     : SrgbStorageViewFormat(pixel_format);
 	const auto view_format         = storage && storage_view_format != vk::Format::eUndefined
 	                                     ? storage_view_format
@@ -605,6 +561,9 @@ TextureBinding RenderExecutor::ResolveTexture(const ShaderRecompiler::IR::ImageR
 	    block_bytes != 0 ? block_bytes : Prospero::NumBytesPerElement(format);
 	desc.info.samples   = samples;
 	desc.info.tile_mode = tile;
+	if (!resource.r128 && descriptor.MetaCompress() && tile != Prospero::TileMode::kDepth) {
+		desc.info.metadata.kind = ImageMetadataKind::Dcc;
+	}
 	if (levels > physical_levels) {
 		if (!ResolveTextureMipView(desc.info, physical_levels,
 		                           !resource.r128 && descriptor.MetaCompress(), view_levels,
@@ -617,8 +576,7 @@ TextureBinding RenderExecutor::ResolveTexture(const ShaderRecompiler::IR::ImageR
 	} else {
 		desc.info.UpdateSize();
 	}
-	const auto alignment =
-	    tile == Prospero::TileMode::kLinear ? 256u : desc.info.tiling.block.block_size;
+	const auto alignment = desc.info.DataAlignment();
 	EXIT_NOT_IMPLEMENTED(desc.info.data.size == 0 || alignment == 0 ||
 	                     (address & (static_cast<uint64_t>(alignment) - 1u)) != 0);
 	if (storage) {
@@ -629,7 +587,6 @@ TextureBinding RenderExecutor::ResolveTexture(const ShaderRecompiler::IR::ImageR
 		(void)TileGetDccSize(width, height, volume ? depth : image_layers,
 		                     desc.info.bytes_per_block, physical_levels, tile, metadata_size,
 		                     std::countr_zero(samples));
-		desc.info.metadata.kind          = ImageMetadataKind::Dcc;
 		desc.info.metadata.range         = {descriptor.MetaAddr() << 8u, metadata_size.size};
 		desc.info.metadata.dcc_alpha_msb = descriptor.DccAlphaPos();
 	}
@@ -685,10 +642,10 @@ static vk::DescriptorBufferInfo NativeUpload(RenderContext&            context,
 }

 void RenderExecutor::BindImage(ImageId id, bool storage) {
-	auto& image = m_context.GetTextureCache().GetImage(id);
-	if (image.info.data.Empty()) {
+	if (!id) {
 		return;
 	}
+	auto& image = m_context.GetTextureCache().GetImage(id);
 	if (image.binding.is_bound) {
 		image.binding.force_general |= image.binding.shader_write != storage;
 	}
@@ -768,6 +725,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;
@@ -785,14 +748,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)};
 		}
 	}
@@ -841,8 +808,9 @@ void RenderExecutor::RebindImages(PreparedBindings& prepared) {
 	auto& texture_cache = m_context.GetTextureCache();
 	// Acquire each view before a later overlapping descriptor can replace its image.
 	for (uint32_t i = 0; i < program.info.images.size(); i++) {
+		if (!images[i].image_id) continue;
 		const auto old_image = texture_cache.m_slot_images.try_get(images[i].image_id);
-		if (old_image == nullptr || (!old_image->registered && !old_image->info.data.Empty()) ||
+		if (old_image == nullptr || !old_image->registered ||
 		    old_image->binding.needs_rebind) {
 			if (old_image != nullptr) {
 				old_image->binding = {};
@@ -873,7 +841,6 @@ void RenderExecutor::RebindImages(PreparedBindings& prepared) {
 		auto&      image   = texture_cache.GetImage(binding.image_id);
 		const bool storage = binding.desc.type == TextureCache::BindingType::Storage;
 		image.usage.storage |= storage;
-		image.usage.texture |= !storage;
 	}
 }

@@ -894,7 +861,7 @@ void RenderExecutor::PrepareGraphicsBindings(std::span<PreparedBindings* const>
 	for (auto& target: colors) {
 		EXIT_IF(!target.image_id);
 		const auto old_image = cache.m_slot_images.try_get(target.image_id);
-		if (old_image == nullptr || (!old_image->registered && !old_image->info.data.Empty()) ||
+		if (old_image == nullptr || !old_image->registered ||
 		    old_image->binding.needs_rebind) {
 			if (old_image != nullptr) {
 				old_image->binding = {};
@@ -974,8 +941,12 @@ 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");
 					}
 				}
@@ -1004,19 +975,14 @@ void RenderExecutor::CommitBindings(CommandBuffer&                     buffer,
 		}

 		for (uint32_t i = 0; i < program.info.images.size(); i++) {
-			auto& image   = m_context.GetTextureCache().GetImage(descriptors.images[i].image_id);
 			auto& binding = descriptors.images[i];
+			if (!binding.image_id) continue;
+			auto& image = m_context.GetTextureCache().GetImage(binding.image_id);
 			const auto&                 view = binding.desc.view_info;
 			const ImageSubresourceRange range {view.base_level, view.level_count, view.base_layer,
 			                                   view.layer_count};
 			const bool storage = binding.desc.type == TextureCache::BindingType::Storage;
-			if (image.info.data.Empty()) {
-				image.Transit(vk::ImageLayout::eGeneral,
-				              storage ? vk::AccessFlagBits2::eShaderRead |
-				                            vk::AccessFlagBits2::eShaderWrite
-				                      : vk::AccessFlagBits2::eShaderRead,
-				              range, vk_buffer);
-			} else if (image.binding.is_target) {
+			if (image.binding.is_target) {
 				const auto layout = image.binding.attachment_layout;
 				EXIT_IF(layout == vk::ImageLayout::eUndefined);
 				if (image.info.IsDepth()) {
@@ -1117,7 +1083,8 @@ 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];
-			EXIT_IF(resource.mip_count != (texture.mip_views.empty() ? 1u : texture.mip_views.size()));
+			EXIT_IF(texture.image_id &&
+			        resource.mip_count != (texture.mip_views.empty() ? 1u : texture.mip_views.size()));
 			const auto kind = *ShaderRecompiler::IR::DescriptorBindingForImage(resource);
 			const auto start = image_starts[ShaderRecompiler::IR::ImageBindingIndex(kind)] +
 			                   resource.descriptor_index;
@@ -1,11 +1,13 @@
 #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"
 #include "common/profiler.h"
 #include "common/threads.h"
+#include "graphics/guest_gpu/gpu_format.h"
 #include "graphics/guest_gpu/hardwareContext.h"
 #include "graphics/host_gpu/renderer/colorRenderTarget.h"
 #include "graphics/host_gpu/renderer/debug.h"
@@ -17,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>
@@ -39,6 +41,8 @@ namespace Libs::Graphics {

 namespace {

+constexpr auto DriverCacheSaveInterval = std::chrono::minutes {5};
+
 uint8_t RemapSourceAlphaFactor(uint8_t factor) {
 	switch (static_cast<Prospero::BlendFactor>(factor)) {
 		case Prospero::BlendFactor::kSrcAlpha:
@@ -78,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);
 }

@@ -148,7 +153,7 @@ void DumpShaderOriginal(const char* stage_name, uint64_t shader_hash,
 }

 bool ValidateShaderSpirv(const char* label, uint64_t shader_hash,
-                         const std::vector<uint32_t>& spirv) {
+                         const std::vector<uint32_t>& spirv, bool maintenance8_enabled) {
 	if (!Config::ShaderValidationEnabled()) {
 		return true;
 	}
@@ -160,7 +165,9 @@ bool ValidateShaderSpirv(const char* label, uint64_t shader_hash,
 		                        static_cast<int>(position.column), static_cast<int>(position.index),
 		                        message);
 	});
-	if (tools.Validate(spirv)) {
+	spvtools::ValidatorOptions validator_options;
+	validator_options.SetAllowOffsetTextureOperand(maintenance8_enabled);
+	if (tools.Validate(spirv.data(), spirv.size(), validator_options)) {
 		return true;
 	}
 	spvtools::SpirvTools disassembler(SPV_ENV_VULKAN_1_2);
@@ -268,6 +275,7 @@ struct PipelineCache::ProgramCache {
 		ShaderRecompiler::IR::ResourceSnapshot       resources;
 		ShaderRecompiler::IR::ResourceSpecialization specialization;
 		std::vector<Permutation>                    permutations;
+		bool                                       skip_dispatch = false;
 	};

 	struct ProgramKeyHash {
@@ -294,16 +302,20 @@ 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)) {
+		if (!ValidateShaderSpirv(options.dump_label, options.shader_hash, result.spirv,
+		                         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);
 		}
 		DumpShaderSpirv(stage_name, options.shader_hash, result.spirv);

-		const auto module = CompileSPV(result.spirv, device);
+		const auto module = CompileSPV(result.spirv, graphics.device);
 		EXIT_IF(module == nullptr);
 		if (options.dump_ir) {
 			LOGF("%s SPIR-V words=%" PRIu64 " wave_size=%u\n", options.dump_label,
@@ -344,10 +356,12 @@ struct PipelineCache::ProgramCache {
 		};
 		if constexpr (std::is_same_v<InputInfo, ShaderComputeInputInfo>) {
 			runtime.workgroup_counts = input_info.workgroup_counts;
+			runtime.indirect_args_address = input_info.indirect_args_address;
 		}
 		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);
@@ -384,11 +398,13 @@ struct PipelineCache::ProgramCache {
 		}
 		ShaderRecompiler::CompileOptions options;
 		options.stage       = stage;
+		options.profile     = &profile;
 		options.shader_hash = params.hash;
 		options.user_data   = user_data;
 		options.back_code      = params.back_code;
 		options.dump_ir     = Config::GetShaderLogDirection() != Config::LogDirection::Silent;
 		options.early_dump  = options.dump_ir;
+		options.emit_debug_info = Config::GraphicsDebugDumpEnabled();
 		options.dump_label  = label;
 		options.input_info  = stage_input;

@@ -405,10 +421,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;
@@ -419,7 +443,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(
@@ -456,26 +480,34 @@ struct PipelineCache::ProgramCache {
 		return permutation.handle;
 	}

-	explicit ProgramCache(vk::Device device): device(device) {
+	explicit ProgramCache(GraphicContext& graphics)
+	    : graphics(graphics),
+	      // 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() {
 		for (const auto& [key, entry]: programs) {
 			(void)key;
 			for (const auto& permutation: entry.permutations) {
-				device.destroyShaderModule(permutation.handle.module, nullptr);
+				graphics.device.destroyShaderModule(permutation.handle.module, nullptr);
 			}
 		}
 	}

 	std::unordered_map<ProgramKey, SourceEntry, ProgramKeyHash> programs;
 	ProgramKey                                                  lookup_key;
-	vk::Device                                                  device;
+	GraphicContext&                                             graphics;
+	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.device)) {
+    : m_graphics(graphics), m_program_cache(std::make_unique<ProgramCache>(graphics)),
+      m_next_driver_cache_save(std::chrono::steady_clock::now() + DriverCacheSaveInterval) {
 	EXIT_NOT_IMPLEMENTED(!Common::Thread::IsMainThread());
 	InitializeDriverCache();
 }
@@ -506,17 +538,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);
@@ -549,7 +570,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 {
@@ -576,17 +597,28 @@ 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");
 	}
 }

-void PipelineCache::Save() {
+void PipelineCache::CheckpointDriverCache() {
 	if (m_driver_cache == nullptr) {
 		return;
 	}
+	m_driver_cache_dirty = true;
+	if (std::chrono::steady_clock::now() >= m_next_driver_cache_save) {
+		Save();
+	}
+}
+
+void PipelineCache::Save() {
+	if (m_driver_cache == nullptr || !m_driver_cache_dirty) {
+		return;
+	}
+	m_next_driver_cache_save = std::chrono::steady_clock::now() + DriverCacheSaveInterval;

 	size_t               size = 0;
 	vk::Result           result;
@@ -606,8 +638,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);
@@ -629,16 +661,20 @@ void PipelineCache::Save() {
 	}
 	const bool flushed = !file.IsInvalid() && file.Flush();
 	file.Close();
-	if (prefix_written != prefix.size() || payload_written != payload.size() || !flushed ||
-	    !Common::File::RenameFile(temp_path, m_driver_cache_path)) {
+	const bool written = prefix_written == prefix.size() && payload_written == payload.size() && flushed;
+	std::error_code error;
+	if (written) {
+		std::filesystem::rename(temp_path, m_driver_cache_path, error);
+	}
+	if (!written || error) {
 		PipelineCacheLog("Vulkan pipeline cache: failed to write {}",
 		                 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_graphics.device.destroyPipelineCache(m_driver_cache, nullptr);
-	m_driver_cache = nullptr;
+	m_driver_cache_dirty = false;
 }

 PipelineCache::GraphicsPrograms PipelineCache::GetGraphicsPrograms(
@@ -654,6 +690,17 @@ PipelineCache::GraphicsPrograms PipelineCache::GetGraphicsPrograms(
 		vertex_params[0] = PrepareProgram(vertex_regs, context, user_config, vertex_info[0]);
 	}
 	const bool mesh_active = vertex_info[0].logical_stage == ShaderType::Mesh;
+	const bool object_primitive_id = context.HasObjectPrimitiveIdOutput();
+	if (object_primitive_id) {
+		// CxObjectIdMode: register source plus a 32-bit pipelined primitive-ID addition.
+		EXIT_NOT_IMPLEMENTED(sh.m_paClObjprimIdCntl != 2u || mesh_active || tess_active ||
+		                     !pixel_active || (sh.m_vgtDrawPayloadCntl & 8u) == 0 ||
+		                     user_config.GetPrimType() != Prospero::PrimitiveType::kTriList);
+		const auto& target = context.GetRenderTarget(7);
+		EXIT_NOT_IMPLEMENTED(Prospero::ResolveRenderTargetFormat(target.info.format,
+		                                                       target.info.channel_type)
+		                         .buffer_format != Prospero::BufferFormat::k32UInt);
+	}
 	if (mesh_active) {
 		EXIT_NOT_IMPLEMENTED(!m_graphics.mesh_shader_enabled);
 		auto& mesh              = vertex_info[0].mesh;
@@ -675,6 +722,29 @@ PipelineCache::GraphicsPrograms PipelineCache::GetGraphicsPrograms(
 	ShaderParams pixel_params;
 	if (pixel_active) {
 		pixel_params      = PrepareProgram(pixel_regs, sh, target_export_mapping, pixel_info);
+		pixel_info.object_primitive_id = object_primitive_id;
+		for (uint32_t export_index = 0; export_index < 8u; export_index++) {
+			const auto slot = ShaderPixelExportTarget(pixel_info.target_shader_mask, export_index);
+			if (slot >= 8u || pixel_info.target_output_mode[export_index] == 0u) continue;
+			const auto& rt = context.GetRenderTarget(slot);
+			const auto target_mask = render_target_mask_slot(context.GetRenderTargetMask(), slot);
+			if (rt.base.addr == 0 || target_mask == 0u) continue;
+			const auto format = Prospero::ResolveRenderTargetFormat(rt.info.format, rt.info.channel_type);
+			if (format.buffer_format != Prospero::BufferFormat::k11_11_10UNorm) continue;
+			const auto mapping = pixel_info.target_export_mapping[export_index];
+			const auto shader_mask = render_target_mask_slot(pixel_info.target_shader_mask, slot);
+			const bool blend = context.GetBlendControl(slot).enable && !rt.info.blend_bypass;
+			if (blend || (mapping.ApplyMask(target_mask) & 7u) != 7u ||
+			    (mapping.ApplyMask(shader_mask) & 7u) != 7u ||
+			    pixel_info.target_output_mode[export_index] != 9u) {
+				EXIT("unsupported packed UNorm render-target export: slot=%u export=%u mode=%u "
+				     "target_mask=0x%x shader_mask=0x%x mapping=0x%02x blend=%u\n",
+				     slot, export_index, pixel_info.target_output_mode[export_index], target_mask,
+				     shader_mask, mapping.packed, static_cast<uint32_t>(blend));
+			}
+			pixel_info.target_packed_unorm_mask |= 1u << export_index;
+			pixel_info.target_round_mode |= static_cast<uint8_t>(rt.info.round_mode) << export_index;
+		}
 		if (!mesh_active && !tess_active && Prospero::IsRectList(user_config.GetPrimType())) {
 			pixel_info.parameter_mode = ShaderPixelParameterMode::Rectangle;
 		} else if (context.GetModeControl().provoking_vtx_last) {
@@ -717,10 +787,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;
@@ -730,9 +802,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 =
-	    mesh_active ? ShaderRecompiler::IR::PushData::MeshDrawDwordCount : 0;
+	uint32_t push_data_cursor = ShaderRecompiler::IR::PushData::GraphicsDrawDwords(
+	    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);
@@ -797,6 +876,9 @@ PipelineCache::Pipeline& PipelineCache::GetGraphicsPipeline(
 		EXIT_IF(!colors[i].image_id || colors[i].desc.view_info.format == vk::Format::eUndefined);
 		static_params.color_mask[slot] = colors[i].export_mapping.ApplyMask(
 		    render_target_mask_slot(ctx.GetRenderTargetMask(), colors[i].target_slot));
+		if (colors[i].desc.info.guest_format == Prospero::BufferFormat::k11_11_10UNorm) {
+			static_params.color_mask[slot] = 1u;
+		}
 		rendering.color_formats[slot] = colors[i].desc.view_info.format;
 		if (attachment_samples == 0) {
 			attachment_samples = colors[i].desc.info.samples;
@@ -896,6 +978,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;
@@ -943,8 +1026,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);
@@ -952,6 +1040,7 @@ PipelineCache::Pipeline& PipelineCache::GetGraphicsPipeline(

 	auto [iter, inserted] = m_graphics_pipelines.emplace(std::move(key), std::move(cached));
 	EXIT_IF(!inserted);
+	CheckpointDriverCache();

 	return *iter->second;
 }
@@ -973,13 +1062,20 @@ 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);

 	auto [iter, inserted] = m_compute_pipelines.emplace(compute_program.id, std::move(cached));
 	EXIT_IF(!inserted);
+	CheckpointDriverCache();

 	return *iter->second;
 }
@@ -8,6 +8,7 @@
 #include "graphics/host_gpu/vulkanCommon.h"
 #include "graphics/shader/shader.h"

+#include <chrono>
 #include <cstddef>
 #include <filesystem>
 #include <memory>
@@ -38,6 +39,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;
@@ -61,7 +63,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 {};
@@ -173,11 +175,14 @@ private:
 	std::unique_ptr<ProgramCache> m_program_cache;
 	vk::PipelineCache             m_driver_cache = nullptr;
 	std::filesystem::path         m_driver_cache_path;
+	std::chrono::steady_clock::time_point m_next_driver_cache_save;
+	bool                          m_driver_cache_dirty = false;
 	std::unordered_map<GraphicsPipelineKey, std::unique_ptr<Pipeline>, GraphicsPipelineKeyHash>
 	                                                        m_graphics_pipelines;
 	std::unordered_map<uint64_t, std::unique_ptr<Pipeline>> m_compute_pipelines;

 	void InitializeDriverCache();
+	void CheckpointDriverCache();
 };

 void LogPipelineTrace(const char* phase, uint64_t vertex_program_id, uint64_t pixel_program_id);
@@ -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;
@@ -19,6 +19,7 @@
 #include "graphics/shader/shader.h"

 #include <algorithm>
+#include <array>
 #include <bit>
 #include <limits>
 #include <span>
@@ -192,9 +193,40 @@ static void AddLayoutBindings(std::vector<vk::DescriptorSetLayoutBinding>& descr

 static void CreateDescriptorLayout(GraphicContext& graphics, PipelineCache::Pipeline& pipeline,
                                   std::span<const vk::DescriptorSetLayoutBinding> bindings) {
-	uint32_t descriptor_count = 0;
+	const auto& limits = graphics.GetPhysicalDeviceProperties().limits;
+	uint64_t descriptor_count = 0;
+	uint64_t storage_buffers = 0;
+	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 (sampler) stage_samplers[stage] += binding.descriptorCount;
+			else
+				stage_resources[stage] += binding.descriptorCount;
+			if (stage_buffers[stage] > limits.maxPerStageDescriptorStorageBuffers ||
+			    stage_resources[stage] > limits.maxPerStageResources ||
+			    stage_samplers[stage] > limits.maxPerStageDescriptorSamplers) {
+				EXIT("native descriptor stage %u exceeds device limits: buffers=%" PRIu64
+				     "/%u resources=%" PRIu64 "/%u samplers=%" PRIu64 "/%u\n", stage, stage_buffers[stage],
+				     limits.maxPerStageDescriptorStorageBuffers, stage_resources[stage], limits.maxPerStageResources,
+				     stage_samplers[stage], limits.maxPerStageDescriptorSamplers);
+			}
+		}
+	}
+	if (storage_buffers > limits.maxDescriptorSetStorageBuffers) {
+		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;

@@ -390,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++) {
@@ -604,6 +638,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 {
@@ -163,8 +173,9 @@ public:
 	                    std::span<PreparedBindings* const> bindings);

 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;
@@ -196,14 +207,16 @@ private:
 	                                               std::span<PreparedBindings* const> stages = {});
 	[[nodiscard]] bool        ResolveColorTargets(CommandBuffer& buffer,
 	                                              uint32_t render_target_slice_offset);
+	[[nodiscard]] bool        ConsumeMetadataColorOperation(CommandBuffer& buffer,
+	                                                       uint32_t render_target_slice_offset);
 	void                      BindImage(ImageId id, bool storage);
 	void                      BindRenderTarget(ImageId id);
 	void                      ResetBindings();
 	[[nodiscard]] bool        TryConsumeComputeMetaClear(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 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"
@@ -22,6 +23,7 @@
 #include "graphics/shader/recompiler/ir/passes/ResourceMaterialization.h"
 #include "graphics/shader/shader.h"
 #include "kernel/eventQueue.h"
+#include "kernel/memory.h"
 #include "kernel/pthread.h"
 #include "libs/errno.h"

@@ -97,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] ||
@@ -111,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))) {
@@ -125,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;
@@ -183,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;
 }
@@ -255,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;
@@ -298,6 +313,7 @@ void RenderExecutor::DispatchDirect(uint64_t submit_id, CommandBuffer& buffer,
 	}
 	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;
@@ -308,8 +324,8 @@ void RenderExecutor::DispatchDirect(uint64_t submit_id, CommandBuffer& buffer,
 	const auto& resources = *input_info.stage.resources;
 	if (resources.specialization_reads.empty() &&
 	    (TryConsumeComputeMetaClear(input_info, buffer) ||
-	     TryConsumeComputeImageClear(input_info, buffer, thread_group_x, thread_group_y,
-	                                 thread_group_z, mode))) {
+	     TryConsumeComputeFill(input_info, buffer, thread_group_x, thread_group_y,
+	                           thread_group_z, mode))) {
 		ResetBindings();
 		return;
 	}
@@ -376,11 +392,12 @@ void RenderExecutor::DispatchDirect(uint64_t submit_id, CommandBuffer& buffer,
 		    });
 		const uint32_t frame_num = static_cast<uint32_t>(m_context.GetGpu().GetFrameNum());
 		LOGF("GraphicsRenderDispatchDirect: frame=%u shader=0x%016" PRIx64
-		     " hash=0x%016" PRIx64 " tick=%" PRIu64
+		     " hash=0x%016" PRIx64 " tick=%" PRIu64 " completed_tick=%" PRIu64
 		     " groups=%ux%ux%u mode=0x%08" PRIx32 " local=%ux%ux%u "
 		     "buffers=%zu textures=%zu sampled=%zu storage=%zu samplers=%zu push=%u\n",
 		     frame_num, sh_ctx.GetCs().cs_regs.data_addr, program.shader_hash,
 		     m_context.GetCommandScheduler().CurrentTick(),
+		     m_context.GetCommandScheduler().GetMasterSemaphore().KnownGpuTick(),
 		     thread_group_x, thread_group_y, thread_group_z, mode,
 		     input_info.threads_num[0], input_info.threads_num[1],
 		     input_info.threads_num[2], program.info.buffers.size(), program.info.images.size(),
@@ -430,7 +447,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);
@@ -446,14 +484,22 @@ void RenderExecutor::DispatchIndirect(uint64_t submit_id, CommandBuffer& buffer,
 	                    static_cast<uint32_t>(args_addr), static_cast<uint32_t>(args_addr >> 32u),
 	                    0, mode, buffer.GetShaders().GetCs().cs_regs.data_addr);
 	Common::LockGuard lock(m_context.GetMutex());
+	vk::DispatchIndirectCommand args {};
+	if (!m_context.GetBufferCache().HasGpuDirtyBytes(args_addr, sizeof(args)) &&
+	    LibKernel::Memory::TryReadBacking(args_addr, &args, sizeof(args)) &&
+	    (args.x == 0 || args.y == 0 || args.z == 0)) {
+		return;
+	}
 	const auto& cs_regs = buffer.GetShaders().GetCs();
 	if (cs_regs.cs_regs.data_addr == 0) {
 		return;
 	}
 	ShaderComputeInputInfo input_info {};
 	input_info.async_compute = async_compute;
+	input_info.indirect_args_address = args_addr;
 	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;
@@ -493,6 +539,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();
 }
@@ -90,6 +90,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) {
@@ -129,9 +130,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;
 }

--- a/Show More
+++ b/Show More