shader: share formatted indirect buffer load emission

ngs2: accept larger custom voice state buffers
Co-authored-by: Frenzied <frenziedsaint@gmail.com>
2026-10-29 16:51:20 +00:00 · 2026-10-29 05:52:04 +01:00 · 2026-10-29 05:52:04 +01:00 · 2026-10-29 05:52:04 +01:00 · 2026-10-29 05:52:04 +01:00 · 2026-10-29 05:52:04 +01:00
78 changed files with 4502 additions and 2194 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 host_input_tests ime_dialog_tests lru_cache_tests virtual_memory_allocation_tests ^
+            audio_out2_port_tests ngs2_sampler_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|host_input|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|ngs2_sampler|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 host_input_tests ime_dialog_tests lru_cache_tests virtual_memory_allocation_tests \
+            --target launcher audio_out2_port_tests ngs2_sampler_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|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)$'
+            -R '^(audio_out2_port|ngs2_sampler|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 host_input_tests ime_dialog_tests lru_cache_tests virtual_memory_allocation_tests \
+              audio_out2_port_tests ngs2_sampler_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|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)$'
+            -R '^(audio_out2_port|ngs2_sampler|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
@@ -42,6 +42,28 @@ if(NOT KYTY_CLANG_CL)
 endif()
 install(FILES "${openxr_SOURCE_DIR}/LICENSE" DESTINATION "licenses/openxr")

+# Noto Sans v2.015 (SIL Open Font License 1.1) is libFont's fallback font.
+set(KYTY_NOTO_SANS_REPOSITORY "https://raw.githubusercontent.com/notofonts/notofonts.github.io")
+set(KYTY_NOTO_SANS_COMMIT 42af67ccd12bc0e31ed85384bf73e39248265374)
+set(KYTY_NOTO_SANS_BASE_URL "${KYTY_NOTO_SANS_REPOSITORY}/${KYTY_NOTO_SANS_COMMIT}/fonts")
+set(KYTY_NOTO_SANS_DIR "${FETCHCONTENT_BASE_DIR}/noto-sans")
+function(kyty_fetch_noto_sans path sha256)
+	get_filename_component(name "${path}" NAME)
+	file(DOWNLOAD "${KYTY_NOTO_SANS_BASE_URL}/${path}" "${KYTY_NOTO_SANS_DIR}/${name}"
+		EXPECTED_HASH SHA256=${sha256})
+endfunction()
+kyty_fetch_noto_sans(LICENSE f2095b08bed08b23a6fe26112fcd679a2bee3f002eef077eb05d215ed1051bd8)
+kyty_fetch_noto_sans(NotoSans/hinted/ttf/NotoSans-Regular.ttf
+	478c558ea716033cd60c03438f628dfa75694dcf6b5f6d505a2f05fd2b4f3823)
+kyty_fetch_noto_sans(NotoSans/hinted/ttf/NotoSans-Bold.ttf
+	1df075a380fc7cb898acf64c1f7b3b4dd780de3caa860178bf929de35817a913)
+kyty_fetch_noto_sans(NotoSans/hinted/ttf/NotoSans-Italic.ttf
+	467e3f89eeca4108bb8710a2b9e0cf2281ac56d5b0609211a83776d0505eecb5)
+kyty_fetch_noto_sans(NotoSans/hinted/ttf/NotoSans-BoldItalic.ttf
+	1b602a9d6353be42c91df097a4857b69fa2696f26703d7a33b54a15d87c2622c)
+install(FILES "${KYTY_NOTO_SANS_DIR}/LICENSE" DESTINATION "licenses/noto-sans")
+set(KYTY_NOTO_SANS_DIR "${KYTY_NOTO_SANS_DIR}" PARENT_SCOPE)
+
 # 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)
@@ -522,6 +522,13 @@ add_executable(ngs2_sampler_tests EXCLUDE_FROM_ALL
 target_link_libraries(ngs2_sampler_tests common fmt::fmt LibAtrac9)
 target_include_directories(ngs2_sampler_tests PRIVATE ${inc_headers})

+add_executable(ajm_resample_tests EXCLUDE_FROM_ALL
+	"${KYTY_TESTS_DIR}/AjmResampleTests.cpp"
+	"${KYTY_SOURCE_DIR}/loader/timer.cpp"
+)
+target_link_libraries(ajm_resample_tests common fmt::fmt FFmpeg::ffmpeg LibAtrac9 opus)
+target_include_directories(ajm_resample_tests PRIVATE ${inc_headers})
+
 add_executable(pad_haptics_tests EXCLUDE_FROM_ALL
 	"${KYTY_TESTS_DIR}/PadHapticsTests.cpp"
 )
@@ -668,6 +675,7 @@ if(BUILD_TESTING)
 	add_test(NAME audio_out2_port COMMAND $<TARGET_FILE:audio_out2_port_tests>)
 	add_test(NAME audio_out_timing COMMAND $<TARGET_FILE:audio_out_timing_tests>)
 	add_test(NAME ngs2_sampler COMMAND $<TARGET_FILE:ngs2_sampler_tests>)
+	add_test(NAME ajm_resample COMMAND $<TARGET_FILE:ajm_resample_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>)
@@ -684,6 +692,16 @@ if(BUILD_TESTING)
 		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_vop3_conversions
+		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --vop3-conversions-only)
+	add_test(NAME shader_recompiler_quadmask
+		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --quadmask-only)
+	add_test(NAME shader_recompiler_integer16_ternary
+		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --integer16-ternary-only)
+	add_test(NAME shader_recompiler_mad_u16
+		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --mad-u16-only)
+	add_test(NAME shader_recompiler_mul_hi24
+		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --mul-hi24-only)
 	add_test(NAME shader_recompiler_alignbyte
 		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --alignbyte-only)
 	add_test(NAME graphics_pipeline_rasterization
@@ -700,8 +718,6 @@ if(BUILD_TESTING)
 		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --stream-buffer-only)
 	add_test(NAME gpu_command_lane
 		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --gpu-command-lane-only)
-	add_test(NAME wait_until_safe_for_rendering
-		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --wait-until-safe-for-rendering-only)
 	add_test(NAME pm4_flip
 		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --flip-only)
 	add_test(NAME depth_bits
@@ -764,6 +780,7 @@ if(BUILD_TESTING)
 		audio_out2_port_tests
 		audio_out_timing_tests
 		ngs2_sampler_tests
+		ajm_resample_tests
 		pad_haptics_tests
 		controller_settings_tests
 		host_input_tests
@@ -815,6 +832,19 @@ add_custom_target(kyty_touch_cursors
 add_dependencies(kyty_emulator kyty_touch_cursors)
 install(DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}/assets/cursors" DESTINATION assets)

+set(KYTY_NOTO_SANS_FONTS
+	"${KYTY_NOTO_SANS_DIR}/NotoSans-Regular.ttf"
+	"${KYTY_NOTO_SANS_DIR}/NotoSans-Bold.ttf"
+	"${KYTY_NOTO_SANS_DIR}/NotoSans-Italic.ttf"
+	"${KYTY_NOTO_SANS_DIR}/NotoSans-BoldItalic.ttf")
+add_custom_target(kyty_noto_sans_fonts
+	COMMAND ${CMAKE_COMMAND} -E make_directory "$<TARGET_FILE_DIR:kyty_emulator>/assets/fonts"
+	COMMAND ${CMAKE_COMMAND} -E copy_if_different ${KYTY_NOTO_SANS_FONTS}
+		"$<TARGET_FILE_DIR:kyty_emulator>/assets/fonts"
+	VERBATIM)
+add_dependencies(kyty_emulator kyty_noto_sans_fonts)
+install(FILES ${KYTY_NOTO_SANS_FONTS} DESTINATION assets/fonts)
+
 target_link_libraries(kyty_emulator ${kyty_emulator_link_libraries})
 if (WIN32)
 	target_link_libraries(kyty_emulator opengl32)
@@ -105,6 +105,10 @@ bool HideCursorEnabled() {
 	return g_config->hide_cursor_enabled;
 }

+bool StretchEnabled() {
+	return g_config->stretch_enabled;
+}
+
 bool VrEnabled() {
 	return g_config->vr_enabled;
 }
@@ -64,6 +64,7 @@ struct ConfigOptions {
 	bool                   fullscreen_enabled          = false;
 	bool                   hdr_enabled                 = false;
 	bool                   hide_cursor_enabled         = false;
+	bool                   stretch_enabled             = false;
 	bool                   splash_enabled              = true;
 	bool                   vr_enabled                  = false;
 	bool                   amd_cpu_enabled             = false;
@@ -125,6 +126,7 @@ int32_t GetGpuIndex();
 bool     FullscreenEnabled();
 bool     HdrEnabled();
 bool     HideCursorEnabled();
+bool     StretchEnabled();
 bool     VrEnabled();
 bool     AmdCpuEnabled();
 bool     IsNetworkConnected();
@@ -99,10 +99,6 @@ static void Init(const Config::ConfigOptions& cfg, const std::filesystem::path&

 	if (Common::File::IsFileExisting(param_json)) {
 		Loader::SystemContentLoadParamSfo(param_json);
-		if (const auto flexible_memory_size = Loader::SystemContentGetFlexibleMemorySize();
-		    flexible_memory_size != 0) {
-			Libs::LibKernel::Memory::SetFlexibleMemorySize(flexible_memory_size);
-		}
 	}
 	if (cfg.splash_enabled) {
 		Libs::Graphics::InitializeSplashScreen(param_json.parent_path());
@@ -46,6 +46,7 @@ constexpr FormatInfo kFormatInfo[] = {
 	{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},
+	{BufferFormat::k2_10_10_10UNorm, 4, 0, 4, true, false},
 	{BufferFormat::k10_10_10_2UNorm, 4, 0, 4, true, false},
 	{BufferFormat::k10_10_10_2UInt, 4, 0, 4, true, true},
 	{BufferFormat::k8_8_8_8UNorm, 4, 0, 4, true, false},
@@ -30,7 +30,6 @@
 #include <mutex>
 #include <semaphore>
 #include <thread>
-#include <utility>
 #include <vector>

 namespace Libs::Graphics {
@@ -92,12 +91,6 @@ void GuestGpu::SendCommand(Common::UniqueFunction<void>&& command) {
 	m_work_available.Signal();
 }

-void GuestGpu::Wake() {
-	Common::LockGuard lock(m_queue_mutex);
-	m_wake_pending = true;
-	m_work_available.Signal();
-}
-
 void GuestGpu::ProcessCommands() {
 	EXIT_IF(!IsGpuThread());
 	while (m_pending_commands.load(std::memory_order_acquire) != 0) {
@@ -443,10 +436,7 @@ void GuestGpu::ThreadRun(void* data) {
 				}
 				if (selected_queue < 0) {
 					gpu->m_processing = false;
-					// A wake that came before this wait must not be lost.
-					if (!std::exchange(gpu->m_wake_pending, false)) {
-						gpu->m_work_available.WaitFor(&gpu->m_queue_mutex, 100);
-					}
+					gpu->m_work_available.WaitFor(&gpu->m_queue_mutex, 100);
 					for (auto& queue: gpu->m_queues) {
 						if (!queue.empty()) {
 							queue.front().blocked = false;
@@ -814,6 +804,27 @@ static void PatchIndirectDrawOffsets(CommandProcessor& cp, IndirectDrawRegisters
 	write(registers.draw_index, draw_index);
 }

+// Writes the GS user-data slots the packet patches; false when the shaders cannot read the
+// offsets from GPU records: a first-index or subdraw-index destination, or another stage's
+// user data.
+static bool IndirectDrawOffsets(IndirectDrawRegisters registers, ShaderIndirectOffsets& offsets) {
+	if (registers.index_offset != Pm4::SH_NOP || registers.draw_index != Pm4::SH_NOP) {
+		return false;
+	}
+	for (const auto [location, slot]: {std::pair {registers.vertex_offset, &offsets.vertex},
+	                                   std::pair {registers.instance_offset, &offsets.instance}}) {
+		if (location == Pm4::SH_NOP) {
+			continue;
+		}
+		const uint32_t index = location - Pm4::SPI_SHADER_USER_DATA_GS_0;
+		if (index >= HW::UserSgprInfo::SGPRS_MAX) {
+			return false;
+		}
+		*slot = static_cast<uint8_t>(index);
+	}
+	return true;
+}
+
 void CommandProcessor::DrawIndirect(uint32_t data_offset, IndirectDrawRegisters registers,
                                    uint32_t draw_initiator, bool indexed) {
 	const auto args_size =
@@ -845,13 +856,10 @@ void CommandProcessor::DrawIndirectMulti(uint32_t data_offset, uint32_t max_coun
 	// 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 && registers.draw_index == Pm4::SH_NOP &&
-	    (!indexed || m_index_buffer_size != 0) &&
+	// The GPU reads the records if the shaders allow; the VS prolog removes Vulkan's bases
+	// from the index VGPRs and copies them only to the user SGPRs the packet patches.
+	DrawIndirectArgs records {args_addr, draw_count, stride_in_bytes};
+	if (IndirectDrawOffsets(registers, records.offsets) && (!indexed || m_index_buffer_size != 0) &&
 	    (indexed ? executor.DrawIndex(m_submit_id, CurrentBuffer(),
 	                                  {.index_count         = m_index_buffer_size,
 	                                   .index_addr          = IndexAddress(0),
@@ -941,11 +949,10 @@ void CommandProcessor::DrawIndexAuto(DrawAutoArgs args) {
 }

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

 void CommandProcessor::ReleaseMem(uint32_t data_sel, uint32_t interrupt_selector, void* address,
@@ -31,8 +31,6 @@ public:
 	[[nodiscard]] bool IsStopping();
 	void               SendCommand(Common::UniqueFunction<void>&& command);
 	void               SendCommandSync(Common::UniqueFunction<void>&& command);
-	// Retries blocked submissions now instead of after the 100 ms stall timeout.
-	void               Wake();

 	// Submitted command memory is borrowed and must remain valid until GPU execution completes.
 	void Submit(std::span<const uint32_t> commands);
@@ -84,7 +82,6 @@ private:
 	bool                                           m_processing        = false;
 	bool                                           m_accepting         = true;
 	bool                                           m_stopping          = false;
-	bool                                           m_wake_pending      = false;
 	bool                                           m_shutdown_complete = false;
 	// Completion callbacks can outlive GuestGpu during renderer shutdown.
 	std::shared_ptr<std::binary_semaphore> m_suspend_point_ready =
@@ -5,6 +5,7 @@

 #include <algorithm>
 #include <cstdint>
+#include <iterator>
 #include <map>

 namespace Libs::Graphics {
@@ -17,14 +18,19 @@ public:
 		if (it != m_ranges.begin() && std::prev(it)->second >= address) {
 			it = std::prev(it);
 		}
-		uint64_t begin = address;
-		uint64_t last  = end;
-		while (it != m_ranges.end() && it->first <= last) {
-			begin = std::min(begin, it->first);
-			last  = std::max(last, it->second);
-			it    = m_ranges.erase(it);
+		if (it == m_ranges.end() || it->first > address) {
+			it = m_ranges.emplace_hint(it, address, end);
+		} else {
+			if (it->second >= end) {
+				return;
+			}
+			it->second = end;
+		}
+		auto next = std::next(it);
+		while (next != m_ranges.end() && next->first <= it->second) {
+			it->second = std::max(it->second, next->second);
+			next       = m_ranges.erase(next);
 		}
-		m_ranges.emplace(begin, last);
 	}

 	void Subtract(uint64_t address, uint64_t size) {
@@ -488,16 +488,15 @@ void TextureCache::RefreshCopySource(ImageId id) {
 }

 bool TextureCache::CopyD16(Image& destination, Image& source) {
-	const bool source_depth      = source.info.IsDepth();
-	const bool destination_depth = destination.info.IsDepth();
-	if (source_depth == destination_depth) {
+	// Converts guest D16 between a four-byte host depth aspect and two-byte color or D16 depth.
+	const bool source_wide = DepthAspectTransferBytes(source.backing.format) == sizeof(uint32_t);
+	if (source_wide == (DepthAspectTransferBytes(destination.backing.format) == sizeof(uint32_t))) {
 		return false;
 	}
-	auto&      depth          = source_depth ? source : destination;
-	auto&      color          = source_depth ? destination : source;
-	const auto transfer_bytes = DepthAspectTransferBytes(depth.backing.format);
+	auto& depth  = source_wide ? source : destination;
+	auto& narrow = source_wide ? destination : source;
 	if (depth.info.bytes_per_block != sizeof(uint16_t) ||
-	    color.info.bytes_per_block != sizeof(uint16_t) || transfer_bytes != sizeof(uint32_t)) {
+	    narrow.info.bytes_per_block != sizeof(uint16_t)) {
 		return false;
 	}
 	EXIT_IF(source.backing.samples != 1 || destination.backing.samples != 1 ||
@@ -507,42 +506,42 @@ bool TextureCache::CopyD16(Image& destination, Image& source) {

 	const auto     layers = depth.info.resources.layers;
 	const uint64_t depth_slice =
-	    static_cast<uint64_t>(depth.info.pitch) * depth.info.extent.height * transfer_bytes;
-	const uint64_t color_slice =
-	    static_cast<uint64_t>(color.info.pitch) * color.info.extent.height * sizeof(uint16_t);
-	EXIT_IF(layers == 0 || depth_slice > UINT64_MAX / layers || color_slice > UINT64_MAX / layers);
-	const auto                       depth_size = depth_slice * layers;
-	const auto                       color_size = color_slice * layers;
-	const std::array                 depth_copies {
+	    static_cast<uint64_t>(depth.info.pitch) * depth.info.extent.height * sizeof(uint32_t);
+	const uint64_t narrow_slice =
+	    static_cast<uint64_t>(narrow.info.pitch) * narrow.info.extent.height * sizeof(uint16_t);
+	EXIT_IF(layers == 0 || depth_slice > UINT64_MAX / layers || narrow_slice > UINT64_MAX / layers);
+	const std::array depth_copies {
 	    vk::BufferImageCopy {0,
 	                         depth.info.pitch,
 	                         depth.info.extent.height,
 	                         {vk::ImageAspectFlagBits::eDepth, 0, 0, layers},
 	                         {},
 	                         depth.info.extent}};
-	const std::array color_copies {
+	const std::array narrow_copies {
 	    vk::BufferImageCopy {0,
-	                         color.info.pitch,
-	                         color.info.extent.height,
-	                         {vk::ImageAspectFlagBits::eColor, 0, 0, layers},
+	                         narrow.info.pitch,
+	                         narrow.info.extent.height,
+	                         {narrow.info.IsDepth() ? vk::ImageAspectFlagBits::eDepth
+	                                                : vk::ImageAspectFlagBits::eColor,
+	                          0, 0, layers},
 	                         {},
-	                         color.info.extent}};
+	                         narrow.info.extent}};

-	auto                         depth_buffer = m_tiler.GetScratchBuffer(depth_size);
-	auto                         color_buffer = m_tiler.GetScratchBuffer(color_size, depth_buffer.buffer);
+	auto depth_buffer  = m_tiler.GetScratchBuffer(depth_slice * layers);
+	auto narrow_buffer = m_tiler.GetScratchBuffer(narrow_slice * layers, depth_buffer.buffer);
 	const TileManager::D16Layout promote_layout {
 	    .width               = depth.info.extent.width,
 	    .height              = depth.info.extent.height,
 	    .layers              = layers,
-	    .source_row_stride   = static_cast<uint64_t>(color.info.pitch) * sizeof(uint16_t),
-	    .target_row_stride   = static_cast<uint64_t>(depth.info.pitch) * transfer_bytes,
-	    .source_slice_stride = color_slice,
+	    .source_row_stride   = static_cast<uint64_t>(narrow.info.pitch) * sizeof(uint16_t),
+	    .target_row_stride   = static_cast<uint64_t>(depth.info.pitch) * sizeof(uint32_t),
+	    .source_slice_stride = narrow_slice,
 	    .target_slice_stride = depth_slice,
 	};
 	const bool d32 = DepthAspectTransferFormat(depth.backing.format) == vk::Format::eD32Sfloat;
-	if (source_depth) {
+	if (source_wide) {
 		source.Download(depth_copies, depth_buffer.buffer, depth_buffer.offset, depth_buffer.size);
-		m_tiler.ConvertD16(depth_buffer, color_buffer, TileManager::D16Direction::Demote, d32,
+		m_tiler.ConvertD16(depth_buffer, narrow_buffer, TileManager::D16Direction::Demote, d32,
 		                   {.width               = promote_layout.width,
 		                    .height              = promote_layout.height,
 		                    .layers              = promote_layout.layers,
@@ -550,11 +549,12 @@ bool TextureCache::CopyD16(Image& destination, Image& source) {
 		                    .target_row_stride   = promote_layout.source_row_stride,
 		                    .source_slice_stride = promote_layout.target_slice_stride,
 		                    .target_slice_stride = promote_layout.source_slice_stride});
-		destination.Upload(color_copies, color_buffer.buffer, color_buffer.offset,
-		                   color_buffer.size);
+		destination.Upload(narrow_copies, narrow_buffer.buffer, narrow_buffer.offset,
+		                   narrow_buffer.size);
 	} else {
-		source.Download(color_copies, color_buffer.buffer, color_buffer.offset, color_buffer.size);
-		m_tiler.ConvertD16(color_buffer, depth_buffer, TileManager::D16Direction::Promote, d32,
+		source.Download(narrow_copies, narrow_buffer.buffer, narrow_buffer.offset,
+		                narrow_buffer.size);
+		m_tiler.ConvertD16(narrow_buffer, depth_buffer, TileManager::D16Direction::Promote, d32,
 		                   promote_layout);
 		destination.Upload(depth_copies, depth_buffer.buffer, depth_buffer.offset,
 		                   depth_buffer.size);
@@ -17,9 +17,15 @@

 namespace Libs::Graphics {

-enum class ColorTransform : uint32_t { None, SwapBgra16, Reverse10_11_11 };
+enum class ColorTransform : uint32_t { None, SwapBgra16, Reverse10_11_11, Rotate2_10_10_10 };

-enum class VideoOutCompression : uint8_t { Uncompressed, Dcc256_256_0, Dcc256_64_64, Unsupported };
+enum class VideoOutCompression : uint8_t {
+	Uncompressed,
+	Dcc256_256_0,
+	Dcc256_64_64,
+	Dcc128_128_128,
+	Unsupported
+};

 enum class ImageMetadataKind : uint8_t { None, Htile, Dcc, Cmask };

@@ -99,6 +105,8 @@ struct ImageInfo {
 		switch (guest_format) {
 			case Prospero::BufferFormat::k10_11_11Float:
 				return ColorTransform::Reverse10_11_11;
+			case Prospero::BufferFormat::k2_10_10_10UNorm:
+				return ColorTransform::Rotate2_10_10_10;
 			default: return ColorTransform::None;
 		}
 	}
@@ -287,8 +295,9 @@ inline bool ImageInfo::IsDepth() const noexcept {
 [[nodiscard]] inline VideoOutCompression
 ClassifyVideoOutCompression(bool compressed, uint64_t metadata_address, uint32_t dcc_control,
                            uint64_t dcc_clear_color) noexcept {
-	constexpr uint32_t DCC_256_256_0 = 0x00000048u;
-	constexpr uint32_t DCC_256_64_64 = 0x00000208u;
+	constexpr uint32_t DCC_256_256_0   = 0x00000048u;
+	constexpr uint32_t DCC_256_64_64   = 0x00000208u;
+	constexpr uint32_t DCC_128_128_128 = 0x00100024u;
 	if (!compressed) {
 		return metadata_address == 0 && dcc_control == 0 && dcc_clear_color == 0
 		           ? VideoOutCompression::Uncompressed
@@ -300,6 +309,7 @@ ClassifyVideoOutCompression(bool compressed, uint64_t metadata_address, uint32_t
 	switch (dcc_control) {
 		case DCC_256_256_0: return VideoOutCompression::Dcc256_256_0;
 		case DCC_256_64_64: return VideoOutCompression::Dcc256_64_64;
+		case DCC_128_128_128: return VideoOutCompression::Dcc128_128_128;
 		default: return VideoOutCompression::Unsupported;
 	}
 }
@@ -32,8 +32,12 @@ struct HostFormatInfo {

 HostFormatInfo ResolveHostFormat(Prospero::BufferFormat guest_format,
                                 Prospero::ChannelOrder order) {
-	if (guest_format == Prospero::BufferFormat::k10_11_11Float) {
-		return {vk::Format::eB10G11R11UfloatPack32, Prospero::ColorMappingBgra};
+	switch (guest_format) {
+		case Prospero::BufferFormat::k10_11_11Float:
+			return {vk::Format::eB10G11R11UfloatPack32, Prospero::ColorMappingBgra};
+		case Prospero::BufferFormat::k2_10_10_10UNorm:
+			return {vk::Format::eA2R10G10B10UnormPack32, Prospero::ColorMappingAbgr};
+		default: break;
 	}
 	if (order == Prospero::ChannelOrder::kAlt) {
 		switch (guest_format) {
@@ -104,7 +104,8 @@ TileManager::Tiling TileManager::Prepare(bool tile, uint64_t tiled_capacity,
 	    block.block_depth == 0 || block.block_size == 0 || info.tiling.texel_width == 0 ||
 	    info.tiling.texel_height == 0 || info.tiled_slice_stride == 0);
 	EXIT_IF((transform == ColorTransform::SwapBgra16 && block.bytes_per_element != 8u) ||
-	        (transform == ColorTransform::Reverse10_11_11 && block.bytes_per_element != 4u));
+	        ((transform == ColorTransform::Reverse10_11_11 ||
+	          transform == ColorTransform::Rotate2_10_10_10) && block.bytes_per_element != 4u));
 	const auto checked_multiply = [](uint64_t left, uint64_t right, uint64_t& result) {
 		return (left == 0 || right <= UINT64_MAX / left) && (result = left * right, true);
 	};
@@ -560,7 +561,8 @@ void TileManager::TransformColor(Result input, Result output, ColorTransform tra
 	uint32_t element_bytes = 0;
 	switch (transform) {
 		case ColorTransform::SwapBgra16: element_bytes = 8; break;
-		case ColorTransform::Reverse10_11_11: element_bytes = 4; break;
+		case ColorTransform::Reverse10_11_11:
+		case ColorTransform::Rotate2_10_10_10: element_bytes = 4; break;
 		default: EXIT("invalid color transform\n");
 	}
 	EXIT_IF(input.size == 0 || input.size % element_bytes != 0 ||
@@ -20,7 +20,7 @@ namespace {

 namespace IR = ShaderRecompiler::IR;

-// Smaller tables are cheap to materialize on each draw and are often transient.
+// Smaller tables are often written for one draw, so they enter on a second draw of a frame.
 constexpr uint64_t MinTableBytes = 64ull * 1024ull;
 constexpr uint32_t ElementMask   = (1u << IR::ImageHeapKindShift) - 1u;

@@ -133,12 +133,16 @@ ImageHeap::~ImageHeap() {
 bool ImageHeap::Acquire(const IR::ImageHeapRequest& request, IR::ImageHeapView& result) {
 	KYTY_PROFILER_FUNCTION();
 	const auto size = request.size & ~uint64_t {3};
-	if (size < MinTableBytes || request.count == 0u || request.step == 0u ||
+	if (request.count == 0u || request.step == 0u ||
 	    ((request.base | request.first | request.step) & 3u) != 0u) {
 		return false;
 	}
 	auto memory = std::ranges::find_if(
 	    m_memories, [&](const auto& candidate) { return candidate->base == request.base; });
+	if (memory == m_memories.end() && size < MinTableBytes &&
+	    m_sightings.try_emplace(request.base, m_draw).first->second == m_draw) {
+		return false;
+	}
 	// A table follows its descriptor's size, as the first program of a draw acquires it.
 	if (memory == m_memories.end() || ((*memory)->size != size && (*memory)->draw != m_draw)) {
 		if (!m_context.IsMapped(request.base, size)) return false;
@@ -593,26 +597,42 @@ void ImageHeap::UnmapMemory(uint64_t address, uint64_t size) {
 	} else {
 		m_unmapped.push_back({address, size});
 	}
-	auto&           scheduler = m_context.GetCommandScheduler();
 	std::scoped_lock lock {m_mutex};
 	std::erase_if(m_memories, [&](const std::unique_ptr<Memory>& memory) {
 		if (address >= memory->base + memory->size || memory->base >= address + size) {
 			return false;
 		}
-		Unwatch(*memory);
-		for (auto& view: memory->views) {
-			std::erase(m_active, view.get());
-			for (const auto& [value, index]: view->elements) {
-				Unreference(m_elements[index]);
-				view->released.push_back(index);
-			}
-			Retire(*view);
-			if (scheduler.Active()) {
-				scheduler.DeferOperation([buffer = std::move(view->buffer)] {});
-			}
-		}
+		Drop(*memory);
 		return true;
 	});
 }

+void ImageHeap::RunGarbageCollector() {
+	// A smaller table rebuilt each frame moves to a new base. Larger ones keep their dead words.
+	std::scoped_lock lock {m_mutex};
+	std::erase_if(m_memories, [&](const std::unique_ptr<Memory>& memory) {
+		if (memory->size >= MinTableBytes || memory->draw > m_suspend_draw) return false;
+		Drop(*memory);
+		return true;
+	});
+	m_suspend_draw = m_draw;
+	m_sightings.clear();
+}
+
+void ImageHeap::Drop(Memory& memory) {
+	auto& scheduler = m_context.GetCommandScheduler();
+	Unwatch(memory);
+	for (auto& view: memory.views) {
+		std::erase(m_active, view.get());
+		for (const auto& [value, index]: view->elements) {
+			Unreference(m_elements[index]);
+			view->released.push_back(index);
+		}
+		Retire(*view);
+		if (scheduler.Active()) {
+			scheduler.DeferOperation([buffer = std::move(view->buffer)] {});
+		}
+	}
+}
+
 } // namespace Libs::Graphics
@@ -26,7 +26,7 @@ struct ImageHeapTestAccess;
 // of a table owns device slot words, one per table dword, that select an array element. CPU
 // writes to a table fault in its snapshot, which each draw refreshes once; a view takes the
 // changes when a draw first acquires it. Changes to the images behind elements bump
-// Image::heap_epoch.
+// Image::heap_epoch. A smaller table leaves at a suspend point if the frame did not acquire it.
 class ImageHeap final: public ShaderRecompiler::IR::ImageHeapSource {
 public:
 	ImageHeap(RenderContext& context, PageManager& page_manager);
@@ -47,6 +47,8 @@ public:
 	// CPU write fault, on any thread.
 	void InvalidateMemory(uint64_t address, uint64_t size);
 	void UnmapMemory(uint64_t address, uint64_t size);
+	// At a suspend point.
+	void RunGarbageCollector();

 private:
 	using DescriptorValue = ShaderRecompiler::IR::DescriptorValue;
@@ -107,6 +109,7 @@ private:
 	template <bool track>
 	void     Watch(const Memory& memory, size_t begin, size_t end);
 	void     Unwatch(const Memory& memory);
+	void     Drop(Memory& memory);
 	void     UpdateProbes(View& view, uint64_t begin, uint64_t end);
 	void     UpdateProbe(View& view, uint64_t probe);
 	void     Assign(View& view, size_t slot, uint32_t entry);
@@ -125,9 +128,11 @@ private:
 	std::mutex                           m_mutex; // Fault-visible Memory state and m_memories.
 	std::vector<std::unique_ptr<Memory>> m_memories;
 	std::vector<View*>                   m_active; // Acquired since the last commit.
-	uint64_t                             m_draw = 1;
+	uint64_t                             m_draw         = 1;
+	uint64_t                             m_suspend_draw = 0; // m_draw at the last suspend point.
 	std::vector<Element>                 m_elements;
 	std::vector<uint32_t>                m_free;
+	std::unordered_map<uint64_t, uint64_t>    m_sightings; // Smaller tables: base, first draw.
 	std::deque<std::pair<uint64_t, uint32_t>> m_retired; // Tick, element.
 	std::vector<std::pair<size_t, size_t>>    m_runs;
 	std::vector<uint32_t>                     m_scratch;
@@ -687,6 +687,10 @@ void PipelineCache::GetGraphicsPrograms(const CommandBuffer& command, DrawRender
 		vertex_params[0] = PrepareProgram(vertex_regs, context, user_config, vertex_info[0]);
 	}
 	const bool mesh_active = vertex_info[0].logical_stage == ShaderType::Mesh;
+	if (!tess_active && !mesh_active) {
+		vertex_info[0].indirect_offsets = state.indirect_offsets;
+		vertex_info[0].indirect_draw    = state.indirect_draw;
+	}
 	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 ||
@@ -137,6 +137,9 @@ struct DrawRenderState {
 	bool                                 ps_active                                = false;
 	std::array<ShaderVertexInputInfo, 3> vertex_info;
 	ShaderPixelInputInfo                 ps_input_info;
+	// Input: the user-data slots a GPU-read indirect draw patches.
+	ShaderIndirectOffsets                indirect_offsets;
+	bool                                 indirect_draw = false;

 	[[nodiscard]] std::span<const ShaderVertexInputInfo> VertexStages() const {
 		return {vertex_info.data(), vertex_info[1].stage ? 3u : 1u};
@@ -45,9 +45,10 @@ enum class DrawOffsetSource : uint8_t {

 // 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;
+	uint64_t              address = 0;
+	uint32_t              count   = 0;
+	uint32_t              stride  = 0;
+	ShaderIndirectOffsets offsets;
 };

 struct DrawIndexArgs {
@@ -147,6 +147,7 @@ void RenderContext::RunGarbageCollector() {
 		m_fault_process_pending = false;
 		m_buffer_cache.ProcessFaultBuffer();
 	}
+	m_image_heap.RunGarbageCollector();
 	m_texture_cache.ProcessDownloadImages();
 	m_texture_cache.RunGarbageCollector();
 	m_buffer_cache.RunGarbageCollector();
@@ -52,7 +52,7 @@ std::pair<int32_t, uint32_t> ResolveDrawOffsets(uint32_t index_offset,
 	                                           const ShaderVertexInputInfo& vs_input_info) {
 	auto     vertex_offset   = static_cast<int32_t>(index_offset);
 	uint32_t instance_offset = 0;
-	if (!vs_input_info.fetch_embedded) {
+	if (!vs_input_info.has_dynamic_vertex_fetch) {
 		return {vertex_offset, instance_offset};
 	}

@@ -1250,24 +1250,36 @@ void RenderExecutor::ExecutePreparedDraw(uint64_t submit_id, CommandBuffer& buff
 	}
 }

-// Vulkan offsets for a prepared draw, or nullopt when its indirect records need CPU counts, as
-// mesh and quad-list draws do. Indirect records carry their own offsets.
+// Vulkan offsets for a prepared draw, or nullopt when its indirect records need the CPU: mesh
+// and quad-list draws for counts, tessellation draws for the index/SGPR prolog. Indirect records
+// carry their own offsets.
 static std::optional<DrawEmitInfo> ResolveDrawEmit(const HW::UserConfig&        ucfg,
                                                   const ShaderVertexInputInfo& vs,
                                                   DrawOffsetSource             source,
                                                   int32_t                      first_vertex,
                                                   uint32_t                     first_instance,
                                                   const DrawIndirectArgs&      indirect) {
+	if ((indirect.address != 0 || source == DrawOffsetSource::IndirectArgs) &&
+	    ucfg.GetIndexOffset() != 0) {
+		static std::once_flag warning_once;
+		std::call_once(warning_once, [&] {
+			Log::WriteToConsoleAndLog(fmt::format(
+			    "GPU: indirect draw ignores GE_INDX_OFFSET=0x{:08x}; geometry may be incorrect.\n",
+			    ucfg.GetIndexOffset()));
+		});
+	}
 	if (indirect.address != 0) {
-		if (vs.stage.module->info.stage == ShaderType::Mesh ||
-		    ucfg.GetPrimType() == Prospero::PrimitiveType::kQuadListLegacy) {
+		if (!vs.indirect_draw ||
+		    vs.stage.module->info.stage == ShaderType::Mesh ||
+		    ucfg.GetPrimType() == Prospero::PrimitiveType::kQuadListLegacy ||
+		    vs.indirect_offsets != indirect.offsets) {
 			return std::nullopt;
 		}
 		return DrawEmitInfo {.indirect = indirect};
 	}
 	if (source == DrawOffsetSource::IndirectArgs) {
 		// Native fetches consume the patched SGPRs; rewritten fetches need Vulkan offsets.
-		return vs.fetch_embedded ? DrawEmitInfo {first_vertex, first_instance} : DrawEmitInfo {};
+		return vs.has_dynamic_vertex_fetch ? DrawEmitInfo {first_vertex, first_instance} : DrawEmitInfo {};
 	}
 	const auto [vertex_offset, instance_offset] = ResolveDrawOffsets(ucfg.GetIndexOffset(), vs);
 	return DrawEmitInfo {vertex_offset + first_vertex, instance_offset};
@@ -1324,7 +1336,8 @@ bool RenderExecutor::DrawIndex(uint64_t submit_id, CommandBuffer& buffer,

 	hw_check(buffer);

-	DrawRenderState state {};
+	DrawRenderState state {.indirect_offsets = args.indirect.offsets,
+	                       .indirect_draw    = args.indirect.address != 0};
 	if (!GetDrawTopology(ucfg, state.topology)) {
 		return true;
 	}
@@ -1429,7 +1442,8 @@ bool RenderExecutor::DrawAuto(uint64_t submit_id, CommandBuffer& buffer, const D
 	const DrawCallInfo draw {CommandBufferDebugOp::DrawIndexAuto,
 	                         args.vertex_count, args.instance_count, args.first_instance};

-	DrawRenderState state {};
+	DrawRenderState state {.indirect_offsets = args.indirect.offsets,
+	                       .indirect_draw    = args.indirect.address != 0};
 	if (!GetDrawTopology(ucfg, state.topology)) {
 		ResetBindings();
 		return true;
@@ -1,4 +1,4 @@
-// ColorTransform values match imageInfo.h. Preserve the floating-point bit fields exactly.
+// ColorTransform values match imageInfo.h. Preserve the component bit fields exactly.
 uvec2 transform_color(uvec2 value, uint transform, bool to_host) {
 	switch (transform) {
 		case 1u:
@@ -15,6 +15,13 @@ uvec2 transform_color(uvec2 value, uint transform, bool to_host) {
 			}
 			break;
 		}
+		case 3u:
+			if (to_host) {
+				value.x = (value.x >> 2u) | (value.x << 30u);
+			} else {
+				value.x = (value.x << 2u) | (value.x >> 30u);
+			}
+			break;
 	}
 	return value;
 }
@@ -533,16 +533,16 @@ void DrawMemoryAndIo(const Snapshot& shown) {
 		    Bytes(guest.pooled_allocated) + " (" + Bytes(guest.pool_committed) + " committed)");
 		Row("Flexible", Bytes(guest.flexible_used) + " / " + Bytes(guest.flexible_total));
 		Row("Extended", Bytes(guest.automatic_allocated));
-		// Past the pool size the guest sees no page-table entries available.
-		const auto entries = std::max(guest.cpu_page_entries, guest.gpu_page_entries);
+		// Past its total a side sees no page-table entries available.
+		const auto percent = std::max(guest.cpu_page_entries * 100 / guest.cpu_page_entries_total,
+		                              guest.gpu_page_entries * 100 / guest.gpu_page_entries_total);
 		Row("Page table (2 MB)",
-		    Fmt("CPU %llu  GPU %llu  / %llu",
+		    Fmt("CPU %llu / %llu  GPU %llu / %llu",
 		        static_cast<unsigned long long>(guest.cpu_page_entries),
+		        static_cast<unsigned long long>(guest.cpu_page_entries_total),
 		        static_cast<unsigned long long>(guest.gpu_page_entries),
-		        static_cast<unsigned long long>(guest.page_entries_total)),
-		    entries > guest.page_entries_total            ? BAD
-		    : entries * 10 > guest.page_entries_total * 9 ? MEDIUM
-		                                                  : GOOD);
+		        static_cast<unsigned long long>(guest.gpu_page_entries_total)),
+		    LevelColor(static_cast<double>(percent), 90, 100));
 		ImGui::EndTable();
 	}
 	if (BeginSection("Emulator process")) {
@@ -221,7 +221,7 @@ public:
 	void WaitForSubmitSlot(VideoOutConfig& cfg);
 	bool Flip(uint32_t micros);
 	void GetFlipStatus(VideoOutConfig& cfg, VideoOutFlipStatus& out);
-	bool IsSafeForRendering(VideoOutConfig& cfg, int index);
+	void Wait(VideoOutConfig& cfg, int index);

 private:
 	enum class RequestState { Reserved, Recording, Ready, Presenting };
@@ -240,7 +240,6 @@ private:
 		Graphics::Presenter::Frame* frame;
 		bool                        premultiplied_alpha;
 	};
-	void SignalDoneLocked();

 	Graphics::Presenter& m_presenter;
 	Common::Mutex        m_mutex;
@@ -251,7 +250,6 @@ private:
 	std::list<Request>   m_cpu_requests;
 	std::list<Request>   m_cancelled_requests;
 	bool                 m_processing      = false;
-	bool                 m_gpu_waiting     = false;
 	uint64_t             m_next_request_id = 1;
 };

@@ -955,7 +953,7 @@ bool FlipQueue::Reserve(VideoOutConfig& cfg, int index, int64_t flip_arg, FlipRe
 				}
 			}
 		}
-		SignalDoneLocked();
+		m_done_cond_var.SignalAll();
 	}
 	if (cfg.flip_status.flipPendingNum >= VIDEO_OUT_FLIP_QUEUE_CAPACITY) {
 		Common::DebugCounters::Add(Common::DebugCounters::Counter::FlipsRejected);
@@ -1035,7 +1033,7 @@ void FlipQueue::Cancel(VideoOutConfig& cfg) {
 			it = queue->erase(it);
 		}
 	}
-	SignalDoneLocked();
+	m_done_cond_var.SignalAll();
 	m_submit_slot_cond_var.SignalAll();
 	m_submit_cond_var.SignalAll();
 	m_mutex.Unlock();
@@ -1207,24 +1205,17 @@ void FlipQueue::WaitForSubmitSlot(VideoOutConfig& cfg) {
 	}
 }

-void FlipQueue::SignalDoneLocked() {
-	m_done_cond_var.SignalAll();
-	// A stalled GPU cannot drain and shut down without re-checking under m_mutex, so it is alive here.
-	if (m_gpu_waiting) {
-		m_gpu_waiting = false;
-		m_presenter.Renderer().GetGpu().Wake();
-	}
-}
-
-// A display buffer is safe for rendering once it is neither displayed nor waiting to be.
-bool FlipQueue::IsSafeForRendering(VideoOutConfig& cfg, int index) {
+void FlipQueue::Wait(VideoOutConfig& cfg, int index) {
 	Common::LockGuard lock(m_mutex);
-	auto matches = [&cfg, index](const auto& r) { return r.cfg == &cfg && r.index == index; };
-	const bool safe = cfg.flip_status.currentBuffer != index &&
-	                  std::none_of(m_requests.begin(), m_requests.end(), matches) &&
-	                  std::none_of(m_cpu_requests.begin(), m_cpu_requests.end(), matches);
-	m_gpu_waiting = !safe;
-	return safe;
+
+	auto has_request = [this, &cfg, index] {
+		auto matches = [&cfg, index](const auto& r) { return r.cfg == &cfg && r.index == index; };
+		return std::any_of(m_requests.begin(), m_requests.end(), matches) ||
+		       std::any_of(m_cpu_requests.begin(), m_cpu_requests.end(), matches);
+	};
+	while (has_request()) {
+		m_done_cond_var.Wait(&m_mutex);
+	}
 }

 bool FlipQueue::Flip(uint32_t micros) {
@@ -1320,7 +1311,7 @@ bool FlipQueue::Flip(uint32_t micros) {
 		it = m_requests.erase(it);
 	}
 	m_processing = false;
-	SignalDoneLocked();
+	m_done_cond_var.SignalAll();
 	m_submit_slot_cond_var.SignalAll();
 	m_mutex.Unlock();
 	for (size_t i = count; i != 0; i--) {
@@ -1703,12 +1694,12 @@ void VideoOutDriver::WaitForSubmitSlot(int handle) {
 	m_impl->GetFlipQueue().WaitForSubmitSlot(*cfg);
 }

-bool VideoOutDriver::IsSafeForRendering(int handle, int index) {
+void VideoOutDriver::WaitFlipDone(int handle, int index) {
 	auto* ctx = m_impl->Get(handle);
 	EXIT_IF(ctx == nullptr);

 	EXIT_NOT_IMPLEMENTED(!IsValidBufferIndex(index));
-	return m_impl->GetFlipQueue().IsSafeForRendering(*ctx, index);
+	m_impl->GetFlipQueue().Wait(*ctx, index);
 }

 KYTY_SYSV_ABI int VideoOutGetFlipStatus(int handle, VideoOutFlipStatus* status) {
@@ -38,7 +38,7 @@ public:
 	void CompleteFlip(uint64_t request_id);
 	void SubmitFlipPreparation(uint64_t request_id);
 	void WaitForSubmitSlot(int handle);
-	[[nodiscard]] bool IsSafeForRendering(int handle, int index);
+	void WaitFlipDone(int handle, int index);

 	[[nodiscard]] Impl& State() noexcept;

@@ -290,7 +290,8 @@ public:
 private:
 	vk::SurfaceFormatKHR SelectFormat();
 	void Destroy();
-	void DrawOverlay(vk::CommandBuffer command, const Presenter::Layer& layer);
+	void DrawOverlay(vk::CommandBuffer command, const Presenter::Layer& layer,
+	                 const vk::Rect2D& rect);

 	WindowContext&   m_window;
 	vk::SwapchainKHR m_handle = nullptr;
@@ -637,7 +638,23 @@ bool Swapchain::PrepareSystemOverlay() {
 	return m_system_overlay->PrepareFrame(m_extent, m_format, ImageCount());
 }

-void Swapchain::DrawOverlay(vk::CommandBuffer command, const Presenter::Layer& layer) {
+static vk::Rect2D FitAspect(const vk::Extent3D& source, const vk::Extent2D& target) {
+	const uint64_t src_w = std::max(source.width, 1u);
+	const uint64_t src_h = std::max(source.height, 1u);
+	uint32_t       w     = target.width;
+	uint32_t       h     = target.height;
+	if (src_w * h > src_h * w) {
+		h = std::max<uint32_t>(1, static_cast<uint32_t>(src_h * w / src_w));
+	} else {
+		w = std::max<uint32_t>(1, static_cast<uint32_t>(src_w * h / src_h));
+	}
+	return {{static_cast<int32_t>((target.width - w) / 2),
+	         static_cast<int32_t>((target.height - h) / 2)},
+	        {w, h}};
+}
+
+void Swapchain::DrawOverlay(vk::CommandBuffer command, const Presenter::Layer& layer,
+                            const vk::Rect2D& rect) {
 	auto device = m_window.graphic_ctx.device;
 	if (m_overlay_pipeline == nullptr) {
 		vk::SamplerCreateInfo sampler {};
@@ -756,12 +773,14 @@ void Swapchain::DrawOverlay(vk::CommandBuffer command, const Presenter::Layer& l
 	rendering.colorAttachmentCount = 1;
 	rendering.pColorAttachments    = &attachment;
 	command.beginRendering(&rendering);
-	const vk::Viewport viewport {
-	    0.0f, 0.0f, static_cast<float>(m_extent.width), static_cast<float>(m_extent.height),
-	    0.0f, 1.0f};
-	const vk::Rect2D scissor {{0, 0}, m_extent};
+	const vk::Viewport viewport {static_cast<float>(rect.offset.x),
+	                             static_cast<float>(rect.offset.y),
+	                             static_cast<float>(rect.extent.width),
+	                             static_cast<float>(rect.extent.height),
+	                             0.0f,
+	                             1.0f};
 	command.setViewport(0, 1, &viewport);
-	command.setScissor(0, 1, &scissor);
+	command.setScissor(0, 1, &rect);
 	command.draw(3, 1, 0, 0);
 	command.endRendering();
 }
@@ -799,6 +818,24 @@ void Swapchain::RecordPresentCommands(CommandBuffer& command, Presenter::Frame*
 	                           vk::PipelineStageFlagBits::eTransfer, vk::DependencyFlags {}, 0,
 	                           nullptr, 0, nullptr, 1, &to_transfer);

+	const auto* aspect_source = source != nullptr ? source : overlay.frame;
+	const auto  game_rect     = !Config::StretchEnabled() && aspect_source != nullptr
+	                               ? FitAspect(aspect_source->image.extent, m_extent)
+	                               : vk::Rect2D {{0, 0}, m_extent};
+	if (source == nullptr || game_rect.extent != m_extent) {
+		const vk::ClearColorValue black {};
+		vk_command.clearColorImage(m_images[m_image_index], vk::ImageLayout::eTransferDstOptimal,
+		                           &black, 1, &to_transfer.subresourceRange);
+		if (source != nullptr) {
+			vk::ImageMemoryBarrier after_clear = to_transfer;
+			after_clear.srcAccessMask          = vk::AccessFlagBits::eTransferWrite;
+			after_clear.oldLayout              = vk::ImageLayout::eTransferDstOptimal;
+			vk_command.pipelineBarrier(vk::PipelineStageFlagBits::eTransfer,
+			                           vk::PipelineStageFlagBits::eTransfer,
+			                           vk::DependencyFlags {}, 0, nullptr, 0, nullptr, 1,
+			                           &after_clear);
+		}
+	}
 	if (source != nullptr) {
 		vk::ImageBlit region {};
 		region.srcSubresource.aspectMask     = vk::ImageAspectFlagBits::eColor;
@@ -812,16 +849,14 @@ void Swapchain::RecordPresentCommands(CommandBuffer& command, Presenter::Frame*
 		region.dstSubresource.mipLevel       = 0;
 		region.dstSubresource.baseArrayLayer = 0;
 		region.dstSubresource.layerCount     = 1;
-		region.dstOffsets[1].x               = static_cast<int>(m_extent.width);
-		region.dstOffsets[1].y               = static_cast<int>(m_extent.height);
+		region.dstOffsets[0].x = game_rect.offset.x;
+		region.dstOffsets[0].y = game_rect.offset.y;
+		region.dstOffsets[1].x = game_rect.offset.x + static_cast<int>(game_rect.extent.width);
+		region.dstOffsets[1].y = game_rect.offset.y + static_cast<int>(game_rect.extent.height);
 		region.dstOffsets[1].z               = 1;
 		vk_command.blitImage(source->image.image, vk::ImageLayout::eTransferSrcOptimal,
 		                     m_images[m_image_index], vk::ImageLayout::eTransferDstOptimal, 1,
 		                     &region, vk::Filter::eLinear);
-	} else {
-		const vk::ClearColorValue black {};
-		vk_command.clearColorImage(m_images[m_image_index], vk::ImageLayout::eTransferDstOptimal,
-		                           &black, 1, &to_transfer.subresourceRange);
 	}

 	vk::ImageMemoryBarrier to_present {};
@@ -848,7 +883,7 @@ void Swapchain::RecordPresentCommands(CommandBuffer& command, Presenter::Frame*
 	                           &to_present);
 	if (draw_attachment) {
 		if (draw_overlay) {
-			DrawOverlay(vk_command, overlay);
+			DrawOverlay(vk_command, overlay, game_rect);
 			if (draw_system_overlay) {
 				to_present.srcAccessMask = vk::AccessFlagBits::eColorAttachmentWrite;
 				to_present.oldLayout     = vk::ImageLayout::eColorAttachmentOptimal;
@@ -51,6 +51,7 @@ struct VulkanExtensions {
 vk::PhysicalDeviceVulkan11Features WindowContext::RequiredVulkan11Features() noexcept {
 	vk::PhysicalDeviceVulkan11Features features {};
 	features.storageBuffer16BitAccess = VK_TRUE;
+	features.shaderDrawParameters     = VK_TRUE;
 	return features;
 }

@@ -278,6 +279,8 @@ static void VulkanFindPhysicalDevice(vk::Instance instance, vk::SurfaceKHR surfa
 		              required_features12.samplerMirrorClampToEdge);
 		check_feature(features11.storageBuffer16BitAccess, "storageBuffer16BitAccess",
 		              WindowContext::RequiredVulkan11Features().storageBuffer16BitAccess);
+		check_feature(features11.shaderDrawParameters, "shaderDrawParameters",
+		              WindowContext::RequiredVulkan11Features().shaderDrawParameters);
 		check_feature(features12.storageBuffer8BitAccess, "storageBuffer8BitAccess",
 		              required_features12.storageBuffer8BitAccess);
 		check_feature(features12.timelineSemaphore, "timelineSemaphore",
@@ -814,7 +814,7 @@ TranslateResult TranslateProgram(const Decoder::Program& decoded, const CompileO
 	const auto& cfg = *selected_cfg;
 	Frontend::EmbeddedFetchPlan embedded_fetch;
 	if ((options.stage == ShaderType::Vertex || options.stage == ShaderType::Local) &&
-	    options.input_info.vertex->fetch_embedded) {
+	    options.input_info.vertex->has_dynamic_vertex_fetch) {
 		embedded_fetch = DetectEmbeddedVertexFetch(
 		    decoded, *options.input_info.vertex, options.user_data_base,
 		    static_cast<uint32_t>(options.user_data.size()), options.wave_size);
@@ -156,11 +156,14 @@ void ReflectTessellationLayout(const Decoder::Program& program, bool local,
 						if (rhs.constant == 0u && third.constant == 8u) value = constant(0u);
 					}
 					break;
-				case Opcode::V_AND_B32:
-					if ((lhs.kind == Kind::PackedControlPoint && rhs == constant(0xffu)) ||
-					    (rhs.kind == Kind::PackedControlPoint && lhs == constant(0xffu)))
-						value = constant(0u);
+				case Opcode::V_AND_B32: {
+					const auto mask = lhs.kind == Kind::PackedControlPoint ? rhs
+					                : rhs.kind == Kind::PackedControlPoint ? lhs
+					                                                       : Address {};
+					if (mask == constant(0xffu)) value = constant(0u);
+					if (mask == constant(0x1f00u)) value = {Kind::Affine, 256u, 0u};
 					break;
+				}
 				case Opcode::V_ADD_NC_U32: value = add(lhs, rhs); break;
 				case Opcode::V_MUL_LO_U32: value = multiply(lhs, rhs); break;
 				case Opcode::V_MUL_U32_U24: value = multiply(low24(lhs), low24(rhs)); break;
@@ -63,7 +63,8 @@ uint32_t EmitBuiltinU32(EmitterState& state, IR::StageInputKind kind, uint32_t c
 	}
 	if (kind == IR::StageInputKind::VertexIndex || kind == IR::StageInputKind::InstanceIndex ||
 	    kind == IR::StageInputKind::InvocationId || kind == IR::StageInputKind::PrimitiveId ||
-	    kind == IR::StageInputKind::Layer || kind == IR::StageInputKind::SampleId) {
+	    kind == IR::StageInputKind::Layer || kind == IR::StageInputKind::SampleId ||
+	    kind == IR::StageInputKind::BaseVertex || kind == IR::StageInputKind::BaseInstance) {
 		const auto value = state.builder.AllocateId();
 		const auto bits  = state.builder.AllocateId();
 		state.builder.AddFunction(spv::OpLoad, TypeI32(state), value, variable);
@@ -118,7 +118,7 @@ struct EmitterState {
 	uint32_t                                         bda_pointer_function    = 0;
 	uint32_t                                         indirect_format_function = 0;
 	uint32_t                                         indirect_format_geometry_function = 0;
-	std::map<std::array<uint32_t, 3>, uint32_t>        indirect_buffer_load_functions;
+	std::map<std::array<uint32_t, 4>, uint32_t>        indirect_buffer_load_functions;
 	uint32_t                                         bvh_intersect_function  = 0;
 	uint32_t                                         gds_variable            = 0;
 	uint32_t                                         gds_length              = 0;
@@ -540,7 +540,6 @@ uint32_t EmitF32BitsToF16RtzBits(EmitterState& state, uint32_t bits);
 void EmitProgram(EmitterState& state);

 void DefineGetBdaPointer(EmitterState& state);
-void DefineIndirectBufferFormats(ValueEmitContext& ctx);
 void DefineIndirectBufferLoads(ValueEmitContext& ctx);
 void DefineBvhIntersect(EmitterState& state);
 uint32_t GetBdaPointer(EmitterState& state, uint32_t address);
@@ -1026,53 +1026,50 @@ uint32_t LoadIndirectBuffer(ValueEmitContext& ctx, const IndirectBufferAccess& b
 	return ConstructU32Composite(state, components, values);
 }

-uint32_t LoadIndirectFormattedX(ValueEmitContext& ctx, const IR::Inst& inst) {
+uint32_t LoadIndirectFormattedX(ValueEmitContext& ctx, const IndirectBufferAccess& buffer) {
 	auto& state = ctx.state;
-	return EmitValueOrZeroIfCondition(state, ctx.Arg(inst, inst.NumArgs() - 1), [&]() {
-		const auto buffer = PrepareIndirectBuffer(state, IndirectBufferArgs(ctx, inst), true);
-		const auto geometry = state.builder.AllocateId();
-		state.builder.AddFunction(spv::OpFunctionCall, TypeU32Composite(state, 4), geometry,
-		                          state.indirect_format_geometry_function, buffer.format);
-		const auto field = [&](uint32_t index) {
-			const auto value = state.builder.AllocateId();
-			state.builder.AddFunction(spv::OpCompositeExtract, TypeU32(state), value, geometry, index);
-			return value;
-		};
-		const auto bytes = field(0);
-		const auto component = Binary(
-		    state, spv::OpUMod, TypeU32(state),
-		    Binary(state, spv::OpISub, TypeU32(state), buffer.selector, ConstantU32(state, 4)), field(1));
-		const auto alignment = Unary(state, spv::OpUConvert, TypeU64(state), field(3));
-		const auto base = Binary(state, spv::OpBitwiseAnd, TypeU64(state), buffer.address,
-		                         Unary(state, spv::OpNot, TypeU64(state), alignment));
-		const auto address = Binary(
-		    state, spv::OpIAdd, TypeU64(state), base,
-		    Unary(state, spv::OpUConvert, TypeU64(state),
-		          Binary(state, spv::OpIMul, TypeU32(state), component, field(2))));
-		const auto memory = AndCondition(
-		    state, Binary(state, spv::OpINotEqual, TypeBool(state), bytes, ConstantU32(state, 0)),
-		    Binary(state, spv::OpUGreaterThanEqual, TypeBool(state), buffer.selector,
-		           ConstantU32(state, 4)));
-		const auto raw = EmitValueOrZeroIfCondition(
-		    state, AndCondition(state, memory, IndirectBufferInBounds(state, buffer, 0u, bytes, true)),
-		    [&]() {
-			    // Format alignment and component spacing keep every selected component within one DWORD.
-			    const auto word = LoadBdaDword(
-			        ctx, Binary(state, spv::OpBitwiseAnd, TypeU64(state), address,
-			                    ConstantU64(state, ~uint64_t {3})));
-			    const auto byte = Binary(
-			        state, spv::OpBitwiseAnd, TypeU32(state),
-			        Unary(state, spv::OpUConvert, TypeU32(state), address), ConstantU32(state, 3));
-			    return Binary(state, spv::OpShiftRightLogical, TypeU32(state), word,
-			                  Binary(state, spv::OpShiftLeftLogical, TypeU32(state), byte,
-			                         ConstantU32(state, 3)));
-		    });
-		const auto result = state.builder.AllocateId();
-		state.builder.AddFunction(spv::OpFunctionCall, TypeU32(state), result,
-		                          state.indirect_format_function, buffer.format, buffer.selector,
-		                          component, raw);
-		return result;
-	});
+	const auto geometry = state.builder.AllocateId();
+	state.builder.AddFunction(spv::OpFunctionCall, TypeU32Composite(state, 4), geometry,
+	                          state.indirect_format_geometry_function, buffer.format);
+	const auto field = [&](uint32_t index) {
+		const auto value = state.builder.AllocateId();
+		state.builder.AddFunction(spv::OpCompositeExtract, TypeU32(state), value, geometry, index);
+		return value;
+	};
+	const auto bytes = field(0);
+	const auto component = Binary(
+	    state, spv::OpUMod, TypeU32(state),
+	    Binary(state, spv::OpISub, TypeU32(state), buffer.selector, ConstantU32(state, 4)), field(1));
+	const auto alignment = Unary(state, spv::OpUConvert, TypeU64(state), field(3));
+	const auto base = Binary(state, spv::OpBitwiseAnd, TypeU64(state), buffer.address,
+	                         Unary(state, spv::OpNot, TypeU64(state), alignment));
+	const auto address = Binary(
+	    state, spv::OpIAdd, TypeU64(state), base,
+	    Unary(state, spv::OpUConvert, TypeU64(state),
+	          Binary(state, spv::OpIMul, TypeU32(state), component, field(2))));
+	const auto memory = AndCondition(
+	    state, Binary(state, spv::OpINotEqual, TypeBool(state), bytes, ConstantU32(state, 0)),
+	    Binary(state, spv::OpUGreaterThanEqual, TypeBool(state), buffer.selector,
+	           ConstantU32(state, 4)));
+	const auto raw = EmitValueOrZeroIfCondition(
+	    state, AndCondition(state, memory, IndirectBufferInBounds(state, buffer, 0u, bytes, true)),
+	    [&]() {
+		    // Format alignment and component spacing keep every selected component within one DWORD.
+		    const auto word = LoadBdaDword(
+		        ctx, Binary(state, spv::OpBitwiseAnd, TypeU64(state), address,
+		                    ConstantU64(state, ~uint64_t {3})));
+		    const auto byte = Binary(
+		        state, spv::OpBitwiseAnd, TypeU32(state),
+		        Unary(state, spv::OpUConvert, TypeU32(state), address), ConstantU32(state, 3));
+		    return Binary(state, spv::OpShiftRightLogical, TypeU32(state), word,
+		                  Binary(state, spv::OpShiftLeftLogical, TypeU32(state), byte,
+		                         ConstantU32(state, 3)));
+	    });
+	const auto result = state.builder.AllocateId();
+	state.builder.AddFunction(spv::OpFunctionCall, TypeU32(state), result,
+	                          state.indirect_format_function, buffer.format, buffer.selector,
+	                          component, raw);
+	return result;
 }

 uint32_t LoadRawBuffer(ValueEmitContext& ctx, const IR::Inst& inst, const IR::MemoryInfo& mem,
@@ -1135,15 +1132,12 @@ uint32_t LoadRawBuffer(ValueEmitContext& ctx, const IR::Inst& inst, const IR::Me
 uint32_t LoadBuffer(ValueEmitContext& ctx, const IR::Inst& inst, uint32_t components) {
 	auto& state = ctx.state;
 	const auto mem = ctx.Memory(inst);
-	if (mem.kind == IR::ResourceKind::IndirectBuffer && mem.formatted) {
-		return LoadIndirectFormattedX(ctx, inst);
-	}
 	return EmitValueOrDefaultIfCondition(
 	    state, ctx.Arg(inst, inst.NumArgs() - 1), TypeU32Composite(state, components),
 	    ConstantU32CompositeZero(state, components), [&]() {
 		    if (mem.kind == IR::ResourceKind::IndirectBuffer) {
 			    const auto function = state.indirect_buffer_load_functions.at(
-			        {mem.data_bits, components, uint32_t(mem.coherent)});
+			        {mem.data_bits, components, uint32_t(mem.coherent), uint32_t(mem.formatted)});
 			    const auto args = IndirectBufferArgs(ctx, inst);
 			    const auto result = state.builder.AllocateId();
 			    state.builder.AddFunction(spv::OpFunctionCall, TypeU32Composite(state, components),
@@ -1231,62 +1225,8 @@ void StoreWideShared(ValueEmitContext& ctx, const IR::Inst& inst, uint32_t compo
 	});
 }

-} // namespace
-
-uint32_t GetBdaPointer(EmitterState& state, uint32_t address) {
-	const auto result = state.builder.AllocateId();
-	state.builder.AddFunction(spv::OpFunctionCall, TypeU64(state), result,
-	                          state.bda_pointer_function, address);
-	return result;
-}
-
-void DefineIndirectBufferLoads(ValueEmitContext& ctx) {
-	auto& state = ctx.state;
-	for (const auto* block: state.program.blocks) {
-		for (const auto& inst: *block) {
-			const auto components = IR::BufferComponentCount(inst.GetOpcode());
-			if (components == 0u || inst.GetOpcode() == IR::ValueOpcode::ReadConstBuffer) continue;
-			const auto& memory = ctx.Memory(inst);
-			if (memory.kind != IR::ResourceKind::IndirectBuffer || memory.formatted ||
-			    !memory.SupportsIndirectBufferLoad(inst.GetOpcode())) continue;
-			auto [entry, inserted] = state.indirect_buffer_load_functions.try_emplace(
-			    std::array {memory.data_bits, components, uint32_t(memory.coherent)}, 0u);
-			if (!inserted) continue;
-			const auto function = entry->second = state.builder.AllocateId();
-			const auto type = TypeU32(state);
-			const auto result_type = TypeU32Composite(state, components);
-			const auto signature = state.builder.Type(
-			    spv::OpTypeFunction, result_type, type, type, type, type, type, type, type, type);
-			const auto name = "load_indirect_buffer_" + std::to_string(memory.data_bits) + "x" +
-			                  std::to_string(components) + (memory.coherent ? "_coherent" : "");
-			state.builder.AddName(function, name.c_str());
-			state.builder.AddFunction(spv::OpFunction, result_type, function,
-			                          spv::FunctionControlMaskNone, signature);
-			std::array<uint32_t, 8> args;
-			for (auto& arg: args) {
-				arg = state.builder.AllocateId();
-				state.builder.AddFunction(spv::OpFunctionParameter, type, arg);
-			}
-			EmitLabel(state, state.builder.AllocateId());
-			const auto result = LoadIndirectBuffer(ctx, PrepareIndirectBuffer(state, args, false),
-			                                       memory.data_bits, components, memory.coherent);
-			state.builder.AddFunction(spv::OpReturnValue, result);
-			state.builder.AddFunction(spv::OpFunctionEnd);
-		}
-	}
-}
-
 void DefineIndirectBufferFormats(ValueEmitContext& ctx) {
 	auto& state = ctx.state;
-	const auto needed = std::ranges::any_of(state.program.blocks, [&](const IR::Block* block) {
-		return std::ranges::any_of(*block, [&](const IR::Inst& inst) {
-			if (inst.GetOpcode() != IR::ValueOpcode::LoadBufferU32) return false;
-			const auto& memory = ctx.Memory(inst);
-			return memory.kind == IR::ResourceKind::IndirectBuffer && memory.formatted;
-		});
-	});
-	if (!needed) return;
-
 	const auto type = TypeU32(state);
 	constexpr auto format_count =
 	    static_cast<uint32_t>(Prospero::BufferFormat::k32_32_32_32Float) + 1u;
@@ -1358,6 +1298,57 @@ void DefineIndirectBufferFormats(ValueEmitContext& ctx) {
 	state.builder.AddFunction(spv::OpFunctionEnd);
 }

+} // namespace
+
+uint32_t GetBdaPointer(EmitterState& state, uint32_t address) {
+	const auto result = state.builder.AllocateId();
+	state.builder.AddFunction(spv::OpFunctionCall, TypeU64(state), result,
+	                          state.bda_pointer_function, address);
+	return result;
+}
+
+void DefineIndirectBufferLoads(ValueEmitContext& ctx) {
+	auto& state = ctx.state;
+	for (const auto* block: state.program.blocks) {
+		for (const auto& inst: *block) {
+			const auto components = IR::BufferComponentCount(inst.GetOpcode());
+			if (components == 0u || inst.GetOpcode() == IR::ValueOpcode::ReadConstBuffer) continue;
+			const auto& memory = ctx.Memory(inst);
+			if (memory.kind != IR::ResourceKind::IndirectBuffer ||
+			    !memory.SupportsIndirectBufferLoad(inst.GetOpcode())) continue;
+			auto [entry, inserted] = state.indirect_buffer_load_functions.try_emplace(
+			    std::array {memory.data_bits, components, uint32_t(memory.coherent),
+			                uint32_t(memory.formatted)}, 0u);
+			if (!inserted) continue;
+			if (memory.formatted && state.indirect_format_function == 0u) {
+				DefineIndirectBufferFormats(ctx);
+			}
+			const auto function = entry->second = state.builder.AllocateId();
+			const auto type = TypeU32(state);
+			const auto result_type = TypeU32Composite(state, components);
+			const auto signature = state.builder.Type(
+			    spv::OpTypeFunction, result_type, type, type, type, type, type, type, type, type);
+			const auto name = "load_indirect_buffer_" + std::to_string(memory.data_bits) + "x" +
+			                  std::to_string(components) + (memory.coherent ? "_coherent" : "") +
+			                  (memory.formatted ? "_formatted" : "");
+			state.builder.AddName(function, name.c_str());
+			state.builder.AddFunction(spv::OpFunction, result_type, function,
+			                          spv::FunctionControlMaskNone, signature);
+			std::array<uint32_t, 8> args;
+			for (auto& arg: args) {
+				arg = state.builder.AllocateId();
+				state.builder.AddFunction(spv::OpFunctionParameter, type, arg);
+			}
+			EmitLabel(state, state.builder.AllocateId());
+			const auto buffer = PrepareIndirectBuffer(state, args, memory.formatted);
+			const auto result = memory.formatted ? LoadIndirectFormattedX(ctx, buffer)
+			    : LoadIndirectBuffer(ctx, buffer, memory.data_bits, components, memory.coherent);
+			state.builder.AddFunction(spv::OpReturnValue, result);
+			state.builder.AddFunction(spv::OpFunctionEnd);
+		}
+	}
+}
+
 void DefineGetBdaPointer(EmitterState& state) {
 	if (!state.program.info.uses_dma) {
 		return;
@@ -476,6 +476,8 @@ uint32_t BuiltInForInput(IR::StageInputKind kind) {
 		case IR::StageInputKind::LocalInvocationId: return spv::BuiltInLocalInvocationId;
 		case IR::StageInputKind::LocalInvocationIndex: return spv::BuiltInLocalInvocationIndex;
 		case IR::StageInputKind::GlobalInvocationId: return spv::BuiltInGlobalInvocationId;
+		case IR::StageInputKind::BaseVertex: return spv::BuiltInBaseVertex;
+		case IR::StageInputKind::BaseInstance: return spv::BuiltInBaseInstance;
 		default: return UINT32_MAX;
 	}
 }
@@ -526,7 +528,9 @@ void DefineInputs(EmitterState& state) {
 			case IR::StageInputKind::PrimitiveId:
 			case IR::StageInputKind::InstanceIndex:
 			case IR::StageInputKind::Layer:
-			case IR::StageInputKind::SampleId: type = TypeI32(state); break;
+			case IR::StageInputKind::SampleId:
+			case IR::StageInputKind::BaseVertex:
+			case IR::StageInputKind::BaseInstance: type = TypeI32(state); break;
 			case IR::StageInputKind::WorkgroupId:
 			case IR::StageInputKind::NumWorkgroups:
 			case IR::StageInputKind::LocalInvocationId:
@@ -802,6 +806,10 @@ void DefineModule(EmitterState& state) {
 	if (InputVariableForKind(state, IR::StageInputKind::SampleId) != 0) {
 		state.builder.RequireCapability(spv::CapabilitySampleRateShading);
 	}
+	if (InputVariableForKind(state, IR::StageInputKind::BaseVertex) != 0 ||
+	    InputVariableForKind(state, IR::StageInputKind::BaseInstance) != 0) {
+		state.builder.RequireCapability(spv::CapabilityDrawParameters);
+	}
 	if (state.program.info.image_gather_extended) {
 		state.builder.RequireCapability(spv::CapabilityImageGatherExtended);
 	}
@@ -675,7 +675,6 @@ void EmitProgram(EmitterState& state) {
 	}
 	DefineGetBdaPointer(state);
 	DefineIndirectBufferLoads(ctx);
-	DefineIndirectBufferFormats(ctx);
 	DefineBvhIntersect(state);
 	if (program.info.uses_swizzle) {
 		ctx.scratch_u32_variable = state.builder.AllocateId();
@@ -29,12 +29,14 @@ uint32_t TessellationPointer(ValueEmitContext& ctx, const IR::Inst& inst) {
 	const auto pointer = state.builder.AllocateId();
 	if (kind == Attribute::Factor) {
 		EXIT_NOT_IMPLEMENTED(!inst.Arg(1).IsImmediate());
-		const auto index = inst.Arg(1).U32() / 4u;
-		const auto outer = index < 3u;
-		EXIT_NOT_IMPLEMENTED(index >= 4u);
+		const auto index       = inst.Arg(1).U32() / 4u;
+		const bool quad        = tess.domain == 2u;
+		const auto outer_count = quad ? 4u : 3u;
+		const bool outer       = index < outer_count;
+		EXIT_NOT_IMPLEMENTED(index >= (quad ? 6u : 4u));
 		state.builder.AddFunction(spv::OpAccessChain, TypePointer(state, storage, TypeF32(state)),
 		                          pointer, outer ? variable : state.tess_inner_variable,
-		                          ConstantU32(state, outer ? index : index - 3u));
+		                          ConstantU32(state, outer ? index : index - outer_count));
 		return pointer;
 	}
 	auto address = ctx.Arg(inst, 1);
@@ -128,20 +130,25 @@ void DefineTessellationInterfaces(EmitterState& state) {

 void DefineTessellationExecutionModes(EmitterState& state) {
 	const auto& tess = state.input_info.vertex->tess;
-	// CxVgtTfParam: triangles, integer or fractional-odd partitioning, clockwise output.
-	EXIT_NOT_IMPLEMENTED(tess.domain != 1u ||
-	                     (tess.partitioning != 0u && tess.partitioning != 2u) ||
-	                     tess.output_topology != 2u);
+	// Triangle or quad domain, integer or fractional partitioning, triangle output.
+	EXIT_NOT_IMPLEMENTED((tess.domain != 1u && tess.domain != 2u) ||
+	                     (tess.partitioning != 0u && tess.partitioning != 2u && tess.partitioning != 3u) ||
+	                     (tess.output_topology != 2u && tess.output_topology != 3u));
 	state.builder.RequireCapability(spv::CapabilityTessellation);
 	if (state.program.stage == ShaderType::TessellationControl) {
 		state.builder.AddExecutionMode(state.main_func, spv::ExecutionModeOutputVertices,
 		                               tess.output_control_points);
 	} else {
-		state.builder.AddExecutionMode(state.main_func, spv::ExecutionModeTriangles);
+		state.builder.AddExecutionMode(state.main_func, tess.domain == 1u
+		                                                   ? spv::ExecutionModeTriangles
+		                                                   : spv::ExecutionModeQuads);
 		state.builder.AddExecutionMode(
 		    state.main_func, tess.partitioning == 0u ? spv::ExecutionModeSpacingEqual
-		                                             : spv::ExecutionModeSpacingFractionalOdd);
-		state.builder.AddExecutionMode(state.main_func, spv::ExecutionModeVertexOrderCw);
+		                     : tess.partitioning == 2u ? spv::ExecutionModeSpacingFractionalOdd
+		                                                : spv::ExecutionModeSpacingFractionalEven);
+		state.builder.AddExecutionMode(state.main_func, tess.output_topology == 2u
+		                                                   ? spv::ExecutionModeVertexOrderCw
+		                                                   : spv::ExecutionModeVertexOrderCcw);
 	}
 }

@@ -70,6 +70,7 @@ constexpr OpcodeMap SOP1_OPCODE_LIST[] = {
    {0x21u, Opcode::S_SWAPPC_B64},
    {0x24u, Opcode::S_AND_SAVEEXEC_B64},
    {0x28u, Opcode::S_ORN2_SAVEEXEC_B64},
+    {0x2cu, Opcode::S_QUADMASK_B32},
    {0x2du, Opcode::S_QUADMASK_B64},
    {0x34u, Opcode::S_ABS_I32},
    {0x37u, Opcode::S_ANDN1_SAVEEXEC_B64},
@@ -517,6 +517,7 @@ std::string InstructionToString(const Instruction& inst) {
 		case Opcode::S_NOT_B64:
 		case Opcode::S_WQM_B32:
 		case Opcode::S_WQM_B64:
+		case Opcode::S_QUADMASK_B32:
 		case Opcode::S_QUADMASK_B64:
 		case Opcode::S_AND_SAVEEXEC_B32:
 		case Opcode::S_ORN2_SAVEEXEC_B32:
@@ -63,6 +63,7 @@ enum class Opcode {
 	S_FLBIT_I32,
 	S_FLBIT_I32_I64,
 	S_BITREPLICATE_B64_B32,
+	S_QUADMASK_B32,
 	S_QUADMASK_B64,
 	S_GETPC_B64,
 	S_SETPC_B64,
@@ -248,7 +249,9 @@ enum class Opcode {
 	V_MIN_I32,
 	V_MAX_I32,
 	V_MUL_I32_I24,
+	V_MUL_HI_I32_I24,
 	V_MUL_U32_U24,
+	V_MUL_HI_U32_U24,
 	V_MUL_LO_U32,
 	V_MUL_HI_U32,
 	V_MUL_LO_I32,
@@ -257,6 +260,7 @@ enum class Opcode {
 	V_CVT_PK_U8_F32,
 	V_MAD_F32,
 	V_MAD_I16,
+	V_MAD_U16,
 	V_MAD_I32_I24,
 	V_MAD_U32_U24,
 	V_MAD_U64_U32,
@@ -277,16 +281,20 @@ enum class Opcode {
 	V_MIN3_I32,
 	V_MIN3_U32,
 	V_MIN3_F16,
+	V_MIN3_I16,
 	V_MIN3_U16,
 	V_MAX3_F32,
 	V_MAX3_I32,
 	V_MAX3_U32,
 	V_MAX3_F16,
+	V_MAX3_I16,
+	V_MAX3_U16,
 	V_MED3_F32,
 	V_MED3_I32,
 	V_MED3_U32,
 	V_MED3_F16,
 	V_MED3_I16,
+	V_MED3_U16,
 	V_SAD_U32,
 	V_ADD3_U32,
 	V_LSHL_ADD_U32,
@@ -354,6 +362,7 @@ enum class Opcode {
 	V_ASHRREV_I32,
 	V_LSHLREV_B64,
 	V_LSHRREV_B64,
+	V_ASHRREV_I64,
 	V_LSHLREV_B16,
 	V_LSHRREV_B16,
 	V_ASHRREV_I16,
@@ -35,7 +35,9 @@ constexpr Vop2OpcodeInfo VOP2_OPCODE_LIST[] = {
    {0x05u, Opcode::V_SUBREV_F32},
    {0x08u, Opcode::V_MUL_F32, Vop2SdwaProfile::Float32},
    {0x09u, Opcode::V_MUL_I32_I24, Vop2SdwaProfile::IntegerPartialDestination},
+    {0x0au, Opcode::V_MUL_HI_I32_I24, Vop2SdwaProfile::IntegerPartialDestination},
    {0x0bu, Opcode::V_MUL_U32_U24, Vop2SdwaProfile::IntegerPartialDestination},
+    {0x0cu, Opcode::V_MUL_HI_U32_U24, Vop2SdwaProfile::IntegerPartialDestination},
    {0x0fu, Opcode::V_MIN_F32},
    {0x10u, Opcode::V_MAX_F32},
    {0x11u, Opcode::V_MIN_I32, Vop2SdwaProfile::IntegerFullDestination},
@@ -157,6 +159,10 @@ constexpr OpcodeMap VOP3_ENCODED_VOP1_OPCODE_LIST[] = {
    {0x04u, Opcode::V_CVT_F64_I32},
    {0x0fu, Opcode::V_CVT_F32_F64},
    {0x10u, Opcode::V_CVT_F64_F32},
+    {0x11u, Opcode::V_CVT_F32_UBYTE0},
+    {0x12u, Opcode::V_CVT_F32_UBYTE1},
+    {0x13u, Opcode::V_CVT_F32_UBYTE2},
+    {0x14u, Opcode::V_CVT_F32_UBYTE3},
    {0x16u, Opcode::V_CVT_F64_U32},
    {0x17u, Opcode::V_TRUNC_F64},
    {0x18u, Opcode::V_CEIL_F64},
@@ -167,6 +173,7 @@ constexpr OpcodeMap VOP3_ENCODED_VOP1_OPCODE_LIST[] = {
    {0x07u, Opcode::V_CVT_U32_F32},
    {0x08u, Opcode::V_CVT_I32_F32},
    {0x0au, Opcode::V_CVT_F16_F32},
+    {0x0bu, Opcode::V_CVT_F32_F16},
    {0x0cu, Opcode::V_CVT_RPI_I32_F32},
    {0x0du, Opcode::V_CVT_FLR_I32_F32},
    {0x0eu, Opcode::V_CVT_OFF_F32_I4},
@@ -334,16 +341,20 @@ constexpr OpcodeMap VOP3_OPCODE_LIST[] = {
    {0x152u, Opcode::V_MIN3_I32},
    {0x153u, Opcode::V_MIN3_U32},
    {0x351u, Opcode::V_MIN3_F16},
+    {0x352u, Opcode::V_MIN3_I16},
    {0x353u, Opcode::V_MIN3_U16},
    {0x154u, Opcode::V_MAX3_F32},
    {0x155u, Opcode::V_MAX3_I32},
    {0x156u, Opcode::V_MAX3_U32},
    {0x354u, Opcode::V_MAX3_F16},
+    {0x355u, Opcode::V_MAX3_I16},
+    {0x356u, Opcode::V_MAX3_U16},
    {0x157u, Opcode::V_MED3_F32},
    {0x158u, Opcode::V_MED3_I32},
    {0x159u, Opcode::V_MED3_U32},
    {0x357u, Opcode::V_MED3_F16},
    {0x358u, Opcode::V_MED3_I16},
+    {0x359u, Opcode::V_MED3_U16},
    {0x15du, Opcode::V_SAD_U32},
    {0x15eu, Opcode::V_CVT_PK_U8_F32},
    {0x178u, Opcode::V_XOR3_B32},
@@ -354,6 +365,7 @@ constexpr OpcodeMap VOP3_OPCODE_LIST[] = {
    {0x16cu, Opcode::V_MUL_HI_I32},
    {0x2ffu, Opcode::V_LSHLREV_B64},
    {0x300u, Opcode::V_LSHRREV_B64},
+    {0x301u, Opcode::V_ASHRREV_I64},
    {0x303u, Opcode::V_ADD_NC_U16},
    {0x304u, Opcode::V_SUB_NC_U16},
    {0x305u, Opcode::V_MUL_LO_U16},
@@ -370,6 +382,7 @@ constexpr OpcodeMap VOP3_OPCODE_LIST[] = {
    {0x311u, Opcode::V_PACK_B32_F16},
    {0x314u, Opcode::V_LSHLREV_B16},
    {0x319u, Opcode::V_SUBREV_I32},
+    {0x340u, Opcode::V_MAD_U16},
    {0x345u, Opcode::V_XAD_U32},
    {0x346u, Opcode::V_LSHL_ADD_U32},
    {0x347u, Opcode::V_ADD_LSHL_U32},
@@ -488,8 +501,10 @@ bool IsNativeVop3F16TernaryOpcode(Opcode opcode) {
 }

 bool IsNativeVop3B16TernaryOpcode(Opcode opcode) {
-	return opcode == Opcode::V_MIN3_U16 || opcode == Opcode::V_MED3_I16 ||
-	       opcode == Opcode::V_MAD_I16;
+	return opcode == Opcode::V_MIN3_I16 || opcode == Opcode::V_MIN3_U16 ||
+	       opcode == Opcode::V_MAX3_I16 || opcode == Opcode::V_MAX3_U16 ||
+	       opcode == Opcode::V_MED3_I16 || opcode == Opcode::V_MED3_U16 ||
+	       opcode == Opcode::V_MAD_I16 || opcode == Opcode::V_MAD_U16;
 }

 bool IsNativeVop3B16BinaryOpcode(Opcode opcode) {
@@ -1335,6 +1350,7 @@ uint32_t NativeVop3SourceCount(Opcode opcode) {
 		case Opcode::V_SUB_NC_I16:
 		case Opcode::V_LSHLREV_B64:
 		case Opcode::V_LSHRREV_B64:
+		case Opcode::V_ASHRREV_I64:
 		case Opcode::V_LSHLREV_B16:
 		case Opcode::V_LSHRREV_B16:
 		case Opcode::V_ASHRREV_I16:
@@ -1523,7 +1539,10 @@ bool SupportsNativeVop3ResultModifiers(Opcode opcode) {

 bool SupportsNativeVop3Clamp(Opcode opcode) {
 	return SupportsNativeVop3ResultModifiers(opcode) || UsesInexactClampControl(opcode) ||
-	       opcode == Opcode::V_MAD_I16;
+	       opcode == Opcode::V_MAD_I16 || opcode == Opcode::V_MAD_U16 ||
+	       opcode == Opcode::V_MIN3_I16 || opcode == Opcode::V_MAX3_I16 ||
+	       opcode == Opcode::V_MAX3_U16 || opcode == Opcode::V_MED3_U16 ||
+	       opcode == Opcode::V_MUL_HI_I32_I24 || opcode == Opcode::V_MUL_HI_U32_U24;
 }

 bool HasUnsupportedNativeVop3Modifiers(Opcode opcode, bool permlane, bool mad_mix,
@@ -1554,7 +1573,17 @@ bool HasUnsupportedNativeVop3Modifiers(Opcode opcode, bool permlane, bool mad_mi
 		return clamp != 0u || omod != 0u || neg != 0u;
 	}
 	switch (opcode) {
+		case Opcode::V_MUL_HI_I32_I24:
+		case Opcode::V_MUL_HI_U32_U24:
+			// The high product fits in 32 bits, so clamp has no effect.
+			return abs != 0u || op_sel != 0u || omod != 0u || neg != 0u;
 		case Opcode::V_LDEXP_F32: return (abs & ~1u) != 0u || op_sel != 0u || (neg & ~1u) != 0u;
+		case Opcode::V_CVT_F32_UBYTE0:
+		case Opcode::V_CVT_F32_UBYTE1:
+		case Opcode::V_CVT_F32_UBYTE2:
+		case Opcode::V_CVT_F32_UBYTE3:
+			// Byte conversions ignore input modifiers and support all output modifiers.
+			return op_sel != 0u;
 		case Opcode::V_INTERP_P1LL_F16:
 		case Opcode::V_INTERP_P2_F16:
 			return abs != 0u || (op_sel & 0x7u) != 0u || clamp != 0u || omod != 0u || neg != 0u;
@@ -20,12 +20,12 @@ void Translator::Integer16Binary(const Decoder::Instruction& inst, IR::ValueOpco
 	Write16Bits(DestinationOperand(inst), result);
 }

-void Translator::V_MAD_I16(const Decoder::Instruction& inst) {
-	const auto lhs    = ReadU16AsU32(inst.src0, true);
-	const auto rhs    = ReadU16AsU32(inst.src1, true);
-	auto       result = ir.IAdd(ir.IMul(lhs, rhs), ReadU16AsU32(inst.src2, true));
+void Translator::Integer16Mad(const Decoder::Instruction& inst, bool sign) {
+	const auto lhs    = ReadU16AsU32(inst.src0, sign);
+	const auto rhs    = ReadU16AsU32(inst.src1, sign);
+	auto       result = ir.IAdd(ir.IMul(lhs, rhs), ReadU16AsU32(inst.src2, sign));
 	if (inst.dst.clamp) {
-		result = Saturate16(result, true);
+		result = Saturate16(result, sign);
 	}
 	Write16Bits(DestinationOperand(inst), result);
 }
@@ -147,6 +147,12 @@ void Translator::SimpleInteger(const Decoder::Instruction& inst, IR::ValueOpcode
 		const auto arg_type = IR::ArgTypeOf(opcode, index);
 		const auto& operand = SourceAt(inst, reverse && index < 2u ? 1u - index : index);
 		args[index]         = ReadOperand(operand, arg_type == IR::Type::Void ? type : arg_type);
+		if (opcode == IR::ValueOpcode::ShiftRightArithmetic64 && index == 0u &&
+		    operand.kind == Decoder::OperandKind::LiteralConstant) {
+			args[index] = ir.ConstructU64(
+			    ReadU32(operand),
+			    IR::U32(IR::Value((operand.value & 0x80000000u) ? 0xffffffffu : 0u)));
+		}
 		if (mask_shift_count && index == 1u) {
 			args[index] = ir.BitwiseAnd(IR::U32(args[index]), IR::U32(IR::Value(31u)));
 		}
@@ -263,7 +269,7 @@ void Translator::V_FFBH_32(const Decoder::Instruction& inst, bool sign) {
 	WriteOperand(DestinationOperand(inst), result);
 }

-void Translator::Integer24(const Decoder::Instruction& inst, bool sign, bool addend) {
+void Translator::Integer24(const Decoder::Instruction& inst, bool sign, bool addend, bool high) {
 	const auto extract24 = [&](IR::U32 value) {
 		return IR::U32(
 		    ir.Emit(sign ? IR::ValueOpcode::BitFieldSExtract : IR::ValueOpcode::BitFieldUExtract,
@@ -271,7 +277,9 @@ void Translator::Integer24(const Decoder::Instruction& inst, bool sign, bool add
 	};
 	const auto lhs    = extract24(ReadU32(inst.src0));
 	const auto rhs    = extract24(ReadU32(inst.src1));
-	auto       result = ir.IMul(lhs, rhs);
+	const auto multiply = high ? (sign ? IR::ValueOpcode::SMulHi : IR::ValueOpcode::UMulHi)
+	                           : IR::ValueOpcode::IMul32;
+	auto result = IR::U32(ir.Emit(multiply, {lhs, rhs}));
 	if (addend) {
 		result = ir.IAdd(result, ReadU32(inst.src2));
 	}
@@ -369,7 +377,7 @@ void Translator::S_BITREPLICATE_B64_B32(const Decoder::Instruction& inst) {
 	WriteOperand(DestinationOperand(inst), ir.ConstructU64(replicate(low), replicate(high)));
 }

-void Translator::S_QUADMASK_B64(const Decoder::Instruction& inst) {
+void Translator::S_QUADMASK(const Decoder::Instruction& inst, bool is64) {
 	const auto compact = [&](IR::U32 value) {
 		auto bits = ir.BitwiseOr(value, ir.ShiftRightLogical(value, IR::U32(IR::Value(1u))));
 		bits      = ir.BitwiseOr(bits, ir.ShiftRightLogical(bits, IR::U32(IR::Value(2u))));
@@ -382,12 +390,17 @@ void Translator::S_QUADMASK_B64(const Decoder::Instruction& inst) {
 		    ir.BitwiseOr(bits, ir.ShiftRightLogical(bits, IR::U32(IR::Value(12u)))),
 		    IR::U32(IR::Value(0xffu)));
 	};
-	const auto source = ReadU32Pair(inst.src0);
-	const auto quads  = ir.BitwiseOr(
-	    compact(source[0]), ir.ShiftLeftLogical(compact(source[1]), IR::U32(IR::Value(8u))));
-	const auto result = ir.ConstructU64(quads, IR::U32(IR::Value(0u)));
-	WriteOperand(DestinationOperand(inst), result);
-	ir.SetScc(IR::U1(ir.Emit(IR::ValueOpcode::INotEqual64, {result, IR::Value(uint64_t {0})})));
+	IR::U32 quads;
+	if (is64) {
+		const auto source = ReadU32Pair(inst.src0);
+		quads = ir.BitwiseOr(
+		    compact(source[0]), ir.ShiftLeftLogical(compact(source[1]), IR::U32(IR::Value(8u))));
+		WriteOperand(DestinationOperand(inst), ir.ConstructU64(quads, IR::U32(IR::Value(0u))));
+	} else {
+		quads = compact(ReadU32(inst.src0));
+		WriteOperand(DestinationOperand(inst), quads);
+	}
+	ir.SetScc(ir.INotEqual(quads, IR::U32(IR::Value(0u))));
 }

 void Translator::BFM_B32(const Decoder::Instruction& inst) {
@@ -248,7 +248,8 @@ void Translator::EmitScalar(const Decoder::Instruction& inst) {
 		case O::S_BITSET0_B64: return S_BITSET_B64(inst, false);
 		case O::S_BITSET1_B64: return S_BITSET_B64(inst, true);
 		case O::S_BITREPLICATE_B64_B32: return S_BITREPLICATE_B64_B32(inst);
-		case O::S_QUADMASK_B64: return S_QUADMASK_B64(inst);
+		case O::S_QUADMASK_B32: return S_QUADMASK(inst, false);
+		case O::S_QUADMASK_B64: return S_QUADMASK(inst, true);
 		case O::S_BFM_B32: return BFM_B32(inst);
 		case O::S_BFM_B64: return S_BFM_B64(inst);
 		case O::S_BFE_U32: return S_BFE_U32(inst, false);
@@ -1165,6 +1165,14 @@ IR::Program TranslateProgram(const Decoder::Program& decoded, const CFG::Graph&
 		                                               : IR::U32(IR::Value(0u));
 		entry_ir.SetExecLo(initial_lo);
 		entry_ir.SetExecHi(initial_hi);
+		const auto active_lane_count = [&] {
+			auto count = IR::U32(entry_ir.Emit(IR::ValueOpcode::BitCount32, {initial_lo}));
+			if (options.wave_size == 64u) {
+				count = entry_ir.IAdd(count, IR::U32(entry_ir.Emit(
+				    IR::ValueOpcode::BitCount32, {initial_hi})));
+			}
+			return count;
+		};
 		if (options.stage == ShaderType::Compute) {
 			const auto* cs = options.input_info.compute;
 			const auto  thread_ids =
@@ -1302,7 +1310,7 @@ IR::Program TranslateProgram(const Decoder::Program& decoded, const CFG::Graph&
 				                      builtin(IR::StageInputKind::WorkgroupId, 1));
 			}
 		} else if (options.stage == ShaderType::Local) {
-			entry_ir.SetScalarReg(static_cast<IR::ScalarReg>(3), IR::U32(IR::Value(64u)));
+			entry_ir.SetScalarReg(static_cast<IR::ScalarReg>(3), active_lane_count());
 			entry_ir.SetVectorReg(static_cast<IR::VectorReg>(2),
 			                      builtin(IR::StageInputKind::VertexIndex));
 			entry_ir.SetVectorReg(static_cast<IR::VectorReg>(3), IR::U32(IR::Value(0u)));
@@ -1380,19 +1388,32 @@ IR::Program TranslateProgram(const Decoder::Program& decoded, const CFG::Graph&
 		} else if (options.stage == ShaderType::Vertex) {
 			// Vulkan owns primitive assembly; each vertex subgroup is one NGG wave.
 			// ESVertCount describes input vertices, independently of lane ordinals.
-			auto vertex_count = IR::U32(entry_ir.Emit(IR::ValueOpcode::BitCount32, {initial_lo}));
-			if (options.wave_size == 64u) {
-				vertex_count = entry_ir.IAdd(vertex_count, IR::U32(entry_ir.Emit(
-				    IR::ValueOpcode::BitCount32, {initial_hi})));
-			}
 			entry_ir.SetScalarReg(static_cast<IR::ScalarReg>(2),
 			                      IR::U32(IR::Value(options.wave_size << 12u)));
 			entry_ir.SetScalarReg(static_cast<IR::ScalarReg>(3),
-			                      entry_ir.BitwiseOr(vertex_count, IR::U32(IR::Value(1u << 28u))));
+			                      entry_ir.BitwiseOr(active_lane_count(), IR::U32(IR::Value(1u << 28u))));
+			// Vulkan adds an indirect record's bases to the indices; the hardware passes them
+			// only in the user SGPRs the packet patches. SH_NOP leaves those SGPRs unchanged;
+			// it does not make the record's bases zero.
+			const auto& offsets = options.input_info.vertex->indirect_offsets;
+			const auto  index   = [&](uint8_t slot, IR::StageInputKind kind,
+			                          IR::StageInputKind base_kind) {
+				if (!options.input_info.vertex->indirect_draw) {
+					return builtin(kind);
+				}
+				const auto base = builtin(base_kind);
+				if (slot < options.user_data_count) {
+					entry_ir.SetScalarReg(static_cast<IR::ScalarReg>(options.user_data_base + slot),
+					                      base);
+				}
+				return entry_ir.ISub(builtin(kind), base);
+			};
 			entry_ir.SetVectorReg(static_cast<IR::VectorReg>(5),
-			                      builtin(IR::StageInputKind::VertexIndex));
+			                      index(offsets.vertex, IR::StageInputKind::VertexIndex,
+			                            IR::StageInputKind::BaseVertex));
 			entry_ir.SetVectorReg(static_cast<IR::VectorReg>(8),
-			                      builtin(IR::StageInputKind::InstanceIndex));
+			                      index(offsets.instance, IR::StageInputKind::InstanceIndex,
+			                            IR::StageInputKind::BaseInstance));
 		}
 	}
 	const bool flush_f32_inputs = options.stage == ShaderType::Compute &&
@@ -167,7 +167,7 @@ private:
 	void Integer16Shift(const Decoder::Instruction& inst, IR::ValueOpcode opcode, bool arithmetic);
 	void Integer16Binary(const Decoder::Instruction& inst, IR::ValueOpcode opcode, bool sign);
 	void Integer16Ternary(const Decoder::Instruction& inst, IR::ValueOpcode opcode, bool sign);
-	void V_MAD_I16(const Decoder::Instruction& inst);
+	void Integer16Mad(const Decoder::Instruction& inst, bool sign);
 	void PackedInteger16Shift(const Decoder::Instruction& inst, IR::ValueOpcode opcode,
 	                          bool arithmetic);
 	void PackedInteger16Binary(const Decoder::Instruction& inst, IR::ValueOpcode opcode);
@@ -187,7 +187,7 @@ private:
 	void V_XOR3_B32(const Decoder::Instruction& inst);
 	void S_FF_I32_B64(const Decoder::Instruction& inst, bool find_zero, bool from_msb, bool sign);
 	void V_FFBH_32(const Decoder::Instruction& inst, bool sign);
-	void Integer24(const Decoder::Instruction& inst, bool sign, bool addend);
+	void Integer24(const Decoder::Instruction& inst, bool sign, bool addend, bool high = false);
 	void V_MAD_U64_U32(const Decoder::Instruction& inst);
 	void V_SAD_U32(const Decoder::Instruction& inst);
 	void V_ADD3_U32(const Decoder::Instruction& inst);
@@ -196,7 +196,7 @@ private:
 	void V_BCNT_U32_B32(const Decoder::Instruction& inst);
 	void V_MBCNT_U32_B32(const Decoder::Instruction& inst, bool low);
 	void S_BITREPLICATE_B64_B32(const Decoder::Instruction& inst);
-	void S_QUADMASK_B64(const Decoder::Instruction& inst);
+	void S_QUADMASK(const Decoder::Instruction& inst, bool is64);
 	void BFM_B32(const Decoder::Instruction& inst);
 	IR::U32 RightMask32(IR::U32 count);
 	void    S_BFM_B64(const Decoder::Instruction& inst);
@@ -375,9 +375,14 @@ void Translator::EmitVector(const Decoder::Instruction& inst) {
 		case O::V_SUB_NC_U16:
 		case O::V_SUB_NC_I16: return Integer16Binary(inst, IR::ValueOpcode::ISub32, false);
 		case O::V_MUL_LO_U16: return Integer16Binary(inst, IR::ValueOpcode::IMul32, false);
-		case O::V_MAD_I16: return V_MAD_I16(inst);
+		case O::V_MAD_I16: return Integer16Mad(inst, true);
+		case O::V_MAD_U16: return Integer16Mad(inst, false);
 		case O::V_MED3_I16: return Integer16Ternary(inst, IR::ValueOpcode::SMedTri32, true);
 		case O::V_MIN3_U16: return Integer16Ternary(inst, IR::ValueOpcode::UMinTri32, false);
+		case O::V_MIN3_I16: return Integer16Ternary(inst, IR::ValueOpcode::SMinTri32, true);
+		case O::V_MAX3_I16: return Integer16Ternary(inst, IR::ValueOpcode::SMaxTri32, true);
+		case O::V_MAX3_U16: return Integer16Ternary(inst, IR::ValueOpcode::UMaxTri32, false);
+		case O::V_MED3_U16: return Integer16Ternary(inst, IR::ValueOpcode::UMedTri32, false);
 		case O::V_MIN_I16: return Integer16Binary(inst, IR::ValueOpcode::SMin32, true);
 		case O::V_MAX_I16: return Integer16Binary(inst, IR::ValueOpcode::SMax32, true);
 		case O::V_MIN_U16: return Integer16Binary(inst, IR::ValueOpcode::UMin32, false);
@@ -565,6 +570,9 @@ void Translator::EmitVector(const Decoder::Instruction& inst) {
 		case O::V_LSHRREV_B64:
 			return SimpleInteger(inst, IR::ValueOpcode::ShiftRightLogical64, IR::Type::U64, true,
 			                     false, false);
+		case O::V_ASHRREV_I64:
+			return SimpleInteger(inst, IR::ValueOpcode::ShiftRightArithmetic64, IR::Type::U64, true,
+			                     false, false);

 		case O::V_XNOR_B32:
 			return ComposedIntegerBinary(inst, IR::ValueOpcode::BitwiseXor32, false, true, false);
@@ -578,6 +586,8 @@ void Translator::EmitVector(const Decoder::Instruction& inst) {
 		case O::V_MAD_U32_U24: return Integer24(inst, false, true);
 		case O::V_MUL_I32_I24: return Integer24(inst, true, false);
 		case O::V_MUL_U32_U24: return Integer24(inst, false, false);
+		case O::V_MUL_HI_I32_I24: return Integer24(inst, true, false, true);
+		case O::V_MUL_HI_U32_U24: return Integer24(inst, false, false, true);
 		case O::V_MAD_U64_U32: return V_MAD_U64_U32(inst);
 		case O::V_SAD_U32: return V_SAD_U32(inst);
 		case O::V_ADD3_U32: return V_ADD3_U32(inst);
@@ -251,6 +251,8 @@ enum class StageInputKind {
 	LocalInvocationId,
 	LocalInvocationIndex,
 	GlobalInvocationId,
+	BaseVertex,
+	BaseInstance,
 	Parameter,
 };

@@ -175,6 +175,8 @@ void AddBuiltinInput(ShaderInfo& info, StageInputKind kind) {
 		case StageInputKind::GlobalInvocationId:
 			AddInput(info, kind, 0, 3, "gl_GlobalInvocationID");
 			break;
+		case StageInputKind::BaseVertex: AddInput(info, kind, 0, 1, "gl_BaseVertex"); break;
+		case StageInputKind::BaseInstance: AddInput(info, kind, 0, 1, "gl_BaseInstance"); break;
 		case StageInputKind::PackedAncillary:
 		case StageInputKind::Parameter: break;
 	}
@@ -444,6 +446,8 @@ void Visit(Program& program, ShaderStageInputInfo input_info, InputUsage& inputs
 				case StageInputKind::InstanceIndex:
 				case StageInputKind::FrontFacing:
 				case StageInputKind::LocalInvocationIndex:
+				case StageInputKind::BaseVertex:
+				case StageInputKind::BaseInstance:
 					if (component != 0u) {
 						return Fail("typed scalar builtin component is out of range");
 					}
@@ -291,10 +291,9 @@ static bool ShaderGetStaticVertexInputInfo(uint64_t shader_addr, const HW::UserS
 	const auto& metadata = data.vertex;

 	if (metadata.vertex_attrib_reg >= 0) {
-		info.fetch_external   = false;
-		info.fetch_embedded   = true;
-		info.fetch_attrib_reg = metadata.vertex_attrib_reg;
-		info.fetch_buffer_reg = metadata.vertex_buffer_reg;
+		info.has_dynamic_vertex_fetch = true;
+		info.fetch_attrib_reg         = metadata.vertex_attrib_reg;
+		info.fetch_buffer_reg         = metadata.vertex_buffer_reg;

 		const auto* attrib = reinterpret_cast<const uint32_t*>(
 		    static_cast<uint64_t>(user_sgpr.value[metadata.vertex_attrib_reg]) |
@@ -413,7 +412,7 @@ static void ShaderGetStaticInputInfoCS(const HW::ComputeShaderInfo& regs,
 void BuildStageStaticKey(const ShaderVertexInputInfo& info, std::vector<uint32_t>& key) {
 	EXIT_IF(info.resources_num < 0 || info.resources_num > ShaderVertexInputInfo::RES_MAX);
 	key.clear();
-	key.push_back(static_cast<uint32_t>(info.fetch_embedded));
+	key.push_back(static_cast<uint32_t>(info.has_dynamic_vertex_fetch));
 	key.push_back(static_cast<uint32_t>(info.fetch_attrib_reg));
 	key.push_back(static_cast<uint32_t>(info.fetch_buffer_reg));
 	key.push_back(info.resources_num);
@@ -422,6 +421,8 @@ void BuildStageStaticKey(const ShaderVertexInputInfo& info, std::vector<uint32_t
 	key.push_back(info.lds_size_dwords);
 	key.push_back(info.pa_cl_vs_out_cntl);
 	key.push_back(info.user_clip_planes);
+	key.push_back(info.indirect_offsets.vertex | (uint32_t {info.indirect_offsets.instance} << 8u) |
+	              (uint32_t {info.indirect_draw} << 16u));
 	key.push_back(static_cast<uint32_t>(info.clip_space.enabled));
 	if (info.clip_space.enabled) {
 		for (const float value: info.clip_space.scale) {
@@ -692,9 +693,7 @@ void ShaderDbgDumpInputInfo(const ShaderVertexInputInfo& info) {

 	LOGF("ShaderDbgDumpInputInfo()\n");

-	LOGF("\t fetch_external = %s\n"
-	     "\t fetch_embedded = %s\n",
-	     info.fetch_external ? "true" : "false", info.fetch_embedded ? "true" : "false");
+	LOGF("\t has_dynamic_vertex_fetch = %s\n", info.has_dynamic_vertex_fetch ? "true" : "false");

 	for (int i = 0; i < info.resources_num; i++) {
 		LOGF("\t input %d\n", i);
@@ -124,6 +124,17 @@ struct ShaderTessellationInputInfo {
 	uint32_t output_topology       = 0;
 };

+// GS user-data slots an indirect draw packet writes its record's base vertex and first instance
+// to. A packet can leave both slots unused even when its record contains nonzero offsets.
+struct ShaderIndirectOffsets {
+	static constexpr uint8_t None = UINT8_MAX;
+
+	uint8_t vertex   = None;
+	uint8_t instance = None;
+
+	bool operator==(const ShaderIndirectOffsets&) const = default;
+};
+
 struct ShaderVertexInputInfo {
 	static constexpr int RES_MAX = 32;

@@ -146,8 +157,11 @@ struct ShaderVertexInputInfo {
 	ShaderClipSpaceTransform clip_space;
 	ShaderMeshInputInfo      mesh;
 	ShaderTessellationInputInfo tess;
-	bool                    fetch_external      = false;
-	bool                    fetch_embedded      = false;
+	ShaderIndirectOffsets   indirect_offsets;
+	// GPU-read records add Vulkan bases even when the packet patches no user SGPRs.
+	bool                    indirect_draw       = false;
+	// Attribute-table metadata enables dynamic fetch detection; loads are matched later.
+	bool                    has_dynamic_vertex_fetch = false;
 };

 struct ShaderComputeInputInfo: ShaderWorkgroupInputInfo {
@@ -26,9 +26,10 @@ struct DebugStats {
 	uint64_t flexible_total      = 0;
 	uint64_t flexible_used       = 0;
 	uint64_t pool_committed      = 0;
-	uint64_t cpu_page_entries    = 0; // 2 MiB page-table entries in use.
-	uint64_t gpu_page_entries    = 0;
-	uint64_t page_entries_total  = 0; // Per side.
+	uint64_t cpu_page_entries       = 0; // 2 MiB page-table entries in use.
+	uint64_t gpu_page_entries       = 0;
+	uint64_t cpu_page_entries_total = 0;
+	uint64_t gpu_page_entries_total = 0;
 };
 [[nodiscard]] DebugStats GetDebugStats();

@@ -125,7 +126,6 @@ static_assert(sizeof(KernelMemoryPoolBlockStats) == 16,
              "KernelMemoryPoolBlockStats struct size is incorrect");

 void                   RegisterCallbacks(callback_func_t alloc_func, callback_func_t free_func);
-void                   SetFlexibleMemorySize(uint64_t size);
 int AllocateDirectMemory(int64_t search_start, int64_t search_end, size_t size, size_t alignment,
                         int memory_type, int64_t* phys_addr_out, bool automatic = false);
 int MapAutomaticMemory(uint64_t vaddr, size_t size, int type, int prot);
@@ -225,6 +225,11 @@ bool     TestPlaceholderRangeIsFree(uint64_t vaddr, uint64_t size);
 bool     TestGuestAddressRangeIsOwned(uint64_t vaddr, uint64_t size);
 bool     TestGuestBackingOutsideAddressSpace();
 uint64_t TestGuestBackingSize();
+uint64_t TestFlexibleBackingHighWater();
+uint64_t TestGuestFreeRangeVisits();
+uint64_t TestDirectMappingCount();
+uint64_t TestFlexibleMappingCount();
+uint64_t TestBackingBytesZeroed();
 bool     TestGuestFreeRangeBounds();
 #endif

@@ -203,6 +203,9 @@ public:
 		}
 		std::lock_guard lock(m_mutex);
 		std::memset(m_backing_base + backing_offset, 0, size);
+#if defined(KYTY_VIRTUAL_MEMORY_ALLOCATION_TESTS)
+		g_test_backing_bytes_zeroed += size;
+#endif
 		return true;
 	}

@@ -287,6 +290,18 @@ public:
 		return true;
 	}

+	bool ProtectViews(uint64_t vaddr, uint64_t size, VirtualMemory::Mode mode) {
+		std::shared_lock lock(m_mutex);
+		bool             ok = true;
+		ForEachSparseRange(
+		    m_maps, vaddr, size,
+		    [&](const auto&, uint64_t start, uint64_t bytes) {
+			    ok = ok && VirtualMemory::Protect(start, bytes, mode);
+		    },
+		    [&](uint64_t, uint64_t) { ok = false; });
+		return ok;
+	}
+
 	bool MapFixed(uint64_t vaddr, uint64_t size, uint64_t backing_offset,
 	              VirtualMemory::Mode mode, FailureReason* failure_reason = nullptr) {
 		if (!IsAvailable() || !RangeOk(backing_offset, size)) {
@@ -324,15 +339,11 @@ public:
 #endif
 	}

-	bool Unmap(uint64_t vaddr, uint64_t size, bool* placeholder_preserved = nullptr) {
+	bool Unmap(uint64_t vaddr, uint64_t size) {
 		if (size == 0 || UINT64_MAX - vaddr < size) {
 			return false;
 		}

-		if (placeholder_preserved != nullptr) {
-			*placeholder_preserved = false;
-		}
-
 		const auto            end = vaddr + size;
 		std::vector<MapEntry> targets;
 		{
@@ -350,42 +361,36 @@ public:
 			}
 		}

-		std::vector<MapEntry> removed;
-		removed.reserve(targets.size());
+		size_t removed  = 0;
+		auto   rollback = [&] {
+			while (removed != 0) {
+				const auto& previous = targets[--removed];
+				EXIT_IF(!MapFixed(previous.vaddr, previous.size, previous.backing_offset,
+				                  previous.mode));
+			}
+		};
 		for (const auto& entry: targets) {
-			bool part_placeholder_preserved = false;
-			if (!UnmapOne(entry.vaddr, entry.size, &part_placeholder_preserved)) {
-				for (auto it = removed.rbegin(); it != removed.rend(); ++it) {
-					EXIT_IF(!MapFixed(it->vaddr, it->size, it->backing_offset, it->mode));
-				}
-				if (placeholder_preserved != nullptr) {
-					*placeholder_preserved = false;
-				}
+			if (!UnmapOne(entry.vaddr, entry.size)) {
+				rollback();
 				return false;
 			}
-			EXIT_IF(!part_placeholder_preserved);
-			removed.push_back(entry);
+			++removed;
 		}

 #if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
 		// Each unmapped file view becomes a separate Windows placeholder. The address-space owner
 		// records this contiguous guest span as one free range, so normalize fragmented backing
 		// views into one placeholder before publishing it there.
-		if (removed.size() > 1 && !CoalescePlaceholders(vaddr, size)) {
-			for (auto it = removed.rbegin(); it != removed.rend(); ++it) {
-				EXIT_IF(!MapFixed(it->vaddr, it->size, it->backing_offset, it->mode));
-			}
+		if (removed > 1 && !CoalescePlaceholders(vaddr, size)) {
+			rollback();
 			return false;
 		}
 #endif

-		if (placeholder_preserved != nullptr) {
-			*placeholder_preserved = true;
-		}
 		return true;
 	}

-	bool UnmapOne(uint64_t vaddr, uint64_t size, bool* placeholder_preserved = nullptr) {
+	bool UnmapOne(uint64_t vaddr, uint64_t size) {
 #if defined(KYTY_VIRTUAL_MEMORY_ALLOCATION_TESTS)
 		if (g_test_backing_store_unmaps_before_failure == 0) {
 			g_test_backing_store_unmaps_before_failure = UINT32_MAX;
@@ -428,9 +433,7 @@ public:
 				m_maps.emplace(old.vaddr, old);
 				return false;
 			}
-			if (placeholder_preserved != nullptr) {
-				*placeholder_preserved = true;
-			}
+
 			return true;
 		}

@@ -451,9 +454,6 @@ public:
 				     ok;
 			}
 			if (ok) {
-				if (placeholder_preserved != nullptr) {
-					*placeholder_preserved = true;
-				}
 				return true;
 			}

@@ -481,9 +481,6 @@ public:
 		std::lock_guard lock(m_mutex);
 		RemoveMappingsUnlocked(vaddr, size);

-		if (placeholder_preserved != nullptr) {
-			*placeholder_preserved = true;
-		}
 		return true;
 #endif
 	}
@@ -634,8 +631,6 @@ private:

 	HANDLE m_handle = nullptr;
 #elif KYTY_PLATFORM == KYTY_PLATFORM_LINUX
-	[[nodiscard]] static uint64_t GetMappableTestSize() { return 0x4000; }
-
 	static int GetProtect(VirtualMemory::Mode mode) {
 		switch (mode) {
 			case VirtualMemory::Mode::Read: return PROT_READ;
@@ -760,7 +755,7 @@ public:
 		if (size == 0 || UINT64_MAX - vaddr < size) {
 			return false;
 		}
-		std::lock_guard lock(m_mutex);
+		std::shared_lock lock(m_mutex);
 		return std::any_of(m_owned.begin(), m_owned.end(), [vaddr, size](const auto& range) {
 			return VirtualRangesOverlap(vaddr, size, range.first, range.second);
 		});
@@ -864,11 +859,7 @@ public:
 		if (!ContainsMappedUnlocked(vaddr, size, MappingType::Backing)) {
 			return false;
 		}
-		bool placeholder_preserved = false;
-		if (!m_backing->Unmap(vaddr, size, &placeholder_preserved)) {
-			return false;
-		}
-		if (!placeholder_preserved) {
+		if (!m_backing->Unmap(vaddr, size)) {
 			return false;
 		}
 		EXIT_IF(!RemoveMappedUnlocked(vaddr, size, MappingType::Backing));
@@ -876,12 +867,13 @@ public:
 		return true;
 	}

+	// Reservations live in VirtualRanges; these free-range checks only read the host placeholders.
 	uint64_t ReserveAligned(uint64_t search_addr, uint64_t size, uint64_t alignment) {
 		if (size == 0) {
 			return 0;
 		}
 		alignment = std::max<uint64_t>(alignment != 0 ? alignment : PageSize(), PageSize());
-		std::lock_guard lock(m_mutex);
+		std::shared_lock lock(m_mutex);
 		if (search_addr == 0) {
 			auto addr = FindFreeAlignedUnlocked(HOST_SYSTEM_MANAGED_MIN,
 			                                    HOST_SYSTEM_MANAGED_MAX + 1u, size, alignment);
@@ -898,7 +890,7 @@ public:
 			return 0;
 		}
 		alignment = std::max<uint64_t>(alignment, PageSize());
-		std::lock_guard lock(m_mutex);
+		std::shared_lock lock(m_mutex);
 		return FindFreeAlignedUnlocked(search_start, search_end, size, alignment);
 	}

@@ -906,19 +898,19 @@ public:
 		if (vaddr == 0 || size == 0 || !IsPageAligned(vaddr) || !IsPageAligned(size)) {
 			return false;
 		}
-		std::lock_guard lock(m_mutex);
+		std::shared_lock lock(m_mutex);
 		return FindContaining(vaddr, size) != m_free.end();
 	}

 #if defined(KYTY_VIRTUAL_MEMORY_ALLOCATION_TESTS)
 	bool TestContainsFree(uint64_t vaddr, uint64_t size) {
-		std::lock_guard lock(m_mutex);
+		std::shared_lock lock(m_mutex);
 		return FindContaining(vaddr, size) != m_free.end();
 	}
 #endif

 	bool ReleaseFree(uint64_t vaddr, uint64_t size) {
-		std::lock_guard lock(m_mutex);
+		std::shared_lock lock(m_mutex);
 		return FindContaining(vaddr, size) != m_free.end();
 	}

@@ -1041,10 +1033,6 @@ private:
 		return (value & (PageSize() - 1u)) == 0;
 	}

-	[[nodiscard]] bool IsGranularityAligned(uint64_t value) const {
-		return (value & (m_granularity - 1u)) == 0;
-	}
-
 	bool ProtectMappedUnlocked(uint64_t vaddr, uint64_t size, VirtualMemory::Mode mode) {
 		const auto end = vaddr + size;
 		auto       it  = m_mapped.upper_bound(vaddr);
@@ -1055,8 +1043,20 @@ private:
 			const auto& region      = it->second;
 			const auto  range_start = std::max(vaddr, region.vaddr);
 			const auto  range_end   = std::min(end, region.vaddr + region.size);
-			if (range_start < range_end &&
-			    !VirtualMemory::Protect(range_start, range_end - range_start, mode)) {
+			if (range_start >= range_end) {
+				continue;
+			}
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+			// Rollback can split a backing view without changing the owner interval.
+			// Windows protection calls must stay within each actual view.
+			if (region.type == MappingType::Backing) {
+				if (!m_backing->ProtectViews(range_start, range_end - range_start, mode)) {
+					return false;
+				}
+				continue;
+			}
+#endif
+			if (!VirtualMemory::Protect(range_start, range_end - range_start, mode)) {
 				return false;
 			}
 		}
@@ -1312,12 +1312,21 @@ private:
 	}

 	uint64_t FindFreeAlignedUnlocked(uint64_t search_start, uint64_t search_end, uint64_t size,
-	                                 uint64_t alignment) {
+	                                 uint64_t alignment) const {
 		if (search_start >= search_end || size > search_end - search_start) {
 			return 0;
 		}
-		for (const auto& [base, range_size]: m_free) {
-			const auto candidate = AlignUp(std::max(base, search_start), alignment);
+		// Include the predecessor because search_start can lie inside a free range.
+		auto it = m_free.upper_bound(search_start);
+		if (it != m_free.begin()) {
+			--it;
+		}
+		for (; it != m_free.end() && it->first < search_end; ++it) {
+#if defined(KYTY_VIRTUAL_MEMORY_ALLOCATION_TESTS)
+			g_test_free_range_visits.fetch_add(1, std::memory_order_relaxed);
+#endif
+			const auto& [base, range_size] = *it;
+			const auto candidate           = AlignUp(std::max(base, search_start), alignment);
 			if (candidate < search_end && size <= search_end - candidate &&
 			    GuestFreeRangeContains(base, range_size, candidate, size)) {
 				return candidate;
@@ -1331,9 +1340,8 @@ private:
 		if (next == m_free.begin()) {
 			return m_free.end();
 		}
-		auto it  = std::prev(next);
-		auto end = it->first + it->second;
-		return (vaddr >= it->first && vaddr + size <= end ? it : m_free.end());
+		auto it = std::prev(next);
+		return GuestFreeRangeContains(it->first, it->second, vaddr, size) ? it : m_free.end();
 	}

 	bool SplitUnlocked(std::map<uint64_t, uint64_t>::iterator it, uint64_t vaddr, uint64_t size) {
@@ -125,6 +125,7 @@ if(APPLE)
 	install(TARGETS kyty_emulator DESTINATION "KytyPS5.app/Contents/MacOS")
 	install(TARGETS kyty_updater DESTINATION "KytyPS5.app/Contents/MacOS")
 	install(FILES "${KYTY_TROPHY_UNLOCK_SOUND}" DESTINATION "KytyPS5.app/Contents/Resources/assets/sounds")
+	install(FILES ${KYTY_NOTO_SANS_FONTS} DESTINATION "KytyPS5.app/Contents/Resources/assets/fonts")
 	install(DIRECTORY "${CMAKE_SOURCE_DIR}/assets/cursors" DESTINATION "KytyPS5.app/Contents/Resources/assets")
 	# Deploy and ad-hoc sign the Qt frameworks and plugins in the bundle.
 	find_program(KYTY_MACDEPLOYQT NAMES macdeployqt PATHS "${Qt6_DIR}/../../../bin" NO_DEFAULT_PATH)
@@ -115,7 +115,17 @@
            </item>
            <item row="0" column="1">
             <widget class="QComboBox" name="comboBox_screen_resolution">
-              <property name="toolTip"><string>Window resolution</string></property>
+              <property name="toolTip"><string>Window resolution as WIDTHxHEIGHT; pick a preset or type a custom size. Non-16:9 sizes such as 21:9 show the game with black bars unless Stretch to window is on.</string></property>
+              <property name="editable"><bool>true</bool></property>
+              <property name="insertPolicy"><enum>QComboBox::NoInsert</enum></property>
+              <item><property name="text"><string notr="true">1280x720</string></property></item>
+              <item><property name="text"><string notr="true">1920x1080</string></property></item>
+              <item><property name="text"><string notr="true">2560x1440</string></property></item>
+              <item><property name="text"><string notr="true">3840x2160</string></property></item>
+              <item><property name="text"><string notr="true">2560x1080</string></property></item>
+              <item><property name="text"><string notr="true">3440x1440</string></property></item>
+              <item><property name="text"><string notr="true">3840x1600</string></property></item>
+              <item><property name="text"><string notr="true">5120x2160</string></property></item>
             </widget>
            </item>
            <item row="1" column="0">
@@ -256,6 +266,12 @@
              <property name="text"><string>Auto-hide cursor</string></property>
             </widget>
            </item>
+            <item>
+             <widget class="QCheckBox" name="checkBox_stretch">
+              <property name="toolTip"><string>Stretch the game image to fill the window instead of keeping its aspect ratio with black bars</string></property>
+              <property name="text"><string>Stretch to window</string></property>
+             </widget>
+            </item>
            <item>
             <widget class="QCheckBox" name="checkBox_perf_overlay">
              <property name="toolTip"><string>Show FPS, frame times, CPU/GPU usage, RAM and memory pressure in the game window. Press F2 (Cmd+P on macOS) to toggle it, and add Shift for the detailed view.</string></property>
@@ -21,10 +21,7 @@ inline QStringList EnumToList() {
 	auto        me    = QMetaEnum::fromType<T>();
 	int         count = me.keyCount();
 	for (int i = 0; i < count; i++) {
-		auto key = QString(me.key(i));
-		ret << (key.startsWith('R') && key.size() > 2 && key.at(1).isDigit()
-		            ? key.remove('R').toLower()
-		            : key);
+		ret << QString(me.key(i));
 	}
 	return ret;
 }
@@ -32,16 +29,12 @@ inline QStringList EnumToList() {
 template <class T>
 T TextToEnum(const QString& text) {
 	auto me = QMetaEnum::fromType<T>();
-	return static_cast<T>(me.keyToValue(
-	    ((text.size() > 1 && text.at(0).isDigit()) ? 'R' + text.toUpper() : text).toUtf8().data()));
+	return static_cast<T>(me.keyToValue(text.toUtf8().constData()));
 }

 template <class T>
 QString EnumToText(T value) {
-	auto me  = QMetaEnum::fromType<T>();
-	auto key = QString(me.valueToKey(static_cast<int>(value)));
-	return (key.startsWith('R') && key.size() > 2 && key.at(1).isDigit() ? key.remove('R').toLower()
-	                                                                     : key);
+	return QString(QMetaEnum::fromType<T>().valueToKey(static_cast<int>(value)));
 }

 struct ControllerSettings {
@@ -80,13 +73,10 @@ public:
 	static constexpr int DEFAULT_CONSOLE_LANGUAGE = 1;
 	static constexpr int MAX_CONSOLE_LANGUAGE     = 29;

-	enum class Resolution {
-		R1280X720,
-		R1920X1080,
-		R2560X1440,
-		R3840X2160,
-	};
-	Q_ENUM(Resolution)
+	static constexpr int MIN_SCREEN_SIZE = 320;
+	static constexpr int MAX_SCREEN_SIZE = 16384;
+
+	static QString NormalizeResolution(const QString& text);

 	enum class ShaderOptimizationType { None, Size, Performance };
 	Q_ENUM(ShaderOptimizationType)
@@ -117,7 +107,7 @@ public:
 	// Controller preferences always come from the global configuration.
 	ControllerSettings controller;

-	Resolution             screen_resolution           = Resolution::R1280X720;
+	QString                screen_resolution           = QStringLiteral("1280x720");
 	QString                user_name                   = "Kyty";
 	int                    user_id                     = Config::DEFAULT_USER_ID;
 	QString                audio_input_device;
@@ -126,6 +116,7 @@ public:
 	bool                   fullscreen_enabled          = false;
 	bool                   hdr_enabled                 = false;
 	bool                   hide_cursor_enabled         = false;
+	bool                   stretch_enabled             = false;
 	bool                   readback_linear_images      = false;
 	bool                   readback_tiled_images       = false;
 	bool                   sync_raw_image_buffers      = false;
@@ -168,6 +159,7 @@ public:
 		fullscreen_enabled          = other.fullscreen_enabled;
 		hdr_enabled                 = other.hdr_enabled;
 		hide_cursor_enabled         = other.hide_cursor_enabled;
+		stretch_enabled             = other.stretch_enabled;
 		readback_linear_images      = other.readback_linear_images;
 		readback_tiled_images       = other.readback_tiled_images;
 		sync_raw_image_buffers      = other.sync_raw_image_buffers;
@@ -46,11 +46,15 @@ private:
 	QLineEdit*                   m_addons_dir_edit        = nullptr;
 	QToolButton*                 m_remove_game_dir_button = nullptr;
 	bool                         m_global_settings        = false;
+	bool                         m_geometry_restored      = false;

 protected:
 	void Init(const Configuration& info);
 	void InitGameDirectories();
 	void AddGameDirectoryItem(const QString& dir);
+	void MakeTabsScrollable();
+	void GrowToFitContents();
+	void ShowTab(QWidget* page);

 	void moveEvent(QMoveEvent* event) override;
 	void resizeEvent(QResizeEvent* event) override;
@@ -1,5 +1,23 @@
 #include "configuration.h"

+#include <QRegularExpression>
+
+QString Configuration::NormalizeResolution(const QString& text) {
+	// Older settings stored the enum key, e.g. "R1920X1080".
+	static const QRegularExpression re(QStringLiteral("^\\s*R?(\\d+)\\s*[xX]\\s*(\\d+)\\s*$"));
+	const auto match = re.match(text);
+	if (!match.hasMatch()) {
+		return {};
+	}
+	const int width  = match.captured(1).toInt();
+	const int height = match.captured(2).toInt();
+	if (width < MIN_SCREEN_SIZE || height < MIN_SCREEN_SIZE || width > MAX_SCREEN_SIZE ||
+	    height > MAX_SCREEN_SIZE) {
+		return {};
+	}
+	return QStringLiteral("%1x%2").arg(width).arg(height);
+}
+
 QVariantMap Configuration::GameSettings() const {
 	QVariantMap values;
 #define KYTY_CFG_SET(n) values.insert(QStringLiteral(#n), QVariant::fromValue(n).toString())
@@ -12,6 +30,7 @@ QVariantMap Configuration::GameSettings() const {
 	KYTY_CFG_SET(fullscreen_enabled);
 	KYTY_CFG_SET(hdr_enabled);
 	KYTY_CFG_SET(hide_cursor_enabled);
+	KYTY_CFG_SET(stretch_enabled);
 	KYTY_CFG_SET(readback_linear_images);
 	KYTY_CFG_SET(readback_tiled_images);
 	KYTY_CFG_SET(sync_raw_image_buffers);
@@ -48,7 +67,8 @@ QVariantMap Configuration::GameSettings() const {
 template <class Settings>
 void Configuration::ReadGameSettingsValues(const Settings& s) {
 #define KYTY_CFG_GET(n) n = s.value(QStringLiteral(#n)).template value<decltype(n)>()
-	KYTY_CFG_GET(screen_resolution);
+	const auto resolution = NormalizeResolution(s.value("screen_resolution").toString());
+	screen_resolution     = resolution.isEmpty() ? QStringLiteral("1280x720") : resolution;
 	user_name          = s.value("user_name", user_name).toString();
 	bool user_id_ok    = false;
 	auto saved_user_id = s.value("user_id", user_id).toInt(&user_id_ok);
@@ -64,6 +84,7 @@ void Configuration::ReadGameSettingsValues(const Settings& s) {
 	KYTY_CFG_GET(fullscreen_enabled);
 	KYTY_CFG_GET(hdr_enabled);
 	KYTY_CFG_GET(hide_cursor_enabled);
+	KYTY_CFG_GET(stretch_enabled);
 	KYTY_CFG_GET(readback_linear_images);
 	KYTY_CFG_GET(readback_tiled_images);
 	KYTY_CFG_GET(sync_raw_image_buffers);
@@ -134,7 +155,9 @@ bool Configuration::SetGameSettings(const QJsonObject& settings, QString& error)
 		if (!it.value().isString()) {
 			return fail(tr("Game setting %1 must be a string.").arg(it.key()));
 		}
-		values.insert(it.key(), it.value().toString());
+		const auto value = it.value().toString();
+		values.insert(it.key(), it.key() == "screen_resolution" ? NormalizeResolution(value)
+		                                                       : value);
 	}

 	ReadGameSettingsValues(values);
@@ -24,6 +24,8 @@
 #include <QMessageBox>
 #include <QPushButton>
 #include <QResizeEvent>
+#include <QScreen>
+#include <QScrollArea>
 #include <QSettings>
 #include <QSlider>
 #include <QSpinBox>
@@ -122,10 +124,10 @@ static QString GameDirectoryKey(const QString& dir) {
 }

 ConfigurationEditDialog::ConfigurationEditDialog(Configuration& info, QWidget* parent)
-    : QDialog(parent, Qt::WindowCloseButtonHint), m_ui(new Ui::ConfigurationEditDialog),
-      m_info(info) {
+    : QDialog(parent, Qt::WindowMaximizeButtonHint | Qt::WindowCloseButtonHint),
+      m_ui(new Ui::ConfigurationEditDialog), m_info(info) {
 	m_ui->setupUi(this);
-	setMinimumWidth(width());
+	MakeTabsScrollable();
 	InitGameDirectories();
 	m_ui->controller_group->setVisible(false);
 	m_ui->settings_button->setVisible(false);
@@ -170,9 +172,11 @@ ConfigurationEditDialog::ConfigurationEditDialog(Configuration& info, QWidget* p

 	setSizeGripEnabled(true);

-	restoreGeometry(g_last_geometry);
-
 	Init(info);
+	m_geometry_restored = restoreGeometry(g_last_geometry);
+	if (!m_geometry_restored) {
+		GrowToFitContents();
+	}
 }

 QByteArray ConfigurationEditDialog::g_last_geometry;
@@ -248,7 +252,7 @@ void ConfigurationEditDialog::Init(const Configuration& info) {
 		                    info.audio_input_device);
 	}
 	microphone->setCurrentIndex(microphone->findData(info.audio_input_device));
-	ListInit(m_ui->comboBox_screen_resolution, info.screen_resolution);
+	m_ui->comboBox_screen_resolution->setCurrentText(info.screen_resolution);
 	ListInit(m_ui->comboBox_present_mode, info.present_mode);
 	m_ui->comboBox_gpu->clear();
 	m_ui->comboBox_gpu->addItem(tr("Auto"));
@@ -285,6 +289,7 @@ void ConfigurationEditDialog::Init(const Configuration& info) {
 	m_ui->checkBox_fullscreen->setChecked(info.fullscreen_enabled);
 	m_ui->checkBox_hdr->setChecked(info.hdr_enabled);
 	m_ui->checkBox_hide_cursor->setChecked(info.hide_cursor_enabled);
+	m_ui->checkBox_stretch->setChecked(info.stretch_enabled);
 	m_ui->checkBox_readback->setChecked(info.readback_linear_images);
 	m_ui->checkBox_readback_tiled->setChecked(info.readback_tiled_images);
 	m_ui->checkBox_sync_raw_image_buffers->setChecked(info.sync_raw_image_buffers);
@@ -413,8 +418,46 @@ void ConfigurationEditDialog::SetGlobalSettings(const QStringList& dirs, const Q
 	m_game_dirs_group->setVisible(true);
 	m_ui->controller_group->setVisible(true);
 	update_game_directory_buttons();
+	if (!m_geometry_restored) {
+		GrowToFitContents();
+	}
+}
+
+void ConfigurationEditDialog::MakeTabsScrollable() {
+	auto* tabs = m_ui->settings_tabs;
+	for (int i = 0; i < tabs->count(); i++) {
+		auto* page   = tabs->widget(i);
+		auto  title  = tabs->tabText(i);
+		auto* scroll = new QScrollArea;
+		scroll->setWidgetResizable(true);
+		scroll->setFrameShape(QFrame::NoFrame);
+		tabs->removeTab(i);
+		scroll->setWidget(page);
+		tabs->insertTab(i, scroll, title);
+	}
+	tabs->setCurrentIndex(0);
+}
+
+void ConfigurationEditDialog::GrowToFitContents() {
 	layout()->activate();
-	resize(size().expandedTo(minimumSizeHint()));
+	auto* tabs = m_ui->settings_tabs;
+	QSize content;
+	for (int i = 0; i < tabs->count(); i++) {
+		content = content.expandedTo(
+		    static_cast<QScrollArea*>(tabs->widget(i))->widget()->minimumSizeHint());
+	}
+	const QSize chrome = size() - tabs->currentWidget()->size();
+	resize(size().expandedTo(content + chrome).boundedTo(screen()->availableGeometry().size()));
+}
+
+void ConfigurationEditDialog::ShowTab(QWidget* page) {
+	auto* tabs = m_ui->settings_tabs;
+	for (int i = 0; i < tabs->count(); i++) {
+		if (tabs->widget(i)->isAncestorOf(page)) {
+			tabs->setCurrentIndex(i);
+			return;
+		}
+	}
 }

 QString ConfigurationEditDialog::GetAddonsDirectory() const {
@@ -474,13 +517,14 @@ static void UpdateInfo(Configuration& info, Ui::ConfigurationEditDialog& ui, boo
 		info.controller.keyboard_player     = ui.comboBox_keyboard_player->currentData().toInt();
 	}
 	info.screen_resolution =
-	    TextToEnum<Configuration::Resolution>(ui.comboBox_screen_resolution->currentText());
+	    Configuration::NormalizeResolution(ui.comboBox_screen_resolution->currentText());
 	info.present_mode =
 	    TextToEnum<Configuration::PresentMode>(ui.comboBox_present_mode->currentText());
 	info.gpu_index                 = ui.comboBox_gpu->currentIndex() - 1;
 	info.fullscreen_enabled        = ui.checkBox_fullscreen->isChecked();
 	info.hdr_enabled               = ui.checkBox_hdr->isChecked();
 	info.hide_cursor_enabled       = ui.checkBox_hide_cursor->isChecked();
+	info.stretch_enabled           = ui.checkBox_stretch->isChecked();
 	info.readback_linear_images    = ui.checkBox_readback->isChecked();
 	info.readback_tiled_images     = ui.checkBox_readback_tiled->isChecked();
 	info.sync_raw_image_buffers    = ui.checkBox_sync_raw_image_buffers->isChecked();
@@ -518,7 +562,7 @@ void ConfigurationEditDialog::save() {
 	for (auto* field: {m_ui->lineEdit_shader_log_folder, m_ui->lineEdit_printf_file,
 	                   m_ui->lineEdit_cmd_dump_folder}) {
 		if (field->isEnabled() && field->text().isEmpty()) {
-			m_ui->settings_tabs->setCurrentWidget(m_ui->debug_tab);
+			ShowTab(m_ui->debug_tab);
 			field->setFocus();
 			break;
 		}
@@ -540,6 +584,16 @@ void ConfigurationEditDialog::save() {
 		                      tr("User ID cannot be 254 (everyone) or 255 (system)"));
 		return;
 	}
+	if (Configuration::NormalizeResolution(m_ui->comboBox_screen_resolution->currentText())
+	        .isEmpty()) {
+		ShowTab(m_ui->settings_tab);
+		m_ui->comboBox_screen_resolution->setFocus();
+		QMessageBox::critical(this, tr("Save failed"),
+		                      tr("Screen resolution must be WIDTHxHEIGHT, each between %1 and %2")
+		                          .arg(Configuration::MIN_SCREEN_SIZE)
+		                          .arg(Configuration::MAX_SCREEN_SIZE));
+		return;
+	}

 	UpdateInfo(m_info, *m_ui, m_global_settings);

@@ -134,7 +134,8 @@ void MainDialogPrivate::Setup(MainDialog* main_dialog) {
 	m_ui->close_terminal->setVisible(false); // The emulator runs without a terminal.
 #endif

-	main_dialog->setWindowFlags(Qt::Dialog /*| Qt::MSWindowsFixedSizeDialogHint*/);
+	main_dialog->setWindowFlags(Qt::Dialog | Qt::WindowMinMaxButtonsHint |
+	                            Qt::WindowCloseButtonHint);

 	connect(main_dialog, &MainDialog::Start, this, &MainDialogPrivate::FindInterpreter,
 	        Qt::QueuedConnection);
@@ -240,7 +241,7 @@ static void AddFlag(QStringList& args, bool enabled, const char* name) {

 static QStringList CreateEmulatorArgs(const Configuration& info) {
 	QStringList args;
-	auto        r = EnumToText(info.screen_resolution).split('x');
+	auto        r = Configuration::NormalizeResolution(info.screen_resolution).split('x');

 	if (r.size() != 2) {
 		return {};
@@ -266,6 +267,7 @@ static QStringList CreateEmulatorArgs(const Configuration& info) {
 	AddFlag(args, info.fullscreen_enabled, "--fullscreen");
 	AddFlag(args, info.hdr_enabled, "--hdr");
 	AddFlag(args, info.hide_cursor_enabled, "--hide-cursor");
+	AddFlag(args, info.stretch_enabled, "--stretch");
 	AddFlag(args, info.readback_linear_images, "--readback-linear-images");
 	AddFlag(args, info.readback_tiled_images, "--readback-tiled-images");
 	AddFlag(args, info.sync_raw_image_buffers, "--sync-raw-image-buffers");
@@ -6,6 +6,7 @@
 #include "libs/ajm/decoder.h"
 #include "libs/ajm/mp3_decoder.h"
 #include "libs/ajm/opus_decoder.h"
+#include "libs/ajm/resampler.h"
 #include "libs/audio.h"
 #include "libs/errno.h"
 #include "libs/libs.h"
@@ -63,12 +64,6 @@ struct AjmSidebandMFrame {
 	uint32_t reserved;
 };

-struct AjmSidebandResampleInfo {
-	float    ratio;
-	int32_t  num_samples;
-	uint32_t reserved[8];
-};
-
 struct AjmGetStatisticsResult {
 	AjmSidebandResult result;
 	float             engine_usage_batch;
@@ -104,6 +99,8 @@ constexpr uint64_t AJM_FLAG_RUN_GET_CODEC_INFO        = 1ull << 11u;
 constexpr uint64_t AJM_FLAG_RUN_MULTIPLE_FRAMES       = 1ull << 12u;
 constexpr uint64_t AJM_FLAG_CONTROL_RESET             = 1ull << 13u;
 constexpr uint64_t AJM_FLAG_CONTROL_INITIALIZE        = 1ull << 14u;
+constexpr uint64_t AJM_FLAG_CONTROL_RESAMPLE          = 1ull << 15u;
+constexpr uint64_t AJM_FLAG_CONTROL_RESAMPLE_EX       = 1ull << 16u;
 constexpr uint64_t AJM_FLAG_SIDEBAND_RESAMPLE_INFO    = 1ull << 44u;
 constexpr uint64_t AJM_FLAG_SIDEBAND_GAPLESS          = 1ull << 45u;
 constexpr uint64_t AJM_FLAG_SIDEBAND_FORMAT           = 1ull << 46u;
@@ -177,6 +174,7 @@ struct AjmInstanceState {
 	uint64_t                    flags   = 0;
 	AjmGaplessState             gapless;
 	std::unique_ptr<AjmDecoder> decoder;
+	AjmResampler                resampler;
 };

 static std::mutex                                     g_ajm_instances_mutex;
@@ -358,7 +356,7 @@ static void AjmWriteSideband(uint64_t flags, void* sideband_output, size_t sideb
 	    offset + AJM_SIDEBAND_RESAMPLE_INFO_SIZE <= sideband_output_size) {
 		auto* info = reinterpret_cast<AjmSidebandResampleInfo*>(
 		    static_cast<uint8_t*>(sideband_output) + offset);
-		info->ratio = 1.0f;
+		*info = decode_result.resample_info;
 		offset += AJM_SIDEBAND_RESAMPLE_INFO_SIZE;
 	}

@@ -425,6 +423,12 @@ static AjmDecodeResult AjmDecodeInstance(uint32_t instance, const void* input, s
 		return result;
 	}

+	// The resampler requires one complete decoded packet per decoder call.
+	if (state->codec == static_cast<uint32_t>(AjmCodec::DecOpus) &&
+	    (state->flags & AJM_INSTANCE_FLAG_RESAMPLE) != 0) {
+		return state->resampler.Decode(*state->decoder, state->flags, input, input_size, output,
+		                               output_size, multiple_frames, &state->gapless);
+	}
 	return state->decoder->Decode(input, input_size, output, output_size, multiple_frames,
 	                              &state->gapless);
 }
@@ -514,6 +518,32 @@ static AjmDecodeResult AjmControlInstance(uint32_t instance, uint64_t flags,
 		sideband_input = static_cast<const uint8_t*>(sideband_input) + AJM_SIDEBAND_GAPLESS_SIZE;
 		sideband_input_size -= AJM_SIDEBAND_GAPLESS_SIZE;
 	}
+	for (const auto flag : {AJM_FLAG_CONTROL_RESAMPLE, AJM_FLAG_CONTROL_RESAMPLE_EX}) {
+		if ((flags & flag) == 0) {
+			continue;
+		}
+		const bool   extended = flag == AJM_FLAG_CONTROL_RESAMPLE_EX;
+		const size_t size     = extended ? 12 : 8;
+		if (sideband_input == nullptr || sideband_input_size < size ||
+		    (state->flags & AJM_INSTANCE_FLAG_RESAMPLE) == 0) {
+			result.result = AJM_RESULT_INVALID_PARAMETER;
+			return result;
+		}
+		float       ratio, change = 0;
+		uint32_t    resample_flags;
+		const auto* data = static_cast<const uint8_t*>(sideband_input);
+		std::memcpy(&ratio, data, sizeof(ratio));
+		if (extended) {
+			std::memcpy(&change, data + sizeof(ratio), sizeof(change));
+		}
+		std::memcpy(&resample_flags, data + size - sizeof(resample_flags), sizeof(resample_flags));
+		result.result = state->resampler.SetParameters(ratio, change, resample_flags);
+		if (result.result != OK) {
+			return result;
+		}
+		sideband_input = data + size;
+		sideband_input_size -= size;
+	}
 	if ((flags & AJM_FLAG_CONTROL_RESET) != 0) {
 		state->gapless.Reset();
 		decoder->Reset();
@@ -524,6 +554,7 @@ static AjmDecodeResult AjmControlInstance(uint32_t instance, uint64_t flags,
 		result = decoder->Initialize(sideband_input, sideband_input_size);
 	}

+	result.resample_info = state->resampler.GetInfo();
 	return result;
 }

@@ -572,6 +603,7 @@ static AjmSidebandFormat AjmGetInstanceFormat(uint32_t          instance,
 	if (out_result != nullptr) {
 		*out_result        = state->decoder->MakeResult();
 		out_result->format = state->decoder->GetFormat();
+		out_result->resample_info = state->resampler.GetInfo();
 	}
 	return state->decoder->GetFormat();
 }
@@ -921,7 +953,13 @@ int KYTY_SYSV_ABI AjmBatchJobSetResampleParameters(AjmBatchInfo* info, uint32_t
                                                   float ratio, uint32_t flags, void* result) {
 	PRINT_NAME();
 	EXIT_NOT_IMPLEMENTED(info == nullptr);
-	AjmWriteResult(result, AJM_SIDEBAND_RESULT_SIZE);
+	const struct {
+		float    ratio;
+		uint32_t flags;
+	} params {ratio, flags};
+	const auto decoded = AjmControlInstance(instance, AJM_FLAG_CONTROL_RESAMPLE, &params,
+	                                        sizeof(params), nullptr, nullptr);
+	AjmWriteResult(result, AJM_SIDEBAND_RESULT_SIZE, decoded);
 	LOGF("\t instance = 0x%08" PRIx32 ", ratio = %f, flags = 0x%08" PRIx32 "\n", instance,
 	     static_cast<double>(ratio), flags);
 	return AjmAppendJob(info, AJM_JOB_SET_RESAMPLE_PARAMETERS_SIZE, "set-resample");
@@ -933,7 +971,14 @@ int KYTY_SYSV_ABI AjmBatchJobSetResampleParametersEx(AjmBatchInfo* info, uint32_
                                                     void* result) {
 	PRINT_NAME();
 	EXIT_NOT_IMPLEMENTED(info == nullptr);
-	AjmWriteResult(result, AJM_SIDEBAND_RESULT_SIZE);
+	const struct {
+		float    ratio;
+		float    change;
+		uint32_t flags;
+	} params {ratio_start, ratio_change_per_sample, flags};
+	const auto decoded = AjmControlInstance(instance, AJM_FLAG_CONTROL_RESAMPLE_EX, &params,
+	                                        sizeof(params), nullptr, nullptr);
+	AjmWriteResult(result, AJM_SIDEBAND_RESULT_SIZE, decoded);
 	LOGF("\t instance = 0x%08" PRIx32 ", ratio_start = %f, ratio_change = %f, flags = 0x%08" PRIx32
 	     "\n",
 	     instance, static_cast<double>(ratio_start), static_cast<double>(ratio_change_per_sample),
@@ -952,7 +997,7 @@ int KYTY_SYSV_ABI AjmBatchJobGetResampleInfo(AjmBatchInfo* info, uint32_t instan
 	if (result != nullptr) {
 		auto* resample  = reinterpret_cast<AjmSidebandResampleInfo*>(static_cast<uint8_t*>(result) +
 		                                                             AJM_SIDEBAND_RESULT_SIZE);
-		resample->ratio = 1.0f;
+		*resample = decode_result.resample_info;
 	}
 	LOGF("\t instance = 0x%08" PRIx32 "\n", instance);
 	return AjmAppendJob(info, AJM_JOB_RUN_SIZE, "get-resample-info");
@@ -56,15 +56,22 @@ enum class AjmSampleEncoding : uint32_t {
 	Float = 2,
 };

+struct AjmSidebandResampleInfo {
+	float    ratio       = 1.0f;
+	int32_t  num_samples = 0;
+	uint32_t reserved[8] {};
+};
+
 struct AjmDecodeResult {
-	int32_t           result                = OK;
-	int32_t           internal_result       = OK;
-	size_t            input_consumed        = 0;
-	size_t            output_written        = 0;
-	uint64_t          total_decoded_samples = 0;
-	uint32_t          frames                = 0;
-	uint32_t          frames_per_packet     = 0;
-	AjmSidebandFormat format {};
+	int32_t                 result                = OK;
+	int32_t                 internal_result       = OK;
+	size_t                  input_consumed        = 0;
+	size_t                  output_written        = 0;
+	uint64_t                total_decoded_samples = 0;
+	uint32_t                frames                = 0;
+	uint32_t                frames_per_packet     = 0;
+	AjmSidebandFormat       format {};
+	AjmSidebandResampleInfo resample_info {};
 };

 constexpr int32_t AJM_RESULT_NOT_INITIALIZED    = 0x00000001;
@@ -0,0 +1,182 @@
+#pragma once
+
+#include "libs/ajm/decoder.h"
+
+#include <algorithm>
+#include <cmath>
+#include <cstring>
+#include <vector>
+
+namespace Libs::Audio::Ajm {
+
+constexpr uint64_t AJM_INSTANCE_FLAG_GAPLESS_LOOP           = 1ull << 10u;
+constexpr uint64_t AJM_INSTANCE_FLAG_RESAMPLE               = 1ull << 12u;
+constexpr uint64_t AJM_INSTANCE_FLAG_SEAMLESS_LOOP          = 1ull << 13u;
+constexpr uint32_t AJM_RESAMPLE_FLAG_RESET                  = 1u;
+constexpr uint32_t AJM_RESAMPLE_FLAG_END                    = 2u;
+constexpr uint32_t AJM_RESAMPLE_FLAG_IGNORE_RATIO_PARAMETER = 8u;
+
+class AjmResampler {
+public:
+	int32_t SetParameters(float ratio, float change, uint32_t flags) {
+		if ((flags & AJM_RESAMPLE_FLAG_IGNORE_RATIO_PARAMETER) == 0 &&
+		    (!std::isfinite(ratio) || ratio < 0 || ratio > 4 || !std::isfinite(change))) {
+			return AJM_RESULT_INVALID_PARAMETER;
+		}
+		if ((flags & AJM_RESAMPLE_FLAG_RESET) != 0) {
+			m_pcm.clear();
+			m_position = 0;
+			m_end      = false;
+		}
+		if ((flags & AJM_RESAMPLE_FLAG_IGNORE_RATIO_PARAMETER) == 0) {
+			m_ratio  = ratio;
+			m_change = std::abs(change) < 0.000001f ? 0 : change;
+		}
+		m_end |= (flags & AJM_RESAMPLE_FLAG_END) != 0;
+		return OK;
+	}
+
+	[[nodiscard]] AjmSidebandResampleInfo GetInfo() const {
+		AjmSidebandResampleInfo info;
+		info.ratio       = static_cast<float>(m_ratio);
+		info.num_samples = m_end && m_position >= Samples()
+		                       ? -1
+		                       : static_cast<int32_t>(std::ceil(Samples() - m_position));
+		return info;
+	}
+
+	AjmDecodeResult Decode(AjmDecoder& decoder, uint64_t flags, const void* input,
+	                       size_t input_size, void* output, size_t output_size,
+	                       bool multiple_frames, AjmGaplessState* gapless) {
+		auto result = decoder.MakeResult();
+		if (result.format.channel_num == 0) {
+			return decoder.Decode(input, input_size, output, output_size, multiple_frames, gapless);
+		}
+		bool stop_decoding = false;
+		for (;;) {
+			const auto format       = decoder.GetFormat();
+			const auto encoding     = static_cast<AjmSampleEncoding>(format.sample_encoding);
+			const auto sample_bytes = format.channel_num * AjmBytesPerSample(encoding);
+			if (!m_pcm.empty() && sample_bytes != m_sample_bytes) {
+				result.result = AJM_RESULT_INVALID_PARAMETER;
+				break;
+			}
+			if (output == nullptr || output_size - result.output_written < sample_bytes) {
+				if (m_position < Samples() || result.input_consumed < input_size) {
+					result.result |= AJM_RESULT_NOT_ENOUGH_ROOM;
+				}
+				break;
+			}
+
+			const auto index     = static_cast<size_t>(m_position);
+			const auto fraction  = m_position - static_cast<double>(index);
+			const bool need_next = fraction != 0 && index + 1 >= Samples();
+			if (index >= Samples() || need_next) {
+				Compact();
+				if (!stop_decoding && input != nullptr && result.input_consumed < input_size &&
+				    (gapless == nullptr || !gapless->IsEnd())) {
+					// A decode job can produce up to six 20 ms Opus frames in one packet.
+					const auto channels = std::max<uint32_t>(format.channel_num, flags & 0x7fu);
+					m_frame.resize(5760 * channels * AjmBytesPerSample(encoding));
+					auto decoded =
+					    decoder.Decode(static_cast<const uint8_t*>(input) + result.input_consumed,
+						               input_size - result.input_consumed, m_frame.data(),
+						               m_frame.size(), false, gapless);
+					result.input_consumed += decoded.input_consumed;
+					result.frames += decoded.frames;
+					result.frames_per_packet     = decoded.frames_per_packet;
+					result.total_decoded_samples = decoded.total_decoded_samples;
+					result.format                = decoded.format;
+					result.result |= decoded.result;
+					result.internal_result = decoded.internal_result;
+					m_sample_bytes = decoded.format.channel_num * AjmBytesPerSample(encoding);
+					m_pcm.insert(m_pcm.end(), m_frame.begin(),
+					             m_frame.begin() + decoded.output_written);
+					stop_decoding =
+					    !multiple_frames || decoded.result != OK || decoded.input_consumed == 0;
+					if (gapless != nullptr && gapless->IsEnd() &&
+					    (flags & AJM_INSTANCE_FLAG_GAPLESS_LOOP) != 0) {
+						gapless->Reset();
+						stop_decoding |= (flags & AJM_INSTANCE_FLAG_SEAMLESS_LOOP) == 0;
+					}
+					continue;
+				}
+				if (!m_end || m_position >= Samples()) {
+					if (!m_end && !stop_decoding && (gapless == nullptr || !gapless->IsEnd())) {
+						result.result |= AJM_RESULT_PARTIAL_INPUT;
+					}
+					break;
+				}
+			}
+
+			const auto  current = static_cast<size_t>(m_position);
+			const auto  next    = std::min(current + 1, Samples() - 1);
+			auto*       dst     = static_cast<uint8_t*>(output) + result.output_written;
+			const auto* a       = m_pcm.data() + current * sample_bytes;
+			const auto* b       = m_pcm.data() + next * sample_bytes;
+			const auto  weight  = m_position - static_cast<double>(current);
+			if (weight == 0) {
+				std::memcpy(dst, a, sample_bytes);
+			} else {
+				switch (encoding) {
+					case AjmSampleEncoding::S16:
+						Interpolate<int16_t>(dst, a, b, format.channel_num, weight);
+						break;
+					case AjmSampleEncoding::S32:
+						Interpolate<int32_t>(dst, a, b, format.channel_num, weight);
+						break;
+					case AjmSampleEncoding::Float:
+						Interpolate<float>(dst, a, b, format.channel_num, weight);
+						break;
+				}
+			}
+			result.output_written += sample_bytes;
+			m_position += m_ratio;
+			m_ratio = std::clamp(m_ratio + m_change, 0.0, 4.0);
+			if (m_ratio == 0 || m_ratio == 4) {
+				m_change = 0;
+			}
+		}
+		Compact();
+		result.resample_info = GetInfo();
+		return result;
+	}
+
+private:
+	template <typename T>
+	static void Interpolate(uint8_t* dst, const uint8_t* a, const uint8_t* b, uint32_t channels,
+	                        double weight) {
+		for (uint32_t channel = 0; channel < channels; ++channel) {
+			T first, second;
+			std::memcpy(&first, a + channel * sizeof(T), sizeof(T));
+			std::memcpy(&second, b + channel * sizeof(T), sizeof(T));
+			const auto value =
+			    static_cast<T>(first + (static_cast<double>(second) - first) * weight);
+			std::memcpy(dst + channel * sizeof(T), &value, sizeof(T));
+		}
+	}
+
+	[[nodiscard]] size_t Samples() const {
+		return m_sample_bytes == 0 ? 0 : m_pcm.size() / m_sample_bytes;
+	}
+
+	void Compact() {
+		const auto consumed = std::min(static_cast<size_t>(m_position), Samples());
+		const auto bytes    = consumed * m_sample_bytes;
+		if (bytes != 0) {
+			std::memmove(m_pcm.data(), m_pcm.data() + bytes, m_pcm.size() - bytes);
+			m_pcm.resize(m_pcm.size() - bytes);
+			m_position -= static_cast<double>(consumed);
+		}
+	}
+
+	std::vector<uint8_t> m_pcm;
+	std::vector<uint8_t> m_frame;
+	size_t               m_sample_bytes = 0;
+	double               m_position     = 0;
+	double               m_ratio        = 1;
+	double               m_change       = 0;
+	bool                 m_end          = false;
+};
+
+} // namespace Libs::Audio::Ajm
@@ -519,7 +519,20 @@ void GameController::SetDevices(int active_id) {
 	m_trigger_state      = {};
 }

+// Input that changes no button or touch contact updates the newest unread sample, so a burst of
+// stick or sensor events between two reads cannot push unread button or touch changes out.
 void GameController::AddState() {
+	if (m_states_num != 0) {
+		const auto newest = (m_first_state + m_states_num - 1) % STATES_MAX;
+		auto&      last   = m_states[newest];
+		if (!m_obtained[newest] && last.buttons == m_state.buttons &&
+		    last.touch[0].down == m_state.touch[0].down &&
+		    last.touch[1].down == m_state.touch[1].down) {
+			last = m_state;
+			return;
+		}
+	}
+
 	if (m_states_num >= STATES_MAX) {
 		m_states_num  = STATES_MAX - 1;
 		m_first_state = (m_first_state + 1) % STATES_MAX;
@@ -13,6 +13,7 @@
 #include <atomic>
 #include <cmath>
 #include <cstring>
+#include <span>
 #include <vector>

 namespace Libs::Audio {
@@ -168,17 +169,12 @@ struct AudioOut2MasteringStatesV2 {

 static std::atomic_uint64_t g_audioout2_next_context {1};
 static std::atomic_uint64_t g_audioout2_next_port {1};
-static AudioOut2UserHandle g_audioout2_next_user = 1;
+static AudioOut2UserHandle  g_audioout2_next_user = 1;

-struct AudioOut2ContextState {
-	bool                   used                = false;
-	AudioOut2ContextHandle handle              = 0;
-	uint32_t               queue_depth         = 4;
-	uint32_t               queued              = 0;
-	uint32_t               num_grains          = 512;
-	uint64_t               last_update         = 0;
-	int                    object_audio_handle = 0;
-};
+static constexpr uint32_t AUDIOOUT2_CHANNEL_PORTS_MAX = 32;
+static constexpr uint32_t AUDIOOUT2_OBJECT_PORTS_MAX  = 512;
+static constexpr uint32_t AUDIOOUT2_PORT_SLOT_BITS    = 10;
+static constexpr uint32_t AUDIOOUT2_CONTEXT_SLOT_BITS = 4;

 struct AudioOut2PortStateEntry {
 	bool                   used          = false;
@@ -196,6 +192,19 @@ struct AudioOut2PortStateEntry {
 	bool                   attributes_locked = false;
 	std::vector<uint8_t>   pcm_data;
 };
+
+struct AudioOut2ContextState {
+	bool                   used                = false;
+	AudioOut2ContextHandle handle              = 0;
+	uint32_t               queue_depth         = 4;
+	uint32_t               queued              = 0;
+	uint32_t               num_grains          = 512;
+	uint64_t               last_update         = 0;
+	int                    object_audio_handle = 0;
+	uint32_t               max_ports           = 0;
+	// Channel slots precede object slots; each partition has its own creation limit.
+	std::vector<AudioOut2PortStateEntry> ports;
+};
 struct AudioOut2SpeakerArrayState {
 	bool                        used          = false;
 	uint32_t                    num_speakers  = 0;
@@ -223,13 +232,16 @@ struct AudioOut2Submission {
 static Common::Mutex                              g_audioout2_state_mutex;
 static std::array<AudioOut2ContextState, 16>      g_audioout2_contexts;
 static std::array<AudioOut2Submission, 16>        g_audioout2_submissions;
-static std::array<AudioOut2PortStateEntry, 256>   g_audioout2_ports;
-static std::vector<AudioOut2UserHandle>          g_audioout2_users;
+static std::vector<AudioOut2UserHandle>           g_audioout2_users;
 static Common::Mutex                              g_audioout2_speaker_array_mutex;
 static std::array<AudioOut2SpeakerArrayState, 32> g_audioout2_speaker_arrays;
 static Common::Mutex                              g_audioout2_latency_mutex;
 static std::array<AudioOut2LatencyState, 16>      g_audioout2_latencies;

+static_assert(AUDIOOUT2_CHANNEL_PORTS_MAX + AUDIOOUT2_OBJECT_PORTS_MAX <=
+              (1u << AUDIOOUT2_PORT_SLOT_BITS));
+static_assert(g_audioout2_contexts.size() == (1u << AUDIOOUT2_CONTEXT_SLOT_BITS));
+
 static constexpr int AUDIO_OUT2_ERROR_NOT_READY                   = -2144960504; /* 0x80268008 */
 static constexpr int AUDIO_OUT2_ERROR_PORT_FULL                   = -2144960494; /* 0x80268012 */
 static constexpr int AUDIO_OUT2_ERROR_INVALID_PARAM               = -2144960511; /* 0x80268001 */
@@ -262,6 +274,11 @@ static AudioOut2Submission& audioout2_submission(const AudioOut2ContextState& co
 	return g_audioout2_submissions[&context - g_audioout2_contexts.data()];
 }

+// Ticks of a sixth of a microsecond: one 48 kHz sample lasts exactly 125 of them.
+static uint64_t audioout2_clock() {
+	return LibKernel::KernelGetProcessTime() * 6;
+}
+
 static uint32_t audioout2_grain_micros(uint32_t grains) {
 	const auto sample_count = (grains == 0 ? 512u : grains);
 	return std::max<uint32_t>((sample_count * 1000000u) / 48000u, 1000u);
@@ -331,49 +348,51 @@ static void audioout2_update_context_locked(AudioOut2ContextState* state) {
 		return;
 	}

-	const auto now = LibKernel::KernelGetProcessTime();
+	const auto now = audioout2_clock();
 	if (state->last_update == 0 || state->queued == 0) {
 		state->last_update = now;
 		return;
 	}
-
-	const auto grain_micros = static_cast<uint64_t>(audioout2_grain_micros(state->num_grains));
-	if (now <= state->last_update || grain_micros == 0) {
+	if (now <= state->last_update) {
 		return;
 	}

-	const auto elapsed = now - state->last_update;
-	const auto drained = std::min<uint64_t>(state->queued, elapsed / grain_micros);
+	const auto grain   = uint64_t {state->num_grains} * 125;
+	const auto drained = std::min<uint64_t>(state->queued, (now - state->last_update) / grain);
 	if (drained == 0) {
 		return;
 	}

 	state->queued -= static_cast<uint32_t>(drained);
-	state->last_update += drained * grain_micros;
+	state->last_update += drained * grain;
 	if (state->queued == 0) {
 		state->last_update = now;
 	}
 }

 static AudioOut2PortStateEntry* audioout2_find_port_locked(AudioOut2PortHandle port) {
-	for (auto& state: g_audioout2_ports) {
-		if (state.used && state.handle == port) {
-			return &state;
-		}
+	const auto context_slot =
+	    (port >> AUDIOOUT2_PORT_SLOT_BITS) & ((1u << AUDIOOUT2_CONTEXT_SLOT_BITS) - 1u);
+	const auto port_slot = port & ((1u << AUDIOOUT2_PORT_SLOT_BITS) - 1u);
+	auto&      context   = g_audioout2_contexts[context_slot];
+	if (context.used && port_slot < context.ports.size()) {
+		auto& state = context.ports[port_slot];
+		if (state.used && state.handle == port) return &state;
 	}

 	return nullptr;
 }

 static size_t audioout2_pcm_size(const AudioOut2PortStateEntry& state) {
-	const auto bytes_per_sample = (state.data_format & 0x7fu) == 1 ? sizeof(int16_t) : sizeof(float);
+	const auto bytes_per_sample =
+	    (state.data_format & 0x7fu) == 1 ? sizeof(int16_t) : sizeof(float);
 	return static_cast<size_t>(state.samples_num) *
 	       audioout2_data_format_channels(state.data_format) * bytes_per_sample;
 }

 static bool audioout2_context_has_queueable_device(const AudioOut2ContextState& context) {
-	for (const auto& state: g_audioout2_ports) {
-		if (state.used && state.context == context.handle && !state.pcm_data.empty()) {
+	for (const auto& state: context.ports) {
+		if (state.used && !state.pcm_data.empty()) {
 			const int handle = audioout2_port_type_is_object(state.port_type)
 			                       ? context.object_audio_handle
 			                       : state.audio_handle;
@@ -427,18 +446,21 @@ static void audioout2_queue_context_audio(AudioOut2ContextHandle ctx, AudioOut2S
 	out.object_mix.clear();
 	std::vector<AudioInternal::OutputParam> params;
 	params.reserve(AudioInternal::OUT_PORTS_MAX);
-	for (const auto& state: g_audioout2_ports) {
-		if (!state.used || state.context != ctx || state.pcm_data.empty()) {
+	// A port's PCM covers one grain: the push consumes it, so a port the game stops feeding
+	// falls silent.
+	for (auto& state: context->ports) {
+		if (!state.used || state.pcm_data.empty()) {
 			continue;
 		}
 		if (audioout2_port_type_is_object(state.port_type)) {
 			audioout2_mix_object(out.object_mix, context->num_grains, state);
 		} else if (state.audio_handle > 0 && params.size() < AudioInternal::OUT_PORTS_MAX) {
 			const auto i = params.size();
-			out.pcm[i]   = state.pcm_data;
+			out.pcm[i].swap(state.pcm_data);
 			out.gains[i] = state.gains;
 			params.push_back({state.audio_handle, out.pcm[i].data(), out.gains[i].data()});
 		}
+		state.pcm_data.clear();
 	}
 	if (!out.object_mix.empty() && params.size() < AudioInternal::OUT_PORTS_MAX) {
 		params.push_back({context->object_audio_handle, out.object_mix.data()});
@@ -467,22 +489,28 @@ int KYTY_SYSV_ABI AudioOut2ContextResetParam(AudioOut2ContextParam* params) {
 	EXIT_NOT_IMPLEMENTED(params == nullptr);

 	std::memset(params, 0, sizeof(AudioOut2ContextParam));
-	params->max_ports              = 256;
-	params->max_object_ports       = 256;
-	params->guarantee_object_ports = 0;
-	params->queue_depth            = 4;
-	params->num_grains             = 512;
-	params->flags                  = 1;
+	params->max_ports   = 8;
+	params->queue_depth = 1;
+	params->num_grains  = 256;

 	return OK;
 }

+static bool audioout2_valid_context_capacity(const AudioOut2ContextParam& params) {
+	return params.max_ports <= AUDIOOUT2_CHANNEL_PORTS_MAX &&
+	       params.max_object_ports <= AUDIOOUT2_OBJECT_PORTS_MAX &&
+	       params.guarantee_object_ports <= params.max_object_ports;
+}
+
 int KYTY_SYSV_ABI AudioOut2ContextQueryMemory(const AudioOut2ContextParam* params,
                                              size_t*                      memory_size) {
 	PRINT_NAME();

 	EXIT_NOT_IMPLEMENTED(params == nullptr);
 	EXIT_NOT_IMPLEMENTED(memory_size == nullptr);
+	if (!audioout2_valid_context_capacity(*params)) {
+		return AUDIO_OUT2_ERROR_INVALID_PARAM;
+	}

 	const auto queue_depth = (params->queue_depth == 0 ? 4u : params->queue_depth);
 	*memory_size           = 0x10000u + static_cast<size_t>(queue_depth) * 0x590u;
@@ -498,6 +526,9 @@ int KYTY_SYSV_ABI AudioOut2ContextCreate(const AudioOut2ContextParam* params, vo

 	EXIT_NOT_IMPLEMENTED(params == nullptr);
 	EXIT_NOT_IMPLEMENTED(ctx == nullptr);
+	if (!audioout2_valid_context_capacity(*params)) {
+		return AUDIO_OUT2_ERROR_INVALID_PARAM;
+	}

 	const auto num_grains = params->num_grains == 0 ? 512u : params->num_grains;
 	const int  object_audio_handle =
@@ -524,8 +555,10 @@ int KYTY_SYSV_ABI AudioOut2ContextCreate(const AudioOut2ContextParam* params, vo
 	state->handle              = *ctx;
 	state->queue_depth         = params->queue_depth == 0 ? 4u : params->queue_depth;
 	state->num_grains          = num_grains;
-	state->last_update         = LibKernel::KernelGetProcessTime();
+	state->last_update         = audioout2_clock();
 	state->object_audio_handle = object_audio_handle;
+	state->max_ports           = params->max_ports;
+	state->ports.resize(params->max_ports + params->max_object_ports);
 	g_audioout2_state_mutex.Unlock();

 	LOGF("\t buffer      = 0x%016" PRIx64 "\n"
@@ -541,8 +574,8 @@ int KYTY_SYSV_ABI AudioOut2ContextDestroy(AudioOut2ContextHandle ctx) {
 	PRINT_NAME();
 	LOGF("\t ctx = 0x%016" PRIx64 "\n", ctx);

-	std::array<int, 257> audio_handles {};
-	size_t audio_handles_num = 0;
+	std::array<int, AUDIOOUT2_CHANNEL_PORTS_MAX + 1> audio_handles {};
+	size_t                                           audio_handles_num = 0;

 	g_audioout2_state_mutex.Lock();
 	auto* state = audioout2_find_context_locked(ctx);
@@ -550,15 +583,12 @@ int KYTY_SYSV_ABI AudioOut2ContextDestroy(AudioOut2ContextHandle ctx) {
 		if (state->object_audio_handle > 0) {
 			audio_handles[audio_handles_num++] = state->object_audio_handle;
 		}
-		*state = AudioOut2ContextState {};
-	}
-	for (auto& port_state: g_audioout2_ports) {
-		if (port_state.used && port_state.context == ctx) {
-			if (port_state.audio_handle > 0 && audio_handles_num < audio_handles.size()) {
+		for (const auto& port_state: state->ports) {
+			if (port_state.used && port_state.audio_handle > 0) {
 				audio_handles[audio_handles_num++] = port_state.audio_handle;
 			}
-			port_state = AudioOut2PortStateEntry {};
 		}
+		*state = AudioOut2ContextState {};
 	}
 	g_audioout2_state_mutex.Unlock();

@@ -586,9 +616,7 @@ int KYTY_SYSV_ABI AudioOut2ContextAdvance(AudioOut2ContextHandle ctx) {
 	g_audioout2_state_mutex.Lock();
 	if (auto* state = audioout2_find_context_locked(ctx); state != nullptr) {
 		audioout2_update_context_locked(state);
-	}
-	for (auto& port: g_audioout2_ports) {
-		if (port.used && port.context == ctx) {
+		for (auto& port: state->ports) {
 			port.attributes_locked = false;
 		}
 	}
@@ -612,7 +640,7 @@ int KYTY_SYSV_ABI AudioOut2ContextPush(AudioOut2ContextHandle ctx, uint32_t bloc
 			sleep_micros = audioout2_grain_micros(state->num_grains);
 			if (state->queued < state->queue_depth || use_device_clock) {
 				if (state->queued == 0) {
-					state->last_update = LibKernel::KernelGetProcessTime();
+					state->last_update = audioout2_clock();
 				}
 				if (state->queued < state->queue_depth) {
 					state->queued++;
@@ -663,21 +691,19 @@ int KYTY_SYSV_ABI AudioOut2PortCreate(AudioOut2ContextHandle ctx, const AudioOut
 	EXIT_NOT_IMPLEMENTED(params == nullptr);
 	EXIT_NOT_IMPLEMENTED(port == nullptr);

-	const auto next_port    = g_audioout2_next_port.fetch_add(1, std::memory_order_relaxed);
 	const auto audio_format = audioout2_data_format_to_audio_format(params->data_format);
 	const auto audio_type   = audioout2_port_type_to_audio_out_type(params->port_type);
+	const bool is_object    = audioout2_port_type_is_object(params->port_type);

 	g_audioout2_state_mutex.Lock();
-	const auto* context_state = audioout2_find_context_locked(ctx);
+	auto* context_state = audioout2_find_context_locked(ctx);
 	if (context_state == nullptr) {
 		g_audioout2_state_mutex.Unlock();
 		return AUDIO_OUT2_ERROR_INVALID_PARAM;
 	}
-	if (audioout2_port_type_is_object(params->port_type) &&
-	    (params->port_type != 0x100 || context_state->object_audio_handle == 0 ||
-	     params->sampling_freq != 48000 ||
-	     (audio_format != AudioInternal::Format::FloatMono &&
-	      audio_format != AudioInternal::Format::Signed16bitMono))) {
+	if (is_object && (params->port_type != 0x100 || params->sampling_freq != 48000 ||
+	                  (audio_format != AudioInternal::Format::FloatMono &&
+	                   audio_format != AudioInternal::Format::Signed16bitMono))) {
 		g_audioout2_state_mutex.Unlock();
 		return AUDIO_OUT2_ERROR_INVALID_PARAM;
 	}
@@ -688,42 +714,45 @@ int KYTY_SYSV_ABI AudioOut2PortCreate(AudioOut2ContextHandle ctx, const AudioOut
 		g_audioout2_state_mutex.Unlock();
 		return AUDIO_OUT2_ERROR_INVALID_PARAM;
 	}
-	AudioOut2PortStateEntry* port_state = nullptr;
-	for (auto& candidate: g_audioout2_ports) {
-		if (!candidate.used) {
-			port_state = &candidate;
-			break;
-		}
-	}
-	if (port_state != nullptr) {
-		*port_state               = AudioOut2PortStateEntry {};
-		port_state->used          = true;
-		port_state->handle        = next_port;
-		port_state->context       = ctx;
-		port_state->user          = params->user_handle;
-		port_state->port_type     = params->port_type;
-		port_state->data_format   = params->data_format;
-		port_state->sampling_freq = params->sampling_freq;
-		port_state->samples_num   = samples_num;
-		port_state->audio_format  = audio_format;
-		port_state->gains.fill(1.0f);
-	}
-	g_audioout2_state_mutex.Unlock();
-
-	if (port_state == nullptr) {
+	auto slots = std::span(context_state->ports);
+	slots =
+	    is_object ? slots.subspan(context_state->max_ports) : slots.first(context_state->max_ports);
+	const auto available = std::find_if(slots.begin(), slots.end(),
+	                                    [](const auto& candidate) { return !candidate.used; });
+	if (available == slots.end()) {
+		g_audioout2_state_mutex.Unlock();
 		return AUDIO_OUT2_ERROR_PORT_FULL;
 	}
+	const auto serial       = g_audioout2_next_port.fetch_add(1, std::memory_order_relaxed);
+	const auto context_slot = static_cast<uint64_t>(context_state - g_audioout2_contexts.data());
+	const auto port_slot    = static_cast<uint64_t>(&*available - context_state->ports.data());
+	const auto next_port    = (serial << (AUDIOOUT2_PORT_SLOT_BITS + AUDIOOUT2_CONTEXT_SLOT_BITS)) |
+	                          (context_slot << AUDIOOUT2_PORT_SLOT_BITS) | port_slot;
+	auto*      port_state   = &*available;
+	*port_state             = AudioOut2PortStateEntry {};
+	port_state->used        = true;
+	port_state->handle      = next_port;
+	port_state->context     = ctx;
+	port_state->user        = params->user_handle;
+	port_state->port_type   = params->port_type;
+	port_state->data_format = params->data_format;
+	port_state->sampling_freq = params->sampling_freq;
+	port_state->samples_num   = samples_num;
+	port_state->audio_format  = audio_format;
+	port_state->gains.fill(1.0f);
+	g_audioout2_state_mutex.Unlock();

 	int audio_handle = 0;

-	if (audio_format != AudioInternal::Format::Unknown &&
-	    !audioout2_port_type_is_object(params->port_type)) {
+	if (audio_format != AudioInternal::Format::Unknown && !is_object) {
 		audio_handle = AudioInternal::AudioOutOpen(audio_type, samples_num, params->sampling_freq,
 		                                           audio_format);
 	}

 	g_audioout2_state_mutex.Lock();
-	const bool reserved = port_state->used && port_state->handle == next_port;
+	// Context destruction may have freed or replaced the slots while the backend opened.
+	port_state          = audioout2_find_port_locked(next_port);
+	const bool reserved = port_state != nullptr;
 	if (reserved) {
 		port_state->audio_handle = audio_handle;
 	}
@@ -735,7 +764,7 @@ int KYTY_SYSV_ABI AudioOut2PortCreate(AudioOut2ContextHandle ctx, const AudioOut

 	*port = next_port;

-	if (next_port <= 16 || (next_port % 600) == 0) {
+	if (serial <= 16 || (serial % 600) == 0) {
 		PRINT_NAME();
 		LOGF("\t ctx           = 0x%016" PRIx64 "\n"
 		     "\t port          = 0x%016" PRIx64 "\n"
@@ -901,21 +930,23 @@ int KYTY_SYSV_ABI AudioOut2UserDestroy(AudioOut2UserHandle handle) {
 	PRINT_NAME();
 	LOGF("\t handle = 0x%016" PRIx64 "\n", static_cast<uint64_t>(handle));
 	Common::LockGuard lock(g_audioout2_state_mutex);
-	const auto it = std::find(g_audioout2_users.begin(), g_audioout2_users.end(), handle);
+	const auto        it = std::find(g_audioout2_users.begin(), g_audioout2_users.end(), handle);
 	if (it == g_audioout2_users.end()) {
 		return AUDIO_OUT2_ERROR_INVALID_PARAM;
 	}
-	if (std::any_of(g_audioout2_ports.begin(), g_audioout2_ports.end(),
-	                [handle](const auto& port) { return port.used && port.user == handle; })) {
-		return AUDIO_OUT2_ERROR_BUSY;
+	for (const auto& context: g_audioout2_contexts) {
+		if (std::any_of(context.ports.begin(), context.ports.end(),
+		                [handle](const auto& port) { return port.used && port.user == handle; })) {
+			return AUDIO_OUT2_ERROR_BUSY;
+		}
 	}
 	g_audioout2_users.erase(it);
 	return OK;
 }

 int KYTY_SYSV_ABI AudioOut2UserGetSupportedAttributes(AudioOut2UserHandle handle,
-                                                       uint32_t* context_attributes,
-                                                       uint32_t* port_attributes) {
+                                                      uint32_t*           context_attributes,
+                                                      uint32_t*           port_attributes) {
 	if (context_attributes == nullptr || port_attributes == nullptr) {
 		return AUDIO_OUT2_ERROR_INVALID_PARAM;
 	}
@@ -38,8 +38,8 @@ struct Global {
 };

 struct Shadow {
-	uintptr_t start = 0;
-	int       count = 0;
+	uintptr_t start         = 0;
+	uint64_t  covered_bytes = 0;
 };

 struct AsanContext {
@@ -50,6 +50,14 @@ struct AsanContext {
 static uint32_t     g_fake_stack = 0;
 static AsanContext* g_ctx        = nullptr;

+// Allocation and release callbacks can partition the same guest range differently.
+static uint64_t ShadowCoverage(uintptr_t page, uintptr_t min_addr, uintptr_t max_addr) {
+	const auto guest_start = (page - 0x10000000000ull) << 3u;
+	const auto guest_end   = guest_start + 0x80000ull - 1;
+	return std::min<uintptr_t>(max_addr, guest_end) - std::max<uintptr_t>(min_addr, guest_start) +
+	       1;
+}
+
 static void AllocShadow(uintptr_t min_addr, uintptr_t max_addr) {
 	EXIT_IF(g_ctx == nullptr);

@@ -67,10 +75,11 @@ static void AllocShadow(uintptr_t min_addr, uintptr_t max_addr) {
 	     static_cast<uint64_t>(min_shadow), static_cast<uint64_t>(size));

 	for (uintptr_t page = min_shadow; page <= max_shadow; page += page_size) {
-		const auto it = std::find_if(g_ctx->shadows.begin(), g_ctx->shadows.end(),
-		                             [page](const Shadow& s) { return s.start == page; });
+		const auto bytes = ShadowCoverage(page, min_addr, max_addr);
+		const auto it    = std::find_if(g_ctx->shadows.begin(), g_ctx->shadows.end(),
+		                                [page](const Shadow& s) { return s.start == page; });
 		if (it != g_ctx->shadows.end()) {
-			it->count++;
+			it->covered_bytes += bytes;
 		} else {
 			if (!Common::VirtualMemory::AllocFixed(page, page_size,
 			                                       Common::VirtualMemory::Mode::ReadWrite)) {
@@ -78,8 +87,8 @@ static void AllocShadow(uintptr_t min_addr, uintptr_t max_addr) {
 			}

 			Shadow s;
-			s.start = page;
-			s.count = 1;
+			s.start         = page;
+			s.covered_bytes = bytes;
 			g_ctx->shadows.push_back(s);
 		}
 	}
@@ -105,14 +114,15 @@ static void FreeShadow(uintptr_t min_addr, uintptr_t max_addr) {
 		const auto it = std::find_if(g_ctx->shadows.begin(), g_ctx->shadows.end(),
 		                             [page](const Shadow& s) { return s.start == page; });
 		if (it != g_ctx->shadows.end()) {
-			it->count--;
-
-			if (it->count <= 0) {
+			const auto bytes = ShadowCoverage(page, min_addr, max_addr);
+			if (bytes >= it->covered_bytes) {
 				if (!Common::VirtualMemory::Free(page)) {
 					EXIT("can't free shadow memory\n");
 				}

 				g_ctx->shadows.erase(it);
+			} else {
+				it->covered_bytes -= bytes;
 			}
 		}
 	}
@@ -126,6 +136,14 @@ static void KernelFree(uintptr_t addr, size_t size) {
 	FreeShadow(addr, addr + size - 1);
 }

+#if defined(KYTY_VIRTUAL_MEMORY_ALLOCATION_TESTS)
+void TestRegisterMemoryCallbacks() {
+	EXIT_IF(g_ctx != nullptr);
+	g_ctx = new AsanContext;
+	LibKernel::Memory::RegisterCallbacks(KernelAlloc, KernelFree);
+}
+#endif
+
 static void KYTY_SYSV_ABI asan_init() {
 	PRINT_NAME();

@@ -2,10 +2,12 @@
 #include "common/assert.h"
 #include "common/common.h"
 #include "common/logging/log.h"
+#include "common/stringUtils.h"
 #include "libs/errno.h"
 #include "libs/libs.h"
 #include "loader/symbolDatabase.h"

+#include <SDL3/SDL.h>
 #include <algorithm>
 #include <array>
 #include <cinttypes>
@@ -675,19 +677,6 @@ static bool font_is_italic(const FontState* font) {
 	       font->render_effect_slant != 0.0f;
 }

-static const char* fallback_font_file_name(bool bold, bool italic) {
-	if (bold && italic) {
-		return "Roboto-BoldItalic.ttf";
-	}
-	if (bold) {
-		return "Roboto-Bold.ttf";
-	}
-	if (italic) {
-		return "Roboto-Italic.ttf";
-	}
-	return "Roboto-Regular.ttf";
-}
-
 static std::vector<uint8_t> read_file_bytes(const std::filesystem::path& path) {
 	std::ifstream file(path, std::ios::binary | std::ios::ate);
 	if (!file) {
@@ -709,49 +698,34 @@ static std::vector<uint8_t> read_file_bytes(const std::filesystem::path& path) {
 	return data;
 }

-static std::vector<std::filesystem::path> fallback_font_candidates(const char* file_name) {
-	const auto rel =
-	    std::filesystem::path("3rdparty") / "tracy" / "profiler" / "src" / "font" / file_name;
-	std::vector<std::filesystem::path> paths;
-
-	auto source_root = std::filesystem::path(__FILE__).parent_path().parent_path().parent_path();
-	paths.push_back(source_root / rel);
-
-	std::error_code ec;
-	auto            cwd = std::filesystem::current_path(ec);
-	if (!ec) {
-		paths.push_back(cwd / rel);
-		paths.push_back(cwd / ".." / rel);
-		paths.push_back(cwd / ".." / ".." / rel);
+static std::vector<uint8_t> load_fallback_font_data(const char* file_name) {
+	const char* base_path = SDL_GetBasePath();
+	if (base_path == nullptr) {
+		LOGF("Font: SDL_GetBasePath failed: %s\n", SDL_GetError());
+		return {};
 	}
-
-	return paths;
-}
-
-static std::vector<uint8_t> load_fallback_font_data(bool bold, bool italic) {
-	for (const auto& path: fallback_font_candidates(fallback_font_file_name(bold, italic))) {
-		auto data = read_file_bytes(path);
-		if (!data.empty()) {
-			return data;
-		}
+	// SDL returns UTF-8; path(const char*) would read it in the Windows ANSI code page.
+	const auto path = Common::PathFromUtf8(base_path) / "assets" / "fonts" / file_name;
+	auto       data = read_file_bytes(path);
+	if (data.empty()) {
+		LOGF("Font: fallback font not found: %s\n", Common::PathToString(path).c_str());
 	}
-
-	return {};
+	return data;
 }

 static const std::vector<uint8_t>& fallback_font_data(bool bold, bool italic) {
-	static const auto regular     = load_fallback_font_data(false, false);
-	static const auto bold_data   = load_fallback_font_data(true, false);
-	static const auto italic_data = load_fallback_font_data(false, true);
-	static const auto bold_italic = load_fallback_font_data(true, true);
+	static const auto regular     = load_fallback_font_data("NotoSans-Regular.ttf");
+	static const auto bold_data   = load_fallback_font_data("NotoSans-Bold.ttf");
+	static const auto italic_data = load_fallback_font_data("NotoSans-Italic.ttf");
+	static const auto bold_italic = load_fallback_font_data("NotoSans-BoldItalic.ttf");

-	if (bold && italic && !bold_italic.empty()) {
+	if (bold && italic) {
 		return bold_italic;
 	}
-	if (bold && !bold_data.empty()) {
+	if (bold) {
 		return bold_data;
 	}
-	if (italic && !italic_data.empty()) {
+	if (italic) {
 		return italic_data;
 	}
 	return regular;
@@ -615,7 +615,12 @@ int KYTY_SYSV_ABI SaveDataDirNameSearch(const SaveDataDirNameSearchCond* cond,
 	Common::LockGuard lock(g_mount_mutex);
 	const std::string title_id = cond->title_id != nullptr ? cond->title_id->data : get_title_id();
 	const auto        root     = save_user_directory(title_id, cond->user_id);
-	std::vector<std::string> dir_list;
+	struct Found {
+		uint64_t    key = 0;
+		std::string name;
+		auto        operator<=>(const Found&) const = default;
+	};
+	std::vector<Found> dir_list;

 	if (Common::File::IsDirectoryExisting(root)) {
 		for (const auto& entry: Common::File::GetDirEntries(root)) {
@@ -625,22 +630,34 @@ int KYTY_SYSV_ABI SaveDataDirNameSearch(const SaveDataDirNameSearchCond* cond,
 				if (cond->dir_name == nullptr || cond->dir_name->data[0] == '\0' ||
 				    dir_name_match(Common::ToLower(entry.name).c_str(),
 				                   Common::ToLower(std::string(cond->dir_name->data)).c_str())) {
-					dir_list.push_back(entry.name);
+					dir_list.push_back({.name = entry.name});
 				}
 			}
 		}
 	}

-	std::sort(dir_list.begin(), dir_list.end(), [](const std::string& a, const std::string& b) {
-		return std::strcmp(a.c_str(), b.c_str()) < 0;
-	});
-
-	if (cond->order == SaveDataSortOrder::Descent) {
-		std::vector<std::string> reversed;
-		for (size_t i = dir_list.size(); i > 0; i--) {
-			reversed.push_back(dir_list[i - 1]);
+	// A save whose metadata cannot be read sorts as 0.
+	for (auto& found: dir_list) {
+		const auto directory = root / found.name;
+		switch (cond->key) {
+			case SaveDataSortKey::UserParam: {
+				SaveDataParam param {};
+				read_save_param(save_param_path(directory), &param);
+				found.key = param.user_param;
+				break;
+			}
+			case SaveDataSortKey::Blocks:
+			case SaveDataSortKey::FreeBlocks: read_save_blocks(directory, &found.key); break;
+			case SaveDataSortKey::Mtime:
+				found.key = static_cast<uint64_t>(newest_save_time(directory));
+				break;
+			default: break;
 		}
-		dir_list = reversed;
+	}
+	// Directory names break ties; descending order reverses both.
+	std::sort(dir_list.begin(), dir_list.end());
+	if (cond->order == SaveDataSortOrder::Descent) {
+		std::reverse(dir_list.begin(), dir_list.end());
 	}

 	auto max_count =
@@ -651,9 +668,9 @@ int KYTY_SYSV_ABI SaveDataDirNameSearch(const SaveDataDirNameSearchCond* cond,

 	for (size_t i = 0; i < max_count; i++) {
 		std::snprintf(result->dir_names[i].data, sizeof(result->dir_names[i].data), "%s",
-		              dir_list[i].c_str());
+		              dir_list[i].name.c_str());
 		if (result->params != nullptr) {
-			const int status = load_save_param(root / dir_list[i], &result->params[i]);
+			const int status = load_save_param(root / dir_list[i].name, &result->params[i]);
 			if (status != OK) {
 				return status;
 			}
@@ -661,7 +678,7 @@ int KYTY_SYSV_ABI SaveDataDirNameSearch(const SaveDataDirNameSearchCond* cond,
 		if (result->infos != nullptr) {
 			auto& info = result->infos[i];
 			info = {};
-			const int status = read_save_blocks(root / dir_list[i], &info.blocks);
+			const int status = read_save_blocks(root / dir_list[i].name, &info.blocks);
 			if (status != OK) {
 				return status;
 			}
@@ -546,6 +546,7 @@ struct Ngs2VoiceInternal {
 	};
 	struct Module {
 		std::vector<uint8_t>     param, work, state;
+		std::vector<float>       channel_matrix;
 		void                     KYTY_SYSV_ABI (*handler)(Ngs2UserFxProcessContext*) = nullptr;
 		std::array<uintptr_t, 3> user_data {};
 		uint32_t                 flags          = 1;
@@ -563,6 +564,7 @@ struct Ngs2VoiceInternal {
 	std::vector<std::vector<float>> matrices;
 	std::vector<Module>             modules;
 	std::vector<float>              samples;
+	std::vector<float>              mix_samples;
 	std::vector<float>              output_matrix;
 	uint32_t                        channels        = 0;
 	uint32_t                        output_id       = 0;
@@ -614,6 +616,27 @@ struct Ngs2VoiceInternal {
 		return samples.data() + size_t(id) * grain * channels;
 	}

+	void MixChannels(const std::vector<float>& matrix, uint32_t buffer, uint32_t grain) {
+		if (matrix.empty()) {
+			return;
+		}
+		EXIT_NOT_IMPLEMENTED(channels > 37);
+		auto* input = Buffer(buffer, grain);
+		mix_samples.assign(size_t(channels) * grain, 0.0f);
+		for (uint32_t dst = 0; dst < channels; ++dst) {
+			for (uint32_t src = 0; src < channels; ++src) {
+				const float level = matrix[dst * 37 + src];
+				if (level == 0.0f) {
+					continue;
+				}
+				for (uint32_t i = 0; i < grain; ++i) {
+					mix_samples[dst * grain + i] += input[src * grain + i] * level;
+				}
+			}
+		}
+		std::copy(mix_samples.begin(), mix_samples.end(), input);
+	}
+
 	void ResetSetupState() {
 		SetEvent(4);
 		std::ranges::fill(ports, Port {});
@@ -1850,6 +1873,10 @@ static void Ngs2RenderVoice(Ngs2VoiceInternal& voice, const std::vector<Ngs2Voic
 			for (size_t m = 0; m < voice.modules.size(); ++m) {
 				const auto& info   = custom.module[m];
 				auto&       module = voice.modules[m];
+				if (info.module_id == 0x17) {
+					voice.MixChannels(module.channel_matrix, info.source_buffer_id, grain);
+					continue;
+				}
 				if (info.module_id != 0x19 && info.module_id != 0x1f) {
 					if (!module.render_warned) {
 						Log::WriteToConsoleAndLog(fmt::sprintf(
@@ -1934,7 +1961,8 @@ int KYTY_SYSV_ABI Ngs2SystemRender(uintptr_t system_handle, const Ngs2RenderBuff
 		EXIT_NOT_IMPLEMENTED(voice->output_id >= num_buffer_info);
 		const auto& output = buffer_info[voice->output_id];
 		EXIT_NOT_IMPLEMENTED(output.buffer == nullptr);
-		EXIT_NOT_IMPLEMENTED((output.waveform_type != 0x12 && output.waveform_type != 0x18) ||
+		EXIT_NOT_IMPLEMENTED(output.waveform_type != 0x12 && output.waveform_type != 0x18);
+		EXIT_NOT_IMPLEMENTED(voice->rack->type == Ngs2RackType::Mastering &&
 		                     output.num_channels != voice->channels);
 		EXIT_NOT_IMPLEMENTED(output.buffer_size < grain * output.num_channels *
 		                                               (output.waveform_type == 0x12 ? sizeof(int16_t)
@@ -1954,6 +1982,9 @@ int KYTY_SYSV_ABI Ngs2SystemRender(uintptr_t system_handle, const Ngs2RenderBuff
 				const float level = voice->output_matrix.empty()
 				                        ? (src == dst ? 1.0f : 0.0f)
 				                        : voice->output_matrix.at(src * output.num_channels + dst);
+				if (level == 0.0f) {
+					continue;
+				}
 				for (uint32_t i = 0; i < grain; ++i) {
 					pcm[i * output.num_channels + dst] += voice->samples[src * grain + i] * level;
 				}
@@ -2452,7 +2483,7 @@ int KYTY_SYSV_ABI Ngs2VoiceControl(uintptr_t voice_handle, const Ngs2VoiceParamH
 				    voice->rack->option.custom_sampler.custom_rack_option.module[index].module_id;
 				auto& module = voice->modules[index];
 				if ((param->id & 0xffffffe0u) != (0x40000000u | module_id << 8u) ||
-				    (module_id != 0x19 && module_id != 0x1f)) {
+				    (module_id != 0x17 && module_id != 0x19 && module_id != 0x1f)) {
 					if (!module.control_warned) {
 						Log::WriteToConsoleAndLog(fmt::sprintf(
 						    "warning: NGS2 custom control 0x%08" PRIx32 " is ignored\n",
@@ -2461,6 +2492,16 @@ int KYTY_SYSV_ABI Ngs2VoiceControl(uintptr_t voice_handle, const Ngs2VoiceParamH
 					}
 					break;
 				}
+				if (module_id == 0x17) {
+					EXIT_NOT_IMPLEMENTED(param->size != sizeof(Ngs2VoiceParamHeader) +
+					                                       37 * 37 * sizeof(float));
+					const auto* levels = reinterpret_cast<const float*>(param + 1);
+					module.channel_matrix.resize(37 * 37);
+					std::transform(levels, levels + module.channel_matrix.size(),
+					               module.channel_matrix.begin(),
+					               [](float level) { return std::clamp(level, -32.0f, 32.0f); });
+					break;
+				}
 				if (module_id == 0x19) {
 					struct FxParam {
 						Ngs2VoiceParamHeader     header;
@@ -2660,24 +2701,22 @@ int KYTY_SYSV_ABI Ngs2VoiceGetState(uintptr_t voice_handle, Ngs2VoiceState* stat
 			submixer->voice_state.state_flags = Ngs2GetStateFlags(voice);
 			break;
 		}
-		case Ngs2RackType::CustomMastering: {
-			const auto configured_size =
-			    voice->rack->option.custom_mastering.custom_rack_option.state_size;
-			EXIT_NOT_IMPLEMENTED(configured_size < sizeof(Ngs2CustomMasteringVoiceState));
-			EXIT_NOT_IMPLEMENTED(state_size != configured_size);
-			std::memset(state, 0, state_size);
-			auto* mastering = reinterpret_cast<Ngs2CustomMasteringVoiceState*>(state);
-			mastering->voice_state.state_flags = Ngs2GetStateFlags(voice);
-			break;
-		}
+		case Ngs2RackType::CustomMastering:
 		case Ngs2RackType::CustomSampler: {
+			const bool sampler = voice->rack->type == Ngs2RackType::CustomSampler;
 			const auto configured_size =
 			    voice->rack->option.custom_sampler.custom_rack_option.state_size;
-			EXIT_NOT_IMPLEMENTED(configured_size < sizeof(Ngs2CustomSamplerVoiceState));
-			EXIT_NOT_IMPLEMENTED(state_size != configured_size);
-			std::memset(state, 0, state_size);
-			auto* sampler = reinterpret_cast<Ngs2CustomSamplerVoiceState*>(state);
-			Ngs2FillSamplerState(*voice, *sampler);
+			EXIT_NOT_IMPLEMENTED(configured_size < (sampler ? sizeof(Ngs2CustomSamplerVoiceState)
+			                                                : sizeof(Ngs2CustomMasteringVoiceState)));
+			if (state_size < configured_size) {
+				return static_cast<int32_t>(0x804a8011u);
+			}
+			std::memset(state, 0, configured_size);
+			if (sampler) {
+				Ngs2FillSamplerState(*voice, *reinterpret_cast<Ngs2CustomSamplerVoiceState*>(state));
+			} else {
+				state->state_flags = Ngs2GetStateFlags(voice);
+			}
 			break;
 		}
 		case Ngs2RackType::Sampler: {
@@ -35,7 +35,7 @@ public:
 	bool     GetParamInt(const char* name, int32_t* value);
 	bool     GetParamString(const char* name, std::string* value);
 	bool     GetParamString(const char* name, char* value, size_t value_size);
-	uint64_t GetFlexibleMemorySize();
+	MemoryParams GetMemoryParams();

 	void DbgPrint();

@@ -61,7 +61,7 @@ private:
 	Common::Mutex      m_mutex;
 	bool               m_opened = false;
 	std::vector<Param> m_params;
-	uint64_t           m_flexible_memory_size = 0;
+	MemoryParams       m_memory_params;
 };

 class PlayGo {
@@ -172,11 +172,21 @@ void Psf::OpenJson(const std::filesystem::path& file_name) {
 	add_int_param("userDefinedParam3", "USER_DEFINED_PARAM_3");
 	add_int_param("userDefinedParam4", "USER_DEFINED_PARAM_4");

+	auto read_size = [](const nlohmann::ordered_json& object, const char* key, uint64_t* size,
+	                    uint64_t unit = 1) {
+		if (auto it = object.find(key); it != object.end()) {
+			*size = it->get<uint64_t>() * unit;
+		}
+	};
 	if (auto kernel = param_json.find("kernel");
 	    kernel != param_json.end() && kernel->is_object()) {
-		if (auto it = kernel->find("flexibleMemorySize"); it != kernel->end()) {
-			m_flexible_memory_size = it->get<uint64_t>();
-		}
+		read_size(*kernel, "flexibleMemorySize", &m_memory_params.flexible_size);
+		read_size(*kernel, "cpuPageTableSize", &m_memory_params.cpu_table_memory);
+		read_size(*kernel, "gpuPageTableSize", &m_memory_params.gpu_table_memory);
+	}
+	if (auto amm = param_json.find("amm"); amm != param_json.end() && amm->is_object()) {
+		read_size(*amm, "pagetableMemorySizeInMib", &m_memory_params.amm_table_memory,
+		          1024ull * 1024ull);
 	}

 	auto localized = param_json.find("localizedParameters");
@@ -227,8 +237,8 @@ void Psf::Close() {
 	Common::LockGuard lock(m_mutex);

 	m_params.clear();
-	m_flexible_memory_size = 0;
-	m_opened               = false;
+	m_memory_params = {};
+	m_opened        = false;
 }

 bool Psf::IsValid() {
@@ -339,10 +349,10 @@ bool Psf::GetParamString(const char* name, char* value, size_t value_size) {
 	return false;
 }

-uint64_t Psf::GetFlexibleMemorySize() {
+MemoryParams Psf::GetMemoryParams() {
 	Common::LockGuard lock(m_mutex);

-	return m_flexible_memory_size;
+	return m_memory_params;
 }

 PlayGo::~PlayGo() {
@@ -646,10 +656,10 @@ bool SystemContentParamSfoGetString(const char* name, char* value, size_t value_
 	return sc->psf.GetParamString(name, value, value_size);
 }

-uint64_t SystemContentGetFlexibleMemorySize() {
+MemoryParams SystemContentGetMemoryParams() {
 	auto* sc = Common::Singleton<SystemContent>::Instance();

-	return sc->psf.GetFlexibleMemorySize();
+	return sc->psf.GetMemoryParams();
 }

 bool SystemContentGetIconPath(std::filesystem::path* path) {
@@ -14,7 +14,16 @@ void     SystemContentLoadParamSfo(const std::filesystem::path& file_name);
 bool     SystemContentParamSfoGetInt(const char* name, int32_t* value);
 bool     SystemContentParamSfoGetString(const char* name, std::string* value);
 bool     SystemContentParamSfoGetString(const char* name, char* value, size_t value_size);
-uint64_t SystemContentGetFlexibleMemorySize();
+// Memory sizes a game's param.json sets, in bytes; 0 where it sets none.
+struct MemoryParams {
+	uint64_t flexible_size    = 0;
+	uint64_t cpu_table_memory = 0;
+	uint64_t gpu_table_memory = 0;
+	uint64_t amm_table_memory = 0;
+};
+
+MemoryParams SystemContentGetMemoryParams();
+
 bool     SystemContentGetIconPath(std::filesystem::path* path);
 bool     SystemContentGetChunksNum(uint32_t* num);

@@ -70,6 +70,8 @@ static void PrintUsage() {
 	::printf("  --hdr                                Enable HDR10 output when supported.\n");
 	::printf(
 	    "  --hide-cursor                        Hide the cursor after 2 s idle. Default: off.\n");
+	::printf(
+	    "  --stretch                            Stretch the image to the window. Default: off.\n");
 	::printf("  --no-splash                          Disable the startup splash screen.\n");
 	::printf("  --vr                                 Enable the virtual VR headset.\n");
 	::printf("  --amd-cpu                            Apply AMD CPU instruction patches.\n");
@@ -199,6 +201,7 @@ static constexpr Flag FLAGS[] = {
    {"--fullscreen", &ConfigOptions::fullscreen_enabled},
    {"--hdr", &ConfigOptions::hdr_enabled},
    {"--hide-cursor", &ConfigOptions::hide_cursor_enabled},
+    {"--stretch", &ConfigOptions::stretch_enabled},
    {"--no-splash", &ConfigOptions::splash_enabled, false},
    {"--vr", &ConfigOptions::vr_enabled},
    {"--amd-cpu", &ConfigOptions::amd_cpu_enabled},
@@ -0,0 +1,360 @@
+#include "libs/ajm.cpp"
+
+#include <array>
+#include <cmath>
+#include <cstdlib>
+#include <fstream>
+#include <opus.h>
+
+using namespace Libs::Audio::Ajm;
+
+namespace {
+
+void Check(bool condition, const char* message) {
+	if (!condition) {
+		std::fprintf(stderr, "AjmResampleTests: %s\n", message);
+		std::abort();
+	}
+}
+
+struct DecodeResult {
+	AjmSidebandResult result;
+	AjmSidebandStream stream;
+	AjmSidebandMFrame frames;
+};
+
+struct Fixture {
+	uint32_t                  instance = 0;
+	uint32_t                  channels;
+	std::array<uint8_t, 4096> batch_data {};
+	AjmBatchInfo              batch {};
+
+	explicit Fixture(uint32_t channels = 1, AjmSampleEncoding encoding = AjmSampleEncoding::S16,
+	                 uint64_t flags = 0x1000)
+	    : channels(channels) {
+		Check(AjmInstanceCreate(1, 24, flags | channels | (uint64_t(encoding) << 7), &instance) ==
+		          OK,
+		      "create Opus instance");
+		Initialize();
+	}
+	~Fixture() { AjmInstanceDestroy(1, instance); }
+
+	AjmBatchInfo* Batch() {
+		batch = {batch_data.data(), 0, batch_data.size(), nullptr, nullptr};
+		return &batch;
+	}
+
+	void Initialize() {
+		AjmDecOpusInitializeParameters params {channels, 48000, 0};
+		AjmSidebandResult              result {};
+		AjmBatchJobInitialize(Batch(), instance, &params, sizeof(params), &result);
+		Check(result.result == OK, "initialize Opus instance");
+	}
+
+	void Parameters(float ratio, uint32_t flags, float change = 0, bool extended = false) {
+		AjmSidebandResult result {};
+		if (extended) {
+			AjmBatchJobSetResampleParametersEx(Batch(), instance, ratio, change, flags, &result);
+		} else {
+			AjmBatchJobSetResampleParameters(Batch(), instance, ratio, flags, &result);
+		}
+		Check(result.result == OK, "set resampling parameters");
+	}
+
+	AjmSidebandResampleInfo Info() {
+		struct {
+			AjmSidebandResult       result;
+			AjmSidebandResampleInfo info;
+		} result {};
+		AjmBatchJobGetResampleInfo(Batch(), instance, &result);
+		Check(result.result.result == OK, "get resampling info");
+		return result.info;
+	}
+
+	void Gapless(uint32_t samples, uint16_t skip) {
+		AjmSidebandGaplessDecode params {samples, skip, 0};
+		AjmSidebandResult        result {};
+		AjmBatchJobSetGaplessDecode(Batch(), instance, &params, 1, &result);
+		Check(result.result == OK, "set gapless parameters");
+	}
+
+	DecodeResult Decode(const void* input, size_t size, void* output, size_t capacity) {
+		DecodeResult result {};
+		AjmBatchJobDecode(Batch(), instance, input, size, output, capacity, &result);
+		Check((result.result.result & ~(AJM_RESULT_NOT_ENOUGH_ROOM | AJM_RESULT_PARTIAL_INPUT)) ==
+		          0,
+		      "decode returned unexpected error");
+		return result;
+	}
+};
+
+std::vector<uint8_t> Encode(uint32_t channels, int packets) {
+	int   error   = 0;
+	auto* encoder = opus_encoder_create(48000, channels, OPUS_APPLICATION_AUDIO, &error);
+	Check(encoder != nullptr && error == OPUS_OK, "create test encoder");
+	std::vector<uint8_t>      stream;
+	std::vector<int16_t>      pcm(960 * channels);
+	std::array<uint8_t, 4096> packet;
+	for (int p = 0; p < packets; ++p) {
+		for (size_t i = 0; i < pcm.size(); ++i) {
+			pcm[i] = static_cast<int16_t>(
+			    12000 * std::sin((i / channels + p * 960) * 0.041 + (i % channels)));
+		}
+		const int size = opus_encode(encoder, pcm.data(), 960, packet.data(), packet.size());
+		Check(size > 0, "encode test packet");
+		stream.push_back(static_cast<uint8_t>(size));
+		stream.push_back(static_cast<uint8_t>(size >> 8));
+		stream.insert(stream.end(), packet.begin(), packet.begin() + size);
+	}
+	opus_encoder_destroy(encoder);
+	return stream;
+}
+
+void SmallOutputs(const std::vector<uint8_t>& packet, uint16_t skip = 0) {
+	Fixture f;
+	f.Gapless(960 - skip, skip);
+	f.Parameters(1, 3);
+	std::array<uint8_t, 1030> output;
+	output.fill(0x5a);
+	auto first = f.Decode(packet.data(), packet.size(), output.data(), 1024);
+	Check(first.stream.size_consumed == static_cast<int>(packet.size()),
+	      "small output must consume packet");
+	Check(first.stream.size_produced == 1024, "small output must produce 1024 bytes");
+	Check(first.frames.num_frames == 1 && first.stream.total_decoded_samples == 960 - skip,
+	      "decoded frame counted once before resampling");
+	Check(f.Info().num_samples == 448 - skip, "report retained source samples");
+	Check(std::all_of(output.begin() + 1024, output.end(), [](auto v) { return v == 0x5a; }),
+	      "do not write past output capacity");
+	const auto last = f.Decode(nullptr, 0, output.data(), 1024);
+	Check(last.result.result == OK && last.stream.size_consumed == 0 &&
+	          last.stream.size_produced == 896 - skip * 2 && last.frames.num_frames == 0 &&
+	          last.stream.total_decoded_samples == first.stream.total_decoded_samples,
+	      "END drains retained PCM without decoding or reporting partial input");
+	Check(f.Decode(nullptr, 0, output.data(), 1024).stream.size_produced == 0,
+	      "END drains only once");
+}
+
+void UnityMatchesWholeDecode(uint32_t channels, AjmSampleEncoding encoding) {
+	const auto stream       = Encode(channels, 6);
+	const auto sample_bytes = channels * AjmBytesPerSample(encoding);
+	Fixture    reference(channels, encoding, 0);
+	reference.Gapless(5328, 312);
+	std::vector<uint8_t> expected(5760 * sample_bytes);
+	const auto           full =
+	    reference.Decode(stream.data(), stream.size(), expected.data(), expected.size());
+	expected.resize(full.stream.size_produced);
+	Check(expected.size() == 5328 * sample_bytes, "reference gapless length");
+
+	Fixture f(channels, encoding);
+	f.Gapless(5328, 312);
+	f.Parameters(1, 3);
+	std::vector<uint8_t>      actual;
+	std::array<uint8_t, 1025> output;
+	size_t                    offset = 0;
+	uint32_t                  frames = 0;
+	for (int i = 0; i < 100; ++i) {
+		const auto decoded = f.Decode(offset < stream.size() ? stream.data() + offset : nullptr,
+		                              stream.size() - offset, output.data(), output.size());
+		offset += decoded.stream.size_consumed;
+		frames += decoded.frames.num_frames;
+		actual.insert(actual.end(), output.begin(), output.begin() + decoded.stream.size_produced);
+		if (decoded.stream.size_produced == 0) {
+			break;
+		}
+	}
+	Check(actual == expected && frames == 6 && offset == stream.size(),
+	      "chunked PCM must exactly match whole decode, without loss or duplication");
+}
+
+void ChunkedRates(const std::vector<uint8_t>& stream, uint32_t samples, uint16_t skip) {
+	for (float ratio: {0.125f, 0.749153f, 1.334840f}) {
+		Fixture reference;
+		reference.Gapless(samples, skip);
+		reference.Parameters(ratio, 3);
+		std::vector<uint8_t> expected(
+		    (static_cast<size_t>(std::ceil(samples / double(ratio))) + 1) * 2);
+		const auto full =
+		    reference.Decode(stream.data(), stream.size(), expected.data(), expected.size());
+		expected.resize(full.stream.size_produced);
+		Check(expected.size() == static_cast<size_t>(std::ceil(samples / double(ratio))) * 2,
+		      "fractional resampling rounds up final output sample count");
+
+		Fixture f;
+		f.Gapless(samples, skip);
+		f.Parameters(ratio, 1);
+		std::vector<uint8_t>      actual;
+		std::array<uint8_t, 1024> output;
+		size_t                    offset = 0;
+		for (int i = 0; i < 1000; ++i) {
+			if (offset == stream.size()) {
+				// The caller can signal END only after submitting the last packet.
+				f.Parameters(0, 2 | 8);
+			}
+			const auto decoded = f.Decode(offset < stream.size() ? stream.data() + offset : nullptr,
+			                              stream.size() - offset, output.data(), output.size());
+			offset += decoded.stream.size_consumed;
+			actual.insert(actual.end(), output.begin(),
+			              output.begin() + decoded.stream.size_produced);
+			if (decoded.stream.size_produced == 0) {
+				break;
+			}
+		}
+		Check(actual == expected && offset == stream.size(),
+		      "non-unity PCM is independent of output chunk boundaries and deferred END");
+	}
+}
+
+void LoopBoundaries() {
+	const auto loop   = Encode(1, 2);
+	auto       stream = loop;
+	stream.insert(stream.end(), loop.begin(), loop.end());
+	for (uint64_t flags: {0x1400u, 0x3400u}) {
+		Fixture f(1, AjmSampleEncoding::S16, flags);
+		f.Gapless(1500, 100);
+		f.Parameters(1, 1);
+		std::array<uint8_t, 8192> pcm;
+		const auto                first = f.Decode(stream.data(), stream.size(), pcm.data(), 1024);
+		Check(first.stream.size_produced == 1024, "loop first partial output");
+		const auto second =
+		    f.Decode(stream.data() + first.stream.size_consumed,
+			         stream.size() - first.stream.size_consumed, pcm.data(), pcm.size());
+		const bool seamless = flags == 0x3400;
+		Check(second.stream.size_produced == (seamless ? 6000 - 1024 : 3000 - 1024),
+		      "non-seamless loop stops job at stream boundary");
+		const auto consumed = first.stream.size_consumed + second.stream.size_consumed;
+		Check(consumed == static_cast<int>(seamless ? stream.size() : loop.size()),
+		      "loop boundary input consumption");
+		if (!seamless) {
+			Check(
+			    f.Decode(stream.data() + consumed, stream.size() - consumed, pcm.data(), pcm.size())
+			            .stream.size_produced == 3000,
+			    "next job resumes gapless loop");
+		}
+	}
+}
+
+void RatesAndControls(const std::vector<uint8_t>& packet) {
+	for (float ratio: {0.125f, 0.5f, 1.0f, 1.25f, 2.0f, 4.0f}) {
+		Fixture f;
+		f.Parameters(ratio, 3);
+		std::vector<int16_t> pcm(8000);
+		const auto result = f.Decode(packet.data(), packet.size(), pcm.data(), pcm.size() * 2);
+		Check(result.stream.size_produced == static_cast<int>(std::ceil(960.0 / ratio)) * 2,
+		      "resampling ratio controls output duration");
+		Check(result.stream.total_decoded_samples == 960,
+		      "ratio does not change decoded sample count");
+	}
+	Fixture f;
+	f.Parameters(1, 1);
+	std::array<int16_t, 512> pcm;
+	f.Decode(packet.data(), packet.size(), pcm.data(), sizeof(pcm));
+	f.Initialize();
+	Check(f.Info().num_samples == 448, "codec initialize must preserve resampler PCM");
+	f.Parameters(0, 0);
+	const auto frozen = f.Decode(nullptr, 0, pcm.data(), sizeof(pcm));
+	Check(
+	    frozen.stream.size_produced == sizeof(pcm) && f.Info().num_samples == 448 &&
+	        std::all_of(pcm.begin(), pcm.end(), [&](auto sample) { return sample == pcm.front(); }),
+	    "zero ratio repeats one sample without advancing input");
+	f.Parameters(1, 0, 0.25f, true);
+	f.Parameters(std::numeric_limits<float>::quiet_NaN(), 8);
+	f.Decode(nullptr, 0, pcm.data(), 4 * sizeof(int16_t));
+	Check(f.Info().ratio == 2.0f, "ignore preserves ratio ramp");
+	f.Parameters(1, 1);
+	Check(f.Info().num_samples == 0, "resample RESET discards retained PCM");
+	Check(f.Decode(nullptr, 0, pcm.data(), sizeof(pcm)).stream.size_produced == 0,
+	      "RESET leaves no stale audio");
+}
+
+void UnbufferedAndSplit(const std::vector<uint8_t>& packet) {
+	Fixture                   raw(1, AjmSampleEncoding::S16, 0);
+	std::array<uint8_t, 1024> pcm;
+	const auto too_small = raw.Decode(packet.data(), packet.size(), pcm.data(), pcm.size());
+	Check(too_small.result.result == AJM_RESULT_NOT_ENOUGH_ROOM &&
+	          too_small.stream.size_consumed == 0 && too_small.stream.size_produced == 0,
+	      "ordinary decoding still requires room for a whole packet");
+	Fixture f;
+	f.Parameters(1, 3);
+	AjmBuffer    input[] {{const_cast<uint8_t*>(packet.data()), 1},
+	                      {const_cast<uint8_t*>(packet.data()) + 1, packet.size() - 1}};
+	AjmBuffer    output[] {{pcm.data(), 257}, {pcm.data() + 257, 767}};
+	DecodeResult result {};
+	AjmBatchJobDecodeSplit(f.Batch(), f.instance, input, 2, output, 2, &result);
+	Check(result.stream.size_produced == 1024 &&
+	          result.stream.size_consumed == static_cast<int>(packet.size()),
+	      "split buffers retain PCM across jobs");
+}
+
+void FormatChangeRequiresReset() {
+	Fixture f(2);
+	f.Parameters(1, 1);
+	const auto                stereo = Encode(2, 1);
+	std::array<uint8_t, 4096> pcm;
+	f.Decode(stereo.data(), stereo.size(), pcm.data(), 1024);
+	f.channels = 1;
+	f.Initialize();
+	DecodeResult result {};
+	AjmBatchJobDecode(f.Batch(), f.instance, nullptr, 0, pcm.data(), pcm.size(), &result);
+	Check(result.result.result == AJM_RESULT_INVALID_PARAMETER && result.stream.size_produced == 0,
+	      "format change cannot reinterpret retained PCM with a different channel stride");
+	f.Parameters(1, 3);
+	const auto mono = Encode(1, 1);
+	Check(f.Decode(mono.data(), mono.size(), pcm.data(), pcm.size()).stream.size_produced == 1920,
+	      "resampler RESET permits a new stream format");
+}
+
+void EndPersistsUntilReset(const std::vector<uint8_t>& packet) {
+	Fixture f;
+	f.Parameters(0.5f, 1);
+	std::array<int16_t, 64> pcm;
+	auto   decoded  = f.Decode(packet.data(), packet.size(), pcm.data(), sizeof(pcm));
+	size_t produced = decoded.stream.size_produced;
+	f.Parameters(0.5f, 2);
+	for (int i = 0; i < 40; ++i) {
+		f.Parameters(0.5f, 0);
+		decoded = f.Decode(nullptr, 0, pcm.data(), sizeof(pcm));
+		produced += decoded.stream.size_produced;
+		Check((decoded.result.result & AJM_RESULT_PARTIAL_INPUT) == 0,
+		      "END remains active across subsequent ratio controls");
+		if (f.Info().num_samples < 0) {
+			break;
+		}
+	}
+	Check(produced == 3840 && f.Info().num_samples < 0,
+	      "fully drained END reports a negative retained sample count");
+	f.Parameters(1, 1);
+	Check(f.Info().num_samples == 0, "RESET clears end-of-stream state");
+}
+
+} // namespace
+
+int main(int argc, char** argv) {
+	const auto packet = Encode(1, 1);
+	SmallOutputs(packet);
+	SmallOutputs(packet, 312);
+	for (uint32_t channels: {1u, 2u}) {
+		for (auto encoding:
+		     {AjmSampleEncoding::S16, AjmSampleEncoding::S32, AjmSampleEncoding::Float}) {
+			UnityMatchesWholeDecode(channels, encoding);
+		}
+	}
+	RatesAndControls(packet);
+	UnbufferedAndSplit(packet);
+	ChunkedRates(Encode(1, 6), 5328, 312);
+	LoopBoundaries();
+	FormatChangeRequiresReset();
+	EndPersistsUntilReset(packet);
+	if (argc >= 2) {
+		std::ifstream        file(argv[1], std::ios::binary);
+		std::vector<uint8_t> real_packet((std::istreambuf_iterator<char>(file)), {});
+		Check(real_packet.size() == 209, "expected real 209-byte packet fixture");
+		SmallOutputs(real_packet, 312);
+	}
+	if (argc >= 3) {
+		std::ifstream        file(argv[2], std::ios::binary);
+		std::vector<uint8_t> real_stream((std::istreambuf_iterator<char>(file)), {});
+		Check(real_stream.size() == 4048, "expected complete real stream fixture");
+		ChunkedRates(real_stream, 25920, 312);
+	}
+	std::puts("AjmResampleTests: passed");
+}
@@ -35,6 +35,8 @@ int                               g_open_waiters = 0;
 bool                              g_block_opens  = false;
 bool                              g_block_outputs = false;
 std::vector<int>                  g_blocked_output;
+std::atomic_uint64_t              g_process_time {0};
+std::atomic_bool                  g_process_time_frozen {false};

 void Check(bool value, const char* text) {
 	if (!value) {
@@ -109,8 +111,10 @@ PortParam MakeParam(uint32_t data_format = 0x200) {
 	return param;
 }

-AudioOut2::AudioOut2ContextHandle CreateContext(uint32_t queue_depth = 4, uint32_t max_object_ports = 0) {
+AudioOut2::AudioOut2ContextHandle
+CreateContext(uint32_t queue_depth = 4, uint32_t max_object_ports = 0, uint32_t max_ports = 8) {
 	ContextParam param {};
+	param.max_ports                           = max_ports;
 	param.max_object_ports                    = max_object_ports;
 	param.queue_depth                         = queue_depth;
 	param.num_grains                          = 512;
@@ -279,13 +283,139 @@ void TestSlotReuse() {
 	AudioOut2::AudioOut2ContextDestroy(context);
 }

+void TestObjectPortCapacityPerContext() {
+	constexpr int port_full    = static_cast<int32_t>(0x80268012u);
+	constexpr int invalid_port = static_cast<int32_t>(0x80268009u);
+	std::array<AudioOut2::AudioOut2ContextHandle, 3> contexts {};
+	std::array<AudioOut2::AudioOut2PortHandle, 3>    last_ports {};
+	auto                                             object = MakeParam(0x100);
+	object.port_type                                        = 0x100;
+	for (size_t i = 0; i < contexts.size(); i++) {
+		contexts[i] = CreateContext(4, 512, 32);
+		CreatePort(contexts[i], MakeParam());
+		for (size_t j = 0; j < 512; j++) {
+			AudioOut2::AudioOut2PortHandle port = 0;
+			Check(AudioOut2::AudioOut2PortCreate(contexts[i], AsParam(&object), &port) == OK,
+			      "context could not create all 512 object ports");
+			last_ports[i] = port;
+		}
+		AudioOut2::AudioOut2PortHandle overflow = UINT64_MAX;
+		Check(AudioOut2::AudioOut2PortCreate(contexts[i], AsParam(&object), &overflow) ==
+		              port_full &&
+		          overflow == UINT64_MAX,
+		      "object capacity did not reject the 513th port without changing its output");
+		CreatePort(contexts[i], MakeParam());
+	}
+
+	for (size_t i = 0; i < contexts.size(); i++) {
+		const auto stale = last_ports[i];
+		AudioOut2::AudioOut2PortDestroy(stale);
+		last_ports[i] = CreatePort(contexts[i], object);
+		Check(last_ports[i] != stale &&
+		          AudioOut2::AudioOut2PortSetAttributes(stale, nullptr, 0) == invalid_port,
+		      "reused object slot accepted its old handle");
+		std::array<float, 512> pcm;
+		pcm.fill(0.25f * static_cast<float>(i + 1));
+		SetPcm(last_ports[i], pcm.data());
+		const auto  output   = PushPcm(contexts[i], 512 * 2 * sizeof(float));
+		const float expected = pcm[0] * 0.70710678f;
+		Check(output.size() == 1, "last object port did not reach its context's mixer");
+		CheckSamples(output[0], {expected, expected}, "object audio crossed context boundaries");
+	}
+
+	AudioOut2::AudioOut2ContextDestroy(contexts[0]);
+	Check(AudioOut2::AudioOut2PortSetAttributes(last_ports[0], nullptr, 0) == invalid_port,
+	      "destroyed context retained its object handles");
+	const auto replacement = CreateContext(4, 512);
+	for (size_t i = 0; i < 512; i++) {
+		CreatePort(replacement, object);
+	}
+	Check(AudioOut2::AudioOut2PortSetAttributes(last_ports[0], nullptr, 0) == invalid_port,
+	      "reused context accepted an old object handle");
+	for (size_t i = 1; i < contexts.size(); i++) {
+		Check(AudioOut2::AudioOut2PortSetAttributes(last_ports[i], nullptr, 0) == OK,
+		      "destroying another context invalidated a live object port");
+		AudioOut2::AudioOut2ContextDestroy(contexts[i]);
+	}
+	AudioOut2::AudioOut2ContextDestroy(replacement);
+	Check(LiveDeviceCount() == 0, "object capacity test leaked audio devices");
+}
+
+void TestContextPortLimits() {
+	constexpr int port_full = static_cast<int32_t>(0x80268012u);
+	ContextParam  param {};
+	param.max_ports                           = 1;
+	param.max_object_ports                    = 2;
+	param.num_grains                          = 512;
+	AudioOut2::AudioOut2ContextHandle context = 0;
+	Check(AudioOut2::AudioOut2ContextCreate(AsParam(&param), nullptr, 0, &context) == OK,
+	      "small context create failed");
+	auto object             = MakeParam(0x100);
+	object.port_type        = 0x100;
+	const auto channel      = CreatePort(context, MakeParam());
+	const auto first_object = CreatePort(context, object);
+	CreatePort(context, object);
+	AudioOut2::AudioOut2PortHandle overflow      = 0;
+	const auto                     channel_param = MakeParam();
+	Check(AudioOut2::AudioOut2PortCreate(context, AsParam(&channel_param), &overflow) ==
+	              port_full &&
+	          AudioOut2::AudioOut2PortCreate(context, AsParam(&object), &overflow) == port_full,
+	      "declared channel/object capacities were ignored");
+	AudioOut2::AudioOut2PortDestroy(channel);
+	Check(AudioOut2::AudioOut2PortCreate(context, AsParam(&object), &overflow) == port_full,
+	      "object port borrowed a channel slot");
+	CreatePort(context, channel_param);
+	AudioOut2::AudioOut2PortDestroy(first_object);
+	Check(AudioOut2::AudioOut2PortCreate(context, AsParam(&channel_param), &overflow) == port_full,
+	      "channel port borrowed an object slot");
+	CreatePort(context, object);
+	AudioOut2::AudioOut2ContextDestroy(context);
+	param.max_ports = param.max_object_ports = 0;
+	Check(AudioOut2::AudioOut2ContextCreate(AsParam(&param), nullptr, 0, &context) == OK,
+	      "empty context create failed");
+	Check(AudioOut2::AudioOut2PortCreate(context, AsParam(&channel_param), &overflow) ==
+	              port_full &&
+	          AudioOut2::AudioOut2PortCreate(context, AsParam(&object), &overflow) == port_full,
+	      "zero port capacity was ignored");
+	AudioOut2::AudioOut2ContextDestroy(context);
+}
+
+void TestContextCapacityParameters() {
+	constexpr int invalid_param = static_cast<int32_t>(0x80268001u);
+	ContextParam  param;
+	std::memset(&param, 0xff, sizeof(param));
+	Check(AudioOut2::AudioOut2ContextResetParam(
+	          reinterpret_cast<AudioOut2::AudioOut2ContextParam*>(&param)) == OK &&
+	          param.max_ports == 8 && param.max_object_ports == 0 &&
+	          param.guarantee_object_ports == 0 && param.num_grains == 256 &&
+	          param.queue_depth == 1 && param.flags == 0 &&
+	          std::all_of(std::begin(param.reserved), std::end(param.reserved),
+	                      [](uint32_t word) { return word == 0; }),
+	      "context reset did not produce valid default parameters");
+	const auto defaults = param;
+	for (size_t i = 0; i < 3; i++) {
+		param = defaults;
+		if (i == 0) param.max_ports = 33;
+		if (i == 1) param.max_object_ports = 513;
+		if (i == 2) param.guarantee_object_ports = 1;
+		size_t                            memory  = SIZE_MAX;
+		AudioOut2::AudioOut2ContextHandle context = UINT64_MAX;
+		Check(AudioOut2::AudioOut2ContextQueryMemory(AsParam(&param), &memory) == invalid_param &&
+		          memory == SIZE_MAX &&
+		          AudioOut2::AudioOut2ContextCreate(AsParam(&param), nullptr, 0, &context) ==
+		              invalid_param &&
+		          context == UINT64_MAX && LiveDeviceCount() == 0,
+		      "invalid capacities allocated resources or modified outputs");
+	}
+}
+
 void TestFullTableRecovers() {
 	const auto                                  context = CreateContext();
 	const auto                                  param   = MakeParam();
 	std::vector<AudioOut2::AudioOut2PortHandle> ports;
-	ports.reserve(256);
+	ports.reserve(8);

-	for (int i = 0; i < 256; i++) {
+	for (int i = 0; i < 8; i++) {
 		AudioOut2::AudioOut2PortHandle port = 0;
 		Check(AudioOut2::AudioOut2PortCreate(context, AsParam(&param), &port) == OK,
 		      "port table filled early");
@@ -322,6 +452,11 @@ void TestConcurrentCreates() {
 		});
 	}
 	WaitForDeviceOpens(thread_count);
+	AudioOut2::AudioOut2PortHandle overflow = 0;
+	Check(AudioOut2::AudioOut2PortCreate(context, AsParam(&param), &overflow) ==
+	              static_cast<int32_t>(0x80268012u) &&
+	          overflow == 0,
+	      "pending creates did not reserve their channel capacity");
 	ReleaseDeviceOpens();
 	for (auto& thread: threads) {
 		thread.join();
@@ -349,17 +484,31 @@ void TestContextDestroyCancelsPendingCreate() {
 	    [&]() { result = AudioOut2::AudioOut2PortCreate(context, AsParam(&param), &port); });
 	WaitForDeviceOpens(1);
 	AudioOut2::AudioOut2ContextDestroy(context);
+	const auto                     replacement        = CreateContext();
+	AudioOut2::AudioOut2PortHandle replacement_port   = 0;
+	int                            replacement_result = OK;
+	std::thread                    replacement_creator([&]() {
+		replacement_result =
+		    AudioOut2::AudioOut2PortCreate(replacement, AsParam(&param), &replacement_port);
+	});
+	WaitForDeviceOpens(2);
 	ReleaseDeviceOpens();
 	creator.join();
+	replacement_creator.join();

 	Check(result != OK, "destroyed context retained a pending port create");
+	Check(replacement_result == OK &&
+	          AudioOut2::AudioOut2PortSetAttributes(replacement_port, nullptr, 0) == OK &&
+	          LiveDeviceCount() == 1,
+	      "cancelled create corrupted the replacement context or leaked a device");
+	AudioOut2::AudioOut2ContextDestroy(replacement);
 	Check(LiveDeviceCount() == 0, "cancelled port create leaked a device");
 }

 void TestSynchronousDevicePushBypassesModelledQueue() {
-	const auto context = CreateContext(1);
-	const auto param   = MakeParam();
-	AudioOut2::AudioOut2PortHandle port = 0;
+	const auto                     context = CreateContext(1);
+	const auto                     param   = MakeParam();
+	AudioOut2::AudioOut2PortHandle port    = 0;
 	Check(AudioOut2::AudioOut2PortCreate(context, AsParam(&param), &port) == OK,
 	      "device port create failed");

@@ -368,6 +517,7 @@ void TestSynchronousDevicePushBypassesModelledQueue() {
 	ResetOutputCalls();

 	Check(AudioOut2::AudioOut2ContextPush(context, 1) == OK, "first sync push failed");
+	SetPcm(port, pcm);
 	Check(AudioOut2::AudioOut2ContextPush(context, 1) == OK,
 	      "device-paced sync push was blocked by modelled queue");
 	const auto calls = OutputCalls();
@@ -447,6 +597,44 @@ void TestHandleWithoutPcmDoesNotBypassQueue() {
 	AudioOut2::AudioOut2ContextDestroy(context);
 }

+void TestPushConsumesPcm() {
+	const auto context = CreateContext();
+	const auto port    = CreatePort(context, MakeParam());
+	uint32_t   pcm[512] {};
+	SetPcm(port, pcm);
+	ResetOutputCalls();
+	Check(AudioOut2::AudioOut2ContextPush(context, 1) == OK, "push with PCM failed");
+	Check(AudioOut2::AudioOut2ContextPush(context, 1) == OK, "push without new PCM failed");
+	SetPcm(port, pcm);
+	SetPcm(port, nullptr);
+	Check(AudioOut2::AudioOut2ContextPush(context, 1) == OK, "push after clearing PCM failed");
+	Check(OutputCalls().size() == 1, "a push replayed PCM the game did not set again");
+
+	AudioOut2::AudioOut2PortDestroy(port);
+	AudioOut2::AudioOut2ContextDestroy(context);
+}
+
+void TestQueueDrainsAtExactGrainRate() {
+	// Three 512-sample grains play in exactly 32000 us.
+	const auto context = CreateContext();
+	g_process_time_frozen = true;
+	const auto start      = g_process_time.load();
+	for (int push = 0; push < 3; push++) {
+		Check(AudioOut2::AudioOut2ContextPush(context, 0) == OK, "async push failed");
+	}
+	const auto queued_after = [&](uint64_t micros) {
+		g_process_time  = start + micros;
+		uint32_t queued = 0;
+		Check(AudioOut2::AudioOut2ContextGetQueueLevel(context, &queued, nullptr) == OK,
+		      "queue-level query failed");
+		return queued;
+	};
+	Check(queued_after(31999) == 1, "the third grain drained before it finished playing");
+	Check(queued_after(32000) == 0, "three grains did not drain in their playing time");
+	g_process_time_frozen = false;
+	AudioOut2::AudioOut2ContextDestroy(context);
+}
+
 void TestPcmCopiedBeforeScratchBufferReuse() {
 	const auto context = CreateContext();
 	const auto param   = MakeParam();
@@ -486,7 +674,6 @@ void TestPortGainAndValidation() {
 	const auto context = CreateContext();
 	const auto port = CreatePort(context, MakeParam());
 	std::vector<float> pcm(512 * 2, 0.25f);
-	SetPcm(port, pcm.data());
 	auto check_output = [&](const std::array<float, 2>& gain, float expected_pcm) {
 		CheckSamples(PushPcm(context, pcm.size() * sizeof(float)).at(0), {expected_pcm},
 		             "gain modified stored PCM");
@@ -494,10 +681,11 @@ void TestPortGainAndValidation() {
 		Check(g_output_gains.size() == 1 && g_output_gains[0] == gain,
 		      "incorrect per-channel gain forwarded to backend");
 	};
+	SetPcm(port, pcm.data());
 	check_output({1.0f, 1.0f}, 0.25f);
 	for (const std::array<float, 2> gain: {std::array {2.0f, 0.5f}, {0.0f, 0.0f}, {1.0f, 1.0f}}) {
 		SetAttribute(port, 1, gain);
-		check_output(gain, 0.25f);
+		SetPcm(port, pcm.data());
 		check_output(gain, 0.25f);
 	}

@@ -507,10 +695,11 @@ void TestPortGainAndValidation() {
 	SetPcm(port, pcm.data());
 	check_output(gain, 0.75f);

-	std::fill(pcm.begin(), pcm.end(), -0.25f);
-	const Pcm replacement {pcm.data()};
+	const std::vector<float> replacement_pcm(pcm.size(), -0.25f);
+	const Pcm                replacement {replacement_pcm.data()};
 	for (const float invalid: {-1.0f, std::numeric_limits<float>::infinity(),
 	                           std::numeric_limits<float>::quiet_NaN()}) {
+		SetPcm(port, pcm.data());
 		const std::array invalid_gain {1.0f, invalid};
 		const Attribute attributes[] {{0, 0, &replacement, sizeof(replacement)},
 		                              {1, 0, &invalid_gain, sizeof(invalid_gain)}};
@@ -523,6 +712,7 @@ void TestPortGainAndValidation() {
 		Check(AudioOut2::AudioOut2PortSetAttributes(port, AsAttribute(&attribute), 1) ==
 		          static_cast<int32_t>(0x80268001u), "malformed gain was accepted");
 	}
+	SetPcm(port, pcm.data());
 	const Attribute duplicates[] {{0, 0, &replacement, sizeof(replacement)},
 	                              {1, 0, &gain, sizeof(gain)}, {1, 0, &gain, sizeof(gain)}};
 	Check(AudioOut2::AudioOut2PortSetAttributes(port, AsAttribute(duplicates), 3) ==
@@ -553,12 +743,10 @@ void TestObjectGainAndAmbisonics() {
 		std::fill(integer.begin(), integer.end(), 0);
 		SetAttribute(object, 1, 0.0f);
 		SetAttribute(object, 8, 65u);
-		for (int push = 0; push < 2; push++) {
-			const auto output = PushPcm(context, output_bytes);
-			Check(output.size() == 1, "object-only push did not output exactly once");
-			CheckSamples(output[0], {0.5f * half_power, 0.5f * half_power},
-			             "object PCM ownership, default gain, or attribute freeze failed");
-		}
+		const auto output = PushPcm(context, output_bytes);
+		Check(output.size() == 1, "object-only push did not output exactly once");
+		CheckSamples(output[0], {0.5f * half_power, 0.5f * half_power},
+		             "object PCM ownership, default gain, or attribute freeze failed");
 		std::fill(floating.begin(), floating.end(), 0.5f);
 		std::fill(integer.begin(), integer.end(), 16384);
 		for (const auto& test: cases) {
@@ -569,10 +757,8 @@ void TestObjectGainAndAmbisonics() {
 			                              {8, 0, &test.channel, sizeof(test.channel)}};
 			Check(AudioOut2::AudioOut2PortSetAttributes(object, AsAttribute(attributes), 3) == OK,
 			      "object attributes could not be set atomically with PCM first");
-			for (int push = 0; push < 2; push++) {
-				CheckSamples(PushPcm(context, output_bytes).at(0), {0.25f * test.left, 0.25f * test.right},
-				             "object gain or ambisonics decoded incorrectly or accumulated across pushes");
-			}
+			CheckSamples(PushPcm(context, output_bytes).at(0), {0.25f * test.left, 0.25f * test.right},
+			             "object gain or ambisonics decoded incorrectly or accumulated across pushes");
 		}
 		AudioOut2::AudioOut2ContextAdvance(context);
 		SetAttribute(object, 8, UINT32_MAX);
@@ -599,22 +785,21 @@ void TestObjectBusRouting() {
 	SetPcm(second, bed.data());
 	auto param = MakeParam(0x100);
 	param.port_type = 0x100;
-	std::vector<float> object_pcm(512, 0.25f);
-	SetPcm(CreatePort(context, param), object_pcm.data());
+	const auto object = CreatePort(context, param);
 	param.user_handle = other_user;
-	SetPcm(CreatePort(context, param), object_pcm.data());
-	std::fill(object_pcm.begin(), object_pcm.end(), 1.0f);
-	SetPcm(CreatePort(other_context, param), object_pcm.data());
+	const auto other_user_object = CreatePort(context, param);
+	std::vector<float> object_pcm(512, 0.25f);
+	SetPcm(object, object_pcm.data());
+	SetPcm(other_user_object, object_pcm.data());
+	SetPcm(CreatePort(other_context, param), std::vector<float>(512, 1.0f).data());
 	constexpr auto bytes = 512 * 2 * sizeof(float);
-	for (int push = 0; push < 2; push++) {
-		const auto output = PushPcm(context, bytes);
-		Check(output.size() == 3, "object bus was missing or queued more than once");
-		CheckSamples(output[0], {0.125f}, "objects modified first bed");
-		CheckSamples(output[1], {0.125f}, "objects modified second bed");
-		CheckSamples(output[2], {0.35355339f}, "MAIN objects failed to mix across users or leaked contexts");
-	}
-	SetPcm(first, nullptr);
-	SetPcm(second, nullptr);
+	const auto output = PushPcm(context, bytes);
+	Check(output.size() == 3, "object bus was missing or queued more than once");
+	CheckSamples(output[0], {0.125f}, "objects modified first bed");
+	CheckSamples(output[1], {0.125f}, "objects modified second bed");
+	CheckSamples(output[2], {0.35355339f}, "MAIN objects failed to mix across users or leaked contexts");
+	SetPcm(object, object_pcm.data());
+	SetPcm(other_user_object, object_pcm.data());
 	CheckSamples(PushPcm(context, bytes).at(0), {0.35355339f}, "object output depended on bed PCM");
 	AudioOut2::AudioOut2ContextDestroy(context);
 	Check(LiveDeviceCount() == 1, "context destroy closed another context's object device");
@@ -731,8 +916,7 @@ uint32_t AudioOutOutputs(const OutputParam* params, uint32_t num, bool blocking)
 namespace Libs::LibKernel {

 uint64_t KYTY_SYSV_ABI KernelGetProcessTime() {
-	static std::atomic_uint64_t now {0};
-	return now.fetch_add(1000);
+	return g_process_time_frozen ? g_process_time.load() : g_process_time.fetch_add(1000);
 }

 } // namespace Libs::LibKernel
@@ -741,6 +925,9 @@ int main() {
 	Check(AudioOut2::AudioOut2UserCreate(1000, &g_user_handle) == OK, "test user create failed");
 	TestUserSupportedAttributes();
 	TestSlotReuse();
+	TestObjectPortCapacityPerContext();
+	TestContextPortLimits();
+	TestContextCapacityParameters();
 	TestFullTableRecovers();
 	TestConcurrentCreates();
 	TestContextDestroyCancelsPendingCreate();
@@ -748,6 +935,8 @@ int main() {
 	TestFloat12ChannelPortOutputsPcm();
 	TestAsynchronousDevicePushKeepsQueueBounded();
 	TestHandleWithoutPcmDoesNotBypassQueue();
+	TestPushConsumesPcm();
+	TestQueueDrainsAtExactGrainRate();
 	TestPcmCopiedBeforeScratchBufferReuse();
 	TestPortGainAndValidation();
 	TestObjectGainAndAmbisonics();
@@ -689,6 +689,27 @@ void TestTriggerEffectState() {
 	check_state(0, 0, "switching controllers retained effect states");
 }

+void TestUnreadPressSurvivesStickBurst() {
+	Controller              controller;
+	std::array<PadData, 64> data {};
+	SetButton(1, PAD_BUTTON_CROSS, true);
+	for (int value = 0; value < 100; value++) {
+		SetAxis(1, Axis::LeftX, value);
+	}
+	SetButton(1, PAD_BUTTON_CROSS, false);
+	for (int value = 100; value < 200; value++) {
+		SetAxis(1, Axis::LeftX, value);
+	}
+	Check(PadRead(1, data.data(), 64) == 2 && data[0].buttons == PAD_BUTTON_CROSS &&
+	          data[0].left_stick_x == 99 && data[1].buttons == 0 && data[1].left_stick_x == 199,
+	      "a stick burst pushed an unread press out of the pad history");
+	SetAxis(1, Axis::LeftX, 50);
+	SetButton(1, PAD_BUTTON_CROSS, true);
+	Check(PadRead(1, data.data(), 64) == 2 && data[0].buttons == 0 &&
+	          data[0].left_stick_x == 50 && data[1].buttons == PAD_BUTTON_CROSS,
+	      "a stick move after a read did not add a sample");
+}
+
 } // namespace

 int main() {
@@ -702,6 +723,7 @@ int main() {
 	TestMaskedTriggersAndValidation();
 	TestIndependentOutputsAndPadSwitch();
 	TestTriggerEffectState();
+	TestUnreadPressSurvivesStickBurst();
 	TestMultiplayerRouting();
 	TestDedicatedKeyboardOutputs();
 	Config::Shutdown();
@@ -751,6 +751,26 @@ void TestCustomSideChainUserFx() {
 	      "user FX did not receive the side-chain buffer after the source channels");
 }

+struct RenderRack {
+	std::unique_ptr<std::byte[]> storage;
+	uintptr_t handle = 0;
+
+	RenderRack(Ngs2Internal& system, uint32_t id, const Ngs2RackOption& option) {
+		Ngs2ContextBufferInfo buffer {};
+		Check(Ngs2RackQueryBufferSize(id, &option, &buffer) == OK, "rack size query failed");
+		storage = std::make_unique<std::byte[]>(buffer.host_buffer_size);
+		buffer.host_buffer = storage.get();
+		Check(Ngs2RackCreate(reinterpret_cast<uintptr_t>(&system), id, &option, &buffer,
+		                     &handle) == OK, "render rack creation failed");
+	}
+
+	~RenderRack() { Check(Ngs2RackDestroy(handle, nullptr) == OK, "render rack cleanup failed"); }
+
+	Ngs2VoiceInternal& Voice() {
+		return *reinterpret_cast<Ngs2VoiceInternal*>(reinterpret_cast<Ngs2RackInternal*>(handle) + 1);
+	}
+};
+
 void TestPcm16RenderOutput() {
 	Ngs2Internal system;
 	system.option                   = Ngs2DefaultSystemOption();
@@ -760,23 +780,11 @@ void TestPcm16RenderOutput() {
 	Ngs2FillDefaultRackOption(0x3000, &mastering_option);
 	sampler_option.common.max_voices = 2;
 	sampler_option.common.max_ports  = 1;
-	alignas(Ngs2RackInternal) alignas(Ngs2VoiceInternal)
-	    std::byte sampler_storage[sizeof(Ngs2RackInternal) + 2 * sizeof(Ngs2VoiceInternal)];
-	alignas(Ngs2RackInternal) alignas(Ngs2VoiceInternal)
-	    std::byte mastering_storage[sizeof(Ngs2RackInternal) + sizeof(Ngs2VoiceInternal)];
-	Ngs2ContextBufferInfo sampler_buffer {sampler_storage, sizeof(sampler_storage)};
-	Ngs2ContextBufferInfo mastering_buffer {mastering_storage, sizeof(mastering_storage)};
-	uintptr_t sampler_handle = 0, mastering_handle = 0;
+	RenderRack sampler(system, 0x1000, sampler_option.common);
+	RenderRack master(system, 0x3000, mastering_option.common);
 	const auto system_handle = reinterpret_cast<uintptr_t>(&system);
-	Check(Ngs2RackCreate(system_handle, 0x1000, &sampler_option.common, &sampler_buffer,
-	                     &sampler_handle) == OK &&
-	          Ngs2RackCreate(system_handle, 0x3000, &mastering_option.common,
-	                         &mastering_buffer, &mastering_handle) == OK,
-	      "PCM16 render rack creation failed");
-	auto* samplers = reinterpret_cast<Ngs2VoiceInternal*>(
-	    reinterpret_cast<Ngs2RackInternal*>(sampler_handle) + 1);
-	auto& mastering = *reinterpret_cast<Ngs2VoiceInternal*>(
-	    reinterpret_cast<Ngs2RackInternal*>(mastering_handle) + 1);
+	auto* samplers = &sampler.Voice();
+	auto& mastering = master.Voice();
 	mastering.channels = 1;
 	mastering.state    = Ngs2VoicePlayState::Playing;
 	const int16_t source[2][4] = {{16384, -16384, 32767, -32768},
@@ -797,9 +805,104 @@ void TestPcm16RenderOutput() {
 	Check(Ngs2SystemRender(system_handle, &output, 1) == OK &&
 	          std::ranges::all_of(output_samples, [](int16_t sample) { return sample == 0; }),
 	      "PCM16 output retained samples from the previous render");
-	Check(Ngs2RackDestroy(mastering_handle, nullptr) == OK &&
-	          Ngs2RackDestroy(sampler_handle, nullptr) == OK,
-	      "PCM16 render rack cleanup failed");
+}
+
+void TestCustomMasteringChannelMixer() {
+	Ngs2Internal system;
+	system.option = Ngs2DefaultSystemOption();
+	system.option.num_grain_samples = 6;
+	Ngs2RackOptionUnion sampler_option {};
+	Ngs2FillDefaultRackOption(0x1000, &sampler_option);
+	sampler_option.common.max_voices = 1;
+	sampler_option.common.max_ports = 1;
+	Ngs2CustomMasteringRackOption master_option {};
+	auto& custom = master_option.custom_rack_option;
+	custom.rack_option.size = sizeof(master_option);
+	custom.rack_option.max_voices = 1;
+	custom.rack_option.max_voice_channels = 6;
+	custom.num_buffers = 1;
+	custom.num_modules = 1;
+	custom.module[0].module_id = 0x17;
+	RenderRack sampler(system, 0x1000, sampler_option.common);
+	RenderRack master(system, 0x4003, custom.rack_option);
+	auto& source = sampler.Voice();
+	auto& output = master.Voice();
+	output.channels = 6;
+	output.SetEvent(1);
+	struct {
+		Ngs2VoiceParamHeader header {5484, 0, 0x40001700};
+		float levels[37][37] {};
+	} matrix;
+	constexpr float center = 0.70710677f;
+	matrix.levels[0][0] = matrix.levels[1][1] = 1.0f;
+	matrix.levels[0][2] = matrix.levels[1][2] = center;
+	matrix.levels[0][4] = matrix.levels[1][5] = center;
+	int16_t impulses[6][6] {};
+	for (uint32_t c = 0; c < 6; ++c) impulses[c][c] = 8192;
+	for (unsigned pass = 0; pass < 2; ++pass) {
+		if (pass == 1) {
+			std::memset(matrix.levels, 0, sizeof(matrix.levels));
+			matrix.levels[0][1] = 64.0f;
+			matrix.levels[1][0] = -64.0f;
+		}
+		Check(Ngs2VoiceControl(reinterpret_cast<uintptr_t>(&output), &matrix.header) == OK,
+		      "custom channel mixer control failed");
+		source.SetupSampler({0x12, 6, 48000});
+		source.blocks.push_back({reinterpret_cast<const uint8_t*>(impulses),
+		                         {0, sizeof(impulses), 0, 0, 6}});
+		source.ports[0].dest = &output;
+		source.SetEvent(1);
+		std::array<float, 14> pcm;
+		pcm.fill(99.0f);
+		Ngs2RenderBufferInfo buffer {pcm.data() + 1, 12 * sizeof(float), 0x18, 2};
+		Check(Ngs2SystemRender(reinterpret_cast<uintptr_t>(&system), &buffer, 1) == OK,
+		      "custom mastering rejected stereo output");
+		const std::array<float, 12> expected = pass == 0
+		    ? std::array<float, 12> {0.25f, 0, 0, 0.25f, center / 4, center / 4,
+		                             0, 0, center / 4, 0, 0, center / 4}
+		    : std::array<float, 12> {0, -8, 8, 0};
+		Check(pcm.front() == 99 && pcm.back() == 99 &&
+		          std::equal(expected.begin(), expected.end(), pcm.begin() + 1),
+		      "custom mixer lost channel routing, in-place input, clamping or output bounds");
+	}
+}
+
+void TestCustomVoiceStateCapacity() {
+	for (auto type: {Ngs2RackType::CustomSampler, Ngs2RackType::CustomMastering}) {
+		Fixture f(48000, 1, Ngs2RackType::CustomSampler);
+		f.rack.type = type;
+		f.rack.option.custom_sampler.custom_rack_option.state_size = 212;
+		f.voice.state_flags = 3;
+		f.voice.decoded_samples = 37;
+		f.voice.decoded_bytes = 128;
+		alignas(Ngs2CustomSamplerVoiceState) std::array<uint8_t, 384> storage;
+		auto* state = reinterpret_cast<Ngs2VoiceState*>(storage.data() + 8);
+		const auto handle = reinterpret_cast<uintptr_t>(&f.voice);
+		for (size_t capacity: {212u, 368u}) {
+			storage.fill(0xa5);
+			Check(Ngs2VoiceGetState(handle, state, capacity) == OK && state->state_flags == 3 &&
+			          state->error_code == 0 &&
+			          std::all_of(storage.begin(), storage.begin() + 8,
+			                      [](uint8_t v) { return v == 0xa5; }) &&
+			          std::all_of(storage.begin() + 220, storage.end(),
+			                      [](uint8_t v) { return v == 0xa5; }),
+			      "oversized custom state buffer was rejected or overwritten past configured state");
+			if (type == Ngs2RackType::CustomSampler) {
+				const auto* sampler = reinterpret_cast<const Ngs2CustomSamplerVoiceState*>(state);
+				Check(sampler->num_decoded_samples == 37 && sampler->decoded_data_size == 128,
+				      "custom sampler state omitted playback progress");
+			}
+			const size_t base_size = type == Ngs2RackType::CustomSampler
+			    ? sizeof(Ngs2CustomSamplerVoiceState) : sizeof(Ngs2CustomMasteringVoiceState);
+			Check(std::all_of(storage.begin() + 8 + base_size, storage.begin() + 220,
+			                  [](uint8_t v) { return v == 0; }),
+			      "configured custom state extension was not initialized");
+		}
+		storage.fill(0xa5);
+		Check(Ngs2VoiceGetState(handle, state, 211) == static_cast<int32_t>(0x804a8011u) &&
+		          std::ranges::all_of(storage, [](uint8_t v) { return v == 0xa5; }),
+		      "undersized custom state buffer was accepted or modified");
+	}
 }

 void TestPitchLoopAndSkip() {
@@ -1700,6 +1803,8 @@ int main() {
 	TestAtrac9Vibration();
 	TestAtrac9Discrete();
 	TestPcm16RenderOutput();
+	TestCustomMasteringChannelMixer();
+	TestCustomVoiceStateCapacity();
 	TestSubmixerSetupAndRouting();
 	TestCustomSideChainUserFx();
 	TestWaveformBlockGeometry();
@@ -660,6 +660,47 @@ void TestClassicSaveParams() {
 	CHECK(params[0].mtime == nested_mtime);
 }

+void TestSearchSortKeys() {
+	Reset("SORTED");
+	const std::array names {DirName("alpha"), DirName("bravo"), DirName("charlie")};
+	constexpr std::array<uint64_t, 3> blocks {96, 48, 48};
+	constexpr std::array<uint32_t, 3> user_params {2, 3, 1};
+	constexpr std::array<int, 3>      hours_ahead {1, 3, 2};
+	for (size_t i = 0; i < names.size(); i++) {
+		struct SaveDataMount3 mount {};
+		mount.user_id    = 1;
+		mount.dir_name   = &names[i];
+		mount.mount_mode = 4;
+		mount.blocks     = blocks[i];
+		SaveDataMountResult result {};
+		CHECK(SaveDataMount3(&mount, &result) == OK);
+		CHECK(SaveDataSetParam(&result.mount_point, 4, &user_params[i], sizeof(user_params[i])) ==
+		      OK);
+		CHECK(SaveDataUmount2(0, &result.mount_point) == OK);
+		fs::last_write_time(fs::path("_SaveData/SORTED") / names[i].data / "sce_sys/param.bin",
+		                    fs::file_time_type::clock::now() + std::chrono::hours(hours_ahead[i]));
+	}
+	const auto search = [](SaveDataSortKey key, SaveDataSortOrder order) {
+		std::array<SceSaveDataDirName, 3> found {};
+		SaveDataDirNameSearchCond         cond {};
+		cond.user_id = 1;
+		cond.key     = key;
+		cond.order   = order;
+		SaveDataDirNameSearchResult result {};
+		result.dir_names     = found.data();
+		result.dir_names_num = found.size();
+		CHECK(SaveDataDirNameSearch(&cond, &result) == OK && result.set_num == 3);
+		return std::string(found[0].data) + " " + found[1].data + " " + found[2].data;
+	};
+	CHECK(search(SaveDataSortKey::DirName, SaveDataSortOrder::Descent) == "charlie bravo alpha");
+	CHECK(search(SaveDataSortKey::UserParam, SaveDataSortOrder::Ascent) == "charlie alpha bravo");
+	CHECK(search(SaveDataSortKey::Blocks, SaveDataSortOrder::Ascent) == "bravo charlie alpha");
+	CHECK(search(SaveDataSortKey::FreeBlocks, SaveDataSortOrder::Descent) == "alpha charlie bravo");
+	CHECK(search(SaveDataSortKey::Mtime, SaveDataSortOrder::Descent) == "bravo charlie alpha");
+	CHECK(fs::remove("_SaveData/SORTED/alpha/sce_sys/blocks.bin"));
+	CHECK(search(SaveDataSortKey::Blocks, SaveDataSortOrder::Ascent) == "alpha bravo charlie");
+}
+
 void RunChild(const fs::path& executable, const char* mode) {
 #ifdef _WIN32
 	CHECK(_spawnl(_P_WAIT, executable.string().c_str(), executable.string().c_str(), mode,
@@ -706,6 +747,7 @@ int main(int argc, char** argv) {
 	TestGuestSaveIsolation();
 	TestSaveAllocations();
 	TestClassicSaveParams();
+	TestSearchSortKeys();
 	CHECK(SaveDataTerminate() == OK);
 	fs::current_path(previous);
 	fs::remove_all(temp);
@@ -1938,7 +1938,7 @@ void TestVertexBufferGrouping() {
  Check(input.fetch_attrib_reg == 2, "a draw reread the vertex metadata of a mapped shader");
  ShaderMapUserData(regs.es_regs.data_addr, mapped);
  (void)PrepareProgram(regs, HW::Context{}, HW::UserConfig{}, input);
-  Check(!input.fetch_embedded, "remapping a shader kept its previous vertex metadata");
+  Check(!input.has_dynamic_vertex_fetch, "remapping a shader kept its previous vertex metadata");
 }

 void TestNggVertexEntryState() {
@@ -2397,6 +2397,50 @@ void TestNewShaderRecompilerVop3LaneReadDestinationEncoding() {
  CheckSpirvBinaryValidates(result.spirv);
 }

+void TestVop3FloatConversionDecode() {
+  namespace Decoder = ShaderRecompiler::Decoder;
+  const std::pair<uint32_t, Decoder::Opcode> conversions[] = {
+      {0x18b, Decoder::Opcode::V_CVT_F32_F16},
+      {0x191, Decoder::Opcode::V_CVT_F32_UBYTE0},
+      {0x192, Decoder::Opcode::V_CVT_F32_UBYTE1},
+      {0x193, Decoder::Opcode::V_CVT_F32_UBYTE2},
+      {0x194, Decoder::Opcode::V_CVT_F32_UBYTE3},
+  };
+  for (const auto &[encoding, opcode] : conversions) {
+    for (const uint32_t source : {261u, 12u, 128u, 255u}) {
+      const uint32_t code[] = {
+          EncodeVop3Word0(encoding, 20, 0, 1, true),
+          EncodeVop3Word1(source, 0, 0) | (2u << 27u) | (1u << 29u),
+          0xa5a53e00u,
+      };
+      Decoder::Instruction inst;
+      Decoder::DecodeInstruction(code, 0, inst);
+      Check(inst.opcode == opcode && inst.family == Decoder::Family::VOP3 &&
+                inst.src_count == 1 && inst.word_count == (source == 255u ? 3u : 2u),
+            "VOP3 conversion lost its opcode, unary source, or literal length");
+      Check(inst.dst.kind == Decoder::OperandKind::Vgpr && inst.dst.reg == 20 &&
+                inst.dst.sdwa_sel == 6 && inst.dst.clamp && inst.dst.omod == 2,
+            "VOP3 conversion lost its full-width destination or output modifiers");
+      const bool half = opcode == Decoder::Opcode::V_CVT_F32_F16;
+      Check(inst.src0.absolute == half && inst.src0.negate == half,
+            "VOP3 conversion must apply half-float modifiers and ignore byte modifiers");
+      Decoder::Operand expected_source;
+      Decoder::DecodeScalarSource(source, 0, expected_source);
+      Check(inst.src0.kind == expected_source.kind && inst.src0.reg == expected_source.reg &&
+                (source != 255u || inst.src0.value == code[2]),
+            "VOP3 conversion changed its source register or literal value");
+    }
+    for (const uint32_t op_sel : {1u, 2u, 4u, 8u}) {
+      const uint32_t code[] = {EncodeVop3Word0(encoding, 20, op_sel),
+                               EncodeVop3Word1(261, 0, 0)};
+      Decoder::Instruction inst;
+      Decoder::DecodeInstruction(code, 0, inst);
+      Check(inst.opcode == Decoder::Opcode::UNSUPPORTED,
+            "VOP3 conversion accepted unsupported OP_SEL bits");
+    }
+  }
+}
+
 void TestNewShaderRecompilerMoreAluFamilies() {
  const uint32_t shader[] = {
      EncodeSopk(0x00, 9, 7), // s_movk_i32 s9, 7
@@ -10550,6 +10594,43 @@ void TestNewShaderRecompilerNativeU64Translation() {
        "captured VOP3 V_LSHLREV_B64 was not present in the decoded dump");
  CheckSpirvBinaryValidates(result.spirv);

+  const uint32_t ashrrev_i64_shader[] = {
+      0xd7010001u, 0x0000d50cu,
+      EncodeSopp(0x01),
+  };
+  Decoder::Instruction decoded_ashrrev_i64;
+  ShaderRecompiler::Decoder::DecodeInstruction(ashrrev_i64_shader, 0u,
+                                               decoded_ashrrev_i64);
+  Check(decoded_ashrrev_i64.opcode == Decoder::Opcode::V_ASHRREV_I64 &&
+            decoded_ashrrev_i64.dst.kind == Decoder::OperandKind::Vgpr &&
+            decoded_ashrrev_i64.dst.reg == 1u &&
+            decoded_ashrrev_i64.src_count == 2u &&
+            decoded_ashrrev_i64.src0.kind == Decoder::OperandKind::Vgpr &&
+            decoded_ashrrev_i64.src0.reg == 12u &&
+            decoded_ashrrev_i64.src1.kind == Decoder::OperandKind::VccLo,
+        "decoder rejected VOP3 V_ASHRREV_I64 fields");
+  result = RecompileForTest(ashrrev_i64_shader, options);
+  Check((result.decoded_dump.find("V_ASHRREV_I64 v1, v12, vcc_lo") != std::string::npos),
+        "VOP3 V_ASHRREV_I64 was not present in the decoded dump");
+  CheckSpirvBinaryValidates(result.spirv);
+
+  const uint32_t ashrrev_captured[] = {0xd7010007u, 0x00020498u, EncodeSopp(0x01)};
+  Decoder::Instruction decoded_ashrrev_captured;
+  ShaderRecompiler::Decoder::DecodeInstruction(ashrrev_captured, 0u,
+                                               decoded_ashrrev_captured);
+  Check(decoded_ashrrev_captured.opcode == Decoder::Opcode::V_ASHRREV_I64 &&
+            decoded_ashrrev_captured.family == Decoder::Family::VOP3 &&
+            decoded_ashrrev_captured.src_count == 2u &&
+            decoded_ashrrev_captured.dst.kind == Decoder::OperandKind::Vgpr &&
+            decoded_ashrrev_captured.dst.reg == 7u &&
+            decoded_ashrrev_captured.src0.kind == Decoder::OperandKind::IntegerInlineConstant &&
+            decoded_ashrrev_captured.src0.value == 24u &&
+            decoded_ashrrev_captured.src1.kind == Decoder::OperandKind::Vgpr &&
+            decoded_ashrrev_captured.src1.reg == 2u,
+        "captured VOP3 V_ASHRREV_I64 opcode 0x301 was not decoded");
+  result = RecompileForTest(ashrrev_captured, options);
+  CheckSpirvBinaryValidates(result.spirv);
+
  const uint32_t cmpx_i64_shader[] = {
      0xd4b5007eu, 0x00020e80u, // v_cmpx_ne_i64 exec, 0, v[7:8]
      EncodeSopp(0x01),
@@ -11575,7 +11656,7 @@ void TestEmbeddedFetchPreservesSharedScalarLoad() {
                             reinterpret_cast<uint64_t>(attributes.data())};
  std::memcpy(user_data.data() + 12, tables, sizeof(tables));
  ShaderVertexInputInfo input{};
-  input.fetch_embedded = true;
+  input.has_dynamic_vertex_fetch = true;
  input.fetch_buffer_reg = 12;
  input.fetch_attrib_reg = 14;
  input.resources_num = 1;
@@ -17121,6 +17202,7 @@ int main() {
  using namespace Libs::Graphics;

  EnsureConfigInitialized();
+  TestVop3FloatConversionDecode();
  TestRayTracingInstructions();
  TestResourceDescriptorClassification();
  TestShaderBufferResourceSize();
Author	SHA1	Message	Date
nmzik	5e837a9bea	shader: share formatted indirect buffer load emission	2026-10-29 05:52:04 +01:00
nmzikandFrenzied	a0b99a8c21	ngs2: accept larger custom voice state buffers Co-authored-by: Frenzied <frenziedsaint@gmail.com>	2026-10-29 05:52:04 +01:00
nmzikandFrenzied	8cfe1c4d34	ngs2: apply custom channel matrices before mastering output Co-authored-by: Frenzied <frenziedsaint@gmail.com>	2026-10-29 05:52:04 +01:00
Çağatay Güçlü	f71da54b9b	shader: support fractional-even tessellation spacing	2026-10-29 05:52:04 +01:00
Çağatay Güçlü	63207739bc	shader: initialize local vertex count from active lanes	2026-10-29 05:52:04 +01:00
Çağatay Güçlü	143f640706	shader: preserve counterclockwise tessellation output	2026-10-29 05:52:04 +01:00
nmzik	54b937c85f	shader: support high products of 24-bit integers	2026-10-28 22:49:21 +01:00
nmzik	22610cb93e	shader: support unsigned 16-bit multiply-add	2026-10-28 22:41:29 +01:00
nmzik	0386f33bf3	shader: support four 16-bit integer ternary operations	2026-10-28 22:37:24 +01:00
nmzik	138d5b86c4	shader: support S_QUADMASK_B32	2026-10-28 22:31:16 +01:00
nmzik	963aaeb3b6	gpu: preserve low-alpha 2/10/10/10 render-target storage	2026-10-28 22:29:50 +01:00
Çağatay Güçlü	77458260f1	shader: reflect masked hull control-point addresses	2026-10-28 22:19:09 +01:00
nmzik	41de2fd64e	shader: support five VOP3 float conversion forms	2026-10-28 22:19:02 +01:00
Çağatay Güçlü	128c938514	shader: support quad tessellation domains and factor routing	2026-10-28 22:09:39 +01:00
Chai T	a8e3c29b2e	shader: implement V_ASHRREV_I64 (VOP3 0x301)	2026-10-28 21:51:28 +01:00
Devran Cosmo Uenal	c528324d27	audio: let the AJM Opus decoder write part of a packet (#1914 ) ajm: buffer and resample Opus output	2026-10-28 21:31:16 +01:00
Mr. Big BritchesandClaude Opus 5.5	c67352d5d6	ci: build and run the NGS2 sampler tests (#1957 ) ngs2_sampler_tests is registered with ctest but no job built it or listed it in the test filter, so the NGS2 unit tests never ran in CI. Each platform now builds the target and includes ngs2_sampler in its ctest filter. Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>	2026-10-28 19:58:06 +01:00
Nawid3333	6c01401862	graphics: accept video-out buffers with 128-byte independent DCC blocks (Gran Turismo 7) (#1945 ) graphics: accept video-out buffers with 128-byte independent DCC blocks	2026-10-28 04:44:09 +01:00
nmzik	a833e1b735	graphics: fix vertex and instance IDs for unpatched indirect draws	2026-10-28 04:33:03 +01:00
nmzik	1be1c6c4f9	graphics: revert `a358b847c` command-processor display waits	2026-10-27 23:15:01 +01:00
M4andnmzik	69fdb03dc2	Launcher display options (#1938 ) * launcher: add minimize and maximize buttons to the main window The main window is a QDialog, which on Windows only gets a close button. * graphics: keep the game's aspect ratio when presenting to the window The guest frame was blitted across the whole swapchain image, so a window or fullscreen display whose shape differs from the frame stretched the picture. Fit the frame inside the window instead and clear the remaining bars to black. The video-out overlay layer is drawn into the same rectangle so it stays aligned with the frame underneath. * launcher: accept custom screen resolutions such as 21:9 The resolution was a fixed enum of four 16:9 sizes. Make the combo box editable so any WIDTHxHEIGHT between 320 and 16384 can be entered, and add ultrawide presets (2560x1080, 3440x1440, 3840x1600, 5120x2160). Settings saved by older launchers stored the enum key ("R1920X1080"); NormalizeResolution still accepts that form, so they load unchanged. * launcher: add a Stretch to window option Presentation now keeps the game's aspect ratio with black bars. Add --stretch to fill the whole window as before, exposed in the launcher as a "Stretch to window" checkbox under General, off by default. * launcher: make the settings dialog resizable and its tabs scrollable The dialog had no maximize button and pinned its minimum width to its initial 900 px, and without scroll areas it could not be made shorter than its content. Add a maximize button (not minimize, since the dialog is modal), drop the width pin, and wrap each tab page in a frameless scroll area. With the scroll areas the layout's minimum size no longer forces the dialog open, so size it explicitly: grow it to fit the largest tab, capped at the available screen, when there is no saved geometry and whenever the global settings add the game folders panel. * launcher: simplify display options and preserve saved settings --------- Co-authored-by: nmzik <Nmzik@mail.ru>	2026-10-27 21:59:56 +01:00
nmzik	2efd31443f	audio: support 512 object ports per AudioOut2 context	2026-10-27 21:41:30 +01:00
nmzikandMajed-allababedi	9eef2e1434	graphics: reuse range-set nodes when merging additions Co-authored-by: Majed-allababedi <60163390+Majed-allababedi@users.noreply.github.com>	2026-10-27 21:41:30 +01:00
nmzik	39faa62ebd	controller: merge stick and sensor events into the newest unread pad sample	2026-10-27 21:15:03 +01:00
nmzik	7ff8aff5e6	audio: drain the AudioOut2 queue at an exact 48 kHz grain rate	2026-10-27 21:15:03 +01:00
nmzik	64a554dea7	audio: a pushed AudioOut2 grain consumes its ports' PCM	2026-10-27 21:15:03 +01:00
nmzik	4b0282e343	savedata: sort directory searches by the requested key	2026-10-27 21:15:03 +01:00
nmzik	ffc81377b2	memory: refactor (#1941 ) * memory: refactor * tests: avoid Windows small macro collision * memory: retain allocation nodes on partial release	2026-10-27 21:14:11 +01:00
bigjjmoneymanandnmzik	0685bffb2f	texture cache: convert D16<->D24/D32 depth copies instead of crashing (#1873 ) gpu: convert guest D16 depth between native D16 and stenciled four-byte host depth Co-authored-by: nmzik <Nmzik@mail.ru>	2026-10-27 10:41:22 +01:00
nmzik	a5b3673a39	graphics: draw patched-offset indirect draws from GPU-read records	2026-10-27 10:13:15 +01:00
nmzik	09753bf8ff	memory: change defaults	2026-10-27 10:13:15 +01:00
nmzik	9d8c458535	gpu: keep reused smaller image tables resident and drop them a frame after their last draw	2026-10-27 10:13:15 +01:00
Mr. Big Britches Claude Opus 5.5 nmzik	358333125d	libFont: ship Noto Sans as the fallback font (#1917 ) * cmake: download Noto Sans at configure time Fetch the hinted Regular, Bold, Italic and Bold Italic of Noto Sans v2.015 and their licence from the Noto project's font distribution (notofonts.github.io) with FetchContent, like the other third-party dependencies. The commit and each file's SHA-256 are pinned in variables. Only these four styles are needed, so they are fetched individually instead of through the release archive (117 MB, every style and build). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * cmake: ship Noto Sans beside the emulator Copy the downloaded fonts to assets/fonts next to kyty_emulator on every build, like the touch cursors and the trophy sound, and install them there with the SIL Open Font License under licenses/noto-sans. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * libFont: fall back to the shipped Noto Sans instead of Tracy's Roboto Fonts stb_truetype cannot read, including every system font set opened with sceFontOpenFontSet, used Roboto from Tracy's source folder, found relative to the source tree or the working directory. Installed builds have neither, so their system fonts drew every character as '?'. Load Noto Sans Regular, Bold, Italic or Bold Italic from assets/fonts next to the emulator instead, and log when a file is missing. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * libFont: fetch Noto Sans with file(DOWNLOAD) and drop the style substitution --------- Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com> Co-authored-by: nmzik <Nmzik@mail.ru>	2026-10-27 07:51:24 +01:00