image: drop optional create flags the driver rejects instead of aborting (#1760)

* image: request extended usage only for storage images

* tests: restore block-texel-view support through the flag check loop order

---------

Co-authored-by: nmzik <Nmzik@mail.ru>

.gitignore

@@ -6,3 +6,4 @@ _Build/vscode-clang/
 _Build/
 cmake-build-debug/
 _PipelineCache/
+CMakeUserPresets.json

.vscode/launch.json

@@ -27,7 +27,7 @@
             "args": [
                 "--game", "D:/gamesps5/PPSA07429",
                 "--printf-direction", "File",
-                "--graphics-debug-dump", "True"
+                "--graphics-debug-dump"
             ],
             "stopAtEntry": false,
             "cwd": "${workspaceFolder}\\_Build\\vscode-clang",

CMakeLists.txt

@@ -657,6 +657,8 @@ if(BUILD_TESTING)
 	add_test(NAME shader_recompiler_compute COMMAND $<TARGET_FILE:shader_recompiler_compute_tests>)
 	add_test(NAME shader_recompiler_tessellation
 		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --tessellation-only)
+	add_test(NAME shader_recompiler_float_min_max
+		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --float-min-max-only)
 	add_test(NAME shader_recompiler_alignbyte
 		COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --alignbyte-only)
 	add_test(NAME graphics_pipeline_rasterization

docs/logo.svg

@@ -0,0 +1,15 @@
+<svg xmlns="http://www.w3.org/2000/svg" viewBox="0 0 512 512"><defs>
+<linearGradient id="body" x1="0" y1="0" x2="0" y2="1"><stop offset="0" stop-color="#ffffff"/><stop offset="1" stop-color="#e3e9f2"/></linearGradient>
+</defs>
+<g transform="translate(256 268) scale(1.16) translate(-256 -268)" stroke-linejoin="round" stroke-linecap="round">
+<path d="M100 196 L120 104 Q124 92 134 99 L214 160 Z" fill="url(#body)"/><path d="M412 196 L392 104 Q388 92 378 99 L298 160 Z" fill="url(#body)"/>
+<path d="M128 164 L137 128 L174 156 Z" fill="#0b0d12"/><path d="M384 164 L375 128 L338 156 Z" fill="#0b0d12"/>
+<g fill="none" stroke="#0b0d12" stroke-width="16"><path d="M100 196 L120 104 Q124 92 134 99 L214 160 Z"/><path d="M412 196 L392 104 Q388 92 378 99 L298 160 Z"/></g>
+<path d="M256 154 C330 154 380 160 406 178 C432 196 447 244 459 302 C471 358 466 398 430 408 C400 416 380 398 360 368 L326 318 C316 304 306 300 290 300 L222 300 C206 300 196 304 186 318 L152 368 C132 398 112 416 82 408 C46 398 41 358 53 302 C65 244 80 196 106 178 C132 160 182 154 256 154 Z" fill="url(#body)"/>
+<path d="M256 154 C330 154 380 160 406 178 C432 196 447 244 459 302 C471 358 466 398 430 408 C400 416 380 398 360 368 L326 318 C316 304 306 300 290 300 L222 300 C206 300 196 304 186 318 L152 368 C132 398 112 416 82 408 C46 398 41 358 53 302 C65 244 80 196 106 178 C132 160 182 154 256 154 Z" fill="none" stroke="#0b0d12" stroke-width="16"/>
+<g fill="none" stroke="#0b0d12" stroke-width="9"><path d="M82 260 C100 274 118 288 128 304 C134 316 121 350 108 388"/><path d="M430 260 C412 274 394 288 384 304 C378 316 391 350 404 388"/></g>
+<path d="M64 322 L82 340 M62 340 L76 354 M448 322 L430 340 M450 340 L436 354" fill="none" stroke="#0b0d12" stroke-width="7"/>
+<path d="M168 230 Q194 202 220 230" fill="none" stroke="#0b0d12" stroke-width="15"/>
+<circle cx="322" cy="220" r="22" fill="#0b0d12"/>
+<path d="M238 252 L274 252 L256 276 Z" fill="#0b0d12" stroke="#0b0d12" stroke-width="8"/>
+</g></svg>

src/common/archive.h

@@ -5,6 +5,7 @@
 
 #include <cstdint>
 #include <filesystem>
+#include <functional>
 #include <memory>
 #include <vector>
 
@@ -45,6 +46,12 @@ private:
 // Accepts either a host archive filename or a virtual path inside the archive.
 [[nodiscard]] std::shared_ptr<ArchiveReader> OpenArchive(const std::filesystem::path& path);
 
+// Packs a host directory into a new ZArchive outside it. progress gets 0 < done <= total and
+// returns false to cancel. A failed or cancelled pack removes the partial archive.
+[[nodiscard]] bool PackArchive(const std::filesystem::path& directory,
+                               const std::filesystem::path& archive,
+                               const std::function<bool(uint64_t done, uint64_t total)>& progress);
+
 } // namespace Common
 
 #endif /* KYTY_COMMON_ARCHIVE_H_ */

src/common/archivePack.cpp

@@ -0,0 +1,84 @@
+// Kept out of archive.cpp so the statically linked emulator does not pull in the writer and the
+// zstd compressor.
+#include "common/archive.h"
+
+#include "common/stringUtils.h"
+
+#include <algorithm>
+#include <zarchive/zarchivewriter.h>
+
+namespace Common {
+
+bool PackArchive(const std::filesystem::path& directory, const std::filesystem::path& archive,
+                 const std::function<bool(uint64_t done, uint64_t total)>& progress) {
+	namespace fs = std::filesystem;
+	struct Output {
+		File file;
+		bool ok = true;
+	} output;
+	// Create the archive before walking, so a re-run started while a cancelled pack is still
+	// walking sees it exists instead of losing its archive to that pack's cleanup.
+	if (!output.file.Create(archive)) {
+		return false;
+	}
+	struct Entry {
+		fs::path path;
+		uint64_t size;
+		bool     is_file;
+	};
+	std::vector<Entry> entries;
+	uint64_t           total = 0;
+	std::error_code    error;
+	// Follow directory links, as the emulator does when it runs the folder.
+	fs::recursive_directory_iterator it(directory, fs::directory_options::follow_directory_symlink,
+	                                    error);
+	for (; !error && it != fs::end(it); it.increment(error)) {
+		const bool is_file = !it->is_directory(error);
+		const auto size    = is_file && !error ? it->file_size(error) : 0;
+		if (error) {
+			break;
+		}
+		entries.push_back({it->path(), size, is_file});
+		total += size;
+	}
+	output.ok = !error;
+	ZArchiveWriter writer(
+	    [](int32_t, void*) {},
+	    [](const void* data, size_t size, void* opaque) {
+		    auto&    out     = *static_cast<Output*>(opaque);
+		    uint32_t written = 0;
+		    out.file.Write(data, static_cast<uint32_t>(size), &written);
+		    out.ok = out.ok && written == size;
+	    },
+	    &output);
+	std::vector<uint8_t> buffer(1u << 20u);
+	uint64_t             done = 0;
+	for (size_t index = 0; output.ok && index < entries.size(); index++) {
+		const auto& entry = entries[index];
+		const auto  name  = PathToGenericString(entry.path.lexically_relative(directory));
+		File        input;
+		output.ok = entry.is_file ? writer.StartNewFile(name.c_str()) &&
+		                                input.Open(entry.path, File::Mode::Read)
+		                          : writer.MakeDir(name.c_str());
+		for (uint64_t left = entry.size; output.ok && left != 0;) {
+			uint32_t read = 0;
+			input.Read(buffer.data(), static_cast<uint32_t>(std::min<uint64_t>(left, buffer.size())),
+			           &read);
+			writer.AppendData(buffer.data(), read);
+			done += read;
+			left -= read;
+			output.ok = output.ok && read != 0 && progress(done, total);
+		}
+	}
+	if (output.ok) {
+		writer.Finalize();
+	}
+	output.ok = output.ok && output.file.Flush();
+	output.file.Close();
+	if (!output.ok) {
+		File::DeleteFile(archive);
+	}
+	return output.ok;
+}
+
+} // namespace Common

src/common/emulatorConfig.cpp

@@ -185,6 +185,10 @@ bool ReadbackLinearImagesEnabled() {
 	return g_config->readback_linear_images;
 }
 
+bool ReadbackTiledImagesEnabled() {
+	return g_config->readback_tiled_images;
+}
+
 bool SyncRawImageBuffersEnabled() {
 	return g_config->sync_raw_image_buffers;
 }

src/common/emulatorConfig.h

@@ -85,6 +85,7 @@ struct ConfigOptions {
 	bool                   gpu_assisted_validation_enabled = false;
 	bool                   renderdoc_enabled           = false;
 	bool                   readback_linear_images      = false;
+	bool                   readback_tiled_images       = false;
 	bool                   sync_raw_image_buffers      = false;
 	bool                   tessellation_enabled        = false;
 	bool                   trophy_enabled              = true;
@@ -150,6 +151,7 @@ bool GpuAssistedValidationEnabled();
 
 bool RenderDocEnabled();
 bool ReadbackLinearImagesEnabled();
+bool ReadbackTiledImagesEnabled();
 bool SyncRawImageBuffersEnabled();
 bool TessellationEnabled();
 bool TrophyEnabled();

src/common/logging/log.h

@@ -51,20 +51,26 @@ inline constexpr auto BrightWhite   = fmt::fg(fmt::terminal_color::bright_white)
 
 } // namespace Color
 
+// Out of line, so callers on small guest stacks (fibers) do not reserve the formatting frame.
+template <typename... Args>
+[[gnu::noinline]] void Printf(fmt::text_style style, const char* format, const Args&... args) {
+	Write(style, fmt::sprintf(format, args...));
+}
+
 } // namespace Log
 
 // NOLINTNEXTLINE(cppcoreguidelines-macro-usage)
 #define LOGF(...)                                                                                  \
 	do {                                                                                           \
 		if (!::Log::IsSilent()) {                                                                  \
-			::Log::Write(::fmt::sprintf(__VA_ARGS__));                                             \
+			::Log::Printf({}, __VA_ARGS__);                                                        \
 		}                                                                                          \
 	} while (false)
 // NOLINTNEXTLINE(cppcoreguidelines-macro-usage)
 #define LOGF_COLOR(style, ...)                                                                     \
 	do {                                                                                           \
 		if (!::Log::IsSilent()) {                                                                  \
-			::Log::Write((style), ::fmt::sprintf(__VA_ARGS__));                                    \
+			::Log::Printf((style), __VA_ARGS__);                                                   \
 		}                                                                                          \
 	} while (false)
 

src/graphics/guest_gpu/command_processor/commandProcessor.h

@@ -104,7 +104,6 @@ public:
 	                       uint32_t value);
 	void PrepareCpuFlip(uint64_t request_id);
 	void SynchronizeGpu();
-	void EmitGlobalBarrier();
 	void TriggerEopEventAtEndOfPipe(uint32_t interrupt_context_id);
 	void DispatchDirect(uint32_t thread_group_x, uint32_t thread_group_y, uint32_t thread_group_z,
 	                    uint32_t mode);

src/graphics/guest_gpu/command_processor/pm4Dispatch.cpp

@@ -73,7 +73,6 @@ constexpr auto MakeContextDispatchTable() {
 	g_hw_ctx_func[Pm4::PA_SC_BINNER_CNTL_0]                   = HwCtxSetPaScExtendedControl;
 	g_hw_ctx_func[Pm4::PA_SC_BINNER_CNTL_1]                   = HwCtxSetPaScExtendedControl;
 	g_hw_ctx_func[Pm4::PA_SC_CONSERVATIVE_RASTERIZATION_CNTL] = HwCtxSetPaScExtendedControl;
-	g_hw_ctx_func[Pm4::DB_ALPHA_TO_MASK]                      = HwCtxSetAlphaToMask;
 	g_hw_ctx_func[Pm4::VGT_DRAW_PAYLOAD_CNTL]                 = HwCtxSetDrawPayloadControl;
 	g_hw_ctx_func[Pm4::VGT_PRIMITIVEID_RESET]                 = HwCtxSetPrimitiveIdReset;
 	g_hw_ctx_func[Pm4::PA_CL_OBJPRIM_ID_CNTL]                 = HwCtxSetObjprimIdControl;

src/graphics/guest_gpu/command_processor/pm4Dispatch.h

@@ -79,7 +79,6 @@ uint32_t HwCtxSetAaMask(CommandProcessor&, uint32_t, uint32_t, const uint32_t*,
 uint32_t HwCtxSetVtxControl(CommandProcessor&, uint32_t, uint32_t, const uint32_t*, uint32_t);
 uint32_t HwCtxSetPaScExtendedControl(CommandProcessor&, uint32_t, uint32_t, const uint32_t*,
                                      uint32_t);
-uint32_t HwCtxSetAlphaToMask(CommandProcessor&, uint32_t, uint32_t, const uint32_t*, uint32_t);
 uint32_t HwCtxSetDrawPayloadControl(CommandProcessor&, uint32_t, uint32_t, const uint32_t*,
                                     uint32_t);
 uint32_t HwCtxSetPrimitiveIdReset(CommandProcessor&, uint32_t, uint32_t, const uint32_t*, uint32_t);

src/graphics/guest_gpu/command_processor/pm4Handlers.cpp

@@ -71,12 +71,6 @@ constexpr uint32_t NormalizeRegisterOffset(uint32_t raw_offset) {
 	return raw_offset & ~RegisterSelectorMask;
 }
 
-bool ReleaseMemGcrNeedsBarrier(uint32_t eop_event_type, uint32_t gcr_cntl) {
-	return eop_event_type != 0x28u ||
-	       (gcr_cntl & (GcrGl2MetadataInvalidate | GcrGl0VectorInvalidate | GcrGl1Invalidate |
-	                    GcrGl2Invalidate | GcrGl2Writeback)) != 0;
-}
-
 uint32_t ReleaseMemCacheActionFromGcr(uint32_t gcr_cntl) {
 	return ((gcr_cntl & GcrGl2Writeback) != 0 ? 0x38u : 0x00u);
 }
@@ -289,8 +283,6 @@ KYTY_HW_CTX_PARSER(HwCtxSetCentroidPriority) {
 static void HwCtxIgnoreAaMaskRegister([[maybe_unused]] uint32_t cmd_offset,
                                       [[maybe_unused]] uint32_t value) {}
 
-static void HwCtxIgnoreAlphaToMaskRegister([[maybe_unused]] uint32_t value) {}
-
 static void HwCtxIgnorePrimitiveIdReset([[maybe_unused]] uint32_t value) {}
 
 static void HwCtxIgnoreVtxControl([[maybe_unused]] uint32_t value) {}
@@ -435,16 +427,6 @@ KYTY_HW_CTX_PARSER(HwCtxSetPointState) {
 	return num_values;
 }
 
-KYTY_HW_CTX_PARSER(HwCtxSetAlphaToMask) {
-	auto num_values = KYTY_PM4_LEN(cmd_id) - 2u;
-
-	for (uint32_t i = 0; i < num_values; i++) {
-		HwCtxIgnoreAlphaToMaskRegister(buffer[i]);
-	}
-
-	return num_values;
-}
-
 KYTY_HW_CTX_PARSER(HwCtxSetDrawPayloadControl) {
 	cp.GetCtx().SetDrawPayloadControl(buffer[0]);
 	return KYTY_PM4_LEN(cmd_id) - 2u;
@@ -2259,30 +2241,14 @@ KYTY_CP_OP_PARSER(CpOpReleaseMem) {
 		}
 	};
 
-	if (data_sel == 0 || interrupt_selector == 4) {
-		if (eop_event_type != 0x28 || gcr_cntl != 0) {
-			cp.EmitGlobalBarrier();
-		}
-
+	// Cache actions only order GPU work, which the renderer's resource barriers already order.
+	if (data_sel == 0 || interrupt_selector == 4 ||
+	    (release_dst == ReleaseMemDstMemory && dst_gpu_addr == nullptr)) {
 		trigger_interrupt();
 
 		return 7;
 	}
 
-	if (release_dst == ReleaseMemDstMemory && dst_gpu_addr == nullptr) {
-		if (eop_event_type != 0x28 || gcr_cntl != 0) {
-			cp.EmitGlobalBarrier();
-		}
-
-		trigger_interrupt();
-
-		return 7;
-	}
-
-	if (ReleaseMemGcrNeedsBarrier(eop_event_type, gcr_cntl)) {
-		cp.EmitGlobalBarrier();
-	}
-
 	auto cache_action = ReleaseMemCacheActionFromGcr(gcr_cntl);
 	cache_policy      = 0;
 
@@ -3487,7 +3453,8 @@ void GraphicsInitJmpTablesCxIndirect() {
 	g_hw_ctx_indirect_func[Pm4::PA_SC_CONSERVATIVE_RASTERIZATION_CNTL] =
 	    [](KYTY_HW_CTX_INDIRECT_ARGS) { HwCtxIgnorePaScExtendedControl(cmd_offset, value); };
 	g_hw_ctx_indirect_func[Pm4::DB_ALPHA_TO_MASK] = [](KYTY_HW_CTX_INDIRECT_ARGS) {
-		HwCtxIgnoreAlphaToMaskRegister(value);
+		cp.GetCtx().SetAlphaToMaskEnable(
+		    KYTY_PM4_GET(value, DB_ALPHA_TO_MASK, ALPHA_TO_MASK_ENABLE) != 0);
 	};
 	g_hw_ctx_indirect_func[Pm4::CB_COLOR_CONTROL] = [](KYTY_HW_CTX_INDIRECT_ARGS) {
 		cp.GetCtx().SetColorControl(DecodeColorControl(value));

src/graphics/guest_gpu/graphicsRun.cpp

@@ -607,7 +607,6 @@ bool GuestGpu::Process(Submission& submission) {
 		case SubmissionType::SuspendPoint:
 			// Age caches at the frame boundary, not once per command list.
 			m_renderer.RunGarbageCollector();
-			cp.EmitGlobalBarrier();
 			m_renderer.GetCommandScheduler().DeferPriorityOperation(
 			    [ready = m_suspend_point_ready] { ready->release(); });
 			cp.BufferFlush();
@@ -1222,22 +1221,6 @@ void CommandProcessor::WriteAtEndOfPipe64(uint32_t cache_policy, uint32_t event_
 	                 interrupt_context_id);
 }
 
-void CommandProcessor::EmitGlobalBarrier() {
-	Common::LockGuard lock(m_renderer.GetMutex());
-
-	vk::MemoryBarrier2 barrier {};
-	barrier.srcStageMask  = vk::PipelineStageFlagBits2::eAllCommands;
-	barrier.srcAccessMask = vk::AccessFlagBits2::eMemoryWrite;
-	barrier.dstStageMask  = vk::PipelineStageFlagBits2::eAllCommands;
-	barrier.dstAccessMask = vk::AccessFlagBits2::eMemoryRead | vk::AccessFlagBits2::eMemoryWrite;
-
-	vk::DependencyInfo dependency {};
-	dependency.memoryBarrierCount = 1;
-	dependency.pMemoryBarriers    = &barrier;
-	GetScheduler().EndRendering();
-	CurrentBuffer().Handle().pipelineBarrier2(dependency);
-}
-
 void CommandProcessor::TriggerEopEventAtEndOfPipe(uint32_t interrupt_context_id) {
 	(void)CurrentBuffer().Handle();
 	GetScheduler().DeferPriorityOperation(
@@ -1257,21 +1240,17 @@ void CommandProcessor::TriggerEvent(uint32_t event_type, uint32_t event_index,
 	}
 
 	const auto valid_cache_event_index = event_index == 0x00000000 || event_index == 0x00000007;
+	// Partial flushes and cache writebacks only order GPU work, which the renderer's resource
+	// barriers already order.
 	switch (event_type) {
 		// CsPartialFlush, GsPartialFlush, PsPartialFlush.
 		case 0x00000007:
 		case 0x0000000f:
-		case 0x00000010: EmitGlobalBarrier(); break;
-		// CbDbDataWritebackInvalidate, CbDataWritebackInvalidate.
+		case 0x00000010: break;
+		// CbDbDataWritebackInvalidate, CbDataWritebackInvalidate, DbDataWritebackInvalidate,
+		// DbMetadataWritebackInvalidate, CbMetadataWritebackInvalidate.
 		case 0x00000016:
 		case 0x00000031:
-			if (!valid_cache_event_index) {
-				EXIT("unknown event type: 0x%08" PRIx32 ", 0x%08" PRIx32 "\n", event_type,
-				     event_index);
-			}
-			EmitGlobalBarrier();
-			break;
-		// DbDataWritebackInvalidate, DbMetadataWritebackInvalidate, CbMetadataWritebackInvalidate.
 		case 0x0000002a:
 		case 0x0000002c:
 		case 0x0000002e:
@@ -1279,7 +1258,6 @@ void CommandProcessor::TriggerEvent(uint32_t event_type, uint32_t event_index,
 				EXIT("unknown event type: 0x%08" PRIx32 ", 0x%08" PRIx32 "\n", event_type,
 				     event_index);
 			}
-			EmitGlobalBarrier();
 			break;
 		case 0x0000000d:
 		case 0x0000000e:

src/graphics/guest_gpu/hardwareContext.h

@@ -576,6 +576,12 @@ struct ShaderRegisters {
 	uint32_t m_paScShaderControl = 0;
 
 	DepthShaderControl db_shader_control;
+	bool               alpha_to_mask_enable = false;
+
+	// DB_SHADER_CONTROL.ALPHA_TO_MASK_DISABLE overrides DB_ALPHA_TO_MASK per shader.
+	[[nodiscard]] bool AlphaToMask() const {
+		return alpha_to_mask_enable && !db_shader_control.alpha_to_mask_disable;
+	}
 
 	[[nodiscard]] uint32_t GetExportCount() const {
 		return 1u + ((m_spiVsOutConfig >> 1u) & 0x1Fu);
@@ -915,6 +921,7 @@ public:
 	void SetDepthShaderControl(const DepthShaderControl& value) {
 		m_sh_regs.db_shader_control = value;
 	}
+	void SetAlphaToMaskEnable(bool value) { m_sh_regs.alpha_to_mask_enable = value; }
 
 	void SetScShaderControl(uint32_t value) { m_sh_regs.m_paScShaderControl = value; }
 

src/graphics/guest_gpu/pm4.h

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

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

@@ -256,6 +256,10 @@ BufferCache::BufferCache(GraphicContext& graphics, CommandScheduler& scheduler,
 	m_gds_buffer.Flush(0, m_gds_buffer.Size());
 	SetVulkanObjectNameF(m_graphics.device, m_bda_pagetable_buffer.Handle(),
 	                     "BDA Page Table Buffer");
+	m_scheduler.BeginCommand();
+	for (auto* buffer: {&m_bda_pagetable_buffer, m_fault_manager.GetFaultBuffer()}) {
+		buffer->Fill(0, buffer->Size(), 0);
+	}
 	const auto null_id =
 	    m_slot_buffers.insert(m_graphics, m_scheduler, MemoryUsage::DeviceLocal, 0, AllFlags, 16);
 	EXIT_IF(null_id != NULL_BUFFER_ID);

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

@@ -34,6 +34,8 @@ namespace Libs::Graphics {
 namespace {
 
 constexpr uint64_t NumFramesBeforeRemoval = 32;
+// Largest GPU-written tiled image (in texels) that is written back to guest memory on submit.
+constexpr uint64_t SmallTiledReadbackTexels = 64 * 64;
 
 [[nodiscard]] bool DecodeColorClear(const TextureCache::ImageDesc& desc, uint8_t code,
                                   vk::ClearColorValue& clear) {
@@ -122,7 +124,8 @@ TextureCache::TextureCache(GraphicContext& graphics, CommandScheduler& scheduler
       m_blit_helper(graphics, scheduler),
       m_tiler(graphics, scheduler, buffer_cache.GetUtilityBuffer(MemoryUsage::Stream)),
       m_buffer_cache(buffer_cache),
-      m_readback_linear_images(Config::ReadbackLinearImagesEnabled()) {
+      m_readback_linear_images(Config::ReadbackLinearImagesEnabled()),
+      m_readback_tiled_images(Config::ReadbackTiledImagesEnabled()) {
 	if (m_graphics.CanReportMemoryUsage()) {
 		constexpr int64_t GiB = 1024ll * 1024 * 1024;
 		const auto        budget =
@@ -401,7 +404,14 @@ void TextureCache::UntrackImageTail(ImageId id) {
 }
 
 void TextureCache::TrackImageDownload(ImageId id, Image& image) {
-	if (m_readback_linear_images && !image.info.IsTiled()) {
+	// Small tiled targets are written back too: games read GPU results such as reductions and
+	// simulation state from them on the CPU (PPSA04263 keeps the origin of its ocean wave grid
+	// in a 1x1 R32G32 target, and its water physics read it to sample the wave heights).
+	const auto& info   = image.info;
+	const auto  texels = static_cast<uint64_t>(info.extent.width) * info.extent.height *
+	                     info.extent.depth * info.resources.layers;
+	if (info.IsTiled() ? m_readback_tiled_images && texels <= SmallTiledReadbackTexels
+	                   : m_readback_linear_images) {
 		if (!image.IsGpuModified()) {
 			EXIT("TextureCache: cannot enroll a non-GPU-owned image for download\n");
 		}
@@ -1956,7 +1966,7 @@ void TextureCache::RunGarbageCollector() {
 }
 
 void TextureCache::ProcessDownloadImages() {
-	if (!m_readback_linear_images) {
+	if (!m_readback_linear_images && !m_readback_tiled_images) {
 		return;
 	}
 	std::scoped_lock lock {m_lock};

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

@@ -180,6 +180,7 @@ private:
 	uint64_t         m_gc_tick                = 0;
 	mutable uint32_t m_image_query_epoch      = 0;
 	bool             m_readback_linear_images = false;
+	bool             m_readback_tiled_images  = false;
 
 	friend struct TextureCacheTestAccess;
 	friend class BufferCache;

src/graphics/host_gpu/renderer/commandScheduler.cpp

@@ -327,7 +327,7 @@ void CommandScheduler::CheckActive() const {
 }
 
 CommandBuffer& CommandScheduler::Current() {
-	CheckActive();
+	EXIT_IF(m_command.IsInvalid());
 	return m_command;
 }
 

src/graphics/host_gpu/renderer/context.cpp

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

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

@@ -30,16 +30,18 @@ namespace {
 	}
 }
 
-[[nodiscard]] vk::ImageCreateFlags ImageCreateFlags(const GraphicContext& graphics,
-                                                   const ImageInfo& info) {
+[[nodiscard]] vk::ImageCreateFlags
+ImageCreateFlags(const GraphicContext& graphics, const ImageInfo& info, vk::ImageUsageFlags usage) {
 	vk::ImageCreateFlags flags {};
 	if (DepthAspectTransferFormat(info.pixel_format) == vk::Format::eUndefined) {
 		flags |= vk::ImageCreateFlagBits::eMutableFormat;
-		flags |= vk::ImageCreateFlagBits::eExtendedUsage;
 		if (info.IsBlock() && graphics.supports_block_texel_view) {
 			flags |= vk::ImageCreateFlagBits::eBlockTexelViewCompatible;
 		}
 	}
+	if (usage & vk::ImageUsageFlagBits::eStorage) {
+		flags |= vk::ImageCreateFlagBits::eExtendedUsage;
+	}
 	if (info.IsVolume()) {
 		flags |= vk::ImageCreateFlagBits::e2DArrayCompatible;
 	}
@@ -59,7 +61,7 @@ namespace {
 			const auto storage = usage | vk::ImageUsageFlagBits::eStorage;
 			if (graphics.GetImageFormatProperties(info.pixel_format, HostImageType(info.type),
 			                                      vk::ImageTiling::eOptimal, storage,
-			                                      ImageCreateFlags(graphics, info),
+			                                      ImageCreateFlags(graphics, info, storage),
 			                                      nullptr) == vk::Result::eSuccess) {
 				usage = storage;
 			} else {
@@ -702,7 +704,8 @@ Image::Image(GraphicContext& graphics, CommandScheduler& scheduler, const ImageI
 	}
 
 	vk::ImageCreateInfo create {};
-	create.flags         = ImageCreateFlags(graphics, info);
+	create.usage         = ImageUsageFlags(graphics, info);
+	create.flags         = ImageCreateFlags(graphics, info, create.usage);
 	create.imageType     = HostImageType(info.type);
 	create.extent        = info.extent;
 	create.mipLevels     = info.resources.levels;
@@ -710,7 +713,6 @@ Image::Image(GraphicContext& graphics, CommandScheduler& scheduler, const ImageI
 	create.format        = info.pixel_format;
 	create.tiling        = vk::ImageTiling::eOptimal;
 	create.initialLayout = vk::ImageLayout::eUndefined;
-	create.usage         = ImageUsageFlags(graphics, info);
 	create.samples       = vulkan_sample_count(info.samples);
 
 	vk::ImageFormatProperties properties {};

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

@@ -941,8 +941,12 @@ void RenderExecutor::CommitBindings(CommandBuffer&                     buffer,
 				for (const auto& resource: program.info.buffers) {
 					if (!resource.written || resource.descriptor_index == UINT32_MAX) continue;
 					const auto& written = writer->buffer_sources[resource.descriptor_index];
-					if (written.size != 0 && ImageRangeOverlaps(address, size,
-					                                          written.address, written.size)) {
+					// ADD_TID and swizzle move fixed-offset writes past write_extent.
+					const bool linear = (resource.packed_stride & ((1u << 14u) | (1u << 20u))) == 0u;
+					const auto written_size =
+					    linear ? std::min<uint64_t>(written.size, resource.write_extent) : written.size;
+					if (written_size != 0 && ImageRangeOverlaps(address, size,
+					                                          written.address, written_size)) {
 						EXIT("scalar resource reads overlap a shader buffer write\n");
 					}
 				}

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

@@ -308,7 +308,7 @@ struct PipelineCache::ProgramCache {
 		auto result = ShaderRecompiler::CompileProgram(std::move(translated), options,
 		                                               specialization, push_data_start_dword);
 		if (!ValidateShaderSpirv(options.dump_label, options.shader_hash, result.spirv,
-		                         options.maintenance8_enabled)) {
+		                         options.profile->maintenance8)) {
 			DumpShaderSpirv(stage_name, options.shader_hash, result.spirv);
 			EXIT("%s failed hash=0x%016" PRIx64 ": SPIR-V validation failed\n", options.dump_label,
 			     options.shader_hash);
@@ -398,7 +398,7 @@ struct PipelineCache::ProgramCache {
 		}
 		ShaderRecompiler::CompileOptions options;
 		options.stage       = stage;
-		options.maintenance8_enabled = graphics.maintenance8_enabled;
+		options.profile     = &profile;
 		options.shader_hash = params.hash;
 		options.user_data   = user_data;
 		options.back_code      = params.back_code;
@@ -480,7 +480,13 @@ struct PipelineCache::ProgramCache {
 		return permutation.handle;
 	}
 
-	explicit ProgramCache(GraphicContext& graphics): graphics(graphics) {
+	explicit ProgramCache(GraphicContext& graphics)
+	    : graphics(graphics),
+	      // AMD GPUs (vendor 0x1002) run NMin/NMax as v_min/v_max, which give the guest's NaN and
+	      // signed-zero results; GLSL.std.450 alone does not order -0 below +0.
+	      profile {.maintenance8 = graphics.maintenance8_enabled,
+		           .native_float_min_max =
+		               graphics.physical_device_properties.vendorID == 0x1002u} {
 		lookup_key.static_state.reserve(MaxStaticKeyWords);
 	}
 	~ProgramCache() {
@@ -495,6 +501,7 @@ struct PipelineCache::ProgramCache {
 	std::unordered_map<ProgramKey, SourceEntry, ProgramKeyHash> programs;
 	ProgramKey                                                  lookup_key;
 	GraphicContext&                                             graphics;
+	const ShaderRecompiler::Spirv::Profile                      profile;
 	uint64_t                                                    next_shader_id = 0;
 };
 
@@ -971,6 +978,7 @@ PipelineCache::Pipeline& PipelineCache::GetGraphicsPipeline(
 	if (static_params.sample_shading_enable && !m_graphics.sample_rate_shading_enabled) {
 		EXIT("Pipeline: sample-rate shading is required but unsupported by the host\n");
 	}
+	static_params.alpha_to_coverage_enable = ps_active && ctx.GetShaderRegisters().AlphaToMask();
 	const bool rect_list = Prospero::IsRectList(command.GetUserConfig().GetPrimType());
 	static_params.cull_back  = !rect_list && mc.cull_back;
 	static_params.cull_front = !rect_list && mc.cull_front;

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

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

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

@@ -92,7 +92,8 @@ bool HasShaderBufferWrites(const ShaderStageRuntime& runtime) {
 	const auto& program   = *runtime.program;
 	const auto& resources = *runtime.resources;
 	EXIT_IF(resources.buffers.size() != program.info.buffers.size());
-	bool has_writes = false;
+	// Address writes reach memory outside the bound buffer descriptors.
+	bool has_writes = program.has_address_writes;
 	for (uint32_t i = 0; i < program.info.buffers.size(); i++) {
 		if (!program.info.buffers[i].written) {
 			continue;

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

@@ -422,6 +422,8 @@ void CreatePipelineInternal(GraphicContext& graphics, PipelineCache::Pipeline& p
 	multisampling.sampleShadingEnable  = static_params.sample_shading_enable ? VK_TRUE : VK_FALSE;
 	multisampling.rasterizationSamples = vulkan_sample_count(static_params.samples);
 	multisampling.minSampleShading     = 1.0f;
+	multisampling.alphaToCoverageEnable =
+	    static_params.alpha_to_coverage_enable ? VK_TRUE : VK_FALSE;
 
 	vk::PipelineColorBlendAttachmentState color_blend_attachment[RENDER_COLOR_ATTACHMENTS_MAX] = {};
 	for (uint32_t i = 0; i < rendering.color_count; i++) {

src/graphics/host_gpu/renderer/renderDraw.cpp

@@ -436,7 +436,7 @@ static bool DrawHasValidVertexShader(const HW::Shader& sh_ctx) {
 static bool PixelShaderHasDepthOrCoverageSideEffects(const HW::ShaderRegisters& sh_regs) {
 	const auto& db = sh_regs.db_shader_control;
 	return db.shader_kill_enable || db.shader_z_export_enable || db.shader_mask_export_enable ||
-	       db.shader_dual_export_enable || db.shader_execute_on_noop;
+	       db.shader_dual_export_enable || db.shader_execute_on_noop || sh_regs.AlphaToMask();
 }
 
 struct DrawRenderState {
@@ -935,6 +935,29 @@ bool RenderExecutor::PrepareDrawRenderState(CommandBuffer& buffer, const DrawCal
 		LogDrawPhase(draw.Name(), "ResolveRenderDepthTarget");
 	}
 	ResolveRenderDepthTarget(buffer, state.depth_info);
+	// TODO(PPSA09243): Vulkan passes take one sample count. Where PS5 attachments mix counts,
+	// keep those matching the raster count instead of resolving multisampled shadows.
+	const auto colors       = std::span {state.color_info, state.color_count};
+	const auto pass_samples = state.depth_info.image_id ? state.depth_info.desc.info.samples
+	                                                    : state.color_info[0].desc.info.samples;
+	if (std::ranges::any_of(colors, [&](const RenderColorInfo& color) {
+		    return color.desc.info.samples != pass_samples;
+	    })) {
+		const auto raster =
+		    render_sample_count(buffer.GetRegisters().GetAaConfig().msaa_num_samples);
+		const auto dropped = std::ranges::remove_if(colors, [&](const RenderColorInfo& color) {
+			return color.desc.info.samples != raster;
+		});
+		state.color_count -= static_cast<uint32_t>(dropped.size());
+		if (state.depth_info.image_id && state.depth_info.desc.info.samples != raster) {
+			state.depth_info = {};
+		}
+		static std::atomic_bool logged = false;
+		if (!logged.exchange(true, std::memory_order_relaxed)) {
+			Log::WriteToConsoleAndLog(
+			    "Warning: dropping attachments whose sample count differs from the raster.\n");
+		}
+	}
 
 	if (state.color_count == 0 && !state.depth_info.image_id && !state.ps_active) {
 		LogFramebufferSkip(draw.Name(), state.color_info[0], state.depth_info, buffer,
@@ -1227,6 +1250,20 @@ void RenderExecutor::ExecutePreparedDraw(uint64_t submit_id, CommandBuffer& buff
 	dynamic_state.SetAttachmentFeedbackLoopAspects(feedback_aspects);
 	dynamic_state.Commit(m_context.GetGraphics(), vk_buffer);
 
+	vk::PipelineStageFlags shader_write_stages = {};
+	for (const auto& stage: vertex_stages) {
+		if (HasShaderBufferWrites(stage.stage)) {
+			shader_write_stages |= ShaderPipelineStages(NativeShaderStage(stage.logical_stage));
+		}
+	}
+	if (state.ps_active && HasShaderBufferWrites(state.ps_input_info.stage)) {
+		shader_write_stages |= vk::PipelineStageFlagBits::eFragmentShader;
+	}
+	if (shader_write_stages) {
+		// As for dispatches, buffer writes wait for earlier accesses to the memory.
+		m_context.GetCommandScheduler().EndRendering();
+		ShaderWriteHazardBarrier(vk_buffer, shader_write_stages);
+	}
 	LogDrawPhase(draw.Name(), "BeginRendering");
 	if (!draw.IsIndexed()) {
 		SetDrawDebugPhase(buffer, submit_id, draw, 0x400u);
@@ -1247,15 +1284,6 @@ void RenderExecutor::ExecutePreparedDraw(uint64_t submit_id, CommandBuffer& buff
 	if (!draw.IsIndexed()) {
 		SetDrawDebugPhase(buffer, submit_id, draw, 0x600u);
 	}
-	vk::PipelineStageFlags shader_write_stages = {};
-	for (const auto& stage: vertex_stages) {
-		if (HasShaderBufferWrites(stage.stage)) {
-			shader_write_stages |= ShaderPipelineStages(NativeShaderStage(stage.logical_stage));
-		}
-	}
-	if (state.ps_active && HasShaderBufferWrites(state.ps_input_info.stage)) {
-		shader_write_stages |= vk::PipelineStageFlagBits::eFragmentShader;
-	}
 	if (shader_write_stages) {
 		m_context.GetCommandScheduler().EndRendering();
 		ShaderWriteBarrier(vk_buffer, shader_write_stages);

src/graphics/shader/recompiler/ShaderRecompiler.cpp

@@ -930,8 +930,8 @@ CompileResult CompileProgram(TranslateResult translated, const CompileOptions& o
 
 	LOGF("%s phase begin: stage=%s hash=0x%016" PRIx64 " SPIR-V EmitProgram\n",
 	     GetDumpLabel(options), StageName(ir.stage), ir.shader_hash);
-	auto spirv = Spirv::EmitProgram(ir, options.input_info, push_data_start_dword,
-	                                options.maintenance8_enabled, options.emit_debug_info);
+	auto spirv = Spirv::EmitProgram(ir, options.input_info, push_data_start_dword, *options.profile,
+	                                options.emit_debug_info);
 	LOGF("%s phase end: stage=%s hash=0x%016" PRIx64 " SPIR-V EmitProgram words=%" PRIu64
 	     " elapsed_ms=%" PRIu64 "\n",
 	     GetDumpLabel(options), StageName(ir.stage), ir.shader_hash,

src/graphics/shader/recompiler/ShaderRecompiler.h

@@ -3,6 +3,7 @@
 
 #include "common/common.h"
 #include "common/stringUtils.h"
+#include "graphics/shader/recompiler/backend/spirv/SpirvEmitter.h"
 #include "graphics/shader/recompiler/ir/passes/ResourceMaterialization.h"
 #include "graphics/shader/recompiler/frontend/decode/ShaderDecoder.h"
 #include "graphics/shader/shader.h"
@@ -23,7 +24,7 @@ struct CompileOptions {
 	bool                        dump_ir                    = true;
 	bool                        early_dump                 = false;
 	bool                        emit_debug_info            = false;
-	bool                        maintenance8_enabled       = false;
+	const Spirv::Profile*       profile                    = &Spirv::DefaultProfile;
 	const char*                 dump_label                 = nullptr;
 	std::span<const uint32_t>   user_data;
 	std::span<const uint32_t>   back_code;

src/graphics/shader/recompiler/Tessellation.cpp

@@ -34,8 +34,8 @@ void ReflectTessellationLayout(const Decoder::Program& program, bool local,
 	using Registers = std::array<Address, IR::NumVectorRegs>;
 
 	const auto control_points = local ? info.input_control_points : info.output_control_points;
-	Registers  registers {};
-	registers[local ? 3u : 1u] =
+	Registers  entry {}, registers;
+	entry[local ? 3u : 1u] =
 	    local ? Address {Kind::Affine, 1u, 0u} : Address {Kind::PackedControlPoint};
 	const auto constant = [](uint32_t value) { return Address {Kind::Affine, 0u, value}; };
 	const auto read     = [&](const Operand& operand) -> Address {
@@ -85,101 +85,114 @@ void ReflectTessellationLayout(const Decoder::Program& program, bool local,
 		           ? value
 		           : Address {};
 	};
-	// A forward branch joins its target with the fallthrough path. Definitions that
-	// differ between the paths are unknown there.
-	std::unordered_map<uint32_t, Registers> joins;
+	// A branch joins its target with the fallthrough path. Definitions that differ between
+	// the paths are unknown there. Loop headers repeat the walk until back edges are stable.
+	std::unordered_map<uint32_t, Registers> joins, loops;
 	const auto merge = [](Registers& into, const Registers& from) {
+		bool changed = false;
 		for (uint32_t index = 0; index < into.size(); index++) {
-			if (into[index] != from[index]) into[index] = {};
+			if (into[index] == from[index]) continue;
+			changed |= into[index] != Address {};
+			into[index] = {};
 		}
+		return changed;
 	};
 
 	uint32_t stride = 0, patch_begin = UINT32_MAX, patch_end = 0;
-	for (const auto& inst: program.instructions) {
-		if (const auto join = joins.find(inst.pc); join != joins.end()) {
-			merge(registers, join->second);
-			joins.erase(join);
-		}
-		if (IsDirectBranch(inst.opcode)) {
-			EXIT_NOT_IMPLEMENTED(inst.branch_target <= inst.pc);
-			if (const auto [join, inserted] = joins.try_emplace(inst.branch_target, registers);
-			    !inserted)
-				merge(join->second, registers);
-		}
-		const bool local_store  = IsDsWriteOpcode(inst.opcode);
-		const bool buffer_store = inst.opcode == Opcode::BUFFER_STORE_DWORD ||
-		                          inst.opcode == Opcode::BUFFER_STORE_DWORDX2 ||
-		                          inst.opcode == Opcode::BUFFER_STORE_DWORDX3 ||
-		                          inst.opcode == Opcode::BUFFER_STORE_DWORDX4;
-		const bool control_store =
-		    !local && buffer_store && inst.src2.kind == OperandKind::Sgpr && inst.src2.reg == 2u;
-		const bool control_read =
-		    !local && (inst.opcode == Opcode::DS_READ_B32 || inst.opcode == Opcode::DS_READ2_B32);
-		if ((local && local_store) || control_store || control_read) {
-			const auto address = read(inst.src0);
-			if (address.kind != Kind::Affine) {
-				EXIT("%s tessellation address is not affine at pc 0x%08x\n", local ? "LS" : "HS",
-				     inst.pc);
+	for (bool repeat = true; repeat;) {
+		repeat    = false;
+		registers = entry;
+		for (const auto& inst: program.instructions) {
+			if (const auto join = joins.find(inst.pc); join != joins.end()) {
+				merge(registers, join->second);
+				joins.erase(join);
 			}
-			if (address.coefficient != 0u) {
-				const auto expected = control_read ? info.ls_stride : stride;
-				EXIT_NOT_IMPLEMENTED((address.coefficient & 3u) != 0u ||
-				                     (expected != 0u && expected != address.coefficient));
-				if (!control_read) stride = address.coefficient;
-			} else if (control_store) {
-				// Control-point-independent HS outputs are patch constants.
-				const auto begin = address.constant + inst.offset;
-				patch_begin      = std::min(patch_begin, begin);
-				patch_end        = std::max(patch_end, begin + 4u * inst.data_dwords);
+			if (const auto loop = loops.find(inst.pc); loop != loops.end())
+				merge(registers, loop->second);
+			if (IsDirectBranch(inst.opcode)) {
+				const bool back = inst.branch_target <= inst.pc;
+				const auto [join, inserted] =
+				    (back ? loops : joins).try_emplace(inst.branch_target, registers);
+				const bool changed = !inserted && merge(join->second, registers);
+				repeat |= back && (inserted || changed);
 			}
-		}
-		// Store vdata is a source, despite occupying the decoder's dst field.
-		if (local_store || buffer_store || inst.dst.kind != OperandKind::Vgpr) continue;
-		const auto lhs   = read(inst.src0);
-		const auto rhs   = read(inst.src1);
-		const auto third = read(inst.src2);
-		Address    value;
-		switch (inst.opcode) {
-			case Opcode::V_BFE_U32:
-				if (lhs.kind == Kind::PackedControlPoint && rhs.kind == Kind::Affine &&
-				    rhs.coefficient == 0u && third.kind == Kind::Affine &&
-				    third.coefficient == 0u) {
-					if (rhs.constant == 8u && third.constant == 5u) value = {Kind::Affine, 1u, 0u};
-					if (rhs.constant == 0u && third.constant == 8u) value = constant(0u);
+			const bool local_store  = IsDsWriteOpcode(inst.opcode);
+			const bool buffer_store = inst.opcode == Opcode::BUFFER_STORE_DWORD ||
+			                          inst.opcode == Opcode::BUFFER_STORE_DWORDX2 ||
+			                          inst.opcode == Opcode::BUFFER_STORE_DWORDX3 ||
+			                          inst.opcode == Opcode::BUFFER_STORE_DWORDX4;
+			const bool control_store = !local && buffer_store &&
+			                           inst.src2.kind == OperandKind::Sgpr && inst.src2.reg == 2u;
+			if ((local && local_store) || control_store) {
+				const auto address = read(inst.src0);
+				if (address.kind != Kind::Affine) {
+					EXIT("%s tessellation address is not affine at pc 0x%08x\n",
+					     local ? "LS" : "HS", inst.pc);
 				}
-				break;
-			case Opcode::V_AND_B32:
-				if ((lhs.kind == Kind::PackedControlPoint && rhs == constant(0xffu)) ||
-				    (rhs.kind == Kind::PackedControlPoint && lhs == constant(0xffu)))
-					value = constant(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;
-			case Opcode::V_MAD_U32_U24: value = add(multiply(low24(lhs), low24(rhs)), third); break;
-			case Opcode::V_LSHL_ADD_U32:
-				if (rhs.kind == Kind::Affine && rhs.coefficient == 0u && rhs.constant < 32u)
-					value = add(multiply(lhs, constant(1u << rhs.constant)), third);
-				break;
-			case Opcode::V_LSHLREV_B32:
-				if (lhs.kind == Kind::Affine && lhs.coefficient == 0u && lhs.constant < 32u)
-					value = multiply(rhs, constant(1u << lhs.constant));
-				break;
-			case Opcode::V_SUB_NC_U32:
-				if (lhs.kind == Kind::Affine && rhs.kind == Kind::Affine &&
-				    lhs.coefficient >= rhs.coefficient && lhs.constant >= rhs.constant)
-					value = affine(lhs.coefficient - rhs.coefficient, lhs.constant - rhs.constant);
-				break;
-			default: break;
+				if (address.coefficient != 0u) {
+					EXIT_NOT_IMPLEMENTED((address.coefficient & 3u) != 0u ||
+					                     (stride != 0u && stride != address.coefficient));
+					stride = address.coefficient;
+				} else if (control_store) {
+					// Control-point-independent HS outputs are patch constants.
+					const auto begin = address.constant + inst.offset;
+					patch_begin      = std::min(patch_begin, begin);
+					patch_end        = std::max(patch_end, begin + 4u * inst.data_dwords);
+				}
+			}
+			// Store vdata is a source, despite occupying the decoder's dst field.
+			if (local_store || buffer_store || inst.dst.kind != OperandKind::Vgpr) continue;
+			const auto lhs   = read(inst.src0);
+			const auto rhs   = read(inst.src1);
+			const auto third = read(inst.src2);
+			Address    value;
+			switch (inst.opcode) {
+				case Opcode::V_BFE_U32:
+					if (lhs.kind == Kind::PackedControlPoint && rhs.kind == Kind::Affine &&
+					    rhs.coefficient == 0u && third.kind == Kind::Affine &&
+					    third.coefficient == 0u) {
+						if (rhs.constant == 8u && third.constant == 5u)
+							value = {Kind::Affine, 1u, 0u};
+						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);
+					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;
+				case Opcode::V_MAD_U32_U24:
+					value = add(multiply(low24(lhs), low24(rhs)), third);
+					break;
+				case Opcode::V_LSHL_ADD_U32:
+					if (rhs.kind == Kind::Affine && rhs.coefficient == 0u && rhs.constant < 32u)
+						value = add(multiply(lhs, constant(1u << rhs.constant)), third);
+					break;
+				case Opcode::V_LSHLREV_B32:
+					if (lhs.kind == Kind::Affine && lhs.coefficient == 0u && lhs.constant < 32u)
+						value = multiply(rhs, constant(1u << lhs.constant));
+					break;
+				case Opcode::V_SUB_NC_U32:
+					if (lhs.kind == Kind::Affine && rhs.kind == Kind::Affine &&
+					    lhs.coefficient >= rhs.coefficient && lhs.constant >= rhs.constant)
+						value =
+						    affine(lhs.coefficient - rhs.coefficient, lhs.constant - rhs.constant);
+					break;
+				default: break;
+			}
+			if (inst.dst.sdwa_sel != 6u || inst.dst.op_sel || inst.dst.omod || inst.dst.clamp ||
+			    inst.dst.dpp)
+				value = {};
+			for (uint32_t index = 0;
+			     index < std::max(inst.data_dwords, 1u) && inst.dst.reg + index < registers.size();
+			     index++) {
+				registers[inst.dst.reg + index] = {};
+			}
+			registers.at(inst.dst.reg) = value;
 		}
-		if (inst.dst.sdwa_sel != 6u || inst.dst.op_sel || inst.dst.omod || inst.dst.clamp ||
-		    inst.dst.dpp)
-			value = {};
-		for (uint32_t index = 0;
-		     index < std::max(inst.data_dwords, 1u) && inst.dst.reg + index < registers.size(); index++) {
-			registers[inst.dst.reg + index] = {};
-		}
-		registers.at(inst.dst.reg) = value;
 	}
 	EXIT_NOT_IMPLEMENTED(!joins.empty() || stride == 0u);
 	(local ? info.ls_stride : info.hs_stride) = stride;
@@ -318,20 +331,22 @@ void LowerTessellationMemory(IR::Program& program, const CompileOptions& options
 					kind = TessellationAttribute::Factor;
 					factors += components;
 				} else {
-					kind = write ? TessellationAttribute::ControlOutput
-					             : TessellationAttribute::EvaluationInput;
-					EXIT_NOT_IMPLEMENTED(write
-					                         ? options.stage != ShaderType::TessellationControl
-					                         : options.stage != ShaderType::TessellationEvaluation);
+					// After its barrier, the HS may read back outputs of other control points.
+					const bool output = options.stage == ShaderType::TessellationControl;
+					EXIT_NOT_IMPLEMENTED(
+					    !output && (write || options.stage != ShaderType::TessellationEvaluation));
+					kind = output ? TessellationAttribute::ControlOutput
+					              : TessellationAttribute::EvaluationInput;
 					if (address.IsImmediate() &&
 					    address.U32() >= tess.hs_stride * tess.output_control_points) {
 						const auto begin = address.U32() + memory.offset;
 						EXIT_NOT_IMPLEMENTED(begin < tess.patch_base ||
 						                     begin + 4u * components >
 						                         tess.patch_base + tess.patch_size);
-						kind = write ? TessellationAttribute::PatchOutput
-						             : TessellationAttribute::PatchInput;
-						patches += components;
+						kind = output ? TessellationAttribute::PatchOutput
+						              : TessellationAttribute::PatchInput;
+						// HS read-backs must not hide misplaced patch stores from the check below.
+						if (write || !output) patches += components;
 					}
 				}
 			} else {

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

@@ -204,7 +204,7 @@ void ValidateNativeProgram(const IR::Program& program, bool lds_storage) {
 } // namespace
 
 std::vector<uint32_t> EmitProgram(IR::Program& program, ShaderStageInputInfo input_info,
-                                  uint32_t push_data_start_dword, bool maintenance8_enabled,
+                                  uint32_t push_data_start_dword, const Profile& profile,
                                   bool emit_debug_info) {
 	using namespace Emitter;
 
@@ -220,7 +220,7 @@ std::vector<uint32_t> EmitProgram(IR::Program& program, ShaderStageInputInfo inp
 	}
 	IR::ValidateProgram(program, true);
 	program.bindings = {.push_data_start_dword = push_data_start_dword};
-	EmitterState state(program, input_info, maintenance8_enabled);
+	EmitterState state(program, input_info, profile);
 	if (emit_debug_info) {
 		char name[64];
 		std::snprintf(name, sizeof(name), "rdna2-%016" PRIx64 "-block-byte-pc-plus-one",

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

@@ -9,11 +9,19 @@
 
 namespace Libs::Graphics::ShaderRecompiler::Spirv {
 
+// Host device capabilities that change the emitted SPIR-V. Built once per device.
+struct Profile {
+	bool maintenance8         = false; // Dynamic texel offsets on image samples.
+	bool native_float_min_max = false; // NMin/NMax give the guest's NaN and signed-zero results.
+};
+
+inline constexpr Profile DefaultProfile {};
+
 // Resets and writes backend IDs and descriptor indices; do not emit the same program concurrently.
 std::vector<uint32_t> EmitProgram(IR::Program& program, ShaderStageInputInfo input_info,
-                                  uint32_t push_data_start_dword = 0,
-                                  bool maintenance8_enabled = false,
-                                  bool emit_debug_info = false);
+                                  uint32_t       push_data_start_dword = 0,
+                                  const Profile& profile               = DefaultProfile,
+                                  bool           emit_debug_info       = false);
 
 } // namespace Libs::Graphics::ShaderRecompiler::Spirv
 

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

@@ -5,18 +5,6 @@
 namespace Libs::Graphics::ShaderRecompiler::Spirv::Emitter {
 namespace {
 
-struct Pair {
-	uint32_t low  = 0;
-	uint32_t high = 0;
-};
-
-Pair ExtractPair(EmitterState& state, uint32_t value) {
-	Pair result {state.builder.AllocateId(), state.builder.AllocateId()};
-	state.builder.AddFunction(spv::OpCompositeExtract, TypeU32(state), result.low, value, 0);
-	state.builder.AddFunction(spv::OpCompositeExtract, TypeU32(state), result.high, value, 1);
-	return result;
-}
-
 uint32_t EmitMinMaxF64(EmitterState& state, uint32_t lhs, uint32_t rhs, bool max_value) {
 	const auto bits_type = TypeU32Vector(state, 2);
 	const auto lhs_bits = Unary(state, spv::OpBitcast, bits_type, lhs);
@@ -32,11 +20,10 @@ uint32_t EmitMinMaxF64(EmitterState& state, uint32_t lhs, uint32_t rhs, bool max
 	result = Select(state, TypeF64(state), equal,
 	                Unary(state, spv::OpBitcast, TypeF64(state), equal_bits), result);
 
-	// Non-IEEE mode selects the other operand for NaN, including rhs when both are NaN.
-	result = Select(state, TypeF64(state), Unary(state, spv::OpIsNan, TypeBool(state), rhs),
-	                lhs, result);
-	return Select(state, TypeF64(state), Unary(state, spv::OpIsNan, TypeBool(state), lhs),
-	              rhs, result);
+	// Non-IEEE mode selects the other operand for NaN, and lhs when both are NaN. The ordered
+	// compares already pick rhs for a NaN lhs.
+	return Select(state, TypeF64(state), Unary(state, spv::OpIsNan, TypeBool(state), rhs),
+	              lhs, result);
 }
 
 uint32_t EmitMulHigh(EmitterState& state, uint32_t lhs, uint32_t rhs, bool signed_value) {
@@ -106,35 +93,23 @@ uint32_t EmitF32ToF16RtzBits(EmitterState& state, uint32_t value) {
 }
 
 uint32_t EmitF32ToU32(EmitterState& state, uint32_t src, bool signed_value) {
-	const auto trunc         = EmitTruncF32Value(state, src);
-	const auto converted_raw = state.builder.AllocateId();
-	if (signed_value) {
-		const auto converted_i = state.builder.AllocateId();
-		state.builder.AddFunction(spv::OpConvertFToS, TypeI32(state), converted_i, trunc);
-		state.builder.AddFunction(spv::OpBitcast, TypeU32(state), converted_raw, converted_i);
-	} else {
-		state.builder.AddFunction(spv::OpConvertFToU, TypeU32(state), converted_raw, trunc);
-	}
-	const auto nan = EmitClassifyF32(state, src).nan;
-	if (signed_value) {
-		const auto below = Binary(state, spv::OpFOrdLessThanEqual, TypeBool(state), src,
-		                          ConstantF32(state, 0xcf000000u));
-		const auto above = Binary(state, spv::OpFOrdGreaterThanEqual, TypeBool(state), src,
-		                          ConstantF32(state, 0x4f000000u));
-		const auto high =
-		    Select(state, TypeU32(state), above, ConstantU32(state, 0x7fffffffu), converted_raw);
-		const auto low =
-		    Select(state, TypeU32(state), below, ConstantU32(state, 0x80000000u), high);
-		return Select(state, TypeU32(state), nan, ConstantU32(state, 0), low);
-	}
-	const auto below =
-	    Binary(state, spv::OpFOrdLessThanEqual, TypeBool(state), src, ConstantF32(state, 0));
-	const auto above = Binary(state, spv::OpFOrdGreaterThanEqual, TypeBool(state), src,
-	                          ConstantF32(state, 0x4f800000u));
-	const auto zero  = Binary(state, spv::OpLogicalOr, TypeBool(state), nan, below);
+	// Clamping first keeps the truncating conversion defined; NaN and values past the largest
+	// float below the integer limit are selected afterwards.
+	const auto clamped =
+	    EmitExt(state, TypeF32(state), GLSLstd450NClamp,
+	            {src, ConstantF32(state, signed_value ? 0xcf000000u : 0u),
+	             ConstantF32(state, signed_value ? 0x4effffffu : 0x4f7fffffu)});
+	const auto converted = signed_value
+	                           ? Unary(state, spv::OpBitcast, TypeU32(state),
+	                                   Unary(state, spv::OpConvertFToS, TypeI32(state), clamped))
+	                           : Unary(state, spv::OpConvertFToU, TypeU32(state), clamped);
+	const auto above     = Binary(state, spv::OpFOrdGreaterThanEqual, TypeBool(state), src,
+	                              ConstantF32(state, signed_value ? 0x4f000000u : 0x4f800000u));
 	const auto high =
-	    Select(state, TypeU32(state), above, ConstantU32(state, 0xffffffffu), converted_raw);
-	return Select(state, TypeU32(state), zero, ConstantU32(state, 0), high);
+	    Select(state, TypeU32(state), above,
+	           ConstantU32(state, signed_value ? 0x7fffffffu : 0xffffffffu), converted);
+	return Select(state, TypeU32(state), Unary(state, spv::OpIsNan, TypeBool(state), src),
+	              ConstantU32(state, 0), high);
 }
 
 } // namespace
@@ -160,18 +135,6 @@ uint32_t EmitFPMad32(EmitterState& state, uint32_t a, uint32_t b, uint32_t c) {
 	return sum;
 }
 
-uint32_t EmitFindUMsb64(EmitterState& state, uint32_t value) {
-	// GLSL FindUMsb is restricted to 32-bit components.
-	const auto unpacked = Unary(state, spv::OpBitcast, TypeU32Vector(state, 2), value);
-	const auto pair =
-	    ExtractPair(state, EmitExt(state, TypeU32Vector(state, 2), GLSLstd450FindUMsb, {unpacked}));
-	const auto high_found =
-	    Binary(state, spv::OpINotEqual, TypeBool(state), pair.high, ConstantU32(state, UINT32_MAX));
-	return Select(state, TypeU32(state), high_found,
-	              Binary(state, spv::OpIAdd, TypeU32(state), pair.high, ConstantU32(state, 32)),
-	              pair.low);
-}
-
 uint32_t EmitConvertU16U32(EmitterState& state, uint32_t arg0) {
 	return EmitNative<spv::OpBitwiseAnd, IR::Type::U16>(state, arg0, ConstantU32(state, 0xffffu));
 }
@@ -235,18 +198,15 @@ uint32_t EmitCompositeExtractU32x2(EmitterState& state, uint32_t arg0, IR::Value
 }
 
 uint32_t EmitPackFloat2x16Rtz(EmitterState& state, uint32_t arg0, uint32_t arg1) {
-	const auto low  = EmitF32ToF16RtzBits(state, arg0);
-	const auto high = Binary(state, spv::OpShiftLeftLogical, TypeU32(state),
-	                         EmitF32ToF16RtzBits(state, arg1), ConstantU32(state, 16));
-	return EmitOrU32(state, low, high);
+	return EmitNative<spv::OpBitFieldInsert, IR::Type::U32>(
+	    state, EmitF32ToF16RtzBits(state, arg0), EmitF32ToF16RtzBits(state, arg1),
+	    ConstantU32(state, 16), ConstantU32(state, 16));
 }
 
 uint32_t EmitFPSaturate32(EmitterState& state, uint32_t arg0) {
-	// The clamp modifier runs in DX10 clamp mode: NaN becomes 0. FClamp leaves NaN undefined.
-	const auto clamped = EmitExt(state, TypeF32(state), GLSLstd450FClamp,
-	                             {arg0, ConstantF32(state, 0), ConstantF32(state, 0x3f800000u)});
-	return Select(state, TypeF32(state), EmitClassifyF32(state, arg0).nan, ConstantF32(state, 0),
-	              clamped);
+	// The clamp modifier runs in DX10 clamp mode: NaN becomes 0, as NClamp's NMax step does.
+	return EmitExt(state, TypeF32(state), GLSLstd450NClamp,
+	               {arg0, ConstantF32(state, 0), ConstantF32(state, 0x3f800000u)});
 }
 
 uint32_t EmitSMulHi(EmitterState& state, uint32_t arg0, uint32_t arg1) {
@@ -257,14 +217,6 @@ uint32_t EmitUMulHi(EmitterState& state, uint32_t arg0, uint32_t arg1) {
 	return EmitMulHigh(state, arg0, arg1, false);
 }
 
-uint32_t EmitBitCount64(EmitterState& state, uint32_t arg0) {
-	// Vulkan bit-count operands remain 32-bit, including when counting a U64.
-	const auto unpacked = Unary(state, spv::OpBitcast, TypeU32Vector(state, 2), arg0);
-	const auto pair =
-	    ExtractPair(state, Unary(state, spv::OpBitCount, TypeU32Vector(state, 2), unpacked));
-	return Binary(state, spv::OpIAdd, TypeU32(state), pair.low, pair.high);
-}
-
 uint32_t EmitSMinTri32(EmitterState& state, uint32_t arg0, uint32_t arg1, uint32_t arg2) {
 	return EmitSMin32(state, arg0, EmitSMin32(state, arg1, arg2));
 }
@@ -341,8 +293,8 @@ uint32_t EmitFPSqrt(EmitterState& state, uint32_t arg0) {
 }
 
 uint32_t EmitFPExp2(EmitterState& state, uint32_t arg0) {
-	return EmitExt(state, TypeF32(state), GLSLstd450Exp2,
-	               {EmitFlushF32DenormToSignedZero(state, arg0)});
+	// A denormal input already rounds to 1.0, the result for its flushed zero.
+	return EmitExt(state, TypeF32(state), GLSLstd450Exp2, {arg0});
 }
 
 uint32_t EmitFPLog2(EmitterState& state, uint32_t arg0) {

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

@@ -17,16 +17,9 @@ uint32_t EmitCompareU32Constant(EmitterState& state, spv::Op opcode, uint32_t va
 }
 
 F32Class EmitClassifyF32Bits(EmitterState& state, uint32_t bits) {
-	F32Class cls;
-	cls.bits            = bits;
-	const auto abs_bits = EmitAndConstant(state, cls.bits, 0x7fffffffu);
-	cls.nan  = EmitCompareU32Constant(state, spv::OpUGreaterThan, abs_bits, 0x7f800000u);
-	cls.zero = EmitCompareU32Constant(state, spv::OpIEqual, abs_bits, 0);
-	return cls;
-}
-
-F32Class EmitClassifyF32(EmitterState& state, uint32_t value) {
-	return EmitClassifyF32Bits(state, EmitBitcastF32ToU32(state, value));
+	const auto abs_bits = EmitAndConstant(state, bits, 0x7fffffffu);
+	return {EmitCompareU32Constant(state, spv::OpUGreaterThan, abs_bits, 0x7f800000u),
+	        EmitCompareU32Constant(state, spv::OpIEqual, abs_bits, 0)};
 }
 
 uint32_t EmitClassMaskBitMatch(EmitterState& state, uint32_t mask, uint32_t bit,
@@ -100,102 +93,88 @@ uint32_t EmitClassMaskF16(EmitterState& state, uint32_t bits, uint32_t mask) {
 	return EmitClassMaskBits(state, bits, mask, true);
 }
 
-static uint32_t EmitMinMaxF32(EmitterState& state, uint32_t lhs, uint32_t rhs,
-                             const F32Class& lhs_class, const F32Class& rhs_class,
-                             bool max_value) {
-	const auto numeric_cond = state.builder.AllocateId();
-	state.builder.AddFunction(max_value ? spv::OpFOrdGreaterThanEqual : spv::OpFOrdLessThan,
-	                          TypeBool(state), numeric_cond, lhs, rhs);
-	const auto ordered_bits =
-	    EmitSelectValueU32(state, numeric_cond, lhs_class.bits, rhs_class.bits);
-
-	const auto both_zero    = EmitLogicalAndBool(state, lhs_class.zero, rhs_class.zero);
-	const auto zero_bits    = max_value ? EmitAndU32(state, lhs_class.bits, rhs_class.bits)
-	                                    : EmitOrU32(state, lhs_class.bits, rhs_class.bits);
-	const auto numeric_bits = EmitSelectValueU32(state, both_zero, zero_bits, ordered_bits);
-
-	const auto rhs_nan_bits =
-	    EmitSelectValueU32(state, rhs_class.nan, lhs_class.bits, numeric_bits);
-	const auto result_bits = EmitSelectValueU32(state, lhs_class.nan, rhs_class.bits, rhs_nan_bits);
-	return EmitBitcastU32ToF32(state, result_bits);
+static uint32_t EmitMinMaxF32(EmitterState& state, uint32_t lhs, uint32_t rhs, uint32_t lhs_bits,
+                              uint32_t rhs_bits, bool max_value) {
+	// PS5 V_MIN/V_MAX_F32 return the other source for a NaN source (S0 when both are NaN)
+	// and order -0 below +0. An ordered compare takes S1 for a NaN S0 and, unlike NMin/NMax,
+	// always returns the exact bits of one source.
+	const auto either    = EmitOrU32(state, lhs_bits, rhs_bits);
+	const auto both_zero = EmitCompareU32Constant(state, spv::OpIEqual,
+	                                              EmitAndConstant(state, either, 0x7fffffffu), 0);
+	const auto zero =
+	    EmitBitcastU32ToF32(state, max_value ? EmitAndU32(state, lhs_bits, rhs_bits) : either);
+	const auto keep_lhs =
+	    Binary(state, max_value ? spv::OpFOrdGreaterThanEqual : spv::OpFOrdLessThan,
+		       TypeBool(state), lhs, rhs);
+	const auto ordered = Select(state, TypeF32(state), keep_lhs, lhs, rhs);
+	const auto numeric = Select(state, TypeF32(state), both_zero, zero, ordered);
+	return Select(state, TypeF32(state), Unary(state, spv::OpIsNan, TypeBool(state), rhs), lhs,
+	              numeric);
 }
 
 uint32_t EmitMinMaxF32Value(EmitterState& state, uint32_t lhs, uint32_t rhs, bool max_value) {
-	const auto lhs_class = EmitClassifyF32(state, lhs);
-	const auto rhs_class = EmitClassifyF32(state, rhs);
-	return EmitMinMaxF32(state, lhs, rhs, lhs_class, rhs_class, max_value);
+	if (state.profile.native_float_min_max) {
+		const auto result = state.builder.AllocateId();
+		state.builder.AddFunction(spv::OpExtInst, TypeF32(state), result, GlslStd450(state),
+		                          max_value ? GLSLstd450NMax : GLSLstd450NMin, lhs, rhs);
+		return result;
+	}
+	return EmitMinMaxF32(state, lhs, rhs, EmitBitcastF32ToU32(state, lhs),
+	                     EmitBitcastF32ToU32(state, rhs), max_value);
 }
 
 uint32_t EmitFPMedTri32(EmitterState& state, uint32_t a, uint32_t b, uint32_t c) {
-	const auto a_class = EmitClassifyF32(state, a);
-	const auto b_class = EmitClassifyF32(state, b);
-	const auto c_class = EmitClassifyF32(state, c);
-	const auto min_ab = EmitMinMaxF32(state, a, b, a_class, b_class, false);
-	const auto min_ab_class = EmitClassifyF32(state, min_ab);
-	const auto min3 = EmitMinMaxF32(state, min_ab, c, min_ab_class, c_class, false);
-	const auto max_ab = EmitMinMaxF32(state, a, b, a_class, b_class, true);
-	const auto max_ab_class = EmitClassifyF32(state, max_ab);
-	const auto max3 = EmitMinMaxF32(state, max_ab, c, max_ab_class, c_class, true);
-	const auto max_bc = EmitMinMaxF32(state, b, c, b_class, c_class, true);
-	const auto max_ac = EmitMinMaxF32(state, a, c, a_class, c_class, true);
-	// RDNA2 compares numerically against MAX3, including ties between signed zeros.
+	const auto a_bits = EmitBitcastF32ToU32(state, a);
+	const auto b_bits = EmitBitcastF32ToU32(state, b);
+	const auto c_bits = EmitBitcastF32ToU32(state, c);
+	// The NaN path is RDNA2 MIN3, which keeps the later source when two are NaN.
+	const auto min_ba = EmitMinMaxF32(state, b, a, b_bits, a_bits, false);
+	const auto min3 =
+	    EmitMinMaxF32(state, c, min_ba, c_bits, EmitBitcastF32ToU32(state, min_ba), false);
+	const auto max_ab = EmitMinMaxF32(state, a, b, a_bits, b_bits, true);
+	const auto max3 =
+	    EmitMinMaxF32(state, max_ab, c, EmitBitcastF32ToU32(state, max_ab), c_bits, true);
+	// RDNA2 compares numerically against MAX3, including ties between signed zeros, then takes
+	// V_MAX_F32 of the other two sources in their original order.
 	const auto a_is_max = Binary(state, spv::OpFOrdEqual, TypeBool(state), max3, a);
 	const auto b_is_max = Binary(state, spv::OpFOrdEqual, TypeBool(state), max3, b);
-	const auto median = Select(state, TypeF32(state), a_is_max, max_bc,
-	                           Select(state, TypeF32(state), b_is_max, max_ac, max_ab));
-	const auto nan_ab = EmitLogicalOrBool(state, a_class.nan, b_class.nan);
-	const auto any_nan = EmitLogicalOrBool(state, nan_ab, c_class.nan);
+	const auto x        = Select(state, TypeF32(state), a_is_max, b, a);
+	const auto y =
+	    Select(state, TypeF32(state), EmitLogicalOrBool(state, a_is_max, b_is_max), c, b);
+	const auto median = EmitMinMaxF32Value(state, x, y, true);
+	const auto is_nan = [&](uint32_t value) {
+		return Unary(state, spv::OpIsNan, TypeBool(state), value);
+	};
+	const auto any_nan =
+	    EmitLogicalOrBool(state, EmitLogicalOrBool(state, is_nan(a), is_nan(b)), is_nan(c));
 	return Select(state, TypeF32(state), any_nan, min3, median);
 }
 
 uint32_t EmitFlushF32DenormToSignedZero(EmitterState& state, uint32_t value) {
-	const auto bits      = state.builder.AllocateId();
-	const auto abs_bits  = state.builder.AllocateId();
-	const auto sign_bits = state.builder.AllocateId();
-	const auto subnormal = state.builder.AllocateId();
-	const auto selected  = state.builder.AllocateId();
-	const auto ret       = state.builder.AllocateId();
-	state.builder.AddFunction(spv::OpBitcast, TypeU32(state), bits, value);
-	state.builder.AddFunction(spv::OpBitwiseAnd, TypeU32(state), abs_bits, bits,
-	                          ConstantU32(state, 0x7fffffffu));
-	state.builder.AddFunction(spv::OpBitwiseAnd, TypeU32(state), sign_bits, bits,
-	                          ConstantU32(state, 0x80000000u));
-	state.builder.AddFunction(spv::OpULessThan, TypeBool(state), subnormal, abs_bits,
-	                          ConstantU32(state, 0x00800000u));
-	state.builder.AddFunction(spv::OpSelect, TypeU32(state), selected, subnormal, sign_bits, bits);
-	state.builder.AddFunction(spv::OpBitcast, TypeF32(state), ret, selected);
-	return ret;
+	// x * 0 keeps the sign of a denormal; NaN and infinity fail the ordered compare.
+	const auto magnitude = state.builder.AllocateId();
+	state.builder.AddFunction(spv::OpExtInst, TypeF32(state), magnitude, GlslStd450(state),
+	                          GLSLstd450FAbs, value);
+	const auto denormal = Binary(state, spv::OpFOrdLessThan, TypeBool(state), magnitude,
+	                             ConstantF32(state, 0x00800000u));
+	return Select(state, TypeF32(state), denormal,
+	              Binary(state, spv::OpFMul, TypeF32(state), value, ConstantF32(state, 0)), value);
 }
 
 uint32_t EmitTrigCycleF32(EmitterState& state, uint32_t src, bool preserve_signed_zero) {
-	const auto fract        = state.builder.AllocateId();
-	const auto bits         = state.builder.AllocateId();
-	const auto abs_bits     = state.builder.AllocateId();
-	const auto large        = state.builder.AllocateId();
-	const auto finite       = state.builder.AllocateId();
-	const auto large_finite = state.builder.AllocateId();
-	const auto reduced      = state.builder.AllocateId();
+	// Finite values of magnitude 2^23 and above are integers, so fract already gives them +0.
+	const auto fract = state.builder.AllocateId();
 	state.builder.AddFunction(spv::OpExtInst, TypeF32(state), fract, GlslStd450(state),
 	                          GLSLstd450Fract, src);
-	state.builder.AddFunction(spv::OpBitcast, TypeU32(state), bits, src);
-	state.builder.AddFunction(spv::OpBitwiseAnd, TypeU32(state), abs_bits, bits,
-	                          ConstantU32(state, 0x7fffffffu));
-	state.builder.AddFunction(spv::OpUGreaterThanEqual, TypeBool(state), large, abs_bits,
-	                          ConstantU32(state, 0x4b000000u));
-	state.builder.AddFunction(spv::OpULessThan, TypeBool(state), finite, abs_bits,
-	                          ConstantU32(state, 0x7f800000u));
-	state.builder.AddFunction(spv::OpLogicalAnd, TypeBool(state), large_finite, large, finite);
-	state.builder.AddFunction(spv::OpSelect, TypeF32(state), reduced, large_finite,
-	                          ConstantF32(state, 0), fract);
 	if (!preserve_signed_zero) {
-		return reduced;
+		return fract;
 	}
-	const auto zero = state.builder.AllocateId();
-	const auto ret  = state.builder.AllocateId();
-	state.builder.AddFunction(spv::OpIEqual, TypeBool(state), zero, abs_bits,
-	                          ConstantU32(state, 0));
-	state.builder.AddFunction(spv::OpSelect, TypeF32(state), ret, zero, src, reduced);
-	return ret;
+	const auto abs_bits = Binary(state, spv::OpBitwiseAnd, TypeU32(state),
+	                             Unary(state, spv::OpBitcast, TypeU32(state), src),
+	                             ConstantU32(state, 0x7fffffffu));
+	return Select(state, TypeF32(state),
+	              Binary(state, spv::OpIEqual, TypeBool(state), abs_bits, ConstantU32(state, 0)),
+	              src, fract);
 }
 
 uint32_t EmitF16BitsToF32(EmitterState& state, uint32_t bits) {

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

@@ -114,51 +114,30 @@ uint32_t EmitBuiltinU32(EmitterState& state, IR::StageInputKind kind, uint32_t c
 
 uint32_t EmitDppWriteCondition(ValueEmitContext& ctx, const IR::DppMoveFlags& flags,
                                uint32_t exec) {
-	auto&      state      = ctx.state;
-	const auto lane       = EmitSubgroupLocalInvocationId(state);
-	const auto bank_shift = state.builder.AllocateId();
-	const auto row_shift  = state.builder.AllocateId();
-	const auto bank       = state.builder.AllocateId();
-	const auto row        = state.builder.AllocateId();
-	const auto bank_bit   = state.builder.AllocateId();
-	const auto row_bit    = state.builder.AllocateId();
-	const auto bank_hit   = state.builder.AllocateId();
-	const auto row_hit    = state.builder.AllocateId();
-	const auto bank_ok    = state.builder.AllocateId();
-	const auto row_ok     = state.builder.AllocateId();
-	const auto masks_ok   = state.builder.AllocateId();
-	state.builder.AddFunction(spv::OpShiftRightLogical, TypeU32(state), bank_shift, lane,
-	                          ConstantU32(state, 2));
-	state.builder.AddFunction(spv::OpShiftRightLogical, TypeU32(state), row_shift, lane,
-	                          ConstantU32(state, 4));
-	state.builder.AddFunction(spv::OpBitwiseAnd, TypeU32(state), bank, bank_shift,
-	                          ConstantU32(state, 3));
-	state.builder.AddFunction(spv::OpBitwiseAnd, TypeU32(state), row, row_shift,
-	                          ConstantU32(state, 3));
-	state.builder.AddFunction(spv::OpShiftLeftLogical, TypeU32(state), bank_bit,
-	                          ConstantU32(state, 1), bank);
-	state.builder.AddFunction(spv::OpShiftLeftLogical, TypeU32(state), row_bit,
-	                          ConstantU32(state, 1), row);
-	state.builder.AddFunction(spv::OpBitwiseAnd, TypeU32(state), bank_hit,
-	                          ConstantU32(state, flags.bank_mask), bank_bit);
-	state.builder.AddFunction(spv::OpBitwiseAnd, TypeU32(state), row_hit,
-	                          ConstantU32(state, flags.row_mask), row_bit);
-	state.builder.AddFunction(spv::OpINotEqual, TypeBool(state), bank_ok, bank_hit,
-	                          ConstantU32(state, 0));
-	state.builder.AddFunction(spv::OpINotEqual, TypeBool(state), row_ok, row_hit,
-	                          ConstantU32(state, 0));
-	state.builder.AddFunction(spv::OpLogicalAnd, TypeBool(state), masks_ok, bank_ok, row_ok);
-	uint32_t writable = masks_ok;
-	if (!flags.bound_control) {
-		const auto target  = EmitDppTargetLane(state, flags);
-		const auto bounded = state.builder.AllocateId();
-		state.builder.AddFunction(spv::OpLogicalAnd, TypeBool(state), bounded, writable,
-		                          target.valid);
-		writable = bounded;
+	auto&    state    = ctx.state;
+	uint32_t writable = exec;
+	uint32_t lane     = 0;
+	// Each mask bit enables writes to one bank (4 lanes) or row (16 lanes); a full mask enables all.
+	const auto enable = [&](uint32_t mask, uint32_t group_shift) {
+		if (mask == 0xfu) return;
+		lane = lane != 0 ? lane : EmitSubgroupLocalInvocationId(state);
+		const auto group = EmitNative<spv::OpBitFieldUExtract, IR::Type::U32>(
+		    state, lane, ConstantU32(state, group_shift), ConstantU32(state, 2));
+		const auto bit = EmitNative<spv::OpBitFieldUExtract, IR::Type::U32>(
+		    state, ConstantU32(state, mask), group, ConstantU32(state, 1));
+		writable = Binary(state, spv::OpLogicalAnd, TypeBool(state), writable,
+		                  Binary(state, spv::OpINotEqual, TypeBool(state), bit, ConstantU32(state, 0)));
+	};
+	enable(flags.bank_mask, 2u);
+	enable(flags.row_mask, 4u);
+	// Only row shifts (0x101-0x10f, 0x111-0x11f) can vacate a lane; other controls always read one.
+	const bool row_shift = !flags.dpp8 && flags.control > 0x100u && flags.control < 0x120u &&
+	                       flags.control != 0x110u;
+	if (!flags.bound_control && row_shift) {
+		writable = Binary(state, spv::OpLogicalAnd, TypeBool(state), writable,
+		                  EmitDppTargetLane(state, flags).valid);
 	}
-	const auto result = state.builder.AllocateId();
-	state.builder.AddFunction(spv::OpLogicalAnd, TypeBool(state), result, exec, writable);
-	return result;
+	return writable;
 }
 
 uint32_t EmitAttribute(EmitterState& state, uint32_t attr, uint32_t chan) {

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

@@ -4,25 +4,16 @@ namespace Libs::Graphics::ShaderRecompiler::Spirv::Emitter {
 
 DppTargetLane EmitDppPermTargetLane(EmitterState& state, uint32_t subid, uint32_t control,
                                     uint32_t lane_bits) {
-	const auto lane_mask  = (1u << lane_bits) - 1u;
-	const auto group_base = state.builder.AllocateId();
-	const auto lane      = state.builder.AllocateId();
-	const auto shift     = state.builder.AllocateId();
-	const auto selected0 = state.builder.AllocateId();
-	const auto selected  = state.builder.AllocateId();
-	const auto target    = state.builder.AllocateId();
-	state.builder.AddFunction(spv::OpBitwiseAnd, TypeU32(state), group_base, subid,
-	                          ConstantU32(state, ~lane_mask));
-	state.builder.AddFunction(spv::OpBitwiseAnd, TypeU32(state), lane, subid,
-	                          ConstantU32(state, lane_mask));
-	state.builder.AddFunction(spv::OpIMul, TypeU32(state), shift, lane,
-	                          ConstantU32(state, lane_bits));
-	state.builder.AddFunction(spv::OpShiftRightLogical, TypeU32(state), selected0,
-	                          ConstantU32(state, control), shift);
-	state.builder.AddFunction(spv::OpBitwiseAnd, TypeU32(state), selected, selected0,
-	                          ConstantU32(state, lane_mask));
-	state.builder.AddFunction(spv::OpBitwiseOr, TypeU32(state), target, group_base, selected);
-	return {target, ConstantBool(state, true)};
+	// Each lane of a group takes its lane_bits-wide source index from control.
+	const auto lane = Binary(state, spv::OpBitwiseAnd, TypeU32(state), subid,
+	                         ConstantU32(state, (1u << lane_bits) - 1u));
+	const auto selected = EmitNative<spv::OpBitFieldUExtract, IR::Type::U32>(
+	    state, ConstantU32(state, control),
+	    Binary(state, spv::OpIMul, TypeU32(state), lane, ConstantU32(state, lane_bits)),
+	    ConstantU32(state, lane_bits));
+	return {EmitNative<spv::OpBitFieldInsert, IR::Type::U32>(
+	            state, subid, selected, ConstantU32(state, 0), ConstantU32(state, lane_bits)),
+	        ConstantBool(state, true)};
 }
 
 DppTargetLane EmitDppRowShiftTargetLane(EmitterState& state, uint32_t subid, uint32_t amount,
@@ -53,42 +44,19 @@ DppTargetLane EmitDppRowShiftTargetLane(EmitterState& state, uint32_t subid, uin
 
 DppTargetLane EmitDppRowRotateRightTargetLane(EmitterState& state, uint32_t subid,
                                               uint32_t amount) {
-	const auto row      = state.builder.AllocateId();
-	const auto lane     = state.builder.AllocateId();
-	const auto in_high  = state.builder.AllocateId();
-	const auto minus    = state.builder.AllocateId();
-	const auto plus     = state.builder.AllocateId();
-	const auto selected = state.builder.AllocateId();
-	const auto target   = state.builder.AllocateId();
-	state.builder.AddFunction(spv::OpBitwiseAnd, TypeU32(state), row, subid,
-	                          ConstantU32(state, 0xfffffff0u));
-	state.builder.AddFunction(spv::OpBitwiseAnd, TypeU32(state), lane, subid,
-	                          ConstantU32(state, 15));
-	state.builder.AddFunction(spv::OpUGreaterThanEqual, TypeBool(state), in_high, lane,
-	                          ConstantU32(state, amount));
-	state.builder.AddFunction(spv::OpISub, TypeU32(state), minus, lane, ConstantU32(state, amount));
-	state.builder.AddFunction(spv::OpIAdd, TypeU32(state), plus, lane,
-	                          ConstantU32(state, 16u - amount));
-	state.builder.AddFunction(spv::OpSelect, TypeU32(state), selected, in_high, minus, plus);
-	state.builder.AddFunction(spv::OpBitwiseOr, TypeU32(state), target, row, selected);
-	return {target, ConstantBool(state, true)};
+	// Rotation wraps within the row, so only the low four lane bits change.
+	const auto rotated = Binary(state, spv::OpISub, TypeU32(state), subid, ConstantU32(state, amount));
+	return {EmitNative<spv::OpBitFieldInsert, IR::Type::U32>(state, subid, rotated,
+	                                                        ConstantU32(state, 0),
+	                                                        ConstantU32(state, 4)),
+	        ConstantBool(state, true)};
 }
 
 DppTargetLane EmitDppMirrorTargetLane(EmitterState& state, uint32_t subid, bool half_row) {
-	const auto base_mask = half_row ? 0xfffffff8u : 0xfffffff0u;
-	const auto lane_mask = half_row ? 7u : 15u;
-	const auto base      = state.builder.AllocateId();
-	const auto lane      = state.builder.AllocateId();
-	const auto mirrored  = state.builder.AllocateId();
-	const auto target    = state.builder.AllocateId();
-	state.builder.AddFunction(spv::OpBitwiseAnd, TypeU32(state), base, subid,
-	                          ConstantU32(state, base_mask));
-	state.builder.AddFunction(spv::OpBitwiseAnd, TypeU32(state), lane, subid,
-	                          ConstantU32(state, lane_mask));
-	state.builder.AddFunction(spv::OpISub, TypeU32(state), mirrored, ConstantU32(state, lane_mask),
-	                          lane);
-	state.builder.AddFunction(spv::OpBitwiseOr, TypeU32(state), target, base, mirrored);
-	return {target, ConstantBool(state, true)};
+	// Mirroring lane i of a row or half row reads lane (size - 1 - i), which flips its low bits.
+	return {Binary(state, spv::OpBitwiseXor, TypeU32(state), subid,
+	               ConstantU32(state, half_row ? 7u : 15u)),
+	        ConstantBool(state, true)};
 }
 
 DppTargetLane EmitDppTargetLane(EmitterState& state, const IR::DppMoveFlags& flags) {
@@ -117,13 +85,10 @@ DppTargetLane EmitDppTargetLane(EmitterState& state, const IR::DppMoveFlags& fla
 	}
 	// row_share:N reads lane N of the caller's row of 16.
 	if (control >= 0x150u && control <= 0x15fu) {
-		const auto row    = state.builder.AllocateId();
-		const auto target = state.builder.AllocateId();
-		state.builder.AddFunction(spv::OpBitwiseAnd, TypeU32(state), row, subid,
-		                          ConstantU32(state, 0xfffffff0u));
-		state.builder.AddFunction(spv::OpBitwiseOr, TypeU32(state), target, row,
-		                          ConstantU32(state, control & 0xfu));
-		return {target, ConstantBool(state, true)};
+		return {EmitNative<spv::OpBitFieldInsert, IR::Type::U32>(
+		            state, subid, ConstantU32(state, control & 0xfu), ConstantU32(state, 0),
+		            ConstantU32(state, 4)),
+		        ConstantBool(state, true)};
 	}
 	if (control >= 0x160u && control <= 0x16fu) {
 		const auto target = state.builder.AllocateId();

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

@@ -1057,8 +1057,8 @@ void EmitImage(ValueEmitContext& ctx, const IR::Inst& inst) {
 			operands[1] = AddressF32(ctx, mem, *address, layout.bias);
 			operand_count = 2;
 		}
-		if (layout.offset != NoImageComponent &&
-		    !offset_value.IsImmediate() && !state.maintenance8_enabled) {
+		if (layout.offset != NoImageComponent && !offset_value.IsImmediate() &&
+		    !state.profile.maintenance8) {
 			ctx.Fail(inst, "requires maintenance8 for a dynamic sample texel offset");
 		}
 		const auto samplers = LoadImageSamplers(ctx, inst);

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

@@ -62,8 +62,6 @@ EMIT_NATIVE(IMul32, OpIMul, U32, uint32_t, uint32_t)
 EMIT_NATIVE(UDiv32, OpUDiv, U32, uint32_t, uint32_t)
 EMIT_NATIVE(IAdd64, OpIAdd, U64, uint32_t, uint32_t)
 EMIT_NATIVE(ISub64, OpISub, U64, uint32_t, uint32_t)
-EMIT_NATIVE(IMul64, OpIMul, U64, uint32_t, uint32_t)
-EMIT_NATIVE(IAddCarry32, OpIAddCarry, U32x2, uint32_t, uint32_t)
 uint32_t EmitSMulHi(EmitterState& state, uint32_t arg0, uint32_t arg1);
 uint32_t EmitUMulHi(EmitterState& state, uint32_t arg0, uint32_t arg1);
 inline constexpr auto EmitIAbs32 = EmitGlsl<GLSLstd450SAbs, IR::Type::U32, uint32_t>;
@@ -92,10 +90,8 @@ EMIT_NATIVE(BitwiseNot32, OpNot, U32, uint32_t)
 EMIT_NATIVE(BitwiseAnd64, OpBitwiseAnd, U64, uint32_t, uint32_t)
 EMIT_NATIVE(BitReverse32, OpBitReverse, U32, uint32_t)
 EMIT_NATIVE(BitCount32, OpBitCount, U32, uint32_t)
-uint32_t EmitBitCount64(EmitterState& state, uint32_t arg0);
 inline constexpr auto EmitFindILsb32 = EmitGlsl<GLSLstd450FindILsb, IR::Type::U32, uint32_t>;
 inline constexpr auto EmitFindUMsb32 = EmitGlsl<GLSLstd450FindUMsb, IR::Type::U32, uint32_t>;
-uint32_t EmitFindUMsb64(EmitterState& state, uint32_t arg0);
 inline constexpr auto EmitSMin32 = EmitGlsl<GLSLstd450SMin, IR::Type::U32, uint32_t, uint32_t>;
 inline constexpr auto EmitSMax32 = EmitGlsl<GLSLstd450SMax, IR::Type::U32, uint32_t, uint32_t>;
 inline constexpr auto EmitUMin32 = EmitGlsl<GLSLstd450UMin, IR::Type::U32, uint32_t, uint32_t>;
@@ -227,10 +223,7 @@ uint32_t              EmitGetUserData(EmitterState& state, IR::ScalarReg reg);
 uint32_t              EmitGetDispatchThreadExtent(ValueEmitContext& ctx, const IR::Inst& inst);
 uint32_t              EmitGetBuiltin(ValueEmitContext& ctx, IR::Value kind, IR::Value index);
 uint32_t              EmitUndefU1(EmitterState& state, const IR::Inst& inst);
-inline constexpr auto EmitUndefU8  = EmitUndefU1;
-inline constexpr auto EmitUndefU16 = EmitUndefU1;
 inline constexpr auto EmitUndefU32 = EmitUndefU1;
-inline constexpr auto EmitUndefU64 = EmitUndefU1;
 uint32_t              EmitDppMoveU32(ValueEmitContext& ctx, const IR::Inst& inst);
 uint32_t              EmitDppUpdateU32(ValueEmitContext& ctx, const IR::Inst& inst);
 uint32_t              EmitWqmU64(EmitterState& state, uint32_t value);

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

@@ -5,6 +5,7 @@
 #include "common/stringUtils.h"
 #include "graphics/shader/recompiler/BufferFormat.h"
 #include "graphics/shader/recompiler/backend/spirv/SpirvBuilder.h"
+#include "graphics/shader/recompiler/backend/spirv/SpirvEmitter.h"
 #include "graphics/shader/recompiler/ir/ShaderIR.h"
 #include "graphics/shader/recompiler/ir/passes/ResourceMaterialization.h"
 
@@ -73,9 +74,9 @@ struct ImageDefinition {
 };
 
 struct EmitterState {
-	EmitterState(IR::Program& program_, ShaderStageInputInfo input_info_, bool maintenance8_enabled_)
+	EmitterState(IR::Program& program_, ShaderStageInputInfo input_info_, const Profile& profile_)
 	    : builder(program_.stage == ShaderType::Mesh ? 0x00010400u : 0x00010300u),
-	      program(program_), input_info(input_info_), maintenance8_enabled(maintenance8_enabled_) {
+	      program(program_), input_info(input_info_), profile(profile_) {
 		if (ShaderWorkgroupInput(program.stage, input_info) != nullptr) {
 			lds_storage_class = spv::StorageClassWorkgroup;
 		}
@@ -88,7 +89,7 @@ struct EmitterState {
 	Builder                                          builder;
 	IR::Program&                                     program;
 	ShaderStageInputInfo                             input_info;
-	bool                                            maintenance8_enabled;
+	const Profile&                                      profile;
 	uint32_t                                        debug_source = 0;
 	uint32_t                                        void_type = 0;
 	uint32_t                                        bool_type = 0;
@@ -285,7 +286,6 @@ struct ImageSampleLayout {
 };
 
 struct F32Class {
-	uint32_t bits = 0;
 	uint32_t nan  = 0;
 	uint32_t zero = 0;
 };
@@ -510,8 +510,6 @@ inline constexpr auto EmitLogicalNotBool = EmitNative<spv::OpLogicalNot, IR::Typ
 
 F32Class EmitClassifyF32Bits(EmitterState& state, uint32_t bits);
 
-F32Class EmitClassifyF32(EmitterState& state, uint32_t value);
-
 uint32_t EmitClassMaskBitMatch(EmitterState& state, uint32_t mask, uint32_t bit,
                                uint32_t class_match);
 
@@ -520,8 +518,6 @@ uint32_t EmitClassMaskF16(EmitterState& state, uint32_t bits, uint32_t mask);
 
 uint32_t EmitMinMaxF32Value(EmitterState& state, uint32_t lhs, uint32_t rhs, bool max_value);
 
-inline constexpr auto EmitTruncF32Value = EmitGlsl<GLSLstd450Trunc, IR::Type::F32, uint32_t>;
-
 uint32_t EmitFlushF32DenormToSignedZero(EmitterState& state, uint32_t value);
 
 uint32_t EmitTrigCycleF32(EmitterState& state, uint32_t src, bool preserve_signed_zero);
@@ -582,6 +578,10 @@ auto EmitIndexSwitch(EmitterState& state, uint32_t index, uint32_t count,
 
 template <typename Fn>
 void EmitIfCondition(EmitterState& state, uint32_t condition, Fn&& fn) {
+	if (condition == ConstantBool(state, true)) {
+		fn();
+		return;
+	}
 	const auto then_label  = state.builder.AllocateId();
 	const auto merge_label = state.builder.AllocateId();
 	state.builder.AddFunction(spv::OpSelectionMerge, merge_label, spv::SelectionControlMaskNone);
@@ -595,6 +595,7 @@ void EmitIfCondition(EmitterState& state, uint32_t condition, Fn&& fn) {
 template <typename Fn>
 uint32_t EmitValueOrDefaultIfCondition(EmitterState& state, uint32_t condition, uint32_t type,
                                        uint32_t default_value, Fn&& fn) {
+	if (condition == ConstantBool(state, true)) return fn();
 	const auto then_label  = state.builder.AllocateId();
 	const auto header     = state.current_label;
 	const auto merge_label = state.builder.AllocateId();

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

@@ -1605,9 +1605,12 @@ void EmitReadConstBuffer(ValueEmitContext& ctx, const IR::Inst& inst) {
 	}
 	auto& state        = ctx.state;
 	mem.kind           = IR::ResourceKind::ScalarBuffer;
-	auto index = Binary(state, spv::OpShiftRightLogical, TypeU32(state), ctx.Arg(inst, 1),
-	                    ConstantU32(state, 2));
-	if (mem.offset >= 4u) {
+	const auto offset = inst.Arg(1).Resolve();
+	auto       index  = offset.IsImmediate()
+	                        ? ConstantU32(state, (offset.U32() >> 2u) + (mem.offset >> 2u))
+	                        : Binary(state, spv::OpShiftRightLogical, TypeU32(state),
+	                                 ctx.Arg(inst, 1), ConstantU32(state, 2));
+	if (!offset.IsImmediate() && mem.offset >= 4u) {
 		index = Binary(state, spv::OpIAdd, TypeU32(state), index, ConstantU32(state, mem.offset >> 2u));
 	}
 	const auto access    = PrepareMemoryResourceAccess(state, mem);

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

@@ -16,10 +16,8 @@ namespace {
 void EmitKillIfBoolFalse(EmitterState& state, uint32_t active) {
 	const auto kill_label  = state.builder.AllocateId();
 	const auto merge_label = state.builder.AllocateId();
-	const auto inactive    = state.builder.AllocateId();
-	state.builder.AddFunction(spv::OpLogicalNot, TypeBool(state), inactive, active);
 	state.builder.AddFunction(spv::OpSelectionMerge, merge_label, spv::SelectionControlMaskNone);
-	state.builder.AddFunction(spv::OpBranchConditional, inactive, kill_label, merge_label);
+	state.builder.AddFunction(spv::OpBranchConditional, active, merge_label, kill_label);
 	EmitLabel(state, kill_label);
 	state.builder.AddFunction(spv::OpKill);
 	EmitLabel(state, merge_label);
@@ -233,14 +231,17 @@ void Invoke(Return (*emit)(Context&, Args...), ValueEmitContext& ctx, const IR::
 }
 
 void EmitDirectInstruction(ValueEmitContext& ctx, const IR::Inst& inst) {
-	if (ctx.half != 0 && (inst.GetOpcode() == IR::ValueOpcode::Ballot ||
-	                     inst.GetOpcode() == IR::ValueOpcode::ReadFirstLane ||
-	                     inst.GetOpcode() == IR::ValueOpcode::SubgroupUMinU32)) {
-		// These operations combine both emulated halves into one whole-wave result.
+	using Op = IR::ValueOpcode;
+	const auto op = inst.GetOpcode();
+	if (ctx.half != 0 &&
+	    (op == Op::Ballot || op == Op::ReadFirstLane || op == Op::SubgroupUMinU32 ||
+	     op == Op::ReadConstBuffer ||
+	     (op == Op::LoadAddressU32 && ctx.Memory(inst).kind == IR::ResourceKind::ScalarAddress))) {
+		// Whole-wave results and SMEM loads (SGPR operands, no EXEC) are identical in both halves.
 		ctx.Define(inst, ctx.other_half->Result(inst));
 		return;
 	}
-	switch (inst.GetOpcode()) {
+	switch (op) {
 #define VALUE_OPCODE(name, ...)                                                                    \
 	case IR::ValueOpcode::name: return Invoke(Emit##name, ctx, inst);
 #include "graphics/shader/recompiler/ir/opcodes/ValueOpcodes.inc"
@@ -492,29 +493,24 @@ uint32_t ValueEmitContext::Ballot(IR::Value predicate) {
 }
 
 uint32_t ValueEmitContext::FirstLane(uint32_t ballot) {
-	if (other_half == nullptr) {
-		const auto result = state.builder.AllocateId();
-		state.builder.AddFunction(spv::OpGroupNonUniformBallotFindLSB, TypeU32(state), result,
-		                          ConstantU32(state, spv::ScopeSubgroup), ballot);
-		return result;
+	// An empty ballot (EXEC zero) gives lane 0, which V_READFIRSTLANE reads on hardware. FindILsb
+	// returns -1 for zero, unlike BallotFindLSB, whose result is then undefined.
+	const auto first = [&](uint32_t word) {
+		const auto bits = state.builder.AllocateId();
+		state.builder.AddFunction(spv::OpCompositeExtract, TypeU32(state), bits, ballot, word);
+		return std::pair {bits, EmitGlsl<GLSLstd450FindILsb, IR::Type::U32>(state, bits)};
+	};
+	const auto [low, low_first] = first(0);
+	if (state.program.wave_size == 32u) {
+		return EmitGlsl<GLSLstd450SMax, IR::Type::U32>(state, low_first, ConstantU32(state, 0));
 	}
-	const auto low        = state.builder.AllocateId();
-	const auto high       = state.builder.AllocateId();
-	const auto low_first  = state.builder.AllocateId();
-	const auto high_first = state.builder.AllocateId();
-	const auto low_active = state.builder.AllocateId();
-	const auto result     = state.builder.AllocateId();
-	state.builder.AddFunction(spv::OpCompositeExtract, TypeU32(state), low, ballot, 0);
-	state.builder.AddFunction(spv::OpCompositeExtract, TypeU32(state), high, ballot, 1);
-	state.builder.AddFunction(spv::OpExtInst, TypeU32(state), low_first, GlslStd450(state),
-	                          GLSLstd450FindILsb, low);
-	state.builder.AddFunction(spv::OpExtInst, TypeU32(state), high_first, GlslStd450(state),
-	                          GLSLstd450FindILsb, high);
-	state.builder.AddFunction(spv::OpINotEqual, TypeBool(state), low_active, low,
-	                          ConstantU32(state, 0));
-	state.builder.AddFunction(spv::OpSelect, TypeU32(state), result, low_active, low_first,
-	                          EmitAddU32(state, high_first, ConstantU32(state, 32)));
-	return result;
+	const auto high_first = first(1).second;
+	const auto high_lane  = EmitGlsl<GLSLstd450SMax, IR::Type::U32>(
+	    state, EmitBinaryU32(state, spv::OpBitwiseOr, high_first, ConstantU32(state, 32)),
+	    ConstantU32(state, 0));
+	return Select(state, TypeU32(state),
+	              Binary(state, spv::OpINotEqual, TypeBool(state), low, ConstantU32(state, 0)),
+	              low_first, high_lane);
 }
 
 uint32_t ValueEmitContext::Shuffle(const IR::Inst& inst, size_t index, uint32_t lane) {

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

@@ -56,8 +56,10 @@ uint32_t TessellationPointer(ValueEmitContext& ctx, const IR::Inst& inst) {
 		state.builder.AddFunction(spv::OpAccessChain, type, pointer, variable, attribute,
 		                          component);
 	} else {
+		// An HS writes only its own control point but may read any other.
 		const auto vertex =
-		    kind == Attribute::ControlOutput
+		    kind == Attribute::ControlOutput &&
+		            inst.GetOpcode() == IR::ValueOpcode::SetTessellationAttribute
 		        ? EmitLaneId(state)
 		        : EmitBinaryU32(state, spv::OpUDiv, address, ConstantU32(state, stride));
 		state.builder.AddFunction(spv::OpAccessChain, type, pointer, variable, vertex, attribute,
@@ -126,7 +128,9 @@ void DefineTessellationInterfaces(EmitterState& state) {
 
 void DefineTessellationExecutionModes(EmitterState& state) {
 	const auto& tess = state.input_info.vertex->tess;
-	EXIT_NOT_IMPLEMENTED(tess.domain != 1u || tess.partitioning != 2u ||
+	// 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);
 	state.builder.RequireCapability(spv::CapabilityTessellation);
 	if (state.program.stage == ShaderType::TessellationControl) {
@@ -134,7 +138,9 @@ void DefineTessellationExecutionModes(EmitterState& state) {
 		                               tess.output_control_points);
 	} else {
 		state.builder.AddExecutionMode(state.main_func, spv::ExecutionModeTriangles);
-		state.builder.AddExecutionMode(state.main_func, spv::ExecutionModeSpacingFractionalOdd);
+		state.builder.AddExecutionMode(
+		    state.main_func, tess.partitioning == 0u ? spv::ExecutionModeSpacingEqual
+		                                             : spv::ExecutionModeSpacingFractionalOdd);
 		state.builder.AddExecutionMode(state.main_func, spv::ExecutionModeVertexOrderCw);
 	}
 }

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

@@ -82,19 +82,17 @@ void Translator::S_SAVEEXEC(const Decoder::Instruction& inst, IR::ValueOpcode op
 }
 
 void Translator::ADD_U32(const Decoder::Instruction& inst, bool vector, bool use_carry_in) {
-	const auto zero     = IR::U32(IR::Value(0u));
-	const auto lhs      = ReadU32(inst.src0);
-	const auto rhs      = ReadU32(inst.src1);
-	const auto carry_in = !use_carry_in ? zero
-	                      : vector      ? ConditionBit(inst.src2)
-	                                    : ConditionBit(ConditionOperand(Decoder::OperandKind::Scc));
-	const auto add0     = ir.Emit(IR::ValueOpcode::IAddCarry32, {lhs, rhs});
-	const auto partial  = ir.CompositeExtract(add0, 0);
-	const auto carry0   = ir.CompositeExtract(add0, 1);
-	const auto add1     = ir.Emit(IR::ValueOpcode::IAddCarry32, {partial, carry_in});
-	const auto result   = ir.CompositeExtract(add1, 0);
-	const auto carry1   = ir.CompositeExtract(add1, 1);
-	const auto carry    = ir.INotEqual(ir.BitwiseOr(carry0, carry1), zero);
+	const auto lhs     = ReadU32(inst.src0);
+	const auto rhs     = ReadU32(inst.src1);
+	const auto partial = ir.IAdd(lhs, rhs);
+	auto       result  = partial;
+	auto       carry   = ir.ULessThan(partial, lhs);
+	if (use_carry_in) {
+		// The carry-in is 0 or 1, so the second add carries out only when it wraps to zero.
+		result = ir.IAdd(partial, vector ? ConditionBit(inst.src2)
+		                                 : ConditionBit(ConditionOperand(Decoder::OperandKind::Scc)));
+		carry  = ir.LogicalOr(carry, ir.ULessThan(result, partial));
+	}
 	WriteOperand(DestinationOperand(inst), result);
 	if (vector) {
 		WriteMask(inst.dst2, ir.LogicalAnd(ir.GetExec(), carry));
@@ -142,29 +140,25 @@ void Translator::S_ABSDIFF_I32(const Decoder::Instruction& inst) {
 }
 
 void Translator::S_ADD_SUB_I32(const Decoder::Instruction& inst, bool subtract) {
-	const auto lhs      = ReadU32(inst.src0);
-	const auto rhs      = ReadU32(inst.src1);
-	const auto result   = subtract ? ir.ISub(lhs, rhs) : ir.IAdd(lhs, rhs);
-	const auto shift    = IR::U32(IR::Value(31u));
-	const auto lhs_sign = ir.ShiftRightLogical(lhs, shift);
-	const auto rhs_sign = ir.ShiftRightLogical(rhs, shift);
-	const auto out_sign = ir.ShiftRightLogical(result, shift);
-	const auto inputs = subtract ? ir.INotEqual(lhs_sign, rhs_sign) : ir.IEqual(lhs_sign, rhs_sign);
-	const auto changed = ir.INotEqual(lhs_sign, out_sign);
+	const auto lhs    = ReadU32(inst.src0);
+	const auto rhs    = ReadU32(inst.src1);
+	const auto result = subtract ? ir.ISub(lhs, rhs) : ir.IAdd(lhs, rhs);
+	// Signed overflow moves the result past lhs the wrong way for the sign of rhs.
+	const auto ge = [&](IR::Value a, IR::Value b) {
+		return IR::U1(ir.Emit(IR::ValueOpcode::SGreaterThanEqual32, {a, b}));
+	};
+	const auto moved = subtract ? ge(lhs, result) : ge(result, lhs);
 	WriteOperand(DestinationOperand(inst), result);
-	ir.SetScc(ir.LogicalAnd(inputs, changed));
+	ir.SetScc(IR::U1(ir.Emit(IR::ValueOpcode::LogicalXor, {moved, ge(rhs, IR::Value(0u))})));
 }
 
 void Translator::S_LSHL_ADD_U32(const Decoder::Instruction& inst, uint32_t shift_amount) {
-	const auto lhs           = ReadU32(inst.src0);
-	const auto shift         = IR::U32(IR::Value(shift_amount));
-	const auto rhs           = ReadU32(inst.src1);
-	const auto shifted       = ir.ShiftLeftLogical(lhs, shift);
-	const auto result        = ir.IAdd(shifted, rhs);
-	const auto add_carry     = ir.ULessThan(result, shifted);
-	const auto inverse_shift = ir.ISub(IR::U32(IR::Value(32u)), shift);
-	const auto shifted_out   = ir.ShiftRightLogical(lhs, inverse_shift);
-	const auto shift_carry   = ir.INotEqual(shifted_out, IR::U32(IR::Value(0u)));
+	const auto lhs         = ReadU32(inst.src0);
+	const auto rhs         = ReadU32(inst.src1);
+	const auto shifted     = ir.ShiftLeftLogical(lhs, IR::U32(IR::Value(shift_amount)));
+	const auto result      = ir.IAdd(shifted, rhs);
+	const auto add_carry   = ir.ULessThan(result, shifted);
+	const auto shift_carry = ir.UGreaterThan(lhs, IR::U32(IR::Value(~0u >> shift_amount)));
 	WriteOperand(DestinationOperand(inst), result);
 	ir.SetScc(ir.LogicalOr(add_carry, shift_carry));
 }

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

@@ -6,43 +6,10 @@ IR::F32 Translator::SelectF32(IR::U1 condition, IR::F32 true_value, IR::F32 fals
 	return IR::F32(ir.Emit(IR::ValueOpcode::SelectF32, {condition, true_value, false_value}));
 }
 
-IR::U32 Translator::ConvertF32ToU32Saturated(IR::F32 value, float upper_bound, float safe_upper,
-                                             uint32_t upper_result) {
-	const auto zero = IR::F32(IR::Value::F32(0.0f));
-	const auto nan  = IR::U1(ir.Emit(IR::ValueOpcode::FPIsNan32, {value}));
-	const auto low  = IR::U1(ir.Emit(IR::ValueOpcode::FPOrdLessThanEqual32, {value, zero}));
-	const auto high = IR::U1(
-	    ir.Emit(IR::ValueOpcode::FPOrdGreaterThanEqual32, {value, IR::Value::F32(upper_bound)}));
-	const auto truncated = IR::F32(ir.Emit(IR::ValueOpcode::FPTrunc32, {value}));
-	const auto safe_low  = SelectF32(ir.LogicalOr(nan, low), zero, truncated);
-	const auto safe      = SelectF32(high, IR::F32(IR::Value::F32(safe_upper)), safe_low);
-	const auto converted = IR::U32(ir.Emit(IR::ValueOpcode::ConvertU32F32, {safe}));
-	return ir.Select(high, IR::U32(IR::Value(upper_result)), converted);
-}
-
-IR::U32 Translator::ConvertF32ToI32Saturated(IR::F32 value, float lower_bound, float upper_bound,
-                                             float safe_upper, uint32_t lower_result,
-                                             uint32_t upper_result) {
-	const auto nan = IR::U1(ir.Emit(IR::ValueOpcode::FPIsNan32, {value}));
-	const auto low = IR::U1(
-	    ir.Emit(IR::ValueOpcode::FPOrdLessThanEqual32, {value, IR::Value::F32(lower_bound)}));
-	const auto high = IR::U1(
-	    ir.Emit(IR::ValueOpcode::FPOrdGreaterThanEqual32, {value, IR::Value::F32(upper_bound)}));
-	const auto truncated    = IR::F32(ir.Emit(IR::ValueOpcode::FPTrunc32, {value}));
-	const auto safe_low     = SelectF32(low, IR::F32(IR::Value::F32(lower_bound)), truncated);
-	const auto safe_high    = SelectF32(high, IR::F32(IR::Value::F32(safe_upper)), safe_low);
-	const auto safe         = SelectF32(nan, IR::F32(IR::Value::F32(0.0f)), safe_high);
-	const auto converted    = IR::U32(ir.Emit(IR::ValueOpcode::ConvertS32F32, {safe}));
-	const auto clamped_high = ir.Select(high, IR::U32(IR::Value(upper_result)), converted);
-	const auto clamped      = ir.Select(low, IR::U32(IR::Value(lower_result)), clamped_high);
-	return ir.Select(nan, IR::U32(IR::Value(0u)), clamped);
-}
-
 void Translator::V_CVT_F32_UBYTE(const Decoder::Instruction& inst, uint32_t byte_index) {
 	const auto source = ReadU32(SourceAt(inst, 0));
-	const auto byte =
-	    ir.BitwiseAnd(ir.ShiftRightLogical(source, IR::U32(IR::Value(byte_index * 8u))),
-	                  IR::U32(IR::Value(0xffu)));
+	const auto byte   = ir.Emit(IR::ValueOpcode::BitFieldUExtract,
+	                            {source, IR::Value(byte_index * 8u), IR::Value(8u)});
 	WriteOperand(DestinationOperand(inst), ir.Emit(IR::ValueOpcode::ConvertF32U32, {byte}));
 }
 
@@ -58,15 +25,12 @@ void Translator::V_CVT_F32_I32(const Decoder::Instruction& inst) {
 
 void Translator::V_CVT_U32_F32(const Decoder::Instruction& inst) {
 	const auto value = IR::F32(ReadOperand(SourceAt(inst, 0), IR::Type::F32));
-	WriteOperand(DestinationOperand(inst),
-	             ConvertF32ToU32Saturated(value, 4294967296.0f, 4294967040.0f, 0xffffffffu));
+	WriteOperand(DestinationOperand(inst), ir.Emit(IR::ValueOpcode::ConvertU32F32, {value}));
 }
 
 void Translator::V_CVT_I32_F32(const Decoder::Instruction& inst) {
 	const auto value = IR::F32(ReadOperand(SourceAt(inst, 0), IR::Type::F32));
-	WriteOperand(DestinationOperand(inst),
-	             ConvertF32ToI32Saturated(value, -2147483648.0f, 2147483648.0f, 2147483520.0f,
-	                                      0x80000000u, 0x7fffffffu));
+	WriteOperand(DestinationOperand(inst), ir.Emit(IR::ValueOpcode::ConvertS32F32, {value}));
 }
 
 void Translator::V_CVT_F16_F32(const Decoder::Instruction& inst) {
@@ -86,31 +50,30 @@ void Translator::V_CVT_F16_16(const Decoder::Instruction& inst, bool signed_valu
 
 void Translator::V_CVT_16_F16(const Decoder::Instruction& inst, bool signed_value) {
 	const auto value = ReadF16AsF32(inst.src0);
+	// The 32-bit conversion already saturates and maps NaN to zero.
 	if (signed_value) {
-		const auto converted =
-		    ConvertF32ToI32Saturated(value, -32768.0f, 32768.0f, 32767.0f, 0xffff8000u, 0x7fffu);
-		Write16Bits(DestinationOperand(inst), ir.BitwiseAnd(converted, IR::U32(IR::Value(0xffffu))));
+		const auto converted = ir.Emit(IR::ValueOpcode::ConvertS32F32, {value});
+		const auto low = ir.Emit(IR::ValueOpcode::SMax32, {converted, IR::Value(0xffff8000u)});
+		Write16Bits(DestinationOperand(inst),
+		            IR::U32(ir.Emit(IR::ValueOpcode::SMin32, {low, IR::Value(0x7fffu)})));
 		return;
 	}
+	const auto converted = ir.Emit(IR::ValueOpcode::ConvertU32F32, {value});
 	Write16Bits(DestinationOperand(inst),
-	         ConvertF32ToU32Saturated(value, 65536.0f, 65535.0f, 0xffffu));
+	            IR::U32(ir.Emit(IR::ValueOpcode::UMin32, {converted, IR::Value(0xffffu)})));
 }
 
 void Translator::V_CVT_RPI_I32_F32(const Decoder::Instruction& inst) {
 	const auto source  = IR::F32(ReadOperand(SourceAt(inst, 0), IR::Type::F32));
 	const auto biased  = IR::F32(ir.Emit(IR::ValueOpcode::FPAdd32, {source, IR::Value::F32(0.5f)}));
 	const auto rounded = IR::F32(ir.Emit(IR::ValueOpcode::FPFloor32, {biased}));
-	WriteOperand(DestinationOperand(inst),
-	             ConvertF32ToI32Saturated(rounded, -2147483648.0f, 2147483648.0f, 2147483520.0f,
-	                                      0x80000000u, 0x7fffffffu));
+	WriteOperand(DestinationOperand(inst), ir.Emit(IR::ValueOpcode::ConvertS32F32, {rounded}));
 }
 
 void Translator::V_CVT_FLR_I32_F32(const Decoder::Instruction& inst) {
 	const auto source  = IR::F32(ReadOperand(SourceAt(inst, 0), IR::Type::F32));
 	const auto rounded = IR::F32(ir.Emit(IR::ValueOpcode::FPFloor32, {source}));
-	WriteOperand(DestinationOperand(inst),
-	             ConvertF32ToI32Saturated(rounded, -2147483648.0f, 2147483648.0f, 2147483520.0f,
-	                                      0x80000000u, 0x7fffffffu));
+	WriteOperand(DestinationOperand(inst), ir.Emit(IR::ValueOpcode::ConvertS32F32, {rounded}));
 }
 
 void Translator::V_FREXP_EXP_I32_F32(const Decoder::Instruction& inst) {
@@ -157,25 +120,24 @@ void Translator::V_CVT_PKNORM_F32(const Decoder::Instruction& inst, IR::ValueOpc
 
 void Translator::V_CVT_PK_U8_F32(const Decoder::Instruction& inst) {
 	const auto source = IR::F32(ReadOperand(SourceAt(inst, 0), IR::Type::F32));
-	const auto byte   = ConvertF32ToU32Saturated(source, 255.0f, 255.0f, 255u);
+	const auto byte =
+	    IR::U32(ir.Emit(IR::ValueOpcode::UMin32,
+		                {ir.Emit(IR::ValueOpcode::ConvertU32F32, {source}), IR::Value(255u)}));
 	const auto index  = ir.BitwiseAnd(ReadU32(SourceAt(inst, 1)), IR::U32(IR::Value(3u)));
 	const auto shift  = ir.ShiftLeftLogical(index, IR::U32(IR::Value(3u)));
-	const auto mask   = ir.ShiftLeftLogical(IR::U32(IR::Value(0xffu)), shift);
-	const auto base   = ir.BitwiseAnd(ReadU32(SourceAt(inst, 2)), ir.BitwiseNot(mask));
-	WriteOperand(DestinationOperand(inst), ir.BitwiseOr(base, ir.ShiftLeftLogical(byte, shift)));
+	WriteOperand(DestinationOperand(inst),
+	             ir.Emit(IR::ValueOpcode::BitFieldInsert,
+	                     {ReadU32(SourceAt(inst, 2)), byte, shift, IR::Value(8u)}));
 }
 
 void Translator::V_PACK_B32_F16(const Decoder::Instruction& inst) {
-	const auto low = Read16LaneBits(inst.src0, false);
-	const auto high =
-	    ir.ShiftLeftLogical(Read16LaneBits(inst.src1, false), IR::U32(IR::Value(16u)));
-	WriteOperand(DestinationOperand(inst), ir.BitwiseOr(low, high));
+	WriteOperand(DestinationOperand(inst),
+	             PackU16Lanes(Read16LaneBits(inst.src0, false), Read16LaneBits(inst.src1, false)));
 }
 
 IR::U32 Translator::PackU16Lanes(IR::U32 low, IR::U32 high) {
-	const auto mask = IR::U32(IR::Value(0xffffu));
-	return ir.BitwiseOr(ir.BitwiseAnd(low, mask),
-	                    ir.ShiftLeftLogical(ir.BitwiseAnd(high, mask), IR::U32(IR::Value(16u))));
+	return IR::U32(
+	    ir.Emit(IR::ValueOpcode::BitFieldInsert, {low, high, IR::Value(16u), IR::Value(16u)}));
 }
 
 } // namespace Libs::Graphics::ShaderRecompiler::Frontend

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

@@ -222,13 +222,10 @@ void Translator::FloatCube(const Decoder::Instruction& inst, uint32_t result_kin
 			break;
 		case 1u: result = select_face(SelectF32(x_neg, z, nz), x, SelectF32(z_neg, nx, x)); break;
 		case 2u: result = select_face(ny, SelectF32(y_neg, nz, z), ny); break;
-		case 3u: {
-			const auto two = IR::F32(IR::Value::F32(2.0f));
-			result         = select_face(IR::F32(ir.Emit(IR::ValueOpcode::FPMul32, {x, two})),
-			                             IR::F32(ir.Emit(IR::ValueOpcode::FPMul32, {y, two})),
-			                             IR::F32(ir.Emit(IR::ValueOpcode::FPMul32, {z, two})));
+		case 3u:
+			result = IR::F32(
+			    ir.Emit(IR::ValueOpcode::FPMul32, {select_face(x, y, z), IR::Value::F32(2.0f)}));
 			break;
-		}
 		default: EXIT("invalid cube result kind");
 	}
 	WriteOperand(DestinationOperand(inst), result);

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

@@ -9,7 +9,7 @@ void Translator::Integer16Shift(const Decoder::Instruction& inst, IR::ValueOpcod
 	const auto value  = ReadU16AsU32(inst.src1, arithmetic);
 	const auto count  = ir.BitwiseAnd(ReadU16AsU32(inst.src0, false), IR::U32(IR::Value(15u)));
 	const auto result = IR::U32(ir.Emit(opcode, {value, count}));
-	Write16Bits(DestinationOperand(inst), ir.BitwiseAnd(result, IR::U32(IR::Value(0xffffu))));
+	Write16Bits(DestinationOperand(inst), result);
 }
 
 void Translator::Integer16Binary(const Decoder::Instruction& inst, IR::ValueOpcode opcode,
@@ -17,7 +17,7 @@ void Translator::Integer16Binary(const Decoder::Instruction& inst, IR::ValueOpco
 	const auto lhs    = ReadU16AsU32(inst.src0, sign);
 	const auto rhs    = ReadU16AsU32(inst.src1, sign);
 	const auto result = IR::U32(ir.Emit(opcode, {lhs, rhs}));
-	Write16Bits(DestinationOperand(inst), ir.BitwiseAnd(result, IR::U32(IR::Value(0xffffu))));
+	Write16Bits(DestinationOperand(inst), result);
 }
 
 void Translator::V_MAD_I16(const Decoder::Instruction& inst) {
@@ -28,7 +28,7 @@ void Translator::V_MAD_I16(const Decoder::Instruction& inst) {
 		result = IR::U32(ir.Emit(IR::ValueOpcode::SMax32, {result, IR::Value(0xffff8000u)}));
 		result = IR::U32(ir.Emit(IR::ValueOpcode::SMin32, {result, IR::Value(0x7fffu)}));
 	}
-	Write16Bits(DestinationOperand(inst), ir.BitwiseAnd(result, IR::U32(IR::Value(0xffffu))));
+	Write16Bits(DestinationOperand(inst), result);
 }
 
 void Translator::Integer16Ternary(const Decoder::Instruction& inst, IR::ValueOpcode opcode,
@@ -221,19 +221,21 @@ void Translator::V_XOR3_B32(const Decoder::Instruction& inst) {
 	WriteOperand(DestinationOperand(inst), result);
 }
 
-void Translator::S_FF_I32_B64(const Decoder::Instruction& inst, bool find_zero) {
+void Translator::S_FF_I32_B64(const Decoder::Instruction& inst, bool find_zero, bool from_msb) {
 	auto source = ReadU32Pair(inst.src0);
-	if (find_zero) {
-		source[0] = ir.BitwiseNot(source[0]);
-		source[1] = ir.BitwiseNot(source[1]);
+	for (auto& word: source) {
+		word = find_zero ? ir.BitwiseNot(word) : word;
 	}
-	const auto low_lsb       = IR::U32(ir.Emit(IR::ValueOpcode::FindILsb32, {source[0]}));
-	const auto high_lsb      = IR::U32(ir.Emit(IR::ValueOpcode::FindILsb32, {source[1]}));
-	const auto high_position = ir.IAdd(high_lsb, IR::U32(IR::Value(32u)));
-	const auto result        = ir.Select(ir.INotEqual(source[0], IR::U32(IR::Value(0u))), low_lsb,
-	                                     ir.Select(ir.INotEqual(source[1], IR::U32(IR::Value(0u))),
-	                                               high_position, IR::U32(IR::Value(0xffffffffu))));
-	WriteOperand(DestinationOperand(inst), result);
+	if (from_msb) {
+		// The leading bit index is the trailing bit index of the bit-reversed 64-bit value.
+		source = {IR::U32(ir.Emit(IR::ValueOpcode::BitReverse32, {source[1]})),
+		          IR::U32(ir.Emit(IR::ValueOpcode::BitReverse32, {source[0]}))};
+	}
+	// FindILsb returns ~0 for zero, so OR-ing 32 keeps a missing high bit at ~0 (-1).
+	const auto low  = ir.Emit(IR::ValueOpcode::FindILsb32, {source[0]});
+	const auto high = ir.BitwiseOr(IR::U32(ir.Emit(IR::ValueOpcode::FindILsb32, {source[1]})),
+	                               IR::U32(IR::Value(32u)));
+	WriteOperand(DestinationOperand(inst), ir.Emit(IR::ValueOpcode::UMin32, {low, high}));
 }
 
 void Translator::V_FFBH_32(const Decoder::Instruction& inst, bool sign) {
@@ -248,17 +250,6 @@ void Translator::V_FFBH_32(const Decoder::Instruction& inst, bool sign) {
 	WriteOperand(DestinationOperand(inst), result);
 }
 
-void Translator::S_FLBIT_I32_B64(const Decoder::Instruction& inst) {
-	const auto source   = ReadU64(inst.src0);
-	const auto msb      = IR::U32(ir.Emit(IR::ValueOpcode::FindUMsb64, {source}));
-	const auto position = ir.ISub(IR::U32(IR::Value(63u)), msb);
-	const auto nonzero =
-	    IR::U1(ir.Emit(IR::ValueOpcode::INotEqual64,
-	                   {source, ir.ConstructU64(IR::U32(IR::Value(0u)), IR::U32(IR::Value(0u)))}));
-	const auto result = ir.Select(nonzero, position, IR::U32(IR::Value(0xffffffffu)));
-	WriteOperand(DestinationOperand(inst), result);
-}
-
 void Translator::Integer24(const Decoder::Instruction& inst, bool sign, bool addend) {
 	const auto extract24 = [&](IR::U32 value) {
 		return IR::U32(
@@ -277,7 +268,7 @@ void Translator::Integer24(const Decoder::Instruction& inst, bool sign, bool add
 void Translator::V_MAD_U64_U32(const Decoder::Instruction& inst) {
 	const auto lhs = ReadU32(inst.src0);
 	const auto rhs = ReadU32(inst.src1);
-	// Keep the 32x32 product in words: native widening IMul64 fails NVIDIA compilation
+	// Keep the 32x32 product in words: a native widening 64-bit multiply fails NVIDIA compilation
 	// in large vertex shaders (Astro's Playroom). The 64-bit addition stays native.
 	const auto product = ir.ConstructU64(
 	    ir.IMul(lhs, rhs), IR::U32(ir.Emit(IR::ValueOpcode::UMulHi, {lhs, rhs})));
@@ -336,8 +327,7 @@ void Translator::V_BCNT_U32_B32(const Decoder::Instruction& inst) {
 void Translator::V_MBCNT_U32_B32(const Decoder::Instruction& inst, bool low) {
 	const auto lane  = IR::U32(ir.Emit(IR::ValueOpcode::LaneId));
 	const auto local = ir.BitwiseAnd(lane, IR::U32(IR::Value(31u)));
-	const auto below =
-	    ir.ISub(ir.ShiftLeftLogical(IR::U32(IR::Value(1u)), local), IR::U32(IR::Value(1u)));
+	const auto below = RightMask32(local);
 	const auto high_lane =
 	    IR::U1(ir.Emit(IR::ValueOpcode::UGreaterThanEqual32, {lane, IR::Value(32u)}));
 	const auto thread_mask = low ? ir.Select(high_lane, IR::U32(IR::Value(0xffffffffu)), below)
@@ -389,9 +379,8 @@ void Translator::S_QUADMASK_B64(const Decoder::Instruction& inst) {
 void Translator::BFM_B32(const Decoder::Instruction& inst) {
 	const auto count  = ir.BitwiseAnd(ReadU32(inst.src0), IR::U32(IR::Value(31u)));
 	const auto offset = ir.BitwiseAnd(ReadU32(inst.src1), IR::U32(IR::Value(31u)));
-	const auto result = ir.Emit(IR::ValueOpcode::BitFieldInsert,
-	                            {IR::Value(0u), IR::Value(0xffffffffu), offset, count});
-	WriteOperand(DestinationOperand(inst), result);
+	// Mask bits shifted past bit 31 are dropped.
+	WriteOperand(DestinationOperand(inst), ir.ShiftLeftLogical(RightMask32(count), offset));
 }
 
 IR::U32 Translator::RightMask32(IR::U32 count) {
@@ -399,15 +388,6 @@ IR::U32 Translator::RightMask32(IR::U32 count) {
 	                       {IR::Value(0u), IR::Value(0xffffffffu), IR::Value(0u), count}));
 }
 
-IR::U64 Translator::RightMask64(IR::U32 count) {
-	const auto below32 = IR::U1(ir.Emit(IR::ValueOpcode::ULessThan32, {count, IR::Value(32u)}));
-	const auto above32 = IR::U1(ir.Emit(IR::ValueOpcode::UGreaterThan32, {count, IR::Value(32u)}));
-	const auto low_count = ir.Select(below32, count, IR::U32(IR::Value(32u)));
-	const auto high_count =
-	    ir.Select(above32, ir.ISub(count, IR::U32(IR::Value(32u))), IR::U32(IR::Value(0u)));
-	return ir.ConstructU64(RightMask32(low_count), RightMask32(high_count));
-}
-
 void Translator::S_BFM_B64(const Decoder::Instruction& inst) {
 	const auto count  = ReadU32(inst.src0);
 	const auto offset = ReadU32(inst.src1);
@@ -439,12 +419,15 @@ void Translator::S_BFE_U64(const Decoder::Instruction& inst) {
 	const auto field  = ReadU32(inst.src1);
 	const auto offset =
 	    IR::U32(ir.Emit(IR::ValueOpcode::BitFieldUExtract, {field, IR::Value(0u), IR::Value(6u)}));
-	const auto raw_count =
+	const auto count =
 	    IR::U32(ir.Emit(IR::ValueOpcode::BitFieldUExtract, {field, IR::Value(16u), IR::Value(7u)}));
-	const auto available = ir.ISub(IR::U32(IR::Value(64u)), offset);
-	const auto count     = IR::U32(ir.Emit(IR::ValueOpcode::UMin32, {raw_count, available}));
-	const auto shifted   = IR::U64(ir.Emit(IR::ValueOpcode::ShiftRightLogical64, {source, offset}));
-	const auto result    = ir.Emit(IR::ValueOpcode::BitwiseAnd64, {shifted, RightMask64(count)});
+	// The shift already clears bits past the source; widths of 64 and more keep every bit.
+	const auto low_count = IR::U32(ir.Emit(IR::ValueOpcode::UMin32, {count, IR::Value(32u)}));
+	const auto high_count =
+	    IR::U32(ir.Emit(IR::ValueOpcode::UMin32, {ir.ISub(count, low_count), IR::Value(32u)}));
+	const auto mask    = ir.ConstructU64(RightMask32(low_count), RightMask32(high_count));
+	const auto shifted = IR::U64(ir.Emit(IR::ValueOpcode::ShiftRightLogical64, {source, offset}));
+	const auto result  = ir.Emit(IR::ValueOpcode::BitwiseAnd64, {shifted, mask});
 	WriteOperand(DestinationOperand(inst), result);
 	ir.SetScc(IR::U1(ir.Emit(IR::ValueOpcode::INotEqual64, {result, IR::Value(uint64_t {0})})));
 }

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

@@ -238,10 +238,10 @@ void Translator::EmitScalar(const Decoder::Instruction& inst) {
 			return ComposedIntegerBinary(inst, IR::ValueOpcode::BitwiseOr32, false, true, true);
 		case O::S_XNOR_B32:
 			return ComposedIntegerBinary(inst, IR::ValueOpcode::BitwiseXor32, false, true, true);
-		case O::S_FF0_I32_B64: return S_FF_I32_B64(inst, true);
-		case O::S_FF1_I32_B64: return S_FF_I32_B64(inst, false);
+		case O::S_FF0_I32_B64: return S_FF_I32_B64(inst, true, false);
+		case O::S_FF1_I32_B64: return S_FF_I32_B64(inst, false, false);
 		case O::S_FLBIT_I32_B32: return V_FFBH_32(inst, false);
-		case O::S_FLBIT_I32_B64: return S_FLBIT_I32_B64(inst);
+		case O::S_FLBIT_I32_B64: return S_FF_I32_B64(inst, false, true);
 
 		case O::S_BITSET0_B32: return S_BITSET_B32(inst, false);
 		case O::S_BITSET1_B32: return S_BITSET_B32(inst, true);

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

@@ -296,15 +296,11 @@ void Translator::WriteRawU32(const Decoder::Operand& operand, IR::U32 value) {
 				value               = ir.ShiftLeftLogical(extended, IR::U32(IR::Value(offset)));
 				break;
 			}
-			case 2: {
-				const uint32_t field_mask = width == 32u ? 0xffffffffu : (1u << width) - 1u;
-				const auto     inserted =
-				    ir.ShiftLeftLogical(ir.BitwiseAnd(value, IR::U32(IR::Value(field_mask))),
-				                        IR::U32(IR::Value(offset)));
-				const auto cleared = ir.BitwiseAnd(ReadRawU32(PlainOperand(operand)),
-				                                   IR::U32(IR::Value(~(field_mask << offset))));
-				value              = ir.BitwiseOr(cleared, inserted);
-			} break;
+			case 2:
+				value = IR::U32(ir.Emit(IR::ValueOpcode::BitFieldInsert,
+				                        {ReadRawU32(PlainOperand(operand)), value, IR::Value(offset),
+				                         IR::Value(width)}));
+				break;
 			default: EXIT("reserved SDWA DST_U mode");
 		}
 	}
@@ -453,7 +449,10 @@ void Translator::Write16Bits(const Decoder::Operand& operand, IR::U32 value) {
 	destination.omod   = 0u;
 	destination.op_sel = false;
 	destination.clamp  = false;
-	WriteRawU32(destination, ir.BitwiseAnd(value, IR::U32(IR::Value(0xffffu))));
+	// A field selector keeps only the field's bits; a full-dword SDWA selector zero-extends.
+	WriteRawU32(destination, destination.sdwa_sel == 6u
+	                             ? ir.BitwiseAnd(value, IR::U32(IR::Value(0xffffu)))
+	                             : value);
 }
 
 void Translator::WriteF16(const Decoder::Operand& operand, IR::F32 value) {
@@ -740,14 +739,16 @@ void Translator::WriteCompareResult(const Decoder::Operand& operand, IR::U1 valu
 void Translator::AddBranchCondition(const CFG::Graph& graph, const CFG::BasicBlock& source,
                                     IR::Block& info) {
 	const auto native_condition = [&](CFG::BranchCondition kind) -> IR::U1 {
+		// A zero kind negates the vote of the nonzero kind that follows it in the enum.
+		const bool zero = kind == CFG::BranchCondition::SccZero ||
+		                  kind == CFG::BranchCondition::VccZero ||
+		                  kind == CFG::BranchCondition::ExecZero;
+		if (zero) kind = static_cast<CFG::BranchCondition>(static_cast<int>(kind) + 1);
 		IR::U1 condition;
 		switch (kind) {
 			case CFG::BranchCondition::Always: condition = IR::U1(IR::Value(true)); break;
-			case CFG::BranchCondition::SccZero: condition = ir.LogicalNot(ir.GetScc()); break;
 			case CFG::BranchCondition::SccNonZero: condition = ir.GetScc(); break;
-			case CFG::BranchCondition::VccZero: condition = ir.LogicalNot(ir.GetVcc()); break;
 			case CFG::BranchCondition::VccNonZero: condition = ir.GetVcc(); break;
-			case CFG::BranchCondition::ExecZero: condition = ir.LogicalNot(ir.GetExec()); break;
 			case CFG::BranchCondition::ExecNonZero: condition = ir.GetExec(); break;
 			case CFG::BranchCondition::ScalarInstruction:
 				EXIT_IF(instruction_branch_condition.IsEmpty());
@@ -755,7 +756,8 @@ void Translator::AddBranchCondition(const CFG::Graph& graph, const CFG::BasicBlo
 				break;
 			default: EXIT("block %u has an invalid native branch condition", source.id);
 		}
-		return IR::U1(ir.Emit(IR::ValueOpcode::ConditionRef, {condition}, kind));
+		const auto vote = IR::U1(ir.Emit(IR::ValueOpcode::ConditionRef, {condition}, kind));
+		return zero ? ir.LogicalNot(vote) : vote;
 	};
 	const auto expression = [&](auto&& self, uint32_t index) -> IR::U1 {
 		const auto& value = graph.expressions.at(index);

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

@@ -113,11 +113,6 @@ private:
 	void DS_ORDERED_COUNT(const Decoder::Instruction& inst);
 
 	IR::F32 SelectF32(IR::U1 condition, IR::F32 true_value, IR::F32 false_value);
-	IR::U32 ConvertF32ToU32Saturated(IR::F32 value, float upper_bound, float safe_upper,
-	                                 uint32_t high_result);
-	IR::U32 ConvertF32ToI32Saturated(IR::F32 value, float lower_bound, float upper_bound,
-	                                 float safe_upper, uint32_t lower_result,
-	                                 uint32_t upper_result);
 	IR::U32 PackU16Lanes(IR::U32 low, IR::U32 high);
 	void EmitCompareResult(const Decoder::Instruction& inst, IR::U1 value, bool scalar, bool cmpx);
 	void EmitCompareConstant(const Decoder::Instruction& inst, bool value, bool scalar, bool cmpx);
@@ -187,9 +182,8 @@ private:
 	void V_AND_OR_B32(const Decoder::Instruction& inst);
 	void V_OR3_B32(const Decoder::Instruction& inst);
 	void V_XOR3_B32(const Decoder::Instruction& inst);
-	void S_FF_I32_B64(const Decoder::Instruction& inst, bool find_zero);
+	void S_FF_I32_B64(const Decoder::Instruction& inst, bool find_zero, bool from_msb);
 	void V_FFBH_32(const Decoder::Instruction& inst, bool sign);
-	void S_FLBIT_I32_B64(const Decoder::Instruction& inst);
 	void Integer24(const Decoder::Instruction& inst, bool sign, bool addend);
 	void V_MAD_U64_U32(const Decoder::Instruction& inst);
 	void V_SAD_U32(const Decoder::Instruction& inst);
@@ -202,7 +196,6 @@ private:
 	void S_QUADMASK_B64(const Decoder::Instruction& inst);
 	void BFM_B32(const Decoder::Instruction& inst);
 	IR::U32 RightMask32(IR::U32 count);
-	IR::U64 RightMask64(IR::U32 count);
 	void    S_BFM_B64(const Decoder::Instruction& inst);
 	void    S_BFE_U32(const Decoder::Instruction& inst, bool sign);
 	void    S_BFE_U64(const Decoder::Instruction& inst);

src/graphics/shader/recompiler/ir/IREmitter.cpp

@@ -6,11 +6,6 @@ IREmitter::IREmitter(Block* value_block): block(value_block) {
 	EXIT_IF(block == nullptr);
 }
 
-void IREmitter::SetBlock(Block* value_block) {
-	EXIT_IF(value_block == nullptr);
-	block = value_block;
-}
-
 Value IREmitter::Emit(ValueOpcode opcode, std::initializer_list<Value> args, uint64_t flags) {
 	auto& inst = block->AppendNewInst(opcode, args, flags);
 	return TypeOf(opcode) == Type::Void ? Value {} : Value(&inst);
@@ -132,16 +127,6 @@ U32 IREmitter::BitCastU32(F32 value) {
 	return U32(Emit(ValueOpcode::BitCastU32F32, {value}));
 }
 
-F16 IREmitter::BitCastF16(U32 value) {
-	const auto half = U16(Emit(ValueOpcode::ConvertU16U32, {value}));
-	return F16(Emit(ValueOpcode::BitCastF16U16, {half}));
-}
-
-U32 IREmitter::BitCastU32(F16 value) {
-	const auto half = U16(Emit(ValueOpcode::BitCastU16F16, {value}));
-	return U32(Emit(ValueOpcode::ConvertU32U16, {half}));
-}
-
 U64 IREmitter::ConstructU64(U32 low, U32 high) {
 	return U64(Emit(ValueOpcode::CompositeConstructU64, {low, high}));
 }

src/graphics/shader/recompiler/ir/IREmitter.h

@@ -12,8 +12,6 @@ class IREmitter {
 public:
 	explicit IREmitter(Block* block);
 
-	void SetBlock(Block* block);
-
 	Value Emit(ValueOpcode opcode, std::initializer_list<Value> args = {}, uint64_t flags = 0);
 
 	template <typename T>
@@ -54,8 +52,6 @@ public:
 
 	F32 BitCastF32(U32 value);
 	U32 BitCastU32(F32 value);
-	F16 BitCastF16(U32 value);
-	U32 BitCastU32(F16 value);
 	U64 ConstructU64(U32 low, U32 high);
 	U32 CompositeExtract(Value composite, uint32_t index);
 

src/graphics/shader/recompiler/ir/ShaderIR.h

@@ -125,6 +125,7 @@ struct BufferResource {
 	uint32_t               descriptor_index   = UINT32_MAX;
 	uint32_t               first_use_pc       = 0;
 	uint32_t               max_byte_extent    = 0;
+	uint32_t               write_extent       = 0; // UINT32_MAX once a write is runtime-addressed or formatted
 	uint32_t               packed_stride      = 0;
 	Prospero::BufferFormat descriptor_format  = Prospero::BufferFormat::kInvalid;
 	uint32_t               descriptor_swizzle = DstSel(4, 5, 6, 7);

src/graphics/shader/recompiler/ir/opcodes/ValueOpcodes.inc

@@ -44,10 +44,7 @@ VALUE_OPCODE(GetM0, U32)
 VALUE_OPCODE(SetM0, Void, U32)
 
 VALUE_OPCODE(UndefU1, U1)
-VALUE_OPCODE(UndefU8, U8)
-VALUE_OPCODE(UndefU16, U16)
 VALUE_OPCODE(UndefU32, U32)
-VALUE_OPCODE(UndefU64, U64)
 
 // Value construction and bit-preserving conversion.
 VALUE_OPCODE(BitCastU16F16, U16, F16)
@@ -100,11 +97,9 @@ VALUE_OPCODE(SelectU1, U1, U1, U1, U1)
 VALUE_OPCODE(SelectF32, F32, U1, F32, F32)
 VALUE_OPCODE(IAdd32, U32, U32, U32)
 VALUE_OPCODE(IAdd64, U64, U64, U64)
-VALUE_OPCODE(IAddCarry32, U32x2, U32, U32)
 VALUE_OPCODE(ISub32, U32, U32, U32)
 VALUE_OPCODE(ISub64, U64, U64, U64)
 VALUE_OPCODE(IMul32, U32, U32, U32)
-VALUE_OPCODE(IMul64, U64, U64, U64)
 VALUE_OPCODE(UDiv32, U32, U32, U32)
 VALUE_OPCODE(SMulHi, U32, U32, U32)
 VALUE_OPCODE(UMulHi, U32, U32, U32)
@@ -122,9 +117,7 @@ VALUE_OPCODE(BitwiseXor32, U32, U32, U32)
 VALUE_OPCODE(BitwiseNot32, U32, U32)
 VALUE_OPCODE(BitReverse32, U32, U32)
 VALUE_OPCODE(BitCount32, U32, U32)
-VALUE_OPCODE(BitCount64, U32, U64)
 VALUE_OPCODE(FindUMsb32, U32, U32)
-VALUE_OPCODE(FindUMsb64, U32, U64)
 VALUE_OPCODE(FindILsb32, U32, U32)
 VALUE_OPCODE(SMin32, U32, U32, U32)
 VALUE_OPCODE(UMin32, U32, U32, U32)

src/graphics/shader/recompiler/ir/passes/ConstantPropagation.cpp

@@ -244,17 +244,6 @@ bool FoldCompositeExtract(Inst& inst, ValueOpcode construct, size_t components)
 		Replace(inst, producer->Arg(component));
 		return true;
 	}
-	if (component < 2u && producer != nullptr &&
-	    producer->GetOpcode() == ValueOpcode::IAddCarry32) {
-		const auto lhs = producer->Arg(0).Resolve();
-		const auto rhs = producer->Arg(1).Resolve();
-		if (IsImmediate(lhs, Type::U32) && IsImmediate(rhs, Type::U32)) {
-			const auto sum = static_cast<uint64_t>(lhs.U32()) + rhs.U32();
-			Replace(inst, Value(component == 0u ? static_cast<uint32_t>(sum)
-			                                    : static_cast<uint32_t>(sum >> 32u)));
-			return true;
-		}
-	}
 	return false;
 }
 
@@ -511,11 +500,6 @@ void FoldInstruction(Block& block, Block::iterator instruction,
 				ReplaceBinaryIdentity(inst, Type::U32, 1u);
 			}
 			return;
-		case ValueOpcode::IMul64:
-			if (!FoldU64(inst, [](uint64_t a, uint64_t b) { return a * b; })) {
-				ReplaceBinaryIdentity(inst, Type::U64, 1u);
-			}
-			return;
 		case ValueOpcode::WqmU64: {
 			const auto value = Arg(inst, 0);
 			if (IsImmediate(value, Type::U64)) {
@@ -667,17 +651,15 @@ void FoldInstruction(Block& block, Block::iterator instruction,
 			}
 			return;
 		}
-		case ValueOpcode::BitCount32: {
+		case ValueOpcode::BitCount32:
+		case ValueOpcode::FindILsb32: {
 			const auto value = Arg(inst, 0);
 			if (IsImmediate(value, Type::U32)) {
-				Replace(inst, Value(static_cast<uint32_t>(std::popcount(value.U32()))));
-			}
-			return;
-		}
-		case ValueOpcode::BitCount64: {
-			const auto value = Arg(inst, 0);
-			if (IsImmediate(value, Type::U64)) {
-				Replace(inst, Value(static_cast<uint32_t>(std::popcount(value.U64()))));
+				const auto bits = value.U32();
+				Replace(inst, Value(static_cast<uint32_t>(
+				    inst.GetOpcode() == ValueOpcode::BitCount32 ? std::popcount(bits)
+				    : bits == 0u                                ? -1
+				                                                : std::countr_zero(bits))));
 			}
 			return;
 		}
@@ -690,7 +672,9 @@ void FoldInstruction(Block& block, Block::iterator instruction,
 			});
 			return;
 		case ValueOpcode::UMin32:
-			FoldU32(inst, [](uint32_t a, uint32_t b) { return std::min(a, b); });
+			if (!FoldU32(inst, [](uint32_t a, uint32_t b) { return std::min(a, b); })) {
+				ReplaceBinaryIdentity(inst, Type::U32, ~0u);
+			}
 			return;
 		case ValueOpcode::UMax32:
 			FoldU32(inst, [](uint32_t a, uint32_t b) { return std::max(a, b); });

src/graphics/shader/recompiler/ir/passes/ResourceTracking.cpp

@@ -56,6 +56,17 @@ uint32_t PossibleU32Bits(Value value, const auto& normalize) {
 				}
 				break;
 			}
+			case ValueOpcode::BitFieldInsert: {
+				const auto offset = inst->Arg(2).Resolve();
+				const auto count  = inst->Arg(3).Resolve();
+				if (offset.IsImmediate() && count.IsImmediate() && offset.U32() < 32u &&
+				    count.U32() <= 32u - offset.U32()) {
+					const auto field = (count.U32() == 32u ? UINT32_MAX : (1u << count.U32()) - 1u)
+					                   << offset.U32();
+					bits = (arg(0) & ~field) | ((arg(1) << offset.U32()) & field);
+				}
+				break;
+			}
 			default: break;
 		}
 		return bits;
@@ -3749,10 +3760,10 @@ private:
 		}
 	}
 
-	uint32_t AddBuffer(uint32_t source, const MemoryInfo& memory, ValueOpcode op, uint32_t pc) {
+	uint32_t AddBuffer(uint32_t source, const MemoryInfo& memory, const Inst& inst, uint32_t pc) {
 		for (uint32_t i = 0; i < m_info.buffers.size(); i++) {
 			if (m_info.buffers[i].source == source) {
-				Merge(m_info.buffers[i], memory, op, pc);
+				Merge(m_info.buffers[i], memory, inst, pc);
 				return i;
 			}
 		}
@@ -3762,13 +3773,14 @@ private:
 		BufferResource resource;
 		resource.source       = source;
 		resource.first_use_pc = pc;
-		Merge(resource, memory, op, pc);
+		Merge(resource, memory, inst, pc);
 		m_info.buffers.push_back(resource);
 		return static_cast<uint32_t>(m_info.buffers.size() - 1);
 	}
 
-	static void Merge(BufferResource& resource, const MemoryInfo& memory, ValueOpcode op,
+	static void Merge(BufferResource& resource, const MemoryInfo& memory, const Inst& inst,
 	                  uint32_t pc) {
+		const auto op            = inst.GetOpcode();
 		const auto access        = BufferAccessOf(op);
 		const bool atomic        = access == BufferAccess::Atomic;
 		const bool write         = access == BufferAccess::Write || atomic;
@@ -3776,6 +3788,13 @@ private:
 		resource.max_byte_extent = std::max(resource.max_byte_extent, ByteExtent(memory));
 		resource.read            = resource.read || !write || atomic;
 		resource.written         = resource.written || write;
+		if (write) {
+			// Formatted stores write the whole format; otherwise zero index, VGPR and SGPR offsets
+			// leave only the immediate offset.
+			const bool fixed = !memory.formatted && inst.Arg(1).Resolve() == Value(0u) &&
+			                   inst.Arg(2).Resolve() == Value(0u) && inst.Arg(3).Resolve() == Value(0u);
+			resource.write_extent = std::max(resource.write_extent, fixed ? ByteExtent(memory) : UINT32_MAX);
+		}
 		resource.atomic          = resource.atomic || atomic;
 		resource.formatted       = resource.formatted || memory.formatted;
 		resource.scalar          = resource.scalar || op == ValueOpcode::ReadConstBuffer ||
@@ -3945,7 +3964,7 @@ private:
 				m_info.uses_dma                         = true;
 				return;
 			}
-			resource = AddBuffer(source, memory, op, flags.pc);
+			resource = AddBuffer(source, memory, inst, flags.pc);
 			if (resource == UINT32_MAX) {
 				Fail(flags.pc, "buffer resource limit exceeded");
 			}

src/graphics/shader/recompiler/ir/passes/SrtWalker.cpp

@@ -134,11 +134,9 @@ bool IsRuntimeUniformOp(ValueOpcode op) {
 		case ValueOpcode::BitFieldSExtract:
 		case ValueOpcode::IAdd32:
 		case ValueOpcode::IAdd64:
-		case ValueOpcode::IAddCarry32:
 		case ValueOpcode::ISub32:
 		case ValueOpcode::ISub64:
 		case ValueOpcode::IMul32:
-		case ValueOpcode::IMul64:
 		case ValueOpcode::UMulHi:
 		case ValueOpcode::UMin32:
 		case ValueOpcode::UMax32:
@@ -259,9 +257,7 @@ private:
 			}
 			return finish(Validate(inst->Arg(1)));
 		}
-		if (op == ValueOpcode::UndefU1 || op == ValueOpcode::UndefU8 ||
-		    op == ValueOpcode::UndefU16 || op == ValueOpcode::UndefU32 ||
-		    op == ValueOpcode::UndefU64 || op == ValueOpcode::Void) {
+		if (op == ValueOpcode::UndefU1 || op == ValueOpcode::UndefU32 || op == ValueOpcode::Void) {
 			return finish(false);
 		}
 		if (op == ValueOpcode::GetUserData) {
@@ -341,9 +337,7 @@ private:
 			const auto* source = inst->NumArgs() == 2 ? inst->Arg(0).ResolveInstruction() : nullptr;
 			const auto  index  = inst->NumArgs() == 2 ? inst->Arg(1).Resolve() : Value {};
 			if (source == nullptr || !index.IsImmediate() || index.GetType() != Type::U32 ||
-			    index.U32() >= 2u ||
-			    (source->GetOpcode() != ValueOpcode::CompositeConstructU32x2 &&
-			     source->GetOpcode() != ValueOpcode::IAddCarry32)) {
+			    index.U32() >= 2u || source->GetOpcode() != ValueOpcode::CompositeConstructU32x2) {
 				return finish(false);
 			}
 		}
@@ -656,18 +650,6 @@ bool SrtWalker::EvaluateExtract(const Inst& inst, uint64_t& result) {
 	if (source->GetOpcode() == ValueOpcode::CompositeConstructU32x2) {
 		return EvaluateWide(source->Arg(component), result);
 	}
-	if (source->GetOpcode() == ValueOpcode::IAddCarry32) {
-		uint64_t lhs = 0;
-		uint64_t rhs = 0;
-		if (!Arg(*source, 0, lhs) || !Arg(*source, 1, rhs)) {
-			return false;
-		}
-		const auto sum =
-		    static_cast<uint64_t>(static_cast<uint32_t>(lhs)) + static_cast<uint32_t>(rhs);
-		result =
-		    component == 0u ? static_cast<uint32_t>(sum) : static_cast<uint32_t>(sum >> 32u);
-		return true;
-	}
 	return false;
 }
 
@@ -832,12 +814,6 @@ bool SrtWalker::EvaluateInst(const Inst& inst, uint64_t& result) {
 				return true;
 			}
 			return false;
-		case ValueOpcode::IMul64:
-			if (binary()) {
-				result = a * b;
-				return true;
-			}
-			return false;
 		case ValueOpcode::UMulHi:
 			if (binary()) {
 				result = (static_cast<uint64_t>(static_cast<uint32_t>(a)) * static_cast<uint32_t>(b)) >> 32u;
@@ -861,12 +837,11 @@ bool SrtWalker::EvaluateInst(const Inst& inst, uint64_t& result) {
 			return false;
 		case ValueOpcode::ConvertU32F32:
 			if (Arg(inst, 0, a)) {
+				// The conversion saturates and maps NaN to zero.
 				const auto value = Float32(a);
-				if (!std::isfinite(value) || value < 0.0f ||
-				    static_cast<double>(value) > UINT32_MAX) {
-					return false;
-				}
-				result = static_cast<uint32_t>(value);
+				result           = std::isnan(value) || value <= 0.0f ? 0u
+				                   : value >= 4294967296.0f ? UINT32_MAX
+				                                            : static_cast<uint32_t>(value);
 				return true;
 			}
 			return false;
@@ -1143,10 +1118,7 @@ bool SrtWalker::EvaluateInst(const Inst& inst, uint64_t& result) {
 			}
 			return false;
 		case ValueOpcode::UndefU1:
-		case ValueOpcode::UndefU8:
-		case ValueOpcode::UndefU16:
-		case ValueOpcode::UndefU32:
-		case ValueOpcode::UndefU64: return false;
+		case ValueOpcode::UndefU32: return false;
 		default: break;
 	}
 	return false;

src/graphics/shader/shader.cpp

@@ -355,10 +355,11 @@ static void ShaderGetStaticInputInfoPS(
 	ps_info.ps_pixel_kill_enable         = sh.db_shader_control.shader_kill_enable;
 	ps_info.ps_depth_export_enable       = sh.db_shader_control.shader_z_export_enable;
 	ps_info.ps_sample_mask_export_enable = sh.db_shader_control.shader_mask_export_enable;
+	// Alpha-to-mask coverage must gate depth writes, so it rules out early fragment tests.
 	ps_info.ps_early_z =
 	    (sh.db_shader_control.shader_z_behavior == 1 && !sh.db_shader_control.shader_kill_enable &&
 	     !sh.db_shader_control.shader_z_export_enable &&
-	     !sh.db_shader_control.shader_mask_export_enable);
+	     !sh.db_shader_control.shader_mask_export_enable && !sh.AlphaToMask());
 	ps_info.ps_execute_on_noop = sh.db_shader_control.shader_execute_on_noop;
 
 	for (uint32_t i = 0; i < ps_info.input_num; i++) {

src/kernel/memoryAddressSpace.inc

@@ -28,9 +28,9 @@ static bool CoalescePlaceholders(uint64_t vaddr, uint64_t size) {
 }
 #endif
 
-template <typename Regions, typename Source, typename Gap>
-static void CopySparseRange(const Regions& regions, uint64_t vaddr, void* data, uint64_t size,
-                            Source&& source, Gap&& gap) {
+template <typename Regions, typename Region, typename Gap>
+static void ForEachSparseRange(const Regions& regions, uint64_t vaddr, uint64_t size,
+                               Region&& region, Gap&& gap) {
 	const auto end     = vaddr + size;
 	auto       mapping = regions.lower_bound(vaddr);
 	if (mapping != regions.begin()) {
@@ -40,23 +40,34 @@ static void CopySparseRange(const Regions& regions, uint64_t vaddr, void* data,
 		}
 	}
 
-	auto* destination = static_cast<uint8_t*>(data);
-	auto  cursor      = vaddr;
+	auto cursor = vaddr;
 	for (; mapping != regions.end() && mapping->second.vaddr < end; ++mapping) {
 		const auto& entry = mapping->second;
 		const auto  start = std::max(vaddr, entry.vaddr);
 		const auto  bytes = std::min(end, entry.vaddr + entry.size) - start;
 		if (cursor < start) {
-			gap(cursor, destination + (cursor - vaddr), start - cursor);
+			gap(cursor, start - cursor);
 		}
-		std::memcpy(destination + (start - vaddr), source(entry, start), bytes);
+		region(entry, start, bytes);
 		cursor = start + bytes;
 	}
 	if (cursor < end) {
-		gap(cursor, destination + (cursor - vaddr), end - cursor);
+		gap(cursor, end - cursor);
 	}
 }
 
+template <typename Regions, typename Source, typename Gap>
+static void CopySparseRange(const Regions& regions, uint64_t vaddr, void* data, uint64_t size,
+                            Source&& source, Gap&& gap) {
+	auto* destination = static_cast<uint8_t*>(data);
+	ForEachSparseRange(
+	    regions, vaddr, size,
+	    [&](const auto& entry, uint64_t start, uint64_t bytes) {
+		    std::memcpy(destination + (start - vaddr), source(entry, start), bytes);
+	    },
+	    [&](uint64_t start, uint64_t bytes) { gap(start, destination + (start - vaddr), bytes); });
+}
+
 class GuestBackingStore {
 public:
 	enum class FailureReason {
@@ -200,36 +211,16 @@ public:
 			return false;
 		}
 
+		// Skip unbacked gaps, mirroring sparse reads that zero-fill them.
 		std::lock_guard lock(m_mutex);
-		const auto      end   = vaddr + size;
-		auto            first = m_maps.upper_bound(vaddr);
-		if (first == m_maps.begin()) {
-			return false;
-		}
-		--first;
-		auto current = vaddr;
-		for (auto mapping = first; current < end; ++mapping) {
-			if (mapping == m_maps.end() || current < mapping->second.vaddr ||
-			    current >= mapping->second.vaddr + mapping->second.size) {
-				return false;
-			}
-			const auto& entry = mapping->second;
-			const auto  bytes = std::min(end, entry.vaddr + entry.size) - current;
-			if (!RangeOk(entry.backing_offset + current - entry.vaddr, bytes)) {
-				return false;
-			}
-			current += bytes;
-		}
-
-		const auto* source = static_cast<const uint8_t*>(data);
-		current = vaddr;
-		for (auto mapping = first; current < end; ++mapping) {
-			const auto& entry  = mapping->second;
-			const auto  bytes  = std::min(end, entry.vaddr + entry.size) - current;
-			const auto  offset = entry.backing_offset + current - entry.vaddr;
-			std::memcpy(m_backing_base + offset, source + (current - vaddr), bytes);
-			current += bytes;
-		}
+		const auto*     source = static_cast<const uint8_t*>(data);
+		ForEachSparseRange(
+		    m_maps, vaddr, size,
+		    [&](const auto& entry, uint64_t start, uint64_t bytes) {
+			    std::memcpy(m_backing_base + (entry.backing_offset + start - entry.vaddr),
+			                source + (start - vaddr), bytes);
+		    },
+		    [](uint64_t, uint64_t) {});
 		return true;
 	}
 
@@ -775,8 +766,14 @@ public:
 		});
 	}
 
+	// Writes the backed bytes of a range that holds no private mappings.
 	bool TryWriteBacking(uint64_t vaddr, const void* data, uint64_t size) {
-		return m_backing->TryWriteBacking(vaddr, data, size);
+		if (!Owns(vaddr, size)) {
+			return false;
+		}
+		std::shared_lock lock(m_mutex);
+		return !OverlapsPrivateUnlocked(vaddr, size) &&
+		       m_backing->TryWriteBacking(vaddr, data, size);
 	}
 
 	bool TryReadBacking(uint64_t vaddr, void* data, uint64_t size) {
@@ -790,15 +787,8 @@ public:
 		}
 		std::shared_lock lock(m_mutex);
 		if constexpr (!copy_private) {
-			auto mapping = m_mapped.upper_bound(vaddr);
-			if (mapping != m_mapped.begin()) {
-				--mapping;
-			}
-			for (; mapping != m_mapped.end() && mapping->second.vaddr < vaddr + size; ++mapping) {
-				const auto& region = mapping->second;
-				if (region.type == MappingType::Private && region.vaddr + region.size > vaddr) {
-					return false;
-				}
+			if (OverlapsPrivateUnlocked(vaddr, size)) {
+				return false;
 			}
 		}
 		return m_backing->TryReadSparseBacking(
@@ -1111,6 +1101,20 @@ private:
 		return true;
 	}
 
+	[[nodiscard]] bool OverlapsPrivateUnlocked(uint64_t vaddr, uint64_t size) const {
+		auto mapping = m_mapped.upper_bound(vaddr);
+		if (mapping != m_mapped.begin()) {
+			--mapping;
+		}
+		for (; mapping != m_mapped.end() && mapping->second.vaddr < vaddr + size; ++mapping) {
+			const auto& region = mapping->second;
+			if (region.type == MappingType::Private && region.vaddr + region.size > vaddr) {
+				return true;
+			}
+		}
+		return false;
+	}
+
 	bool AddMappedUnlocked(uint64_t vaddr, uint64_t size, MappingType type) {
 		if (size == 0 || UINT64_MAX - vaddr < size) {
 			return false;

src/kernel/pthread.cpp

@@ -3389,6 +3389,11 @@ int KYTY_SYSV_ABI PthreadCancel(Pthread thread) {
 		return KERNEL_ERROR_EINVAL;
 	}
 
+	// A joined thread's native handle is gone; the guest may still hold its handle.
+	if (thread->free) {
+		return KERNEL_ERROR_ESRCH;
+	}
+
 	int result = pthread_cancel(thread->p);
 
 	LOGF("\tthread cancel: %s, %d\n", thread->name.c_str(), result);
@@ -3492,10 +3497,8 @@ int KYTY_SYSV_ABI PthreadGetprio(Pthread thread, int* prio) {
 
 	EXIT_NOT_IMPLEMENTED(prio == nullptr);
 
-	sched_param native_param {};
-	int         native_policy = 0;
-	if (pthread_getschedparam(thread->p, &native_policy, &native_param) != 0) {
-		return KERNEL_ERROR_EINVAL;
+	if (thread->free) {
+		return KERNEL_ERROR_ESRCH;
 	}
 
 	*prio = thread->attr->guest_priority;
@@ -3510,6 +3513,11 @@ int KYTY_SYSV_ABI PthreadSetprio(Pthread thread, int prio) {
 		return KERNEL_ERROR_ESRCH;
 	}
 
+	// A joined thread's native handle is gone; the guest may still hold its handle.
+	if (thread->free) {
+		return KERNEL_ERROR_ESRCH;
+	}
+
 	sched_param param {};
 	int         pol = 0;
 

src/launcher/forms/configuration_edit_dialog.ui

@@ -168,12 +168,18 @@
              </widget>
             </item>
             <item row="6" column="1">
+             <widget class="QCheckBox" name="checkBox_readback_tiled">
+              <property name="toolTip"><string>Read back small writable tiled images</string></property>
+              <property name="text"><string>Enable tiled readback</string></property>
+             </widget>
+            </item>
+            <item row="7" column="1">
              <widget class="QCheckBox" name="checkBox_tessellation">
               <property name="toolTip"><string>Draw tessellated geometry. Disabled by default; unsupported tessellation may cause crashes.</string></property>
               <property name="text"><string>Enable tessellation support (experimental)</string></property>
              </widget>
             </item>
-            <item row="7" column="1">
+            <item row="8" column="1">
              <widget class="QCheckBox" name="checkBox_hdr">
               <property name="toolTip"><string>Enable HDR10 output for supported games. Requires HDR enabled on the host display.</string></property>
               <property name="text"><string>Enable HDR display support</string></property>

src/launcher/icons/kyty-small.svg

@@ -0,0 +1,15 @@
+<svg xmlns="http://www.w3.org/2000/svg" viewBox="0 0 512 512"><defs>
+<linearGradient id="body" x1="0" y1="0" x2="0" y2="1"><stop offset="0" stop-color="#ffffff"/><stop offset="1" stop-color="#e3e9f2"/></linearGradient>
+</defs>
+<g transform="translate(256 268) scale(1.16) translate(-256 -268)" stroke-linejoin="round" stroke-linecap="round">
+<path d="M100 196 L120 104 Q124 92 134 99 L214 160 Z" fill="url(#body)"/><path d="M412 196 L392 104 Q388 92 378 99 L298 160 Z" fill="url(#body)"/>
+<path d="M116 181 L137 121 L190 157 Z" fill="#0b0d12"/><path d="M396 181 L375 121 L322 157 Z" fill="#0b0d12"/>
+<g fill="none" stroke="#0b0d12" stroke-width="20"><path d="M100 196 L120 104 Q124 92 134 99 L214 160 Z"/><path d="M412 196 L392 104 Q388 92 378 99 L298 160 Z"/></g>
+<path d="M256 154 C330 154 380 160 406 178 C432 196 447 244 459 302 C471 358 466 398 430 408 C400 416 380 398 360 368 L326 318 C316 304 306 300 290 300 L222 300 C206 300 196 304 186 318 L152 368 C132 398 112 416 82 408 C46 398 41 358 53 302 C65 244 80 196 106 178 C132 160 182 154 256 154 Z" fill="url(#body)"/>
+<path d="M256 154 C330 154 380 160 406 178 C432 196 447 244 459 302 C471 358 466 398 430 408 C400 416 380 398 360 368 L326 318 C316 304 306 300 290 300 L222 300 C206 300 196 304 186 318 L152 368 C132 398 112 416 82 408 C46 398 41 358 53 302 C65 244 80 196 106 178 C132 160 182 154 256 154 Z" fill="none" stroke="#0b0d12" stroke-width="20"/>
+<g fill="none" stroke="#0b0d12" stroke-width="14"><path d="M82 260 C100 274 118 288 128 304 C134 316 121 350 108 388"/><path d="M430 260 C412 274 394 288 384 304 C378 316 391 350 404 388"/></g>
+<path d="M72 316 L92 336 M70 336 L86 352 M440 316 L420 336 M442 336 L426 352" fill="none" stroke="#0b0d12" stroke-width="10"/>
+<path d="M179 230 Q194 210 209 230" fill="none" stroke="#0b0d12" stroke-width="16"/>
+<circle cx="318" cy="222" r="21" fill="#0b0d12"/>
+<path d="M240 256 L272 256 L256 278 Z" fill="#0b0d12" stroke="#0b0d12" stroke-width="8"/>
+</g></svg>

src/launcher/icons/kyty.svg

@@ -0,0 +1,15 @@
+<svg xmlns="http://www.w3.org/2000/svg" viewBox="0 0 512 512"><defs>
+<linearGradient id="body" x1="0" y1="0" x2="0" y2="1"><stop offset="0" stop-color="#ffffff"/><stop offset="1" stop-color="#e3e9f2"/></linearGradient>
+</defs>
+<g transform="translate(256 268) scale(1.16) translate(-256 -268)" stroke-linejoin="round" stroke-linecap="round">
+<path d="M100 196 L120 104 Q124 92 134 99 L214 160 Z" fill="url(#body)"/><path d="M412 196 L392 104 Q388 92 378 99 L298 160 Z" fill="url(#body)"/>
+<path d="M128 164 L137 128 L174 156 Z" fill="#0b0d12"/><path d="M384 164 L375 128 L338 156 Z" fill="#0b0d12"/>
+<g fill="none" stroke="#0b0d12" stroke-width="16"><path d="M100 196 L120 104 Q124 92 134 99 L214 160 Z"/><path d="M412 196 L392 104 Q388 92 378 99 L298 160 Z"/></g>
+<path d="M256 154 C330 154 380 160 406 178 C432 196 447 244 459 302 C471 358 466 398 430 408 C400 416 380 398 360 368 L326 318 C316 304 306 300 290 300 L222 300 C206 300 196 304 186 318 L152 368 C132 398 112 416 82 408 C46 398 41 358 53 302 C65 244 80 196 106 178 C132 160 182 154 256 154 Z" fill="url(#body)"/>
+<path d="M256 154 C330 154 380 160 406 178 C432 196 447 244 459 302 C471 358 466 398 430 408 C400 416 380 398 360 368 L326 318 C316 304 306 300 290 300 L222 300 C206 300 196 304 186 318 L152 368 C132 398 112 416 82 408 C46 398 41 358 53 302 C65 244 80 196 106 178 C132 160 182 154 256 154 Z" fill="none" stroke="#0b0d12" stroke-width="16"/>
+<g fill="none" stroke="#0b0d12" stroke-width="9"><path d="M82 260 C100 274 118 288 128 304 C134 316 121 350 108 388"/><path d="M430 260 C412 274 394 288 384 304 C378 316 391 350 404 388"/></g>
+<path d="M64 322 L82 340 M62 340 L76 354 M448 322 L430 340 M450 340 L436 354" fill="none" stroke="#0b0d12" stroke-width="7"/>
+<path d="M168 230 Q194 202 220 230" fill="none" stroke="#0b0d12" stroke-width="15"/>
+<circle cx="322" cy="220" r="22" fill="#0b0d12"/>
+<path d="M238 252 L274 252 L256 276 Z" fill="#0b0d12" stroke="#0b0d12" stroke-width="8"/>
+</g></svg>

src/launcher/include/configuration.h

@@ -126,6 +126,7 @@ public:
 	bool                   hdr_enabled                 = false;
 	bool                   hide_cursor_enabled         = false;
 	bool                   readback_linear_images      = false;
+	bool                   readback_tiled_images       = false;
 	bool                   sync_raw_image_buffers      = false;
 	bool                   tessellation_enabled        = false;
 	bool                   trophy_enabled              = true;
@@ -165,6 +166,7 @@ public:
 		hdr_enabled                 = other.hdr_enabled;
 		hide_cursor_enabled         = other.hide_cursor_enabled;
 		readback_linear_images      = other.readback_linear_images;
+		readback_tiled_images       = other.readback_tiled_images;
 		sync_raw_image_buffers      = other.sync_raw_image_buffers;
 		tessellation_enabled        = other.tessellation_enabled;
 		trophy_enabled              = other.trophy_enabled;

src/launcher/include/configurationListWidget.h

@@ -68,6 +68,7 @@ protected slots:
 	void list_itemDoubleClicked(QTreeWidgetItem* witem, int column);
 	void open_game_folder();
 	void remove_save_data();
+	void compress_game();
 	void filter_configurations(const QString& text);
 
 private:

src/launcher/resources.qrc

@@ -7,6 +7,8 @@
   <file alias="grid-view.svg">icons/grid-view.svg</file>
   <file alias="remove-configuration.svg">icons/remove-configuration.svg</file>
   <file alias="trophy.svg">icons/trophy.svg</file>
+  <file alias="kyty.svg">icons/kyty.svg</file>
+  <file alias="kyty-small.svg">icons/kyty-small.svg</file>
   <file alias="hidden-trophy.png">../../assets/images/hidden-trophy.png</file>
  </qresource>
  <qresource prefix="/styles">

src/launcher/src/configuration.cpp

@@ -13,6 +13,7 @@ QVariantMap Configuration::GameSettings() const {
 	KYTY_CFG_SET(hdr_enabled);
 	KYTY_CFG_SET(hide_cursor_enabled);
 	KYTY_CFG_SET(readback_linear_images);
+	KYTY_CFG_SET(readback_tiled_images);
 	KYTY_CFG_SET(sync_raw_image_buffers);
 	KYTY_CFG_SET(tessellation_enabled);
 	KYTY_CFG_SET(trophy_enabled);
@@ -62,6 +63,7 @@ void Configuration::ReadGameSettingsValues(const Settings& s) {
 	KYTY_CFG_GET(hdr_enabled);
 	KYTY_CFG_GET(hide_cursor_enabled);
 	KYTY_CFG_GET(readback_linear_images);
+	KYTY_CFG_GET(readback_tiled_images);
 	KYTY_CFG_GET(sync_raw_image_buffers);
 	KYTY_CFG_GET(tessellation_enabled);
 	trophy_enabled   = s.value("trophy_enabled", trophy_enabled).toBool();

src/launcher/src/configurationEditDialog.cpp

@@ -286,6 +286,7 @@ void ConfigurationEditDialog::Init(const Configuration& info) {
 	m_ui->checkBox_hdr->setChecked(info.hdr_enabled);
 	m_ui->checkBox_hide_cursor->setChecked(info.hide_cursor_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);
 	m_ui->checkBox_tessellation->setChecked(info.tessellation_enabled);
 	m_ui->checkBox_skip_notice_screen->setChecked(info.skip_notice_screen);
@@ -479,6 +480,7 @@ static void UpdateInfo(Configuration& info, Ui::ConfigurationEditDialog& ui, boo
 	info.hdr_enabled               = ui.checkBox_hdr->isChecked();
 	info.hide_cursor_enabled       = ui.checkBox_hide_cursor->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();
 	info.tessellation_enabled      = ui.checkBox_tessellation->isChecked();
 	info.skip_notice_screen        = ui.checkBox_skip_notice_screen->isChecked();

src/launcher/src/configurationListWidget.cpp

@@ -27,6 +27,7 @@
 #include <QFile>
 #include <QFileDialog>
 #include <QFileInfo>
+#include <QFutureWatcher>
 #include <QHeaderView>
 #include <QIcon>
 #include <QInputDialog>
@@ -42,6 +43,8 @@
 #include <QPalette>
 #include <QPointer>
 #include <QProcess>
+#include <QProgressDialog>
+#include <QPromise>
 #include <QPushButton>
 #include <QRegularExpression>
 #include <QSaveFile>
@@ -54,6 +57,7 @@
 #include <QToolButton>
 #include <QTreeWidget>
 #include <QUrl>
+#include <QtConcurrentRun>
 #include <QtCore>
 
 #ifdef __linux__
@@ -62,6 +66,7 @@
 #include <QDBusPendingCallWatcher>
 #endif
 
+#include <atomic>
 #include <memory>
 #include <vector>
 
@@ -1171,6 +1176,71 @@ void ConfigurationListWidget::remove_save_data() {
 	}
 }
 
+void ConfigurationListWidget::compress_game() {
+	auto* item = static_cast<ConfigurationItem*>(m_ui->cfgs_list->currentItem());
+	if (item == nullptr) {
+		return;
+	}
+
+	const auto base    = item->GetInfo().basedir;
+	const auto archive = base + QStringLiteral(".zar");
+	const auto title   = tr("Compress to ZAR");
+	if (QFileInfo::exists(archive)) {
+		QMessageBox::warning(this, title,
+		                     tr("%1 already exists.").arg(QDir::toNativeSeparators(archive)));
+		return;
+	}
+
+	// Cancel through a flag, not the future, so a pack that already succeeded keeps its result.
+	auto  canceled = std::make_shared<std::atomic_bool>(false);
+	auto* progress = new QProgressDialog(
+	    tr("Compressing %1...").arg(QDir::toNativeSeparators(base)), tr("Cancel"), 0, 1000, this);
+	auto* watcher  = new QFutureWatcher<bool>(this);
+	// Stay modal through finalization and fsync, which follow the last progress step.
+	progress->setAutoReset(false);
+	connect(progress, &QProgressDialog::canceled, this, [canceled] { *canceled = true; });
+	connect(watcher, &QFutureWatcher<bool>::progressValueChanged, progress,
+	        &QProgressDialog::setValue);
+	connect(watcher, &QFutureWatcher<bool>::finished, this,
+	        [this, progress, watcher, canceled, base, archive, title] {
+		        // finished can arrive inside the modal dialog's setValue event processing.
+		        progress->deleteLater();
+		        watcher->deleteLater();
+		        if (!watcher->result()) {
+			        if (!*canceled) {
+				        QMessageBox::warning(
+				            this, title,
+				            tr("Could not compress %1.").arg(QDir::toNativeSeparators(base)));
+			        }
+			        return;
+		        }
+		        if (QMessageBox::Yes ==
+		                QMessageBox::question(this, title,
+		                                      tr("Created %1.\n\nDelete the original folder %2?")
+		                                          .arg(QDir::toNativeSeparators(archive),
+		                                               QDir::toNativeSeparators(base)),
+		                                      QMessageBox::Yes | QMessageBox::No, QMessageBox::No) &&
+		            !QDir(base).removeRecursively()) {
+			        QMessageBox::warning(
+			            this, title, tr("Could not remove %1.").arg(QDir::toNativeSeparators(base)));
+		        }
+		        ScanGameDirectory();
+	        });
+	progress->open();
+	watcher->setFuture(QtConcurrent::run(
+	    [directory = GameContent::ToPath(base), path = GameContent::ToPath(archive),
+	     canceled](QPromise<bool>& promise) {
+		    promise.setProgressRange(0, 1000);
+		    // A pack cancelled while queued must not touch a path a later pack may already own.
+		    promise.addResult(
+		        !*canceled &&
+		        Common::PackArchive(directory, path, [&](uint64_t done, uint64_t total) {
+			        promise.setProgressValue(static_cast<int>(done * 1000 / total));
+			        return !*canceled;
+		        }));
+	    }));
+}
+
 void ConfigurationListWidget::filter_configurations(const QString& text) {
 	const auto query              = text.trimmed();
 	const bool has_query          = !query.isEmpty();
@@ -1268,6 +1338,8 @@ void ConfigurationListWidget::ShowContextMenu(QTreeWidgetItem* witem) {
 		        }
 	        });
 	action_addcont->setEnabled(item != nullptr);
+	QAction* action_compress =
+	    menu.addAction(tr("Compress to ZAR..."), this, SLOT(compress_game()));
 	QAction* action_remove_save_data =
 	    menu.addAction(style()->standardIcon(QStyle::SP_DialogDiscardButton),
 	                   tr("Remove save data..."), this, SLOT(remove_save_data()));
@@ -1291,6 +1363,7 @@ void ConfigurationListWidget::ShowContextMenu(QTreeWidgetItem* witem) {
 		const auto& base = item->GetInfo().basedir;
 		action_open_folder->setDisabled(!QDir(base).exists() && !GameContent::IsArchive(base));
 		action_view_trophies->setDisabled(!has_trophy_data);
+		action_compress->setDisabled(item->IsRunning() || !QDir(base).exists());
 		action_remove_save_data->setDisabled(item->IsRunning() || save_data_dirs.isEmpty());
 		action_edit->setDisabled(item->IsRunning());
 		action_delete->setDisabled(item->IsRunning() || !item->GetInfo().custom_settings);
@@ -1298,6 +1371,7 @@ void ConfigurationListWidget::ShowContextMenu(QTreeWidgetItem* witem) {
 		action_run->setDisabled(true);
 		action_open_folder->setDisabled(true);
 		action_view_trophies->setDisabled(true);
+		action_compress->setDisabled(true);
 		action_remove_save_data->setDisabled(true);
 		action_edit->setDisabled(true);
 		action_delete->setDisabled(true);

src/launcher/src/main.cpp

@@ -2,9 +2,21 @@
 #include "launcherTheme.h"
 
 #include <QApplication>
+#include <QIcon>
 
 int main(int argc, char* argv[]) {
 	QApplication a(argc, argv);
+	// Title bars and taskbars get the bolder small-size artwork.
+	const QIcon small_art(QStringLiteral(":/icons/kyty-small.svg"));
+	const QIcon large_art(QStringLiteral(":/icons/kyty.svg"));
+	QIcon       icon;
+	for (int size: {16, 24, 32, 48, 64}) {
+		icon.addPixmap(small_art.pixmap(QSize(size, size), 1.0));
+	}
+	for (int size: {128, 256}) {
+		icon.addPixmap(large_art.pixmap(QSize(size, size), 1.0));
+	}
+	QApplication::setWindowIcon(icon);
 	LauncherTheme::Initialize(a);
 
 	MainDialog w;

src/launcher/src/mainDialog.cpp

@@ -244,8 +244,10 @@ void MainDialogPrivate::FindInterpreter() {
 	}
 }
 
-static QString BoolArg(bool value) {
-	return value ? QStringLiteral("true") : QStringLiteral("false");
+static void AddFlag(QStringList& args, bool enabled, const char* name) {
+	if (enabled) {
+		args << name;
+	}
 }
 
 static QStringList CreateEmulatorArgs(const Configuration& info) {
@@ -273,54 +275,38 @@ static QStringList CreateEmulatorArgs(const Configuration& info) {
 	if (info.gpu_index >= 0) {
 		args << "--gpu" << QString::number(info.gpu_index);
 	}
-	if (info.fullscreen_enabled) {
-		args << "--fullscreen";
-	}
-	if (info.hdr_enabled) {
-		args << "--hdr";
-	}
-	if (info.hide_cursor_enabled) {
-		args << "--hide-cursor";
-	}
-	args << "--readback-linear-images" << BoolArg(info.readback_linear_images);
-	args << "--sync-raw-image-buffers" << BoolArg(info.sync_raw_image_buffers);
-	args << "--trophy-notifications" << BoolArg(info.trophy_enabled);
+	AddFlag(args, info.fullscreen_enabled, "--fullscreen");
+	AddFlag(args, info.hdr_enabled, "--hdr");
+	AddFlag(args, info.hide_cursor_enabled, "--hide-cursor");
+	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");
+	AddFlag(args, !info.trophy_enabled, "--no-trophy-notifications");
 	args << "--trophy-duration" << QString::number(info.trophy_duration_seconds);
-	args << "--perf-overlay" << BoolArg(info.perf_overlay_enabled);
-	args << "--skip-notice-screen" << BoolArg(info.skip_notice_screen);
-	args << "--network-connected" << BoolArg(info.network_connected);
-	if (info.tessellation_enabled) {
-		args << "--tessellation";
-	}
+	AddFlag(args, info.perf_overlay_enabled, "--perf-overlay");
+	AddFlag(args, info.skip_notice_screen, "--skip-notice-screen");
+	AddFlag(args, info.network_connected, "--network-connected");
+	AddFlag(args, info.tessellation_enabled, "--tessellation");
 	args << "--vblank-frequency" << QString::number(info.vblank_frequency);
 	args << "--console-language" << QString::number(info.console_language);
-	args << "--vulkan-validation" << BoolArg(info.vulkan_validation_enabled);
-	args << "--shader-validation" << BoolArg(info.shader_validation_enabled);
+	AddFlag(args, info.vulkan_validation_enabled, "--vulkan-validation");
+	AddFlag(args, info.shader_validation_enabled, "--shader-validation");
 	args << "--shader-optimization-type" << EnumToText(info.shader_optimization_type);
 	args << "--shader-log-direction" << EnumToText(info.shader_log_direction);
 	args << "--shader-log-folder" << info.shader_log_folder;
-	args << "--command-buffer-dump" << BoolArg(info.command_buffer_dump_enabled);
+	AddFlag(args, info.command_buffer_dump_enabled, "--command-buffer-dump");
 	args << "--command-buffer-dump-folder" << info.command_buffer_dump_folder;
 	args << "--printf-direction" << EnumToText(info.printf_direction);
 	args << "--printf-output-file" << info.printf_output_file;
-	if (info.profiler_enabled) {
-		args << "--profile";
-	}
-	args << "--spirv-debug-printf" << "false";
-	if (info.amd_cpu_enabled) {
-		args << "--amd-cpu";
-	}
+	AddFlag(args, info.profiler_enabled, "--profile");
+	AddFlag(args, info.amd_cpu_enabled, "--amd-cpu");
 #if defined(_WIN32)
-	if (info.red_zone_protection_enabled) {
-		args << "--redzone";
-	}
+	AddFlag(args, info.red_zone_protection_enabled, "--redzone");
 #endif
 	for (const auto& binding: info.host_input_mapping) {
 		args << "--keymap" << binding;
 	}
-	if (info.renderdoc_enabled) {
-		args << "--rd";
-	}
+	AddFlag(args, info.renderdoc_enabled, "--rd");
 
 	QString game = info.basedir;
 	if (!info.elf.isEmpty() && !GameContent::IsArchive(info.basedir)) {

src/launcher/src/patchesDialog.cpp

@@ -105,7 +105,17 @@ PatchesDialog::PatchesDialog(const Configuration& game, QWidget* parent)
 	auto* local_buttons = new QDialogButtonBox(local);
 	auto* reload        = local_buttons->addButton(tr("Reload"), QDialogButtonBox::ActionRole);
 	m_apply = local_buttons->addButton(tr("Apply selection"), QDialogButtonBox::ActionRole);
+	auto* enable_all    = local_buttons->addButton(tr("Enable all"), QDialogButtonBox::ActionRole);
+	auto* disable_all   = local_buttons->addButton(tr("Disable all"), QDialogButtonBox::ActionRole);
 	local_layout->addWidget(local_buttons);
+	auto set_all = [this](Qt::CheckState state) {
+		for (int i = 0; i < m_patches->count(); ++i)
+		{
+			m_patches->item(i)->setCheckState(state);
+		}
+	};
+	connect(enable_all, &QPushButton::clicked, this, [set_all] { set_all(Qt::Checked); });
+	connect(disable_all, &QPushButton::clicked, this, [set_all] { set_all(Qt::Unchecked); });
 	auto* buttons = new QDialogButtonBox(QDialogButtonBox::Close, this);
 	layout->addWidget(buttons);
 	connect(buttons, &QDialogButtonBox::rejected, this, &QDialog::close);

src/launcher/src/updateChecker.cpp

@@ -4,6 +4,7 @@
 
 #include <QAbstractButton>
 #include <QApplication>
+#include <QDateTime>
 #include <QDebug>
 #include <QDesktopServices>
 #include <QDir>
@@ -96,7 +97,27 @@ bool UpdateChecker::IsSupported() {
 
 void UpdateChecker::ShowPreviousResult(QWidget* parent) {
 	const auto directory = InstallDirectory();
-	QFile      result(QDir(directory).filePath(QStringLiteral(".kyty-update-result")));
+	if (!directory.isEmpty()) {
+		// The helper cannot delete its own executable on Windows, so its workspace is
+		// removed by the next launcher start.
+		const QDir install(directory);
+		for (const auto& name:
+		     install.entryList({QStringLiteral(".kyty-update-??????")},
+		                       QDir::Dirs | QDir::Hidden | QDir::NoSymLinks | QDir::NoDotAndDotDot)) {
+			QDir       workspace(install.filePath(name));
+			QDateTime  newest = QFileInfo(workspace.path()).lastModified();
+			for (const auto& entry: workspace.entryInfoList(QDir::Files | QDir::Dirs | QDir::Hidden |
+			                                                QDir::NoDotAndDotDot)) {
+				newest = qMax(newest, entry.lastModified());
+			}
+			// A workspace touched recently may belong to a running download or helper. Windows
+			// lists a file still being downloaded with a stale time, so allow a generous margin.
+			if (newest.secsTo(QDateTime::currentDateTime()) > 60 * 60) {
+				workspace.removeRecursively();
+			}
+		}
+	}
+	QFile result(QDir(directory).filePath(QStringLiteral(".kyty-update-result")));
 	if (directory.isEmpty() || !result.open(QIODevice::ReadOnly)) {
 		return;
 	}

src/libs/agc.cpp

@@ -1439,14 +1439,6 @@ int KYTY_SYSV_ABI AgcCreateInterpolantMapping2(ShaderRegister* regs, const Shade
 int KYTY_SYSV_ABI AgcGetDataPacketPayloadAddress(uint32_t** addr, uint32_t* cmd, int type) {
 	PRINT_NAME();
 
-	static std::atomic<uint32_t> log_count {0};
-	if (log_count.fetch_add(1) < 64) {
-		LOGF("\t addr = 0x%016" PRIx64 "\n"
-		     "\t cmd  = 0x%016" PRIx64 "\n"
-		     "\t type = %d\n",
-		     reinterpret_cast<uint64_t>(addr), reinterpret_cast<uint64_t>(cmd), type);
-	}
-
 	EXIT_NOT_IMPLEMENTED(addr == nullptr);
 
 	if (type != 0) {
@@ -1463,20 +1455,10 @@ int KYTY_SYSV_ABI AgcGetDataPacketPayloadAddress(uint32_t** addr, uint32_t* cmd,
 int KYTY_SYSV_ABI AgcGetDataPacketPayloadRange(MemoryRange* range, uint32_t* cmd, int type) {
 	PRINT_NAME();
 
-	static std::atomic<uint32_t> log_count {0};
-	const bool                   should_log = (log_count.fetch_add(1) < 64);
-	if (should_log) {
-		LOGF("\t range = 0x%016" PRIx64 "\n"
-		     "\t cmd   = 0x%016" PRIx64 "\n"
-		     "\t type  = %d\n",
-		     reinterpret_cast<uint64_t>(range), reinterpret_cast<uint64_t>(cmd), type);
-	}
-
 	EXIT_NOT_IMPLEMENTED(range == nullptr);
 	EXIT_NOT_IMPLEMENTED(cmd == nullptr);
 
 	const auto cmd_id = cmd[0];
-	const auto len    = KYTY_PM4_LEN(cmd_id);
 
 	if (type != 0) {
 		range->m_base = cmd + 2;
@@ -1491,13 +1473,6 @@ int KYTY_SYSV_ABI AgcGetDataPacketPayloadRange(MemoryRange* range, uint32_t* cmd
 		}
 	}
 
-	if (should_log) {
-		LOGF("\t cmd_id = 0x%08" PRIx32 ", len = %" PRIu32 "\n"
-		     "\t base   = 0x%016" PRIx64 "\n"
-		     "\t size   = 0x%016" PRIx64 "\n",
-		     cmd_id, len, reinterpret_cast<uint64_t>(range->m_base), range->m_size);
-	}
-
 	return OK;
 }
 
@@ -1746,15 +1721,6 @@ uint32_t* KYTY_SYSV_ABI AgcCbDispatch(CommandBuffer* buf, uint32_t thread_group_
                                       uint32_t modifier) {
 	PRINT_NAME();
 
-	static std::atomic<uint32_t> log_count {0};
-	if (log_count.fetch_add(1) < 64) {
-		LOGF("\t thread_group_x = %" PRIu32 "\n"
-		     "\t thread_group_y = %" PRIu32 "\n"
-		     "\t thread_group_z = %" PRIu32 "\n"
-		     "\t modifier       = 0x%08" PRIx32 "\n",
-		     thread_group_x, thread_group_y, thread_group_z, modifier);
-	}
-
 	if (buf == nullptr) {
 		return nullptr;
 	}
@@ -1862,16 +1828,6 @@ uint32_t* KYTY_SYSV_ABI AgcCbSetShRegisterRangeDirect(CommandBuffer* buf, uint32
                                                       const uint32_t* values, uint32_t num_values) {
 	PRINT_NAME();
 
-	static std::atomic<uint32_t> log_count {0};
-	if (log_count.fetch_add(1) < 64) {
-		LOGF("\t buf        = 0x%016" PRIx64 "\n"
-		     "\t offset     = %" PRIx32 "\n"
-		     "\t values     = 0x%016" PRIx64 "\n"
-		     "\t num_values = %" PRIu32 "\n",
-		     reinterpret_cast<uint64_t>(buf), offset, reinterpret_cast<uint64_t>(values),
-		     num_values);
-	}
-
 	EXIT_NOT_IMPLEMENTED(buf == nullptr);
 
 	auto* cmd = buf->AllocateDW(num_values + 2);
@@ -1902,12 +1858,6 @@ uint32_t KYTY_SYSV_ABI AgcCbSetShRegisterRangeDirectGetSize(uint32_t num_values)
 uint32_t* KYTY_SYSV_ABI AgcCbSetShRegistersDirect(CommandBuffer*                 buf,
                                                   const volatile ShaderRegister* regs,
                                                   uint32_t                       num_regs) {
-	PRINT_NAME();
-
-	AgcTrace("\t regs     = 0x%016" PRIx64 "\n"
-	     "\t num_regs = %" PRIu32 "\n",
-	     reinterpret_cast<uint64_t>(regs), num_regs);
-
 	if (num_regs == 0) {
 		return nullptr;
 	}
@@ -2047,23 +1997,6 @@ uint32_t* KYTY_SYSV_ABI AgcCbReleaseMem(CommandBuffer* buf, uint8_t action, uint
                                         uint8_t interrupt, uint32_t interrupt_ctx_id) {
 	PRINT_NAME();
 
-	static std::atomic<uint32_t> log_count {0};
-	if (log_count.fetch_add(1) < 64) {
-		LOGF("\t action           = 0x%02" PRIx8 "\n"
-		     "\t gcr_cntl         = 0x%04" PRIx16 "\n"
-		     "\t dst              = %" PRIu8 "\n"
-		     "\t cache_policy     = 0x%02" PRIx8 "\n"
-		     "\t address          = 0x%016" PRIx64 "\n"
-		     "\t data_sel         = 0x%02" PRIx8 "\n"
-		     "\t data             = 0x%016" PRIx64 "\n"
-		     "\t gds_offset       = %" PRIu16 "\n"
-		     "\t gds_size         = %" PRIu16 "\n"
-		     "\t interrupt        = 0x%02" PRIx8 "\n"
-		     "\t interrupt_ctx_id = %" PRIu32 "\n",
-		     action, gcr_cntl, dst, cache_policy, reinterpret_cast<uint64_t>(address), data_sel,
-		     data, gds_offset, gds_size, interrupt, interrupt_ctx_id);
-	}
-
 	EXIT_NOT_IMPLEMENTED(buf == nullptr);
 	EXIT_NOT_IMPLEMENTED(dst > 1);
 	EXIT_NOT_IMPLEMENTED(data_sel != 0 && data_sel != 1 && data_sel != 2 && data_sel != 3 &&
@@ -2950,13 +2883,6 @@ uint32_t* KYTY_SYSV_ABI AgcDcbEventWrite(CommandBuffer* buf, uint8_t event_type,
                                          const volatile void* address) {
 	PRINT_NAME();
 
-	static std::atomic<uint32_t> log_count {0};
-	if (log_count.fetch_add(1) < 64) {
-		LOGF("\t event_type = 0x%02" PRIx8 "\n"
-		     "\t address    = 0x%016" PRIx64 "\n",
-		     event_type, reinterpret_cast<uint64_t>(address));
-	}
-
 	EXIT_NOT_IMPLEMENTED(buf == nullptr);
 	EXIT_NOT_IMPLEMENTED(event_type > 0x3fu);
 
@@ -3011,17 +2937,6 @@ uint32_t* KYTY_SYSV_ABI AgcDcbAcquireMem(CommandBuffer* buf, uint8_t engine, uin
                                          uint64_t size_bytes, uint32_t poll_cycles) {
 	PRINT_NAME();
 
-	static std::atomic<uint32_t> log_count {0};
-	if (log_count.fetch_add(1) < 64) {
-		LOGF("\t engine      = 0x%02" PRIx8 "\n"
-		     "\t cb_db_op    = 0x%08" PRIx32 "\n"
-		     "\t gcr_cntl    = 0x%08" PRIx32 "\n"
-		     "\t base        = 0x%016" PRIx64 "\n"
-		     "\t size_bytes  = 0x%016" PRIx64 "\n"
-		     "\t poll_cycles = %" PRIu32 "\n",
-		     engine, cb_db_op, gcr_cntl, reinterpret_cast<uint64_t>(base), size_bytes, poll_cycles);
-	}
-
 	EXIT_NOT_IMPLEMENTED(buf == nullptr);
 
 	auto* cmd = buf->AllocateDW(8);

src/libs/audio.cpp

@@ -841,6 +841,12 @@ struct AudioOutPortState {
 	uint64_t reserved2[2];
 };
 
+struct AudioOutSystemState {
+	float    loudness;
+	uint8_t  reserved8[4];
+	uint64_t reserved64[3];
+};
+
 int KYTY_SYSV_ABI AudioOutInit() {
 	PRINT_NAME();
 
@@ -952,6 +958,19 @@ int KYTY_SYSV_ABI AudioOutGetPortState(int handle, AudioOutPortState* state) {
 	return OK;
 }
 
+int KYTY_SYSV_ABI AudioOutGetSystemState(AudioOutSystemState* state) {
+	PRINT_NAME();
+
+	if (state == nullptr) {
+		return AUDIO_OUT_ERROR_INVALID_POINTER;
+	}
+
+	// Loudness is undefined until measurement starts; no measurement is emulated.
+	*state = {};
+
+	return OK;
+}
+
 int KYTY_SYSV_ABI AudioOutSetVolume(int handle, uint32_t flag, int* vol) {
 	PRINT_NAME();
 

src/libs/audio.h

@@ -18,6 +18,7 @@ namespace AudioOut {
 
 struct AudioOutOutputParam;
 struct AudioOutPortState;
+struct AudioOutSystemState;
 
 int KYTY_SYSV_ABI AudioOutInit();
 int KYTY_SYSV_ABI AudioOutOpen(int user_id, int type, int index, uint32_t len, uint32_t freq,
@@ -27,6 +28,7 @@ int KYTY_SYSV_ABI AudioOutOutputs(AudioOutOutputParam* param, uint32_t num);
 int KYTY_SYSV_ABI AudioOutOutput(int handle, const void* ptr);
 int KYTY_SYSV_ABI AudioOutClose(int handle);
 int KYTY_SYSV_ABI AudioOutGetPortState(int handle, AudioOutPortState* state);
+int KYTY_SYSV_ABI AudioOutGetSystemState(AudioOutSystemState* state);
 
 } // namespace AudioOut
 

src/libs/libAudio.cpp

@@ -131,6 +131,7 @@ LIB_DEFINE(InitAudio_1_AudioOut) {
 	LIB_FUNC("QOQtbeDqsT4", AudioOut::AudioOutOutput);
 	LIB_FUNC("s1--uE9mBFw", AudioOut::AudioOutClose);
 	LIB_FUNC("GrQ9s4IrNaQ", AudioOut::AudioOutGetPortState);
+	LIB_FUNC("R5hemoKKID8", AudioOut::AudioOutGetSystemState);
 }
 
 } // namespace LibAudioOut

src/libs/libAudio2.cpp

@@ -178,7 +178,6 @@ struct AudioOut2ContextState {
 	uint32_t               num_grains          = 512;
 	uint64_t               last_update         = 0;
 	int                    object_audio_handle = 0;
-	std::vector<float>     object_mix;
 };
 
 struct AudioOut2PortStateEntry {
@@ -212,8 +211,18 @@ struct AudioOut2LatencyState {
 	uint32_t latency_us = 0;
 };
 
+// Per-context copy of what a push hands the backend. The backend can sleep while pacing, so a push
+// holds only its context's mutex across it; closing the context's handles waits on that mutex.
+struct AudioOut2Submission {
+	Common::Mutex                                                   mutex;
+	std::array<std::vector<uint8_t>, AudioInternal::OUT_PORTS_MAX>  pcm;
+	std::array<std::array<float, 16>, AudioInternal::OUT_PORTS_MAX> gains;
+	std::vector<float>                                              object_mix;
+};
+
 static Common::Mutex                              g_audioout2_state_mutex;
 static std::array<AudioOut2ContextState, 16>      g_audioout2_contexts;
+static std::array<AudioOut2Submission, 16>        g_audioout2_submissions;
 static std::array<AudioOut2PortStateEntry, 256>   g_audioout2_ports;
 static std::vector<AudioOut2UserHandle>          g_audioout2_users;
 static Common::Mutex                              g_audioout2_speaker_array_mutex;
@@ -249,6 +258,10 @@ static AudioOut2ContextState* audioout2_find_context_locked(AudioOut2ContextHand
 	return nullptr;
 }
 
+static AudioOut2Submission& audioout2_submission(const AudioOut2ContextState& context) {
+	return g_audioout2_submissions[&context - g_audioout2_contexts.data()];
+}
+
 static uint32_t audioout2_grain_micros(uint32_t grains) {
 	const auto sample_count = (grains == 0 ? 512u : grains);
 	return std::max<uint32_t>((sample_count * 1000000u) / 48000u, 1000u);
@@ -372,11 +385,10 @@ static bool audioout2_context_has_queueable_device(const AudioOut2ContextState&
 	return false;
 }
 
-static void audioout2_mix_object(AudioOut2ContextState&         context,
+static void audioout2_mix_object(std::vector<float>& mix, uint32_t num_grains,
                                  const AudioOut2PortStateEntry& object) {
-	auto& mix = context.object_mix;
 	if (mix.empty()) {
-		mix.resize(static_cast<size_t>(context.num_grains) * 2, 0.0f);
+		mix.resize(static_cast<size_t>(num_grains) * 2, 0.0f);
 	}
 
 	// Center ordinary objects. First-order stereo ambisonics uses left/right-facing
@@ -393,7 +405,7 @@ static void audioout2_mix_object(AudioOut2ContextState&         context,
 	}
 	left *= object.gains[0];
 	right *= object.gains[0];
-	for (uint32_t frame = 0; frame < context.num_grains; frame++) {
+	for (uint32_t frame = 0; frame < num_grains; frame++) {
 		const float sample =
 		    object.audio_format == AudioInternal::Format::FloatMono
 		        ? reinterpret_cast<const float*>(object.pcm_data.data())[frame]
@@ -403,15 +415,16 @@ static void audioout2_mix_object(AudioOut2ContextState&         context,
 	}
 }
 
-static void audioout2_queue_context_audio(AudioOut2ContextHandle ctx, bool blocking) {
-	// The backend consumes the buffers synchronously, but can block while pacing SDL's queue.
-	// Keep both PCM storage and port handles alive until it returns.
-	Common::LockGuard lock(g_audioout2_state_mutex);
-	auto*             context = audioout2_find_context_locked(ctx);
+static void audioout2_queue_context_audio(AudioOut2ContextHandle ctx, AudioOut2Submission& out,
+                                          bool blocking) {
+	Common::LockGuard submit_lock(out.mutex);
+	g_audioout2_state_mutex.Lock();
+	auto* context = audioout2_find_context_locked(ctx);
 	if (context == nullptr) {
+		g_audioout2_state_mutex.Unlock();
 		return;
 	}
-	context->object_mix.clear();
+	out.object_mix.clear();
 	std::vector<AudioInternal::OutputParam> params;
 	params.reserve(AudioInternal::OUT_PORTS_MAX);
 	for (const auto& state: g_audioout2_ports) {
@@ -419,14 +432,18 @@ static void audioout2_queue_context_audio(AudioOut2ContextHandle ctx, bool block
 			continue;
 		}
 		if (audioout2_port_type_is_object(state.port_type)) {
-			audioout2_mix_object(*context, state);
+			audioout2_mix_object(out.object_mix, context->num_grains, state);
 		} else if (state.audio_handle > 0 && params.size() < AudioInternal::OUT_PORTS_MAX) {
-			params.push_back({state.audio_handle, state.pcm_data.data(), state.gains.data()});
+			const auto i = params.size();
+			out.pcm[i]   = state.pcm_data;
+			out.gains[i] = state.gains;
+			params.push_back({state.audio_handle, out.pcm[i].data(), out.gains[i].data()});
 		}
 	}
-	if (!context->object_mix.empty() && params.size() < AudioInternal::OUT_PORTS_MAX) {
-		params.push_back({context->object_audio_handle, context->object_mix.data()});
+	if (!out.object_mix.empty() && params.size() < AudioInternal::OUT_PORTS_MAX) {
+		params.push_back({context->object_audio_handle, out.object_mix.data()});
 	}
+	g_audioout2_state_mutex.Unlock();
 	if (!params.empty()) {
 		(void)AudioInternal::AudioOutOutputs(params.data(), static_cast<uint32_t>(params.size()),
 		                                     blocking);
@@ -528,7 +545,8 @@ int KYTY_SYSV_ABI AudioOut2ContextDestroy(AudioOut2ContextHandle ctx) {
 	size_t audio_handles_num = 0;
 
 	g_audioout2_state_mutex.Lock();
-	if (auto* state = audioout2_find_context_locked(ctx); state != nullptr) {
+	auto* state = audioout2_find_context_locked(ctx);
+	if (state != nullptr) {
 		if (state->object_audio_handle > 0) {
 			audio_handles[audio_handles_num++] = state->object_audio_handle;
 		}
@@ -544,6 +562,10 @@ int KYTY_SYSV_ABI AudioOut2ContextDestroy(AudioOut2ContextHandle ctx) {
 	}
 	g_audioout2_state_mutex.Unlock();
 
+	if (state != nullptr) {
+		// Wait for an in-flight push that may still use the handles.
+		Common::LockGuard wait(audioout2_submission(*state).mutex);
+	}
 	for (size_t i = 0; i < audio_handles_num; i++) {
 		audioout2_close_audio_handle(audio_handles[i]);
 	}
@@ -595,8 +617,9 @@ int KYTY_SYSV_ABI AudioOut2ContextPush(AudioOut2ContextHandle ctx, uint32_t bloc
 				if (state->queued < state->queue_depth) {
 					state->queued++;
 				}
+				auto& submission = audioout2_submission(*state);
 				g_audioout2_state_mutex.Unlock();
-				audioout2_queue_context_audio(ctx, blocking != 0);
+				audioout2_queue_context_audio(ctx, submission, blocking != 0);
 				return OK;
 			}
 		}
@@ -729,14 +752,20 @@ int KYTY_SYSV_ABI AudioOut2PortDestroy(AudioOut2PortHandle port) {
 	PRINT_NAME();
 	LOGF("\t port = 0x%016" PRIx64 "\n", port);
 
-	int audio_handle = 0;
+	int                          audio_handle = 0;
+	const AudioOut2ContextState* context      = nullptr;
 	g_audioout2_state_mutex.Lock();
 	if (auto* state = audioout2_find_port_locked(port); state != nullptr) {
 		audio_handle = state->audio_handle;
+		context      = audioout2_find_context_locked(state->context);
 		*state       = AudioOut2PortStateEntry {};
 	}
 	g_audioout2_state_mutex.Unlock();
 
+	if (audio_handle > 0 && context != nullptr) {
+		// Wait for an in-flight push that may still use the handle.
+		Common::LockGuard wait(audioout2_submission(*context).mutex);
+	}
 	audioout2_close_audio_handle(audio_handle);
 
 	return OK;

src/libs/ngs2.cpp

@@ -541,7 +541,6 @@ struct Ngs2VoiceInternal {
 	struct Block {
 		const uint8_t*    data;
 		Ngs2WaveformBlock info;
-		uint32_t          cursor         = 0;
 		size_t            data_cursor    = 0;
 		uint32_t          repeated_count = 0;
 	};
@@ -595,10 +594,15 @@ struct Ngs2VoiceInternal {
 			return waveform_end;
 		}
 		const auto& block = blocks.front();
-		return block.data + (decoder != nullptr
-		                         ? block.data_cursor
-		                         : (uint64_t(block.info.num_skip_samples) + block.cursor) *
-		                               channels * sizeof(int16_t));
+		if (decoder != nullptr) {
+			return block.data + block.data_cursor;
+		}
+		// PCM never reads skipped samples, so the next read starts past a pending skip.
+		const auto skip = block.info.num_samples != 0 && block.data_cursor == 0
+		                      ? block.info.num_skip_samples
+		                      : skip_remaining;
+		return block.data + std::min(block.data_cursor + size_t(skip) * channels * sizeof(int16_t),
+		                             block.info.data_size);
 	}
 
 	// Custom racks keep num_buffers consecutive channel-major buffers; other racks keep one.
@@ -687,7 +691,6 @@ struct Ngs2VoiceInternal {
 		const bool clear     = (flags & 4u) != 0;
 		const bool only_data = (flags & 2u) != 0;
 		const bool join      = only_data && (flags & 0x10u) == 0;
-		EXIT_NOT_IMPLEMENTED(only_data && decoder == nullptr);
 		const auto retained = clear ? 0 : blocks.size();
 		Block      tail     = retained == 0 ? Block {} : blocks.back();
 		std::vector<Block> pending;
@@ -703,10 +706,6 @@ struct Ngs2VoiceInternal {
 				info.num_skip_samples = 0;
 				info.num_samples      = 0;
 			}
-			EXIT_NOT_IMPLEMENTED(decoder == nullptr &&
-			                     (uint64_t(info.num_skip_samples) + info.num_samples) *
-			                             channels * sizeof(int16_t) >
-			                         info.data_size);
 			Block block {data + info.data_offset, info};
 			auto* previous = pending.empty() ? (retained == 0 ? nullptr : &tail)
 			                                 : &pending.back();
@@ -1546,9 +1545,8 @@ static bool Ngs2FinishBlock(Ngs2VoiceInternal& voice) {
 			--block.info.num_repeats;
 		}
 		++block.repeated_count;
-		block.cursor = 0;
+		block.data_cursor = 0;
 		if (voice.decoder != nullptr) {
-			block.data_cursor = 0;
 			voice.compressed_input.clear();
 			voice.decoder->Reset();
 		}
@@ -1573,37 +1571,8 @@ static bool Ngs2FinishBlock(Ngs2VoiceInternal& voice) {
 	return voice.source_generation == generation;
 }
 
-static bool Ngs2DecodePcmFrame(Ngs2VoiceInternal& voice, std::vector<float>& frame,
-                               uint32_t& frame_samples) {
-	// PCM has no encoded frame size. Use one system grain as a bounded staging
-	// quantum, keeping source decoding separate from resampled output consumption.
-	const auto quantum = std::max(voice.rack->ngs->option.num_grain_samples, 1u);
-	frame.resize(size_t(quantum) * voice.channels);
-	frame_samples = 0;
-	while (voice.state == Ngs2VoicePlayState::Playing && !voice.blocks.empty() &&
-	       frame_samples < quantum) {
-		auto& block = voice.blocks.front();
-		const auto count = std::min(quantum - frame_samples, block.info.num_samples - block.cursor);
-		const auto* pcm = reinterpret_cast<const int16_t*>(voice.WaveformData());
-		for (size_t i = 0; i < size_t(count) * voice.channels; ++i) {
-			frame[size_t(frame_samples) * voice.channels + i] = pcm[i] / 32768.0f;
-		}
-		block.cursor += count;
-		frame_samples += count;
-		voice.decoded_samples += count;
-		voice.decoded_bytes += size_t(count) * voice.channels * sizeof(int16_t);
-		if (block.cursor == block.info.num_samples && !Ngs2FinishBlock(voice)) {
-			return false;
-		}
-	}
-	return frame_samples != 0;
-}
-
 static bool Ngs2DecodeFrame(Ngs2VoiceInternal& voice, std::vector<float>& frame,
                             uint32_t& frame_samples) {
-	if (voice.decoder == nullptr) {
-		return Ngs2DecodePcmFrame(voice, frame, frame_samples);
-	}
 	frame_samples = 0;
 	while (voice.state == Ngs2VoicePlayState::Playing && !voice.blocks.empty()) {
 		auto& block = voice.blocks.front();
@@ -1617,6 +1586,41 @@ static bool Ngs2DecodeFrame(Ngs2VoiceInternal& voice, std::vector<float>& frame,
 			}
 			continue;
 		}
+		if (voice.decoder == nullptr) {
+			// PCM has no encoded frame size, so stage at most one system grain.
+			const auto quantum     = std::max(voice.rack->ngs->option.num_grain_samples, 1u);
+			const auto frame_bytes = voice.channels * sizeof(int16_t);
+			const auto available   = (block.info.data_size - block.data_cursor) / frame_bytes;
+			const auto skip        = std::min<size_t>(voice.skip_remaining, available);
+			const auto count =
+			    std::min<size_t>({available - skip, voice.duration_remaining, quantum});
+			frame.resize(size_t(quantum) * voice.channels);
+			const auto* pcm = reinterpret_cast<const int16_t*>(block.data + block.data_cursor) +
+			                  skip * voice.channels;
+			for (size_t i = 0; i < count * voice.channels; ++i) {
+				frame[i] = pcm[i] / 32768.0f;
+			}
+			voice.skip_remaining -= skip;
+			voice.duration_remaining -= count;
+			block.data_cursor += (skip + count) * frame_bytes;
+			voice.decoded_bytes += (skip + count) * frame_bytes;
+			voice.decoded_samples += count;
+			frame_samples = static_cast<uint32_t>(count);
+			if (voice.duration_remaining == 0 ||
+			    block.info.data_size - block.data_cursor < frame_bytes) {
+				// A repeating block restarts its duration, so it must hold all of it.
+				EXIT_NOT_IMPLEMENTED(voice.duration_remaining != 0 &&
+				                     (block.data_cursor != block.info.data_size ||
+				                      block.info.num_repeats != 0));
+				if (!Ngs2FinishBlock(voice)) {
+					return false;
+				}
+			}
+			if (count != 0) {
+				return true;
+			}
+			continue;
+		}
 		const auto skip =
 		    std::min(voice.skip_remaining,
 		             static_cast<uint32_t>(frame.size() / voice.channels));

src/main.cpp

@@ -9,6 +9,7 @@
 #include "emulator.h"
 #include "kytyGitVersion.h"
 
+#include <algorithm>
 #include <charconv>
 #include <cstdio>
 #include <filesystem>
@@ -71,35 +72,31 @@ static void PrintUsage() {
 	::printf("  --amd-cpu                            Apply AMD CPU instruction patches.\n");
 	::printf("  --vblank-frequency <num>             Virtual vblank frequency. Default: 60.\n");
 	::printf("  --console-language <0-29>            Console language. Default: 1 (English US).\n");
-	::printf("  --vulkan-validation <true|false>     Enable Vulkan validation.\n");
-	::printf("  --network-connected <true|false>     Use the host local network (default: false).\n");
-	::printf("  --gpu-assisted-validation <t|f>      Bounds-check shader accesses on the GPU.\n"
+	::printf("  --vulkan-validation                  Enable Vulkan validation.\n");
+	::printf("  --network-connected                  Use the host local network.\n");
+	::printf("  --gpu-assisted-validation            Bounds-check shader accesses on the GPU.\n"
 	         "                                       Implies --vulkan-validation; very slow.\n");
-	::printf("  --shader-validation <true|false>     Enable shader validation.\n");
+	::printf("  --shader-validation                  Enable shader validation.\n");
 	::printf("  --tessellation                      Draw tessellation patches; skipped by default.\n");
 	::printf("  --shader-optimization-type <value>   None, Size, or Performance.\n");
 	::printf("  --shader-log-direction <value>       Silent, Console, or File.\n");
 	::printf("  --shader-log-folder <path>           Shader log output folder.\n");
-	::printf("  --command-buffer-dump <true|false>   Enable command buffer dumps.\n");
+	::printf("  --command-buffer-dump                Enable command buffer dumps.\n");
 	::printf("  --command-buffer-dump-folder <path>  Command buffer dump folder.\n");
-	::printf("  --graphics-debug-dump <true|false>   Enable graphics debug dumps.\n");
+	::printf("  --graphics-debug-dump                Enable graphics debug dumps.\n");
 	::printf("  --printf-direction <value>           Silent, Console, or File.\n");
 	::printf("  --printf-output-file <path>          Guest printf output file.\n");
 	::printf("  --profile                            Enable the Tracy profiler.\n");
-	::printf("  --spirv-debug-printf <true|false>    Enable SPIR-V debug printf.\n");
-	::printf(
-	    "  --readback-linear-images <true|false> Read back writable linear images on submit.\n");
-	::printf(
-	    "  --sync-raw-image-buffers <true|false> Synchronize raw reads of GPU images. Default: false.\n");
-	::printf(
-	    "  --trophy-notifications <true|false>   Show trophy unlock toasts and play their sound.\n");
+	::printf("  --spirv-debug-printf                 Enable SPIR-V debug printf.\n");
+	::printf("  --readback-linear-images             Read back writable linear images on submit.\n");
+	::printf("  --readback-tiled-images              Read back small writable tiled images on submit.\n");
+	::printf("  --sync-raw-image-buffers             Synchronize raw reads of GPU images.\n");
+	::printf("  --no-trophy-notifications            Disable trophy unlock toasts and their sound.\n");
 	::printf(
 	    "  --trophy-duration <seconds>          Trophy toast duration, 5-10. Default: 6.\n");
-	::printf("  --perf-overlay <true|false>           Show the performance panel (F2 or Cmd+P).\n");
+	::printf("  --perf-overlay                       Show the performance panel (F2 or Cmd+P).\n");
 	::printf("  --playgo-hack                       Use the supplied PlayGo stub fallback.\n");
-	::printf(
-	    "  --skip-notice-screen <true|false>    Skip startup logos and notices in supported games.\n"
-	    "                                      Default: false.\n");
+	::printf("  --skip-notice-screen                 Skip startup logos and notices in supported games.\n");
 #if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
 	::printf("  --redzone                            Protect the guest SysV red zone.\n");
 #endif
@@ -117,22 +114,6 @@ static bool NextArg(int argc, char* argv[], int& index, std::string& out) {
 	return true;
 }
 
-static bool ParseBool(const std::string& value, bool& out) {
-	if (Common::EqualNoCase(value, "true") || value == "1" || Common::EqualNoCase(value, "yes") ||
-	    Common::EqualNoCase(value, "on")) {
-		out = true;
-		return true;
-	}
-
-	if (Common::EqualNoCase(value, "false") || value == "0" || Common::EqualNoCase(value, "no") ||
-	    Common::EqualNoCase(value, "off")) {
-		out = false;
-		return true;
-	}
-
-	return false;
-}
-
 template <typename E>
 static bool ParseEnum(const std::string& value, E& out) {
 	auto enum_value = magic_enum::enum_cast<E>(value.c_str());
@@ -200,6 +181,43 @@ static bool ParseUserId(const std::string& value, int32_t& out) {
 	return true;
 }
 
+using Config::ConfigOptions;
+
+struct Flag {
+	std::string_view name;
+	bool ConfigOptions::* field;
+	bool                  value = true;
+};
+
+static constexpr Flag FLAGS[] = {
+    {"--rd", &ConfigOptions::renderdoc_enabled},
+    {"--fullscreen", &ConfigOptions::fullscreen_enabled},
+    {"--hdr", &ConfigOptions::hdr_enabled},
+    {"--hide-cursor", &ConfigOptions::hide_cursor_enabled},
+    {"--no-splash", &ConfigOptions::splash_enabled, false},
+    {"--vr", &ConfigOptions::vr_enabled},
+    {"--amd-cpu", &ConfigOptions::amd_cpu_enabled},
+    {"--network-connected", &ConfigOptions::network_connected},
+    {"--vulkan-validation", &ConfigOptions::vulkan_validation_enabled},
+    {"--gpu-assisted-validation", &ConfigOptions::gpu_assisted_validation_enabled},
+    {"--shader-validation", &ConfigOptions::shader_validation_enabled},
+    {"--tessellation", &ConfigOptions::tessellation_enabled},
+    {"--command-buffer-dump", &ConfigOptions::command_buffer_dump_enabled},
+    {"--graphics-debug-dump", &ConfigOptions::graphics_debug_dump_enabled},
+    {"--profile", &ConfigOptions::profiler_enabled},
+    {"--spirv-debug-printf", &ConfigOptions::spirv_debug_printf_enabled},
+    {"--readback-linear-images", &ConfigOptions::readback_linear_images},
+    {"--readback-tiled-images", &ConfigOptions::readback_tiled_images},
+    {"--sync-raw-image-buffers", &ConfigOptions::sync_raw_image_buffers},
+    {"--no-trophy-notifications", &ConfigOptions::trophy_enabled, false},
+    {"--perf-overlay", &ConfigOptions::perf_overlay_enabled},
+    {"--playgo-hack", &ConfigOptions::playgo_hack_enabled},
+    {"--skip-notice-screen", &ConfigOptions::skip_notice_screen},
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+    {"--redzone", &ConfigOptions::red_zone_protection_enabled},
+#endif
+};
+
 static bool ParseArgs(int argc, char* argv[], RunOptions& options, bool& show_help) {
 	show_help = false;
 
@@ -212,63 +230,11 @@ static bool ParseArgs(int argc, char* argv[], RunOptions& options, bool& show_he
 			continue;
 		}
 
-		if (arg == "--rd") {
-			options.config.renderdoc_enabled = true;
+		if (const auto* flag = std::ranges::find(FLAGS, arg, &Flag::name); flag != std::end(FLAGS)) {
+			options.config.*flag->field = flag->value;
 			continue;
 		}
 
-		if (arg == "--fullscreen") {
-			options.config.fullscreen_enabled = true;
-			continue;
-		}
-
-		if (arg == "--hdr") {
-			options.config.hdr_enabled = true;
-			continue;
-		}
-
-		if (arg == "--hide-cursor") {
-			options.config.hide_cursor_enabled = true;
-			continue;
-		}
-
-		if (arg == "--no-splash") {
-			options.config.splash_enabled = false;
-			continue;
-		}
-
-		if (arg == "--vr") {
-			options.config.vr_enabled = true;
-			continue;
-		}
-
-		if (arg == "--amd-cpu") {
-			options.config.amd_cpu_enabled = true;
-			continue;
-		}
-
-		if (arg == "--playgo-hack") {
-			options.config.playgo_hack_enabled = true;
-			continue;
-		}
-
-		if (arg == "--tessellation") {
-			options.config.tessellation_enabled = true;
-			continue;
-		}
-
-		if (arg == "--profile") {
-			options.config.profiler_enabled = true;
-			continue;
-		}
-
-#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
-		if (arg == "--redzone") {
-			options.config.red_zone_protection_enabled = true;
-			continue;
-		}
-#endif
-
 		if (!arg.starts_with("--")) {
 			::printf("game input must be provided with --game\n");
 			return false;
@@ -411,26 +377,6 @@ static bool ParseArgs(int argc, char* argv[], RunOptions& options, bool& show_he
 				::printf("invalid console language: %s\n", value.c_str());
 				return false;
 			}
-		} else if (arg == "--network-connected") {
-			if (!ParseBool(value, options.config.network_connected)) {
-				::printf("invalid boolean for %s: %s\n", arg.c_str(), value.c_str());
-				return false;
-			}
-		} else if (arg == "--vulkan-validation") {
-			if (!ParseBool(value, options.config.vulkan_validation_enabled)) {
-				::printf("invalid boolean for %s: %s\n", arg.c_str(), value.c_str());
-				return false;
-			}
-		} else if (arg == "--gpu-assisted-validation") {
-			if (!ParseBool(value, options.config.gpu_assisted_validation_enabled)) {
-				::printf("invalid boolean for %s: %s\n", arg.c_str(), value.c_str());
-				return false;
-			}
-		} else if (arg == "--shader-validation") {
-			if (!ParseBool(value, options.config.shader_validation_enabled)) {
-				::printf("invalid boolean for %s: %s\n", arg.c_str(), value.c_str());
-				return false;
-			}
 		} else if (arg == "--shader-optimization-type") {
 			if (!ParseEnum(value, options.config.shader_optimization_type)) {
 				::printf("invalid shader optimization type: %s\n", value.c_str());
@@ -443,18 +389,8 @@ static bool ParseArgs(int argc, char* argv[], RunOptions& options, bool& show_he
 			}
 		} else if (arg == "--shader-log-folder") {
 			options.config.shader_log_folder = Common::PathFromUtf8(value);
-		} else if (arg == "--command-buffer-dump") {
-			if (!ParseBool(value, options.config.command_buffer_dump_enabled)) {
-				::printf("invalid boolean for %s: %s\n", arg.c_str(), value.c_str());
-				return false;
-			}
 		} else if (arg == "--command-buffer-dump-folder") {
 			options.config.command_buffer_dump_folder = Common::PathFromUtf8(value);
-		} else if (arg == "--graphics-debug-dump") {
-			if (!ParseBool(value, options.config.graphics_debug_dump_enabled)) {
-				::printf("invalid boolean for %s: %s\n", arg.c_str(), value.c_str());
-				return false;
-			}
 		} else if (arg == "--printf-direction") {
 			if (!ParseEnum(value, options.config.printf_direction)) {
 				::printf("invalid printf direction: %s\n", value.c_str());
@@ -462,36 +398,6 @@ static bool ParseArgs(int argc, char* argv[], RunOptions& options, bool& show_he
 			}
 		} else if (arg == "--printf-output-file") {
 			options.config.printf_output_file = Common::PathFromUtf8(value);
-		} else if (arg == "--spirv-debug-printf") {
-			if (!ParseBool(value, options.config.spirv_debug_printf_enabled)) {
-				::printf("invalid boolean for %s: %s\n", arg.c_str(), value.c_str());
-				return false;
-			}
-		} else if (arg == "--readback-linear-images") {
-			if (!ParseBool(value, options.config.readback_linear_images)) {
-				::printf("invalid boolean for %s: %s\n", arg.c_str(), value.c_str());
-				return false;
-			}
-		} else if (arg == "--sync-raw-image-buffers") {
-			if (!ParseBool(value, options.config.sync_raw_image_buffers)) {
-				::printf("invalid boolean for %s: %s\n", arg.c_str(), value.c_str());
-				return false;
-			}
-		} else if (arg == "--trophy-notifications") {
-			if (!ParseBool(value, options.config.trophy_enabled)) {
-				::printf("invalid boolean for %s: %s\n", arg.c_str(), value.c_str());
-				return false;
-			}
-		} else if (arg == "--skip-notice-screen") {
-			if (!ParseBool(value, options.config.skip_notice_screen)) {
-				::printf("invalid boolean for %s: %s\n", arg.c_str(), value.c_str());
-				return false;
-			}
-		} else if (arg == "--perf-overlay") {
-			if (!ParseBool(value, options.config.perf_overlay_enabled)) {
-				::printf("invalid boolean for %s: %s\n", arg.c_str(), value.c_str());
-				return false;
-			}
 		} else if (arg == "--keymap") {
 			const auto split = value.find('=');
 			if (split == std::string::npos || split == 0 || split + 1 == value.size()) {

tests/ArchiveFileTests.cpp

@@ -12,6 +12,7 @@
 #include <cstdlib>
 #include <cstring>
 #include <filesystem>
+#include <fstream>
 #include <limits>
 #include <string>
 #include <thread>
@@ -143,6 +144,72 @@ void CheckMalformedArchives(const std::filesystem::path &source,
         "recover by reading an unaffected compressed block");
 }
 
+void WriteHostFile(const std::filesystem::path &path, const void *data, size_t size) {
+  std::ofstream stream(path, std::ios::binary);
+  stream.write(static_cast<const char *>(data), static_cast<std::streamsize>(size));
+  Check(stream.good(), "write pack source file");
+}
+
+void CheckPackArchive(const std::filesystem::path &root, const std::vector<uint8_t> &payload) {
+  const auto source = root / "pack-source";
+  const auto unicode = Common::PathFromUtf8(ArchiveTests::UnicodeFilename);
+  std::error_code error;
+  Check(std::filesystem::create_directories(source / "assets/subdir", error) &&
+            std::filesystem::create_directories(source / "empty", error),
+        "create pack source tree");
+  WriteHostFile(source / "eboot.bin", ArchiveTests::Eboot.data(), ArchiveTests::Eboot.size());
+  WriteHostFile(source / "assets/subdir/data.bin", payload.data(), payload.size());
+  WriteHostFile(source / unicode, nullptr, 0);
+  uint64_t expected = ArchiveTests::Eboot.size() + payload.size();
+#if !defined(_WIN32)
+  std::filesystem::create_directory_symlink("assets", source / "linked", error);
+  Check(!error, "link a pack source directory");
+  expected += payload.size();
+#endif
+
+  const auto archive = root / "packed.zar";
+  uint64_t done = 0;
+  uint64_t total = 0;
+  Check(Common::PackArchive(source, archive,
+                            [&](uint64_t packed, uint64_t size) {
+                              done = packed;
+                              total = size;
+                              return true;
+                            }),
+        "pack a host directory");
+  Check(done == expected && total == expected, "report every packed byte");
+
+  const auto packed = Common::MakeArchivePath(archive);
+  const auto check_payload = [&](const std::filesystem::path &member) {
+    Common::File file(packed / member, Common::File::Mode::Read);
+    const auto data = file.ReadWholeBuffer();
+    Check(data.size() == payload.size() &&
+              std::memcmp(data.data(), payload.data(), payload.size()) == 0,
+          "round-trip packed file contents");
+  };
+  check_payload("assets/subdir/data.bin");
+#if !defined(_WIN32)
+  check_payload("linked/subdir/data.bin");
+#endif
+  const auto empty_file = Common::File::GetInfo(packed / unicode);
+  Check(empty_file && empty_file->is_file && empty_file->size == 0 &&
+            Common::File::IsDirectoryExisting(packed / "empty"),
+        "keep empty files and directories");
+
+  const auto failed = root / "failed.zar";
+  const auto rejected = [&](const std::filesystem::path &directory, bool proceed) {
+    return !Common::PackArchive(directory, failed, [=](uint64_t, uint64_t) { return proceed; }) &&
+           !std::filesystem::exists(failed, error);
+  };
+  Check(rejected(source, false), "cancelling removes the partial archive");
+  Check(rejected(root / "missing", true), "reject a missing source directory");
+#if defined(__linux__)
+  // Archive names compare case-insensitively, so such a folder cannot be packed faithfully.
+  WriteHostFile(source / "EBOOT.BIN", nullptr, 0);
+  Check(rejected(source, true), "reject names that differ only in case");
+#endif
+}
+
 #if !defined(_WIN32) || defined(_GLIBCXX_GCC_GTHR_POSIX_H)
 void CheckCancellation(const std::filesystem::path &path,
                        const std::vector<uint8_t> &payload) {
@@ -395,6 +462,8 @@ int main() {
             std::memcmp(replaced_data.data(), replacement.data(), replacement.size()) == 0,
         "reopening after the final reader closes sees replacement contents");
 
+  CheckPackArchive(temporary.Path(), payload);
+
   std::printf("ArchiveFileTests: all cases passed\n");
   return 0;
 }

tests/AudioOut2PortTests.cpp

@@ -5,6 +5,7 @@
 #include <algorithm>
 #include <array>
 #include <atomic>
+#include <chrono>
 #include <cmath>
 #include <condition_variable>
 #include <cstdint>
@@ -32,6 +33,8 @@ size_t                            g_capture_bytes = 0;
 int                               g_next_device  = 1;
 int                               g_open_waiters = 0;
 bool                              g_block_opens  = false;
+bool                              g_block_outputs = false;
+std::vector<int>                  g_blocked_output;
 
 void Check(bool value, const char* text) {
 	if (!value) {
@@ -190,6 +193,22 @@ void ReleaseDeviceOpens() {
 	g_device_cv.notify_all();
 }
 
+void BlockOutputs() {
+	std::lock_guard lock(g_device_mutex);
+	g_block_outputs = true;
+}
+
+void WaitForBlockedOutput() {
+	std::unique_lock lock(g_device_mutex);
+	g_device_cv.wait(lock, []() { return !g_blocked_output.empty(); });
+}
+
+void ReleaseOutputs() {
+	std::lock_guard lock(g_device_mutex);
+	g_block_outputs = false;
+	g_device_cv.notify_all();
+}
+
 int LiveDeviceCount() {
 	std::lock_guard lock(g_device_mutex);
 	return static_cast<int>(g_live_devices.size());
@@ -605,6 +624,46 @@ void TestObjectBusRouting() {
 	Check(LiveDeviceCount() == 0, "object bus device leaked");
 }
 
+void TestBlockedPushOnlyHoldsItsContext() {
+	const auto blocked      = CreateContext(4, 1);
+	const auto other        = CreateContext();
+	const auto blocked_port = CreatePort(blocked, MakeParam());
+	const auto other_port   = CreatePort(other, MakeParam());
+	auto       object_param = MakeParam(0x100);
+	object_param.port_type  = 0x100;
+	std::vector<float> pcm(512 * 2, 0.25f);
+	SetPcm(blocked_port, pcm.data());
+	SetPcm(CreatePort(blocked, object_param), pcm.data());
+	CaptureOutputPcm(pcm.size() * sizeof(float));
+
+	BlockOutputs();
+	std::thread pusher([&]() { Check(AudioOut2::AudioOut2ContextPush(blocked, 1) == OK, "blocked push failed"); });
+	WaitForBlockedOutput();
+	std::fill(pcm.begin(), pcm.end(), 0.5f);
+	SetPcm(blocked_port, pcm.data());
+	SetAttribute(blocked_port, 1, std::array {0.5f, 0.5f});
+	SetPcm(other_port, pcm.data());
+	Check(AudioOut2::AudioOut2ContextPush(other, 1) == OK, "push stalled behind another context");
+	// Give each destroyer time to close a device of the blocked output; the backend fails if it does.
+	std::thread port_destroyer([&]() { AudioOut2::AudioOut2PortDestroy(blocked_port); });
+	std::this_thread::sleep_for(std::chrono::milliseconds(20));
+	std::thread context_destroyer([&]() { AudioOut2::AudioOut2ContextDestroy(blocked); });
+	std::this_thread::sleep_for(std::chrono::milliseconds(20));
+	ReleaseOutputs();
+	pusher.join();
+	port_destroyer.join();
+	context_destroyer.join();
+
+	const auto output = OutputPcm();
+	Check(output.size() == 3, "pushes did not reach the backend once each");
+	CheckSamples(output[0], {0.5f}, "other context missed its PCM");
+	CheckSamples(output[1], {0.25f}, "PCM changed while its push was in flight");
+	Check(g_output_gains[0] == std::array {1.0f, 1.0f}, "gains changed while their push was in flight");
+	Check(LiveDeviceCount() == 1, "destroyed port or context kept its device");
+	CaptureOutputPcm(0);
+	AudioOut2::AudioOut2ContextDestroy(other);
+}
+
 } // namespace
 
 namespace Libs::Audio::AudioInternal {
@@ -625,6 +684,8 @@ int AudioOutOpen(int type, uint32_t /*samples_num*/, uint32_t /*freq*/, Format /
 void AudioOutClose(int handle) {
 	std::lock_guard lock(g_device_mutex);
 	const auto      it = std::find(g_live_devices.begin(), g_live_devices.end(), handle);
+	Check(std::find(g_blocked_output.begin(), g_blocked_output.end(), handle) == g_blocked_output.end(),
+	      "device closed during its output");
 	if (it != g_live_devices.end()) {
 		g_live_devices.erase(it);
 	}
@@ -642,8 +703,16 @@ bool AudioOutHasDevice(int handle) {
 }
 
 uint32_t AudioOutOutputs(const OutputParam* params, uint32_t num, bool blocking) {
-	std::lock_guard lock(g_device_mutex);
+	std::unique_lock lock(g_device_mutex);
 	g_output_blocking.push_back(blocking);
+	if (g_block_outputs && g_blocked_output.empty()) {
+		for (uint32_t i = 0; i < num; i++) {
+			g_blocked_output.push_back(params[i].handle);
+		}
+		g_device_cv.notify_all();
+		g_device_cv.wait(lock, []() { return !g_block_outputs; });
+		g_blocked_output.clear();
+	}
 	g_output_gains.clear();
 	if (g_capture_bytes != 0) {
 		for (uint32_t i = 0; i < num; i++) {
@@ -683,6 +752,7 @@ int main() {
 	TestPortGainAndValidation();
 	TestObjectGainAndAmbisonics();
 	TestObjectBusRouting();
+	TestBlockedPushOnlyHoldsItsContext();
 	Check(AudioOut2::AudioOut2UserDestroy(g_user_handle) == OK, "test user destroy failed");
 	std::printf("AudioOut2PortTests: all cases passed\n");
 	return 0;

tests/Ngs2SamplerTests.cpp

@@ -124,7 +124,7 @@ void TestWaveformQueueCapacity() {
 			for (uint32_t i = 0; i < capacity; ++i) {
 				Check(f.voice.blocks[i].data == reinterpret_cast<const uint8_t*>(pcm) &&
 				          f.voice.blocks[i].info.user_data == i + 1 &&
-				          f.voice.blocks[i].cursor == (i == 0 ? 1u : 0u),
+				          f.voice.blocks[i].data_cursor == (i == 0 ? sizeof(int16_t) : 0u),
 				      "rejected waveform request replaced an outstanding block");
 			}
 			f.Render(3);
@@ -846,13 +846,57 @@ void TestStarvationAndPause() {
 	f.Queue(pcm, 2, 0);
 	f.voice.SetEvent(16);
 	f.Render(2);
-	Check(!f.voice.has_samples && f.voice.blocks.front().cursor == 0, "paused voice consumed data");
+	Check(!f.voice.has_samples && f.voice.blocks.front().data_cursor == 0,
+	      "paused voice consumed data");
 	f.voice.SetEvent(32);
 	f.Render(2);
 	Check(f.voice.samples[0] == 0.5f && f.voice.samples[1] == 0.25f,
 	      "stream failed to resume after starvation/pause");
 }
 
+void TestPcmByteUnitStreaming() {
+	Fixture                 f(48000);
+	const int16_t           pcm[] = {0, 4096, 8192, 12288, 16384, 20480, 24576, 28672};
+	const Ngs2WaveformBlock first {0, 4, 0, 1, 5, 0, 0}, rest {4, 12, 0, 0, 99, 0, 0};
+	Check(f.QueueBlocks(pcm, &first, 1, 1) == OK && f.QueueBlocks(pcm, &rest, 1, 0x12) == OK &&
+	          f.voice.WaveformData() == reinterpret_cast<const uint8_t*>(pcm + 1),
+	      "byte-unit stream was rejected or its read address ignored the skipped frame");
+	f.Render(8);
+	Check(f.voice.samples == std::vector<float> {0.125f, 0.25f, 0.375f, 0.5f, 0.625f, 0, 0, 0} &&
+	          f.voice.blocks.empty() && f.voice.decoded_samples == 5,
+	      "byte-unit stream did not play exactly its duration across data-only blocks");
+}
+
+void TestPcmSkipAcrossBlocks() {
+	Fixture                 f(48000);
+	const int16_t           pcm[] = {1, 2, 3, 8192, 16384};
+	const Ngs2WaveformBlock first {0, 4, 0, 3, 2, 0, 0}, rest {4, 6, 0, 0, 0, 0, 0};
+	Check(f.QueueBlocks(pcm, &first, 1, 1) == OK, "skipping stream block was rejected");
+	f.Render(1);
+	Check(!f.voice.has_samples && f.voice.blocks.empty(), "skip-only block produced samples");
+	Check(f.QueueBlocks(pcm, &rest, 1, 2) == OK &&
+	          f.voice.WaveformData() == reinterpret_cast<const uint8_t*>(pcm + 3),
+	      "read address ignored the skip remaining from the previous block");
+	f.Render(2);
+	Check(f.voice.samples == std::vector<float> {0.25f, 0.5f} && f.voice.decoded_samples == 2 &&
+	          f.voice.decoded_bytes == sizeof(pcm),
+	      "skip spanning blocks played skipped frames or miscounted decoded data");
+}
+
+void TestPcmBlocksEndWithTheirData() {
+	// The Last Oricru queues blocks holding fewer frames than they declare, even none.
+	Fixture                 f(48000, 2);
+	const int16_t           pcm[] = {4096, -4096, 8192, -8192, 12288, -12288};
+	const Ngs2WaveformBlock blocks[] = {
+	    {0, 0, 0, 0, 16385, 0, 0}, {0, 8, 0, 0, 4, 0, 0}, {8, 4, 0, 0, 3, 0, 0}};
+	Check(f.QueueBlocks(pcm, blocks, 3, 0) == OK, "partially filled blocks were rejected");
+	f.Render(5);
+	Check(f.voice.samples == std::vector<float> {0.125f, 0.25f, 0.375f, 0, 0, -0.125f, -0.25f,
+	                                             -0.375f, 0, 0} &&
+	          f.voice.blocks.empty() && f.voice.decoded_samples == 3,
+	      "partially filled blocks did not play their data back to back");
+}
+
 void TestMonoRateAndRouting() {
 	Fixture              f(22050);
 	std::vector<int16_t> pcm(441, 16384);
@@ -930,7 +974,7 @@ void TestSamplerReuseResetsRouting() {
 		const int16_t pcm[] = {8192, 16384, -8192, -16384};
 		f.Queue(pcm, 2, 0);
 		f.Render(1);
-		Check(!f.voice.has_samples && f.voice.blocks.front().cursor == 0,
+		Check(!f.voice.has_samples && f.voice.blocks.front().data_cursor == 0,
 		      "sampler setup started the new waveform without a play event");
 		f.voice.ports[0].dest = &mastering;
 		f.voice.SetEvent(1);
@@ -1216,7 +1260,7 @@ void TestLoopCallbacksAndExit(Ngs2RackType type) {
 	active.Queue(pcm + 2, 1, 0);
 	active.Render(6);
 	Check(log.events.size() == 1 && active.voice.decoded_samples == 2 &&
-	          active.voice.blocks.front().cursor == 0 && active.voice.samples[2] == 0,
+	          active.voice.blocks.front().data_cursor == 0 && active.voice.samples[2] == 0,
 	      "loop callback pause consumed another iteration");
 	Ngs2VoiceParamHeader exit_loop {sizeof(Ngs2VoiceParamHeader), 0,
 	                                type == Ngs2RackType::Sampler ? 0x10000004u : 0x40010004u};
@@ -1423,7 +1467,7 @@ void TestCallbackStopsConsumption() {
 			      "completion callback failed to interrupt playback");
 			Check(f.voice.samples[0] == 0.0f && f.voice.samples[1] == 0.0f &&
 			          f.voice.decoded_samples == 1 && f.voice.blocks.size() == 1 &&
-			          f.voice.blocks.front().cursor == 0,
+			          f.voice.blocks.front().data_cursor == 0,
 			      "callback pause/stop consumed the next queued block");
 			if (event == 16) {
 				f.voice.SetEvent(32);
@@ -1523,7 +1567,7 @@ void TestCallbackReplacesSource() {
 				f.Render(4);
 				Check(!f.voice.has_samples && f.voice.samples.size() == (setup ? 4u : 8u) &&
 				          std::ranges::all_of(f.voice.samples, [](float value) { return value == 0; }) &&
-				          f.voice.blocks.size() == 1 && f.voice.blocks.front().cursor == 0,
+				          f.voice.blocks.size() == 1 && f.voice.blocks.front().data_cursor == 0,
 				      "callback source replacement used an old frame or channel layout");
 				f.Render(4);
 				const std::vector<float> expected = setup
@@ -1684,6 +1728,9 @@ int main() {
 	TestPitchLoopAndSkip();
 	TestPitchClamps();
 	TestStarvationAndPause();
+	TestPcmByteUnitStreaming();
+	TestPcmSkipAcrossBlocks();
+	TestPcmBlocksEndWithTheirData();
 	TestMonoRateAndRouting();
 	TestMatrixOutputChannels();
 	TestCustomPcmStillPlays();

tests/ScalarProvenanceTests.cpp

@@ -237,12 +237,9 @@ void TestShaderBaseAndUserData() {
 
 void TestCarryAndBitFields() {
   Fixture fixture;
-  const auto carry =
-      fixture.Emit(ValueOpcode::IAddCarry32, {Value(0xffffffffu), Value(2u)});
-  const auto low =
-      fixture.Emit(ValueOpcode::CompositeExtractU32x2, {carry, Value(0u)});
-  const auto high =
-      fixture.Emit(ValueOpcode::CompositeExtractU32x2, {carry, Value(1u)});
+  const auto low = fixture.Emit(ValueOpcode::IAdd32, {Value(0xffffffffu), Value(2u)});
+  const auto carry = fixture.Emit(ValueOpcode::ULessThan32, {low, Value(0xffffffffu)});
+  const auto high = fixture.Emit(ValueOpcode::SelectU32, {carry, Value(1u), Value(0u)});
   const auto inserted =
       fixture.Emit(ValueOpcode::BitFieldInsert,
                    {Value(0u), Value(0x89abcdefu), Value(0u), Value(32u)});

tests/ShaderRecompilerComputeTests.cpp

@@ -1313,6 +1313,9 @@ struct BdaMapping {
   u32 backing_offset = 0;
 };
 
+using ShaderProfile = ShaderRecompiler::Spirv::Profile;
+constexpr ShaderProfile Maintenance8Profile{.maintenance8 = true};
+
 struct TestCase {
   const char *name = "";
   std::vector<u32> code;
@@ -1340,6 +1343,7 @@ struct TestCase {
   std::vector<u32> expected_storage_image_r32ui;
   std::vector<std::string> required_spirv;
   std::vector<std::string> forbidden_spirv;
+  const ShaderProfile *profile = &Maintenance8Profile;
   ShaderComputeInputInfo compute_info = [] {
     ShaderComputeInputInfo info{};
     info.lds_size_dwords = 1024;
@@ -1672,7 +1676,7 @@ CompiledShader CompileCase(const TestCase &test, u32 host_subgroup_size = 64) {
   }
   ShaderRecompiler::CompileOptions options;
   options.stage = ShaderType::Compute;
-  options.maintenance8_enabled = true;
+  options.profile = test.profile;
   options.dump_ir = true;
   auto compute_info = test.compute_info;
   compute_info.host_subgroup_size = host_subgroup_size;
@@ -1745,7 +1749,7 @@ CompiledShader CompileCase(const TestCase &test, u32 host_subgroup_size = 64) {
     result.program.shader_info_complete = false;
     ShaderRecompiler::IR::CollectShaderInfo(result.program, options.input_info);
     result.spirv = ShaderRecompiler::Spirv::EmitProgram(
-        result.program, options.input_info, 0u, options.maintenance8_enabled);
+        result.program, options.input_info, 0u, *options.profile);
     const auto shader_data = result.program.bindings.descriptor_counts[
         static_cast<size_t>(ShaderRecompiler::IR::DescriptorBindingKind::ShaderData)];
     Require(test.name, "binding layout",
@@ -1862,7 +1866,7 @@ CompiledShader CompileFragmentCase(const GraphicsCase &test) {
 
   ShaderRecompiler::CompileOptions options;
   options.stage = ShaderType::Pixel;
-  options.maintenance8_enabled = true;
+  options.profile = &Maintenance8Profile;
   options.dump_ir = false;
   options.input_info.pixel = &pixel_info;
   options.user_data = user_data;
@@ -2001,6 +2005,12 @@ class VulkanHarness {
 public:
   VulkanHarness() { Init(); }
   ~VulkanHarness() { Destroy(); }
+  // Host float min/max is used on AMD GPUs, as in Kyty.
+  [[nodiscard]] bool SupportsNativeFloatMinMax() const {
+    vk::PhysicalDeviceProperties properties;
+    m_physical_device.getProperties(&properties);
+    return properties.vendorID == 0x1002u;
+  }
 
   VulkanHarness(const VulkanHarness &) = delete;
   VulkanHarness &operator=(const VulkanHarness &) = delete;
@@ -10799,7 +10809,7 @@ public:
               ShaderRecompiler::CompileOptions options;
               options.wave_size = wave_size;
               options.dump_ir = false;
-              options.maintenance8_enabled = true;
+              options.profile = &Maintenance8Profile;
               input.host_subgroup_size = SubgroupSize();
               options.input_info.compute = &input;
               options.user_data = user_data;
@@ -22996,6 +23006,38 @@ TestCase ScalarZeroBitOps() {
   return test;
 }
 
+// V_READFIRSTLANE reads the first EXEC lane, or lane 0 when EXEC is zero; v1 = 100 + lane.
+TestCase VectorReadFirstLaneExecMasks(u32 wave_size) {
+  using O = ShaderOpcode;
+  TestCase test;
+  test.name = wave_size == 32 ? "VectorReadFirstLaneExecMasksWave32"
+                              : "VectorReadFirstLaneExecMasksWave64";
+  auto &code = test.code;
+  AppendVMovU32(&code, 1, 100);
+  code.push_back(EncodeVop2(0x25, 1, Vgpr(0), 1));
+  code.push_back(EncodeSop1(0x04, 10, 126)); // s_mov_b64 s[10:11], exec
+  const uint64_t masks[] = {0, 1ull << 5, 0x80000000ull, 1ull << 32, 1ull << 63, 0xffff000000000000ull};
+  for (const uint64_t mask : masks) {
+    if (wave_size == 32 && mask > UINT32_MAX) continue;
+    AppendSMovLiteral(&code, 126, u32(mask));
+    if (wave_size == 64) AppendSMovLiteral(&code, 127, u32(mask >> 32));
+    code.push_back(EncodeVop1(0x02, 20, Vgpr(1)));
+    code.push_back(EncodeSop1(0x04, 126, 10));
+    AppendStoreSgpr(&code, 20, u32(test.expected.size()));
+    test.expected.push_back(100 + (mask == 0 ? 0 : std::countr_zero(mask)));
+  }
+  AppendEnd(&code);
+  test.initial.resize(test.expected.size());
+  test.opcodes = {O::V_MOV_B32, O::V_ADD_NC_U32, O::S_MOV_B32, O::S_MOV_B64,
+                  O::V_READFIRSTLANE_B32, O::BUFFER_STORE_DWORD, O::S_ENDPGM};
+  test.compute_info.wave_size = wave_size;
+  test.compute_info.threads_num[0] = wave_size;
+  test.compute_info.threads_num[1] = test.compute_info.threads_num[2] = 1;
+  test.compute_info.thread_ids_num = 1;
+  test.has_compute_info = true;
+  return test;
+}
+
 TestCase ScalarSextMasks(u32 bits, u32 wave_size) {
   using O = ShaderOpcode;
   TestCase test;
@@ -24666,7 +24708,7 @@ TestCase Vop1SdwaMovByteDestinations() {
   test.opcodes = {O::BUFFER_LOAD_DWORD, O::V_MOV_B32, O::BUFFER_STORE_DWORD,
                   O::S_ENDPGM};
   test.decoded_counts = {{"V_MOV_B32 v1.sdwa(sel=2,sext=0), 0", 1}};
-  test.required_spirv = {"OpBitFieldInsert", "OpBitFieldSExtract", "OpBitwiseOr"};
+  test.required_spirv = {"OpBitFieldInsert", "OpBitFieldSExtract"};
   return test;
 }
 
@@ -24783,7 +24825,7 @@ TestCase Vop2SdwaAddNcCapturedHighWordDestination() {
   test.expected = {0x00010002u, 0x0006c3d4u};
   test.opcodes = {O::BUFFER_LOAD_DWORD, O::V_READFIRSTLANE_B32, O::V_MOV_B32,
                   O::V_ADD_NC_U32,      O::BUFFER_STORE_DWORD,  O::S_ENDPGM};
-  test.required_spirv = {"OpIAdd", "OpShiftLeftLogical", "OpBitwiseOr"};
+  test.required_spirv = {"OpIAdd", "OpBitFieldInsert"};
   return test;
 }
 
@@ -25099,10 +25141,8 @@ TestCase Vop3CvtPkI16I32Captured() {
   test.opcodes = {O::BUFFER_LOAD_DWORD, O::V_MOV_B32, O::V_CVT_PK_I16_I32,
                   O::BUFFER_STORE_DWORD, O::S_ENDPGM};
   test.decoded_counts = {{"0x00000014: V_CVT_PK_I16_I32 v5, v7, v16\n", 1}};
-  test.ir_counts = {{" = BitwiseAnd32 ", 2},
-                    {" = ShiftLeftLogical32 ", 1},
-                    {" = BitwiseOr32 ", 1}};
-  test.required_spirv = {"OpBitwiseAnd", "OpShiftLeftLogical", "OpBitwiseOr"};
+  test.ir_counts = {{" = BitFieldInsert ", 1}};
+  test.required_spirv = {"OpBitFieldInsert"};
   return test;
 }
 
@@ -25614,6 +25654,31 @@ TestCase VectorBfeI32SignExtendsField() {
   return test;
 }
 
+TestCase VectorBfmDropsMaskBitsPastBit31() {
+  using O = ShaderOpcode;
+
+  std::vector<u32> code;
+  for (u32 vgpr = 0; vgpr < 4; vgpr++) {
+    AppendVMovU32(&code, 30, vgpr * sizeof(u32));
+    AppendBufferLoadDword(&code, vgpr, 30);
+  }
+  AppendVop3(&code, 0x363, 10, Vgpr(0), Vgpr(1));
+  AppendVop3(&code, 0x363, 11, Vgpr(2), Vgpr(3));
+  AppendStoreVgpr(&code, 10, 0);
+  AppendStoreVgpr(&code, 11, 1);
+  AppendEnd(&code);
+
+  TestCase test;
+  test.name = "VectorBfmDropsMaskBitsPastBit31";
+  test.code = std::move(code);
+  // Size and offset use bits [4:0]; mask bits shifted past bit 31 are dropped.
+  test.initial = {24, 16, 0x3fu, 0x21u};
+  test.expected = {0xffff0000u, 0xfffffffeu};
+  test.opcodes = {O::V_MOV_B32, O::BUFFER_LOAD_DWORD, O::V_BFM_B32,
+                  O::BUFFER_STORE_DWORD, O::S_ENDPGM};
+  return test;
+}
+
 TestCase VectorAlignByteUsesTwoBitByteOffset() {
   using O = ShaderOpcode;
 
@@ -25796,7 +25861,7 @@ TestCase VectorAddcWritesPerLaneCarryOut() {
                   O::S_ENDPGM};
   test.decoded_counts = {
       {"V_ADDC_U32 v2, vcc_lo, 0, v26.sdwa(sel=4,sext=1), vcc_lo", 1}};
-  test.ir_counts = {{" = BitFieldSExtract ", 1}, {" = IAddCarry32 ", 2}};
+  test.ir_counts = {{" = BitFieldSExtract ", 1}, {" = ULessThan32 ", 3}};
   test.required_spirv = {"OpBitFieldSExtract"};
   test.compute_info.threads_num[0] = 4;
   test.compute_info.threads_num[1] = 1;
@@ -27265,6 +27330,35 @@ TestCase VectorDppRowShare(u32 wave_size) {
   return test;
 }
 
+// row_mirror, row_half_mirror and row_ror read lanes within the row; v1 = 100 + lane.
+TestCase VectorDppMirrorAndRotate(u32 wave_size) {
+  using O = ShaderOpcode;
+  TestCase test;
+  test.name = wave_size == 32 ? "VectorDppMirrorAndRotateWave32" : "VectorDppMirrorAndRotateWave64";
+  auto &code = test.code;
+  AppendVMovU32(&code, 1, 100);
+  code.push_back(EncodeVop2(0x25, 1, Vgpr(0), 1));
+  for (const u32 control : {0x140u, 0x141u, 0x121u, 0x125u, 0x12fu}) {
+    code.push_back(EncodeVop1(0x01, 2, 250));
+    code.push_back(EncodeVop1Dpp(1, control));
+    AppendStoreVgprAtLaneDwordOffset(&code, 2, 0, u32(test.expected.size()));
+    for (u32 lane = 0; lane < wave_size; ++lane) {
+      const u32 source = control == 0x140u   ? (lane | 15u) - (lane & 15u)
+                         : control == 0x141u ? (lane | 7u) - (lane & 7u)
+                                             : (lane & ~15u) | ((lane - (control & 15u)) & 15u);
+      test.expected.push_back(100 + source);
+    }
+  }
+  AppendEnd(&code);
+  test.opcodes = {O::V_MOV_B32, O::V_ADD_NC_U32, O::BUFFER_STORE_DWORD, O::S_ENDPGM};
+  test.compute_info.wave_size = wave_size;
+  test.compute_info.threads_num[0] = wave_size;
+  test.compute_info.threads_num[1] = test.compute_info.threads_num[2] = 1;
+  test.compute_info.thread_ids_num = 1;
+  test.has_compute_info = true;
+  return test;
+}
+
 TestCase VectorDppBankMaskPreservesDestination() {
   using O = ShaderOpcode;
 
@@ -27863,6 +27957,68 @@ TestCase VectorMinMaxF32NanAndSignedZeroEdges() {
            O::V_MAX3_F32, O::BUFFER_STORE_DWORD, O::S_ENDPGM}};
 }
 
+// MIN/MAX/MIN3/MAX3 on runtime inputs, so no compiler can fold them. The expected bits follow
+// the emulated rules: a NaN operand yields the other (S0 for two NaNs) and signed zeros are
+// ordered.
+// Denormal operands are left out; their expected results are not pinned down here.
+TestCase VectorMinMaxF32RuntimeEdges(u32 wave) {
+  using O = ShaderOpcode;
+  const auto nan = [](u32 v) { return (v & 0x7fffffffu) > 0x7f800000u; };
+  const auto zero = [](u32 v) { return (v & 0x7fffffffu) == 0u; };
+  const auto value = [](u32 v) { return std::bit_cast<float>(v); };
+  const auto min2 = [&](u32 a, u32 b) {
+    if (nan(b)) return a;
+    if (nan(a)) return b;
+    if (zero(a) && zero(b)) return a | b;
+    return value(a) < value(b) ? a : b;
+  };
+  const auto max2 = [&](u32 a, u32 b) {
+    if (nan(b)) return a;
+    if (nan(a)) return b;
+    if (zero(a) && zero(b)) return a & b;
+    return value(a) >= value(b) ? a : b;
+  };
+  constexpr u32 cases[][3] = {
+      {0x40000000u, 0x40400000u, 0x3f800000u}, {0x7fc00000u, 0x40000000u, 0xbf800000u},
+      {0x40000000u, 0x7fc00000u, 0x3f800000u}, {0x7fa00001u, 0x40000000u, 0x40400000u},
+      {0x40000000u, 0x7fa00001u, 0xc0000000u}, {0x80000000u, 0x00000000u, 0x80000000u},
+      {0x00000000u, 0x80000000u, 0x00000000u}, {0x80000000u, 0x80000000u, 0x00000000u},
+      {0xff800000u, 0x7f800000u, 0x00000000u}, {0x7fc00000u, 0x00000000u, 0x80000000u},
+      {0xffc05678u, 0x3f800000u, 0x7fa00001u}, {0xbf800000u, 0xc0000000u, 0x7fc01234u},
+      {0x7fc01234u, 0xffc05678u, 0x7fa00001u}, {0xffc05678u, 0x7fc01234u, 0x3f800000u}};
+  TestCase test;
+  test.name = "VectorMinMaxF32RuntimeEdges";
+  test.has_compute_info = true;
+  test.compute_info.wave_size = wave;
+  test.compute_info.float_mode = 0xf0; // Preserve FP32 input/output denormals.
+  test.compute_info.threads_num[0] = test.compute_info.threads_num[1] =
+      test.compute_info.threads_num[2] = 1u;
+  for (const auto &values : cases)
+    test.initial.insert(test.initial.end(), values, values + 3);
+  test.expected = test.initial;
+  auto &code = test.code;
+  for (u32 index = 0; index < std::size(cases); ++index) {
+    for (u32 operand = 0; operand < 3; ++operand) {
+      AppendVMovU32(&code, 30, (index * 3u + operand) * 4u);
+      AppendBufferLoadDword(&code, operand, 30);
+    }
+    const auto [a, b, c] = cases[index];
+    code.push_back(EncodeVop2(0x0f, 10, Vgpr(0), 1));
+    code.push_back(EncodeVop2(0x10, 11, Vgpr(0), 1));
+    AppendVop3(&code, 0x151, 12, Vgpr(0), Vgpr(1), Vgpr(2));
+    AppendVop3(&code, 0x154, 13, Vgpr(0), Vgpr(1), Vgpr(2));
+    for (u32 result = 0; result < 4; ++result)
+      AppendStoreVgpr(&code, 10 + result, static_cast<u32>(test.expected.size() + result));
+    test.expected.insert(test.expected.end(),
+                         {min2(a, b), max2(a, b), min2(min2(a, b), c), max2(max2(a, b), c)});
+  }
+  AppendEnd(&code);
+  test.initial.resize(test.expected.size());
+  test.opcodes = {O::V_MOV_B32,  O::BUFFER_LOAD_DWORD, O::V_MIN_F32,         O::V_MAX_F32,
+                  O::V_MIN3_F32, O::V_MAX3_F32,        O::BUFFER_STORE_DWORD, O::S_ENDPGM};
+  return test;
+}
+
 TestCase VectorMed3F32NanUsesMin3Path(u32 wave) {
   using O = ShaderOpcode;
   // Runtime inputs keep every MED3 in the emitted module; compare raw result bits.
@@ -27914,8 +28070,8 @@ TestCase VectorMed3F32NanUsesMin3Path(u32 wave) {
   test.opcodes = {O::V_MOV_B32, O::BUFFER_LOAD_DWORD, O::V_MED3_F32,
                   O::BUFFER_STORE_DWORD, O::S_ENDPGM};
   test.ir_counts = {{"FPMedTri32", std::size(cases)}};
-  // Measured 17,593/36,607 words; keep classification growth below a single MED3.
-  test.max_spirv_words = wave == 32u ? 17600u : 36620u;
+  // Measured 16,140/33,705 words; keep min/max expansion growth below a single MED3.
+  test.max_spirv_words = wave == 32u ? 16150u : 33720u;
   return test;
 }
 
@@ -28640,7 +28796,7 @@ TestCase VectorFractF64CapturedAndEdges() {
 
 TestCase VectorMinMaxF64CapturedAndEdges() {
   using O = ShaderOpcode;
-  // MODE.IEEE=0: select the numeric operand, or the second NaN unchanged.
+  // MODE.IEEE=0: select the numeric operand, or the first NaN unchanged.
   constexpr uint64_t positive_zero = 0, negative_zero = 0x8000000000000000ull;
   constexpr uint64_t one = 0x3ff0000000000000ull, minus_two = 0xc000000000000000ull;
   constexpr uint64_t qnan = 0x7ff8123456789abcull, snan = 0xfff0000000000042ull;
@@ -28654,7 +28810,7 @@ TestCase VectorMinMaxF64CapturedAndEdges() {
       {positive_zero, positive_zero, positive_zero, positive_zero},
       {qnan, one, one, one}, {one, qnan, one, one},
       {snan, minus_two, minus_two, minus_two}, {minus_two, snan, minus_two, minus_two},
-      {qnan, snan, snan, snan}, {snan, qnan, qnan, qnan},
+      {qnan, snan, qnan, qnan}, {snan, qnan, snan, snan},
       {infinity, one, one, infinity},
       {infinity | negative_zero, one, infinity | negative_zero, one},
       {0x3fe0000000000000ull, one, 0x3fe0000000000000ull, one},
@@ -29106,23 +29262,28 @@ TestCase VectorRcpIflagF32IntegerReciprocal() {
   return test;
 }
 
+// Runtime inputs keep the driver from folding; finite values from 2^23 up are whole periods.
 TestCase VectorSinCosMaxFiniteSpecialCases() {
   using O = ShaderOpcode;
-
-  std::vector<u32> code;
-  AppendVMovLiteral(&code, 0, 0xff7fffffu);
-  code.push_back(EncodeVop1(0x35, 1, Vgpr(0)));
-  code.push_back(EncodeVop1(0x36, 2, Vgpr(0)));
-  AppendStoreVgpr(&code, 1, 0);
-  AppendStoreVgpr(&code, 2, 1);
-  AppendEnd(&code);
-
-  return {"VectorSinCosMaxFiniteSpecialCases",
-          code,
-          {},
-          {0x00000000u, 0x3f800000u},
-          {O::V_MOV_B32, O::V_SIN_F32, O::V_COS_F32, O::BUFFER_STORE_DWORD,
-           O::S_ENDPGM}};
+  constexpr u32 inputs[] = {0xff7fffffu, 0x7f7fffffu, 0x4b000001u, 0xcb800000u};
+  TestCase test;
+  test.name = "VectorSinCosMaxFiniteSpecialCases";
+  test.initial.assign(std::begin(inputs), std::end(inputs));
+  test.expected = test.initial;
+  for (u32 i = 0; i < std::size(inputs); i++) {
+    AppendVMovU32(&test.code, 30, i * 4u);
+    AppendBufferLoadDword(&test.code, 0, 30);
+    test.code.push_back(EncodeVop1(0x35, 1, Vgpr(0)));
+    test.code.push_back(EncodeVop1(0x36, 2, Vgpr(0)));
+    AppendStoreVgpr(&test.code, 1, u32(test.expected.size()));
+    AppendStoreVgpr(&test.code, 2, u32(test.expected.size() + 1));
+    test.expected.insert(test.expected.end(), {0x00000000u, 0x3f800000u});
+  }
+  AppendEnd(&test.code);
+  test.initial.resize(test.expected.size());
+  test.opcodes = {O::V_MOV_B32, O::BUFFER_LOAD_DWORD, O::V_SIN_F32, O::V_COS_F32,
+                  O::BUFFER_STORE_DWORD, O::S_ENDPGM};
+  return test;
 }
 
 TestCase VectorCompareF32ExposureGuard() {
@@ -31084,7 +31245,7 @@ TestCase ScalarBranchRestoresInactiveLanes(u32 wave_size, u32 threads, bool uppe
                   upper_nonzero ? O::S_CBRANCH_EXECNZ : O::S_CBRANCH_EXECZ,
                   O::V_READLANE_B32, O::V_LSHLREV_B32,
                   O::BUFFER_STORE_DWORD, O::S_ENDPGM};
-  test.required_spirv = {upper_nonzero ? "OpGroupNonUniformAny" : "OpGroupNonUniformAll"};
+  test.required_spirv = {"OpGroupNonUniformAny"};
   test.compute_info.wave_size = wave_size;
   test.compute_info.threads_num[0] = threads;
   test.compute_info.thread_ids_num = 1;
@@ -38144,7 +38305,8 @@ void CheckDynamicSampleOffsets(VulkanHarness &vulkan, u32 wave_size) {
           "dynamic offsets expanded samples or lost signed six-bit components");
 #if KYTY_PLATFORM != KYTY_PLATFORM_WINDOWS
   if (wave_size == 32u) {
-    Require(test.name, "default feature gate", !ShaderRecompiler::CompileOptions{}.maintenance8_enabled,
+    Require(test.name, "default feature gate",
+            !ShaderRecompiler::CompileOptions{}.profile->maintenance8,
             "maintenance8 must reflect the enabled device feature");
     ExpectFatal("DynamicSampleOffsetsRequireMaintenance8", [&] {
       ShaderRecompiler::Spirv::EmitProgram(compiled.program, {.compute = &test.compute_info});
@@ -42912,6 +43074,7 @@ std::vector<TestCase> MakeCases() {
   AddCase(VectorVop3LshlrevB64Captured);
   AddCase(VectorVop3IntegerOps);
   AddCase(VectorBfeI32SignExtendsField);
+  AddCase(VectorBfmDropsMaskBitsPastBit31);
   AddCase(VectorAlignByteUsesTwoBitByteOffset);
   AddCase(VectorCarryAndBitCountOps);
   AddCase(ScalarVector64BitFloatInlineConstants);
@@ -42961,6 +43124,10 @@ std::vector<TestCase> MakeCases() {
   AddCase(VectorDppRowXmask);
   cases.push_back(VectorDppRowShare(32));
   cases.push_back(VectorDppRowShare(64));
+  cases.push_back(VectorDppMirrorAndRotate(32));
+  cases.push_back(VectorDppMirrorAndRotate(64));
+  cases.push_back(VectorReadFirstLaneExecMasks(32));
+  cases.push_back(VectorReadFirstLaneExecMasks(64));
   cases.push_back(ComputeWave64Minimum(false));
   cases.push_back(ComputeWave64Minimum(true));
   AddCase(VectorDppBankMaskPreservesDestination);
@@ -42983,6 +43150,8 @@ std::vector<TestCase> MakeCases() {
   AddCase(VectorMinMaxF32NanAndSignedZeroEdges);
   cases.push_back(VectorMed3F32NanUsesMin3Path(32));
   cases.push_back(VectorMed3F32NanUsesMin3Path(64));
+  cases.push_back(VectorMinMaxF32RuntimeEdges(32));
+  cases.push_back(VectorMinMaxF32RuntimeEdges(64));
   AddCase(VectorFloatConversionOps);
   AddCase(VectorFrexpF32Edges);
   AddCase(CvtF32ToIntSaturatesNaNAndOutOfRange);
@@ -43885,7 +44054,8 @@ void CheckTessellationProgram(const char *name, u32 ls_stride, u32 hs_stride) {
 
 // PPSA25468: CP0 writes patch constants below the 32 KiB offchip ring end
 // through mul24 addresses that reuse data VGPRs under nested EXEC masks, and
-// the TES reads them back.
+// the TES reads them back. As in God of War Ragnarok, the HS reads its next LS
+// record, loops on a v_cmpx header and CP0 reads back another CP's outputs.
 void CheckTessellationPatchConstants() {
   using namespace ShaderRecompiler;
   using Attribute = IR::TessellationAttribute;
@@ -43899,7 +44069,7 @@ void CheckTessellationPatchConstants() {
   local.push_back(EncodeDs1(0, 2, 20));
   local.push_back(EncodeSop1(0x20, 0, 6));
 
-  const auto make_control = [&](bool divergent) {
+  const auto make_control = [&](bool divergent, bool carried) {
     std::vector<u32> control;
     AppendVop3(&control, 0x365, 0, 193u, InlineU32(0));
     AppendVop3(&control, 0x366, 0, 193u, Vgpr(0));
@@ -43912,7 +44082,10 @@ void CheckTessellationPatchConstants() {
     control.push_back(0xffu);
     control.push_back(EncodeVop2(0x0b, 0, 255u, 16)); // v_mul_u32_u24 v0, hs_stride, v16
     control.push_back(hs_stride);
-    AppendVop3(&control, 0x169, 1, Vgpr(16), 255u); // v_mul_lo_u32 v1, v16, ls_stride
+    control.push_back(EncodeVopc(0xc4, InlineU32(2), 16));    // v_cmp_gt_u32 vcc, 2, v16
+    control.push_back(EncodeVop2(0x25, 1, InlineU32(1), 16)); // v_add_nc_u32 v1, 1, v16
+    control.push_back(EncodeVop2(0x01, 1, InlineU32(0), 1));  // v_cndmask_b32 v1, 0, v1, vcc
+    AppendVop3(&control, 0x169, 1, Vgpr(1), 255u); // v_mul_lo_u32 v1, v1, ls_stride
     control.push_back(ls_stride);
     AppendVop3(&control, 0x143, 17, 255u, Vgpr(50), Vgpr(0));
     control.push_back(3 * hs_stride);
@@ -43932,12 +44105,33 @@ void CheckTessellationPatchConstants() {
       control.push_back(EncodeVop2(0x25, 17, InlineU32(4), 17)); // v_add_nc_u32 v17, 4, v17
     }
     control[skip_scale] = EncodeSopp(0x06, control.size() - skip_scale - 1u);
+    control.push_back(EncodeSop1(0x04, 6, 126));       // s_mov_b64 s[6:7], exec
+    control.push_back(EncodeSMovB32(5, InlineU32(0))); // s_mov_b32 s5, 0
+    const auto loop = control.size();
+    control.push_back(EncodeVopc(0xd1, 5, 16)); // v_cmpx_lt_u32 s5, v16
+    const auto exit_loop = control.size();
+    control.push_back(0u); // s_cbranch_execz
     control.push_back(EncodeMubuf0(0x1c, 24));
     control.push_back(EncodeMubuf1(12, 2, 17, 2));
+    // Only the back edge shows a loop-carried v17 at the store above.
+    if (carried) {
+      control.push_back(EncodeVop2(0x25, 17, InlineU32(4), 17)); // v_add_nc_u32 v17, 4, v17
+    }
+    control.push_back(EncodeVop2(0x25, 13, InlineU32(1), 13)); // v_add_nc_u32 v13, 1, v13
+    control.push_back(EncodeSop2(0x00, 5, 5, InlineU32(1)));   // s_add_u32 s5, s5, 1
+    control.push_back(EncodeSopp(0x02, loop - control.size() - 1u)); // s_branch
+    control[exit_loop] = EncodeSopp(0x08, control.size() - exit_loop - 1u);
+    control.push_back(EncodeSop1(0x04, 126, 6)); // s_mov_b64 exec, s[6:7]
+    control.push_back(EncodeSopp(0x0a));         // s_barrier
     // CP0 reuses data VGPR v10 for the ring address 0x8000 - (patch + 1) * 0xf0.
     control.push_back(EncodeVopc(0xd4, InlineU32(1), 16)); // v_cmpx_gt_u32 1, v16
     const auto skip_patch = control.size();
     control.push_back(0u); // s_cbranch_execz
+    control.push_back(EncodeVop2(0x0b, 11, 255u, 50)); // v_mul_u32_u24 v11, 3 * hs_stride, v50
+    control.push_back(3 * hs_stride);
+    // buffer_load_dwordx2 v[11:12], v11, s[8:11], s2 offen offset:hs_stride + 16 (CP1)
+    control.push_back(EncodeMubuf0(0x0d, hs_stride + 16));
+    control.push_back(EncodeMubuf1(11, 2, 11, 2));
     control.push_back(EncodeVop2(0x25, 10, InlineU32(1), 50)); // v_add_nc_u32 v10, 1, v50
     control.push_back(EncodeVop2(0x0b, 10, 255u, 10));         // v_mul_u32_u24 v10, 0xf0, v10
     control.push_back(patch_stride);
@@ -43959,7 +44153,7 @@ void CheckTessellationPatchConstants() {
     AppendEnd(&control);
     return control;
   };
-  const auto control = make_control(false);
+  const auto control = make_control(false, false);
 
   std::vector<u32> evaluation;
   AppendVop3(&evaluation, 0x169, 24, 255u, Vgpr(7));
@@ -43982,13 +44176,14 @@ void CheckTessellationPatchConstants() {
   ShaderTessellationInputInfo tess{.input_control_points = 3,
                                   .output_control_points = 3,
                                   .domain = 1,
-                                  .partitioning = 2,
+                                  .partitioning = 0,
                                   .output_topology = 2};
   AnalyzeTessellationPrograms(local, control, tess);
   Require(name, "reflected interface",
           tess.ls_stride == ls_stride && tess.hs_stride == hs_stride &&
               tess.patch_base == 0x7f10u && tess.patch_size == 28u,
-          "b128 LS records, joined HS paths and CP0 patch stores must reflect the layout");
+          "b128 LS records, selected LS reads, joined and looping HS paths and CP0 patch "
+          "stores must reflect the layout");
 
   auto ls = CompileTessellationStage(name, local, ShaderType::Local, tess);
   auto hs = CompileTessellationStage(name, control, ShaderType::TessellationControl, tess);
@@ -43999,6 +44194,10 @@ void CheckTessellationPatchConstants() {
   Require(name, "predicated patch outputs",
           hs.addresses[Attribute::PatchOutput] == std::set<u32>{0x7f10, 0x7f20, 0x7f24, 0x7f28},
           "CP0 mul24 addresses under nested EXEC masks must reach patch constants");
+  Require(name, "control point read-back",
+          hs.addresses[Attribute::ControlOutput] == std::set<u32>{hs_stride + 16, hs_stride + 20} &&
+              hs.assembly.find("OpUDiv %uint %uint_96 %uint_80") != std::string::npos,
+          "CP0 must read CP1's outputs by address instead of its own invocation's");
   Require(name, "patch inputs",
           ds.addresses[Attribute::PatchInput] == std::set<u32>{0x7f20, 0x7f24} &&
               ds.addresses[Attribute::EvaluationInput] ==
@@ -44008,12 +44207,16 @@ void CheckTessellationPatchConstants() {
           HasPatchInterface(hs.assembly, "Output", 2, 5) &&
               HasPatchInterface(ds.assembly, "Input", 2, 5),
           "HS and TES patch variables must share their Location and type");
+  Require(name, "integer partitioning", ds.assembly.find("SpacingEqual") != std::string::npos,
+          "integer partitioning must map to equal spacing");
   std::printf("[host]    %-32s ok\n", name);
 #if KYTY_PLATFORM != KYTY_PLATFORM_WINDOWS
-  ExpectFatal("TessellationDivergentJoin", [&] {
-    ShaderTessellationInputInfo info{.input_control_points = 3, .output_control_points = 3};
-    AnalyzeTessellationPrograms(local, make_control(true), info);
-  });
+  for (const bool carried : {false, true}) {
+    ExpectFatal(carried ? "TessellationLoopCarriedAddress" : "TessellationDivergentJoin", [&] {
+      ShaderTessellationInputInfo info{.input_control_points = 3, .output_control_points = 3};
+      AnalyzeTessellationPrograms(local, make_control(!carried, carried), info);
+    });
+  }
 #endif
 }
 
@@ -44725,6 +44928,26 @@ void CheckImageTransitionState(RenderContext &renderer) {
           ImageTestAccess::CopyRows(copy_capacity + 4, 1, copy_capacity) == 0,
       "buffered image copy does not split at the fixed scratch capacity");
 
+  // RADV rejects 3D 128-bit BC images with block-texel views plus extended usage.
+  auto volume_info = MakeInfo(vk::Format::eBc7UnormBlock, 1, 1);
+  volume_info.guest_format = Prospero::BufferFormat::kBc7UNorm;
+  volume_info.type = Prospero::ImageType::kColor3D;
+  volume_info.extent.depth = 4;
+  volume_info.bytes_per_block = 16;
+  for (const bool supported : {false, context.supports_block_texel_view}) {
+    context.supports_block_texel_view = supported;
+    Image volume(context, scheduler, volume_info);
+    const auto flags = volume.backing.flags;
+    Require(name, "block volume create flags",
+            (flags & vk::ImageCreateFlagBits::e2DArrayCompatible) &&
+                static_cast<bool>(flags & vk::ImageCreateFlagBits::eBlockTexelViewCompatible) ==
+                    supported &&
+                static_cast<bool>(flags & vk::ImageCreateFlagBits::eExtendedUsage) ==
+                    static_cast<bool>(volume.backing.usage &
+                                      vk::ImageUsageFlagBits::eStorage),
+            "block volume requested extended usage without storage");
+  }
+
   Image image(context, scheduler, MakeInfo(vk::Format::eR8Unorm, 2, 3));
   auto barriers =
       image.GetBarriers(vk::ImageLayout::eShaderReadOnlyOptimal,
@@ -48906,6 +49129,26 @@ int main(int argc, char **argv) {
     RunCase(&vulkan, CvtPkrtzF16F32SdwaAndOutputModifiers());
     return 0;
   }
+  if (argc == 2 && std::strcmp(argv[1], "--float-min-max-only") == 0) {
+    // Host min/max must keep the emulated results on AMD GPUs; other GPUs only
+    // compile it.
+    static constexpr ShaderProfile native_profile{.maintenance8 = true,
+                                                  .native_float_min_max = true};
+    VulkanHarness vulkan;
+    const bool supported = vulkan.SupportsNativeFloatMinMax();
+    for (bool native : {false, true}) {
+      for (auto test :
+           {VectorMinMaxF32NanAndSignedZeroEdges(),
+            VectorMinMaxF32RuntimeEdges(32), VectorMinMaxF32RuntimeEdges(64)}) {
+        auto &spirv = native ? test.required_spirv : test.forbidden_spirv;
+        spirv = {" NMin ", " NMax "};
+        test.profile = native ? &native_profile : &Maintenance8Profile;
+        test.compile_only = native && !supported;
+        RunCase(&vulkan, test);
+      }
+    }
+    return 0;
+  }
   if (argc == 2 && std::strcmp(argv[1], "--thread-dimensions-only") == 0) {
     VulkanHarness vulkan;
     CheckComputeThreadDimensions(vulkan);
@@ -48995,6 +49238,10 @@ int main(int argc, char **argv) {
     RunCase(&vulkan, VectorDppRowXmask());
     RunCase(&vulkan, VectorDppRowShare(32));
     RunCase(&vulkan, VectorDppRowShare(64));
+    RunCase(&vulkan, VectorDppMirrorAndRotate(32));
+    RunCase(&vulkan, VectorDppMirrorAndRotate(64));
+    RunCase(&vulkan, VectorReadFirstLaneExecMasks(32));
+    RunCase(&vulkan, VectorReadFirstLaneExecMasks(64));
     RunCase(&vulkan, VectorDppBankMaskPreservesDestination());
     RunCase(&vulkan, VectorDppBoundsControlZeroPreservesDestination());
     return 0;

tests/VirtualMemoryAllocationTests.cpp

@@ -1596,20 +1596,23 @@ void TestDirectMapQueryOffsetAndPartialMunmap() {
 	ExpectUnmapped(test, base + SceKernelPageSize);
 
 	constexpr uint64_t transaction_sentinel = 0x5452414e53414354ull; // "TRANSACT"
-	constexpr uint64_t rejected_write       = 0x4e4f504152544941ull; // "NOPARTIA"
+	constexpr uint64_t sparse_write         = 0x4e4f504152544941ull; // "NOPARTIA"
+	constexpr uint64_t backed_prefix_mask   = 0xffffffffull;
 	const auto         crossing_address     = base + SceKernelPageSize - sizeof(uint32_t);
 	std::memcpy(reinterpret_cast<void*>(alias_base + SceKernelPageSize - sizeof(uint32_t)),
 	            &transaction_sentinel, sizeof(transaction_sentinel));
 	Check(test,
-	      !Libs::LibKernel::Memory::TryWriteBacking(crossing_address, &rejected_write,
-	                                                sizeof(rejected_write)),
-	      "TryWriteBacking should reject a range crossing an unmapped span");
-	uint64_t backing_after_rejected_write = 0;
-	std::memcpy(&backing_after_rejected_write,
+	      Libs::LibKernel::Memory::TryWriteBacking(crossing_address, &sparse_write,
+	                                               sizeof(sparse_write)),
+	      "TryWriteBacking should write a range crossing an unmapped span");
+	uint64_t backing_after_sparse_write = 0;
+	std::memcpy(&backing_after_sparse_write,
 	            reinterpret_cast<const void*>(alias_base + SceKernelPageSize - sizeof(uint32_t)),
-	            sizeof(backing_after_rejected_write));
-	Check(test, backing_after_rejected_write == transaction_sentinel,
-	      "failed backing writes must not modify a validated prefix");
+	            sizeof(backing_after_sparse_write));
+	Check(test,
+	      backing_after_sparse_write == ((sparse_write & backed_prefix_mask) |
+	                                     (transaction_sentinel & ~backed_prefix_mask)),
+	      "sparse backing writes must fill the backed prefix and skip the unmapped span");
 	uint64_t rejected_read = transaction_sentinel;
 	Check(test,
 	      !Libs::LibKernel::Memory::TryReadBacking(crossing_address, &rejected_read,

tests/shaderCfgTests.cpp

@@ -10429,14 +10429,12 @@ void TestNewShaderRecompilerNativeU64Translation() {
       ir.Emit(IR::ValueOpcode::BitwiseAnd64,
               {base, IR::Value(uint64_t{0xff00ff00ff00ff00ull})});
   const auto sum = ir.Emit(IR::ValueOpcode::IAdd64, {masked, base});
-  const auto product = ir.Emit(IR::ValueOpcode::IMul64, {sum, base});
   const auto shifted =
-      ir.Emit(IR::ValueOpcode::ShiftRightArithmetic64, {product, lane});
+      ir.Emit(IR::ValueOpcode::ShiftRightArithmetic64, {sum, lane});
   ir.Emit(IR::ValueOpcode::IEqual64, {shifted, base});
   ir.Emit(IR::ValueOpcode::INotEqual64, {shifted, base});
   ir.Emit(IR::ValueOpcode::ULessThan64, {shifted, base});
   ir.Emit(IR::ValueOpcode::SLessThan64, {shifted, base});
-  ir.Emit(IR::ValueOpcode::BitCount64, {shifted});
   for (const uint32_t count : {1u, 31u, 32u, 33u, 63u}) {
     ir.Emit(IR::ValueOpcode::ShiftLeftLogical64, {base, IR::Value(count)});
     ir.Emit(IR::ValueOpcode::ShiftRightLogical64, {base, IR::Value(count)});
@@ -10451,7 +10449,6 @@ void TestNewShaderRecompilerNativeU64Translation() {
   Check(source.find("OpCapability Int64") != std::string::npos &&
             source.find("OpTypeInt 64 0") != std::string::npos &&
             SpirvSourceHasInstructionUsing(source, "OpIAdd", "%ulong") &&
-            SpirvSourceHasInstructionUsing(source, "OpIMul", "%ulong") &&
             SpirvSourceHasInstructionUsing(source, "OpBitwiseAnd", "%ulong"),
         "U64 arithmetic did not emit native 64-bit scalar instructions");
   Check(SpirvInstructionOpcodeCount(spirv, spv::OpIAddCarry) == 0u &&
@@ -10535,6 +10532,7 @@ void TestComputeDispatchWaveSize() {
 
 void TestNewShaderRecompilerBufferLoadsGuardedByExec() {
   const uint32_t shader[] = {
+      EncodeVopc(0xd4, 0 + 256, 1), // v_cmpx_gt_u32 v0, v1
       EncodeMubuf0(0x0c),
       EncodeMubuf1(0, 0, 1), // buffer_load_dword
                              // v0
@@ -15147,7 +15145,7 @@ void TestTypedDescriptorRealCarryAndScalarLoads() {
             ValueDependsOn(carry_source->dwords[1],
                            ShaderRecompiler::IR::ValueOpcode::GetShaderBase) &&
             ValueDependsOn(carry_source->dwords[1],
-                           ShaderRecompiler::IR::ValueOpcode::IAddCarry32),
+                           ShaderRecompiler::IR::ValueOpcode::ULessThan32),
         "real s_add_u32/s_addc_u32 translation lost SCC carry provenance");
   std::array<uint32_t, 64> carry_user_data{};
   const uint64_t shader_base = 0x00000012fffffffbull;
@@ -16350,8 +16348,8 @@ void CheckSpirvBlockSourceMapping(TestCompileResult &result) {
   Check(SpirvInstructionOpcodeCount(result.spirv, spv::OpString) == 0 &&
             SpirvInstructionOpcodeCount(result.spirv, spv::OpLine) == 0,
         "default emission unexpectedly contains source mapping");
-  const auto mapped = Spirv::EmitProgram(result.program, options.input_info, 0,
-                                         false, true);
+  const auto mapped =
+      Spirv::EmitProgram(result.program, options.input_info, 0, {}, true);
   CheckSpirvBinaryValidates(mapped);
   uint32_t source = 0;
   uint32_t label = 0;