shader: share formatted indirect buffer load emission

Co-authored-by: Frenzied <frenziedsaint@gmail.com>

.github/workflows/build.yml

@@ -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

CMakeLists.txt

@@ -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

src/common/emulatorConfig.cpp

@@ -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;
 }

src/common/emulatorConfig.h

@@ -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();

src/graphics/guest_gpu/gpu_format.cpp

@@ -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},

src/graphics/guest_gpu/graphicsRun.cpp

@@ -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);
-					}
 					for (auto& queue: gpu->m_queues) {
 						if (!queue.empty()) {
 							queue.front().blocked = false;
@@ -866,9 +856,8 @@ 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);
-	// The GPU reads the records if the shaders allow; the VS prolog loads the patched SGPRs from
+	// The GPU reads the records if the shaders allow; the VS prolog removes Vulkan's bases
-	// Vulkan's bases. With no SGPR patched, Vulkan applies the offsets, which matches while they
+	// from the index VGPRs and copies them only to the user SGPRs the packet patches.
-	// are zero.
 	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(),
@@ -960,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.
+	BufferFlush();
-	if (!m_renderer.GetVideoOut().IsSafeForRendering(static_cast<int>(video_out_handle),
+
-	                                                 static_cast<int>(display_buffer_index))) {
+	m_renderer.GetVideoOut().WaitFlipDone(static_cast<int>(video_out_handle),
-		SuspendPm4();
+	                                      static_cast<int>(display_buffer_index));
-	}
 }
 
 void CommandProcessor::ReleaseMem(uint32_t data_sel, uint32_t interrupt_selector, void* address,

src/graphics/guest_gpu/graphicsRun.h

@@ -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 =

src/graphics/host_gpu/rangeSet.h

@@ -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;
+		if (it == m_ranges.end() || it->first > address) {
-		uint64_t last  = end;
+			it = m_ranges.emplace_hint(it, address, end);
-		while (it != m_ranges.end() && it->first <= last) {
+		} else {
-			begin = std::min(begin, it->first);
+			if (it->second >= end) {
-			last  = std::max(last, it->second);
+				return;
-			it    = m_ranges.erase(it);
+			}
+			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) {

src/graphics/host_gpu/renderer/cache/textureCache.cpp

@@ -488,16 +488,15 @@ void TextureCache::RefreshCopySource(ImageId id) {
 }
 
 bool TextureCache::CopyD16(Image& destination, Image& source) {
-	const bool source_depth      = source.info.IsDepth();
+	// Converts guest D16 between a four-byte host depth aspect and two-byte color or D16 depth.
-	const bool destination_depth = destination.info.IsDepth();
+	const bool source_wide = DepthAspectTransferBytes(source.backing.format) == sizeof(uint32_t);
-	if (source_depth == destination_depth) {
+	if (source_wide == (DepthAspectTransferBytes(destination.backing.format) == sizeof(uint32_t))) {
 		return false;
 	}
-	auto&      depth          = source_depth ? source : destination;
+	auto& depth  = source_wide ? source : destination;
-	auto&      color          = source_depth ? destination : source;
+	auto& narrow = source_wide ? destination : source;
-	const auto transfer_bytes = DepthAspectTransferBytes(depth.backing.format);
 	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,12 +506,10 @@ 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;
+	    static_cast<uint64_t>(depth.info.pitch) * depth.info.extent.height * sizeof(uint32_t);
-	const uint64_t color_slice =
+	const uint64_t narrow_slice =
-	    static_cast<uint64_t>(color.info.pitch) * color.info.extent.height * sizeof(uint16_t);
+	    static_cast<uint64_t>(narrow.info.pitch) * narrow.info.extent.height * sizeof(uint16_t);
-	EXIT_IF(layers == 0 || depth_slice > UINT64_MAX / layers || color_slice > UINT64_MAX / layers);
+	EXIT_IF(layers == 0 || depth_slice > UINT64_MAX / layers || narrow_slice > UINT64_MAX / layers);
-	const auto                       depth_size = depth_slice * layers;
-	const auto                       color_size = color_slice * layers;
 	const std::array depth_copies {
 	    vk::BufferImageCopy {0,
 	                         depth.info.pitch,
@@ -520,29 +517,31 @@ bool TextureCache::CopyD16(Image& destination, Image& source) {
 	                         {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,
+	                         narrow.info.pitch,
-	                         color.info.extent.height,
+	                         narrow.info.extent.height,
-	                         {vk::ImageAspectFlagBits::eColor, 0, 0, layers},
+	                         {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 depth_buffer  = m_tiler.GetScratchBuffer(depth_slice * layers);
-	auto                         color_buffer = m_tiler.GetScratchBuffer(color_size, depth_buffer.buffer);
+	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),
+	    .source_row_stride   = static_cast<uint64_t>(narrow.info.pitch) * sizeof(uint16_t),
-	    .target_row_stride   = static_cast<uint64_t>(depth.info.pitch) * transfer_bytes,
+	    .target_row_stride   = static_cast<uint64_t>(depth.info.pitch) * sizeof(uint32_t),
-	    .source_slice_stride = color_slice,
+	    .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,
+		destination.Upload(narrow_copies, narrow_buffer.buffer, narrow_buffer.offset,
-		                   color_buffer.size);
+		                   narrow_buffer.size);
 	} else {
-		source.Download(color_copies, color_buffer.buffer, color_buffer.offset, color_buffer.size);
+		source.Download(narrow_copies, narrow_buffer.buffer, narrow_buffer.offset,
-		m_tiler.ConvertD16(color_buffer, depth_buffer, TileManager::D16Direction::Promote, d32,
+		                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);

src/graphics/host_gpu/renderer/image/imageInfo.h

@@ -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;
 		}
 	}
@@ -289,6 +297,7 @@ ClassifyVideoOutCompression(bool compressed, uint64_t metadata_address, uint32_t
                             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_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;
 	}
 }

src/graphics/host_gpu/renderer/image/textureCommon.cpp

@@ -32,8 +32,12 @@ struct HostFormatInfo {
 
 HostFormatInfo ResolveHostFormat(Prospero::BufferFormat guest_format,
                                  Prospero::ChannelOrder order) {
-	if (guest_format == Prospero::BufferFormat::k10_11_11Float) {
+	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) {

src/graphics/host_gpu/renderer/image/tiler.cpp

@@ -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 ||

src/graphics/host_gpu/renderer/pipeline/pipelineCache.cpp

@@ -689,6 +689,7 @@ void PipelineCache::GetGraphicsPrograms(const CommandBuffer& command, DrawRender
 	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.

src/graphics/host_gpu/renderer/pipeline/pipelineCache.h

@@ -139,6 +139,7 @@ struct DrawRenderState {
 	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};

src/graphics/host_gpu/renderer/renderDraw.cpp

@@ -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};
 	}
 
@@ -1251,16 +1251,26 @@ void RenderExecutor::ExecutePreparedDraw(uint64_t submit_id, CommandBuffer& buff
 }
 
 // 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 patched SGPRs. Indirect records carry
+// and quad-list draws for counts, tessellation draws for the index/SGPR prolog. Indirect records
-// their own offsets.
+// 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 ||
+		if (!vs.indirect_draw ||
+		    vs.stage.module->info.stage == ShaderType::Mesh ||
 		    ucfg.GetPrimType() == Prospero::PrimitiveType::kQuadListLegacy ||
 		    vs.indirect_offsets != indirect.offsets) {
 			return std::nullopt;
@@ -1269,7 +1279,7 @@ static std::optional<DrawEmitInfo> ResolveDrawEmit(const HW::UserConfig&
 	}
 	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};
@@ -1326,7 +1336,8 @@ bool RenderExecutor::DrawIndex(uint64_t submit_id, CommandBuffer& buffer,
 
 	hw_check(buffer);
 
-	DrawRenderState state {.indirect_offsets = args.indirect.offsets};
+	DrawRenderState state {.indirect_offsets = args.indirect.offsets,
+	                       .indirect_draw    = args.indirect.address != 0};
 	if (!GetDrawTopology(ucfg, state.topology)) {
 		return true;
 	}
@@ -1431,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 {.indirect_offsets = args.indirect.offsets};
+	DrawRenderState state {.indirect_offsets = args.indirect.offsets,
+	                       .indirect_draw    = args.indirect.address != 0};
 	if (!GetDrawTopology(ucfg, state.topology)) {
 		ResetBindings();
 		return true;

src/graphics/host_gpu/shaders/gpu_tiler_color.inc

@@ -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;
 }

src/graphics/presentation/videoOut.cpp

@@ -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() {
+void FlipQueue::Wait(VideoOutConfig& cfg, int index) {
-	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) {
 	Common::LockGuard lock(m_mutex);
+
+	auto has_request = [this, &cfg, index] {
 		auto matches = [&cfg, index](const auto& r) { return r.cfg == &cfg && r.index == index; };
-	const bool safe = cfg.flip_status.currentBuffer != index &&
+		return std::any_of(m_requests.begin(), m_requests.end(), matches) ||
-	                  std::none_of(m_requests.begin(), m_requests.end(), matches) &&
+		       std::any_of(m_cpu_requests.begin(), m_cpu_requests.end(), matches);
-	                  std::none_of(m_cpu_requests.begin(), m_cpu_requests.end(), matches);
+	};
-	m_gpu_waiting = !safe;
+	while (has_request()) {
-	return safe;
+		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) {

src/graphics/presentation/videoOut.h

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

src/graphics/presentation/window/swapchain.cpp

@@ -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 {
+	const vk::Viewport viewport {static_cast<float>(rect.offset.x),
-	    0.0f, 0.0f, static_cast<float>(m_extent.width), static_cast<float>(m_extent.height),
+	                             static_cast<float>(rect.offset.y),
-	    0.0f, 1.0f};
+	                             static_cast<float>(rect.extent.width),
-	const vk::Rect2D scissor {{0, 0}, m_extent};
+	                             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[0].x = game_rect.offset.x;
-		region.dstOffsets[1].y               = static_cast<int>(m_extent.height);
+		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;

src/graphics/shader/recompiler/ShaderRecompiler.cpp

@@ -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);

src/graphics/shader/recompiler/Tessellation.cpp

@@ -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:
+				case Opcode::V_AND_B32: {
-					if ((lhs.kind == Kind::PackedControlPoint && rhs == constant(0xffu)) ||
+					const auto mask = lhs.kind == Kind::PackedControlPoint ? rhs
-					    (rhs.kind == Kind::PackedControlPoint && lhs == constant(0xffu)))
+					                : rhs.kind == Kind::PackedControlPoint ? lhs
-						value = constant(0u);
+					                                                       : 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;

src/graphics/shader/recompiler/backend/spirv/spirvEmitterInternal.h

@@ -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);

src/graphics/shader/recompiler/backend/spirv/spirvEmitterMemory.cpp

@@ -1026,10 +1026,8 @@ 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);
@@ -1072,7 +1070,6 @@ uint32_t LoadIndirectFormattedX(ValueEmitContext& ctx, const IR::Inst& inst) {
 	                          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;

src/graphics/shader/recompiler/backend/spirv/spirvEmitterProgram.cpp

@@ -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();

src/graphics/shader/recompiler/backend/spirv/spirvEmitterTessellation.cpp

@@ -30,11 +30,13 @@ uint32_t TessellationPointer(ValueEmitContext& ctx, const IR::Inst& inst) {
 	if (kind == Attribute::Factor) {
 		EXIT_NOT_IMPLEMENTED(!inst.Arg(1).IsImmediate());
 		const auto index       = inst.Arg(1).U32() / 4u;
-		const auto outer = index < 3u;
+		const bool quad        = tess.domain == 2u;
-		EXIT_NOT_IMPLEMENTED(index >= 4u);
+		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.
+	// Triangle or quad domain, integer or fractional partitioning, triangle output.
-	EXIT_NOT_IMPLEMENTED(tess.domain != 1u ||
+	EXIT_NOT_IMPLEMENTED((tess.domain != 1u && tess.domain != 2u) ||
-	                     (tess.partitioning != 0u && tess.partitioning != 2u) ||
+	                     (tess.partitioning != 0u && tess.partitioning != 2u && tess.partitioning != 3u) ||
-	                     tess.output_topology != 2u);
+	                     (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);
+		                     : tess.partitioning == 2u ? spv::ExecutionModeSpacingFractionalOdd
-		state.builder.AddExecutionMode(state.main_func, spv::ExecutionModeVertexOrderCw);
+		                                                : spv::ExecutionModeSpacingFractionalEven);
+		state.builder.AddExecutionMode(state.main_func, tess.output_topology == 2u
+		                                                   ? spv::ExecutionModeVertexOrderCw
+		                                                   : spv::ExecutionModeVertexOrderCcw);
 	}
 }
 

src/graphics/shader/recompiler/frontend/decode/ScalarAluOps.cpp

@@ -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},

src/graphics/shader/recompiler/frontend/decode/ShaderDecoder.cpp

@@ -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:

src/graphics/shader/recompiler/frontend/decode/ShaderDecoder.h

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

src/graphics/shader/recompiler/frontend/decode/VectorAluOps.cpp

@@ -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 ||
+	return opcode == Opcode::V_MIN3_I16 || opcode == Opcode::V_MIN3_U16 ||
-	       opcode == Opcode::V_MAD_I16;
+	       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;

src/graphics/shader/recompiler/frontend/translate/Integer.cpp

@@ -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) {
+void Translator::Integer16Mad(const Decoder::Instruction& inst, bool sign) {
-	const auto lhs    = ReadU16AsU32(inst.src0, true);
+	const auto lhs    = ReadU16AsU32(inst.src0, sign);
-	const auto rhs    = ReadU16AsU32(inst.src1, true);
+	const auto rhs    = ReadU16AsU32(inst.src1, sign);
-	auto       result = ir.IAdd(ir.IMul(lhs, rhs), ReadU16AsU32(inst.src2, true));
+	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)));
 	};
+	IR::U32 quads;
+	if (is64) {
 		const auto source = ReadU32Pair(inst.src0);
-	const auto quads  = ir.BitwiseOr(
+		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), ir.ConstructU64(quads, IR::U32(IR::Value(0u))));
-	WriteOperand(DestinationOperand(inst), result);
+	} else {
-	ir.SetScc(IR::U1(ir.Emit(IR::ValueOpcode::INotEqual64, {result, IR::Value(uint64_t {0})})));
+		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) {

src/graphics/shader/recompiler/frontend/translate/Scalar.cpp

@@ -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);

src/graphics/shader/recompiler/frontend/translate/Translate.cpp

@@ -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,21 +1388,17 @@ 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.
+			// 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 (offsets == ShaderIndirectOffsets {}) {
+				if (!options.input_info.vertex->indirect_draw) {
 					return builtin(kind);
 				}
 				const auto base = builtin(base_kind);

src/graphics/shader/recompiler/frontend/translate/Translator.h

@@ -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);

src/graphics/shader/recompiler/frontend/translate/Vector.cpp

@@ -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);

src/graphics/shader/shader.cpp

@@ -291,8 +291,7 @@ 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.has_dynamic_vertex_fetch = true;
-		info.fetch_embedded   = true;
 		info.fetch_attrib_reg         = metadata.vertex_attrib_reg;
 		info.fetch_buffer_reg         = metadata.vertex_buffer_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,7 +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));
+	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) {
@@ -693,9 +693,7 @@ void ShaderDbgDumpInputInfo(const ShaderVertexInputInfo& info) {
 
 	LOGF("ShaderDbgDumpInputInfo()\n");
 
-	LOGF("\t fetch_external = %s\n"
+	LOGF("\t has_dynamic_vertex_fetch = %s\n", info.has_dynamic_vertex_fetch ? "true" : "false");
-	     "\t fetch_embedded = %s\n",
-	     info.fetch_external ? "true" : "false", info.fetch_embedded ? "true" : "false");
 
 	for (int i = 0; i < info.resources_num; i++) {
 		LOGF("\t input %d\n", i);

src/graphics/shader/shader.h

@@ -125,8 +125,7 @@ struct ShaderTessellationInputInfo {
 };
 
 // GS user-data slots an indirect draw packet writes its record's base vertex and first instance
-// to. Like the hardware index VGPRs, the vertex shader's indices exclude both bases once either
+// to. A packet can leave both slots unused even when its record contains nonzero offsets.
-// slot is set.
 struct ShaderIndirectOffsets {
 	static constexpr uint8_t None = UINT8_MAX;
 
@@ -159,8 +158,10 @@ struct ShaderVertexInputInfo {
 	ShaderMeshInputInfo      mesh;
 	ShaderTessellationInputInfo tess;
 	ShaderIndirectOffsets   indirect_offsets;
-	bool                    fetch_external      = false;
+	// GPU-read records add Vulkan bases even when the packet patches no user SGPRs.
-	bool                    fetch_embedded      = false;
+	bool                    indirect_draw       = false;
+	// Attribute-table metadata enables dynamic fetch detection; loads are matched later.
+	bool                    has_dynamic_vertex_fetch = false;
 };
 
 struct ShaderComputeInputInfo: ShaderWorkgroupInputInfo {

src/kernel/memory.h

@@ -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
 

src/kernel/memoryAddressSpace.inc

@@ -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;
+		size_t removed  = 0;
-		removed.reserve(targets.size());
+		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)) {
-			if (!UnmapOne(entry.vaddr, entry.size, &part_placeholder_preserved)) {
+				rollback();
-				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;
-				}
 				return false;
 			}
-			EXIT_IF(!part_placeholder_preserved);
+			++removed;
-			removed.push_back(entry);
 		}
 
 #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)) {
+		if (removed > 1 && !CoalescePlaceholders(vaddr, size)) {
-			for (auto it = removed.rbegin(); it != removed.rend(); ++it) {
+			rollback();
-				EXIT_IF(!MapFixed(it->vaddr, it->size, it->backing_offset, it->mode));
-			}
 			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)) {
-		if (!m_backing->Unmap(vaddr, size, &placeholder_preserved)) {
-			return false;
-		}
-		if (!placeholder_preserved) {
 			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 &&
+			if (range_start >= range_end) {
-			    !VirtualMemory::Protect(range_start, range_end - range_start, mode)) {
+				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,11 +1312,20 @@ 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) {
+		// 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)) {
@@ -1332,8 +1341,7 @@ private:
 			return m_free.end();
 		}
 		auto it = std::prev(next);
-		auto end = it->first + it->second;
+		return GuestFreeRangeContains(it->first, it->second, vaddr, size) ? it : m_free.end();
-		return (vaddr >= it->first && vaddr + size <= end ? it : m_free.end());
 	}
 
 	bool SplitUnlocked(std::map<uint64_t, uint64_t>::iterator it, uint64_t vaddr, uint64_t size) {

src/launcher/forms/configuration_edit_dialog.ui

@@ -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>

src/launcher/include/configuration.h

@@ -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 << QString(me.key(i));
-		ret << (key.startsWith('R') && key.size() > 2 && key.at(1).isDigit()
-		            ? key.remove('R').toLower()
-		            : key);
 	}
 	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(
+	return static_cast<T>(me.keyToValue(text.toUtf8().constData()));
-	    ((text.size() > 1 && text.at(0).isDigit()) ? 'R' + text.toUpper() : text).toUtf8().data()));
 }
 
 template <class T>
 QString EnumToText(T value) {
-	auto me  = QMetaEnum::fromType<T>();
+	return QString(QMetaEnum::fromType<T>().valueToKey(static_cast<int>(value)));
-	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);
 }
 
 struct ControllerSettings {
@@ -80,13 +73,10 @@ public:
 	static constexpr int DEFAULT_CONSOLE_LANGUAGE = 1;
 	static constexpr int MAX_CONSOLE_LANGUAGE     = 29;
 
-	enum class Resolution {
+	static constexpr int MIN_SCREEN_SIZE = 320;
-		R1280X720,
+	static constexpr int MAX_SCREEN_SIZE = 16384;
-		R1920X1080,
+
-		R2560X1440,
+	static QString NormalizeResolution(const QString& text);
-		R3840X2160,
-	};
-	Q_ENUM(Resolution)
 
 	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;

src/launcher/include/configurationEditDialog.h

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

src/launcher/src/configuration.cpp

@@ -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);

src/launcher/src/configurationEditDialog.cpp

@@ -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),
+    : QDialog(parent, Qt::WindowMaximizeButtonHint | Qt::WindowCloseButtonHint),
-      m_info(info) {
+      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);
 

src/launcher/src/mainDialog.cpp

@@ -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");

src/libs/ajm.cpp

@@ -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");

src/libs/ajm/decoder.h

@@ -56,6 +56,12 @@ 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;
@@ -65,6 +71,7 @@ struct AjmDecodeResult {
 	uint32_t                frames                = 0;
 	uint32_t                frames_per_packet     = 0;
 	AjmSidebandFormat       format {};
+	AjmSidebandResampleInfo resample_info {};
 };
 
 constexpr int32_t AJM_RESULT_NOT_INITIALIZED    = 0x00000001;

src/libs/ajm/resampler.h

@@ -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

src/libs/controller.cpp

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

src/libs/libAudio2.cpp

@@ -13,6 +13,7 @@
 #include <atomic>
 #include <cmath>
 #include <cstring>
+#include <span>
 #include <vector>
 
 namespace Libs::Audio {
@@ -170,15 +171,10 @@ 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;
 
-struct AudioOut2ContextState {
+static constexpr uint32_t AUDIOOUT2_CHANNEL_PORTS_MAX = 32;
-	bool                   used                = false;
+static constexpr uint32_t AUDIOOUT2_OBJECT_PORTS_MAX  = 512;
-	AudioOut2ContextHandle handle              = 0;
+static constexpr uint32_t AUDIOOUT2_PORT_SLOT_BITS    = 10;
-	uint32_t               queue_depth         = 4;
+static constexpr uint32_t AUDIOOUT2_CONTEXT_SLOT_BITS = 4;
-	uint32_t               queued              = 0;
-	uint32_t               num_grains          = 512;
-	uint64_t               last_update         = 0;
-	int                    object_audio_handle = 0;
-};
 
 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 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;
 	}
-
+	if (now <= state->last_update) {
-	const auto grain_micros = static_cast<uint64_t>(audioout2_grain_micros(state->num_grains));
-	if (now <= state->last_update || grain_micros == 0) {
 		return;
 	}
 
-	const auto elapsed = now - state->last_update;
+	const auto grain   = uint64_t {state->num_grains} * 125;
-	const auto drained = std::min<uint64_t>(state->queued, elapsed / grain_micros);
+	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) {
+	const auto context_slot =
-		if (state.used && state.handle == port) {
+	    (port >> AUDIOOUT2_PORT_SLOT_BITS) & ((1u << AUDIOOUT2_CONTEXT_SLOT_BITS) - 1u);
-			return &state;
+	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) {
+	for (const auto& state: context.ports) {
-		if (state.used && state.context == context.handle && !state.pcm_data.empty()) {
+		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) {
+	// A port's PCM covers one grain: the push consumes it, so a port the game stops feeding
-		if (!state.used || state.context != ctx || state.pcm_data.empty()) {
+	// 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_ports   = 8;
-	params->max_object_ports       = 256;
+	params->queue_depth = 1;
-	params->guarantee_object_ports = 0;
+	params->num_grains  = 256;
-	params->queue_depth            = 4;
-	params->num_grains             = 512;
-	params->flags                  = 1;
 
 	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,7 +574,7 @@ int KYTY_SYSV_ABI AudioOut2ContextDestroy(AudioOut2ContextHandle ctx) {
 	PRINT_NAME();
 	LOGF("\t ctx = 0x%016" PRIx64 "\n", ctx);
 
-	std::array<int, 257> audio_handles {};
+	std::array<int, AUDIOOUT2_CHANNEL_PORTS_MAX + 1> audio_handles {};
 	size_t                                           audio_handles_num = 0;
 
 	g_audioout2_state_mutex.Lock();
@@ -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 (const auto& port_state: state->ports) {
-	}
+			if (port_state.used && port_state.audio_handle > 0) {
-	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()) {
 				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: state->ports) {
-	for (auto& port: g_audioout2_ports) {
-		if (port.used && port.context == ctx) {
 			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,19 +691,17 @@ 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) &&
+	if (is_object && (params->port_type != 0x100 || params->sampling_freq != 48000 ||
-	    (params->port_type != 0x100 || context_state->object_audio_handle == 0 ||
-	     params->sampling_freq != 48000 ||
 	                  (audio_format != AudioInternal::Format::FloatMono &&
 	                   audio_format != AudioInternal::Format::Signed16bitMono))) {
 		g_audioout2_state_mutex.Unlock();
@@ -688,14 +714,21 @@ int KYTY_SYSV_ABI AudioOut2PortCreate(AudioOut2ContextHandle ctx, const AudioOut
 		g_audioout2_state_mutex.Unlock();
 		return AUDIO_OUT2_ERROR_INVALID_PARAM;
 	}
-	AudioOut2PortStateEntry* port_state = nullptr;
+	auto slots = std::span(context_state->ports);
-	for (auto& candidate: g_audioout2_ports) {
+	slots =
-		if (!candidate.used) {
+	    is_object ? slots.subspan(context_state->max_ports) : slots.first(context_state->max_ports);
-			port_state = &candidate;
+	const auto available = std::find_if(slots.begin(), slots.end(),
-			break;
+	                                    [](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);
-	if (port_state != nullptr) {
+	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;
@@ -707,23 +740,19 @@ int KYTY_SYSV_ABI AudioOut2PortCreate(AudioOut2ContextHandle ctx, const AudioOut
 	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) {
-		return AUDIO_OUT2_ERROR_PORT_FULL;
-	}
-
 	int audio_handle = 0;
 
-	if (audio_format != AudioInternal::Format::Unknown &&
+	if (audio_format != AudioInternal::Format::Unknown && !is_object) {
-	    !audioout2_port_type_is_object(params->port_type)) {
 		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"
@@ -905,10 +934,12 @@ int KYTY_SYSV_ABI AudioOut2UserDestroy(AudioOut2UserHandle 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(),
+	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;
 }

src/libs/libDbgAsan.cpp

@@ -39,7 +39,7 @@ struct Global {
 
 struct Shadow {
 	uintptr_t start         = 0;
-	int       count = 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 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)) {
@@ -79,7 +88,7 @@ static void AllocShadow(uintptr_t min_addr, uintptr_t max_addr) {
 
 			Shadow s;
 			s.start         = page;
-			s.count = 1;
+			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--;
+			const auto bytes = ShadowCoverage(page, min_addr, max_addr);
-
+			if (bytes >= it->covered_bytes) {
-			if (it->count <= 0) {
 				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();
 

src/libs/libSaveData.cpp

@@ -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) {
+	// A save whose metadata cannot be read sorts as 0.
-		return std::strcmp(a.c_str(), b.c_str()) < 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;
+		}
+	}
+	// Directory names break ties; descending order reverses both.
+	std::sort(dir_list.begin(), dir_list.end());
 	if (cond->order == SaveDataSortOrder::Descent) {
-		std::vector<std::string> reversed;
+		std::reverse(dir_list.begin(), dir_list.end());
-		for (size_t i = dir_list.size(); i > 0; i--) {
-			reversed.push_back(dir_list[i - 1]);
-		}
-		dir_list = reversed;
 	}
 
 	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;
 			}

src/libs/ngs2.cpp

@@ -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: {
+		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::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(configured_size < (sampler ? sizeof(Ngs2CustomSamplerVoiceState)
-			EXIT_NOT_IMPLEMENTED(state_size != configured_size);
+			                                                : sizeof(Ngs2CustomMasteringVoiceState)));
-			std::memset(state, 0, state_size);
+			if (state_size < configured_size) {
-			auto* sampler = reinterpret_cast<Ngs2CustomSamplerVoiceState*>(state);
+				return static_cast<int32_t>(0x804a8011u);
-			Ngs2FillSamplerState(*voice, *sampler);
+			}
+			std::memset(state, 0, configured_size);
+			if (sampler) {
+				Ngs2FillSamplerState(*voice, *reinterpret_cast<Ngs2CustomSamplerVoiceState*>(state));
+			} else {
+				state->state_flags = Ngs2GetStateFlags(voice);
+			}
 			break;
 		}
 		case Ngs2RackType::Sampler: {

src/main.cpp

@@ -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},

tests/AjmResampleTests.cpp

@@ -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");
+}

tests/AudioOut2PortTests.cpp

@@ -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,10 +484,24 @@ 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");
 }
 
@@ -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 std::vector<float> replacement_pcm(pcm.size(), -0.25f);
-	const Pcm replacement {pcm.data()};
+	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");
-		}
 		std::fill(floating.begin(), floating.end(), 0.5f);
 		std::fill(integer.begin(), integer.end(), 16384);
 		for (const auto& test: cases) {
@@ -569,11 +757,9 @@ 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");
 		}
-		}
 		AudioOut2::AudioOut2ContextAdvance(context);
 		SetAttribute(object, 8, UINT32_MAX);
 		SetAttribute(object, 1, 0.0f);
@@ -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);
+	const auto object = CreatePort(context, param);
-	SetPcm(CreatePort(context, param), object_pcm.data());
 	param.user_handle = other_user;
-	SetPcm(CreatePort(context, param), object_pcm.data());
+	const auto other_user_object = CreatePort(context, param);
-	std::fill(object_pcm.begin(), object_pcm.end(), 1.0f);
+	std::vector<float> object_pcm(512, 0.25f);
-	SetPcm(CreatePort(other_context, param), object_pcm.data());
+	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(object, object_pcm.data());
-	SetPcm(first, nullptr);
+	SetPcm(other_user_object, object_pcm.data());
-	SetPcm(second, nullptr);
 	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 g_process_time_frozen ? g_process_time.load() : g_process_time.fetch_add(1000);
-	return now.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();

tests/ControllerSettingsTests.cpp

@@ -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();

tests/Ngs2SamplerTests.cpp

@@ -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)
+	RenderRack sampler(system, 0x1000, sampler_option.common);
-	    std::byte sampler_storage[sizeof(Ngs2RackInternal) + 2 * sizeof(Ngs2VoiceInternal)];
+	RenderRack master(system, 0x3000, mastering_option.common);
-	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;
 	const auto system_handle = reinterpret_cast<uintptr_t>(&system);
-	Check(Ngs2RackCreate(system_handle, 0x1000, &sampler_option.common, &sampler_buffer,
+	auto* samplers = &sampler.Voice();
-	                     &sampler_handle) == OK &&
+	auto& mastering = master.Voice();
-	          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);
 	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();

tests/SaveDataMemoryTests.cpp

@@ -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);

tests/shaderCfgTests.cpp

@@ -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();