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
-	// Vulkan's bases. With no SGPR patched, Vulkan applies the offsets, which matches while they
-	// are zero.
+	// The GPU reads the records if the shaders allow; the VS prolog removes Vulkan's bases
+	// from the index VGPRs and copies them only to the user SGPRs the packet patches.
 	DrawIndirectArgs records {args_addr, draw_count, stride_in_bytes};
 	if (IndirectDrawOffsets(registers, records.offsets) && (!indexed || m_index_buffer_size != 0) &&
 	    (indexed ? executor.DrawIndex(m_submit_id, CurrentBuffer(),
@@ -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.
-	if (!m_renderer.GetVideoOut().IsSafeForRendering(static_cast<int>(video_out_handle),
-	                                                 static_cast<int>(display_buffer_index))) {
-		SuspendPm4();
-	}
+	BufferFlush();
+
+	m_renderer.GetVideoOut().WaitFlipDone(static_cast<int>(video_out_handle),
+	                                      static_cast<int>(display_buffer_index));
 }
 
 void CommandProcessor::ReleaseMem(uint32_t data_sel, uint32_t interrupt_selector, void* address,

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/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
-// their own offsets.
+// and quad-list draws for counts, tessellation draws for the index/SGPR prolog. Indirect records
+// carry their own offsets.
 static std::optional<DrawEmitInfo> ResolveDrawEmit(const HW::UserConfig&        ucfg,
                                                    const ShaderVertexInputInfo& vs,
                                                    DrawOffsetSource             source,
                                                    int32_t                      first_vertex,
                                                    uint32_t                     first_instance,
                                                    const DrawIndirectArgs&      indirect) {
+	if ((indirect.address != 0 || source == DrawOffsetSource::IndirectArgs) &&
+	    ucfg.GetIndexOffset() != 0) {
+		static std::once_flag warning_once;
+		std::call_once(warning_once, [&] {
+			Log::WriteToConsoleAndLog(fmt::format(
+			    "GPU: indirect draw ignores GE_INDX_OFFSET=0x{:08x}; geometry may be incorrect.\n",
+			    ucfg.GetIndexOffset()));
+		});
+	}
 	if (indirect.address != 0) {
-		if (vs.stage.module->info.stage == ShaderType::Mesh ||
+		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() {
-	m_done_cond_var.SignalAll();
-	// A stalled GPU cannot drain and shut down without re-checking under m_mutex, so it is alive here.
-	if (m_gpu_waiting) {
-		m_gpu_waiting = false;
-		m_presenter.Renderer().GetGpu().Wake();
-	}
-}
-
-// A display buffer is safe for rendering once it is neither displayed nor waiting to be.
-bool FlipQueue::IsSafeForRendering(VideoOutConfig& cfg, int index) {
+void FlipQueue::Wait(VideoOutConfig& cfg, int index) {
 	Common::LockGuard lock(m_mutex);
+
+	auto 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 &&
-	                  std::none_of(m_requests.begin(), m_requests.end(), matches) &&
-	                  std::none_of(m_cpu_requests.begin(), m_cpu_requests.end(), matches);
-	m_gpu_waiting = !safe;
-	return safe;
+		return std::any_of(m_requests.begin(), m_requests.end(), matches) ||
+		       std::any_of(m_cpu_requests.begin(), m_cpu_requests.end(), matches);
+	};
+	while (has_request()) {
+		m_done_cond_var.Wait(&m_mutex);
+	}
 }
 
 bool FlipQueue::Flip(uint32_t micros) {
@@ -1320,7 +1311,7 @@ bool FlipQueue::Flip(uint32_t micros) {
 		it = m_requests.erase(it);
 	}
 	m_processing = false;
-	SignalDoneLocked();
+	m_done_cond_var.SignalAll();
 	m_submit_slot_cond_var.SignalAll();
 	m_mutex.Unlock();
 	for (size_t i = count; i != 0; i--) {
@@ -1703,12 +1694,12 @@ void VideoOutDriver::WaitForSubmitSlot(int handle) {
 	m_impl->GetFlipQueue().WaitForSubmitSlot(*cfg);
 }
 
-bool VideoOutDriver::IsSafeForRendering(int handle, int index) {
+void VideoOutDriver::WaitFlipDone(int handle, int index) {
 	auto* ctx = m_impl->Get(handle);
 	EXIT_IF(ctx == nullptr);
 
 	EXIT_NOT_IMPLEMENTED(!IsValidBufferIndex(index));
-	return m_impl->GetFlipQueue().IsSafeForRendering(*ctx, index);
+	m_impl->GetFlipQueue().Wait(*ctx, index);
 }
 
 KYTY_SYSV_ABI int VideoOutGetFlipStatus(int handle, VideoOutFlipStatus* status) {

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 {
-	    0.0f, 0.0f, static_cast<float>(m_extent.width), static_cast<float>(m_extent.height),
-	    0.0f, 1.0f};
-	const vk::Rect2D scissor {{0, 0}, m_extent};
+	const vk::Viewport viewport {static_cast<float>(rect.offset.x),
+	                             static_cast<float>(rect.offset.y),
+	                             static_cast<float>(rect.extent.width),
+	                             static_cast<float>(rect.extent.height),
+	                             0.0f,
+	                             1.0f};
 	command.setViewport(0, 1, &viewport);
-	command.setScissor(0, 1, &scissor);
+	command.setScissor(0, 1, &rect);
 	command.draw(3, 1, 0, 0);
 	command.endRendering();
 }
@@ -799,6 +818,24 @@ void Swapchain::RecordPresentCommands(CommandBuffer& command, Presenter::Frame*
 	                           vk::PipelineStageFlagBits::eTransfer, vk::DependencyFlags {}, 0,
 	                           nullptr, 0, nullptr, 1, &to_transfer);
 
+	const auto* aspect_source = source != nullptr ? source : overlay.frame;
+	const auto  game_rect     = !Config::StretchEnabled() && aspect_source != nullptr
+	                               ? FitAspect(aspect_source->image.extent, m_extent)
+	                               : vk::Rect2D {{0, 0}, m_extent};
+	if (source == nullptr || game_rect.extent != m_extent) {
+		const vk::ClearColorValue black {};
+		vk_command.clearColorImage(m_images[m_image_index], vk::ImageLayout::eTransferDstOptimal,
+		                           &black, 1, &to_transfer.subresourceRange);
+		if (source != nullptr) {
+			vk::ImageMemoryBarrier after_clear = to_transfer;
+			after_clear.srcAccessMask          = vk::AccessFlagBits::eTransferWrite;
+			after_clear.oldLayout              = vk::ImageLayout::eTransferDstOptimal;
+			vk_command.pipelineBarrier(vk::PipelineStageFlagBits::eTransfer,
+			                           vk::PipelineStageFlagBits::eTransfer,
+			                           vk::DependencyFlags {}, 0, nullptr, 0, nullptr, 1,
+			                           &after_clear);
+		}
+	}
 	if (source != nullptr) {
 		vk::ImageBlit region {};
 		region.srcSubresource.aspectMask     = vk::ImageAspectFlagBits::eColor;
@@ -812,16 +849,14 @@ void Swapchain::RecordPresentCommands(CommandBuffer& command, Presenter::Frame*
 		region.dstSubresource.mipLevel       = 0;
 		region.dstSubresource.baseArrayLayer = 0;
 		region.dstSubresource.layerCount     = 1;
-		region.dstOffsets[1].x               = static_cast<int>(m_extent.width);
-		region.dstOffsets[1].y               = static_cast<int>(m_extent.height);
+		region.dstOffsets[0].x = game_rect.offset.x;
+		region.dstOffsets[0].y = game_rect.offset.y;
+		region.dstOffsets[1].x = game_rect.offset.x + static_cast<int>(game_rect.extent.width);
+		region.dstOffsets[1].y = game_rect.offset.y + static_cast<int>(game_rect.extent.height);
 		region.dstOffsets[1].z               = 1;
 		vk_command.blitImage(source->image.image, vk::ImageLayout::eTransferSrcOptimal,
 		                     m_images[m_image_index], vk::ImageLayout::eTransferDstOptimal, 1,
 		                     &region, vk::Filter::eLinear);
-	} else {
-		const vk::ClearColorValue black {};
-		vk_command.clearColorImage(m_images[m_image_index], vk::ImageLayout::eTransferDstOptimal,
-		                           &black, 1, &to_transfer.subresourceRange);
 	}
 
 	vk::ImageMemoryBarrier to_present {};
@@ -848,7 +883,7 @@ void Swapchain::RecordPresentCommands(CommandBuffer& command, Presenter::Frame*
 	                           &to_present);
 	if (draw_attachment) {
 		if (draw_overlay) {
-			DrawOverlay(vk_command, overlay);
+			DrawOverlay(vk_command, overlay, game_rect);
 			if (draw_system_overlay) {
 				to_present.srcAccessMask = vk::AccessFlagBits::eColorAttachmentWrite;
 				to_present.oldLayout     = vk::ImageLayout::eColorAttachmentOptimal;

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:
-					if ((lhs.kind == Kind::PackedControlPoint && rhs == constant(0xffu)) ||
-					    (rhs.kind == Kind::PackedControlPoint && lhs == constant(0xffu)))
-						value = constant(0u);
+				case Opcode::V_AND_B32: {
+					const auto mask = lhs.kind == Kind::PackedControlPoint ? rhs
+					                : rhs.kind == Kind::PackedControlPoint ? lhs
+					                                                       : Address {};
+					if (mask == constant(0xffu)) value = constant(0u);
+					if (mask == constant(0x1f00u)) value = {Kind::Affine, 256u, 0u};
 					break;
+				}
 				case Opcode::V_ADD_NC_U32: value = add(lhs, rhs); break;
 				case Opcode::V_MUL_LO_U32: value = multiply(lhs, rhs); break;
 				case Opcode::V_MUL_U32_U24: value = multiply(low24(lhs), low24(rhs)); break;

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;
-		EXIT_NOT_IMPLEMENTED(index >= 4u);
+		const bool quad        = tess.domain == 2u;
+		const auto outer_count = quad ? 4u : 3u;
+		const bool outer       = index < outer_count;
+		EXIT_NOT_IMPLEMENTED(index >= (quad ? 6u : 4u));
 		state.builder.AddFunction(spv::OpAccessChain, TypePointer(state, storage, TypeF32(state)),
 		                          pointer, outer ? variable : state.tess_inner_variable,
-		                          ConstantU32(state, outer ? index : index - 3u));
+		                          ConstantU32(state, outer ? index : index - outer_count));
 		return pointer;
 	}
 	auto address = ctx.Arg(inst, 1);
@@ -128,20 +130,25 @@ void DefineTessellationInterfaces(EmitterState& state) {
 
 void DefineTessellationExecutionModes(EmitterState& state) {
 	const auto& tess = state.input_info.vertex->tess;
-	// CxVgtTfParam: triangles, integer or fractional-odd partitioning, clockwise output.
-	EXIT_NOT_IMPLEMENTED(tess.domain != 1u ||
-	                     (tess.partitioning != 0u && tess.partitioning != 2u) ||
-	                     tess.output_topology != 2u);
+	// Triangle or quad domain, integer or fractional partitioning, triangle output.
+	EXIT_NOT_IMPLEMENTED((tess.domain != 1u && tess.domain != 2u) ||
+	                     (tess.partitioning != 0u && tess.partitioning != 2u && tess.partitioning != 3u) ||
+	                     (tess.output_topology != 2u && tess.output_topology != 3u));
 	state.builder.RequireCapability(spv::CapabilityTessellation);
 	if (state.program.stage == ShaderType::TessellationControl) {
 		state.builder.AddExecutionMode(state.main_func, spv::ExecutionModeOutputVertices,
 		                               tess.output_control_points);
 	} else {
-		state.builder.AddExecutionMode(state.main_func, spv::ExecutionModeTriangles);
+		state.builder.AddExecutionMode(state.main_func, tess.domain == 1u
+		                                                   ? spv::ExecutionModeTriangles
+		                                                   : spv::ExecutionModeQuads);
 		state.builder.AddExecutionMode(
 		    state.main_func, tess.partitioning == 0u ? spv::ExecutionModeSpacingEqual
-		                                             : spv::ExecutionModeSpacingFractionalOdd);
-		state.builder.AddExecutionMode(state.main_func, spv::ExecutionModeVertexOrderCw);
+		                     : tess.partitioning == 2u ? spv::ExecutionModeSpacingFractionalOdd
+		                                                : spv::ExecutionModeSpacingFractionalEven);
+		state.builder.AddExecutionMode(state.main_func, tess.output_topology == 2u
+		                                                   ? spv::ExecutionModeVertexOrderCw
+		                                                   : spv::ExecutionModeVertexOrderCcw);
 	}
 }
 

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 ||
-	       opcode == Opcode::V_MAD_I16;
+	return opcode == Opcode::V_MIN3_I16 || opcode == Opcode::V_MIN3_U16 ||
+	       opcode == Opcode::V_MAX3_I16 || opcode == Opcode::V_MAX3_U16 ||
+	       opcode == Opcode::V_MED3_I16 || opcode == Opcode::V_MED3_U16 ||
+	       opcode == Opcode::V_MAD_I16 || opcode == Opcode::V_MAD_U16;
 }
 
 bool IsNativeVop3B16BinaryOpcode(Opcode opcode) {
@@ -1335,6 +1350,7 @@ uint32_t NativeVop3SourceCount(Opcode opcode) {
 		case Opcode::V_SUB_NC_I16:
 		case Opcode::V_LSHLREV_B64:
 		case Opcode::V_LSHRREV_B64:
+		case Opcode::V_ASHRREV_I64:
 		case Opcode::V_LSHLREV_B16:
 		case Opcode::V_LSHRREV_B16:
 		case Opcode::V_ASHRREV_I16:
@@ -1523,7 +1539,10 @@ bool SupportsNativeVop3ResultModifiers(Opcode opcode) {
 
 bool SupportsNativeVop3Clamp(Opcode opcode) {
 	return SupportsNativeVop3ResultModifiers(opcode) || UsesInexactClampControl(opcode) ||
-	       opcode == Opcode::V_MAD_I16;
+	       opcode == Opcode::V_MAD_I16 || opcode == Opcode::V_MAD_U16 ||
+	       opcode == Opcode::V_MIN3_I16 || opcode == Opcode::V_MAX3_I16 ||
+	       opcode == Opcode::V_MAX3_U16 || opcode == Opcode::V_MED3_U16 ||
+	       opcode == Opcode::V_MUL_HI_I32_I24 || opcode == Opcode::V_MUL_HI_U32_U24;
 }
 
 bool HasUnsupportedNativeVop3Modifiers(Opcode opcode, bool permlane, bool mad_mix,
@@ -1554,7 +1573,17 @@ bool HasUnsupportedNativeVop3Modifiers(Opcode opcode, bool permlane, bool mad_mi
 		return clamp != 0u || omod != 0u || neg != 0u;
 	}
 	switch (opcode) {
+		case Opcode::V_MUL_HI_I32_I24:
+		case Opcode::V_MUL_HI_U32_U24:
+			// The high product fits in 32 bits, so clamp has no effect.
+			return abs != 0u || op_sel != 0u || omod != 0u || neg != 0u;
 		case Opcode::V_LDEXP_F32: return (abs & ~1u) != 0u || op_sel != 0u || (neg & ~1u) != 0u;
+		case Opcode::V_CVT_F32_UBYTE0:
+		case Opcode::V_CVT_F32_UBYTE1:
+		case Opcode::V_CVT_F32_UBYTE2:
+		case Opcode::V_CVT_F32_UBYTE3:
+			// Byte conversions ignore input modifiers and support all output modifiers.
+			return op_sel != 0u;
 		case Opcode::V_INTERP_P1LL_F16:
 		case Opcode::V_INTERP_P2_F16:
 			return abs != 0u || (op_sel & 0x7u) != 0u || clamp != 0u || omod != 0u || neg != 0u;

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) {
-	const auto lhs    = ReadU16AsU32(inst.src0, true);
-	const auto rhs    = ReadU16AsU32(inst.src1, true);
-	auto       result = ir.IAdd(ir.IMul(lhs, rhs), ReadU16AsU32(inst.src2, true));
+void Translator::Integer16Mad(const Decoder::Instruction& inst, bool sign) {
+	const auto lhs    = ReadU16AsU32(inst.src0, sign);
+	const auto rhs    = ReadU16AsU32(inst.src1, sign);
+	auto       result = ir.IAdd(ir.IMul(lhs, rhs), ReadU16AsU32(inst.src2, sign));
 	if (inst.dst.clamp) {
-		result = Saturate16(result, true);
+		result = Saturate16(result, sign);
 	}
 	Write16Bits(DestinationOperand(inst), result);
 }
@@ -147,6 +147,12 @@ void Translator::SimpleInteger(const Decoder::Instruction& inst, IR::ValueOpcode
 		const auto arg_type = IR::ArgTypeOf(opcode, index);
 		const auto& operand = SourceAt(inst, reverse && index < 2u ? 1u - index : index);
 		args[index]         = ReadOperand(operand, arg_type == IR::Type::Void ? type : arg_type);
+		if (opcode == IR::ValueOpcode::ShiftRightArithmetic64 && index == 0u &&
+		    operand.kind == Decoder::OperandKind::LiteralConstant) {
+			args[index] = ir.ConstructU64(
+			    ReadU32(operand),
+			    IR::U32(IR::Value((operand.value & 0x80000000u) ? 0xffffffffu : 0u)));
+		}
 		if (mask_shift_count && index == 1u) {
 			args[index] = ir.BitwiseAnd(IR::U32(args[index]), IR::U32(IR::Value(31u)));
 		}
@@ -263,7 +269,7 @@ void Translator::V_FFBH_32(const Decoder::Instruction& inst, bool sign) {
 	WriteOperand(DestinationOperand(inst), result);
 }
 
-void Translator::Integer24(const Decoder::Instruction& inst, bool sign, bool addend) {
+void Translator::Integer24(const Decoder::Instruction& inst, bool sign, bool addend, bool high) {
 	const auto extract24 = [&](IR::U32 value) {
 		return IR::U32(
 		    ir.Emit(sign ? IR::ValueOpcode::BitFieldSExtract : IR::ValueOpcode::BitFieldUExtract,
@@ -271,7 +277,9 @@ void Translator::Integer24(const Decoder::Instruction& inst, bool sign, bool add
 	};
 	const auto lhs    = extract24(ReadU32(inst.src0));
 	const auto rhs    = extract24(ReadU32(inst.src1));
-	auto       result = ir.IMul(lhs, rhs);
+	const auto multiply = high ? (sign ? IR::ValueOpcode::SMulHi : IR::ValueOpcode::UMulHi)
+	                           : IR::ValueOpcode::IMul32;
+	auto result = IR::U32(ir.Emit(multiply, {lhs, rhs}));
 	if (addend) {
 		result = ir.IAdd(result, ReadU32(inst.src2));
 	}
@@ -369,7 +377,7 @@ void Translator::S_BITREPLICATE_B64_B32(const Decoder::Instruction& inst) {
 	WriteOperand(DestinationOperand(inst), ir.ConstructU64(replicate(low), replicate(high)));
 }
 
-void Translator::S_QUADMASK_B64(const Decoder::Instruction& inst) {
+void Translator::S_QUADMASK(const Decoder::Instruction& inst, bool is64) {
 	const auto compact = [&](IR::U32 value) {
 		auto bits = ir.BitwiseOr(value, ir.ShiftRightLogical(value, IR::U32(IR::Value(1u))));
 		bits      = ir.BitwiseOr(bits, ir.ShiftRightLogical(bits, IR::U32(IR::Value(2u))));
@@ -382,12 +390,17 @@ void Translator::S_QUADMASK_B64(const Decoder::Instruction& inst) {
 		    ir.BitwiseOr(bits, ir.ShiftRightLogical(bits, IR::U32(IR::Value(12u)))),
 		    IR::U32(IR::Value(0xffu)));
 	};
+	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), result);
-	ir.SetScc(IR::U1(ir.Emit(IR::ValueOpcode::INotEqual64, {result, IR::Value(uint64_t {0})})));
+		WriteOperand(DestinationOperand(inst), ir.ConstructU64(quads, IR::U32(IR::Value(0u))));
+	} else {
+		quads = compact(ReadU32(inst.src0));
+		WriteOperand(DestinationOperand(inst), quads);
+	}
+	ir.SetScc(ir.INotEqual(quads, IR::U32(IR::Value(0u))));
 }
 
 void Translator::BFM_B32(const Decoder::Instruction& inst) {

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.fetch_embedded   = true;
+		info.has_dynamic_vertex_fetch = 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"
-	     "\t fetch_embedded = %s\n",
-	     info.fetch_external ? "true" : "false", info.fetch_embedded ? "true" : "false");
+	LOGF("\t has_dynamic_vertex_fetch = %s\n", info.has_dynamic_vertex_fetch ? "true" : "false");
 
 	for (int i = 0; i < info.resources_num; i++) {
 		LOGF("\t input %d\n", i);

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
-// slot is set.
+// to. A packet can leave both slots unused even when its record contains nonzero offsets.
 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;
-	bool                    fetch_embedded      = false;
+	// GPU-read records add Vulkan bases even when the packet patches no user SGPRs.
+	bool                    indirect_draw       = false;
+	// Attribute-table metadata enables dynamic fetch detection; loads are matched later.
+	bool                    has_dynamic_vertex_fetch = false;
 };
 
 struct ShaderComputeInputInfo: ShaderWorkgroupInputInfo {

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 << (key.startsWith('R') && key.size() > 2 && key.at(1).isDigit()
-		            ? key.remove('R').toLower()
-		            : key);
+		ret << QString(me.key(i));
 	}
 	return ret;
 }
@@ -32,16 +29,12 @@ inline QStringList EnumToList() {
 template <class T>
 T TextToEnum(const QString& text) {
 	auto me = QMetaEnum::fromType<T>();
-	return static_cast<T>(me.keyToValue(
-	    ((text.size() > 1 && text.at(0).isDigit()) ? 'R' + text.toUpper() : text).toUtf8().data()));
+	return static_cast<T>(me.keyToValue(text.toUtf8().constData()));
 }
 
 template <class T>
 QString EnumToText(T value) {
-	auto me  = QMetaEnum::fromType<T>();
-	auto key = QString(me.valueToKey(static_cast<int>(value)));
-	return (key.startsWith('R') && key.size() > 2 && key.at(1).isDigit() ? key.remove('R').toLower()
-	                                                                     : key);
+	return QString(QMetaEnum::fromType<T>().valueToKey(static_cast<int>(value)));
 }
 
 struct ControllerSettings {
@@ -80,13 +73,10 @@ public:
 	static constexpr int DEFAULT_CONSOLE_LANGUAGE = 1;
 	static constexpr int MAX_CONSOLE_LANGUAGE     = 29;
 
-	enum class Resolution {
-		R1280X720,
-		R1920X1080,
-		R2560X1440,
-		R3840X2160,
-	};
-	Q_ENUM(Resolution)
+	static constexpr int MIN_SCREEN_SIZE = 320;
+	static constexpr int MAX_SCREEN_SIZE = 16384;
+
+	static QString NormalizeResolution(const QString& text);
 
 	enum class ShaderOptimizationType { None, Size, Performance };
 	Q_ENUM(ShaderOptimizationType)
@@ -117,7 +107,7 @@ public:
 	// Controller preferences always come from the global configuration.
 	ControllerSettings controller;
 
-	Resolution             screen_resolution           = Resolution::R1280X720;
+	QString                screen_resolution           = QStringLiteral("1280x720");
 	QString                user_name                   = "Kyty";
 	int                    user_id                     = Config::DEFAULT_USER_ID;
 	QString                audio_input_device;
@@ -126,6 +116,7 @@ public:
 	bool                   fullscreen_enabled          = false;
 	bool                   hdr_enabled                 = false;
 	bool                   hide_cursor_enabled         = false;
+	bool                   stretch_enabled             = false;
 	bool                   readback_linear_images      = false;
 	bool                   readback_tiled_images       = false;
 	bool                   sync_raw_image_buffers      = false;
@@ -168,6 +159,7 @@ public:
 		fullscreen_enabled          = other.fullscreen_enabled;
 		hdr_enabled                 = other.hdr_enabled;
 		hide_cursor_enabled         = other.hide_cursor_enabled;
+		stretch_enabled             = other.stretch_enabled;
 		readback_linear_images      = other.readback_linear_images;
 		readback_tiled_images       = other.readback_tiled_images;
 		sync_raw_image_buffers      = other.sync_raw_image_buffers;

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),
-      m_info(info) {
+    : QDialog(parent, Qt::WindowMaximizeButtonHint | Qt::WindowCloseButtonHint),
+      m_ui(new Ui::ConfigurationEditDialog), m_info(info) {
 	m_ui->setupUi(this);
-	setMinimumWidth(width());
+	MakeTabsScrollable();
 	InitGameDirectories();
 	m_ui->controller_group->setVisible(false);
 	m_ui->settings_button->setVisible(false);
@@ -170,9 +172,11 @@ ConfigurationEditDialog::ConfigurationEditDialog(Configuration& info, QWidget* p
 
 	setSizeGripEnabled(true);
 
-	restoreGeometry(g_last_geometry);
-
 	Init(info);
+	m_geometry_restored = restoreGeometry(g_last_geometry);
+	if (!m_geometry_restored) {
+		GrowToFitContents();
+	}
 }
 
 QByteArray ConfigurationEditDialog::g_last_geometry;
@@ -248,7 +252,7 @@ void ConfigurationEditDialog::Init(const Configuration& info) {
 		                    info.audio_input_device);
 	}
 	microphone->setCurrentIndex(microphone->findData(info.audio_input_device));
-	ListInit(m_ui->comboBox_screen_resolution, info.screen_resolution);
+	m_ui->comboBox_screen_resolution->setCurrentText(info.screen_resolution);
 	ListInit(m_ui->comboBox_present_mode, info.present_mode);
 	m_ui->comboBox_gpu->clear();
 	m_ui->comboBox_gpu->addItem(tr("Auto"));
@@ -285,6 +289,7 @@ void ConfigurationEditDialog::Init(const Configuration& info) {
 	m_ui->checkBox_fullscreen->setChecked(info.fullscreen_enabled);
 	m_ui->checkBox_hdr->setChecked(info.hdr_enabled);
 	m_ui->checkBox_hide_cursor->setChecked(info.hide_cursor_enabled);
+	m_ui->checkBox_stretch->setChecked(info.stretch_enabled);
 	m_ui->checkBox_readback->setChecked(info.readback_linear_images);
 	m_ui->checkBox_readback_tiled->setChecked(info.readback_tiled_images);
 	m_ui->checkBox_sync_raw_image_buffers->setChecked(info.sync_raw_image_buffers);
@@ -413,8 +418,46 @@ void ConfigurationEditDialog::SetGlobalSettings(const QStringList& dirs, const Q
 	m_game_dirs_group->setVisible(true);
 	m_ui->controller_group->setVisible(true);
 	update_game_directory_buttons();
+	if (!m_geometry_restored) {
+		GrowToFitContents();
+	}
+}
+
+void ConfigurationEditDialog::MakeTabsScrollable() {
+	auto* tabs = m_ui->settings_tabs;
+	for (int i = 0; i < tabs->count(); i++) {
+		auto* page   = tabs->widget(i);
+		auto  title  = tabs->tabText(i);
+		auto* scroll = new QScrollArea;
+		scroll->setWidgetResizable(true);
+		scroll->setFrameShape(QFrame::NoFrame);
+		tabs->removeTab(i);
+		scroll->setWidget(page);
+		tabs->insertTab(i, scroll, title);
+	}
+	tabs->setCurrentIndex(0);
+}
+
+void ConfigurationEditDialog::GrowToFitContents() {
 	layout()->activate();
-	resize(size().expandedTo(minimumSizeHint()));
+	auto* tabs = m_ui->settings_tabs;
+	QSize content;
+	for (int i = 0; i < tabs->count(); i++) {
+		content = content.expandedTo(
+		    static_cast<QScrollArea*>(tabs->widget(i))->widget()->minimumSizeHint());
+	}
+	const QSize chrome = size() - tabs->currentWidget()->size();
+	resize(size().expandedTo(content + chrome).boundedTo(screen()->availableGeometry().size()));
+}
+
+void ConfigurationEditDialog::ShowTab(QWidget* page) {
+	auto* tabs = m_ui->settings_tabs;
+	for (int i = 0; i < tabs->count(); i++) {
+		if (tabs->widget(i)->isAncestorOf(page)) {
+			tabs->setCurrentIndex(i);
+			return;
+		}
+	}
 }
 
 QString ConfigurationEditDialog::GetAddonsDirectory() const {
@@ -474,13 +517,14 @@ static void UpdateInfo(Configuration& info, Ui::ConfigurationEditDialog& ui, boo
 		info.controller.keyboard_player     = ui.comboBox_keyboard_player->currentData().toInt();
 	}
 	info.screen_resolution =
-	    TextToEnum<Configuration::Resolution>(ui.comboBox_screen_resolution->currentText());
+	    Configuration::NormalizeResolution(ui.comboBox_screen_resolution->currentText());
 	info.present_mode =
 	    TextToEnum<Configuration::PresentMode>(ui.comboBox_present_mode->currentText());
 	info.gpu_index                 = ui.comboBox_gpu->currentIndex() - 1;
 	info.fullscreen_enabled        = ui.checkBox_fullscreen->isChecked();
 	info.hdr_enabled               = ui.checkBox_hdr->isChecked();
 	info.hide_cursor_enabled       = ui.checkBox_hide_cursor->isChecked();
+	info.stretch_enabled           = ui.checkBox_stretch->isChecked();
 	info.readback_linear_images    = ui.checkBox_readback->isChecked();
 	info.readback_tiled_images     = ui.checkBox_readback_tiled->isChecked();
 	info.sync_raw_image_buffers    = ui.checkBox_sync_raw_image_buffers->isChecked();
@@ -518,7 +562,7 @@ void ConfigurationEditDialog::save() {
 	for (auto* field: {m_ui->lineEdit_shader_log_folder, m_ui->lineEdit_printf_file,
 	                   m_ui->lineEdit_cmd_dump_folder}) {
 		if (field->isEnabled() && field->text().isEmpty()) {
-			m_ui->settings_tabs->setCurrentWidget(m_ui->debug_tab);
+			ShowTab(m_ui->debug_tab);
 			field->setFocus();
 			break;
 		}
@@ -540,6 +584,16 @@ void ConfigurationEditDialog::save() {
 		                      tr("User ID cannot be 254 (everyone) or 255 (system)"));
 		return;
 	}
+	if (Configuration::NormalizeResolution(m_ui->comboBox_screen_resolution->currentText())
+	        .isEmpty()) {
+		ShowTab(m_ui->settings_tab);
+		m_ui->comboBox_screen_resolution->setFocus();
+		QMessageBox::critical(this, tr("Save failed"),
+		                      tr("Screen resolution must be WIDTHxHEIGHT, each between %1 and %2")
+		                          .arg(Configuration::MIN_SCREEN_SIZE)
+		                          .arg(Configuration::MAX_SCREEN_SIZE));
+		return;
+	}
 
 	UpdateInfo(m_info, *m_ui, m_global_settings);
 

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/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: {
-			const auto configured_size =
-			    voice->rack->option.custom_mastering.custom_rack_option.state_size;
-			EXIT_NOT_IMPLEMENTED(configured_size < sizeof(Ngs2CustomMasteringVoiceState));
-			EXIT_NOT_IMPLEMENTED(state_size != configured_size);
-			std::memset(state, 0, state_size);
-			auto* mastering = reinterpret_cast<Ngs2CustomMasteringVoiceState*>(state);
-			mastering->voice_state.state_flags = Ngs2GetStateFlags(voice);
-			break;
-		}
+		case Ngs2RackType::CustomMastering:
 		case Ngs2RackType::CustomSampler: {
+			const bool sampler = voice->rack->type == Ngs2RackType::CustomSampler;
 			const auto configured_size =
 			    voice->rack->option.custom_sampler.custom_rack_option.state_size;
-			EXIT_NOT_IMPLEMENTED(configured_size < sizeof(Ngs2CustomSamplerVoiceState));
-			EXIT_NOT_IMPLEMENTED(state_size != configured_size);
-			std::memset(state, 0, state_size);
-			auto* sampler = reinterpret_cast<Ngs2CustomSamplerVoiceState*>(state);
-			Ngs2FillSamplerState(*voice, *sampler);
+			EXIT_NOT_IMPLEMENTED(configured_size < (sampler ? sizeof(Ngs2CustomSamplerVoiceState)
+			                                                : sizeof(Ngs2CustomMasteringVoiceState)));
+			if (state_size < configured_size) {
+				return static_cast<int32_t>(0x804a8011u);
+			}
+			std::memset(state, 0, configured_size);
+			if (sampler) {
+				Ngs2FillSamplerState(*voice, *reinterpret_cast<Ngs2CustomSamplerVoiceState*>(state));
+			} else {
+				state->state_flags = Ngs2GetStateFlags(voice);
+			}
 			break;
 		}
 		case Ngs2RackType::Sampler: {

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

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