Add Nix flake and development shell (#645)

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

.github/workflows/build.yml

@@ -304,15 +304,9 @@ 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 \

.gitmodules

@@ -22,7 +22,7 @@
 	url = https://github.com/GPUOpen-LibrariesAndSDKs/VulkanMemoryAllocator.git
 [submodule "3rdparty/ffmpeg-core"]
 	path = 3rdparty/ffmpeg-core
-	url = https://github.com/shadps4-emu/ext-ffmpeg-core.git
+	url = https://github.com/KytyPS5/ext-ffmpeg-core.git
 	shallow = true
 [submodule "3rdparty/fmt"]
 	path = 3rdparty/fmt

3rdparty/CMakeLists.txt

@@ -88,13 +88,6 @@ 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).

README.md

@@ -165,14 +165,13 @@ The finished application and its runtime dependencies will be placed in
 
 ### Building on Linux
 
-Install the toolchain, system FFmpeg, and the libraries the bundled SDL2 needs. Without the audio, Wayland and
+Install the toolchain 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 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
@@ -193,14 +192,37 @@ 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. The emulator still requires
-the system FFmpeg runtime libraries matching those used to build it.
+`_Build/linux/install` runs without a matching system Qt. FFmpeg is linked statically
+from the pinned [KytyPS5 FFmpeg core](https://github.com/KytyPS5/ext-ffmpeg-core)
+release, including VP9 and WebM support. System FFmpeg packages are not required.
 
 As on Windows, the MSVC compiler is not used; Clang is required. `cl.exe` is rejected at configure
 time.
 
 The CMake source root is the repository root.
 
+### Building on NixOS
+
+A development shell provides Clang, CMake, Ninja, Qt 6, the Vulkan headers, and the SDL2 backend
+libraries. Enter it and configure exactly as on other Linux distributions; the shell exports
+`CMAKE_PREFIX_PATH` and `QT_PLUGIN_PATH`, so the `-DCMAKE_PREFIX_PATH="$Qt6_DIR"` argument is not
+needed:
+
+```bash
+nix-shell # or: nix develop
+git submodule update --init --recursive
+
+cmake -S . -B _Build/linux -G Ninja -DCMAKE_BUILD_TYPE=Release \
+  -DCMAKE_C_COMPILER=clang -DCMAKE_CXX_COMPILER=clang++
+
+cmake --build _Build/linux --target launcher --parallel
+cmake --install _Build/linux --prefix _Build/linux/install
+```
+
+The configure step downloads the FFmpeg prebuilts and the `xbyak`/`zydis` sources, so it needs
+network access; a fully sandboxed `nix build` would require vendoring those inputs. A Vulkan 1.3
+driver must be available at runtime (on NixOS, `hardware.graphics.enable = true`).
+
 ### Building on macOS
 
 macOS builds target x86-64 and run under Rosetta 2 on Apple Silicon, so the PS5's x86-64 game

flake.lock

@@ -0,0 +1,27 @@
+{
+  "nodes": {
+    "nixpkgs": {
+      "locked": {
+        "lastModified": 1789426278,
+        "narHash": "sha256-GnotcvKtbTMnoVx6G4E7ZdUX6487Dky74dqq3pU7iRk=",
+        "owner": "NixOS",
+        "repo": "nixpkgs",
+        "rev": "efe6f071ede9d21c37462d2d6682d5e670099684",
+        "type": "github"
+      },
+      "original": {
+        "owner": "NixOS",
+        "ref": "nixos-unstable",
+        "repo": "nixpkgs",
+        "type": "github"
+      }
+    },
+    "root": {
+      "inputs": {
+        "nixpkgs": "nixpkgs"
+      }
+    }
+  },
+  "root": "root",
+  "version": 7
+}

flake.nix

@@ -0,0 +1,20 @@
+{
+  description = "KytyPS5 development shell";
+
+  inputs.nixpkgs.url = "github:NixOS/nixpkgs/nixos-unstable";
+
+  outputs = { self, nixpkgs }:
+    let
+      systems = [ "x86_64-linux" ];
+      forAllSystems = nixpkgs.lib.genAttrs systems;
+    in
+    {
+      devShells = forAllSystems (system:
+        let
+          pkgs = import nixpkgs { inherit system; };
+        in
+        {
+          default = import ./nix/devshell.nix { inherit pkgs; };
+        });
+    };
+}

nix/devshell.nix

@@ -0,0 +1,84 @@
+{ pkgs }:
+
+let
+  # Build-time dependencies. Every -dev package the bundled SDL2 probes for must
+  # be listed here, otherwise SDL2 quietly drops the audio/Wayland/udev backends.
+  buildDeps = with pkgs; [
+    qt6.qtbase # Concurrent, Network, Widgets (launcher)
+    vulkan-headers
+    vulkan-loader
+    vulkan-tools
+    vulkan-validation-layers
+    libglvnd
+    # X11 / input
+    libX11
+    libXcursor
+    libXext
+    libXfixes
+    libXi
+    libXrandr
+    libXScrnSaver
+    libxkbcommon
+    # Wayland
+    wayland
+    wayland-protocols
+    libdecor
+    # audio + device/dbus
+    alsa-lib
+    libpulseaudio
+    systemd # libudev
+    dbus
+  ];
+
+  # Libraries the built binaries need at runtime.
+  runtimeLibs = with pkgs; [
+    stdenv.cc.cc.lib # libstdc++ / libgcc_s
+    libgcc.lib
+    qt6.qtbase # libQt6Concurrent/Network/Widgets
+    vulkan-loader
+    libglvnd
+    alsa-lib
+    libpulseaudio
+    libX11
+    libXcursor
+    libXext
+    libXfixes
+    libXi
+    libXrandr
+    libXScrnSaver
+    libxkbcommon
+    wayland
+    libdecor
+    systemd
+    dbus
+  ];
+in
+pkgs.mkShell {
+  name = "kytyps5";
+
+  nativeBuildInputs = with pkgs; [
+    clang
+    lld
+    cmake
+    ninja
+    pkg-config
+    git
+    glslang # provides glslangValidator for the bundled shaders
+  ];
+
+  buildInputs = buildDeps;
+
+  shellHook = ''
+    export CMAKE_PREFIX_PATH="${pkgs.qt6.qtbase}:''${CMAKE_PREFIX_PATH:-}"
+    export QT_PLUGIN_PATH="${pkgs.qt6.qtbase}/lib/qt-6/plugins:${pkgs.qt6.qtbase}/lib/qt6/plugins:''${QT_PLUGIN_PATH:-}"
+    export LD_LIBRARY_PATH="${pkgs.lib.makeLibraryPath runtimeLibs}:''${LD_LIBRARY_PATH:-}"
+
+    if [ -z "''${KYTY_SHELL_QUIET:-}" ]; then
+      echo "KytyPS5 dev shell"
+      echo "  cmake -S . -B _Build/linux -G Ninja -DCMAKE_BUILD_TYPE=Release \\"
+      echo "    -DCMAKE_C_COMPILER=clang -DCMAKE_CXX_COMPILER=clang++"
+      echo "  cmake --build _Build/linux --target launcher --parallel"
+      echo "  cmake --install _Build/linux --prefix _Build/linux/install"
+    fi
+  '';
+}

shell.nix

@@ -0,0 +1,4 @@
+let
+  pkgs = import <nixpkgs> { };
+in
+import ./nix/devshell.nix { inherit pkgs; }

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/graphics/shader/recompiler/frontend/decode/ScalarAluOps.cpp

@@ -105,6 +105,7 @@ constexpr OpcodeMap SOPP_OPCODE_LIST[] = {
     {0x10u, Opcode::S_SENDMSG},
     {0x12u, Opcode::S_TRAP},
     {0x16u, Opcode::S_TTRACEDATA},
+    {0x17u, Opcode::S_CBRANCH_CDBGSYS},
     {0x20u, Opcode::S_INST_PREFETCH},
     {0x23u, Opcode::S_WAITCNT_DEPCTR},
 };

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

@@ -187,6 +187,8 @@ std::string FormatExp(const Instruction& inst) {
 
 bool IsConditionalBranch(Opcode opcode) {
 	switch (opcode) {
+		// DevKit NGG should be disabled.
+		case Opcode::S_CBRANCH_CDBGSYS: return false;
 		case Opcode::S_CBRANCH_SCC0:
 		case Opcode::S_CBRANCH_SCC1:
 		case Opcode::S_CBRANCH_VCCZ:
@@ -554,6 +556,7 @@ std::string InstructionToString(const Instruction& inst) {
 		case Opcode::S_CBRANCH_VCCNZ:
 		case Opcode::S_CBRANCH_EXECZ:
 		case Opcode::S_CBRANCH_EXECNZ:
+		case Opcode::S_CBRANCH_CDBGSYS:
 			return WithUnsupportedReason(inst, fmt::format("0x{:08x}: {} 0x{:08x}", inst.pc,
 			                                               magic_enum::enum_name(inst.opcode),
 			                                               inst.branch_target));

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

@@ -581,6 +581,7 @@ enum class Opcode {
 	S_CBRANCH_VCCNZ,
 	S_CBRANCH_EXECZ,
 	S_CBRANCH_EXECNZ,
+	S_CBRANCH_CDBGSYS,
 	S_SENDMSG,
 	S_SETREG_B32,
 	S_SLEEP,

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

@@ -227,6 +227,7 @@ bool Translator::EmitScalar(const Decoder::Instruction& inst) {
 		case O::S_CBRANCH_VCCNZ:
 		case O::S_CBRANCH_EXECZ:
 		case O::S_CBRANCH_EXECNZ:
+		case O::S_CBRANCH_CDBGSYS:
 		case O::S_ENDPGM: return true;
 		default: return 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);

src/graphics/shader/shader.cpp

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

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

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/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/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/libKernel.cpp

@@ -29,11 +29,9 @@
 #include <cstdio>
 #include <cstdlib>
 #include <cstring>
-#include <mutex>
 #include <random>
 #include <string>
 #include <thread>
-#include <unordered_map>
 #include <vector>
 
 #if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
@@ -43,9 +41,11 @@
 #include <windows.h>
 #elif defined(__APPLE__)
 #include <csignal>
+#include <pthread.h>
 #include <sys/ucontext.h>
 #else
 #include <csignal>
+#include <pthread.h>
 #include <ucontext.h>
 #endif
 
@@ -2356,132 +2356,71 @@ static bool FiberIsValid(const FiberObject* fiber) {
 	       fiber->magic_end == FIBER_MAGIC_END;
 }
 
-static thread_local FiberObject*                  g_current_fiber       = nullptr;
-static thread_local FiberObject*                  g_thread_return_fiber = nullptr;
-static thread_local FiberObject*                  g_starting_fiber      = nullptr;
-static thread_local FiberObject*                  g_pending_idle_fiber  = nullptr;
-static thread_local FiberCpuContext               g_thread_fiber_context {};
-static std::mutex                                 g_fiber_owner_mutex;
-static std::unordered_map<FiberObject*, uint64_t> g_fiber_owner_thread;
-static std::unordered_map<uint64_t, FiberObject*> g_fiber_current_by_thread;
+struct FiberThreadContext {
+	FiberObject*    current_fiber;
+	FiberObject*    returned_fiber;
+	FiberObject*    pending_idle_fiber;
+	FiberCpuContext cpu_context;
+};
+
+static const auto g_fiber_context_key = [] {
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+	const auto key = TlsAlloc();
+	if (key == TLS_OUT_OF_INDEXES) {
+		std::abort();
+	}
+#else
+	pthread_key_t key;
+	if (pthread_key_create(&key, nullptr) != 0) {
+		std::abort();
+	}
+#endif
+	return key;
+}();
+
+// Native lookup must stay fresh when a fiber resumes on another host thread.
+static FiberThreadContext* FiberGetThreadContext() {
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+	return static_cast<FiberThreadContext*>(TlsGetValue(g_fiber_context_key));
+#else
+	return static_cast<FiberThreadContext*>(pthread_getspecific(g_fiber_context_key));
+#endif
+}
+
+static void FiberSetThreadContext(FiberThreadContext* context) {
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+	if (TlsSetValue(g_fiber_context_key, context) == 0) {
+#else
+	if (pthread_setspecific(g_fiber_context_key, context) != 0) {
+#endif
+		std::abort();
+	}
+}
 
 static uint32_t FiberLoadState(const FiberObject* fiber) {
 	auto& state = const_cast<uint32_t&>(fiber->state);
 	return std::atomic_ref<uint32_t>(state).load(std::memory_order_acquire);
 }
 
-static uint64_t FiberCurrentHostThreadId() {
-#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
-	return static_cast<uint64_t>(::GetCurrentThreadId());
-#else
-	return std::hash<std::thread::id> {}(std::this_thread::get_id());
-#endif
-}
-
-static void FiberSetOwner(FiberObject* fiber) {
-	std::lock_guard lock(g_fiber_owner_mutex);
-	g_fiber_owner_thread[fiber] = FiberCurrentHostThreadId();
-}
-
-static void FiberClearOwner(FiberObject* fiber) {
-	std::lock_guard lock(g_fiber_owner_mutex);
-	g_fiber_owner_thread.erase(fiber);
-}
-
-static uint64_t FiberGetOwner(FiberObject* fiber) {
-	std::lock_guard lock(g_fiber_owner_mutex);
-	auto            it = g_fiber_owner_thread.find(fiber);
-	return it != g_fiber_owner_thread.end() ? it->second : 0;
-}
-
-static void FiberSetCurrentFiber(FiberObject* fiber) {
-	g_current_fiber = fiber;
-
-	std::lock_guard lock(g_fiber_owner_mutex);
-	const auto      thread_id = FiberCurrentHostThreadId();
-	if (fiber != nullptr) {
-		g_fiber_current_by_thread[thread_id] = fiber;
-	} else {
-		g_fiber_current_by_thread.erase(thread_id);
-	}
-}
-
 static void FiberStoreState(FiberObject* fiber, uint32_t state) {
 	std::atomic_ref<uint32_t>(fiber->state).store(state, std::memory_order_release);
-	if (state == FIBER_STATE_RUNNING) {
-		FiberSetOwner(fiber);
-	} else {
-		FiberClearOwner(fiber);
-	}
 }
 
-static void FiberDeferIdle(FiberObject* fiber) {
-	g_pending_idle_fiber = fiber;
-}
-
-static void FiberCommitDeferredIdle() {
-	auto* fiber          = g_pending_idle_fiber;
-	g_pending_idle_fiber = nullptr;
+static void FiberCompleteSwitch(FiberObject* current) {
+	auto* context = FiberGetThreadContext();
+	context->current_fiber = current;
+	// Only publish IDLE after switching away from the departing fiber's stack.
+	auto* fiber = context->pending_idle_fiber;
+	context->pending_idle_fiber = nullptr;
 	if (fiber != nullptr) {
 		FiberStoreState(fiber, FIBER_STATE_IDLE);
 	}
 }
 
 static bool FiberCompareExchangeState(FiberObject* fiber, uint32_t expected, uint32_t desired) {
-	const bool ok = std::atomic_ref<uint32_t>(fiber->state)
-	                    .compare_exchange_strong(expected, desired, std::memory_order_acq_rel,
-	                                             std::memory_order_acquire);
-	if (ok) {
-		if (desired == FIBER_STATE_RUNNING) {
-			FiberSetOwner(fiber);
-		} else {
-			FiberClearOwner(fiber);
-		}
-	}
-	return ok;
-}
-
-static bool FiberWaitAndEnterRunning(FiberObject* fiber, uint32_t* observed_state) {
-	const auto start = std::chrono::steady_clock::now();
-	uint32_t   spin  = 0;
-	auto&      state = fiber->state;
-	auto       ref   = std::atomic_ref<uint32_t>(state);
-
-	for (;;) {
-		uint32_t expected = FIBER_STATE_IDLE;
-		if (ref.compare_exchange_strong(expected, FIBER_STATE_RUNNING, std::memory_order_acq_rel,
-		                                std::memory_order_acquire)) {
-			FiberSetOwner(fiber);
-			return true;
-		}
-		if (observed_state != nullptr) {
-			*observed_state = expected;
-		}
-		if (expected != FIBER_STATE_RUNNING) {
-			return false;
-		}
-		if (std::chrono::steady_clock::now() - start >= std::chrono::milliseconds(5)) {
-			return false;
-		}
-		if (spin++ < 64) {
-			std::this_thread::yield();
-		} else {
-			std::this_thread::sleep_for(std::chrono::microseconds(100));
-		}
-	}
-}
-
-static bool FiberRepairStaleRunningOnThisThread(FiberObject* fiber, uint32_t observed_state) {
-	if (observed_state != FIBER_STATE_RUNNING || fiber == g_current_fiber) {
-		return false;
-	}
-
-	const auto owner = FiberGetOwner(fiber);
-	if (owner != 0) {
-		return false;
-	}
-
-	return FiberCompareExchangeState(fiber, FIBER_STATE_RUNNING, FIBER_STATE_IDLE);
+	return std::atomic_ref<uint32_t>(fiber->state)
+	    .compare_exchange_strong(expected, desired, std::memory_order_acq_rel,
+	                             std::memory_order_acquire);
 }
 
 static void FiberSetContextValid(FiberObject* fiber, bool valid) {
@@ -2538,45 +2477,36 @@ static void FiberRestoreContext(FiberCpuContext* ctx, uint64_t ret) {
 }
 #endif
 
-[[noreturn]] static void FiberStartTrampoline();
+[[noreturn]] static KYTY_SYSV_ABI void FiberStartTrampoline(FiberObject* fiber);
 
-[[noreturn]] static void FiberStartOnGuestStack(FiberObject* fiber) {
+[[noreturn]] static KYTY_SYSV_ABI void FiberStartOnGuestStack(FiberObject* fiber) {
 	FiberCpuContext ctx {};
 	const auto      stack_top = reinterpret_cast<uintptr_t>(fiber->addr_context) +
 	                            static_cast<uintptr_t>(fiber->size_context);
 	auto            rsp       = (stack_top & ~static_cast<uintptr_t>(0x0f));
-#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
-	rsp -= 4u * sizeof(uint64_t);
-#endif
 	rsp -= sizeof(uint64_t);
 	*reinterpret_cast<uint64_t*>(rsp) = 0;
 
 	ctx.rsp = rsp;
 	ctx.rip = reinterpret_cast<uint64_t>(&FiberStartTrampoline);
+	ctx.rdi = reinterpret_cast<uint64_t>(fiber);
 
-	g_starting_fiber = fiber;
 	FiberRestoreContext(&ctx, 1);
 }
 
-[[noreturn]] static void FiberStartTrampoline() {
-	auto* fiber      = g_starting_fiber;
-	g_starting_fiber = nullptr;
-	if (fiber == nullptr) {
-		EXIT("Fiber start without current fiber\n");
-	}
-
-	FiberCommitDeferredIdle();
-	FiberSetCurrentFiber(fiber);
+[[noreturn]] static KYTY_SYSV_ABI void FiberStartTrampoline(FiberObject* fiber) {
+	FiberCompleteSwitch(fiber);
 
 	fiber->entry(fiber->arg_on_initialize, fiber->arg_on_run);
 
 	FiberStoreState(fiber, FIBER_STATE_TERMINATED);
 	FiberSetContextValid(fiber, false);
 	fiber->arg_on_return  = 0;
-	g_thread_return_fiber = fiber;
-	FiberSetCurrentFiber(nullptr);
+	auto* context = FiberGetThreadContext();
+	context->returned_fiber = fiber;
+	context->current_fiber = nullptr;
 
-	FiberRestoreContext(&g_thread_fiber_context, 1);
+	FiberRestoreContext(&context->cpu_context, 1);
 }
 
 int32_t KYTY_SYSV_ABI FiberInitialize(FiberObject* fiber, const char* name, FiberEntry entry,
@@ -2686,34 +2616,39 @@ int32_t KYTY_SYSV_ABI FiberRun(FiberObject* fiber, uint64_t arg_on_run, uint64_t
 	if (!FiberIsValid(fiber)) {
 		return FIBER_ERROR_INVALID;
 	}
-	if (g_current_fiber != nullptr) {
+	if (FiberGetThreadContext() != nullptr) {
 		return FIBER_ERROR_PERMISSION;
 	}
 	if (!FiberCompareExchangeState(fiber, FIBER_STATE_IDLE, FIBER_STATE_RUNNING)) {
 		return FIBER_ERROR_STATE;
 	}
 
-	fiber->arg_on_run     = arg_on_run;
-	fiber->arg_on_return  = 0;
-	g_thread_return_fiber = nullptr;
+	fiber->arg_on_run    = arg_on_run;
+	fiber->arg_on_return = 0;
+	// Register on the host stack; switching fibers only reads the native TLS slot.
+	FiberThreadContext context {};
+	FiberSetThreadContext(&context);
 
-	if (FiberSaveContext(&g_thread_fiber_context) == 0) {
+	if (FiberSaveContext(&context.cpu_context) == 0) {
 		if (fiber->context_valid) {
 			FiberRestoreContext(&fiber->saved_context, 1);
 		}
 		FiberStartOnGuestStack(fiber);
 	}
 
-	FiberCommitDeferredIdle();
-	FiberSetCurrentFiber(nullptr);
-	auto* returned_fiber  = (g_thread_return_fiber != nullptr ? g_thread_return_fiber : fiber);
-	g_thread_return_fiber = nullptr;
+	auto* returned_fiber  = (context.returned_fiber != nullptr ? context.returned_fiber : fiber);
+	const auto return_value = returned_fiber->arg_on_return;
+	const auto return_code =
+	    FiberLoadState(returned_fiber) == FIBER_STATE_TERMINATED ? FIBER_ERROR_STATE : OK;
+	// Copy results before another thread can acquire and resume the fiber.
+	FiberCompleteSwitch(nullptr);
+	FiberSetThreadContext(nullptr);
 
 	if (arg_on_return != nullptr) {
-		*arg_on_return = returned_fiber->arg_on_return;
+		*arg_on_return = return_value;
 	}
 
-	return (FiberLoadState(returned_fiber) == FIBER_STATE_TERMINATED ? FIBER_ERROR_STATE : OK);
+	return return_code;
 }
 
 int32_t KYTY_SYSV_ABI FiberSwitch(FiberObject* fiber, uint64_t arg_on_run,
@@ -2723,38 +2658,30 @@ int32_t KYTY_SYSV_ABI FiberSwitch(FiberObject* fiber, uint64_t arg_on_run,
 	if (!FiberIsValid(fiber)) {
 		return FIBER_ERROR_INVALID;
 	}
-	if (g_current_fiber == nullptr) {
+	auto* context = FiberGetThreadContext();
+	if (context == nullptr) {
 		return FIBER_ERROR_PERMISSION;
 	}
 
-	for (;;) {
-		uint32_t observed_state = 0;
-		if (FiberWaitAndEnterRunning(fiber, &observed_state)) {
-			break;
-		}
-		if (FiberRepairStaleRunningOnThisThread(fiber, observed_state)) {
-			continue;
-		}
+	if (!FiberCompareExchangeState(fiber, FIBER_STATE_IDLE, FIBER_STATE_RUNNING)) {
 		return FIBER_ERROR_STATE;
 	}
 
-	auto* caller = g_current_fiber;
+	auto* caller = context->current_fiber;
 
 	fiber->arg_on_run    = arg_on_run;
 	fiber->arg_on_return = 0;
 
 	if (FiberSaveContext(&caller->saved_context) == 0) {
 		FiberSetContextValid(caller, true);
-		FiberDeferIdle(caller);
+		context->pending_idle_fiber = caller;
 		if (fiber->context_valid) {
 			FiberRestoreContext(&fiber->saved_context, 1);
 		}
 		FiberStartOnGuestStack(fiber);
 	}
 
-	FiberCommitDeferredIdle();
-	FiberSetCurrentFiber(caller);
-	FiberStoreState(caller, FIBER_STATE_RUNNING);
+	FiberCompleteSwitch(caller);
 
 	if (arg_on_return != nullptr) {
 		*arg_on_return = caller->arg_on_run;
@@ -2770,7 +2697,8 @@ int32_t KYTY_SYSV_ABI FiberGetSelf(FiberObject** fiber) {
 		return FIBER_ERROR_NULL;
 	}
 
-	*fiber = g_current_fiber;
+	const auto* context = FiberGetThreadContext();
+	*fiber = context != nullptr ? context->current_fiber : nullptr;
 
 	return OK;
 }
@@ -2778,23 +2706,22 @@ int32_t KYTY_SYSV_ABI FiberGetSelf(FiberObject** fiber) {
 int32_t KYTY_SYSV_ABI FiberReturnToThread(uint64_t arg_on_return, uint64_t* arg_on_run) {
 	PRINT_NAME();
 
-	if (g_current_fiber == nullptr) {
+	auto* context = FiberGetThreadContext();
+	if (context == nullptr) {
 		return FIBER_ERROR_PERMISSION;
 	}
 
-	auto* fiber          = g_current_fiber;
+	auto* fiber          = context->current_fiber;
 	fiber->arg_on_return = arg_on_return;
 
 	if (FiberSaveContext(&fiber->saved_context) == 0) {
 		FiberSetContextValid(fiber, true);
-		g_thread_return_fiber = fiber;
-		FiberDeferIdle(fiber);
-		FiberRestoreContext(&g_thread_fiber_context, 1);
+		context->returned_fiber = fiber;
+		context->pending_idle_fiber = fiber;
+		FiberRestoreContext(&context->cpu_context, 1);
 	}
 
-	FiberCommitDeferredIdle();
-	FiberSetCurrentFiber(fiber);
-	FiberStoreState(fiber, FIBER_STATE_RUNNING);
+	FiberCompleteSwitch(fiber);
 	if (arg_on_run != nullptr) {
 		*arg_on_run = fiber->arg_on_run;
 	}

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/VirtualMemoryAllocationTests.cpp

@@ -10,18 +10,32 @@
 #include "loader/redZonePatcher.h"
 #include "loader/runtimeLinker.h"
 #include "loader/systemContent.h"
+#include "loader/x64InstructionEmulator.h"
 
+#include <algorithm>
 #include <array>
+#include <atomic>
+#include <bit>
+#include <chrono>
+#include <cmath>
 #include <cinttypes>
 #include <cstdint>
 #include <cstdio>
 #include <cstring>
 #include <filesystem>
 #include <string>
+#include <thread>
 #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
 
@@ -35,6 +49,19 @@
 #endif
 #endif
 
+namespace Libs::Fiber {
+struct FiberObject;
+struct FiberOptParam;
+using FiberEntry = KYTY_SYSV_ABI void (*)(uint64_t, uint64_t);
+int32_t KYTY_SYSV_ABI FiberInitialize(FiberObject*, const char*, FiberEntry, uint64_t, void*,
+                                     uint64_t, const FiberOptParam*, uint32_t);
+int32_t KYTY_SYSV_ABI FiberFinalize(FiberObject*);
+int32_t KYTY_SYSV_ABI FiberRun(FiberObject*, uint64_t, uint64_t*);
+int32_t KYTY_SYSV_ABI FiberSwitch(FiberObject*, uint64_t, uint64_t*);
+int32_t KYTY_SYSV_ABI FiberGetSelf(FiberObject**);
+int32_t KYTY_SYSV_ABI FiberReturnToThread(uint64_t, uint64_t*);
+} // namespace Libs::Fiber
+
 namespace {
 
 using Libs::LibKernel::Memory::VirtualQueryInfo;
@@ -209,10 +236,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 +2563,449 @@ 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
+
+#if defined(__x86_64__) || defined(_M_X64)
+constexpr int32_t FiberErrorState = -2141650938; // SCE_FIBER_ERROR_STATE
+constexpr int32_t FiberErrorPermission = -2141650939;
+
+struct FiberRoundTrip {
+	Libs::Fiber::FiberObject* first;
+	Libs::Fiber::FiberObject* second;
+	std::atomic<bool> running {false};
+	std::atomic<bool> release {false};
+	int first_errors = 0;
+	int second_errors = 0;
+	uint32_t first_visits = 0;
+	uint32_t second_visits = 0;
+	uint64_t first_arg = 0;
+	uint64_t second_arg = 0;
+};
+
+[[noreturn]] void KYTY_SYSV_ABI FirstFiberEntry(uint64_t initial, uint64_t arg) {
+	auto& data = *reinterpret_cast<FiberRoundTrip*>(initial);
+	for (;;) {
+		Libs::Fiber::FiberObject* self = nullptr;
+		data.first_errors |= Libs::Fiber::FiberGetSelf(&self);
+		data.first_errors |= self != data.first;
+		data.first_errors |= Libs::Fiber::FiberSwitch(self, 0, nullptr) != FiberErrorState;
+		if (data.first_errors != 0) {
+			__builtin_trap();
+		}
+		++data.first_visits;
+		data.first_arg = arg;
+		if (arg == 20) {
+			data.running.store(true, std::memory_order_release);
+			while (!data.release.load(std::memory_order_acquire)) {
+				asm volatile("pause");
+			}
+		}
+		data.first_errors |= Libs::Fiber::FiberSwitch(data.second, arg + 1, &arg);
+		if (data.first_errors != 0) {
+			__builtin_trap();
+		}
+		data.first_errors |= Libs::Fiber::FiberReturnToThread(arg + 1, &arg);
+	}
+}
+
+[[noreturn]] void KYTY_SYSV_ABI SecondFiberEntry(uint64_t initial, uint64_t arg) {
+	auto& data = *reinterpret_cast<FiberRoundTrip*>(initial);
+	for (;;) {
+		Libs::Fiber::FiberObject* self = nullptr;
+		data.second_errors |= Libs::Fiber::FiberGetSelf(&self);
+		data.second_errors |= self != data.second;
+		if (data.second_errors != 0) {
+			__builtin_trap();
+		}
+		++data.second_visits;
+		data.second_arg = arg;
+		data.second_errors |= Libs::Fiber::FiberSwitch(data.first, arg + 1, &arg);
+		if (data.second_errors != 0) {
+			__builtin_trap();
+		}
+	}
+}
+
+void TestSmallFiberStacksAndMigration() {
+	const char* test = "SmallFiberStacksAndMigration";
+	// 256-byte objects, 8-byte object alignment, 16-byte context
+	// alignment, and a 512-byte minimum context. The game supplies 2048 bytes.
+	for (const size_t stack_size: {512u, 2048u}) {
+		alignas(8) std::array<uint8_t, 256> first_object {};
+		alignas(8) std::array<uint8_t, 256> second_object {};
+		constexpr size_t GuardSize = 4096;
+		alignas(16) std::array<uint8_t, GuardSize + 2048 + 64> first_stack;
+		alignas(16) std::array<uint8_t, GuardSize + 2048 + 64> second_stack;
+		first_stack.fill(0xa5);
+		second_stack.fill(0xa5);
+		FiberRoundTrip data {reinterpret_cast<Libs::Fiber::FiberObject*>(first_object.data()),
+		                     reinterpret_cast<Libs::Fiber::FiberObject*>(second_object.data())};
+		CheckOk(test, Libs::Fiber::FiberInitialize(data.first, "first", FirstFiberEntry,
+		    reinterpret_cast<uint64_t>(&data), first_stack.data() + GuardSize, stack_size,
+		    nullptr, 0x0a000000), "initialize first fiber");
+		CheckOk(test, Libs::Fiber::FiberInitialize(data.second, "second", SecondFiberEntry,
+		    reinterpret_cast<uint64_t>(&data), second_stack.data() + GuardSize, stack_size,
+		    nullptr, 0x0a000000), "initialize second fiber");
+		const auto thread_context_is_clear = [&] {
+			Libs::Fiber::FiberObject* self = data.first;
+			return Libs::Fiber::FiberGetSelf(&self) == OK && self == nullptr &&
+			       Libs::Fiber::FiberSwitch(data.first, 0, nullptr) == FiberErrorPermission &&
+			       Libs::Fiber::FiberReturnToThread(0, nullptr) == FiberErrorPermission;
+		};
+		Check(test, thread_context_is_clear(), "thread retained a fiber context before its first run");
+		uint64_t first_base = 0;
+		uint64_t second_base = 0;
+		std::memcpy(&first_base, first_stack.data() + GuardSize, sizeof(first_base));
+		std::memcpy(&second_base, second_stack.data() + GuardSize, sizeof(second_base));
+		const auto check_stacks = [&] {
+			for (const auto* stack: {&first_stack, &second_stack}) {
+				const auto is_guard = [](uint8_t byte) { return byte == 0xa5; };
+				Check(test, std::all_of(stack->begin(), stack->begin() + GuardSize, is_guard) &&
+				                std::all_of(stack->begin() + GuardSize + stack_size, stack->end(), is_guard),
+				      "fiber wrote outside its supplied stack");
+			}
+			Check(test, std::memcmp(&first_base, first_stack.data() + GuardSize, sizeof(first_base)) == 0 &&
+			                std::memcmp(&second_base, second_stack.data() + GuardSize, sizeof(second_base)) == 0,
+			      "fiber overwrote the bottom of its stack");
+		};
+		uint64_t returned = 0;
+		CheckOk(test, Libs::Fiber::FiberRun(data.first, 10, &returned), "run first fiber");
+		Check(test, returned == 13 && data.first_visits == 1 && data.second_visits == 1 &&
+		                data.first_arg == 10 && data.second_arg == 11,
+		      "switch/return arguments were not preserved");
+		Check(test, thread_context_is_clear(), "run retained its thread context after returning");
+		check_stacks();
+
+		std::atomic<bool> done {false};
+		int worker_result = -1;
+		int worker_repeat_result = -1;
+		bool worker_context_clear = false;
+		uint64_t worker_returned = 0;
+		uint64_t worker_repeat_returned = 0;
+		std::thread worker([&] {
+			worker_context_clear = thread_context_is_clear();
+			worker_result = Libs::Fiber::FiberRun(data.first, 20, &worker_returned);
+			worker_context_clear = thread_context_is_clear() && worker_context_clear;
+			worker_repeat_result = Libs::Fiber::FiberRun(data.first, 30, &worker_repeat_returned);
+			worker_context_clear = thread_context_is_clear() && worker_context_clear;
+			done.store(true, std::memory_order_release);
+		});
+		const auto deadline = std::chrono::steady_clock::now() + std::chrono::seconds(5);
+		while (!data.running.load(std::memory_order_acquire) &&
+		       !done.load(std::memory_order_acquire) && std::chrono::steady_clock::now() < deadline) {
+			std::this_thread::yield();
+		}
+		const bool running = data.running.load(std::memory_order_acquire);
+		Libs::Fiber::FiberObject* main_self = data.first;
+		const int main_self_result = Libs::Fiber::FiberGetSelf(&main_self);
+		const int busy_run = running ? Libs::Fiber::FiberRun(data.first, 30, nullptr) : 0;
+		const int busy_finalize = running ? Libs::Fiber::FiberFinalize(data.first) : 0;
+		data.release.store(true, std::memory_order_release);
+		worker.join();
+		Check(test, thread_context_is_clear() && main_self_result == OK && main_self == nullptr,
+		      "migrated fiber changed the original thread's current fiber");
+		Check(test, running && busy_run == FiberErrorState && busy_finalize == FiberErrorState,
+		      "running fiber was not exclusively owned");
+		CheckOk(test, worker_result, "resume fiber on another thread");
+		CheckOk(test, worker_repeat_result, "repeat run on another thread");
+		Check(test, worker_context_clear, "worker retained a context outside a fiber run");
+		Check(test, worker_returned == 23 && worker_repeat_returned == 33 &&
+		                data.first_visits == 3 && data.second_visits == 3 &&
+		                data.first_arg == 30 && data.second_arg == 31 &&
+		                data.first_errors == 0 && data.second_errors == 0,
+		      "migration lost fiber identity, arguments, or resumable contexts");
+		CheckOk(test, Libs::Fiber::FiberRun(data.first, 40, &returned), "migrate fiber back to original thread");
+		Check(test, returned == 43 && data.first_visits == 4 && data.second_visits == 4 &&
+		                data.first_arg == 40 && data.second_arg == 41 &&
+		                data.first_errors == 0 && data.second_errors == 0,
+		      "return migration reused an expired thread context");
+		Check(test, thread_context_is_clear(), "repeated run retained its thread context after returning");
+		check_stacks();
+		CheckOk(test, Libs::Fiber::FiberFinalize(data.first), "finalize first suspended fiber");
+		CheckOk(test, Libs::Fiber::FiberFinalize(data.second), "finalize second suspended fiber");
+		std::printf("[host]    %s stack=%zu ok\n", test, stack_size);
+	}
+}
+#endif
+
 } // namespace
 
 int main(int argc, char** argv) {
 	InitSubsystems();
+#if defined(__x86_64__) || defined(_M_X64)
+	if (argc == 2 && std::strcmp(argv[1], "--fiber-only") == 0) {
+		RunTest(TestSmallFiberStacksAndMigration);
+		return g_failed_tests == 0 ? 0 : 1;
+	}
+#endif
+#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(__x86_64__) || defined(_M_X64)
+	RunTest(TestSmallFiberStacksAndMigration);
+#endif
+#if defined(__linux__) || KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+	RunTest(TestPackedReciprocalSquareRoot);
+#endif
 	RunTest(TestWindowsGuestRedZoneStaticPatcher);
 	RunTest(TestProsperoArgumentAndInfoSizeContracts);
 	RunTest(TestGuestAddressSpaceOwnsReservationsBeforeBacking);

tests/shaderCfgTests.cpp

@@ -1795,6 +1795,39 @@ void TestSopkCompareImmediateExtension() {
   }
 }
 
+void TestDisabledSystemDebugBranch() {
+  using namespace ShaderRecompiler;
+  // Relocate PPSA08709 MS pc 0x530 to zero, retaining its displacement to the
+  // system validation helper immediately after the main shader's S_ENDPGM.
+  std::array<uint32_t, 274> shader;
+  shader.fill(EncodeSopp(0x00));
+  shader[0] = 0xbf97010fu;
+  shader[1] = EncodeVop1(0x01, 1, 129);
+  shader[2] = EncodeMubuf0(0x1c, 0, false);
+  shader[3] = EncodeMubuf1(1, 0, 0);
+  shader[271] = EncodeSopp(0x01);
+  shader[272] = 0xffffffffu; // System-only helper must not be decoded.
+  shader[273] = 0xffffffffu;
+  Decoder::Program decoded;
+  Decoder::DecodeProgram(shader, decoded);
+  Check(decoded.instructions.front().opcode == Decoder::Opcode::S_CBRANCH_CDBGSYS &&
+            decoded.instructions.front().branch_target == 0x440u &&
+            decoded.instructions.back().pc == 0x43cu &&
+            decoded.instructions.back().opcode == Decoder::Opcode::S_ENDPGM,
+        "disabled system debug branch reached its post-ENDPGM validation helper");
+  auto graph = CFG::BuildGraph(decoded);
+  Check(graph.blocks.size() == 1u && graph.FindBlockByPc(0x440u) == nullptr &&
+            graph.blocks.front().terminator.kind == CFG::TerminatorKind::Return,
+        "disabled system debug branch changed normal shader control flow");
+  auto options = MakeCompileOptions(ShaderType::Compute);
+  options.dump_ir = true;
+  auto result = RecompileForTest(shader, options);
+  Check(Common::ContainsStr(result.decoded_dump, "S_CBRANCH_CDBGSYS 0x00000440") &&
+            Common::ContainsStr(result.ir_dump, "StoreBufferU32"),
+        "disabled system debug branch lost its identity or fallthrough write");
+  CheckSpirvBinaryValidates(result.spirv);
+}
+
 void TestNewShaderRecompilerRdna2ScalarOpcodes() {
   const uint32_t shader[] = {
       EncodeSMovB32(2, 135),         // s2 = 7
@@ -9606,6 +9639,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() {
@@ -9791,6 +9833,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,
@@ -13448,6 +13502,7 @@ int main() {
   TestNewShaderRecompilerBufferLoadsGuardedByExec();
   TestNewShaderRecompilerBufferAtomicsGuardedByBounds();
   TestCapturedBufferAtomicsX2();
+  TestDisabledSystemDebugBranch();
   TestNewShaderRecompilerPixelImageSampleLodSelection();
   TestNewShaderRecompilerBranchConditionForms();
   TestNewShaderRecompilerSetpcBranch();