loader: replace SSE4a EXTRQ with native trampolines on hosts without SSE4a (#839)

loader: add native EXTRQ patches with trap fallback

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

3rdparty/CMakeLists.txt

@@ -25,7 +25,7 @@ FetchContent_MakeAvailable(opus)
 set(BUILD_TESTING "${kyty_build_testing_saved}")
 install(FILES "${opus_SOURCE_DIR}/COPYING" DESTINATION "licenses/opus")
 
-# The Windows guest red-zone patcher needs a decoder/encoder and a small
+# The guest instruction patcher needs a decoder/encoder and a small
 # trampoline assembler. Pin the same revisions used by shadPS4's implementation.
 set(XBYAK_BUILD_TESTS OFF CACHE BOOL "" FORCE)
 FetchContent_Declare(
@@ -48,6 +48,7 @@ FetchContent_Declare(
 )
 
 FetchContent_MakeAvailable(xbyak zydis)
+target_compile_definitions(xbyak INTERFACE XBYAK_NO_EXCEPTION)
 
 # ZArchive provides random-access reads for compressed game dumps. Build only
 # the reader/writer library and its zstd dependency; the command-line tool is

CMakeLists.txt

@@ -1,7 +1,7 @@
 cmake_minimum_required(VERSION 3.18...4.3)
 
 if(CMAKE_VERSION VERSION_GREATER_EQUAL 4.0)
-	set(CMAKE_POLICY_VERSION_MINIMUM 3.5 CACHE STRING "Minimum policy version for bundled dependencies")
+	set(CMAKE_POLICY_VERSION_MINIMUM 3.10 CACHE STRING "Minimum policy version for bundled dependencies")
 endif()
 
 get_property(KYTY_GENERATOR_IS_MULTI_CONFIG GLOBAL PROPERTY GENERATOR_IS_MULTI_CONFIG)

src/common/emulatorConfig.cpp

@@ -23,6 +23,7 @@ void Load(const ConfigOptions& cfg) {
 	EXIT_IF(g_config == nullptr);
 	EXIT_IF(cfg.user_name.empty() || cfg.user_name.size() > MAX_USER_NAME_LENGTH);
 	EXIT_IF(!IsConfiguredUserIdValid(cfg.user_id));
+	EXIT_IF(cfg.controller_speaker_volume > 100 || cfg.controller_vibration_intensity > 100);
 
 	*g_config = cfg;
 }
@@ -51,6 +52,14 @@ const std::optional<ControllerColor>& GetControllerColor() {
 	return g_config->controller_color;
 }
 
+uint32_t GetControllerSpeakerVolume() {
+	return g_config->controller_speaker_volume;
+}
+
+uint32_t GetControllerVibrationIntensity() {
+	return g_config->controller_vibration_intensity;
+}
+
 PresentMode GetPresentMode() {
 	return g_config->present_mode;
 }
@@ -63,6 +72,10 @@ bool FullscreenEnabled() {
 	return g_config->fullscreen_enabled;
 }
 
+bool HideCursorEnabled() {
+	return g_config->hide_cursor_enabled;
+}
+
 bool VrEnabled() {
 	return g_config->vr_enabled;
 }

src/common/emulatorConfig.h

@@ -48,9 +48,12 @@ struct ConfigOptions {
 	int32_t                user_id                     = DEFAULT_USER_ID;
 	std::string            audio_input_device;
 	std::optional<ControllerColor> controller_color;
+	uint32_t               controller_speaker_volume      = 100;
+	uint32_t               controller_vibration_intensity = 100;
 	PresentMode            present_mode                = PresentMode::Mailbox;
 	int32_t                gpu_index                   = -1;
 	bool                   fullscreen_enabled          = false;
+	bool                   hide_cursor_enabled         = false;
 	bool                   vr_enabled                  = false;
 	bool                   amd_cpu_enabled             = false;
 	uint32_t               vblank_frequency            = 60;
@@ -86,9 +89,12 @@ const std::string& GetUserName();
 int32_t  GetUserId();
 const std::string& GetAudioInputDevice();
 const std::optional<ControllerColor>& GetControllerColor();
+uint32_t GetControllerSpeakerVolume();
+uint32_t GetControllerVibrationIntensity();
 PresentMode GetPresentMode();
 int32_t GetGpuIndex();
 bool     FullscreenEnabled();
+bool     HideCursorEnabled();
 bool     VrEnabled();
 bool     AmdCpuEnabled();
 uint32_t GetVblankFrequency();

src/graphics/guest_gpu/gpu_format.cpp

@@ -61,8 +61,8 @@ constexpr FormatInfo kFormatInfo[] = {
 	{BufferFormat::k32_32_32_32UInt, 16, 0, 16, true, true},
 	{BufferFormat::k32_32_32_32SInt, 16, 0, 16, true, false, true},
 	{BufferFormat::k32_32_32_32Float, 16, 0, 16, true, false},
-	{BufferFormat::k8Srgb, 1, 0, 0, true, false},
-	{BufferFormat::k8_8Srgb, 2, 0, 0, true, false},
+	{BufferFormat::k8Srgb, 1, 0, 1, true, false},
+	{BufferFormat::k8_8Srgb, 2, 0, 2, true, false},
 	{BufferFormat::k8_8_8_8Srgb, 4, 0, 4, true, false},
 	{BufferFormat::k9_9_9_5Float, 4, 0, 0, true, false},
 	{BufferFormat::k5_6_5UNorm, 2, 0, 2, true, false},

src/graphics/host_gpu/graphicContext.h

@@ -29,6 +29,7 @@ struct GraphicContext {
 	bool                               memory_budget_ext_enabled             = false;
 	bool                               compute_subgroup_size_control_enabled = false;
 	bool                               sample_rate_shading_enabled           = false;
+	bool                               shader_image_int64_atomics_enabled    = false;
 	// bool fp64_denorm_preserve = false; // Temporarily disabled.
 	bool                               attachment_feedback_loop_enabled      = false;
 	bool                               provoking_vertex_last_enabled         = false;

src/graphics/host_gpu/regionManager.h

@@ -114,6 +114,9 @@ public:
 	void ForEachModifiedRange(uint64_t vaddr, uint64_t size, Func&& func) {
 		const auto [start, end] = GetPageRange(vaddr, size);
 		auto&      bits         = GetBits<source>();
+		if (bits.FirstRangeFrom(start).first >= end) {
+			return;
+		}
 		RegionBits mask(bits, start, end);
 		if constexpr (clear) {
 			bits.UnsetRange(start, end);

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

@@ -791,13 +791,6 @@ TextureCache::OverlapResult TextureCache::ResolveOverlap(const ImageInfo& reques
 		if (const auto depth_id = ResolveDepthOverlap(requested, binding, cached_id)) {
 			return {depth_id};
 		}
-		if (requested.IsBlock() && !cached.info.IsBlock()) {
-			return {ExpandImage(requested, cached_id)};
-		}
-		if (requested.data.size == cached.info.data.size &&
-		    (requested.IsVolume() || cached.info.IsVolume())) {
-			return {ExpandImage(requested, cached_id)};
-		}
 		// Equal pitch does not imply equal mip placement: a changed extent can move
 		// a level into or out of the mip tail. These are separate guest layouts.
 		if (requested.tile_mode != cached.info.tile_mode ||
@@ -808,6 +801,19 @@ TextureCache::OverlapResult TextureCache::ResolveOverlap(const ImageInfo& reques
 			}
 			return {merged_id};
 		}
+		if (requested.IsBlock() && !cached.info.IsBlock()) {
+			return {ExpandImage(requested, cached_id)};
+		}
+		// Volume depth is not an array-layer count. A larger depth can retain the
+		// same block-slice layout while requiring a larger native image.
+		if ((requested.IsVolume() || cached.info.IsVolume()) &&
+		    (requested.data.size == cached.info.data.size ||
+		     (requested.type == cached.info.type && requested.resources == cached.info.resources &&
+		      requested.extent.width == cached.info.extent.width &&
+		      requested.extent.height == cached.info.extent.height &&
+		      requested.extent.depth > cached.info.extent.depth))) {
+			return {ExpandImage(requested, cached_id)};
+		}
 		// PPSA08394
 		// A view cannot change the native image type or grow its extent.
 		if (requested.data.size == cached.info.data.size &&
@@ -878,9 +884,7 @@ TextureCache::OverlapResult TextureCache::ResolveOverlap(const ImageInfo& reques
 ImageId TextureCache::ExpandImage(const ImageInfo& info, ImageId source_id) {
 	RefreshCopySource(source_id);
 	const auto expanded_id = InsertImage(info);
-	auto&      expanded    = m_slot_images[expanded_id];
 	auto&      source      = m_slot_images[source_id];
-	expanded.usage         = source.usage;
 	if (source.binding.is_bound || source.binding.is_target) {
 		source.binding.needs_rebind = true;
 	}
@@ -926,22 +930,12 @@ TextureCache::BuildTextureTransfer(const Image& image, BindingType binding,
 	const bool  upload           = direction == TransferDirection::Upload;
 	const bool  render_target    = binding == BindingType::RenderTarget;
 	const bool  video_out        = binding == BindingType::VideoOut;
-	auto        format           = info.guest_format;
 	uint32_t    layers           = info.TransferLayers();
 	bool        volume           = info.IsVolume();
-	bool        allow_depth_tile = upload;
 	const char* owner            = "TextureCache readback";
 
 	TextureTransfer transfer;
-	transfer.swap_bgra16 = info.bgra16 && (!upload || render_target || video_out);
-	if (render_target) {
-		format = ImageOps::RenderTargetTransferFormat(info.bytes_per_block);
-	}
-	if (video_out) {
-		allow_depth_tile = false;
-	} else if (render_target || binding == BindingType::Storage) {
-		allow_depth_tile = true;
-	}
+	transfer.swap_bgra16 = info.bgra16;
 	if (upload) {
 		if ((render_target || video_out) &&
 		    (info.resources.layers == 0 || info.data.size % info.resources.layers != 0 ||
@@ -963,9 +957,9 @@ TextureCache::BuildTextureTransfer(const Image& image, BindingType binding,
 		}
 	}
 
-	transfer.layout  = TextureCalcUploadLayout(format, info.extent.width, info.extent.height,
-	                                       info.resources.levels, layers, info.tile_mode,
-	                                       info.data.size, allow_depth_tile, volume, owner);
+	transfer.layout = TextureCalcUploadLayout(info.guest_format, info.extent.width,
+	                                         info.extent.height, info.resources.levels, layers,
+	                                         info.tile_mode, info.data.size, volume, owner);
 	transfer.regions = TextureBuildImageCopies(transfer.layout);
 	if (info.IsDepth()) {
 		for (auto& region: transfer.regions) {
@@ -1133,7 +1127,7 @@ void TextureCache::MaterializeColorClear(ImageId id, const ImageDesc& desc,
 		image.info.metadata = desc.info.metadata;
 		// Native color metadata must not retain a reused HTile/CMask/FMask clear flag.
 		m_surface_metas.erase(range.address);
-		if (range.size == 0 || desc.info.resources.levels != 1 || image.info.resources.levels != 1) {
+		if (range.size == 0 || desc.info.resources.levels != 1) {
 			return;
 		}
 	}
@@ -1181,10 +1175,12 @@ void TextureCache::MaterializeColorClear(ImageId id, const ImageDesc& desc,
 			           {vk::ImageAspectFlagBits::eColor, view.base_level, view.level_count,
 			            image_first + slice, 1}, clear);
 		}
-		// Native expanded keys own consumption. Existing buffer tracking publishes this CPU
-		// write to future GPU readers; FillBuffer can fault and must run outside the texture lock.
+		// Publish the conversion's expanded keys without treating them as guest writes
+		// to overlapping image data. Invalidate the buffer before updating its backing.
 		if (desc.type != BindingType::VideoOut) {
-			m_buffer_cache.FillBuffer(address, slice_size, UINT32_MAX, false);
+			std::fill(bytes.begin(), bytes.end(), uint8_t {0xff});
+			m_buffer_cache.InvalidateMemory(address, slice_size);
+			LibKernel::Memory::WriteBacking(address, bytes.data(), bytes.size());
 		}
 	}
 }
@@ -1850,7 +1846,7 @@ void TextureCache::InvalidateMemoryFromGPU(uint64_t address, uint64_t size) {
 	std::scoped_lock lock {m_lock};
 	for (const auto id: FindImagesInRegion(address, size, true)) {
 		auto& image = m_slot_images[id];
-		if (!image.Overlaps(address, size)) {
+		if (image.info.data.address != address) {
 			continue;
 		}
 		if (image.IsGpuModified()) {

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

@@ -340,7 +340,9 @@ void Image::CopyImage(Image& source) {
 	EXIT_IF(source.backing.samples != backing.samples);
 	m_scheduler.EndRendering();
 	const uint32_t levels     = std::min(source.backing.mip_levels, backing.mip_levels);
-	const uint32_t base_depth = backing.image_type == vk::ImageType::e3D
+	const uint32_t base_depth = source.backing.image_type == backing.image_type
+	                                ? std::min(source.backing.extent.depth, backing.extent.depth)
+	                            : backing.image_type == vk::ImageType::e3D
 	                                ? backing.extent.depth
 	                                : source.backing.extent.depth;
 	const auto     source_aspect =

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

@@ -358,10 +358,8 @@ vk::ImageView Image::FindView(const ImageViewInfo& view_info) {
 	}
 
 	vk::ImageViewUsageCreateInfo usage {};
-	usage.usage = image.usage;
-	if (!is_storage) {
-		usage.usage &= ~vk::ImageUsageFlagBits::eStorage;
-	}
+	usage.usage = is_storage ? vk::ImageUsageFlagBits::eStorage
+	                         : image.usage & ~vk::ImageUsageFlagBits::eStorage;
 	vk::ImageViewMinLodCreateInfoEXT min_lod {};
 	if (normalized.min_lod != 0) {
 		min_lod.minLod = static_cast<float>(normalized.base_level) +

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

@@ -90,8 +90,11 @@ IsSupportedSampledDepthResource(const ShaderRecompiler::IR::ImageResource& resou
 	if (resource.numeric_class != Prospero::TextureNumericClass::Float && resource.depth_compare) {
 		return false;
 	}
-	return resource.mip_mode == ShaderRecompiler::IR::ImageMipMode::None && resource.read &&
-	       !resource.written && !resource.atomic;
+	return ((resource.mip_mode == ShaderRecompiler::IR::ImageMipMode::None &&
+	         resource.mip_count == 1u) ||
+	        (resource.mip_mode == ShaderRecompiler::IR::ImageMipMode::Dynamic &&
+	         resource.mip_count != 0u)) &&
+	       resource.read && !resource.written && !resource.atomic;
 }
 
 inline void ValidateStorageColorView(vk::Format image_format, vk::Format view_format,
@@ -114,7 +117,7 @@ IsSupportedStorageImageResource(const ShaderRecompiler::IR::ImageResource& resou
 	        resource.dimension == ShaderRecompiler::Decoder::ImageDimension::Dim2DArray) &&
 	       ((resource.mip_mode == ShaderRecompiler::IR::ImageMipMode::None &&
 	         resource.mip_count == 1u) ||
-	        (resource.mip_mode == ShaderRecompiler::IR::ImageMipMode::DynamicStorage &&
+	        (resource.mip_mode == ShaderRecompiler::IR::ImageMipMode::Dynamic &&
 	         resource.mip_count != 0u)) &&
 	       resource.written &&
 	       (!resource.atomic ||

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

@@ -175,8 +175,7 @@ SurfaceFormatInfo TextureGetSurfaceFormatInfo(Prospero::BufferFormat format) {
 TextureUploadLayout TextureCalcUploadLayout(Prospero::BufferFormat format, uint32_t width,
                                             uint32_t height, uint32_t levels, uint32_t depth,
                                             Prospero::TileMode tile_mode, uint64_t upload_size,
-                                            bool allow_depth_tile, bool volume_texture,
-                                            const char* owner) {
+                                            bool volume_texture, const char* owner) {
 	TextureUploadLayout layout {};
 	layout.surface.description = {
 	    format,
@@ -204,8 +203,7 @@ TextureUploadLayout TextureCalcUploadLayout(Prospero::BufferFormat format, uint3
 			     static_cast<uint32_t>(format));
 		}
 	} else {
-		if ((tile_mode == Prospero::TileMode::kDepth && !allow_depth_tile) ||
-		    !TileGetTiledTextureLayout(description, layout.surface)) {
+		if (!TileGetTiledTextureLayout(description, layout.surface)) {
 			EXIT("%s: unsupported typed tiled upload: fmt=%u tile=%u "
 			     "size=%" PRIu64 " extent=%ux%u levels=%u\n",
 			     owner, static_cast<uint32_t>(format), static_cast<uint32_t>(tile_mode),

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

@@ -47,8 +47,7 @@ RenderTargetFormatInfo TextureGetRenderTargetFormat(Prospero::ChannelLayout layo
 TextureUploadLayout    TextureCalcUploadLayout(Prospero::BufferFormat format, uint32_t width,
                                                uint32_t height, uint32_t levels, uint32_t depth,
                                                Prospero::TileMode tile_mode, uint64_t upload_size,
-                                               bool allow_depth_tile, bool volume_texture,
-                                               const char* owner);
+                                               bool volume_texture, const char* owner);
 std::vector<vk::BufferImageCopy> TextureBuildImageCopies(const TextureUploadLayout& layout);
 bool TextureBuildGpuTileInfos(uint64_t tiled_size, const std::vector<vk::BufferImageCopy>& regions,
                               const TextureUploadLayout& layout, uint32_t levels,

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

@@ -312,13 +312,14 @@ void ValidateStorageTexture(const ShaderRecompiler::IR::ImageResource& resource,
 	                              resource.written && !resource.read && !resource.atomic;
 	const auto numeric_class = Prospero::SampledTextureNumericClass(format);
 	const bool raw_float_atomic = format == Prospero::BufferFormat::k32Float && uint_resource &&
-	                              resource.atomic;
+	                              resource.atomic && !resource.atomic64;
 	const bool format_ok =
 	    raw_sint_storage || raw_float_atomic ||
 	    (numeric_class != Prospero::TextureNumericClass::Unsupported &&
 	     numeric_class != Prospero::TextureNumericClass::Sint &&
 	     uint_resource == (numeric_class == Prospero::TextureNumericClass::Uint) &&
-	     (!resource.atomic || format == Prospero::BufferFormat::k32UInt));
+	     (!resource.atomic || format == (resource.atomic64 ? Prospero::BufferFormat::k32_32UInt
+	                                                       : Prospero::BufferFormat::k32UInt)));
 	if (resource_ok && descriptor_ok && encoding_ok && format_ok && size != 0) {
 		return;
 	}
@@ -354,7 +355,8 @@ static TextureCache::ImageDesc NullTextureDesc(const ShaderRecompiler::IR::Image
 			desc.info.guest_format = Prospero::BufferFormat::k32Float;
 			break;
 		case Prospero::TextureNumericClass::Uint:
-			desc.info.guest_format = Prospero::BufferFormat::k32UInt;
+			desc.info.guest_format = resource.atomic64 ? Prospero::BufferFormat::k32_32UInt
+			                                         : Prospero::BufferFormat::k32UInt;
 			break;
 		case Prospero::TextureNumericClass::Sint:
 			desc.info.guest_format = Prospero::BufferFormat::k32SInt;
@@ -365,10 +367,10 @@ static TextureCache::ImageDesc NullTextureDesc(const ShaderRecompiler::IR::Image
 	desc.info.type            = Prospero::ImageType::kColor2D;
 	desc.info.extent          = {1, 1, 1};
 	desc.info.resources       = {1, 1};
-	desc.info.bytes_per_block = 4;
+	desc.info.bytes_per_block = Prospero::NumBytesPerElement(desc.info.guest_format);
 	desc.info.samples         = 1;
 	desc.info.mip_layout[0]   = {0, 0, 1, 1};
-	desc.view_info.format     = desc.info.pixel_format;
+	desc.view_info.format     = resource.atomic64 ? vk::Format::eR64Uint : desc.info.pixel_format;
 	desc.view_info.type       = vk::ImageViewType::e2D;
 	desc.view_info.aspect     = vk::ImageAspectFlagBits::eColor;
 	desc.view_info.usage      = binding == TextureCache::BindingType::Storage
@@ -538,6 +540,9 @@ static bool ResolveTextureMipView(const TileSurfaceDescription& description, boo
 
 TextureBinding RenderExecutor::ResolveTexture(const ShaderRecompiler::IR::ImageResource&   resource,
                                               const ShaderRecompiler::IR::DescriptorValue& value) {
+	if (resource.atomic64 && !m_context.GetGraphics().shader_image_int64_atomics_enabled) {
+		EXIT("64-bit image atomics require shaderImageInt64Atomics\n");
+	}
 	auto descriptor = DecodeNativeDescriptor<ShaderTextureResource>(value);
 	const bool storage = resource.written;
 	if (storage) {
@@ -563,7 +568,7 @@ TextureBinding RenderExecutor::ResolveTexture(const ShaderRecompiler::IR::ImageR
 	const auto physical_levels = multisampled ? 1u : static_cast<uint32_t>(max_mip) + 1u;
 	// IMAGE_STORE addresses BASE_LEVEL; only IMAGE_STORE_MIP selects other view mips.
 	const bool single_storage_mip =
-	    storage && resource.mip_mode != ShaderRecompiler::IR::ImageMipMode::DynamicStorage;
+	    storage && resource.mip_mode != ShaderRecompiler::IR::ImageMipMode::Dynamic;
 	const auto view_levels = multisampled || single_storage_mip
 	                             ? 1u
 	                             : static_cast<uint32_t>(last_level - base_level) + 1u;
@@ -650,8 +655,9 @@ TextureBinding RenderExecutor::ResolveTexture(const ShaderRecompiler::IR::ImageR
 			pixel_format = depth_format->depth_attachment_format;
 		}
 	}
-	const auto storage_view_format = storage && (resource.atomic ||
-	                                            format == Prospero::BufferFormat::k32SInt)
+	const auto storage_view_format = resource.atomic64 ? vk::Format::eR64Uint
+	                                 : storage && (resource.atomic ||
+	                                               format == Prospero::BufferFormat::k32SInt)
 	                                     ? vk::Format::eR32Uint
 	                                     : SrgbStorageViewFormat(pixel_format);
 	const auto view_format         = storage && storage_view_format != vk::Format::eUndefined
@@ -785,7 +791,8 @@ void RenderExecutor::PrepareBindings(const ShaderStageRuntime& runtime,
 	}
 	prepared.samplers.reserve(program.info.samplers.size());
 	for (uint32_t i = 0; i < program.info.samplers.size(); i++) {
-		prepared.samplers.push_back(NativeSampler(m_context, program, i, snapshot.samplers[i]));
+		prepared.samplers.push_back(NativeSampler(
+		    m_context, program, i, snapshot.samplers[program.info.samplers[i].snapshot_index]));
 	}
 	prepared.shader_data.reserve(program.bindings.ShaderDataDwords());
 	for (const auto reg: program.bindings.user_data_registers) {
@@ -885,7 +892,7 @@ void RenderExecutor::RebindImages(PreparedBindings& prepared) {
 		auto& binding = images[i];
 		binding.mip_views.clear();
 		const auto& resource = program.info.images[i];
-		if (resource.mip_mode == ShaderRecompiler::IR::ImageMipMode::DynamicStorage) {
+		if (resource.mip_mode == ShaderRecompiler::IR::ImageMipMode::Dynamic) {
 			EXIT_IF(resource.mip_count == 0u ||
 			        resource.mip_count != binding.desc.view_info.level_count);
 			binding.mip_views.reserve(resource.mip_count);
@@ -893,6 +900,8 @@ void RenderExecutor::RebindImages(PreparedBindings& prepared) {
 				auto desc = binding.desc;
 				desc.view_info.base_level += mip;
 				desc.view_info.level_count = 1;
+				// The shader selects the mip after applying the guest minimum LOD.
+				desc.view_info.min_lod = 0;
 				binding.mip_views.push_back(texture_cache.FindTexture(binding.image_id, desc));
 			}
 			binding.image_view = binding.mip_views.front();

src/graphics/host_gpu/renderer/renderCompute.cpp

@@ -321,7 +321,8 @@ void RenderExecutor::DispatchDirect(uint64_t submit_id, CommandBuffer& buffer,
 			     static_cast<uint32_t>(r.TileMode()));
 		}
 		for (uint32_t i = 0; i < program.info.samplers.size(); i++) {
-			const auto r = DecodeNativeDescriptor<ShaderSamplerResource>(resources.samplers[i]);
+			const auto r = DecodeNativeDescriptor<ShaderSamplerResource>(
+			    resources.samplers[program.info.samplers[i].snapshot_index]);
 			LOGF("  CS sampler[%u]: source=%u clamp=%u/%u/%u filter=%u/%u/%u mip=%u "
 			     "lod=%u-%u bias=%d\n",
 			     i, program.info.samplers[i].source, static_cast<uint32_t>(r.ClampX()),

src/graphics/presentation/systemOverlay.cpp

@@ -29,9 +29,9 @@ namespace Libs::Graphics {
 
 namespace {
 
-namespace CoreIme   = Libs::Ime;
-namespace DialogIme = Libs::Dialog::ImeDialog;
-namespace ErrorDialog = Libs::Dialog::ErrorDialog;
+namespace CoreIme      = Libs::Ime;
+namespace DialogIme    = Libs::Dialog::ImeDialog;
+namespace SystemDialog = Libs::Dialog::SystemDialog;
 
 namespace Ime {
 
@@ -104,7 +104,7 @@ bool HostQueueExternalInput(uint64_t generation, ExternalInput input) {
 
 } // namespace Ime
 
-enum class OverlayKind : uint8_t { None, Ime, Error };
+enum class OverlayKind : uint8_t { None, Ime, Dialog };
 
 struct OverlaySession {
 	OverlayKind kind       = OverlayKind::None;
@@ -114,14 +114,14 @@ struct OverlaySession {
 };
 
 struct OverlaySnapshot {
-	OverlaySession            session;
-	Ime::HostSnapshot         ime;
-	ErrorDialog::HostSnapshot error;
+	OverlaySession             session;
+	Ime::HostSnapshot          ime;
+	SystemDialog::HostSnapshot dialog;
 };
 
 bool GetOverlaySnapshot(OverlaySnapshot* snapshot) {
-	if (ErrorDialog::GetHostSnapshot(&snapshot->error)) {
-		snapshot->session = {OverlayKind::Error, snapshot->error.generation};
+	if (SystemDialog::GetHostSnapshot(&snapshot->dialog)) {
+		snapshot->session = {OverlayKind::Dialog, snapshot->dialog.generation};
 		return true;
 	}
 	if (Ime::GetHostSnapshot(&snapshot->ime)) {
@@ -234,7 +234,7 @@ void RefreshVisibility() {
 	bool capture_keyboard   = false;
 	bool text_input         = false;
 	bool multiline          = false;
-	if (snapshot.session.kind == OverlayKind::Error) {
+	if (snapshot.session.kind == OverlayKind::Dialog) {
 		capture_controller = true;
 		capture_keyboard   = true;
 	} else if (snapshot.session.kind == OverlayKind::Ime) {
@@ -429,14 +429,14 @@ void InitializeSystemOverlayInput(SDL_Window* window) {
 	}
 	CoreIme::SetVisibilityCallback(OnCoreVisibilityChanged);
 	DialogIme::SetVisibilityCallback(OnDialogVisibilityChanged);
-	ErrorDialog::SetVisibilityCallback(RefreshVisibility);
+	SystemDialog::SetVisibilityCallback(RefreshVisibility);
 	RefreshVisibility();
 }
 
 void ShutdownSystemOverlayInput() {
 	CoreIme::SetVisibilityCallback(nullptr);
 	DialogIme::SetVisibilityCallback(nullptr);
-	ErrorDialog::SetVisibilityCallback(nullptr);
+	SystemDialog::SetVisibilityCallback(nullptr);
 	{
 		std::scoped_lock lock(g_visibility_mutex);
 		g_input_lifecycle_active = false;
@@ -466,9 +466,9 @@ void ShutdownSystemOverlayInput() {
 SystemOverlayVisualState GetSystemOverlayVisualState() noexcept {
 	const auto core   = CoreIme::GetVisualState();
 	const auto dialog = DialogIme::GetVisualState();
-	const auto error  = ErrorDialog::GetVisualState();
-	return {core.active || dialog.active || error.active,
-	        core.revision + dialog.revision + error.revision};
+	const auto system = SystemDialog::GetVisualState();
+	return {core.active || dialog.active || system.active,
+	        core.revision + dialog.revision + system.revision};
 }
 
 bool ProcessSystemOverlayInput(const SDL_Event& event) {
@@ -521,10 +521,11 @@ bool ProcessSystemOverlayInput(const SDL_Event& event) {
 	if (keyboard_event && !g_input_keyboard) {
 		return false;
 	}
-	if (keyboard_event && session.kind == OverlayKind::Error) {
+	if (keyboard_event && session.kind == OverlayKind::Dialog) {
 		if (event.type == SDL_EVENT_KEY_DOWN && !event.key.repeat &&
-		    (event.key.key == SDLK_RETURN || event.key.key == SDLK_KP_ENTER)) {
-			ErrorDialog::HostAccept(generation);
+		    (event.key.key == SDLK_RETURN || event.key.key == SDLK_KP_ENTER ||
+		     event.key.key == SDLK_ESCAPE)) {
+			SystemDialog::HostClose(generation);
 		}
 		return true;
 	}
@@ -904,7 +905,7 @@ struct SystemOverlay::Impl {
 		ImGui::End();
 	}
 
-	void DrawError(const ErrorDialog::HostSnapshot& snapshot, vk::Extent2D extent) {
+	void DrawDialog(const SystemDialog::HostSnapshot& snapshot, vk::Extent2D extent) {
 		const ImVec2 display(static_cast<float>(extent.width), static_cast<float>(extent.height));
 		const float  scale = std::max(std::min(display.x / 1280.0f, display.y / 720.0f), 0.5f);
 		ImGui::GetBackgroundDrawList()->AddRectFilled({0.0f, 0.0f}, display,
@@ -922,17 +923,23 @@ struct SystemOverlay::Impl {
 		constexpr ImGuiWindowFlags flags = ImGuiWindowFlags_NoDecoration | ImGuiWindowFlags_NoMove |
 		                                   ImGuiWindowFlags_NoSavedSettings |
 		                                   ImGuiWindowFlags_AlwaysAutoResize;
-		ImGui::Begin("##SystemError", nullptr, flags);
-		ImGui::TextUnformatted("Error");
+		const bool                 signin = snapshot.kind == SystemDialog::Kind::Signin;
+		ImGui::Begin("##SystemDialog", nullptr, flags);
+		ImGui::TextUnformatted(signin ? "Sign in to PlayStation Network" : "Error");
 		ImGui::Separator();
 		const auto error_code = static_cast<uint32_t>(snapshot.error_code);
-		ImGui::TextWrapped("%s", error_code == 0x80550006u
+		ImGui::TextWrapped("%s", signin ? "PlayStation Network sign-in is unavailable."
+		                         : error_code == 0x80550006u
 		                             ? "You are not signed in to PlayStation Network."
 		                             : "An error has occurred.");
-		ImGui::TextDisabled("Error code: 0x%08X", error_code);
+		if (!signin) {
+			ImGui::TextDisabled("Error code: 0x%08X", error_code);
+		}
 		const float button_width = std::min(140.0f * scale, ImGui::GetContentRegionAvail().x);
 		ImGui::SetCursorPosX((ImGui::GetWindowSize().x - button_width) * 0.5f);
-		const bool accepted = ImGui::Button("OK", {button_width, 44.0f * scale});
+		const bool accepted =
+		    ImGui::Button(signin ? "Cancel" : "OK", {button_width, 44.0f * scale}) ||
+		    (signin && ImGui::IsKeyPressed(ImGuiKey_GamepadFaceRight, false));
 		if (focus_pending) {
 			ImGui::SetItemDefaultFocus();
 			focus_pending = false;
@@ -941,7 +948,7 @@ struct SystemOverlay::Impl {
 		ImGui::PopFont();
 		ImGui::PopStyleVar(2);
 		if (accepted) {
-			ErrorDialog::HostAccept(snapshot.generation);
+			SystemDialog::HostClose(snapshot.generation);
 		}
 	}
 
@@ -955,7 +962,8 @@ struct SystemOverlay::Impl {
 		if (session != snapshot.session) {
 			session       = snapshot.session;
 			focus_pending = true;
-			shift         = (snapshot.ime.option & Ime::OPTION_NO_AUTO_CAPITALIZE) == 0;
+			shift         = session.kind == OverlayKind::Ime &&
+			                (snapshot.ime.option & Ime::OPTION_NO_AUTO_CAPITALIZE) == 0;
 			symbol_mode   = false;
 			panel_offset  = {};
 			right_stick   = {};
@@ -981,8 +989,8 @@ struct SystemOverlay::Impl {
 			ImGui::EndFrame();
 			return false;
 		}
-		if (snapshot.session.kind == OverlayKind::Error) {
-			DrawError(snapshot.error, frame_extent);
+		if (snapshot.session.kind == OverlayKind::Dialog) {
+			DrawDialog(snapshot.dialog, frame_extent);
 		} else {
 			DrawIme(snapshot.ime, frame_extent);
 		}

src/graphics/presentation/window/hostInput.cpp

@@ -72,6 +72,7 @@ struct Binding {
 };
 
 constexpr int              MOUSE_POLL_INTERVAL_MS = 33;
+constexpr uint64_t         CURSOR_IDLE_HIDE_MS    = 2000;
 constexpr std::string_view MOUSE_SENSITIVITY      = "MouseSensitivity=";
 
 struct MouseJoystickState {
@@ -81,7 +82,8 @@ struct MouseJoystickState {
 };
 
 MouseJoystickState g_mouse;
-SDL_Window*        g_mouse_window = nullptr;
+SDL_Window*        g_mouse_window   = nullptr;
+uint64_t           g_cursor_hide_at = 0;
 
 std::size_t ControlFromName(std::string_view name) {
 	const auto info = std::find_if(CONTROL_INFO.begin(), CONTROL_INFO.end(),
@@ -380,11 +382,31 @@ int PollMouse(uint64_t now_ms) {
 	return MOUSE_POLL_INTERVAL_MS;
 }
 
+bool IsCursorActivity(const SDL_Event& event) {
+	SDL_MouseID which;
+	switch (event.type) {
+		case SDL_EVENT_MOUSE_MOTION:
+			if (event.motion.xrel == 0.0f && event.motion.yrel == 0.0f) {
+				return false;
+			}
+			which = event.motion.which;
+			break;
+		case SDL_EVENT_MOUSE_BUTTON_DOWN:
+		case SDL_EVENT_MOUSE_BUTTON_UP: which = event.button.which; break;
+		case SDL_EVENT_MOUSE_WHEEL: which = event.wheel.which; break;
+		default: return false;
+	}
+	return which != SDL_TOUCH_MOUSEID && which != SDL_PEN_MOUSEID;
+}
+
 } // namespace
 
 void HostInputInit(SDL_Window* window) {
 	GetInputMap();
 	g_mouse_window = window;
+	if (Config::HideCursorEnabled()) {
+		g_cursor_hide_at = SDL_GetTicks() + CURSOR_IDLE_HIDE_MS;
+	}
 }
 
 void HostInputShutdown() {
@@ -393,7 +415,11 @@ void HostInputShutdown() {
 		CenterMouseStick();
 		g_mouse = {};
 	}
-	g_mouse_window = nullptr;
+	g_mouse_window   = nullptr;
+	g_cursor_hide_at = 0;
+	if (Config::HideCursorEnabled()) {
+		SDL_ShowCursor();
+	}
 }
 
 void HostInputKey(int key_code, bool down) {
@@ -429,20 +455,36 @@ void HostInputToggleMouseToJoystick() {
 }
 
 bool HostInputWaitEvent(SDL_Event* event) {
-	bool has_event;
+	int timeout = -1;
 	if (!g_mouse.enabled || SDL_GetKeyboardFocus() != g_mouse_window) {
 		g_mouse.next_poll = 0;
 		CenterMouseStick();
-		has_event = SDL_WaitEvent(event);
-		if (!has_event) {
-			EXIT("%s\n", SDL_GetError());
-		}
 	} else {
 		if (g_mouse.next_poll == 0) {
 			SDL_GetRelativeMouseState(nullptr, nullptr);
 			g_mouse.next_poll = SDL_GetTicks() + MOUSE_POLL_INTERVAL_MS;
 		}
-		has_event = SDL_WaitEventTimeout(event, PollMouse(SDL_GetTicks()));
+		timeout = PollMouse(SDL_GetTicks());
+	}
+
+	if (g_cursor_hide_at != 0) {
+		const auto now_ms = SDL_GetTicks();
+		const int  cursor_timeout =
+		    now_ms < g_cursor_hide_at ? static_cast<int>(g_cursor_hide_at - now_ms) : 0;
+		timeout = timeout < 0 ? cursor_timeout : std::min(timeout, cursor_timeout);
+	}
+	const bool has_event = SDL_WaitEventTimeout(event, timeout);
+
+	if (Config::HideCursorEnabled()) {
+		const auto now_ms = SDL_GetTicks();
+		if (has_event && !g_mouse.enabled && IsCursorActivity(*event) &&
+		    SDL_GetWindowFromEvent(event) == g_mouse_window) {
+			SDL_ShowCursor();
+			g_cursor_hide_at = now_ms + CURSOR_IDLE_HIDE_MS;
+		} else if (g_cursor_hide_at != 0 && now_ms >= g_cursor_hide_at) {
+			SDL_HideCursor();
+			g_cursor_hide_at = 0;
+		}
 	}
 
 	if (has_event && event->type == SDL_EVENT_WINDOW_FOCUS_LOST &&

src/graphics/presentation/window/vulkanWindow.cpp

@@ -460,7 +460,15 @@ static vk::Device VulkanCreateDevice(GraphicContext& graphics,
 		provoking_vertex.pNext = supported_features2.pNext;
 		supported_features2.pNext = &provoking_vertex;
 	}
+	const bool image_atomic_int64_extension =
+	    HasExtension(device_extensions, VK_EXT_SHADER_IMAGE_ATOMIC_INT64_EXTENSION_NAME);
+	vk::PhysicalDeviceShaderImageAtomicInt64FeaturesEXT image_atomic_int64 {};
+	if (image_atomic_int64_extension) {
+		image_atomic_int64.pNext = supported_features2.pNext;
+		supported_features2.pNext = &image_atomic_int64;
+	}
 	physical_device.getFeatures2(&supported_features2);
+	graphics.shader_image_int64_atomics_enabled = image_atomic_int64.shaderImageInt64Atomics;
 
 	auto features12 = WindowContext::RequiredVulkan12Features();
 	features12.shaderSharedInt64Atomics = supported_features12.shaderSharedInt64Atomics;
@@ -605,6 +613,11 @@ static vk::Device VulkanCreateDevice(GraphicContext& graphics,
 		provoking_vertex.transformFeedbackPreservesProvokingVertex = VK_FALSE;
 		create_info.pNext = &provoking_vertex;
 	}
+	if (graphics.shader_image_int64_atomics_enabled) {
+		image_atomic_int64.pNext = const_cast<void*>(create_info.pNext);
+		image_atomic_int64.sparseImageInt64Atomics = VK_FALSE;
+		create_info.pNext = &image_atomic_int64;
+	}
 	create_info.pQueueCreateInfos       = &queue_create_info;
 	create_info.queueCreateInfoCount    = 1;
 	create_info.enabledExtensionCount   = static_cast<uint32_t>(device_extensions.size());
@@ -969,6 +982,7 @@ void WindowContext::CreateVulkan() {
 			graphic_ctx.memory_budget_ext_enabled = true;
 		}
 		for (const auto* extension: {VK_EXT_ROBUSTNESS_2_EXTENSION_NAME,
+		                             VK_EXT_SHADER_IMAGE_ATOMIC_INT64_EXTENSION_NAME,
 		                             VK_EXT_PROVOKING_VERTEX_EXTENSION_NAME,
 		                             VK_EXT_MESH_SHADER_EXTENSION_NAME,
 		                             VK_KHR_WORKGROUP_MEMORY_EXPLICIT_LAYOUT_EXTENSION_NAME,

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

@@ -41,7 +41,7 @@ void ValidateNativeProgram(const IR::Program& program) {
 			Fail(program, "native shader plan has an invalid image class");
 		}
 		present[static_cast<size_t>(*kind)] = true;
-		const auto dynamic = program.info.images[i].mip_mode == IR::ImageMipMode::DynamicStorage;
+		const auto dynamic = program.info.images[i].mip_mode == IR::ImageMipMode::Dynamic;
 		const auto count   = dynamic ? program.info.images[i].mip_count : 1u;
 		if (count == 0u || (!dynamic && program.info.images[i].mip_count != 1u)) {
 			Fail(program, "native shader plan has an invalid image mip descriptor count");

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

@@ -124,13 +124,21 @@ uint32_t ImageType(EmitterState& state, const IR::ImageResource& image) {
 		sampled = 2;
 		if (image.atomic) {
 			EXIT_IF(image.numeric_class != Prospero::TextureNumericClass::Uint);
-			format = spv::ImageFormatR32ui;
+			format = image.atomic64 ? spv::ImageFormatR64ui : spv::ImageFormatR32ui;
+			if (image.atomic64) {
+				state.builder.RequireExtension("SPV_EXT_shader_image_int64");
+				state.builder.RequireCapability(spv::CapabilityInt64);
+				state.builder.RequireCapability(spv::CapabilityInt64Atomics);
+				state.builder.RequireCapability(spv::CapabilityInt64ImageEXT);
+			}
 		}
 	} else {
 		EXIT("invalid image resource class");
 	}
 	const auto& info = ImageDimensionInfoFor(image.dimension);
-	return state.builder.Type(spv::OpTypeImage, ImageScalarType(state, image.numeric_class),
+	const auto scalar_type = image.atomic64 ? TypeScalarU64(state)
+	                                        : ImageScalarType(state, image.numeric_class);
+	return state.builder.Type(spv::OpTypeImage, scalar_type,
 	                          info.spirv_dimension, image.depth_compare ? 1u : 0u, info.arrayed,
 	                          info.multisampled, sampled, format);
 }
@@ -144,20 +152,11 @@ uint32_t ImageViewSizeType(EmitterState& state, ImageDimension dimension) {
 	}
 }
 
-uint32_t LoadSampledImageDescriptor(EmitterState& state, uint32_t resource) {
-	const auto& image_resource = state.program.info.images.at(resource);
-	EXIT_IF(image_resource.resource_class != IR::ImageResourceClass::Sampled);
-	const auto kind = IR::DescriptorBindingForImage(image_resource);
-	EXIT_IF(!kind.has_value());
-	const auto array_index  = ResourceForDescriptor(state, *kind, resource);
-	const auto variable     = state.image_variables[IR::ImageBindingIndex(*kind)];
-	const auto pointer_type = state.builder.Type(
-	    spv::OpTypePointer, spv::StorageClassUniformConstant, ImageType(state, image_resource));
-	const auto pointer =
-	    DescriptorElementPointer(state, pointer_type, variable, array_index, *kind, resource,
-	                             "sampled image descriptor array was not emitted");
+uint32_t LoadImageDescriptor(EmitterState& state, uint32_t resource, uint32_t mip) {
+	const auto pointer = ImageDescriptorPointer(state, resource, mip);
 	const auto image = state.builder.AllocateId();
-	state.builder.AddFunction(spv::OpLoad, ImageType(state, image_resource), image, pointer);
+	state.builder.AddFunction(spv::OpLoad, ImageType(state, state.program.info.images.at(resource)),
+	                          image, pointer);
 	return image;
 }
 
@@ -175,10 +174,10 @@ uint32_t LoadSamplerDescriptor(EmitterState& state, uint32_t sampler) {
 	return sampler_id;
 }
 
-uint32_t MakeSampledImage(EmitterState& state, uint32_t resource, uint32_t sampler) {
+uint32_t MakeSampledImage(EmitterState& state, uint32_t resource, uint32_t sampler_id, uint32_t mip) {
 	const auto& image_resource = state.program.info.images.at(resource);
-	const auto  image          = LoadSampledImageDescriptor(state, resource);
-	const auto  sampler_id     = LoadSamplerDescriptor(state, sampler);
+	EXIT_IF(image_resource.resource_class != IR::ImageResourceClass::Sampled);
+	const auto  image          = LoadImageDescriptor(state, resource, mip);
 	const auto  sampled_image = state.builder.AllocateId();
 	const auto  sampled_type =
 	    state.builder.Type(spv::OpTypeSampledImage, ImageType(state, image_resource));
@@ -186,17 +185,17 @@ uint32_t MakeSampledImage(EmitterState& state, uint32_t resource, uint32_t sampl
 	return sampled_image;
 }
 
-uint32_t StorageImageDescriptorPointer(EmitterState& state, uint32_t resource) {
+uint32_t ImageDescriptorPointer(EmitterState& state, uint32_t resource, uint32_t mip) {
 	const auto& image = state.program.info.images.at(resource);
-	EXIT_IF(image.resource_class != IR::ImageResourceClass::Storage);
+	EXIT_IF(mip >= image.mip_count);
 	const auto kind = IR::DescriptorBindingForImage(image);
 	EXIT_IF(!kind.has_value());
-	const auto array_index  = ResourceForDescriptor(state, *kind, resource);
+	const auto array_index  = ResourceForDescriptor(state, *kind, resource) + mip;
 	const auto pointer_type = state.builder.Type(
 	    spv::OpTypePointer, spv::StorageClassUniformConstant, ImageType(state, image));
 	const auto variable = state.image_variables[IR::ImageBindingIndex(*kind)];
 	return DescriptorElementPointer(state, pointer_type, variable, array_index, *kind, resource,
-	                                "storage image descriptor array was not emitted");
+	                                "image descriptor array was not emitted");
 }
 
 void EmitStorageImageWrite(EmitterState& state, uint32_t resource, uint32_t mip_lod, uint32_t coord,
@@ -206,46 +205,15 @@ void EmitStorageImageWrite(EmitterState& state, uint32_t resource, uint32_t mip_
 	if (!image.atomic) {
 		state.builder.RequireCapability(spv::CapabilityStorageImageWriteWithoutFormat);
 	}
-	const auto kind = IR::DescriptorBindingForImage(image);
-	EXIT_IF(!kind.has_value());
-	const auto array_index = ResourceForDescriptor(state, *kind, resource);
-	const auto image_type  = ImageType(state, image);
-	const auto pointer_type =
-	    state.builder.Type(spv::OpTypePointer, spv::StorageClassUniformConstant, image_type);
-	const auto variable = state.image_variables[IR::ImageBindingIndex(*kind)];
-	const auto LoadAt   = [&](uint32_t array_index) {
-		const auto pointer =
-		    DescriptorElementPointer(state, pointer_type, variable, array_index, *kind, resource,
-		                             "storage image descriptor array was not emitted");
-		const auto descriptor = state.builder.AllocateId();
-		state.builder.AddFunction(spv::OpLoad, image_type, descriptor, pointer);
-		return descriptor;
+	const auto EmitWrite = [&](uint32_t mip) {
+		state.builder.AddFunction(spv::OpImageWrite, LoadImageDescriptor(state, resource, mip),
+		                          coord, texel);
 	};
-	if (image.mip_mode != IR::ImageMipMode::DynamicStorage) {
-		state.builder.AddFunction(spv::OpImageWrite, LoadAt(array_index), coord, texel);
+	if (image.mip_mode != IR::ImageMipMode::Dynamic) {
+		EmitWrite(0);
 		return;
 	}
-	if (image.mip_count == 0u) {
-		ExitDescriptorBindingFailure(state, *kind, resource,
-		                             "dynamic storage image has no mip descriptors");
-	}
-
-	const auto            merge_label = state.builder.AllocateId();
-	std::vector<uint32_t> labels(image.mip_count);
-	std::vector<uint32_t> words {spv::OpSwitch, mip_lod, merge_label};
-	for (uint32_t mip = 0; mip < image.mip_count; mip++) {
-		labels[mip] = state.builder.AllocateId();
-		words.push_back(mip);
-		words.push_back(labels[mip]);
-	}
-	state.builder.AddFunction(spv::OpSelectionMerge, merge_label, spv::SelectionControlMaskNone);
-	state.builder.AddFunction(words);
-	for (uint32_t mip = 0; mip < image.mip_count; mip++) {
-		EmitLabel(state, labels[mip]);
-		state.builder.AddFunction(spv::OpImageWrite, LoadAt(array_index + mip), coord, texel);
-		state.builder.AddFunction(spv::OpBranch, merge_label);
-	}
-	EmitLabel(state, merge_label);
+	EmitImageMipSwitch(state, mip_lod, image.mip_count, 0, EmitWrite);
 }
 
 spv::ExecutionModel ExecutionModelForStage(ShaderType stage) {

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

@@ -3,9 +3,7 @@
 #include "graphics/shader/recompiler/frontend/decode/ImageOps.h"
 
 #include <algorithm>
-#include <atomic>
 #include <bit>
-#include <cstdio>
 
 namespace Libs::Graphics::ShaderRecompiler::Spirv::Emitter {
 namespace {
@@ -86,6 +84,46 @@ uint32_t ZeroF32(EmitterState& state) {
 	return ConstantF32(state, 0);
 }
 
+uint32_t GatherMip(ValueEmitContext& ctx, const IR::Inst& inst, const IR::MemoryInfo& mem,
+                   const IR::Inst& address, const ImageSampleLayout& layout, uint32_t mip_count) {
+	auto& state = ctx.state;
+	const auto Extract = [&](uint32_t word, uint32_t first, uint32_t count) {
+		return EmitBitFieldUExtract(state, word, ConstantU32(state, first),
+		                           ConstantU32(state, count));
+	};
+	const auto LodFixed = [&](uint32_t word, uint32_t first) {
+		return Binary(state, spv::OpFMul, TypeF32(state),
+		              Unary(state, spv::OpConvertUToF, TypeF32(state), Extract(word, first, 12)),
+		              ConstantF32Value(state, 1.0f / 256.0f));
+	};
+	const auto sampler_control = ctx.Arg(inst, 6);
+	const auto preclamp = Binary(state, spv::OpINotEqual, TypeBool(state),
+	                            Extract(sampler_control, 28, 1), ConstantU32(state, 0));
+	const auto half = ConstantF32Value(state, 0.5f);
+	const auto zero = ZeroF32(state);
+	const auto base = Unary(state, spv::OpConvertUToF, TypeF32(state),
+	                        Extract(ctx.Arg(inst, 4), 12, 4));
+	const auto minimum = Binary(state, spv::OpFSub, TypeF32(state),
+	                            LodFixed(ctx.Arg(inst, 3), 8), base);
+	const auto last = ConstantF32Value(state, static_cast<float>(mip_count - 1u));
+	auto lod = Binary(state, spv::OpFAdd, TypeF32(state),
+	                  AddressF32(ctx, mem, address, layout.lod),
+	                  Select(state, TypeF32(state), preclamp, half, zero));
+	lod = EmitGlsl<GLSLstd450FClamp, IR::Type::F32>(
+	    state, lod, LodFixed(ctx.Arg(inst, 5), 0), LodFixed(ctx.Arg(inst, 5), 12));
+	// T# minimum LOD is in physical mip space and applies after the S# view-relative clamp.
+	lod = EmitGlsl<GLSLstd450FMax, IR::Type::F32>(state, lod, minimum);
+	lod = EmitGlsl<GLSLstd450FClamp, IR::Type::F32>(state, lod, zero, last);
+	lod = Binary(state, spv::OpFAdd, TypeF32(state), lod,
+	             Select(state, TypeF32(state), preclamp, zero, half));
+	// The clamped value is nonnegative, so conversion performs the integer mip rounding.
+	const auto mip = EmitUMin32(state, Unary(state, spv::OpConvertFToU, TypeU32(state), lod),
+	                           ConstantU32(state, mip_count - 1u));
+	const auto mip_none = Binary(state, spv::OpIEqual, TypeBool(state),
+	                            Extract(sampler_control, 26, 2), ConstantU32(state, 0));
+	return Select(state, TypeU32(state), mip_none, ConstantU32(state, 0), mip);
+}
+
 uint32_t CubeAxis(EmitterState& state, uint32_t value) {
 	return Binary(state, spv::OpFSub, TypeF32(state), value, ConstantF32(state, 0x3f800000u));
 }
@@ -252,7 +290,7 @@ uint32_t QueryDimensions(ValueEmitContext& ctx, const IR::MemoryInfo& mem,
                          const IR::Inst& address) {
 	const auto  dimension = ctx.state.program.info.images.at(mem.resource).dimension;
 	const auto& info      = ImageDimensionInfoFor(dimension);
-	const auto  image     = LoadSampledImageDescriptor(ctx.state, mem.resource);
+	const auto  image     = LoadImageDescriptor(ctx.state, mem.resource);
 	const auto  size      = ctx.state.builder.AllocateId();
 	if (info.multisampled != 0u) {
 		ctx.state.builder.AddFunction(spv::OpImageQuerySize,
@@ -433,7 +471,7 @@ uint32_t EmitOneDimensionalGatherLz(ValueEmitContext& ctx, const IR::MemoryInfo&
                                     uint32_t coord, Prospero::TextureNumericClass numeric_class) {
 	auto& state = ctx.state;
 	state.builder.RequireCapability(spv::CapabilityImageQuery);
-	const auto image = LoadSampledImageDescriptor(state, mem.resource);
+	const auto image = LoadImageDescriptor(state, mem.resource);
 	const auto width = state.builder.AllocateId();
 	state.builder.AddFunction(spv::OpImageQuerySizeLod, TypeU32(state), width, image,
 	                          ConstantU32(state, 0));
@@ -446,7 +484,8 @@ uint32_t EmitOneDimensionalGatherLz(ValueEmitContext& ctx, const IR::MemoryInfo&
 	                                 Binary(state, spv::OpFMul, TypeF32(state), coord, width_f32),
 	                                 ConstantF32(state, 0x3f000000u)));
 
-	const auto sampled     = MakeSampledImage(state, mem.resource, mem.sampler);
+	const auto sampled = MakeSampledImage(state, mem.resource,
+	                                     LoadSamplerDescriptor(state, mem.sampler));
 	const auto vector_type = ImageVectorType(state, numeric_class, 4);
 	const auto scalar_type = ImageScalarType(state, numeric_class);
 	const auto component =
@@ -558,7 +597,8 @@ void EmitImage(ValueEmitContext& ctx, const IR::Inst& inst) {
 	if (op == IR::ValueOpcode::ImageQueryLod) {
 		state.builder.RequireCapability(spv::CapabilityImageQuery);
 		const auto dimension = image.dimension;
-		const auto sampled   = MakeSampledImage(state, mem.resource, mem.sampler);
+		const auto sampled = MakeSampledImage(state, mem.resource,
+		                                     LoadSamplerDescriptor(state, mem.sampler));
 		const auto lod       = state.builder.AllocateId();
 		state.builder.AddFunction(
 		    spv::OpImageQueryLod, TypeF32Vector(state, 2), lod, sampled,
@@ -588,7 +628,7 @@ void EmitImage(ValueEmitContext& ctx, const IR::Inst& inst) {
 		    EmitValueOrDefaultIfCondition(
 		        state, condition, TypeU32Vector(state, 4), ConstantU32CompositeZero(state, 4),
 		        [&]() {
-			        const auto descriptor = LoadSampledImageDescriptor(state, mem.resource);
+			        const auto descriptor = LoadImageDescriptor(state, mem.resource);
 			        const auto color      = state.builder.AllocateId();
 			        const auto coord      = CoordU32(ctx, mem, *address, dimension);
 			        if (dimension_info.multisampled != 0u) {
@@ -612,7 +652,7 @@ void EmitImage(ValueEmitContext& ctx, const IR::Inst& inst) {
 		const auto dimension  = image.dimension;
 		EmitIfCondition(state, ctx.Arg(inst, 3), [&]() {
 			const auto mip_lod =
-			    state.program.info.images[mem.resource].mip_mode == IR::ImageMipMode::DynamicStorage
+			    state.program.info.images[mem.resource].mip_mode == IR::ImageMipMode::Dynamic
 			        ? LodU32(ctx, mem, *address, dimension)
 			        : 0u;
 			const auto coord = CoordU32(ctx, mem, *address, dimension);
@@ -635,13 +675,6 @@ void EmitImage(ValueEmitContext& ctx, const IR::Inst& inst) {
 		if (op == IR::ValueOpcode::ImageGatherRaw) {
 			const auto coord = CoordF32(ctx, mem, *address, layout.coord,
 			                            dimension_info.coordinate_components, image.cube);
-			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) ||
@@ -658,44 +691,48 @@ void EmitImage(ValueEmitContext& ctx, const IR::Inst& inst) {
 				ctx.Fail(inst, "has an unsupported 1D-array gather");
 				return;
 			}
-			const auto            sampled = MakeSampledImage(state, mem.resource, mem.sampler);
-			const auto            sample  = state.builder.AllocateId();
-			std::vector<uint32_t> words;
+			const auto result_numeric_class = dref ? Prospero::TextureNumericClass::Float
+			                                       : numeric_class;
+			const auto result_type = ImageVectorType(state, result_numeric_class, 4);
+			uint32_t component_or_dref;
 			if (dref) {
-				auto dref_value = ZeroF32(state);
-				if (layout.dref != NoImageComponent) {
-					dref_value = AddressF32(ctx, mem, *address, layout.dref);
-				}
-				words = {spv::OpImageDrefGather,
-				         TypeF32Vector(state, 4),
-				         sample,
-				         sampled,
-				         coord,
-				         dref_value};
+				component_or_dref = layout.dref != NoImageComponent
+				                        ? AddressF32(ctx, mem, *address, layout.dref)
+				                        : ZeroF32(state);
 			} else {
-				uint32_t component = 0;
-				if (ImageConversionFormat(state, mem).format == Prospero::BufferFormat::kInvalid) {
-					component = ImageGatherComponent(mem.dmask);
-				}
-				words = {spv::OpImageGather,
-				         ImageVectorType(state, numeric_class, 4),
-				         sample,
-				         sampled,
-				         coord,
-				         ConstantU32(state, component)};
+				component_or_dref = ConstantU32(
+				    state, ImageConversionFormat(state, mem).format == Prospero::BufferFormat::kInvalid
+				               ? ImageGatherComponent(mem.dmask) : 0u);
 			}
+			uint32_t operand_mask = 0;
+			uint32_t offset = 0;
 			if (HasFlag(mem, Decoder::ImageSampleFlagGatherHorizontal)) {
-				words.push_back(spv::ImageOperandsConstOffsetsMask);
-				words.push_back(HorizontalOffsets(state, dimension));
+				operand_mask = spv::ImageOperandsConstOffsetsMask;
+				offset = HorizontalOffsets(state, dimension);
 			} else if (layout.offset != NoImageComponent) {
-				words.push_back(spv::ImageOperandsOffsetMask);
-				words.push_back(PackedOffset(ctx, mem, *address, layout, dimension));
-			}
-			state.builder.AddFunction(words);
-			auto result_numeric_class = numeric_class;
-			if (dref) {
-				result_numeric_class = Prospero::TextureNumericClass::Float;
+				operand_mask = spv::ImageOperandsOffsetMask;
+				offset = PackedOffset(ctx, mem, *address, layout, dimension);
 			}
+			const auto sampler_id = LoadSamplerDescriptor(state, mem.sampler);
+			const auto EmitGather = [&](uint32_t mip) {
+				const auto sampled = MakeSampledImage(state, mem.resource, sampler_id, mip);
+				const auto sample = state.builder.AllocateId();
+				std::vector<uint32_t> words {
+				    static_cast<uint32_t>(dref ? spv::OpImageDrefGather : spv::OpImageGather), result_type,
+				    sample, sampled, coord, component_or_dref};
+				if (operand_mask != 0u) {
+					words.push_back(operand_mask);
+					words.push_back(offset);
+				}
+				state.builder.AddFunction(words);
+				return sample;
+			};
+			const auto sample = image.mip_mode == IR::ImageMipMode::Dynamic && image.mip_count > 1u
+			                        ? EmitImageMipSwitch(
+			                              state, GatherMip(ctx, inst, mem, *address, layout,
+			                                               image.mip_count),
+			                              image.mip_count, result_type, EmitGather)
+			                        : EmitGather(0);
 			ctx.Define(inst, ResultVector(ctx, UnpackImageGather(ctx, mem, sample),
 			                              result_numeric_class, false, mem, true));
 			return;
@@ -738,13 +775,14 @@ void EmitImage(ValueEmitContext& ctx, const IR::Inst& inst) {
 			operand_mask |= spv::ImageOperandsBiasMask;
 			operands.push_back(AddressF32(ctx, mem, *address, layout.bias));
 		}
+		const auto sampler_id = LoadSamplerDescriptor(state, mem.sampler);
 		const auto EmitSample = [&](uint32_t resource) {
 			const auto& candidate = state.program.info.images[resource];
 			const auto coord =
 			    CoordF32(ctx, mem, *address, layout.coord,
 			             ImageDimensionInfoFor(candidate.dimension).coordinate_components,
 			             candidate.cube);
-			const auto            sampled = MakeSampledImage(state, resource, mem.sampler);
+			const auto            sampled = MakeSampledImage(state, resource, sampler_id);
 			const auto            sample  = state.builder.AllocateId();
 			std::vector<uint32_t> sample_operands {result_type, sample, sampled, coord};
 			if (dref) {
@@ -865,12 +903,16 @@ void EmitImage(ValueEmitContext& ctx, const IR::Inst& inst) {
 	}
 	if (image_info.access == IR::ImageAccess::Atomic) {
 		const auto dimension = image.dimension;
-		ctx.Define(inst, EmitValueOrZeroIfCondition(state, ctx.Arg(inst, inst.NumArgs() - 1), [&]() {
+		const auto result_type = TypeId(state, inst.GetType());
+		const auto zero = image.atomic64 ? ConstantU64(state, 0) : ConstantU32(state, 0);
+		ctx.Define(inst, EmitValueOrDefaultIfCondition(
+		                     state, ctx.Arg(inst, inst.NumArgs() - 1), result_type, zero, [&]() {
 			           const auto pointer      = state.builder.AllocateId();
 			           const auto pointer_type = state.builder.Type(
-			               spv::OpTypePointer, spv::StorageClassImage, TypeU32(state));
+			               spv::OpTypePointer, spv::StorageClassImage,
+			               image.atomic64 ? TypeScalarU64(state) : TypeU32(state));
 			           state.builder.AddFunction(spv::OpImageTexelPointer, pointer_type, pointer,
-			                                     StorageImageDescriptorPointer(state, mem.resource),
+			                                     ImageDescriptorPointer(state, mem.resource),
 			                                     CoordU32(ctx, mem, *address, dimension),
 			                                     ConstantU32(state, 0));
 			           if (op == IR::ValueOpcode::ImageAtomicFMin32 ||

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

@@ -317,6 +317,7 @@ inline constexpr auto EmitImageAtomicSMin32    = EmitImage;
 inline constexpr auto EmitImageAtomicUMin32    = EmitImage;
 inline constexpr auto EmitImageAtomicSMax32    = EmitImage;
 inline constexpr auto EmitImageAtomicUMax32    = EmitImage;
+inline constexpr auto EmitImageAtomicUMax64    = EmitImage;
 inline constexpr auto EmitImageAtomicAnd32     = EmitImage;
 inline constexpr auto EmitImageAtomicOr32      = EmitImage;
 inline constexpr auto EmitImageAtomicXor32     = EmitImage;

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

@@ -19,6 +19,7 @@
 #include <map>
 #include <spirv/unified1/GLSL.std.450.h>
 #include <string>
+#include <type_traits>
 #include <unordered_map>
 #include <utility>
 #include <vector>
@@ -314,13 +315,14 @@ uint32_t ImageType(EmitterState& state, const IR::ImageResource& image);
 
 uint32_t ImageViewSizeType(EmitterState& state, ImageDimension dimension);
 
-uint32_t LoadSampledImageDescriptor(EmitterState& state, uint32_t resource);
+uint32_t ImageDescriptorPointer(EmitterState& state, uint32_t resource, uint32_t mip = 0);
+
+uint32_t LoadImageDescriptor(EmitterState& state, uint32_t resource, uint32_t mip = 0);
 
 uint32_t LoadSamplerDescriptor(EmitterState& state, uint32_t sampler);
 
-uint32_t MakeSampledImage(EmitterState& state, uint32_t resource, uint32_t sampler);
-
-uint32_t StorageImageDescriptorPointer(EmitterState& state, uint32_t resource);
+uint32_t MakeSampledImage(EmitterState& state, uint32_t resource, uint32_t sampler_id,
+                          uint32_t mip = 0);
 
 void EmitStorageImageWrite(EmitterState& state, uint32_t resource, uint32_t mip_lod, uint32_t coord,
                            uint32_t texel);
@@ -514,6 +516,45 @@ uint32_t ConstantDeviceAddress(EmitterState& state, uint64_t value);
 uint32_t DeviceAddressFromWords(EmitterState& state, uint32_t low, uint32_t high);
 
 // These templates accept local lambdas from several emitter translation units.
+template <typename Fn>
+auto EmitImageMipSwitch(EmitterState& state, uint32_t mip_lod, uint32_t mip_count,
+                        uint32_t result_type, Fn&& emit) {
+	constexpr bool has_result = !std::is_void_v<std::invoke_result_t<Fn, uint32_t>>;
+	EXIT_IF(mip_count == 0u);
+	const auto merge_label = state.builder.AllocateId();
+	std::vector<uint32_t> labels(mip_count);
+	std::vector<uint32_t> words {spv::OpSwitch, mip_lod, merge_label};
+	for (uint32_t mip = 0; mip < mip_count; mip++) {
+		labels[mip] = state.builder.AllocateId();
+		words.push_back(mip);
+		words.push_back(labels[mip]);
+	}
+	if constexpr (has_result) {
+		words[2] = labels.front();
+	}
+	state.builder.AddFunction(spv::OpSelectionMerge, merge_label, spv::SelectionControlMaskNone);
+	state.builder.AddFunction(words);
+	std::vector<uint32_t> phi;
+	if constexpr (has_result) {
+		phi = {spv::OpPhi, result_type, state.builder.AllocateId()};
+	}
+	for (uint32_t mip = 0; mip < mip_count; mip++) {
+		EmitLabel(state, labels[mip]);
+		if constexpr (has_result) {
+			phi.push_back(emit(mip));
+			phi.push_back(state.current_label);
+		} else {
+			emit(mip);
+		}
+		state.builder.AddFunction(spv::OpBranch, merge_label);
+	}
+	EmitLabel(state, merge_label);
+	if constexpr (has_result) {
+		state.builder.AddFunction(phi);
+		return phi[2];
+	}
+}
+
 template <typename Fn>
 void EmitIfCondition(EmitterState& state, uint32_t condition, Fn&& fn) {
 	const auto then_label  = state.builder.AllocateId();

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

@@ -1044,17 +1044,11 @@ uint32_t EmitBufferAtomic64(ValueEmitContext& ctx, const IR::Inst& inst) {
 		    return EmitValueOrDefaultIfCondition(
 		        state, EmitMemoryElementInBounds(state, resource, index), TypeU64(state),
 		        ConstantU64(state, 0), [&]() {
-			        const auto value = Unary(state, spv::OpBitcast, TypeScalarU64(state),
-			                                 ctx.Arg(inst, inst.NumArgs() - 2));
-			        const auto old   = state.builder.AllocateId();
-			        state.builder.AddFunction(
-			            SpirvAtomicOpcode(inst.GetOpcode()), TypeScalarU64(state), old,
-			            EmitStorageBufferElementPointer(state, resource, index,
-			                                            TypeStorageBufferU64ElementPointer(state)),
-			            ConstantU32(state, spv::ScopeDevice),
-			            ConstantU32(state, spv::MemorySemanticsMaskNone), value);
+			        const auto pointer = EmitStorageBufferElementPointer(
+			            state, resource, index, TypeStorageBufferU64ElementPointer(state));
+			        const auto old = EmitAtomicOperation(ctx, inst, pointer, spv::ScopeDevice);
 			        EmitAtomicMemoryBarrier(state, mem.kind);
-			        return Unary(state, spv::OpBitcast, TypeU64(state), old);
+			        return old;
 		        });
 	    });
 }

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

@@ -330,6 +330,7 @@ spv::Op SpirvAtomicOpcode(IR::ValueOpcode opcode) {
 		case IR::ValueOpcode::BufferAtomicSMax32:
 		case IR::ValueOpcode::SharedAtomicSMax32: return spv::OpAtomicSMax;
 		case IR::ValueOpcode::ImageAtomicUMax32:
+		case IR::ValueOpcode::ImageAtomicUMax64:
 		case IR::ValueOpcode::BufferAtomicUMax32:
 		case IR::ValueOpcode::SharedAtomicUMax32: return spv::OpAtomicUMax;
 		case IR::ValueOpcode::ImageAtomicAnd32:
@@ -352,6 +353,8 @@ uint32_t EmitAtomicOperation(ValueEmitContext& ctx, const IR::Inst& inst, uint32
                              uint32_t scope) {
 	const auto opcode = SpirvAtomicOpcode(inst.GetOpcode());
 	const auto old    = ctx.state.builder.AllocateId();
+	const bool wide   = inst.GetType() == IR::Type::U64;
+	const auto type   = wide ? TypeScalarU64(ctx.state) : TypeU32(ctx.state);
 	if (opcode == spv::OpAtomicCompareExchange) {
 		const auto desired    = ctx.Arg(inst, inst.NumArgs() - 3);
 		const auto comparator = ctx.Arg(inst, inst.NumArgs() - 2);
@@ -360,12 +363,13 @@ uint32_t EmitAtomicOperation(ValueEmitContext& ctx, const IR::Inst& inst, uint32
 		    ConstantU32(ctx.state, scope), ConstantU32(ctx.state, spv::MemorySemanticsMaskNone),
 		    ConstantU32(ctx.state, spv::MemorySemanticsMaskNone), desired, comparator);
 	} else {
-		const auto value = ctx.Arg(inst, inst.NumArgs() - 2);
-		ctx.state.builder.AddFunction(opcode, TypeU32(ctx.state), old, pointer,
+		auto value = ctx.Arg(inst, inst.NumArgs() - 2);
+		if (wide) value = Unary(ctx.state, spv::OpBitcast, type, value);
+		ctx.state.builder.AddFunction(opcode, type, old, pointer,
 		                              ConstantU32(ctx.state, scope),
 		                              ConstantU32(ctx.state, spv::MemorySemanticsMaskNone), value);
 	}
-	return old;
+	return wide ? Unary(ctx.state, spv::OpBitcast, TypeU64(ctx.state), old) : old;
 }
 
 void EmitAtomicMemoryBarrier(EmitterState& state, IR::ResourceKind kind) {

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

@@ -360,9 +360,19 @@ void DecodeMimg(uint32_t pc, std::span<const uint32_t> code, uint32_t word_index
 		SetUnsupported(inst, Family::MIMG, opcode,
 		               "MIMG image gather requires exactly one dmask bit");
 	}
-	if (inst.opcode == Opcode::IMAGE_ATOMIC_CMPSWAP && inst.dmask != 0x3u) {
-		SetUnsupported(inst, Family::MIMG, opcode,
-		               "MIMG image compare-and-swap requires 32-bit DMASK 0x3");
+	if (atomic != nullptr) {
+		const bool compare_swap = inst.opcode == Opcode::IMAGE_ATOMIC_CMPSWAP;
+		const auto mask32       = compare_swap ? 0x3u : 0x1u;
+		const auto mask64       = compare_swap ? 0xfu : 0x3u;
+		if (inst.dmask != mask32 && inst.dmask != mask64) {
+			SetUnsupported(inst, Family::MIMG, opcode, "MIMG image atomic has invalid DMASK");
+		} else if (inst.dmask == mask64) {
+			inst.data_bits = 64u;
+			if (inst.opcode != Opcode::IMAGE_ATOMIC_UMAX) {
+				SetUnsupported(inst, Family::MIMG, opcode,
+				               "MIMG 64-bit image atomic opcode is not implemented");
+			}
+		}
 	}
 	const bool supports_d16 = sample != nullptr || gather != nullptr || opcode == 0x00u ||
 	                          opcode == 0x01u || opcode == 0x08u || opcode == 0x09u;

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

@@ -568,7 +568,9 @@ void Translator::IMAGE_ATOMIC(const Decoder::Instruction& inst, IR::ValueOpcode
 		                  ir.GetExec()},
 		                 flags);
 	} else {
-		result = ir.Emit(opcode, {resource, address, ReadU32(data), ir.GetExec()}, flags);
+		const IR::Value value =
+		    memory.data_bits == 64u ? IR::Value(ReadU64(data)) : IR::Value(ReadU32(data));
+		result = ir.Emit(opcode, {resource, address, value, ir.GetExec()}, flags);
 	}
 	if (inst.glc) {
 		WriteOperand(inst.dst, result);
@@ -706,8 +708,16 @@ void Translator::IMAGE_GATHER(const Decoder::Instruction& inst) {
 	const auto resource = GetImageResource(memory);
 	const auto sampler  = GetSamplerResource(memory);
 	const auto address  = MakeImageAddress(inst, MemorySourceAt(inst, 0));
-	const auto result   = ir.Emit(IR::ValueOpcode::ImageGatherRaw, {resource, sampler, address},
-	                              AddMemoryInfo(memory, inst.pc));
+	const bool has_lod = (memory.image_sample_flags & Decoder::ImageSampleFlagLod) != 0u;
+	// LOD selection consumes these words on the GPU after descriptor handles are stripped.
+	const auto result = ir.Emit(
+	    IR::ValueOpcode::ImageGatherRaw,
+	    {resource, sampler, address,
+	     has_lod ? resource.Instruction()->Arg(1) : IR::Value(0u),
+	     has_lod ? resource.Instruction()->Arg(3) : IR::Value(0u),
+	     has_lod ? sampler.Instruction()->Arg(1) : IR::Value(0u),
+	     has_lod ? sampler.Instruction()->Arg(2) : IR::Value(0u)},
+	    AddMemoryInfo(memory, inst.pc));
 	for (uint32_t index = 0; index < memory.data_dwords; index++) {
 		WriteOperand(OffsetOperand(inst.dst, index),
 		             ir.Emit(IR::ValueOpcode::CompositeExtractU32x4, {result, IR::Value(index)}));
@@ -1052,7 +1062,8 @@ void Translator::EmitMemory(const Decoder::Instruction& inst) {
 		case Decoder::Opcode::IMAGE_ATOMIC_SMAX:
 			return IMAGE_ATOMIC(inst, IR::ValueOpcode::ImageAtomicSMax32);
 		case Decoder::Opcode::IMAGE_ATOMIC_UMAX:
-			return IMAGE_ATOMIC(inst, IR::ValueOpcode::ImageAtomicUMax32);
+			return IMAGE_ATOMIC(inst, inst.data_bits == 64u ? IR::ValueOpcode::ImageAtomicUMax64
+			                                                : IR::ValueOpcode::ImageAtomicUMax32);
 		case Decoder::Opcode::IMAGE_ATOMIC_AND:
 			return IMAGE_ATOMIC(inst, IR::ValueOpcode::ImageAtomicAnd32);
 		case Decoder::Opcode::IMAGE_ATOMIC_OR:

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

@@ -114,7 +114,7 @@ struct BufferResource {
 	bool operator==(const BufferResource& other) const = default;
 };
 
-enum class ImageMipMode { None, DynamicStorage };
+enum class ImageMipMode { None, Dynamic };
 
 constexpr uint32_t ShaderImageIdentitySwizzle = 0x00000facu;
 
@@ -133,6 +133,7 @@ struct ImageResource {
 	bool                          read              = false;
 	bool                          written           = false;
 	bool                          atomic            = false;
+	bool                          atomic64          = false;
 	bool                          depth_compare     = false;
 	bool                          cube              = false;
 	bool                          r128              = false;
@@ -147,9 +148,12 @@ struct ImageResource {
 struct SamplerResource {
 	uint32_t source                = 0;
 	uint32_t first_use_pc          = 0;
+	// Native filtering/border variants share the original sampler's runtime descriptor.
+	uint32_t snapshot_index        = 0;
 	bool     force_point_filtering = false;
 	bool     depth_compare         = false;
 	bool     integer_border        = false;
+	bool     gather_lod            = false;
 
 	bool operator==(const SamplerResource& other) const = default;
 };
@@ -278,7 +282,7 @@ struct StageOutput {
 inline constexpr uint32_t FirstImageBinding           = 1u;
 inline constexpr uint32_t FirstComparisonImageBinding = 22u;
 inline constexpr uint32_t FirstStorageImageBinding    = 29u;
-inline constexpr uint32_t ImageBindingCount           = 43u;
+inline constexpr uint32_t ImageBindingCount           = 48u;
 
 enum class DescriptorBindingKind : uint32_t {
 	Buffers  = 0u,
@@ -291,8 +295,8 @@ enum class DescriptorBindingKind : uint32_t {
 	Count,
 };
 
-static_assert(static_cast<uint32_t>(DescriptorBindingKind::Samplers) == 44u);
-static_assert(static_cast<uint32_t>(DescriptorBindingKind::Count) == 50u);
+static_assert(static_cast<uint32_t>(DescriptorBindingKind::Samplers) == 49u);
+static_assert(static_cast<uint32_t>(DescriptorBindingKind::Count) == 55u);
 
 struct PushData {
 	static constexpr uint32_t DwordCount = 32;
@@ -369,7 +373,7 @@ DescriptorBindingForImage(const ImageResource& image) {
 			if (image.numeric_class != Prospero::TextureNumericClass::Uint) {
 				return std::nullopt;
 			}
-			base = AtomicUintBinding;
+			base = AtomicUintBinding + (image.atomic64 ? 5u : 0u);
 		} else {
 			switch (image.numeric_class) {
 				case Prospero::TextureNumericClass::Float: base = StorageFloatBinding; break;

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

@@ -227,6 +227,7 @@ ImageOpcodeInfo ImageOpcodeInfoOf(ValueOpcode opcode) {
 		case ValueOpcode::ImageAtomicUMin32:
 		case ValueOpcode::ImageAtomicSMax32:
 		case ValueOpcode::ImageAtomicUMax32:
+		case ValueOpcode::ImageAtomicUMax64:
 		case ValueOpcode::ImageAtomicAnd32:
 		case ValueOpcode::ImageAtomicOr32:
 		case ValueOpcode::ImageAtomicXor32:

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

@@ -297,7 +297,7 @@ VALUE_OPCODE(ImageQueryLod, U32x4, ImageResource, SamplerResource, ImageAddress)
 VALUE_OPCODE(ImageRead, U32x4, ImageResource, ImageAddress, U1)
 VALUE_OPCODE(ImageWrite, Void, ImageResource, ImageAddress, U32x4, U1)
 VALUE_OPCODE(ImageSampleRaw, U32x4, ImageResource, SamplerResource, ImageAddress)
-VALUE_OPCODE(ImageGatherRaw, U32x4, ImageResource, SamplerResource, ImageAddress)
+VALUE_OPCODE(ImageGatherRaw, U32x4, ImageResource, SamplerResource, ImageAddress, U32, U32, U32, U32)
 VALUE_OPCODE(ImageAtomicCompareSwap32, U32, ImageResource, ImageAddress, U32, U32, U1)
 VALUE_OPCODE(ImageAtomicSwap32, U32, ImageResource, ImageAddress, U32, U1)
 VALUE_OPCODE(ImageAtomicIAdd32, U32, ImageResource, ImageAddress, U32, U1)
@@ -305,6 +305,7 @@ VALUE_OPCODE(ImageAtomicSMin32, U32, ImageResource, ImageAddress, U32, U1)
 VALUE_OPCODE(ImageAtomicUMin32, U32, ImageResource, ImageAddress, U32, U1)
 VALUE_OPCODE(ImageAtomicSMax32, U32, ImageResource, ImageAddress, U32, U1)
 VALUE_OPCODE(ImageAtomicUMax32, U32, ImageResource, ImageAddress, U32, U1)
+VALUE_OPCODE(ImageAtomicUMax64, U64, ImageResource, ImageAddress, U64, U1)
 VALUE_OPCODE(ImageAtomicAnd32, U32, ImageResource, ImageAddress, U32, U1)
 VALUE_OPCODE(ImageAtomicOr32, U32, ImageResource, ImageAddress, U32, U1)
 VALUE_OPCODE(ImageAtomicXor32, U32, ImageResource, ImageAddress, U32, U1)

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

@@ -126,7 +126,7 @@ void AllocateBindings(Program& program, uint32_t push_data_start_dword) {
 			EXIT("shader binding layout failed: image %u has an unmapped binding class", i);
 		}
 		auto&      resources = image_groups[group];
-		const auto dynamic   = program.info.images[i].mip_mode == ImageMipMode::DynamicStorage;
+		const auto dynamic   = program.info.images[i].mip_mode == ImageMipMode::Dynamic;
 		const auto count     = dynamic ? program.info.images[i].mip_count : 1u;
 		if (count == 0u || (!dynamic && program.info.images[i].mip_count != 1u)) {
 			EXIT("shader binding layout failed: image %u has invalid specialized mip count %u", i,

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

@@ -109,8 +109,7 @@ Prospero::BufferFormat ImageConversionFormat(Prospero::BufferFormat format) {
 
 enum class SamplerClass : uint8_t { Float, Integer, PointInteger };
 
-template <typename Image>
-SamplerClass ClassifySampler(const Image& image) {
+SamplerClass ClassifySampler(const ImageResource& image) {
 	if (image.numeric_class == Prospero::TextureNumericClass::Sint ||
 	    image.conversion_format != Prospero::BufferFormat::kInvalid) {
 		return SamplerClass::PointInteger;
@@ -124,8 +123,8 @@ bool DescriptorIsCube(const DescriptorValue& descriptor) {
 	       Prospero::ImageType::kCube;
 }
 
-uint32_t StorageMipCount(const ImageResource& image, const DescriptorValue& descriptor) {
-	if (image.mip_mode != ImageMipMode::DynamicStorage || NullImageDescriptor(descriptor)) {
+uint32_t ImageMipCount(const ImageResource& image, const DescriptorValue& descriptor) {
+	if (image.mip_mode != ImageMipMode::Dynamic || NullImageDescriptor(descriptor)) {
 		return 1;
 	}
 	const auto base = (descriptor.dwords[3] >> 12u) & 0xfu;
@@ -418,9 +417,6 @@ private:
 	uint32_t                                    count = 0;
 };
 
-template <typename Images>
-bool BuildSamplerPlan(const ShaderInfo& base, const Images& images, SamplerPlan& plan);
-
 static bool BuildResourceSpecialization(const ResourcePlan& program, ResourceSnapshot& snapshot,
                                         ResourceSpecialization& specialization) {
 	for (uint32_t i = 0; i < specialization.images.size(); i++) {
@@ -435,10 +431,10 @@ static bool BuildResourceSpecialization(const ResourcePlan& program, ResourceSna
 		    (base.atomic && base.resource_class != ImageResourceClass::Storage)) {
 			return SpecializationFail(fmt::format("image resource {} has an invalid class", i));
 		}
-		image.mip_count = StorageMipCount(base, descriptor);
+		image.mip_count = ImageMipCount(base, descriptor);
 		if (image.mip_count == 0u) {
 			return SpecializationFail(
-			    fmt::format("storage image descriptor {} has an invalid mip range", i));
+			    fmt::format("image descriptor {} has an invalid mip range", i));
 		}
 		if (NullImageDescriptor(descriptor)) {
 			image.numeric_class = base.atomic ? Prospero::TextureNumericClass::Uint
@@ -460,8 +456,10 @@ static bool BuildResourceSpecialization(const ResourcePlan& program, ResourceSna
 		image.cube      = DescriptorIsCube(descriptor);
 		const auto format =
 		    static_cast<Prospero::BufferFormat>((descriptor.dwords[1] >> 20u) & 0x1ffu);
-		if (base.atomic && format != Prospero::BufferFormat::k32UInt &&
-		    format != Prospero::BufferFormat::k32Float) {
+		if (base.atomic &&
+		    (base.atomic64 ? format != Prospero::BufferFormat::k32_32UInt
+		                   : format != Prospero::BufferFormat::k32UInt &&
+		                         format != Prospero::BufferFormat::k32Float)) {
 			return SpecializationFail(
 			    fmt::format("atomic image descriptor {} uses unsupported format {}", i,
 			                static_cast<uint32_t>(format)));
@@ -561,30 +559,21 @@ static bool BuildResourceSpecialization(const ResourcePlan& program, ResourceSna
 			}
 		}
 	}
-	SamplerPlan sampler_plan;
-	if (!BuildSamplerPlan(program.info, specialization.images, sampler_plan)) {
-		return SpecializationFail("specialized sampler layout exceeds its resource limit");
-	}
-	for (uint32_t index = static_cast<uint32_t>(program.info.samplers.size());
-	     index < sampler_plan.sampler_count; index++) {
-		snapshot.samplers.push_back(snapshot.samplers[sampler_plan.bindings[index].source]);
-	}
 	ImageRemap(specialization).Apply(snapshot.images);
 	return true;
 }
 
-template <typename Images>
-bool BuildSamplerPlan(const ShaderInfo& base, const Images& images, SamplerPlan& plan) {
+bool BuildSamplerPlan(const ShaderInfo& base, SamplerPlan& plan) {
 	if (base.samplers.size() > plan.mapping.size()) {
 		return false;
 	}
 	std::array<uint8_t, ShaderInfo::MaxSamplers> usage {};
 	plan.sampler_count = static_cast<uint32_t>(base.samplers.size());
 	for (const auto& pair: base.sampled_pairs) {
-		if (pair.image >= images.size() || pair.sampler >= base.samplers.size()) {
+		if (pair.image >= base.images.size() || pair.sampler >= base.samplers.size()) {
 			return false;
 		}
-		usage[pair.sampler] |= 1u << static_cast<uint32_t>(ClassifySampler(images[pair.image]));
+		usage[pair.sampler] |= 1u << static_cast<uint32_t>(ClassifySampler(base.images[pair.image]));
 	}
 	for (uint32_t index = 0; index < base.samplers.size(); index++) {
 		auto& mapping = plan.mapping[index];
@@ -1075,6 +1064,16 @@ bool MaterializeResources(const ResourcePlan& program, const SrtRuntime& runtime
 		if (!evaluate(program.info.samplers[i].source, snapshot.samplers[i])) {
 			return false;
 		}
+		if (program.info.samplers[i].gather_lod) {
+			const auto control = snapshot.samplers[i].dwords[2];
+			const auto filter = (control >> 26u) & 3u;
+			// MipNone always selects the base level. Explicit point gathers currently require
+			// encoded-zero primary and secondary bias; linear primary-mip selection is unsupported.
+			if (filter > 1u || (filter == 1u && (control & 0xfffffu) != 0u)) {
+				return SpecializationFail(
+				    "explicit-LOD gather requires mip filtering None or Point with zero LOD biases");
+			}
+		}
 	}
 	snapshot.user_data.assign(runtime.user_data.begin(), runtime.user_data.end());
 	return BuildResourceSpecialization(program, snapshot, specialization);
@@ -1122,7 +1121,7 @@ void ApplyResourceSpecialization(Program& program, const ResourceSpecialization&
 	}
 
 	SamplerPlan sampler_plan;
-	EXIT_IF(!BuildSamplerPlan(program.info, images, sampler_plan));
+	EXIT_IF(!BuildSamplerPlan(program.info, sampler_plan));
 	auto& samplers      = program.info.samplers;
 	auto& sampled_pairs = program.info.sampled_pairs;
 	const auto original_sampler_count = samplers.size();
@@ -1132,6 +1131,7 @@ void ApplyResourceSpecialization(Program& program, const ResourceSpecialization&
 		if (index >= original_sampler_count) {
 			samplers.push_back(samplers[binding.source]);
 		}
+		samplers[index].snapshot_index = binding.source;
 		samplers[index].force_point_filtering = binding.type == SamplerClass::PointInteger;
 		samplers[index].integer_border        = binding.type != SamplerClass::Float;
 	}

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

@@ -1725,15 +1725,20 @@ private:
 
 	uint32_t AddImage(uint32_t source, const MemoryInfo& memory, ValueOpcode op, uint32_t pc) {
 		const auto resource_class = ImageOpcodeInfoOf(op).resource_class;
-		const auto mip   = resource_class == ImageResourceClass::Storage && memory.image_has_mip
-		                       ? ImageMipMode::DynamicStorage
-		                       : ImageMipMode::None;
+		const bool atomic64 = ImageOpcodeInfoOf(op).access == ImageAccess::Atomic &&
+		                      memory.data_bits == 64u;
+		const bool dynamic_mip =
+		    (resource_class == ImageResourceClass::Storage && memory.image_has_mip) ||
+		    (op == ValueOpcode::ImageGatherRaw &&
+		     (memory.image_sample_flags & Decoder::ImageSampleFlagLod) != 0u);
+		const auto mip = dynamic_mip ? ImageMipMode::Dynamic : ImageMipMode::None;
 		const bool depth = (memory.image_sample_flags & Decoder::ImageSampleFlagCompare) != 0;
 		for (uint32_t i = 0; i < m_info.images.size(); i++) {
 			auto& image = m_info.images[i];
 			if (image.source == source && image.resource_class == resource_class &&
 			    image.dimension == memory.image_dimension && image.mip_mode == mip &&
-			    image.depth_compare == depth && image.r128 == memory.image_r128) {
+			    image.depth_compare == depth && image.r128 == memory.image_r128 &&
+			    image.atomic64 == atomic64) {
 				Merge(image, op, pc);
 				return i;
 			}
@@ -1749,6 +1754,7 @@ private:
 		image.mip_mode       = mip;
 		image.depth_compare  = depth;
 		image.r128           = memory.image_r128;
+		image.atomic64       = atomic64;
 		Merge(image, op, pc);
 		m_info.images.push_back(image);
 		return static_cast<uint32_t>(m_info.images.size() - 1);
@@ -1935,6 +1941,9 @@ private:
 			if (sampler == UINT32_MAX) {
 				Fail(flags.pc, "sampler resource limit exceeded");
 			}
+			m_info.samplers[sampler].gather_lod |=
+			    op == ValueOpcode::ImageGatherRaw &&
+			    (memory.image_sample_flags & Decoder::ImageSampleFlagLod) != 0u;
 			AddHandlePatch(sampler_handle, sampler, flags.pc);
 			AddSampledPair(resource, sampler, flags.pc);
 		}

src/kernel/eventQueue.cpp

@@ -138,7 +138,13 @@ int KernelEqueuePrivate::GetTriggeredEvents(KernelEvent* ev, int num) {
 
 void KernelEqueuePrivate::TriggerExpiredTimers(uint64_t now_ns) {
 	for (auto& event: m_events) {
-		if (!event.triggered && event.deadline_ns != 0 && event.deadline_ns <= now_ns) {
+		if (event.deadline_ns != 0 && event.deadline_ns <= now_ns) {
+			if (event.event.filter == KERNEL_EVFILT_TIMER) {
+				const auto count = event.interval_ns == 0 ? !event.triggered :
+				    1 + (now_ns - event.deadline_ns) / event.interval_ns;
+				event.event.data += static_cast<intptr_t>(count);
+				event.deadline_ns += count * event.interval_ns;
+			}
 			event.triggered = true;
 		}
 	}
@@ -209,7 +215,9 @@ int KernelEqueuePrivate::AddEvent(const KernelEqueueEvent& event) {
 		                       return e.event.ident == ident && e.event.filter == filter;
 	                       });
 	if (it != m_events.end()) {
+		TriggerExpiredTimers(MonotonicTimeNs());
 		it->deadline_ns = event.deadline_ns;
+		it->interval_ns = event.interval_ns;
 		it->event.udata = event.event.udata;
 		for (auto& pending: it->pending_events) {
 			pending.udata = event.event.udata;
@@ -511,6 +519,22 @@ int KYTY_SYSV_ABI KernelDeleteUserEvent(KernelEqueue eq, int id) {
 	return KernelDeleteEvent(eq, static_cast<uintptr_t>(id), KERNEL_EVFILT_USER);
 }
 
+static int AddTimerEvent(KernelEqueue eq, int id, uint64_t delay_ns, bool periodic, void* udata) {
+	const auto now_ns = MonotonicTimeNs();
+	KernelEqueueEvent event {};
+	event.deadline_ns  = delay_ns <= UINT64_MAX - now_ns ? now_ns + delay_ns : UINT64_MAX;
+	event.interval_ns  = periodic ? delay_ns : 0;
+	event.event.ident  = static_cast<uintptr_t>(id);
+	event.event.filter = periodic ? KERNEL_EVFILT_TIMER : KERNEL_EVFILT_HRTIMER;
+	event.event.flags  = EV_ADD | (periodic ? EV_CLEAR : EV_ONESHOT);
+	event.event.udata  = udata;
+	return KernelAddEvent(eq, event);
+}
+
+int KYTY_SYSV_ABI KernelAddTimerEvent(KernelEqueue eq, int id, KernelUseconds usec, void* udata) {
+	return AddTimerEvent(eq, id, static_cast<uint64_t>(usec) * 1000, true, udata);
+}
+
 int KYTY_SYSV_ABI KernelAddHRTimerEvent(KernelEqueue eq, int id, const KernelTimespec* ts,
                                         void* udata) {
 	if (ts == nullptr) {
@@ -524,17 +548,7 @@ int KYTY_SYSV_ABI KernelAddHRTimerEvent(KernelEqueue eq, int id, const KernelTim
 
 	const auto delay_ns =
 	    static_cast<uint64_t>(ts->tv_sec) * 1000000000ull + static_cast<uint64_t>(ts->tv_nsec);
-	const auto now_ns = MonotonicTimeNs();
-
-	KernelEqueueEvent event {};
-	event.deadline_ns  = delay_ns <= UINT64_MAX - now_ns ? now_ns + delay_ns : UINT64_MAX;
-	event.event.ident  = static_cast<uintptr_t>(id);
-	event.event.filter = KERNEL_EVFILT_HRTIMER;
-	event.event.flags  = EV_ADD | EV_ONESHOT;
-	event.event.fflags = 0;
-	event.event.data   = 0;
-	event.event.udata  = udata;
-	return KernelAddEvent(eq, event);
+	return AddTimerEvent(eq, id, delay_ns, false, udata);
 }
 
 int KYTY_SYSV_ABI KernelDeleteHRTimerEvent(KernelEqueue eq, int id) {

src/kernel/eventQueue.h

@@ -51,6 +51,7 @@ struct KernelFilter {
 struct KernelEqueueEvent {
 	bool                    triggered   = false;
 	uint64_t                deadline_ns = 0;
+	uint64_t                interval_ns = 0;
 	KernelEvent             event;
 	KernelFilter            filter;
 	std::deque<KernelEvent> pending_events;
@@ -72,6 +73,8 @@ int KYTY_SYSV_ABI KernelAddUserEventEdge(KernelEqueue eq, int id);
 int KYTY_SYSV_ABI KernelTriggerUserEvent(KernelEqueue eq, int id, void* udata);
 int KYTY_SYSV_ABI KernelTriggerUserEventForAll(int id, void* udata);
 int KYTY_SYSV_ABI KernelDeleteUserEvent(KernelEqueue eq, int id);
+int KYTY_SYSV_ABI KernelAddTimerEvent(KernelEqueue eq, int id, KernelUseconds usec,
+                                      void* udata);
 int KYTY_SYSV_ABI KernelAddHRTimerEvent(KernelEqueue eq, int id, const KernelTimespec* ts,
                                         void* udata);
 int KYTY_SYSV_ABI KernelDeleteHRTimerEvent(KernelEqueue eq, int id);

src/kernel/pthread.cpp

@@ -4208,6 +4208,10 @@ int KYTY_SYSV_ABI pthread_attr_setschedpolicy(LibKernel::PthreadAttr* attr, int
 	return POSIX_PTHREAD_CALL(LibKernel::PthreadAttrSetschedpolicy(attr, policy));
 }
 
+int KYTY_SYSV_ABI pthread_attr_setstack(LibKernel::PthreadAttr* attr, void* addr, size_t size) {
+	return POSIX_PTHREAD_CALL(LibKernel::PthreadAttrSetstack(attr, addr, size));
+}
+
 int KYTY_SYSV_ABI pthread_attr_setstacksize(LibKernel::PthreadAttr* attr, size_t stack_size) {
 	// PRINT_NAME();
 

src/kernel/pthread.h

@@ -249,6 +249,7 @@ int KYTY_SYSV_ABI pthread_attr_setinheritsched(LibKernel::PthreadAttr* attr, int
 int KYTY_SYSV_ABI pthread_attr_setschedparam(LibKernel::PthreadAttr*            attr,
                                              const LibKernel::KernelSchedParam* param);
 int KYTY_SYSV_ABI pthread_attr_setschedpolicy(LibKernel::PthreadAttr* attr, int policy);
+int KYTY_SYSV_ABI pthread_attr_setstack(LibKernel::PthreadAttr* attr, void* addr, size_t size);
 int KYTY_SYSV_ABI pthread_attr_setstacksize(LibKernel::PthreadAttr* attr, size_t stack_size);
 int KYTY_SYSV_ABI pthread_cond_broadcast(LibKernel::PthreadCond* cond);
 int KYTY_SYSV_ABI pthread_cond_signal(LibKernel::PthreadCond* cond);

src/launcher/forms/configuration_edit_dialog.ui

@@ -23,7 +23,7 @@
    <property name="verticalSpacing">
     <number>8</number>
    </property>
-   <item row="0" column="0">
+   <item row="0" column="0" colspan="2">
     <layout class="QGridLayout" name="sectionsLayout" columnstretch="1,1">
      <property name="horizontalSpacing">
       <number>8</number>
@@ -72,27 +72,6 @@
           <property name="toolTip"><string>Language</string></property>
          </widget>
         </item>
-        <item row="3" column="0">
-         <widget class="QLabel" name="label_controller_color">
-          <property name="text"><string>DualSense lightbar:</string></property>
-         </widget>
-        </item>
-        <item row="3" column="1">
-         <layout class="QHBoxLayout" name="controllerColorLayout">
-          <item>
-           <widget class="QPushButton" name="button_controller_color">
-            <property name="toolTip"><string>Choose and preview a fixed LED color. Bluetooth DualSense may take several seconds to update.</string></property>
-            <property name="text"><string>Custom</string></property>
-           </widget>
-          </item>
-          <item>
-           <widget class="QPushButton" name="button_controller_color_reset">
-            <property name="toolTip"><string>Let the game control the lightbar</string></property>
-            <property name="text"><string>Default</string></property>
-           </widget>
-          </item>
-         </layout>
-        </item>
        </layout>
       </widget>
      </item>
@@ -194,6 +173,12 @@
           <property name="text"><string>Enable tessellation support (experimental)</string></property>
          </widget>
         </item>
+        <item row="8" column="1">
+         <widget class="QCheckBox" name="checkBox_hide_cursor">
+          <property name="toolTip"><string>Hide the host cursor after 2 seconds without mouse activity</string></property>
+          <property name="text"><string>Auto-hide cursor</string></property>
+         </widget>
+        </item>
        </layout>
       </widget>
      </item>
@@ -316,7 +301,82 @@
      </item>
     </layout>
    </item>
-   <item row="2" column="0">
+   <item row="1" column="1">
+    <widget class="QGroupBox" name="controller_group">
+     <property name="title"><string>Controller</string></property>
+     <layout class="QFormLayout" name="controllerLayout">
+      <property name="fieldGrowthPolicy"><enum>QFormLayout::AllNonFixedFieldsGrow</enum></property>
+      <item row="0" column="0">
+       <widget class="QLabel" name="label_controller_color">
+        <property name="text"><string>DualSense lightbar:</string></property>
+       </widget>
+      </item>
+      <item row="0" column="1">
+       <layout class="QHBoxLayout" name="controllerColorLayout">
+        <item>
+         <widget class="QPushButton" name="button_controller_color">
+          <property name="toolTip"><string>Choose and preview a fixed LED color. Bluetooth DualSense may take several seconds to update.</string></property>
+          <property name="text"><string>Custom</string></property>
+         </widget>
+        </item>
+        <item>
+         <widget class="QPushButton" name="button_controller_color_reset">
+          <property name="toolTip"><string>Let the game control the lightbar</string></property>
+          <property name="text"><string>Default</string></property>
+         </widget>
+        </item>
+       </layout>
+      </item>
+      <item row="1" column="0">
+       <widget class="QLabel" name="label_controller_vibration">
+        <property name="text"><string>Vibration intensity:</string></property>
+       </widget>
+      </item>
+      <item row="1" column="1">
+       <layout class="QHBoxLayout" name="controllerVibrationLayout">
+        <item>
+         <widget class="QSlider" name="slider_controller_vibration">
+          <property name="toolTip"><string>Scale DualSense vibration over USB and Bluetooth</string></property>
+          <property name="maximum"><number>100</number></property>
+          <property name="value"><number>100</number></property>
+          <property name="orientation"><enum>Qt::Horizontal</enum></property>
+         </widget>
+        </item>
+        <item>
+         <widget class="QLabel" name="label_controller_vibration_value">
+          <property name="minimumSize"><size><width>40</width><height>0</height></size></property>
+          <property name="text"><string>100%</string></property>
+         </widget>
+        </item>
+       </layout>
+      </item>
+      <item row="2" column="0">
+       <widget class="QLabel" name="label_controller_volume">
+        <property name="text"><string>Speaker volume:</string></property>
+       </widget>
+      </item>
+      <item row="2" column="1">
+       <layout class="QHBoxLayout" name="controllerVolumeLayout">
+        <item>
+         <widget class="QSlider" name="slider_controller_volume">
+          <property name="toolTip"><string>Scale the DualSense controller speaker over USB and Bluetooth</string></property>
+          <property name="maximum"><number>100</number></property>
+          <property name="value"><number>100</number></property>
+          <property name="orientation"><enum>Qt::Horizontal</enum></property>
+         </widget>
+        </item>
+        <item>
+         <widget class="QLabel" name="label_controller_volume_value">
+          <property name="minimumSize"><size><width>40</width><height>0</height></size></property>
+          <property name="text"><string>100%</string></property>
+         </widget>
+        </item>
+       </layout>
+      </item>
+     </layout>
+    </widget>
+   </item>
+   <item row="2" column="0" colspan="2">
     <layout class="QHBoxLayout" name="buttonLayout">
      <item>
      <widget class="QPushButton" name="clear_button">

src/launcher/include/configuration.h

@@ -46,6 +46,30 @@ QString EnumToText(T value) {
 	                                                                     : key);
 }
 
+struct ControllerSettings {
+	QString color;
+	int     speaker_volume      = 100;
+	int     vibration_intensity = 100;
+
+	void WriteSettings(QSettings* s) const {
+		s->setValue("controller_color", color);
+		s->setValue("controller_speaker_volume", speaker_volume);
+		s->setValue("controller_vibration_intensity", vibration_intensity);
+	}
+
+	void ReadSettings(QSettings* s) {
+		const QColor saved_color(s->value("controller_color").toString());
+		color = saved_color.isValid() ? saved_color.name(QColor::HexRgb) : QString {};
+		const auto read_percent = [s](const char* key) {
+			bool      ok    = false;
+			const int value = s->value(key, 100).toInt(&ok);
+			return ok ? qBound(0, value, 100) : 100;
+		};
+		speaker_volume      = read_percent("controller_speaker_volume");
+		vibration_intensity = read_percent("controller_vibration_intensity");
+	}
+};
+
 class Configuration: public QObject {
 	Q_OBJECT
 
@@ -85,14 +109,17 @@ public:
 	GameStatus game_status     = GameStatus::Unknown;
 	QString    game_comment;
 
+	// Controller preferences always come from the global configuration.
+	ControllerSettings controller;
+
 	Resolution             screen_resolution           = Resolution::R1280X720;
 	QString                user_name                   = "Kyty";
 	int                    user_id                     = Config::DEFAULT_USER_ID;
 	QString                audio_input_device;
-	QString                controller_color;
 	PresentMode            present_mode                = PresentMode::Mailbox;
 	int                    gpu_index                   = -1;
 	bool                   fullscreen_enabled          = false;
+	bool                   hide_cursor_enabled         = false;
 	bool                   readback_linear_images      = false;
 	bool                   tessellation_enabled        = false;
 	int                    vblank_frequency            = 60;
@@ -121,10 +148,10 @@ public:
 		user_name                   = other.user_name;
 		user_id                     = other.user_id;
 		audio_input_device          = other.audio_input_device;
-		controller_color            = other.controller_color;
 		present_mode                = other.present_mode;
 		gpu_index                   = other.gpu_index;
 		fullscreen_enabled          = other.fullscreen_enabled;
+		hide_cursor_enabled         = other.hide_cursor_enabled;
 		readback_linear_images      = other.readback_linear_images;
 		tessellation_enabled        = other.tessellation_enabled;
 		vblank_frequency            = other.vblank_frequency;
@@ -168,10 +195,10 @@ public:
 		KYTY_CFG_SET(user_name);
 		KYTY_CFG_SET(user_id);
 		KYTY_CFG_SET(audio_input_device);
-		KYTY_CFG_SET(controller_color);
 		KYTY_CFG_SET(present_mode);
 		KYTY_CFG_SET(gpu_index);
 		KYTY_CFG_SET(fullscreen_enabled);
+		KYTY_CFG_SET(hide_cursor_enabled);
 		KYTY_CFG_SET(readback_linear_images);
 		KYTY_CFG_SET(tessellation_enabled);
 		KYTY_CFG_SET(vblank_frequency);
@@ -208,14 +235,13 @@ public:
 		                         ? saved_user_id
 		                         : Config::DEFAULT_USER_ID;
 		audio_input_device = s->value("audio_input_device", audio_input_device).toString();
-		const QColor color(s->value("controller_color", controller_color).toString());
-		controller_color = color.isValid() ? color.name(QColor::HexRgb) : QString {};
 		KYTY_CFG_GET(present_mode);
 		gpu_index = s->value("gpu_index", -1).toInt();
 		if (EnumToText(present_mode).isEmpty()) {
 			present_mode = PresentMode::Mailbox;
 		}
 		KYTY_CFG_GET(fullscreen_enabled);
+		KYTY_CFG_GET(hide_cursor_enabled);
 		KYTY_CFG_GET(readback_linear_images);
 		KYTY_CFG_GET(tessellation_enabled);
 		vblank_frequency = s->value("vblank_frequency", vblank_frequency).toInt();

src/launcher/include/configurationEditDialog.h

@@ -28,7 +28,7 @@ public:
 	static void WriteSettings(QSettings& s);
 	static void ReadSettings(QSettings& s);
 
-	void                      SetGameDirectories(const QStringList& dirs);
+	void                      SetGlobalSettings(const QStringList& dirs);
 	[[nodiscard]] QStringList GetGameDirectories() const;
 
 signals:
@@ -40,7 +40,7 @@ private:
 	QGroupBox*                   m_game_dirs_group        = nullptr;
 	QListWidget*                 m_game_dirs_list         = nullptr;
 	QToolButton*                 m_remove_game_dir_button = nullptr;
-	bool                         m_show_game_dirs         = false;
+	bool                         m_global_settings        = false;
 
 protected:
 	void Init(const Configuration& info);

src/launcher/include/configurationListWidget.h

@@ -34,7 +34,7 @@ public:
 
 	[[nodiscard]] const QString& GetSettingsFile() const { return m_settings_file; }
 	[[nodiscard]] const QString& GetGlobalControllerColor() const {
-		return m_global_info.controller_color;
+		return m_global_info.controller.color;
 	}
 	[[nodiscard]] std::unique_ptr<Configuration>
 	CreateConfiguration(const ConfigurationItem& item) const;

src/launcher/src/configurationEditDialog.cpp

@@ -23,6 +23,7 @@
 #include <QPushButton>
 #include <QResizeEvent>
 #include <QSettings>
+#include <QSlider>
 #include <QSpinBox>
 #include <QStyle>
 #include <QToolButton>
@@ -39,7 +40,7 @@
 
 constexpr char SETTINGS_CFG_DIALOG[]               = "ConfigurationEditDialog";
 constexpr char SETTINGS_CFG_LAST_GEOMETRY[]        = "geometry";
-constexpr int  GLOBAL_SETTINGS_GAME_DIRS_MIN_WIDTH = 560;
+constexpr int  GLOBAL_SETTINGS_GAME_DIRS_MIN_WIDTH = 400;
 
 static void UpdateControllerColorButton(QPushButton* button, const QString& hex) {
 	const QColor color(hex);
@@ -123,6 +124,7 @@ ConfigurationEditDialog::ConfigurationEditDialog(Configuration& info, QWidget* p
       m_info(info) {
 	m_ui->setupUi(this);
 	InitGameDirectories();
+	m_ui->controller_group->setVisible(false);
 
 	connect(m_ui->ok_button, &QPushButton::clicked, this, &ConfigurationEditDialog::save);
 	connect(m_ui->cancel_button, &QPushButton::clicked, this, &QDialog::reject);
@@ -142,6 +144,12 @@ ConfigurationEditDialog::ConfigurationEditDialog(Configuration& info, QWidget* p
 		UpdateControllerColorButton(m_ui->button_controller_color, {});
 		emit PreviewControllerColor({});
 	});
+	connect(m_ui->slider_controller_vibration, &QSlider::valueChanged, this, [this](int value) {
+		m_ui->label_controller_vibration_value->setText(tr("%1%").arg(value));
+	});
+	connect(m_ui->slider_controller_volume, &QSlider::valueChanged, this, [this](int value) {
+		m_ui->label_controller_volume_value->setText(tr("%1%").arg(value));
+	});
 	connect(m_ui->comboBox_shader_log_direction, &QComboBox::currentTextChanged, this,
 	        [this](const QString& text) {
 		        auto log = TextToEnum<Configuration::LogDirection>(text);
@@ -201,7 +209,9 @@ void ConfigurationEditDialog::Init(const Configuration& info) {
 	m_ui->lineEdit_user_name->setMaxLength(static_cast<int>(Config::MAX_USER_NAME_LENGTH));
 	m_ui->lineEdit_user_name->setText(info.user_name);
 	m_ui->spinBox_user_id->setValue(info.user_id);
-	UpdateControllerColorButton(m_ui->button_controller_color, info.controller_color);
+	UpdateControllerColorButton(m_ui->button_controller_color, info.controller.color);
+	m_ui->slider_controller_vibration->setValue(info.controller.vibration_intensity);
+	m_ui->slider_controller_volume->setValue(info.controller.speaker_volume);
 	auto* microphone = m_ui->comboBox_audio_input_device;
 	microphone->clear();
 	microphone->addItem(tr("None"), QString {});
@@ -263,6 +273,7 @@ void ConfigurationEditDialog::Init(const Configuration& info) {
 	    info.gpu_index >= 0 && info.gpu_index < m_ui->comboBox_gpu->count() - 1 ? info.gpu_index + 1
 	                                                                            : 0);
 	m_ui->checkBox_fullscreen->setChecked(info.fullscreen_enabled);
+	m_ui->checkBox_hide_cursor->setChecked(info.hide_cursor_enabled);
 	m_ui->checkBox_readback->setChecked(info.readback_linear_images);
 	m_ui->checkBox_tessellation->setChecked(info.tessellation_enabled);
 	m_ui->spinBox_vblank_frequency->setValue(info.vblank_frequency);
@@ -342,8 +353,8 @@ void ConfigurationEditDialog::InitGameDirectories() {
 	update_game_directory_buttons();
 }
 
-void ConfigurationEditDialog::SetGameDirectories(const QStringList& dirs) {
-	m_show_game_dirs = true;
+void ConfigurationEditDialog::SetGlobalSettings(const QStringList& dirs) {
+	m_global_settings = true;
 	m_game_dirs_list->clear();
 	m_game_dirs_group->setMinimumWidth(GLOBAL_SETTINGS_GAME_DIRS_MIN_WIDTH);
 
@@ -352,6 +363,7 @@ void ConfigurationEditDialog::SetGameDirectories(const QStringList& dirs) {
 	}
 
 	m_game_dirs_group->setVisible(true);
+	m_ui->controller_group->setVisible(true);
 	update_game_directory_buttons();
 	layout()->activate();
 	resize(size().expandedTo(minimumSizeHint()));
@@ -359,7 +371,7 @@ void ConfigurationEditDialog::SetGameDirectories(const QStringList& dirs) {
 
 QStringList ConfigurationEditDialog::GetGameDirectories() const {
 	QStringList dirs;
-	if (!m_show_game_dirs) {
+	if (!m_global_settings) {
 		return dirs;
 	}
 
@@ -399,17 +411,22 @@ void ConfigurationEditDialog::resizeEvent(QResizeEvent* event) {
 	g_last_geometry = saveGeometry();
 }
 
-static void UpdateInfo(Configuration& info, Ui::ConfigurationEditDialog& ui) {
+static void UpdateInfo(Configuration& info, Ui::ConfigurationEditDialog& ui, bool global_settings) {
 	info.user_name = ui.lineEdit_user_name->text().trimmed();
 	info.user_id   = ui.spinBox_user_id->value();
 	info.audio_input_device = ui.comboBox_audio_input_device->currentData().toString();
-	info.controller_color   = ui.button_controller_color->property("controllerColor").toString();
+	if (global_settings) {
+		info.controller.color = ui.button_controller_color->property("controllerColor").toString();
+		info.controller.vibration_intensity = ui.slider_controller_vibration->value();
+		info.controller.speaker_volume      = ui.slider_controller_volume->value();
+	}
 	info.screen_resolution =
 	    TextToEnum<Configuration::Resolution>(ui.comboBox_screen_resolution->currentText());
 	info.present_mode =
 	    TextToEnum<Configuration::PresentMode>(ui.comboBox_present_mode->currentText());
 	info.gpu_index                 = ui.comboBox_gpu->currentIndex() - 1;
 	info.fullscreen_enabled        = ui.checkBox_fullscreen->isChecked();
+	info.hide_cursor_enabled       = ui.checkBox_hide_cursor->isChecked();
 	info.readback_linear_images    = ui.checkBox_readback->isChecked();
 	info.tessellation_enabled      = ui.checkBox_tessellation->isChecked();
 	info.vblank_frequency          = ui.spinBox_vblank_frequency->value();
@@ -453,7 +470,7 @@ void ConfigurationEditDialog::save() {
 		return;
 	}
 
-	UpdateInfo(m_info, *m_ui);
+	UpdateInfo(m_info, *m_ui, m_global_settings);
 
 	emit accept();
 }
@@ -461,9 +478,9 @@ void ConfigurationEditDialog::save() {
 void ConfigurationEditDialog::clear() {
 	Configuration default_info;
 	Init(default_info);
-	emit PreviewControllerColor({});
 
-	if (m_show_game_dirs) {
+	if (m_global_settings) {
+		emit PreviewControllerColor({});
 		m_game_dirs_list->clear();
 		update_game_directory_buttons();
 	}

src/launcher/src/configurationListWidget.cpp

@@ -314,6 +314,7 @@ void ConfigurationListWidget::WriteSettings() {
 	s->remove(CONF_GLOBAL);
 	s->beginGroup(CONF_GLOBAL);
 	m_global_info.WriteSettings(s.get());
+	m_global_info.controller.WriteSettings(s.get());
 	s->endGroup();
 
 	s->remove(CONF_SECTION_NAME);
@@ -357,6 +358,7 @@ void ConfigurationListWidget::ReadSettings() {
 	if (!s->childKeys().isEmpty()) {
 		m_global_info.ReadSettings(s.get());
 	}
+	m_global_info.controller.ReadSettings(s.get());
 	s->endGroup();
 
 	qDeleteAll(m_custom_infos);
@@ -409,6 +411,7 @@ ConfigurationListWidget::CreateConfiguration(const ConfigurationItem& item) cons
 	const auto* custom = m_custom_infos.value(item.GetInfo().game_path);
 	info->CopyGameInfoFrom(item.GetInfo());
 	info->CopyEmulatorSettingsFrom(custom != nullptr ? *custom : m_global_info);
+	info->controller = m_global_info.controller;
 	if (custom != nullptr && !custom->elf.isEmpty()) {
 		info->elf = custom->elf;
 	}
@@ -787,16 +790,18 @@ void ConfigurationListWidget::delete_configuartion() {
 void ConfigurationListWidget::edit_global_settings() {
 	Configuration info;
 	info.CopyEmulatorSettingsFrom(m_global_info);
+	info.controller = m_global_info.controller;
 	info.name = tr("Global settings");
 
 	ConfigurationEditDialog dlg(info, this);
 	dlg.setWindowTitle(tr("Global settings"));
-	dlg.SetGameDirectories(m_game_dirs);
+	dlg.SetGlobalSettings(m_game_dirs);
 	connect(&dlg, &ConfigurationEditDialog::PreviewControllerColor, this,
 	        &ConfigurationListWidget::PreviewControllerColor);
 
 	if (dlg.exec() == QDialog::Accepted) {
 		m_global_info.CopyEmulatorSettingsFrom(info);
+		m_global_info.controller     = info.controller;
 		const auto game_dirs         = NormalizeGameDirectories(dlg.GetGameDirectories());
 		const bool game_dirs_changed = game_dirs != m_game_dirs;
 		m_game_dirs                  = game_dirs;

src/launcher/src/mainDialog.cpp

@@ -228,9 +228,11 @@ static QStringList CreateEmulatorArgs(const Configuration& info) {
 	if (!info.audio_input_device.isEmpty()) {
 		args << "--mic" << info.audio_input_device;
 	}
-	if (!info.controller_color.isEmpty()) {
-		args << "--controller-color" << info.controller_color;
+	if (!info.controller.color.isEmpty()) {
+		args << "--controller-color" << info.controller.color;
 	}
+	args << "--controller-volume" << QString::number(info.controller.speaker_volume);
+	args << "--controller-vibration" << QString::number(info.controller.vibration_intensity);
 	args << "--present-mode" << EnumToText(info.present_mode);
 	if (info.gpu_index >= 0) {
 		args << "--gpu" << QString::number(info.gpu_index);
@@ -238,6 +240,9 @@ static QStringList CreateEmulatorArgs(const Configuration& info) {
 	if (info.fullscreen_enabled) {
 		args << "--fullscreen";
 	}
+	if (info.hide_cursor_enabled) {
+		args << "--hide-cursor";
+	}
 	args << "--readback-linear-images" << BoolArg(info.readback_linear_images);
 	if (info.tessellation_enabled) {
 		args << "--tessellation";
@@ -529,8 +534,7 @@ void MainDialogPrivate::Update() {
 		m_lightbar.Stop();
 		return;
 	}
-	m_lightbar.SetColor(item != nullptr ? m_ui->widget->CreateConfiguration(*item)->controller_color
-	                                    : m_ui->widget->GetGlobalControllerColor());
+	m_lightbar.SetColor(m_ui->widget->GetGlobalControllerColor());
 }
 
 #include "mainDialog.moc"

src/libs/controller.cpp

@@ -223,7 +223,13 @@ void GameController::CycleSetting(Setting setting) {
 
 float GameController::GetSettingScale(Setting setting) const {
 	const auto step = m_setting_steps[static_cast<size_t>(setting)].load(std::memory_order_relaxed);
-	return setting == Setting::SpeakerVolume ? SPEAKER_VOLUME[step] : INTENSITY[step];
+	if (setting == Setting::SpeakerVolume) {
+		return SPEAKER_VOLUME[step] * (Config::GetControllerSpeakerVolume() / 100.0f);
+	}
+	if (setting == Setting::VibrationIntensity) {
+		return INTENSITY[step] * (Config::GetControllerVibrationIntensity() / 100.0f);
+	}
+	return INTENSITY[step];
 }
 
 static uint8_t Scale(uint8_t value, float scale) {

src/libs/dialog.cpp

@@ -14,6 +14,81 @@
 // NOLINTNEXTLINE(modernize-concat-nested-namespaces)
 namespace Libs::Dialog {
 
+namespace {
+
+constexpr int DIALOG_STATUS_RUNNING  = 2;
+constexpr int DIALOG_STATUS_FINISHED = 3;
+
+struct DialogState {
+	int      status     = 0;
+	uint64_t generation = 0;
+};
+
+std::mutex  g_dialog_mutex;
+DialogState g_error;
+DialogState g_signin;
+uint64_t    g_dialog_generation        = 0;
+uint64_t    g_dialog_revision          = 0;
+int32_t     g_error_code               = 0;
+void (*g_dialog_visibility_callback)() = nullptr;
+
+void SetDialogStatus(DialogState& dialog, int status, std::unique_lock<std::mutex>& lock) {
+	const bool visibility_changed =
+	    (dialog.status == DIALOG_STATUS_RUNNING) != (status == DIALOG_STATUS_RUNNING);
+	dialog.status = status;
+	if (visibility_changed) {
+		g_dialog_revision++;
+		if (status == DIALOG_STATUS_RUNNING) {
+			dialog.generation = ++g_dialog_generation;
+		}
+	}
+	const auto callback = visibility_changed ? g_dialog_visibility_callback : nullptr;
+	lock.unlock();
+	if (callback != nullptr) {
+		callback();
+	}
+}
+
+} // namespace
+
+namespace SystemDialog {
+
+bool GetHostSnapshot(HostSnapshot* snapshot) {
+	std::scoped_lock lock(g_dialog_mutex);
+	if (g_signin.status == DIALOG_STATUS_RUNNING) {
+		*snapshot = {Kind::Signin, g_signin.generation, 0};
+		return true;
+	}
+	if (g_error.status == DIALOG_STATUS_RUNNING) {
+		*snapshot = {Kind::Error, g_error.generation, g_error_code};
+		return true;
+	}
+	return false;
+}
+
+VisualState GetVisualState() noexcept {
+	std::scoped_lock lock(g_dialog_mutex);
+	const bool       signin = g_signin.status == DIALOG_STATUS_RUNNING;
+	return {g_error.status == DIALOG_STATUS_RUNNING || signin, signin, g_dialog_revision};
+}
+
+void SetVisibilityCallback(void (*callback)()) {
+	std::scoped_lock lock(g_dialog_mutex);
+	g_dialog_visibility_callback = callback;
+}
+
+bool HostClose(uint64_t generation) {
+	std::unique_lock lock(g_dialog_mutex);
+	auto&            dialog = g_signin.status == DIALOG_STATUS_RUNNING ? g_signin : g_error;
+	if (dialog.status != DIALOG_STATUS_RUNNING || generation != dialog.generation) {
+		return false;
+	}
+	SetDialogStatus(dialog, DIALOG_STATUS_FINISHED, lock);
+	return true;
+}
+
+} // namespace SystemDialog
+
 namespace CommonDialog {
 
 LIB_NAME("CommonDialog", "CommonDialog");
@@ -207,6 +282,7 @@ constexpr int SIGNIN_ERROR_NOT_INITIALIZED     = static_cast<int>(0x81350001u);
 constexpr int SIGNIN_ERROR_ALREADY_INITIALIZED = static_cast<int>(0x81350002u);
 constexpr int SIGNIN_ERROR_PARAM_INVALID       = static_cast<int>(0x81350003u);
 constexpr int SIGNIN_ERROR_INVALID_STATE       = static_cast<int>(0x81350005u);
+constexpr int SIGNIN_ERROR_INVALID_USER_ID     = static_cast<int>(0x81350007u);
 
 struct SigninDialogResult {
 	int32_t result;
@@ -222,70 +298,78 @@ struct SigninDialogParam {
 static_assert(sizeof(SigninDialogResult) == 16);
 static_assert(sizeof(SigninDialogParam) == 16);
 
-static int g_signin_status = SIGNIN_STATUS_NONE;
-
 int KYTY_SYSV_ABI SigninDialogInitialize() {
 	PRINT_NAME();
-	if (g_signin_status != SIGNIN_STATUS_NONE) {
+	std::scoped_lock lock(g_dialog_mutex);
+	if (g_signin.status != SIGNIN_STATUS_NONE) {
 		return SIGNIN_ERROR_ALREADY_INITIALIZED;
 	}
-	g_signin_status = SIGNIN_STATUS_INITIALIZED;
+	g_signin.status = SIGNIN_STATUS_INITIALIZED;
 	return OK;
 }
 
 int KYTY_SYSV_ABI SigninDialogTerminate() {
 	PRINT_NAME();
-	if (g_signin_status == SIGNIN_STATUS_NONE) {
+	std::unique_lock lock(g_dialog_mutex);
+	if (g_signin.status == SIGNIN_STATUS_NONE) {
 		return SIGNIN_ERROR_NOT_INITIALIZED;
 	}
-	g_signin_status = SIGNIN_STATUS_NONE;
+	SetDialogStatus(g_signin, SIGNIN_STATUS_NONE, lock);
 	return OK;
 }
 
 int KYTY_SYSV_ABI SigninDialogOpen(const void* param) {
 	PRINT_NAME();
-	if (g_signin_status != SIGNIN_STATUS_INITIALIZED && g_signin_status != SIGNIN_STATUS_FINISHED) {
+	std::unique_lock lock(g_dialog_mutex);
+	if (g_signin.status == SIGNIN_STATUS_NONE) {
+		return SIGNIN_ERROR_NOT_INITIALIZED;
+	}
+	if (g_signin.status != SIGNIN_STATUS_INITIALIZED && g_signin.status != SIGNIN_STATUS_FINISHED) {
 		return SIGNIN_ERROR_INVALID_STATE;
 	}
-	if (param == nullptr ||
-	    static_cast<const SigninDialogParam*>(param)->size != sizeof(SigninDialogParam)) {
+	if (param == nullptr) {
 		return SIGNIN_ERROR_PARAM_INVALID;
 	}
-	g_signin_status = SIGNIN_STATUS_RUNNING;
+	const auto* p = static_cast<const SigninDialogParam*>(param);
+	if (p->size != sizeof(SigninDialogParam) || p->reserved[0] != 0 || p->reserved[1] != 0) {
+		return SIGNIN_ERROR_PARAM_INVALID;
+	}
+	if (p->user_id != Config::GetUserId()) {
+		return SIGNIN_ERROR_INVALID_USER_ID;
+	}
+	SetDialogStatus(g_signin, SIGNIN_STATUS_RUNNING, lock);
 	return OK;
 }
 
 int KYTY_SYSV_ABI SigninDialogClose() {
 	PRINT_NAME();
-	if (g_signin_status == SIGNIN_STATUS_NONE) {
+	std::unique_lock lock(g_dialog_mutex);
+	if (g_signin.status == SIGNIN_STATUS_NONE) {
 		return SIGNIN_ERROR_NOT_INITIALIZED;
 	}
-	if (g_signin_status != SIGNIN_STATUS_RUNNING && g_signin_status != SIGNIN_STATUS_FINISHED) {
-		return SIGNIN_ERROR_INVALID_STATE;
-	}
-	g_signin_status = SIGNIN_STATUS_FINISHED;
+	SetDialogStatus(g_signin, SIGNIN_STATUS_FINISHED, lock);
 	return OK;
 }
 
 int KYTY_SYSV_ABI SigninDialogUpdateStatus() {
 	PRINT_NAME();
-	if (g_signin_status == SIGNIN_STATUS_RUNNING) {
-		g_signin_status = SIGNIN_STATUS_FINISHED;
-	}
-	return g_signin_status;
+	std::scoped_lock lock(g_dialog_mutex);
+	return g_signin.status;
 }
 
 int KYTY_SYSV_ABI SigninDialogGetStatus() {
 	PRINT_NAME();
-	return g_signin_status;
+	std::scoped_lock lock(g_dialog_mutex);
+	return g_signin.status;
 }
 
 int KYTY_SYSV_ABI SigninDialogGetResult(void* result) {
 	PRINT_NAME();
-	if (g_signin_status == SIGNIN_STATUS_NONE) {
+	std::scoped_lock lock(g_dialog_mutex);
+	if (g_signin.status == SIGNIN_STATUS_NONE) {
 		return SIGNIN_ERROR_NOT_INITIALIZED;
 	}
-	if (g_signin_status != SIGNIN_STATUS_FINISHED) {
+	if (g_signin.status != SIGNIN_STATUS_FINISHED) {
 		return SIGNIN_ERROR_INVALID_STATE;
 	}
 	if (result == nullptr) {
@@ -661,64 +745,15 @@ struct ErrorDialogParam {
 
 static_assert(sizeof(ErrorDialogParam) == 16);
 
-static std::mutex g_error_mutex;
-static int        g_error_status     = STATUS_NONE;
-static uint64_t   g_error_generation = 0;
-static uint64_t   g_error_revision   = 0;
-static int32_t    g_error_code       = 0;
-static void (*g_error_visibility_callback)() = nullptr;
-
-static void SetStatus(int status, std::unique_lock<std::mutex>& lock) {
-	const bool visibility_changed =
-	    (g_error_status == STATUS_RUNNING) != (status == STATUS_RUNNING);
-	g_error_status = status;
-	if (visibility_changed) {
-		g_error_revision++;
-	}
-	const auto callback = visibility_changed ? g_error_visibility_callback : nullptr;
-	lock.unlock();
-	if (callback != nullptr) {
-		callback();
-	}
-}
-
-bool GetHostSnapshot(HostSnapshot* snapshot) {
-	std::scoped_lock lock(g_error_mutex);
-	if (g_error_status != STATUS_RUNNING) {
-		return false;
-	}
-	*snapshot = {g_error_generation, g_error_code};
-	return true;
-}
-
-VisualState GetVisualState() noexcept {
-	std::scoped_lock lock(g_error_mutex);
-	return {g_error_status == STATUS_RUNNING, g_error_revision};
-}
-
-void SetVisibilityCallback(void (*callback)()) {
-	std::scoped_lock lock(g_error_mutex);
-	g_error_visibility_callback = callback;
-}
-
-bool HostAccept(uint64_t generation) {
-	std::unique_lock lock(g_error_mutex);
-	if (g_error_status != STATUS_RUNNING || generation != g_error_generation) {
-		return false;
-	}
-	SetStatus(STATUS_FINISHED, lock);
-	return true;
-}
-
 int KYTY_SYSV_ABI ErrorDialogInitialize() {
 	PRINT_NAME();
 
-	std::scoped_lock lock(g_error_mutex);
-	if (g_error_status != STATUS_NONE) {
+	std::scoped_lock lock(g_dialog_mutex);
+	if (g_error.status != STATUS_NONE) {
 		return ERROR_ALREADY_INITIALIZED;
 	}
 
-	g_error_status = STATUS_INITIALIZED;
+	g_error.status = STATUS_INITIALIZED;
 
 	return OK;
 }
@@ -726,8 +761,8 @@ int KYTY_SYSV_ABI ErrorDialogInitialize() {
 int KYTY_SYSV_ABI ErrorDialogOpen(const void* param) {
 	PRINT_NAME();
 
-	std::unique_lock lock(g_error_mutex);
-	if (g_error_status != STATUS_INITIALIZED && g_error_status != STATUS_FINISHED) {
+	std::unique_lock lock(g_dialog_mutex);
+	if (g_error.status != STATUS_INITIALIZED && g_error.status != STATUS_FINISHED) {
 		return ERROR_INVALID_STATE;
 	}
 	if (param == nullptr) {
@@ -745,8 +780,7 @@ int KYTY_SYSV_ABI ErrorDialogOpen(const void* param) {
 	     p->size, static_cast<uint32_t>(p->error_code), p->user_id);
 
 	g_error_code = p->error_code;
-	g_error_generation++;
-	SetStatus(STATUS_RUNNING, lock);
+	SetDialogStatus(g_error, STATUS_RUNNING, lock);
 
 	return OK;
 }
@@ -754,12 +788,12 @@ int KYTY_SYSV_ABI ErrorDialogOpen(const void* param) {
 int KYTY_SYSV_ABI ErrorDialogClose() {
 	PRINT_NAME();
 
-	std::unique_lock lock(g_error_mutex);
-	if (g_error_status != STATUS_RUNNING) {
+	std::unique_lock lock(g_dialog_mutex);
+	if (g_error.status != STATUS_RUNNING) {
 		return ERROR_INVALID_STATE;
 	}
 
-	SetStatus(STATUS_FINISHED, lock);
+	SetDialogStatus(g_error, STATUS_FINISHED, lock);
 
 	return OK;
 }
@@ -767,12 +801,12 @@ int KYTY_SYSV_ABI ErrorDialogClose() {
 int KYTY_SYSV_ABI ErrorDialogTerminate() {
 	PRINT_NAME();
 
-	std::unique_lock lock(g_error_mutex);
-	if (g_error_status == STATUS_NONE) {
+	std::unique_lock lock(g_dialog_mutex);
+	if (g_error.status == STATUS_NONE) {
 		return ERROR_NOT_INITIALIZED;
 	}
 
-	SetStatus(STATUS_NONE, lock);
+	SetDialogStatus(g_error, STATUS_NONE, lock);
 
 	return OK;
 }
@@ -780,15 +814,15 @@ int KYTY_SYSV_ABI ErrorDialogTerminate() {
 int KYTY_SYSV_ABI ErrorDialogUpdateStatus() {
 	PRINT_NAME();
 
-	std::scoped_lock lock(g_error_mutex);
-	return g_error_status;
+	std::scoped_lock lock(g_dialog_mutex);
+	return g_error.status;
 }
 
 int KYTY_SYSV_ABI ErrorDialogGetStatus() {
 	PRINT_NAME();
 
-	std::scoped_lock lock(g_error_mutex);
-	return g_error_status;
+	std::scoped_lock lock(g_dialog_mutex);
+	return g_error.status;
 }
 
 } // namespace ErrorDialog

src/libs/dialog.h

@@ -68,22 +68,30 @@ int KYTY_SYSV_ABI MsgDialogProgressBarSetMsg(int target, const char* msg);
 
 } // namespace MsgDialog
 
-namespace ErrorDialog {
+namespace SystemDialog {
+
+enum class Kind { Error, Signin };
 
 struct HostSnapshot {
+	Kind     kind;
 	uint64_t generation;
 	int32_t  error_code;
 };
 
 struct VisualState {
 	bool     active;
+	bool     background;
 	uint64_t revision;
 };
 
 bool        GetHostSnapshot(HostSnapshot* snapshot);
 VisualState GetVisualState() noexcept;
 void        SetVisibilityCallback(void (*callback)());
-bool        HostAccept(uint64_t generation);
+bool        HostClose(uint64_t generation);
+
+} // namespace SystemDialog
+
+namespace ErrorDialog {
 
 int KYTY_SYSV_ABI ErrorDialogInitialize();
 int KYTY_SYSV_ABI ErrorDialogOpen(const void* param);

src/libs/libKernel.cpp

@@ -2189,6 +2189,7 @@ LIB_DEFINE(InitLibKernel_1_Posix) {
 	LIB_FUNC("wtkt-teR1so", Posix::pthread_attr_init);
 	LIB_FUNC("zHchY8ft5pk", Posix::pthread_attr_destroy);
 	LIB_FUNC("vQm4fDEsWi8", Posix::pthread_attr_getstack);
+	LIB_FUNC("-SrbXpGR1f0", Posix::pthread_attr_setstack);
 	LIB_FUNC("2Q0z6rnBrTE", Posix::pthread_attr_setstacksize);
 	LIB_FUNC("Ucsu-OK+els", Posix::pthread_attr_get_np);
 	LIB_FUNC("RtLRV-pBTTY", Posix::pthread_attr_getschedpolicy);
@@ -2827,6 +2828,7 @@ constexpr int32_t KERNEL_AIO_STATE_SUBMITTED  = 1;
 constexpr int32_t KERNEL_AIO_STATE_PROCESSING = 2;
 constexpr int32_t KERNEL_AIO_STATE_COMPLETED  = 3;
 constexpr int32_t KERNEL_AIO_STATE_ABORTED    = 4;
+constexpr int32_t KERNEL_AIO_STATE_NOTIFIED   = 0x10000;
 constexpr int32_t KERNEL_AIO_MAX_QUEUE        = 512;
 constexpr int32_t KERNEL_AIO_MAX_REQUESTS     = 128;
 
@@ -2939,10 +2941,31 @@ int KYTY_SYSV_ABI KernelAioPollRequest(int32_t id, int32_t* state) {
 	}
 
 	if (!kernel_aio_is_valid_id(id)) {
-		return LibKernel::KERNEL_ERROR_EINVAL;
+		*state = LibKernel::KERNEL_ERROR_ESRCH;
+		return OK;
 	}
 
-	*state = g_kernel_aio_state[id].load(std::memory_order_acquire);
+	auto current = g_kernel_aio_state[id].load(std::memory_order_acquire);
+	while (current == KERNEL_AIO_STATE_COMPLETED || current == KERNEL_AIO_STATE_ABORTED) {
+		if (g_kernel_aio_state[id].compare_exchange_weak(current, current | KERNEL_AIO_STATE_NOTIFIED,
+		                                                std::memory_order_acq_rel)) {
+			break;
+		}
+	}
+	*state = current == 0 ? LibKernel::KERNEL_ERROR_ESRCH : current;
+	return OK;
+}
+
+int KYTY_SYSV_ABI KernelAioPollRequests(int32_t* ids, int32_t num, int32_t* states) {
+	if (ids == nullptr || states == nullptr) {
+		return LibKernel::KERNEL_ERROR_EFAULT;
+	}
+	if (num <= 0 || num > KERNEL_AIO_MAX_REQUESTS) {
+		return LibKernel::KERNEL_ERROR_EINVAL;
+	}
+	for (int32_t i = 0; i < num; i++) {
+		KernelAioPollRequest(ids[i], &states[i]);
+	}
 	return OK;
 }
 
@@ -2954,7 +2977,7 @@ int KYTY_SYSV_ABI KernelAioWaitRequest(int32_t id, int32_t* state, uint32_t* use
 	}
 
 	if (!kernel_aio_is_valid_id(id)) {
-		return LibKernel::KERNEL_ERROR_EINVAL;
+		return KernelAioPollRequest(id, state);
 	}
 
 	uint32_t waited  = 0;
@@ -2970,8 +2993,7 @@ int KYTY_SYSV_ABI KernelAioWaitRequest(int32_t id, int32_t* state, uint32_t* use
 		current = g_kernel_aio_state[id].load(std::memory_order_acquire);
 	}
 
-	*state = current;
-	return OK;
+	return KernelAioPollRequest(id, state);
 }
 
 int KYTY_SYSV_ABI KernelAioDeleteRequest(int32_t id, int32_t* ret) {
@@ -2982,12 +3004,35 @@ int KYTY_SYSV_ABI KernelAioDeleteRequest(int32_t id, int32_t* ret) {
 	}
 
 	if (!kernel_aio_is_valid_id(id)) {
-		return LibKernel::KERNEL_ERROR_EINVAL;
+		*ret = LibKernel::KERNEL_ERROR_ESRCH;
+		return OK;
 	}
 
-	g_kernel_aio_state[id].store(KERNEL_AIO_STATE_ABORTED, std::memory_order_release);
-	*ret = OK;
+	auto current = g_kernel_aio_state[id].load(std::memory_order_acquire);
+	for (;;) {
+		const auto state = current & ~KERNEL_AIO_STATE_NOTIFIED;
+		if (state != KERNEL_AIO_STATE_COMPLETED && state != KERNEL_AIO_STATE_ABORTED) {
+			*ret = current == 0 ? LibKernel::KERNEL_ERROR_ESRCH : LibKernel::KERNEL_ERROR_EBUSY;
+			break;
+		}
+		if (g_kernel_aio_state[id].compare_exchange_weak(current, 0, std::memory_order_acq_rel)) {
+			*ret = OK;
+			break;
+		}
+	}
+	return OK;
+}
 
+int KYTY_SYSV_ABI KernelAioDeleteRequests(int32_t* ids, int32_t num, int32_t* rets) {
+	if (ids == nullptr || rets == nullptr) {
+		return LibKernel::KERNEL_ERROR_EFAULT;
+	}
+	if (num <= 0 || num > KERNEL_AIO_MAX_REQUESTS) {
+		return LibKernel::KERNEL_ERROR_EINVAL;
+	}
+	for (int32_t i = 0; i < num; i++) {
+		KernelAioDeleteRequest(ids[i], &rets[i]);
+	}
 	return OK;
 }
 
@@ -3103,6 +3148,7 @@ LIB_DEFINE(InitLibKernel_1_Equeue) {
 	LIB_FUNC("WDszmSbWuDk", EventQueue::KernelAddUserEventEdge);
 	LIB_FUNC("F6e0kwo4cnk", EventQueue::KernelTriggerUserEvent);
 	LIB_FUNC("LJDwdSNTnDg", EventQueue::KernelDeleteUserEvent);
+	LIB_FUNC("57ZK+ODEXWY", EventQueue::KernelAddTimerEvent);
 	LIB_FUNC("R74tt43xP6k", EventQueue::KernelAddHRTimerEvent);
 	LIB_FUNC("J+LF6LwObXU", EventQueue::KernelDeleteHRTimerEvent);
 	LIB_FUNC("bBfz7kMF2Ho", EventQueue::KernelAddAmprEvent);
@@ -3325,8 +3371,10 @@ LIB_DEFINE(InitLibKernel_1) {
 	LIB_FUNC("kbw4UHHSYy0", LibKernel::pthread_cxa_finalize);
 	LIB_FUNC("lLMT9vJAck0", LibKernel::clock_gettime);
 	LIB_FUNC("5TgME6AYty4", KernelAioDeleteRequest);
+	LIB_FUNC("Ft3EtsZzAoY", KernelAioDeleteRequests);
 	LIB_FUNC("HgX7+AORI58", KernelAioSubmitReadCommands);
 	LIB_FUNC("2pOuoWoCxdk", KernelAioPollRequest);
+	LIB_FUNC("o7O4z3jwKzo", KernelAioPollRequests);
 	LIB_FUNC("KOF-oJbQVvc", KernelAioWaitRequest);
 	LIB_FUNC("XQ8C8y+de+E", KernelAioSubmitWriteCommands);
 	LIB_FUNC("nu4a0-arQis", KernelAioInitializeParam);

src/libs/libNet.cpp

@@ -3138,9 +3138,21 @@ namespace LibNpWebApi2 {
 
 LIB_VERSION("NpWebApi2", 1, "NpWebApi2", 1, 1);
 
-constexpr int NP_WEBAPI2_ERROR_INVALID_ARGUMENT  = -2141899774; /* 0x80553402 */
-constexpr int NP_WEBAPI2_ERROR_REQUEST_NOT_FOUND = -2141899770; /* 0x80553406 */
-constexpr int NP_WEBAPI2_ERROR_NOT_SIGNED_IN     = -2141899769; /* 0x80553407 */
+constexpr int NP_WEBAPI2_ERROR_INVALID_ARGUMENT       = -2141899774; /* 0x80553402 */
+constexpr int NP_WEBAPI2_ERROR_INVALID_LIB_CONTEXT_ID = -2141899773; /* 0x80553403 */
+constexpr int NP_WEBAPI2_ERROR_LIB_CONTEXT_NOT_FOUND  = -2141899772; /* 0x80553404 */
+constexpr int NP_WEBAPI2_ERROR_REQUEST_NOT_FOUND      = -2141899770; /* 0x80553406 */
+constexpr int NP_WEBAPI2_ERROR_NOT_SIGNED_IN          = -2141899769; /* 0x80553407 */
+
+struct NpWebApi2MemoryPoolStats {
+	size_t  pool_size;
+	size_t  max_inuse_size;
+	size_t  current_inuse_size;
+	int32_t reserved;
+};
+
+static std::mutex            g_np_webapi2_context_mutex;
+static std::map<int, size_t> g_np_webapi2_contexts;
 
 struct NpWebApi2ResponseInformationOption {
 	int32_t http_status;
@@ -3171,7 +3183,30 @@ static int KYTY_SYSV_ABI NpWebApi2Initialize(int lib_http_ctx_id, size_t pool_si
 
 	static int id = 0;
 
-	return ++id;
+	if (pool_size > std::numeric_limits<size_t>::max() - 0x3fff) {
+		return NP_WEBAPI2_ERROR_INVALID_ARGUMENT;
+	}
+	std::lock_guard lock(g_np_webapi2_context_mutex);
+	g_np_webapi2_contexts[++id] = (pool_size + 0x3fff) & ~size_t {0x3fff};
+	return id;
+}
+
+static int KYTY_SYSV_ABI NpWebApi2GetMemoryPoolStats(int lib_ctx_id,
+                                                   NpWebApi2MemoryPoolStats* stats) {
+	if (stats == nullptr) {
+		return NP_WEBAPI2_ERROR_INVALID_ARGUMENT;
+	}
+	if (lib_ctx_id <= 0) {
+		return NP_WEBAPI2_ERROR_INVALID_LIB_CONTEXT_ID;
+	}
+	std::lock_guard lock(g_np_webapi2_context_mutex);
+	const auto context = g_np_webapi2_contexts.find(lib_ctx_id);
+	if (context == g_np_webapi2_contexts.end()) {
+		return NP_WEBAPI2_ERROR_LIB_CONTEXT_NOT_FOUND;
+	}
+	// Requests currently use host storage, not allocations from the library pool.
+	*stats = {context->second, 0, 0, 0};
+	return OK;
 }
 
 static int KYTY_SYSV_ABI NpWebApi2PushEventCreateHandle(int lib_ctx_id) {
@@ -3438,11 +3473,16 @@ static int KYTY_SYSV_ABI NpWebApi2Terminate(int lib_ctx_id) {
 
 	LOGF("\t lib_ctx_id = %d\n", lib_ctx_id);
 
-	return 0;
+	if (lib_ctx_id <= 0) {
+		return NP_WEBAPI2_ERROR_INVALID_LIB_CONTEXT_ID;
+	}
+	std::lock_guard lock(g_np_webapi2_context_mutex);
+	return g_np_webapi2_contexts.erase(lib_ctx_id) != 0 ? OK : NP_WEBAPI2_ERROR_LIB_CONTEXT_NOT_FOUND;
 }
 
 LIB_DEFINE(InitNet_1_NpWebApi2) {
 	LIB_FUNC("+o9816YQhqQ", LibNpWebApi2::NpWebApi2Initialize);
+	LIB_FUNC("Xweb+naPZ8Y", LibNpWebApi2::NpWebApi2GetMemoryPoolStats);
 	LIB_FUNC("WV1GwM32NgY", LibNpWebApi2::NpWebApi2PushEventCreateHandle);
 	LIB_FUNC("sk54bi6FtYM", LibNpWebApi2::NpWebApi2CreateUserContext);
 	LIB_FUNC("9X9+cneTGUU", LibNpWebApi2::NpWebApi2DeleteUserContext);

src/libs/libSystemService.cpp

@@ -163,8 +163,10 @@ static int KYTY_SYSV_ABI SystemServiceGetStatus(SystemServiceStatus* status) {
 		return SYSTEM_SERVICE_ERROR_PARAMETER;
 	}
 
-	*status                      = SystemServiceStatus();
-	status->is_system_ui_overlaid = Dialog::ErrorDialog::GetVisualState().active;
+	const auto dialog                  = Dialog::SystemDialog::GetVisualState();
+	*status                            = SystemServiceStatus();
+	status->is_system_ui_overlaid      = dialog.active;
+	status->is_in_background_execution = dialog.background;
 
 	return OK;
 }

src/loader/guestInstructionPatcher.cpp

@@ -1,12 +1,11 @@
 // SPDX-FileCopyrightText: Copyright 2026 shadPS4 Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
-#include "loader/redZonePatcher.h"
+#include "loader/guestInstructionPatcher.h"
 
 #include "common/assert.h"
 #include "common/logging/log.h"
 #include "common/virtualMemory.h"
-#include "loader/x64InstructionEmulator.h"
 
 #include <Zydis/Zydis.h>
 #include <algorithm>
@@ -22,7 +21,7 @@
 #include <set>
 #include <unordered_set>
 #include <vector>
-#if defined(_WIN32)
+#if !defined(__APPLE__)
 #include <xbyak/xbyak.h>
 #include <xbyak/xbyak_util.h>
 #endif
@@ -34,7 +33,7 @@
 #undef max
 #endif
 
-#if defined(_WIN32)
+#if !defined(__APPLE__)
 using namespace Xbyak::util;
 #endif
 
@@ -53,7 +52,7 @@ using s64 = int64_t;
 
 constexpr size_t NearJumpSize = 5;
 
-#if defined(_WIN32)
+#if !defined(__APPLE__)
 
 struct PatchModule {
 	std::mutex           mutex {};
@@ -62,7 +61,7 @@ struct PatchModule {
 	std::set<u8*>        patched;
 	Xbyak::CodeGenerator patch_gen;
 	Xbyak::CodeGenerator trampoline_gen;
-	bool                 trampoline_exhausted = false;
+	bool                 trampoline_exhaustion_reported = false;
 
 	PatchModule(u8* module_ptr, u64 module_size, u8* trampoline_ptr, u64 trampoline_size)
 	    : start(module_ptr), end(module_ptr + module_size), patch_gen(module_size, module_ptr),
@@ -81,14 +80,14 @@ static PatchModule* GetContainingModule(const void* ptr) {
 	return address >= module->start && address < module->end ? module : nullptr;
 }
 
-static bool HandleTrampolineError(PatchModule* module, const Xbyak::Error& error) {
-	if (static_cast<int>(error) != Xbyak::ERR_CODE_IS_TOO_BIG) {
+static bool HandleTrampolineError(PatchModule* module, int error) {
+	if (error != Xbyak::ERR_CODE_IS_TOO_BIG) {
 		return false;
 	}
-	if (!module->trampoline_exhausted) {
-		LOGF("Windows guest red-zone trampoline space exhausted for module %p\n",
+	if (!module->trampoline_exhaustion_reported) {
+		LOGF("Guest instruction trampoline space exhausted for module %p\n",
 		     static_cast<void*>(module->start));
-		module->trampoline_exhausted = true;
+		module->trampoline_exhaustion_reported = true;
 	}
 	return true;
 }
@@ -103,9 +102,6 @@ static ZydisDecoder& GetDecoder() {
 	return decoder;
 }
 
-#endif
-#if defined(_WIN32)
-
 namespace {
 
 constexpr size_t GuestRedZoneSize = 128;
@@ -136,11 +132,21 @@ struct DecodedFunction {
 	bool                                        requires_conservative_red_zone_tracking {};
 };
 
+enum class InstructionReplacement { None, ReciprocalSquareRoot, ExtractQ };
+
 struct InstructionRewrite {
-	bool protect_red_zone {};
-	bool protected_indirect_call {};
+	bool                   protect_red_zone {};
+	bool                   protected_indirect_call {};
+	InstructionReplacement replacement {};
 };
 
+InstructionPatchCounts& ReplacementCounts(GuestInstructionPatchResult& result,
+                                          InstructionReplacement       replacement) {
+	ASSERT(replacement != InstructionReplacement::None);
+	return replacement == InstructionReplacement::ReciprocalSquareRoot ? result.reciprocal_sqrt
+	                                                                   : result.extrq;
+}
+
 bool IsStackPointerRegister(ZydisRegister reg) {
 	return reg != ZYDIS_REGISTER_NONE &&
 	       ZydisRegisterGetLargestEnclosing(ZYDIS_MACHINE_MODE_LONG_64, reg) == ZYDIS_REGISTER_RSP;
@@ -266,7 +272,7 @@ DecodedCodeInstruction DecodeCodeInstruction(uintptr_t address, uintptr_t end) {
 		const s64 access_start = operand.mem.disp.value;
 		const s64 access_size  = std::max<s64>(operand.size / 8, 1);
 		const s64 range_start  = std::max(access_start, -static_cast<s64>(GuestRedZoneSize));
-		const s64 range_end    = std::min(access_start + access_size, 0LL);
+		const s64 range_end    = std::min<s64>(access_start + access_size, 0);
 		for (s64 offset = range_start; offset < range_end; ++offset) {
 			const size_t bit = static_cast<size_t>(offset + static_cast<s64>(GuestRedZoneSize));
 			if ((operand.actions & ZYDIS_OPERAND_ACTION_MASK_READ) != 0) {
@@ -738,9 +744,9 @@ bool GenerateProtectedIndirectCall(const DecodedCodeInstruction& decoded,
 	return true;
 }
 
-void CollectRedZoneMemoryInstructions(const DecodedFunction& function,
+void CollectRedZoneMemoryInstructions(const DecodedFunction&                   function,
                                       std::map<uintptr_t, InstructionRewrite>& rewrite_sites,
-                                      RedZonePatchResult& result) {
+                                      GuestInstructionPatchResult&             result) {
 	if (!function.uses_red_zone) {
 		return;
 	}
@@ -777,98 +783,262 @@ void CollectRedZoneMemoryInstructions(const DecodedFunction& function,
 	}
 }
 
-struct ReciprocalSquareRootSite {
-	uintptr_t address;
-	u8        length;
-	bool      requires_red_zone_protection;
-};
+// Compute 1/sqrt in double precision, rounded to float, with denormals treated as
+// signed zero. Preserve guest MXCSR, RFLAGS and the scratch register's full YMM.
+void GenerateReciprocalSquareRoot(const DecodedCodeInstruction& decoded,
+                                  Xbyak::CodeGenerator&         generator) {
+	const int destination   = decoded.operands[0].reg.value - ZYDIS_REGISTER_XMM0;
+	const int source        = decoded.operands[1].reg.value - ZYDIS_REGISTER_XMM0;
+	int       scratch_index = 0;
+	while (scratch_index == destination || scratch_index == source) {
+		++scratch_index;
+	}
+	const Xbyak::Ymm scratch(scratch_index);
+	Xbyak::Label     one;
+	Xbyak::Label     done;
+	constexpr size_t SpillSize = 48;
 
-void CollectReciprocalSquareRoots(const DecodedFunction& function,
-                                 std::map<uintptr_t, InstructionRewrite>& rewrite_sites,
-                                 std::vector<ReciprocalSquareRootSite>& sites) {
+	// Guest code uses the SysV red zone on both hosts; put our spills below it.
+	// LEA/MOV/SIMD leave RFLAGS intact, including when the guest red zone is live.
+	generator.lea(rsp, ptr[rsp - GuestRedZoneSize - SpillSize]);
+	generator.vmovdqu(ptr[rsp], scratch);
+	generator.stmxcsr(ptr[rsp + 32]);
+	generator.mov(dword[rsp + 36], 0x1fc0); // round to nearest, exceptions masked, DAZ
+	generator.ldmxcsr(ptr[rsp + 36]);
+	generator.vcvtps2pd(scratch, Xbyak::Xmm(source));
+	generator.vsqrtpd(scratch, scratch);
+	generator.vbroadcastsd(Xbyak::Ymm(destination), ptr[rip + one]);
+	generator.vdivpd(scratch, Xbyak::Ymm(destination), scratch);
+	generator.vcvtpd2ps(Xbyak::Xmm(destination), scratch); // Clears upper YMM lanes.
+	generator.ldmxcsr(ptr[rsp + 32]);
+	generator.vmovdqu(scratch, ptr[rsp]);
+	generator.lea(rsp, ptr[rsp + GuestRedZoneSize + SpillSize]);
+	generator.jmp(done);
+	// A failed label reference can have an incomplete displacement. Do not bind it.
+	if (Xbyak::GetError() != 0) {
+		return;
+	}
+	generator.L(one);
+	generator.dq(0x3ff0000000000000ULL);
+	generator.L(done);
+}
+
+bool IsSupportedExtractQ(const DecodedCodeInstruction& decoded) {
+	if (decoded.instruction.mnemonic != ZYDIS_MNEMONIC_EXTRQ) {
+		return false;
+	}
+	const auto is_xmm = [](const ZydisDecodedOperand& operand) {
+		return operand.type == ZYDIS_OPERAND_TYPE_REGISTER &&
+		       operand.reg.value >= ZYDIS_REGISTER_XMM0 &&
+		       operand.reg.value <= ZYDIS_REGISTER_XMM15;
+	};
+	const bool immediate = decoded.instruction.operand_count_visible == 3 &&
+	                       decoded.operands[1].type == ZYDIS_OPERAND_TYPE_IMMEDIATE &&
+	                       decoded.operands[2].type == ZYDIS_OPERAND_TYPE_IMMEDIATE;
+	if (!is_xmm(decoded.operands[0]) ||
+	    (!immediate &&
+	     !(decoded.instruction.operand_count_visible == 2 && is_xmm(decoded.operands[1])))) {
+		return false;
+	}
+	// Match the trap handler's encoding support: 66 [REX] 0F 78/79 ModRM [imm8 imm8].
+	const auto*  code   = reinterpret_cast<const u8*>(decoded.address);
+	const size_t opcode = (code[1] & 0xf0u) == 0x40u ? 2 : 1;
+	return code[0] == 0x66 && code[opcode] == 0x0f &&
+	       code[opcode + 1] == (immediate ? 0x78 : 0x79) &&
+	       decoded.instruction.length == opcode + (immediate ? 5 : 3);
+}
+
+void GenerateExtractQ(const DecodedCodeInstruction& decoded, Xbyak::CodeGenerator& generator) {
+	const Xbyak::Xmm destination(decoded.operands[0].reg.value - ZYDIS_REGISTER_XMM0);
+	// Keep spills below the guest red zone on both hosts, and restore all modified flags/GPRs.
+	generator.lea(rsp, ptr[rsp - GuestRedZoneSize]);
+	generator.pushfq();
+	generator.push(rax);
+	generator.push(rcx);
+	generator.push(rdx);
+	if (decoded.operands[1].type == ZYDIS_OPERAND_TYPE_IMMEDIATE) {
+		const u32 length = decoded.operands[1].imm.value.u & 63u;
+		const u32 index  = decoded.operands[2].imm.value.u & 63u;
+		const u64 mask   = length == 0 ? UINT64_MAX : (u64 {1} << length) - 1;
+		generator.movq(rax, destination);
+		if (index != 0) {
+			generator.shr(rax, index);
+		}
+		generator.mov(rdx, mask);
+	} else {
+		const Xbyak::Xmm source(decoded.operands[1].reg.value - ZYDIS_REGISTER_XMM0);
+		generator.movq(rax, source);
+		generator.mov(ecx, eax);
+		generator.neg(ecx);
+		generator.mov(rdx, UINT64_MAX);
+		// CL shifts use six bits: -length gives 64-length, with zero retaining all 64 bits.
+		generator.shr(rdx, cl);
+		generator.shr(rax, 8);
+		generator.mov(ecx, eax);
+		generator.movq(rax, destination);
+		generator.shr(rax, cl);
+	}
+	generator.and_(rax, rdx);
+	// Match the emulator: clear the upper XMM half, preserving the upper YMM half.
+	generator.movq(destination, rax);
+	generator.pop(rdx);
+	generator.pop(rcx);
+	generator.pop(rax);
+	generator.popfq();
+	generator.lea(rsp, ptr[rsp + GuestRedZoneSize]);
+}
+
+void CollectAmdInstructions(const DecodedFunction&                   function,
+                            std::map<uintptr_t, InstructionRewrite>& rewrite_sites,
+                            GuestInstructionPatchResult&             result) {
+	static const bool host_lacks_sse4a = !Xbyak::util::Cpu().has(Xbyak::util::Cpu::tSSE4a);
 	for (const auto& [address, decoded]: function.instructions) {
-		if (!X64InstructionEmulator::IsReciprocalSquareRoot(decoded.instruction,
-		                                                    decoded.operands.data())) {
+		const auto&            instruction = decoded.instruction;
+		InstructionReplacement replacement {};
+		if (instruction.mnemonic == ZYDIS_MNEMONIC_VRSQRTPS &&
+		    instruction.encoding == ZYDIS_INSTRUCTION_ENCODING_VEX &&
+		    instruction.raw.vex.offset == 0 && decoded.operands[0].size == 128 &&
+		    decoded.operands[1].type == ZYDIS_OPERAND_TYPE_REGISTER) {
+			replacement = InstructionReplacement::ReciprocalSquareRoot;
+		} else if (host_lacks_sse4a && IsSupportedExtractQ(decoded)) {
+			replacement = InstructionReplacement::ExtractQ;
+		} else {
 			continue;
 		}
-		const bool protect_red_zone = decoded.red_zone_live.any();
-		sites.push_back({address, decoded.instruction.length, protect_red_zone});
-		if (protect_red_zone) {
-			rewrite_sites[address].protect_red_zone = true;
-		}
+		++ReplacementCounts(result, replacement).found;
+		rewrite_sites[address].replacement = replacement;
 	}
 }
 
-uint64_t ApplyReciprocalSquareRootPatches(const PatchModule& module,
-                                        std::span<const ReciprocalSquareRootSite> sites,
-                                        uint64_t trampoline_addr, uint64_t trampoline_size) {
-	// Validate every required relocation before introducing any traps.
-	for (const auto& site: sites) {
-		if (site.requires_red_zone_protection &&
-		    !module.patched.contains(reinterpret_cast<u8*>(site.address))) {
-			EXIT("Cannot preserve the guest red zone at emulated VRSQRTPS 0x%016" PRIx64 "\n",
-			     static_cast<u64>(site.address));
-		}
+void MarkInstructionTrap(u8* code, const ZydisDecodedInstruction& instruction,
+                         InstructionReplacement replacement) {
+	if (replacement == InstructionReplacement::ReciprocalSquareRoot) {
+		// Clear a reserved VEX.vvvv bit, retaining both register operands for the handler.
+		code[instruction.raw.vex.size - 1] &= ~0x08u;
 	}
-
-	uint64_t patched = 0;
-	for (const auto& site: sites) {
-		if (!module.patched.contains(reinterpret_cast<u8*>(site.address))) {
-			patched += X64InstructionEmulator::PatchReciprocalSquareRoots(site.address, site.length);
-		}
-	}
-	return patched +
-	       X64InstructionEmulator::PatchReciprocalSquareRoots(trampoline_addr, trampoline_size);
+	// EXTRQ already raises #UD on hosts without SSE4a; keep its original encoding.
 }
 
-void RelocateRedZoneInstructions(PatchModule* module, const DecodedFunction& function,
-                                const std::map<uintptr_t, InstructionRewrite>& rewrite_sites,
-                                RedZonePatchResult& result) {
+void GenerateInstructionTrap(const DecodedCodeInstruction& decoded,
+                             InstructionReplacement replacement, Xbyak::CodeGenerator& generator) {
+	std::array<u8, ZYDIS_MAX_INSTRUCTION_LENGTH> code {};
+	std::memcpy(code.data(), reinterpret_cast<const void*>(decoded.address),
+	            decoded.instruction.length);
+	MarkInstructionTrap(code.data(), decoded.instruction, replacement);
+	generator.db(code.data(), decoded.instruction.length);
+}
+
+bool NeedsTrapRedZoneProtection(const DecodedCodeInstruction& decoded) {
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+	return decoded.red_zone_live.any();
+#else
+	return false;
+#endif
+}
+
+void TrapUnrelocatedInstructions(const PatchModule& module, const DecodedFunction& function,
+                                 const std::map<uintptr_t, InstructionRewrite>& rewrite_sites,
+                                 GuestInstructionPatchResult&                   result) {
+	for (const auto& [address, rewrite]: rewrite_sites) {
+		if (rewrite.replacement == InstructionReplacement::None ||
+		    module.patched.contains(reinterpret_cast<u8*>(address))) {
+			continue;
+		}
+		const auto& decoded = function.instructions.at(address);
+		if (NeedsTrapRedZoneProtection(decoded)) {
+			if (rewrite.replacement == InstructionReplacement::ExtractQ) {
+				// Unlike VRSQRTPS, leaving EXTRQ unchanged would still trap and corrupt the red
+				// zone.
+				EXIT("AMD CPU compatibility: cannot safely trap EXTRQ at %p (guest red zone is "
+				     "live)\n",
+				     reinterpret_cast<void*>(address));
+			}
+			continue;
+		}
+		MarkInstructionTrap(reinterpret_cast<u8*>(address), decoded.instruction,
+		                    rewrite.replacement);
+		++ReplacementCounts(result, rewrite.replacement).trapped;
+	}
+}
+
+void RelocateGuestInstructions(PatchModule* module, const DecodedFunction& function,
+                               const std::map<uintptr_t, InstructionRewrite>& rewrite_sites,
+                               GuestInstructionPatchResult&                   result) {
 	struct RelocationSpan {
 		std::vector<const DecodedCodeInstruction*> instructions;
 		uintptr_t                                  patch_start {};
 		uintptr_t                                  continuation {};
 		size_t                                     patch_size {};
+		bool                                       trap_replacements {};
 	};
 
-	const auto emit_span = [&](const RelocationSpan& span) -> std::optional<size_t> {
-		const size_t trampoline_offset = module->trampoline_gen.getSize();
-		if (module->trampoline_exhausted) {
-			return std::nullopt;
-		}
-		try {
+	const auto emit_span = [&](RelocationSpan& span) -> std::optional<size_t> {
+		auto&        generator         = module->trampoline_gen;
+		const size_t trampoline_offset = generator.getSize();
+		const bool   has_replacements =
+		    std::ranges::any_of(span.instructions, [&](const auto* decoded) {
+			    const auto rewrite = rewrite_sites.find(decoded->address);
+			    return rewrite != rewrite_sites.end() &&
+			           rewrite->second.replacement != InstructionReplacement::None;
+		    });
+		for (bool trap_replacements: {false, true}) {
+			span.trap_replacements = trap_replacements;
+			Xbyak::ClearError();
+			bool encoded = true;
 			for (const auto* decoded: span.instructions) {
-				const auto rewrite = rewrite_sites.find(decoded->address);
+				const auto rewrite     = rewrite_sites.find(decoded->address);
+				const auto replacement = rewrite != rewrite_sites.end()
+				                             ? rewrite->second.replacement
+				                             : InstructionReplacement::None;
 				const bool protected_indirect_call =
 				    rewrite != rewrite_sites.end() && rewrite->second.protected_indirect_call;
-				const bool protect_red_zone = rewrite != rewrite_sites.end() &&
-				                              rewrite->second.protect_red_zone &&
-				                              !protected_indirect_call;
+				const bool protect_red_zone =
+				    (rewrite != rewrite_sites.end() && rewrite->second.protect_red_zone &&
+				     !protected_indirect_call) ||
+				    (replacement != InstructionReplacement::None && trap_replacements &&
+				     NeedsTrapRedZoneProtection(*decoded));
 				if (protect_red_zone) {
-					module->trampoline_gen.lea(rsp, ptr[rsp - GuestRedZoneSize]);
+					generator.lea(rsp, ptr[rsp - GuestRedZoneSize]);
 				}
-				if (protected_indirect_call) {
-					if (!GenerateProtectedIndirectCall(*decoded, module->trampoline_gen)) {
-						module->trampoline_gen.setSize(trampoline_offset);
-						return std::nullopt;
+				if (replacement != InstructionReplacement::None) {
+					if (trap_replacements) {
+						GenerateInstructionTrap(*decoded, replacement, generator);
+					} else if (replacement == InstructionReplacement::ReciprocalSquareRoot) {
+						GenerateReciprocalSquareRoot(*decoded, generator);
+					} else {
+						GenerateExtractQ(*decoded, generator);
 					}
-				} else if (!EncodeRelocatedInstruction(*decoded, module->trampoline_gen)) {
-					module->trampoline_gen.setSize(trampoline_offset);
-					return std::nullopt;
+				} else if (protected_indirect_call) {
+					encoded = GenerateProtectedIndirectCall(*decoded, generator);
+				} else {
+					encoded = EncodeRelocatedInstruction(*decoded, generator);
 				}
 				if (protect_red_zone && !decoded->replaces_stack_pointer) {
-					module->trampoline_gen.lea(rsp, ptr[rsp + GuestRedZoneSize]);
+					generator.lea(rsp, ptr[rsp + GuestRedZoneSize]);
+				}
+				if (!encoded || Xbyak::GetError() != 0) {
+					break;
 				}
 			}
-			module->trampoline_gen.jmp(reinterpret_cast<void*>(span.continuation));
-		} catch (const Xbyak::Error& error) {
-			module->trampoline_gen.setSize(trampoline_offset);
-			if (HandleTrampolineError(module, error)) {
-				return std::nullopt;
+			if (encoded && Xbyak::GetError() == 0) {
+				generator.jmp(reinterpret_cast<void*>(span.continuation));
+				if (Xbyak::GetError() == 0) {
+					return trampoline_offset;
+				}
+			}
+			const int error = Xbyak::GetError();
+			// Roll back before retrying this uncommitted span with smaller trap replacements.
+			generator.reset();
+			generator.setSize(trampoline_offset);
+			if (error != 0 && !HandleTrampolineError(module, error)) {
+				EXIT("Guest instruction trampoline generation failed: %s\n",
+				     Xbyak::ConvertErrorToString(error));
+			}
+			if (error != Xbyak::ERR_CODE_IS_TOO_BIG || !has_replacements) {
+				break;
 			}
-			throw;
 		}
-		return trampoline_offset;
+		return std::nullopt;
 	};
 
 	const auto record_rewrites = [&](const RelocationSpan& span) {
@@ -881,6 +1051,14 @@ void RelocateRedZoneInstructions(PatchModule* module, const DecodedFunction& fun
 			if (rewrite->second.protect_red_zone && decoded->accesses_memory) {
 				++result.patched_memory_instruction_count;
 			}
+			if (rewrite->second.replacement != InstructionReplacement::None) {
+				auto& counts = ReplacementCounts(result, rewrite->second.replacement);
+				if (span.trap_replacements) {
+					++counts.trapped;
+				} else {
+					++counts.native;
+				}
+			}
 		}
 	};
 
@@ -983,6 +1161,7 @@ void RelocateRedZoneInstructions(PatchModule* module, const DecodedFunction& fun
 		patch_gen.setSize(span.patch_start - reinterpret_cast<uintptr_t>(patch_gen.getCode()));
 		patch_gen.jmp(trampoline, Xbyak::CodeGenerator::LabelType::T_NEAR);
 		patch_gen.nop(span.patch_size - NearJumpSize);
+		EXIT_IF(Xbyak::GetError() != 0);
 
 		record_rewrites(span);
 		patched_spans.emplace_back(span.patch_start, span.continuation);
@@ -1021,7 +1200,7 @@ void RelocateRedZoneInstructions(PatchModule* module, const DecodedFunction& fun
 			continue;
 		}
 
-		const RelocationSpan site_span {
+		RelocationSpan site_span {
 		    .instructions = {&site_instruction->second},
 		    .patch_start  = site,
 		    .continuation = site + site_instruction->second.instruction.length,
@@ -1049,12 +1228,14 @@ void RelocateRedZoneInstructions(PatchModule* module, const DecodedFunction& fun
 			patch_gen.reset();
 			patch_gen.setSize(*relay_slot - reinterpret_cast<uintptr_t>(patch_gen.getCode()));
 			patch_gen.jmp(site_trampoline, Xbyak::CodeGenerator::LabelType::T_NEAR);
+			EXIT_IF(Xbyak::GetError() != 0);
 
 			patch_gen.reset();
 			patch_gen.setSize(site - reinterpret_cast<uintptr_t>(patch_gen.getCode()));
 			patch_gen.jmp(reinterpret_cast<void*>(*relay_slot),
 			              Xbyak::CodeGenerator::LabelType::T_SHORT);
 			patch_gen.nop(site_span.patch_size - ShortJumpSize);
+			EXIT_IF(Xbyak::GetError() != 0);
 
 			record_rewrites(site_span);
 			patched_spans.emplace_back(site_span.patch_start, site_span.continuation);
@@ -1187,6 +1368,7 @@ void RelocateRedZoneInstructions(PatchModule* module, const DecodedFunction& fun
 			patch_gen.reset();
 			patch_gen.setSize(relay_address - reinterpret_cast<uintptr_t>(patch_gen.getCode()));
 			patch_gen.jmp(site_trampoline, Xbyak::CodeGenerator::LabelType::T_NEAR);
+			EXIT_IF(Xbyak::GetError() != 0);
 			std::erase(relay_slots, relay_address);
 			std::erase(short_relay_slots, relay_address);
 		} else {
@@ -1201,6 +1383,7 @@ void RelocateRedZoneInstructions(PatchModule* module, const DecodedFunction& fun
 			patch_gen.jmp(host_trampoline, Xbyak::CodeGenerator::LabelType::T_NEAR);
 			patch_gen.jmp(site_trampoline, Xbyak::CodeGenerator::LabelType::T_NEAR);
 			patch_gen.nop(host_span->patch_size - NearJumpSize * 2);
+			EXIT_IF(Xbyak::GetError() != 0);
 		}
 
 		uintptr_t jump_target = relay_address;
@@ -1210,6 +1393,7 @@ void RelocateRedZoneInstructions(PatchModule* module, const DecodedFunction& fun
 			                  reinterpret_cast<uintptr_t>(patch_gen.getCode()));
 			patch_gen.jmp(reinterpret_cast<void*>(jump_target),
 			              Xbyak::CodeGenerator::LabelType::T_SHORT);
+			EXIT_IF(Xbyak::GetError() != 0);
 			jump_target       = final_short_jump;
 			const auto parent = relay_parent.find(final_short_jump);
 			ASSERT(parent != relay_parent.end());
@@ -1223,6 +1407,7 @@ void RelocateRedZoneInstructions(PatchModule* module, const DecodedFunction& fun
 		patch_gen.jmp(reinterpret_cast<void*>(jump_target),
 		              Xbyak::CodeGenerator::LabelType::T_SHORT);
 		patch_gen.nop(site_span.patch_size - ShortJumpSize);
+		EXIT_IF(Xbyak::GetError() != 0);
 
 		if (host_span) {
 			record_rewrites(*host_span);
@@ -1240,10 +1425,13 @@ void RelocateRedZoneInstructions(PatchModule* module, const DecodedFunction& fun
 }
 } // namespace
 
-RedZonePatchResult PatchGuestInstructions(u64 segment_addr, u64 segment_size,
-                                          std::span<const uintptr_t> function_starts,
-                                          bool protect_memory, bool emulate_rsqrt) {
-	RedZonePatchResult result {};
+GuestInstructionPatchResult PatchGuestInstructions(u64 segment_addr, u64 segment_size,
+                                                   std::span<const uintptr_t> function_starts,
+                                                   bool protect_memory, bool emulate_amd) {
+	GuestInstructionPatchResult result {};
+	if (!protect_memory && !emulate_amd) {
+		return result;
+	}
 	auto*              module = GetContainingModule(reinterpret_cast<void*>(segment_addr));
 	if (module == nullptr || function_starts.empty()) {
 		return result;
@@ -1262,8 +1450,11 @@ RedZonePatchResult PatchGuestInstructions(u64 segment_addr, u64 segment_size,
 	starts.erase(unique_end, starts.end());
 
 	std::unique_lock lock {module->mutex};
-	const size_t trampoline_begin = module->trampoline_gen.getSize();
-	std::vector<ReciprocalSquareRootSite> reciprocal_sqrt_sites;
+	const size_t     trampoline_begin = module->trampoline_gen.getSize();
+	bool analyze_red_zone = protect_memory;
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+	analyze_red_zone |= emulate_amd;
+#endif
 	for (size_t function_index = 0; function_index < starts.size(); ++function_index) {
 		const uintptr_t function_start = starts[function_index];
 		const uintptr_t function_end =
@@ -1274,7 +1465,9 @@ RedZonePatchResult PatchGuestInstructions(u64 segment_addr, u64 segment_size,
 
 		++result.function_count;
 		auto function = DecodeFunction(function_start, function_end, segment_addr, segment_end);
-		AnalyzeRedZoneLiveness(function);
+		if (analyze_red_zone) {
+			AnalyzeRedZoneLiveness(function);
+		}
 		result.instruction_count += function.instructions.size();
 
 		std::map<uintptr_t, InstructionRewrite> rewrite_sites;
@@ -1286,22 +1479,17 @@ RedZonePatchResult PatchGuestInstructions(u64 segment_addr, u64 segment_size,
 		if (protect_memory) {
 			CollectRedZoneMemoryInstructions(function, rewrite_sites, result);
 		}
-		if (emulate_rsqrt) {
-			CollectReciprocalSquareRoots(function, rewrite_sites, reciprocal_sqrt_sites);
+		if (emulate_amd) {
+			CollectAmdInstructions(function, rewrite_sites, result);
 		}
 		if (!rewrite_sites.empty()) {
-			RelocateRedZoneInstructions(module, function, rewrite_sites, result);
+			RelocateGuestInstructions(module, function, rewrite_sites, result);
+			TrapUnrelocatedInstructions(*module, function, rewrite_sites, result);
 		}
 	}
-	// Preserve valid instruction encodings throughout CFG analysis and relocation.
-	// Only now turn reciprocal roots into traps, including the relocated copies.
 	const auto trampoline_addr =
 	    reinterpret_cast<u64>(module->trampoline_gen.getCode()) + trampoline_begin;
 	const auto trampoline_size = module->trampoline_gen.getSize() - trampoline_begin;
-	if (emulate_rsqrt) {
-		result.reciprocal_sqrt_instruction_count = ApplyReciprocalSquareRootPatches(
-		    *module, reciprocal_sqrt_sites, trampoline_addr, trampoline_size);
-	}
 	Common::VirtualMemory::FlushInstructionCache(segment_addr, segment_size);
 	if (trampoline_size != 0) {
 		Common::VirtualMemory::FlushInstructionCache(trampoline_addr, trampoline_size);
@@ -1311,7 +1499,8 @@ RedZonePatchResult PatchGuestInstructions(u64 segment_addr, u64 segment_size,
 
 #else
 
-RedZonePatchResult PatchGuestInstructions(u64, u64, std::span<const uintptr_t>, bool, bool) {
+GuestInstructionPatchResult PatchGuestInstructions(u64, u64, std::span<const uintptr_t>, bool,
+                                                   bool) {
 	return {};
 }
 
@@ -1492,9 +1681,9 @@ bool DecodeEhFrameFunctionStarts(uint64_t eh_frame_header_addr, uint64_t eh_fram
 	return true;
 }
 
-void RegisterRedZonePatchModule(void* module_ptr, uint64_t module_size, void* trampoline_area_ptr,
-                                uint64_t trampoline_area_size) {
-#if defined(_WIN32)
+void RegisterGuestInstructionPatchModule(void* module_ptr, uint64_t module_size,
+                                         void* trampoline_area_ptr, uint64_t trampoline_area_size) {
+#if !defined(__APPLE__)
 	EXIT_IF(module_ptr == nullptr || module_size == 0 || trampoline_area_ptr == nullptr ||
 	        trampoline_area_size == 0);
 	const auto module_addr = reinterpret_cast<u64>(module_ptr);
@@ -1511,8 +1700,8 @@ void RegisterRedZonePatchModule(void* module_ptr, uint64_t module_size, void* tr
 #endif
 }
 
-void UnregisterRedZonePatchModule(void* module_ptr) {
-#if defined(_WIN32)
+void UnregisterGuestInstructionPatchModule(void* module_ptr) {
+#if !defined(__APPLE__)
 	g_patch_modules.erase(reinterpret_cast<u64>(module_ptr));
 #else
 	(void)module_ptr;

src/loader/guestInstructionPatcher.h

@@ -0,0 +1,58 @@
+// SPDX-FileCopyrightText: Copyright 2026 shadPS4 Emulator Project
+// SPDX-License-Identifier: GPL-2.0-or-later
+
+#ifndef KYTY_LOADER_GUEST_INSTRUCTION_PATCHER_H_
+#define KYTY_LOADER_GUEST_INSTRUCTION_PATCHER_H_
+
+#include "common/common.h"
+
+#include <span>
+#include <vector>
+
+namespace Loader {
+
+struct InstructionPatchCounts {
+	uint64_t found   = 0;
+	uint64_t native  = 0;
+	uint64_t trapped = 0;
+
+	uint64_t Skipped() const { return found - native - trapped; }
+
+	InstructionPatchCounts& operator+=(const InstructionPatchCounts& other) {
+		found += other.found;
+		native += other.native;
+		trapped += other.trapped;
+		return *this;
+	}
+};
+
+struct GuestInstructionPatchResult {
+	uint64_t               function_count                           = 0;
+	uint64_t               instruction_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;
+	uint64_t               stack_dependent_memory_instruction_count = 0;
+	uint64_t               control_flow_memory_instruction_count    = 0;
+	uint64_t               unrelocatable_memory_instruction_count   = 0;
+	uint64_t               indirect_red_zone_function_count         = 0;
+	InstructionPatchCounts reciprocal_sqrt;
+	InstructionPatchCounts extrq;
+};
+
+void RegisterGuestInstructionPatchModule(void* module_ptr, uint64_t module_size,
+                                         void* trampoline_area_ptr, uint64_t trampoline_area_size);
+void UnregisterGuestInstructionPatchModule(void* module_ptr);
+
+// Apply enabled instruction fixes using native trampolines or safe trap fallbacks.
+GuestInstructionPatchResult PatchGuestInstructions(uint64_t segment_addr, uint64_t segment_size,
+                                                   std::span<const uintptr_t> function_starts,
+                                                   bool protect_memory, bool emulate_amd);
+
+bool DecodeEhFrameFunctionStarts(uint64_t eh_frame_header_addr, uint64_t eh_frame_header_size,
+                                 std::vector<uintptr_t>* function_starts);
+
+} // namespace Loader
+
+#endif /* KYTY_LOADER_GUEST_INSTRUCTION_PATCHER_H_ */

src/loader/redZonePatcher.h

@@ -1,42 +0,0 @@
-// SPDX-FileCopyrightText: Copyright 2026 shadPS4 Emulator Project
-// SPDX-License-Identifier: GPL-2.0-or-later
-
-#ifndef KYTY_LOADER_RED_ZONE_PATCHER_H_
-#define KYTY_LOADER_RED_ZONE_PATCHER_H_
-
-#include "common/common.h"
-
-#include <span>
-#include <vector>
-
-namespace Loader {
-
-struct RedZonePatchResult {
-	uint64_t function_count                           = 0;
-	uint64_t instruction_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;
-	uint64_t stack_dependent_memory_instruction_count = 0;
-	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);
-
-// 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);
-
-} // namespace Loader
-
-#endif /* KYTY_LOADER_RED_ZONE_PATCHER_H_ */

src/loader/runtimeLinker.cpp

@@ -17,8 +17,8 @@
 #include "kernel/pthread.h"
 #include "loader/elf.h"
 #include "loader/gamePatch.h"
+#include "loader/guestInstructionPatcher.h"
 #include "loader/jit.h"
-#include "loader/redZonePatcher.h"
 #include "loader/symbolDatabase.h"
 #include "loader/x64InstructionEmulator.h"
 
@@ -1812,38 +1812,39 @@ 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         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);
-		program->mapped_size += RED_ZONE_TRAMPOLINE_SIZE;
-	}
+#if !defined(__APPLE__)
+	const bool emulate_amd = Config::AmdCpuEnabled();
+#else
+	const bool emulate_amd = false;
 #endif
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+	const bool protect_memory_faults = Config::RedZoneProtectionEnabled();
+#else
+	const bool protect_memory_faults = false;
+#endif
+	const bool patch_guest_instructions = protect_memory_faults || emulate_amd;
+
+	constexpr uint64_t INSTRUCTION_TRAMPOLINE_SIZE = 8u * 1024u * 1024u;
+	if (patch_guest_instructions) {
+		EXIT_IF(INSTRUCTION_TRAMPOLINE_SIZE > UINT64_MAX - program->mapped_size);
+		program->mapped_size += INSTRUCTION_TRAMPOLINE_SIZE;
+	}
 
 	program->base_vaddr = Libs::LibKernel::Memory::AllocateProgramMemory(
 	    g_desired_base_addr, program->mapped_size, Common::VirtualMemory::Mode::ExecuteReadWrite,
 	    Common::PathToString(program->file_name.filename()).c_str());
 	EXIT_IF(program->base_vaddr == 0);
 
-#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
-	if (use_red_zone_protection) {
-		program->red_zone_trampoline_vaddr = program->base_vaddr + program->base_size_aligned;
-		program->red_zone_trampoline_size  = RED_ZONE_TRAMPOLINE_SIZE;
-		RegisterRedZonePatchModule(reinterpret_cast<void*>(program->base_vaddr),
-		                           program->base_size_aligned,
-		                           reinterpret_cast<void*>(program->red_zone_trampoline_vaddr),
-		                           program->red_zone_trampoline_size);
+	if (patch_guest_instructions) {
+		const auto trampoline_addr           = program->base_vaddr + program->base_size_aligned;
+		program->instruction_trampoline_size = INSTRUCTION_TRAMPOLINE_SIZE;
+		RegisterGuestInstructionPatchModule(
+		    reinterpret_cast<void*>(program->base_vaddr), program->base_size_aligned,
+		    reinterpret_cast<void*>(trampoline_addr), program->instruction_trampoline_size);
 	}
-#endif
 	if (!is_shared) {
-		program->tls.handler_vaddr = program->base_vaddr + program->base_size_aligned;
-#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
-		program->tls.handler_vaddr += program->red_zone_trampoline_size;
-#endif
+		program->tls.handler_vaddr =
+		    program->base_vaddr + program->base_size_aligned + program->instruction_trampoline_size;
 	}
 
 	g_desired_base_addr += CODE_BASE_INCR * (1 + program->mapped_size / CODE_BASE_INCR);
@@ -1863,10 +1864,8 @@ void RuntimeLinker::LoadProgramToMemory(Program* program) {
 	}
 
 	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
 
 	for (Elf64_Half i = 0; i < ehdr->e_phnum; i++) {
 		if (phdr[i].p_memsz != 0 && (phdr[i].p_type == PT_LOAD || phdr[i].p_type == PT_OS_RELRO)) {
@@ -1925,52 +1924,57 @@ void RuntimeLinker::LoadProgramToMemory(Program* program) {
 			program->proc_param_vaddr = phdr[i].p_vaddr + program->base_vaddr;
 		}
 
-#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
-		if (use_red_zone_protection && phdr[i].p_type == PT_GNU_EH_FRAME) {
+		if (patch_guest_instructions && phdr[i].p_type == PT_GNU_EH_FRAME) {
 			eh_frame_header_addr = phdr[i].p_vaddr + program->base_vaddr;
 			eh_frame_header_size = phdr[i].p_memsz;
 		}
-#endif
 	}
 
-#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
-	std::vector<uintptr_t> function_starts;
-	if (use_red_zone_protection) {
-		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());
+	if (patch_guest_instructions) {
+		std::vector<uintptr_t> function_starts;
+		const bool             have_function_starts =
+		    DecodeEhFrameFunctionStarts(eh_frame_header_addr, eh_frame_header_size,
+		                                &function_starts) &&
+		    !function_starts.empty();
+		const auto module_name = Common::PathToString(program->file_name.filename());
+		if (!have_function_starts) {
+			Log::WriteToConsoleAndLog(
+			    fmt::format("{}: {} not patched (function boundaries unavailable)\n",
+			                emulate_amd ? "AMD CPU compatibility" : "Guest red-zone protection",
+			                module_name));
 		}
-	}
-#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 =
-			    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",
-			     Common::PathToString(program->file_name.filename()).c_str(), result.function_count,
-			     result.red_zone_function_count, result.memory_instruction_count,
-			     result.patched_memory_instruction_count, result.short_memory_instruction_count,
-			     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;
+		GuestInstructionPatchResult totals {};
+		for (const auto& [segment_addr, segment_size]: executable_segments) {
+			const auto result = PatchGuestInstructions(segment_addr, segment_size, function_starts,
+			                                           protect_memory_faults, emulate_amd);
+			totals.reciprocal_sqrt += result.reciprocal_sqrt;
+			totals.extrq += result.extrq;
+			if (protect_memory_faults) {
+				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",
+				     module_name.c_str(),
+				     result.function_count, result.red_zone_function_count,
+				     result.memory_instruction_count, result.patched_memory_instruction_count,
+				     result.short_memory_instruction_count,
+				     result.stack_dependent_memory_instruction_count,
+				     result.control_flow_memory_instruction_count,
+				     result.unrelocatable_memory_instruction_count);
+			}
 		}
-#else
-		if (emulate_rsqrt) {
-			reciprocal_sqrt_count =
-			    X64InstructionEmulator::PatchReciprocalSquareRoots(segment_addr, segment_size);
-			Common::VirtualMemory::FlushInstructionCache(segment_addr, segment_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 (emulate_amd && have_function_starts) {
+			const auto  found   = totals.reciprocal_sqrt.found + totals.extrq.found;
+			const auto  skipped = totals.reciprocal_sqrt.Skipped() + totals.extrq.Skipped();
+			const char* status  = found == 0         ? "no matching instructions"
+			                      : skipped == found ? "not patched"
+			                      : skipped != 0     ? "partially patched"
+			                                         : "patched";
+			Log::WriteToConsoleAndLog(fmt::format(
+			    "AMD CPU compatibility: {} {} (VRSQRTPS: native={}, trapped={}, skipped={}; "
+			    "EXTRQ: native={}, trapped={}, skipped={})\n",
+			    module_name, status, totals.reciprocal_sqrt.native, totals.reciprocal_sqrt.trapped,
+			    totals.reciprocal_sqrt.Skipped(), totals.extrq.native, totals.extrq.trapped,
+			    totals.extrq.Skipped()));
 		}
 	}
 
@@ -1999,11 +2003,9 @@ void RuntimeLinker::DeleteProgram(Program* p) {
 
 	if (program->base_vaddr != 0 || program->mapped_size != 0) {
 		EXIT_IF(program->base_vaddr == 0 || program->mapped_size == 0);
-#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
-		if (program->red_zone_trampoline_size != 0) {
-			UnregisterRedZonePatchModule(reinterpret_cast<void*>(program->base_vaddr));
+		if (program->instruction_trampoline_size != 0) {
+			UnregisterGuestInstructionPatchModule(reinterpret_cast<void*>(program->base_vaddr));
 		}
-#endif
 		EXIT_IF(
 		    !Libs::LibKernel::Memory::FreeGuestMemory(program->base_vaddr, program->mapped_size));
 	}

src/loader/runtimeLinker.h

@@ -123,10 +123,7 @@ struct Program {
 	uint64_t                     base_size         = 0;
 	uint64_t                     base_size_aligned = 0;
 	uint64_t                     mapped_size       = 0;
-#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
-	uint64_t red_zone_trampoline_vaddr = 0;
-	uint64_t red_zone_trampoline_size  = 0;
-#endif
+	uint64_t instruction_trampoline_size = 0;
 	std::unique_ptr<SymbolDatabase> export_symbols;
 	std::unique_ptr<SymbolDatabase> import_symbols;
 	ThreadLocalStorage              tls;

src/loader/x64InstructionEmulator.cpp

@@ -2,7 +2,6 @@
 
 #include "common/common.h"
 
-#include <Zydis/Zydis.h>
 #include <bit>
 #include <cstring>
 #if !defined(__APPLE__)
@@ -713,46 +712,6 @@ static bool TryEmulateReciprocalSquareRoot(Context& context) {
 
 #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;
@@ -769,11 +728,8 @@ bool TryEmulate(void* native_context) {
 	Context context {static_cast<ucontext_t*>(native_context)};
 #endif
 #if !defined(__APPLE__)
-	if (TryEmulateReciprocalSquareRoot(context)) {
-		return true;
-	}
-	return TryEmulateMonitorxMwaitx(context) || TryEmulateSse4a(context) ||
-	       TryEmulateShaNi(context);
+	return TryEmulateReciprocalSquareRoot(context) || TryEmulateMonitorxMwaitx(context) ||
+	       TryEmulateSse4a(context) || TryEmulateShaNi(context);
 #else
 	return TryEmulateSse4a(context);
 #endif

src/loader/x64InstructionEmulator.h

@@ -1,14 +1,8 @@
 #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

@@ -54,11 +54,15 @@ static void PrintUsage() {
 	         Config::DEFAULT_USER_ID);
 	::printf("  --mic <name>                        Capture from this microphone; omit for silence.\n");
 	::printf("  --controller-color <#RRGGBB>        Override the controller lightbar color.\n");
+	::printf("  --controller-volume <0-100>         DualSense speaker volume. Default: 100.\n");
+	::printf("  --controller-vibration <0-100>      DualSense vibration intensity. Default: 100.\n");
 	::printf(
 	    "  --present-mode <value>               Fifo, Mailbox, or Immediate. Default: Mailbox.\n");
 	::printf(
 	    "  --gpu <index>                        Vulkan physical device index. Default: auto.\n");
 	::printf("  --fullscreen                         Run in borderless desktop fullscreen.\n");
+	::printf(
+	    "  --hide-cursor                        Hide the cursor after 2 s idle. Default: off.\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");
@@ -203,6 +207,11 @@ static bool ParseArgs(int argc, char* argv[], RunOptions& options, bool& show_he
 			continue;
 		}
 
+		if (arg == "--hide-cursor") {
+			options.config.hide_cursor_enabled = true;
+			continue;
+		}
+
 		if (arg == "--vr") {
 			options.config.vr_enabled = true;
 			continue;
@@ -324,6 +333,18 @@ static bool ParseArgs(int argc, char* argv[], RunOptions& options, bool& show_he
 				return false;
 			}
 			options.config.controller_color = color;
+		} else if (arg == "--controller-volume") {
+			if (!ParseUint32(value, options.config.controller_speaker_volume) ||
+			    options.config.controller_speaker_volume > 100) {
+				::printf("invalid controller volume: %s\n", value.c_str());
+				return false;
+			}
+		} else if (arg == "--controller-vibration") {
+			if (!ParseUint32(value, options.config.controller_vibration_intensity) ||
+			    options.config.controller_vibration_intensity > 100) {
+				::printf("invalid controller vibration intensity: %s\n", value.c_str());
+				return false;
+			}
 		} else if (arg == "--present-mode") {
 			if (!ParseEnum(value, options.config.present_mode)) {
 				::printf("invalid present mode: %s\n", value.c_str());

tests/ControllerSettingsTests.cpp

@@ -111,8 +111,10 @@ struct Controller {
 		haptics_handles_rumble = false;
 		Initialize();
 		Connect(1);
-		Check(GetSettingScale(Setting::SpeakerVolume) == 1.0f &&
-		          GetSettingScale(Setting::VibrationIntensity) == 1.0f &&
+		Check(GetSettingScale(Setting::SpeakerVolume) ==
+		              Config::GetControllerSpeakerVolume() / 100.0f &&
+		          GetSettingScale(Setting::VibrationIntensity) ==
+		              Config::GetControllerVibrationIntensity() / 100.0f &&
 		          GetSettingScale(Setting::TriggerEffectIntensity) == 1.0f,
 		      "controller initialization retained old settings");
 		rumble.clear();
@@ -161,6 +163,48 @@ void TestSettingCycles() {
 	CycleSetting(Setting::TriggerEffectIntensity);
 }
 
+void TestGlobalControllerLevels() {
+	Config::ConfigOptions options;
+	options.controller_speaker_volume      = 50;
+	options.controller_vibration_intensity = 25;
+	Config::Load(options);
+	{
+		Controller controller;
+		Check(GetSettingScale(Setting::SpeakerVolume) == 0.5f &&
+		          GetSettingScale(Setting::VibrationIntensity) == 0.25f &&
+		          GetSettingScale(Setting::TriggerEffectIntensity) == 1.0f,
+		      "global controller levels did not scale their outputs independently");
+		SetRumble(200, 100);
+		Check(rumble.back().large == 50 * 257 && rumble.back().small == 25 * 257,
+		      "global vibration level did not scale rumble");
+		CycleSetting(Setting::SpeakerVolume);
+		CycleSetting(Setting::VibrationIntensity);
+		Check(GetSettingScale(Setting::SpeakerVolume) == 0.0f &&
+		          GetSettingScale(Setting::VibrationIntensity) == 0.0f,
+		      "cycle hotkeys did not mute globally scaled outputs");
+		CycleSetting(Setting::SpeakerVolume);
+		CycleSetting(Setting::VibrationIntensity);
+		Check(GetSettingScale(Setting::SpeakerVolume) == 0.01f &&
+		          GetSettingScale(Setting::VibrationIntensity) == 0.0825f,
+		      "cycle hotkeys did not multiply the global controller levels");
+		Check(rumble.back().large == 17 * 257 && rumble.back().small == 8 * 257,
+		      "cycling vibration did not apply the combined intensity to cached rumble");
+	}
+	options.controller_speaker_volume      = 0;
+	options.controller_vibration_intensity = 0;
+	Config::Load(options);
+	{
+		Controller controller;
+		Check(GetSettingScale(Setting::SpeakerVolume) == 0.0f &&
+		          GetSettingScale(Setting::VibrationIntensity) == 0.0f,
+		      "zero global controller levels did not mute outputs");
+		SetRumble(200, 100);
+		Check(rumble.back().large == 0 && rumble.back().small == 0,
+		      "zero global vibration level did not mute motors");
+	}
+	Config::Load(Config::ConfigOptions {});
+}
+
 void TestVibrationLifetime() {
 	Controller controller;
 	SetRumble(255, 1);
@@ -287,6 +331,7 @@ void TestIndependentOutputsAndPadSwitch() {
 int main() {
 	Config::Initialize();
 	TestSettingCycles();
+	TestGlobalControllerLevels();
 	// Each following test initializes another controller and requires strong defaults.
 	TestVibrationLifetime();
 	TestMaskedTriggersAndValidation();

tests/EventQueueLifetimeTests.cpp

@@ -3,6 +3,7 @@
 
 #include <algorithm>
 #include <atomic>
+#include <chrono>
 #include <cstdio>
 #include <cstdlib>
 #include <memory>
@@ -27,6 +28,69 @@ void CheckConcurrentResult(int result, const char* text) {
 	Check(result == OK || result == KERNEL_ERROR_EBADF || result == KERNEL_ERROR_ENOENT, text);
 }
 
+void TestPeriodicTimerEvents() {
+	using namespace std::chrono_literals;
+	using Libs::LibKernel::KERNEL_ERROR_ETIMEDOUT;
+	Check(EventQueue::KernelAddTimerEvent(EventQueue::KERNEL_EQUEUE_INVALID, 10, 1000, nullptr) ==
+	          KERNEL_ERROR_EBADF,
+	      "timer rejects invalid queue");
+	EventQueue::KernelEqueue queue = EventQueue::KERNEL_EQUEUE_INVALID;
+	Check(EventQueue::KernelCreateEqueue(&queue, "periodic-timer") == OK, "create timer queue");
+	auto* user_data = reinterpret_cast<void*>(0x1234);
+	Check(EventQueue::KernelAddTimerEvent(queue, 10, 1000, user_data) == OK, "add periodic timer");
+	std::this_thread::sleep_for(5ms);
+	EventQueue::KernelEvent events[2] {};
+	int out = 0;
+	Libs::LibKernel::KernelUseconds timeout = 0;
+	Check(EventQueue::KernelWaitEqueue(queue, events, 2, &out, &timeout) == OK && out == 1,
+	      "periodic occurrences are aggregated into one event");
+	Check(events[0].ident == 10 && events[0].filter == EventQueue::KERNEL_EVFILT_TIMER &&
+	          events[0].data >= 5 && events[0].udata == user_data,
+	      "timer reports identity, userdata and elapsed occurrence count");
+	timeout = 100000;
+	Check(EventQueue::KernelWaitEqueue(queue, events, 2, &out, &timeout) == OK && out == 1 &&
+	          events[0].data >= 1,
+	      "periodic timer remains registered after delivery");
+
+	std::this_thread::sleep_for(5ms);
+	user_data = reinterpret_cast<void*>(0x5678);
+	Check(EventQueue::KernelAddTimerEvent(queue, 10, 3000000000u, user_data) == OK,
+	      "update periodic timer interval");
+	timeout = 0;
+	Check(EventQueue::KernelWaitEqueue(queue, events, 2, &out, &timeout) == OK && out == 1 &&
+	          events[0].data >= 5 && events[0].udata == user_data,
+	      "timer update preserves accumulated occurrences and updates userdata");
+	Check(EventQueue::KernelWaitEqueue(queue, events, 2, &out, &timeout) == KERNEL_ERROR_ETIMEDOUT,
+	      "timer update replaces the repeating interval");
+	Check(EventQueue::KernelDeleteEvent(queue, 10, EventQueue::KERNEL_EVFILT_TIMER) == OK,
+	      "delete periodic timer");
+
+	Check(EventQueue::KernelAddTimerEvent(queue, 10, 0, user_data) == OK,
+	      "add immediate periodic timer");
+	for (int i = 0; i < 2; ++i) {
+		Check(EventQueue::KernelAddTimerEvent(queue, 10, 0, user_data) == OK,
+		      "update pending immediate timer");
+	}
+	for (int i = 0; i < 2; ++i) {
+		Check(EventQueue::KernelWaitEqueue(queue, events, 2, &out, &timeout) == OK && out == 1 &&
+		          events[0].data == 1,
+		      "zero interval remains immediately ready after delivery");
+	}
+	Check(EventQueue::KernelDeleteEvent(queue, 10, EventQueue::KERNEL_EVFILT_TIMER) == OK,
+	      "delete immediate timer");
+
+	const Libs::LibKernel::KernelTimespec delay {0, 1000000};
+	Check(EventQueue::KernelAddHRTimerEvent(queue, 10, &delay, user_data) == OK,
+	      "add high resolution timer");
+	timeout = 100000;
+	Check(EventQueue::KernelWaitEqueue(queue, events, 2, &out, &timeout) == OK && out == 1 &&
+	          events[0].filter == EventQueue::KERNEL_EVFILT_HRTIMER && events[0].udata == user_data,
+	      "shared timer scheduling delivers high resolution event");
+	Check(EventQueue::KernelDeleteHRTimerEvent(queue, 10) == KERNEL_ERROR_ENOENT,
+	      "high resolution timer is removed after one delivery");
+	Check(EventQueue::KernelDeleteEqueue(queue) == OK, "delete timer queue");
+}
+
 void CountDeletedEvent(EventQueue::KernelEqueue, EventQueue::KernelEqueueEvent* event) {
 	auto* count = static_cast<std::atomic_uint32_t*>(event->filter.data);
 	count->fetch_add(1, std::memory_order_relaxed);
@@ -431,6 +495,7 @@ void TestConcurrentDelete() {
 } // namespace
 
 int main() {
+	TestPeriodicTimerEvents();
 	TestDuplicateAddPreservesEventState();
 	TestCallbackStateOutlivesPort();
 	TestCallbackOwnsPayload();

tests/KernelFileSystemTests.cpp

@@ -35,6 +35,10 @@ namespace Libs::LibKernelApr {
 void InitLibKernel_1_Apr(Loader::SymbolDatabase *symbols);
 }
 
+namespace Libs {
+void InitLibKernel_1(Loader::SymbolDatabase *symbols);
+}
+
 namespace Libs::LibAmpr {
 void InitAmpr_1(Loader::SymbolDatabase *symbols);
 }
@@ -43,6 +47,10 @@ namespace Libs::LibNet {
 void InitNet_1_Net(Loader::SymbolDatabase *symbols);
 }
 
+namespace Libs::LibNpWebApi2 {
+void InitNet_1_NpWebApi2(Loader::SymbolDatabase *symbols);
+}
+
 namespace {
 
 namespace FileSystem = Libs::LibKernel::FileSystem;
@@ -132,6 +140,119 @@ void TestSaveOpenVisibility() {
   Check(FileSystem::KernelClose(truncated) == OK, "close truncated save file");
 }
 
+void TestAioBatches() {
+  namespace Kernel = Libs::LibKernel;
+  Loader::SymbolDatabase symbols;
+  Libs::InitLibKernel_1(&symbols);
+  const auto find = [&](const char *nid) {
+    const auto *symbol = symbols.FindByNid(nid, Loader::SymbolType::Func);
+    Check(symbol != nullptr, "AIO exports resolve");
+    return symbol->vaddr;
+  };
+  struct Result { int64_t return_value; uint32_t state; };
+  struct Request {
+    int64_t offset;
+    size_t size;
+    void *buffer;
+    Result *result;
+    int32_t fd;
+  };
+  using Submit = int (KYTY_SYSV_ABI *)(Request *, int32_t, int32_t, int32_t *);
+  using Batch = int (KYTY_SYSV_ABI *)(int32_t *, int32_t, int32_t *);
+  using Single = int (KYTY_SYSV_ABI *)(int32_t, int32_t *);
+  using Wait = int (KYTY_SYSV_ABI *)(int32_t, int32_t *, uint32_t *);
+  const auto submit = reinterpret_cast<Submit>(find("HgX7+AORI58"));
+  const auto poll = reinterpret_cast<Batch>(find("o7O4z3jwKzo"));
+  const auto erase = reinterpret_cast<Batch>(find("Ft3EtsZzAoY"));
+  const auto poll_one = reinterpret_cast<Single>(find("2pOuoWoCxdk"));
+  const auto erase_one = reinterpret_cast<Single>(find("5TgME6AYty4"));
+  const auto wait = reinterpret_cast<Wait>(find("KOF-oJbQVvc"));
+  constexpr char Payload[] = "AIO payload";
+  const int fd = FileSystem::KernelOpen("/savedata0/aio.dat", 0x602, 0777);
+  Check(fd >= 3 && FileSystem::KernelWrite(fd, Payload, sizeof(Payload)) == sizeof(Payload),
+        "create AIO read fixture");
+  std::array<int32_t, 3> ids {};
+  std::array<char, 4> buffer {};
+  Result result {};
+  Request request {2, buffer.size(), buffer.data(), &result, fd};
+  for (int i = 0; i < 2; ++i) {
+    Check(submit(&request, 1, 2, &ids[i]) == OK &&
+              result.return_value == buffer.size() && result.state == 3 &&
+              std::memcmp(buffer.data(), Payload + 2, buffer.size()) == 0,
+          "AIO submission reads bytes at the requested offset");
+  }
+  ids[2] = -1;
+  std::array<int32_t, 3> states {-1, -1, -1};
+  Check(poll(ids.data(), ids.size(), states.data()) == OK &&
+            states[0] == 3 && states[1] == 3 && states[2] == Kernel::KERNEL_ERROR_ESRCH,
+        "batch poll writes every state and per-request invalid-ID error");
+  Check(poll_one(ids[0], &states[0]) == OK && states[0] == (3 | 0x10000) &&
+            poll(ids.data(), 2, states.data()) == OK &&
+            states[0] == (3 | 0x10000) && states[1] == (3 | 0x10000) &&
+            wait(ids[0], &states[0], nullptr) == OK && states[0] == (3 | 0x10000),
+        "single and batch polls and waits share completion notification state");
+  Check(erase(ids.data(), ids.size(), states.data()) == OK &&
+            states[0] == OK && states[1] == OK && states[2] == Kernel::KERNEL_ERROR_ESRCH,
+        "batch deletion writes per-request results");
+  Check(poll_one(ids[0], &states[0]) == OK && states[0] == Kernel::KERNEL_ERROR_ESRCH &&
+            erase_one(ids[1], &states[1]) == OK && states[1] == Kernel::KERNEL_ERROR_ESRCH,
+        "deleted AIO IDs are invalid for single-request APIs");
+  for (const auto batch : {poll, erase}) {
+    states.fill(42);
+    Check(batch(nullptr, 1, states.data()) == Kernel::KERNEL_ERROR_EFAULT &&
+              batch(ids.data(), 1, nullptr) == Kernel::KERNEL_ERROR_EFAULT &&
+              batch(ids.data(), 0, states.data()) == Kernel::KERNEL_ERROR_EINVAL &&
+              batch(ids.data(), 129, states.data()) == Kernel::KERNEL_ERROR_EINVAL &&
+              states == std::array<int32_t, 3> {42, 42, 42},
+          "invalid batch arguments do not modify outputs");
+    std::array<int32_t, 128> invalid_ids {};
+    std::array<int32_t, 128> errors {};
+    Check(batch(invalid_ids.data(), invalid_ids.size(), errors.data()) == OK &&
+              std::all_of(errors.begin(), errors.end(), [](int32_t error) {
+                return error == Kernel::KERNEL_ERROR_ESRCH;
+              }),
+          "maximum-sized batch returns all invalid-ID errors");
+  }
+  Check(FileSystem::KernelClose(fd) == OK, "close AIO read fixture");
+}
+
+void TestNpWebApi2Memory() {
+  Loader::SymbolDatabase symbols;
+  Libs::LibNpWebApi2::InitNet_1_NpWebApi2(&symbols);
+  const auto find = [&](const char *nid) {
+    const auto *symbol = symbols.FindByNid(nid, Loader::SymbolType::Func);
+    Check(symbol != nullptr, "NpWebApi2 library and memory exports resolve");
+    return symbol->vaddr;
+  };
+  struct Stats { size_t pool, maximum, current; int32_t reserved; };
+  static_assert(sizeof(Stats) == 32 && offsetof(Stats, reserved) == 24);
+  using Initialize = int (KYTY_SYSV_ABI *)(int, size_t);
+  using GetStats = int (KYTY_SYSV_ABI *)(int, Stats *);
+  using Terminate = int (KYTY_SYSV_ABI *)(int);
+  const auto initialize = reinterpret_cast<Initialize>(find("+o9816YQhqQ"));
+  const auto get_stats = reinterpret_cast<GetStats>(find("Xweb+naPZ8Y"));
+  const auto terminate = reinterpret_cast<Terminate>(find("bEvXpcEk200"));
+  const int first = initialize(1, 65537);
+  const int second = initialize(1, 16384);
+  Stats stats {1, 2, 3, 4};
+  Check(first > 0 && second > 0 && first != second &&
+            get_stats(first, &stats) == OK && stats.pool == 81920 &&
+            stats.maximum == 0 && stats.current == 0 && stats.reserved == 0 &&
+            get_stats(second, &stats) == OK && stats.pool == 16384,
+        "NpWebApi2 statistics retain each context's rounded pool capacity");
+  Check(get_stats(first, nullptr) == static_cast<int32_t>(0x80553402) &&
+            get_stats(0, &stats) == static_cast<int32_t>(0x80553403) &&
+            initialize(1, std::numeric_limits<size_t>::max()) == static_cast<int32_t>(0x80553402),
+        "NpWebApi2 statistics validate output, context ID, and pool rounding");
+  Check(terminate(first) == OK &&
+            get_stats(first, &stats) == static_cast<int32_t>(0x80553404) &&
+            terminate(first) == static_cast<int32_t>(0x80553404) &&
+            terminate(-1) == static_cast<int32_t>(0x80553403) &&
+            get_stats(second, &stats) == OK && stats.pool == 16384 &&
+            terminate(second) == OK,
+        "NpWebApi2 termination removes only the selected library context");
+}
+
 void CheckMountRoot(const std::filesystem::path &root) {
   Common::File cache;
   Check(cache.Create(root / "rpf.cache"), "create directory listing fixture");
@@ -836,10 +957,12 @@ int main(int, char**) {
   CheckAprPaths(temporary.Path());
   FileSystem::Mount(temporary.Path(), "/savedata0");
   TestSaveOpenVisibility();
+  TestAioBatches();
   CheckSaveRename(temporary.Path(), "first-save");
   CheckSaveRename(temporary.Path(), "replacement-save");
   FileSystem::Shutdown();
   CheckSocketWakeup();
+  TestNpWebApi2Memory();
   graphics.reset();
   subsystems.Destroy();
 

tests/MemoryTrackerTests.cpp

@@ -402,6 +402,100 @@ void TestCpuDirtyUpload() {
   Release(memory);
 }
 
+void TestCleanUploadPreservesOwnership() {
+  constexpr auto page_size = Libs::Graphics::TRACKER_PAGE_SIZE;
+  constexpr auto region_size = Libs::Graphics::TRACKER_REGION_SIZE;
+  TrackerHarness harness;
+  auto &tracker = harness.tracker;
+  auto *memory = AllocateFixedGuestRange(region_size * 2, region_size);
+  const auto address = reinterpret_cast<uint64_t>(memory);
+  tracker.ForEachUploadRange(address, region_size * 2, false,
+                             [](uint64_t, uint64_t) noexcept {},
+                             []() noexcept {});
+
+  for (const auto start : {address + 63 * page_size,
+                           address + region_size - page_size}) {
+    const auto before = start - page_size;
+    const auto after = start + 2 * page_size;
+    tracker.MarkRegionAsCpuModified(before, page_size);
+    tracker.MarkRegionAsCpuModified(after, page_size);
+    uint32_t ranges = 0;
+    uint32_t completions = 0;
+    ResetProtectionLog();
+    tracker.ForEachUploadRange(
+        start, 2 * page_size, false,
+        [&](uint64_t, uint64_t) noexcept { ranges++; },
+        [&]() noexcept { completions++; });
+    Check(ranges == 0 && completions == 1 && g_protection_calls == 0 &&
+              !tracker.IsRegionCpuModified(start, 2 * page_size),
+          "clean read-only upload changed dirty state or protection");
+
+    tracker.ForEachUploadRange(
+        start, 2 * page_size, true,
+        [&](uint64_t, uint64_t) noexcept { ranges++; },
+        [&]() noexcept { completions++; });
+    Check(ranges == 0 && completions == 2 &&
+              tracker.IsRegionGpuModified(start, page_size) &&
+              tracker.IsRegionGpuModified(start + page_size, page_size) &&
+              Protection(reinterpret_cast<void *>(start)) == PAGE_NOACCESS &&
+              Protection(reinterpret_cast<void *>(start + page_size)) ==
+                  PAGE_NOACCESS &&
+              tracker.IsRegionCpuModified(before, page_size) &&
+              tracker.IsRegionCpuModified(after, page_size) &&
+              !tracker.IsRegionGpuModified(before, page_size) &&
+              !tracker.IsRegionGpuModified(after, page_size) &&
+              IsWritable(reinterpret_cast<void *>(before)) &&
+              IsWritable(reinterpret_cast<void *>(after)),
+          "clean written upload lost GPU ownership or changed dirty neighbors");
+    tracker.UnmarkRegionAsGpuModified(start, 2 * page_size);
+  }
+  tracker.UntrackMemory(address, region_size * 2);
+  Release(memory);
+}
+
+void BenchmarkCleanUploads() {
+  constexpr uint64_t size = 256ull * 1024 * 1024;
+  TrackerHarness harness;
+  auto &tracker = harness.tracker;
+  auto *memory = AllocateFixedGuestRange(size, 0);
+  const auto address = reinterpret_cast<uint64_t>(memory);
+  tracker.ForEachUploadRange(address, size, false,
+                             [](uint64_t, uint64_t) noexcept {},
+                             []() noexcept {});
+  std::puts("chunk_bytes,sweeps,calls,elapsed_ns,ns_per_sweep,ns_per_call");
+  for (const uint64_t chunk : {size, Libs::Graphics::TRACKER_REGION_SIZE,
+                               uint64_t{64 * 1024}}) {
+    uint64_t ranges = 0;
+    uint64_t completions = 0;
+    uint64_t sweeps = 0;
+    ResetProtectionLog();
+    const auto start = std::chrono::steady_clock::now();
+    std::chrono::nanoseconds elapsed{};
+    do {
+      for (uint64_t offset = 0; offset < size; offset += chunk) {
+        tracker.ForEachUploadRange(
+            address + offset, chunk, false,
+            [&](uint64_t, uint64_t) noexcept { ranges++; },
+            [&]() noexcept { completions++; });
+      }
+      sweeps++;
+      elapsed = std::chrono::steady_clock::now() - start;
+    } while (elapsed < std::chrono::milliseconds(250));
+    Check(ranges == 0 && completions == sweeps * (size / chunk) &&
+              g_protection_calls == 0,
+          "clean upload benchmark changed ranges, completion, or protection");
+    std::printf("%llu,%llu,%llu,%lld,%.2f,%.2f\n",
+                static_cast<unsigned long long>(chunk),
+                static_cast<unsigned long long>(sweeps),
+                static_cast<unsigned long long>(completions),
+                static_cast<long long>(elapsed.count()),
+                static_cast<double>(elapsed.count()) / sweeps,
+                static_cast<double>(elapsed.count()) / completions);
+  }
+  tracker.UntrackMemory(address, size);
+  Release(memory);
+}
+
 void TestRangeInvalidation() {
   TrackerHarness harness;
   auto &tracker = harness.tracker;
@@ -1027,11 +1121,16 @@ int main(int argc, char **argv) {
   if (argc == 3 && std::strcmp(argv[1], "--death") == 0) {
     RunDeathCase(argv[2]);
   }
+  if (argc == 2 && std::strcmp(argv[1], "--benchmark-clean-upload") == 0) {
+    BenchmarkCleanUploads();
+    return 0;
+  }
   TestGuestRange();
   TestRangeSet();
   TestQueriesDoNotRequireMappedOwnership();
   TestConcurrentRegionPublication();
   TestCpuDirtyUpload();
+  TestCleanUploadPreservesOwnership();
   TestRangeInvalidation();
   TestGpuReacquisitionAfterInvalidation();
   TestGpuDirtyBits();

tests/ResourceMaterializationTests.cpp

@@ -95,29 +95,33 @@ Libs::Graphics::ShaderRecompiler::IR::ResourcePlan UserDataBufferPlan() {
   return ExtractResourcePlan(program);
 }
 
-Libs::Graphics::ShaderRecompiler::IR::ResourcePlan MixedSamplerPlan() {
+Libs::Graphics::ShaderRecompiler::IR::Program MixedSamplerProgram() {
   using namespace Libs::Graphics::ShaderRecompiler::IR;
   Program program;
   program.stage = Libs::Graphics::ShaderType::Compute;
   program.srt_plan_complete = true;
   program.resource_tracking_complete = true;
-  AddValueBlock(program);
+  auto &block = AddValueBlock(program);
 
-  const auto AddSource = [&program](uint32_t dword_count, uint32_t first) {
+  const auto AddSource = [&program](uint32_t dword_count) {
     DescriptorSource source;
     source.dword_count = dword_count;
-    source.dwords[0] = Value(first);
-    for (uint32_t i = 1; i < dword_count; i++) {
+    for (uint32_t i = 0; i < dword_count; i++) {
       source.dwords[i] = Value(0u);
     }
     program.descriptor_sources.push_back(source);
     return static_cast<uint32_t>(program.descriptor_sources.size() - 1u);
   };
 
-  const auto image0 = AddSource(8, 0);
-  const auto image1 = AddSource(8, 0);
-  const auto sampler0 = AddSource(4, 0x11111111u);
-  const auto sampler1 = AddSource(4, 0x22222222u);
+  const auto image0 = AddSource(8);
+  const auto image1 = AddSource(8);
+  const auto sampler0 = AddSource(4);
+  const auto sampler1 = AddSource(4);
+  for (uint32_t index = 0; index < 2; ++index) {
+    auto &value = block.AppendNewInst(
+        ValueOpcode::GetUserData, {Value(static_cast<ScalarReg>(index))});
+    program.descriptor_sources[sampler0 + index].dwords[0] = Value(&value);
+  }
   program.info.images.push_back(
       {.source = image0,
        .resource_class = ImageResourceClass::Sampled,
@@ -137,7 +141,7 @@ Libs::Graphics::ShaderRecompiler::IR::ResourcePlan MixedSamplerPlan() {
   program.info.sampled_pairs.push_back({.image = 0, .sampler = 0});
   program.info.sampled_pairs.push_back({.image = 0, .sampler = 1});
   program.info.sampled_pairs.push_back({.image = 1, .sampler = 1});
-  return ExtractResourcePlan(program);
+  return program;
 }
 
 void TestMappedSrtUsesDirectReaderByDefault() {
@@ -373,18 +377,35 @@ void TestFiniteImageRefreshReusesScalarReads() {
         "finite image refresh grew reusable resource storage after warmup");
 }
 
-void TestMixedSamplerDuplicatesTheCorrectSnapshot() {
+void TestMixedSamplerVariantsShareRuntimeDescriptor() {
   using namespace Libs::Graphics::ShaderRecompiler::IR;
-  auto plan = MixedSamplerPlan();
+  auto program = MixedSamplerProgram();
+  auto plan = ExtractResourcePlan(program);
+  std::array<uint32_t, 2> user_data{0x11111111u, 0x22222222u};
+  const SrtRuntime runtime{.user_data = user_data};
   ResourceSnapshot snapshot;
   ResourceSpecialization specialization;
-  Check(MaterializeResources(plan, {}, snapshot, specialization),
+  Check(MaterializeResources(plan, runtime, snapshot, specialization),
         "mixed sampler materialization failed");
-  Check(snapshot.samplers.size() == 3,
-        "mixed sampler materialization appended unrelated samplers");
-  Check(snapshot.samplers[2] == snapshot.samplers[1] &&
-            snapshot.samplers[2] != snapshot.samplers[0],
-        "point sampler variant duplicated the wrong runtime descriptor");
+  ApplyResourceSpecialization(program, specialization);
+  const auto &samplers = program.info.samplers;
+  Check(snapshot.samplers.size() == 2 && samplers.size() == 3 &&
+            samplers[0].snapshot_index == 0 && samplers[1].snapshot_index == 1 &&
+            samplers[2].snapshot_index == 1 &&
+            samplers[0].source == plan.info.samplers[0].source &&
+            samplers[1].source == plan.info.samplers[1].source &&
+            samplers[2].source == samplers[1].source &&
+            !samplers[1].force_point_filtering && samplers[2].force_point_filtering &&
+            program.info.sampled_pairs[2].sampler == 2,
+        "native sampler variants lost their source identity or binding order");
+  const auto capacity = snapshot.samplers.capacity();
+  user_data[1] = 0x33333333u;
+  Check(MaterializeResources(plan, runtime, snapshot, specialization) &&
+            snapshot.samplers.size() == 2 && snapshot.samplers.capacity() == capacity &&
+            snapshot.samplers[samplers[0].snapshot_index].dwords[0] == user_data[0] &&
+            snapshot.samplers[samplers[1].snapshot_index].dwords[0] == user_data[1] &&
+            snapshot.samplers[samplers[2].snapshot_index].dwords[0] == user_data[1],
+        "sampler variants retained stale or duplicated descriptors after refresh");
 }
 
 } // namespace
@@ -410,7 +431,7 @@ int main() {
   TestWrittenDescriptorUsesStrictReaderOnce();
   TestFailedMaterializationRejectsStage();
   TestFiniteImageRefreshReusesScalarReads();
-  TestMixedSamplerDuplicatesTheCorrectSnapshot();
+  TestMixedSamplerVariantsShareRuntimeDescriptor();
   std::puts("ResourceMaterializationTests: all cases passed");
   return 0;
 }

tests/ResourceTrackingTests.cpp

@@ -2032,6 +2032,58 @@ void TestFmaskLoadSpecialization() {
         "rebinding FMASK as a texture reused the metadata specialization");
 }
 
+void TestGatherLodSamplerValidation() {
+  for (const bool explicit_lod : {false, true}) {
+    for (const bool reverse : {false, true}) {
+      Fixture fixture;
+      std::array<Value, 8> image_words;
+      std::array<Value, 4> sampler_words;
+      for (uint32_t i = 0; i < image_words.size(); i++)
+        image_words[i] = fixture.UserData(i);
+      for (uint32_t i = 0; i < sampler_words.size(); i++)
+        sampler_words[i] = fixture.UserData(i + 8);
+      for (uint32_t i = 0; i < 2; i++) {
+        const bool lod = explicit_lod && i == (reverse ? 0u : 1u);
+        MemoryInfo memory;
+        memory.kind = ResourceKind::Image;
+        memory.image_dimension = Decoder::ImageDimension::Dim2D;
+        memory.image_sample_flags = lod ? Decoder::ImageSampleFlagLod
+                                        : Decoder::ImageSampleFlagLevelZero;
+        fixture.Emit(ValueOpcode::ImageGatherRaw,
+                     {fixture.Image(image_words), fixture.Sampler(sampler_words),
+                      fixture.ImageAddress(), image_words[1], image_words[3],
+                      sampler_words[1], sampler_words[2]},
+                     fixture.AddMemory(memory, i * 4));
+      }
+      fixture.PlanAndTrack();
+      auto plan = ExtractResourcePlan(fixture.program);
+      Check(plan.info.samplers.size() == 1 &&
+                plan.info.samplers[0].gather_lod == explicit_lod,
+            "shared sampler lost explicit gather validation or restricted LZ gather");
+      std::array<uint32_t, 12> user_data{};
+      user_data[0] = 0x1000u;
+      user_data[1] = static_cast<uint32_t>(
+          Libs::Graphics::Prospero::BufferFormat::k32Float) << 20u;
+      user_data[3] = Libs::Graphics::DstSel(4, 5, 6, 7) |
+          (static_cast<uint32_t>(Libs::Graphics::Prospero::ImageType::kColor2D) << 28u);
+      user_data[9] = 0xfff000u;
+      constexpr std::array sampler_controls{
+          0u, 0x4001u, 1u << 26u, (1u << 26u) | 256u,
+          (1u << 26u) | (1u << 14u), 2u << 26u, 3u << 26u};
+      for (const auto control : sampler_controls) {
+        user_data[10] = control;
+        ResourceSnapshot snapshot;
+        ResourceSpecialization specialization;
+        const bool supported = !explicit_lod || (control >> 26u) == 0 ||
+                               control == (1u << 26u);
+        Check(MaterializeResources(plan, {.user_data = user_data}, snapshot,
+                                   specialization) == supported,
+              "explicit gather accepted an unsupported sampler or rejected a valid one");
+      }
+    }
+  }
+}
+
 void TestDynamicStorageMipTracking() {
   Fixture fixture;
   std::array<Value, 8> image_words;
@@ -2067,7 +2119,7 @@ void TestDynamicStorageMipTracking() {
   const auto &images = fixture.program.info.images;
   Check(images.size() == 2 && images[0].mip_mode == ImageMipMode::None &&
             images[0].mip_count == 1 &&
-            images[1].mip_mode == ImageMipMode::DynamicStorage &&
+            images[1].mip_mode == ImageMipMode::Dynamic &&
             images[1].mip_count == 1,
         "storage mip writes did not share one dynamic logical resource");
   Check(plain.first.Instruction()->Flags<uint32_t>() == 0 &&
@@ -3139,15 +3191,15 @@ void TestShaderInfoAndBindingLayout() {
 void TestImageBindingAbi() {
   using NumericClass = Libs::Graphics::Prospero::TextureNumericClass;
 
-  Check(ImageBindingCount == 43u &&
+  Check(ImageBindingCount == 48u &&
             static_cast<uint32_t>(DescriptorBindingKind::Buffers) == 0u &&
-            static_cast<uint32_t>(DescriptorBindingKind::Samplers) == 44u &&
-            static_cast<uint32_t>(DescriptorBindingKind::Gds) == 45u &&
-            static_cast<uint32_t>(DescriptorBindingKind::BdaPagetable) == 46u &&
-            static_cast<uint32_t>(DescriptorBindingKind::FaultBuffer) == 47u &&
-            static_cast<uint32_t>(DescriptorBindingKind::FlattenedSrt) == 48u &&
-            static_cast<uint32_t>(DescriptorBindingKind::ShaderData) == 49u &&
-            static_cast<uint32_t>(DescriptorBindingKind::Count) == 50u,
+            static_cast<uint32_t>(DescriptorBindingKind::Samplers) == 49u &&
+            static_cast<uint32_t>(DescriptorBindingKind::Gds) == 50u &&
+            static_cast<uint32_t>(DescriptorBindingKind::BdaPagetable) == 51u &&
+            static_cast<uint32_t>(DescriptorBindingKind::FaultBuffer) == 52u &&
+            static_cast<uint32_t>(DescriptorBindingKind::FlattenedSrt) == 53u &&
+            static_cast<uint32_t>(DescriptorBindingKind::ShaderData) == 54u &&
+            static_cast<uint32_t>(DescriptorBindingKind::Count) == 55u,
         "native descriptor binding anchors changed");
 
   const std::array sampled_dimensions{
@@ -3170,12 +3222,13 @@ void TestImageBindingAbi() {
   uint32_t index = 0;
   const auto CheckBinding =
       [&](ImageResourceClass resource_class, NumericClass numeric_class,
-          Decoder::ImageDimension dimension, bool atomic, bool comparison = false) {
+          Decoder::ImageDimension dimension, bool atomic, bool comparison = false, bool atomic64 = false) {
         ImageResource image;
         image.resource_class = resource_class;
         image.numeric_class = numeric_class;
         image.dimension = dimension;
         image.atomic = atomic;
+        image.atomic64 = atomic64;
         image.depth_compare = comparison;
         const auto kind = DescriptorBindingForImage(image);
         Check(kind.has_value() &&
@@ -3210,6 +3263,10 @@ void TestImageBindingAbi() {
     CheckBinding(ImageResourceClass::Storage, NumericClass::Uint, dimension,
                  true);
   }
+  for (const auto dimension : storage_dimensions) {
+    CheckBinding(ImageResourceClass::Storage, NumericClass::Uint, dimension,
+                 true, false, true);
+  }
   Check(index == ImageBindingCount, "image descriptor ABI case count changed");
 
   const auto Invalid = [](ImageResource image) {
@@ -3407,6 +3464,7 @@ int main() {
     Run("images and samplers", TestImagesSamplersAndAliases);
     Run("SampleAdjust sampler scratch", TestSampleAdjustSamplerScratch);
     Run("FMASK load specialization", TestFmaskLoadSpecialization);
+    Run("gather LOD sampler validation", TestGatherLodSamplerValidation);
     Run("dynamic storage mips", TestDynamicStorageMipTracking);
     Run("invariant indirect images", TestInvariantIndirectImageMaterialization);
     Run("guarded direct image table", TestGuardedDirectImageTable);

tests/VirtualMemoryAllocationTests.cpp

@@ -7,7 +7,7 @@
 #include "kernel/memory.h"
 #include "kernel/pthread.h"
 #include "libs/errno.h"
-#include "loader/redZonePatcher.h"
+#include "loader/guestInstructionPatcher.h"
 #include "loader/runtimeLinker.h"
 #include "loader/systemContent.h"
 #include "loader/x64InstructionEmulator.h"
@@ -20,6 +20,7 @@
 #include <cmath>
 #include <cinttypes>
 #include <cstdint>
+#include <cstddef>
 #include <cstdio>
 #include <cstring>
 #include <filesystem>
@@ -31,6 +32,7 @@
 #include <csignal>
 #include <immintrin.h>
 #include <xbyak/xbyak.h>
+#include <xbyak/xbyak_util.h>
 #endif
 
 #if defined(__linux__)
@@ -232,15 +234,15 @@ void TestWindowsGuestRedZoneStaticPatcher() {
 	DWORD old_protection = 0;
 	Check(test, VirtualProtect(g_red_zone_fault_page, 0x1000, PAGE_NOACCESS, &old_protection) != FALSE,
 	      "failed to reset fault page protection");
-	Loader::RegisterRedZonePatchModule(reinterpret_cast<void*>(mapping), CODE_SIZE,
-	                                   reinterpret_cast<void*>(mapping + CODE_SIZE),
-	                                   TRAMPOLINE_SIZE);
+	Loader::RegisterGuestInstructionPatchModule(
+	    reinterpret_cast<void*>(mapping), CODE_SIZE, 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::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));
+	Loader::UnregisterGuestInstructionPatchModule(reinterpret_cast<void*>(mapping));
 	RemoveVectoredExceptionHandler(handler);
 	VirtualFree(g_red_zone_fault_page, 0, MEM_RELEASE);
 	g_red_zone_fault_page = nullptr;
@@ -2827,9 +2829,9 @@ void TestModuleRelocationUsesWritableHostMapping() {
 }
 
 #if defined(__linux__) || KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
-volatile sig_atomic_t g_rsqrt_traps = 0;
+volatile sig_atomic_t g_instruction_traps = 0;
 
-bool EmulateReciprocalSquareRootContext(void* context) {
+bool EmulateInstructionContext(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);
@@ -2838,83 +2840,92 @@ bool EmulateReciprocalSquareRootContext(void* context) {
 	if (!emulated || !preserved_mxcsr) {
 		return false;
 	}
-	g_rsqrt_traps = g_rsqrt_traps + 1;
+	g_instruction_traps = g_instruction_traps + 1;
 	return true;
 }
 
 #if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
-LONG CALLBACK ReciprocalSquareRootHandler(EXCEPTION_POINTERS* exception) {
+LONG CALLBACK InstructionHandler(EXCEPTION_POINTERS* exception) {
 	if (exception->ExceptionRecord->ExceptionCode != EXCEPTION_ILLEGAL_INSTRUCTION ||
-	    !EmulateReciprocalSquareRootContext(exception->ContextRecord)) {
+	    !EmulateInstructionContext(exception->ContextRecord)) {
 		return EXCEPTION_CONTINUE_SEARCH;
 	}
 	return EXCEPTION_CONTINUE_EXECUTION;
 }
 #else
-void ReciprocalSquareRootHandler(int, siginfo_t*, void* context) {
-	if (!EmulateReciprocalSquareRootContext(context)) {
+void InstructionHandler(int, siginfo_t*, void* context) {
+	if (!EmulateInstructionContext(context)) {
 		_exit(190);
 	}
 }
 #endif
 
+struct InstructionTestScope {
+	uint64_t mapping;
+	uint64_t size;
+	uint32_t mxcsr = _mm_getcsr();
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+	void* handler = nullptr;
+#else
+	struct sigaction previous {};
+	bool             installed = false;
+#endif
+	~InstructionTestScope() {
+		_mm_setcsr(mxcsr);
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+		if (handler != nullptr) {
+			RemoveVectoredExceptionHandler(handler);
+		}
+#else
+		if (installed) {
+			sigaction(SIGILL, &previous, nullptr);
+		}
+#endif
+		Loader::UnregisterGuestInstructionPatchModule(reinterpret_cast<void*>(mapping));
+		Libs::LibKernel::Memory::FreeGuestMemory(mapping, size);
+	}
+
+	void InstallHandler(const char* test) {
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+		handler = AddVectoredExceptionHandler(1, InstructionHandler);
+		Check(test, handler != nullptr, "failed to install instruction handler");
+#else
+		struct sigaction action {};
+		action.sa_sigaction = InstructionHandler;
+		action.sa_flags     = SA_SIGINFO;
+		sigemptyset(&action.sa_mask);
+		installed = sigaction(SIGILL, &action, &previous) == 0;
+		Check(test, installed, "failed to install instruction handler");
+#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
+	InstructionTestScope restore {mapping, allocation_size};
+	restore.InstallHandler(test);
 	using GuestFunction = void(KYTY_SYSV_ABI*)(const uint32_t*, uint32_t*);
 	const auto function = reinterpret_cast<GuestFunction>(mapping);
+	const std::array<uintptr_t, 1> function_starts {mapping};
+	const auto register_module = [&](uint64_t trampoline_size) {
+		Loader::RegisterGuestInstructionPatchModule(
+		    reinterpret_cast<void*>(mapping), code_size,
+		    reinterpret_cast<void*>(mapping + code_size), trampoline_size);
+	};
 	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 {};
+	std::array<uint32_t, 74> output {};
 	input.fill(0x3f800000);
-	for (const auto registers: {std::array {2, 1}, std::array {10, 9}, std::array {1, 1}}) {
+	for (const auto registers: {std::array {2, 1}, std::array {10, 9}, std::array {1, 1},
+	                            std::array {0, 1}, std::array {2, 0}, std::array {0, 0}}) {
 		const auto destination = registers[0];
 		const auto source = registers[1];
 		input[0] = 0x3f800000;
@@ -2924,8 +2935,8 @@ void TestPackedReciprocalSquareRoot() {
 			expected = 0x3f000000;
 		}
 		Xbyak::CodeGenerator code(code_size, reinterpret_cast<void*>(mapping));
-		if (source == 9) {
-			code.vmovups(Xbyak::Ymm(1), code.ptr[code.rdi + 32]);
+		for (int reg = 0; reg < 3; ++reg) {
+			code.vmovups(Xbyak::Ymm(reg), code.ptr[code.rdi + 32]);
 		}
 		code.vmovups(Xbyak::Ymm(source), code.ptr[code.rdi]);
 		if (destination != source) {
@@ -2935,18 +2946,32 @@ void TestPackedReciprocalSquareRoot() {
 			code.mov(code.rax, red_zone_sentinel ^ offset);
 			code.mov(code.qword[code.rsp - offset], code.rax);
 		}
+		code.cmp(code.rdi, code.rdi);
+		code.stc();
 		code.vrsqrtps(Xbyak::Xmm(destination), Xbyak::Xmm(source));
+		const uint64_t instruction_count = destination == 2 && source == 1 ? 2 : 1;
+		if (instruction_count == 2) {
+			// Adjacent four-byte instructions share one relocation span.
+			code.vrsqrtps(Xbyak::Xmm(destination), Xbyak::Xmm(source));
+		}
+		code.lea(code.rsp, code.ptr[code.rsp - 128]);
+		code.pushfq();
+		code.pop(code.rax);
+		code.lea(code.rsp, code.ptr[code.rsp + 128]);
+		code.mov(code.qword[code.rsi + 288], code.rax);
 		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));
+		for (int reg = 0; reg < 3; ++reg) {
+			code.vmovups(code.ptr[code.rsi + 192 + reg * 32], Xbyak::Ymm(reg));
+		}
 		code.vzeroupper();
 		code.ret();
-#if defined(__linux__)
+		Check(test, Xbyak::GetError() == 0, "failed to generate reciprocal-root fixture");
 		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());
@@ -2965,88 +2990,193 @@ void TestPackedReciprocalSquareRoot() {
 		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");
+		register_module(code_size);
+		const std::vector<uint8_t> original_trampoline(code.getCode() + code_size,
+		                                              code.getCode() + allocation_size);
+		const auto disabled = Loader::PatchGuestInstructions(
+		    mapping, code.getSize(), function_starts, false, false);
+		Check(test, Xbyak::GetError() == 0 && disabled.reciprocal_sqrt.found == 0 &&
+		                disabled.reciprocal_sqrt.native == 0 &&
+		                disabled.reciprocal_sqrt.trapped == 0 &&
+		                std::equal(original.begin(), original.end(), code.getCode()) &&
+		                std::equal(original_trampoline.begin(), original_trampoline.end(),
+		                           code.getCode() + code_size),
+		      "disabled instruction patching modified executable code or trampolines");
+		if (destination == 2 && source == 1) {
+			// Fail initially, after the RIP-relative constant reference, and inside its literal.
+			for (const uint64_t capacity: {1, 48, 80}) {
+				register_module(capacity);
+				const auto exhausted = Loader::PatchGuestInstructions(
+				    mapping, code.getSize(), function_starts, false, true);
 #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");
+				const uint64_t expected_traps = capacity == 1 ? 0 : instruction_count;
 #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");
+				const uint64_t expected_traps = instruction_count;
 #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);
+				Check(test, exhausted.reciprocal_sqrt.found == instruction_count &&
+				                exhausted.reciprocal_sqrt.native == 0 &&
+				                exhausted.reciprocal_sqrt.trapped == expected_traps,
+				      "exhausted reciprocal-root generation lost fallback coverage");
+				Check(test, Xbyak::GetError() == 0,
+				      "failed trampoline generation leaked Xbyak's thread-local error");
+				if (expected_traps == 0) {
+					Check(test, std::equal(original.begin(), original.end(), code.getCode()),
+					      "unsafe fallback modified a live guest red zone site");
+				}
+				const auto traps_before = g_instruction_traps;
 				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");
+				Check(test, static_cast<uint64_t>(g_instruction_traps - traps_before) == expected_traps &&
+				                red_zone_intact() &&
+				                output[0] == (expected_traps ? expected : std::bit_cast<uint32_t>(native)),
+				      "exhaustion fallback changed guest state or failed to execute");
+				std::memcpy(reinterpret_cast<void*>(mapping), original.data(), original.size());
 			}
+			register_module(code_size);
+		}
+		// Memory protection alone must retain the host's original reciprocal-root result.
+		const bool protect_memory = source == 9;
+		if (protect_memory) {
+			const auto memory_only = Loader::PatchGuestInstructions(
+			    mapping, code.getSize(), function_starts, true, false);
+			Check(test, Xbyak::GetError() == 0 &&
+			                memory_only.reciprocal_sqrt.found == 0 &&
+			                memory_only.reciprocal_sqrt.native == 0 &&
+			                memory_only.reciprocal_sqrt.trapped == 0 &&
+			                memory_only.patched_memory_instruction_count > 0,
+			      "memory protection enabled reciprocal-root replacement");
+			function(input.data(), output.data());
+			Check(test, output[0] == std::bit_cast<uint32_t>(native) && red_zone_intact(),
+			      "memory protection changed native reciprocal-root behavior");
+			std::memcpy(reinterpret_cast<void*>(mapping), original.data(), original.size());
+			register_module(code_size);
+		}
+		for (const bool trap_fallback: {false, true}) {
+			if (trap_fallback) {
+				std::memcpy(reinterpret_cast<void*>(mapping), original.data(), original.size());
+				register_module(48);
+			}
+			const bool protect_memory_sites = protect_memory && !trap_fallback;
+			const auto patched = Loader::PatchGuestInstructions(
+			    mapping, code.getSize(), function_starts, protect_memory_sites, true);
+			Check(test, Xbyak::GetError() == 0 &&
+			                patched.reciprocal_sqrt.found == instruction_count &&
+			                patched.reciprocal_sqrt.native ==
+			                    (trap_fallback ? 0 : instruction_count) &&
+			                patched.reciprocal_sqrt.trapped ==
+			                    (trap_fallback ? instruction_count : 0) &&
+			                patched.unrelocatable_memory_instruction_count == 0 &&
+			                (protect_memory_sites ? patched.patched_memory_instruction_count > 0
+			                                      : patched.patched_memory_instruction_count == 0),
+			      "instruction pass lost reciprocal root or memory patch coverage");
+			const auto traps_before = g_instruction_traps;
+			function(input.data(), output.data());
+			Check(test, static_cast<uint64_t>(g_instruction_traps - traps_before) ==
+			                (trap_fallback ? instruction_count : 0),
+			      "reciprocal-root execution did not match its native/fallback patch count");
+			Check(test, (output[72] & 0x8d5u) == 0x45u, "patched instruction corrupted RFLAGS");
+			for (int reg = 0; reg < 3; ++reg) {
+				if (reg != destination && reg != source) {
+					Check(test, std::equal(input.begin() + 8, input.end(), output.begin() + 48 + reg * 8),
+					      "patched instruction corrupted an unrelated YMM register");
+				}
+			}
+			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 %s native=%08x patched=%08x\n",
+			            destination, source, trap_fallback ? "fallback" : "replacement",
+			            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 replacement 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);
 		}
-		input.fill(0x3f800000);
 	}
-#if defined(__linux__)
-	std::array<uint8_t, 16> unknown {0x0f, 0x0b};
+	// An indirect branch prevents borrowing neighboring bytes for this four-byte site.
+	Xbyak::CodeGenerator unsupported_code(code_size, reinterpret_cast<void*>(mapping));
+	Xbyak::Label continuation;
+	unsupported_code.vmovups(unsupported_code.ymm1, unsupported_code.ptr[unsupported_code.rdi]);
+	unsupported_code.lea(unsupported_code.rax,
+	                     unsupported_code.ptr[unsupported_code.rip + continuation]);
+	const auto trap_offset = unsupported_code.getSize();
+	unsupported_code.vrsqrtps(unsupported_code.xmm1, unsupported_code.xmm1);
+	unsupported_code.jmp(unsupported_code.rax);
+	unsupported_code.L(continuation);
+	unsupported_code.vmovups(unsupported_code.ptr[unsupported_code.rsi], unsupported_code.ymm1);
+	unsupported_code.vzeroupper();
+	unsupported_code.ret();
+	Check(test, Xbyak::GetError() == 0, "failed to generate unsupported-relocation fixture");
+	std::vector<uint8_t> expected_fallback(unsupported_code.getCode(), unsupported_code.getCurr());
+	expected_fallback[trap_offset + 1] &= ~0x08u;
+	register_module(code_size);
+	const auto unsupported = Loader::PatchGuestInstructions(
+	    mapping, unsupported_code.getSize(), function_starts, false, true);
+	Check(test, Xbyak::GetError() == 0 && unsupported.reciprocal_sqrt.found == 1 &&
+	                unsupported.reciprocal_sqrt.native == 0 &&
+	                unsupported.reciprocal_sqrt.trapped == 1 &&
+	                std::equal(expected_fallback.begin(), expected_fallback.end(),
+	                           unsupported_code.getCode()),
+	      "unsafe native relocation did not select the in-place trap fallback");
+	const auto traps_before = g_instruction_traps;
+	function(input.data(), output.data());
+	Check(test, g_instruction_traps == traps_before + 1 && output[0] == 0x3f800000 &&
+	                std::all_of(output.begin() + 4, output.begin() + 8,
+	                            [](uint32_t lane) { return lane == 0; }),
+	      "unsafe native relocation fallback did not execute correctly");
+#if KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+	CONTEXT context {};
+	auto& instruction_pointer = context.Rip;
+#else
 	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;
+	auto& instruction_pointer = context.uc_mcontext.gregs[REG_RIP];
+#endif
+	std::array<uint8_t, 16> unknown {0x0f, 0x0b};
+	instruction_pointer = reinterpret_cast<uintptr_t>(unknown.data());
+	Check(test, !Loader::X64InstructionEmulator::TryEmulate(&context), "unrecognized UD2 was swallowed");
+	std::array<uint8_t, 16> unmarked_rsqrt {0xc5, 0xf8, 0x52, 0xd1};
+	instruction_pointer = reinterpret_cast<uintptr_t>(unmarked_rsqrt.data());
+	Check(test, !Loader::X64InstructionEmulator::TryEmulate(&context) &&
+	                instruction_pointer == reinterpret_cast<uintptr_t>(unmarked_rsqrt.data()),
+	      "unmarked native reciprocal root was incorrectly emulated");
+#if defined(__linux__)
 	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");
@@ -3098,6 +3228,316 @@ void TestPackedReciprocalSquareRoot() {
 #endif
 	std::printf("[host]    %-48s ok\n", test);
 }
+
+void TestPackedBitFieldExtract() {
+	const char*        test      = "PackedBitFieldExtract";
+	constexpr uint64_t code_size = 0x80000;
+	const auto         mapping   = Libs::LibKernel::Memory::AllocateProgramMemory(
+	    0x904000000, code_size * 2, Common::VirtualMemory::Mode::ExecuteReadWrite, "extrq_test");
+	Check(test, mapping != 0, "failed to allocate EXTRQ test code");
+	InstructionTestScope restore {mapping, code_size * 2};
+	restore.InstallHandler(test);
+	const bool host_sse4a      = Xbyak::util::Cpu().has(Xbyak::util::Cpu::tSSE4a);
+	const auto register_module = [&](uint64_t capacity) {
+		Loader::RegisterGuestInstructionPatchModule(reinterpret_cast<void*>(mapping), code_size,
+		                                            reinterpret_cast<void*>(mapping + code_size),
+		                                            capacity);
+	};
+	// Xbyak has no EXTRQ mnemonic. Emit the two canonical SSE4a encodings directly.
+	const auto extrq = [](Xbyak::CodeGenerator& code, int destination, int source,
+	                      uint8_t length = 0, uint8_t index = 0) {
+		code.db(0x66);
+		const int rex = source < 0 ? (destination >> 3) : ((destination >> 3) << 2) | (source >> 3);
+		if (rex != 0) {
+			code.db(0x40 | rex);
+		}
+		code.db(0x0f);
+		code.db(source < 0 ? 0x78 : 0x79);
+		code.db(source < 0 ? 0xc0 | (destination & 7)
+		                   : 0xc0 | ((destination & 7) << 3) | (source & 7));
+		if (source < 0) {
+			code.db(length);
+			code.db(index);
+		}
+	};
+	// Copy individual source bits so the oracle is independent of the emitter's shift/mask.
+	const auto extract = [](uint64_t value, uint32_t length, uint32_t index) {
+		const uint32_t count  = (length & 63u) == 0 ? 64 : length & 63u;
+		const uint32_t first  = index & 63u;
+		uint64_t       result = 0;
+		for (uint32_t bit = 0; bit < count && first + bit < 64; ++bit) {
+			result |= ((value >> (first + bit)) & 1u) << bit;
+		}
+		return result;
+	};
+	struct Result {
+		std::array<uint64_t, 4>  destination;
+		std::array<uint64_t, 4>  source;
+		std::array<uint64_t, 13> gprs;
+		uint64_t                 flags;
+		uint64_t                 stack_before;
+		uint64_t                 stack_after;
+		std::array<uint64_t, 16> red_zone;
+	};
+	constexpr uint64_t             gpr_sentinel      = 0x8192a3b4c5d6e700ull;
+	constexpr uint64_t             red_zone_sentinel = 0xa1b2c3d4e5f60000ull;
+	std::array<uint64_t, 12>       input {0,
+	                                      0xdeadbeef12345678,
+	                                      0x8877665544332211,
+	                                      0x1122334455667788,
+	                                      0,
+	                                      0x123456789abcdef0,
+	                                      0x13579bdf2468ace0,
+	                                      0xfdb97531eca86420,
+	                                      0x3f8000003f800000,
+	                                      0x3f8000003f800000,
+	                                      0x3f8000003f800000,
+	                                      0x3f8000003f800000};
+	const std::array<uintptr_t, 1> function_starts {mapping};
+	using GuestFunction = void(KYTY_SYSV_ABI*)(const uint64_t*, void*);
+	const auto function = reinterpret_cast<GuestFunction>(mapping);
+	for (const auto registers: {std::array {1, 8}, std::array {8, 9}, std::array {3, 3},
+	                            std::array {15, 0}, std::array {8, -1}}) {
+		const auto           destination = registers[0];
+		const auto           source      = registers[1];
+		Xbyak::CodeGenerator code(code_size, reinterpret_cast<void*>(mapping));
+		const std::array     saved {code.rbx, code.rbp, code.r12, code.r13, code.r14, code.r15};
+		const std::array gprs {code.rax, code.rcx, code.rdx, code.rbx, code.rbp, code.r8, code.r9,
+		                       code.r10, code.r11, code.r12, code.r13, code.r14, code.r15};
+		for (const auto& reg: saved) {
+			code.push(reg);
+		}
+		code.vmovups(Xbyak::Ymm(destination), code.ptr[code.rdi]);
+		if (source >= 0 && source != destination) {
+			code.vmovups(Xbyak::Ymm(source), 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.mov(code.eax, 0x7fffffff);
+		code.add(code.eax, 1);
+		code.stc(); // CF, PF, AF, SF and OF set; ZF clear.
+		for (size_t index = 0; index < gprs.size(); ++index) {
+			code.mov(gprs[index], gpr_sentinel | index);
+		}
+		code.mov(code.qword[code.rsi + offsetof(Result, stack_before)], code.rsp);
+		extrq(code, destination, source, 0xc8, 0xc4);
+		code.mov(code.qword[code.rsi + offsetof(Result, stack_after)], code.rsp);
+		for (size_t index = 0; index < gprs.size(); ++index) {
+			code.mov(code.qword[code.rsi + offsetof(Result, gprs) + index * 8], gprs[index]);
+		}
+		code.lea(code.rsp, code.ptr[code.rsp - 128]);
+		code.pushfq();
+		code.pop(code.rax);
+		code.lea(code.rsp, code.ptr[code.rsp + 128]);
+		code.mov(code.qword[code.rsi + offsetof(Result, flags)], code.rax);
+		for (uint32_t offset = 8; offset <= 128; offset += 8) {
+			code.mov(code.rax, code.qword[code.rsp - offset]);
+			code.mov(code.qword[code.rsi + offsetof(Result, red_zone) + offset - 8], code.rax);
+		}
+		code.vmovups(code.ptr[code.rsi + offsetof(Result, destination)], Xbyak::Ymm(destination));
+		code.vmovups(code.ptr[code.rsi + offsetof(Result, source)],
+		             Xbyak::Ymm(source < 0 ? destination : source));
+		code.vzeroupper();
+		for (auto reg = saved.rbegin(); reg != saved.rend(); ++reg) {
+			code.pop(*reg);
+		}
+		code.ret();
+		Check(test, Xbyak::GetError() == 0, "failed to generate EXTRQ state fixture");
+		const std::vector<uint8_t> original(code.getCode(), code.getCurr());
+		register_module(code_size);
+		const std::vector<uint8_t> trampoline(code.getCode() + code_size,
+		                                      code.getCode() + code_size * 2);
+		const auto                 disabled =
+		    Loader::PatchGuestInstructions(mapping, code.getSize(), function_starts, false, false);
+		Check(test,
+		      disabled.extrq.found == 0 &&
+		          std::equal(original.begin(), original.end(), code.getCode()) &&
+		          std::equal(trampoline.begin(), trampoline.end(), code.getCode() + code_size),
+		      "disabled AMD patching changed EXTRQ code or trampolines");
+		for (const bool fallback: {false, true}) {
+			std::memcpy(reinterpret_cast<void*>(mapping), original.data(), original.size());
+			register_module(fallback ? 32 : code_size);
+			const auto counts = Loader::PatchGuestInstructions(mapping, code.getSize(),
+			                                                   function_starts, false, true)
+			                        .extrq;
+			if (host_sse4a) {
+				Check(test,
+				      counts.found == 0 &&
+				          std::equal(original.begin(), original.end(), code.getCode()),
+				      "SSE4a host received unnecessary EXTRQ patches");
+				std::printf("[host]    %-48s skipped (native SSE4a host)\n", test);
+				return;
+			}
+			std::printf("[host]    EXTRQ xmm%d,%d %s found=%" PRIu64 " native=%" PRIu64
+			            " trapped=%" PRIu64 "\n",
+			            destination, source, fallback ? "fallback" : "replacement", counts.found,
+			            counts.native, counts.trapped);
+			Check(test,
+			      Xbyak::GetError() == 0 && counts.found == 1 && counts.Skipped() == 0 &&
+			          counts.native == (fallback ? 0u : 1u) &&
+			          counts.trapped == (fallback ? 1u : 0u),
+			      "EXTRQ native/fallback patch counts differ");
+			const auto traps_before = g_instruction_traps;
+			uint64_t   calls        = 0;
+			for (uint32_t length = 0; length < 64; ++length) {
+				for (uint32_t index = 0; index < 64; ++index) {
+					if (source < 0 && (length != 8 || index != 4)) {
+						continue;
+					}
+					for (const uint64_t pattern: {UINT64_MAX, uint64_t {0x0123456789abcdef}}) {
+						const uint64_t controls = (pattern & ~uint64_t {0xffff}) |
+						                          ((index | 0xc0u) << 8) | (length | 0xc0u);
+						input[0]                = source == destination ? controls : pattern;
+						input[4]                = controls;
+						Result output {};
+						_mm_setcsr(0x5fa5);
+						function(input.data(), &output);
+						const auto mxcsr = _mm_getcsr();
+						_mm_setcsr(restore.mxcsr);
+						++calls;
+						Check(test,
+						      output.destination[0] == extract(input[0], length, index) &&
+						          output.destination[1] == 0 && output.destination[2] == input[2] &&
+						          output.destination[3] == input[3],
+						      "EXTRQ result or destination lanes differ");
+						Check(test,
+						      source < 0 || source == destination
+						          ? output.source == output.destination
+						          : std::equal(output.source.begin(), output.source.end(),
+						                       input.begin() + 4),
+						      "EXTRQ changed its separate source register");
+						Check(test,
+						      mxcsr == 0x5fa5 && (output.flags & 0x8d5u) == 0x895u &&
+						          output.stack_before == output.stack_after,
+						      "EXTRQ changed MXCSR, RFLAGS or RSP");
+						for (size_t reg = 0; reg < output.gprs.size(); ++reg) {
+							Check(test, output.gprs[reg] == (gpr_sentinel | reg),
+							      "EXTRQ corrupted a guest GPR");
+						}
+						for (size_t slot = 0; slot < output.red_zone.size(); ++slot) {
+							Check(test,
+							      output.red_zone[slot] == (red_zone_sentinel | ((slot + 1) * 8)),
+							      "EXTRQ corrupted the guest red zone");
+						}
+					}
+				}
+			}
+			Check(test,
+			      static_cast<uint64_t>(g_instruction_traps - traps_before) ==
+			          (fallback ? calls : 0),
+			      "EXTRQ execution did not match native/fallback selection");
+		}
+	}
+
+	// Cover every immediate length/index pair without re-decoding thousands of large fixtures.
+	Xbyak::CodeGenerator   code(code_size, reinterpret_cast<void*>(mapping));
+	std::vector<uintptr_t> starts;
+	for (uint32_t length = 0; length < 64; ++length) {
+		for (uint32_t index = 0; index < 64; ++index) {
+			starts.push_back(mapping + code.getSize());
+			code.vmovups(code.ymm8, code.ptr[code.rdi]);
+			extrq(code, 8, -1, length | 0xc0, index | 0xc0);
+			code.vmovups(code.ptr[code.rsi], code.ymm8);
+			code.vzeroupper();
+			code.ret();
+		}
+	}
+	Check(test, Xbyak::GetError() == 0, "failed to generate exhaustive immediate EXTRQ fixture");
+	const std::vector<uint8_t> original(code.getCode(), code.getCurr());
+	for (const bool fallback: {false, true}) {
+		std::memcpy(reinterpret_cast<void*>(mapping), original.data(), original.size());
+		register_module(fallback ? 1 : code_size);
+		const auto counts =
+		    Loader::PatchGuestInstructions(mapping, code.getSize(), starts, false, true).extrq;
+		Check(test,
+		      Xbyak::GetError() == 0 && counts.found == starts.size() && counts.Skipped() == 0 &&
+		          counts.native == (fallback ? 0 : starts.size()) &&
+		          counts.trapped == (fallback ? starts.size() : 0),
+		      "immediate EXTRQ coverage was lost");
+		const auto traps_before = g_instruction_traps;
+		for (size_t entry = 0; entry < starts.size(); ++entry) {
+			for (const uint64_t pattern: {UINT64_MAX, uint64_t {0x0123456789abcdef}}) {
+				input[0] = pattern;
+				std::array<uint64_t, 4> output {};
+				reinterpret_cast<GuestFunction>(starts[entry])(input.data(), output.data());
+				Check(test,
+				      output[0] == extract(pattern, entry / 64, entry % 64) && output[1] == 0 &&
+				          output[2] == input[2] && output[3] == input[3],
+				      "immediate EXTRQ bit slice differs");
+			}
+		}
+		Check(test,
+		      static_cast<uint64_t>(g_instruction_traps - traps_before) ==
+		          (fallback ? starts.size() * 2 : 0),
+		      "immediate EXTRQ did not execute its selected path");
+	}
+
+	// An indirect branch prevents relocating this four-byte register form; retain its natural trap.
+	code.reset();
+	Xbyak::Label continuation;
+	code.vmovups(code.ymm1, code.ptr[code.rdi]);
+	code.vmovups(code.ymm2, code.ptr[code.rdi + 32]);
+	code.lea(code.rax, code.ptr[code.rip + continuation]);
+	extrq(code, 1, 2);
+	code.jmp(code.rax);
+	code.L(continuation);
+	code.vmovups(code.ptr[code.rsi], code.ymm1);
+	code.vzeroupper();
+	code.ret();
+	Check(test, Xbyak::GetError() == 0, "failed to generate unrelocatable EXTRQ fixture");
+	const std::vector<uint8_t> unsafe_original(code.getCode(), code.getCurr());
+	register_module(code_size);
+	const auto unsafe =
+	    Loader::PatchGuestInstructions(mapping, code.getSize(), function_starts, false, true).extrq;
+	Check(test,
+	      unsafe.found == 1 && unsafe.native == 0 && unsafe.trapped == 1 &&
+	          std::equal(unsafe_original.begin(), unsafe_original.end(), code.getCode()),
+	      "unrelocatable EXTRQ did not retain its natural trap");
+	input[0] = 0x0123456789abcdef;
+	input[4] = 0xffffc4c8;
+	std::array<uint64_t, 8> output {};
+	const auto              traps_before = g_instruction_traps;
+	function(input.data(), output.data());
+	Check(test, g_instruction_traps == traps_before + 1 && output[0] == extract(input[0], 8, 4),
+	      "unrelocatable EXTRQ fallback did not execute");
+
+	// Neighboring replacement kinds must share the same relocation/fallback span safely.
+	code.reset();
+	code.vmovups(code.ymm1, code.ptr[code.rdi]);
+	code.vmovups(code.ymm2, code.ptr[code.rdi + 32]);
+	code.vmovups(code.ymm3, code.ptr[code.rdi + 64]);
+	extrq(code, 1, 2);
+	code.vrsqrtps(code.xmm4, code.xmm3);
+	code.vmovups(code.ptr[code.rsi], code.ymm1);
+	code.vmovups(code.ptr[code.rsi + 32], code.ymm4);
+	code.vzeroupper();
+	code.ret();
+	Check(test, Xbyak::GetError() == 0, "failed to generate mixed AMD instruction fixture");
+	const std::vector<uint8_t> mixed_original(code.getCode(), code.getCurr());
+	for (const bool fallback: {false, true}) {
+		std::memcpy(reinterpret_cast<void*>(mapping), mixed_original.data(), mixed_original.size());
+		register_module(fallback ? 48 : code_size);
+		const auto mixed =
+		    Loader::PatchGuestInstructions(mapping, code.getSize(), function_starts, false, true);
+		for (const auto counts: {mixed.extrq, mixed.reciprocal_sqrt}) {
+			Check(test,
+			      counts.found == 1 && counts.native == (fallback ? 0u : 1u) &&
+			          counts.trapped == (fallback ? 1u : 0u),
+			      "mixed AMD instruction counts differ");
+		}
+		const auto before = g_instruction_traps;
+		function(input.data(), output.data());
+		Check(test,
+		      g_instruction_traps == before + (fallback ? 2 : 0) &&
+		          output[0] == extract(input[0], 8, 4) && output[4] == 0x3f8000003f800000 &&
+		          output[5] == 0x3f8000003f800000 && output[6] == 0 && output[7] == 0,
+		      "mixed native/fallback AMD instructions lost their results");
+	}
+	std::printf("[host]    %-48s ok (all 4096 length/index pairs)\n", test);
+}
+
 #endif
 
 #if defined(__x86_64__) || defined(_M_X64)
@@ -3274,6 +3714,10 @@ int main(int argc, char** argv) {
 	}
 #endif
 #if defined(__linux__) || KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
+	if (argc == 2 && std::strcmp(argv[1], "--extrq-only") == 0) {
+		RunTest(TestPackedBitFieldExtract);
+		return g_failed_tests == 0 ? 0 : 1;
+	}
 	if (argc == 2 && std::strcmp(argv[1], "--rsqrt-only") == 0) {
 		RunTest(TestPackedReciprocalSquareRoot);
 		return g_failed_tests == 0 ? 0 : 1;
@@ -3289,6 +3733,7 @@ int main(int argc, char** argv) {
 #endif
 #if defined(__linux__) || KYTY_PLATFORM == KYTY_PLATFORM_WINDOWS
 	RunTest(TestPackedReciprocalSquareRoot);
+	RunTest(TestPackedBitFieldExtract);
 #endif
 	RunTest(TestWindowsGuestRedZoneStaticPatcher);
 	RunTest(TestProsperoArgumentAndInfoSizeContracts);

tests/shaderCfgTests.cpp

@@ -1309,7 +1309,8 @@ void TestSpirvRequirementsAnalysis() {
   block->AppendNewInst(ValueOpcode::ImageQueryLod,
                        {Value(0u), Value(0u), Value(0u)});
   block->AppendNewInst(ValueOpcode::ImageGatherRaw,
-                       {Value(0u), Value(0u), Value(0u)});
+                       {Value(0u), Value(0u), Value(0u), Value(0u),
+                        Value(0u), Value(0u), Value(0u)});
   auto &shared = block->AppendNewInst(ValueOpcode::LoadSharedU32,
                                       {Value(0u), Value(true)});
   shared.SetFlags(MemoryFlags{.index = 0});
@@ -4269,6 +4270,38 @@ void TestScalarAshrI64Decoder() {
         "S_ASHR_I64 is missing from the decoded dump");
 }
 
+void TestImageAtomicWidthDecoder() {
+  using namespace ShaderRecompiler::Decoder;
+  // Captured PPSA29343 pixel shader: image_atomic_umax v[2:3], v[0:1], s[0:7].
+  const uint32_t captured[] = {0xf05c0308u, 0x00000200u};
+  Instruction decoded;
+  DecodeInstruction(captured, 0, decoded);
+  Check(decoded.opcode == Opcode::IMAGE_ATOMIC_UMAX &&
+            decoded.data_bits == 64u && decoded.data_dwords == 2u &&
+            decoded.dst.reg == 2u && decoded.dmask == 3u && !decoded.glc,
+        "captured image atomic lost its 64-bit data width");
+
+  for (const auto opcode : {0x10u, 0x17u}) {
+    for (uint32_t mask = 0; mask < 16u; ++mask) {
+      const uint32_t words[] = {EncodeMimg0(opcode, mask), captured[1]};
+      DecodeInstruction(words, 0, decoded);
+      const bool compare_swap = opcode == 0x10u;
+      const bool supported = mask == (compare_swap ? 3u : 1u) ||
+                             (!compare_swap && mask == 3u);
+      Check((decoded.opcode != Opcode::UNSUPPORTED) == supported,
+            "image atomic accepted an invalid or unsupported width mask");
+      if (supported) {
+        Check(decoded.data_bits == (!compare_swap && mask == 3u ? 64u : 32u),
+              "image atomic DMASK selected the wrong data width");
+      }
+    }
+  }
+  const uint32_t unsupported[] = {EncodeMimg0(0x11u, 3u), captured[1]};
+  DecodeInstruction(unsupported, 0, decoded);
+  Check(decoded.opcode == Opcode::UNSUPPORTED && decoded.data_bits == 64u,
+        "unsupported 64-bit image atomic silently decoded as 32-bit");
+}
+
 void TestNewShaderDecoderArchitecture() {
   using namespace ShaderRecompiler::Decoder;
 
@@ -14301,6 +14334,7 @@ int main() {
   // ShaderRecompilerComputeTests; keep the distinct decoder contract checks
   // here.
   TestScalarAshrI64Decoder();
+  TestImageAtomicWidthDecoder();
   TestNewShaderDecoderArchitecture();
   TestImageAddressOperands();
   TestSopkCompareImmediateExtension();