loader: emulate immediate SSE4a instructions on macOS (#635)

* loader: emulate immediate SSE4a instructions on macOS

* tests: simplify macOS SSE4a test integration

Keep the Darwin target, CTest registration, and aggregate dependency in
one conditional block without a separate target list. Preserve the
standalone target when BUILD_TESTING is disabled.

Lowercase the test filename and diagnostic prefix, and drop unrelated
formatting changes from the instruction emulator.

Credit Brandon Strong's related SSE4a work in
https://github.com/KytyPS5/KytyPS5/pull/467.

Co-authored-by: Brandon Strong <brandostrong@gmail.com>

---------

Co-authored-by: nmzik <Nmzik@mail.ru>
Co-authored-by: Brandon Strong <brandostrong@gmail.com>

.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).

CMakeLists.txt

@@ -601,6 +601,21 @@ if(BUILD_TESTING)
 	)
 endif()
 
+# Test Darwin signal-context restoration for SSE4a instructions.
+if(APPLE AND (CMAKE_OSX_ARCHITECTURES STREQUAL "x86_64" OR
+		(NOT CMAKE_OSX_ARCHITECTURES AND CMAKE_SYSTEM_PROCESSOR MATCHES "^(x86_64|AMD64)$")))
+	add_executable(macos_sse4a_tests EXCLUDE_FROM_ALL
+		"${KYTY_TESTS_DIR}/macosSse4aTests.cpp"
+		"${KYTY_SOURCE_DIR}/loader/x64InstructionEmulator.cpp"
+	)
+	target_link_libraries(macos_sse4a_tests common fmt::fmt Zydis::Zydis)
+	target_include_directories(macos_sse4a_tests PRIVATE ${inc_headers})
+	if(BUILD_TESTING)
+		add_test(NAME macos_sse4a COMMAND $<TARGET_FILE:macos_sse4a_tests>)
+		set_tests_properties(macos_sse4a PROPERTIES TIMEOUT 15)
+		add_dependencies(kyty_tests macos_sse4a_tests)
+	endif()
+endif()
 
 add_executable(kyty_emulator "${KYTY_SOURCE_DIR}/main.cpp" ${kyty_emulator_src})
 configure_macos_guest_address_space(kyty_emulator)

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/guest_gpu/graphicsRun.cpp

@@ -386,8 +386,12 @@ void CommandProcessor::WriteReferenceClock(uint64_t dst_address, uint32_t num_by
 	}
 	const auto value = Sync::ReadReferenceClock();
 	std::memcpy(reinterpret_cast<void*>(dst_address), &value, num_bytes);
-	LOGF("\t copy_data reference clock: dst=0x%016" PRIx64 " value=0x%016" PRIx64 " size=%u\n",
-	     dst_address, value, num_bytes);
+	static std::atomic<uint32_t> clock_log_count {0};
+	if (clock_log_count.fetch_add(1) < 64) {
+		LOGF("\t copy_data reference clock: dst=0x%016" PRIx64 " value=0x%016" PRIx64
+		     " size=%u\n",
+		     dst_address, value, num_bytes);
+	}
 }
 
 void CommandProcessor::DmaData(uint8_t engine, uint8_t dst_sel, uint8_t dst_cache_policy,

src/graphics/guest_gpu/tile.h

@@ -86,6 +86,7 @@ struct TileMipLayout {
 	uint32_t padded_height = 0;
 	uint32_t tail_x        = 0;
 	uint32_t tail_y        = 0;
+	bool     operator==(const TileMipLayout&) const = default;
 };
 
 struct TileSurfaceLayout {

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

@@ -208,6 +208,10 @@ bool TextureCache::SameBacking(const ImageInfo& cached, const ImageInfo& request
 	if (cached.extent != requested.extent) {
 		return false;
 	}
+	if (cached.resources.levels < requested.resources.levels ||
+	    cached.resources.layers < requested.resources.layers) {
+		return false;
+	}
 	if (cached.samples != requested.samples) {
 		return false;
 	}
@@ -843,6 +847,16 @@ TextureCache::OverlapResult TextureCache::ResolveOverlap(const ImageInfo& reques
 		    ImageViewOps::FormatsCompatible(cached.info.pixel_format, requested.pixel_format)) {
 			return {ExpandImage(requested, cached_id)};
 		}
+		// PS5 mip tails can expose more levels without increasing the guest allocation.
+		if (requested.pixel_format == cached.info.pixel_format &&
+		    requested.type == cached.info.type && requested.resources > cached.info.resources &&
+		    (requested.data.size > cached.info.data.size ||
+		     (requested.data.size == cached.info.data.size &&
+		      requested.extent == cached.info.extent &&
+		      cached.info.resources.levels > 1 &&
+		      requested.resources.layers == cached.info.resources.layers))) {
+			return {ExpandImage(requested, cached_id)};
+		}
 		if (requested.pixel_format != cached.info.pixel_format ||
 		    requested.data.size <= cached.info.data.size) {
 			const auto result_id = merged_id ? merged_id : cached_id;
@@ -852,9 +866,6 @@ TextureCache::OverlapResult TextureCache::ResolveOverlap(const ImageInfo& reques
 			            ? result_id
 			            : ImageId {}};
 		}
-		if (requested.type == cached.info.type && requested.resources > cached.info.resources) {
-			return {ExpandImage(requested, cached_id)};
-		}
 		EXIT("TextureCache: unresolvable equal-address image overlap, address=0x%016" PRIx64
 		     " requested=%ux%u "
 		     "cached=%ux%u requested_size=0x%016" PRIx64 " cached_size=0x%016" PRIx64
@@ -1226,7 +1237,8 @@ void TextureCache::AssociateStencil(ImageId depth_id, GuestRange stencil) {
 	ImageId association {};
 	for (const auto id: FindImagesInRegion(stencil.address, stencil.size, false)) {
 		const auto owner = m_slot_images.try_get(id);
-		if (owner != nullptr && owner->info.data.address == stencil.address) {
+		if (owner != nullptr && owner->info.data == stencil &&
+		    owner->info.extent == depth.info.extent) {
 			association = id;
 		}
 	}

src/graphics/host_gpu/renderer/depthRenderTarget.cpp

@@ -159,8 +159,8 @@ static TextureCache::ImageDesc MakeDepthTargetDesc(const CommandBuffer& buffer,
 			DepthFatal("invalid depth view: base=%u last=%u", z.depth_view.slice_start,
 			           z.depth_view.slice_max);
 	}
-	if (z.z_info.expclear_enabled || z.stencil_info.expclear_enabled ||
-	    z.z_info.partially_resident ||
+	// EXPCLEAR permits an HTile acceleration state; the host attachment is already expanded.
+	if (z.z_info.partially_resident ||
 	    z.stencil_info.partially_resident || z.z_info.max_mip_level != 0 ||
 	    z.depth_view.current_mip_level != 0 || unsupported_shading_rate_encoding ||
 	    depth_address == 0 || (depth_address & 0xffffu) != 0) {

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

@@ -159,6 +159,7 @@ struct ImageViewInfo {
 	vk::ImageAspectFlags aspect      = vk::ImageAspectFlagBits::eColor;
 	uint32_t             base_level  = 0;
 	uint32_t             level_count = 1;
+	uint32_t             min_lod     = 0; // U4.8 clamp relative to base_level.
 	uint32_t             base_layer  = 0;
 	uint32_t             layer_count = 1;
 	vk::ComponentMapping mapping     = {};
@@ -167,9 +168,10 @@ struct ImageViewInfo {
 	[[nodiscard]] bool operator==(const ImageViewInfo& rhs) const noexcept {
 		return format == rhs.format && type == rhs.type && aspect == rhs.aspect &&
 		       base_level == rhs.base_level && level_count == rhs.level_count &&
-		       base_layer == rhs.base_layer && layer_count == rhs.layer_count &&
-		       mapping.r == rhs.mapping.r && mapping.g == rhs.mapping.g &&
-		       mapping.b == rhs.mapping.b && mapping.a == rhs.mapping.a && usage == rhs.usage;
+		       min_lod == rhs.min_lod && base_layer == rhs.base_layer &&
+		       layer_count == rhs.layer_count && mapping.r == rhs.mapping.r &&
+		       mapping.g == rhs.mapping.g && mapping.b == rhs.mapping.b &&
+		       mapping.a == rhs.mapping.a && usage == rhs.usage;
 	}
 };
 

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

@@ -362,6 +362,12 @@ vk::ImageView Image::FindView(const ImageViewInfo& view_info) {
 	if (!is_storage) {
 		usage.usage &= ~vk::ImageUsageFlagBits::eStorage;
 	}
+	vk::ImageViewMinLodCreateInfoEXT min_lod {};
+	if (normalized.min_lod != 0) {
+		min_lod.minLod = static_cast<float>(normalized.base_level) +
+		                 static_cast<float>(normalized.min_lod) / 256.0f;
+		usage.pNext    = &min_lod;
+	}
 	vk::ImageViewCreateInfo create {};
 	create.pNext                           = &usage;
 	create.image                           = image.image;

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

@@ -293,13 +293,8 @@ static bool IsSupportedStorageTextureDescriptor(const ShaderRecompiler::IR::Imag
 	const bool supported_swizzle =
 	    IsValidImageSwizzle(swizzle) &&
 	    (swizzle == DstSel(4, 5, 6, 7) || !resource.read || resource.atomic);
-	const auto max_mip = resource.r128 ? descriptor.LastLevel() : descriptor.MaxMip();
-	const auto view_last_level =
-	    resource.mip_mode == ShaderRecompiler::IR::ImageMipMode::DynamicStorage
-	        ? descriptor.LastLevel()
-	        : std::min(descriptor.LastLevel(), max_mip);
 	return (is_1d || is_1d_array || is_2d || is_2d_array || is_3d) && supported_tile &&
-	       descriptor.BaseLevel() <= view_last_level && view_last_level <= max_mip &&
+	       descriptor.BaseLevel() <= descriptor.LastLevel() &&
 	       descriptor.MinLod() == 0 && supported_swizzle && descriptor.BCSwizzle() == 0 &&
 	       !descriptor.MsaaDepth();
 }
@@ -457,6 +452,14 @@ static ImageViewInfo TextureViewInfo(const ShaderRecompiler::IR::ImageResource&
 	view.aspect      = vk::ImageAspectFlagBits::eColor;
 	view.base_level  = descriptor.BaseLevel();
 	view.level_count = view_levels;
+	if (descriptor.MinLod() > descriptor.LastLevel() * 256u) {
+		EXIT("texture minimum LOD exceeds last mip level: min_lod=%u last_level=%u\n",
+		     descriptor.MinLod(), descriptor.LastLevel());
+	}
+	const auto base_lod = view.base_level * 256u;
+	if (descriptor.MinLod() > base_lod) {
+		view.min_lod = descriptor.MinLod() - base_lod;
+	}
 	view.usage = storage ? vk::ImageUsageFlagBits::eStorage : vk::ImageUsageFlagBits::eSampled;
 	view.mapping =
 	    storage || surface_format.conversion_format != Prospero::BufferFormat::kInvalid
@@ -507,6 +510,21 @@ static ImageViewInfo TextureViewInfo(const ShaderRecompiler::IR::ImageResource&
 	return view;
 }
 
+static bool TextureViewPreservesMipLayout(const TileSurfaceDescription& description,
+                                          uint32_t                      view_levels) {
+	TileSurfaceLayout physical {};
+	TileSurfaceLayout view {};
+	auto              view_description = description;
+	view_description.levels            = view_levels;
+	return TileGetTiledTextureLayout(description, physical) &&
+	       TileGetTiledTextureLayout(view_description, view) &&
+	       physical.first_tail_level == view.first_tail_level &&
+	       physical.block_slice_size == view.block_slice_size &&
+	       physical.total_size == view.total_size &&
+	       std::equal(std::begin(physical.mips), std::begin(physical.mips) + description.levels,
+	                  std::begin(view.mips));
+}
+
 TextureBinding RenderExecutor::ResolveTexture(const ShaderRecompiler::IR::ImageResource&   resource,
                                               const ShaderRecompiler::IR::DescriptorValue& value) {
 	auto descriptor = DecodeNativeDescriptor<ShaderTextureResource>(value);
@@ -523,24 +541,22 @@ TextureBinding RenderExecutor::ResolveTexture(const ShaderRecompiler::IR::ImageR
 		return {id, nullptr, std::move(desc)};
 	}
 
-	const auto address      = descriptor.Base40();
-	const auto width        = static_cast<uint32_t>(descriptor.Width5()) + 1u;
-	const auto height       = static_cast<uint32_t>(descriptor.Height5()) + 1u;
-	const auto base_level   = descriptor.BaseLevel();
-	const auto last_level   = descriptor.LastLevel();
-	const auto type         = TextureType(descriptor);
-	const bool multisampled = IsMultisampledTexture(type);
-	const auto max_mip      = resource.r128 ? last_level : descriptor.MaxMip();
-	const auto levels       = multisampled ? 1u : static_cast<uint32_t>(max_mip) + 1u;
-	const bool dynamic_storage =
-	    storage && resource.mip_mode == ShaderRecompiler::IR::ImageMipMode::DynamicStorage;
-	const auto view_last_level =
-	    !multisampled && !dynamic_storage ? std::min(last_level, max_mip) : last_level;
+	const auto address         = descriptor.Base40();
+	const auto width           = static_cast<uint32_t>(descriptor.Width5()) + 1u;
+	const auto height          = static_cast<uint32_t>(descriptor.Height5()) + 1u;
+	const auto base_level      = descriptor.BaseLevel();
+	const auto last_level      = descriptor.LastLevel();
+	const auto type            = TextureType(descriptor);
+	const bool multisampled    = IsMultisampledTexture(type);
+	const auto max_mip         = resource.r128 ? last_level : descriptor.MaxMip();
+	const auto physical_levels = multisampled ? 1u : static_cast<uint32_t>(max_mip) + 1u;
+	const auto levels =
+	    multisampled ? 1u : std::max(physical_levels, static_cast<uint32_t>(last_level) + 1u);
 	const auto tile       = descriptor.TileMode();
 	const bool depth_tile = tile == Prospero::TileMode::kDepth;
 	const bool msaa_tile  = depth_tile || tile == Prospero::TileMode::kRenderTarget;
 	const bool msaa_array = type == Prospero::ImageType::kColor2DMsaaArray;
-	if ((!multisampled && (base_level > view_last_level || view_last_level >= levels)) ||
+	if ((!multisampled && base_level > last_level) ||
 	    (multisampled &&
 	     (base_level != 0 || last_level == 0 || last_level > 3 || max_mip != last_level ||
 	      !msaa_tile || (descriptor.MsaaDepth() && !depth_tile) ||
@@ -559,7 +575,7 @@ TextureBinding RenderExecutor::ResolveTexture(const ShaderRecompiler::IR::ImageR
 	}
 	const auto samples = multisampled ? 1u << last_level : 1u;
 	const auto view_levels =
-	    multisampled ? 1u : static_cast<uint32_t>(view_last_level - base_level) + 1u;
+	    multisampled ? 1u : static_cast<uint32_t>(last_level - base_level) + 1u;
 	const auto depth          = static_cast<uint32_t>(descriptor.Depth()) + 1u;
 	const auto format         = descriptor.Format();
 	const auto surface_format = TextureGetSurfaceFormatInfo(format);
@@ -578,7 +594,19 @@ TextureBinding RenderExecutor::ResolveTexture(const ShaderRecompiler::IR::ImageR
 	                             type == Prospero::ImageType::kColor2DArray ||
 	                             type == Prospero::ImageType::kColor2DMsaaArray;
 	const auto    image_layers = layered ? depth : 1u;
-	uint32_t      pitch        = 0;
+	if (levels > physical_levels) {
+		const TileSurfaceDescription physical {
+		    format, tile, volume ? TileSurfaceDimension::Dim3D : TileSurfaceDimension::Dim2D,
+		    width, height, volume ? depth : 1u, physical_levels, image_layers};
+		// Texture mip views take precedence over the resource count, but must keep its storage layout.
+		if (!TextureViewPreservesMipLayout(physical, levels)) {
+			EXIT("unsupported texture mip view changes physical layout: base=%u last=%u max=%u "
+			     "extent=%ux%ux%u tile=%u\n",
+			     base_level, last_level, max_mip, width, height, depth,
+			     static_cast<uint32_t>(tile));
+		}
+	}
+	uint32_t      pitch = 0;
 	TileSizeAlign size {};
 	if (multisampled) {
 		const auto bytes = Prospero::NumBytesPerElement(format);
@@ -591,8 +619,8 @@ TextureBinding RenderExecutor::ResolveTexture(const ShaderRecompiler::IR::ImageR
 		size.size *= image_layers;
 	} else {
 		pitch = TileGetTexturePitch(format, width, tile);
-		TileGetTextureTotalSize(format, width, height, volume ? depth : image_layers, levels, tile,
-		                        volume, size);
+		TileGetTextureTotalSize(format, width, height, volume ? depth : image_layers,
+		                        physical_levels, tile, volume, size);
 	}
 	EXIT_NOT_IMPLEMENTED(size.size == 0 || size.align == 0 ||
 	                     (address & (static_cast<uint64_t>(size.align) - 1u)) != 0);
@@ -629,7 +657,7 @@ TextureBinding RenderExecutor::ResolveTexture(const ShaderRecompiler::IR::ImageR
 	    !desc.info.IsDepth()) {
 		TileSizeAlign metadata_size {};
 		(void)TileGetDccSize(width, height, volume ? depth : image_layers,
-		                     desc.info.bytes_per_block, levels, tile, metadata_size,
+		                     desc.info.bytes_per_block, physical_levels, tile, metadata_size,
 		                     std::countr_zero(samples));
 		desc.info.metadata.kind          = ImageMetadataKind::Dcc;
 		desc.info.metadata.range         = {descriptor.MetaAddr() << 8u, metadata_size.size};

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

@@ -666,13 +666,16 @@ PipelineCache::Pipeline& PipelineCache::GetGraphicsPipeline(
 	key.ps_shader_id            = ps_id;
 	auto& static_params         = key.static_params;
 	auto& rendering             = key.rendering;
-	rendering.color_count       = color_count;
+	rendering.color_count       = 0;
 	uint32_t attachment_samples = 0;
 	for (uint32_t i = 0; i < color_count; i++) {
+		const auto slot = colors[i].target_slot;
+		EXIT_IF(slot >= RENDER_COLOR_ATTACHMENTS_MAX);
+		rendering.color_count = std::max(rendering.color_count, slot + 1);
 		EXIT_IF(!colors[i].image_id || colors[i].desc.view_info.format == vk::Format::eUndefined);
-		static_params.color_mask[i] = colors[i].export_mapping.ApplyMask(
+		static_params.color_mask[slot] = colors[i].export_mapping.ApplyMask(
 		    render_target_mask_slot(ctx.GetRenderTargetMask(), colors[i].target_slot));
-		rendering.color_formats[i] = colors[i].desc.view_info.format;
+		rendering.color_formats[slot] = colors[i].desc.view_info.format;
 		if (attachment_samples == 0) {
 			attachment_samples = colors[i].desc.info.samples;
 		} else if (attachment_samples != colors[i].desc.info.samples) {
@@ -681,14 +684,14 @@ PipelineCache::Pipeline& PipelineCache::GetGraphicsPipeline(
 		}
 		const auto& rt                        = ctx.GetRenderTarget(colors[i].target_slot);
 		const auto& bc                        = ctx.GetBlendControl(colors[i].target_slot);
-		static_params.color_srcblend[i]       = bc.color_srcblend;
-		static_params.color_comb_fcn[i]       = bc.color_comb_fcn;
-		static_params.color_destblend[i]      = bc.color_destblend;
-		static_params.alpha_srcblend[i]       = bc.alpha_srcblend;
-		static_params.alpha_comb_fcn[i]       = bc.alpha_comb_fcn;
-		static_params.alpha_destblend[i]      = bc.alpha_destblend;
-		static_params.separate_alpha_blend[i] = bc.separate_alpha_blend;
-		static_params.blend_enable[i]         = bc.enable && !rt.info.blend_bypass;
+		static_params.color_srcblend[slot]       = bc.color_srcblend;
+		static_params.color_comb_fcn[slot]       = bc.color_comb_fcn;
+		static_params.color_destblend[slot]      = bc.color_destblend;
+		static_params.alpha_srcblend[slot]       = bc.alpha_srcblend;
+		static_params.alpha_comb_fcn[slot]       = bc.alpha_comb_fcn;
+		static_params.alpha_destblend[slot]      = bc.alpha_destblend;
+		static_params.separate_alpha_blend[slot] = bc.separate_alpha_blend;
+		static_params.blend_enable[slot]         = bc.enable && !rt.info.blend_bypass;
 	}
 	const bool with_depth =
 	    depth.desc.view_info.format != vk::Format::eUndefined && static_cast<bool>(depth.image_id);

src/graphics/host_gpu/renderer/renderCompute.cpp

@@ -206,6 +206,17 @@ void RenderExecutor::DispatchDirect(uint64_t submit_id, CommandBuffer& buffer,
 	                    thread_group_x, thread_group_y, thread_group_z, mode,
 	                    sh_ctx.GetCs().cs_regs.data_addr);
 
+	if (thread_group_x == 0 || thread_group_y == 0 || thread_group_z == 0) {
+		static std::atomic<uint32_t> log_count {0};
+		if (log_count.fetch_add(1, std::memory_order_relaxed) < 32) {
+			LOGF("GraphicsRenderDispatchDirect: skipping zero-sized dispatch groups=%ux%ux%u "
+			     "mode=0x%08" PRIx32 " shader=0x%016" PRIx64 "\n",
+			     thread_group_x, thread_group_y, thread_group_z, mode,
+			     sh_ctx.GetCs().cs_regs.data_addr);
+		}
+		return;
+	}
+
 	Common::LockGuard lock(m_context.GetMutex());
 	if (sh_ctx.GetCs().cs_regs.data_addr == 0) {
 		LOGF("GraphicsRenderDispatchDirect: temporary: ignoring dispatch with null CS shader, "
@@ -347,17 +358,6 @@ void RenderExecutor::DispatchDirect(uint64_t submit_id, CommandBuffer& buffer,
 		}
 	}
 
-	if (thread_group_x == 0 || thread_group_y == 0 || thread_group_z == 0) {
-		static std::atomic<uint32_t> log_count {0};
-		if (log_count.fetch_add(1, std::memory_order_relaxed) < 32) {
-			LOGF("GraphicsRenderDispatchDirect: skipping zero-sized dispatch groups=%ux%ux%u "
-			     "mode=0x%08" PRIx32 " shader=0x%016" PRIx64 "\n",
-			     thread_group_x, thread_group_y, thread_group_z, mode,
-			     sh_ctx.GetCs().cs_regs.data_addr);
-		}
-		return;
-	}
-
 	buffer.EndRendering();
 	auto& pipeline =
 	    m_context.GetPipelineCache().GetComputePipeline(input_info, compute_program);

src/graphics/host_gpu/renderer/renderDraw.cpp

@@ -15,6 +15,7 @@
 #include "graphics/host_gpu/renderer/colorRenderTarget.h"
 #include "graphics/host_gpu/renderer/debug.h"
 #include "graphics/host_gpu/renderer/depthRenderTarget.h"
+#include "graphics/host_gpu/renderer/image/textureCommon.h"
 #include "graphics/host_gpu/renderer/pipeline/pipelineCache.h"
 #include "graphics/host_gpu/renderer/pipeline/shaderResourceBarrier.h"
 #include "graphics/host_gpu/renderer/render.h"
@@ -308,24 +309,12 @@ static void LogDrawInputState(const CommandBuffer& buffer, const RenderColorInfo
 
 static void SetGraphicsDynamicParams(const CommandBuffer& buffer, vk::CommandBuffer vk_buffer,
                                      const ShaderVertexInputInfo& vs_input_info,
-                                     const RenderColorInfo* colors, uint32_t color_count,
-                                     const RenderDepthInfo& depth) {
+                                     const RenderDepthInfo& depth, const RenderState& rendering) {
 	KYTY_PROFILER_FUNCTION();
 
-	EXIT_IF(colors == nullptr);
 	const auto& ctx = buffer.GetRegisters();
-
-	const auto&  vp = ctx.GetScreenViewport();
-	vk::Extent2D framebuffer_extent {};
-	if (color_count > 0 && colors[0].image_id) {
-		framebuffer_extent = colors[0].Extent();
-	} else if (depth.image_id) {
-		framebuffer_extent = {depth.desc.info.extent.width, depth.desc.info.extent.height};
-	} else {
-		const auto& limits = buffer.GetGraphics().GetPhysicalDeviceProperties().limits;
-		framebuffer_extent = {limits.maxFramebufferWidth, limits.maxFramebufferHeight};
-	}
-
+	const auto&        vp  = ctx.GetScreenViewport();
+	const vk::Extent2D framebuffer_extent {rendering.width, rendering.height};
 	const auto& outputs = vs_input_info.stage.program->info.outputs;
 	const bool  indexed_viewports =
 	    std::any_of(outputs.begin(), outputs.end(), [](const auto& output) {
@@ -420,15 +409,11 @@ static void SetGraphicsDynamicParams(const CommandBuffer& buffer, vk::CommandBuf
 	// eColorWriteEnableEXT dynamic state and relies on the static colorWriteMask instead.
 #else
 	vk::Bool32 enable[RENDER_COLOR_ATTACHMENTS_MAX] = {};
-	// Color-control operation selects special color-buffer paths, not the normal component write
-	// mask. Attachment availability therefore follows the target write mask.
-	for (uint32_t i = 0; i < color_count; i++) {
-		enable[i] = render_target_mask_slot(ctx.GetRenderTargetMask(), colors[i].target_slot) != 0
-		                ? VK_TRUE
-		                : VK_FALSE;
+	for (uint32_t slot = 0; slot < rendering.num_color_attachments; slot++) {
+		enable[slot] = rendering.color_attachments[slot].image_view != nullptr;
 	}
-	if (color_count != 0) {
-		vk_buffer.setColorWriteEnableEXT(color_count, enable);
+	if (rendering.num_color_attachments != 0) {
+		vk_buffer.setColorWriteEnableEXT(rendering.num_color_attachments, enable);
 	}
 #endif
 }
@@ -474,7 +459,7 @@ RenderState RenderExecutor::AcquireRenderTargets(CommandBuffer& buffer, RenderCo
 	state.width                 = std::numeric_limits<uint32_t>::max();
 	state.height                = std::numeric_limits<uint32_t>::max();
 	state.num_layers            = std::numeric_limits<uint32_t>::max();
-	state.num_color_attachments = color_count;
+	state.num_color_attachments = 0;
 	uint32_t attachment_samples = 0;
 	for (uint32_t i = 0; i < color_count; i++) {
 		auto& target = colors[i];
@@ -516,7 +501,8 @@ RenderState RenderExecutor::AcquireRenderTargets(CommandBuffer& buffer, RenderCo
 		state.width             = std::min(state.width, extent.width);
 		state.height            = std::min(state.height, extent.height);
 		state.num_layers        = std::min(state.num_layers, view.layer_count);
-		auto& attachment        = state.color_attachments[i];
+		state.num_color_attachments = std::max(state.num_color_attachments, target.target_slot + 1);
+		auto& attachment            = state.color_attachments[target.target_slot];
 		attachment.image_view   = image_view;
 		attachment.image_layout = layout;
 	}
@@ -891,42 +877,6 @@ static bool ResolvePrimitiveRestart(const CommandBuffer& buffer,
 	return false;
 }
 
-bool RenderExecutor::PrepareDrawRenderState(CommandBuffer& buffer, const DrawCallInfo& draw,
-                                            uint32_t            render_target_slice_offset,
-	                                        DrawRenderState& state) {
-	auto& ctx = buffer.GetRegisters();
-
-	if (ResolveColorTargets(buffer, render_target_slice_offset)) {
-		return false;
-	}
-	if (draw.IsIndexed()) {
-		LogDrawPhase(draw.Name(), "ResolveRenderColorTarget");
-	}
-	for (uint32_t slot = 0; slot < RENDER_COLOR_ATTACHMENTS_MAX; slot++) {
-		if (slot == 0 || (render_target_mask_slot(ctx.GetRenderTargetMask(), slot) != 0 &&
-		                  ctx.GetRenderTarget(slot).base.addr != 0)) {
-			ResolveRenderColorTarget(buffer, state.color_info[state.color_count],
-			                         render_target_slice_offset, slot);
-			if (state.color_info[state.color_count].image_id) {
-				state.color_count++;
-			}
-		}
-	}
-	if (draw.IsIndexed()) {
-		LogDrawPhase(draw.Name(), "ResolveRenderDepthTarget");
-	}
-	ResolveRenderDepthTarget(buffer, state.depth_info);
-
-	state.ps_active       = DrawHasActivePixelShader(buffer);
-	if (state.color_count == 0 && !state.depth_info.image_id && !state.ps_active) {
-		LogFramebufferSkip(draw.Name(), state.color_info[0], state.depth_info, buffer,
-		                   draw.index_count, 0);
-		return false;
-	}
-
-	return true;
-}
-
 static void RefreshShaders(CommandBuffer& buffer, const DrawCallInfo& draw,
                            DrawRenderState& state) {
 	auto& ctx    = buffer.GetRegisters();
@@ -940,8 +890,14 @@ static void RefreshShaders(CommandBuffer& buffer, const DrawCallInfo& draw,
 	state.ps_input_info = {};
 	std::array<Prospero::ColorComponentMapping, RENDER_COLOR_ATTACHMENTS_MAX>
 	    target_export_mapping {};
-	for (uint32_t i = 0; i < state.color_count; i++) {
-		target_export_mapping[state.color_info[i].target_slot] = state.color_info[i].export_mapping;
+	for (uint32_t slot = 0; slot < RENDER_COLOR_ATTACHMENTS_MAX; slot++) {
+		const auto& rt = ctx.GetRenderTarget(slot);
+		if (rt.base.addr != 0 && render_target_mask_slot(ctx.GetRenderTargetMask(), slot) != 0) {
+			target_export_mapping[slot] =
+			    TextureGetRenderTargetFormat(rt.info.format, rt.info.channel_type,
+			                                 rt.info.channel_order)
+			        .export_mapping;
+		}
 	}
 	auto& pipeline_cache = buffer.GetContext().GetPipelineCache();
 	if (draw.IsIndexed()) {
@@ -952,6 +908,45 @@ static void RefreshShaders(CommandBuffer& buffer, const DrawCallInfo& draw,
 	    target_export_mapping, state.ps_active, state.vertex_info, state.ps_input_info);
 }
 
+bool RenderExecutor::PrepareDrawRenderState(CommandBuffer& buffer, const DrawCallInfo& draw,
+                                            uint32_t            render_target_slice_offset,
+	                                        DrawRenderState& state) {
+	state.ps_active = DrawHasActivePixelShader(buffer);
+	RefreshShaders(buffer, draw, state);
+	uint32_t mrt_mask = 0;
+	if (state.ps_active) {
+		for (const auto& output: state.ps_input_info.stage.program->info.outputs) {
+			if (output.kind == ShaderRecompiler::IR::StageOutputKind::Mrt) {
+				mrt_mask |= 1u << output.index;
+			}
+		}
+	}
+	if (draw.IsIndexed()) {
+		LogDrawPhase(draw.Name(), "ResolveRenderColorTarget");
+	}
+	for (uint32_t slot = 0; slot < RENDER_COLOR_ATTACHMENTS_MAX; slot++) {
+		if ((mrt_mask & (1u << slot)) != 0) {
+			ResolveRenderColorTarget(buffer, state.color_info[state.color_count],
+			                         render_target_slice_offset, slot);
+			if (state.color_info[state.color_count].image_id) {
+				state.color_count++;
+			}
+		}
+	}
+	if (draw.IsIndexed()) {
+		LogDrawPhase(draw.Name(), "ResolveRenderDepthTarget");
+	}
+	ResolveRenderDepthTarget(buffer, state.depth_info);
+
+	if (state.color_count == 0 && !state.depth_info.image_id && !state.ps_active) {
+		LogFramebufferSkip(draw.Name(), state.color_info[0], state.depth_info, buffer,
+		                   draw.index_count, 0);
+		return false;
+	}
+
+	return true;
+}
+
 static PreparedIndexBuffer PrepareIndexBuffer(CommandBuffer&               buffer,
                                               const DrawIndexBufferSource& source) {
 	PreparedIndexBuffer prepared;
@@ -1073,7 +1068,13 @@ void RenderExecutor::ExecutePreparedDraw(uint64_t submit_id, CommandBuffer& buff
 	uint32_t   mesh_groups = 0;
 	if (mesh_active) {
 		const auto& mesh = state.vertex_info[0].mesh;
-		if (primitive_restart_enable || mesh.primitives_per_group == 0) {
+		static std::atomic_bool restart_warned = false;
+		if (primitive_restart_enable && !restart_warned.exchange(true, std::memory_order_relaxed)) {
+			std::printf("Warning: primitive restart is not implemented for mesh shaders; "
+			            "continuing draw (primitive=%u indexed=%u)\n",
+			            static_cast<uint32_t>(ucfg.GetPrimType()), draw.IsIndexed());
+		}
+		if (mesh.primitives_per_group == 0) {
 			EXIT("unsupported mesh draw: primitive=%u indexed=%u restart=%u\n",
 			     static_cast<uint32_t>(ucfg.GetPrimType()), draw.IsIndexed(), primitive_restart_enable);
 		}
@@ -1160,8 +1161,7 @@ void RenderExecutor::ExecutePreparedDraw(uint64_t submit_id, CommandBuffer& buff
 		CommitIndexBuffer(vk_buffer, index_binding);
 	}
 
-	SetGraphicsDynamicParams(buffer, vk_buffer, vertex_stages.back(), state.color_info,
-	                         state.color_count, state.depth_info);
+	SetGraphicsDynamicParams(buffer, vk_buffer, vertex_stages.back(), state.depth_info, rendering);
 	if (m_context.GetGraphics().attachment_feedback_loop_enabled) {
 		vk_buffer.setAttachmentFeedbackLoopEnableEXT(
 		    rendering.depth_stencil_attachment.image_layout ==
@@ -1226,7 +1226,8 @@ void RenderExecutor::DrawIndex(uint64_t submit_id, CommandBuffer& buffer,
 		return;
 	}
 
-	if (ConsumeMetadataColorOperation(buffer) || DepthStencilCopy(buffer)) {
+	if (ConsumeMetadataColorOperation(buffer) || DepthStencilCopy(buffer) ||
+	    ResolveColorTargets(buffer, args.render_target_slice_offset)) {
 		ResetBindings();
 		return;
 	}
@@ -1300,8 +1301,6 @@ void RenderExecutor::DrawIndex(uint64_t submit_id, CommandBuffer& buffer,
 		return;
 	}
 
-	RefreshShaders(buffer, draw, state);
-
 	LogDrawStateIfNeeded(buffer, draw, state, args.index_type_and_size,
 	                     args.index_addr);
 
@@ -1338,7 +1337,8 @@ void RenderExecutor::DrawAuto(uint64_t submit_id, CommandBuffer& buffer, const D
 		return;
 	}
 
-	if (ConsumeMetadataColorOperation(buffer) || DepthStencilCopy(buffer)) {
+	if (ConsumeMetadataColorOperation(buffer) || DepthStencilCopy(buffer) ||
+	    ResolveColorTargets(buffer, args.render_target_slice_offset)) {
 		ResetBindings();
 		return;
 	}
@@ -1367,18 +1367,16 @@ void RenderExecutor::DrawAuto(uint64_t submit_id, CommandBuffer& buffer, const D
 	const DrawCallInfo draw {CommandBufferDebugOp::DrawIndexAuto,
 	                         args.vertex_count, args.instance_count, args.first_instance};
 
-	DrawRenderState state {};
-	if (!PrepareDrawRenderState(buffer, draw, args.render_target_slice_offset, state)) {
-		ResetBindings();
-		return;
-	}
-
 	vk::PrimitiveTopology topology = vk::PrimitiveTopology::ePointList;
 	if (!GetDrawTopology(ucfg, true, topology)) {
 		ResetBindings();
 		return;
 	}
-	RefreshShaders(buffer, draw, state);
+	DrawRenderState state {};
+	if (!PrepareDrawRenderState(buffer, draw, args.render_target_slice_offset, state)) {
+		ResetBindings();
+		return;
+	}
 
 	const bool rect_list = ucfg.GetPrimType() == Prospero::PrimitiveType::kRectList;
 	if (rect_list && state.vertex_info[0].buffers_num == 0 &&

src/graphics/presentation/window/vulkanWindow.cpp

@@ -201,6 +201,8 @@ static void VulkanFindPhysicalDevice(vk::Instance instance, vk::SurfaceKHR surfa
 		vk::PhysicalDeviceVulkan13Features features13 {};
 
 		vk::PhysicalDeviceColorWriteEnableFeaturesEXT color_write_ext {};
+		vk::PhysicalDeviceImageViewMinLodFeaturesEXT  image_view_min_lod {};
+		color_write_ext.pNext = &image_view_min_lod;
 
 		vk::PhysicalDeviceDepthClipEnableFeaturesEXT depth_clip_enable {};
 		depth_clip_enable.pNext = &color_write_ext;
@@ -235,6 +237,10 @@ static void VulkanFindPhysicalDevice(vk::Instance instance, vk::SurfaceKHR surfa
 			skip_device = true;
 #endif
 		}
+		if (image_view_min_lod.minLod != VK_TRUE) {
+			LOGF("image view minLod is not supported\n");
+			skip_device = true;
+		}
 
 		if (depth_clip_control.depthClipControl != VK_TRUE) {
 			LOGF("depthClipControl is not supported\n");
@@ -514,11 +520,14 @@ static vk::Device VulkanCreateDevice(GraphicContext& graphics,
 	depth_clip_enable.depthClipEnable = VK_TRUE;
 
 	vk::PhysicalDeviceDepthClipControlFeaturesEXT depth_clip_control {};
+	vk::PhysicalDeviceImageViewMinLodFeaturesEXT  image_view_min_lod {};
+	image_view_min_lod.minLod = VK_TRUE;
+	depth_clip_control.pNext  = &image_view_min_lod;
 	// MoltenVK lacks VK_EXT_depth_clip_enable and VK_EXT_color_write_enable, so drop those
 	// feature structs from the chain on macOS (the renderer falls back to default depth
 	// clipping and static color-write masks).
 #if !defined(__APPLE__)
-	depth_clip_control.pNext = &depth_clip_enable;
+	image_view_min_lod.pNext = &depth_clip_enable;
 #endif
 	depth_clip_control.depthClipControl = VK_TRUE;
 
@@ -984,7 +993,8 @@ void WindowContext::CreateVulkan() {
 
 	std::vector<const char*> device_extensions = {
 	    VK_KHR_SWAPCHAIN_EXTENSION_NAME, VK_EXT_DEPTH_CLIP_CONTROL_EXTENSION_NAME,
-	    VK_KHR_PUSH_DESCRIPTOR_EXTENSION_NAME, "VK_KHR_maintenance1"};
+	    VK_EXT_IMAGE_VIEW_MIN_LOD_EXTENSION_NAME, VK_KHR_PUSH_DESCRIPTOR_EXTENSION_NAME,
+	    "VK_KHR_maintenance1"};
 
 #if defined(__APPLE__)
 	// MoltenVK lacks VK_EXT_depth_clip_enable and VK_EXT_color_write_enable; the renderer

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

@@ -3,7 +3,9 @@
 #include "graphics/shader/recompiler/frontend/decode/ImageOps.h"
 
 #include <algorithm>
+#include <atomic>
 #include <bit>
+#include <cstdio>
 
 namespace Libs::Graphics::ShaderRecompiler::Spirv::Emitter {
 namespace {
@@ -646,6 +648,13 @@ void EmitImage(ValueEmitContext& ctx, const IR::Inst& inst) {
 		const auto coord =
 		    CoordF32(ctx, mem, *address, layout.coord, dimension_info.coordinate_components);
 		if (op == IR::ValueOpcode::ImageGatherRaw) {
+			if (HasFlag(mem, Decoder::ImageSampleFlagLod)) {
+				static std::atomic_flag warned = ATOMIC_FLAG_INIT;
+				if (!warned.test_and_set(std::memory_order_relaxed)) {
+					std::fputs("Warning: approximating IMAGE_GATHER4_L at mip level 0; explicit LOD is ignored.\n",
+					           stderr);
+				}
+			}
 			if (dimension == ImageDimension::Dim1D) {
 				if (dref || !HasFlag(mem, Decoder::ImageSampleFlagLevelZero) ||
 				    HasFlag(mem, Decoder::ImageSampleFlagOffset) ||

src/graphics/shader/recompiler/frontend/cfg/ShaderCFG.cpp

@@ -1206,13 +1206,12 @@ uint32_t FindSelectionMerge(const Graph& graph, const BasicBlock& block) {
 					return false_target;
 				}
 				// A return can leave a selection without reaching its merge. Keep the
-				// continuing arm as the merge instead of joining live state with a return.
-				if (graph.Dominates(block.id, false_target) &&
-				    HasLinearPathToTerminal(graph, true_target)) {
+				// continuing arm as the merge. Shared continuations receive a dedicated
+				// gateway in SplitSharedMergeBlock before structured control is emitted.
+				if (HasLinearPathToTerminal(graph, true_target)) {
 					return false_target;
 				}
-				if (graph.Dominates(block.id, true_target) &&
-				    HasLinearPathToTerminal(graph, false_target)) {
+				if (HasLinearPathToTerminal(graph, false_target)) {
 					return true_target;
 				}
 			}

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

@@ -175,6 +175,7 @@ constexpr MimgSampleInfo MIMG_SAMPLE_OPCODE_LIST[] = {
 };
 
 constexpr MimgGatherInfo MIMG_GATHER_OPCODE_LIST[] = {
+    {0x44u, Opcode::IMAGE_GATHER4_L, ImageSampleFlagLod},
     {0x47u, Opcode::IMAGE_GATHER4_LZ, ImageSampleFlagLevelZero},
     {0x48u, Opcode::IMAGE_GATHER4_C, ImageSampleFlagCompare},
     {0x4fu, Opcode::IMAGE_GATHER4_C_LZ,

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

@@ -24,7 +24,8 @@ constexpr OpcodeMap SOP2_OPCODE_LIST[] = {
     {0x1cu, Opcode::S_XNOR_B32},        {0x1du, Opcode::S_XNOR_B64},
     {0x1eu, Opcode::S_LSHL_B32},        {0x1fu, Opcode::S_LSHL_B64},
     {0x20u, Opcode::S_LSHR_B32},        {0x21u, Opcode::S_LSHR_B64},
-    {0x22u, Opcode::S_ASHR_I32},        {0x24u, Opcode::S_BFM_B32},
+    {0x22u, Opcode::S_ASHR_I32},        {0x23u, Opcode::S_ASHR_I64},
+    {0x24u, Opcode::S_BFM_B32},
     {0x25u, Opcode::S_BFM_B64},         {0x26u, Opcode::S_MUL_I32},
     {0x27u, Opcode::S_BFE_U32},         {0x28u, Opcode::S_BFE_I32},
     {0x29u, Opcode::S_BFE_U64},         {0x2cu, Opcode::S_ABSDIFF_I32},
@@ -105,6 +106,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},
 };
@@ -136,6 +138,10 @@ void DecodeSop1(uint32_t pc, std::span<const uint32_t> code, uint32_t word_index
 	inst.family    = Family::SOP1;
 	inst.opcode_id = opcode;
 	inst.opcode    = Detail::LookupOpcode(SOP1_OPS, opcode);
+	if (opcode == 0x21u && sdst == 125u) {
+		// S_SWAPPC_B64 with NULL discards the return PC, so it is a plain jump.
+		inst.opcode = Opcode::S_SETPC_B64;
+	}
 	SetRawWords(inst, code, word_index, 1);
 
 	if (inst.opcode == Opcode::UNSUPPORTED) {

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

@@ -144,6 +144,7 @@ std::string FormatMimg(const Instruction& inst) {
 	}
 	switch (inst.opcode) {
 		case Opcode::IMAGE_SAMPLE:
+		case Opcode::IMAGE_GATHER4_L:
 		case Opcode::IMAGE_GATHER4_LZ:
 		case Opcode::IMAGE_GATHER4_C:
 		case Opcode::IMAGE_GATHER4_C_LZ:
@@ -187,6 +188,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 +557,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));
@@ -577,6 +581,7 @@ std::string InstructionToString(const Instruction& inst) {
 		case Opcode::IMAGE_LOAD_MIP:
 		case Opcode::IMAGE_GET_RESINFO:
 		case Opcode::IMAGE_GET_LOD:
+		case Opcode::IMAGE_GATHER4_L:
 		case Opcode::IMAGE_GATHER4_LZ:
 		case Opcode::IMAGE_GATHER4_C:
 		case Opcode::IMAGE_GATHER4_C_LZ:

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

@@ -106,6 +106,7 @@ enum class Opcode {
 	S_LSHR_B32,
 	S_LSHR_B64,
 	S_ASHR_I32,
+	S_ASHR_I64,
 	S_MUL_I32,
 	S_MUL_HI_U32,
 	S_MUL_HI_I32,
@@ -559,6 +560,7 @@ enum class Opcode {
 	IMAGE_ATOMIC_OR,
 	IMAGE_ATOMIC_XOR,
 	IMAGE_SAMPLE,
+	IMAGE_GATHER4_L,
 	IMAGE_GATHER4_LZ,
 	IMAGE_GATHER4_C,
 	IMAGE_GATHER4_C_LZ,
@@ -581,6 +583,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/Integer.cpp

@@ -168,6 +168,21 @@ bool Translator::SimpleInteger(const Decoder::Instruction& inst, IR::ValueOpcode
 	return true;
 }
 
+bool Translator::S_ASHR_I64(const Decoder::Instruction& inst) {
+	// Signed 64-bit sources sign-extend literals; generic B64 operands zero-extend them.
+	const auto source =
+	    inst.src0.kind == Decoder::OperandKind::LiteralConstant
+	        ? ir.ConstructU64(
+	              ReadU32(inst.src0),
+	              IR::U32(IR::Value((inst.src0.value & 0x80000000u) ? 0xffffffffu : 0u)))
+	        : ReadU64(inst.src0);
+	const auto result =
+	    ir.Emit(IR::ValueOpcode::ShiftRightArithmetic64, {source, ReadU32(inst.src1)});
+	WriteOperand(inst.dst, result);
+	ir.SetScc(IR::U1(ir.Emit(IR::ValueOpcode::INotEqual64, {result, IR::Value(uint64_t {0})})));
+	return true;
+}
+
 bool Translator::ComposedIntegerBinary(const Decoder::Instruction& inst, IR::ValueOpcode opcode,
                                        bool negate_rhs, bool negate_result, bool update_scc) {
 	const auto lhs = ReadU32(inst.src0);

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

@@ -1034,6 +1034,7 @@ bool Translator::EmitMemory(const Decoder::Instruction& inst) {
 		case Decoder::Opcode::IMAGE_STORE:
 		case Decoder::Opcode::IMAGE_STORE_MIP: return IMAGE_STORE(inst);
 		case Decoder::Opcode::IMAGE_SAMPLE: return IMAGE_SAMPLE(inst);
+		case Decoder::Opcode::IMAGE_GATHER4_L:
 		case Decoder::Opcode::IMAGE_GATHER4_LZ:
 		case Decoder::Opcode::IMAGE_GATHER4_C:
 		case Decoder::Opcode::IMAGE_GATHER4_C_LZ:

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

@@ -179,6 +179,7 @@ bool Translator::EmitScalar(const Decoder::Instruction& inst) {
 		case O::S_LSHR_B64:
 			return SimpleInteger(inst, IR::ValueOpcode::ShiftRightLogical64, IR::Type::U64, false,
 			                     false, true);
+		case O::S_ASHR_I64: return S_ASHR_I64(inst);
 
 		case O::S_ANDN2_B32:
 			return ComposedIntegerBinary(inst, IR::ValueOpcode::BitwiseAnd32, true, false, true);
@@ -227,6 +228,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/recompiler/frontend/translate/Translator.h

@@ -176,6 +176,7 @@ private:
 	                     bool negate_result);
 	bool SimpleInteger(const Decoder::Instruction& inst, IR::ValueOpcode opcode, IR::Type type,
 	                   bool reverse, bool mask_shift_count, bool update_scc);
+	bool S_ASHR_I64(const Decoder::Instruction& inst);
 	bool ComposedIntegerBinary(const Decoder::Instruction& inst, IR::ValueOpcode opcode,
 	                           bool negate_rhs, bool negate_result, bool update_scc);
 	bool V_AND_OR_B32(const Decoder::Instruction& inst);

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

@@ -112,6 +112,10 @@ ResourcePlan& ResourcePlan::operator=(ResourcePlan&& other) noexcept {
 }
 
 Program::~Program() {
+	// Planning expressions can refer to block values but outlive block storage in the base class.
+	for (auto& inst: value_storage) {
+		inst.Invalidate();
+	}
 	// Values may cross block boundaries. Detach all arguments before any block starts destroying
 	// its instruction storage so reverse-use links always point to live definitions.
 	for (auto* block: blocks) {
@@ -184,6 +188,22 @@ Value ResolveInvariantPhi(const ResourcePlan& program, Value value) {
 	return invariant;
 }
 
+bool HasShaderMemoryWrites(const Program& program) {
+	for (const auto* block: program.blocks) {
+		for (const auto& inst: *block) {
+			const auto op     = inst.GetOpcode();
+			const auto buffer = BufferAccessOf(op);
+			const auto image  = ImageOpcodeInfoOf(op).access;
+			if (buffer == BufferAccess::Write || buffer == BufferAccess::Atomic ||
+			    image == ImageAccess::Write || image == ImageAccess::Atomic ||
+			    AddressOpcodeInfoOf(op).access == AddressAccess::Write) {
+				return true;
+			}
+		}
+	}
+	return false;
+}
+
 void ValidateProgram(const Program& program, bool require_ssa) {
 	if (program.blocks.size() != program.block_info.size() ||
 	    program.blocks.size() != program.block_storage.size()) {

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

@@ -563,6 +563,7 @@ struct Program: ResourcePlan {
 };
 
 std::string ProgramToString(const Program& program);
+bool        HasShaderMemoryWrites(const Program& program);
 
 void  ValidateProgram(const Program& program, bool require_ssa);
 void  ResolveControlFlowIdentities(Program& program);

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

@@ -647,7 +647,8 @@ bool BuildSamplerPlan(const ShaderInfo& base, const Images& images, SamplerPlan&
 }
 
 static std::vector<ResourceBlock> ResourceControlFlow(const Program& program) {
-	if (program.blocks.size() != program.block_info.size()) {
+	// Any shader write may alias a scalar predicate read, including on a later loop visit.
+	if (program.blocks.size() != program.block_info.size() || HasShaderMemoryWrites(program)) {
 		return {};
 	}
 	std::unordered_map<uint32_t, uint32_t> indices;
@@ -687,12 +688,6 @@ static std::vector<ResourceBlock> ResourceControlFlow(const Program& program) {
 			const auto op     = inst.GetOpcode();
 			const auto buffer = BufferAccessOf(op);
 			const auto image  = ImageOpcodeInfoOf(op);
-			// Any shader write may alias a scalar predicate read, including on a later loop visit.
-			if (buffer == BufferAccess::Write || buffer == BufferAccess::Atomic ||
-			    image.access == ImageAccess::Write || image.access == ImageAccess::Atomic ||
-			    AddressOpcodeInfoOf(op).access == AddressAccess::Write) {
-				return {};
-			}
 			if (buffer == BufferAccess::None && image.access == ImageAccess::None) {
 				continue;
 			}

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

@@ -88,6 +88,7 @@ public:
 		m_info.samplers.clear();
 		m_info.sampled_pairs.clear();
 		m_info.uses_dma = false;
+		m_shader_writes = HasShaderMemoryWrites(program);
 	}
 
 	void Run() {
@@ -168,15 +169,85 @@ private:
 		std::abort();
 	}
 
+	Value LowerDescriptorPhi(Value value) {
+		value           = value.Resolve();
+		const auto* phi = value.TryInstruction();
+		if (m_shader_writes || phi == nullptr || phi->GetOpcode() != ValueOpcode::Phi ||
+		    phi->NumArgs() != 2u || phi->NumPhiBlocks() != 2u || phi->GetType() != Type::U32 ||
+		    m_program.blocks.size() != m_program.block_info.size()) {
+			return value;
+		}
+		for (const auto& [original, selected]: m_descriptor_selections) {
+			if (original == phi) {
+				return selected;
+			}
+		}
+		const auto* merge  = phi->Parent();
+		const auto* branch = phi->PhiBlock(0);
+		if (merge == nullptr || branch == nullptr || phi->PhiBlock(1) == nullptr ||
+		    branch == phi->PhiBlock(1)) {
+			return value;
+		}
+		if (branch->ImmSuccessors().size() != 2u) {
+			if (branch->ImmPredecessors().size() != 1u) {
+				return value;
+			}
+			branch = branch->ImmPredecessors()[0];
+		}
+		if (branch == merge || branch->ImmSuccessors().size() != 2u) {
+			return value;
+		}
+		std::array<uint32_t, 2> target_ids;
+		for (uint32_t arm = 0; arm < 2; arm++) {
+			const auto* incoming = phi->PhiBlock(arm);
+			if (incoming == merge ||
+			    (incoming != branch &&
+			     (incoming->ImmPredecessors().size() != 1u ||
+			      incoming->ImmPredecessors()[0] != branch ||
+			      incoming->ImmSuccessors().size() != 1u ||
+			      incoming->ImmSuccessors()[0] != merge))) {
+				return value;
+			}
+			const auto* target = incoming == branch ? merge : incoming;
+			const auto  it     = std::ranges::find(m_program.blocks, target);
+			if (it == m_program.blocks.end()) {
+				return value;
+			}
+			target_ids[arm] = m_program.block_info[it - m_program.blocks.begin()].id;
+		}
+		const auto branch_it = std::ranges::find(m_program.blocks, branch);
+		if (branch_it == m_program.blocks.end()) {
+			return value;
+		}
+		const auto& info = m_program.block_info[branch_it - m_program.blocks.begin()];
+		const auto& term = info.terminator;
+		if (term.kind != CFG::TerminatorKind::ConditionalBranch ||
+		    !((term.true_block == target_ids[0] && term.false_block == target_ids[1]) ||
+		      (term.false_block == target_ids[0] && term.true_block == target_ids[1])) ||
+		    !ValidateRuntimeValue(m_program, info.condition, RuntimeValueType::Integer) ||
+		    !ValidateRuntimeValue(m_program, phi->Arg(0)) ||
+		    !ValidateRuntimeValue(m_program, phi->Arg(1))) {
+			return value;
+		}
+		// Retain a host expression; replacing the GPU Phi would break SSA dominance.
+		const auto true_arg = term.true_block == target_ids[0] ? 0u : 1u;
+		auto&      selected = m_program.value_storage.emplace_back(ValueOpcode::SelectU32);
+		selected.SetArg(0, info.condition);
+		selected.SetArg(1, phi->Arg(true_arg));
+		selected.SetArg(2, phi->Arg(true_arg ^ 1u));
+		m_descriptor_selections.emplace_back(phi, Value(&selected));
+		return Value(&selected);
+	}
+
 	void MakeSource(const Inst& handle, uint32_t width, bool sampler, bool sample_adjust,
-	                DescriptorSource& descriptor, uint32_t pc) const {
+	                DescriptorSource& descriptor, uint32_t pc) {
 		if (handle.NumArgs() != width) {
 			Fail(pc, fmt::format("{} has {} descriptor dwords, expected {}",
 			                     ValueOpcodeName(handle.GetOpcode()), handle.NumArgs(), width));
 		}
 		descriptor.dword_count = width;
 		for (uint32_t i = 0; i < width; i++) {
-			descriptor.dwords[i] = handle.Arg(i).Resolve();
+			descriptor.dwords[i] = LowerDescriptorPhi(handle.Arg(i));
 		}
 		if (sample_adjust) {
 			descriptor.dwords[3] = CanonicalizeSampleAdjustDword3(descriptor.dwords[3]);
@@ -744,12 +815,14 @@ private:
 		}
 	}
 
-	Program&                       m_program;
-	ShaderInfo                     m_info;
-	std::vector<DescriptorSource>  m_sources;
-	std::vector<HandlePatch>       m_handle_patches;
-	std::vector<MemoryPatch>       m_memory_patches;
-	std::vector<IndirectImagePlan> m_indirect_images;
+	Program&                                   m_program;
+	ShaderInfo                                 m_info;
+	std::vector<DescriptorSource>              m_sources;
+	std::vector<HandlePatch>                   m_handle_patches;
+	std::vector<MemoryPatch>                   m_memory_patches;
+	std::vector<IndirectImagePlan>             m_indirect_images;
+	std::vector<std::pair<const Inst*, Value>> m_descriptor_selections;
+	bool                                       m_shader_writes = false;
 };
 
 } // namespace

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

@@ -16,6 +16,13 @@ namespace {
 
 constexpr uint64_t AddressMask = 0x0000ffffffffffffull;
 
+SrtRuntime CleanRuntime(SrtRuntime runtime) {
+	runtime.read_memory = runtime.read_specialization_memory != nullptr
+	                          ? runtime.read_specialization_memory
+	                          : +[](void*, uint64_t, uint32_t*) { return false; };
+	return runtime;
+}
+
 const char* StageName(ShaderType stage) {
 	switch (stage) {
 		case ShaderType::Vertex: return "vertex";
@@ -646,8 +653,10 @@ private:
 			case ValueOpcode::GetShaderBase: result = m_runtime.shader_base; return true;
 			case ValueOpcode::Phi: return EvaluatePhi(inst, result);
 			case ValueOpcode::ReadFirstLane: {
-				Evaluator active(m_program, m_runtime, m_clean_flat_slots, m_clean_evaluator,
-				                 inst.Arg(1));
+				const auto clean_runtime = CleanRuntime(m_runtime);
+				Evaluator  clean_active(m_program, clean_runtime, {}, nullptr, inst.Arg(1));
+				Evaluator  active(m_program, m_runtime, m_clean_flat_slots, &clean_active,
+				                  inst.Arg(1));
 				return active.EvaluateWide(inst.Arg(0), result);
 			}
 			case ValueOpcode::BitCastU32F32:
@@ -892,12 +901,13 @@ private:
 			}
 			case ValueOpcode::SelectU32:
 			case ValueOpcode::SelectU1:
-			case ValueOpcode::SelectF32:
-				if (ternary()) {
-					result = a != 0u ? b : c;
-					return true;
+			case ValueOpcode::SelectF32: {
+				auto& predicate = m_clean_evaluator != nullptr ? *m_clean_evaluator : *this;
+				if (predicate.EvaluateWide(inst.Arg(0), a)) {
+					return Arg(inst, a != 0u ? 1u : 2u, result);
 				}
 				return false;
+			}
 			case ValueOpcode::IEqual32:
 				if (binary()) {
 					result = static_cast<uint32_t>(a) == static_cast<uint32_t>(b);
@@ -977,7 +987,7 @@ bool EvaluateRuntimeSourcesImpl(const ResourcePlan& program, std::span<const uin
                                 const SrtRuntime& runtime, std::vector<DescriptorValue>& results,
                                 std::vector<uint32_t>& flat, bool evaluate_flat,
                                 std::span<const uint8_t> clean_flat_slots,
-                                std::vector<uint8_t>& active_sources) {
+                                std::vector<uint8_t>&    active_sources) {
 	if (!program.srt_plan_complete) {
 		return false;
 	}
@@ -985,8 +995,7 @@ bool EvaluateRuntimeSourcesImpl(const ResourcePlan& program, std::span<const uin
 	    runtime.read_specialization_memory == nullptr) {
 		return false;
 	}
-	SrtRuntime clean_runtime  = runtime;
-	clean_runtime.read_memory = runtime.read_specialization_memory;
+	const auto           clean_runtime = CleanRuntime(runtime);
 	Evaluator            clean_evaluator(program, clean_runtime);
 	Evaluator            evaluator(program, runtime, clean_flat_slots, &clean_evaluator);
 	std::vector<uint8_t> active;
@@ -1081,11 +1090,8 @@ bool EvaluateUniformValues(const ResourcePlan& program, std::span<const Value> v
 	if (values.size() != results.size()) {
 		return false;
 	}
-	auto clean = runtime;
-	clean.read_memory = runtime.read_specialization_memory != nullptr
-	                        ? runtime.read_specialization_memory
-	                        : +[](void*, uint64_t, uint32_t*) { return false; };
-	Evaluator evaluator(program, clean);
+	const auto clean = CleanRuntime(runtime);
+	Evaluator  evaluator(program, clean);
 	for (size_t i = 0; i < values.size(); ++i) {
 		if (!evaluator.Evaluate(values[i], results[i])) {
 			return false;

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

@@ -213,7 +213,7 @@ void ConfigurationEditDialog::Init(const Configuration& info) {
 #endif
 	QVulkanInstance instance;
 	instance.setApiVersion(QVersionNumber(1, 3, 0));
-#if !defined(__APPLE__)
+#if !defined(__APPLE__) && QT_VERSION >= QT_VERSION_CHECK(6, 6, 0)
 	instance.setFlags(QVulkanInstance::NoPortabilityDrivers);
 #endif
 	if (instance.create()) {
@@ -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/agc.cpp

@@ -262,35 +262,25 @@ struct CommandBuffer {
 		     reserved_dw);
 	}
 
-	[[nodiscard]] KYTY_SYSV_ABI uint32_t GetAvailableSizeDW() const {
-		if (cursor_up == nullptr || cursor_down == nullptr || cursor_down <= cursor_up) {
-			return 0;
-		}
-
-		auto available = static_cast<uint64_t>(cursor_down - cursor_up);
+	[[nodiscard]] KYTY_SYSV_ABI uint64_t GetAvailableSizeDW() const {
+		// Interpret the signed cursor distance as an unsigned 64-bit DWORD count.
+		const auto distance = static_cast<int64_t>(reinterpret_cast<uintptr_t>(cursor_down) -
+		                                           reinterpret_cast<uintptr_t>(cursor_up));
+		const auto available = static_cast<uint64_t>(distance >> 2);
 		if (available <= reserved_dw) {
 			return 0;
 		}
-		if (available - reserved_dw > UINT32_MAX) {
-			LOGF_COLOR(
-			    Log::Color::Red,
-			    "\t command buffer has suspiciously large free space: cursor_up = 0x%016" PRIx64
-			    ", cursor_down = 0x%016" PRIx64 ", reserved_dw = %" PRIu32 "\n",
-			    reinterpret_cast<uint64_t>(cursor_up), reinterpret_cast<uint64_t>(cursor_down),
-			    reserved_dw);
-			return UINT32_MAX;
-		}
-		return static_cast<uint32_t>(available - reserved_dw);
+		return available - reserved_dw;
 	}
 
 	KYTY_SYSV_ABI bool ReserveDW(uint32_t num_dw) {
-		uint32_t remaining = GetAvailableSizeDW();
+		const uint64_t remaining = GetAvailableSizeDW();
 		if (num_dw > remaining) {
 			if (callback == nullptr) {
 				LOGF_COLOR(
 				    Log::Color::Red,
 				    "\t command buffer exhausted and has no grow callback: requested = %" PRIu32
-				    ", remaining = %" PRIu32 ", reserved_dw = %" PRIu32 "\n",
+				    ", remaining = %" PRIu64 ", reserved_dw = %" PRIu32 "\n",
 				    num_dw, remaining, reserved_dw);
 				DbgDump();
 				return false;
@@ -300,7 +290,7 @@ struct CommandBuffer {
 			if (!result) {
 				LOGF_COLOR(Log::Color::Red,
 				           "\t command buffer grow callback failed: requested = %" PRIu32
-				           ", remaining = %" PRIu32 ", reserved_dw = %" PRIu32 "\n",
+				           ", remaining = %" PRIu64 ", reserved_dw = %" PRIu32 "\n",
 				           num_dw, remaining, reserved_dw);
 				DbgDump();
 				return false;
@@ -308,7 +298,7 @@ struct CommandBuffer {
 			if (GetAvailableSizeDW() < num_dw) {
 				LOGF_COLOR(Log::Color::Red,
 				           "\t command buffer grow callback did not provide enough space: "
-				           "requested = %" PRIu32 ", remaining = %" PRIu32
+				           "requested = %" PRIu32 ", remaining = %" PRIu64
 				           ", reserved_dw = %" PRIu32 "\n",
 				           num_dw, GetAvailableSizeDW(), reserved_dw);
 				DbgDump();

src/libs/libJson2.cpp

@@ -41,6 +41,10 @@ struct JsonArray {
 	std::vector<JsonValue*>* impl;
 };
 
+struct JsonArrayIterator {
+	JsonValue* const* position;
+};
+
 struct JsonObject {
 	std::map<std::string, JsonValue*>* impl;
 };
@@ -62,6 +66,7 @@ struct JsonValue {
 };
 
 static_assert(sizeof(JsonArray) == 8);
+static_assert(sizeof(JsonArrayIterator) == 8);
 static_assert(sizeof(JsonValue) == 32 && offsetof(JsonValue, type) == 28);
 
 struct JsonInitParameter2 {
@@ -751,7 +756,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) {
@@ -906,6 +915,51 @@ static size_t KYTY_SYSV_ABI JsonArraySize(const JsonArray* self) {
 	return JsonArrayImpl(self)->size();
 }
 
+// The non-trivial guest iterator is returned through an explicit result pointer.
+static JsonArrayIterator* KYTY_SYSV_ABI JsonArrayBegin(JsonArrayIterator* result,
+                                                    const JsonArray* self) {
+	PRINT_NAME();
+
+	result->position = JsonArrayImpl(self)->data();
+	return result;
+}
+
+static JsonArrayIterator* KYTY_SYSV_ABI JsonArrayEnd(JsonArrayIterator* result,
+                                                  const JsonArray* self) {
+	PRINT_NAME();
+
+	const auto* impl = JsonArrayImpl(self);
+	result->position = impl->data();
+	if (!impl->empty()) {
+		result->position += impl->size();
+	}
+	return result;
+}
+
+static bool KYTY_SYSV_ABI JsonArrayIteratorNotEqual(const JsonArrayIterator* self,
+                                                 const JsonArrayIterator* other) {
+	PRINT_NAME();
+
+	return self->position != other->position;
+}
+
+static JsonValue* KYTY_SYSV_ABI JsonArrayIteratorDereference(const JsonArrayIterator* self) {
+	PRINT_NAME();
+
+	return *self->position;
+}
+
+static JsonArrayIterator* KYTY_SYSV_ABI JsonArrayIteratorIncrement(JsonArrayIterator* self) {
+	PRINT_NAME();
+
+	++self->position;
+	return self;
+}
+
+static void KYTY_SYSV_ABI JsonArrayIteratorDtor(JsonArrayIterator*) {
+	PRINT_NAME();
+}
+
 static JsonObject* KYTY_SYSV_ABI JsonObjectCtor(JsonObject* self) {
 	PRINT_NAME();
 
@@ -1080,6 +1134,12 @@ LIB_DEFINE(InitNet_1_Json2) {
 	LIB_FUNC("3qrge7L-AU4", LibJson2::JsonValueGetReal);
 	LIB_FUNC("sOmU4vnx3s0", LibJson2::JsonValueDoubleCtor);
 	LIB_FUNC("rQGJeNjOuUk", LibJson2::JsonArraySize);
+	LIB_FUNC("bcH5EnFE2xY", LibJson2::JsonArrayBegin);
+	LIB_FUNC("WXF2ihRF+B8", LibJson2::JsonArrayEnd);
+	LIB_FUNC("5AZPp99ogrc", LibJson2::JsonArrayIteratorNotEqual);
+	LIB_FUNC("wcgr5mte7T8", LibJson2::JsonArrayIteratorDereference);
+	LIB_FUNC("w5+VCznos5E", LibJson2::JsonArrayIteratorIncrement);
+	LIB_FUNC("9yLjn46Ypfs", LibJson2::JsonArrayIteratorDtor);
 	LIB_FUNC("5yHuiWXo2gg", LibJson2::JsonValueSetBool);
 	LIB_FUNC("QxVVYhP-mvg", LibJson2::JsonValueSetInt);
 	LIB_FUNC("SIe1ZmW7e7s", LibJson2::JsonValueSetUInt);

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,11 +2,19 @@
 
 #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
-#elif !defined(__APPLE__)
+#elif defined(__APPLE__)
+#include <sys/ucontext.h>
+#else
 #include <sched.h>
 #include <ucontext.h>
 #endif
@@ -400,109 +408,170 @@ static bool ExecuteShaNiInsn(const ShaNiInsn& insn, const XmmWords& src2, const
 	}
 }
 
+// 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] = {};
+		}
+	}
+#elif defined(__APPLE__)
+	ucontext_t* native;
+
+	[[nodiscard]] uint64_t Rip() const {
+		return static_cast<uint64_t>(native->uc_mcontext->__ss.__rip);
+	}
+	void Advance(size_t length) {
+		native->uc_mcontext->__ss.__rip += static_cast<uint64_t>(length);
+	}
+	// Darwin names the XMM file __fpu_xmm0..__fpu_xmm15 instead of exposing an array.
+	[[nodiscard]] void* Xmm(uint8_t index) const {
+		auto* fs = &native->uc_mcontext->__fs;
+		switch (index) {
+			case 0: return &fs->__fpu_xmm0;
+			case 1: return &fs->__fpu_xmm1;
+			case 2: return &fs->__fpu_xmm2;
+			case 3: return &fs->__fpu_xmm3;
+			case 4: return &fs->__fpu_xmm4;
+			case 5: return &fs->__fpu_xmm5;
+			case 6: return &fs->__fpu_xmm6;
+			case 7: return &fs->__fpu_xmm7;
+			case 8: return &fs->__fpu_xmm8;
+			case 9: return &fs->__fpu_xmm9;
+			case 10: return &fs->__fpu_xmm10;
+			case 11: return &fs->__fpu_xmm11;
+			case 12: return &fs->__fpu_xmm12;
+			case 13: return &fs->__fpu_xmm13;
+			case 14: return &fs->__fpu_xmm14;
+			case 15: return &fs->__fpu_xmm15;
+			default: return nullptr;
+		}
+	}
+#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 !defined(__APPLE__)
+
+static bool TryEmulateShaNi(Context& context) {
+	const auto* rip = reinterpret_cast<const uint8_t*>(context.Rip());
+	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   = context.Xmm(dest_index);
+	auto*         xmm0       = context.Xmm(0);
+	if (dest_xmm == nullptr || xmm0 == nullptr) {
+		return false;
+	}
+
+	XmmWords dest {};
+	XmmWords src2 {};
+	XmmWords 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   = context.Xmm(src_index);
+		if (src_xmm == nullptr) {
+			return false;
+		}
+		std::memcpy(&src2, src_xmm, sizeof(src2));
+	} else {
+		uint64_t    gpr[16] {};
+		const void* source = nullptr;
+		context.LoadGprs(gpr);
+		if (!ResolveShaNiMemoryAddress(rip, insn, gpr, source)) {
+			return false;
+		}
+		std::memcpy(&src2, source, sizeof(src2));
+	}
+
+	if (!ExecuteShaNiInsn(insn, src2, xmm0_words, dest)) {
+		return false;
+	}
+
+	std::memcpy(dest_xmm, &dest, sizeof(dest));
+	context.Advance(insn.length);
+	return true;
 }
 
 #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);
-	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 {};
-	LoadXmmWordsWin(dest_xmm, dest);
-	LoadXmmWordsWin(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;
-		}
-		LoadXmmWordsWin(src_xmm, src2);
-	} else {
-		uint64_t    gpr[16] {};
-		const void* source = nullptr;
-		LoadContextGprsWin(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;
-	}
-
-	StoreXmmWordsWin(dest_xmm, dest);
-	context->Rip += 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,277 +580,212 @@ 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) {
+	if (rip[offset] != 0x0f) {
+		return false;
+	}
+	const bool register_extract = prefix == 0x66 && rip[offset + 1] == 0x79;
+	if (rip[offset + 1] != 0x78 && !register_extract) {
 		return false;
 	}
 
-	auto modrm = rip[offset + 2];
+	const uint8_t 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];
+	const uint8_t reg = ((modrm >> 3u) & 0x07u) | ((rex & 0x04u) << 1u);
+	const uint8_t rm  = (modrm & 0x07u) | ((rex & 0x01u) << 3u);
 
-	// 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;
+	// Immediate EXTRQ encodes its destination in r/m; the two-register form uses reg.
+	uint8_t dest_index = reg;
+	if (prefix == 0x66 && !register_extract) {
+		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;
 	}
-
-	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);
+	uint64_t dest[2] {};
+	uint64_t source = 0;
+	std::memcpy(dest, dest_xmm, sizeof(dest));
+	std::memcpy(&source, src_xmm, sizeof(source));
+	uint8_t length             = 0;
+	uint8_t index              = 0;
+	size_t  instruction_length = offset + 3;
+	if (register_extract) {
+		length = static_cast<uint8_t>(source);
+		index  = static_cast<uint8_t>(source >> 8u);
 	} 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));
+		length = rip[offset + 3];
+		index  = rip[offset + 4];
+		instruction_length += 2;
 	}
-
-	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];
-	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.
 	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[0] = ExtractBitField(dest[0], length, index);
+		dest[1] = 0;
+	} else {
+		dest[0] = InsertBitField(dest[0], source, length, index);
 	}
-
-	auto* dst = GetContextXmm(context, reg);
-	auto* src = GetContextXmm(context, rm);
-	if (dst == nullptr || src == nullptr) {
-		return false;
-	}
-
-	SetXmmLow(dst, InsertBitField(GetXmmLow(dst), GetXmmLow(src), length, index));
-	rip_reg += static_cast<greg_t>(offset + 5);
+	std::memcpy(dest_xmm, dest, sizeof(dest));
+	context.Advance(instruction_length);
 	return true;
 }
 
-static bool TryEmulateMonitorxMwaitx(ucontext_t* context) {
-	if (context == nullptr) {
-		return false;
-	}
+#if !defined(__APPLE__)
 
-	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 (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)};
+#elif defined(__APPLE__)
+	auto* saved_context = static_cast<ucontext_t*>(native_context);
+	if (saved_context->uc_mcontext == nullptr) {
+		return false;
+	}
+	Context context {saved_context};
+#else
+	Context context {static_cast<ucontext_t*>(native_context)};
+#endif
+#if !defined(__APPLE__)
+	if (TryEmulateReciprocalSquareRoot(context)) {
+		return true;
+	}
 	return TryEmulateMonitorxMwaitx(context) || TryEmulateSse4a(context) ||
 	       TryEmulateShaNi(context);
 #else
-	(void)native_context;
-	return false;
+	return TryEmulateSse4a(context);
 #endif
 }
 

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

@@ -1208,6 +1208,156 @@ void TestPhiValidation() {
         "control-dependent descriptor phi was not rejected transactionally");
 }
 
+ResourcePlan ConditionalSamplerPlan(bool diamond, bool reverse, bool reverse_phi,
+                                    bool nonuniform = false,
+                                    bool writable = false) {
+  namespace CFG = Libs::Graphics::ShaderRecompiler::CFG;
+  Fixture fixture(ShaderType::Pixel);
+  auto *entry = fixture.block;
+  auto *initial = diamond ? fixture.AddBlock() : entry;
+  auto *alternate = fixture.AddBlock();
+  auto *merge = fixture.AddBlock();
+  const uint32_t alternate_id = diamond ? 2u : 1u;
+  const uint32_t merge_id = alternate_id + 1;
+  const uint32_t initial_target = diamond ? 1u : merge_id;
+  entry->AddBranch(alternate);
+  entry->AddBranch(diamond ? initial : merge);
+  alternate->AddBranch(merge);
+  if (diamond) {
+    initial->AddBranch(merge);
+    fixture.program.block_info[1].terminator = {
+        .kind = CFG::TerminatorKind::Branch, .true_block = merge_id};
+  }
+  fixture.program.block_info[0].terminator = {
+      .kind = CFG::TerminatorKind::ConditionalBranch,
+      .true_block = reverse ? alternate_id : initial_target,
+      .false_block = reverse ? initial_target : alternate_id};
+  fixture.program.block_info[alternate_id].terminator = {
+      .kind = CFG::TerminatorKind::Branch, .true_block = merge_id};
+  fixture.program.block_info[merge_id].terminator.kind =
+      CFG::TerminatorKind::Return;
+  const auto control =
+      fixture.Buffer({Value(0x2000u), Value(0u), Value(200u), Value(0u)});
+  MemoryInfo scalar;
+  scalar.kind = ResourceKind::ScalarBuffer;
+  auto flag = fixture.Emit(ValueOpcode::ReadConstBuffer, {control, Value(196u)},
+                           fixture.AddMemory(scalar, 0x498));
+  if (nonuniform) {
+    flag = fixture.Emit(ValueOpcode::LaneId);
+  }
+  fixture.program.block_info[0].condition =
+      fixture.Emit(ValueOpcode::IEqual32, {flag, Value(0u)});
+  if (writable) {
+    MemoryInfo memory;
+    memory.kind = ResourceKind::Buffer;
+    fixture.Emit(
+        ValueOpcode::StoreBufferU32,
+        {control, Value(0u), Value(0u), Value(0u), Value(1u), Value(true)},
+        fixture.AddMemory(memory, 0x170));
+  }
+
+  std::array<Value, 4> sampler_words;
+  for (uint32_t word = 0; word < sampler_words.size(); ++word) {
+    const auto read = [&](Block *block, uint32_t address, uint32_t pc) {
+      const auto handle = fixture.Emit(ValueOpcode::GetAddressResource,
+                                       {Value(address), Value(0u)}, 0, block);
+      MemoryInfo memory;
+      memory.kind = ResourceKind::ScalarAddress;
+      memory.offset = word * 4;
+      return fixture.Emit(ValueOpcode::LoadAddressU32,
+                          {handle, Value(0u), Value(0u), Value(true)},
+                          fixture.AddMemory(memory, pc), block);
+    };
+    // PS 2190adcc312b2e6e selects SRT+448 or SRT+480 before its sample at
+    // 0x4d4.
+    const auto first = read(initial, 0x1000 + 448, 0x4c8);
+    const auto second = read(alternate, 0x1000 + 480, 0x4bc);
+    auto &phi = merge->AppendNewInst(ValueOpcode::Phi, {},
+                                     static_cast<uint64_t>(Type::U32));
+    if (reverse_phi) {
+      phi.AddPhiOperand(alternate, second);
+      phi.AddPhiOperand(initial, first);
+    } else {
+      phi.AddPhiOperand(initial, first);
+      phi.AddPhiOperand(alternate, second);
+    }
+    sampler_words[word] = Value(&phi);
+  }
+  fixture.block = merge;
+  const auto image =
+      fixture.Image({Value(0u), Value(0u), Value(0u), Value(0u), Value(0u),
+                     Value(0u), Value(0u), Value(0u)});
+  const auto address = fixture.ImageAddress();
+  for (uint32_t use = 0; use < 2; ++use) {
+    const auto sampler = fixture.Sampler(sampler_words);
+    MemoryInfo sample;
+    sample.kind = ResourceKind::Image;
+    sample.image_dimension = Decoder::ImageDimension::Dim2D;
+    const auto result =
+        fixture.Emit(ValueOpcode::ImageSampleRaw, {image, sampler, address},
+                     fixture.AddMemory(sample, 0x4d4));
+    fixture.Emit(ValueOpcode::ReferenceU32,
+                 {fixture.Emit(ValueOpcode::CompositeExtractU32x4,
+                               {result, Value(0u)})});
+  }
+  fixture.PlanAndTrack();
+  Check(fixture.program.value_storage.size() == 4,
+        "repeated sampler uses retained duplicate planning selections");
+  Check(sampler_words[0].ResolveInstruction()->GetOpcode() == ValueOpcode::Phi,
+        "host descriptor selection changed the GPU Phi");
+  EliminateDeadCode(fixture.program.blocks);
+  ValidateProgram(fixture.program, true);
+  return ExtractResourcePlan(fixture.program);
+}
+
+void TestConditionalSamplerPhi() {
+  for (const bool diamond : {false, true}) {
+    for (const bool reverse : {false, true}) {
+      for (const bool reverse_phi : {false, true}) {
+        auto plan = ConditionalSamplerPlan(diamond, reverse, reverse_phi);
+        const auto source = plan.info.samplers[0].source;
+        LinearTestMemory memory;
+        for (uint32_t word = 448 / 4; word < (480 + 16) / 4; ++word) {
+          memory.words[word] = 0x400u + word;
+        }
+        const SrtRuntime runtime{.read_memory = ReadLinearTestMemory,
+                                 .userdata = &memory,
+                                 .read_specialization_memory =
+                                     ReadLinearTestMemory};
+        for (uint32_t flag = 0; flag < 2; ++flag) {
+          memory.words[(0x1000 + 196) / 4] = flag;
+          const uint32_t first = (flag != 0u) != reverse ? 480 / 4 : 448 / 4;
+          memory.fail_address = 0x1000 + (first == 448 / 4 ? 480u : 448u);
+          DescriptorValue selected;
+          Check(EvaluateDescriptorSource(plan, source, runtime, selected),
+                "conditional sampler did not survive detached plan lifetime");
+          for (uint32_t word = 0; word < 4; ++word) {
+            Check(selected.dwords[word] == memory.words[first + word],
+                  "conditional sampler chose the wrong incoming descriptor");
+          }
+        }
+        DescriptorValue selected;
+        auto no_clean_reader = runtime;
+        no_clean_reader.read_specialization_memory = nullptr;
+        Check(
+            !EvaluateDescriptorSource(plan, source, no_clean_reader, selected),
+            "conditional sampler used unchecked memory for its predicate");
+        memory.fail_address = 0x2000 + 196;
+        Check(!EvaluateDescriptorSource(plan, source, runtime, selected),
+              "conditional sampler ignored unavailable coherent predicate "
+              "memory");
+      }
+    }
+    CheckFatal([&] { ConditionalSamplerPlan(diamond, false, false, true); },
+               "not a valid runtime value",
+               "nonuniform sampler selection was accepted");
+    CheckFatal(
+        [&] { ConditionalSamplerPlan(diamond, false, false, false, true); },
+        "not a valid runtime value",
+        "shader-written sampler predicate was accepted");
+  }
+}
+
 void TestLoopCycleEnteredThroughRuntimeValue() {
   Fixture fixture;
   auto *entry = fixture.block;
@@ -1865,6 +2015,7 @@ int main() {
     Run("SRT runtime", TestSrtFlatteningAndRuntimeMemoization);
     Run("dynamic SRT", TestDynamicSrtReadRemainsExplicit);
     Run("phi validation", TestPhiValidation);
+    Run("conditional sampler phi", TestConditionalSamplerPhi);
     Run("runtime-rooted loop", TestLoopCycleEnteredThroughRuntimeValue);
     Run("invariant loop phi", TestInvariantLoopPhi);
     Run("DMA address materialization", TestDmaAddressMaterialization);

tests/ShaderRecompilerComputeTests.cpp

@@ -379,6 +379,11 @@ struct TextureCacheTestAccess {
 };
 
 struct RenderExecutorTestAccess {
+  static void DrawAuto(RenderExecutor &executor, CommandBuffer &command,
+                       const DrawAutoArgs &args) {
+    executor.DrawAuto(0, command, args);
+  }
+
   static bool TryConsumeComputeImageClear(RenderExecutor &executor,
       const ShaderComputeInputInfo &input, CommandBuffer &command,
       uint32_t x, uint32_t y, uint32_t z, uint32_t mode) {
@@ -4476,6 +4481,17 @@ public:
     TextureCacheTestAccess::RegisterHtileMeta(texture_cache, read_only_meta);
     TextureCacheTestAccess::RegisterHtileMeta(texture_cache, read_write_meta);
     TextureCacheTestAccess::RegisterHtileMeta(texture_cache, write_only_meta);
+
+    // An empty dispatch must not inspect even an unreadable shader address.
+    shaders.SetCsShader({.data_addr = 1});
+    constexpr std::array<std::array<uint32_t, 3>, 3> empty_dispatches{{
+        {0, 1, 1}, {1, 0, 1}, {1, 1, 0}}};
+    for (const auto mode : {0x41u, 0x61u}) {
+      for (const auto &groups : empty_dispatches) {
+        context.GetRenderExecutor().DispatchDirect(
+            0, scheduler.Current(), groups[0], groups[1], groups[2], mode);
+      }
+    }
     Require(name, "HTile fixture",
             texture_cache.IsMeta(read_only_meta) &&
                 texture_cache.IsMeta(read_write_meta) &&
@@ -5652,7 +5668,7 @@ public:
       auto stencil_storage_desc = MakeLinearDesc(
           base + second_stencil_offset, ms_stencil_size, vk::Format::eR8Uint,
           Prospero::BufferFormat::k8UInt, Prospero::ImageType::kColor2D,
-          {256, 256, 1}, 256, 1, 1);
+          ms_depth_desc.info.extent, 1, 1, 1);
       stencil_storage_desc.type = BindingType::Storage;
       stencil_storage_desc.view_info.usage = vk::ImageUsageFlagBits::eStorage;
       const auto stencil_storage_image =
@@ -7107,6 +7123,64 @@ public:
                                                 ms_stencil_size, false),
               "unmapping depth retained its re-associated stencil proxy");
 
+      auto reused_stencil_color = MakeLinearDesc(
+          base + 0x2200000, 0x40000, vk::Format::eR8Unorm,
+          Prospero::BufferFormat::k8UNorm, Prospero::ImageType::kColor2D,
+          {480, 408, 1}, 1, 1, 1);
+      const auto reused_stencil_color_id =
+          texture_cache.FindImage(reused_stencil_color);
+      auto reused_stencil_depth = MakeLinearDesc(
+          base + 0x2240000, 0x10000, vk::Format::eD32SfloatS8Uint,
+          Prospero::BufferFormat::k32Float, Prospero::ImageType::kColor2D,
+          {128, 128, 1}, 1, 4, 1);
+      reused_stencil_depth.type = BindingType::DepthTarget;
+      reused_stencil_depth.info.stencil = {reused_stencil_color.info.data.address,
+                                          0x10000};
+      reused_stencil_depth.view_info.aspect =
+          vk::ImageAspectFlagBits::eDepth | vk::ImageAspectFlagBits::eStencil;
+      reused_stencil_depth.view_info.usage =
+          vk::ImageUsageFlagBits::eDepthStencilAttachment;
+      const auto reused_stencil_depth_id =
+          texture_cache.FindImage(reused_stencil_depth);
+      TextureCacheTestAccess::AssociateStencil(
+          texture_cache, reused_stencil_depth_id,
+          reused_stencil_depth.info.stencil);
+      const auto reused_stencil_proxy = texture_cache.FindImageFromRange(
+          reused_stencil_depth.info.stencil.address,
+          reused_stencil_depth.info.stencil.size, false);
+      Require(name, "stencil address reuse preserves color",
+              reused_stencil_proxy &&
+                  reused_stencil_proxy != reused_stencil_color_id &&
+                  !texture_cache.GetImage(reused_stencil_color_id).depth_id &&
+                  texture_cache.GetImage(reused_stencil_proxy).depth_id ==
+                      reused_stencil_depth_id &&
+                  texture_cache.GetImage(reused_stencil_proxy).info.data ==
+                      reused_stencil_depth.info.stencil &&
+                  texture_cache.GetImage(reused_stencil_proxy).info.extent ==
+                      reused_stencil_depth.info.extent,
+              "a smaller stencil plane took ownership of an existing R8 image");
+      vk::ClearValue reused_stencil_color_clear{};
+      reused_stencil_color_clear.color.float32 =
+          std::array{1.0f, 0.0f, 0.0f, 0.0f};
+      TextureCacheTestAccess::ClearImage(
+          texture_cache, scheduler.Current(), reused_stencil_color_id,
+          {vk::ImageAspectFlagBits::eColor, 0, 1, 0, 1},
+          reused_stencil_color_clear);
+      Require(name, "reused stencil plane clear",
+              texture_cache.ClearImageFromBuffer(
+                  scheduler.Current(), reused_stencil_depth.info.stencil.address,
+                  reused_stencil_depth.info.stencil.size, 0x5a5a5a5au) &&
+                  texture_cache.GetImage(reused_stencil_depth_id).IsGpuModified(),
+              "the separate stencil association did not resolve to depth");
+      Require(name, "color clear after stencil address reuse",
+              ReadCachedTexel(name, context, reused_stencil_color_id, {},
+                              {4, 1, 1}) == std::vector<u32>{0xffffffffu},
+              "stencil address reuse changed the native R8 clear contents");
+      texture_cache.UnmapMemory(reused_stencil_depth.info.data.address,
+                                reused_stencil_depth.info.data.size);
+      texture_cache.UnmapMemory(reused_stencil_color.info.data.address,
+                                reused_stencil_color.info.data.size);
+
       constexpr uint64_t exact_image_offset = 0x334100;
       constexpr uint64_t dirty_sibling_offset = 0x334200;
       constexpr uint32_t dirty_sibling_value = 0xc001d00du;
@@ -8583,7 +8657,8 @@ public:
 
   std::vector<u32> ReadCachedTexel(const char *name, RenderContext &context,
                                  ImageId id, vk::Offset3D offset = {},
-                                 vk::Extent3D extent = {1, 1, 1}, uint32_t layer = 0) {
+                                 vk::Extent3D extent = {1, 1, 1}, uint32_t layer = 0,
+                                 uint32_t mip = 0) {
     auto &scheduler = context.GetCommandScheduler();
     auto &image = context.GetTextureCache().GetImage(id);
     const auto bytes = image.info.bytes_per_block * extent.width *
@@ -8596,7 +8671,7 @@ public:
                   scheduler.Current().Handle());
     vk::BufferImageCopy copy{};
     copy.imageSubresource = {image.info.IsDepth() ? vk::ImageAspectFlagBits::eDepth
-                                                : vk::ImageAspectFlagBits::eColor, 0, layer, 1};
+                                                : vk::ImageAspectFlagBits::eColor, mip, layer, 1};
     copy.imageOffset = offset;
     copy.imageExtent = extent;
     scheduler.Current().Handle().copyImageToBuffer(
@@ -9591,6 +9666,9 @@ public:
     constexpr uint64_t allocation_alignment = 0x10000;
     constexpr uint64_t depth_address = base + 0x40000;
     constexpr uint64_t stencil_address = base + 0x70000;
+    constexpr uint64_t mipped_storage_address = base + 0x100000;
+    constexpr uint32_t tail_backing_value = 0x2468ace0u;
+    constexpr uint32_t tail_gpu_value = 0x13579bdfu;
     EnsureRuntimeContext();
 
     int64_t direct_offset = -1;
@@ -9607,6 +9685,8 @@ public:
                 mapped == reinterpret_cast<void *>(base),
             "descriptor discovery fixed mapping failed");
     std::memset(mapped, 0, allocation_size);
+    std::fill_n(reinterpret_cast<uint32_t *>(mipped_storage_address),
+                0x10000 / 4, tail_backing_value);
 
     {
       RenderContext context(m_runtime_context);
@@ -9873,17 +9953,17 @@ public:
                 storage_descriptor.dwords.begin());
       storage_descriptor.dword_count = 8;
       auto mipped_storage = storage;
-      constexpr uint64_t mipped_storage_address = base + 0x100000;
       const auto encoded_mipped_storage_address = mipped_storage_address >> 8u;
       mipped_storage.fields[0] =
           static_cast<uint32_t>(encoded_mipped_storage_address);
       mipped_storage.fields[1] =
           static_cast<uint32_t>(encoded_mipped_storage_address >> 32u) |
-          (static_cast<uint32_t>(stencil_format) << 20u) | (3u << 30u);
-      mipped_storage.fields[2] = 1u | (7u << 14u);
+          (static_cast<uint32_t>(Prospero::BufferFormat::k32UInt) << 20u) |
+          (3u << 30u);
+      mipped_storage.fields[2] = 15u | (63u << 14u);
       mipped_storage.fields[3] =
           DstSel(4, 5, 6, 7) | (1u << 12u) | (3u << 16u) |
-          (static_cast<uint32_t>(linear) << 20u) |
+          (static_cast<uint32_t>(Prospero::TileMode::kStandard64KB) << 20u) |
           (static_cast<uint32_t>(Prospero::ImageType::kColor2D) << 28u);
       mipped_storage.fields[5] = 0x00700000u | (3u << 4u);
       ShaderRecompiler::IR::DescriptorValue mipped_storage_descriptor{};
@@ -9900,6 +9980,14 @@ public:
                 std::end(overwide_mipped_storage.fields),
                 overwide_mipped_storage_descriptor.dwords.begin());
       overwide_mipped_storage_descriptor.dword_count = 8;
+#if KYTY_PLATFORM != KYTY_PLATFORM_WINDOWS
+      ExpectFatal("MipViewPhysicalLayoutChange", [&] {
+        auto linear_view = overwide_mipped_storage_descriptor;
+        linear_view.dwords[3] &= ~(0x1fu << 20u);
+        (void)RenderExecutorTestAccess::ResolveTexture(executor, storage_resource,
+                                                       linear_view);
+      });
+#endif
       auto mipped_storage_resource = storage_resource;
       mipped_storage_resource.mip_mode =
           ShaderRecompiler::IR::ImageMipMode::DynamicStorage;
@@ -9914,6 +10002,11 @@ public:
       auto plain_mipped_storage_binding =
           RenderExecutorTestAccess::ResolveTexture(executor, storage_resource,
                                                    mipped_storage_descriptor);
+      vk::ClearValue tail_clear{};
+      tail_clear.color.uint32[0] = tail_gpu_value;
+      TextureCacheTestAccess::ClearImage(
+          texture_cache, scheduler.Current(), plain_mipped_storage_binding.image_id,
+          {vk::ImageAspectFlagBits::eColor, 0, 4, 0, 1}, tail_clear);
       auto mipped_storage_binding = RenderExecutorTestAccess::ResolveTexture(
           executor, mipped_storage_resource, mipped_storage_descriptor);
       auto overwide_mipped_storage_binding =
@@ -9922,6 +10015,16 @@ public:
       auto sampled_overwide_resolved = RenderExecutorTestAccess::ResolveTexture(
           executor, sampled_overwide_resource,
           overwide_mipped_storage_descriptor);
+      const auto tail0 = ReadCachedTexel(name, context, overwide_mipped_storage_binding.image_id);
+      const auto tail3 = ReadCachedTexel(name, context, overwide_mipped_storage_binding.image_id,
+                                        {}, {1, 1, 1}, 0, 3);
+      const auto tail4 = ReadCachedTexel(name, context, overwide_mipped_storage_binding.image_id,
+                                        {}, {1, 1, 1}, 0, 4);
+      Require(name, "expanded tail preserves GPU and backing data",
+              tail0 == std::vector<u32>{tail_gpu_value} &&
+                  tail3 == std::vector<u32>{tail_gpu_value} &&
+                  tail4 == std::vector<u32>{tail_backing_value},
+              "expanding a shared mip tail lost GPU writes or the added mip's backing bytes");
       PreparedBindings mipped_prepared{};
       ShaderRecompiler::IR::CompiledShaderInfo mipped_program{};
       mipped_program.info.images.push_back(storage_resource);
@@ -9983,23 +10086,129 @@ public:
                   overwide_mipped_storage.MaxMip() == 3 &&
                   overwide_mipped_binding.image_id ==
                       plain_mipped_binding.image_id &&
-                  overwide_mipped_binding.desc.info.resources.levels == 4 &&
+                  overwide_mipped_binding.desc.info.resources.levels == 5 &&
                   overwide_mipped_binding.desc.view_info.base_level == 1 &&
-                  overwide_mipped_binding.desc.view_info.level_count == 3 &&
+                  overwide_mipped_binding.desc.view_info.level_count == 4 &&
                   overwide_mipped_binding.mip_views.empty() &&
                   overwide_mipped_binding.image_view ==
                       plain_mipped_binding.image_view,
-              "fixed storage view was not intersected with its physical mip "
-              "range before Vulkan acquisition");
+              "fixed storage view lost addressable mips in the allocated tail");
       Require(name, "over-wide sampled mip view",
               sampled_overwide_binding.image_id ==
                       plain_mipped_binding.image_id &&
-                  sampled_overwide_binding.desc.info.resources.levels == 4 &&
+                  sampled_overwide_binding.desc.info.resources.levels == 5 &&
                   sampled_overwide_binding.desc.view_info.base_level == 1 &&
-                  sampled_overwide_binding.desc.view_info.level_count == 3 &&
+                  sampled_overwide_binding.desc.view_info.level_count == 4 &&
                   sampled_overwide_binding.mip_views.empty() &&
                   sampled_overwide_binding.image_view != nullptr,
-              "sampled view was not intersected with its physical mip range");
+              "sampled view lost addressable mips in the allocated tail");
+      // Streaming T# clamps apply after S# max LOD, without changing the view
+      // base.
+      {
+        constexpr uint64_t lod_address = base + 0x170000;
+        auto lod_descriptor = mipped_storage_descriptor;
+        lod_descriptor.dwords[0] = static_cast<uint32_t>(lod_address >> 8u);
+        lod_descriptor.dwords[1] =
+            static_cast<uint32_t>(lod_address >> 40u) |
+            (static_cast<uint32_t>(Prospero::BufferFormat::k32Float) << 20u) |
+            (3u << 30u);
+        lod_descriptor.dwords[3] &= ~(0xfu << 12u);
+        ShaderSamplerResource lod_sampler{
+            {0, 0,
+             static_cast<uint32_t>(Prospero::SamplerMipFilter::kLinear) << 26u,
+             0}};
+        TestCase lod_test;
+        lod_test.name = "TextureMinLodAfterSamplerClamp";
+        lod_test.has_user_data = true;
+        lod_test.image_descriptor_swizzle = mipped_storage.DstSelXYZW();
+        std::copy_n(lod_descriptor.dwords.begin(), 8,
+                    lod_test.user_data.begin());
+        std::copy_n(lod_sampler.fields, 4, lod_test.user_data.begin() + 8);
+        lod_test.user_data[50] = sizeof(uint32_t);
+        lod_test.opcodes = {ShaderOpcode::V_MOV_B32, ShaderOpcode::IMAGE_SAMPLE,
+                            ShaderOpcode::BUFFER_STORE_DWORD,
+                            ShaderOpcode::S_ENDPGM};
+        lod_test.required_spirv = {"OpImageSampleExplicitLod"};
+        AppendVMovLiteral(&lod_test.code, 20, std::bit_cast<uint32_t>(0.5f));
+        AppendVMovLiteral(&lod_test.code, 21, std::bit_cast<uint32_t>(0.5f));
+        lod_test.code.push_back(EncodeMimg0(0x27, 1));
+        lod_test.code.push_back(EncodeMimg1(0, 20, 0, 2));
+        AppendStoreVgpr(&lod_test.code, 0, 0);
+        AppendEnd(&lod_test.code);
+        const auto lod_program = CompileCase(lod_test, SubgroupSize());
+#if KYTY_PLATFORM != KYTY_PLATFORM_WINDOWS
+        ExpectFatal("TextureMinLodBeyondView", [&] {
+          auto invalid_lod = lod_descriptor;
+          invalid_lod.dwords[1] |= 1024u << 8u;
+          (void)RenderExecutorTestAccess::ResolveTexture(
+              executor, lod_program.program.info.images[0], invalid_lod);
+        });
+#endif
+        const auto lod_binding = RenderExecutorTestAccess::ResolveTexture(
+            executor, lod_program.program.info.images[0], lod_descriptor);
+        for (uint32_t mip = 0; mip < 4; ++mip) {
+          vk::ClearValue clear{};
+          clear.color.float32[0] = static_cast<float>(1u << mip);
+          TextureCacheTestAccess::ClearImage(
+              texture_cache, scheduler.Current(), lod_binding.image_id,
+              {vk::ImageAspectFlagBits::eColor, mip, 1, 0, 1}, clear);
+        }
+        const auto sampler = context.GetSamplerCache().GetSampler(lod_sampler);
+        auto output = CreateStorageBuffer(lod_test.name, {}, 1);
+        struct LodCase {
+          uint32_t base_level;
+          uint32_t min_lod;
+          float expected;
+          float integer_min_lod_expected;
+        };
+        constexpr std::array lod_cases{
+            LodCase{0, 0, 1.0f, 1.0f},   LodCase{0, 256, 2.0f, 2.0f},
+            LodCase{0, 384, 3.0f, 2.0f}, LodCase{1, 384, 3.0f, 2.0f},
+            LodCase{1, 512, 4.0f, 4.0f}, LodCase{1, 0, 2.0f, 2.0f},
+            LodCase{0, 256, 2.0f, 2.0f}};
+        std::array<vk::ImageView, lod_cases.size()> views{};
+        for (size_t index = 0; index < lod_cases.size(); ++index) {
+          const auto &test = lod_cases[index];
+          lod_descriptor.dwords[1] =
+              (lod_descriptor.dwords[1] & ~0xfff00u) | (test.min_lod << 8u);
+          lod_descriptor.dwords[3] =
+              (lod_descriptor.dwords[3] & ~(0xfu << 12u)) |
+              (test.base_level << 12u);
+          std::copy_n(lod_descriptor.dwords.begin(), 8,
+                      lod_test.user_data.begin());
+          const auto binding = RenderExecutorTestAccess::ResolveTexture(
+              executor, lod_program.program.info.images[0], lod_descriptor);
+          views[index] =
+              texture_cache.FindTexture(binding.image_id, binding.desc);
+          auto &image = texture_cache.GetImage(binding.image_id);
+          image.Transit(vk::ImageLayout::eShaderReadOnlyOptimal,
+                        vk::AccessFlagBits2::eShaderRead, {},
+                        scheduler.Current().Handle());
+          Image sampled;
+          sampled.view = views[index];
+          sampled.layout = image.backing.state.layout;
+          scheduler.Finish();
+          Dispatch(lod_test, lod_program, output, nullptr, &sampled, nullptr,
+                   nullptr, sampler);
+          const auto result = ReadBuffer(lod_test.name, output, 1)[0];
+          // Vulkan permits flooring imageViewMinLod instead of retaining its
+          // fraction.
+          Require(lod_test.name, "sampled mip",
+                  result == std::bit_cast<uint32_t>(test.expected) ||
+                      result == std::bit_cast<uint32_t>(
+                                    test.integer_min_lod_expected),
+                  "texture clamp was lost, applied before sampler max LOD, or "
+                  "rebased incorrectly");
+        }
+        Require(lod_test.name, "view identity",
+                views[0] != views[1] && views[1] != views[2] &&
+                    views[2] != views[3] && views[1] == views.back(),
+                "different texture clamps aliased or identical clamped views "
+                "were not reused");
+        DestroyBuffer(&output);
+        RenderExecutorTestAccess::ResetBindings(executor);
+      }
+
       auto srgb_storage = storage;
       constexpr auto srgb_format =
           static_cast<uint32_t>(Prospero::BufferFormat::k8_8_8_8Srgb);
@@ -12120,7 +12329,7 @@ public:
     return ret;
   }
 
-  vk::Sampler CreateNearestSampler(const char *shader_name) {
+  vk::Sampler CreateNearestSampler(const char *shader_name, u32 mip_levels = 1) {
     vk::SamplerCreateInfo sampler_info{};
     sampler_info.sType = vk::StructureType::eSamplerCreateInfo;
     sampler_info.magFilter = vk::Filter::eNearest;
@@ -12130,7 +12339,7 @@ public:
     sampler_info.addressModeV = vk::SamplerAddressMode::eClampToEdge;
     sampler_info.addressModeW = vk::SamplerAddressMode::eClampToEdge;
     sampler_info.minLod = 0.0f;
-    sampler_info.maxLod = 0.0f;
+    sampler_info.maxLod = static_cast<float>(mip_levels - 1u);
     vk::Sampler sampler = nullptr;
     RequireVk(shader_name, "dispatch",
               m_device.createSampler(&sampler_info, nullptr, &sampler),
@@ -13056,8 +13265,6 @@ public:
       check_stencil("masked compare passes", false, 0xb8, 0xb9);
       check_stencil("masked compare fails", false, 0xb0, 0xb0);
 
-      DestroyBuffer(&stencil_readback);
-
       // SPI_PS_IN_CONTROL must select the native wave width through the actual
       // program cache. In wave32 a low-word compare preserves scalar VCC_HI.
       static const auto native_vertex = [] {
@@ -13209,6 +13416,117 @@ public:
                 "zero positions drew an extra triangle, or a valid position was culled");
       }
 
+      // A NULL-export stencil pass must ignore smaller stale color targets,
+      // materialize the whole pending HTile clear, and still execute pixel discard.
+      static const auto null_pixels = [] {
+        std::array<std::vector<u32>, 2> result;
+        for (uint32_t discard = 0; discard < result.size(); discard++) {
+          auto &code = result[discard];
+          if (discard != 0) {
+            code.push_back(EncodeSop1(0x04, 126, InlineU32(0)));
+          }
+          code.push_back(EncodeExp0(0x09, 0, true, false, true));
+          code.push_back(EncodeExp1(0, 0, 0, 0));
+          AppendEnd(&code);
+        }
+        return result;
+      }();
+      registers.Reset();
+      registers.SetViewportTransformControl(0x300);
+      registers.SetViewportScaleOffset(0, extent / 2, extent / 2,
+                                       extent / 2, extent / 2, 1, 0);
+      registers.SetViewportZMax(0, 1);
+      registers.SetScreenScissor(0, 0, extent, extent);
+      registers.SetWindowScissor(0, 0, extent, extent, false);
+      registers.SetGenericScissor(0, 0, extent, extent, false);
+      registers.SetViewportScissor(0, 0, 0, extent, extent, false);
+      registers.SetRenderTargetMask(0xf00f);
+      registers.SetShaderMask(0xf00f);
+      registers.SetPsInControl(0x8000);
+      registers.SetDepthShaderControl({.shader_kill_enable = true});
+      for (const auto slot : {0u, 3u}) {
+        registers.SetColorBase(slot,
+            {.addr = depth_address + (slot == 0 ? 0x38000 : 0x10000)});
+        registers.SetColorInfo(slot,
+            {.format = Prospero::ChannelLayout::k32_32_32_32,
+             .channel_type = Prospero::ChannelType::kFloat,
+             .channel_order = Prospero::ChannelOrder::kStandard});
+        const auto side = slot == 0 ? extent / 4 : extent;
+        registers.SetColorAttrib2(slot, {.height = side - 1, .width = side - 1});
+        registers.SetColorAttrib3(slot,
+            {.tile_mode = Prospero::TileMode::kLinear, .dimension = 1});
+        registers.SetTargetOutputMode(slot, 4);
+      }
+      stencil_target.z_info.htile_acceleration = true;
+      stencil_target.htile_data_base_addr = depth_address + 0x30000;
+      registers.SetDepthRenderTarget(stencil_target);
+      registers.SetDepthClearValue(1);
+      stencil_control = {};
+      stencil_control.stencil_enable = true;
+      stencil_control.stencilfunc = static_cast<uint8_t>(vk::CompareOp::eAlways);
+      registers.SetDepthControl(stencil_control);
+      registers.SetStencilControl({3, 3, 3, 3, 3, 3});
+      registers.SetStencilMask({0x80, 0, 0x80, 1, 0x80, 0, 0x80, 1});
+      user_config.SetPrimitiveType(Prospero::PrimitiveType::kTriList);
+      shaders.SetEsShaderBase(vertex_address);
+      for (uint32_t discard = 0; discard < null_pixels.size(); discard++) {
+        const auto &code = null_pixels[discard];
+        const auto address = reinterpret_cast<uint64_t>(code.data());
+        ShaderMapUserData(address,
+            {.type = Prospero::ShaderBinaryType::kPs,
+             .code_size_bytes = static_cast<uint32_t>(code.size() * sizeof(u32))});
+        shaders.SetPsShaderBase(address);
+        RenderExecutorTestAccess::ResolveRenderDepthTarget(executor, scheduler.Current(), depth);
+        (void)cache.FindDepthTarget(depth.image_id, depth.desc);
+        vk::ClearValue zero{};
+        TextureCacheTestAccess::ClearImage(cache, scheduler.Current(), depth.image_id,
+            {vk::ImageAspectFlagBits::eDepth | vk::ImageAspectFlagBits::eStencil,
+             0, 1, 0, 1}, zero);
+        Require(name, "pending stencil-pass HTile clear",
+                cache.ClearMeta(stencil_target.htile_data_base_addr),
+                "the stencil pass did not retain its HTile metadata");
+        RenderExecutorTestAccess::DrawAuto(
+            executor, scheduler.Current(), {.vertex_count = 3, .instance_count = 1});
+        const auto depth_pixels = ReadCachedTexel(name, context, depth.image_id,
+                                                  {}, {extent, extent, 1});
+        Require(name, "NULL-export full depth clear",
+                std::ranges::all_of(depth_pixels, [](u32 v) { return v == 0x3f800000u; }),
+                "an unexported color target clipped the pending depth clear");
+        const auto stencil = read_stencil();
+        const auto expected = discard == 0 ? 0x80808080u : 0u;
+        Require(name, "NULL-export stencil coverage and discard",
+                std::ranges::all_of(stencil, [=](u32 v) { return v == expected; }),
+                "a stale color target clipped stencil coverage or the NULL-export shader was skipped");
+      }
+
+      // With MRT0 still bound, an MRT3-only export must retain location3 in
+      // rendering attachments, pipeline formats, blend masks and dynamic write enables.
+      static const auto sparse_pixel = [] {
+        auto code = native_pixel;
+        *std::ranges::find(code, EncodeExp0(0x00, 0xf)) = EncodeExp0(0x03, 0xf);
+        return code;
+      }();
+      const auto sparse_address = reinterpret_cast<uint64_t>(sparse_pixel.data());
+      ShaderMapUserData(sparse_address,
+          {.type = Prospero::ShaderBinaryType::kPs,
+           .code_size_bytes = static_cast<uint32_t>(sparse_pixel.size() * sizeof(u32))});
+      shaders.SetPsShaderBase(sparse_address);
+      registers.SetDepthControl({});
+      registers.SetDepthShaderControl({});
+      RenderExecutorTestAccess::DrawAuto(
+            executor, scheduler.Current(), {.vertex_count = 3, .instance_count = 1});
+      RenderColorInfo sparse_color{};
+      RenderExecutorTestAccess::ResolveRenderColorTarget(
+          executor, scheduler.Current(), sparse_color, 3);
+      const auto sparse_pixels = ReadCachedTexel(name, context, sparse_color.image_id,
+                                                {}, {extent, extent, 1});
+      for (size_t component = 0; component < sparse_pixels.size(); component++) {
+        const auto expected = component % 4 == 3 ? 0x3f800000u : 0x3e800000u;
+        Require(name, "sparse MRT3 output", sparse_pixels[component] == expected,
+                "MRT3 was compacted to a different slot or clipped by unused MRT0");
+      }
+      DestroyBuffer(&stencil_readback);
+
       vertex_shader = owned_vertex_shader;
       pixel_shader = owned_pixel_shader;
     }
@@ -14662,9 +14980,11 @@ private:
     available_feedback_dynamic.pNext = &available_feedback_layout;
     vk::PhysicalDeviceProvokingVertexFeaturesEXT available_provoking_vertex{};
     available_provoking_vertex.pNext = &available_feedback_dynamic;
+    vk::PhysicalDeviceImageViewMinLodFeaturesEXT available_min_lod{};
+    available_min_lod.pNext = &available_provoking_vertex;
     vk::PhysicalDeviceFeatures2 available_features2{};
     available_features2.sType = vk::StructureType::ePhysicalDeviceFeatures2;
-    available_features2.pNext = &available_provoking_vertex;
+    available_features2.pNext = &available_min_lod;
     m_physical_device.getFeatures2(&available_features2);
     Require("VulkanHarness", "dispatch",
             available_features.shaderStorageImageWriteWithoutFormat == true,
@@ -14690,6 +15010,8 @@ private:
     Require("VulkanHarness", "dispatch",
             available_features12.bufferDeviceAddress == true,
             "bufferDeviceAddress is not supported");
+    Require("VulkanHarness", "dispatch", available_min_lod.minLod == true,
+            "image view minimum LOD is not supported");
     Require("VulkanHarness", "graphics", available_features12.shaderOutputLayer == true,
             "vertex layer output is not supported");
     Require("VulkanHarness", "graphics", available_features.fillModeNonSolid &&
@@ -14748,7 +15070,10 @@ private:
     vk::PhysicalDeviceProvokingVertexFeaturesEXT provoking_vertex{};
     provoking_vertex.pNext = &feedback_dynamic;
     provoking_vertex.provokingVertexLast = available_provoking_vertex.provokingVertexLast;
-    device_info.pNext = &provoking_vertex;
+    vk::PhysicalDeviceImageViewMinLodFeaturesEXT min_lod{};
+    min_lod.pNext = &provoking_vertex;
+    min_lod.minLod = true;
+    device_info.pNext = &min_lod;
     vk::PhysicalDeviceFeatures device_features{};
     device_features.shaderStorageImageWriteWithoutFormat = true;
     device_features.shaderImageGatherExtended = true;
@@ -14765,7 +15090,8 @@ private:
         VK_EXT_COLOR_WRITE_ENABLE_EXTENSION_NAME,
         VK_EXT_ATTACHMENT_FEEDBACK_LOOP_LAYOUT_EXTENSION_NAME,
         VK_EXT_ATTACHMENT_FEEDBACK_LOOP_DYNAMIC_STATE_EXTENSION_NAME,
-        VK_EXT_PROVOKING_VERTEX_EXTENSION_NAME};
+        VK_EXT_PROVOKING_VERTEX_EXTENSION_NAME,
+        VK_EXT_IMAGE_VIEW_MIN_LOD_EXTENSION_NAME};
     device_info.enabledExtensionCount = std::size(device_extensions);
     device_info.ppEnabledExtensionNames = device_extensions;
     RequireVk("VulkanHarness", "dispatch",
@@ -15216,7 +15542,7 @@ void RunCase(VulkanHarness *vulkan, const TestCase &test) {
         vk::ImageLayout::eGeneral);
   }
   if (needs_sampler) {
-    sampler = vulkan->CreateNearestSampler(test.name);
+    sampler = vulkan->CreateNearestSampler(test.name, sampled_image.mip_levels);
   }
 
   vulkan->Dispatch(test, compiled, buffer, needs_gds ? &gds_buffer : nullptr,
@@ -16871,6 +17197,251 @@ TestCase Scalar64BitOps() {
            O::V_MOV_B32, O::BUFFER_STORE_DWORD, O::S_ENDPGM}};
 }
 
+uint64_t ReferenceAshrI64(uint64_t bits, u32 count) {
+  const auto shift = count & 63u;
+  // Complementing negative values expresses sign fill using unsigned shifts.
+  return (bits & (uint64_t{1} << 63u)) != 0 ? ~(~bits >> shift) : bits >> shift;
+}
+
+TestCase ScalarAshrI64Edges(bool dynamic) {
+  using O = ShaderOpcode;
+  TestCase test;
+  test.name =
+      dynamic ? "ScalarAshrI64DynamicEdges" : "ScalarAshrI64ConstantEdges";
+  constexpr uint64_t sources[] = {0,
+                                  1,
+                                  0x80000000ull,
+                                  0x100000000ull,
+                                  0x7fffffffffffffffull,
+                                  0x8000000000000000ull,
+                                  0xffffffffffffffffull,
+                                  0x0123456789abcdefull,
+                                  0xfedcba9876543210ull};
+  constexpr u32 counts[] = {0,  1,  7,  31, 32,  33,         63,
+                            64, 65, 95, 96, 127, 0xffffffffu};
+  test.initial.assign(std::begin(counts), std::end(counts));
+  for (const auto source : sources) {
+    test.initial.push_back(static_cast<u32>(source));
+    test.initial.push_back(static_cast<u32>(source >> 32u));
+  }
+  test.expected = test.initial;
+  auto &code = test.code;
+  const auto load_scalar = [&](u32 reg, u32 index) {
+    AppendVMovU32(&code, 30, index * sizeof(u32));
+    AppendBufferLoadDword(&code, 1, 30);
+    code.push_back(EncodeVop1(0x02, reg, Vgpr(1)));
+  };
+  for (u32 source_index = 0; source_index < std::size(sources);
+       ++source_index) {
+    const auto source = sources[source_index];
+    // Runtime sources exercise both constant-count and dynamic-count emission.
+    load_scalar(20, static_cast<u32>(std::size(counts)) + source_index * 2u);
+    load_scalar(21,
+                static_cast<u32>(std::size(counts)) + source_index * 2u + 1u);
+    for (u32 index = 0; index < std::size(counts); ++index) {
+      const auto expected = ReferenceAshrI64(source, counts[index]);
+      if (dynamic) {
+        load_scalar(22, index);
+      }
+      // Start SCC at the opposite value so every result must update it.
+      code.push_back(
+          EncodeSopc(0x06, InlineU32(0), InlineU32(expected != 0 ? 1u : 0u)));
+      const auto count_src = dynamic                ? 22u
+                             : counts[index] <= 64u ? InlineU32(counts[index])
+                                                    : 255u;
+      code.push_back(EncodeSop2(0x23, 24, 20, count_src));
+      if (!dynamic && count_src == 255u) {
+        code.push_back(counts[index]);
+      }
+      code.push_back(EncodeSMovB32(26, 253));
+      const auto output = static_cast<u32>(test.expected.size());
+      AppendStoreSgprPair(&code, 24, output);
+      AppendStoreSgpr(&code, 26, output + 2u);
+      test.expected.push_back(static_cast<u32>(expected));
+      test.expected.push_back(static_cast<u32>(expected >> 32u));
+      test.expected.push_back(expected != 0 ? 1u : 0u);
+    }
+  }
+  AppendEnd(&code);
+  test.initial.resize(test.expected.size());
+  test.opcodes = {O::BUFFER_LOAD_DWORD,  O::V_READFIRSTLANE_B32,
+                  O::S_MOV_B32,          O::S_CMP_EQ_U32,
+                  O::S_ASHR_I64,         O::V_MOV_B32,
+                  O::BUFFER_STORE_DWORD, O::S_ENDPGM};
+  test.decoded_counts = {
+      {"S_ASHR_I64", std::size(sources) * std::size(counts)}};
+  test.required_spirv = {"OpShiftRightArithmetic"};
+  test.forbidden_spirv = {"OpTypeInt 64"};
+  return test;
+}
+
+TestCase ScalarAshrI64OperandsAndAliases() {
+  using O = ShaderOpcode;
+  TestCase test;
+  test.name = "ScalarAshrI64OperandsAndAliases";
+  test.initial = {33}; // Runtime shift for signed literal source tests.
+  test.expected = test.initial;
+  auto &code = test.code;
+  AppendVMovU32(&code, 30, 0);
+  AppendBufferLoadDword(&code, 1, 30);
+  code.push_back(EncodeVop1(0x02, 22, Vgpr(1)));
+  size_t instructions = 0;
+  const auto check_result = [&](u32 dst, u32 src0, u32 src1, uint64_t value,
+                                u32 count, std::optional<u32> literal = {}) {
+    code.push_back(EncodeSop2(0x23, dst, src0, src1));
+    if (literal) {
+      code.push_back(*literal);
+    }
+    code.push_back(EncodeSMovB32(26, 253));
+    const auto output = static_cast<u32>(test.expected.size());
+    AppendStoreSgprPair(&code, dst, output);
+    AppendStoreSgpr(&code, 26, output + 2u);
+    const auto expected = ReferenceAshrI64(value, count);
+    test.expected.push_back(static_cast<u32>(expected));
+    test.expected.push_back(static_cast<u32>(expected >> 32u));
+    test.expected.push_back(expected != 0 ? 1u : 0u);
+    ++instructions;
+  };
+  for (const u32 literal : {0u, 0x7fffffffu, 0x80000000u, 0xffffffffu}) {
+    const uint64_t value =
+        literal | ((literal & 0x80000000u) != 0 ? 0xffffffff00000000ull : 0ull);
+    for (const u32 count : {0u, 1u, 31u, 32u, 63u, 64u}) {
+      check_result(24, 255, InlineU32(count), value, count, literal);
+    }
+    check_result(24, 255, 22, value, 33, literal);
+  }
+  for (const u32 source : {128u, 129u, 192u, 193u, 208u, 125u}) {
+    const uint64_t value = source == 125u   ? 0ull
+                           : source <= 192u ? source - 128u
+                                            : uint64_t{192} - source;
+    for (const u32 count : {0u, 1u, 32u, 63u}) {
+      check_result(24, source, InlineU32(count), value, count);
+    }
+  }
+  for (const u32 scc : {0u, 1u}) {
+    // SCC may be either source and must be read before it is overwritten.
+    code.push_back(EncodeSopc(0x06, InlineU32(scc), InlineU32(1)));
+    check_result(24, 253, InlineU32(1), scc, 1);
+    AppendSMovLiteral(&code, 20, 1);
+    AppendSMovLiteral(&code, 21, 0);
+    code.push_back(EncodeSopc(0x06, InlineU32(scc), InlineU32(1)));
+    check_result(24, 20, 253, 1, scc);
+  }
+  constexpr uint64_t value = 0xfedcba9889abcdefull;
+  for (const u32 count : {0u, 1u, 31u, 32u, 33u, 63u, 65u}) {
+    AppendSMovLiteral(&code, 20, static_cast<u32>(value));
+    AppendSMovLiteral(&code, 21, static_cast<u32>(value >> 32u));
+    AppendSMovLiteral(&code, 23, count);
+    check_result(20, 20, 23, value, count); // In-place source/destination pair.
+    for (const u32 count_reg : {24u, 25u}) {
+      AppendSMovLiteral(&code, 20, static_cast<u32>(value));
+      AppendSMovLiteral(&code, 21, static_cast<u32>(value >> 32u));
+      AppendSMovLiteral(&code, count_reg, count);
+      check_result(24, 20, count_reg, value,
+                   count); // Count aliases destination.
+    }
+  }
+  for (const u32 count_reg : {20u, 21u}) {
+    const uint64_t alias_value =
+        count_reg == 20u ? 0x8000000000000021ull : 0x0000002100000001ull;
+    AppendSMovLiteral(&code, 20, static_cast<u32>(alias_value));
+    AppendSMovLiteral(&code, 21, static_cast<u32>(alias_value >> 32u));
+    check_result(20, 20, count_reg, alias_value, 33);
+  }
+  // Both operands may refer to the single literal word in the instruction.
+  check_result(24, 255, 255, 0xffffffffffffffffull, 0xffffffffu, 0xffffffffu);
+  for (const u32 source : {0u, 1u}) {
+    // Discarding the destination must still update SCC and leave M0 untouched.
+    AppendSMovLiteral(&code, 124, 0x5a13abcdu);
+    AppendSMovLiteral(&code, 20, 0);
+    AppendSMovLiteral(&code, 21, source);
+    code.push_back(EncodeSopc(0x06, InlineU32(0), InlineU32(source)));
+    code.push_back(EncodeSop2(0x23, 125, 20, InlineU32(0)));
+    code.push_back(EncodeSMovB32(26, 253));
+    const auto output = static_cast<u32>(test.expected.size());
+    AppendStoreSgpr(&code, 26, output);
+    AppendStoreSgpr(&code, 124, output + 1u);
+    AppendStoreSgprPair(&code, 20, output + 2u);
+    test.expected.insert(test.expected.end(), {source, 0x5a13abcdu, 0, source});
+    ++instructions;
+  }
+  AppendEnd(&code);
+  test.initial.resize(test.expected.size());
+  test.opcodes = {O::BUFFER_LOAD_DWORD,  O::V_READFIRSTLANE_B32,
+                  O::S_MOV_B32,          O::S_CMP_EQ_U32,
+                  O::S_ASHR_I64,         O::V_MOV_B32,
+                  O::BUFFER_STORE_DWORD, O::S_ENDPGM};
+  test.decoded_counts = {{"S_ASHR_I64", instructions}};
+  test.forbidden_spirv = {"OpTypeInt 64"};
+  return test;
+}
+
+TestCase ScalarAshrI64Masks(u32 wave_size) {
+  using O = ShaderOpcode;
+  TestCase test;
+  test.name =
+      wave_size == 64 ? "ScalarAshrI64MasksWave64" : "ScalarAshrI64MasksWave32";
+  auto &code = test.code;
+  struct Case {
+    uint64_t source;
+    u32 count;
+  };
+  constexpr Case cases[] = {
+      {0x8000000100000001ull, 1},  {0x8000000100000001ull, 32},
+      {0x8000000100000001ull, 63}, {0x0000000100000000ull, 0},
+      {0x0000000100000000ull, 32}, {1, 1}};
+  for (const u32 mask_reg : {106u, 126u}) {
+    for (const auto &item : cases) {
+      const auto result = ReferenceAshrI64(item.source, item.count);
+      AppendVMovU32(&code, 1, 1);
+      AppendVMovU32(&code, 2, 0);
+      AppendSMovLiteral(&code, mask_reg, static_cast<u32>(item.source));
+      AppendSMovLiteral(&code, mask_reg + 1u,
+                        static_cast<u32>(item.source >> 32u));
+      code.push_back(
+          EncodeSopc(0x06, InlineU32(0), InlineU32(result != 0 ? 1u : 0u)));
+      code.push_back(
+          EncodeSop2(0x23, mask_reg, mask_reg, InlineU32(item.count)));
+      code.push_back(EncodeSMovB32(20, mask_reg));
+      code.push_back(EncodeSMovB32(21, mask_reg + 1u));
+      code.push_back(EncodeSMovB32(22, 253));
+      if (mask_reg == 106u) {
+        code.push_back(EncodeVop2(0x01, 2, InlineU32(0), 1));
+      } else {
+        code.push_back(EncodeVop1(0x01, 2, InlineU32(1)));
+      }
+      code.push_back(EncodeSop1(0x04, 126, 193)); // Restore EXEC for readback.
+      const u32 expected[] = {static_cast<u32>(result),
+                              static_cast<u32>(result >> 32u),
+                              result != 0 ? 1u : 0u};
+      for (u32 i = 0; i < std::size(expected); ++i) {
+        AppendStoreSgprAtLaneDwordOffset(
+            &code, 20 + i, 0, static_cast<u32>(test.expected.size()));
+        test.expected.insert(test.expected.end(), wave_size, expected[i]);
+      }
+      AppendStoreVgprAtLaneDwordOffset(&code, 2, 0,
+                                       static_cast<u32>(test.expected.size()));
+      for (u32 lane = 0; lane < wave_size; ++lane) {
+        test.expected.push_back(static_cast<u32>((result >> lane) & 1u));
+      }
+    }
+  }
+  AppendEnd(&code);
+  test.initial.resize(test.expected.size());
+  test.opcodes = {O::S_MOV_B32,     O::S_MOV_B64,    O::S_CMP_EQ_U32,
+                  O::S_ASHR_I64,    O::V_MOV_B32,    O::V_CNDMASK_B32,
+                  O::V_LSHLREV_B32, O::V_ADD_NC_U32, O::BUFFER_STORE_DWORD,
+                  O::S_ENDPGM};
+  test.decoded_counts = {{"S_ASHR_I64 vcc_lo, vcc_lo", std::size(cases)},
+                         {"S_ASHR_I64 exec_lo, exec_lo", std::size(cases)}};
+  test.forbidden_spirv = {"OpTypeInt 64"};
+  test.compute_info.threads_num[0] = wave_size;
+  test.compute_info.threads_num[1] = test.compute_info.threads_num[2] = 1;
+  test.compute_info.thread_ids_num = 1;
+  test.compute_info.wave_size = wave_size;
+  test.has_compute_info = true;
+  return test;
+}
 TestCase ScalarConditionalMoveB64() {
   using O = ShaderOpcode;
   TestCase test;
@@ -24733,6 +25304,34 @@ TestCase ImageSampleAndGather() {
   return test;
 }
 
+TestCase ImageGatherLodApproximatesLevelZero() {
+  using O = ShaderOpcode;
+  std::vector<u32> code;
+  AppendVMovLiteral(&code, 20, 0x3f000000u);
+  AppendVMovLiteral(&code, 21, 0x3f000000u);
+  AppendVMovLiteral(&code, 22, 0x3f800000u);
+  code.push_back(EncodeMimg0(0x44, 0x1));
+  code.push_back(EncodeMimg1(0, 20));
+  code.push_back(EncodeMimg0(0x24, 0x1));
+  code.push_back(EncodeMimg1(4, 20));
+  for (u32 i = 0; i < 5u; i++) {
+    AppendStoreVgpr(&code, i, i);
+  }
+  AppendEnd(&code);
+
+  TestCase test;
+  test.name = "ImageGatherLodApproximatesLevelZero";
+  test.code = std::move(code);
+  test.opcodes = {O::V_MOV_B32, O::IMAGE_GATHER4_L, O::IMAGE_SAMPLE,
+                  O::BUFFER_STORE_DWORD, O::S_ENDPGM};
+  // Explicit sampling proves mip 1 is accessible; the gather must still read mip 0.
+  test.expected = {0x3f800000u, 0x3f800000u, 0x3f800000u, 0x3f800000u, 0x41000000u};
+  test.sampled_image_rgba_mips = {std::vector<u32>(4 * 4 * 4, 0x3f800000u),
+                                 std::vector<u32>(2 * 2 * 4, 0x41000000u)};
+  test.required_spirv = {"OpImageGather", "OpImageSampleExplicitLod"};
+  return test;
+}
+
 TestCase ImageD16GatherPacksHalfPairs() {
   using O = ShaderOpcode;
 
@@ -26057,6 +26656,11 @@ std::vector<TestCase> MakeCases() {
   AddCase(ScalarBfeI32CapturedRawSignExtends);
   AddCase(BitfieldExtractWidthPastEndEdges);
   AddCase(Scalar64BitOps);
+  cases.push_back(ScalarAshrI64Edges(false));
+  cases.push_back(ScalarAshrI64Edges(true));
+  AddCase(ScalarAshrI64OperandsAndAliases);
+  cases.push_back(ScalarAshrI64Masks(32));
+  cases.push_back(ScalarAshrI64Masks(64));
   AddCase(VectorDynamicU64ShiftEdges);
   AddCase(ScalarConditionalMoveB64);
   AddCase(ScalarConditionalMoveB64PreservesMasks);
@@ -26307,6 +26911,7 @@ std::vector<TestCase> MakeCases() {
   AddCase(ImageGetResinfoDmaskWidthHeight);
   AddCase(ImageGetResinfoDmaskMipLevels);
   AddCase(ImageSampleAndGather);
+  AddCase(ImageGatherLodApproximatesLevelZero);
   AddCase(ImageD16GatherPacksHalfPairs);
   AddCase(ImageSampleA16SamplerCoordsOnGpu);
   AddCase(ImageSampleOpcodeAliasUsesNormalCoords);
@@ -27972,12 +28577,6 @@ ShaderTextureResource AtomicStorageTextureDescriptor() {
     descriptor.fields[3] =
         (descriptor.fields[3] & ~(0x1fu << 20u)) |
         (static_cast<uint32_t>(Prospero::TileMode::kStandard256B) << 20u);
-  } else if (std::strcmp(kind, "base-mip-out-of-resource") == 0) {
-    descriptor.fields[3] |= (1u << 12u) | (1u << 16u);
-  } else if (std::strcmp(kind, "dynamic-mip-out-of-resource") == 0) {
-    resource.mip_mode = ShaderRecompiler::IR::ImageMipMode::DynamicStorage;
-    resource.mip_count = 2;
-    descriptor.fields[3] |= 1u << 16u;
   } else if (std::strcmp(kind, "inverted-mip-range") == 0) {
     descriptor.fields[3] |= 1u << 12u;
     descriptor.fields[5] |= 1u << 4u;
@@ -28449,8 +29048,6 @@ void CheckBasicStorageTextureDescriptor() {
   for (const char *kind : {"resource",
                            "type",
                            "standard256b-volume",
-                           "base-mip-out-of-resource",
-                           "dynamic-mip-out-of-resource",
                            "inverted-mip-range",
                            "swizzle",
                            "linear-rgb1-read",
@@ -30827,6 +31424,15 @@ int main(int argc, char **argv) {
   std::setvbuf(stdout, nullptr, _IONBF, 0);
   EnsureConfigInitialized();
   CheckLeastRecentlyUsedCacheOrdering();
+  if (argc == 2 && std::strcmp(argv[1], "--s-ashr-i64-only") == 0) {
+    VulkanHarness vulkan;
+    RunCase(&vulkan, ScalarAshrI64Edges(false));
+    RunCase(&vulkan, ScalarAshrI64Edges(true));
+    RunCase(&vulkan, ScalarAshrI64OperandsAndAliases());
+    RunCase(&vulkan, ScalarAshrI64Masks(32));
+    RunCase(&vulkan, ScalarAshrI64Masks(64));
+    return 0;
+  }
   if (argc == 2 && std::strcmp(argv[1], "--tessellation-only") == 0) {
     CheckTessellationPrograms();
     CheckEmbeddedFetchVertexOffset();
@@ -30906,6 +31512,7 @@ int main(int argc, char **argv) {
     RunCase(&vulkan, BranchVccnzUsesWaveMask());
     RunCase(&vulkan, BranchVccnzUsesCarryProducedWaveMask());
     RunCase(&vulkan, ImageSampleAndGather());
+    RunCase(&vulkan, ImageGatherLodApproximatesLevelZero());
     RunCase(&vulkan, SharedReturnKeepsSelectedValues());
     return 0;
   }

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,470 @@ 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[1].element[0] = 0x1234;
+	fpstate._xmm[8].element[0] = 0xffffc4c8;
+	emulate({0x66, 0x41, 0x0f, 0x79, 0xc8}, 5); // observed fault: extrq xmm1, xmm8
+	Check(test, fpstate._xmm[1].element[0] == 0x23 && fpstate._xmm[8].element[0] == 0xffffc4c8,
+	      "register EXTRQ lost its source controls, ignored bits, or separate destination");
+	fpstate._xmm[11].element[0] = 0x89abcdef;
+	fpstate._xmm[11].element[1] = 0x01234567;
+	fpstate._xmm[0].element[0] = 0x2010;
+	emulate({0x66, 0x44, 0x0f, 0x79, 0xd8}, 5); // extrq xmm11, xmm0
+	Check(test, fpstate._xmm[11].element[0] == 0x4567 && fpstate._xmm[11].element[1] == 0,
+	      "register EXTRQ lost its extended destination or 64-bit extraction");
+	fpstate._xmm[8].element[0] = 0x89abcdef;
+	fpstate._xmm[8].element[1] = 0x01234567;
+	fpstate._xmm[9].element[0] = 0;
+	emulate({0x66, 0x45, 0x0f, 0x79, 0xc1}, 5); // extrq xmm8, xmm9
+	Check(test, fpstate._xmm[8].element[0] == 0x89abcdef && fpstate._xmm[8].element[1] == 0x01234567,
+	      "register EXTRQ did not interpret zero length as 64 bits");
+	fpstate._xmm[3].element[0] = 0xab0408;
+	emulate({0x66, 0x0f, 0x79, 0xdb}, 4); // extrq xmm3, xmm3
+	Check(test, fpstate._xmm[3].element[0] == 0x40,
+	      "register EXTRQ overwrote aliased controls before reading them");
+	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/macosSse4aTests.cpp

@@ -0,0 +1,219 @@
+#include "common/hostException.h"
+#include "loader/x64InstructionEmulator.h"
+
+#include <atomic>
+#include <bit>
+#include <cstddef>
+#include <cstdint>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <sys/mman.h>
+#include <sys/ucontext.h>
+#include <unistd.h>
+
+namespace {
+
+using Common::HostException::ExceptionInfo;
+using Common::HostException::ExceptionType;
+
+// The SSE4a bytes are the point of this test, so blobs are spelled out byte by
+// byte instead of going through an encoder. Every blob is a complete routine:
+// the instruction under test, a store of the result, and ret. Encodings were
+// disassembled with otool to confirm the operand roles: EXTRQ immediate is
+// unary (ModRM.r/m is both source and destination, the reg field is unused),
+// INSERTQ immediate reads ModRM.r/m into the reg field.
+constexpr uint8_t kExtrqXmm2Len40Index0[] = {
+	0x66, 0x0f, 0x78, 0xc2, 0x28, 0x00, // extrq xmm2, 40, 0
+	0xf3, 0x0f, 0x7f, 0x17,             // movdqu %xmm2, (%rdi)
+	0xc3,                               // ret
+};
+
+constexpr uint8_t kExtrqXmm2Len64Index0[] = {
+	0x66, 0x0f, 0x78, 0xc2, 0x40, 0x00, // extrq xmm2, 64, 0
+	0xf3, 0x0f, 0x7f, 0x17,             // movdqu %xmm2, (%rdi)
+	0xc3,                               // ret
+};
+
+constexpr uint8_t kInsertqXmm2Xmm0Len32Index0[] = {
+	0xf2, 0x0f, 0x78, 0xd0, 0x20, 0x00, // insertq xmm2, xmm0, 32, 0
+	0xf3, 0x0f, 0x7f, 0x17,             // movdqu %xmm2, (%rdi)
+	0xc3,                               // ret
+};
+
+// REX.B moves the operand into the high register file. xmm10 is not an argument
+// register, so the blob loads it from memory. EXTRQ has a single operand, so
+// only the destination needs loading.
+constexpr uint8_t kExtrqXmm10Len40Index0[] = {
+	0xf3, 0x44, 0x0f, 0x6f, 0x16,             // movdqu (%rsi), %xmm10
+	0x66, 0x41, 0x0f, 0x78, 0xc2, 0x28, 0x00, // extrq xmm10, 40, 0
+	0xf3, 0x44, 0x0f, 0x7f, 0x17,             // movdqu %xmm10, (%rdi)
+	0xc3,                                     // ret
+};
+
+// REX.R keeps the destination in xmm10 while REX.B puts the source in xmm9, and
+// the inserted field sits at a nonzero index.
+constexpr uint8_t kInsertqXmm10Xmm9Len16Index8[] = {
+	0xf3, 0x44, 0x0f, 0x6f, 0x0e,             // movdqu (%rsi), %xmm9
+	0xf3, 0x44, 0x0f, 0x6f, 0x12,             // movdqu (%rdx), %xmm10
+	0xf2, 0x45, 0x0f, 0x78, 0xd1, 0x10, 0x08, // insertq xmm10, xmm9, 16, 8
+	0xf3, 0x44, 0x0f, 0x7f, 0x17,             // movdqu %xmm10, (%rdi)
+	0xc3,                                     // ret
+};
+
+constexpr uint8_t kUd2[] = {
+	0x0f, 0x0b, // ud2
+	0xc3,       // ret
+};
+
+constexpr size_t kPageSize = 4096;
+
+std::atomic<int>    g_illegal_hits {0};
+std::atomic<bool>   g_refused {false};
+std::atomic<size_t> g_refuse_skip {0};
+
+static_assert(decltype(g_illegal_hits)::is_always_lock_free);
+static_assert(decltype(g_refused)::is_always_lock_free);
+static_assert(decltype(g_refuse_skip)::is_always_lock_free);
+
+void Check(bool value, const char *text) {
+	if (!value) {
+		std::fprintf(stderr, "macosSse4aTests: failed: %s\n", text);
+		std::abort();
+	}
+}
+
+bool Handler(const ExceptionInfo &info) {
+	if (info.type != ExceptionType::IllegalInstruction) {
+		return false;
+	}
+	g_illegal_hits.fetch_add(1, std::memory_order_relaxed);
+	if (Loader::X64InstructionEmulator::TryEmulate(info.native_context)) {
+		return true;
+	}
+	// A test that expects a refusal arms g_refuse_skip so the host-illegal
+	// instruction can be stepped over. Any other refusal is outside the
+	// emulator's contract, and returning there would re-run the faulting
+	// instruction forever, so end the run instead.
+	const size_t skip = g_refuse_skip.load(std::memory_order_relaxed);
+	if (skip == 0) {
+		static const char message[] = "macosSse4aTests: failed: unexpected emulator refusal\n";
+		::write(STDERR_FILENO, message, sizeof(message) - 1);
+		::_exit(1);
+	}
+	g_refused.store(true, std::memory_order_relaxed);
+	auto *saved_context = static_cast<ucontext_t *>(info.native_context);
+	saved_context->uc_mcontext->__ss.__rip += static_cast<uint64_t>(skip);
+	return true;
+}
+
+template <typename Fn, size_t N>
+Fn MapCode(const uint8_t (&code)[N]) {
+	void *page = ::mmap(nullptr, kPageSize, PROT_READ | PROT_WRITE, MAP_PRIVATE | MAP_ANON, -1, 0);
+	Check(page != MAP_FAILED, "map test code page");
+	std::memcpy(page, code, N);
+	Check(::mprotect(page, kPageSize, PROT_READ | PROT_EXEC) == 0, "make test code page executable");
+	return reinterpret_cast<Fn>(page);
+}
+
+// Blobs receive their inputs through the SysV ABI, so the xmm argument registers
+// are reachable as ordinary function parameters.
+using BlobFn     = void (*)(uint64_t *out, double xmm0, double xmm1, double xmm2);
+using LoadBlobFn = void (*)(uint64_t *out, const uint64_t *in);
+using PairBlobFn = void (*)(uint64_t *out, const uint64_t *src, const uint64_t *dest);
+using VoidBlobFn = void (*)();
+
+int IllegalHits() {
+	return g_illegal_hits.load(std::memory_order_relaxed);
+}
+
+void TestRefusals() {
+	Check(!Loader::X64InstructionEmulator::TryEmulate(nullptr), "TryEmulate refuses a null native context");
+	ucontext_t saved_context {};
+	saved_context.uc_mcontext = nullptr;
+	Check(!Loader::X64InstructionEmulator::TryEmulate(&saved_context), "TryEmulate refuses a saved context without mcontext");
+}
+
+void TestExtrqImmediate() {
+	auto     fn      = MapCode<BlobFn>(kExtrqXmm2Len40Index0);
+	uint64_t out[2] {};
+	// The pattern is already in xmm2; the emulator keeps bits [39:0] and zeroes
+	// the upper half.
+	const double pattern = std::bit_cast<double>(0xaabbccddeeff1122ull);
+	const int    before  = IllegalHits();
+	fn(out, 0.0, 0.0, pattern);
+	Check(IllegalHits() > before, "host raises SIGILL for EXTRQ");
+	Check(out[0] == 0x000000ddeeff1122ull, "extrq xmm2, 40, 0 keeps bits [39:0] of xmm2");
+	Check(out[1] == 0, "extrq xmm2, 40, 0 zero-extends the upper 64 bits");
+}
+
+void TestExtrqFullWidth() {
+	auto     fn        = MapCode<BlobFn>(kExtrqXmm2Len64Index0);
+	uint64_t out[2] {};
+	const double pattern = std::bit_cast<double>(0x0123456789abcdefull);
+	const int    before  = IllegalHits();
+	fn(out, 0.0, 0.0, pattern);
+	Check(IllegalHits() > before, "host raises SIGILL for a full width EXTRQ");
+	Check(out[0] == 0x0123456789abcdefull, "extrq xmm2, 64, 0 keeps the whole low 64 bits");
+	Check(out[1] == 0, "extrq xmm2, 64, 0 zero-extends the upper 64 bits");
+}
+
+void TestInsertqImmediate() {
+	auto     fn        = MapCode<BlobFn>(kInsertqXmm2Xmm0Len32Index0);
+	uint64_t out[2] {};
+	const double source = std::bit_cast<double>(0x1122334455667788ull);
+	const double dest   = std::bit_cast<double>(0x0123456789abcdefull);
+	const int    before = IllegalHits();
+	fn(out, source, 0.0, dest);
+	Check(IllegalHits() > before, "host raises SIGILL for INSERTQ");
+	Check(out[0] == 0x0123456755667788ull, "insertq xmm2, xmm0, 32, 0 inserts bits [31:0] of xmm0");
+}
+
+void TestRexHighRegisterExtrq() {
+	auto           fn     = MapCode<LoadBlobFn>(kExtrqXmm10Len40Index0);
+	uint64_t       out[2] {};
+	const uint64_t in[2] {0xaabbccddeeff1122ull, 0};
+	const int      before = IllegalHits();
+	fn(out, in);
+	Check(IllegalHits() > before, "host raises SIGILL for a REX.B EXTRQ");
+	Check(out[0] == 0x000000ddeeff1122ull, "REX.B destination xmm10 is emulated through the saved context");
+	Check(out[1] == 0, "REX.B destination xmm10 zero-extends the upper 64 bits");
+}
+
+void TestRexHighRegisterInsertqWithIndex() {
+	auto           fn     = MapCode<PairBlobFn>(kInsertqXmm10Xmm9Len16Index8);
+	uint64_t       out[2] {};
+	const uint64_t src[2] {0x8877665544332211ull, 0xdeadbeefcafebabeull};
+	const uint64_t dest[2] {0xaaaaaaaaaaaaaaaaull, 0xbbbbbbbbbbbbbbbbull};
+	const int      before = IllegalHits();
+	fn(out, src, dest);
+	Check(IllegalHits() > before, "host raises SIGILL for a REX.R/REX.B INSERTQ");
+	Check(out[0] == 0xaaaaaaaaaa2211aaull, "insertq xmm10, xmm9, 16, 8 inserts bits [15:0] at index 8");
+	Check(out[1] == 0xbbbbbbbbbbbbbbbbull, "insertq xmm10, xmm9, 16, 8 leaves the upper 64 bits untouched");
+}
+
+void TestUnknownInstructionRefused() {
+	g_refused.store(false, std::memory_order_relaxed);
+	g_refuse_skip.store(2, std::memory_order_relaxed);
+	auto      fn     = MapCode<VoidBlobFn>(kUd2);
+	const int before = IllegalHits();
+	fn();
+	Check(IllegalHits() > before, "host raises SIGILL for the unknown instruction");
+	Check(g_refused.load(std::memory_order_relaxed), "TryEmulate refuses an unknown instruction");
+	g_refuse_skip.store(0, std::memory_order_relaxed);
+}
+
+} // namespace
+
+int main() {
+	Check(Common::HostException::InstallHandler(Handler), "install host exception handler");
+	TestRefusals();
+	TestExtrqImmediate();
+	TestExtrqFullWidth();
+	TestInsertqImmediate();
+	TestRexHighRegisterExtrq();
+	TestRexHighRegisterInsertqWithIndex();
+	TestUnknownInstructionRefused();
+	std::printf("macosSse4aTests: all passed\n");
+	return 0;
+}

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
@@ -4097,6 +4130,56 @@ void CheckNewDecoderUnsupported(const uint32_t *shader, uint32_t words,
 #endif
 }
 
+void TestScalarAshrI64Decoder() {
+  using namespace ShaderRecompiler::Decoder;
+
+  const uint32_t shader[] = {
+      0x91860204u, // s_ashr_i64 s[6:7], s[4:5], s2
+      EncodeSop2(0x23, 8, 255, 129), 0xfffffffeu,
+      EncodeSop2(0x23, 10, 8, 255), 0xffffffffu,
+      EncodeSop2(0x23, 12, 255, 255), 0xffffffc0u,
+      EncodeSopp(0x01),
+  };
+  Program program;
+  ShaderRecompiler::Decoder::DecodeProgram(shader, program);
+  Check(program.instructions.size() == 5u,
+        "S_ASHR_I64 decoder lost instruction boundaries around literals");
+  const auto &reg = program.instructions[0];
+  Check(reg.family == Family::SOP2 && reg.opcode == Opcode::S_ASHR_I64 &&
+            reg.opcode_id == 0x23u && reg.pc == 0u && reg.word_count == 1u &&
+            reg.src_count == 2u && reg.dst.kind == OperandKind::Sgpr &&
+            reg.dst.reg == 6u && reg.src0.kind == OperandKind::Sgpr &&
+            reg.src0.reg == 4u && reg.src1.kind == OperandKind::Sgpr &&
+            reg.src1.reg == 2u,
+        "S_ASHR_I64 register encoding was decoded incorrectly");
+  const auto &source_literal = program.instructions[1];
+  const auto &count_literal = program.instructions[2];
+  const auto &shared_literal = program.instructions[3];
+  for (const auto *inst : {&source_literal, &count_literal, &shared_literal}) {
+    Check(inst->opcode == Opcode::S_ASHR_I64 && inst->word_count == 2u &&
+              inst->src_count == 2u,
+          "S_ASHR_I64 literal encoding was decoded incorrectly");
+  }
+  Check(source_literal.pc == 4u &&
+            source_literal.src0.kind == OperandKind::LiteralConstant &&
+            source_literal.src0.value == 0xfffffffeu &&
+            source_literal.src1.kind == OperandKind::IntegerInlineConstant &&
+            source_literal.src1.value == 1u && count_literal.pc == 12u &&
+            count_literal.src1.kind == OperandKind::LiteralConstant &&
+            count_literal.src1.value == 0xffffffffu &&
+            shared_literal.pc == 20u &&
+            shared_literal.src0.kind == OperandKind::LiteralConstant &&
+            shared_literal.src1.kind == OperandKind::LiteralConstant &&
+            shared_literal.src0.value == 0xffffffc0u &&
+            shared_literal.src1.value == 0xffffffc0u &&
+            program.instructions[4].pc == 28u &&
+            program.instructions[4].opcode == Opcode::S_ENDPGM,
+        "S_ASHR_I64 decoder mishandled source, count, or shared literals");
+  Check(Common::ContainsStr(ProgramToString(program),
+                            "S_ASHR_I64 s6, s4, s2"),
+        "S_ASHR_I64 is missing from the decoded dump");
+}
+
 void TestNewShaderDecoderArchitecture() {
   using namespace ShaderRecompiler::Decoder;
 
@@ -8729,6 +8812,54 @@ void TestNewShaderRecompilerCfgNestedEarlyExitSharedTerminal() {
   CheckSpirvBinaryValidates(result.spirv);
 }
 
+void TestNewShaderRecompilerCfgEarlyReturnSharedLoopContinuation() {
+  // PS d6fb5f22c689ceff: an outer EXEC skip and an inner SCC exit share the
+  // continuing arm, which contains another loop and reaches a different return.
+  const uint32_t shader[] = {
+      EncodeSopc(0x06, 0, 128),
+      EncodeSopp(0x04, 7), // outer -> shared loop or preceding work
+      EncodeSMovB32(2, 129),
+      EncodeSopc(0x06, 2, 128), // preceding work loop
+      EncodeSopp(0x04, 2),      // loop -> early-return condition or body
+      EncodeSop2(0x01, 2, 2, 129),
+      EncodeSopp(0x02, -4),
+      EncodeSopc(0x06, 1, 128),
+      EncodeSopp(0x04, 5),      // inner -> private return or shared loop
+      EncodeSopc(0x06, 3, 128), // shared loop header
+      EncodeSopp(0x04, 2),      // loop -> normal return or body
+      EncodeSop2(0x01, 3, 3, 129),
+      EncodeSopp(0x02, -4),  // body -> loop header
+      0xbf810000u,           // normal return
+      EncodeSMovB32(4, 130), // private return epilogue
+      0xbf810000u,
+  };
+
+  ShaderRecompiler::Decoder::Program decoded;
+  ShaderRecompiler::Decoder::DecodeProgram(std::span{shader}, decoded);
+  auto graph = ShaderRecompiler::CFG::BuildGraph(decoded);
+  const auto original_coverage =
+      CfgInstructionCoverage(graph, decoded.instructions.size());
+  Check(ShaderRecompiler::CFG::Structurize(graph),
+        graph.unsupported_reason.c_str());
+  Check(CfgInstructionCoverage(graph, decoded.instructions.size()) ==
+                original_coverage &&
+            graph.natural_loops.size() == 2u,
+        "shared-continuation gateway changed semantic blocks or the loops");
+  Check(std::ranges::none_of(graph.blocks,
+                             [](const auto &block) {
+                               return block.terminator.goto_variable !=
+                                      UINT32_MAX;
+                             }),
+        "private early return introduced unnecessary routing state");
+
+  auto options = MakeCompileOptions(ShaderType::Compute);
+  auto result = RecompileForTest(shader, options);
+  Check(!result.program.dispatcher_fallback &&
+            SpirvInstructionOpcodeCount(result.spirv, 251) == 0u,
+        "early return with a shared loop continuation selected the dispatcher");
+  CheckSpirvBinaryValidates(result.spirv);
+}
+
 void TestNewShaderRecompilerCfgSharedTerminalEarlyExit() {
   const uint32_t shader[] = {
       EncodeSopc(0x06, 0, 0), // outer early-exit condition
@@ -9354,7 +9485,7 @@ void TestNewShaderRecompilerSetpcBranch() {
 
 void TestFusedShaderHandoffPreservesRegisters() {
   using namespace ShaderRecompiler;
-  const uint32_t front[] = {
+  uint32_t front[] = {
       EncodeSMovB32(12, 255), 0x1003u, // three vertices and one primitive
       EncodeSop1(0x20, 0, 6), // merged-stage handoff through s[6:7]
       0xffffffffu,            // front shader metadata must not be decoded
@@ -9373,20 +9504,23 @@ void TestFusedShaderHandoffPreservesRegisters() {
   options.stage = ShaderType::Mesh;
   options.input_info.vertex = &input;
   options.back_code = back;
-  auto translated = TranslateProgram(front, options);
-  uint32_t allocations = 0;
-  for (const auto* block: translated.program.blocks) {
-    for (const auto& inst: *block) {
-      if (inst.GetOpcode() != IR::ValueOpcode::MeshAllocate) {
-        continue;
+  for (const auto handoff: {EncodeSop1(0x20, 0, 6), 0xbefd2106u}) {
+    front[2] = handoff; // SETPC or captured SWAPPC with NULL destination
+    auto translated = TranslateProgram(front, options);
+    uint32_t allocations = 0;
+    for (const auto* block: translated.program.blocks) {
+      for (const auto& inst: *block) {
+        if (inst.GetOpcode() != IR::ValueOpcode::MeshAllocate) {
+          continue;
+        }
+        const auto value = inst.Arg(0).Resolve();
+        Check(value.IsImmediate() && value.U32() == 0x1003u,
+              "fused back shader lost the front shader's scalar register value");
+        allocations++;
       }
-      const auto value = inst.Arg(0).Resolve();
-      Check(value.IsImmediate() && value.U32() == 0x1003u,
-            "fused back shader lost the front shader's scalar register value");
-      allocations++;
     }
+    Check(allocations == 1u, "fused shader omitted the back shader allocation");
   }
-  Check(allocations == 1u, "fused shader omitted the back shader allocation");
 }
 
 void TestMeshExportStorage() {
@@ -9606,6 +9740,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 +9934,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,
@@ -13388,6 +13543,7 @@ int main() {
   // Opcode semantics and optimized SPIR-V are exercised by
   // ShaderRecompilerComputeTests; keep the distinct decoder contract checks
   // here.
+  TestScalarAshrI64Decoder();
   TestNewShaderDecoderArchitecture();
   TestImageAddressOperands();
   TestSopkCompareImmediateExtension();
@@ -13437,6 +13593,7 @@ int main() {
   TestNewShaderRecompilerCfgSharedRegionBeforeEarlyBreakLoop();
   TestNewShaderRecompilerCfgOverlappingEarlyExitLadder();
   TestNewShaderRecompilerCfgNestedEarlyExitSharedTerminal();
+  TestNewShaderRecompilerCfgEarlyReturnSharedLoopContinuation();
   TestNewShaderRecompilerCfgSharedTerminalEarlyExit();
   TestNewShaderRecompilerCfgPrunesUnreachableSelectionEntry();
   TestNewShaderRecompilerCfgFailedStructurizationPreservesGraph();
@@ -13448,6 +13605,7 @@ int main() {
   TestNewShaderRecompilerBufferLoadsGuardedByExec();
   TestNewShaderRecompilerBufferAtomicsGuardedByBounds();
   TestCapturedBufferAtomicsX2();
+  TestDisabledSystemDebugBranch();
   TestNewShaderRecompilerPixelImageSampleLodSelection();
   TestNewShaderRecompilerBranchConditionForms();
   TestNewShaderRecompilerSetpcBranch();