Protect Windows guest red zones when emulating VRSQRTPS

Keep the fan center across subgroups and reuse the existing index fetch.

.github/workflows/build.yml

@@ -6,7 +6,12 @@ on:
     branches: [ main, master ]
   pull_request:
 
+concurrency:
+  group: ${{ github.workflow }}-${{ github.event.pull_request.number || github.ref }}
+  cancel-in-progress: ${{ github.event_name == 'pull_request' }}
+
 env:
+  SCCACHE_GHA_ENABLED: "true"
   KYTY_BUILD_ORIGIN: ${{ github.repository != 'KytyPS5/KytyPS5' && 'Fork' || github.event_name == 'push' && 'Official' || 'Source' }}
   KYTY_BUILD_REPOSITORY: ${{ github.repository }}
 
@@ -78,6 +83,9 @@ jobs:
           arch: win64_msvc2022_64
           cache: true
 
+      - name: Setup sccache
+        uses: mozilla-actions/sccache-action@v0.0.11
+
       - name: Verify toolchain
         shell: cmd
         run: |
@@ -100,6 +108,8 @@ jobs:
             -DKYTY_RELEASE_TAG=${{ needs.metadata.outputs.release_tag }} ^
             -DCMAKE_C_COMPILER=clang-cl ^
             -DCMAKE_CXX_COMPILER=clang-cl ^
+            -DCMAKE_C_COMPILER_LAUNCHER=sccache ^
+            -DCMAKE_CXX_COMPILER_LAUNCHER=sccache ^
             -DCMAKE_PREFIX_PATH="%Qt6_DIR%"
 
       - name: Build
@@ -156,6 +166,9 @@ jobs:
           arch: clang_64
           cache: true
 
+      - name: Setup sccache
+        uses: mozilla-actions/sccache-action@v0.0.11
+
       - name: Verify toolchain
         shell: bash
         run: |
@@ -180,6 +193,8 @@ jobs:
             -DCMAKE_OSX_ARCHITECTURES=x86_64 \
             -DCMAKE_C_COMPILER=clang \
             -DCMAKE_CXX_COMPILER=clang++ \
+            -DCMAKE_C_COMPILER_LAUNCHER=sccache \
+            -DCMAKE_CXX_COMPILER_LAUNCHER=sccache \
             -DCMAKE_PREFIX_PATH="$QT_ROOT_DIR"
 
       - name: Build
@@ -289,9 +304,15 @@ jobs:
             clang \
             glslang-tools \
             libasound2-dev \
+            libavcodec-dev \
+            libavfilter-dev \
+            libavformat-dev \
+            libavutil-dev \
             libdbus-1-dev \
             libgl1-mesa-dev \
             libpulse-dev \
+            libswresample-dev \
+            libswscale-dev \
             libudev-dev \
             libwayland-dev \
             libx11-dev \
@@ -304,6 +325,7 @@ jobs:
             libxss-dev \
             lld \
             ninja-build \
+            pkg-config \
             wayland-protocols
 
       - name: Install Qt 6.10.3
@@ -315,6 +337,9 @@ jobs:
           arch: linux_gcc_64
           cache: true
 
+      - name: Setup sccache
+        uses: mozilla-actions/sccache-action@v0.0.11
+
       - name: Verify toolchain
         shell: bash
         run: |
@@ -339,6 +364,8 @@ jobs:
             -DKYTY_RELEASE_TAG="${{ needs.metadata.outputs.release_tag }}" \
             -DCMAKE_C_COMPILER=clang \
             -DCMAKE_CXX_COMPILER=clang++ \
+            -DCMAKE_C_COMPILER_LAUNCHER=sccache \
+            -DCMAKE_CXX_COMPILER_LAUNCHER=sccache \
             -DCMAKE_PREFIX_PATH="$Qt6_DIR" 2>&1 | tee _Build/configure.log
           exit "${PIPESTATUS[0]}"
 

3rdparty/CMakeLists.txt

@@ -88,6 +88,13 @@ set(SPDLOG_FMT_EXTERNAL ON CACHE BOOL "" FORCE)
 set(SPDLOG_NO_EXCEPTIONS ON CACHE BOOL "" FORCE)
 add_subdirectory(spdlog EXCLUDE_FROM_ALL)
 
+if (NOT TARGET FFmpeg::ffmpeg AND CMAKE_SYSTEM_NAME STREQUAL "Linux")
+	find_package(PkgConfig REQUIRED)
+	pkg_check_modules(FFMPEG REQUIRED IMPORTED_TARGET GLOBAL
+		libavformat libavcodec libswscale libavutil libavfilter libswresample)
+	add_library(FFmpeg::ffmpeg ALIAS PkgConfig::FFMPEG)
+endif()
+
 if (NOT TARGET FFmpeg::ffmpeg)
 	# On macOS the target arch is driven by CMAKE_OSX_ARCHITECTURES, not the host
 	# CMAKE_SYSTEM_PROCESSOR (which reports arm64 even for an x86_64/Rosetta build).

CMakeLists.txt

@@ -317,24 +317,10 @@ endif()
 
 set(kyty_emulator_link_libraries common Vulkan::Headers spirv-tools-opt spirv-tools SDL2-static imgui xxhash FFmpeg::ffmpeg fmt::fmt nlohmann_json::nlohmann_json LibAtrac9 kyty_git_version Zydis::Zydis xbyak::xbyak)
 
-# Linux system libraries required by the static FFmpeg archive.
+# Linux system libraries.
 if(LINUX)
 	find_package(Threads REQUIRED)
 	list(APPEND kyty_emulator_link_libraries m ${CMAKE_DL_LIBS} Threads::Threads)
-
-	# Optional FFmpeg dependencies.
-	find_package(ZLIB)
-	if(ZLIB_FOUND)
-		list(APPEND kyty_emulator_link_libraries ZLIB::ZLIB)
-	endif()
-	find_library(KYTY_BZ2_LIBRARY bz2)
-	if(KYTY_BZ2_LIBRARY)
-		list(APPEND kyty_emulator_link_libraries ${KYTY_BZ2_LIBRARY})
-	endif()
-	find_library(KYTY_LZMA_LIBRARY lzma)
-	if(KYTY_LZMA_LIBRARY)
-		list(APPEND kyty_emulator_link_libraries ${KYTY_LZMA_LIBRARY})
-	endif()
 endif()
 
 set(inc_headers

README.md

@@ -165,13 +165,14 @@ The finished application and its runtime dependencies will be placed in
 
 ### Building on Linux
 
-Install the toolchain and the libraries the bundled SDL2 needs. Without the audio, Wayland and
+Install the toolchain, system FFmpeg, and the libraries the bundled SDL2 needs. Without the audio, Wayland and
 udev development packages SDL2 quietly configures itself without those backends, and the resulting
 build has no working sound and no gamepad hotplug:
 
 ```bash
 sudo apt-get install --no-install-recommends \
-  clang lld ninja-build cmake git glslang-tools \
+  clang lld ninja-build cmake git glslang-tools pkg-config \
+  libavformat-dev libavcodec-dev libswscale-dev libavutil-dev libavfilter-dev libswresample-dev \
   libgl1-mesa-dev libx11-dev libxcursor-dev libxext-dev libxfixes-dev \
   libxi-dev libxrandr-dev libxss-dev libxkbcommon-dev \
   libasound2-dev libpulse-dev libudev-dev libdbus-1-dev libwayland-dev wayland-protocols
@@ -192,7 +193,8 @@ cmake --install _Build/linux --prefix _Build/linux/install
 ```
 
 The install step copies the Qt libraries and plugins next to the binaries, so
-`_Build/linux/install` runs without a matching system Qt.
+`_Build/linux/install` runs without a matching system Qt. The emulator still requires
+the system FFmpeg runtime libraries matching those used to build it.
 
 As on Windows, the MSVC compiler is not used; Clang is required. `cl.exe` is rejected at configure
 time.

src/common/emulatorConfig.cpp

@@ -63,6 +63,10 @@ bool VrEnabled() {
 	return g_config->vr_enabled;
 }
 
+bool AmdCpuEnabled() {
+	return g_config->amd_cpu_enabled;
+}
+
 uint32_t GetVblankFrequency() {
 	return std::clamp(g_config->vblank_frequency, 30u, 360u);
 }

src/common/emulatorConfig.h

@@ -48,6 +48,7 @@ struct ConfigOptions {
 	int32_t                gpu_index                   = -1;
 	bool                   fullscreen_enabled          = false;
 	bool                   vr_enabled                  = false;
+	bool                   amd_cpu_enabled             = false;
 	uint32_t               vblank_frequency            = 60;
 	uint32_t               console_language            = DEFAULT_CONSOLE_LANGUAGE;
 	bool                   vulkan_validation_enabled   = false;
@@ -83,6 +84,7 @@ PresentMode GetPresentMode();
 int32_t GetGpuIndex();
 bool     FullscreenEnabled();
 bool     VrEnabled();
+bool     AmdCpuEnabled();
 uint32_t GetVblankFrequency();
 uint32_t GetConsoleLanguage();
 bool     VulkanValidationEnabled();

src/emulator.cpp

@@ -196,6 +196,10 @@ void Run(const RunOptions& options) {
 	ClearDebugTextureFolder();
 
 	PrintSystemInfo();
+	std::string title_id;
+	if (Loader::SystemContentParamSfoGetString("TITLE_ID", &title_id) && !title_id.empty()) {
+		Log::WriteToConsoleAndLog(fmt::format("Title ID: {}\n", title_id));
+	}
 
 	int ok = atexit(KytyClose);
 	EXIT_NOT_IMPLEMENTED(ok != 0);

src/graphics/guest_gpu/command_processor/commandProcessor.h

@@ -32,11 +32,12 @@ private:
 
 	struct BufferCursor {
 		std::span<const uint32_t> commands;
-		uint32_t                  offset_dw           = 0;
-		uint32_t                  deferred_advance_dw = 0;
+		uint32_t                  offset_dw = 0;
 	};
 
 	std::vector<BufferCursor> m_buffer_stack;
+	std::span<const uint32_t> m_next_buffer;
+	bool                      m_chain         = false;
 	bool                      m_suspended     = false;
 	bool                      m_made_progress = false;
 };
@@ -129,7 +130,7 @@ public:
 	[[nodiscard]] bool ShouldSkipPredicatedPackets() const { return m_predicate_skip; }
 
 	Pm4ProcessResult Process(Pm4Execution& execution, std::span<const uint32_t> commands);
-	void             ProcessIndirectBuffer(std::span<const uint32_t> commands);
+	void             ProcessIndirectBuffer(std::span<const uint32_t> commands, bool chain);
 
 	void SetFlip(const FlipInfo& flip) { m_flip = flip; }
 
@@ -143,7 +144,7 @@ private:
 	                      uint32_t cache_action, uint32_t event_index, uint32_t event_write_source,
 	                      void* dst_gpu_addr, T value, uint32_t interrupt_selector,
 	                      uint32_t interrupt_context_id);
-	void ProcessPm4(Pm4Execution& execution, size_t stop_depth);
+	void ProcessPm4(Pm4Execution& execution);
 	void SuspendPm4();
 	CommandScheduler&   GetScheduler() const { return m_renderer.GetCommandScheduler(); }
 	CommandBuffer&      CurrentBuffer() { return GetScheduler().Current(); }

src/graphics/guest_gpu/command_processor/pm4Handlers.cpp

@@ -1495,10 +1495,10 @@ KYTY_CP_OP_PARSER(CpOpBranch) {
 	     reinterpret_cast<uint64_t>(else_buffer), else_num_dw);
 
 	if (take_then) {
-		cp.ProcessIndirectBuffer({then_buffer, then_num_dw});
+		cp.ProcessIndirectBuffer({then_buffer, then_num_dw}, true);
 	} else if (mode == 2 && else_num_dw != 0) {
 		EXIT_NOT_IMPLEMENTED(else_buffer == nullptr);
-		cp.ProcessIndirectBuffer({else_buffer, else_num_dw});
+		cp.ProcessIndirectBuffer({else_buffer, else_num_dw}, true);
 	}
 
 	return payload_dw;
@@ -1959,7 +1959,7 @@ KYTY_CP_OP_PARSER(CpOpIndirectBuffer) {
 
 	GraphicsDbgDumpDcb("ci", indirect_num_dw, indirect_buffer);
 
-	cp.ProcessIndirectBuffer({indirect_buffer, indirect_num_dw});
+	cp.ProcessIndirectBuffer({indirect_buffer, indirect_num_dw}, (control & (1u << 20u)) != 0);
 
 	return 3;
 }

src/graphics/guest_gpu/graphicsRun.cpp

@@ -672,20 +672,16 @@ Pm4ProcessResult CommandProcessor::Process(Pm4Execution&             execution,
 		Pm4Execution*     previous_execution;
 	} execution_scope(*this, execution);
 
-	ProcessPm4(execution, 0);
+	ProcessPm4(execution);
 	return execution.m_buffer_stack.empty() ? Pm4ProcessResult::Complete
 	                                        : Pm4ProcessResult::Blocked;
 }
 
-void CommandProcessor::ProcessIndirectBuffer(std::span<const uint32_t> commands) {
+void CommandProcessor::ProcessIndirectBuffer(std::span<const uint32_t> commands, bool chain) {
 	EXIT_IF(g_current_execution == nullptr);
-	if (commands.empty()) {
-		return;
-	}
-	auto&      execution  = *g_current_execution;
-	const auto stop_depth = execution.m_buffer_stack.size();
-	execution.m_buffer_stack.push_back({commands});
-	ProcessPm4(execution, stop_depth);
+	EXIT_IF(!g_current_execution->m_next_buffer.empty());
+	g_current_execution->m_next_buffer = commands;
+	g_current_execution->m_chain       = chain;
 }
 
 void CommandProcessor::SuspendPm4() {
@@ -693,21 +689,13 @@ void CommandProcessor::SuspendPm4() {
 	g_current_execution->m_suspended = true;
 }
 
-void CommandProcessor::ProcessPm4(Pm4Execution& execution, size_t stop_depth) {
-	while (execution.m_buffer_stack.size() > stop_depth) {
+void CommandProcessor::ProcessPm4(Pm4Execution& execution) {
+	while (!execution.m_buffer_stack.empty()) {
 		if (g_gpu_state != nullptr) {
 			g_gpu_state->ProcessCommands();
 		}
-		const auto buffer_index = execution.m_buffer_stack.size() - 1;
-		auto&      cursor       = execution.m_buffer_stack[buffer_index];
+		auto& cursor = execution.m_buffer_stack.back();
 		EXIT_IF(cursor.offset_dw > cursor.commands.size());
-		if (cursor.deferred_advance_dw != 0) {
-			EXIT_IF(cursor.deferred_advance_dw > cursor.commands.size() - cursor.offset_dw);
-			cursor.offset_dw += cursor.deferred_advance_dw;
-			cursor.deferred_advance_dw = 0;
-			execution.m_made_progress  = true;
-			continue;
-		}
 		if (cursor.offset_dw == cursor.commands.size()) {
 			execution.m_buffer_stack.pop_back();
 			continue;
@@ -783,14 +771,19 @@ void CommandProcessor::ProcessPm4(Pm4Execution& execution, size_t stop_depth) {
 		    handler(*this, packet_header & ~1u, packet + 1, remaining_dw, total_dw) + 1;
 		EXIT_IF(packet_dw > remaining_dw);
 		if (execution.m_suspended) {
-			if (execution.m_buffer_stack.size() > buffer_index + 1) {
-				execution.m_buffer_stack[buffer_index].deferred_advance_dw = packet_dw;
-			}
 			return;
 		}
-		EXIT_IF(execution.m_buffer_stack.size() != buffer_index + 1);
-		execution.m_buffer_stack[buffer_index].offset_dw += packet_dw;
+		cursor.offset_dw += packet_dw;
 		execution.m_made_progress = true;
+		if (!execution.m_next_buffer.empty()) {
+			// Chains and taken branches reuse the fetcher; only calls retain a return cursor.
+			if (execution.m_chain) {
+				cursor = {execution.m_next_buffer};
+			} else {
+				execution.m_buffer_stack.push_back({execution.m_next_buffer});
+			}
+			execution.m_next_buffer = {};
+		}
 	}
 }
 
@@ -825,36 +818,55 @@ void CommandProcessor::SetNumInstances(uint32_t num_instances) {
 
 void CommandProcessor::SetPredication(uint32_t condition, uint32_t op, uint32_t wait_op,
                                       const volatile void* address, uint32_t count_in_dwords) {
-	if (wait_op != 0) {
-		BufferFlushAndWait();
-	}
-
 	(void)count_in_dwords;
+	uint64_t value = 0;
 
 	switch (op) {
-		case 0x00: {
+		case 0x00:
 			m_predicate_skip = false;
-		} break;
-		case 0x03: {
+			return;
+		case 0x01: {
 			EXIT_NOT_IMPLEMENTED(address == nullptr);
-
-			auto value = *reinterpret_cast<const volatile uint64_t*>(address);
-
-			switch (condition) {
-				case 0x00: m_predicate_skip = (value != 0); break;
-				case 0x01: m_predicate_skip = (value == 0); break;
-				default: EXIT("unknown predication condition: 0x%08" PRIx32 "\n", condition);
-			}
-			static std::atomic<uint32_t> log_count {0};
-			if (log_count.fetch_add(1) < 128) {
-				LOGF("\t bool predication: addr=0x%016" PRIx64 ", value=0x%016" PRIx64
-				     ", condition=%" PRIu32 ", skip=%u, wait_op=%" PRIu32 "\n",
-				     reinterpret_cast<uint64_t>(address), value, condition,
-				     m_predicate_skip ? 1u : 0u, wait_op);
+			// One begin/end pair per DB; bit 63 marks each counter ready.
+			constexpr uint64_t ready_bit = 1ull << 63u;
+			const auto* results = reinterpret_cast<const volatile uint64_t*>(address);
+			for (uint32_t db = 0; db < 16u; db++) {
+				const auto begin = results[db * 2u];
+				const auto end   = results[db * 2u + 1u];
+				if ((begin & end & ready_bit) == 0) {
+					if (wait_op == 0) {
+						SuspendPm4();
+					} else {
+						m_predicate_skip = false;
+					}
+					return;
+				}
+				value += end - begin;
 			}
 		} break;
+		case 0x03:
+			if (wait_op != 0) {
+				BufferFlushAndWait();
+			}
+			EXIT_NOT_IMPLEMENTED(address == nullptr);
+			value = *reinterpret_cast<const volatile uint64_t*>(address);
+			break;
 		default: EXIT("unknown predication op: 0x%08" PRIx32 "\n", op);
 	}
+	switch (condition) {
+		case 0x00: m_predicate_skip = (value != 0); break;
+		case 0x01: m_predicate_skip = (value == 0); break;
+		default: EXIT("unknown predication condition: 0x%08" PRIx32 "\n", condition);
+	}
+	if (op == 0x03) {
+		static std::atomic<uint32_t> log_count {0};
+		if (log_count.fetch_add(1) < 128) {
+			LOGF("\t bool predication: addr=0x%016" PRIx64 ", value=0x%016" PRIx64
+			     ", condition=%" PRIu32 ", skip=%u, wait_op=%" PRIu32 "\n",
+			     reinterpret_cast<uint64_t>(address), value, condition,
+			     m_predicate_skip ? 1u : 0u, wait_op);
+		}
+	}
 }
 
 void CommandProcessor::DrawIndex(DrawIndexArgs args) {
@@ -1090,17 +1102,6 @@ void CommandProcessor::DispatchDirect(uint32_t thread_group_x, uint32_t thread_g
 		// local_x        = std::max(cs.num_thread_x, 1u);
 		// local_y        = std::max(cs.num_thread_y, 1u);
 		// local_z        = std::max(cs.num_thread_z, 1u);
-		if (cs.wave_size == 64u) {
-			static std::atomic_bool logged_wave64_shader {false};
-			if (!logged_wave64_shader.exchange(true, std::memory_order_relaxed)) {
-				LOGF("warning: executing wave64 compute shader cs=0x%016" PRIx64 "\n",
-				     cs.data_addr);
-				std::printf("warning: executing wave64 compute shader cs=0x%016" PRIx64 "\n",
-				            cs.data_addr);
-				std::fflush(stdout);
-			}
-		}
-
 		m_renderer.GetRenderExecutor().DispatchDirect(m_submit_id, CurrentBuffer(), thread_group_x,
 		                                              thread_group_y, thread_group_z, mode);
 	}

src/graphics/presentation/window/hostInput.cpp

@@ -220,6 +220,8 @@ uint32_t DefaultKeyboardButton(int key_code) {
 		case SDLK_l: return Controller::PAD_BUTTON_CIRCLE;
 		case SDLK_q: return Controller::PAD_BUTTON_L1;
 		case SDLK_e: return Controller::PAD_BUTTON_R1;
+		case SDLK_z: return Controller::PAD_BUTTON_L2;
+		case SDLK_c: return Controller::PAD_BUTTON_R2;
 		case SDLK_LSHIFT: return Controller::PAD_BUTTON_L3;
 		case SDLK_LCTRL: return Controller::PAD_BUTTON_R3;
 		case SDLK_RETURN:

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

@@ -41,7 +41,9 @@ uint32_t EmitBuiltinU32(EmitterState& state, IR::StageInputKind kind, uint32_t c
 		return EmitAddU32(state, local,
 		                  EmitBinaryU32(state, spv::OpIMul, group, ConstantU32(state, size)));
 	}
-	const auto variable = InputVariableForKind(state, kind);
+	const bool centroid = kind == IR::StageInputKind::BaryCoordSmoothCentroid;
+	const auto variable = InputVariableForKind(
+	    state, centroid ? IR::StageInputKind::BaryCoordSmooth : kind);
 	if (variable == 0) {
 		return ConstantU32(state, 0);
 	}
@@ -74,15 +76,24 @@ uint32_t EmitBuiltinU32(EmitterState& state, IR::StageInputKind kind, uint32_t c
 		state.builder.AddFunction(spv::OpBitcast, TypeU32(state), bits, value);
 		return bits;
 	}
-	if (kind == IR::StageInputKind::BaryCoordSmooth ||
+	if (centroid || kind == IR::StageInputKind::BaryCoordSmooth ||
 	    kind == IR::StageInputKind::BaryCoordNoPerspective) {
-		const auto pointer = state.builder.AllocateId();
 		const auto value   = state.builder.AllocateId();
 		const auto bits    = state.builder.AllocateId();
-		state.builder.AddFunction(spv::OpAccessChain,
-		                          TypePointer(state, spv::StorageClassInput, TypeF32(state)),
-		                          pointer, variable, ConstantU32(state, component + 1u));
-		state.builder.AddFunction(spv::OpLoad, TypeF32(state), value, pointer);
+		if (centroid) {
+			const auto coordinates = state.builder.AllocateId();
+			state.builder.RequireCapability(spv::CapabilityInterpolationFunction);
+			state.builder.AddFunction(spv::OpExtInst, TypeF32Vector(state, 3), coordinates,
+			                          GlslStd450(state), GLSLstd450InterpolateAtCentroid, variable);
+			state.builder.AddFunction(spv::OpCompositeExtract, TypeF32(state), value,
+			                          coordinates, component + 1u);
+		} else {
+			const auto pointer = state.builder.AllocateId();
+			state.builder.AddFunction(spv::OpAccessChain,
+			                          TypePointer(state, spv::StorageClassInput, TypeF32(state)),
+			                          pointer, variable, ConstantU32(state, component + 1u));
+			state.builder.AddFunction(spv::OpLoad, TypeF32(state), value, pointer);
+		}
 		state.builder.AddFunction(spv::OpBitcast, TypeU32(state), bits, value);
 		return bits;
 	}

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

@@ -388,6 +388,19 @@ void DefineInputs(EmitterState& state) {
 		}
 	}
 	for (auto& input: state.inputs) {
+		if (state.program.stage == ShaderType::Pixel &&
+		    input.kind == IR::StageInputKind::Parameter) {
+			const auto location = PixelParameterLocation(state, input.location);
+			const auto alias = std::ranges::find_if(state.inputs, [&](const InputBinding& other) {
+				return other.kind == IR::StageInputKind::Parameter && other.variable_id != 0 &&
+				       PixelParameterLocation(state, other.location) == location;
+			});
+			if (alias != state.inputs.end()) {
+				EXIT_IF(alias->per_vertex != input.per_vertex);
+				input.variable_id = alias->variable_id;
+				continue;
+			}
+		}
 		uint32_t type = TypeU32(state);
 		switch (input.kind) {
 			case IR::StageInputKind::VertexIndex:

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

@@ -516,7 +516,7 @@ struct Vop1SdwaRule {
 };
 
 constexpr Vop1SdwaRule VOP1_SDWA_RULES[] = {
-    {Opcode::V_MOV_B32, SdwaSelBytes() | SdwaSelWords() | SdwaSelFull(), SdwaSelWords(),
+    {Opcode::V_MOV_B32, SdwaSelBytes() | SdwaSelWords() | SdwaSelFull(), SdwaSelBytes() | SdwaSelWords(),
      SdwaSelWords() | SdwaSelFull(), false},
     {Opcode::V_CVT_F32_U32, SdwaSelBytes() | SdwaSelWords() | SdwaSelFull(), 0, 0, false},
     {Opcode::V_CVT_F32_I32, SdwaSelBytes() | SdwaSelWords() | SdwaSelFull(), 0, 0, false},

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

@@ -1102,26 +1102,32 @@ IR::Program TranslateProgram(const Decoder::Program& decoded, const CFG::Graph&
 			    entry_ir.BitwiseOr(wave_info, entry_ir.BitwiseOr(entry_ir.ShiftLeftLogical(
 			                                                         primitive_count, u32(8)),
 			                                                     vertex_count)));
-			// GS adjacency addresses local ES records in LDS. Strip winding alternates
-			// with the global primitive number, including across subgroup boundaries.
-			const auto parity = mesh.input_primitive ==
-			                            static_cast<uint32_t>(Prospero::PrimitiveType::kTriStrip)
-			                        ? entry_ir.BitwiseAnd(entry_ir.IAdd(primitive_chunk, local), u32(1))
-			                        : u32(0);
+			// GS adjacency addresses local ES records in LDS. Fans retain the draw's
+			// center in every subgroup; strip winding follows the global primitive.
 			const auto vertex = entry_ir.IMul(local, step);
-			const auto first  = entry_ir.IAdd(vertex, parity);
-			const auto second = mesh.InputPrimitiveSize() >= 2u
-			                        ? entry_ir.ISub(entry_ir.IAdd(vertex, u32(1)), parity)
-			                        : u32(0);
-			const auto third = mesh.InputPrimitiveSize() == 3u
-			                       ? entry_ir.IAdd(vertex, u32(2))
-			                       : u32(0);
+			auto       first  = vertex;
+			auto       second = u32(0);
+			auto       third  = u32(0);
+			if (mesh.InputPrimitiveSize() >= 2u) {
+				second = entry_ir.IAdd(vertex, u32(1));
+			}
+			if (mesh.InputPrimitiveSize() == 3u) {
+				third = entry_ir.IAdd(vertex, u32(2));
+			}
+			auto input_vertex = entry_ir.IAdd(chunk, local);
+			if (mesh.input_primitive == static_cast<uint32_t>(Prospero::PrimitiveType::kTriFan)) {
+				first = u32(0);
+				input_vertex = entry_ir.Select(entry_ir.IEqual(local, u32(0)), u32(0), input_vertex);
+			} else if (mesh.input_primitive == static_cast<uint32_t>(Prospero::PrimitiveType::kTriStrip)) {
+				const auto parity = entry_ir.BitwiseAnd(entry_ir.IAdd(primitive_chunk, local), u32(1));
+				first = entry_ir.IAdd(first, parity);
+				second = entry_ir.ISub(second, parity);
+			}
 			entry_ir.SetVectorReg(static_cast<IR::VectorReg>(0),
 			                      entry_ir.BitwiseOr(entry_ir.ShiftLeftLogical(first, u32(2)),
 			                                         entry_ir.ShiftLeftLogical(second, u32(18))));
 			entry_ir.SetVectorReg(static_cast<IR::VectorReg>(1),
 			                      entry_ir.ShiftLeftLogical(third, u32(2)));
-			const auto input_vertex = entry_ir.IAdd(chunk, local);
 			const auto index_bytes  = draw(3);
 			const auto indexed      = entry_ir.INotEqual(index_bytes, u32(0));
 			const auto index_low    = draw(4);
@@ -1184,13 +1190,15 @@ IR::Program TranslateProgram(const Decoder::Program& decoded, const CFG::Graph&
 			                      builtin(IR::StageInputKind::PrimitiveId));
 		} else if (options.stage == ShaderType::Pixel) {
 			const auto* ps = options.input_info.pixel;
-			if (ps->ps_perspective_center_vgpr != UINT32_MAX) {
-				entry_ir.SetVectorReg(static_cast<IR::VectorReg>(ps->ps_perspective_center_vgpr),
-				                      builtin(IR::StageInputKind::BaryCoordSmooth, 0));
-				entry_ir.SetVectorReg(
-				    static_cast<IR::VectorReg>(ps->ps_perspective_center_vgpr + 1u),
-				    builtin(IR::StageInputKind::BaryCoordSmooth, 1));
-			}
+			const auto barycentric_pair = [&](uint32_t reg, IR::StageInputKind kind) {
+				if (reg != UINT32_MAX) {
+					entry_ir.SetVectorReg(static_cast<IR::VectorReg>(reg), builtin(kind, 0));
+					entry_ir.SetVectorReg(static_cast<IR::VectorReg>(reg + 1u), builtin(kind, 1));
+				}
+			};
+			barycentric_pair(ps->ps_perspective_center_vgpr, IR::StageInputKind::BaryCoordSmooth);
+			barycentric_pair(ps->ps_perspective_centroid_vgpr,
+			                 IR::StageInputKind::BaryCoordSmoothCentroid);
 			uint32_t reg = ps->ps_system_input_base;
 			if (ps->ps_pos_x) {
 				entry_ir.SetVectorReg(static_cast<IR::VectorReg>(reg++),

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

@@ -171,6 +171,7 @@ enum class StageInputKind {
 	Layer,
 	SampleId,
 	BaryCoordSmooth,
+	BaryCoordSmoothCentroid,
 	BaryCoordNoPerspective,
 	WorkgroupId,
 	LocalInvocationId,

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

@@ -143,6 +143,7 @@ void ValidateValueReferences(const Program& program, ShaderStageInputInfo input_
 							}
 							break;
 						case StageInputKind::BaryCoordSmooth:
+						case StageInputKind::BaryCoordSmoothCentroid:
 						case StageInputKind::BaryCoordNoPerspective:
 							if (component >= 2u) {
 								return Fail("typed barycentric component is out of range");
@@ -225,6 +226,18 @@ void CollectPixelInputs(const Program& program, const ShaderPixelInputInfo* pixe
 			}
 		}
 	}
+	// Aliases of a vertex output share one SPIR-V interface variable. If any
+	// alias reads raw vertices, interpolate the other aliases from those too.
+	for (uint32_t input = 0; input < pixel->input_num; input++) {
+		for (uint32_t alias = 0; alias < pixel->input_num; alias++) {
+			if (ShaderPixelParameterMappedLocation(*pixel, input) ==
+			        ShaderPixelParameterMappedLocation(*pixel, alias) &&
+			    ShaderPixelParameterIsFlat(*pixel, input) ==
+			        ShaderPixelParameterIsFlat(*pixel, alias)) {
+				per_vertex[input] = per_vertex[input] || per_vertex[alias];
+			}
+		}
+	}
 	for (uint32_t input = 0; input < pixel->input_num; input++) {
 		AddInput(info, StageInputKind::Parameter, input, 4, fmt::format("in_param_{}", input),
 		         per_vertex[input]);
@@ -287,7 +300,8 @@ void CollectBuiltinInputs(const Program& program, ShaderInfo& info) {
 				case StageInputKind::Layer: AddInput(info, kind, 0, 1, "gl_Layer"); break;
 				case StageInputKind::SampleId: AddInput(info, kind, 0, 1, "gl_SampleID"); break;
 				case StageInputKind::BaryCoordSmooth:
-					AddInput(info, kind, 0, 3, "gl_BaryCoordKHR");
+				case StageInputKind::BaryCoordSmoothCentroid:
+					AddInput(info, StageInputKind::BaryCoordSmooth, 0, 3, "gl_BaryCoordKHR");
 					break;
 				case StageInputKind::BaryCoordNoPerspective:
 					AddInput(info, kind, 0, 3, "gl_BaryCoordNoPerspKHR");

src/graphics/shader/rectListShader.cpp

@@ -40,12 +40,15 @@ std::vector<Parameter> GetParameters(const ShaderVertexInputInfo& vertex_info,
 	}
 
 	std::vector<Parameter> parameters;
+	std::array<bool, 32> output_locations {};
 	for (const auto input: active_inputs) {
 		const auto input_location = ShaderPixelParameterMappedLocation(*pixel_info, input);
-		if ((vertex_info.stage.program->param_export_mask & (1u << input_location)) != 0) {
-			parameters.push_back({input_location,
-			                      ShaderPixelParameterLocation(*pixel_info, active_inputs, input),
+		const auto output_location = ShaderPixelParameterLocation(*pixel_info, active_inputs, input);
+		if ((vertex_info.stage.program->param_export_mask & (1u << input_location)) != 0 &&
+		    !output_locations[output_location]) {
+			parameters.push_back({input_location, output_location,
 			                      ShaderPixelParameterIsFlat(*pixel_info, input)});
+			output_locations[output_location] = true;
 		}
 	}
 	return parameters;

src/graphics/shader/shader.cpp

@@ -584,6 +584,9 @@ static void ShaderGetStaticInputInfoPS(
 	if ((active_inputs & 0x00000002u) != 0) {
 		ps_info.ps_perspective_center_vgpr = (active_inputs & 0x00000001u) != 0 ? 2u : 0u;
 	}
+	if ((active_inputs & 0x00000004u) != 0) {
+		ps_info.ps_perspective_centroid_vgpr = 2u * std::popcount(active_inputs & 0x3u);
+	}
 	for (uint32_t i = 0; i < data.num_input_semantics && i < ps_info.input_num && i < 32u; i++) {
 		const auto& semantic = data.input_semantics[i];
 		if (semantic.is_custom != 0 && semantic.is_f16 == 0) {
@@ -707,6 +710,7 @@ void BuildStageStaticKey(const ShaderPixelInputInfo& info, std::vector<uint32_t>
 	key.push_back(info.ps_system_input_base);
 	key.push_back(info.custom_interpolation_mask);
 	key.push_back(info.ps_perspective_center_vgpr);
+	key.push_back(info.ps_perspective_centroid_vgpr);
 	key.push_back(static_cast<uint32_t>(info.ps_pos_x));
 	key.push_back(static_cast<uint32_t>(info.ps_pos_y));
 	key.push_back(static_cast<uint32_t>(info.ps_pos_z));
@@ -793,6 +797,7 @@ ShaderParams PrepareProgram(const HW::VertexShaderInfo& regs, const HW::Context&
 	const auto& group = user_config.GetGeControl();
 	if ((user_config.GetPrimType() != Prospero::PrimitiveType::kPointList &&
 	     user_config.GetPrimType() != Prospero::PrimitiveType::kLineList &&
+	     user_config.GetPrimType() != Prospero::PrimitiveType::kTriFan &&
 	     user_config.GetPrimType() != Prospero::PrimitiveType::kTriStrip &&
 	     user_config.GetPrimType() != Prospero::PrimitiveType::kTriList) ||
 	    sh.m_vgtGsOutPrimType != 2u || sh.m_vgtGsMaxVertOut < 3u ||
@@ -958,6 +963,7 @@ void ShaderDbgDumpInputInfo(const ShaderPixelInputInfo& info) {
 	     "\t ps_system_input_base = %u\n"
 	     "\t custom_interpolation_mask = 0x%08" PRIx32 "\n"
 	     "\t ps_perspective_center_vgpr = %" PRIu32 "\n"
+	     "\t ps_perspective_centroid_vgpr = %" PRIu32 "\n"
 	     "\t ps_pos_x             = %s\n"
 	     "\t ps_pos_y             = %s\n"
 	     "\t ps_pos_z             = %s\n"
@@ -970,7 +976,8 @@ void ShaderDbgDumpInputInfo(const ShaderPixelInputInfo& info) {
 	     "\t ps_early_z           = %s\n"
 	     "\t ps_execute_on_noop   = %s\n",
 	     info.input_num, info.ps_system_input_base, info.custom_interpolation_mask,
-	     info.ps_perspective_center_vgpr, info.ps_pos_x ? "true" : "false",
+	     info.ps_perspective_center_vgpr, info.ps_perspective_centroid_vgpr,
+	     info.ps_pos_x ? "true" : "false",
 	     info.ps_pos_y ? "true" : "false", info.ps_pos_z ? "true" : "false",
 	     info.ps_pos_w ? "true" : "false", info.ps_front_face ? "true" : "false",
 	     info.ps_ancillary ? "true" : "false",

src/graphics/shader/shader.h

@@ -97,8 +97,11 @@ struct ShaderMeshInputInfo: ShaderWorkgroupInputInfo {
 		}
 	}
 	[[nodiscard]] constexpr uint32_t InputPrimitiveStep() const {
-		return input_primitive == static_cast<uint32_t>(Prospero::PrimitiveType::kTriStrip)
-		           ? 1u : InputPrimitiveSize();
+		switch (static_cast<Prospero::PrimitiveType>(input_primitive)) {
+			case Prospero::PrimitiveType::kTriFan:
+			case Prospero::PrimitiveType::kTriStrip: return 1u;
+			default: return InputPrimitiveSize();
+		}
 	}
 	[[nodiscard]] constexpr uint32_t InputPrimitiveCount(uint32_t vertices) const {
 		const auto size = InputPrimitiveSize();
@@ -157,6 +160,7 @@ struct ShaderPixelInputInfo {
 	uint32_t                                       ps_system_input_base         = 0;
 	uint32_t                                       custom_interpolation_mask    = 0;
 	uint32_t                                       ps_perspective_center_vgpr   = UINT32_MAX;
+	uint32_t                                       ps_perspective_centroid_vgpr = UINT32_MAX;
 	uint8_t                                        target_output_mode[8]        = {};
 	std::array<Prospero::ColorComponentMapping, 8> target_export_mapping        = {};
 	uint32_t                                       scratch_size_dwords          = 0;

src/graphics/shader/shaderPixelParameter.cpp

@@ -19,19 +19,31 @@ uint32_t ShaderPixelParameterLocation(const ShaderPixelInputInfo& info,
                                       std::span<const uint32_t> active_inputs, uint32_t input) {
 	std::array<bool, 32> used_locations {};
 	for (const auto active_input: active_inputs) {
-		auto location = ShaderPixelParameterMappedLocation(info, active_input);
-		if (location < used_locations.size() && used_locations[location]) {
-			location = active_input;
-			while (location < used_locations.size() && used_locations[location]) {
-				location++;
+		used_locations[ShaderPixelParameterMappedLocation(info, active_input)] = true;
+	}
+	std::array<uint32_t, 64> group_locations;
+	group_locations.fill(UINT32_MAX);
+	for (const auto active_input: active_inputs) {
+		const auto mapped = ShaderPixelParameterMappedLocation(info, active_input);
+		const auto group  = mapped * 2u + ShaderPixelParameterIsFlat(info, active_input);
+		auto&      location = group_locations[group];
+		if (location == UINT32_MAX) {
+			location = mapped;
+			// Smooth and custom interpolation read the same vertex output. Only
+			// differing flat/smooth rectangle outputs need separate locations.
+			if (group_locations[group ^ 1u] != UINT32_MAX) {
+				location = 0;
+				while (location < used_locations.size() && used_locations[location]) {
+					location++;
+				}
+				EXIT_NOT_IMPLEMENTED(location >= used_locations.size());
 			}
-			EXIT_NOT_IMPLEMENTED(location >= used_locations.size());
+			used_locations[location] = true;
 		}
 
 		if (active_input == input) {
 			return location;
 		}
-		used_locations[location] = true;
 	}
 	return ShaderPixelParameterMappedLocation(info, input);
 }

src/launcher/forms/configuration_edit_dialog.ui

@@ -271,21 +271,19 @@
        <property name="title">
         <string>Compatibility</string>
        </property>
-       <layout class="QHBoxLayout" name="compatibilityLayout">
+       <layout class="QVBoxLayout" name="compatibilityLayout">
+        <item>
+         <widget class="QCheckBox" name="checkBox_amd_cpu">
+          <property name="toolTip"><string>Patch CPU instructions that differ between AMD and the host CPU.</string></property>
+          <property name="text"><string>AMD CPU patch (EXPERIMENTAL)</string></property>
+         </widget>
+        </item>
         <item>
          <widget class="QCheckBox" name="checkBox_red_zone_protection">
           <property name="toolTip"><string>Helps prevent crashes in some games on Windows</string></property>
           <property name="text"><string>Windows SysV red zone crash protection (Experimental)</string></property>
          </widget>
         </item>
-        <item>
-         <spacer name="compatibilitySpacer">
-          <property name="orientation"><enum>Qt::Horizontal</enum></property>
-          <property name="sizeHint" stdset="0">
-           <size><width>40</width><height>20</height></size>
-          </property>
-         </spacer>
-        </item>
        </layout>
       </widget>
      </item>

src/launcher/include/configuration.h

@@ -105,6 +105,7 @@ public:
 	QString                printf_output_file          = "_kyty.txt";
 	bool                   profiler_enabled            = false;
 	bool                   renderdoc_enabled           = false;
+	bool                   amd_cpu_enabled             = false;
 #if defined(_WIN32)
 	bool red_zone_protection_enabled = false;
 #endif
@@ -134,6 +135,7 @@ public:
 		printf_output_file          = other.printf_output_file;
 		profiler_enabled            = other.profiler_enabled;
 		renderdoc_enabled           = other.renderdoc_enabled;
+		amd_cpu_enabled             = other.amd_cpu_enabled;
 #if defined(_WIN32)
 		red_zone_protection_enabled = other.red_zone_protection_enabled;
 #endif
@@ -180,6 +182,7 @@ public:
 		KYTY_CFG_SET(printf_output_file);
 		KYTY_CFG_SET(profiler_enabled);
 		KYTY_CFG_SET(renderdoc_enabled);
+		KYTY_CFG_SET(amd_cpu_enabled);
 #if defined(_WIN32)
 		KYTY_CFG_SET(red_zone_protection_enabled);
 #endif
@@ -223,6 +226,7 @@ public:
 		KYTY_CFG_GET(printf_output_file);
 		KYTY_CFG_GET(profiler_enabled);
 		KYTY_CFG_GET(renderdoc_enabled);
+		amd_cpu_enabled = s->value("amd_cpu_enabled", false).toBool();
 #if defined(_WIN32)
 		red_zone_protection_enabled =
 		    s->value("red_zone_protection_enabled", red_zone_protection_enabled).toBool();

src/launcher/src/configurationEditDialog.cpp

@@ -239,6 +239,10 @@ void ConfigurationEditDialog::Init(const Configuration& info) {
 	m_ui->checkBox_shader_validation->setChecked(info.shader_validation_enabled);
 	m_ui->checkBox_vulkan_validation->setChecked(info.vulkan_validation_enabled);
 	m_ui->checkBox_renderdoc_capture->setChecked(info.renderdoc_enabled);
+	m_ui->checkBox_amd_cpu->setChecked(info.amd_cpu_enabled);
+#if defined(__APPLE__)
+	m_ui->checkBox_amd_cpu->setVisible(false);
+#endif
 #if defined(_WIN32)
 	m_ui->checkBox_red_zone_protection->setChecked(info.red_zone_protection_enabled);
 #else
@@ -379,6 +383,7 @@ static void UpdateInfo(Configuration& info, Ui::ConfigurationEditDialog& ui) {
 	info.vulkan_validation_enabled = ui.checkBox_vulkan_validation->isChecked();
 	info.shader_validation_enabled = ui.checkBox_shader_validation->isChecked();
 	info.renderdoc_enabled         = ui.checkBox_renderdoc_capture->isChecked();
+	info.amd_cpu_enabled           = ui.checkBox_amd_cpu->isChecked();
 #if defined(_WIN32)
 	info.red_zone_protection_enabled = ui.checkBox_red_zone_protection->isChecked();
 #endif

src/launcher/src/inputMappingDialog.cpp

@@ -46,8 +46,8 @@ constexpr PadControl PAD_CONTROLS[] = {
     {"Square", "Square", "K"},
     {"L1", "L1", "Q"},
     {"R1", "R1", "E"},
-    {"L2", "L2", ""},
-    {"R2", "R2", ""},
+    {"L2", "L2", "Z"},
+    {"R2", "R2", "C"},
     {"L3", "L3", "Left Shift"},
     {"R3", "R3", "Left Ctrl"},
     {"Options", "Options", "Return"},

src/launcher/src/mainDialog.cpp

@@ -254,6 +254,9 @@ static QStringList CreateEmulatorArgs(const Configuration& info) {
 		args << "--profile";
 	}
 	args << "--spirv-debug-printf" << "false";
+	if (info.amd_cpu_enabled) {
+		args << "--amd-cpu";
+	}
 #if defined(_WIN32)
 	if (info.red_zone_protection_enabled) {
 		args << "--redzone";

src/libs/avPlayer.cpp

@@ -384,6 +384,16 @@ static bool codec_is_supported_for_source(AvPlayerSourceType source_type, AVMedi
 	}
 }
 
+static const AVCodec* find_decoder(AVCodecID codec_id) {
+	const auto* decoder = avcodec_find_decoder(codec_id);
+	if (decoder == nullptr && codec_id == AV_CODEC_ID_VP9) {
+		Log::WriteToConsoleAndLog(
+		    "WARNING: AvPlayer: The game requested VP9 video, but FFmpeg has no VP9 "
+		    "decoder. Use FFmpeg with VP9 decoding enabled.\n");
+	}
+	return decoder;
+}
+
 struct PacketDeleter {
 	void operator()(AVPacket* packet) const { av_packet_free(&packet); }
 };
@@ -1064,7 +1074,7 @@ private:
 			     static_cast<int>(s->codecpar->codec_id), static_cast<uint32_t>(source_type));
 			return false;
 		}
-		if (avcodec_find_decoder(s->codecpar->codec_id) == nullptr) {
+		if (find_decoder(s->codecpar->codec_id) == nullptr) {
 			LOGF("\t decoder not found: stream=%d codec=%d\n", id,
 			     static_cast<int>(s->codecpar->codec_id));
 			return false;
@@ -1073,7 +1083,7 @@ private:
 	}
 	bool OpenCodec(int id, AVCodecContext** out) {
 		auto* s   = fmt->streams[id];
-		auto* dec = avcodec_find_decoder(s->codecpar->codec_id);
+		auto* dec = find_decoder(s->codecpar->codec_id);
 		if (dec == nullptr) {
 			LOGF("\t avcodec_find_decoder failed: stream=%d codec=%d\n", id,
 			     static_cast<int>(s->codecpar->codec_id));

src/libs/libFont.cpp

@@ -38,7 +38,12 @@
 #endif
 
 extern "C" {
+#include <libavutil/version.h>
+#if LIBAVUTIL_VERSION_MAJOR < 61
 extern const uint8_t avpriv_vga16_font[4096];
+#else
+const uint8_t* avpriv_vga16_font_get(void);
+#endif
 }
 
 namespace Libs {
@@ -950,10 +955,15 @@ static void init_image(std::array<uint8_t, FONT_BITMAP_MAX_DIM * FONT_BITMAP_MAX
 
 	const auto width  = static_cast<uint32_t>(scaled_font_width(font));
 	const auto height = static_cast<uint32_t>(scaled_font_height(font));
+#if LIBAVUTIL_VERSION_MAJOR < 61
+	const auto* vga_font = avpriv_vga16_font;
+#else
+	const auto* vga_font = avpriv_vga16_font_get();
+#endif
 
 	for (uint32_t y = 0; y < height; y++) {
 		const uint32_t src_y = std::min<uint32_t>((y * 16u) / height, 15u);
-		const uint8_t  row   = avpriv_vga16_font[code * 16 + src_y];
+		const uint8_t  row   = vga_font[code * 16 + src_y];
 		for (uint32_t x = 0; x < width; x++) {
 			const uint32_t src_x = std::min<uint32_t>((x * 8u) / width, 7u);
 			if ((row & (0x80u >> src_x)) != 0) {

src/libs/libJson2.cpp

@@ -751,7 +751,11 @@ static const int64_t* KYTY_SYSV_ABI JsonValueGetInteger(const JsonValue* self) {
 	PRINT_NAME();
 
 	static const int64_t zero = 0;
-	return (self != nullptr && self->type == JsonValueTypeInteger ? &self->integer : &zero);
+	if (self != nullptr &&
+	    (self->type == JsonValueTypeInteger || self->type == JsonValueTypeUInteger)) {
+		return &self->integer;
+	}
+	return &zero;
 }
 
 static const double* KYTY_SYSV_ABI JsonValueGetReal(const JsonValue* self) {

src/libs/videoDec2Decoder.cpp

@@ -92,6 +92,11 @@ public:
 	[[nodiscard]] bool Initialize() {
 		const AVCodec* decoder = avcodec_find_decoder(GetAvCodecId(m_config.codec_type));
 		if (decoder == nullptr) {
+			if (m_config.codec_type == CODEC_TYPE_VP9) {
+				Log::WriteToConsoleAndLog(
+				    "WARNING: Videodec2: The game requested VP9 video, but FFmpeg has no VP9 "
+				    "decoder. Use FFmpeg with VP9 decoding enabled.\n");
+			}
 			LOGF("Videodec2: FFmpeg decoder is unavailable for codec type %u\n",
 			     m_config.codec_type);
 			return false;

src/loader/redZonePatcher.h

@@ -14,7 +14,7 @@ namespace Loader {
 struct RedZonePatchResult {
 	uint64_t function_count                           = 0;
 	uint64_t instruction_count                        = 0;
-	uint64_t red_zone_function_count                  = 0;
+	uint64_t red_zone_function_count                   = 0;
 	uint64_t memory_instruction_count                 = 0;
 	uint64_t short_memory_instruction_count           = 0;
 	uint64_t patched_memory_instruction_count         = 0;
@@ -22,14 +22,17 @@ struct RedZonePatchResult {
 	uint64_t control_flow_memory_instruction_count    = 0;
 	uint64_t unrelocatable_memory_instruction_count   = 0;
 	uint64_t indirect_red_zone_function_count         = 0;
+	uint64_t reciprocal_sqrt_instruction_count        = 0;
 };
 
 void RegisterRedZonePatchModule(void* module_ptr, uint64_t module_size, void* trampoline_area_ptr,
                                 uint64_t trampoline_area_size);
 void UnregisterRedZonePatchModule(void* module_ptr);
 
-RedZonePatchResult PatchRedZoneMemoryInstructions(uint64_t segment_addr, uint64_t segment_size,
-                                                  std::span<const uintptr_t> function_starts);
+// Analyze and protect fault sites before replacing reciprocal roots with traps.
+RedZonePatchResult PatchGuestInstructions(uint64_t segment_addr, uint64_t segment_size,
+                                          std::span<const uintptr_t> function_starts,
+                                          bool protect_memory, bool emulate_rsqrt);
 
 bool DecodeEhFrameFunctionStarts(uint64_t eh_frame_header_addr, uint64_t eh_frame_header_size,
                                  std::vector<uintptr_t>* function_starts);

src/loader/runtimeLinker.cpp

@@ -2014,9 +2014,11 @@ void RuntimeLinker::LoadProgramToMemory(Program* program) {
 	uint64_t tls_handler_size = is_shared ? 0 : Jit::SafeCall::GetSize();
 	EXIT_IF(tls_handler_size > UINT64_MAX - program->base_size_aligned);
 	program->mapped_size = program->base_size_aligned + tls_handler_size;
+	const bool emulate_rsqrt = Config::AmdCpuEnabled();
 
 #if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
-	const bool         use_red_zone_protection  = Config::RedZoneProtectionEnabled();
+	const bool         protect_memory_faults   = Config::RedZoneProtectionEnabled();
+	const bool         use_red_zone_protection = protect_memory_faults || emulate_rsqrt;
 	constexpr uint64_t RED_ZONE_TRAMPOLINE_SIZE = 8u * 1024u * 1024u;
 	if (use_red_zone_protection) {
 		EXIT_IF(RED_ZONE_TRAMPOLINE_SIZE > UINT64_MAX - program->mapped_size);
@@ -2062,9 +2064,8 @@ void RuntimeLinker::LoadProgramToMemory(Program* program) {
 		EXIT("Failed to install the required vectored exception handler\n");
 	}
 
-	// program->elf->SetBaseVAddr(program->base_vaddr);
-#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
 	std::vector<std::pair<uint64_t, uint64_t>> executable_segments;
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
 	uint64_t                                   eh_frame_header_addr = 0;
 	uint64_t                                   eh_frame_header_size = 0;
 #endif
@@ -2090,11 +2091,7 @@ void RuntimeLinker::LoadProgramToMemory(Program* program) {
 
 			if (Common::VirtualMemory::IsExecute(mode)) {
 				PatchProgram(program, segment_addr, segment_memory_size);
-#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
-				if (use_red_zone_protection) {
-					executable_segments.emplace_back(segment_addr, segment_file_size);
-				}
-#endif
+				executable_segments.emplace_back(segment_addr, segment_file_size);
 			}
 
 			if (!skip_protect) {
@@ -2139,16 +2136,22 @@ void RuntimeLinker::LoadProgramToMemory(Program* program) {
 	}
 
 #if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+	std::vector<uintptr_t> function_starts;
 	if (use_red_zone_protection) {
-		std::vector<uintptr_t> function_starts;
 		if (!DecodeEhFrameFunctionStarts(eh_frame_header_addr, eh_frame_header_size,
 		                                 &function_starts)) {
 			LOGF("Windows guest red-zone patching could not decode function boundaries for %s\n",
 			     Common::PathToString(program->file_name).c_str());
 		}
-		for (const auto& [segment_addr, segment_size]: executable_segments) {
+	}
+#endif
+	for (const auto& [segment_addr, segment_size]: executable_segments) {
+		uint64_t reciprocal_sqrt_count = 0;
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+		if (use_red_zone_protection) {
 			const auto result =
-			    PatchRedZoneMemoryInstructions(segment_addr, segment_size, function_starts);
+			    PatchGuestInstructions(segment_addr, segment_size, function_starts,
+			                           protect_memory_faults, emulate_rsqrt);
 			LOGF("Windows guest red-zone patching: %s, functions=%" PRIu64 ", red_zone=%" PRIu64
 			     ", memory=%" PRIu64 ", patched=%" PRIu64 ", short=%" PRIu64 ", stack=%" PRIu64
 			     ", control=%" PRIu64 ", unrelocatable=%" PRIu64 "\n",
@@ -2158,12 +2161,20 @@ void RuntimeLinker::LoadProgramToMemory(Program* program) {
 			     result.stack_dependent_memory_instruction_count,
 			     result.control_flow_memory_instruction_count,
 			     result.unrelocatable_memory_instruction_count);
+			reciprocal_sqrt_count = result.reciprocal_sqrt_instruction_count;
+		}
+#else
+		if (emulate_rsqrt) {
+			reciprocal_sqrt_count =
+			    X64InstructionEmulator::PatchReciprocalSquareRoots(segment_addr, segment_size);
 			Common::VirtualMemory::FlushInstructionCache(segment_addr, segment_size);
 		}
-		Common::VirtualMemory::FlushInstructionCache(program->red_zone_trampoline_vaddr,
-		                                             program->red_zone_trampoline_size);
-	}
 #endif
+		if (reciprocal_sqrt_count != 0) {
+			LOGF("Guest VRSQRTPS emulation: %s, instructions=%" PRIu64 "\n",
+			     Common::PathToString(program->file_name.filename()).c_str(), reciprocal_sqrt_count);
+		}
+	}
 
 	if (!is_shared) {
 		SetupTlsHandler(program);

src/loader/x64InstructionEmulator.cpp

@@ -2,7 +2,13 @@
 
 #include "common/common.h"
 
+#include <Zydis/Zydis.h>
+#include <bit>
 #include <cstring>
+#if !defined(__APPLE__)
+#include <emmintrin.h>
+#include <xmmintrin.h>
+#endif
 
 #if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
 #include <windows.h> // IWYU pragma: keep
@@ -400,57 +406,89 @@ static bool ExecuteShaNiInsn(const ShaNiInsn& insn, const XmmWords& src2, const
 	}
 }
 
+#if !defined(__APPLE__)
+
+// Keep instruction semantics shared; only access to the saved host context differs.
+struct Context {
 #if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+	PCONTEXT native;
 
-static void LoadXmmWordsWin(const M128A* xmm, XmmWords& out) {
-	out.w[0] = static_cast<uint32_t>(xmm->Low);
-	out.w[1] = static_cast<uint32_t>(xmm->Low >> 32u);
-	out.w[2] = static_cast<uint32_t>(xmm->High);
-	out.w[3] = static_cast<uint32_t>(xmm->High >> 32u);
-}
+	[[nodiscard]] uint64_t Rip() const { return native->Rip; }
+	void                   Advance(size_t length) { native->Rip += length; }
+	[[nodiscard]] void*    Xmm(uint8_t index) const { return &native->Xmm0 + index; }
 
-static void StoreXmmWordsWin(M128A* xmm, const XmmWords& in) {
-	xmm->Low  = static_cast<uint64_t>(in.w[0]) | (static_cast<uint64_t>(in.w[1]) << 32u);
-	xmm->High = static_cast<uint64_t>(in.w[2]) | (static_cast<uint64_t>(in.w[3]) << 32u);
-}
+	void LoadGprs(uint64_t (&gpr)[16]) const {
+		const uint64_t registers[] = {native->Rax, native->Rcx, native->Rdx, native->Rbx,
+		                              native->Rsp, native->Rbp, native->Rsi, native->Rdi,
+		                              native->R8,  native->R9,  native->R10, native->R11,
+		                              native->R12, native->R13, native->R14, native->R15};
+		std::memcpy(gpr, registers, sizeof(gpr));
+	}
 
+	void ClearUpperYmm(uint8_t index) const {
+		if ((native->ContextFlags & CONTEXT_XSTATE) != CONTEXT_XSTATE) {
+			return;
+		}
+		DWORD64 features = 0;
+		if (!GetXStateFeaturesMask(native, &features) || (features & XSTATE_MASK_AVX) == 0) {
+			return; // An absent AVX component restores zeroes.
+		}
+		DWORD size = 0;
+		auto* ymm = static_cast<M128A*>(LocateXStateFeature(native, XSTATE_AVX, &size));
+		if (ymm != nullptr && size >= (index + 1u) * sizeof(M128A)) {
+			ymm[index] = {};
+		}
+	}
+#else
+	ucontext_t* native;
+
+	[[nodiscard]] uint64_t Rip() const {
+		return static_cast<uint64_t>(native->uc_mcontext.gregs[REG_RIP]);
+	}
+	void Advance(size_t length) {
+		native->uc_mcontext.gregs[REG_RIP] += static_cast<greg_t>(length);
+	}
+	[[nodiscard]] void* Xmm(uint8_t index) const {
+		if (native->uc_mcontext.fpregs == nullptr) {
+			return nullptr;
+		}
+		return native->uc_mcontext.fpregs->_xmm[index].element;
+	}
+
+	void LoadGprs(uint64_t (&gpr)[16]) const {
+		constexpr int registers[] = {REG_RAX, REG_RCX, REG_RDX, REG_RBX, REG_RSP, REG_RBP,
+		                             REG_RSI, REG_RDI, REG_R8,  REG_R9,  REG_R10, REG_R11,
+		                             REG_R12, REG_R13, REG_R14, REG_R15};
+		for (size_t i = 0; i < 16; ++i) {
+			gpr[i] = static_cast<uint64_t>(native->uc_mcontext.gregs[registers[i]]);
+		}
+	}
+
+	void ClearUpperYmm(uint8_t index) const {
+		// Linux signal frames use the standard XSAVE layout. An absent AVX component
+		// already restores the architectural initial value (all zeroes).
+		auto*    state    = reinterpret_cast<uint8_t*>(native->uc_mcontext.fpregs);
+		uint32_t magic    = 0;
+		uint32_t size     = 0;
+		uint64_t features = 0;
+		std::memcpy(&magic, state + 464, sizeof(magic));
+		if (magic != 0x46505853) {
+			return;
+		}
+		std::memcpy(&size, state + 480, sizeof(size));
+		if (size < 832) {
+			return;
+		}
+		std::memcpy(&features, state + 512, sizeof(features));
+		if ((features & 4) != 0) {
+			std::memset(state + 576 + index * 16, 0, 16);
+		}
+	}
 #endif
+};
 
-#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
-
-static M128A* GetContextXmm(PCONTEXT context, uint8_t index) {
-	if (context == nullptr || index >= 16) {
-		return nullptr;
-	}
-
-	return &context->Xmm0 + index;
-}
-
-static void LoadContextGprsWin(PCONTEXT context, uint64_t (&gpr)[16]) {
-	gpr[0]  = context->Rax;
-	gpr[1]  = context->Rcx;
-	gpr[2]  = context->Rdx;
-	gpr[3]  = context->Rbx;
-	gpr[4]  = context->Rsp;
-	gpr[5]  = context->Rbp;
-	gpr[6]  = context->Rsi;
-	gpr[7]  = context->Rdi;
-	gpr[8]  = context->R8;
-	gpr[9]  = context->R9;
-	gpr[10] = context->R10;
-	gpr[11] = context->R11;
-	gpr[12] = context->R12;
-	gpr[13] = context->R13;
-	gpr[14] = context->R14;
-	gpr[15] = context->R15;
-}
-
-static bool TryEmulateShaNi(PCONTEXT context) {
-	if (context == nullptr) {
-		return false;
-	}
-
-	const auto* rip = reinterpret_cast<const uint8_t*>(context->Rip);
+static bool TryEmulateShaNi(Context& context) {
+	const auto* rip = reinterpret_cast<const uint8_t*>(context.Rip());
 	ShaNiInsn   insn {};
 	if (!DecodeShaNiInsn(rip, insn)) {
 		return false;
@@ -458,8 +496,8 @@ static bool TryEmulateShaNi(PCONTEXT context) {
 
 	const uint8_t modrm_byte = rip[insn.modrm_offset];
 	const uint8_t dest_index = ShaNiRegIndex(modrm_byte, insn.rex, true);
-	auto*         dest_xmm   = GetContextXmm(context, dest_index);
-	auto*         xmm0       = GetContextXmm(context, 0);
+	auto*         dest_xmm   = context.Xmm(dest_index);
+	auto*         xmm0       = context.Xmm(0);
 	if (dest_xmm == nullptr || xmm0 == nullptr) {
 		return false;
 	}
@@ -467,20 +505,20 @@ static bool TryEmulateShaNi(PCONTEXT context) {
 	XmmWords dest {};
 	XmmWords src2 {};
 	XmmWords xmm0_words {};
-	LoadXmmWordsWin(dest_xmm, dest);
-	LoadXmmWordsWin(xmm0, xmm0_words);
+	std::memcpy(&dest, dest_xmm, sizeof(dest));
+	std::memcpy(&xmm0_words, xmm0, sizeof(xmm0_words));
 
 	if (ShaNiModrmIsRegister(modrm_byte)) {
 		const uint8_t src_index = ShaNiRegIndex(modrm_byte, insn.rex, false);
-		auto*         src_xmm   = GetContextXmm(context, src_index);
+		auto*         src_xmm   = context.Xmm(src_index);
 		if (src_xmm == nullptr) {
 			return false;
 		}
-		LoadXmmWordsWin(src_xmm, src2);
+		std::memcpy(&src2, src_xmm, sizeof(src2));
 	} else {
 		uint64_t    gpr[16] {};
 		const void* source = nullptr;
-		LoadContextGprsWin(context, gpr);
+		context.LoadGprs(gpr);
 		if (!ResolveShaNiMemoryAddress(rip, insn, gpr, source)) {
 			return false;
 		}
@@ -491,18 +529,13 @@ static bool TryEmulateShaNi(PCONTEXT context) {
 		return false;
 	}
 
-	StoreXmmWordsWin(dest_xmm, dest);
-	context->Rip += insn.length;
+	std::memcpy(dest_xmm, &dest, sizeof(dest));
+	context.Advance(insn.length);
 	return true;
 }
 
-static bool TryEmulateSse4a(PCONTEXT context) {
-	if (context == nullptr) {
-		return false;
-	}
-
-	const auto* rip = reinterpret_cast<const uint8_t*>(context->Rip);
-
+static bool TryEmulateSse4a(Context& context) {
+	const auto*   rip    = reinterpret_cast<const uint8_t*>(context.Rip());
 	const uint8_t prefix = rip[0];
 	if (prefix != 0x66 && prefix != 0xf2) {
 		return false;
@@ -511,210 +544,13 @@ static bool TryEmulateSse4a(PCONTEXT context) {
 	size_t  offset = 1;
 	uint8_t rex    = 0;
 	if ((rip[offset] & 0xf0u) == 0x40u) {
-		rex = rip[offset];
-		offset++;
+		rex = rip[offset++];
 	}
-
 	if (rip[offset] != 0x0f || rip[offset + 1] != 0x78) {
 		return false;
 	}
 
-	auto modrm = rip[offset + 2];
-	if ((modrm & 0xc0u) != 0xc0u) {
-		return false;
-	}
-
-	const uint8_t reg    = ((modrm >> 3u) & 0x07u) | ((rex & 0x04u) << 1u);
-	const uint8_t rm     = (modrm & 0x07u) | ((rex & 0x01u) << 3u);
-	const uint8_t length = rip[offset + 3];
-	const uint8_t index  = rip[offset + 4];
-
-	// AMD SSE4a immediate-form EXTRQ/INSERTQ. PS5 code can execute these natively on AMD hardware,
-	// while Intel hosts raise an illegal-instruction exception.
-	if (prefix == 0x66) {
-		auto* dst = GetContextXmm(context, rm);
-		if (dst == nullptr) {
-			return false;
-		}
-
-		dst->Low  = ExtractBitField(dst->Low, length, index);
-		dst->High = 0;
-		context->Rip += offset + 5;
-		return true;
-	}
-
-	auto* dst = GetContextXmm(context, reg);
-	auto* src = GetContextXmm(context, rm);
-	if (dst == nullptr || src == nullptr) {
-		return false;
-	}
-
-	dst->Low = InsertBitField(dst->Low, src->Low, length, index);
-	context->Rip += offset + 5;
-	return true;
-}
-
-static bool TryEmulateMonitorxMwaitx(PCONTEXT context) {
-	if (context == nullptr) {
-		return false;
-	}
-
-	const auto* rip = reinterpret_cast<const uint8_t*>(context->Rip);
-	if (rip[0] != 0x0f || rip[1] != 0x01 || (rip[2] != 0xfa && rip[2] != 0xfb)) {
-		return false;
-	}
-
-	// AMD MONITORX/MWAITX are used by PS5 code in wait loops. Intel hosts can raise an illegal-
-	// instruction exception, so approximate them as a no-op/yield pair.
-	if (rip[2] == 0xfb) {
-		SwitchToThread();
-	}
-	context->Rip += 3;
-	return true;
-}
-
-#elif !defined(__APPLE__)
-
-// Linux signal contexts expose registers through ucontext_t.
-
-static uint32_t* GetContextXmm(ucontext_t* context, uint8_t index) {
-	if (context == nullptr || index >= 16) {
-		return nullptr;
-	}
-
-	auto* fpregs = context->uc_mcontext.fpregs;
-	if (fpregs == nullptr) {
-		return nullptr;
-	}
-
-	return static_cast<uint32_t*>(fpregs->_xmm[index].element);
-}
-
-static void LoadContextGprsLin(ucontext_t* context, uint64_t (&gpr)[16]) {
-	gpr[0]  = static_cast<uint64_t>(context->uc_mcontext.gregs[REG_RAX]);
-	gpr[1]  = static_cast<uint64_t>(context->uc_mcontext.gregs[REG_RCX]);
-	gpr[2]  = static_cast<uint64_t>(context->uc_mcontext.gregs[REG_RDX]);
-	gpr[3]  = static_cast<uint64_t>(context->uc_mcontext.gregs[REG_RBX]);
-	gpr[4]  = static_cast<uint64_t>(context->uc_mcontext.gregs[REG_RSP]);
-	gpr[5]  = static_cast<uint64_t>(context->uc_mcontext.gregs[REG_RBP]);
-	gpr[6]  = static_cast<uint64_t>(context->uc_mcontext.gregs[REG_RSI]);
-	gpr[7]  = static_cast<uint64_t>(context->uc_mcontext.gregs[REG_RDI]);
-	gpr[8]  = static_cast<uint64_t>(context->uc_mcontext.gregs[REG_R8]);
-	gpr[9]  = static_cast<uint64_t>(context->uc_mcontext.gregs[REG_R9]);
-	gpr[10] = static_cast<uint64_t>(context->uc_mcontext.gregs[REG_R10]);
-	gpr[11] = static_cast<uint64_t>(context->uc_mcontext.gregs[REG_R11]);
-	gpr[12] = static_cast<uint64_t>(context->uc_mcontext.gregs[REG_R12]);
-	gpr[13] = static_cast<uint64_t>(context->uc_mcontext.gregs[REG_R13]);
-	gpr[14] = static_cast<uint64_t>(context->uc_mcontext.gregs[REG_R14]);
-	gpr[15] = static_cast<uint64_t>(context->uc_mcontext.gregs[REG_R15]);
-}
-
-static void LoadXmmWordsLin(const uint32_t* xmm, XmmWords& out) {
-	out.w[0] = xmm[0];
-	out.w[1] = xmm[1];
-	out.w[2] = xmm[2];
-	out.w[3] = xmm[3];
-}
-
-static void StoreXmmWordsLin(uint32_t* xmm, const XmmWords& in) {
-	xmm[0] = in.w[0];
-	xmm[1] = in.w[1];
-	xmm[2] = in.w[2];
-	xmm[3] = in.w[3];
-}
-
-static bool TryEmulateShaNi(ucontext_t* context) {
-	if (context == nullptr) {
-		return false;
-	}
-
-	auto&       rip_reg = context->uc_mcontext.gregs[REG_RIP];
-	const auto* rip     = reinterpret_cast<const uint8_t*>(rip_reg);
-	ShaNiInsn   insn {};
-	if (!DecodeShaNiInsn(rip, insn)) {
-		return false;
-	}
-
-	const uint8_t modrm_byte = rip[insn.modrm_offset];
-	const uint8_t dest_index = ShaNiRegIndex(modrm_byte, insn.rex, true);
-	auto*         dest_xmm   = GetContextXmm(context, dest_index);
-	auto*         xmm0       = GetContextXmm(context, 0);
-	if (dest_xmm == nullptr || xmm0 == nullptr) {
-		return false;
-	}
-
-	XmmWords dest {};
-	XmmWords src2 {};
-	XmmWords xmm0_words {};
-	LoadXmmWordsLin(dest_xmm, dest);
-	LoadXmmWordsLin(xmm0, xmm0_words);
-
-	if (ShaNiModrmIsRegister(modrm_byte)) {
-		const uint8_t src_index = ShaNiRegIndex(modrm_byte, insn.rex, false);
-		auto*         src_xmm   = GetContextXmm(context, src_index);
-		if (src_xmm == nullptr) {
-			return false;
-		}
-		LoadXmmWordsLin(src_xmm, src2);
-	} else {
-		uint64_t    gpr[16] {};
-		const void* source = nullptr;
-		LoadContextGprsLin(context, gpr);
-		if (!ResolveShaNiMemoryAddress(rip, insn, gpr, source)) {
-			return false;
-		}
-		std::memcpy(&src2, source, sizeof(src2));
-	}
-
-	if (!ExecuteShaNiInsn(insn, src2, xmm0_words, dest)) {
-		return false;
-	}
-
-	StoreXmmWordsLin(dest_xmm, dest);
-	rip_reg += static_cast<greg_t>(insn.length);
-	return true;
-}
-
-static uint64_t GetXmmLow(const uint32_t* xmm) {
-	return static_cast<uint64_t>(xmm[0]) | (static_cast<uint64_t>(xmm[1]) << 32u);
-}
-
-static void SetXmmLow(uint32_t* xmm, uint64_t value) {
-	xmm[0] = static_cast<uint32_t>(value);
-	xmm[1] = static_cast<uint32_t>(value >> 32u);
-}
-
-static void SetXmmHigh(uint32_t* xmm, uint64_t value) {
-	xmm[2] = static_cast<uint32_t>(value);
-	xmm[3] = static_cast<uint32_t>(value >> 32u);
-}
-
-static bool TryEmulateSse4a(ucontext_t* context) {
-	if (context == nullptr) {
-		return false;
-	}
-
-	auto& rip_reg = context->uc_mcontext.gregs[REG_RIP];
-
-	const auto* rip = reinterpret_cast<const uint8_t*>(rip_reg);
-
-	const uint8_t prefix = rip[0];
-	if (prefix != 0x66 && prefix != 0xf2) {
-		return false;
-	}
-
-	size_t  offset = 1;
-	uint8_t rex    = 0;
-	if ((rip[offset] & 0xf0u) == 0x40u) {
-		rex = rip[offset];
-		offset++;
-	}
-
-	if (rip[offset] != 0x0f || rip[offset + 1] != 0x78) {
-		return false;
-	}
-
-	auto modrm = rip[offset + 2];
+	const uint8_t modrm = rip[offset + 2];
 	if ((modrm & 0xc0u) != 0xc0u) {
 		return false;
 	}
@@ -725,58 +561,170 @@ static bool TryEmulateSse4a(ucontext_t* context) {
 	const uint8_t index  = rip[offset + 4];
 
 	// AMD SSE4a immediate-form EXTRQ/INSERTQ.
+	uint8_t dest_index = reg;
 	if (prefix == 0x66) {
-		auto* dst = GetContextXmm(context, rm);
-		if (dst == nullptr) {
-			return false;
-		}
-
-		SetXmmLow(dst, ExtractBitField(GetXmmLow(dst), length, index));
-		SetXmmHigh(dst, 0);
-		rip_reg += static_cast<greg_t>(offset + 5);
-		return true;
+		dest_index = rm;
 	}
-
-	auto* dst = GetContextXmm(context, reg);
-	auto* src = GetContextXmm(context, rm);
-	if (dst == nullptr || src == nullptr) {
+	auto* dest_xmm = context.Xmm(dest_index);
+	auto* src_xmm  = context.Xmm(rm);
+	if (dest_xmm == nullptr || src_xmm == nullptr) {
 		return false;
 	}
-
-	SetXmmLow(dst, InsertBitField(GetXmmLow(dst), GetXmmLow(src), length, index));
-	rip_reg += static_cast<greg_t>(offset + 5);
+	uint64_t dest[2] {};
+	uint64_t source = 0;
+	std::memcpy(dest, dest_xmm, sizeof(dest));
+	std::memcpy(&source, src_xmm, sizeof(source));
+	if (prefix == 0x66) {
+		dest[0] = ExtractBitField(dest[0], length, index);
+		dest[1] = 0;
+	} else {
+		dest[0] = InsertBitField(dest[0], source, length, index);
+	}
+	std::memcpy(dest_xmm, dest, sizeof(dest));
+	context.Advance(offset + 5);
 	return true;
 }
 
-static bool TryEmulateMonitorxMwaitx(ucontext_t* context) {
-	if (context == nullptr) {
-		return false;
-	}
-
-	auto& rip_reg = context->uc_mcontext.gregs[REG_RIP];
-
-	const auto* rip = reinterpret_cast<const uint8_t*>(rip_reg);
+static bool TryEmulateMonitorxMwaitx(Context& context) {
+	const auto* rip = reinterpret_cast<const uint8_t*>(context.Rip());
 	if (rip[0] != 0x0f || rip[1] != 0x01 || (rip[2] != 0xfa && rip[2] != 0xfb)) {
 		return false;
 	}
 
 	// Approximate AMD MONITORX/MWAITX as no-op/yield.
 	if (rip[2] == 0xfb) {
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+		SwitchToThread();
+#else
 		::sched_yield();
+#endif
 	}
-	rip_reg += 3;
+	context.Advance(3);
+	return true;
+}
+
+static uint32_t ReciprocalSquareRoot(uint32_t bits) {
+	const uint32_t magnitude = bits & 0x7fffffffu;
+	const uint32_t exponent  = magnitude & 0x7f800000u;
+	if (exponent == 0) {
+		// RSQRT treats denormals as signed zero regardless of MXCSR.DAZ.
+		return (bits & 0x80000000u) | 0x7f800000u;
+	}
+	if (magnitude > 0x7f800000u) {
+		return bits | 0x00400000u; // Quiet NaNs without raising an exception.
+	}
+	if ((bits & 0x80000000u) != 0) {
+		return 0xffc00000u;
+	}
+	if (magnitude == 0x7f800000u) {
+		return 0;
+	}
+
+	// A deterministic accurate estimate meets the instruction's relative-error
+	// bound without relying on the host vendor's approximation table.
+	const __m128d input  = _mm_set_sd(static_cast<double>(std::bit_cast<float>(bits)));
+	const __m128d result = _mm_div_sd(_mm_set_sd(1.0), _mm_sqrt_sd(input, input));
+	return std::bit_cast<uint32_t>(_mm_cvtss_f32(_mm_cvtsd_ss(_mm_setzero_ps(), result)));
+}
+
+static bool TryEmulateReciprocalSquareRoot(Context& context) {
+	const auto* rip            = reinterpret_cast<const uint8_t*>(context.Rip());
+	size_t      prefix_size    = 0;
+	uint8_t     dest_extension = 0;
+	uint8_t     src_extension  = 0;
+	if (rip[0] == 0xc5 && (rip[1] & 0x7fu) == 0x70u) {
+		prefix_size    = 2;
+		dest_extension = (~rip[1] & 0x80u) >> 4u;
+	} else if (rip[0] == 0xc4 && (rip[1] & 0x1fu) == 1 && (rip[2] & 0x7fu) == 0x70u) {
+		prefix_size    = 3;
+		dest_extension = (~rip[1] & 0x80u) >> 4u;
+		src_extension  = (~rip[1] & 0x20u) >> 2u;
+	} else {
+		return false;
+	}
+	if (rip[prefix_size] != 0x52 || (rip[prefix_size + 1] & 0xc0u) != 0xc0u) {
+		return false;
+	}
+
+	const uint8_t modrm    = rip[prefix_size + 1];
+	const uint8_t dest     = ((modrm >> 3u) & 7u) | dest_extension;
+	const uint8_t source   = (modrm & 7u) | src_extension;
+	auto*         dest_xmm = context.Xmm(dest);
+	auto*         src_xmm  = context.Xmm(source);
+	if (dest_xmm == nullptr || src_xmm == nullptr) {
+		return false;
+	}
+	XmmWords result {};
+	std::memcpy(&result, src_xmm, sizeof(result));
+	// RSQRT ignores the rounding mode and never changes guest exception flags.
+	// Mask host exceptions while calculating, then restore the handler's state.
+	const uint32_t mxcsr = _mm_getcsr();
+	_mm_setcsr(0x1f80);
+	for (auto& word: result.w) {
+		word = ReciprocalSquareRoot(word);
+	}
+	_mm_setcsr(mxcsr);
+	std::memcpy(dest_xmm, &result, sizeof(result));
+	context.ClearUpperYmm(dest);
+	context.Advance(prefix_size + 2);
 	return true;
 }
 
 #endif
 
+bool IsReciprocalSquareRoot(const ZydisDecodedInstruction& instruction,
+                            const ZydisDecodedOperand* operands) {
+	return instruction.mnemonic == ZYDIS_MNEMONIC_VRSQRTPS &&
+	       instruction.encoding == ZYDIS_INSTRUCTION_ENCODING_VEX &&
+	       instruction.raw.vex.offset == 0 && operands[0].size == 128 &&
+	       operands[1].type == ZYDIS_OPERAND_TYPE_REGISTER;
+}
+
+uint64_t PatchReciprocalSquareRoots(uint64_t address, uint64_t size) {
+	uint64_t patched = 0;
+#if !defined(__APPLE__)
+	ZydisDecoder decoder {};
+	if (!ZYAN_SUCCESS(
+	        ZydisDecoderInit(&decoder, ZYDIS_MACHINE_MODE_LONG_64, ZYDIS_STACK_WIDTH_64))) {
+		return 0;
+	}
+	for (uint64_t offset = 0; offset < size;) {
+		auto*                   code = reinterpret_cast<uint8_t*>(address + offset);
+		ZydisDecodedInstruction instruction {};
+		ZydisDecodedOperand     operands[ZYDIS_MAX_OPERAND_COUNT] {};
+		if (!ZYAN_SUCCESS(
+		        ZydisDecoderDecodeFull(&decoder, code, size - offset, &instruction, operands))) {
+			++offset;
+			continue;
+		}
+		if (IsReciprocalSquareRoot(instruction, operands)) {
+			// vvvv is reserved (must be 1111b). Clear one bit to route this
+			// otherwise intact instruction through the illegal-instruction emulator.
+			code[instruction.raw.vex.size - 1] &= ~0x08u;
+			++patched;
+		}
+		offset += instruction.length;
+	}
+#else
+	(void)address;
+	(void)size;
+#endif
+	return patched;
+}
+
 bool TryEmulate(void* native_context) {
+#if !defined(__APPLE__)
+	if (native_context == nullptr) {
+		return false;
+	}
 #if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
-	auto* context = static_cast<PCONTEXT>(native_context);
-	return TryEmulateMonitorxMwaitx(context) || TryEmulateSse4a(context) ||
-	       TryEmulateShaNi(context);
-#elif !defined(__APPLE__)
-	auto* context = static_cast<ucontext_t*>(native_context);
+	Context context {static_cast<PCONTEXT>(native_context)};
+#else
+	Context context {static_cast<ucontext_t*>(native_context)};
+#endif
+	if (TryEmulateReciprocalSquareRoot(context)) {
+		return true;
+	}
 	return TryEmulateMonitorxMwaitx(context) || TryEmulateSse4a(context) ||
 	       TryEmulateShaNi(context);
 #else

src/loader/x64InstructionEmulator.h

@@ -1,8 +1,14 @@
 #ifndef KYTY_LOADER_X64_INSTRUCTION_EMULATOR_H_
 #define KYTY_LOADER_X64_INSTRUCTION_EMULATOR_H_
 
+#include <Zydis/DecoderTypes.h>
+#include <cstdint>
+
 namespace Loader::X64InstructionEmulator {
 
+[[nodiscard]] bool IsReciprocalSquareRoot(const ZydisDecodedInstruction& instruction,
+                                         const ZydisDecodedOperand* operands);
+uint64_t           PatchReciprocalSquareRoots(uint64_t address, uint64_t size);
 [[nodiscard]] bool TryEmulate(void* native_context);
 
 } // namespace Loader::X64InstructionEmulator

src/main.cpp

@@ -56,6 +56,7 @@ static void PrintUsage() {
 	    "  --gpu <index>                        Vulkan physical device index. Default: auto.\n");
 	::printf("  --fullscreen                         Run in borderless desktop fullscreen.\n");
 	::printf("  --vr                                 Enable the virtual VR headset.\n");
+	::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");
@@ -168,6 +169,11 @@ static bool ParseArgs(int argc, char* argv[], RunOptions& options, bool& show_he
 			continue;
 		}
 
+		if (arg == "--amd-cpu") {
+			options.config.amd_cpu_enabled = true;
+			continue;
+		}
+
 		if (arg == "--playgo-hack") {
 			options.config.playgo_hack_enabled = true;
 			continue;

tests/ShaderRecompilerComputeTests.cpp

@@ -1219,6 +1219,8 @@ struct GraphicsCase {
   bool pixel_position_w = false;
   float vertex_clip_w = 1.0f;
   bool pixel_depth_export = false;
+  u32 pixel_perspective_centroid_vgpr = UINT32_MAX;
+  u32 pixel_custom_interpolation_mask = 0;
 };
 
 struct CompiledShader {
@@ -1274,6 +1276,43 @@ size_t CountText(const std::string &text, const std::string &needle) {
   return count;
 }
 
+void CheckPixelParameterAliases() {
+  constexpr const char *name = "PixelParameterAliases";
+  ShaderPixelInputInfo pixel{};
+  pixel.input_num = 2;
+  const std::array<uint32_t, 2> pair = {0, 1};
+  for (const auto flat : {0u, 0x400u}) {
+    pixel.interpolator_settings[0] = flat | 3u;
+    pixel.interpolator_settings[1] = flat | 3u;
+    Require(name, "same-mode nonidentity alias",
+            ShaderPixelParameterLocation(pixel, pair, 0) == 3u &&
+                ShaderPixelParameterLocation(pixel, pair, 1) == 3u,
+            "aliases of one export with the same interpolation mode must "
+            "share that physical location");
+  }
+  pixel.interpolator_settings[0] = 3u;
+  pixel.interpolator_settings[1] = 0x423u;
+  pixel.custom_interpolation_mask = 2u;
+  Require(name, "smooth/custom alias",
+          ShaderPixelParameterLocation(pixel, pair, 0) == 3u &&
+              ShaderPixelParameterLocation(pixel, pair, 1) == 3u,
+          "custom interpolation must retain its shared vertex export");
+
+  pixel.input_num = 3;
+  pixel.custom_interpolation_mask = 0;
+  pixel.interpolator_settings[0] = 0x400u;
+  pixel.interpolator_settings[1] = 0u;
+  pixel.interpolator_settings[2] = 1u;
+  const std::array<uint32_t, 3> active = {0, 1, 2};
+  Require(name, "reserved physical locations",
+          ShaderPixelParameterLocation(pixel, active, 0) == 0u &&
+              ShaderPixelParameterLocation(pixel, active, 1) == 2u &&
+              ShaderPixelParameterLocation(pixel, active, 2) == 1u,
+          "a synthesized opposite-mode alias must not occupy another "
+          "vertex export's physical location");
+  std::printf("[host]    %-32s ok\n", name);
+}
+
 void CheckRectListShaders() {
   constexpr const char *name = "RectListShaders";
 
@@ -1355,6 +1394,25 @@ void CheckRectListShaders() {
           "duplicate pixel mappings must share one vertex input and keep "
           "distinct patch outputs");
 
+  for (const auto flat : {0u, 0x400u}) {
+    pixel.interpolator_settings[0] = flat;
+    pixel.interpolator_settings[1] = flat;
+    const auto shared = BuildRectListShaders(vertex, &pixel);
+    ValidateSpirv(name, shared.control);
+    ValidateSpirv(name, shared.evaluation);
+    Require(name, "same-mode alias disassembly",
+            tools.Disassemble(shared.control, &control_text) &&
+                tools.Disassemble(shared.evaluation, &evaluation_text),
+            "failed to disassemble shared rectangle-list parameters");
+    Require(name, "same-mode interface deduplication",
+            CountText(control_text, " Location 0") == 2 &&
+                CountText(evaluation_text, " Location 0") == 2 &&
+                control_text.find(" Location 1") == std::string::npos &&
+                evaluation_text.find(" Location 1") == std::string::npos,
+            "same-mode aliases must produce one input and one output at "
+            "the shared location in each rectangle-list stage");
+  }
+
   const auto position_only = BuildRectListShaders(vertex, nullptr);
   ValidateSpirv(name, position_only.control);
   ValidateSpirv(name, position_only.evaluation);
@@ -1565,6 +1623,8 @@ CompiledShader CompileFragmentCase(const GraphicsCase &test) {
   pixel_info.ps_pos_w = test.pixel_position_w;
   pixel_info.ps_depth_export_enable = test.pixel_depth_export;
   pixel_info.ps_system_input_base = 2;
+  pixel_info.ps_perspective_centroid_vgpr = test.pixel_perspective_centroid_vgpr;
+  pixel_info.custom_interpolation_mask = test.pixel_custom_interpolation_mask;
   for (u32 i = 0; i < std::size(pixel_info.interpolator_settings); i++) {
     pixel_info.interpolator_settings[i] = i;
   }
@@ -17352,6 +17412,45 @@ TestCase Vop1SdwaNotCapturedByte0Source() {
   return test;
 }
 
+TestCase Vop1SdwaMovByteDestinations() {
+  using O = ShaderOpcode;
+
+  std::vector<u32> code;
+  AppendBufferLoadDword(&code, 3, 30);
+  AppendVMovU32(&code, 30, 4);
+  AppendBufferLoadDword(&code, 2, 30);
+  code.push_back(EncodeVop1(0x01, 1, Vgpr(3)));
+  code.push_back(0x7e0202f9u);
+  code.push_back(0x00861280u); // Captured v_mov_b32 v1.byte2, 0, preserve.
+  AppendStoreVgpr(&code, 1, 2);
+
+  // RDNA2 table 88: select any byte, then pad, sign extend, or preserve.
+  for (u32 dst_u = 0; dst_u < 3; dst_u++) {
+    for (u32 dst_sel = 0; dst_sel < 4; dst_sel++) {
+      code.push_back(EncodeVop1(0x01, 1, Vgpr(3)));
+      code.push_back(EncodeVop1(0x01, 1, 249));
+      code.push_back(EncodeVop1Sdwa(2, dst_sel, dst_u));
+      AppendStoreVgpr(&code, 1, 3 + dst_u * 4 + dst_sel);
+    }
+  }
+  AppendEnd(&code);
+
+  TestCase test;
+  test.name = "Vop1SdwaMovByteDestinations";
+  test.code = std::move(code);
+  test.initial = {0xa1b2c3d4u, 0x12345680u};
+  test.expected = {0xa1b2c3d4u, 0x12345680u, 0xa100c3d4u,
+                   0x00000080u, 0x00008000u, 0x00800000u, 0x80000000u,
+                   0xffffff80u, 0xffff8000u, 0xff800000u, 0x80000000u,
+                   0xa1b2c380u, 0xa1b280d4u, 0xa180c3d4u, 0x80b2c3d4u};
+  test.initial.resize(test.expected.size());
+  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"};
+  return test;
+}
+
 TestCase Vop2SdwaSubNcExactByte2Destination() {
   using O = ShaderOpcode;
 
@@ -25514,6 +25613,91 @@ GraphicsCase GraphicsPositionWExport() {
   return test;
 }
 
+GraphicsCase GraphicsPackedHalfCentroid() {
+  GraphicsCase test;
+  test.name = "GraphicsPackedHalfCentroid";
+  test.pixel_perspective_centroid_vgpr = 0;
+  test.pixel_custom_interpolation_mask = 1;
+  test.pixel_interpolator_settings = {0x420u};
+  // Captured ab810715011baef3 interpolation, with attr3.x remapped to attr0.x.
+  test.fragment_code = {
+      0xc8120002u, 0xc80e0000u,             // raw vertex0/vertex1 packed halves
+      0xcc204007u, 0x9c1206f2u,             // low(vertex1) - low(vertex0)
+      0xcc207005u, 0x9c1206f2u,             // high(vertex1) - high(vertex0)
+      0xc80a0001u,                         // raw vertex2
+      0xcc204006u, 0x9c1204f2u,
+      0xcc207003u, 0x9c1204f2u,
+      0xcc20400eu, 0x04120f00u,             // low = I * delta10 + vertex0
+      0xcc20600fu, 0x04120b00u,             // high = I * delta10 + vertex0
+      0x3e1c0306u, 0x3e1e0303u,             // += J * delta20
+      EncodeExp0(0x00, 0xf), EncodeExp1(14, 15, 14, 15)};
+  AppendEnd(&test.fragment_code);
+  // At the pixel center the vertex weights are (5/8, 1/8, 1/4).
+  // Packed values are (0,1), (1,2), (3,3): expected result (0.875, 1.625).
+  test.vertices = {
+      0xbf800000u, 0xbf800000u, 0x3c000000u, 0, 0, 0,
+      0x40e00000u, 0xbf800000u, 0x40003c00u, 0, 0, 0,
+      0xbf800000u, 0x40400000u, 0x42004200u, 0, 0, 0};
+  test.expected_pixel = {0x3f600000u, 0x3fd00000u, 0x3f600000u, 0x3fd00000u};
+  test.opcodes = {ShaderOpcode::V_INTERP_MOV_F32, ShaderOpcode::V_FMA_F32, ShaderOpcode::V_MAC_F32,
+                  ShaderOpcode::EXP, ShaderOpcode::S_ENDPGM};
+  return test;
+}
+
+GraphicsCase GraphicsPackedHalfInputAlias(bool second_weights) {
+  auto test = GraphicsPackedHalfCentroid();
+  test.name = second_weights ? "GraphicsPackedHalfAliasSecondWeights"
+                            : "GraphicsPackedHalfInputAlias";
+  test.pixel_custom_interpolation_mask = 2u;
+  test.pixel_interpolator_settings = {0u, 0x420u};
+  // The head shader reads one exported vec4 through two logical inputs:
+  // ordinary interpolation and raw per-vertex packed-half interpolation.
+  // Move the captured raw attr0.x loads to logical attr1.x, still mapped to 0.
+  for (const auto word : {0u, 1u, 6u}) {
+    test.fragment_code[word] |= 1u << 10u;
+  }
+  test.fragment_code.resize(test.fragment_code.size() - 3u);
+  test.fragment_code.push_back(EncodeVintrp(0x00, 16, 0, 1, 0));
+  test.fragment_code.push_back(EncodeVintrp(0x01, 16, 0, 1, 1));
+  AppendVMovLiteral(&test.fragment_code, 17, 0x3f800000u);
+  test.fragment_code.push_back(EncodeExp0(0x00, 0xf));
+  test.fragment_code.push_back(EncodeExp1(14, 15, 16, 17));
+  AppendEnd(&test.fragment_code);
+  // The smooth component (2,4,8) must vary independently of packed UVs.
+  test.vertices[3] = 0x40000000u;
+  test.vertices[9] = 0x40800000u;
+  test.vertices[15] = 0x41000000u;
+  if (second_weights) {
+    // Change the right vertex from x=7 to x=3: weights become (1/2,1/4,1/4).
+    test.vertices[6] = 0x40400000u;
+    test.expected_pixel = {0x3f800000u, 0x3fe00000u, 0x40800000u, 0x3f800000u};
+  } else {
+    test.expected_pixel = {0x3f600000u, 0x3fd00000u, 0x40700000u, 0x3f800000u};
+  }
+  test.opcodes.insert(test.opcodes.end(),
+                      {ShaderOpcode::V_INTERP_P1_F32, ShaderOpcode::V_INTERP_P2_F32,
+                       ShaderOpcode::V_MOV_B32});
+  return test;
+}
+
+GraphicsCase GraphicsSmoothRawInputAlias() {
+  auto test = GraphicsPackedHalfInputAlias(false);
+  test.name = "GraphicsSmoothRawInputAlias";
+  test.pixel_perspective_centroid_vgpr = UINT32_MAX;
+  // No explicit barycentric builtin: alias promotion must discover that the
+  // ordinary input needs weights while the raw alias returns vertex zero.
+  test.fragment_code = {EncodeVintrp(0x00, 0, 0, 1, 0),
+                        EncodeVintrp(0x01, 0, 0, 1, 1),
+                        EncodeVintrp(0x02, 1, 1, 1, 2),
+                        EncodeExp0(0x00, 0xf), EncodeExp1(0, 1, 0, 1)};
+  AppendEnd(&test.fragment_code);
+  test.expected_pixel = {0x40700000u, 0x40000000u, 0x40700000u, 0x40000000u};
+  test.opcodes = {ShaderOpcode::V_INTERP_P1_F32, ShaderOpcode::V_INTERP_P2_F32,
+                  ShaderOpcode::V_INTERP_MOV_F32, ShaderOpcode::EXP,
+                  ShaderOpcode::S_ENDPGM};
+  return test;
+}
+
 GraphicsCase GraphicsAncillaryLayer(bool front_face) {
   GraphicsCase test;
   test.name = front_face ? "GraphicsAncillaryAfterFrontFace" : "GraphicsAncillaryLayer";
@@ -25891,6 +26075,7 @@ std::vector<TestCase> MakeCases() {
   AddCase(VectorFfbhI32NativeAndVop3OnGpu);
   AddCase(Vop1SdwaFfblCapturedHighWordSource);
   AddCase(Vop1SdwaNotCapturedByte0Source);
+  AddCase(Vop1SdwaMovByteDestinations);
   AddCase(Vop2SdwaSubNcExactByte2Destination);
   AddCase(Vop2SdwaAddNcCapturedHighWordDestination);
   AddCase(Vop2SdwaAshrrevCapturedWord0SignExtends);
@@ -26156,6 +26341,10 @@ std::vector<GraphicsCase> MakeGraphicsCases() {
   return {
       GraphicsInterpolationExport(),
       GraphicsPositionWExport(),
+      GraphicsPackedHalfCentroid(),
+      GraphicsPackedHalfInputAlias(false),
+      GraphicsPackedHalfInputAlias(true),
+      GraphicsSmoothRawInputAlias(),
       GraphicsAncillaryLayer(false),
       GraphicsAncillaryLayer(true),
       GraphicsAncillarySampleId(),
@@ -29374,6 +29563,106 @@ void CheckPm4SyntheticOcclusionCounterDump(RenderContext &renderer) {
   std::printf("[host]    %-32s ok\n", "Pm4SyntheticOcclusionCounterDump");
 }
 
+void CheckPm4Predication(RenderContext &renderer) {
+  constexpr const char *name = "Pm4Predication";
+  constexpr uint64_t ready = 1ull << 63u;
+  GraphicsInitJmpTables();
+  CommandProcessor processor(renderer, 0);
+  processor.BufferInit();
+  alignas(16) std::array<uint64_t, 32> query{};
+  alignas(16) uint64_t boolean = 0;
+  uint32_t predicated = 0;
+  uint32_t unconditional = 0;
+  const auto commands = [&](uint32_t op, uint32_t condition, uint32_t wait,
+                            const void *source) {
+    const auto address = reinterpret_cast<uint64_t>(source);
+    const auto predicated_address = reinterpret_cast<uint64_t>(&predicated);
+    const auto unconditional_address = reinterpret_cast<uint64_t>(&unconditional);
+    return std::array<uint32_t, 14>{
+        // Captured AGC form: c0022000 00010100 <query low> <query high>.
+        0xc0022000u, (op << 16u) | (wait << 12u) | (condition << 8u),
+        static_cast<uint32_t>(address), static_cast<uint32_t>(address >> 32u),
+        KYTY_PM4(5, Pm4::IT_WRITE_DATA, 0) | 1u, 0,
+        static_cast<uint32_t>(predicated_address),
+        static_cast<uint32_t>(predicated_address >> 32u), 11,
+        KYTY_PM4(5, Pm4::IT_WRITE_DATA, 0), 0,
+        static_cast<uint32_t>(unconditional_address),
+        static_cast<uint32_t>(unconditional_address >> 32u), 22};
+  };
+  const auto check = [&](const char *stage, uint32_t op, uint32_t condition,
+                         uint32_t wait, const void *source, bool skip) {
+    predicated = unconditional = 0;
+    const auto packet = commands(op, condition, wait, source);
+    Pm4Execution execution;
+    Require(name, stage,
+            processor.Process(execution, packet) == Pm4ProcessResult::Complete &&
+                processor.ShouldSkipPredicatedPackets() == skip &&
+                predicated == (skip ? 0u : 11u) && unconditional == 22,
+            "predicate polarity, tagged packet execution, or packet consumption is wrong");
+  };
+  struct Case {
+    uint64_t delta;
+    uint32_t condition;
+    bool skip;
+  };
+  constexpr Case cases[]{{1, 1, false}, {1, 0, true},
+                         {0, 1, true}, {0, 0, false}};
+  for (const auto &test : cases) {
+    query.fill(ready | 0x344u);
+    for (size_t db = 0; db < 16; db++) {
+      query[db * 2 + 1] += test.delta;
+    }
+    boolean = test.delta;
+    for (uint32_t wait : {0u, 1u}) {
+      check("ready Z-pass", 1, test.condition, wait, query.data(), test.skip);
+      check("bool preservation", 3, test.condition, wait, &boolean, test.skip);
+    }
+  }
+  query.fill(ready | 0x344u);
+  query[31]++;
+  check("only last DB visible", 1, 1, 0, query.data(), false);
+  check("inverse last DB visibility", 1, 0, 0, query.data(), true);
+  check("clear", 0, 0, 0, nullptr, false);
+
+  // Every begin AND end counter must be ready; no-wait must discard old skip state.
+  for (size_t missing = 0; missing < query.size(); missing++) {
+    query.fill(ready | 0x344u);
+    query[missing] &= ~ready;
+    for (uint32_t condition : {0u, 1u}) {
+      boolean = 0;
+      processor.SetPredication(1, 3, 0, &boolean, 0);
+      check("unavailable no-wait", 1, condition, 1, query.data(), false);
+    }
+  }
+
+  // Retry the same packet when its last DB's begin or end arrives later.
+  for (size_t missing : {30u, 31u}) {
+    query.fill(ready | 0x344u);
+    query[31]++;
+    query[missing] &= ~ready;
+    boolean = 0;
+    processor.SetPredication(1, 3, 0, &boolean, 0);
+    predicated = unconditional = 0;
+    const auto packet = commands(1, 1, 0, query.data());
+    Pm4Execution execution;
+    for (int retry = 0; retry < 2; retry++) {
+      Require(name, "pending wait suspends",
+              processor.Process(execution, packet) == Pm4ProcessResult::Blocked &&
+                  !execution.MadeProgress() && predicated == 0 && unconditional == 0 &&
+                  processor.ShouldSkipPredicatedPackets(),
+              "pending query advanced its packet or changed the prior predicate");
+    }
+    query[missing] |= ready;
+    Require(name, "ready wait resumes",
+            processor.Process(execution, packet) == Pm4ProcessResult::Complete &&
+                execution.MadeProgress() && predicated == 11 && unconditional == 22 &&
+                !processor.ShouldSkipPredicatedPackets(),
+            "ready query did not reevaluate and execute the suspended packet's suffix");
+  }
+  processor.BufferWait();
+  std::printf("[host]    %-32s ok\n", name);
+}
+
 void CheckPm4StencilInfoValueLane(RenderContext &renderer) {
   CommandProcessor processor(renderer, 0);
   constexpr std::array<uint32_t, 2> payload{0x00100801u, 0x28000000u};
@@ -30295,6 +30584,89 @@ void CheckPm4ContextStateOperations(RenderContext &renderer) {
   std::printf("[host]    %-32s ok\n", "Pm4ContextState");
 }
 
+void CheckPm4IndirectControlFlow(RenderContext &renderer) {
+  constexpr const char *name = "Pm4IndirectControlFlow";
+  GraphicsInitJmpTables();
+  CommandProcessor processor(renderer, 0);
+  uint32_t selected = 0;
+  uint32_t returned = 0;
+  const auto address = [](const void *value) {
+    return reinterpret_cast<uint64_t>(value);
+  };
+  const auto write = [&](uint32_t *destination, uint32_t value) {
+    return std::array<uint32_t, 5>{
+        KYTY_PM4(5, Pm4::IT_WRITE_DATA, 0), 0,
+        static_cast<uint32_t>(address(destination)),
+        static_cast<uint32_t>(address(destination) >> 32u), value};
+  };
+  const auto indirect = [&](std::span<const uint32_t> target, bool chain) {
+    uint32_t control = 0x0f200000u | static_cast<uint32_t>(target.size());
+    if (chain) {
+      control |= 1u << 20u;
+    }
+    return std::array<uint32_t, 4>{
+        KYTY_PM4(4, Pm4::IT_INDIRECT_BUFFER, 0),
+        static_cast<uint32_t>(address(target.data())),
+        static_cast<uint32_t>(address(target.data()) >> 32u), control};
+  };
+  const auto then_commands = write(&selected, 11);
+  const auto else_commands = write(&selected, 33);
+  const auto branch_suffix = write(&selected, 44);
+  const auto caller_suffix = write(&returned, 22);
+  alignas(8) uint64_t condition = 0;
+  std::array<uint32_t, 19> branch{
+      KYTY_PM4(14, Pm4::IT_INDIRECT_BUFFER, 0), 0,
+      static_cast<uint32_t>(address(&condition)),
+      static_cast<uint32_t>(address(&condition) >> 32u), UINT32_MAX,
+      UINT32_MAX, 1, 0,
+      static_cast<uint32_t>(address(then_commands.data())),
+      static_cast<uint32_t>(address(then_commands.data()) >> 32u),
+      static_cast<uint32_t>(then_commands.size()),
+      static_cast<uint32_t>(address(else_commands.data())),
+      static_cast<uint32_t>(address(else_commands.data()) >> 32u),
+      static_cast<uint32_t>(else_commands.size())};
+  std::copy(branch_suffix.begin(), branch_suffix.end(), branch.begin() + 14);
+  std::array<uint32_t, 9> caller{};
+  const auto call_branch = indirect(branch, false);
+  std::copy(call_branch.begin(), call_branch.end(), caller.begin());
+  std::copy(caller_suffix.begin(), caller_suffix.end(), caller.begin() + 4);
+  for (uint32_t mode : {1u, 2u}) {
+    branch[1] = mode | (3u << 8u);
+    for (uint64_t value : {0ull, 1ull}) {
+      condition = value;
+      selected = returned = 0;
+      uint32_t expected = 44;
+      if (value == 1) {
+        expected = 11;
+      } else if (mode == 2) {
+        expected = 33;
+      }
+      Pm4Execution execution;
+      Require(name, "conditional fetcher replacement",
+              processor.Process(execution, caller) == Pm4ProcessResult::Complete &&
+                  selected == expected && returned == 22,
+              "taken branch resumed its discarded suffix or lost the caller's return");
+    }
+  }
+
+  // This stream exceeds native recursion capacity, while chains need one fetcher cursor.
+  std::vector<std::array<uint32_t, 4>> links(65536);
+  std::span<const uint32_t> target = then_commands;
+  for (auto &link : links) {
+    link = indirect(target, true);
+    target = link;
+  }
+  const auto call_chain = indirect(target, false);
+  std::copy(call_chain.begin(), call_chain.end(), caller.begin());
+  selected = returned = 0;
+  Pm4Execution execution;
+  Require(name, "long chain and call return",
+          processor.Process(execution, caller) == Pm4ProcessResult::Complete &&
+              selected == 11 && returned == 22,
+          "long chain did not complete and return to the original caller");
+  std::printf("[host]    %-32s ok\n", name);
+}
+
 void CheckPm4WaitResume(RenderContext &renderer) {
   GraphicsInitJmpTables();
   CommandProcessor processor(renderer, 0);
@@ -30324,7 +30696,7 @@ void CheckPm4WaitResume(RenderContext &renderer) {
   nested[0] = KYTY_PM4(4, Pm4::IT_INDIRECT_BUFFER, 0);
   nested[1] = static_cast<uint32_t>(address(child.data()));
   nested[2] = static_cast<uint32_t>(address(child.data()) >> 32u);
-  nested[3] = 0x0f200000u | static_cast<uint32_t>(child.size());
+  nested[3] = 0x0f300000u | static_cast<uint32_t>(child.size());
 
   std::array<uint32_t, 14> commands{};
   commands[0] = KYTY_PM4(5, Pm4::IT_WRITE_DATA, 0);
@@ -30347,6 +30719,10 @@ void CheckPm4WaitResume(RenderContext &renderer) {
           processor.Process(execution, commands) == Pm4ProcessResult::Blocked &&
               prefix == 11 && child_observation == 0 && suffix == 0,
           "blocked indirect wait did not preserve its command position");
+  Require("Pm4WaitResume", "still blocked",
+          processor.Process(execution, commands) == Pm4ProcessResult::Blocked &&
+              !execution.MadeProgress() && suffix == 0,
+          "blocked wait advanced or replayed its caller");
 
   label = 1;
   child[4] = 1;
@@ -30538,6 +30914,21 @@ int main(int argc, char **argv) {
     RunGraphicsCase(&vulkan, GraphicsPositionWExport());
     return 0;
   }
+  if (argc == 2 && std::strcmp(argv[1], "--centroid-only") == 0) {
+    VulkanHarness vulkan;
+    RunGraphicsCase(&vulkan, GraphicsPackedHalfCentroid());
+    return 0;
+  }
+  if (argc == 2 && std::strcmp(argv[1], "--pixel-alias-only") == 0) {
+    CheckPixelParameterAliases();
+    CheckRectListShaders();
+    VulkanHarness vulkan;
+    RunGraphicsCase(&vulkan, GraphicsPackedHalfCentroid());
+    RunGraphicsCase(&vulkan, GraphicsPackedHalfInputAlias(false));
+    RunGraphicsCase(&vulkan, GraphicsPackedHalfInputAlias(true));
+    RunGraphicsCase(&vulkan, GraphicsSmoothRawInputAlias());
+    return 0;
+  }
   if (argc == 2 && std::strcmp(argv[1], "--clip-control-only") == 0) {
     CheckClipControlDepthClipState();
     return 0;
@@ -30561,6 +30952,11 @@ int main(int argc, char **argv) {
     CheckPm4SyntheticOcclusionCounterDump(vulkan.RuntimeRenderer());
     return 0;
   }
+  if (argc == 2 && std::strcmp(argv[1], "--predication-only") == 0) {
+    VulkanHarness vulkan;
+    CheckPm4Predication(vulkan.RuntimeRenderer());
+    return 0;
+  }
   if (argc == 2 && std::strcmp(argv[1], "--descriptor-heap-only") == 0) {
     VulkanHarness vulkan;
     vulkan.CheckDescriptorHeapLargeSet();
@@ -30609,6 +31005,10 @@ int main(int argc, char **argv) {
   if (argc == 2 && std::strcmp(argv[1], "--gpu-command-lane-only") == 0) {
     VulkanHarness vulkan;
     vulkan.CheckGpuCommandLane();
+    CheckPm4IndirectControlFlow(vulkan.RuntimeRenderer());
+    CheckPm4WaitResume(vulkan.RuntimeRenderer());
+    CheckPm4RewindResume(vulkan.RuntimeRenderer());
+    CheckPm4CeCompletion(vulkan.RuntimeRenderer());
     return 0;
   }
   if (argc == 2 && std::strcmp(argv[1], "--alignbyte-only") == 0) {
@@ -30616,6 +31016,11 @@ int main(int argc, char **argv) {
     RunCase(&vulkan, VectorAlignByteUsesFiveBitByteOffset());
     return 0;
   }
+  if (argc == 2 && std::strcmp(argv[1], "--sdwa-mov-only") == 0) {
+    VulkanHarness vulkan;
+    RunCase(&vulkan, Vop1SdwaMovByteDestinations());
+    return 0;
+  }
   if (argc == 2 && std::strcmp(argv[1], "--sdwa-ashr-only") == 0) {
     VulkanHarness vulkan;
     RunCase(&vulkan, Vop2SdwaAshrrevCapturedWord0SignExtends());
@@ -30864,6 +31269,7 @@ int main(int argc, char **argv) {
   CheckVulkan13FeatureRequirements();
   CheckPm4AcquireMemNoOp(vulkan.RuntimeRenderer());
   CheckPm4SyntheticOcclusionCounterDump(vulkan.RuntimeRenderer());
+  CheckPm4Predication(vulkan.RuntimeRenderer());
   CheckPm4StencilInfoValueLane(vulkan.RuntimeRenderer());
   CheckPm4NativeTargetGeometryRegisters(vulkan.RuntimeRenderer());
   CheckPm4PrivateAgcShaderRegisters(vulkan.RuntimeRenderer());
@@ -30878,12 +31284,14 @@ int main(int argc, char **argv) {
   CheckAgcWaitPackets(vulkan.RuntimeRenderer());
   CheckAgcDrawIndirectMultiPacket(vulkan.RuntimeRenderer());
   CheckPm4ContextStateOperations(vulkan.RuntimeRenderer());
+  CheckPm4IndirectControlFlow(vulkan.RuntimeRenderer());
   CheckPm4WaitResume(vulkan.RuntimeRenderer());
   CheckPm4RewindResume(vulkan.RuntimeRenderer());
   CheckPm4CeCompletion(vulkan.RuntimeRenderer());
   CheckEmbeddedFetchVertexOffset();
   CheckEmbeddedFetchLaneSpill();
   CheckTessellationPrograms();
+  CheckPixelParameterAliases();
   CheckRectListShaders();
   CheckIndirectImageKeySwitch();
   CheckPs5GameExampleImageClearRuntimeShape();

tests/VirtualMemoryAllocationTests.cpp

@@ -10,8 +10,12 @@
 #include "loader/redZonePatcher.h"
 #include "loader/runtimeLinker.h"
 #include "loader/systemContent.h"
+#include "loader/x64InstructionEmulator.h"
 
+#include <algorithm>
 #include <array>
+#include <bit>
+#include <cmath>
 #include <cinttypes>
 #include <cstdint>
 #include <cstdio>
@@ -20,8 +24,15 @@
 #include <string>
 #include <vector>
 
+#if defined(__linux__) || KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+#include <csignal>
+#include <immintrin.h>
+#include <xbyak/xbyak.h>
+#endif
+
 #if defined(__linux__)
 #include <sys/uio.h>
+#include <ucontext.h>
 #include <unistd.h>
 #endif
 
@@ -209,10 +220,8 @@ void TestWindowsGuestRedZoneStaticPatcher() {
 	                                   reinterpret_cast<void*>(mapping + CODE_SIZE),
 	                                   TRAMPOLINE_SIZE);
 	const std::array<uintptr_t, 1> function_starts = {static_cast<uintptr_t>(mapping)};
-	const auto result = Loader::PatchRedZoneMemoryInstructions(
-	    mapping, code.size(), function_starts);
-	Check(test, Common::VirtualMemory::FlushInstructionCache(mapping, CODE_SIZE + TRAMPOLINE_SIZE),
-	      "failed to flush patched test code");
+	const auto result = Loader::PatchGuestInstructions(
+	    mapping, code.size(), function_starts, true, false);
 	const bool patched_preserved = function(static_cast<const uint64_t*>(g_red_zone_fault_page)) == 1;
 
 	Loader::UnregisterRedZonePatchModule(reinterpret_cast<void*>(mapping));
@@ -2538,15 +2547,277 @@ void TestModuleRelocationUsesWritableHostMapping() {
 	std::printf("[host]    %-48s ok\n", test);
 }
 
+#if defined(__linux__) || KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+volatile sig_atomic_t g_rsqrt_traps = 0;
+
+bool EmulateReciprocalSquareRootContext(void* context) {
+	const auto host_mxcsr = _mm_getcsr();
+	_mm_setcsr(0x7fe1); // Exercise the handler under a distinct rounding mode and raised flags.
+	const bool emulated        = Loader::X64InstructionEmulator::TryEmulate(context);
+	const bool preserved_mxcsr = _mm_getcsr() == 0x7fe1;
+	_mm_setcsr(host_mxcsr);
+	if (!emulated || !preserved_mxcsr) {
+		return false;
+	}
+	g_rsqrt_traps = g_rsqrt_traps + 1;
+	return true;
+}
+
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+LONG CALLBACK ReciprocalSquareRootHandler(EXCEPTION_POINTERS* exception) {
+	if (exception->ExceptionRecord->ExceptionCode != EXCEPTION_ILLEGAL_INSTRUCTION ||
+	    !EmulateReciprocalSquareRootContext(exception->ContextRecord)) {
+		return EXCEPTION_CONTINUE_SEARCH;
+	}
+	return EXCEPTION_CONTINUE_EXECUTION;
+}
+#else
+void ReciprocalSquareRootHandler(int, siginfo_t*, void* context) {
+	if (!EmulateReciprocalSquareRootContext(context)) {
+		_exit(190);
+	}
+}
+#endif
+
+void TestPackedReciprocalSquareRoot() {
+	const char* test = "PackedReciprocalSquareRoot";
+	constexpr uint64_t code_size = 0x4000;
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+	constexpr uint64_t allocation_size = code_size * 2;
+#else
+	constexpr uint64_t allocation_size = code_size;
+#endif
+	const auto mapping = Libs::LibKernel::Memory::AllocateProgramMemory(
+	    0x902000000, allocation_size, Common::VirtualMemory::Mode::ExecuteReadWrite, "rsqrt_test");
+	Check(test, mapping != 0, "failed to allocate instruction test code");
+	struct RestoreState {
+		uint64_t mapping;
+		uint64_t size;
+		uint32_t mxcsr;
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+		void* handler = nullptr;
+#else
+		struct sigaction previous {};
+		bool installed = false;
+#endif
+		~RestoreState() {
+			_mm_setcsr(mxcsr);
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+			if (handler != nullptr) {
+				RemoveVectoredExceptionHandler(handler);
+			}
+			Loader::UnregisterRedZonePatchModule(reinterpret_cast<void*>(mapping));
+#else
+			if (installed) {
+				sigaction(SIGILL, &previous, nullptr);
+			}
+#endif
+			Libs::LibKernel::Memory::FreeGuestMemory(mapping, size);
+		}
+	} restore {mapping, allocation_size, _mm_getcsr()};
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+	restore.handler = AddVectoredExceptionHandler(1, ReciprocalSquareRootHandler);
+	Check(test, restore.handler != nullptr, "failed to install instruction handler");
+#else
+	struct sigaction action {};
+	action.sa_sigaction = ReciprocalSquareRootHandler;
+	action.sa_flags = SA_SIGINFO;
+	sigemptyset(&action.sa_mask);
+	restore.installed = sigaction(SIGILL, &action, &restore.previous) == 0;
+	Check(test, restore.installed, "failed to install instruction handler");
+#endif
+	using GuestFunction = void(KYTY_SYSV_ABI*)(const uint32_t*, uint32_t*);
+	const auto function = reinterpret_cast<GuestFunction>(mapping);
+	const auto refine = [](float estimate) {
+		return (estimate * 0.5f) * std::fma(-estimate, estimate, 3.0f);
+	};
+	constexpr uint64_t red_zone_sentinel = 0x1122334455667788ull;
+	std::array<uint32_t, 16> input {};
+	std::array<uint32_t, 48> output {};
+	input.fill(0x3f800000);
+	for (const auto registers: {std::array {2, 1}, std::array {10, 9}, std::array {1, 1}}) {
+		const auto destination = registers[0];
+		const auto source = registers[1];
+		input[0] = 0x3f800000;
+		uint32_t expected = 0x3f800000;
+		if (source == 9) {
+			input[0] = 0x40800000;
+			expected = 0x3f000000;
+		}
+		Xbyak::CodeGenerator code(code_size, reinterpret_cast<void*>(mapping));
+		if (source == 9) {
+			code.vmovups(Xbyak::Ymm(1), code.ptr[code.rdi + 32]);
+		}
+		code.vmovups(Xbyak::Ymm(source), code.ptr[code.rdi]);
+		if (destination != source) {
+			code.vmovups(Xbyak::Ymm(destination), code.ptr[code.rdi + 32]);
+		}
+		for (uint32_t offset = 8; offset <= 128; offset += 8) {
+			code.mov(code.rax, red_zone_sentinel ^ offset);
+			code.mov(code.qword[code.rsp - offset], code.rax);
+		}
+		code.vrsqrtps(Xbyak::Xmm(destination), Xbyak::Xmm(source));
+		for (uint32_t offset = 8; offset <= 128; offset += 8) {
+			code.mov(code.rax, code.qword[code.rsp - offset]);
+			code.mov(code.qword[code.rsi + 64 + offset - 8], code.rax);
+		}
+		code.vmovups(code.ptr[code.rsi], Xbyak::Ymm(destination));
+		code.vmovups(code.ptr[code.rsi + 32], Xbyak::Ymm(source));
+		code.vzeroupper();
+		code.ret();
+#if defined(__linux__)
+		const std::vector<uint8_t> original(code.getCode(), code.getCurr());
+#endif
+		Check(test, Common::VirtualMemory::FlushInstructionCache(mapping, code.getSize()),
+		      "failed to flush generated instruction test code");
+		function(input.data(), output.data());
+		const float native = std::bit_cast<float>(output[0]);
+		const auto red_zone_intact = [&] {
+			for (uint32_t offset = 8; offset <= 128; offset += 8) {
+				const auto index = 16 + (offset - 8) / sizeof(uint32_t);
+				const auto value = static_cast<uint64_t>(output[index]) |
+				                   (static_cast<uint64_t>(output[index + 1]) << 32);
+				if (value != (red_zone_sentinel ^ offset)) {
+					return false;
+				}
+			}
+			return true;
+		};
+		Check(test, red_zone_intact(), "native instruction corrupted the guest red zone");
+		Check(test, std::isfinite(native) && std::abs(native - std::bit_cast<float>(expected)) < 0.001f,
+		      "native reciprocal root is outside its error bound");
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+		Loader::RegisterRedZonePatchModule(reinterpret_cast<void*>(mapping), code_size,
+		                                   reinterpret_cast<void*>(mapping + code_size), code_size);
+		const std::array<uintptr_t, 1> function_starts {mapping};
+		// The extended-register case also relocates ordinary memory accesses while
+		// red-zone data is live, exercising both enabled patchers in one function.
+		const bool protect_memory = source == 9;
+		const auto patched = Loader::PatchGuestInstructions(
+		    mapping, code.getSize(), function_starts, protect_memory, true);
+		Check(test, patched.reciprocal_sqrt_instruction_count == 1 &&
+		                patched.unrelocatable_memory_instruction_count == 0 &&
+		                (protect_memory ? patched.patched_memory_instruction_count > 0
+		                                : patched.patched_memory_instruction_count == 0),
+		      "instruction pass lost reciprocal root or memory patch coverage");
+#else
+		Check(test, Loader::X64InstructionEmulator::PatchReciprocalSquareRoots(mapping, code.getSize()) == 1,
+		      "instruction pass did not patch exactly one packed reciprocal root");
+		const auto* patched = reinterpret_cast<const uint8_t*>(mapping);
+		size_t changed = 0;
+		for (size_t i = 0; i < original.size(); ++i) {
+			changed += original[i] != patched[i];
+		}
+		Check(test, changed == 1, "instruction pass changed unrelated code bytes");
+		Check(test, Common::VirtualMemory::FlushInstructionCache(mapping, code.getSize()),
+		      "failed to flush patched instruction test code");
+#endif
+		const auto before = g_rsqrt_traps;
+		function(input.data(), output.data());
+		Check(test, g_rsqrt_traps == before + 1, "patched instruction did not execute its handler");
+		Check(test, red_zone_intact(), "patched instruction corrupted the guest red zone");
+		Check(test, output[0] == expected, "patched instruction read or wrote the wrong register");
+		if (destination != source) {
+			Check(test, std::equal(input.begin(), input.begin() + 4, output.begin() + 8),
+			      "source lower XMM lanes were corrupted");
+		}
+		if (source != 9) {
+			Check(test, refine(refine(std::bit_cast<float>(output[0]))) == 1.0f,
+			      "identity quaternion normalization drifted below one");
+		}
+		for (size_t i = 0; i < 4; ++i) {
+			Check(test, output[4 + i] == 0, "128-bit VEX destination retained upper YMM lanes");
+			const auto expected_upper = destination == source ? 0u : input[4 + i];
+			Check(test, output[12 + i] == expected_upper, "source upper YMM lanes were corrupted");
+		}
+		std::printf("[host]    rsqrt xmm%d,xmm%d native=%08x patched=%08x\n",
+		            destination, source, std::bit_cast<uint32_t>(native), output[0]);
+		if (destination != 2) {
+			continue;
+		}
+		constexpr std::array<std::array<uint32_t, 8>, 4> cases {{
+		    {0x00000000, 0x80000000, 0x00000001, 0x80000001,
+		     0x7f800000, 0xff800000, 0x7f800000, 0xff800000},
+		    {0x7f800000, 0xff800000, 0xbf800000, 0x7fc12345,
+		     0x00000000, 0xffc00000, 0xffc00000, 0x7fc12345},
+		    {0x7f812345, 0xff812345, 0x40800000, 0x40000000,
+		     0x7fc12345, 0xffc12345, 0x3f000000, 0x3f3504f3},
+		    {0x00800000, 0x7f7fffff, 0x3f800000, 0x41800000,
+		     0x5f000000, 0x1f800000, 0x3f800000, 0x3e800000},
+		}};
+		for (const auto& values: cases) {
+			std::copy_n(values.begin(), 4, input.begin());
+			for (uint32_t controls = 0; controls < 8; ++controls) {
+				const uint32_t mxcsr = 0x1fa1u | ((controls & 3u) << 13u) | ((controls & 4u) << 4u);
+				_mm_setcsr(mxcsr);
+				function(input.data(), output.data());
+				const auto result_mxcsr = _mm_getcsr();
+				_mm_setcsr(restore.mxcsr);
+				Check(test, red_zone_intact(), "special-value trap corrupted the guest red zone");
+				Check(test, result_mxcsr == mxcsr, "instruction changed MXCSR controls or exception flags");
+				Check(test, std::equal(output.begin(), output.begin() + 4, values.begin() + 4),
+				      "special values or round-independent reciprocal roots differ from ISA semantics");
+			}
+		}
+		input.fill(0x3f800000);
+	}
+#if defined(__linux__)
+	std::array<uint8_t, 16> unknown {0x0f, 0x0b};
+	ucontext_t context {};
+	context.uc_mcontext.gregs[REG_RIP] = reinterpret_cast<greg_t>(unknown.data());
+	Check(test, !Loader::X64InstructionEmulator::TryEmulate(&context), "unrecognized UD2 was swallowed");
+	_libc_fpstate fpstate {};
+	context.uc_mcontext.fpregs = &fpstate;
+	const auto emulate = [&](std::array<uint8_t, 16> instruction, size_t length) {
+		context.uc_mcontext.gregs[REG_RIP] = reinterpret_cast<greg_t>(instruction.data());
+		Check(test, Loader::X64InstructionEmulator::TryEmulate(&context), "existing instruction emulation was rejected");
+		Check(test, context.uc_mcontext.gregs[REG_RIP] == reinterpret_cast<greg_t>(instruction.data() + length),
+		      "existing instruction emulation advanced RIP incorrectly");
+	};
+	fpstate._xmm[8].element[0] = 0x1234;
+	fpstate._xmm[8].element[2] = 0xdeadbeef;
+	emulate({0x66, 0x41, 0x0f, 0x78, 0xc0, 8, 4}, 7); // extrq xmm8, 8, 4
+	Check(test, fpstate._xmm[8].element[0] == 0x23 && fpstate._xmm[8].element[2] == 0,
+	      "SSE4a extraction lost its extended register or upper-half semantics");
+	fpstate._xmm[8].element[0] = 0x1111;
+	fpstate._xmm[8].element[2] = 0xdeadbeef;
+	fpstate._xmm[9].element[0] = 0xab;
+	emulate({0xf2, 0x45, 0x0f, 0x78, 0xc1, 8, 4}, 7); // insertq xmm8, xmm9, 8, 4
+	Check(test, fpstate._xmm[8].element[0] == 0x1ab1 && fpstate._xmm[8].element[2] == 0xdeadbeef,
+	      "SSE4a insertion corrupted its destination lanes");
+	std::array<uint32_t, 8> sha_source {0, 0, 0, 0, 5, 6, 7, 8};
+	context.uc_mcontext.gregs[REG_RDI] = reinterpret_cast<greg_t>(sha_source.data());
+	for (uint32_t lane = 0; lane < 4; ++lane) {
+		fpstate._xmm[8].element[lane] = lane + 1;
+	}
+	emulate({0x44, 0x0f, 0x38, 0xc9, 0x47, 0x10}, 6); // sha1msg1 xmm8, [rdi+16]
+	constexpr std::array<uint32_t, 4> sha_expected {6, 10, 2, 6};
+	Check(test, std::equal(sha_expected.begin(), sha_expected.end(), fpstate._xmm[8].element),
+	      "SHA emulation lost its memory operand or extended destination register");
+	emulate({0x0f, 0x01, 0xfa}, 3); // monitorx
+#endif
+	std::printf("[host]    %-48s ok\n", test);
+}
+#endif
+
 } // namespace
 
 int main(int argc, char** argv) {
 	InitSubsystems();
+#if defined(__linux__) || KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+	if (argc == 2 && std::strcmp(argv[1], "--rsqrt-only") == 0) {
+		RunTest(TestPackedReciprocalSquareRoot);
+		return g_failed_tests == 0 ? 0 : 1;
+	}
+#endif
 	if (argc == 2 && std::strcmp(argv[1], "--red-zone-patcher-only") == 0) {
 		RunTest(TestWindowsGuestRedZoneStaticPatcher);
 		return g_failed_tests == 0 ? 0 : 1;
 	}
 
+#if defined(__linux__) || KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+	RunTest(TestPackedReciprocalSquareRoot);
+#endif
 	RunTest(TestWindowsGuestRedZoneStaticPatcher);
 	RunTest(TestProsperoArgumentAndInfoSizeContracts);
 	RunTest(TestGuestAddressSpaceOwnsReservationsBeforeBacking);

tests/shaderCfgTests.cpp

@@ -6103,6 +6103,53 @@ void TestCustomVintrpMovTranslation() {
   CheckSpirvBinaryValidates(mixed_linear_result.spirv);
 }
 
+void TestPerspectiveCentroidInputs() {
+  constexpr std::array cases{std::array{4u, UINT32_MAX, 0u},
+                             std::array{5u, UINT32_MAX, 2u},
+                             std::array{6u, 0u, 2u}, std::array{7u, 2u, 4u}};
+  for (const auto &[inputs, center, centroid] : cases) {
+    const uint32_t other = center == UINT32_MAX ? centroid : center;
+    const std::vector<uint32_t> shader = {
+        EncodeExp0(0x00, 0xf), EncodeExp1(centroid, centroid + 1, other, other + 1),
+        EncodeSopp(0x01)};
+    HW::PixelShaderInfo regs{};
+    regs.ps_regs.data_addr = reinterpret_cast<uint64_t>(shader.data());
+    ShaderMappedData mapped{};
+    mapped.code_size_bytes = shader.size() * sizeof(uint32_t);
+    ShaderMapUserData(regs.ps_regs.data_addr, mapped);
+    HW::ShaderRegisters sh{};
+    sh.ps_input_ena = sh.ps_input_addr = inputs;
+    const std::array<Prospero::ColorComponentMapping, 8> mappings{};
+    ShaderPixelInputInfo pixel{};
+    (void)PrepareProgram(regs, sh, mappings, pixel);
+    Check(pixel.ps_perspective_centroid_vgpr == centroid &&
+              pixel.ps_perspective_center_vgpr == center &&
+              pixel.ps_system_input_base == centroid + 2,
+          "perspective-centroid pair did not follow enabled sample/center inputs");
+    const auto key = MakeStageStaticKey(pixel);
+    pixel.ps_perspective_centroid_vgpr = UINT32_MAX;
+    Check(key != MakeStageStaticKey(pixel), "centroid input is missing from shader key");
+    pixel.ps_perspective_centroid_vgpr = centroid;
+    auto options = MakeCompileOptions(ShaderType::Pixel);
+    options.input_info.pixel = &pixel;
+    const auto result = RecompileForTest(shader, options);
+    const auto source = DisassembleSpirvBinary(result.spirv);
+    Check(SpirvContainsCapability(result.spirv, 52u) &&
+              SpirvHasDecorationValue(result.spirv, 11u, 5286u) &&
+              Common::ContainsStr(source, "InterpolateAtCentroid %gl_BaryCoordKHR") &&
+              SpirvSourceHasInstructionUsing(source, "OpCompositeExtract", " 1") &&
+              SpirvSourceHasInstructionUsing(source, "OpCompositeExtract", " 2"),
+          "centroid I/J did not evaluate BaryCoordKHR Y/Z at the centroid");
+    Check(!SpirvHasDecorationValueWithDecoration(result.spirv, 11u, 5286u, 16u),
+          "centroid input changed the shared center builtin's interpolation location");
+    if (center != UINT32_MAX) {
+      Check(SpirvSourceHasInstructionUsing(source, "OpLoad", "%float "),
+            "center pair lost its ordinary barycentric loads when centroid was enabled");
+    }
+    CheckSpirvBinaryValidates(result.spirv);
+  }
+}
+
 void TestPsInputCountRegisterDecode() {
   HW::Context context;
   // NUM_INTERP is 3 while bit 14 is an independent control flag that must be
@@ -9559,6 +9606,15 @@ void TestMergedShaderUserDataSnapshot() {
                          params.user_data.begin() + 8),
           "monolithic NGG shader used stale GS-back state or lost its s8 user data");
   }
+  context.SetMaxOutputPerSubgroup(256);
+  context.SetGsMaxVertOut(8);
+  user_config.SetPrimitiveType(Prospero::PrimitiveType::kTriFan);
+  user_config.SetGeControl({32, 32});
+  ShaderVertexInputInfo fan_input{};
+  PrepareProgram(regs, context, user_config, fan_input);
+  Check(fan_input.mesh.primitives_per_group == 30 && fan_input.mesh.vertices_per_group == 32 &&
+            fan_input.mesh.max_vertices == 256 && fan_input.mesh.max_primitives == 192,
+        "captured triangle-fan GS configuration lost its subgroup assembly limits");
 }
 
 void TestEmbeddedFetchPreservesSharedScalarLoad() {
@@ -9744,6 +9800,18 @@ void TestMeshInputAssembly() {
        0x40000305, 1, 2, 3, 0, 15, false},
       {Prospero::PrimitiveType::kTriStrip, 5, 8, 0, 1, 0, 0, 11,
        0x40000305, 2, 1, 3, 0, 12, false},
+      {Prospero::PrimitiveType::kTriFan, 32, 6, 0, 3, 2, 0x1002, 0,
+       0x40000406, 0, 4, 5, 8, 0x0123, true},
+      {Prospero::PrimitiveType::kTriFan, 5, 8, 1, 0, 2, 0x1002, 0,
+       0x40000305, 0, 1, 2, 0, 0xabcd, true},
+      {Prospero::PrimitiveType::kTriFan, 5, 8, 1, 1, 0, 0, 11,
+       0x40000305, 0, 2, 3, 0, 15, false},
+      {Prospero::PrimitiveType::kTriFan, 5, 7, 1, 3, 2, 0x1002, 11,
+       0x40000204, 0, 4, 5, 12, 0xabd8, true},
+      {Prospero::PrimitiveType::kTriFan, 5, 7, 1, 4, 2, 0x1002, 0,
+       0x40000204, 0, 5, 6, 16, 0, false},
+      {Prospero::PrimitiveType::kTriFan, 5, 8, 1, 64, 4, 0x1000, 0,
+       0x41000000, 0, 65, 66, 268, 0, false},
       {Prospero::PrimitiveType::kLineList, 5, 17, 3, 0, 0, 0, 11,
        0x40000204, 0, 1, 0, 0, 23, false},
       {Prospero::PrimitiveType::kLineList, 6, 17, 2, 4, 2, 0x1002, 11,
@@ -13444,6 +13512,7 @@ int main() {
   TestNewShaderRecompilerNativeBindingPlan();
   TestNewShaderRecompilerStageInputInfo();
   TestCustomVintrpMovTranslation();
+  TestPerspectiveCentroidInputs();
   TestGraphicsCreateInterpolantMapping();
   TestNewShaderRecompilerPixelPipelineEntry();
   TestComputeLdsAllocationIdentity();