Missing var

Signed-off-by: Caio Oliveira <caiooliveirafarias0@gmail.com>

src/common/host_memory.cpp

@@ -578,7 +578,10 @@ public:
 #endif
         int flags = (fd > 0 ? MAP_SHARED : MAP_PRIVATE) | MAP_FIXED;
         void* ret = mmap(virtual_base + virtual_offset, length, prot_flags, flags, fd, host_offset);
-        ASSERT_MSG(ret != MAP_FAILED, "mmap: {}", strerror(errno));
+        ASSERT_MSG(ret != MAP_FAILED,
+               "mmap(virt_off=0x{:X}, host_off=0x{:X}, len=0x{:X}, virt_size=0x{:X}, backing_size=0x{:X}, perms=0x{:X}) failed: {}",
+               virtual_offset, host_offset, length, virtual_size, backing_size,
+               static_cast<u32>(perms), strerror(errno));
     }
 
     void Unmap(size_t virtual_offset, size_t length) {

src/common/settings.h

@@ -536,7 +536,7 @@ struct Values {
                                                 Category::RendererExtensions,
                                                 Specialization::Scalar};
 
-    SwitchableSetting<bool> vertex_input_dynamic_state{linkage, true, "vertex_input_dynamic_state", Category::RendererExtensions};
+    SwitchableSetting<bool> vertex_input_dynamic_state{linkage, false, "vertex_input_dynamic_state", Category::RendererExtensions};
     SwitchableSetting<bool> provoking_vertex{linkage, false, "provoking_vertex", Category::RendererExtensions};
     SwitchableSetting<bool> descriptor_indexing{linkage, false, "descriptor_indexing", Category::RendererExtensions};
 

src/common/slot_vector.h

@@ -130,6 +130,17 @@ public:
         ResetStorageBit(id.index);
     }
 
+    [[nodiscard]] bool Contains(SlotId id) const noexcept {
+        if (!id) {
+            return false;
+        }
+        const size_t word = id.index / 64;
+        if (word >= stored_bitset.size()) {
+            return false;
+        }
+        return ((stored_bitset[word] >> (id.index % 64)) & 1) != 0;
+    }
+
     [[nodiscard]] Iterator begin() noexcept {
         const auto it = std::ranges::find_if(stored_bitset, [](u64 value) { return value != 0; });
         if (it == stored_bitset.end()) {

src/core/hle/service/hle_ipc.cpp

@@ -1,8 +1,12 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2018 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
 #include <algorithm>
 #include <array>
+#include <cstring>
 #include <sstream>
 
 #include <boost/range/algorithm_ext/erase.hpp>
@@ -187,7 +191,7 @@ void HLERequestContext::ParseCommandBuffer(u32_le* src_cmdbuf, bool incoming) {
         buffer_w_descriptors.push_back(rp.PopRaw<IPC::BufferDescriptorABW>());
     }
 
-    const auto buffer_c_offset = rp.GetCurrentOffset() + command_header->data_size;
+    buffer_c_offset = rp.GetCurrentOffset() + command_header->data_size;
 
     if (!command_header->IsTipc()) {
         // Padding to align to 16 bytes
@@ -294,7 +298,15 @@ Result HLERequestContext::WriteToOutgoingCommandBuffer() {
     }
 
     // Write the domain objects to the command buffer, these go after the raw untranslated data.
-    // TODO(Subv): This completely ignores C buffers.
+
+    if (buffer_c_offset != 0 && !buffer_c_descriptors.empty()) {
+        constexpr u32 WORDS_PER_DESCRIPTOR = sizeof(IPC::BufferDescriptorC) / sizeof(u32);
+        u32 descriptor_offset = buffer_c_offset;
+        for (const auto& descriptor : buffer_c_descriptors) {
+            std::memcpy(&cmd_buf[descriptor_offset], &descriptor, sizeof(descriptor));
+            descriptor_offset += WORDS_PER_DESCRIPTOR;
+        }
+    }
 
     if (GetManager()->IsDomain()) {
         current_offset = domain_offset - static_cast<u32>(outgoing_domain_objects.size());
@@ -393,10 +405,14 @@ std::size_t HLERequestContext::WriteBuffer(const void* buffer, std::size_t size,
     const bool is_buffer_b{BufferDescriptorB().size() > buffer_index &&
                            BufferDescriptorB()[buffer_index].Size()};
     const std::size_t buffer_size{GetWriteBufferSize(buffer_index)};
+    if (buffer_size == 0) {
+        LOG_WARNING(Core, "WriteBuffer target index {} has zero capacity", buffer_index);
+        return 0;
+    }
     if (size > buffer_size) {
-        LOG_CRITICAL(Core, "size ({:016X}) is greater than buffer_size ({:016X})", size,
+        LOG_WARNING(Core, "size ({:016X}) is greater than buffer_size ({:016X}); clamping",
-                     buffer_size);
+                    size, buffer_size);
-        size = buffer_size; // TODO(bunnei): This needs to be HW tested
+        size = buffer_size;
     }
 
     if (is_buffer_b) {
@@ -418,15 +434,25 @@ std::size_t HLERequestContext::WriteBuffer(const void* buffer, std::size_t size,
 
 std::size_t HLERequestContext::WriteBufferB(const void* buffer, std::size_t size,
                                             std::size_t buffer_index) const {
-    if (buffer_index >= BufferDescriptorB().size() || size == 0) {
+    if (buffer_index >= BufferDescriptorB().size()) {
+        LOG_WARNING(Core, "WriteBufferB invalid buffer index {}", buffer_index);
+        return 0;
+    }
+
+    if (size == 0) {
+        LOG_WARNING(Core, "skip empty buffer write (B)");
         return 0;
     }
 
     const auto buffer_size{BufferDescriptorB()[buffer_index].Size()};
+    if (buffer_size == 0) {
+        LOG_WARNING(Core, "WriteBufferB target index {} has zero capacity", buffer_index);
+        return 0;
+    }
     if (size > buffer_size) {
-        LOG_CRITICAL(Core, "size ({:016X}) is greater than buffer_size ({:016X})", size,
+        LOG_WARNING(Core, "size ({:016X}) is greater than buffer_size ({:016X}); clamping",
-                     buffer_size);
+                    size, buffer_size);
-        size = buffer_size; // TODO(bunnei): This needs to be HW tested
+        size = buffer_size;
     }
 
     memory.WriteBlock(BufferDescriptorB()[buffer_index].Address(), buffer, size);
@@ -435,15 +461,25 @@ std::size_t HLERequestContext::WriteBufferB(const void* buffer, std::size_t size
 
 std::size_t HLERequestContext::WriteBufferC(const void* buffer, std::size_t size,
                                             std::size_t buffer_index) const {
-    if (buffer_index >= BufferDescriptorC().size() || size == 0) {
+    if (buffer_index >= BufferDescriptorC().size()) {
+        LOG_WARNING(Core, "WriteBufferC invalid buffer index {}", buffer_index);
+        return 0;
+    }
+
+    if (size == 0) {
+        LOG_WARNING(Core, "skip empty buffer write (C)");
         return 0;
     }
 
     const auto buffer_size{BufferDescriptorC()[buffer_index].Size()};
+    if (buffer_size == 0) {
+        LOG_WARNING(Core, "WriteBufferC target index {} has zero capacity", buffer_index);
+        return 0;
+    }
     if (size > buffer_size) {
-        LOG_CRITICAL(Core, "size ({:016X}) is greater than buffer_size ({:016X})", size,
+        LOG_WARNING(Core, "size ({:016X}) is greater than buffer_size ({:016X}); clamping",
-                     buffer_size);
+                    size, buffer_size);
-        size = buffer_size; // TODO(bunnei): This needs to be HW tested
+        size = buffer_size;
     }
 
     memory.WriteBlock(BufferDescriptorC()[buffer_index].Address(), buffer, size);
@@ -473,12 +509,20 @@ std::size_t HLERequestContext::GetWriteBufferSize(std::size_t buffer_index) cons
         ASSERT_OR_EXECUTE_MSG(
             BufferDescriptorB().size() > buffer_index, { return 0; },
             "BufferDescriptorB invalid buffer_index {}", buffer_index);
-        return BufferDescriptorB()[buffer_index].Size();
+        const auto size = BufferDescriptorB()[buffer_index].Size();
+        if (size == 0) {
+            LOG_WARNING(Core, "BufferDescriptorB index {} has zero size", buffer_index);
+        }
+        return size;
     } else {
         ASSERT_OR_EXECUTE_MSG(
             BufferDescriptorC().size() > buffer_index, { return 0; },
             "BufferDescriptorC invalid buffer_index {}", buffer_index);
-        return BufferDescriptorC()[buffer_index].Size();
+        const auto size = BufferDescriptorC()[buffer_index].Size();
+        if (size == 0) {
+            LOG_WARNING(Core, "BufferDescriptorC index {} has zero size", buffer_index);
+        }
+        return size;
     }
     return 0;
 }

src/core/hle/service/hle_ipc.h

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2018 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
@@ -422,6 +425,7 @@ private:
     u32 data_payload_offset{};
     u32 handles_offset{};
     u32 domain_offset{};
+    u32 buffer_c_offset{};
 
     std::weak_ptr<SessionRequestManager> manager{};
     bool is_deferred{false};

src/shader_recompiler/backend/spirv/emit_spirv.cpp

@@ -11,6 +11,7 @@
 #include <vector>
 #include <spirv-tools/optimizer.hpp>
 
+#include "common/logging/log.h"
 #include "common/settings.h"
 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
@@ -439,15 +440,23 @@ void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ct
         ctx.AddExtension("SPV_KHR_shader_draw_parameters");
         ctx.AddCapability(spv::Capability::DrawParameters);
     }
-    if ((info.uses_subgroup_vote || info.uses_subgroup_invocation_id ||
+    const bool stage_supports_warp = profile.SupportsWarpIntrinsics(ctx.stage);
-         info.uses_subgroup_shuffles) &&
+    const bool needs_warp_intrinsics = info.uses_subgroup_vote ||
-        profile.support_vote) {
+                                       info.uses_subgroup_invocation_id ||
+                                       info.uses_subgroup_shuffles;
+
+    if (needs_warp_intrinsics && profile.support_vote && stage_supports_warp) {
         ctx.AddCapability(spv::Capability::GroupNonUniformBallot);
         ctx.AddCapability(spv::Capability::GroupNonUniformShuffle);
         if (!profile.warp_size_potentially_larger_than_guest) {
             // vote ops are only used when not taking the long path
             ctx.AddCapability(spv::Capability::GroupNonUniformVote);
         }
+    } else if (needs_warp_intrinsics && !stage_supports_warp) {
+        LOG_WARNING(Shader,
+                    "Warp intrinsics requested in stage {} but the device does not report subgroup "
+                    "support; falling back to scalar approximations",
+                    static_cast<u32>(ctx.stage));
     }
     if (info.uses_int64_bit_atomics && profile.support_int64_atomics) {
         ctx.AddCapability(spv::Capability::Int64Atomics);

src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp

@@ -491,9 +491,24 @@ void EmitSetPatch(EmitContext& ctx, IR::Patch patch, Id value) {
 }
 
 void EmitSetFragColor(EmitContext& ctx, u32 index, u32 component, Id value) {
+    const AttributeType output_type{ctx.runtime_info.color_output_types[index]};
+    Id pointer_type{ctx.output_f32};
+    Id store_value{value};
+    switch (output_type) {
+    case AttributeType::SignedInt:
+        pointer_type = ctx.output_s32;
+        store_value = ctx.OpBitcast(ctx.S32[1], value);
+        break;
+    case AttributeType::UnsignedInt:
+        pointer_type = ctx.output_u32;
+        store_value = ctx.OpBitcast(ctx.U32[1], value);
+        break;
+    default:
+        break;
+    }
     const Id component_id{ctx.Const(component)};
-    const Id pointer{ctx.OpAccessChain(ctx.output_f32, ctx.frag_color.at(index), component_id)};
+    const Id pointer{ctx.OpAccessChain(pointer_type, ctx.frag_color.at(index), component_id)};
-    ctx.OpStore(pointer, value);
+    ctx.OpStore(pointer, store_value);
 }
 
 void EmitSetSampleMask(EmitContext& ctx, Id value) {

src/shader_recompiler/backend/spirv/emit_spirv_image.cpp

@@ -195,6 +195,41 @@ Id Texture(EmitContext& ctx, IR::TextureInstInfo info, [[maybe_unused]] const IR
     }
 }
 
+Id TextureColorResultType(EmitContext& ctx, const TextureDefinition& def) {
+    switch (def.component_type) {
+    case SamplerComponentType::Float:
+    case SamplerComponentType::Depth:
+        return ctx.F32[4];
+    case SamplerComponentType::Sint:
+        return ctx.S32[4];
+    case SamplerComponentType::Stencil:
+        return ctx.U32[4];
+    case SamplerComponentType::Uint:
+        return ctx.U32[4];
+    }
+    throw InvalidArgument("Invalid sampler component type {}", def.component_type);
+}
+
+Id TextureSampleResultToFloat(EmitContext& ctx, const TextureDefinition& def, Id color) {
+    switch (def.component_type) {
+    case SamplerComponentType::Float:
+    case SamplerComponentType::Depth:
+        return color;
+    case SamplerComponentType::Sint:
+        return ctx.OpConvertSToF(ctx.F32[4], color);
+    case SamplerComponentType::Stencil:
+        {
+            const Id converted{ctx.OpConvertUToF(ctx.F32[4], color)};
+            const Id inv255{ctx.Const(1.0f / 255.0f)};
+            const Id scale{ctx.ConstantComposite(ctx.F32[4], inv255, inv255, inv255, inv255)};
+            return ctx.OpFMul(ctx.F32[4], converted, scale);
+        }
+    case SamplerComponentType::Uint:
+        return ctx.OpConvertUToF(ctx.F32[4], color);
+    }
+    throw InvalidArgument("Invalid sampler component type {}", def.component_type);
+}
+
 Id TextureImage(EmitContext& ctx, IR::TextureInstInfo info, const IR::Value& index) {
     if (!index.IsImmediate() || index.U32() != 0) {
         throw NotImplementedException("Indirect image indexing");
@@ -449,31 +484,39 @@ Id EmitBoundImageWrite(EmitContext&) {
 Id EmitImageSampleImplicitLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords,
                               Id bias_lc, const IR::Value& offset) {
     const auto info{inst->Flags<IR::TextureInstInfo>()};
+    const TextureDefinition& def{ctx.textures.at(info.descriptor_index)};
+    const Id color_type{TextureColorResultType(ctx, def)};
+    const Id texture{Texture(ctx, info, index)};
+    Id color{};
     if (ctx.stage == Stage::Fragment) {
         const ImageOperands operands(ctx, info.has_bias != 0, false, info.has_lod_clamp != 0,
                                      bias_lc, offset);
-        return Emit(&EmitContext::OpImageSparseSampleImplicitLod,
+        color = Emit(&EmitContext::OpImageSparseSampleImplicitLod,
-                    &EmitContext::OpImageSampleImplicitLod, ctx, inst, ctx.F32[4],
+                     &EmitContext::OpImageSampleImplicitLod, ctx, inst, color_type, texture,
-                    Texture(ctx, info, index), coords, operands.MaskOptional(), operands.Span());
+                     coords, operands.MaskOptional(), operands.Span());
     } else {
         // We can't use implicit lods on non-fragment stages on SPIR-V. Maxwell hardware behaves as
         // if the lod was explicitly zero.  This may change on Turing with implicit compute
         // derivatives
         const Id lod{ctx.Const(0.0f)};
         const ImageOperands operands(ctx, false, true, info.has_lod_clamp != 0, lod, offset);
-        return Emit(&EmitContext::OpImageSparseSampleExplicitLod,
+        color = Emit(&EmitContext::OpImageSparseSampleExplicitLod,
-                    &EmitContext::OpImageSampleExplicitLod, ctx, inst, ctx.F32[4],
+                     &EmitContext::OpImageSampleExplicitLod, ctx, inst, color_type, texture,
-                    Texture(ctx, info, index), coords, operands.Mask(), operands.Span());
+                     coords, operands.Mask(), operands.Span());
     }
+    return TextureSampleResultToFloat(ctx, def, color);
 }
 
 Id EmitImageSampleExplicitLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords,
                               Id lod, const IR::Value& offset) {
     const auto info{inst->Flags<IR::TextureInstInfo>()};
+    const TextureDefinition& def{ctx.textures.at(info.descriptor_index)};
+    const Id color_type{TextureColorResultType(ctx, def)};
     const ImageOperands operands(ctx, false, true, false, lod, offset);
-    return Emit(&EmitContext::OpImageSparseSampleExplicitLod,
+    const Id color{Emit(&EmitContext::OpImageSparseSampleExplicitLod,
-                &EmitContext::OpImageSampleExplicitLod, ctx, inst, ctx.F32[4],
+                        &EmitContext::OpImageSampleExplicitLod, ctx, inst, color_type,
-                Texture(ctx, info, index), coords, operands.Mask(), operands.Span());
+                        Texture(ctx, info, index), coords, operands.Mask(), operands.Span())};
+    return TextureSampleResultToFloat(ctx, def, color);
 }
 
 Id EmitImageSampleDrefImplicitLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index,
@@ -509,13 +552,18 @@ Id EmitImageSampleDrefExplicitLod(EmitContext& ctx, IR::Inst* inst, const IR::Va
 Id EmitImageGather(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords,
                    const IR::Value& offset, const IR::Value& offset2) {
     const auto info{inst->Flags<IR::TextureInstInfo>()};
+    const TextureDefinition& def{ctx.textures.at(info.descriptor_index)};
+    const Id color_type{TextureColorResultType(ctx, def)};
     const ImageOperands operands(ctx, offset, offset2);
+    const Id texture{Texture(ctx, info, index)};
     if (ctx.profile.need_gather_subpixel_offset) {
         coords = ImageGatherSubpixelOffset(ctx, info, TextureImage(ctx, info, index), coords);
     }
-    return Emit(&EmitContext::OpImageSparseGather, &EmitContext::OpImageGather, ctx, inst,
+    const Id color{
-                ctx.F32[4], Texture(ctx, info, index), coords, ctx.Const(info.gather_component),
+        Emit(&EmitContext::OpImageSparseGather, &EmitContext::OpImageGather, ctx, inst, color_type,
-                operands.MaskOptional(), operands.Span());
+             texture, coords, ctx.Const(info.gather_component), operands.MaskOptional(),
+             operands.Span())};
+    return TextureSampleResultToFloat(ctx, def, color);
 }
 
 Id EmitImageGatherDref(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords,
@@ -533,6 +581,9 @@ Id EmitImageGatherDref(EmitContext& ctx, IR::Inst* inst, const IR::Value& index,
 Id EmitImageFetch(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords, Id offset,
                   Id lod, Id ms) {
     const auto info{inst->Flags<IR::TextureInstInfo>()};
+    const TextureDefinition* def =
+        info.type == TextureType::Buffer ? nullptr : &ctx.textures.at(info.descriptor_index);
+    const Id result_type{def ? TextureColorResultType(ctx, *def) : ctx.F32[4]};
     AddOffsetToCoordinates(ctx, info, coords, offset);
     if (info.type == TextureType::Buffer) {
         lod = Id{};
@@ -542,8 +593,13 @@ Id EmitImageFetch(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id c
         lod = Id{};
     }
     const ImageOperands operands(lod, ms);
-    return Emit(&EmitContext::OpImageSparseFetch, &EmitContext::OpImageFetch, ctx, inst, ctx.F32[4],
+    Id color{Emit(&EmitContext::OpImageSparseFetch, &EmitContext::OpImageFetch, ctx, inst,
-                TextureImage(ctx, info, index), coords, operands.MaskOptional(), operands.Span());
+                  result_type, TextureImage(ctx, info, index), coords, operands.MaskOptional(),
+                  operands.Span())};
+    if (def) {
+        color = TextureSampleResultToFloat(ctx, *def, color);
+    }
+    return color;
 }
 
 Id EmitImageQueryDimensions(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id lod,
@@ -588,14 +644,17 @@ Id EmitImageQueryLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, I
 Id EmitImageGradient(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords,
                      Id derivatives, const IR::Value& offset, Id lod_clamp) {
     const auto info{inst->Flags<IR::TextureInstInfo>()};
+    const TextureDefinition& def{ctx.textures.at(info.descriptor_index)};
+    const Id color_type{TextureColorResultType(ctx, def)};
     const auto operands = info.num_derivatives == 3
                               ? ImageOperands(ctx, info.has_lod_clamp != 0, derivatives,
                                               ctx.Def(offset), {}, lod_clamp)
                               : ImageOperands(ctx, info.has_lod_clamp != 0, derivatives,
                                               info.num_derivatives, offset, lod_clamp);
-    return Emit(&EmitContext::OpImageSparseSampleExplicitLod,
+    const Id color{Emit(&EmitContext::OpImageSparseSampleExplicitLod,
-                &EmitContext::OpImageSampleExplicitLod, ctx, inst, ctx.F32[4],
+                        &EmitContext::OpImageSampleExplicitLod, ctx, inst, color_type,
-                Texture(ctx, info, index), coords, operands.Mask(), operands.Span());
+                        Texture(ctx, info, index), coords, operands.Mask(), operands.Span())};
+    return TextureSampleResultToFloat(ctx, def, color);
 }
 
 Id EmitImageRead(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords) {

src/shader_recompiler/backend/spirv/emit_spirv_warp.cpp

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
@@ -78,9 +81,25 @@ Id AddPartitionBase(EmitContext& ctx, Id thread_id) {
     const Id partition_base{ctx.OpShiftLeftLogical(ctx.U32[1], partition_idx, ctx.Const(5u))};
     return ctx.OpIAdd(ctx.U32[1], thread_id, partition_base);
 }
+
+bool SupportsWarpIntrinsics(const EmitContext& ctx) {
+    return ctx.profile.SupportsWarpIntrinsics(ctx.stage);
+}
+
+void SetAlwaysInBounds(EmitContext& ctx, IR::Inst* inst) {
+    SetInBoundsFlag(inst, ctx.true_value);
+}
+
+Id FallbackBallotMask(EmitContext& ctx, Id pred) {
+    const Id full_mask{ctx.Const(0xFFFFFFFFu)};
+    return ctx.OpSelect(ctx.U32[1], pred, full_mask, ctx.u32_zero_value);
+}
 } // Anonymous namespace
 
 Id EmitLaneId(EmitContext& ctx) {
+    if (!SupportsWarpIntrinsics(ctx)) {
+        return ctx.u32_zero_value;
+    }
     const Id id{GetThreadId(ctx)};
     if (!ctx.profile.warp_size_potentially_larger_than_guest) {
         return id;
@@ -89,6 +108,9 @@ Id EmitLaneId(EmitContext& ctx) {
 }
 
 Id EmitVoteAll(EmitContext& ctx, Id pred) {
+    if (!SupportsWarpIntrinsics(ctx)) {
+        return pred;
+    }
     if (!ctx.profile.warp_size_potentially_larger_than_guest) {
         return ctx.OpGroupNonUniformAll(ctx.U1, SubgroupScope(ctx), pred);
     }
@@ -102,6 +124,9 @@ Id EmitVoteAll(EmitContext& ctx, Id pred) {
 }
 
 Id EmitVoteAny(EmitContext& ctx, Id pred) {
+    if (!SupportsWarpIntrinsics(ctx)) {
+        return pred;
+    }
     if (!ctx.profile.warp_size_potentially_larger_than_guest) {
         return ctx.OpGroupNonUniformAny(ctx.U1, SubgroupScope(ctx), pred);
     }
@@ -115,6 +140,9 @@ Id EmitVoteAny(EmitContext& ctx, Id pred) {
 }
 
 Id EmitVoteEqual(EmitContext& ctx, Id pred) {
+    if (!SupportsWarpIntrinsics(ctx)) {
+        return pred;
+    }
     if (!ctx.profile.warp_size_potentially_larger_than_guest) {
         return ctx.OpGroupNonUniformAllEqual(ctx.U1, SubgroupScope(ctx), pred);
     }
@@ -129,6 +157,9 @@ Id EmitVoteEqual(EmitContext& ctx, Id pred) {
 }
 
 Id EmitSubgroupBallot(EmitContext& ctx, Id pred) {
+    if (!SupportsWarpIntrinsics(ctx)) {
+        return FallbackBallotMask(ctx, pred);
+    }
     const Id ballot{ctx.OpGroupNonUniformBallot(ctx.U32[4], SubgroupScope(ctx), pred)};
     if (!ctx.profile.warp_size_potentially_larger_than_guest) {
         return ctx.OpCompositeExtract(ctx.U32[1], ballot, 0U);
@@ -137,27 +168,46 @@ Id EmitSubgroupBallot(EmitContext& ctx, Id pred) {
 }
 
 Id EmitSubgroupEqMask(EmitContext& ctx) {
+    if (!SupportsWarpIntrinsics(ctx)) {
+        return ctx.u32_zero_value;
+    }
     return LoadMask(ctx, ctx.subgroup_mask_eq);
 }
 
 Id EmitSubgroupLtMask(EmitContext& ctx) {
+    if (!SupportsWarpIntrinsics(ctx)) {
+        return ctx.u32_zero_value;
+    }
     return LoadMask(ctx, ctx.subgroup_mask_lt);
 }
 
 Id EmitSubgroupLeMask(EmitContext& ctx) {
+    if (!SupportsWarpIntrinsics(ctx)) {
+        return ctx.u32_zero_value;
+    }
     return LoadMask(ctx, ctx.subgroup_mask_le);
 }
 
 Id EmitSubgroupGtMask(EmitContext& ctx) {
+    if (!SupportsWarpIntrinsics(ctx)) {
+        return ctx.u32_zero_value;
+    }
     return LoadMask(ctx, ctx.subgroup_mask_gt);
 }
 
 Id EmitSubgroupGeMask(EmitContext& ctx) {
+    if (!SupportsWarpIntrinsics(ctx)) {
+        return ctx.u32_zero_value;
+    }
     return LoadMask(ctx, ctx.subgroup_mask_ge);
 }
 
 Id EmitShuffleIndex(EmitContext& ctx, IR::Inst* inst, Id value, Id index, Id clamp,
                     Id segmentation_mask) {
+    if (!SupportsWarpIntrinsics(ctx)) {
+        SetAlwaysInBounds(ctx, inst);
+        return value;
+    }
     const Id not_seg_mask{ctx.OpNot(ctx.U32[1], segmentation_mask)};
     const Id thread_id{EmitLaneId(ctx)};
     const Id min_thread_id{ComputeMinThreadId(ctx, thread_id, segmentation_mask)};
@@ -177,6 +227,10 @@ Id EmitShuffleIndex(EmitContext& ctx, IR::Inst* inst, Id value, Id index, Id cla
 
 Id EmitShuffleUp(EmitContext& ctx, IR::Inst* inst, Id value, Id index, Id clamp,
                  Id segmentation_mask) {
+    if (!SupportsWarpIntrinsics(ctx)) {
+        SetAlwaysInBounds(ctx, inst);
+        return value;
+    }
     const Id thread_id{EmitLaneId(ctx)};
     const Id max_thread_id{GetMaxThreadId(ctx, thread_id, clamp, segmentation_mask)};
     Id src_thread_id{ctx.OpISub(ctx.U32[1], thread_id, index)};
@@ -192,6 +246,10 @@ Id EmitShuffleUp(EmitContext& ctx, IR::Inst* inst, Id value, Id index, Id clamp,
 
 Id EmitShuffleDown(EmitContext& ctx, IR::Inst* inst, Id value, Id index, Id clamp,
                    Id segmentation_mask) {
+    if (!SupportsWarpIntrinsics(ctx)) {
+        SetAlwaysInBounds(ctx, inst);
+        return value;
+    }
     const Id thread_id{EmitLaneId(ctx)};
     const Id max_thread_id{GetMaxThreadId(ctx, thread_id, clamp, segmentation_mask)};
     Id src_thread_id{ctx.OpIAdd(ctx.U32[1], thread_id, index)};
@@ -207,6 +265,10 @@ Id EmitShuffleDown(EmitContext& ctx, IR::Inst* inst, Id value, Id index, Id clam
 
 Id EmitShuffleButterfly(EmitContext& ctx, IR::Inst* inst, Id value, Id index, Id clamp,
                         Id segmentation_mask) {
+    if (!SupportsWarpIntrinsics(ctx)) {
+        SetAlwaysInBounds(ctx, inst);
+        return value;
+    }
     const Id thread_id{EmitLaneId(ctx)};
     const Id max_thread_id{GetMaxThreadId(ctx, thread_id, clamp, segmentation_mask)};
     Id src_thread_id{ctx.OpBitwiseXor(ctx.U32[1], thread_id, index)};

src/shader_recompiler/backend/spirv/spirv_emit_context.cpp

@@ -28,27 +28,41 @@ enum class Operation {
     FPMax,
 };
 
-Id ImageType(EmitContext& ctx, const TextureDescriptor& desc) {
+Id ComponentScalarType(EmitContext& ctx, SamplerComponentType component_type) {
+    switch (component_type) {
+    case SamplerComponentType::Float:
+    case SamplerComponentType::Depth:
+        return ctx.F32[1];
+    case SamplerComponentType::Sint:
+        return ctx.S32[1];
+    case SamplerComponentType::Stencil:
+        return ctx.U32[1];
+    case SamplerComponentType::Uint:
+        return ctx.U32[1];
+    }
+    throw InvalidArgument("Invalid sampler component type {}", component_type);
+}
+
+Id ImageType(EmitContext& ctx, const TextureDescriptor& desc, Id sampled_type) {
     const spv::ImageFormat format{spv::ImageFormat::Unknown};
-    const Id type{ctx.F32[1]};
     const bool depth{desc.is_depth};
     const bool ms{desc.is_multisample};
     switch (desc.type) {
     case TextureType::Color1D:
-        return ctx.TypeImage(type, spv::Dim::Dim1D, depth, false, false, 1, format);
+        return ctx.TypeImage(sampled_type, spv::Dim::Dim1D, depth, false, false, 1, format);
     case TextureType::ColorArray1D:
-        return ctx.TypeImage(type, spv::Dim::Dim1D, depth, true, false, 1, format);
+        return ctx.TypeImage(sampled_type, spv::Dim::Dim1D, depth, true, false, 1, format);
     case TextureType::Color2D:
     case TextureType::Color2DRect:
-        return ctx.TypeImage(type, spv::Dim::Dim2D, depth, false, ms, 1, format);
+        return ctx.TypeImage(sampled_type, spv::Dim::Dim2D, depth, false, ms, 1, format);
     case TextureType::ColorArray2D:
-        return ctx.TypeImage(type, spv::Dim::Dim2D, depth, true, ms, 1, format);
+        return ctx.TypeImage(sampled_type, spv::Dim::Dim2D, depth, true, ms, 1, format);
     case TextureType::Color3D:
-        return ctx.TypeImage(type, spv::Dim::Dim3D, depth, false, false, 1, format);
+        return ctx.TypeImage(sampled_type, spv::Dim::Dim3D, depth, false, false, 1, format);
     case TextureType::ColorCube:
-        return ctx.TypeImage(type, spv::Dim::Cube, depth, false, false, 1, format);
+        return ctx.TypeImage(sampled_type, spv::Dim::Cube, depth, false, false, 1, format);
     case TextureType::ColorArrayCube:
-        return ctx.TypeImage(type, spv::Dim::Cube, depth, true, false, 1, format);
+        return ctx.TypeImage(sampled_type, spv::Dim::Cube, depth, true, false, 1, format);
     case TextureType::Buffer:
         break;
     }
@@ -549,6 +563,7 @@ void EmitContext::DefineCommonTypes(const Info& info) {
 
     output_f32 = Name(TypePointer(spv::StorageClass::Output, F32[1]), "output_f32");
     output_u32 = Name(TypePointer(spv::StorageClass::Output, U32[1]), "output_u32");
+    output_s32 = Name(TypePointer(spv::StorageClass::Output, S32[1]), "output_s32");
 
     if (info.uses_int8 && profile.support_int8) {
         AddCapability(spv::Capability::Int8);
@@ -1362,7 +1377,8 @@ void EmitContext::DefineImageBuffers(const Info& info, u32& binding) {
 void EmitContext::DefineTextures(const Info& info, u32& binding, u32& scaling_index) {
     textures.reserve(info.texture_descriptors.size());
     for (const TextureDescriptor& desc : info.texture_descriptors) {
-        const Id image_type{ImageType(*this, desc)};
+        const Id result_type{ComponentScalarType(*this, desc.component_type)};
+        const Id image_type{ImageType(*this, desc, result_type)};
         const Id sampled_type{TypeSampledImage(image_type)};
         const Id pointer_type{TypePointer(spv::StorageClass::UniformConstant, sampled_type)};
         const Id desc_type{DescType(*this, sampled_type, pointer_type, desc.count)};
@@ -1375,8 +1391,10 @@ void EmitContext::DefineTextures(const Info& info, u32& binding, u32& scaling_in
             .sampled_type = sampled_type,
             .pointer_type = pointer_type,
             .image_type = image_type,
+            .result_type = result_type,
             .count = desc.count,
             .is_multisample = desc.is_multisample,
+            .component_type = desc.component_type,
         });
         if (profile.supported_spirv >= 0x00010400) {
             interfaces.push_back(id);
@@ -1419,6 +1437,7 @@ void EmitContext::DefineImages(const Info& info, u32& binding, u32& scaling_inde
 void EmitContext::DefineInputs(const IR::Program& program) {
     const Info& info{program.info};
     const VaryingState loads{info.loads.mask | info.passthrough.mask};
+    const bool stage_supports_warp = profile.SupportsWarpIntrinsics(stage);
 
     if (info.uses_workgroup_id) {
         workgroup_id = DefineInput(*this, U32[3], false, spv::BuiltIn::WorkgroupId);
@@ -1442,7 +1461,7 @@ void EmitContext::DefineInputs(const IR::Program& program) {
     if (info.uses_is_helper_invocation) {
         is_helper_invocation = DefineInput(*this, U1, false, spv::BuiltIn::HelperInvocation);
     }
-    if (info.uses_subgroup_mask) {
+    if (info.uses_subgroup_mask && stage_supports_warp) {
         subgroup_mask_eq = DefineInput(*this, U32[4], false, spv::BuiltIn::SubgroupEqMaskKHR);
         subgroup_mask_lt = DefineInput(*this, U32[4], false, spv::BuiltIn::SubgroupLtMaskKHR);
         subgroup_mask_le = DefineInput(*this, U32[4], false, spv::BuiltIn::SubgroupLeMaskKHR);
@@ -1456,9 +1475,10 @@ void EmitContext::DefineInputs(const IR::Program& program) {
             Decorate(subgroup_mask_ge, spv::Decoration::Flat);
         }
     }
-    if (info.uses_fswzadd || info.uses_subgroup_invocation_id || info.uses_subgroup_shuffles ||
+    if (stage_supports_warp &&
-        (profile.warp_size_potentially_larger_than_guest &&
+        (info.uses_fswzadd || info.uses_subgroup_invocation_id || info.uses_subgroup_shuffles ||
-         (info.uses_subgroup_vote || info.uses_subgroup_mask))) {
+         (profile.warp_size_potentially_larger_than_guest &&
+          (info.uses_subgroup_vote || info.uses_subgroup_mask)))) {
         AddCapability(spv::Capability::GroupNonUniform);
         subgroup_local_invocation_id =
             DefineInput(*this, U32[1], false, spv::BuiltIn::SubgroupLocalInvocationId);
@@ -1680,7 +1700,18 @@ void EmitContext::DefineOutputs(const IR::Program& program) {
             if (!info.stores_frag_color[index] && !profile.need_declared_frag_colors) {
                 continue;
             }
-            frag_color[index] = DefineOutput(*this, F32[4], std::nullopt);
+            const AttributeType output_type{runtime_info.color_output_types[index]};
+            const Id vec_type = [&, output_type]() -> Id {
+                switch (output_type) {
+                case AttributeType::SignedInt:
+                    return S32[4];
+                case AttributeType::UnsignedInt:
+                    return U32[4];
+                default:
+                    return F32[4];
+                }
+            }();
+            frag_color[index] = DefineOutput(*this, vec_type, std::nullopt);
             Decorate(frag_color[index], spv::Decoration::Location, index);
             Name(frag_color[index], fmt::format("frag_color{}", index));
         }

src/shader_recompiler/backend/spirv/spirv_emit_context.h

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
@@ -36,8 +39,10 @@ struct TextureDefinition {
     Id sampled_type;
     Id pointer_type;
     Id image_type;
+    Id result_type;
     u32 count;
     bool is_multisample;
+    SamplerComponentType component_type;
 };
 
 struct TextureBufferDefinition {
@@ -244,6 +249,7 @@ public:
 
     Id output_f32{};
     Id output_u32{};
+    Id output_s32{};
 
     Id image_buffer_type{};
     Id image_u32{};

src/shader_recompiler/environment.h

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
@@ -22,6 +25,8 @@ public:
 
     [[nodiscard]] virtual TextureType ReadTextureType(u32 raw_handle) = 0;
 
+    [[nodiscard]] virtual SamplerComponentType ReadTextureComponentType(u32 raw_handle) = 0;
+
     [[nodiscard]] virtual TexturePixelFormat ReadTexturePixelFormat(u32 raw_handle) = 0;
 
     [[nodiscard]] virtual bool IsTexturePixelFormatInteger(u32 raw_handle) = 0;

src/shader_recompiler/ir_opt/texture_pass.cpp

@@ -396,6 +396,10 @@ bool IsTexturePixelFormatInteger(Environment& env, const ConstBufferAddr& cbuf)
     return env.IsTexturePixelFormatInteger(GetTextureHandle(env, cbuf));
 }
 
+SamplerComponentType ReadTextureComponentType(Environment& env, const ConstBufferAddr& cbuf) {
+    return env.ReadTextureComponentType(GetTextureHandle(env, cbuf));
+}
+
 class Descriptors {
 public:
     explicit Descriptors(TextureBufferDescriptors& texture_buffer_descriptors_,
@@ -433,7 +437,9 @@ public:
 
     u32 Add(const TextureDescriptor& desc) {
         const u32 index{Add(texture_descriptors, desc, [&desc](const auto& existing) {
-            return desc.type == existing.type && desc.is_depth == existing.is_depth &&
+            return desc.type == existing.type &&
+                   desc.component_type == existing.component_type &&
+                   desc.is_depth == existing.is_depth &&
                    desc.has_secondary == existing.has_secondary &&
                    desc.cbuf_index == existing.cbuf_index &&
                    desc.cbuf_offset == existing.cbuf_offset &&
@@ -666,10 +672,12 @@ void TexturePass(Environment& env, IR::Program& program, const HostTranslateInfo
                     .secondary_shift_left = cbuf.secondary_shift_left,
                     .count = cbuf.count,
                     .size_shift = DESCRIPTOR_SIZE_SHIFT,
+                    .component_type = ReadTextureComponentType(env, cbuf),
                 });
             } else {
                 index = descriptors.Add(TextureDescriptor{
                     .type = flags.type,
+                    .component_type = ReadTextureComponentType(env, cbuf),
                     .is_depth = flags.is_depth != 0,
                     .is_multisample = is_multisample,
                     .has_secondary = cbuf.has_secondary,

src/shader_recompiler/profile.h

@@ -1,9 +1,15 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
 #pragma once
 
+#include <limits>
+
 #include "common/common_types.h"
+#include "shader_recompiler/stage.h"
 
 namespace Shader {
 
@@ -46,6 +52,8 @@ struct Profile {
     bool support_multi_viewport{};
     bool support_geometry_streams{};
 
+    u32 warp_stage_support_mask{std::numeric_limits<u32>::max()};
+
     bool warp_size_potentially_larger_than_guest{};
 
     bool lower_left_origin_mode{};
@@ -90,6 +98,11 @@ struct Profile {
     u64 min_ssbo_alignment{};
 
     u32 max_user_clip_distances{};
+
+    bool SupportsWarpIntrinsics(Stage stage) const {
+        const u32 bit = 1u << static_cast<u32>(stage);
+        return (warp_stage_support_mask & bit) != 0;
+    }
 };
 
 } // namespace Shader

src/shader_recompiler/runtime_info.h

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
@@ -80,6 +83,7 @@ struct TransformFeedbackVarying {
 
 struct RuntimeInfo {
     std::array<AttributeType, 32> generic_input_types{};
+    std::array<AttributeType, 8> color_output_types{};
     VaryingState previous_stage_stores;
     std::map<IR::Attribute, IR::Attribute> previous_stage_legacy_stores_mapping;
 

src/shader_recompiler/shader_info.h

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
@@ -151,6 +154,14 @@ enum class ImageFormat : u32 {
     R32G32B32A32_UINT,
 };
 
+enum class SamplerComponentType : u8 {
+    Float,
+    Sint,
+    Uint,
+    Depth,
+    Stencil,
+};
+
 enum class Interpolation {
     Smooth,
     Flat,
@@ -183,6 +194,7 @@ struct TextureBufferDescriptor {
     u32 secondary_shift_left;
     u32 count;
     u32 size_shift;
+    SamplerComponentType component_type;
 
     auto operator<=>(const TextureBufferDescriptor&) const = default;
 };
@@ -204,6 +216,7 @@ using ImageBufferDescriptors = boost::container::small_vector<ImageBufferDescrip
 
 struct TextureDescriptor {
     TextureType type;
+    SamplerComponentType component_type;
     bool is_depth;
     bool is_multisample;
     bool has_secondary;

src/video_core/buffer_cache/buffer_cache.h

@@ -407,6 +407,12 @@ void BufferCache<P>::SetComputeUniformBufferState(u32 mask,
 
 template <class P>
 void BufferCache<P>::UnbindGraphicsStorageBuffers(size_t stage) {
+    if constexpr (requires { runtime.ShouldLimitDynamicStorageBuffers(); }) {
+        if (runtime.ShouldLimitDynamicStorageBuffers()) {
+            channel_state->total_graphics_storage_buffers -=
+                static_cast<u32>(std::popcount(channel_state->enabled_storage_buffers[stage]));
+        }
+    }
     channel_state->enabled_storage_buffers[stage] = 0;
     channel_state->written_storage_buffers[stage] = 0;
 }
@@ -414,8 +420,26 @@ void BufferCache<P>::UnbindGraphicsStorageBuffers(size_t stage) {
 template <class P>
 bool BufferCache<P>::BindGraphicsStorageBuffer(size_t stage, size_t ssbo_index, u32 cbuf_index,
                                                u32 cbuf_offset, bool is_written) {
+    const bool already_enabled =
+        ((channel_state->enabled_storage_buffers[stage] >> ssbo_index) & 1U) != 0;
+    if constexpr (requires { runtime.ShouldLimitDynamicStorageBuffers(); }) {
+        if (runtime.ShouldLimitDynamicStorageBuffers() && !already_enabled) {
+            const u32 max_bindings = runtime.GetMaxDynamicStorageBuffers();
+            if (channel_state->total_graphics_storage_buffers >= max_bindings) {
+                LOG_WARNING(HW_GPU,
+                            "Skipping graphics storage buffer {} due to driver limit {}",
+                            ssbo_index, max_bindings);
+                return false;
+            }
+        }
+    }
     channel_state->enabled_storage_buffers[stage] |= 1U << ssbo_index;
     channel_state->written_storage_buffers[stage] |= (is_written ? 1U : 0U) << ssbo_index;
+    if constexpr (requires { runtime.ShouldLimitDynamicStorageBuffers(); }) {
+        if (runtime.ShouldLimitDynamicStorageBuffers() && !already_enabled) {
+            ++channel_state->total_graphics_storage_buffers;
+        }
+    }
 
     const auto& cbufs = maxwell3d->state.shader_stages[stage];
     const GPUVAddr ssbo_addr = cbufs.const_buffers[cbuf_index].address + cbuf_offset;
@@ -446,6 +470,12 @@ void BufferCache<P>::BindGraphicsTextureBuffer(size_t stage, size_t tbo_index, G
 
 template <class P>
 void BufferCache<P>::UnbindComputeStorageBuffers() {
+    if constexpr (requires { runtime.ShouldLimitDynamicStorageBuffers(); }) {
+        if (runtime.ShouldLimitDynamicStorageBuffers()) {
+            channel_state->total_compute_storage_buffers -=
+                static_cast<u32>(std::popcount(channel_state->enabled_compute_storage_buffers));
+        }
+    }
     channel_state->enabled_compute_storage_buffers = 0;
     channel_state->written_compute_storage_buffers = 0;
     channel_state->image_compute_texture_buffers = 0;
@@ -459,8 +489,26 @@ void BufferCache<P>::BindComputeStorageBuffer(size_t ssbo_index, u32 cbuf_index,
                   ssbo_index);
         return;
     }
+    const bool already_enabled =
+        ((channel_state->enabled_compute_storage_buffers >> ssbo_index) & 1U) != 0;
+    if constexpr (requires { runtime.ShouldLimitDynamicStorageBuffers(); }) {
+        if (runtime.ShouldLimitDynamicStorageBuffers() && !already_enabled) {
+            const u32 max_bindings = runtime.GetMaxDynamicStorageBuffers();
+            if (channel_state->total_compute_storage_buffers >= max_bindings) {
+                LOG_WARNING(HW_GPU,
+                            "Skipping compute storage buffer {} due to driver limit {}",
+                            ssbo_index, max_bindings);
+                return;
+            }
+        }
+    }
     channel_state->enabled_compute_storage_buffers |= 1U << ssbo_index;
     channel_state->written_compute_storage_buffers |= (is_written ? 1U : 0U) << ssbo_index;
+    if constexpr (requires { runtime.ShouldLimitDynamicStorageBuffers(); }) {
+        if (runtime.ShouldLimitDynamicStorageBuffers() && !already_enabled) {
+            ++channel_state->total_compute_storage_buffers;
+        }
+    }
 
     const auto& launch_desc = kepler_compute->launch_description;
     if (((launch_desc.const_buffer_enable_mask >> cbuf_index) & 1) == 0) {
@@ -793,9 +841,23 @@ void BufferCache<P>::BindHostGraphicsUniformBuffer(size_t stage, u32 index, u32
     const u32 size = (std::min)(binding.size, (*channel_state->uniform_buffer_sizes)[stage][index]);
     Buffer& buffer = slot_buffers[binding.buffer_id];
     TouchBuffer(buffer, binding.buffer_id);
-    const bool use_fast_buffer = binding.buffer_id != NULL_BUFFER_ID &&
+    const bool has_host_buffer = binding.buffer_id != NULL_BUFFER_ID;
-                                 size <= channel_state->uniform_buffer_skip_cache_size &&
+    const u32 offset = has_host_buffer ? buffer.Offset(device_addr) : 0;
-                                 !memory_tracker.IsRegionGpuModified(device_addr, size);
+    const bool needs_alignment_stream = [&]() {
+        if constexpr (IS_OPENGL) {
+            return false;
+        } else {
+            if (!has_host_buffer) {
+                return false;
+            }
+            const u32 alignment = runtime.GetUniformBufferAlignment();
+            return alignment > 1 && (offset % alignment) != 0;
+        }
+    }();
+    const bool use_fast_buffer = needs_alignment_stream ||
+                                 (has_host_buffer &&
+                                  size <= channel_state->uniform_buffer_skip_cache_size &&
+                                  !memory_tracker.IsRegionGpuModified(device_addr, size));
     if (use_fast_buffer) {
         if constexpr (IS_OPENGL) {
             if (runtime.HasFastBufferSubData()) {
@@ -834,7 +896,6 @@ void BufferCache<P>::BindHostGraphicsUniformBuffer(size_t stage, u32 index, u32
     if (!needs_bind) {
         return;
     }
-    const u32 offset = buffer.Offset(device_addr);
     if constexpr (IS_OPENGL) {
         // Mark the index as dirty if offset doesn't match
         const bool is_copy_bind = offset != 0 && !runtime.SupportsNonZeroUniformOffset();
@@ -951,9 +1012,30 @@ void BufferCache<P>::BindHostComputeUniformBuffers() {
         TouchBuffer(buffer, binding.buffer_id);
         const u32 size =
             (std::min)(binding.size, (*channel_state->compute_uniform_buffer_sizes)[index]);
+        const bool has_host_buffer = binding.buffer_id != NULL_BUFFER_ID;
+        const u32 offset = has_host_buffer ? buffer.Offset(binding.device_addr) : 0;
+        const bool needs_alignment_stream = [&]() {
+            if constexpr (IS_OPENGL) {
+                return false;
+            } else {
+                if (!has_host_buffer) {
+                    return false;
+                }
+                const u32 alignment = runtime.GetUniformBufferAlignment();
+                return alignment > 1 && (offset % alignment) != 0;
+            }
+        }();
+        if constexpr (!IS_OPENGL) {
+            if (needs_alignment_stream) {
+                const std::span<u8> span =
+                    runtime.BindMappedUniformBuffer(0, binding_index, size);
+                device_memory.ReadBlockUnsafe(binding.device_addr, span.data(), size);
+                return;
+            }
+        }
+
         SynchronizeBuffer(buffer, binding.device_addr, size);
 
-        const u32 offset = buffer.Offset(binding.device_addr);
         buffer.MarkUsage(offset, size);
         if constexpr (NEEDS_BIND_UNIFORM_INDEX) {
             runtime.BindComputeUniformBuffer(binding_index, buffer, offset, size);

src/video_core/buffer_cache/buffer_cache_base.h

@@ -8,6 +8,7 @@
 
 #include <algorithm>
 #include <array>
+#include <bit>
 #include <functional>
 #include <memory>
 #include <mutex>
@@ -132,6 +133,9 @@ public:
     u32 enabled_compute_storage_buffers = 0;
     u32 written_compute_storage_buffers = 0;
 
+    u32 total_graphics_storage_buffers = 0;
+    u32 total_compute_storage_buffers = 0;
+
     std::array<u32, NUM_STAGES> enabled_texture_buffers{};
     std::array<u32, NUM_STAGES> written_texture_buffers{};
     std::array<u32, NUM_STAGES> image_texture_buffers{};

src/video_core/compatible_formats.cpp

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2020 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
@@ -235,6 +238,9 @@ constexpr Table MakeViewTable() {
     EnableRange(view, VIEW_CLASS_ASTC_10x10_RGBA);
     EnableRange(view, VIEW_CLASS_ASTC_12x10_RGBA);
     EnableRange(view, VIEW_CLASS_ASTC_12x12_RGBA);
+    Enable(view, PixelFormat::D24_UNORM_S8_UINT, PixelFormat::S8_UINT);
+    Enable(view, PixelFormat::S8_UINT_D24_UNORM, PixelFormat::S8_UINT);
+    Enable(view, PixelFormat::D32_FLOAT_S8_UINT, PixelFormat::S8_UINT);
     return view;
 }
 

src/video_core/fence_manager.h

@@ -87,6 +87,10 @@ public:
         }
         pending_operations.emplace_back(std::move(uncommitted_operations));
         QueueFence(new_fence);
+        if (!new_fence->IsStubbed()) {
+            std::scoped_lock lock{texture_cache.mutex};
+            texture_cache.CommitPendingGpuAccesses(new_fence->WaitTick());
+        }
         if (!delay_fence) {
             func();
         }
@@ -178,7 +182,7 @@ private:
                     return;
                 }
             }
-            PopAsyncFlushes();
+            PopAsyncFlushes(current_fence->WaitTick());
             auto operations = std::move(pending_operations.front());
             pending_operations.pop_front();
             for (auto& operation : operations) {
@@ -213,7 +217,7 @@ private:
             if (!current_fence->IsStubbed()) {
                 WaitFence(current_fence);
             }
-            PopAsyncFlushes();
+            PopAsyncFlushes(current_fence->WaitTick());
             for (auto& operation : current_operations) {
                 operation();
             }
@@ -236,10 +240,11 @@ private:
                query_cache.HasUncommittedFlushes();
     }
 
-    void PopAsyncFlushes() {
+    void PopAsyncFlushes(u64 completed_tick) {
         {
             std::scoped_lock lock{buffer_cache.mutex, texture_cache.mutex};
             texture_cache.PopAsyncFlushes();
+            texture_cache.CompleteGpuAccesses(completed_tick);
             buffer_cache.PopAsyncFlushes();
         }
         query_cache.PopAsyncFlushes();

src/video_core/query_cache/query_cache.h

@@ -211,6 +211,12 @@ void QueryCacheBase<Traits>::CounterClose(QueryType counter_type) {
     streamer->CloseCounter();
 }
 
+template <typename Traits>
+bool QueryCacheBase<Traits>::HasStreamer(QueryType counter_type) const {
+    const size_t index = static_cast<size_t>(counter_type);
+    return impl->streamers[index] != nullptr;
+}
+
 template <typename Traits>
 void QueryCacheBase<Traits>::CounterReset(QueryType counter_type) {
     size_t index = static_cast<size_t>(counter_type);

src/video_core/query_cache/query_cache_base.h

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2023 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-3.0-or-later
 
@@ -92,6 +95,8 @@ public:
 
     void CounterReset(QueryType counter_type);
 
+    [[nodiscard]] bool HasStreamer(QueryType counter_type) const;
+
     void CounterClose(QueryType counter_type);
 
     void CounterReport(GPUVAddr addr, QueryType counter_type, QueryPropertiesFlags flags,

src/video_core/renderer_opengl/gl_buffer_cache.h

@@ -198,6 +198,10 @@ public:
         return device.CanReportMemoryUsage();
     }
 
+    u32 GetUniformBufferAlignment() const {
+        return static_cast<u32>(device.GetUniformBufferAlignment());
+    }
+
     u32 GetStorageBufferAlignment() const {
         return static_cast<u32>(device.GetShaderStorageBufferAlignment());
     }

src/video_core/renderer_opengl/gl_compute_pipeline.cpp

@@ -7,13 +7,50 @@
 #include <cstring>
 #include <bit>
 #include <numeric>
+#include <optional>
 #include "common/cityhash.h"
 #include "common/settings.h" // for enum class Settings::ShaderBackend
 #include "video_core/renderer_opengl/gl_compute_pipeline.h"
 #include "video_core/renderer_opengl/gl_shader_manager.h"
 #include "video_core/renderer_opengl/gl_shader_util.h"
+#include "video_core/surface.h"
 
 namespace OpenGL {
+namespace {
+
+std::optional<VideoCore::Surface::PixelFormatNumeric>
+NumericFromComponentType(Shader::SamplerComponentType component_type) {
+    using VideoCore::Surface::PixelFormatNumeric;
+    switch (component_type) {
+    case Shader::SamplerComponentType::Float:
+        return PixelFormatNumeric::Float;
+    case Shader::SamplerComponentType::Sint:
+        return PixelFormatNumeric::Sint;
+    case Shader::SamplerComponentType::Uint:
+        return PixelFormatNumeric::Uint;
+    default:
+        return std::nullopt;
+    }
+}
+
+VideoCore::Surface::PixelFormat ResolveTexelBufferFormat(
+    VideoCore::Surface::PixelFormat format, Shader::SamplerComponentType component_type) {
+    const auto desired_numeric = NumericFromComponentType(component_type);
+    if (!desired_numeric) {
+        return format;
+    }
+    const auto current_numeric = VideoCore::Surface::GetPixelFormatNumericType(format);
+    if (*desired_numeric == current_numeric) {
+        return format;
+    }
+    if (const auto variant =
+            VideoCore::Surface::FindPixelFormatVariant(format, *desired_numeric)) {
+        return *variant;
+    }
+    return format;
+}
+
+} // Anonymous namespace
 
 using Shader::ImageBufferDescriptor;
 using Tegra::Texture::TexturePair;
@@ -174,8 +211,12 @@ void ComputePipeline::Configure() {
                 is_written = desc.is_written;
             }
             ImageView& image_view{texture_cache.GetImageView(views[texbuf_index].id)};
+            auto buffer_format = image_view.format;
+            if constexpr (!is_image) {
+                buffer_format = ResolveTexelBufferFormat(buffer_format, desc.component_type);
+            }
             buffer_cache.BindComputeTextureBuffer(texbuf_index, image_view.GpuAddr(),
-                                                  image_view.BufferSize(), image_view.format,
+                                                  image_view.BufferSize(), buffer_format,
                                                   is_written, is_image);
             ++texbuf_index;
         }
@@ -205,7 +246,8 @@ void ComputePipeline::Configure() {
     for (const auto& desc : info.texture_descriptors) {
         for (u32 index = 0; index < desc.count; ++index) {
             ImageView& image_view{texture_cache.GetImageView((views_it++)->id)};
-            textures[texture_binding] = image_view.Handle(desc.type);
+            textures[texture_binding] =
+                image_view.SampledView(desc.type, desc.component_type);
             if (texture_cache.IsRescaling(image_view)) {
                 texture_scaling_mask |= 1u << texture_binding;
             }

src/video_core/renderer_opengl/gl_fence_manager.h

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2020 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
@@ -25,6 +28,10 @@ public:
 
     void Wait();
 
+    [[nodiscard]] u64 WaitTick() const noexcept {
+        return 0;
+    }
+
 private:
     OGLSync sync_object;
 };

src/video_core/renderer_opengl/gl_graphics_pipeline.cpp

@@ -6,6 +6,7 @@
 
 #include <algorithm>
 #include <array>
+#include <optional>
 #include <string>
 #include <vector>
 #include <bit>
@@ -18,6 +19,7 @@
 #include "video_core/renderer_opengl/gl_shader_util.h"
 #include "video_core/renderer_opengl/gl_state_tracker.h"
 #include "video_core/shader_notify.h"
+#include "video_core/surface.h"
 #include "video_core/texture_cache/texture_cache.h"
 
 #if defined(_MSC_VER) && defined(NDEBUG)
@@ -39,6 +41,38 @@ using VideoCommon::ImageId;
 constexpr u32 MAX_TEXTURES = 64;
 constexpr u32 MAX_IMAGES = 8;
 
+std::optional<VideoCore::Surface::PixelFormatNumeric>
+NumericFromComponentType(Shader::SamplerComponentType component_type) {
+    using VideoCore::Surface::PixelFormatNumeric;
+    switch (component_type) {
+    case Shader::SamplerComponentType::Float:
+        return PixelFormatNumeric::Float;
+    case Shader::SamplerComponentType::Sint:
+        return PixelFormatNumeric::Sint;
+    case Shader::SamplerComponentType::Uint:
+        return PixelFormatNumeric::Uint;
+    default:
+        return std::nullopt;
+    }
+}
+
+VideoCore::Surface::PixelFormat ResolveTexelBufferFormat(
+    VideoCore::Surface::PixelFormat format, Shader::SamplerComponentType component_type) {
+    const auto desired_numeric = NumericFromComponentType(component_type);
+    if (!desired_numeric) {
+        return format;
+    }
+    const auto current_numeric = VideoCore::Surface::GetPixelFormatNumericType(format);
+    if (*desired_numeric == current_numeric) {
+        return format;
+    }
+    if (const auto variant =
+            VideoCore::Surface::FindPixelFormatVariant(format, *desired_numeric)) {
+        return *variant;
+    }
+    return format;
+}
+
 GLenum Stage(size_t stage_index) {
     switch (stage_index) {
     case 0:
@@ -397,8 +431,12 @@ bool GraphicsPipeline::ConfigureImpl(bool is_indexed) {
                     is_written = desc.is_written;
                 }
                 ImageView& image_view{texture_cache.GetImageView(texture_buffer_it->id)};
+                auto buffer_format = image_view.format;
+                if constexpr (!is_image) {
+                    buffer_format = ResolveTexelBufferFormat(buffer_format, desc.component_type);
+                }
                 buffer_cache.BindGraphicsTextureBuffer(stage, index, image_view.GpuAddr(),
-                                                       image_view.BufferSize(), image_view.format,
+                                                       image_view.BufferSize(), buffer_format,
                                                        is_written, is_image);
                 ++index;
                 ++texture_buffer_it;
@@ -483,7 +521,8 @@ bool GraphicsPipeline::ConfigureImpl(bool is_indexed) {
         for (const auto& desc : info.texture_descriptors) {
             for (u32 index = 0; index < desc.count; ++index) {
                 ImageView& image_view{texture_cache.GetImageView((views_it++)->id)};
-                textures[texture_binding] = image_view.Handle(desc.type);
+                textures[texture_binding] =
+                    image_view.SampledView(desc.type, desc.component_type);
                 if (texture_cache.IsRescaling(image_view)) {
                     texture_scaling_mask |= 1u << stage_texture_binding;
                 }

src/video_core/renderer_opengl/gl_shader_cache.cpp

@@ -54,7 +54,7 @@ using VideoCommon::LoadPipelines;
 using VideoCommon::SerializePipeline;
 using Context = ShaderContext::Context;
 
-constexpr u32 CACHE_VERSION = 13;
+constexpr u32 CACHE_VERSION = 14;
 
 template <typename Container>
 auto MakeSpan(Container& container) {
@@ -220,6 +220,7 @@ ShaderCache::ShaderCache(Tegra::MaxwellDeviceMemoryManager& device_memory_,
           .support_gl_sparse_textures = device.HasSparseTexture2(),
           .support_gl_derivative_control = device.HasDerivativeControl(),
           .support_geometry_streams = true,
+          .warp_stage_support_mask = 0xFFFFFFFFu,
 
           .warp_size_potentially_larger_than_guest = device.IsWarpSizePotentiallyLargerThanGuest(),
 

src/video_core/renderer_opengl/gl_texture_cache.cpp

@@ -692,6 +692,15 @@ bool TextureCacheRuntime::HasNativeASTC() const noexcept {
     return device.HasASTC();
 }
 
+bool TextureCacheRuntime::SupportsLinearFilter(VideoCore::Surface::PixelFormat format) const noexcept {
+    using VideoCore::Surface::GetFormatType;
+    using VideoCore::Surface::IsPixelFormatInteger;
+    if (IsPixelFormatInteger(format)) {
+        return false;
+    }
+    return GetFormatType(format) == VideoCore::Surface::SurfaceType::ColorTexture;
+}
+
 Image::Image(TextureCacheRuntime& runtime_, const VideoCommon::ImageInfo& info_, GPUVAddr gpu_addr_,
              VAddr cpu_addr_)
     : VideoCommon::ImageBase(info_, gpu_addr_, cpu_addr_), runtime{&runtime_} {
@@ -1229,6 +1238,13 @@ GLuint ImageView::StorageView(Shader::TextureType texture_type, Shader::ImageFor
     return view;
 }
 
+GLuint ImageView::SampledView(Shader::TextureType view_type,
+                              Shader::SamplerComponentType /*component_type*/) {
+    // OpenGL swizzles already configure depth/stencil selection per TIC entry,
+    // so fall back to the default view handle.
+    return Handle(view_type);
+}
+
 void ImageView::SetupView(Shader::TextureType view_type) {
     views[static_cast<size_t>(view_type)] = MakeView(view_type, internal_format);
 }

src/video_core/renderer_opengl/gl_texture_cache.h

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2019 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
@@ -13,6 +16,7 @@
 #include "video_core/renderer_opengl/gl_resource_manager.h"
 #include "video_core/renderer_opengl/gl_staging_buffer_pool.h"
 #include "video_core/renderer_opengl/util_shaders.h"
+#include "video_core/surface.h"
 #include "video_core/texture_cache/image_view_base.h"
 #include "video_core/texture_cache/texture_cache_base.h"
 
@@ -129,6 +133,8 @@ public:
         return false;
     }
 
+    bool SupportsLinearFilter(VideoCore::Surface::PixelFormat format) const noexcept;
+
     bool HasBrokenTextureViewFormats() const noexcept {
         return has_broken_texture_view_formats;
     }
@@ -137,6 +143,8 @@ public:
 
     void TickFrame() {}
 
+    void WaitForGpuTick(u64) {}
+
     StateTracker& GetStateTracker() {
         return state_tracker;
     }
@@ -264,6 +272,9 @@ public:
     [[nodiscard]] GLuint StorageView(Shader::TextureType texture_type,
                                      Shader::ImageFormat image_format);
 
+    [[nodiscard]] GLuint SampledView(Shader::TextureType view_type,
+                                     Shader::SamplerComponentType component_type);
+
     [[nodiscard]] GLuint Handle(Shader::TextureType handle_type) const noexcept {
         return views[static_cast<size_t>(handle_type)];
     }

src/video_core/renderer_vulkan/fixed_pipeline_state.h

@@ -27,6 +27,7 @@ struct DynamicFeatures {
     bool has_extended_dynamic_state_2_patch_control_points;
     bool has_extended_dynamic_state_3_blend;
     bool has_extended_dynamic_state_3_enables;
+    bool has_dual_source_blend;
     bool has_dynamic_vertex_input;
 };
 

src/video_core/renderer_vulkan/maxwell_to_vk.cpp

@@ -173,7 +173,7 @@ FormatInfo SurfaceFormat(const Device& device, FormatType format_type, bool with
     SURFACE_FORMAT_ELEM(VK_FORMAT_R16G16_SINT, usage_attachable | usage_storage, R16G16_SINT) \
     SURFACE_FORMAT_ELEM(VK_FORMAT_R16G16_SNORM, usage_attachable | usage_storage, R16G16_SNORM) \
     SURFACE_FORMAT_ELEM(VK_FORMAT_R32G32B32_SFLOAT, 0, R32G32B32_FLOAT) \
-    SURFACE_FORMAT_ELEM(VK_FORMAT_A8B8G8R8_SRGB_PACK32, usage_attachable, A8B8G8R8_SRGB) \
+    SURFACE_FORMAT_ELEM(VK_FORMAT_A8B8G8R8_SRGB_PACK32, usage_attachable | usage_storage, A8B8G8R8_SRGB) \
     SURFACE_FORMAT_ELEM(VK_FORMAT_R8G8_UNORM, usage_attachable | usage_storage, R8G8_UNORM) \
     SURFACE_FORMAT_ELEM(VK_FORMAT_R8G8_SNORM, usage_attachable | usage_storage, R8G8_SNORM) \
     SURFACE_FORMAT_ELEM(VK_FORMAT_R8G8_SINT, usage_attachable | usage_storage, R8G8_SINT) \
@@ -185,7 +185,7 @@ FormatInfo SurfaceFormat(const Device& device, FormatType format_type, bool with
     SURFACE_FORMAT_ELEM(VK_FORMAT_ASTC_8x8_UNORM_BLOCK, 0, ASTC_2D_8X8_UNORM) \
     SURFACE_FORMAT_ELEM(VK_FORMAT_ASTC_8x5_UNORM_BLOCK, 0, ASTC_2D_8X5_UNORM) \
     SURFACE_FORMAT_ELEM(VK_FORMAT_ASTC_5x4_UNORM_BLOCK, 0, ASTC_2D_5X4_UNORM) \
-    SURFACE_FORMAT_ELEM(VK_FORMAT_B8G8R8A8_SRGB, usage_attachable, B8G8R8A8_SRGB) \
+    SURFACE_FORMAT_ELEM(VK_FORMAT_B8G8R8A8_SRGB, usage_attachable | usage_storage, B8G8R8A8_SRGB) \
     SURFACE_FORMAT_ELEM(VK_FORMAT_BC1_RGBA_SRGB_BLOCK, 0, BC1_RGBA_SRGB) \
     SURFACE_FORMAT_ELEM(VK_FORMAT_BC2_SRGB_BLOCK, 0, BC2_SRGB) \
     SURFACE_FORMAT_ELEM(VK_FORMAT_BC3_SRGB_BLOCK, 0, BC3_SRGB) \

src/video_core/renderer_vulkan/pipeline_helper.h

@@ -189,12 +189,16 @@ inline void PushImageDescriptors(TextureCache& texture_cache,
             const VideoCommon::ImageViewId image_view_id{(views++)->id};
             const VideoCommon::SamplerId sampler_id{*(samplers++)};
             ImageView& image_view{texture_cache.GetImageView(image_view_id)};
-            const VkImageView vk_image_view{image_view.Handle(desc.type)};
+            const VkImageView vk_image_view{
+                image_view.SampledView(desc.type, desc.component_type)};
             const Sampler& sampler{texture_cache.GetSampler(sampler_id)};
-            const bool use_fallback_sampler{sampler.HasAddedAnisotropy() &&
+            const bool supports_linear_filter{
-                                            !image_view.SupportsAnisotropy()};
+                texture_cache.SupportsLinearFilter(image_view.format)};
-            const VkSampler vk_sampler{use_fallback_sampler ? sampler.HandleWithDefaultAnisotropy()
+            const bool supports_depth_compare_sampling{
-                                                            : sampler.Handle()};
+                image_view.SupportsDepthCompareSampling()};
+            const VkSampler vk_sampler{
+                sampler.SelectHandle(supports_linear_filter, image_view.SupportsAnisotropy(),
+                                     supports_depth_compare_sampling)};
             guest_descriptor_queue.AddSampledImage(vk_image_view, vk_sampler);
             rescaling.PushTexture(texture_cache.IsRescaling(image_view));
         }

src/video_core/renderer_vulkan/present/layer.cpp

@@ -280,7 +280,6 @@ void Layer::UpdateRawImage(const Tegra::FramebufferConfig& framebuffer, size_t i
         Tegra::Texture::UnswizzleTexture(
             mapped_span.subspan(image_offset, linear_size), std::span(host_ptr, tiled_size),
             bytes_per_pixel, framebuffer.width, framebuffer.height, 1, block_height_log2, 0);
-        buffer.Flush();  // Ensure host writes are visible before the GPU copy.
     }
 
     const VkBufferImageCopy copy{

src/video_core/renderer_vulkan/vk_buffer_cache.cpp

@@ -7,6 +7,7 @@
 #include <algorithm>
 #include <array>
 #include <cstring>
+#include <limits>
 #include <span>
 #include <vector>
 
@@ -333,6 +334,13 @@ BufferCacheRuntime::BufferCacheRuntime(const Device& device_, MemoryAllocator& m
       staging_pool{staging_pool_}, guest_descriptor_queue{guest_descriptor_queue_},
       quad_index_pass(device, scheduler, descriptor_pool, staging_pool,
                       compute_pass_descriptor_queue) {
+    const VkDriverIdKHR driver_id = device.GetDriverID();
+    limit_dynamic_storage_buffers = driver_id == VK_DRIVER_ID_QUALCOMM_PROPRIETARY ||
+                                    driver_id == VK_DRIVER_ID_MESA_TURNIP ||
+                                    driver_id == VK_DRIVER_ID_ARM_PROPRIETARY;
+    if (limit_dynamic_storage_buffers) {
+        max_dynamic_storage_buffers = device.GetMaxDescriptorSetStorageBuffersDynamic();
+    }
     if (device.GetDriverID() != VK_DRIVER_ID_QUALCOMM_PROPRIETARY) {
         // TODO: FixMe: Uint8Pass compute shader does not build on some Qualcomm drivers.
         uint8_pass = std::make_unique<Uint8Pass>(device, scheduler, descriptor_pool, staging_pool,
@@ -368,6 +376,10 @@ bool BufferCacheRuntime::CanReportMemoryUsage() const {
     return device.CanReportMemoryUsage();
 }
 
+u32 BufferCacheRuntime::GetUniformBufferAlignment() const {
+    return static_cast<u32>(device.GetUniformBufferAlignment());
+}
+
 u32 BufferCacheRuntime::GetStorageBufferAlignment() const {
     return static_cast<u32>(device.GetStorageBufferAlignment());
 }
@@ -542,7 +554,15 @@ void BufferCacheRuntime::BindVertexBuffer(u32 index, VkBuffer buffer, u32 offset
     if (index >= device.GetMaxVertexInputBindings()) {
         return;
     }
-    if (device.IsExtExtendedDynamicStateSupported()) {
+    if (!device.HasNullDescriptor() && buffer == VK_NULL_HANDLE) {
+        ReserveNullBuffer();
+        buffer = *null_buffer;
+        offset = 0;
+        size = std::numeric_limits<u32>::max();
+    }
+    // Use BindVertexBuffers2EXT only if EDS1 is supported AND VIDS is not active
+    // When VIDS is active, the pipeline doesn't declare VERTEX_INPUT_BINDING_STRIDE as dynamic
+    if (device.IsExtExtendedDynamicStateSupported() && !device.IsExtVertexInputDynamicStateSupported()) {
         scheduler.Record([index, buffer, offset, size, stride](vk::CommandBuffer cmdbuf) {
             const VkDeviceSize vk_offset = buffer != VK_NULL_HANDLE ? offset : 0;
             const VkDeviceSize vk_size = buffer != VK_NULL_HANDLE ? size : VK_WHOLE_SIZE;
@@ -582,7 +602,8 @@ void BufferCacheRuntime::BindVertexBuffers(VideoCommon::HostBindings<Buffer>& bi
     if (binding_count == 0) {
         return;
     }
-    if (device.IsExtExtendedDynamicStateSupported()) {
+    // Use BindVertexBuffers2EXT only if EDS1 is supported AND VIDS is not active
+    if (device.IsExtExtendedDynamicStateSupported() && !device.IsExtVertexInputDynamicStateSupported()) {
         scheduler.Record([bindings_ = std::move(bindings),
                           buffer_handles_ = std::move(buffer_handles),
                           binding_count](vk::CommandBuffer cmdbuf) {
@@ -639,27 +660,50 @@ void BufferCacheRuntime::BindTransformFeedbackBuffers(VideoCommon::HostBindings<
 }
 
 void BufferCacheRuntime::ReserveNullBuffer() {
+    const VkBufferUsageFlags expected_usage = NullBufferUsageFlags();
+    if (null_buffer && null_buffer_usage_flags != expected_usage) {
+        RefreshNullBuffer();
+    }
     if (!null_buffer) {
         null_buffer = CreateNullBuffer();
     }
 }
 
+VkBufferUsageFlags BufferCacheRuntime::NullBufferUsageFlags() const {
+    VkBufferUsageFlags usage = VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT |
+                              VK_BUFFER_USAGE_UNIFORM_TEXEL_BUFFER_BIT |
+                              VK_BUFFER_USAGE_STORAGE_TEXEL_BUFFER_BIT |
+                              VK_BUFFER_USAGE_UNIFORM_BUFFER_BIT | VK_BUFFER_USAGE_STORAGE_BUFFER_BIT |
+                              VK_BUFFER_USAGE_VERTEX_BUFFER_BIT | VK_BUFFER_USAGE_INDEX_BUFFER_BIT |
+                              VK_BUFFER_USAGE_INDIRECT_BUFFER_BIT;
+    if (device.IsExtTransformFeedbackSupported()) {
+        usage |= VK_BUFFER_USAGE_TRANSFORM_FEEDBACK_BUFFER_BIT_EXT;
+    }
+    return usage;
+}
+
+void BufferCacheRuntime::RefreshNullBuffer() {
+    if (!null_buffer) {
+        return;
+    }
+    scheduler.Finish();
+    null_buffer.reset();
+    null_buffer = CreateNullBuffer();
+}
+
 vk::Buffer BufferCacheRuntime::CreateNullBuffer() {
     VkBufferCreateInfo create_info{
         .sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO,
         .pNext = nullptr,
         .flags = 0,
         .size = 4,
-        .usage = VK_BUFFER_USAGE_VERTEX_BUFFER_BIT | VK_BUFFER_USAGE_INDEX_BUFFER_BIT |
+        .usage = NullBufferUsageFlags(),
-                 VK_BUFFER_USAGE_TRANSFER_DST_BIT | VK_BUFFER_USAGE_INDIRECT_BUFFER_BIT,
         .sharingMode = VK_SHARING_MODE_EXCLUSIVE,
         .queueFamilyIndexCount = 0,
         .pQueueFamilyIndices = nullptr,
     };
-    if (device.IsExtTransformFeedbackSupported()) {
-        create_info.usage |= VK_BUFFER_USAGE_TRANSFORM_FEEDBACK_BUFFER_BIT_EXT;
-    }
     vk::Buffer ret = memory_allocator.CreateBuffer(create_info, MemoryUsage::DeviceLocal);
+    null_buffer_usage_flags = create_info.usage;
     if (device.HasDebuggingToolAttached()) {
         ret.SetObjectNameEXT("Null buffer");
     }

src/video_core/renderer_vulkan/vk_buffer_cache.h

@@ -6,6 +6,8 @@
 
 #pragma once
 
+#include <limits>
+
 #include "video_core/buffer_cache/buffer_cache_base.h"
 #include "video_core/buffer_cache/memory_tracker_base.h"
 #include "video_core/buffer_cache/usage_tracker.h"
@@ -94,6 +96,8 @@ public:
 
     bool CanReportMemoryUsage() const;
 
+    u32 GetUniformBufferAlignment() const;
+
     u32 GetStorageBufferAlignment() const;
 
     [[nodiscard]] StagingBufferRef UploadStagingBuffer(size_t size);
@@ -127,6 +131,9 @@ public:
 
     void BindTransformFeedbackBuffers(VideoCommon::HostBindings<Buffer>& bindings);
 
+    /// Forces destruction and recreation of the shared null buffer so new usage flags take effect.
+    void RefreshNullBuffer();
+
     std::span<u8> BindMappedUniformBuffer([[maybe_unused]] size_t stage,
                                           [[maybe_unused]] u32 binding_index,
                                           u32 size) {
@@ -149,6 +156,14 @@ public:
         guest_descriptor_queue.AddTexelBuffer(buffer.View(offset, size, format));
     }
 
+    bool ShouldLimitDynamicStorageBuffers() const {
+        return limit_dynamic_storage_buffers;
+    }
+
+    u32 GetMaxDynamicStorageBuffers() const {
+        return max_dynamic_storage_buffers;
+    }
+
 private:
     void BindBuffer(VkBuffer buffer, u32 offset, u32 size) {
         guest_descriptor_queue.AddBuffer(buffer, offset, size);
@@ -156,6 +171,7 @@ private:
 
     void ReserveNullBuffer();
     vk::Buffer CreateNullBuffer();
+    VkBufferUsageFlags NullBufferUsageFlags() const;
 
     const Device& device;
     MemoryAllocator& memory_allocator;
@@ -167,9 +183,13 @@ private:
     std::shared_ptr<QuadStripIndexBuffer> quad_strip_index_buffer;
 
     vk::Buffer null_buffer;
+    VkBufferUsageFlags null_buffer_usage_flags = 0;
 
     std::unique_ptr<Uint8Pass> uint8_pass;
     QuadIndexedPass quad_index_pass;
+
+    bool limit_dynamic_storage_buffers = false;
+    u32 max_dynamic_storage_buffers = std::numeric_limits<u32>::max();
 };
 
 struct BufferCacheParams {

src/video_core/renderer_vulkan/vk_compute_pass.cpp

@@ -462,10 +462,14 @@ QueriesPrefixScanPass::QueriesPrefixScanPass(
           device_, descriptor_pool_, QUERIES_SCAN_DESCRIPTOR_SET_BINDINGS,
           QUERIES_SCAN_DESCRIPTOR_UPDATE_TEMPLATE, QUERIES_SCAN_BANK_INFO,
           COMPUTE_PUSH_CONSTANT_RANGE<sizeof(QueriesPrefixScanPushConstants)>,
-          device_.IsSubgroupFeatureSupported(VK_SUBGROUP_FEATURE_BASIC_BIT) &&
+          device_.IsSubgroupFeatureSupported(VK_SUBGROUP_FEATURE_BASIC_BIT,
-                  device_.IsSubgroupFeatureSupported(VK_SUBGROUP_FEATURE_ARITHMETIC_BIT) &&
+                                             VK_SHADER_STAGE_COMPUTE_BIT) &&
-                  device_.IsSubgroupFeatureSupported(VK_SUBGROUP_FEATURE_SHUFFLE_BIT) &&
+                  device_.IsSubgroupFeatureSupported(VK_SUBGROUP_FEATURE_ARITHMETIC_BIT,
-                  device_.IsSubgroupFeatureSupported(VK_SUBGROUP_FEATURE_SHUFFLE_RELATIVE_BIT)
+                                                      VK_SHADER_STAGE_COMPUTE_BIT) &&
+                  device_.IsSubgroupFeatureSupported(VK_SUBGROUP_FEATURE_SHUFFLE_BIT,
+                                                      VK_SHADER_STAGE_COMPUTE_BIT) &&
+                  device_.IsSubgroupFeatureSupported(VK_SUBGROUP_FEATURE_SHUFFLE_RELATIVE_BIT,
+                                                      VK_SHADER_STAGE_COMPUTE_BIT)
               ? std::span<const u32>(QUERIES_PREFIX_SCAN_SUM_COMP_SPV)
               : std::span<const u32>(QUERIES_PREFIX_SCAN_SUM_NOSUBGROUPS_COMP_SPV)),
       scheduler{scheduler_}, compute_pass_descriptor_queue{compute_pass_descriptor_queue_} {}

src/video_core/renderer_vulkan/vk_compute_pipeline.cpp

@@ -18,14 +18,49 @@
 #include "video_core/renderer_vulkan/vk_scheduler.h"
 #include "video_core/renderer_vulkan/vk_update_descriptor.h"
 #include "video_core/shader_notify.h"
+#include "video_core/surface.h"
+#include "video_core/texture_cache/texture_cache.h"
 #include "video_core/vulkan_common/vulkan_device.h"
 #include "video_core/vulkan_common/vulkan_wrapper.h"
+#include <optional>
 
 namespace Vulkan {
 
 using Shader::ImageBufferDescriptor;
 using Shader::Backend::SPIRV::RESCALING_LAYOUT_WORDS_OFFSET;
 using Tegra::Texture::TexturePair;
+using VideoCore::Surface::PixelFormat;
+using VideoCore::Surface::PixelFormatNumeric;
+
+static std::optional<PixelFormatNumeric> NumericFromComponentType(
+    Shader::SamplerComponentType component_type) {
+    switch (component_type) {
+    case Shader::SamplerComponentType::Float:
+        return PixelFormatNumeric::Float;
+    case Shader::SamplerComponentType::Sint:
+        return PixelFormatNumeric::Sint;
+    case Shader::SamplerComponentType::Uint:
+        return PixelFormatNumeric::Uint;
+    default:
+        return std::nullopt;
+    }
+}
+
+static PixelFormat ResolveTexelBufferFormat(PixelFormat format,
+                                     Shader::SamplerComponentType component_type) {
+    const auto desired_numeric = NumericFromComponentType(component_type);
+    if (!desired_numeric) {
+        return format;
+    }
+    const auto current_numeric = VideoCore::Surface::GetPixelFormatNumericType(format);
+    if (*desired_numeric == current_numeric) {
+        return format;
+    }
+    if (const auto variant = VideoCore::Surface::FindPixelFormatVariant(format, *desired_numeric)) {
+        return *variant;
+    }
+    return format;
+}
 
 ComputePipeline::ComputePipeline(const Device& device_, vk::PipelineCache& pipeline_cache_,
                                  DescriptorPool& descriptor_pool,
@@ -182,8 +217,12 @@ void ComputePipeline::Configure(Tegra::Engines::KeplerCompute& kepler_compute,
                 is_written = desc.is_written;
             }
             ImageView& image_view = texture_cache.GetImageView(views[index].id);
+            VideoCore::Surface::PixelFormat buffer_format = image_view.format;
+            if constexpr (!is_image) {
+                buffer_format = ResolveTexelBufferFormat(buffer_format, desc.component_type);
+            }
             buffer_cache.BindComputeTextureBuffer(index, image_view.GpuAddr(),
-                                                  image_view.BufferSize(), image_view.format,
+                                                  image_view.BufferSize(), buffer_format,
                                                   is_written, is_image);
             ++index;
         }

src/video_core/renderer_vulkan/vk_fence_manager.h

@@ -34,6 +34,10 @@ public:
 
     void Wait();
 
+    [[nodiscard]] u64 WaitTick() const noexcept {
+        return wait_tick;
+    }
+
 private:
     Scheduler& scheduler;
     u64 wait_tick = 0;

src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp

@@ -5,6 +5,8 @@
 // SPDX-License-Identifier: GPL-2.0-or-later
 
 #include <algorithm>
+#include <array>
+#include <optional>
 #include <iostream>
 #include <span>
 
@@ -23,7 +25,9 @@
 #include "video_core/renderer_vulkan/vk_texture_cache.h"
 #include "video_core/renderer_vulkan/vk_update_descriptor.h"
 #include "video_core/shader_notify.h"
+#include "video_core/texture_cache/samples_helper.h"
 #include "video_core/texture_cache/texture_cache.h"
+#include "video_core/surface.h"
 #include "video_core/vulkan_common/vulkan_device.h"
 
 #if defined(_MSC_VER) && defined(NDEBUG)
@@ -44,10 +48,83 @@ using Tegra::Texture::TexturePair;
 using VideoCore::Surface::PixelFormat;
 using VideoCore::Surface::PixelFormatFromDepthFormat;
 using VideoCore::Surface::PixelFormatFromRenderTargetFormat;
+using VideoCore::Surface::PixelFormatNumeric;
 
 constexpr size_t NUM_STAGES = Maxwell::MaxShaderStage;
 constexpr size_t MAX_IMAGE_ELEMENTS = 64;
 
+std::optional<PixelFormatNumeric> NumericFromComponentType(
+    Shader::SamplerComponentType component_type) {
+    switch (component_type) {
+    case Shader::SamplerComponentType::Float:
+        return PixelFormatNumeric::Float;
+    case Shader::SamplerComponentType::Sint:
+        return PixelFormatNumeric::Sint;
+    case Shader::SamplerComponentType::Uint:
+        return PixelFormatNumeric::Uint;
+    default:
+        return std::nullopt;
+    }
+}
+
+PixelFormat ResolveTexelBufferFormat(PixelFormat format,
+                                     Shader::SamplerComponentType component_type) {
+    const auto desired_numeric = NumericFromComponentType(component_type);
+    if (!desired_numeric) {
+        return format;
+    }
+    const auto current_numeric = VideoCore::Surface::GetPixelFormatNumericType(format);
+    if (*desired_numeric == current_numeric) {
+        return format;
+    }
+    if (const auto variant = VideoCore::Surface::FindPixelFormatVariant(format, *desired_numeric)) {
+        return *variant;
+    }
+    return format;
+}
+
+bool UsesDualSourceFactor(Maxwell::Blend::Factor factor) {
+    switch (factor) {
+    case Maxwell::Blend::Factor::Source1Color_D3D:
+    case Maxwell::Blend::Factor::Source1Color_GL:
+    case Maxwell::Blend::Factor::OneMinusSource1Color_D3D:
+    case Maxwell::Blend::Factor::OneMinusSource1Color_GL:
+    case Maxwell::Blend::Factor::Source1Alpha_D3D:
+    case Maxwell::Blend::Factor::Source1Alpha_GL:
+    case Maxwell::Blend::Factor::OneMinusSource1Alpha_D3D:
+    case Maxwell::Blend::Factor::OneMinusSource1Alpha_GL:
+        return true;
+    default:
+        return false;
+    }
+}
+
+Maxwell::Blend::Factor FallbackDualSourceFactor(Maxwell::Blend::Factor factor) {
+    switch (factor) {
+    case Maxwell::Blend::Factor::Source1Color_D3D:
+    case Maxwell::Blend::Factor::Source1Color_GL:
+        return Maxwell::Blend::Factor::SourceColor_D3D;
+    case Maxwell::Blend::Factor::OneMinusSource1Color_D3D:
+    case Maxwell::Blend::Factor::OneMinusSource1Color_GL:
+        return Maxwell::Blend::Factor::OneMinusSourceColor_D3D;
+    case Maxwell::Blend::Factor::Source1Alpha_D3D:
+    case Maxwell::Blend::Factor::Source1Alpha_GL:
+        return Maxwell::Blend::Factor::SourceAlpha_D3D;
+    case Maxwell::Blend::Factor::OneMinusSource1Alpha_D3D:
+    case Maxwell::Blend::Factor::OneMinusSource1Alpha_GL:
+        return Maxwell::Blend::Factor::OneMinusSourceAlpha_D3D;
+    default:
+        return factor;
+    }
+}
+
+bool AttachmentUsesDualSource(const FixedPipelineState::BlendingAttachment& blend) {
+    return UsesDualSourceFactor(blend.SourceRGBFactor()) ||
+           UsesDualSourceFactor(blend.DestRGBFactor()) ||
+           UsesDualSourceFactor(blend.SourceAlphaFactor()) ||
+           UsesDualSourceFactor(blend.DestAlphaFactor());
+}
+
 DescriptorLayoutBuilder MakeBuilder(const Device& device, std::span<const Shader::Info> infos) {
     DescriptorLayoutBuilder builder{device};
     for (size_t index = 0; index < infos.size(); ++index) {
@@ -417,8 +494,12 @@ bool GraphicsPipeline::ConfigureImpl(bool is_indexed) {
                     is_written = desc.is_written;
                 }
                 ImageView& image_view{texture_cache.GetImageView(texture_buffer_it->id)};
+                VideoCore::Surface::PixelFormat buffer_format = image_view.format;
+                if constexpr (!is_image) {
+                    buffer_format = ResolveTexelBufferFormat(buffer_format, desc.component_type);
+                }
                 buffer_cache.BindGraphicsTextureBuffer(stage, index, image_view.GpuAddr(),
-                                                       image_view.BufferSize(), image_view.format,
+                                                       image_view.BufferSize(), buffer_format,
                                                        is_written, is_image);
                 ++index;
                 ++texture_buffer_it;
@@ -747,11 +828,13 @@ void GraphicsPipeline::MakePipeline(VkRenderPass render_pass) {
 
     const bool supports_alpha_output = fragment_has_color0_output;
     const bool alpha_to_one_supported = device.SupportsAlphaToOne();
-    const VkPipelineMultisampleStateCreateInfo multisample_ci{
+    const auto msaa_mode = key.state.msaa_mode.Value();
+    const VkSampleCountFlagBits vk_samples = MaxwellToVK::MsaaMode(msaa_mode);
+    VkPipelineMultisampleStateCreateInfo multisample_ci{
         .sType = VK_STRUCTURE_TYPE_PIPELINE_MULTISAMPLE_STATE_CREATE_INFO,
         .pNext = nullptr,
         .flags = 0,
-        .rasterizationSamples = MaxwellToVK::MsaaMode(key.state.msaa_mode),
+        .rasterizationSamples = vk_samples,
         .sampleShadingEnable = Settings::values.sample_shading.GetValue() > 0 ? VK_TRUE : VK_FALSE,
         .minSampleShading = f32(Settings::values.sample_shading.GetValue()) / 100.0f,
         .pSampleMask = nullptr,
@@ -760,6 +843,8 @@ void GraphicsPipeline::MakePipeline(VkRenderPass render_pass) {
         .alphaToOneEnable = supports_alpha_output && alpha_to_one_supported &&
                            key.state.alpha_to_one_enabled != 0 ? VK_TRUE : VK_FALSE,
     };
+
+
     const VkPipelineDepthStencilStateCreateInfo depth_stencil_ci{
         .sType = VK_STRUCTURE_TYPE_PIPELINE_DEPTH_STENCIL_STATE_CREATE_INFO,
         .pNext = nullptr,
@@ -781,6 +866,12 @@ void GraphicsPipeline::MakePipeline(VkRenderPass render_pass) {
     }
     static_vector<VkPipelineColorBlendAttachmentState, Maxwell::NumRenderTargets> cb_attachments;
     const size_t num_attachments{NumAttachments(key.state)};
+    const bool supports_dual_source_blend = device.SupportsDualSourceBlend();
+    const u32 max_dual_source_attachments = supports_dual_source_blend
+                                                ? device.MaxFragmentDualSrcAttachments()
+                                                : 0;
+    u32 granted_dual_source_attachments = 0;
+    bool logged_dual_source_warning = false;
     for (size_t index = 0; index < num_attachments; ++index) {
         static constexpr std::array mask_table{
             VK_COLOR_COMPONENT_R_BIT,
@@ -794,13 +885,30 @@ void GraphicsPipeline::MakePipeline(VkRenderPass render_pass) {
         for (size_t i = 0; i < mask_table.size(); ++i) {
             write_mask |= mask[i] ? mask_table[i] : 0;
         }
+        const bool attachment_uses_dual_source = AttachmentUsesDualSource(blend);
+        const bool allow_dual_source = attachment_uses_dual_source && supports_dual_source_blend &&
+                                       granted_dual_source_attachments < max_dual_source_attachments;
+        if (allow_dual_source) {
+            ++granted_dual_source_attachments;
+        } else if (attachment_uses_dual_source && !logged_dual_source_warning) {
+            LOG_WARNING(Render_Vulkan,
+                        "Dual-source blend factors exceed device limit (maxFragmentDualSrcAttachments={}), falling back to single-source factors",
+                        max_dual_source_attachments);
+            logged_dual_source_warning = true;
+        }
+        const auto sanitize_factor = [&](Maxwell::Blend::Factor factor) {
+            if (allow_dual_source || !UsesDualSourceFactor(factor)) {
+                return factor;
+            }
+            return FallbackDualSourceFactor(factor);
+        };
         cb_attachments.push_back({
             .blendEnable = blend.enable != 0,
-            .srcColorBlendFactor = MaxwellToVK::BlendFactor(blend.SourceRGBFactor()),
+            .srcColorBlendFactor = MaxwellToVK::BlendFactor(sanitize_factor(blend.SourceRGBFactor())),
-            .dstColorBlendFactor = MaxwellToVK::BlendFactor(blend.DestRGBFactor()),
+            .dstColorBlendFactor = MaxwellToVK::BlendFactor(sanitize_factor(blend.DestRGBFactor())),
             .colorBlendOp = MaxwellToVK::BlendEquation(blend.EquationRGB()),
-            .srcAlphaBlendFactor = MaxwellToVK::BlendFactor(blend.SourceAlphaFactor()),
+            .srcAlphaBlendFactor = MaxwellToVK::BlendFactor(sanitize_factor(blend.SourceAlphaFactor())),
-            .dstAlphaBlendFactor = MaxwellToVK::BlendFactor(blend.DestAlphaFactor()),
+            .dstAlphaBlendFactor = MaxwellToVK::BlendFactor(sanitize_factor(blend.DestAlphaFactor())),
             .alphaBlendOp = MaxwellToVK::BlendEquation(blend.EquationAlpha()),
             .colorWriteMask = write_mask,
         });

src/video_core/renderer_vulkan/vk_pipeline_cache.cpp

@@ -36,6 +36,7 @@
 #include "video_core/renderer_vulkan/vk_scheduler.h"
 #include "video_core/renderer_vulkan/vk_shader_util.h"
 #include "video_core/renderer_vulkan/vk_update_descriptor.h"
+#include "video_core/surface.h"
 #include "video_core/shader_cache.h"
 #include "video_core/shader_environment.h"
 #include "video_core/shader_notify.h"
@@ -105,6 +106,41 @@ Shader::CompareFunction MaxwellToCompareFunction(Maxwell::ComparisonOp compariso
     return {};
 }
 
+Shader::AttributeType RenderTargetAttributeType(Tegra::RenderTargetFormat format) {
+    if (format == Tegra::RenderTargetFormat::NONE) {
+        return Shader::AttributeType::Float;
+    }
+    const auto pixel_format{
+        VideoCore::Surface::PixelFormatFromRenderTargetFormat(format)};
+    if (!VideoCore::Surface::IsPixelFormatInteger(pixel_format)) {
+        return Shader::AttributeType::Float;
+    }
+    if (VideoCore::Surface::IsPixelFormatSignedInteger(pixel_format)) {
+        return Shader::AttributeType::SignedInt;
+    }
+    return Shader::AttributeType::UnsignedInt;
+}
+
+VkShaderStageFlagBits StageToVkStage(Shader::Stage stage) {
+    switch (stage) {
+    case Shader::Stage::VertexA:
+    case Shader::Stage::VertexB:
+        return VK_SHADER_STAGE_VERTEX_BIT;
+    case Shader::Stage::TessellationControl:
+        return VK_SHADER_STAGE_TESSELLATION_CONTROL_BIT;
+    case Shader::Stage::TessellationEval:
+        return VK_SHADER_STAGE_TESSELLATION_EVALUATION_BIT;
+    case Shader::Stage::Geometry:
+        return VK_SHADER_STAGE_GEOMETRY_BIT;
+    case Shader::Stage::Fragment:
+        return VK_SHADER_STAGE_FRAGMENT_BIT;
+    case Shader::Stage::Compute:
+        return VK_SHADER_STAGE_COMPUTE_BIT;
+    default:
+        return VK_SHADER_STAGE_VERTEX_BIT;
+    }
+}
+
 Shader::AttributeType CastAttributeType(const FixedPipelineState::VertexAttribute& attr) {
     if (attr.enabled == 0) {
         return Shader::AttributeType::Disabled;
@@ -238,6 +274,10 @@ Shader::RuntimeInfo MakeRuntimeInfo(std::span<const Shader::IR::Program> program
         info.alpha_test_func = MaxwellToCompareFunction(
             key.state.UnpackComparisonOp(key.state.alpha_test_func.Value()));
         info.alpha_test_reference = std::bit_cast<float>(key.state.alpha_test_ref);
+        for (size_t index = 0; index < Maxwell::NumRenderTargets; ++index) {
+            const auto format = static_cast<Tegra::RenderTargetFormat>(key.state.color_formats[index]);
+            info.color_output_types[index] = RenderTargetAttributeType(format);
+        }
         break;
     default:
         break;
@@ -278,8 +318,8 @@ size_t GetTotalPipelineWorkers() {
     const size_t max_core_threads =
         std::max<size_t>(static_cast<size_t>(std::thread::hardware_concurrency()), 2ULL) - 1ULL;
 #ifdef ANDROID
-    // Leave at least a few cores free in android
+    // Leave at least one core free on Android to reduce thermal pressure.
-    constexpr size_t free_cores = 3ULL;
+    constexpr size_t free_cores = 1ULL;
     if (max_core_threads <= free_cores) {
         return 1ULL;
     }
@@ -326,6 +366,7 @@ PipelineCache::PipelineCache(Tegra::MaxwellDeviceMemoryManager& device_memory_,
               "VkPipelineBuilder"),
       serialization_thread(1, "VkPipelineSerialization") {
     const auto& float_control{device.FloatControlProperties()};
+    const bool float_controls_supported{device.IsKhrShaderFloatControlsSupported()};
     const VkDriverId driver_id{device.GetDriverID()};
     profile = Shader::Profile{
         .supported_spirv = device.SupportedSpirvVersion(),
@@ -335,20 +376,24 @@ PipelineCache::PipelineCache(Tegra::MaxwellDeviceMemoryManager& device_memory_,
         .support_int16 = device.IsShaderInt16Supported(),
         .support_int64 = device.IsShaderInt64Supported(),
         .support_vertex_instance_id = false,
-        .support_float_controls = device.IsKhrShaderFloatControlsSupported(),
+        .support_float_controls = float_controls_supported,
-        .support_separate_denorm_behavior =
+        .support_separate_denorm_behavior = float_controls_supported &&
             float_control.denormBehaviorIndependence == VK_SHADER_FLOAT_CONTROLS_INDEPENDENCE_ALL,
-        .support_separate_rounding_mode =
+        .support_separate_rounding_mode = float_controls_supported &&
             float_control.roundingModeIndependence == VK_SHADER_FLOAT_CONTROLS_INDEPENDENCE_ALL,
-        .support_fp16_denorm_preserve = float_control.shaderDenormPreserveFloat16 != VK_FALSE,
+        .support_fp16_denorm_preserve = float_controls_supported &&
-        .support_fp32_denorm_preserve = float_control.shaderDenormPreserveFloat32 != VK_FALSE,
+            float_control.shaderDenormPreserveFloat16 != VK_FALSE,
-        .support_fp16_denorm_flush = float_control.shaderDenormFlushToZeroFloat16 != VK_FALSE,
+        .support_fp32_denorm_preserve = float_controls_supported &&
-        .support_fp32_denorm_flush = float_control.shaderDenormFlushToZeroFloat32 != VK_FALSE,
+            float_control.shaderDenormPreserveFloat32 != VK_FALSE,
-        .support_fp16_signed_zero_nan_preserve =
+        .support_fp16_denorm_flush = float_controls_supported &&
+            float_control.shaderDenormFlushToZeroFloat16 != VK_FALSE,
+        .support_fp32_denorm_flush = float_controls_supported &&
+            float_control.shaderDenormFlushToZeroFloat32 != VK_FALSE,
+        .support_fp16_signed_zero_nan_preserve = float_controls_supported &&
             float_control.shaderSignedZeroInfNanPreserveFloat16 != VK_FALSE,
-        .support_fp32_signed_zero_nan_preserve =
+        .support_fp32_signed_zero_nan_preserve = float_controls_supported &&
             float_control.shaderSignedZeroInfNanPreserveFloat32 != VK_FALSE,
-        .support_fp64_signed_zero_nan_preserve =
+        .support_fp64_signed_zero_nan_preserve = float_controls_supported &&
             float_control.shaderSignedZeroInfNanPreserveFloat64 != VK_FALSE,
         .support_explicit_workgroup_layout = device.IsKhrWorkgroupMemoryExplicitLayoutSupported(),
         .support_vote = device.IsSubgroupFeatureSupported(VK_SUBGROUP_FEATURE_VOTE_BIT),
@@ -404,6 +449,27 @@ PipelineCache::PipelineCache(Tegra::MaxwellDeviceMemoryManager& device_memory_,
         .support_conditional_barrier = device.SupportsConditionalBarriers(),
     };
 
+    profile.warp_stage_support_mask = 0;
+    static constexpr std::array kAllStages{
+        Shader::Stage::VertexA,          Shader::Stage::VertexB,
+        Shader::Stage::TessellationControl, Shader::Stage::TessellationEval,
+        Shader::Stage::Geometry,         Shader::Stage::Fragment,
+        Shader::Stage::Compute,
+    };
+    for (const auto stage : kAllStages) {
+        const auto vk_stage = StageToVkStage(stage);
+        if (device.SupportsWarpIntrinsics(vk_stage)) {
+            profile.warp_stage_support_mask |= 1u << static_cast<u32>(stage);
+        }
+    }
+    profile.support_vote = profile.warp_stage_support_mask != 0;
+
+    if (!profile.SupportsWarpIntrinsics(Shader::Stage::Fragment)) {
+        LOG_WARNING(Render_Vulkan,
+                    "Fragment shaders lack subgroup support on this driver; warp intrinsics will be "
+                    "approximated and visual artifacts may remain");
+    }
+
     if (device.GetMaxVertexInputAttributes() < Maxwell::NumVertexAttributes) {
         LOG_WARNING(Render_Vulkan, "maxVertexInputAttributes is too low: {} < {}",
                     device.GetMaxVertexInputAttributes(), Maxwell::NumVertexAttributes);
@@ -435,6 +501,11 @@ PipelineCache::PipelineCache(Tegra::MaxwellDeviceMemoryManager& device_memory_,
     
     dynamic_features.has_extended_dynamic_state_3_blend = 
         device.IsExtExtendedDynamicState3BlendingSupported();
+    dynamic_features.has_dual_source_blend = device.SupportsDualSourceBlend();
+    if (!dynamic_features.has_dual_source_blend) {
+        LOG_WARNING(Render_Vulkan, "Dual-source blending unsupported, disabling dynamic blend");
+        dynamic_features.has_extended_dynamic_state_3_blend = false;
+    }
     dynamic_features.has_extended_dynamic_state_3_enables = 
         device.IsExtExtendedDynamicState3EnablesSupported();
     
@@ -451,6 +522,13 @@ PipelineCache::~PipelineCache() {
     }
 }
 
+void PipelineCache::DrainPendingBuilds() {
+    if (!device.HasBrokenParallelShaderCompiling()) {
+        return;
+    }
+    workers.WaitForRequests();
+}
+
 GraphicsPipeline* PipelineCache::CurrentGraphicsPipeline() {
 
     if (!RefreshStages(graphics_key.unique_hashes)) {
@@ -481,12 +559,17 @@ ComputePipeline* PipelineCache::CurrentComputePipeline() {
         .shared_memory_size = qmd.shared_alloc,
         .workgroup_size{qmd.block_dim_x, qmd.block_dim_y, qmd.block_dim_z},
     };
-    const auto [pair, is_new]{compute_cache.try_emplace(key)};
+    const auto [pair, inserted]{compute_cache.try_emplace(key)};
     auto& pipeline{pair->second};
-    if (!is_new) {
+    if (!pipeline) {
-        return pipeline.get();
+        auto [slot, should_build] = AcquireComputeBuildSlot(key);
+        if (!should_build) {
+            WaitForBuildCompletion(slot);
+        } else {
+            pipeline = CreateComputePipeline(key, shader);
+            ReleaseComputeBuildSlot(key, slot);
+        }
     }
-    pipeline = CreateComputePipeline(key, shader);
     return pipeline.get();
 }
 
@@ -602,13 +685,20 @@ void PipelineCache::LoadDiskResources(u64 title_id, std::stop_token stop_loading
 }
 
 GraphicsPipeline* PipelineCache::CurrentGraphicsPipelineSlowPath() {
-    const auto [pair, is_new]{graphics_cache.try_emplace(graphics_key)};
+    const auto [pair, inserted]{graphics_cache.try_emplace(graphics_key)};
     auto& pipeline{pair->second};
-    if (is_new) {
-        pipeline = CreateGraphicsPipeline();
-    }
     if (!pipeline) {
-        return nullptr;
+        const auto key = pair->first;
+        auto [slot, should_build] = AcquireGraphicsBuildSlot(key);
+        if (!should_build) {
+            WaitForBuildCompletion(slot);
+        } else {
+            pipeline = CreateGraphicsPipeline();
+            ReleaseGraphicsBuildSlot(key, slot);
+        }
+        if (!pipeline) {
+            return nullptr;
+        }
     }
     if (current_pipeline) {
         current_pipeline->AddTransition(pipeline.get());
@@ -631,6 +721,7 @@ GraphicsPipeline* PipelineCache::BuiltPipeline(GraphicsPipeline* pipeline) const
     if (draw_state.index_buffer.count <= 6 || draw_state.vertex_buffer.count <= 6) {
         return pipeline;
     }
+    scheduler.KeepAliveTick();
     return nullptr;
 }
 
@@ -734,6 +825,10 @@ std::unique_ptr<GraphicsPipeline> PipelineCache::CreateGraphicsPipeline(
     }
     LOG_ERROR(Render_Vulkan, "{}", exception.what());
     return nullptr;
+} catch (const vk::Exception& exception) {
+    LOG_ERROR(Render_Vulkan, "Failed to create graphics pipeline 0x{:016x}: {}", key.Hash(),
+              exception.what());
+    return nullptr;
 }
 
 std::unique_ptr<GraphicsPipeline> PipelineCache::CreateGraphicsPipeline() {
@@ -797,6 +892,19 @@ std::unique_ptr<ComputePipeline> PipelineCache::CreateComputePipeline(
     }
 
     auto program{TranslateProgram(pools.inst, pools.block, env, cfg, host_info)};
+    const VkDriverIdKHR driver_id = device.GetDriverID();
+    const bool needs_shared_mem_clamp =
+        driver_id == VK_DRIVER_ID_QUALCOMM_PROPRIETARY ||
+        driver_id == VK_DRIVER_ID_ARM_PROPRIETARY;
+    const u32 max_shared_memory = device.GetMaxComputeSharedMemorySize();
+    if (needs_shared_mem_clamp && program.shared_memory_size > max_shared_memory) {
+        LOG_WARNING(Render_Vulkan,
+                    "Compute shader 0x{:016x} requests {}KB shared memory but device max is {}KB - clamping",
+                    key.unique_hash,
+                    program.shared_memory_size / 1024,
+                    max_shared_memory / 1024);
+        program.shared_memory_size = max_shared_memory;
+    }
     const std::vector<u32> code{EmitSPIRV(profile, program, this->optimize_spirv_output)};
     device.SaveShader(code);
     vk::ShaderModule spv_module{BuildShader(device, code)};
@@ -812,6 +920,10 @@ std::unique_ptr<ComputePipeline> PipelineCache::CreateComputePipeline(
 } catch (const Shader::Exception& exception) {
     LOG_ERROR(Render_Vulkan, "{}", exception.what());
     return nullptr;
+} catch (const vk::Exception& exception) {
+    LOG_ERROR(Render_Vulkan, "Failed to create compute pipeline 0x{:016x}: {}", key.Hash(),
+              exception.what());
+    return nullptr;
 }
 
 void PipelineCache::SerializeVulkanPipelineCache(const std::filesystem::path& filename,
@@ -909,4 +1021,68 @@ vk::PipelineCache PipelineCache::LoadVulkanPipelineCache(const std::filesystem::
     }
 }
 
+auto PipelineCache::AcquireGraphicsBuildSlot(const GraphicsPipelineCacheKey& key)
+    -> std::pair<InFlightPipelinePtr, bool> {
+    std::scoped_lock lock(graphics_inflight_mutex);
+    auto [it, inserted] = graphics_inflight_builds.try_emplace(key);
+    if (inserted || !it->second) {
+        it->second = std::make_shared<InFlightPipelineBuild>();
+        return {it->second, true};
+    }
+    return {it->second, false};
+}
+
+auto PipelineCache::AcquireComputeBuildSlot(const ComputePipelineCacheKey& key)
+    -> std::pair<InFlightPipelinePtr, bool> {
+    std::scoped_lock lock(compute_inflight_mutex);
+    auto [it, inserted] = compute_inflight_builds.try_emplace(key);
+    if (inserted || !it->second) {
+        it->second = std::make_shared<InFlightPipelineBuild>();
+        return {it->second, true};
+    }
+    return {it->second, false};
+}
+
+void PipelineCache::ReleaseGraphicsBuildSlot(const GraphicsPipelineCacheKey& key,
+                                             const InFlightPipelinePtr& slot) {
+    if (!slot) {
+        return;
+    }
+    {
+        std::scoped_lock slot_lock(slot->mutex);
+        slot->building = false;
+    }
+    slot->cv.notify_all();
+    std::scoped_lock map_lock(graphics_inflight_mutex);
+    auto it = graphics_inflight_builds.find(key);
+    if (it != graphics_inflight_builds.end() && it->second == slot) {
+        graphics_inflight_builds.erase(it);
+    }
+}
+
+void PipelineCache::ReleaseComputeBuildSlot(const ComputePipelineCacheKey& key,
+                                            const InFlightPipelinePtr& slot) {
+    if (!slot) {
+        return;
+    }
+    {
+        std::scoped_lock slot_lock(slot->mutex);
+        slot->building = false;
+    }
+    slot->cv.notify_all();
+    std::scoped_lock map_lock(compute_inflight_mutex);
+    auto it = compute_inflight_builds.find(key);
+    if (it != compute_inflight_builds.end() && it->second == slot) {
+        compute_inflight_builds.erase(it);
+    }
+}
+
+void PipelineCache::WaitForBuildCompletion(const InFlightPipelinePtr& slot) const {
+    if (!slot) {
+        return;
+    }
+    std::unique_lock lock(slot->mutex);
+    slot->cv.wait(lock, [&] { return !slot->building; });
+}
+
 } // namespace Vulkan

src/video_core/renderer_vulkan/vk_pipeline_cache.h

@@ -5,8 +5,10 @@
 
 #include <array>
 #include <cstddef>
+#include <condition_variable>
 #include <filesystem>
 #include <memory>
+#include <mutex>
 #include <type_traits>
 #include <unordered_map>
 #include <vector>
@@ -113,7 +115,17 @@ public:
     void LoadDiskResources(u64 title_id, std::stop_token stop_loading,
                            const VideoCore::DiskResourceLoadCallback& callback);
 
+    void DrainPendingBuilds();
+
 private:
+    struct InFlightPipelineBuild {
+        std::mutex mutex;
+        std::condition_variable cv;
+        bool building{true};
+    };
+
+    using InFlightPipelinePtr = std::shared_ptr<InFlightPipelineBuild>;
+
     [[nodiscard]] GraphicsPipeline* CurrentGraphicsPipelineSlowPath();
 
     [[nodiscard]] GraphicsPipeline* BuiltPipeline(GraphicsPipeline* pipeline) const noexcept;
@@ -140,6 +152,14 @@ private:
     vk::PipelineCache LoadVulkanPipelineCache(const std::filesystem::path& filename,
                                               u32 expected_cache_version);
 
+    std::pair<InFlightPipelinePtr, bool> AcquireGraphicsBuildSlot(
+        const GraphicsPipelineCacheKey& key);
+    std::pair<InFlightPipelinePtr, bool> AcquireComputeBuildSlot(
+        const ComputePipelineCacheKey& key);
+    void ReleaseGraphicsBuildSlot(const GraphicsPipelineCacheKey& key, const InFlightPipelinePtr& slot);
+    void ReleaseComputeBuildSlot(const ComputePipelineCacheKey& key, const InFlightPipelinePtr& slot);
+    void WaitForBuildCompletion(const InFlightPipelinePtr& slot) const;
+
     const Device& device;
     Scheduler& scheduler;
     DescriptorPool& descriptor_pool;
@@ -158,6 +178,11 @@ private:
     std::unordered_map<ComputePipelineCacheKey, std::unique_ptr<ComputePipeline>> compute_cache;
     std::unordered_map<GraphicsPipelineCacheKey, std::unique_ptr<GraphicsPipeline>> graphics_cache;
 
+    std::mutex graphics_inflight_mutex;
+    std::unordered_map<GraphicsPipelineCacheKey, InFlightPipelinePtr> graphics_inflight_builds;
+    std::mutex compute_inflight_mutex;
+    std::unordered_map<ComputePipelineCacheKey, InFlightPipelinePtr> compute_inflight_builds;
+
     ShaderPools main_pools;
 
     Shader::Profile profile;

src/video_core/renderer_vulkan/vk_query_cache.cpp

@@ -42,7 +42,8 @@ public:
     static constexpr size_t BANK_SIZE = 256;
     static constexpr size_t QUERY_SIZE = 8;
     explicit SamplesQueryBank(const Device& device_, size_t index_)
-        : BankBase(BANK_SIZE), device{device_}, index{index_} {
+        : BankBase(BANK_SIZE), device{device_}, index{index_},
+          supports_host_query_reset{device_.SupportsHostQueryReset()} {
         const auto& dev = device.GetLogical();
         query_pool = dev.CreateQueryPool({
             .sType = VK_STRUCTURE_TYPE_QUERY_POOL_CREATE_INFO,
@@ -60,8 +61,10 @@ public:
     void Reset() override {
         ASSERT(references == 0);
         VideoCommon::BankBase::Reset();
-        const auto& dev = device.GetLogical();
+        if (supports_host_query_reset) {
-        dev.ResetQueryPool(*query_pool, 0, BANK_SIZE);
+            const auto& dev = device.GetLogical();
+            dev.ResetQueryPool(*query_pool, 0, BANK_SIZE);
+        }
         host_results.fill(0ULL);
         next_bank = 0;
     }
@@ -99,6 +102,7 @@ public:
 private:
     const Device& device;
     const size_t index;
+    const bool supports_host_query_reset;
     vk::QueryPool query_pool;
     std::array<u64, BANK_SIZE> host_results;
 };
@@ -1202,21 +1206,24 @@ struct QueryCacheRuntimeImpl {
         hcr_setup.pNext = nullptr;
         hcr_setup.flags = 0;
 
-        conditional_resolve_pass = std::make_unique<ConditionalRenderingResolvePass>(
+        if (device.IsExtConditionalRendering()) {
-            device, scheduler, descriptor_pool, compute_pass_descriptor_queue);
+            conditional_resolve_pass = std::make_unique<ConditionalRenderingResolvePass>(
+                device, scheduler, descriptor_pool, compute_pass_descriptor_queue);
 
-        const VkBufferCreateInfo buffer_ci = {
+            const VkBufferCreateInfo buffer_ci = {
-            .sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO,
+                .sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO,
-            .pNext = nullptr,
+                .pNext = nullptr,
-            .flags = 0,
+                .flags = 0,
-            .size = sizeof(u32),
+                .size = sizeof(u32),
-            .usage = VK_BUFFER_USAGE_TRANSFER_DST_BIT | VK_BUFFER_USAGE_STORAGE_BUFFER_BIT |
+                .usage = VK_BUFFER_USAGE_TRANSFER_DST_BIT | VK_BUFFER_USAGE_STORAGE_BUFFER_BIT |
-                     VK_BUFFER_USAGE_CONDITIONAL_RENDERING_BIT_EXT,
+                         VK_BUFFER_USAGE_CONDITIONAL_RENDERING_BIT_EXT,
-            .sharingMode = VK_SHARING_MODE_EXCLUSIVE,
+                .sharingMode = VK_SHARING_MODE_EXCLUSIVE,
-            .queueFamilyIndexCount = 0,
+                .queueFamilyIndexCount = 0,
-            .pQueueFamilyIndices = nullptr,
+                .pQueueFamilyIndices = nullptr,
-        };
+            };
-        hcr_resolve_buffer = memory_allocator.CreateBuffer(buffer_ci, MemoryUsage::DeviceLocal);
+            hcr_resolve_buffer =
+                memory_allocator.CreateBuffer(buffer_ci, MemoryUsage::DeviceLocal);
+        }
     }
 
     VideoCore::RasterizerInterface* rasterizer;

src/video_core/renderer_vulkan/vk_rasterizer.cpp

@@ -215,6 +215,10 @@ void RasterizerVulkan::PrepareDraw(bool is_indexed, Func&& draw_func) {
     FlushWork();
     gpu_memory->FlushCaching();
 
+    if (device.HasBrokenParallelShaderCompiling()) {
+        pipeline_cache.DrainPendingBuilds();
+    }
+
     GraphicsPipeline* const pipeline{pipeline_cache.CurrentGraphicsPipeline()};
     if (!pipeline) {
         return;
@@ -996,7 +1000,6 @@ void RasterizerVulkan::UpdateDynamicStates() {
     UpdateDepthBounds(regs);
     UpdateStencilFaces(regs);
     UpdateLineWidth(regs);
-
     // EDS1: CullMode, DepthCompare, FrontFace, StencilOp, DepthBoundsTest, DepthTest, DepthWrite, StencilTest
     if (device.IsExtExtendedDynamicStateSupported()) {
         UpdateCullMode(regs);
@@ -1236,116 +1239,138 @@ void RasterizerVulkan::UpdateBlendConstants(Tegra::Engines::Maxwell3D::Regs& reg
     if (!device.UsesAdvancedCoreDynamicState()) {
         return;
     }
-    const std::array blend_color = {regs.blend_color.r, regs.blend_color.g, regs.blend_color.b,
+    const std::array<float, 4> blend_constants{
-                                    regs.blend_color.a};
+        regs.blend_color.r,
-    scheduler.Record(
+        regs.blend_color.g,
-        [blend_color](vk::CommandBuffer cmdbuf) { cmdbuf.SetBlendConstants(blend_color.data()); });
+        regs.blend_color.b,
+        regs.blend_color.a,
+    };
+    scheduler.Record([blend_constants](vk::CommandBuffer cmdbuf) {
+        cmdbuf.SetBlendConstants(blend_constants.data());
+    });
 }
 
 void RasterizerVulkan::UpdateDepthBounds(Tegra::Engines::Maxwell3D::Regs& regs) {
     if (!state_tracker.TouchDepthBounds()) {
         return;
     }
-    if (!device.UsesAdvancedCoreDynamicState() || !device.IsDepthBoundsSupported()) {
+    if (!device.IsDepthBoundsSupported()) {
         return;
     }
-    scheduler.Record([min = regs.depth_bounds[0], max = regs.depth_bounds[1]](
+    if (!device.UsesAdvancedCoreDynamicState()) {
-                         vk::CommandBuffer cmdbuf) { cmdbuf.SetDepthBounds(min, max); });
+        return;
+    }
+    const bool unrestricted = device.IsExtDepthRangeUnrestrictedSupported();
+    const float min_depth = unrestricted ? regs.depth_bounds[0]
+                                         : std::clamp(regs.depth_bounds[0], 0.0f, 1.0f);
+    const float max_depth = unrestricted ? regs.depth_bounds[1]
+                                         : std::clamp(regs.depth_bounds[1], 0.0f, 1.0f);
+    scheduler.Record([min_depth, max_depth](vk::CommandBuffer cmdbuf) {
+        cmdbuf.SetDepthBounds(min_depth, max_depth);
+    });
 }
 
 void RasterizerVulkan::UpdateStencilFaces(Tegra::Engines::Maxwell3D::Regs& regs) {
-    if (!state_tracker.TouchStencilProperties()) {
+    const bool two_sided = regs.stencil_two_side_enable != 0;
+    const bool update_properties = state_tracker.TouchStencilProperties();
+    const bool update_side = state_tracker.TouchStencilSide(two_sided);
+    const bool refs_dirty = state_tracker.TouchStencilReference();
+    const bool write_dirty = state_tracker.TouchStencilWriteMask();
+    const bool compare_dirty = state_tracker.TouchStencilCompare();
+
+    const bool update_references = update_properties || update_side || refs_dirty;
+    const bool update_write_masks = update_properties || update_side || write_dirty;
+    const bool update_compare_masks = update_properties || update_side || compare_dirty;
+
+    if (!update_references && !update_write_masks && !update_compare_masks) {
+        state_tracker.ClearStencilReset();
         return;
     }
+
     if (!device.UsesAdvancedCoreDynamicState()) {
         state_tracker.ClearStencilReset();
         return;
     }
-    bool update_references = state_tracker.TouchStencilReference();
+
-    bool update_write_mask = state_tracker.TouchStencilWriteMask();
-    bool update_compare_masks = state_tracker.TouchStencilCompare();
-    if (state_tracker.TouchStencilSide(regs.stencil_two_side_enable != 0)) {
-        update_references = true;
-        update_write_mask = true;
-        update_compare_masks = true;
-    }
     if (update_references) {
-        [&]() {
+        if (two_sided) {
-            if (regs.stencil_two_side_enable) {
+            const bool front_dirty =
-                if (!state_tracker.CheckStencilReferenceFront(regs.stencil_front_ref) &&
+                state_tracker.CheckStencilReferenceFront(regs.stencil_front_ref);
-                    !state_tracker.CheckStencilReferenceBack(regs.stencil_back_ref)) {
+            const bool back_dirty =
-                    return;
+                state_tracker.CheckStencilReferenceBack(regs.stencil_back_ref);
-                }
+            if (front_dirty || back_dirty) {
-            } else {
+                scheduler.Record([front_ref = regs.stencil_front_ref,
-                if (!state_tracker.CheckStencilReferenceFront(regs.stencil_front_ref)) {
+                                  back_ref = regs.stencil_back_ref,
-                    return;
+                                  is_two_sided = two_sided](vk::CommandBuffer cmdbuf) {
-                }
+                    const bool set_back = is_two_sided && front_ref != back_ref;
+                    cmdbuf.SetStencilReference(set_back ? VK_STENCIL_FACE_FRONT_BIT
+                                                        : VK_STENCIL_FACE_FRONT_AND_BACK,
+                                               front_ref);
+                    if (set_back) {
+                        cmdbuf.SetStencilReference(VK_STENCIL_FACE_BACK_BIT, back_ref);
+                    }
+                });
             }
-            scheduler.Record([front_ref = regs.stencil_front_ref, back_ref = regs.stencil_back_ref,
+        } else if (state_tracker.CheckStencilReferenceFront(regs.stencil_front_ref)) {
-                              two_sided = regs.stencil_two_side_enable](vk::CommandBuffer cmdbuf) {
+            const u32 reference = regs.stencil_front_ref;
-                const bool set_back = two_sided && front_ref != back_ref;
+            scheduler.Record([reference](vk::CommandBuffer cmdbuf) {
-                // Front face
+                cmdbuf.SetStencilReference(VK_STENCIL_FACE_FRONT_AND_BACK, reference);
-                cmdbuf.SetStencilReference(set_back ? VK_STENCIL_FACE_FRONT_BIT
-                                                    : VK_STENCIL_FACE_FRONT_AND_BACK,
-                                           front_ref);
-                if (set_back) {
-                    cmdbuf.SetStencilReference(VK_STENCIL_FACE_BACK_BIT, back_ref);
-                }
             });
-        }();
+        }
     }
-    if (update_write_mask) {
+
-        [&]() {
+    if (update_write_masks) {
-            if (regs.stencil_two_side_enable) {
+        if (two_sided) {
-                if (!state_tracker.CheckStencilWriteMaskFront(regs.stencil_front_mask) &&
+            if (state_tracker.CheckStencilWriteMaskFront(regs.stencil_front_mask) ||
-                    !state_tracker.CheckStencilWriteMaskBack(regs.stencil_back_mask)) {
+                state_tracker.CheckStencilWriteMaskBack(regs.stencil_back_mask)) {
-                    return;
+                scheduler.Record([
-                }
+                                     front_write_mask = regs.stencil_front_mask,
-            } else {
+                                     back_write_mask = regs.stencil_back_mask,
-                if (!state_tracker.CheckStencilWriteMaskFront(regs.stencil_front_mask)) {
+                                     is_two_sided = regs.stencil_two_side_enable
-                    return;
+                                 ](vk::CommandBuffer cmdbuf) {
-                }
+                    const bool set_back = is_two_sided && front_write_mask != back_write_mask;
+                    cmdbuf.SetStencilWriteMask(set_back ? VK_STENCIL_FACE_FRONT_BIT
+                                                        : VK_STENCIL_FACE_FRONT_AND_BACK,
+                                               front_write_mask);
+                    if (set_back) {
+                        cmdbuf.SetStencilWriteMask(VK_STENCIL_FACE_BACK_BIT, back_write_mask);
+                    }
+                });
             }
-            scheduler.Record([front_write_mask = regs.stencil_front_mask,
+        } else if (state_tracker.CheckStencilWriteMaskFront(regs.stencil_front_mask)) {
-                              back_write_mask = regs.stencil_back_mask,
+            const u32 front_write_mask = regs.stencil_front_mask;
-                              two_sided = regs.stencil_two_side_enable](vk::CommandBuffer cmdbuf) {
+            scheduler.Record([front_write_mask](vk::CommandBuffer cmdbuf) {
-                const bool set_back = two_sided && front_write_mask != back_write_mask;
+                cmdbuf.SetStencilWriteMask(VK_STENCIL_FACE_FRONT_AND_BACK, front_write_mask);
-                // Front face
-                cmdbuf.SetStencilWriteMask(set_back ? VK_STENCIL_FACE_FRONT_BIT
-                                                    : VK_STENCIL_FACE_FRONT_AND_BACK,
-                                           front_write_mask);
-                if (set_back) {
-                    cmdbuf.SetStencilWriteMask(VK_STENCIL_FACE_BACK_BIT, back_write_mask);
-                }
             });
-        }();
+        }
     }
+
     if (update_compare_masks) {
-        [&]() {
+        if (two_sided) {
-            if (regs.stencil_two_side_enable) {
+            if (state_tracker.CheckStencilCompareMaskFront(regs.stencil_front_func_mask) ||
-                if (!state_tracker.CheckStencilCompareMaskFront(regs.stencil_front_func_mask) &&
+                state_tracker.CheckStencilCompareMaskBack(regs.stencil_back_func_mask)) {
-                    !state_tracker.CheckStencilCompareMaskBack(regs.stencil_back_func_mask)) {
+                scheduler.Record([
-                    return;
+                                     front_test_mask = regs.stencil_front_func_mask,
-                }
+                                     back_test_mask = regs.stencil_back_func_mask,
-            } else {
+                                     is_two_sided = regs.stencil_two_side_enable
-                if (!state_tracker.CheckStencilCompareMaskFront(regs.stencil_front_func_mask)) {
+                                 ](vk::CommandBuffer cmdbuf) {
-                    return;
+                    const bool set_back = is_two_sided && front_test_mask != back_test_mask;
-                }
+                    cmdbuf.SetStencilCompareMask(set_back ? VK_STENCIL_FACE_FRONT_BIT
+                                                          : VK_STENCIL_FACE_FRONT_AND_BACK,
+                                                 front_test_mask);
+                    if (set_back) {
+                        cmdbuf.SetStencilCompareMask(VK_STENCIL_FACE_BACK_BIT, back_test_mask);
+                    }
+                });
             }
-            scheduler.Record([front_test_mask = regs.stencil_front_func_mask,
+        } else if (state_tracker.CheckStencilCompareMaskFront(regs.stencil_front_func_mask)) {
-                              back_test_mask = regs.stencil_back_func_mask,
+            const u32 front_test_mask = regs.stencil_front_func_mask;
-                              two_sided = regs.stencil_two_side_enable](vk::CommandBuffer cmdbuf) {
+            scheduler.Record([front_test_mask](vk::CommandBuffer cmdbuf) {
-                const bool set_back = two_sided && front_test_mask != back_test_mask;
+                cmdbuf.SetStencilCompareMask(VK_STENCIL_FACE_FRONT_AND_BACK, front_test_mask);
-                // Front face
-                cmdbuf.SetStencilCompareMask(set_back ? VK_STENCIL_FACE_FRONT_BIT
-                                                      : VK_STENCIL_FACE_FRONT_AND_BACK,
-                                             front_test_mask);
-                if (set_back) {
-                    cmdbuf.SetStencilCompareMask(VK_STENCIL_FACE_BACK_BIT, back_test_mask);
-                }
             });
-        }();
+        }
     }
+
     state_tracker.ClearStencilReset();
 }
 
@@ -1368,65 +1393,15 @@ void RasterizerVulkan::UpdateCullMode(Tegra::Engines::Maxwell3D::Regs& regs) {
     });
 }
 
-void RasterizerVulkan::UpdateDepthBoundsTestEnable(Tegra::Engines::Maxwell3D::Regs& regs) {
+void RasterizerVulkan::UpdateConservativeRasterizationMode(
-    if (!state_tracker.TouchDepthBoundsTestEnable()) {
+    Tegra::Engines::Maxwell3D::Regs& regs) {
-        return;
-    }
-    bool enabled = regs.depth_bounds_enable;
-    if (enabled && !device.IsDepthBoundsSupported()) {
-        LOG_WARNING(Render_Vulkan, "Depth bounds is enabled but not supported");
-        enabled = false;
-    }
-    scheduler.Record([enable = enabled](vk::CommandBuffer cmdbuf) {
-        cmdbuf.SetDepthBoundsTestEnableEXT(enable);
-    });
-}
-
-void RasterizerVulkan::UpdateDepthTestEnable(Tegra::Engines::Maxwell3D::Regs& regs) {
-    if (!state_tracker.TouchDepthTestEnable()) {
-        return;
-    }
-    scheduler.Record([enable = regs.depth_test_enable](vk::CommandBuffer cmdbuf) {
-        cmdbuf.SetDepthTestEnableEXT(enable);
-    });
-}
-
-void RasterizerVulkan::UpdateDepthWriteEnable(Tegra::Engines::Maxwell3D::Regs& regs) {
-    if (!state_tracker.TouchDepthWriteEnable()) {
-        return;
-    }
-    scheduler.Record([enable = regs.depth_write_enabled](vk::CommandBuffer cmdbuf) {
-        cmdbuf.SetDepthWriteEnableEXT(enable);
-    });
-}
-
-void RasterizerVulkan::UpdatePrimitiveRestartEnable(Tegra::Engines::Maxwell3D::Regs& regs) {
-    if (!state_tracker.TouchPrimitiveRestartEnable()) {
-        return;
-    }
-    scheduler.Record([enable = regs.primitive_restart.enabled](vk::CommandBuffer cmdbuf) {
-        cmdbuf.SetPrimitiveRestartEnableEXT(enable);
-    });
-}
-
-void RasterizerVulkan::UpdateRasterizerDiscardEnable(Tegra::Engines::Maxwell3D::Regs& regs) {
-    if (!state_tracker.TouchRasterizerDiscardEnable()) {
-        return;
-    }
-    scheduler.Record([disable = regs.rasterize_enable](vk::CommandBuffer cmdbuf) {
-        cmdbuf.SetRasterizerDiscardEnableEXT(disable == 0);
-    });
-}
-
-void RasterizerVulkan::UpdateConservativeRasterizationMode(Tegra::Engines::Maxwell3D::Regs& regs) {
     if (!state_tracker.TouchConservativeRasterizationMode()) {
         return;
     }
-
+    if (!device.SupportsDynamicState3ConservativeRasterizationMode() ||
-    if (!device.SupportsDynamicState3ConservativeRasterizationMode()) {
+        !device.IsExtConservativeRasterizationSupported()) {
         return;
     }
-
     scheduler.Record([enable = regs.conservative_raster_enable](vk::CommandBuffer cmdbuf) {
         cmdbuf.SetConservativeRasterizationModeEXT(
             enable ? VK_CONSERVATIVE_RASTERIZATION_MODE_OVERESTIMATE_EXT
@@ -1485,6 +1460,25 @@ void RasterizerVulkan::UpdateLineRasterizationMode(Tegra::Engines::Maxwell3D::Re
     });
 }
 
+void RasterizerVulkan::UpdatePrimitiveRestartEnable(Tegra::Engines::Maxwell3D::Regs& regs) {
+    if (!state_tracker.TouchPrimitiveRestartEnable()) {
+        return;
+    }
+    scheduler.Record([enable = regs.primitive_restart.enabled != 0](vk::CommandBuffer cmdbuf) {
+        cmdbuf.SetPrimitiveRestartEnableEXT(enable);
+    });
+}
+
+void RasterizerVulkan::UpdateRasterizerDiscardEnable(Tegra::Engines::Maxwell3D::Regs& regs) {
+    if (!state_tracker.TouchRasterizerDiscardEnable()) {
+        return;
+    }
+    const bool discard = regs.rasterize_enable == 0;
+    scheduler.Record([discard](vk::CommandBuffer cmdbuf) {
+        cmdbuf.SetRasterizerDiscardEnableEXT(discard);
+    });
+}
+
 void RasterizerVulkan::UpdateDepthBiasEnable(Tegra::Engines::Maxwell3D::Regs& regs) {
     if (!state_tracker.TouchDepthBiasEnable()) {
         return;
@@ -1593,6 +1587,36 @@ void RasterizerVulkan::UpdateDepthCompareOp(Tegra::Engines::Maxwell3D::Regs& reg
     });
 }
 
+void RasterizerVulkan::UpdateDepthBoundsTestEnable(Tegra::Engines::Maxwell3D::Regs& regs) {
+    if (!state_tracker.TouchDepthBoundsTestEnable()) {
+        return;
+    }
+    if (!device.IsDepthBoundsSupported()) {
+        return;
+    }
+    scheduler.Record([enable = regs.depth_bounds_enable != 0](vk::CommandBuffer cmdbuf) {
+        cmdbuf.SetDepthBoundsTestEnableEXT(enable);
+    });
+}
+
+void RasterizerVulkan::UpdateDepthTestEnable(Tegra::Engines::Maxwell3D::Regs& regs) {
+    if (!state_tracker.TouchDepthTestEnable()) {
+        return;
+    }
+    scheduler.Record([enable = regs.depth_test_enable != 0](vk::CommandBuffer cmdbuf) {
+        cmdbuf.SetDepthTestEnableEXT(enable);
+    });
+}
+
+void RasterizerVulkan::UpdateDepthWriteEnable(Tegra::Engines::Maxwell3D::Regs& regs) {
+    if (!state_tracker.TouchDepthWriteEnable()) {
+        return;
+    }
+    scheduler.Record([enable = regs.depth_write_enabled != 0](vk::CommandBuffer cmdbuf) {
+        cmdbuf.SetDepthWriteEnableEXT(enable);
+    });
+}
+
 void RasterizerVulkan::UpdateFrontFace(Tegra::Engines::Maxwell3D::Regs& regs) {
     if (!state_tracker.TouchFrontFace()) {
         return;

src/video_core/renderer_vulkan/vk_rasterizer.h

@@ -25,6 +25,7 @@
 #include "video_core/renderer_vulkan/vk_staging_buffer_pool.h"
 #include "video_core/renderer_vulkan/vk_texture_cache.h"
 #include "video_core/renderer_vulkan/vk_update_descriptor.h"
+#include "video_core/texture_cache/samples_helper.h"
 #include "video_core/vulkan_common/vulkan_memory_allocator.h"
 #include "video_core/vulkan_common/vulkan_wrapper.h"
 
@@ -168,7 +169,6 @@ private:
     void UpdateDepthBounds(Tegra::Engines::Maxwell3D::Regs& regs);
     void UpdateStencilFaces(Tegra::Engines::Maxwell3D::Regs& regs);
     void UpdateLineWidth(Tegra::Engines::Maxwell3D::Regs& regs);
-
     void UpdateCullMode(Tegra::Engines::Maxwell3D::Regs& regs);
     void UpdateDepthBoundsTestEnable(Tegra::Engines::Maxwell3D::Regs& regs);
     void UpdateDepthTestEnable(Tegra::Engines::Maxwell3D::Regs& regs);

src/video_core/renderer_vulkan/vk_scheduler.cpp

@@ -4,6 +4,7 @@
 // SPDX-FileCopyrightText: Copyright 2019 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
+#include <chrono>
 #include <memory>
 #include <mutex>
 #include <thread>
@@ -78,6 +79,14 @@ void Scheduler::WaitWorker() {
     std::scoped_lock el{execution_mutex};
 }
 
+void Scheduler::KeepAliveTick() {
+    const auto now = Clock::now();
+    if (now - last_submission_time < KEEPALIVE_INTERVAL) {
+        return;
+    }
+    Flush();
+}
+
 void Scheduler::DispatchWork() {
     if (chunk->Empty()) {
         return;
@@ -271,6 +280,7 @@ u64 Scheduler::SubmitExecution(VkSemaphore signal_semaphore, VkSemaphore wait_se
     });
     chunk->MarkSubmit();
     DispatchWork();
+    last_submission_time = Clock::now();
     return signal_value;
 }
 
@@ -295,8 +305,13 @@ void Scheduler::EndRenderPass()
             return;
         }
 
-        query_cache->CounterEnable(VideoCommon::QueryType::ZPassPixelCount64, false);
+        if (query_cache) {
-        query_cache->NotifySegment(false);
+            query_cache->CounterEnable(VideoCommon::QueryType::ZPassPixelCount64, false);
+            if (query_cache->HasStreamer(VideoCommon::QueryType::StreamingByteCount)) {
+                query_cache->CounterEnable(VideoCommon::QueryType::StreamingByteCount, false);
+            }
+            query_cache->NotifySegment(false);
+        }
 
         Record([num_images = num_renderpass_images,
                        images = renderpass_images,

src/video_core/renderer_vulkan/vk_scheduler.h

@@ -6,6 +6,7 @@
 
 #pragma once
 
+#include <chrono>
 #include <condition_variable>
 #include <cstddef>
 #include <functional>
@@ -52,6 +53,9 @@ public:
     /// safe to touch worker resources.
     void WaitWorker();
 
+    /// Submits a tiny chunk of work if recent GPU submissions are stale.
+    void KeepAliveTick();
+
     /// Sends currently recorded work to the worker thread.
     void DispatchWork();
 
@@ -128,6 +132,8 @@ public:
     std::mutex submit_mutex;
 
 private:
+    using Clock = std::chrono::steady_clock;
+
     class Command {
     public:
         virtual ~Command() = default;
@@ -250,6 +256,9 @@ private:
 
     State state;
 
+    Clock::time_point last_submission_time{Clock::time_point::min()};
+    static constexpr std::chrono::milliseconds KEEPALIVE_INTERVAL{4};
+
     u32 num_renderpass_images = 0;
     std::array<VkImage, 9> renderpass_images{};
     std::array<VkImageSubresourceRange, 9> renderpass_image_ranges{};

src/video_core/renderer_vulkan/vk_staging_buffer_pool.cpp

@@ -5,6 +5,7 @@
 // SPDX-License-Identifier: GPL-3.0-or-later
 
 #include <algorithm>
+#include <optional>
 #include <utility>
 #include <vector>
 
@@ -49,6 +50,7 @@ size_t GetStreamBufferSize(const Device& device) {
 }
 } // Anonymous namespace
 
+
 StagingBufferPool::StagingBufferPool(const Device& device_, MemoryAllocator& memory_allocator_,
                                      Scheduler& scheduler_)
     : device{device_}, memory_allocator{memory_allocator_}, scheduler{scheduler_},
@@ -74,13 +76,16 @@ StagingBufferPool::StagingBufferPool(const Device& device_, MemoryAllocator& mem
     }
     stream_pointer = stream_buffer.Mapped();
     ASSERT_MSG(!stream_pointer.empty(), "Stream buffer must be host visible!");
+
 }
 
 StagingBufferPool::~StagingBufferPool() = default;
 
 StagingBufferRef StagingBufferPool::Request(size_t size, MemoryUsage usage, bool deferred) {
-    if (!deferred && usage == MemoryUsage::Upload && size <= region_size) {
+    if (!deferred && usage == MemoryUsage::Upload) {
-        return GetStreamBuffer(size);
+        if (size <= region_size) {
+            return GetStreamBuffer(size);
+        }
     }
     return GetStagingBuffer(size, usage, deferred);
 }
@@ -142,6 +147,7 @@ StagingBufferRef StagingBufferPool::GetStreamBuffer(size_t size) {
 }
 
 bool StagingBufferPool::AreRegionsActive(size_t region_begin, size_t region_end) const {
+    scheduler.GetMasterSemaphore().Refresh();
     const u64 gpu_tick = scheduler.GetMasterSemaphore().KnownGpuTick();
     return std::any_of(sync_ticks.begin() + region_begin, sync_ticks.begin() + region_end,
                        [gpu_tick](u64 sync_tick) { return gpu_tick < sync_tick; });

src/video_core/renderer_vulkan/vk_staging_buffer_pool.h

@@ -1,9 +1,13 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2022 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-3.0-or-later
 
 #pragma once
 
 #include <climits>
+#include <optional>
 #include <vector>
 
 #include "common/common_types.h"

src/video_core/renderer_vulkan/vk_texture_cache.h

@@ -3,6 +3,7 @@
 
 #pragma once
 
+#include <optional>
 #include <span>
 
 #include "video_core/texture_cache/texture_cache_base.h"
@@ -56,12 +57,16 @@ public:
 
     void TickFrame();
 
+    void WaitForGpuTick(u64 tick);
+
     u64 GetDeviceLocalMemory() const;
 
     u64 GetDeviceMemoryUsage() const;
 
     bool CanReportMemoryUsage() const;
 
+    std::optional<size_t> GetSamplerHeapBudget() const;
+
     void BlitImage(Framebuffer* dst_framebuffer, ImageView& dst, ImageView& src,
                    const Region2D& dst_region, const Region2D& src_region,
                    Tegra::Engines::Fermi2D::Filter filter,
@@ -108,10 +113,15 @@ public:
         return view_formats[static_cast<std::size_t>(format)];
     }
 
+    bool RequiresBlockCompatibleViewFormats(PixelFormat format) const noexcept {
+        return requires_block_view_formats[static_cast<std::size_t>(format)];
+    }
+
     void BarrierFeedbackLoop();
 
     bool IsFormatDitherable(VideoCore::Surface::PixelFormat format);
     bool IsFormatScalable(VideoCore::Surface::PixelFormat format);
+    bool SupportsLinearFilter(VideoCore::Surface::PixelFormat format) const;
 
     VkFormat GetSupportedFormat(VkFormat requested_format, VkFormatFeatureFlags required_features) const;
 
@@ -125,6 +135,7 @@ public:
     std::unique_ptr<MSAACopyPass> msaa_copy_pass;
     const Settings::ResolutionScalingInfo& resolution;
     std::array<std::vector<VkFormat>, VideoCore::Surface::MaxPixelFormat> view_formats;
+    std::array<bool, VideoCore::Surface::MaxPixelFormat> requires_block_view_formats{};
 
     static constexpr size_t indexing_slots = 8 * sizeof(size_t);
     std::array<vk::Buffer, indexing_slots> buffers{};
@@ -171,6 +182,30 @@ public:
         return (this->*current_image).UsageFlags();
     }
 
+    void TrackGpuReadTick(u64 tick) noexcept {
+        TrackPendingReadTick(tick);
+    }
+
+    void TrackGpuWriteTick(u64 tick) noexcept {
+        TrackPendingWriteTick(tick);
+    }
+
+    void CompleteGpuReadTick(u64 completed_tick) noexcept {
+        ClearPendingReadTick(completed_tick);
+    }
+
+    void CompleteGpuWriteTick(u64 completed_tick) noexcept {
+        ClearPendingWriteTick(completed_tick);
+    }
+
+    [[nodiscard]] std::optional<u64> PendingGpuReadTick() const noexcept {
+        return PendingReadTick();
+    }
+
+    [[nodiscard]] std::optional<u64> PendingGpuWriteTick() const noexcept {
+        return PendingWriteTick();
+    }
+
     /// Returns true when the image is already initialized and mark it as initialized
     [[nodiscard]] bool ExchangeInitialization() noexcept {
         return std::exchange(initialized, true);
@@ -233,11 +268,20 @@ public:
 
     [[nodiscard]] VkImageView ColorView();
 
+    [[nodiscard]] VkImageView SampledView(Shader::TextureType texture_type,
+                                          Shader::SamplerComponentType component_type);
+
     [[nodiscard]] VkImageView StorageView(Shader::TextureType texture_type,
                                           Shader::ImageFormat image_format);
 
     [[nodiscard]] bool IsRescaled() const noexcept;
 
+    [[nodiscard]] bool SupportsDepthCompareSampling() const noexcept;
+
+    [[nodiscard]] bool Is3DImage() const noexcept {
+        return is_3d_image;
+    }
+
     [[nodiscard]] VkImageView Handle(Shader::TextureType texture_type) const noexcept {
         return *image_views[static_cast<size_t>(texture_type)];
     }
@@ -266,23 +310,36 @@ private:
     struct StorageViews {
         std::array<vk::ImageView, Shader::NUM_TEXTURE_TYPES> signeds;
         std::array<vk::ImageView, Shader::NUM_TEXTURE_TYPES> unsigneds;
+        std::array<vk::ImageView, Shader::NUM_TEXTURE_TYPES> typeless;
     };
 
+    static constexpr size_t NUMERIC_VIEW_TYPES = 3;
+
+    [[nodiscard]] Shader::TextureType BaseTextureType() const noexcept;
+    [[nodiscard]] std::optional<u32> LayerCountOverride(Shader::TextureType texture_type) const noexcept;
+    [[nodiscard]] VkImageView DepthView(Shader::TextureType texture_type);
+    [[nodiscard]] VkImageView StencilView(Shader::TextureType texture_type);
+
     [[nodiscard]] vk::ImageView MakeView(VkFormat vk_format, VkImageAspectFlags aspect_mask);
+    [[nodiscard]] vk::ImageView MakeView(VkFormat vk_format, VkImageAspectFlags aspect_mask,
+                                        Shader::TextureType texture_type);
 
     const Device* device = nullptr;
     const SlotVector<Image>* slot_images = nullptr;
 
     std::array<vk::ImageView, Shader::NUM_TEXTURE_TYPES> image_views;
     std::unique_ptr<StorageViews> storage_views;
-    vk::ImageView depth_view;
+    std::array<vk::ImageView, Shader::NUM_TEXTURE_TYPES> depth_views;
-    vk::ImageView stencil_view;
+    std::array<vk::ImageView, Shader::NUM_TEXTURE_TYPES> stencil_views;
+    std::array<std::array<vk::ImageView, Shader::NUM_TEXTURE_TYPES>, NUMERIC_VIEW_TYPES>
+        sampled_component_views;
     vk::ImageView color_view;
     vk::Image null_image;
     VkImage image_handle = VK_NULL_HANDLE;
     VkImageView render_target = VK_NULL_HANDLE;
     VkSampleCountFlagBits samples = VK_SAMPLE_COUNT_1_BIT;
     u32 buffer_size = 0;
+    bool is_3d_image = false;
 };
 
 class ImageAlloc : public VideoCommon::ImageAllocBase {};
@@ -303,9 +360,19 @@ public:
         return static_cast<bool>(sampler_default_anisotropy);
     }
 
+    [[nodiscard]] VkSampler SelectHandle(bool supports_linear_filter,
+                                         bool supports_anisotropy,
+                                         bool allow_depth_compare) const noexcept;
+
 private:
     vk::Sampler sampler;
     vk::Sampler sampler_default_anisotropy;
+    vk::Sampler sampler_force_point;
+    vk::Sampler sampler_compare_disabled;
+    vk::Sampler sampler_default_anisotropy_compare_disabled;
+    vk::Sampler sampler_force_point_compare_disabled;
+    bool uses_linear_filter = false;
+    bool depth_compare_enabled = false;
 };
 
 class Framebuffer {

src/video_core/shader_environment.cpp

@@ -70,6 +70,102 @@ static Shader::TexturePixelFormat ConvertTexturePixelFormat(const Tegra::Texture
                                    entry.a_type, entry.srgb_conversion));
 }
 
+namespace {
+
+[[nodiscard]] bool UsesSwizzleSource(const Tegra::Texture::TICEntry& entry,
+                                     Tegra::Texture::SwizzleSource source) {
+    const std::array swizzles{entry.x_source.Value(), entry.y_source.Value(),
+                              entry.z_source.Value(), entry.w_source.Value()};
+    return std::ranges::any_of(swizzles, [source](Tegra::Texture::SwizzleSource current) {
+        return current == source;
+    });
+}
+
+[[nodiscard]] std::optional<Shader::SamplerComponentType> DepthStencilComponentFromSwizzle(
+    const Tegra::Texture::TICEntry& entry, VideoCore::Surface::PixelFormat pixel_format) {
+    using Tegra::Texture::SwizzleSource;
+    const bool uses_r = UsesSwizzleSource(entry, SwizzleSource::R);
+    const bool uses_g = UsesSwizzleSource(entry, SwizzleSource::G);
+
+    switch (pixel_format) {
+    case VideoCore::Surface::PixelFormat::D24_UNORM_S8_UINT:
+    case VideoCore::Surface::PixelFormat::D32_FLOAT_S8_UINT:
+        if (uses_r != uses_g) {
+            return uses_r ? Shader::SamplerComponentType::Depth
+                          : Shader::SamplerComponentType::Stencil;
+        }
+        break;
+    case VideoCore::Surface::PixelFormat::S8_UINT_D24_UNORM:
+        if (uses_r != uses_g) {
+            return uses_r ? Shader::SamplerComponentType::Stencil
+                          : Shader::SamplerComponentType::Depth;
+        }
+        break;
+    default:
+        break;
+    }
+
+    return std::nullopt;
+}
+
+} // Anonymous namespace
+
+static Shader::SamplerComponentType ConvertSamplerComponentType(
+    const Tegra::Texture::TICEntry& entry) {
+    const auto pixel_format = PixelFormatFromTextureInfo(entry.format, entry.r_type, entry.g_type,
+                                                        entry.b_type, entry.a_type,
+                                                        entry.srgb_conversion);
+    const auto surface_type = VideoCore::Surface::GetFormatType(pixel_format);
+    if (entry.depth_texture != 0 || surface_type == VideoCore::Surface::SurfaceType::Depth) {
+        return Shader::SamplerComponentType::Depth;
+    }
+    if (surface_type == VideoCore::Surface::SurfaceType::Stencil) {
+        return Shader::SamplerComponentType::Stencil;
+    }
+    if (surface_type == VideoCore::Surface::SurfaceType::DepthStencil) {
+        if (const auto inferred = DepthStencilComponentFromSwizzle(entry, pixel_format)) {
+            return *inferred;
+        }
+        return entry.depth_texture != 0 ? Shader::SamplerComponentType::Depth
+                                        : Shader::SamplerComponentType::Stencil;
+    }
+
+    const auto accumulate = [](const Tegra::Texture::ComponentType component,
+                               bool& has_signed, bool& has_unsigned) {
+        switch (component) {
+        case Tegra::Texture::ComponentType::SINT:
+            has_signed = true;
+            break;
+        case Tegra::Texture::ComponentType::UINT:
+            has_unsigned = true;
+            break;
+        default:
+            break;
+        }
+    };
+
+    bool has_signed{};
+    bool has_unsigned{};
+    accumulate(entry.r_type, has_signed, has_unsigned);
+    accumulate(entry.g_type, has_signed, has_unsigned);
+    accumulate(entry.b_type, has_signed, has_unsigned);
+    accumulate(entry.a_type, has_signed, has_unsigned);
+
+    if (has_signed && !has_unsigned) {
+        return Shader::SamplerComponentType::Sint;
+    }
+    if (has_unsigned && !has_signed) {
+        return Shader::SamplerComponentType::Uint;
+    }
+    if (has_signed) {
+        return Shader::SamplerComponentType::Sint;
+    }
+    if (has_unsigned) {
+        return Shader::SamplerComponentType::Uint;
+    }
+    return Shader::SamplerComponentType::Float;
+}
+
 static std::string_view StageToPrefix(Shader::Stage stage) {
     switch (stage) {
     case Shader::Stage::VertexB:
@@ -200,6 +296,7 @@ void GenericEnvironment::Serialize(std::ofstream& file) const {
     const u64 code_size{static_cast<u64>(CachedSizeBytes())};
     const u64 num_texture_types{static_cast<u64>(texture_types.size())};
     const u64 num_texture_pixel_formats{static_cast<u64>(texture_pixel_formats.size())};
+    const u64 num_texture_component_types{static_cast<u64>(texture_component_types.size())};
     const u64 num_cbuf_values{static_cast<u64>(cbuf_values.size())};
     const u64 num_cbuf_replacement_values{static_cast<u64>(cbuf_replacements.size())};
 
@@ -207,6 +304,8 @@ void GenericEnvironment::Serialize(std::ofstream& file) const {
         .write(reinterpret_cast<const char*>(&num_texture_types), sizeof(num_texture_types))
         .write(reinterpret_cast<const char*>(&num_texture_pixel_formats),
                sizeof(num_texture_pixel_formats))
+        .write(reinterpret_cast<const char*>(&num_texture_component_types),
+               sizeof(num_texture_component_types))
         .write(reinterpret_cast<const char*>(&num_cbuf_values), sizeof(num_cbuf_values))
         .write(reinterpret_cast<const char*>(&num_cbuf_replacement_values),
                sizeof(num_cbuf_replacement_values))
@@ -223,6 +322,10 @@ void GenericEnvironment::Serialize(std::ofstream& file) const {
         file.write(reinterpret_cast<const char*>(&key), sizeof(key))
             .write(reinterpret_cast<const char*>(&type), sizeof(type));
     }
+    for (const auto& [key, component] : texture_component_types) {
+        file.write(reinterpret_cast<const char*>(&key), sizeof(key))
+            .write(reinterpret_cast<const char*>(&component), sizeof(component));
+    }
     for (const auto& [key, format] : texture_pixel_formats) {
         file.write(reinterpret_cast<const char*>(&key), sizeof(key))
             .write(reinterpret_cast<const char*>(&format), sizeof(format));
@@ -277,7 +380,17 @@ std::optional<u64> GenericEnvironment::TryFindSize() {
 Tegra::Texture::TICEntry GenericEnvironment::ReadTextureInfo(GPUVAddr tic_addr, u32 tic_limit,
                                                              bool via_header_index, u32 raw) {
     const auto handle{Tegra::Texture::TexturePair(raw, via_header_index)};
-    ASSERT(handle.first <= tic_limit);
+     if (handle.first > tic_limit) {
+        static std::atomic<size_t> oob_count{0};
+        const size_t n = ++oob_count;
+        if (n <= 4 || (n & 63) == 0) {
+            LOG_WARNING(Shader,
+                        "TIC handle {} exceeds limit {} (via_header_index={}) — returning empty",
+                        handle.first, tic_limit, via_header_index);
+        }
+        return {};
+    }
+
     const GPUVAddr descriptor_addr{tic_addr + handle.first * sizeof(Tegra::Texture::TICEntry)};
     Tegra::Texture::TICEntry entry;
     gpu_memory->ReadBlock(descriptor_addr, &entry, sizeof(entry));
@@ -374,6 +487,21 @@ Shader::TextureType GraphicsEnvironment::ReadTextureType(u32 handle) {
         ReadTextureInfo(regs.tex_header.Address(), regs.tex_header.limit, via_header_index, handle);
     const Shader::TextureType result{ConvertTextureType(entry)};
     texture_types.emplace(handle, result);
+    texture_component_types.emplace(handle, ConvertSamplerComponentType(entry));
+    return result;
+}
+
+Shader::SamplerComponentType GraphicsEnvironment::ReadTextureComponentType(u32 handle) {
+    const auto it{texture_component_types.find(handle)};
+    if (it != texture_component_types.end()) {
+        return it->second;
+    }
+    const auto& regs{maxwell3d->regs};
+    const bool via_header_index{regs.sampler_binding == Maxwell::SamplerBinding::ViaHeaderBinding};
+    auto entry =
+        ReadTextureInfo(regs.tex_header.Address(), regs.tex_header.limit, via_header_index, handle);
+    const Shader::SamplerComponentType result{ConvertSamplerComponentType(entry)};
+    texture_component_types.emplace(handle, result);
     return result;
 }
 
@@ -430,6 +558,20 @@ Shader::TextureType ComputeEnvironment::ReadTextureType(u32 handle) {
     auto entry = ReadTextureInfo(regs.tic.Address(), regs.tic.limit, qmd.linked_tsc != 0, handle);
     const Shader::TextureType result{ConvertTextureType(entry)};
     texture_types.emplace(handle, result);
+    texture_component_types.emplace(handle, ConvertSamplerComponentType(entry));
+    return result;
+}
+
+Shader::SamplerComponentType ComputeEnvironment::ReadTextureComponentType(u32 handle) {
+    const auto it{texture_component_types.find(handle)};
+    if (it != texture_component_types.end()) {
+        return it->second;
+    }
+    const auto& regs{kepler_compute->regs};
+    const auto& qmd{kepler_compute->launch_description};
+    auto entry = ReadTextureInfo(regs.tic.Address(), regs.tic.limit, qmd.linked_tsc != 0, handle);
+    const Shader::SamplerComponentType result{ConvertSamplerComponentType(entry)};
+    texture_component_types.emplace(handle, result);
     return result;
 }
 
@@ -455,12 +597,15 @@ void FileEnvironment::Deserialize(std::ifstream& file) {
     u64 code_size{};
     u64 num_texture_types{};
     u64 num_texture_pixel_formats{};
+    u64 num_texture_component_types{};
     u64 num_cbuf_values{};
     u64 num_cbuf_replacement_values{};
     file.read(reinterpret_cast<char*>(&code_size), sizeof(code_size))
         .read(reinterpret_cast<char*>(&num_texture_types), sizeof(num_texture_types))
-        .read(reinterpret_cast<char*>(&num_texture_pixel_formats),
+          .read(reinterpret_cast<char*>(&num_texture_pixel_formats),
               sizeof(num_texture_pixel_formats))
+          .read(reinterpret_cast<char*>(&num_texture_component_types),
+              sizeof(num_texture_component_types))
         .read(reinterpret_cast<char*>(&num_cbuf_values), sizeof(num_cbuf_values))
         .read(reinterpret_cast<char*>(&num_cbuf_replacement_values),
               sizeof(num_cbuf_replacement_values))
@@ -480,6 +625,13 @@ void FileEnvironment::Deserialize(std::ifstream& file) {
             .read(reinterpret_cast<char*>(&type), sizeof(type));
         texture_types.emplace(key, type);
     }
+    for (size_t i = 0; i < num_texture_component_types; ++i) {
+        u32 key;
+        Shader::SamplerComponentType component;
+        file.read(reinterpret_cast<char*>(&key), sizeof(key))
+            .read(reinterpret_cast<char*>(&component), sizeof(component));
+        texture_component_types.emplace(key, component);
+    }
     for (size_t i = 0; i < num_texture_pixel_formats; ++i) {
         u32 key;
         Shader::TexturePixelFormat format;
@@ -534,6 +686,15 @@ u32 FileEnvironment::ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) {
     return it->second;
 }
 
+Shader::SamplerComponentType FileEnvironment::ReadTextureComponentType(u32 handle) {
+    const auto it{texture_component_types.find(handle)};
+    if (it == texture_component_types.end()) {
+        LOG_WARNING(Render_Vulkan, "Texture component descriptor {:08x} not found", handle);
+        return Shader::SamplerComponentType::Float;
+    }
+    return it->second;
+}
+
 Shader::TextureType FileEnvironment::ReadTextureType(u32 handle) {
     const auto it{texture_types.find(handle)};
     if (it == texture_types.end()) {

src/video_core/shader_environment.h

@@ -80,6 +80,7 @@ protected:
 
     std::vector<u64> code;
     std::unordered_map<u32, Shader::TextureType> texture_types;
+    std::unordered_map<u32, Shader::SamplerComponentType> texture_component_types;
     std::unordered_map<u32, Shader::TexturePixelFormat> texture_pixel_formats;
     std::unordered_map<u64, u32> cbuf_values;
     std::unordered_map<u64, Shader::ReplaceConstant> cbuf_replacements;
@@ -116,6 +117,8 @@ public:
 
     Shader::TextureType ReadTextureType(u32 handle) override;
 
+    Shader::SamplerComponentType ReadTextureComponentType(u32 handle) override;
+
     Shader::TexturePixelFormat ReadTexturePixelFormat(u32 handle) override;
 
     bool IsTexturePixelFormatInteger(u32 handle) override;
@@ -142,6 +145,8 @@ public:
 
     Shader::TextureType ReadTextureType(u32 handle) override;
 
+    Shader::SamplerComponentType ReadTextureComponentType(u32 handle) override;
+
     Shader::TexturePixelFormat ReadTexturePixelFormat(u32 handle) override;
 
     bool IsTexturePixelFormatInteger(u32 handle) override;
@@ -176,6 +181,8 @@ public:
 
     [[nodiscard]] Shader::TextureType ReadTextureType(u32 handle) override;
 
+    [[nodiscard]] Shader::SamplerComponentType ReadTextureComponentType(u32 handle) override;
+
     [[nodiscard]] Shader::TexturePixelFormat ReadTexturePixelFormat(u32 handle) override;
 
     [[nodiscard]] bool IsTexturePixelFormatInteger(u32 handle) override;
@@ -202,6 +209,7 @@ public:
 private:
     std::vector<u64> code;
     std::unordered_map<u32, Shader::TextureType> texture_types;
+    std::unordered_map<u32, Shader::SamplerComponentType> texture_component_types;
     std::unordered_map<u32, Shader::TexturePixelFormat> texture_pixel_formats;
     std::unordered_map<u64, u32> cbuf_values;
     std::unordered_map<u64, Shader::ReplaceConstant> cbuf_replacements;

src/video_core/surface.cpp

@@ -4,6 +4,8 @@
 // SPDX-FileCopyrightText: 2014 Citra Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
+#include <optional>
+
 #include "common/common_types.h"
 #include "common/math_util.h"
 #include "common/settings.h"
@@ -408,6 +410,126 @@ bool IsPixelFormatSignedInteger(PixelFormat format) {
     }
 }
 
+namespace {
+
+struct NumericVariantSet {
+    PixelFormat float_format = PixelFormat::Invalid;
+    PixelFormat uint_format = PixelFormat::Invalid;
+    PixelFormat sint_format = PixelFormat::Invalid;
+
+    [[nodiscard]] std::optional<PixelFormat> Select(PixelFormatNumeric numeric) const {
+        PixelFormat candidate = PixelFormat::Invalid;
+        switch (numeric) {
+        case PixelFormatNumeric::Float:
+            candidate = float_format;
+            break;
+        case PixelFormatNumeric::Uint:
+            candidate = uint_format;
+            break;
+        case PixelFormatNumeric::Sint:
+            candidate = sint_format;
+            break;
+        }
+        if (candidate == PixelFormat::Invalid) {
+            return std::nullopt;
+        }
+        return candidate;
+    }
+};
+
+constexpr NumericVariantSet MakeVariant(PixelFormat float_format, PixelFormat uint_format,
+                                        PixelFormat sint_format) {
+    return NumericVariantSet{
+        .float_format = float_format,
+        .uint_format = uint_format,
+        .sint_format = sint_format,
+    };
+}
+
+std::optional<NumericVariantSet> LookupNumericVariantSet(PixelFormat format) {
+    switch (format) {
+    case PixelFormat::R8_UNORM:
+    case PixelFormat::R8_SNORM:
+    case PixelFormat::R8_UINT:
+    case PixelFormat::R8_SINT:
+        return MakeVariant(PixelFormat::R8_UNORM, PixelFormat::R8_UINT, PixelFormat::R8_SINT);
+    case PixelFormat::R16_FLOAT:
+    case PixelFormat::R16_UNORM:
+    case PixelFormat::R16_SNORM:
+    case PixelFormat::R16_UINT:
+    case PixelFormat::R16_SINT:
+        return MakeVariant(PixelFormat::R16_FLOAT, PixelFormat::R16_UINT, PixelFormat::R16_SINT);
+    case PixelFormat::R32_FLOAT:
+    case PixelFormat::R32_UINT:
+    case PixelFormat::R32_SINT:
+        return MakeVariant(PixelFormat::R32_FLOAT, PixelFormat::R32_UINT, PixelFormat::R32_SINT);
+    case PixelFormat::R8G8_UNORM:
+    case PixelFormat::R8G8_SNORM:
+    case PixelFormat::R8G8_UINT:
+    case PixelFormat::R8G8_SINT:
+        return MakeVariant(PixelFormat::R8G8_UNORM, PixelFormat::R8G8_UINT, PixelFormat::R8G8_SINT);
+    case PixelFormat::R16G16_FLOAT:
+    case PixelFormat::R16G16_UNORM:
+    case PixelFormat::R16G16_SNORM:
+    case PixelFormat::R16G16_UINT:
+    case PixelFormat::R16G16_SINT:
+        return MakeVariant(PixelFormat::R16G16_FLOAT, PixelFormat::R16G16_UINT,
+                           PixelFormat::R16G16_SINT);
+    case PixelFormat::R32G32_FLOAT:
+    case PixelFormat::R32G32_UINT:
+    case PixelFormat::R32G32_SINT:
+        return MakeVariant(PixelFormat::R32G32_FLOAT, PixelFormat::R32G32_UINT,
+                           PixelFormat::R32G32_SINT);
+    case PixelFormat::R16G16B16A16_FLOAT:
+    case PixelFormat::R16G16B16A16_UNORM:
+    case PixelFormat::R16G16B16A16_SNORM:
+    case PixelFormat::R16G16B16A16_UINT:
+    case PixelFormat::R16G16B16A16_SINT:
+        return MakeVariant(PixelFormat::R16G16B16A16_FLOAT, PixelFormat::R16G16B16A16_UINT,
+                           PixelFormat::R16G16B16A16_SINT);
+    case PixelFormat::R32G32B32A32_FLOAT:
+    case PixelFormat::R32G32B32A32_UINT:
+    case PixelFormat::R32G32B32A32_SINT:
+        return MakeVariant(PixelFormat::R32G32B32A32_FLOAT, PixelFormat::R32G32B32A32_UINT,
+                           PixelFormat::R32G32B32A32_SINT);
+    case PixelFormat::A8B8G8R8_UNORM:
+    case PixelFormat::A8B8G8R8_SNORM:
+    case PixelFormat::A8B8G8R8_SRGB:
+    case PixelFormat::A8B8G8R8_UINT:
+    case PixelFormat::A8B8G8R8_SINT:
+        return MakeVariant(PixelFormat::A8B8G8R8_UNORM, PixelFormat::A8B8G8R8_UINT,
+                           PixelFormat::A8B8G8R8_SINT);
+    case PixelFormat::A2B10G10R10_UNORM:
+    case PixelFormat::A2B10G10R10_UINT:
+        return MakeVariant(PixelFormat::A2B10G10R10_UNORM, PixelFormat::A2B10G10R10_UINT,
+                           PixelFormat::Invalid);
+    default:
+        return std::nullopt;
+    }
+}
+
+} // Anonymous namespace
+
+PixelFormatNumeric GetPixelFormatNumericType(PixelFormat format) {
+    if (IsPixelFormatInteger(format)) {
+        return IsPixelFormatSignedInteger(format) ? PixelFormatNumeric::Sint
+                                                  : PixelFormatNumeric::Uint;
+    }
+    return PixelFormatNumeric::Float;
+}
+
+std::optional<PixelFormat> FindPixelFormatVariant(PixelFormat format,
+                                                  PixelFormatNumeric target_numeric) {
+    const auto variants = LookupNumericVariantSet(format);
+    if (!variants) {
+        return std::nullopt;
+    }
+    if (const auto candidate = variants->Select(target_numeric)) {
+        return candidate;
+    }
+    return std::nullopt;
+}
+
 size_t PixelComponentSizeBitsInteger(PixelFormat format) {
     switch (format) {
     case PixelFormat::A8B8G8R8_SINT:

src/video_core/surface.h

@@ -1,11 +1,13 @@
 // SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
 // SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: 2014 Citra Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
 #pragma once
 
 #include <climits>
+#include <optional>
 #include <utility>
 #include "common/assert.h"
 #include "common/common_types.h"
@@ -202,6 +204,17 @@ bool IsPixelFormatSRGB(PixelFormat format);
 bool IsPixelFormatInteger(PixelFormat format);
 bool IsPixelFormatSignedInteger(PixelFormat format);
 size_t PixelComponentSizeBitsInteger(PixelFormat format);
+
+enum class PixelFormatNumeric {
+    Float,
+    Uint,
+    Sint,
+};
+
+PixelFormatNumeric GetPixelFormatNumericType(PixelFormat format);
+std::optional<PixelFormat> FindPixelFormatVariant(PixelFormat format,
+                                                  PixelFormatNumeric target_numeric);
+
 std::pair<u32, u32> GetASTCBlockSize(PixelFormat format);
 u64 TranscodedAstcSize(u64 base_size, PixelFormat format);
 

src/video_core/texture_cache/format_lookup_table.cpp

@@ -1,3 +1,6 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2019 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
@@ -21,6 +24,27 @@ constexpr auto FLOAT = ComponentType::FLOAT;
 constexpr bool LINEAR = false;
 constexpr bool SRGB = true;
 
+constexpr TextureFormat SanitizeFormat(TextureFormat format) {
+    if (format == static_cast<TextureFormat>(0)) {
+        return TextureFormat::A8B8G8R8;
+    }
+    return format;
+}
+
+constexpr ComponentType SanitizeComponent(ComponentType component) {
+    if (component == static_cast<ComponentType>(0)) {
+        return ComponentType::UNORM;
+    }
+    switch (component) {
+    case ComponentType::SNORM_FORCE_FP16:
+        return ComponentType::SNORM;
+    case ComponentType::UNORM_FORCE_FP16:
+        return ComponentType::UNORM;
+    default:
+        return component;
+    }
+}
+
 constexpr u32 Hash(TextureFormat format, ComponentType red_component, ComponentType green_component,
                    ComponentType blue_component, ComponentType alpha_component, bool is_srgb) {
     u32 hash = is_srgb ? 1 : 0;
@@ -41,6 +65,11 @@ constexpr u32 Hash(TextureFormat format, ComponentType component, bool is_srgb =
 PixelFormat PixelFormatFromTextureInfo(TextureFormat format, ComponentType red, ComponentType green,
                                        ComponentType blue, ComponentType alpha,
                                        bool is_srgb) noexcept {
+    format = SanitizeFormat(format);
+    red = SanitizeComponent(red);
+    green = SanitizeComponent(green);
+    blue = SanitizeComponent(blue);
+    alpha = SanitizeComponent(alpha);
     switch (Hash(format, red, green, blue, alpha, is_srgb)) {
     case Hash(TextureFormat::A8B8G8R8, UNORM):
         return PixelFormat::A8B8G8R8_UNORM;

src/video_core/texture_cache/image_base.h

@@ -1,8 +1,12 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2020 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
 #pragma once
 
+#include <algorithm>
 #include <array>
 #include <optional>
 #include <vector>
@@ -58,6 +62,50 @@ struct ImageBase {
     explicit ImageBase(const ImageInfo& info, GPUVAddr gpu_addr, VAddr cpu_addr);
     explicit ImageBase(const NullImageParams&);
 
+    void TrackPendingReadTick(u64 tick) noexcept {
+        if (pending_read_tick) {
+            *pending_read_tick = std::max(*pending_read_tick, tick);
+        } else {
+            pending_read_tick = tick;
+        }
+    }
+
+    void TrackPendingWriteTick(u64 tick) noexcept {
+        if (pending_write_tick) {
+            *pending_write_tick = std::max(*pending_write_tick, tick);
+        } else {
+            pending_write_tick = tick;
+        }
+    }
+
+    void ClearPendingReadTick(u64 completed_tick) noexcept {
+        if (pending_read_tick && completed_tick >= *pending_read_tick) {
+            pending_read_tick.reset();
+        }
+    }
+
+    void ClearPendingWriteTick(u64 completed_tick) noexcept {
+        if (pending_write_tick && completed_tick >= *pending_write_tick) {
+            pending_write_tick.reset();
+        }
+    }
+
+    [[nodiscard]] bool HasPendingReadTick() const noexcept {
+        return pending_read_tick.has_value();
+    }
+
+    [[nodiscard]] bool HasPendingWriteTick() const noexcept {
+        return pending_write_tick.has_value();
+    }
+
+    [[nodiscard]] std::optional<u64> PendingReadTick() const noexcept {
+        return pending_read_tick;
+    }
+
+    [[nodiscard]] std::optional<u64> PendingWriteTick() const noexcept {
+        return pending_write_tick;
+    }
+
     [[nodiscard]] std::optional<SubresourceBase> TryFindBase(GPUVAddr other_addr) const noexcept;
 
     [[nodiscard]] ImageViewId FindView(const ImageViewInfo& view_info) const noexcept;
@@ -115,6 +163,9 @@ struct ImageBase {
     std::vector<AliasedImage> aliased_images;
     std::vector<ImageId> overlapping_images;
     ImageMapId map_view_id{};
+
+    std::optional<u64> pending_read_tick;
+    std::optional<u64> pending_write_tick;
 };
 
 struct ImageMapView {

src/video_core/texture_cache/samples_helper.h

@@ -1,15 +1,22 @@
+// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
 // SPDX-FileCopyrightText: Copyright 2020 yuzu Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
 #pragma once
 
+#include <cstddef>
 #include <utility>
 
 #include "common/assert.h"
+#include "common/common_types.h"
 #include "video_core/textures/texture.h"
 
 namespace VideoCommon {
 
+constexpr inline std::size_t MaxSampleLocationSlots = 16;
+
 [[nodiscard]] inline std::pair<int, int> SamplesLog2(int num_samples) {
     switch (num_samples) {
     case 1:
@@ -95,4 +102,24 @@ namespace VideoCommon {
     return 1;
 }
 
+// NVN guarantees up to sixteen programmable sample slots shared across a repeating pixel grid
+// (per the 0.7.0 addon release notes), so the grid dimensions shrink as MSAA increases.
+[[nodiscard]] inline std::pair<u32, u32> SampleLocationGridSize(Tegra::Texture::MsaaMode msaa_mode) {
+    const int samples = NumSamples(msaa_mode);
+    switch (samples) {
+    case 1:
+        return {4u, 4u};
+    case 2:
+        return {4u, 2u};
+    case 4:
+        return {2u, 2u};
+    case 8:
+        return {2u, 1u};
+    case 16:
+        return {1u, 1u};
+    }
+    ASSERT_MSG(false, "Unsupported sample count for grid size={}", samples);
+    return {1u, 1u};
+}
+
 } // namespace VideoCommon

src/video_core/texture_cache/texture_cache.h

@@ -6,6 +6,8 @@
 
 #pragma once
 
+#include <limits>
+#include <optional>
 #include <unordered_set>
 #include <boost/container/small_vector.hpp>
 
@@ -538,6 +540,7 @@ void TextureCache<P>::WriteMemory(DAddr cpu_addr, size_t size) {
         if (True(image.flags & ImageFlagBits::CpuModified)) {
             return;
         }
+        EnsureImageReady(image, ImageAccessType::Write);
         image.flags |= ImageFlagBits::CpuModified;
         if (True(image.flags & ImageFlagBits::Tracked)) {
             UntrackImage(image, image_id);
@@ -548,11 +551,12 @@ void TextureCache<P>::WriteMemory(DAddr cpu_addr, size_t size) {
 template <class P>
 void TextureCache<P>::DownloadMemory(DAddr cpu_addr, size_t size) {
     boost::container::small_vector<ImageId, 16> images;
-    ForEachImageInRegion(cpu_addr, size, [&images](ImageId image_id, ImageBase& image) {
+    ForEachImageInRegion(cpu_addr, size, [this, &images](ImageId image_id, ImageBase& image) {
         if (!image.IsSafeDownload()) {
             return;
         }
         image.flags &= ~ImageFlagBits::GpuModified;
+        EnsureImageReady(this->slot_images[image_id], ImageAccessType::Read);
         images.push_back(image_id);
     });
     if (images.empty()) {
@@ -604,6 +608,7 @@ void TextureCache<P>::UnmapMemory(DAddr cpu_addr, size_t size) {
     ForEachImageInRegion(cpu_addr, size, [&](ImageId id, Image&) { deleted_images.push_back(id); });
     for (const ImageId id : deleted_images) {
         Image& image = slot_images[id];
+        EnsureImageReady(image, ImageAccessType::Write);
         if (True(image.flags & ImageFlagBits::Tracked)) {
             UntrackImage(image, id);
         }
@@ -619,6 +624,7 @@ void TextureCache<P>::UnmapGPUMemory(size_t as_id, GPUVAddr gpu_addr, size_t siz
                             [&](ImageId id, Image&) { deleted_images.push_back(id); });
     for (const ImageId id : deleted_images) {
         Image& image = slot_images[id];
+        EnsureImageReady(image, ImageAccessType::Write);
         if (False(image.flags & ImageFlagBits::CpuModified)) {
             image.flags |= ImageFlagBits::CpuModified;
             if (True(image.flags & ImageFlagBits::Tracked)) {
@@ -1736,11 +1742,89 @@ SamplerId TextureCache<P>::FindSampler(const TSCEntry& config) {
     }
     const auto [pair, is_new] = channel_state->samplers.try_emplace(config);
     if (is_new) {
+        EnforceSamplerBudget();
         pair->second = slot_samplers.insert(runtime, config);
     }
     return pair->second;
 }
 
+template <class P>
+std::optional<size_t> TextureCache<P>::QuerySamplerBudget() const {
+    if constexpr (requires { runtime.GetSamplerHeapBudget(); }) {
+        return runtime.GetSamplerHeapBudget();
+    } else {
+        return std::nullopt;
+    }
+}
+
+template <class P>
+void TextureCache<P>::EnforceSamplerBudget() {
+    const auto budget = QuerySamplerBudget();
+    if (!budget) {
+        return;
+    }
+    if (slot_samplers.size() < *budget) {
+        return;
+    }
+    if (!channel_state) {
+        return;
+    }
+    if (last_sampler_gc_frame == frame_tick) {
+        return;
+    }
+    last_sampler_gc_frame = frame_tick;
+    TrimInactiveSamplers(*budget);
+}
+
+template <class P>
+void TextureCache<P>::TrimInactiveSamplers(size_t budget) {
+    if (channel_state->samplers.empty()) {
+        return;
+    }
+    static constexpr size_t SAMPLER_GC_SLACK = 1024;
+    auto mark_active = [](auto& set, SamplerId id) {
+        if (!id || id == CORRUPT_ID || id == NULL_SAMPLER_ID) {
+            return;
+        }
+        set.insert(id);
+    };
+    std::unordered_set<SamplerId> active;
+    active.reserve(channel_state->graphics_sampler_ids.size() +
+                   channel_state->compute_sampler_ids.size());
+    for (const SamplerId id : channel_state->graphics_sampler_ids) {
+        mark_active(active, id);
+    }
+    for (const SamplerId id : channel_state->compute_sampler_ids) {
+        mark_active(active, id);
+    }
+
+    size_t removed = 0;
+    auto& sampler_map = channel_state->samplers;
+    for (auto it = sampler_map.begin(); it != sampler_map.end();) {
+        const SamplerId sampler_id = it->second;
+        if (!sampler_id || sampler_id == CORRUPT_ID) {
+            it = sampler_map.erase(it);
+            continue;
+        }
+        if (active.find(sampler_id) != active.end()) {
+            ++it;
+            continue;
+        }
+        slot_samplers.erase(sampler_id);
+        it = sampler_map.erase(it);
+        ++removed;
+        if (slot_samplers.size() + SAMPLER_GC_SLACK <= budget) {
+            break;
+        }
+    }
+
+    if (removed != 0) {
+        LOG_WARNING(HW_GPU,
+                    "Sampler cache exceeded {} entries on this driver; reclaimed {} inactive samplers",
+                    budget, removed);
+    }
+}
+
 template <class P>
 ImageViewId TextureCache<P>::FindColorBuffer(size_t index) {
     const auto& regs = maxwell3d->regs;
@@ -2343,6 +2427,9 @@ void TextureCache<P>::SynchronizeAliases(ImageId image_id) {
 template <class P>
 void TextureCache<P>::PrepareImage(ImageId image_id, bool is_modification, bool invalidate) {
     Image& image = slot_images[image_id];
+    EnsureImageReady(image, is_modification ? ImageAccessType::Write : ImageAccessType::Read);
+    TrackGpuImageAccess(image_id, is_modification ? ImageAccessType::Write : ImageAccessType::Read);
+    runtime.TransitionImageLayout(image);
     if (invalidate) {
         image.flags &= ~(ImageFlagBits::CpuModified | ImageFlagBits::GpuModified);
         if (False(image.flags & ImageFlagBits::Tracked)) {
@@ -2358,6 +2445,85 @@ void TextureCache<P>::PrepareImage(ImageId image_id, bool is_modification, bool
     lru_cache.Touch(image.lru_index, frame_tick);
 }
 
+template <class P>
+void TextureCache<P>::TrackGpuImageAccess(ImageId image_id, ImageAccessType access) {
+    staged_gpu_accesses.push_back({image_id, access});
+}
+
+template <class P>
+void TextureCache<P>::CommitPendingGpuAccesses(u64 fence_value) {
+    if (staged_gpu_accesses.empty()) {
+        return;
+    }
+    if (fence_value == 0) {
+        return;
+    }
+    auto& batch = committed_gpu_accesses.emplace_back(std::move(staged_gpu_accesses));
+    staged_gpu_accesses.clear();
+    committed_gpu_ticks.push_back(fence_value);
+    for (const PendingImageAccess& access : batch) {
+        ImageBase& image = slot_images[access.image_id];
+        if (access.access == ImageAccessType::Read) {
+            image.TrackPendingReadTick(fence_value);
+        } else {
+            image.TrackPendingWriteTick(fence_value);
+        }
+    }
+}
+
+template <class P>
+void TextureCache<P>::CompleteGpuAccesses(u64 completed_fence) {
+    if (completed_fence == 0) {
+        return;
+    }
+    while (!committed_gpu_ticks.empty() && committed_gpu_ticks.front() <= completed_fence) {
+        auto accesses = std::move(committed_gpu_accesses.front());
+        committed_gpu_accesses.pop_front();
+        committed_gpu_ticks.pop_front();
+        for (const PendingImageAccess& access : accesses) {
+            if (!slot_images.Contains(access.image_id)) {
+                continue;
+            }
+            ImageBase& image = slot_images[access.image_id];
+            if (access.access == ImageAccessType::Read) {
+                image.ClearPendingReadTick(completed_fence);
+            } else {
+                image.ClearPendingWriteTick(completed_fence);
+            }
+        }
+    }
+}
+
+template <class P>
+void TextureCache<P>::EnsureImageReady(ImageBase& image, ImageAccessType access) {
+    auto wait_tick = [this](std::optional<u64> tick) -> std::optional<u64> {
+        if (!tick) {
+            return std::nullopt;
+        }
+        runtime.WaitForGpuTick(*tick);
+        return tick;
+    };
+
+    if (access == ImageAccessType::Write) {
+        if (const auto tick = image.PendingReadTick()) {
+            if (const auto waited = wait_tick(tick)) {
+                image.ClearPendingReadTick(*waited);
+            }
+        }
+        if (const auto tick = image.PendingWriteTick()) {
+            if (const auto waited = wait_tick(tick)) {
+                image.ClearPendingWriteTick(*waited);
+            }
+        }
+    } else {
+        if (const auto tick = image.PendingWriteTick()) {
+            if (const auto waited = wait_tick(tick)) {
+                image.ClearPendingWriteTick(*waited);
+            }
+        }
+    }
+}
+
 template <class P>
 void TextureCache<P>::PrepareImageView(ImageViewId image_view_id, bool is_modification,
                                        bool invalidate) {
@@ -2375,6 +2541,8 @@ template <class P>
 void TextureCache<P>::CopyImage(ImageId dst_id, ImageId src_id, std::vector<ImageCopy> copies) {
     Image& dst = slot_images[dst_id];
     Image& src = slot_images[src_id];
+    EnsureImageReady(dst, ImageAccessType::Write);
+    EnsureImageReady(src, ImageAccessType::Read);
     const bool is_rescaled = True(src.flags & ImageFlagBits::Rescaled);
     if (is_rescaled) {
         ASSERT(True(dst.flags & ImageFlagBits::Rescaled));
@@ -2391,20 +2559,30 @@ void TextureCache<P>::CopyImage(ImageId dst_id, ImageId src_id, std::vector<Imag
             }
         }
     }
+    const auto TrackCopyAccesses = [this, dst_id, src_id]() {
+        TrackGpuImageAccess(dst_id, ImageAccessType::Write);
+        TrackGpuImageAccess(src_id, ImageAccessType::Read);
+    };
     const auto dst_format_type = GetFormatType(dst.info.format);
     const auto src_format_type = GetFormatType(src.info.format);
     if (src_format_type == dst_format_type) {
         if constexpr (HAS_EMULATED_COPIES) {
             if (!runtime.CanImageBeCopied(dst, src)) {
-                return runtime.EmulateCopyImage(dst, src, copies);
+                runtime.EmulateCopyImage(dst, src, copies);
+                TrackCopyAccesses();
+                return;
             }
         }
-        return runtime.CopyImage(dst, src, copies);
+        runtime.CopyImage(dst, src, copies);
+        TrackCopyAccesses();
+        return;
     }
     UNIMPLEMENTED_IF(dst.info.type != ImageType::e2D);
     UNIMPLEMENTED_IF(src.info.type != ImageType::e2D);
     if (runtime.ShouldReinterpret(dst, src)) {
-        return runtime.ReinterpretImage(dst, src, copies);
+        runtime.ReinterpretImage(dst, src, copies);
+        TrackCopyAccesses();
+        return;
     }
     for (const ImageCopy& copy : copies) {
         UNIMPLEMENTED_IF(copy.dst_subresource.num_layers != 1);
@@ -2457,6 +2635,7 @@ void TextureCache<P>::CopyImage(ImageId dst_id, ImageId src_id, std::vector<Imag
 
         runtime.ConvertImage(dst_framebuffer, dst_view, src_view);
     }
+    TrackCopyAccesses();
 }
 
 template <class P>

src/video_core/texture_cache/texture_cache_base.h

@@ -182,6 +182,11 @@ public:
     /// Return a reference to the given sampler id
     [[nodiscard]] Sampler& GetSampler(SamplerId id) noexcept;
 
+    /// Returns true when the runtime format supports linear filtering
+    [[nodiscard]] bool SupportsLinearFilter(PixelFormat format) const noexcept {
+        return runtime.SupportsLinearFilter(format);
+    }
+
     /// Refresh the state for graphics image view and sampler descriptors
     void SynchronizeGraphicsDescriptors();
 
@@ -258,6 +263,15 @@ public:
     /// Prepare an image to be used
     void PrepareImage(ImageId image_id, bool is_modification, bool invalidate);
 
+    /// Track that an image participates in upcoming GPU work with the given access type
+    void TrackGpuImageAccess(ImageId image_id, ImageAccessType access);
+
+    /// Notify the cache that tracked GPU work has been submitted with the specified fence value
+    void CommitPendingGpuAccesses(u64 fence_value);
+
+    /// Notify the cache that a fence value has completed so tracked accesses can be released
+    void CompleteGpuAccesses(u64 completed_fence);
+
     std::recursive_mutex mutex;
 
 private:
@@ -408,6 +422,8 @@ private:
     /// Execute copies from one image to the other, even if they are incompatible
     void CopyImage(ImageId dst_id, ImageId src_id, std::vector<ImageCopy> copies);
 
+    void EnsureImageReady(ImageBase& image, ImageAccessType access);
+
     /// Bind an image view as render target, downloading resources preemtively if needed
     void BindRenderTarget(ImageViewId* old_id, ImageViewId new_id);
 
@@ -429,6 +445,9 @@ private:
 
     void QueueAsyncDecode(Image& image, ImageId image_id);
     void TickAsyncDecode();
+    void EnforceSamplerBudget();
+    void TrimInactiveSamplers(size_t budget);
+    std::optional<size_t> QuerySamplerBudget() const;
 
     Runtime& runtime;
 
@@ -462,6 +481,11 @@ private:
         Common::SlotId object_id;
     };
 
+    struct PendingImageAccess {
+        ImageId image_id;
+        ImageAccessType access;
+    };
+
     Common::SlotVector<Image> slot_images;
     Common::SlotVector<ImageMapView> slot_map_views;
     Common::SlotVector<ImageView> slot_image_views;
@@ -477,6 +501,9 @@ private:
     std::vector<AsyncBuffer> uncommitted_async_buffers;
     std::deque<std::vector<AsyncBuffer>> async_buffers;
     std::deque<AsyncBuffer> async_buffers_death_ring;
+    std::vector<PendingImageAccess> staged_gpu_accesses;
+    std::deque<std::vector<PendingImageAccess>> committed_gpu_accesses;
+    std::deque<u64> committed_gpu_ticks;
 
     struct LRUItemParams {
         using ObjectType = ImageId;
@@ -500,6 +527,7 @@ private:
 
     u64 modification_tick = 0;
     u64 frame_tick = 0;
+    u64 last_sampler_gc_frame = (std::numeric_limits<u64>::max)();
 
     Common::ThreadWorker texture_decode_worker{1, "TextureDecoder"};
     std::vector<std::unique_ptr<AsyncDecodeContext>> async_decodes;

src/video_core/texture_cache/types.h

@@ -155,4 +155,9 @@ struct SwizzleParameters {
     s32 level;
 };
 
+enum class ImageAccessType : u8 {
+    Read,
+    Write,
+};
+
 } // namespace VideoCommon

src/video_core/vulkan_common/vulkan.h

@@ -29,9 +29,6 @@
 #ifndef VK_KHR_MAINTENANCE_8_EXTENSION_NAME
 #define VK_KHR_MAINTENANCE_8_EXTENSION_NAME "VK_KHR_maintenance8"
 #endif
-#ifndef VK_KHR_MAINTENANCE_9_EXTENSION_NAME
-#define VK_KHR_MAINTENANCE_9_EXTENSION_NAME "VK_KHR_maintenance9"
-#endif
 
 // Sanitize macros
 #undef CreateEvent

src/video_core/vulkan_common/vulkan_device.cpp

@@ -7,6 +7,7 @@
 #include <algorithm>
 #include <bitset>
 #include <chrono>
+#include <mutex>
 #include <optional>
 #include <thread>
 #include <unordered_set>
@@ -92,8 +93,57 @@ constexpr std::array VK_FORMAT_A4B4G4R4_UNORM_PACK16{
     VK_FORMAT_UNDEFINED,
 };
 
+constexpr std::array B10G11R11_UFLOAT_PACK32{
+    VK_FORMAT_R16G16B16A16_SFLOAT,
+    VK_FORMAT_R16G16B16A16_UNORM,
+    VK_FORMAT_UNDEFINED,
+};
+
+constexpr std::array R16G16B16A16_SFLOAT{
+    VK_FORMAT_A8B8G8R8_UNORM_PACK32,
+    VK_FORMAT_B8G8R8A8_UNORM,
+    VK_FORMAT_UNDEFINED,
+};
+
+constexpr std::array R16G16B16A16_UNORM{
+    VK_FORMAT_A8B8G8R8_UNORM_PACK32,
+    VK_FORMAT_B8G8R8A8_UNORM,
+    VK_FORMAT_UNDEFINED,
+};
+
 } // namespace Alternatives
 
+struct UnsupportedFormatKey {
+    VkFormat format;
+    VkFormatFeatureFlags usage;
+    FormatType type;
+
+    bool operator==(const UnsupportedFormatKey&) const noexcept = default;
+};
+
+struct UnsupportedFormatKeyHash {
+    size_t operator()(const UnsupportedFormatKey& key) const noexcept {
+        size_t seed = std::hash<int>{}(static_cast<int>(key.format));
+        seed ^= static_cast<size_t>(key.usage) + 0x9e3779b97f4a7c15ULL + (seed << 6) + (seed >> 2);
+        seed ^= static_cast<size_t>(key.type) + 0x9e3779b97f4a7c15ULL + (seed << 6) + (seed >> 2);
+        return seed;
+    }
+};
+
+bool ShouldLogUnsupportedFormat(VkFormat format, VkFormatFeatureFlags usage, FormatType type) {
+    static std::mutex mutex;
+    static std::unordered_set<UnsupportedFormatKey, UnsupportedFormatKeyHash> logged_keys;
+    const UnsupportedFormatKey key{format, usage, type};
+    std::scoped_lock lock{mutex};
+    const auto [it, inserted] = logged_keys.insert(key);
+    return inserted;
+}
+
+[[maybe_unused]] constexpr VkShaderStageFlags GraphicsStageMask =
+    VK_SHADER_STAGE_VERTEX_BIT | VK_SHADER_STAGE_TESSELLATION_CONTROL_BIT |
+    VK_SHADER_STAGE_TESSELLATION_EVALUATION_BIT | VK_SHADER_STAGE_GEOMETRY_BIT |
+    VK_SHADER_STAGE_FRAGMENT_BIT;
+
 template <typename T>
 void SetNext(void**& next, T& data) {
     *next = &data;
@@ -124,6 +174,12 @@ constexpr const VkFormat* GetFormatAlternatives(VkFormat format) {
         return Alternatives::VK_FORMAT_R32G32B32_SFLOAT.data();
     case VK_FORMAT_A4B4G4R4_UNORM_PACK16_EXT:
         return Alternatives::VK_FORMAT_A4B4G4R4_UNORM_PACK16.data();
+    case VK_FORMAT_B10G11R11_UFLOAT_PACK32:
+        return Alternatives::B10G11R11_UFLOAT_PACK32.data();
+    case VK_FORMAT_R16G16B16A16_SFLOAT:
+        return Alternatives::R16G16B16A16_SFLOAT.data();
+    case VK_FORMAT_R16G16B16A16_UNORM:
+        return Alternatives::R16G16B16A16_UNORM.data();
     default:
         return nullptr;
     }
@@ -379,10 +435,10 @@ void Device::RemoveExtension(bool& extension, const std::string& extension_name)
     loaded_extensions.erase(extension_name);
 }
 
-void Device::RemoveExtensionIfUnsuitable(bool is_suitable, const std::string& extension_name) {
+void Device::RemoveExtensionIfUnsuitable(bool& extension, const std::string& extension_name) {
-    if (loaded_extensions.contains(extension_name) && !is_suitable) {
+    if (loaded_extensions.contains(extension_name) && !extension) {
         LOG_WARNING(Render_Vulkan, "Removing unsuitable extension {}", extension_name);
-        this->RemoveExtension(is_suitable, extension_name);
+        this->RemoveExtension(extension, extension_name);
     }
 }
 
@@ -403,11 +459,12 @@ void Device::RemoveExtensionFeature(bool& extension, Feature& feature,
 }
 
 template <typename Feature>
-void Device::RemoveExtensionFeatureIfUnsuitable(bool is_suitable, Feature& feature,
+void Device::RemoveExtensionFeatureIfUnsuitable(bool& extension, Feature& feature,
                                                 const std::string& extension_name) {
-    if (loaded_extensions.contains(extension_name) && !is_suitable) {
+    if (loaded_extensions.contains(extension_name) && !extension) {
-        LOG_WARNING(Render_Vulkan, "Removing features for unsuitable extension {}", extension_name);
+        LOG_WARNING(Render_Vulkan,
-        this->RemoveExtensionFeature(is_suitable, feature, extension_name);
+                    "Removing features for unsuitable extension {}", extension_name);
+        this->RemoveExtensionFeature(extension, feature, extension_name);
     }
 }
 
@@ -501,9 +558,16 @@ Device::Device(VkInstance instance_, vk::PhysicalDevice physical_, VkSurfaceKHR
         LOG_WARNING(Render_Vulkan,
                     "Qualcomm drivers require scaled vertex format emulation; forcing fallback");
 
-        LOG_WARNING(Render_Vulkan,
+        if (extensions.shader_float_controls) {
-                    "Disabling shader float controls and 64-bit integer features on Qualcomm proprietary drivers");
+            LOG_WARNING(Render_Vulkan,
-        RemoveExtension(extensions.shader_float_controls, VK_KHR_SHADER_FLOAT_CONTROLS_EXTENSION_NAME);
+                        "Qualcomm drivers: VK_KHR_shader_float_controls is unstable; disabling usage");
+            RemoveExtension(extensions.shader_float_controls,
+                            VK_KHR_SHADER_FLOAT_CONTROLS_EXTENSION_NAME);
+        } else {
+            LOG_INFO(Render_Vulkan,
+                     "Qualcomm drivers: VK_KHR_shader_float_controls already unavailable");
+        }
+        disable_shader_float_controls_usage = true;
         RemoveExtensionFeature(extensions.shader_atomic_int64, features.shader_atomic_int64,
                                VK_KHR_SHADER_ATOMIC_INT64_EXTENSION_NAME);
         features.shader_atomic_int64.shaderBufferInt64Atomics = false;
@@ -600,10 +664,55 @@ Device::Device(VkInstance instance_, vk::PhysicalDevice physical_, VkSurfaceKHR
         }
     }
 
+    const bool needs_mobile_alignment_clamp = is_qualcomm || is_arm;
+
     if (is_qualcomm) {
-        const u32 version = (properties.properties.driverVersion << 3) >> 3;
+        has_broken_parallel_compiling = true;
-        if (version < VK_MAKE_API_VERSION(0, 255, 615, 512)) {
+        LOG_WARNING(Render_Vulkan,
-            has_broken_parallel_compiling = true;
+                    "Adreno drivers exhibit instability with parallel shader compilation; "
+                    "forcing single-threaded pipeline builds");
+        if (extensions.push_descriptor) {
+            LOG_WARNING(Render_Vulkan,
+                        "Qualcomm driver mishandles push descriptors; disabling "
+                        "VK_KHR_push_descriptor to avoid incomplete descriptor writes");
+            RemoveExtension(extensions.push_descriptor, VK_KHR_PUSH_DESCRIPTOR_EXTENSION_NAME);
+        }
+    }
+
+    if (needs_mobile_alignment_clamp) {
+        const char* driver_label = is_qualcomm ? "Qualcomm" : "ARM";
+
+        constexpr VkDeviceSize MIN_UNIFORM_ALIGNMENT = 256;
+        const VkDeviceSize reported_uniform_alignment =
+            properties.properties.limits.minUniformBufferOffsetAlignment;
+        if (reported_uniform_alignment < MIN_UNIFORM_ALIGNMENT) {
+            uniform_buffer_alignment_minimum = MIN_UNIFORM_ALIGNMENT;
+            LOG_WARNING(Render_Vulkan,
+                        "{} driver reports {}-byte minUniformBufferOffsetAlignment; clamping to {}",
+                        driver_label, reported_uniform_alignment, uniform_buffer_alignment_minimum);
+        }
+
+        constexpr VkDeviceSize MIN_STORAGE_ALIGNMENT = 64;
+        const VkDeviceSize reported_storage_alignment =
+            properties.properties.limits.minStorageBufferOffsetAlignment;
+        if (reported_storage_alignment < MIN_STORAGE_ALIGNMENT) {
+            storage_buffer_alignment_minimum = MIN_STORAGE_ALIGNMENT;
+            LOG_WARNING(Render_Vulkan,
+                        "{} driver reports {}-byte minStorageBufferOffsetAlignment; clamping to {}",
+                        driver_label, reported_storage_alignment, storage_buffer_alignment_minimum);
+        }
+
+        const size_t sampler_limit = properties.properties.limits.maxSamplerAllocationCount;
+        if (sampler_limit > 0) {
+            constexpr size_t MIN_SAMPLER_BUDGET = 1024U;
+            const size_t reserved = sampler_limit / 4U;
+            const size_t derived_budget =
+                (std::max)(MIN_SAMPLER_BUDGET, sampler_limit - reserved);
+            sampler_heap_budget = derived_budget;
+            LOG_WARNING(Render_Vulkan,
+                        "{} driver reports max {} samplers; reserving {} (25%) and "
+                        "allowing Eden to use {} (75%) to avoid heap exhaustion",
+                        driver_label, sampler_limit, reserved, sampler_heap_budget);
         }
     }
 
@@ -662,7 +771,7 @@ Device::Device(VkInstance instance_, vk::PhysicalDevice physical_, VkSurfaceKHR
         properties.properties.limits.maxVertexInputBindings = 32;
     }
 
-    const auto dyna_state = Settings::values.dyna_state.GetValue();
+    const auto dyna_state = u32(Settings::values.dyna_state.GetValue());
 
     // Base dynamic states (VIEWPORT, SCISSOR, DEPTH_BIAS, etc.) are ALWAYS active in vk_graphics_pipeline.cpp
     // This slider controls EXTENDED dynamic states with accumulative levels per Vulkan specs:
@@ -672,7 +781,7 @@ Device::Device(VkInstance instance_, vk::PhysicalDevice physical_, VkSurfaceKHR
     //   Level 3 = Core + VK_EXT_extended_dynamic_state + VK_EXT_extended_dynamic_state2 + VK_EXT_extended_dynamic_state3
     
     switch (dyna_state) {
-    case Settings::ExtendedDynamicState::Disabled:
+    case 0:
         // Level 0: Disable all extended dynamic state extensions
         RemoveExtensionFeature(extensions.extended_dynamic_state, features.extended_dynamic_state, 
                               VK_EXT_EXTENDED_DYNAMIC_STATE_EXTENSION_NAME);
@@ -683,7 +792,7 @@ Device::Device(VkInstance instance_, vk::PhysicalDevice physical_, VkSurfaceKHR
         dynamic_state3_blending = false;
         dynamic_state3_enables = false;
         break;
-    case Settings::ExtendedDynamicState::EDS1:
+    case 1:
         // Level 1: Enable EDS1, disable EDS2 and EDS3
         RemoveExtensionFeature(extensions.extended_dynamic_state2, features.extended_dynamic_state2, 
                               VK_EXT_EXTENDED_DYNAMIC_STATE_2_EXTENSION_NAME);
@@ -692,14 +801,14 @@ Device::Device(VkInstance instance_, vk::PhysicalDevice physical_, VkSurfaceKHR
         dynamic_state3_blending = false;
         dynamic_state3_enables = false;
         break;
-    case Settings::ExtendedDynamicState::EDS2:
+    case 2:
         // Level 2: Enable EDS1 + EDS2, disable EDS3
         RemoveExtensionFeature(extensions.extended_dynamic_state3, features.extended_dynamic_state3, 
                               VK_EXT_EXTENDED_DYNAMIC_STATE_3_EXTENSION_NAME);
         dynamic_state3_blending = false;
         dynamic_state3_enables = false;
         break;
-    case Settings::ExtendedDynamicState::EDS3:
+    case 3:
     default:
         // Level 3: Enable all (EDS1 + EDS2 + EDS3)
         break;
@@ -752,32 +861,63 @@ VkFormat Device::GetSupportedFormat(VkFormat wanted_format, VkFormatFeatureFlags
     if (IsFormatSupported(wanted_format, wanted_usage, format_type)) {
         return wanted_format;
     }
-    // The wanted format is not supported by hardware, search for alternatives
+
     const VkFormat* alternatives = GetFormatAlternatives(wanted_format);
     if (alternatives == nullptr) {
-        LOG_ERROR(Render_Vulkan,
+        if (ShouldLogUnsupportedFormat(wanted_format, wanted_usage, format_type)) {
-                  "Format={} with usage={} and type={} has no defined alternatives and host "
+            LOG_ERROR(Render_Vulkan,
-                  "hardware does not support it",
+                      "Format={} with usage={} and type={} has no defined alternatives and host "
-                  wanted_format, wanted_usage, format_type);
+                      "hardware does not support it",
+                      wanted_format, wanted_usage, format_type);
+        }
         return wanted_format;
     }
 
-    std::size_t i = 0;
+#if defined(ANDROID)
-    for (VkFormat alternative = *alternatives; alternative; alternative = alternatives[++i]) {
+    const auto downgrade_r16_to_ldr = [&](VkFormat candidate) -> VkFormat {
+        if (wanted_format != VK_FORMAT_B10G11R11_UFLOAT_PACK32) {
+            return candidate;
+        }
+        if (candidate != VK_FORMAT_R16G16B16A16_SFLOAT &&
+            candidate != VK_FORMAT_R16G16B16A16_UNORM) {
+            return candidate;
+        }
+        static constexpr std::array ldr_candidates{
+            VK_FORMAT_A8B8G8R8_UNORM_PACK32,
+            VK_FORMAT_A8B8G8R8_SRGB_PACK32,
+            VK_FORMAT_B8G8R8A8_UNORM,
+            VK_FORMAT_B8G8R8A8_SRGB,
+        };
+        for (const VkFormat format : ldr_candidates) {
+            if (IsFormatSupported(format, wanted_usage, format_type)) {
+                return format;
+            }
+        }
+        return candidate;
+    };
+#endif
+
+    for (std::size_t i = 0; VkFormat alternative = alternatives[i]; ++i) {
         if (!IsFormatSupported(alternative, wanted_usage, format_type)) {
             continue;
         }
+#if defined(ANDROID)
+        const VkFormat selected = downgrade_r16_to_ldr(alternative);
+#else
+        const VkFormat selected = alternative;
+#endif
         LOG_DEBUG(Render_Vulkan,
                   "Emulating format={} with alternative format={} with usage={} and type={}",
-                  wanted_format, alternative, wanted_usage, format_type);
+                  wanted_format, selected, wanted_usage, format_type);
-        return alternative;
+        return selected;
     }
 
-    // No alternatives found, panic
+    if (ShouldLogUnsupportedFormat(wanted_format, wanted_usage, format_type)) {
-    LOG_ERROR(Render_Vulkan,
+        LOG_ERROR(Render_Vulkan,
-              "Format={} with usage={} and type={} is not supported by the host hardware and "
+                  "Format={} with usage={} and type={} is not supported by the host hardware and "
-              "doesn't support any of the alternatives",
+                  "doesn't support any of the alternatives",
-              wanted_format, wanted_usage, format_type);
+                  wanted_format, wanted_usage, format_type);
+    }
     return wanted_format;
 }
 
@@ -924,11 +1064,8 @@ bool Device::ShouldBoostClocks() const {
 }
 
 bool Device::HasTimelineSemaphore() const {
-    if (GetDriverID() == VK_DRIVER_ID_QUALCOMM_PROPRIETARY ||
+    const VkDriverIdKHR driver_id = GetDriverID();
-        GetDriverID() == VK_DRIVER_ID_MESA_TURNIP) {
+    if (driver_id == VK_DRIVER_ID_MESA_TURNIP || driver_id == VK_DRIVER_ID_QUALCOMM_PROPRIETARY) {
-        // Timeline semaphores do not work properly on all Qualcomm drivers.
-        // They generally work properly with Turnip drivers, but are problematic on some devices
-        // (e.g. ZTE handsets with Snapdragon 870).
         return false;
     }
     return features.timeline_semaphore.timelineSemaphore;
@@ -1018,11 +1155,13 @@ bool Device::GetSuitability(bool requires_swapchain) {
     // Vulkan 1.2 and 1.3 features
     if (instance_version >= VK_API_VERSION_1_2) {
         features_1_2.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_VULKAN_1_2_FEATURES;
-        features_1_3.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_VULKAN_1_3_FEATURES;
+        features_1_2.pNext = nullptr;
-
-        features_1_2.pNext = &features_1_3;
-
         *next = &features_1_2;
+        if (instance_version >= VK_API_VERSION_1_3) {
+            features_1_3.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_VULKAN_1_3_FEATURES;
+            features_1_3.pNext = nullptr;
+            features_1_2.pNext = &features_1_3;
+        }
     }
 
 // Test all features we know about. If the feature is not available in core at our
@@ -1168,6 +1307,17 @@ bool Device::GetSuitability(bool requires_swapchain) {
     // Driver detection variables for workarounds in GetSuitability
     const VkDriverId driver_id = properties.driver.driverID;
     const bool is_intel_windows = driver_id == VK_DRIVER_ID_INTEL_PROPRIETARY_WINDOWS;
+    const bool is_turnip = driver_id == VK_DRIVER_ID_MESA_TURNIP;
+    const bool is_qualcomm = driver_id == VK_DRIVER_ID_QUALCOMM_PROPRIETARY;
+
+    // VK_EXT_CUSTOM_BORDER_COLOR
+    
+    if (extensions.custom_border_color && (is_qualcomm || is_turnip)) {
+        const char* driver_name = is_turnip ? "Turnip" : "Adreno";
+        LOG_WARNING(Render_Vulkan,
+                    "{} driver previously required custom border color disablement; leaving VK_EXT_custom_border_color enabled for testing",
+                    driver_name);
+    }
     
     // VK_EXT_extended_dynamic_state2 below this will appear drivers that need workarounds.
     
@@ -1373,6 +1523,14 @@ void Device::RemoveUnsuitableExtensions() {
         features.transform_feedback.transformFeedback &&
         properties.transform_feedback.maxTransformFeedbackBuffers > 0 &&
         properties.transform_feedback.transformFeedbackQueries;
+
+    if (extensions.transform_feedback) {
+        features.transform_feedback.transformFeedback = VK_TRUE;
+        features.transform_feedback.geometryStreams = VK_TRUE;
+    } else {
+        features.transform_feedback.transformFeedback = VK_FALSE;
+        features.transform_feedback.geometryStreams = VK_FALSE;
+    }
     RemoveExtensionFeatureIfUnsuitable(extensions.transform_feedback, features.transform_feedback,
                                        VK_EXT_TRANSFORM_FEEDBACK_EXTENSION_NAME);
     if (extensions.transform_feedback) {
@@ -1390,12 +1548,13 @@ void Device::RemoveUnsuitableExtensions() {
 
     // VK_EXT_multi_draw
     extensions.multi_draw = features.multi_draw.multiDraw;
-
     if (extensions.multi_draw) {
+        features.multi_draw.multiDraw = VK_TRUE;
         LOG_INFO(Render_Vulkan, "VK_EXT_multi_draw: maxMultiDrawCount={}",
                  properties.multi_draw.maxMultiDrawCount);
+    } else {
+        features.multi_draw.multiDraw = VK_FALSE;
     }
-
     RemoveExtensionFeatureIfUnsuitable(extensions.multi_draw, features.multi_draw,
                                        VK_EXT_MULTI_DRAW_EXTENSION_NAME);
 
@@ -1462,7 +1621,6 @@ void Device::RemoveUnsuitableExtensions() {
 
     // VK_KHR_maintenance5
     extensions.maintenance5 = features.maintenance5.maintenance5;
-
     if (extensions.maintenance5) {
         LOG_INFO(Render_Vulkan, "VK_KHR_maintenance5 properties: polygonModePointSize={} "
                                 "depthStencilSwizzleOne={} earlyFragmentTests={} nonStrictWideLines={}",
@@ -1472,7 +1630,6 @@ void Device::RemoveUnsuitableExtensions() {
                  properties.maintenance5.earlyFragmentSampleMaskTestBeforeSampleCounting,
                  properties.maintenance5.nonStrictWideLinesUseParallelogram);
     }
-
     RemoveExtensionFeatureIfUnsuitable(extensions.maintenance5, features.maintenance5,
                                        VK_KHR_MAINTENANCE_5_EXTENSION_NAME);
 
@@ -1489,9 +1646,33 @@ void Device::RemoveUnsuitableExtensions() {
     extensions.maintenance8 = loaded_extensions.contains(VK_KHR_MAINTENANCE_8_EXTENSION_NAME);
     RemoveExtensionIfUnsuitable(extensions.maintenance8, VK_KHR_MAINTENANCE_8_EXTENSION_NAME);
 
-    // VK_KHR_maintenance9 (proposed for Vulkan 1.4, no features)
+}
-    extensions.maintenance9 = loaded_extensions.contains(VK_KHR_MAINTENANCE_9_EXTENSION_NAME);
+
-    RemoveExtensionIfUnsuitable(extensions.maintenance9, VK_KHR_MAINTENANCE_9_EXTENSION_NAME);
+bool Device::SupportsSubgroupStage(VkShaderStageFlags stage_mask) const {
+    if (stage_mask == 0) {
+        return true;
+    }
+    const VkShaderStageFlags supported = properties.subgroup_properties.supportedStages;
+    return (supported & stage_mask) == stage_mask;
+}
+
+bool Device::IsSubgroupFeatureSupported(VkSubgroupFeatureFlagBits feature,
+                                        VkShaderStageFlags stage_mask) const {
+    if ((properties.subgroup_properties.supportedOperations & feature) == 0) {
+        return false;
+    }
+    return SupportsSubgroupStage(stage_mask);
+}
+
+bool Device::SupportsWarpIntrinsics(VkShaderStageFlagBits stage) const {
+    constexpr VkSubgroupFeatureFlags required_ops =
+        VK_SUBGROUP_FEATURE_BASIC_BIT | VK_SUBGROUP_FEATURE_VOTE_BIT |
+        VK_SUBGROUP_FEATURE_ARITHMETIC_BIT | VK_SUBGROUP_FEATURE_BALLOT_BIT |
+        VK_SUBGROUP_FEATURE_SHUFFLE_BIT | VK_SUBGROUP_FEATURE_SHUFFLE_RELATIVE_BIT;
+    if ((properties.subgroup_properties.supportedOperations & required_ops) != required_ops) {
+        return false;
+    }
+    return SupportsSubgroupStage(stage);
 }
 
 void Device::SetupFamilies(VkSurfaceKHR surface) {
@@ -1529,6 +1710,13 @@ void Device::SetupFamilies(VkSurfaceKHR surface) {
     }
 }
 
+std::optional<size_t> Device::GetSamplerHeapBudget() const {
+    if (sampler_heap_budget == 0) {
+        return std::nullopt;
+    }
+    return sampler_heap_budget;
+}
+
 u64 Device::GetDeviceMemoryUsage() const {
     VkPhysicalDeviceMemoryBudgetPropertiesEXT budget;
     budget.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_MEMORY_BUDGET_PROPERTIES_EXT;

src/video_core/vulkan_common/vulkan_device.h

@@ -6,6 +6,8 @@
 
 #pragma once
 
+#include <algorithm>
+#include <optional>
 #include <set>
 #include <span>
 #include <string>
@@ -97,7 +99,6 @@ VK_DEFINE_HANDLE(VmaAllocator)
     EXTENSION(KHR, MAINTENANCE_3, maintenance3)                                                    \
     EXTENSION(KHR, MAINTENANCE_7, maintenance7)                                                    \
     EXTENSION(KHR, MAINTENANCE_8, maintenance8)                                                    \
-    EXTENSION(KHR, MAINTENANCE_9, maintenance9)                                                    \
     EXTENSION(NV, DEVICE_DIAGNOSTICS_CONFIG, device_diagnostics_config)                            \
     EXTENSION(NV, GEOMETRY_SHADER_PASSTHROUGH, geometry_shader_passthrough)                        \
     EXTENSION(NV, VIEWPORT_ARRAY2, viewport_array2)                                                \
@@ -303,12 +304,14 @@ public:
 
     /// Returns uniform buffer alignment requirement.
     VkDeviceSize GetUniformBufferAlignment() const {
-        return properties.properties.limits.minUniformBufferOffsetAlignment;
+        return (std::max)(properties.properties.limits.minUniformBufferOffsetAlignment,
+                          uniform_buffer_alignment_minimum);
     }
 
     /// Returns storage alignment requirement.
     VkDeviceSize GetStorageBufferAlignment() const {
-        return properties.properties.limits.minStorageBufferOffsetAlignment;
+        return (std::max)(properties.properties.limits.minStorageBufferOffsetAlignment,
+                          storage_buffer_alignment_minimum);
     }
 
     /// Returns the maximum range for storage buffers.
@@ -326,11 +329,22 @@ public:
         return properties.properties.limits.maxComputeSharedMemorySize;
     }
 
+    /// Returns the maximum number of dynamic storage buffer descriptors per set.
+    u32 GetMaxDescriptorSetStorageBuffersDynamic() const {
+        return properties.properties.limits.maxDescriptorSetStorageBuffersDynamic;
+    }
+
+    /// Returns the maximum number of dynamic uniform buffer descriptors per set.
+    u32 GetMaxDescriptorSetUniformBuffersDynamic() const {
+        return properties.properties.limits.maxDescriptorSetUniformBuffersDynamic;
+    }
+
     /// Returns float control properties of the device.
     const VkPhysicalDeviceFloatControlsPropertiesKHR& FloatControlProperties() const {
         return properties.float_controls;
     }
 
+
     /// Returns true if ASTC is natively supported.
     bool IsOptimalAstcSupported() const {
         return features.features.textureCompressionASTC_LDR;
@@ -346,6 +360,11 @@ public:
         return GetDriverID() != VK_DRIVER_ID_QUALCOMM_PROPRIETARY;
     }
 
+    /// Returns true if vkResetQueryPool can be used from the host.
+    bool SupportsHostQueryReset() const {
+        return features.host_query_reset.hostQueryReset;
+    }
+
     /// Returns true if the device supports float64 natively.
     bool IsFloat64Supported() const {
         return features.features.shaderFloat64;
@@ -376,11 +395,26 @@ public:
         return properties.subgroup_size_control.requiredSubgroupSizeStages & stage;
     }
 
-    /// Returns true if the device supports the provided subgroup feature.
+    /// Returns true if the device supports the provided subgroup feature for the given stages.
-    bool IsSubgroupFeatureSupported(VkSubgroupFeatureFlagBits feature) const {
+    bool IsSubgroupFeatureSupported(VkSubgroupFeatureFlagBits feature,
-        return properties.subgroup_properties.supportedOperations & feature;
+                                    VkShaderStageFlags stage_mask = 0) const;
+
+    /// Returns true if the device reports subgroup support for the provided shader stages.
+    bool SupportsSubgroupStage(VkShaderStageFlags stage_mask) const;
+
+    /// Returns the set of stages that report subgroup support.
+    VkShaderStageFlags GetSubgroupSupportedStages() const {
+        return properties.subgroup_properties.supportedStages;
     }
 
+    /// Returns the set of subgroup operations reported by the driver.
+    VkSubgroupFeatureFlags GetSubgroupSupportedOperations() const {
+        return properties.subgroup_properties.supportedOperations;
+    }
+
+    /// Returns true if the driver can execute all warp intrinsics for the given shader stage.
+    bool SupportsWarpIntrinsics(VkShaderStageFlagBits stage) const;
+
     /// Returns the maximum number of push descriptors.
     u32 MaxPushDescriptors() const {
         return properties.push_descriptor.maxPushDescriptors;
@@ -463,7 +497,7 @@ public:
 
     /// Returns true if VK_KHR_shader_float_controls is enabled.
     bool IsKhrShaderFloatControlsSupported() const {
-        return extensions.shader_float_controls;
+        return extensions.shader_float_controls && !disable_shader_float_controls_usage;
     }
 
     /// Returns true if the device supports VK_KHR_workgroup_memory_explicit_layout.
@@ -654,6 +688,16 @@ public:
         return features.features.alphaToOne != VK_FALSE;
     }
 
+    bool SupportsDualSourceBlend(u32 required_dual_source_attachments = 1) const {
+        const u32 max_dual = properties.properties.limits.maxFragmentDualSrcAttachments;
+        return features.features.dualSrcBlend != VK_FALSE &&
+               max_dual >= required_dual_source_attachments;
+    }
+
+    u32 MaxFragmentDualSrcAttachments() const {
+        return properties.properties.limits.maxFragmentDualSrcAttachments;
+    }
+
     bool SupportsDynamicState3DepthClampEnable() const {
         return dynamic_state3_depth_clamp_enable;
     }
@@ -687,9 +731,10 @@ public:
         return u32(Settings::values.dyna_state.GetValue()) > 0;
     }
 
-    /// Returns true if the device supports VK_EXT_vertex_input_dynamic_state.
+    /// Returns true if VK_EXT_vertex_input_dynamic_state is enabled on the device.
     bool IsExtVertexInputDynamicStateSupported() const {
-        return extensions.vertex_input_dynamic_state;
+        return extensions.vertex_input_dynamic_state &&
+               features.vertex_input_dynamic_state.vertexInputDynamicState;
     }
 
     /// Returns true if the device supports VK_EXT_shader_demote_to_helper_invocation
@@ -720,10 +765,11 @@ public:
 
     /// Returns the minimum supported version of SPIR-V.
     u32 SupportedSpirvVersion() const {
-        if (instance_version >= VK_API_VERSION_1_3) {
+        const bool has_float_controls = extensions.shader_float_controls;
+        if (instance_version >= VK_API_VERSION_1_3 && has_float_controls) {
             return 0x00010600U;
         }
-        if (extensions.spirv_1_4) {
+        if (extensions.spirv_1_4 && has_float_controls) {
             return 0x00010400U;
         }
         return 0x00010300U;
@@ -749,6 +795,8 @@ public:
         return has_broken_parallel_compiling;
     }
 
+    std::optional<size_t> GetSamplerHeapBudget() const;
+
     /// Returns the vendor name reported from Vulkan.
     std::string_view GetVendorName() const {
         return properties.driver.driverName;
@@ -883,11 +931,6 @@ public:
         return extensions.maintenance8;
     }
 
-    /// Returns true if the device supports VK_KHR_maintenance9.
-    bool IsKhrMaintenance9Supported() const {
-        return extensions.maintenance9;
-    }
-
     [[nodiscard]] static constexpr bool CheckBrokenCompute(VkDriverId driver_id,
                                                            u32 driver_version) {
         if (driver_id == VK_DRIVER_ID_INTEL_PROPRIETARY_WINDOWS) {
@@ -925,13 +968,13 @@ private:
     void RemoveUnsuitableExtensions();
 
     void RemoveExtension(bool& extension, const std::string& extension_name);
-    void RemoveExtensionIfUnsuitable(bool is_suitable, const std::string& extension_name);
+    void RemoveExtensionIfUnsuitable(bool& extension, const std::string& extension_name);
 
     template <typename Feature>
     void RemoveExtensionFeature(bool& extension, Feature& feature,
                                 const std::string& extension_name);
     template <typename Feature>
-    void RemoveExtensionFeatureIfUnsuitable(bool is_suitable, Feature& feature,
+    void RemoveExtensionFeatureIfUnsuitable(bool& extension, Feature& feature,
                                             const std::string& extension_name);
 
     /// Sets up queue families.
@@ -1035,6 +1078,7 @@ private:
     bool cant_blit_msaa{};                     ///< Does not support MSAA<->MSAA blitting.
     bool must_emulate_scaled_formats{};        ///< Requires scaled vertex format emulation
     bool must_emulate_bgr565{};                ///< Emulates BGR565 by swizzling RGB565 format.
+    bool disable_shader_float_controls_usage{}; ///< True when VK_KHR_shader_float_controls cannot be safely used.
     bool dynamic_state3_blending{};            ///< Has blending features of dynamic_state3.
     bool dynamic_state3_enables{};             ///< Has at least one enable feature of dynamic_state3.
     bool dynamic_state3_depth_clamp_enable{};
@@ -1045,6 +1089,9 @@ private:
     bool dynamic_state3_alpha_to_coverage{};
     bool dynamic_state3_alpha_to_one{};
     bool supports_conditional_barriers{};      ///< Allows barriers in conditional control flow.
+    size_t sampler_heap_budget{};              ///< Sampler budget for buggy drivers (0 = unlimited).
+    VkDeviceSize uniform_buffer_alignment_minimum{}; ///< Minimum enforced UBO alignment.
+    VkDeviceSize storage_buffer_alignment_minimum{}; ///< Minimum enforced SSBO alignment.
     u64 device_access_memory{};                ///< Total size of device local memory in bytes.
     u32 sets_per_pool{};                       ///< Sets per Description Pool
     NvidiaArchitecture nvidia_arch{NvidiaArchitecture::Arch_AmpereOrNewer};