Missing var

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

src/common/settings.h

@@ -600,13 +600,7 @@ struct Values {
                                                 Category::RendererExtensions,
                                                 Specialization::Scalar};
 
-    SwitchableSetting<bool> vertex_input_dynamic_state{linkage,
-#if defined (ANDROID)
-                                                       false,
-#else
-                                                       true,
-#endif
-                                                       "vertex_input_dynamic_state", Category::RendererExtensions};
+    SwitchableSetting<bool> vertex_input_dynamic_state{linkage, false, "vertex_input_dynamic_state", Category::RendererExtensions};
 
     Setting<bool> renderer_debug{linkage, false, "debug", Category::RendererDebug};
     Setting<bool> renderer_shader_feedback{linkage, false, "shader_feedback",

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

@@ -6,6 +6,7 @@
 
 #include <algorithm>
 #include <array>
+#include <cstring>
 #include <sstream>
 
 #include <boost/range/algorithm_ext/erase.hpp>
@@ -190,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
@@ -297,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());
@@ -396,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,
-                     buffer_size);
-        size = buffer_size; // TODO(bunnei): This needs to be HW tested
+        LOG_WARNING(Core, "size ({:016X}) is greater than buffer_size ({:016X}); clamping",
+                    size, buffer_size);
+        size = buffer_size;
     }
 
     if (is_buffer_b) {
@@ -421,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,
-                     buffer_size);
-        size = buffer_size; // TODO(bunnei): This needs to be HW tested
+        LOG_WARNING(Core, "size ({:016X}) is greater than buffer_size ({:016X}); clamping",
+                    size, buffer_size);
+        size = buffer_size;
     }
 
     memory.WriteBlock(BufferDescriptorB()[buffer_index].Address(), buffer, size);
@@ -438,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,
-                     buffer_size);
-        size = buffer_size; // TODO(bunnei): This needs to be HW tested
+        LOG_WARNING(Core, "size ({:016X}) is greater than buffer_size ({:016X}); clamping",
+                    size, buffer_size);
+        size = buffer_size;
     }
 
     memory.WriteBlock(BufferDescriptorC()[buffer_index].Address(), buffer, size);
@@ -476,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

@@ -10,6 +10,7 @@
 #include <utility>
 #include <vector>
 
+#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"
@@ -430,15 +431,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 ||
-         info.uses_subgroup_shuffles) &&
-        profile.support_vote) {
+    const bool stage_supports_warp = profile.SupportsWarpIntrinsics(ctx.stage);
+    const bool needs_warp_intrinsics = info.uses_subgroup_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

@@ -483,6 +483,21 @@ 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 AttributeType type{ctx.runtime_info.color_output_types[index]};
     if (type == AttributeType::Float) {

src/shader_recompiler/backend/spirv/emit_spirv_image.cpp

@@ -40,7 +40,10 @@ public:
 
     explicit ImageOperands(EmitContext& ctx, const IR::Value& offset, const IR::Value& offset2) {
         if (offset2.IsEmpty()) {
-            AddOffset(ctx, offset, ImageGatherOffsetAllowed);
+            if (offset.IsEmpty()) {
+                return;
+            }
+            Add(spv::ImageOperandsMask::Offset, ctx.Def(offset));
             return;
         }
         const std::array values{offset.InstRecursive(), offset2.InstRecursive()};
@@ -196,6 +199,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 (info.type == TextureType::Buffer) {
         const TextureBufferDefinition& def{ctx.texture_buffers.at(info.descriptor_index)};
@@ -348,8 +386,8 @@ void AddOffsetToCoordinates(EmitContext& ctx, const IR::TextureInstInfo& info, I
         break;
     }
     case TextureType::ColorArray2D:
-        offset = ctx.OpCompositeConstruct(ctx.U32[3], ctx.OpCompositeExtract(ctx.U32[1], offset, 0),
-                                          ctx.OpCompositeExtract(ctx.U32[1], offset, 1),
+        offset = ctx.OpCompositeConstruct(ctx.U32[3], ctx.OpCompositeExtract(ctx.U32[1], coords, 0),
+                                          ctx.OpCompositeExtract(ctx.U32[1], coords, 1),
                                           ctx.u32_zero_value);
         [[fallthrough]];
     case TextureType::Color3D: {
@@ -463,38 +501,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,
-                    &EmitContext::OpImageSampleImplicitLod, ctx, inst, ctx.F32[4],
-                    Texture(ctx, info, index), coords, operands.MaskOptional(), operands.Span());
+        color = Emit(&EmitContext::OpImageSparseSampleImplicitLod,
+                     &EmitContext::OpImageSampleImplicitLod, ctx, inst, color_type, texture,
+                     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,
-                    &EmitContext::OpImageSampleExplicitLod, ctx, inst, ctx.F32[4],
-                    Texture(ctx, info, index), coords, operands.Mask(), operands.Span());
+        color = Emit(&EmitContext::OpImageSparseSampleExplicitLod,
+                     &EmitContext::OpImageSampleExplicitLod, ctx, inst, color_type, texture,
+                     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);
-
-    Id result = Emit(&EmitContext::OpImageSparseSampleExplicitLod,
-                     &EmitContext::OpImageSampleExplicitLod, ctx, inst, ctx.F32[4],
-                     Texture(ctx, info, index), coords, operands.Mask(), operands.Span());
-#ifdef ANDROID
-    if (Settings::values.fix_bloom_effects.GetValue()) {
-        result = ctx.OpVectorTimesScalar(ctx.F32[4], result, ctx.Const(0.98f));
-    }
-#endif
-    return result;
+    const Id color{Emit(&EmitContext::OpImageSparseSampleExplicitLod,
+                        &EmitContext::OpImageSampleExplicitLod, ctx, inst, color_type,
+                        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,
@@ -530,13 +569,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,
-                ctx.F32[4], Texture(ctx, info, index), coords, ctx.Const(info.gather_component),
-                operands.MaskOptional(), operands.Span());
+    const Id color{
+        Emit(&EmitContext::OpImageSparseGather, &EmitContext::OpImageGather, ctx, inst, color_type,
+             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,
@@ -554,6 +598,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{};
@@ -563,8 +610,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],
-                TextureImage(ctx, info, index), coords, operands.MaskOptional(), operands.Span());
+    Id color{Emit(&EmitContext::OpImageSparseFetch, &EmitContext::OpImageFetch, ctx, inst,
+                  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,
@@ -609,14 +661,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,
-                &EmitContext::OpImageSampleExplicitLod, ctx, inst, ctx.F32[4],
-                Texture(ctx, info, index), coords, operands.Mask(), operands.Span());
+    const Id color{Emit(&EmitContext::OpImageSparseSampleExplicitLod,
+                        &EmitContext::OpImageSampleExplicitLod, ctx, inst, color_type,
+                        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);
@@ -1360,7 +1375,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)};
@@ -1373,8 +1389,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);
@@ -1415,6 +1433,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);
@@ -1438,7 +1457,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);
@@ -1452,9 +1471,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 ||
-        (profile.warp_size_potentially_larger_than_guest &&
-         (info.uses_subgroup_vote || info.uses_subgroup_mask))) {
+    if (stage_supports_warp &&
+        (info.uses_fswzadd || info.uses_subgroup_invocation_id || info.uses_subgroup_shuffles ||
+         (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);
@@ -1689,6 +1709,22 @@ void EmitContext::DefineOutputs(const IR::Program& program) {
                 Decorate(frag_color[index], spv::Decoration::Location, index);
                 Name(frag_color[index], fmt::format("frag_color{}", index));
             }
+#if 0
+            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));
+#endif
         }
         if (info.stores_frag_depth) {
             frag_depth = DefineOutput(*this, F32[1], std::nullopt);

src/shader_recompiler/backend/spirv/spirv_emit_context.h

@@ -39,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 {
@@ -249,6 +251,7 @@ public:
 
     Id output_f32{};
     Id output_u32{};
+    Id output_s32{};
 
     Id image_buffer_type{};
     Id image_u32{};

src/shader_recompiler/environment.h

@@ -77,6 +77,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

@@ -418,6 +418,10 @@ u32 GetTextureHandle(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_,
@@ -455,7 +459,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 &&
@@ -690,10 +696,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

@@ -84,6 +84,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

@@ -154,6 +154,14 @@ enum class ImageFormat : u32 {
     R32G32B32A32_UINT,
 };
 
+enum class SamplerComponentType : u8 {
+    Float,
+    Sint,
+    Uint,
+    Depth,
+    Stencil,
+};
+
 enum class Interpolation {
     Smooth,
     Flat,
@@ -186,6 +194,7 @@ struct TextureBufferDescriptor {
     u32 secondary_shift_left;
     u32 count;
     u32 size_shift;
+    SamplerComponentType component_type;
 
     auto operator<=>(const TextureBufferDescriptor&) const = default;
 };
@@ -207,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/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

@@ -95,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_compute_pipeline.cpp

@@ -7,13 +7,50 @@
 #include <cstring>
 #include <bit>
 #include <numeric>
+#include <optional>
 #include "common/cityhash.h"
 #include "common/settings.h"
 #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;
@@ -175,8 +212,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;
         }
@@ -210,7 +251,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:
@@ -395,8 +429,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;
@@ -481,7 +519,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

@@ -219,6 +219,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

@@ -693,6 +693,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_} {
@@ -1227,6 +1236,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

@@ -16,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"
 
@@ -133,6 +134,8 @@ public:
         return false;
     }
 
+    bool SupportsLinearFilter(VideoCore::Surface::PixelFormat format) const noexcept;
+
     bool HasBrokenTextureViewFormats() const noexcept {
         return has_broken_texture_view_formats;
     }
@@ -141,6 +144,8 @@ public:
 
     void TickFrame() {}
 
+    void WaitForGpuTick(u64) {}
+
     StateTracker& GetStateTracker() {
         return state_tracker;
     }
@@ -270,6 +275,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/blit_image.cpp

@@ -37,7 +37,7 @@ namespace Vulkan {
 using VideoCommon::ImageViewType;
 
 namespace {
-
+    
 [[nodiscard]] VkImageAspectFlags AspectMaskFromFormat(VideoCore::Surface::PixelFormat format) {
     using VideoCore::Surface::SurfaceType;
     switch (VideoCore::Surface::GetFormatType(format)) {
@@ -516,16 +516,16 @@ BlitImageHelper::BlitImageHelper(const Device& device_, Scheduler& scheduler_,
           nullptr, PUSH_CONSTANT_RANGE<VK_SHADER_STAGE_FRAGMENT_BIT, sizeof(float) * 4>))),
       full_screen_vert(BuildShader(device, FULL_SCREEN_TRIANGLE_VERT_SPV)),
       blit_color_to_color_frag(BuildShader(device, BLIT_COLOR_FLOAT_FRAG_SPV)),
-      blit_depth_stencil_frag(device.IsExtShaderStencilExportSupported()
-                             ? BuildShader(device, VULKAN_BLIT_DEPTH_STENCIL_FRAG_SPV)
+      blit_depth_stencil_frag(device.IsExtShaderStencilExportSupported() 
+                             ? BuildShader(device, VULKAN_BLIT_DEPTH_STENCIL_FRAG_SPV) 
                              : vk::ShaderModule{}),
       clear_color_vert(BuildShader(device, VULKAN_COLOR_CLEAR_VERT_SPV)),
       clear_color_frag(BuildShader(device, VULKAN_COLOR_CLEAR_FRAG_SPV)),
       clear_stencil_frag(BuildShader(device, VULKAN_DEPTHSTENCIL_CLEAR_FRAG_SPV)),
       convert_depth_to_float_frag(BuildShader(device, CONVERT_DEPTH_TO_FLOAT_FRAG_SPV)),
       convert_float_to_depth_frag(BuildShader(device, CONVERT_FLOAT_TO_DEPTH_FRAG_SPV)),
-      convert_abgr8_to_d24s8_frag(device.IsExtShaderStencilExportSupported()
-                                 ? BuildShader(device, CONVERT_ABGR8_TO_D24S8_FRAG_SPV)
+      convert_abgr8_to_d24s8_frag(device.IsExtShaderStencilExportSupported() 
+                                 ? BuildShader(device, CONVERT_ABGR8_TO_D24S8_FRAG_SPV) 
                                  : vk::ShaderModule{}),
       convert_abgr8_to_d32f_frag(BuildShader(device, CONVERT_ABGR8_TO_D32F_FRAG_SPV)),
       convert_d32f_to_abgr8_frag(BuildShader(device, CONVERT_D32F_TO_ABGR8_FRAG_SPV)),

src/video_core/renderer_vulkan/fixed_pipeline_state.h

@@ -30,6 +30,7 @@ struct DynamicFeatures {
     bool has_dynamic_state3_depth_clamp_enable;
     bool has_dynamic_state3_logic_op_enable;
     bool has_dynamic_state3_line_stipple_enable;
+    bool has_dual_source_blend;
     bool has_dynamic_vertex_input;
     bool has_provoking_vertex;
     bool has_provoking_vertex_first_mode;

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

@@ -193,12 +193,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() &&
-                                            !image_view.SupportsAnisotropy()};
-            const VkSampler vk_sampler{use_fallback_sampler ? sampler.HandleWithDefaultAnisotropy()
-                                                            : sampler.Handle()};
+            const bool supports_linear_filter{
+                texture_cache.SupportsLinearFilter(image_view.format)};
+            const bool supports_depth_compare_sampling{
+                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

@@ -283,7 +283,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>
 
@@ -335,8 +336,8 @@ BufferCacheRuntime::BufferCacheRuntime(const Device& device_, MemoryAllocator& m
       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_ARM_PROPRIETARY;
+    limit_dynamic_storage_buffers = driver_id == VK_DRIVER_ID_QUALCOMM_PROPRIETARY
+        || driver_id == VK_DRIVER_ID_ARM_PROPRIETARY;
     if (limit_dynamic_storage_buffers) {
         max_dynamic_storage_buffers = device.GetMaxDescriptorSetStorageBuffersDynamic();
     }    
@@ -556,7 +557,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;
@@ -596,7 +605,9 @@ 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
+    // When VIDS is active, the pipeline doesn't declare VERTEX_INPUT_BINDING_STRIDE as dynamic
+    if (device.IsExtExtendedDynamicStateSupported() && !device.IsExtVertexInputDynamicStateSupported()) {
         scheduler.Record([bindings_ = std::move(bindings), buffer_handles_ = std::move(buffer_handles), binding_count](vk::CommandBuffer cmdbuf) {
             cmdbuf.BindVertexBuffers2EXT(bindings_.min_index, binding_count, buffer_handles_.data(), bindings_.offsets.data(), bindings_.sizes.data(), bindings_.strides.data());
         });
@@ -650,27 +661,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 |
-                 VK_BUFFER_USAGE_TRANSFER_DST_BIT | VK_BUFFER_USAGE_INDIRECT_BUFFER_BIT,
+        .usage = NullBufferUsageFlags(),
         .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

@@ -131,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) {
@@ -168,6 +171,7 @@ private:
 
     void ReserveNullBuffer();
     vk::Buffer CreateNullBuffer();
+    VkBufferUsageFlags NullBufferUsageFlags() const;
 
     const Device& device;
     MemoryAllocator& memory_allocator;
@@ -179,6 +183,7 @@ 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;

src/video_core/renderer_vulkan/vk_compute_pass.cpp

@@ -473,10 +473,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_ARITHMETIC_BIT) &&
-                  device_.IsSubgroupFeatureSupported(VK_SUBGROUP_FEATURE_SHUFFLE_BIT) &&
-                  device_.IsSubgroupFeatureSupported(VK_SUBGROUP_FEATURE_SHUFFLE_RELATIVE_BIT)
+          device_.IsSubgroupFeatureSupported(VK_SUBGROUP_FEATURE_BASIC_BIT,
+                                             VK_SHADER_STAGE_COMPUTE_BIT) &&
+                  device_.IsSubgroupFeatureSupported(VK_SUBGROUP_FEATURE_ARITHMETIC_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

@@ -19,16 +19,51 @@
 #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 "video_core/gpu_logging/gpu_logging.h"
 #include "common/settings.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,
@@ -193,8 +228,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>
 
@@ -24,7 +26,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"
 #include "video_core/gpu_logging/gpu_logging.h"
 #include "common/settings.h"
@@ -47,10 +51,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 INLINE_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) {
@@ -426,8 +503,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;
@@ -780,11 +861,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,
@@ -793,6 +876,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,
@@ -814,6 +899,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,
@@ -827,13 +918,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()),
-            .dstColorBlendFactor = MaxwellToVK::BlendFactor(blend.DestRGBFactor()),
+            .srcColorBlendFactor = MaxwellToVK::BlendFactor(sanitize_factor(blend.SourceRGBFactor())),
+            .dstColorBlendFactor = MaxwellToVK::BlendFactor(sanitize_factor(blend.DestRGBFactor())),
             .colorBlendOp = MaxwellToVK::BlendEquation(blend.EquationRGB()),
-            .srcAlphaBlendFactor = MaxwellToVK::BlendFactor(blend.SourceAlphaFactor()),
-            .dstAlphaBlendFactor = MaxwellToVK::BlendFactor(blend.DestAlphaFactor()),
+            .srcAlphaBlendFactor = MaxwellToVK::BlendFactor(sanitize_factor(blend.SourceAlphaFactor())),
+            .dstAlphaBlendFactor = MaxwellToVK::BlendFactor(sanitize_factor(blend.DestAlphaFactor())),
             .alphaBlendOp = MaxwellToVK::BlendEquation(blend.EquationAlpha()),
             .colorWriteMask = write_mask,
         });
@@ -855,6 +963,18 @@ void GraphicsPipeline::MakePipeline(VkRenderPass render_pass) {
         VK_DYNAMIC_STATE_STENCIL_WRITE_MASK, VK_DYNAMIC_STATE_STENCIL_REFERENCE,
         VK_DYNAMIC_STATE_LINE_WIDTH,
     };
+
+    if (device.UsesAdvancedCoreDynamicState()) {
+        static constexpr std::array core_dynamic_states{
+            VK_DYNAMIC_STATE_BLEND_CONSTANTS,
+            VK_DYNAMIC_STATE_DEPTH_BOUNDS,
+            VK_DYNAMIC_STATE_STENCIL_COMPARE_MASK,
+            VK_DYNAMIC_STATE_STENCIL_WRITE_MASK,
+            VK_DYNAMIC_STATE_STENCIL_REFERENCE,
+        };
+        dynamic_states.insert(dynamic_states.end(), core_dynamic_states.begin(),
+                              core_dynamic_states.end());
+    }
     if (key.state.extended_dynamic_state) {
         static constexpr std::array extended{
             VK_DYNAMIC_STATE_CULL_MODE_EXT,
@@ -867,20 +987,20 @@ void GraphicsPipeline::MakePipeline(VkRenderPass render_pass) {
             VK_DYNAMIC_STATE_STENCIL_OP_EXT,
         };
         dynamic_states.insert(dynamic_states.end(), extended.begin(), extended.end());
-
+        
         // VK_DYNAMIC_STATE_VERTEX_INPUT_BINDING_STRIDE_EXT is part of EDS1
         // Only use it if VIDS is not active (VIDS replaces it with full vertex input control)
         if (!key.state.dynamic_vertex_input) {
             dynamic_states.push_back(VK_DYNAMIC_STATE_VERTEX_INPUT_BINDING_STRIDE_EXT);
         }
     }
-
+    
     // VK_DYNAMIC_STATE_VERTEX_INPUT_EXT (VIDS) - Independent from EDS
     // Provides full dynamic vertex input control, replaces VERTEX_INPUT_BINDING_STRIDE
     if (key.state.dynamic_vertex_input) {
         dynamic_states.push_back(VK_DYNAMIC_STATE_VERTEX_INPUT_EXT);
     }
-
+    
     // EDS2 - Core (3 states)
     if (key.state.extended_dynamic_state_2) {
         static constexpr std::array extended2{
@@ -890,12 +1010,12 @@ void GraphicsPipeline::MakePipeline(VkRenderPass render_pass) {
         };
         dynamic_states.insert(dynamic_states.end(), extended2.begin(), extended2.end());
     }
-
+    
     // EDS2 - LogicOp (granular)
     if (key.state.extended_dynamic_state_2_logic_op) {
         dynamic_states.push_back(VK_DYNAMIC_STATE_LOGIC_OP_EXT);
     }
-
+    
     // EDS3 - Blending (composite: 3 states)
     if (key.state.extended_dynamic_state_3_blend) {
         static constexpr std::array extended3{
@@ -905,7 +1025,7 @@ void GraphicsPipeline::MakePipeline(VkRenderPass render_pass) {
         };
         dynamic_states.insert(dynamic_states.end(), extended3.begin(), extended3.end());
     }
-
+    
     // EDS3 - Enables (composite: per-feature)
     if (key.state.extended_dynamic_state_3_enables) {
         if (device.SupportsDynamicState3DepthClampEnable()) {

src/video_core/renderer_vulkan/vk_pipeline_cache.cpp

@@ -37,6 +37,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"
@@ -108,6 +109,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;
@@ -325,8 +361,7 @@ 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 one core free on Android. Previously we reserved two, but
-    // shipping builds benefit from one extra compilation worker.
+    // Leave at least one core free on Android to reduce thermal pressure.
     constexpr size_t free_cores = 1ULL;
     if (max_core_threads <= free_cores) {
         return 1ULL;
@@ -373,6 +408,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(),
@@ -382,20 +418,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_separate_denorm_behavior =
+        .support_float_controls = float_controls_supported,
+        .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_fp32_denorm_preserve = float_control.shaderDenormPreserveFloat32 != VK_FALSE,
-        .support_fp16_denorm_flush = float_control.shaderDenormFlushToZeroFloat16 != VK_FALSE,
-        .support_fp32_denorm_flush = float_control.shaderDenormFlushToZeroFloat32 != VK_FALSE,
-        .support_fp16_signed_zero_nan_preserve =
+        .support_fp16_denorm_preserve = float_controls_supported &&
+            float_control.shaderDenormPreserveFloat16 != VK_FALSE,
+        .support_fp32_denorm_preserve = float_controls_supported &&
+            float_control.shaderDenormPreserveFloat32 != VK_FALSE,
+        .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),
@@ -451,6 +491,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);
@@ -461,28 +522,33 @@ PipelineCache::PipelineCache(Tegra::MaxwellDeviceMemoryManager& device_memory_,
     }
 
     LOG_INFO(Render_Vulkan, "DynamicState setting value: {}", u32(Settings::values.dyna_state.GetValue()));
-
+    
     dynamic_features = {};
-
+    
     // User granularity enforced in vulkan_device.cpp switch statement:
     //   Level 0: Core Dynamic States only
     //   Level 1: Core + EDS1
     //   Level 2: Core + EDS1 + EDS2 (accumulative)
     //   Level 3: Core + EDS1 + EDS2 + EDS3 (accumulative)
     // Here we only verify if extensions were successfully loaded by the device
-
-    dynamic_features.has_extended_dynamic_state =
+    
+    dynamic_features.has_extended_dynamic_state = 
         device.IsExtExtendedDynamicStateSupported();
-
-    dynamic_features.has_extended_dynamic_state_2 =
+    
+    dynamic_features.has_extended_dynamic_state_2 = 
         device.IsExtExtendedDynamicState2Supported();
-    dynamic_features.has_extended_dynamic_state_2_logic_op =
+    dynamic_features.has_extended_dynamic_state_2_logic_op = 
         device.IsExtExtendedDynamicState2ExtrasSupported();
     dynamic_features.has_extended_dynamic_state_2_patch_control_points = false;
-
-    dynamic_features.has_extended_dynamic_state_3_blend =
+    
+    dynamic_features.has_extended_dynamic_state_3_blend = 
         device.IsExtExtendedDynamicState3BlendingSupported();
-    dynamic_features.has_extended_dynamic_state_3_enables =
+    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();
     dynamic_features.has_dynamic_state3_depth_clamp_enable =
         device.SupportsDynamicState3DepthClampEnable();
@@ -492,8 +558,8 @@ PipelineCache::PipelineCache(Tegra::MaxwellDeviceMemoryManager& device_memory_,
         device.SupportsDynamicState3LineStippleEnable();
 
     // VIDS: Independent toggle (not affected by dyna_state levels)
-    dynamic_features.has_dynamic_vertex_input =
-        device.IsExtVertexInputDynamicStateSupported() &&
+    dynamic_features.has_dynamic_vertex_input = 
+        device.IsExtVertexInputDynamicStateSupported() && 
         Settings::values.vertex_input_dynamic_state.GetValue();
 
     dynamic_features.has_provoking_vertex = device.IsExtProvokingVertexSupported();
@@ -512,6 +578,13 @@ PipelineCache::~PipelineCache() {
     }
 }
 
+void PipelineCache::DrainPendingBuilds() {
+    if (!device.HasBrokenParallelShaderCompiling()) {
+        return;
+    }
+    workers.WaitForRequests();
+}
+
 GraphicsPipeline* PipelineCache::CurrentGraphicsPipeline() {
 
     if (!RefreshStages(graphics_key.unique_hashes)) {
@@ -542,12 +615,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) {
-        return pipeline.get();
+    if (!pipeline) {
+        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();
 }
 
@@ -671,13 +749,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());
@@ -700,6 +785,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;
 }
 
@@ -814,6 +900,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() {
@@ -915,6 +1005,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,
@@ -1012,4 +1106,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

@@ -8,8 +8,10 @@
 
 #include <array>
 #include <cstddef>
+#include <condition_variable>
 #include <filesystem>
 #include <memory>
+#include <mutex>
 #include <type_traits>
 #include <ankerl/unordered_dense.h>
 #include <vector>
@@ -116,7 +118,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;
@@ -143,6 +155,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;
@@ -160,6 +180,11 @@ private:
     ankerl::unordered_dense::map<ComputePipelineCacheKey, std::unique_ptr<ComputePipeline>> compute_cache;
     ankerl::unordered_dense::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();
-        dev.ResetQueryPool(*query_pool, 0, BANK_SIZE);
+        if (supports_host_query_reset) {
+            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;
 };

src/video_core/renderer_vulkan/vk_rasterizer.cpp

@@ -235,6 +235,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;
@@ -1028,7 +1032,6 @@ void RasterizerVulkan::UpdateDynamicStates() {
     UpdateDepthBounds(regs);
     UpdateStencilFaces(regs);
     UpdateLineWidth(regs);
-
     // EDS1: CullMode, DepthCompare, FrontFace, StencilOp, DepthBoundsTest, DepthTest, DepthWrite, StencilTest
     if (device.IsExtExtendedDynamicStateSupported()) {
         UpdateCullMode(regs);
@@ -1042,19 +1045,19 @@ void RasterizerVulkan::UpdateDynamicStates() {
             UpdateStencilTestEnable(regs);
         }
     }
-
+    
     // EDS2: PrimitiveRestart, RasterizerDiscard, DepthBias enable/disable
     if (device.IsExtExtendedDynamicState2Supported()) {
         UpdatePrimitiveRestartEnable(regs);
         UpdateRasterizerDiscardEnable(regs);
         UpdateDepthBiasEnable(regs);
     }
-
+    
     // EDS2 Extras: LogicOp operation selection
     if (device.IsExtExtendedDynamicState2ExtrasSupported()) {
         UpdateLogicOp(regs);
     }
-
+    
     // EDS3 Enables: LogicOpEnable, DepthClamp, LineStipple, ConservativeRaster
     if (device.IsExtExtendedDynamicState3EnablesSupported()) {
         using namespace Tegra::Engines;
@@ -1079,12 +1082,12 @@ void RasterizerVulkan::UpdateDynamicStates() {
         UpdateAlphaToCoverageEnable(regs);
         UpdateAlphaToOneEnable(regs);
     }
-
+    
     // EDS3 Blending: ColorBlendEnable, ColorBlendEquation, ColorWriteMask
     if (device.IsExtExtendedDynamicState3BlendingSupported()) {
         UpdateBlending(regs);
     }
-
+    
     // Vertex Input Dynamic State: Independent from EDS levels
     if (device.IsExtVertexInputDynamicStateSupported()) {
         if (auto* gp = pipeline_cache.CurrentGraphicsPipeline(); gp && gp->HasDynamicVertexInput()) {
@@ -1121,7 +1124,7 @@ void RasterizerVulkan::HandleTransformFeedback() {
         UNIMPLEMENTED_IF(regs.IsShaderConfigEnabled(Maxwell::ShaderType::TessellationInit) ||
                          regs.IsShaderConfigEnabled(Maxwell::ShaderType::Tessellation));
     }
-}
+} 
 
 void RasterizerVulkan::UpdateViewportsState(Tegra::Engines::Maxwell3D::Regs& regs) {
     if (!state_tracker.TouchViewports()) {
@@ -1281,109 +1284,141 @@ void RasterizerVulkan::UpdateBlendConstants(Tegra::Engines::Maxwell3D::Regs& reg
     if (!state_tracker.TouchBlendConstants()) {
         return;
     }
-    const std::array blend_color = {regs.blend_color.r, regs.blend_color.g, regs.blend_color.b,
-                                    regs.blend_color.a};
-    scheduler.Record(
-        [blend_color](vk::CommandBuffer cmdbuf) { cmdbuf.SetBlendConstants(blend_color.data()); });
+    if (!device.UsesAdvancedCoreDynamicState()) {
+        return;
+    }
+    const std::array<float, 4> blend_constants{
+        regs.blend_color.r,
+        regs.blend_color.g,
+        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;
     }
-    scheduler.Record([min = regs.depth_bounds[0], max = regs.depth_bounds[1]](
-                         vk::CommandBuffer cmdbuf) { cmdbuf.SetDepthBounds(min, max); });
+    if (!device.IsDepthBoundsSupported()) {
+        return;
+    }
+    if (!device.UsesAdvancedCoreDynamicState()) {
+        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;
     }
-    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 (!device.UsesAdvancedCoreDynamicState()) {
+        state_tracker.ClearStencilReset();
+        return;
     }
+
     if (update_references) {
-        [&]() {
-            if (regs.stencil_two_side_enable) {
-                if (!state_tracker.CheckStencilReferenceFront(regs.stencil_front_ref) &&
-                    !state_tracker.CheckStencilReferenceBack(regs.stencil_back_ref)) {
-                    return;
-                }
-            } else {
-                if (!state_tracker.CheckStencilReferenceFront(regs.stencil_front_ref)) {
-                    return;
-                }
+        if (two_sided) {
+            const bool front_dirty =
+                state_tracker.CheckStencilReferenceFront(regs.stencil_front_ref);
+            const bool back_dirty =
+                state_tracker.CheckStencilReferenceBack(regs.stencil_back_ref);
+            if (front_dirty || back_dirty) {
+                scheduler.Record([front_ref = regs.stencil_front_ref,
+                                  back_ref = regs.stencil_back_ref,
+                                  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,
-                              two_sided = regs.stencil_two_side_enable](vk::CommandBuffer cmdbuf) {
-                const bool set_back = two_sided && front_ref != back_ref;
-                // Front face
-                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);
-                }
+        } else if (state_tracker.CheckStencilReferenceFront(regs.stencil_front_ref)) {
+            const u32 reference = regs.stencil_front_ref;
+            scheduler.Record([reference](vk::CommandBuffer cmdbuf) {
+                cmdbuf.SetStencilReference(VK_STENCIL_FACE_FRONT_AND_BACK, reference);
             });
-        }();
+        }
     }
-    if (update_write_mask) {
-        [&]() {
-            if (regs.stencil_two_side_enable) {
-                if (!state_tracker.CheckStencilWriteMaskFront(regs.stencil_front_mask) &&
-                    !state_tracker.CheckStencilWriteMaskBack(regs.stencil_back_mask)) {
-                    return;
-                }
-            } else {
-                if (!state_tracker.CheckStencilWriteMaskFront(regs.stencil_front_mask)) {
-                    return;
-                }
+
+    if (update_write_masks) {
+        if (two_sided) {
+            if (state_tracker.CheckStencilWriteMaskFront(regs.stencil_front_mask) ||
+                state_tracker.CheckStencilWriteMaskBack(regs.stencil_back_mask)) {
+                scheduler.Record([
+                                     front_write_mask = regs.stencil_front_mask,
+                                     back_write_mask = regs.stencil_back_mask,
+                                     is_two_sided = regs.stencil_two_side_enable
+                                 ](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,
-                              back_write_mask = regs.stencil_back_mask,
-                              two_sided = regs.stencil_two_side_enable](vk::CommandBuffer cmdbuf) {
-                const bool set_back = two_sided && front_write_mask != back_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);
-                }
+        } else if (state_tracker.CheckStencilWriteMaskFront(regs.stencil_front_mask)) {
+            const u32 front_write_mask = regs.stencil_front_mask;
+            scheduler.Record([front_write_mask](vk::CommandBuffer cmdbuf) {
+                cmdbuf.SetStencilWriteMask(VK_STENCIL_FACE_FRONT_AND_BACK, front_write_mask);
             });
-        }();
+        }
     }
+
     if (update_compare_masks) {
-        [&]() {
-            if (regs.stencil_two_side_enable) {
-                if (!state_tracker.CheckStencilCompareMaskFront(regs.stencil_front_func_mask) &&
-                    !state_tracker.CheckStencilCompareMaskBack(regs.stencil_back_func_mask)) {
-                    return;
-                }
-            } else {
-                if (!state_tracker.CheckStencilCompareMaskFront(regs.stencil_front_func_mask)) {
-                    return;
-                }
+        if (two_sided) {
+            if (state_tracker.CheckStencilCompareMaskFront(regs.stencil_front_func_mask) ||
+                state_tracker.CheckStencilCompareMaskBack(regs.stencil_back_func_mask)) {
+                scheduler.Record([
+                                     front_test_mask = regs.stencil_front_func_mask,
+                                     back_test_mask = regs.stencil_back_func_mask,
+                                     is_two_sided = regs.stencil_two_side_enable
+                                 ](vk::CommandBuffer cmdbuf) {
+                    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,
-                              back_test_mask = regs.stencil_back_func_mask,
-                              two_sided = regs.stencil_two_side_enable](vk::CommandBuffer cmdbuf) {
-                const bool set_back = two_sided && front_test_mask != back_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);
-                }
+        } else if (state_tracker.CheckStencilCompareMaskFront(regs.stencil_front_func_mask)) {
+            const u32 front_test_mask = regs.stencil_front_func_mask;
+            scheduler.Record([front_test_mask](vk::CommandBuffer cmdbuf) {
+                cmdbuf.SetStencilCompareMask(VK_STENCIL_FACE_FRONT_AND_BACK, front_test_mask);
             });
-        }();
+        }
     }
+
     state_tracker.ClearStencilReset();
 }
 
@@ -1406,75 +1441,15 @@ void RasterizerVulkan::UpdateCullMode(Tegra::Engines::Maxwell3D::Regs& regs) {
     });
 }
 
-void RasterizerVulkan::UpdateDepthBoundsTestEnable(Tegra::Engines::Maxwell3D::Regs& regs) {
-    if (!state_tracker.TouchDepthBoundsTestEnable()) {
-        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;
-    }
-
-    bool enable = regs.primitive_restart.enabled != 0;
-    if (device.IsMoltenVK()) {
-        enable = true;
-    } else if (enable) {
-        const auto topology =
-            MaxwellToVK::PrimitiveTopology(device, maxwell3d->draw_manager->GetDrawState().topology);
-        enable = IsPrimitiveRestartSupported(device, topology);
-    }
-
-    scheduler.Record([enable](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) {
+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
@@ -1533,6 +1508,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;
@@ -1641,6 +1635,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"
 
@@ -169,7 +170,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>
@@ -92,6 +93,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;
@@ -300,6 +309,7 @@ u64 Scheduler::SubmitExecution(VkSemaphore signal_semaphore, VkSemaphore wait_se
     });
     chunk->MarkSubmit();
     DispatchWork();
+    last_submission_time = Clock::now();
     return signal_value;
 }
 
@@ -333,7 +343,9 @@ void Scheduler::EndRenderPass()
         }
 
         query_cache->CounterEnable(VideoCommon::QueryType::ZPassPixelCount64, false);
-        query_cache->NotifySegment(false);
+        if (supports_transform_feedback) {
+            query_cache->CounterEnable(VideoCommon::QueryType::StreamingByteCount, 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>
@@ -53,6 +54,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();
 
@@ -159,6 +163,8 @@ public:
     std::mutex submit_mutex;
 
 private:
+    using Clock = std::chrono::steady_clock;
+
     class Command {
     public:
         virtual ~Command() = default;
@@ -281,6 +287,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) {
-        return GetStreamBuffer(size);
+    if (!deferred && usage == MemoryUsage::Upload) {
+        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.cpp

@@ -38,6 +38,7 @@ namespace Vulkan {
 
 using Tegra::Engines::Fermi2D;
 using Tegra::Texture::SwizzleSource;
+using Tegra::Texture::TextureFilter;
 using Tegra::Texture::TextureMipmapFilter;
 using VideoCommon::BufferImageCopy;
 using VideoCommon::ImageFlagBits;
@@ -45,12 +46,32 @@ using VideoCommon::ImageInfo;
 using VideoCommon::ImageType;
 using VideoCommon::SubresourceRange;
 using VideoCore::Surface::BytesPerBlock;
+using VideoCore::Surface::DefaultBlockHeight;
+using VideoCore::Surface::DefaultBlockWidth;
 using VideoCore::Surface::HasAlpha;
 using VideoCore::Surface::IsPixelFormatASTC;
 using VideoCore::Surface::IsPixelFormatInteger;
+using VideoCore::Surface::IsPixelFormatSRGB;
 using VideoCore::Surface::SurfaceType;
+using VideoCore::Surface::PixelFormat;
+using VideoCore::Surface::PixelFormatNumeric;
 
 namespace {
+
+[[nodiscard]] std::optional<PixelFormatNumeric>
+ComponentNumericType(Shader::SamplerComponentType component) {
+    switch (component) {
+    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;
+    }
+}
+
 constexpr VkBorderColor ConvertBorderColor(const std::array<float, 4>& color) {
     if (color == std::array<float, 4>{0, 0, 0, 0}) {
         return VK_BORDER_COLOR_FLOAT_TRANSPARENT_BLACK;
@@ -103,6 +124,19 @@ constexpr VkBorderColor ConvertBorderColor(const std::array<float, 4>& color) {
     }
 }
 
+[[nodiscard]] VkFormat LinearizeSrgbFormat(VkFormat format) {
+    switch (format) {
+    case VK_FORMAT_A8B8G8R8_SRGB_PACK32:
+        return VK_FORMAT_A8B8G8R8_UNORM_PACK32;
+    case VK_FORMAT_R8G8B8A8_SRGB:
+        return VK_FORMAT_R8G8B8A8_UNORM;
+    case VK_FORMAT_B8G8R8A8_SRGB:
+        return VK_FORMAT_B8G8R8A8_UNORM;
+    default:
+        return format;
+    }
+}
+
 [[nodiscard]] VkImageUsageFlags ImageUsageFlags(const MaxwellToVK::FormatInfo& info,
                                                 PixelFormat format) {
     VkImageUsageFlags usage = VK_IMAGE_USAGE_TRANSFER_SRC_BIT | VK_IMAGE_USAGE_TRANSFER_DST_BIT |
@@ -139,13 +173,19 @@ constexpr VkBorderColor ConvertBorderColor(const std::array<float, 4>& color) {
     if (info.type == ImageType::e3D) {
         flags |= VK_IMAGE_CREATE_2D_ARRAY_COMPATIBLE_BIT;
     }
+
+    VkFormat base_format = format_info.format;
+    if (IsPixelFormatSRGB(info.format)) {
+        base_format = LinearizeSrgbFormat(base_format);
+    }
+
     const auto [samples_x, samples_y] = VideoCommon::SamplesLog2(info.num_samples);
     return VkImageCreateInfo{
         .sType = VK_STRUCTURE_TYPE_IMAGE_CREATE_INFO,
         .pNext = nullptr,
         .flags = flags,
         .imageType = ConvertImageType(info.type),
-        .format = format_info.format,
+        .format = base_format,
         .extent{
             .width = info.size.width >> samples_x,
             .height = info.size.height >> samples_y,
@@ -164,7 +204,8 @@ constexpr VkBorderColor ConvertBorderColor(const std::array<float, 4>& color) {
 }
 
 [[nodiscard]] vk::Image MakeImage(const Device& device, const MemoryAllocator& allocator,
-                                  const ImageInfo& info, std::span<const VkFormat> view_formats) {
+                                  const ImageInfo& info, std::span<const VkFormat> view_formats,
+                                  bool needs_block_compatible_views) {
     if (info.type == ImageType::Buffer) {
         return vk::Image{};
     }
@@ -324,31 +365,6 @@ void SanitizeDepthStencilSwizzle(std::array<SwizzleSource, 4>& swizzle,
     return VK_IMAGE_VIEW_TYPE_2D;
 }
 
-[[nodiscard]] VkImageViewType ImageViewType(VideoCommon::ImageViewType type) {
-    switch (type) {
-    case VideoCommon::ImageViewType::e1D:
-        return VK_IMAGE_VIEW_TYPE_1D;
-    case VideoCommon::ImageViewType::e2D:
-    case VideoCommon::ImageViewType::Rect:
-        return VK_IMAGE_VIEW_TYPE_2D;
-    case VideoCommon::ImageViewType::Cube:
-        return VK_IMAGE_VIEW_TYPE_CUBE;
-    case VideoCommon::ImageViewType::e3D:
-        return VK_IMAGE_VIEW_TYPE_3D;
-    case VideoCommon::ImageViewType::e1DArray:
-        return VK_IMAGE_VIEW_TYPE_1D_ARRAY;
-    case VideoCommon::ImageViewType::e2DArray:
-        return VK_IMAGE_VIEW_TYPE_2D_ARRAY;
-    case VideoCommon::ImageViewType::CubeArray:
-        return VK_IMAGE_VIEW_TYPE_CUBE_ARRAY;
-    case VideoCommon::ImageViewType::Buffer:
-        ASSERT_MSG(false, "Texture buffers can't be image views");
-        return VK_IMAGE_VIEW_TYPE_1D;
-    }
-    ASSERT_MSG(false, "Invalid image view type={}", type);
-    return VK_IMAGE_VIEW_TYPE_2D;
-}
-
 [[nodiscard]] VkImageSubresourceLayers MakeImageSubresourceLayers(
     VideoCommon::SubresourceLayers subresource, VkImageAspectFlags aspect_mask) {
     return VkImageSubresourceLayers{
@@ -386,6 +402,17 @@ void SanitizeDepthStencilSwizzle(std::array<SwizzleSource, 4>& swizzle,
     };
 }
 
+[[nodiscard]] VkImageResolve MakeImageResolve(const VideoCommon::ImageCopy& copy,
+                                              VkImageAspectFlags aspect_mask) noexcept {
+    return VkImageResolve{
+        .srcSubresource = MakeImageSubresourceLayers(copy.src_subresource, aspect_mask),
+        .srcOffset = MakeOffset3D(copy.src_offset),
+        .dstSubresource = MakeImageSubresourceLayers(copy.dst_subresource, aspect_mask),
+        .dstOffset = MakeOffset3D(copy.dst_offset),
+        .extent = MakeExtent3D(copy.extent),
+    };
+}
+
 [[nodiscard]] VkBufferImageCopy MakeBufferImageCopy(const VideoCommon::ImageCopy& copy, bool is_src,
                                                     VkImageAspectFlags aspect_mask) noexcept {
     return VkBufferImageCopy{
@@ -470,6 +497,15 @@ TransformBufferCopies(std::span<const VideoCommon::BufferCopy> copies, size_t bu
     };
 }
 
+[[nodiscard]] VkImageSubresourceRange Expand3DSubresourceRange(VkImageSubresourceRange range,
+                                                               bool is_3d_image) {
+    if (is_3d_image) {
+        range.baseArrayLayer = 0;
+        range.layerCount = VK_REMAINING_ARRAY_LAYERS;
+    }
+    return range;
+}
+
 [[nodiscard]] VkImageSubresourceRange MakeSubresourceRange(const ImageView* image_view) {
     SubresourceRange range = image_view->range;
     if (True(image_view->flags & VideoCommon::ImageViewFlagBits::Slice)) {
@@ -478,7 +514,8 @@ TransformBufferCopies(std::span<const VideoCommon::BufferCopy> copies, size_t bu
         range.base.layer = 0;
         range.extent.layers = 1;
     }
-    return MakeSubresourceRange(ImageAspectMask(image_view->format), range);
+    return Expand3DSubresourceRange(MakeSubresourceRange(ImageAspectMask(image_view->format), range),
+                                    image_view->Is3DImage());
 }
 
 [[nodiscard]] VkImageSubresourceLayers MakeSubresourceLayers(const ImageView* image_view) {
@@ -559,7 +596,7 @@ struct RangedBarrierRange {
     }
 };
 void CopyBufferToImage(vk::CommandBuffer cmdbuf, VkBuffer src_buffer, VkImage image,
-                       VkImageAspectFlags aspect_mask, bool is_initialized,
+                       VkImageAspectFlags aspect_mask, bool is_initialized, bool is_3d_image,
                        std::span<const VkBufferImageCopy> copies) {
     static constexpr VkAccessFlags WRITE_ACCESS_FLAGS =
                                            VK_ACCESS_SHADER_WRITE_BIT | VK_ACCESS_COLOR_ATTACHMENT_WRITE_BIT |
@@ -573,7 +610,8 @@ void CopyBufferToImage(vk::CommandBuffer cmdbuf, VkBuffer src_buffer, VkImage im
     for (const auto& region : copies) {
         range.AddLayers(region.imageSubresource);
     }
-    const VkImageSubresourceRange subresource_range = range.SubresourceRange(aspect_mask);
+    VkImageSubresourceRange subresource_range = range.SubresourceRange(aspect_mask);
+    subresource_range = Expand3DSubresourceRange(subresource_range, is_3d_image);
 
     const VkImageMemoryBarrier read_barrier{
             .sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER,
@@ -732,7 +770,7 @@ void TryTransformSwizzleIfNeeded(PixelFormat format, std::array<SwizzleSource, 4
 
 void BlitScale(Scheduler& scheduler, VkImage src_image, VkImage dst_image, const ImageInfo& info,
                VkImageAspectFlags aspect_mask, const Settings::ResolutionScalingInfo& resolution,
-               bool up_scaling = true) {
+               bool supports_linear_filter, bool up_scaling = true) {
     const bool is_2d = info.type == ImageType::e2D;
     const auto resources = info.resources;
     const VkExtent2D extent{
@@ -741,7 +779,8 @@ void BlitScale(Scheduler& scheduler, VkImage src_image, VkImage dst_image, const
     };
     // Depth and integer formats must use NEAREST filter for blits.
     const bool is_color{aspect_mask == VK_IMAGE_ASPECT_COLOR_BIT};
-    const bool is_bilinear{is_color && !IsPixelFormatInteger(info.format)};
+    const bool is_bilinear = supports_linear_filter && is_color &&
+                             !IsPixelFormatInteger(info.format);
     const VkFilter vk_filter = is_bilinear ? VK_FILTER_LINEAR : VK_FILTER_NEAREST;
 
     scheduler.RequestOutsideRenderPassOperationContext();
@@ -887,22 +926,42 @@ TextureCacheRuntime::TextureCacheRuntime(const Device& device_, Scheduler& sched
     if (device.IsStorageImageMultisampleSupported()) {
         msaa_copy_pass.emplace(device, scheduler, descriptor_pool, staging_buffer_pool, compute_pass_descriptor_queue);
     }
-    if (!device.IsKhrImageFormatListSupported()) {
-        return;
-    }
     for (size_t index_a = 0; index_a < VideoCore::Surface::MaxPixelFormat; index_a++) {
         const auto image_format = static_cast<PixelFormat>(index_a);
+        const auto image_surface_type = VideoCore::Surface::GetFormatType(image_format);
+        auto& image_view_formats = view_formats[index_a];
+        const auto image_block_width = DefaultBlockWidth(image_format);
+        const auto image_block_height = DefaultBlockHeight(image_format);
+        const auto image_bytes_per_block = BytesPerBlock(image_format);
+        bool needs_block_view = false;
         if (IsPixelFormatASTC(image_format) && !device.IsOptimalAstcSupported()) {
-            view_formats[index_a].push_back(VK_FORMAT_A8B8G8R8_UNORM_PACK32);
+            image_view_formats.push_back(VK_FORMAT_A8B8G8R8_UNORM_PACK32);
+            needs_block_view = true;
         }
         for (size_t index_b = 0; index_b < VideoCore::Surface::MaxPixelFormat; index_b++) {
             const auto view_format = static_cast<PixelFormat>(index_b);
-            if (VideoCore::Surface::IsViewCompatible(image_format, view_format, false, true)) {
-                const auto view_info =
-                    MaxwellToVK::SurfaceFormat(device, FormatType::Optimal, true, view_format);
-                view_formats[index_a].push_back(view_info.format);
+            if (!VideoCore::Surface::IsViewCompatible(image_format, view_format, false, true)) {
+                continue;
+            }
+            if (image_surface_type == SurfaceType::DepthStencil &&
+                VideoCore::Surface::GetFormatType(view_format) == SurfaceType::Stencil &&
+                view_format != image_format) {
+                continue;
+            }
+            const auto view_info =
+                MaxwellToVK::SurfaceFormat(device, FormatType::Optimal, true, view_format);
+            if (std::find(image_view_formats.begin(), image_view_formats.end(), view_info.format) ==
+                image_view_formats.end()) {
+                image_view_formats.push_back(view_info.format);
+            }
+            if (!needs_block_view) {
+                const bool different_block = image_bytes_per_block != BytesPerBlock(view_format) ||
+                                            image_block_width != DefaultBlockWidth(view_format) ||
+                                            image_block_height != DefaultBlockHeight(view_format);
+                needs_block_view = different_block;
             }
         }
+        requires_block_view_formats[index_a] = needs_block_view;
     }
 
     if (Settings::values.gpu_unswizzle_enabled.GetValue()) {
@@ -961,6 +1020,13 @@ void TextureCacheRuntime::FreeDeferredStagingBuffer(StagingBufferRef& ref) {
     staging_buffer_pool.FreeDeferred(ref);
 }
 
+void TextureCacheRuntime::WaitForGpuTick(u64 tick) {
+    if (tick == 0) {
+        return;
+    }
+    scheduler.Wait(tick);
+}
+
 bool TextureCacheRuntime::ShouldReinterpret(Image& dst, Image& src) {
     if (VideoCore::Surface::GetFormatType(dst.info.format) ==
             VideoCore::Surface::SurfaceType::DepthStencil &&
@@ -1146,12 +1212,17 @@ void TextureCacheRuntime::BlitImage(Framebuffer* dst_framebuffer, ImageView& dst
     const VkImageAspectFlags aspect_mask = ImageAspectMask(src.format);
     const bool is_dst_msaa = dst.Samples() != VK_SAMPLE_COUNT_1_BIT;
     const bool is_src_msaa = src.Samples() != VK_SAMPLE_COUNT_1_BIT;
+    const bool supports_linear_filter = SupportsLinearFilter(src.format);
+    const auto effective_filter =
+        (filter == Fermi2D::Filter::Bilinear && supports_linear_filter)
+            ? Fermi2D::Filter::Bilinear
+            : Fermi2D::Filter::Point;
     if (aspect_mask != ImageAspectMask(dst.format)) {
         UNIMPLEMENTED_MSG("Incompatible blit from format {} to {}", src.format, dst.format);
         return;
     }
     if (aspect_mask == VK_IMAGE_ASPECT_COLOR_BIT && !is_src_msaa && !is_dst_msaa) {
-        blit_image_helper.BlitColor(dst_framebuffer, src, dst_region, src_region, filter,
+        blit_image_helper.BlitColor(dst_framebuffer, src, dst_region, src_region, effective_filter,
                                     operation);
         return;
     }
@@ -1174,7 +1245,7 @@ void TextureCacheRuntime::BlitImage(Framebuffer* dst_framebuffer, ImageView& dst
         if (!can_blit_depth_stencil) {
             UNIMPLEMENTED_IF(is_src_msaa || is_dst_msaa);
             blit_image_helper.BlitDepthStencil(dst_framebuffer, src, dst_region, src_region,
-                                               filter, operation);
+                                               effective_filter, operation);
             return;
         }
     }
@@ -1186,7 +1257,7 @@ void TextureCacheRuntime::BlitImage(Framebuffer* dst_framebuffer, ImageView& dst
     const VkImageSubresourceLayers dst_layers = MakeSubresourceLayers(&dst);
     const VkImageSubresourceLayers src_layers = MakeSubresourceLayers(&src);
     const bool is_msaa_to_msaa = is_src_msaa && is_dst_msaa;
-
+    
     // NVIDIA 510+ and Intel crash on MSAA->MSAA blits (scaling operations)
     // Fall back to 3D helpers for MSAA scaling
     if (is_msaa_to_msaa && device.CantBlitMSAA()) {
@@ -1196,8 +1267,8 @@ void TextureCacheRuntime::BlitImage(Framebuffer* dst_framebuffer, ImageView& dst
     }
     const bool is_resolve = is_src_msaa && !is_dst_msaa;
     scheduler.RequestOutsideRenderPassOperationContext();
-    scheduler.Record([filter, dst_region, src_region, dst_image, src_image, dst_layers, src_layers,
-                      aspect_mask, is_resolve](vk::CommandBuffer cmdbuf) {
+    scheduler.Record([effective_filter, dst_region, src_region, dst_image, src_image, dst_layers,
+                      src_layers, aspect_mask, is_resolve](vk::CommandBuffer cmdbuf) {
         const std::array read_barriers{
             VkImageMemoryBarrier{
                 .sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER,
@@ -1268,7 +1339,7 @@ void TextureCacheRuntime::BlitImage(Framebuffer* dst_framebuffer, ImageView& dst
                                 VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL,
                                 MakeImageResolve(dst_region, src_region, dst_layers, src_layers));
         } else {
-            const bool is_linear = filter == Fermi2D::Filter::Bilinear;
+            const bool is_linear = effective_filter == Fermi2D::Filter::Bilinear;
             const VkFilter vk_filter = is_linear ? VK_FILTER_LINEAR : VK_FILTER_NEAREST;
             cmdbuf.BlitImage(
                 src_image, VK_IMAGE_LAYOUT_GENERAL, dst_image, VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL,
@@ -1447,6 +1518,20 @@ bool TextureCacheRuntime::IsFormatScalable(PixelFormat format) {
     }
 }
 
+bool TextureCacheRuntime::SupportsLinearFilter(PixelFormat format) const {
+    if (IsPixelFormatInteger(format)) {
+        return false;
+    }
+    if (VideoCore::Surface::GetFormatType(format) != SurfaceType::ColorTexture) {
+        return false;
+    }
+    const auto vk_format =
+        MaxwellToVK::SurfaceFormat(device, FormatType::Optimal, false, format).format;
+    constexpr VkFormatFeatureFlags linear_filter_feature =
+        VK_FORMAT_FEATURE_SAMPLED_IMAGE_FILTER_LINEAR_BIT;
+    return device.IsFormatSupported(vk_format, linear_filter_feature, FormatType::Optimal);
+}
+
 void TextureCacheRuntime::CopyImage(Image& dst, Image& src,
                                     std::span<const VideoCommon::ImageCopy> copies) {
     // As per the size-compatible formats section of vulkan, copy manually via ReinterpretImage
@@ -1475,14 +1560,21 @@ void TextureCacheRuntime::CopyImage(Image& dst, Image& src,
     });
     const VkImage dst_image = dst.Handle();
     const VkImage src_image = src.Handle();
+    const bool dst_is_3d = dst.info.type == ImageType::e3D;
+    const bool src_is_3d = src.info.type == ImageType::e3D;
     scheduler.RequestOutsideRenderPassOperationContext();
-    scheduler.Record([dst_image, src_image, aspect_mask, vk_copies](vk::CommandBuffer cmdbuf) {
+    scheduler.Record([dst_image, src_image, aspect_mask, vk_copies, dst_is_3d,
+                      src_is_3d](vk::CommandBuffer cmdbuf) {
         RangedBarrierRange dst_range;
         RangedBarrierRange src_range;
         for (const VkImageCopy& copy : vk_copies) {
             dst_range.AddLayers(copy.dstSubresource);
             src_range.AddLayers(copy.srcSubresource);
         }
+        const VkImageSubresourceRange dst_subresource =
+            Expand3DSubresourceRange(dst_range.SubresourceRange(aspect_mask), dst_is_3d);
+        const VkImageSubresourceRange src_subresource =
+            Expand3DSubresourceRange(src_range.SubresourceRange(aspect_mask), src_is_3d);
         const std::array pre_barriers{
             VkImageMemoryBarrier{
                 .sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER,
@@ -1496,7 +1588,7 @@ void TextureCacheRuntime::CopyImage(Image& dst, Image& src,
                 .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
                 .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
                 .image = src_image,
-                .subresourceRange = src_range.SubresourceRange(aspect_mask),
+                .subresourceRange = src_subresource,
             },
             VkImageMemoryBarrier{
                 .sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER,
@@ -1510,7 +1602,7 @@ void TextureCacheRuntime::CopyImage(Image& dst, Image& src,
                 .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
                 .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
                 .image = dst_image,
-                .subresourceRange = dst_range.SubresourceRange(aspect_mask),
+                .subresourceRange = dst_subresource,
             },
         };
         const std::array post_barriers{
@@ -1524,7 +1616,7 @@ void TextureCacheRuntime::CopyImage(Image& dst, Image& src,
                 .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
                 .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
                 .image = src_image,
-                .subresourceRange = src_range.SubresourceRange(aspect_mask),
+                .subresourceRange = src_subresource,
             },
             VkImageMemoryBarrier{
                 .sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER,
@@ -1541,7 +1633,7 @@ void TextureCacheRuntime::CopyImage(Image& dst, Image& src,
                 .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
                 .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
                 .image = dst_image,
-                .subresourceRange = dst_range.SubresourceRange(aspect_mask),
+                .subresourceRange = dst_subresource,
             },
         };
         cmdbuf.PipelineBarrier(
@@ -1562,8 +1654,115 @@ void TextureCacheRuntime::CopyImage(Image& dst, Image& src,
 void TextureCacheRuntime::CopyImageMSAA(Image& dst, Image& src,
                                         std::span<const VideoCommon::ImageCopy> copies) {
     const bool msaa_to_non_msaa = src.info.num_samples > 1 && dst.info.num_samples == 1;
+    if (msaa_to_non_msaa) {
+        const VkImageAspectFlags aspect_mask = dst.AspectMask();
+        ASSERT(aspect_mask == src.AspectMask());
+        boost::container::small_vector<VkImageResolve, 16> vk_resolves(copies.size());
+        std::ranges::transform(copies, vk_resolves.begin(), [aspect_mask](const auto& copy) {
+            return MakeImageResolve(copy, aspect_mask);
+        });
+        const VkImage dst_image = dst.Handle();
+        const VkImage src_image = src.Handle();
+        const bool dst_is_3d = dst.info.type == ImageType::e3D;
+        const bool src_is_3d = src.info.type == ImageType::e3D;
+        scheduler.RequestOutsideRenderPassOperationContext();
+        scheduler.Record([dst_image, src_image, aspect_mask, vk_resolves, dst_is_3d,
+                          src_is_3d](vk::CommandBuffer cmdbuf) {
+            RangedBarrierRange dst_range;
+            RangedBarrierRange src_range;
+            for (const VkImageResolve& region : vk_resolves) {
+                dst_range.AddLayers(region.dstSubresource);
+                src_range.AddLayers(region.srcSubresource);
+            }
+            const VkImageSubresourceRange dst_subresource =
+                Expand3DSubresourceRange(dst_range.SubresourceRange(aspect_mask), dst_is_3d);
+            const VkImageSubresourceRange src_subresource =
+                Expand3DSubresourceRange(src_range.SubresourceRange(aspect_mask), src_is_3d);
+            const std::array pre_barriers{
+                VkImageMemoryBarrier{
+                    .sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER,
+                    .pNext = nullptr,
+                    .srcAccessMask = VK_ACCESS_SHADER_WRITE_BIT |
+                                     VK_ACCESS_COLOR_ATTACHMENT_WRITE_BIT |
+                                     VK_ACCESS_DEPTH_STENCIL_ATTACHMENT_WRITE_BIT |
+                                     VK_ACCESS_TRANSFER_WRITE_BIT,
+                    .dstAccessMask = VK_ACCESS_TRANSFER_READ_BIT,
+                    .oldLayout = VK_IMAGE_LAYOUT_GENERAL,
+                    .newLayout = VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL,
+                    .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
+                    .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
+                    .image = src_image,
+                    .subresourceRange = src_subresource,
+                },
+                VkImageMemoryBarrier{
+                    .sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER,
+                    .pNext = nullptr,
+                    .srcAccessMask = VK_ACCESS_SHADER_WRITE_BIT |
+                                     VK_ACCESS_COLOR_ATTACHMENT_WRITE_BIT |
+                                     VK_ACCESS_DEPTH_STENCIL_ATTACHMENT_WRITE_BIT |
+                                     VK_ACCESS_TRANSFER_WRITE_BIT,
+                    .dstAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT,
+                    .oldLayout = VK_IMAGE_LAYOUT_GENERAL,
+                    .newLayout = VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL,
+                    .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
+                    .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
+                    .image = dst_image,
+                    .subresourceRange = dst_subresource,
+                },
+            };
+            const std::array post_barriers{
+                VkImageMemoryBarrier{
+                    .sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER,
+                    .pNext = nullptr,
+                    .srcAccessMask = 0,
+                    .dstAccessMask = 0,
+                    .oldLayout = VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL,
+                    .newLayout = VK_IMAGE_LAYOUT_GENERAL,
+                    .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
+                    .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
+                    .image = src_image,
+                    .subresourceRange = src_subresource,
+                },
+                VkImageMemoryBarrier{
+                    .sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER,
+                    .pNext = nullptr,
+                    .srcAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT,
+                    .dstAccessMask = VK_ACCESS_SHADER_READ_BIT | VK_ACCESS_SHADER_WRITE_BIT |
+                                     VK_ACCESS_COLOR_ATTACHMENT_READ_BIT |
+                                     VK_ACCESS_COLOR_ATTACHMENT_WRITE_BIT |
+                                     VK_ACCESS_DEPTH_STENCIL_ATTACHMENT_READ_BIT |
+                                     VK_ACCESS_DEPTH_STENCIL_ATTACHMENT_WRITE_BIT |
+                                     VK_ACCESS_TRANSFER_READ_BIT |
+                                     VK_ACCESS_TRANSFER_WRITE_BIT,
+                    .oldLayout = VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL,
+                    .newLayout = VK_IMAGE_LAYOUT_GENERAL,
+                    .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
+                    .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
+                    .image = dst_image,
+                    .subresourceRange = dst_subresource,
+                },
+            };
+            cmdbuf.PipelineBarrier(
+                VK_PIPELINE_STAGE_LATE_FRAGMENT_TESTS_BIT | VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT |
+                    VK_PIPELINE_STAGE_COLOR_ATTACHMENT_OUTPUT_BIT |
+                    VK_PIPELINE_STAGE_TRANSFER_BIT,
+                VK_PIPELINE_STAGE_TRANSFER_BIT,
+                0, nullptr, nullptr, pre_barriers);
+            cmdbuf.ResolveImage(src_image, VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL, dst_image,
+                                VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL,
+                                VideoCommon::FixSmallVectorADL(vk_resolves));
+            cmdbuf.PipelineBarrier(
+                VK_PIPELINE_STAGE_TRANSFER_BIT,
+                VK_PIPELINE_STAGE_LATE_FRAGMENT_TESTS_BIT | VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT |
+                    VK_PIPELINE_STAGE_COLOR_ATTACHMENT_OUTPUT_BIT |
+                    VK_PIPELINE_STAGE_TRANSFER_BIT,
+                0, nullptr, nullptr, post_barriers);
+        });
+        return;
+    }
     if (msaa_copy_pass) {
-        return msaa_copy_pass->CopyImage(dst, src, copies, msaa_to_non_msaa);
+        msaa_copy_pass->CopyImage(dst, src, copies, msaa_to_non_msaa);
+        return;
     }
     UNIMPLEMENTED_MSG("Copying images with different samples is not supported.");
 }
@@ -1588,9 +1787,11 @@ void TextureCacheRuntime::TickFrame() {}
 
 Image::Image(TextureCacheRuntime& runtime_, const ImageInfo& info_, GPUVAddr gpu_addr_,
              VAddr cpu_addr_)
-    : VideoCommon::ImageBase(info_, gpu_addr_, cpu_addr_), scheduler{&runtime_.scheduler},
-      runtime{&runtime_}, original_image(MakeImage(runtime_.device, runtime_.memory_allocator, info,
-                                                   runtime->ViewFormats(info.format))),
+        : VideoCommon::ImageBase(info_, gpu_addr_, cpu_addr_), scheduler{&runtime_.scheduler},
+            runtime{&runtime_},
+            original_image(MakeImage(runtime_.device, runtime_.memory_allocator, info,
+                                                             runtime_.ViewFormats(info.format),
+                                                             runtime_.RequiresBlockCompatibleViewFormats(info.format))),
       aspect_mask(ImageAspectMask(info.format)) {
     if (IsPixelFormatASTC(info.format) && !runtime->device.IsOptimalAstcSupported()) {
         switch (Settings::values.accelerate_astc.GetValue()) {
@@ -1628,6 +1829,8 @@ Image::Image(TextureCacheRuntime& runtime_, const ImageInfo& info_, GPUVAddr gpu
                 MakeStorageView(device, level, *original_image, VK_FORMAT_A8B8G8R8_UNORM_PACK32);
         }
     }
+
+    runtime->TransitionImageLayout(*this);
 }
 
 Image::Image(const VideoCommon::NullImageParams& params) : VideoCommon::ImageBase{params} {}
@@ -1721,9 +1924,10 @@ void Image::UploadMemory(VkBuffer buffer, VkDeviceSize offset,
         const VkImage temp_vk_image = *temp_wrapper->original_image;
         const VkImageAspectFlags vk_aspect_mask = temp_wrapper->aspect_mask;
 
-        scheduler->Record([src_buffer, temp_vk_image, vk_aspect_mask, vk_copies,
+        const bool temp_is_3d = temp_info.type == ImageType::e3D;
+        scheduler->Record([src_buffer, temp_vk_image, vk_aspect_mask, vk_copies, temp_is_3d,
                            keep = temp_wrapper](vk::CommandBuffer cmdbuf) {
-            CopyBufferToImage(cmdbuf, src_buffer, temp_vk_image, vk_aspect_mask, false, VideoCommon::FixSmallVectorADL(vk_copies));
+            CopyBufferToImage(cmdbuf, src_buffer, temp_vk_image, vk_aspect_mask, false, temp_is_3d, VideoCommon::FixSmallVectorADL(vk_copies));
         });
 
         // Use MSAACopyPass to convert from non-MSAA to MSAA
@@ -1760,9 +1964,10 @@ void Image::UploadMemory(VkBuffer buffer, VkDeviceSize offset,
     const VkImageAspectFlags vk_aspect_mask = aspect_mask;
     const bool was_initialized = std::exchange(initialized, true);
 
-    scheduler->Record([src_buffer, vk_image, vk_aspect_mask, was_initialized,
+    const bool is_3d_image = info.type == ImageType::e3D;
+    scheduler->Record([src_buffer, vk_image, vk_aspect_mask, was_initialized, is_3d_image,
                        vk_copies](vk::CommandBuffer cmdbuf) {
-        CopyBufferToImage(cmdbuf, src_buffer, vk_image, vk_aspect_mask, was_initialized, VideoCommon::FixSmallVectorADL(vk_copies));
+        CopyBufferToImage(cmdbuf, src_buffer, vk_image, vk_aspect_mask, was_initialized, is_3d_image, VideoCommon::FixSmallVectorADL(vk_copies));
     });
 
     if (is_rescaled) {
@@ -2005,7 +2210,8 @@ bool Image::ScaleUp(bool ignore) {
         scaled_info.size.width = scaled_width;
         scaled_info.size.height = scaled_height;
         scaled_image = MakeImage(runtime->device, runtime->memory_allocator, scaled_info,
-                                 runtime->ViewFormats(info.format));
+                     runtime->ViewFormats(info.format),
+                     runtime->RequiresBlockCompatibleViewFormats(info.format));
         ignore = false;
     }
     current_image = &Image::scaled_image;
@@ -2018,7 +2224,9 @@ bool Image::ScaleUp(bool ignore) {
     if (NeedsScaleHelper()) {
         return BlitScaleHelper(true);
     } else {
-        BlitScale(*scheduler, *original_image, *scaled_image, info, aspect_mask, resolution);
+        const bool supports_linear_filter = runtime->SupportsLinearFilter(info.format);
+        BlitScale(*scheduler, *original_image, *scaled_image, info, aspect_mask, resolution,
+                  supports_linear_filter);
     }
     return true;
 }
@@ -2043,7 +2251,9 @@ bool Image::ScaleDown(bool ignore) {
     if (NeedsScaleHelper()) {
         return BlitScaleHelper(false);
     } else {
-        BlitScale(*scheduler, *scaled_image, *original_image, info, aspect_mask, resolution, false);
+        const bool supports_linear_filter = runtime->SupportsLinearFilter(info.format);
+        BlitScale(*scheduler, *scaled_image, *original_image, info, aspect_mask, resolution,
+                  supports_linear_filter, false);
     }
     return true;
 }
@@ -2052,9 +2262,21 @@ bool Image::BlitScaleHelper(bool scale_up) {
     using namespace VideoCommon;
     static constexpr auto BLIT_OPERATION = Tegra::Engines::Fermi2D::Operation::SrcCopy;
     const bool is_color{aspect_mask == VK_IMAGE_ASPECT_COLOR_BIT};
-    const bool is_bilinear{is_color && !IsPixelFormatInteger(info.format)};
-    const auto operation = is_bilinear ? Tegra::Engines::Fermi2D::Filter::Bilinear
-                                       : Tegra::Engines::Fermi2D::Filter::Point;
+    const bool has_depth = (aspect_mask & VK_IMAGE_ASPECT_DEPTH_BIT) != 0;
+    const bool has_stencil = (aspect_mask & VK_IMAGE_ASPECT_STENCIL_BIT) != 0;
+    const bool is_depth_only = has_depth && !has_stencil;
+    const bool is_stencil_only = has_stencil && !has_depth;
+    if ((is_depth_only || is_stencil_only) && info.num_samples == 1) {
+        const VkImage src_image = scale_up ? *original_image : *scaled_image;
+        const VkImage dst_image = scale_up ? *scaled_image : *original_image;
+        BlitScale(*scheduler, src_image, dst_image, info, aspect_mask, runtime->resolution,
+                  /*supports_linear_filter=*/false, scale_up);
+        return true;
+    }
+    const bool supports_linear_filter = runtime->SupportsLinearFilter(info.format);
+    const bool is_bilinear = is_color && supports_linear_filter;
+    const auto filter_mode = is_bilinear ? Tegra::Engines::Fermi2D::Filter::Bilinear
+                                         : Tegra::Engines::Fermi2D::Filter::Point;
 
     const bool is_2d = info.type == ImageType::e2D;
     const auto& resolution = runtime->resolution;
@@ -2088,13 +2310,11 @@ bool Image::BlitScaleHelper(bool scale_up) {
     if (aspect_mask == VK_IMAGE_ASPECT_COLOR_BIT) {
         if (!blit_framebuffer)
             blit_framebuffer.emplace(*runtime, view_ptr, nullptr, extent, scale_up);
-        runtime->blit_image_helper.BlitColor(&*blit_framebuffer, *blit_view,
-            dst_region, src_region, operation, BLIT_OPERATION);
+        runtime->blit_image_helper.BlitColor(&blit_framebuffer.value(), *blit_view, dst_region, src_region, filter_mode, BLIT_OPERATION);
     } else if (aspect_mask == (VK_IMAGE_ASPECT_DEPTH_BIT | VK_IMAGE_ASPECT_STENCIL_BIT)) {
         if (!blit_framebuffer)
             blit_framebuffer.emplace(*runtime, nullptr, view_ptr, extent, scale_up);
-        runtime->blit_image_helper.BlitDepthStencil(&*blit_framebuffer, *blit_view,
-            dst_region, src_region, operation, BLIT_OPERATION);
+        runtime->blit_image_helper.BlitDepthStencil(&blit_framebuffer.value(), *blit_view, dst_region, src_region, filter_mode, BLIT_OPERATION);
     } else {
         // TODO: Use helper blits where applicable
         flags &= ~ImageFlagBits::Rescaled;
@@ -2122,7 +2342,8 @@ ImageView::ImageView(TextureCacheRuntime& runtime, const VideoCommon::ImageViewI
                      ImageId image_id_, Image& image)
     : VideoCommon::ImageViewBase{info, image.info, image_id_, image.gpu_addr},
       device{&runtime.device}, image_handle{image.Handle()},
-      samples(ConvertSampleCount(image.info.num_samples)) {
+    samples(ConvertSampleCount(image.info.num_samples)),
+    is_3d_image(image.info.type == ImageType::e3D) {
     using Shader::TextureType;
 
     const VkImageAspectFlags aspect_mask = ImageViewAspectMask(info);
@@ -2166,11 +2387,11 @@ ImageView::ImageView(TextureCacheRuntime& runtime, const VideoCommon::ImageViewI
         },
         .subresourceRange = MakeSubresourceRange(aspect_mask, info.range),
     };
-    const auto create = [&](TextureType tex_type, std::optional<u32> num_layers) {
+    const auto create = [&](TextureType tex_type) {
         VkImageViewCreateInfo ci{create_info};
         ci.viewType = ImageViewType(tex_type);
-        if (num_layers) {
-            ci.subresourceRange.layerCount = *num_layers;
+        if (const auto override_layers = LayerCountOverride(tex_type)) {
+            ci.subresourceRange.layerCount = *override_layers;
         }
         vk::ImageView handle = device->GetLogical().CreateImageView(ci);
         if (device->HasDebuggingToolAttached()) {
@@ -2181,25 +2402,29 @@ ImageView::ImageView(TextureCacheRuntime& runtime, const VideoCommon::ImageViewI
     switch (info.type) {
     case VideoCommon::ImageViewType::e1D:
     case VideoCommon::ImageViewType::e1DArray:
-        create(TextureType::Color1D, 1);
-        create(TextureType::ColorArray1D, std::nullopt);
+        create(TextureType::Color1D);
+        create(TextureType::ColorArray1D);
         render_target = Handle(TextureType::ColorArray1D);
         break;
     case VideoCommon::ImageViewType::e2D:
     case VideoCommon::ImageViewType::e2DArray:
+        create(TextureType::Color2D);
+        create(TextureType::ColorArray2D);
+        render_target = Handle(Shader::TextureType::ColorArray2D);
+        break;
     case VideoCommon::ImageViewType::Rect:
-        create(TextureType::Color2D, 1);
-        create(TextureType::ColorArray2D, std::nullopt);
+        create(TextureType::Color2DRect);
+        create(TextureType::ColorArray2D);
         render_target = Handle(Shader::TextureType::ColorArray2D);
         break;
     case VideoCommon::ImageViewType::e3D:
-        create(TextureType::Color3D, std::nullopt);
+        create(TextureType::Color3D);
         render_target = Handle(Shader::TextureType::Color3D);
         break;
     case VideoCommon::ImageViewType::Cube:
     case VideoCommon::ImageViewType::CubeArray:
-        create(TextureType::ColorCube, 6);
-        create(TextureType::ColorArrayCube, std::nullopt);
+        create(TextureType::ColorCube);
+        create(TextureType::ColorArrayCube);
         break;
     case VideoCommon::ImageViewType::Buffer:
         ASSERT(false);
@@ -2216,7 +2441,8 @@ ImageView::ImageView(TextureCacheRuntime& runtime, const VideoCommon::ImageViewI
 ImageView::ImageView(TextureCacheRuntime& runtime, const VideoCommon::ImageInfo& info,
                      const VideoCommon::ImageViewInfo& view_info, GPUVAddr gpu_addr_)
     : VideoCommon::ImageViewBase{info, view_info, gpu_addr_},
-      buffer_size{VideoCommon::CalculateGuestSizeInBytes(info)} {}
+    buffer_size{VideoCommon::CalculateGuestSizeInBytes(info)},
+    is_3d_image(info.type == ImageType::e3D) {}
 
 ImageView::ImageView(TextureCacheRuntime& runtime, const VideoCommon::NullImageViewParams& params)
     : VideoCommon::ImageViewBase{params}, device{&runtime.device} {
@@ -2228,7 +2454,7 @@ ImageView::ImageView(TextureCacheRuntime& runtime, const VideoCommon::NullImageV
     ImageInfo info{};
     info.format = PixelFormat::A8B8G8R8_UNORM;
 
-    null_image = MakeImage(*device, runtime.memory_allocator, info, {});
+    null_image = MakeImage(*device, runtime.memory_allocator, info, {}, false);
     image_handle = *null_image;
     for (u32 i = 0; i < Shader::NUM_TEXTURE_TYPES; i++) {
         image_views[i] = MakeView(VK_FORMAT_A8B8G8R8_UNORM_PACK32, VK_IMAGE_ASPECT_COLOR_BIT);
@@ -2237,28 +2463,81 @@ ImageView::ImageView(TextureCacheRuntime& runtime, const VideoCommon::NullImageV
 
 ImageView::~ImageView() = default;
 
+Shader::TextureType ImageView::BaseTextureType() const noexcept {
+    using VideoCommon::ImageViewType;
+    switch (type) {
+    case ImageViewType::e1D:
+        return Shader::TextureType::Color1D;
+    case ImageViewType::e1DArray:
+        return Shader::TextureType::ColorArray1D;
+    case ImageViewType::e2D:
+        return Shader::TextureType::Color2D;
+    case ImageViewType::e2DArray:
+        return Shader::TextureType::ColorArray2D;
+    case ImageViewType::Rect:
+        return Shader::TextureType::Color2DRect;
+    case ImageViewType::e3D:
+        return Shader::TextureType::Color3D;
+    case ImageViewType::Cube:
+        return Shader::TextureType::ColorCube;
+    case ImageViewType::CubeArray:
+        return Shader::TextureType::ColorArrayCube;
+    case ImageViewType::Buffer:
+        break;
+    }
+    return Shader::TextureType::Color2D;
+}
+
+std::optional<u32> ImageView::LayerCountOverride(Shader::TextureType texture_type) const noexcept {
+    switch (texture_type) {
+    case Shader::TextureType::Color1D:
+    case Shader::TextureType::Color2D:
+    case Shader::TextureType::Color2DRect:
+        return 1u;
+    case Shader::TextureType::ColorCube:
+        return 6u;
+    default:
+        return std::nullopt;
+    }
+}
+
 VkImageView ImageView::DepthView() {
+    return DepthView(BaseTextureType());
+}
+
+VkImageView ImageView::DepthView(Shader::TextureType texture_type) {
     if (!image_handle) {
         return VK_NULL_HANDLE;
     }
-    if (depth_view) {
-        return *depth_view;
+    auto& view = depth_views[static_cast<size_t>(texture_type)];
+    if (view) {
+        return *view;
     }
     const auto& info = MaxwellToVK::SurfaceFormat(*device, FormatType::Optimal, true, format);
-    depth_view = MakeView(info.format, VK_IMAGE_ASPECT_DEPTH_BIT);
-    return *depth_view;
+    view = MakeView(info.format, VK_IMAGE_ASPECT_DEPTH_BIT, texture_type);
+    return *view;
 }
 
 VkImageView ImageView::StencilView() {
+    return StencilView(BaseTextureType());
+}
+
+VkImageView ImageView::StencilView(Shader::TextureType texture_type) {
     if (!image_handle) {
         return VK_NULL_HANDLE;
     }
-    if (stencil_view) {
-        return *stencil_view;
+    auto& view = stencil_views[static_cast<size_t>(texture_type)];
+    if (view) {
+        return *view;
     }
-    const auto& info = MaxwellToVK::SurfaceFormat(*device, FormatType::Optimal, true, format);
-    stencil_view = MakeView(info.format, VK_IMAGE_ASPECT_STENCIL_BIT);
-    return *stencil_view;
+    const auto surface_type = VideoCore::Surface::GetFormatType(format);
+    PixelFormat view_format = format;
+    if (surface_type == SurfaceType::Stencil) {
+        view_format = PixelFormat::S8_UINT;
+    }
+    const auto& info = MaxwellToVK::SurfaceFormat(*device, FormatType::Optimal, true, view_format);
+    view = MakeView(info.format, VK_IMAGE_ASPECT_STENCIL_BIT, texture_type);
+    return *view;
 }
 
 VkImageView ImageView::ColorView() {
@@ -2272,20 +2551,74 @@ VkImageView ImageView::ColorView() {
     return *color_view;
 }
 
+VkImageView ImageView::SampledView(Shader::TextureType texture_type,
+                                   Shader::SamplerComponentType component_type) {
+    using VideoCore::Surface::GetFormatType;
+    using VideoCore::Surface::SurfaceType;
+
+    const SurfaceType surface_type = GetFormatType(format);
+    switch (component_type) {
+    case Shader::SamplerComponentType::Depth:
+        if (surface_type == SurfaceType::Depth || surface_type == SurfaceType::DepthStencil) {
+            return DepthView(texture_type);
+        }
+        break;
+    case Shader::SamplerComponentType::Stencil:
+        if (surface_type == SurfaceType::Stencil || surface_type == SurfaceType::DepthStencil) {
+            return StencilView(texture_type);
+        }
+        break;
+    default:
+        break;
+    }
+    const auto desired_numeric = ComponentNumericType(component_type);
+    if (!desired_numeric) {
+        return Handle(texture_type);
+    }
+    const PixelFormatNumeric current_numeric =
+        VideoCore::Surface::GetPixelFormatNumericType(format);
+    if (*desired_numeric == current_numeric) {
+        return Handle(texture_type);
+    }
+    const auto remapped_format =
+        VideoCore::Surface::FindPixelFormatVariant(format, *desired_numeric);
+    if (!remapped_format) {
+        return Handle(texture_type);
+    }
+    auto& cached_view = sampled_component_views[static_cast<size_t>(*desired_numeric)]
+                                            [static_cast<size_t>(texture_type)];
+    if (cached_view) {
+        return *cached_view;
+    }
+    const auto& info =
+        MaxwellToVK::SurfaceFormat(*device, FormatType::Optimal, true, *remapped_format);
+    cached_view = MakeView(info.format, VK_IMAGE_ASPECT_COLOR_BIT, texture_type);
+    return *cached_view;
+}
+
 VkImageView ImageView::StorageView(Shader::TextureType texture_type,
                                    Shader::ImageFormat image_format) {
-    if (image_handle) {
-        if (image_format == Shader::ImageFormat::Typeless) {
-            return Handle(texture_type);
+    if (!image_handle) {
+        return VK_NULL_HANDLE;
+    }
+    const bool is_signed{image_format == Shader::ImageFormat::R8_SINT ||
+                         image_format == Shader::ImageFormat::R16_SINT};
+    if (!storage_views) {
+        storage_views.emplace();
+    }
+    if (image_format == Shader::ImageFormat::Typeless) {
+        auto& view = storage_views->typeless[static_cast<size_t>(texture_type)];
+        if (view) {
+            return *view;
         }
-        const bool is_signed = image_format == Shader::ImageFormat::R8_SINT
-            || image_format == Shader::ImageFormat::R16_SINT;
-        if (!storage_views)
-            storage_views.emplace();
-        auto& views{is_signed ? storage_views->signeds : storage_views->unsigneds};
-        auto& view{views[size_t(texture_type)]};
-        if (!view)
-            view = MakeView(Format(image_format), VK_IMAGE_ASPECT_COLOR_BIT);
+        const auto& format_info =
+            MaxwellToVK::SurfaceFormat(*device, FormatType::Optimal, true, format);
+        view = MakeView(format_info.format, VK_IMAGE_ASPECT_COLOR_BIT, texture_type);
+        return *view;
+    }
+    auto& views{is_signed ? storage_views->signeds : storage_views->unsigneds};
+    auto& view{views[static_cast<size_t>(texture_type)]};
+    if (view) {
         return *view;
     }
     return VK_NULL_HANDLE;
@@ -2295,13 +2628,24 @@ bool ImageView::IsRescaled() const noexcept {
     return (*slot_images)[image_id].IsRescaled();
 }
 
+bool ImageView::SupportsDepthCompareSampling() const noexcept {
+    const auto surface_type = VideoCore::Surface::GetFormatType(format);
+    return surface_type == VideoCore::Surface::SurfaceType::Depth ||
+           surface_type == VideoCore::Surface::SurfaceType::DepthStencil;
+}
+
 vk::ImageView ImageView::MakeView(VkFormat vk_format, VkImageAspectFlags aspect_mask) {
-    return device->GetLogical().CreateImageView({
+    return MakeView(vk_format, aspect_mask, BaseTextureType());
+}
+
+vk::ImageView ImageView::MakeView(VkFormat vk_format, VkImageAspectFlags aspect_mask,
+                                  Shader::TextureType texture_type) {
+    VkImageViewCreateInfo ci{
         .sType = VK_STRUCTURE_TYPE_IMAGE_VIEW_CREATE_INFO,
         .pNext = nullptr,
         .flags = 0,
         .image = image_handle,
-        .viewType = ImageViewType(type),
+        .viewType = ImageViewType(texture_type),
         .format = vk_format,
         .components{
             .r = VK_COMPONENT_SWIZZLE_IDENTITY,
@@ -2310,7 +2654,11 @@ vk::ImageView ImageView::MakeView(VkFormat vk_format, VkImageAspectFlags aspect_
             .a = VK_COMPONENT_SWIZZLE_IDENTITY,
         },
         .subresourceRange = MakeSubresourceRange(aspect_mask, range),
-    });
+    };
+    if (const auto override_layers = LayerCountOverride(texture_type)) {
+        ci.subresourceRange.layerCount = *override_layers;
+    }
+    return device->GetLogical().CreateImageView(ci);
 }
 
 Sampler::Sampler(TextureCacheRuntime& runtime, const Tegra::Texture::TSCEntry& tsc) {
@@ -2350,21 +2698,33 @@ Sampler::Sampler(TextureCacheRuntime& runtime, const Tegra::Texture::TSCEntry& t
     // Some games have samplers with garbage. Sanitize them here.
     const f32 max_anisotropy = std::clamp(tsc.MaxAnisotropy(), 1.0f, 16.0f);
 
-    const auto create_sampler = [&](const f32 anisotropy) {
+    depth_compare_enabled = tsc.depth_compare_enabled != 0;
+
+    const auto create_sampler = [&](const f32 anisotropy, bool force_point_filter,
+                                    bool enable_compare) {
+        const VkFilter mag_filter = force_point_filter ? VK_FILTER_NEAREST
+                                                       : MaxwellToVK::Sampler::Filter(tsc.mag_filter);
+        const VkFilter min_filter = force_point_filter ? VK_FILTER_NEAREST
+                                                       : MaxwellToVK::Sampler::Filter(tsc.min_filter);
+        const VkSamplerMipmapMode mipmap_mode = force_point_filter
+                                                    ? VK_SAMPLER_MIPMAP_MODE_NEAREST
+                                                    : MaxwellToVK::Sampler::MipmapMode(tsc.mipmap_filter);
+        const float resolved_anisotropy = force_point_filter ? 1.0f : anisotropy;
+        const bool enable_anisotropy = resolved_anisotropy > 1.0f && !force_point_filter;
         return device.GetLogical().CreateSampler(VkSamplerCreateInfo{
             .sType = VK_STRUCTURE_TYPE_SAMPLER_CREATE_INFO,
             .pNext = pnext,
             .flags = 0,
-            .magFilter = MaxwellToVK::Sampler::Filter(tsc.mag_filter),
-            .minFilter = MaxwellToVK::Sampler::Filter(tsc.min_filter),
-            .mipmapMode = MaxwellToVK::Sampler::MipmapMode(tsc.mipmap_filter),
+            .magFilter = mag_filter,
+            .minFilter = min_filter,
+            .mipmapMode = mipmap_mode,
             .addressModeU = MaxwellToVK::Sampler::WrapMode(device, tsc.wrap_u, tsc.mag_filter),
             .addressModeV = MaxwellToVK::Sampler::WrapMode(device, tsc.wrap_v, tsc.mag_filter),
             .addressModeW = MaxwellToVK::Sampler::WrapMode(device, tsc.wrap_p, tsc.mag_filter),
             .mipLodBias = tsc.LodBias(),
-            .anisotropyEnable = static_cast<VkBool32>(anisotropy > 1.0f ? VK_TRUE : VK_FALSE),
-            .maxAnisotropy = anisotropy,
-            .compareEnable = tsc.depth_compare_enabled,
+            .anisotropyEnable = static_cast<VkBool32>(enable_anisotropy ? VK_TRUE : VK_FALSE),
+            .maxAnisotropy = enable_anisotropy ? resolved_anisotropy : 1.0f,
+            .compareEnable = enable_compare ? VK_TRUE : VK_FALSE,
             .compareOp = MaxwellToVK::Sampler::DepthCompareFunction(tsc.depth_compare_func),
             .minLod = tsc.mipmap_filter == TextureMipmapFilter::None ? 0.0f : tsc.MinLod(),
             .maxLod = tsc.mipmap_filter == TextureMipmapFilter::None ? 0.25f : tsc.MaxLod(),
@@ -2374,12 +2734,54 @@ Sampler::Sampler(TextureCacheRuntime& runtime, const Tegra::Texture::TSCEntry& t
         });
     };
 
-    sampler = create_sampler(max_anisotropy);
+    sampler = create_sampler(max_anisotropy, false, depth_compare_enabled);
+    if (depth_compare_enabled) {
+        sampler_compare_disabled = create_sampler(max_anisotropy, false, false);
+    }
 
     const f32 max_anisotropy_default = static_cast<f32>(1U << tsc.max_anisotropy);
     if (max_anisotropy > max_anisotropy_default) {
-        sampler_default_anisotropy = create_sampler(max_anisotropy_default);
+        sampler_default_anisotropy =
+            create_sampler(max_anisotropy_default, false, depth_compare_enabled);
+        if (depth_compare_enabled) {
+            sampler_default_anisotropy_compare_disabled =
+                create_sampler(max_anisotropy_default, false, false);
+        }
     }
+
+    uses_linear_filter = (tsc.mag_filter == TextureFilter::Linear) ||
+                         (tsc.min_filter == TextureFilter::Linear) || (max_anisotropy > 1.0f);
+    if (uses_linear_filter) {
+        sampler_force_point = create_sampler(1.0f, true, depth_compare_enabled);
+        if (depth_compare_enabled) {
+            sampler_force_point_compare_disabled = create_sampler(1.0f, true, false);
+        }
+    }
+}
+
+VkSampler Sampler::SelectHandle(bool supports_linear_filter, bool supports_anisotropy,
+                                bool allow_depth_compare) const noexcept {
+    const auto select_compare_variant = [&](const vk::Sampler& compare_sampler,
+                                            const vk::Sampler& non_compare_sampler) {
+        if (!depth_compare_enabled || allow_depth_compare) {
+            return *compare_sampler;
+        }
+        if (non_compare_sampler) {
+            return *non_compare_sampler;
+        }
+        return *compare_sampler;
+    };
+
+    const bool needs_linear_fallback = uses_linear_filter && !supports_linear_filter;
+    if (needs_linear_fallback && sampler_force_point) {
+        return select_compare_variant(sampler_force_point, sampler_force_point_compare_disabled);
+    }
+    const bool needs_aniso_fallback = HasAddedAnisotropy() && !supports_anisotropy;
+    if (needs_aniso_fallback) {
+        return select_compare_variant(sampler_default_anisotropy,
+                                      sampler_default_anisotropy_compare_disabled);
+    }
+    return select_compare_variant(sampler, sampler_compare_disabled);
 }
 
 Framebuffer::Framebuffer(TextureCacheRuntime& runtime, std::span<ImageView*, NUM_RT> color_buffers,

src/video_core/renderer_vulkan/vk_texture_cache.h

@@ -6,6 +6,7 @@
 
 #pragma once
 
+#include <optional>
 #include <span>
 
 #include "video_core/texture_cache/texture_cache_base.h"
@@ -59,6 +60,8 @@ public:
 
     void TickFrame();
 
+    void WaitForGpuTick(u64 tick);
+
     u64 GetDeviceLocalMemory() const;
 
     u64 GetDeviceMemoryUsage() const;
@@ -114,10 +117,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;
 
@@ -136,6 +144,7 @@ public:
     std::optional<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{};
@@ -267,6 +276,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);
@@ -338,11 +371,16 @@ public:
 
     [[nodiscard]] VkImageView ColorView();
 
-    [[nodiscard]] VkImageView StorageView(Shader::TextureType texture_type,
-                                          Shader::ImageFormat image_format);
+    [[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)];
     }
@@ -371,15 +409,26 @@ 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;
     };
 
-    [[nodiscard]] vk::ImageView MakeView(VkFormat vk_format, VkImageAspectFlags aspect_mask);
+    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::optional<StorageViews> storage_views;
+    std::array<vk::ImageView, Shader::NUM_TEXTURE_TYPES> depth_views;
+    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 depth_view;
     vk::ImageView stencil_view;
     vk::ImageView color_view;
@@ -388,6 +437,7 @@ private:
     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 {};
@@ -405,12 +455,20 @@ public:
     }
 
     [[nodiscard]] bool HasAddedAnisotropy() const noexcept {
-        return static_cast<bool>(sampler_default_anisotropy);
+        return 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;
 };
 
 struct TextureCacheParams {

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;
     ankerl::unordered_dense::map<u32, Shader::TextureType> texture_types;
+    ankerl::unordered_dense::map<u32, Shader::SamplerComponentType> texture_component_types;
     ankerl::unordered_dense::map<u32, Shader::TexturePixelFormat> texture_pixel_formats;
     ankerl::unordered_dense::map<u64, u32> cbuf_values;
     ankerl::unordered_dense::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;
     ankerl::unordered_dense::map<u32, Shader::TextureType> texture_types;
+    ankerl::unordered_dense::map<u32, Shader::SamplerComponentType> texture_component_types;
     ankerl::unordered_dense::map<u32, Shader::TexturePixelFormat> texture_pixel_formats;
     ankerl::unordered_dense::map<u64, u32> cbuf_values;
     ankerl::unordered_dense::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 2026 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

@@ -24,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;
@@ -44,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

@@ -680,6 +680,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);
@@ -690,11 +691,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()) {
@@ -746,6 +748,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);
         }
@@ -761,6 +764,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)) {
@@ -2039,6 +2043,7 @@ 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);
         EnforceSamplerBudget();
     }
@@ -2726,6 +2731,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)) {
@@ -2743,6 +2751,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) {
@@ -2760,6 +2847,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));
@@ -2776,20 +2865,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);
@@ -2842,6 +2941,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

@@ -195,6 +195,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();
 
@@ -271,6 +276,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:
@@ -421,6 +435,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);
 
@@ -494,6 +510,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;
@@ -509,6 +530,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;

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_device.cpp

@@ -7,6 +7,7 @@
 #include <algorithm>
 #include <bitset>
 #include <chrono>
+#include <mutex>
 #include <optional>
 #include <thread>
 #include <ankerl/unordered_dense.h>
@@ -25,6 +26,7 @@
 #include "video_core/vulkan_common/vulkan_device.h"
 #include "video_core/vulkan_common/vulkan_wrapper.h"
 #include "video_core/gpu_logging/gpu_logging.h"
+#include "vulkan/vulkan_core.h"
 
 #if defined(ANDROID) && defined(ARCHITECTURE_arm64)
 #include <adrenotools/bcenabler.h>
@@ -95,8 +97,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;
@@ -127,6 +178,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;
     }
@@ -382,10 +439,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) {
-    if (loaded_extensions.contains(extension_name) && !is_suitable) {
+void Device::RemoveExtensionIfUnsuitable(bool& extension, const std::string& extension_name) {
+    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);
     }
 }
 
@@ -406,11 +463,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) {
-        LOG_WARNING(Render_Vulkan, "Removing features for unsuitable extension {}", extension_name);
-        this->RemoveExtensionFeature(is_suitable, feature, extension_name);
+    if (loaded_extensions.contains(extension_name) && !extension) {
+        LOG_WARNING(Render_Vulkan,
+                    "Removing features for unsuitable extension {}", extension_name);
+        this->RemoveExtensionFeature(extension, feature, extension_name);
     }
 }
 
@@ -491,9 +549,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,
-                    "Disabling shader float controls and 64-bit integer features on Qualcomm proprietary drivers");
-        RemoveExtension(extensions.shader_float_controls, VK_KHR_SHADER_FLOAT_CONTROLS_EXTENSION_NAME);
+        if (extensions.shader_float_controls) {
+            LOG_WARNING(Render_Vulkan,
+                        "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;
@@ -561,7 +626,7 @@ Device::Device(VkInstance instance_, vk::PhysicalDevice physical_, VkSurfaceKHR
                      is_blit_depth24_stencil8_supported);
             LOG_INFO(Render_Vulkan, "  D32S8 hardware blit support: {}",
                      is_blit_depth32_stencil8_supported);
-
+            
             if (!is_blit_depth24_stencil8_supported && !is_blit_depth32_stencil8_supported) {
                 LOG_WARNING(Render_Vulkan,
                             "NVIDIA: Neither shader export nor hardware blits available for "
@@ -590,7 +655,44 @@ Device::Device(VkInstance instance_, vk::PhysicalDevice physical_, VkSurfaceKHR
         }
     }
 
+    const bool needs_mobile_alignment_clamp = is_qualcomm || is_arm;
+
     if (is_qualcomm) {
+        has_broken_parallel_compiling = true;
+        LOG_WARNING(Render_Vulkan,
+                    "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;
@@ -599,9 +701,9 @@ Device::Device(VkInstance instance_, vk::PhysicalDevice physical_, VkSurfaceKHR
                 (std::max)(MIN_SAMPLER_BUDGET, sampler_limit - reserved);
             sampler_heap_budget = derived_budget;
             LOG_WARNING(Render_Vulkan,
-                        "Qualcomm driver reports max {} samplers; reserving {} (25%) and "
+                        "{} driver reports max {} samplers; reserving {} (25%) and "
                         "allowing Eden to use {} (75%) to avoid heap exhaustion",
-                        sampler_limit, reserved, sampler_heap_budget);
+                        driver_label, sampler_limit, reserved, sampler_heap_budget);
         }
     }
 
@@ -622,7 +724,7 @@ Device::Device(VkInstance instance_, vk::PhysicalDevice physical_, VkSurfaceKHR
     }
 
     if (is_intel_windows) {
-        LOG_WARNING(Render_Vulkan,
+        LOG_WARNING(Render_Vulkan, 
                     "Intel proprietary drivers do not support MSAA->MSAA image blits. "
                     "MSAA scaling will use 3D helpers. MSAA resolves work normally.");
         cant_blit_msaa = true;
@@ -646,7 +748,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:
@@ -654,36 +756,36 @@ Device::Device(VkInstance instance_, vk::PhysicalDevice physical_, VkSurfaceKHR
     //   Level 1 = Core + VK_EXT_extended_dynamic_state
     //   Level 2 = Core + VK_EXT_extended_dynamic_state + VK_EXT_extended_dynamic_state2
     //   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,
+        RemoveExtensionFeature(extensions.extended_dynamic_state, features.extended_dynamic_state, 
                               VK_EXT_EXTENDED_DYNAMIC_STATE_EXTENSION_NAME);
-        RemoveExtensionFeature(extensions.extended_dynamic_state2, features.extended_dynamic_state2,
+        RemoveExtensionFeature(extensions.extended_dynamic_state2, features.extended_dynamic_state2, 
                               VK_EXT_EXTENDED_DYNAMIC_STATE_2_EXTENSION_NAME);
-        RemoveExtensionFeature(extensions.extended_dynamic_state3, features.extended_dynamic_state3,
+        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::EDS1:
+    case 1:
         // Level 1: Enable EDS1, disable EDS2 and EDS3
-        RemoveExtensionFeature(extensions.extended_dynamic_state2, features.extended_dynamic_state2,
+        RemoveExtensionFeature(extensions.extended_dynamic_state2, features.extended_dynamic_state2, 
                               VK_EXT_EXTENDED_DYNAMIC_STATE_2_EXTENSION_NAME);
-        RemoveExtensionFeature(extensions.extended_dynamic_state3, features.extended_dynamic_state3,
+        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::EDS2:
+    case 2:
         // Level 2: Enable EDS1 + EDS2, disable EDS3
-        RemoveExtensionFeature(extensions.extended_dynamic_state3, features.extended_dynamic_state3,
+        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;
@@ -740,32 +842,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,
-                  "Format={} with usage={} and type={} has no defined alternatives and host "
-                  "hardware does not support it",
-                  wanted_format, wanted_usage, format_type);
+        if (ShouldLogUnsupportedFormat(wanted_format, wanted_usage, format_type)) {
+            LOG_ERROR(Render_Vulkan,
+                      "Format={} with usage={} and type={} has no defined alternatives and host "
+                      "hardware does not support it",
+                      wanted_format, wanted_usage, format_type);
+        }
         return wanted_format;
     }
 
-    std::size_t i = 0;
-    for (VkFormat alternative = *alternatives; alternative; alternative = alternatives[++i]) {
+#if defined(ANDROID)
+    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);
-        return alternative;
+                  wanted_format, selected, wanted_usage, format_type);
+        return selected;
     }
 
-    // No alternatives found, panic
-    LOG_ERROR(Render_Vulkan,
-              "Format={} with usage={} and type={} is not supported by the host hardware and "
-              "doesn't support any of the alternatives",
-              wanted_format, wanted_usage, format_type);
+    if (ShouldLogUnsupportedFormat(wanted_format, wanted_usage, format_type)) {
+        LOG_ERROR(Render_Vulkan,
+                  "Format={} with usage={} and type={} is not supported by the host hardware and "
+                  "doesn't support any of the alternatives",
+                  wanted_format, wanted_usage, format_type);
+    }
     return wanted_format;
 }
 
@@ -863,7 +996,8 @@ bool Device::ShouldBoostClocks() const {
 }
 
 bool Device::HasTimelineSemaphore() const {
-    if (GetDriverID() == VK_DRIVER_ID_MESA_TURNIP) {
+    const VkDriverIdKHR driver_id = GetDriverID();
+    if (driver_id == VK_DRIVER_ID_MESA_TURNIP || driver_id == VK_DRIVER_ID_QUALCOMM_PROPRIETARY) {
         return false;
     }
     return features.timeline_semaphore.timelineSemaphore;
@@ -968,11 +1102,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 = &features_1_3;
-
+        features_1_2.pNext = nullptr;
         *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
@@ -1114,24 +1250,34 @@ bool Device::GetSuitability(bool requires_swapchain) {
     }
 
     // VK_DYNAMIC_STATE
-
+    
     // Driver detection variables for workarounds in GetSuitability
     const VkDriverId driver_id = properties.driver.driverID;
 
-    // VK_EXT_extended_dynamic_state2 below this will appear drivers that need workarounds.
+    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_extended_dynamic_state3 below this will appear drivers that need workarounds.
+    // VK_EXT_CUSTOM_BORDER_COLOR
+    if (extensions.custom_border_color && (driver_id == VK_DRIVER_ID_QUALCOMM_PROPRIETARY || 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.
 
     // Samsung: Broken extendedDynamicState3ColorBlendEquation
     // Disable blend equation dynamic state, force static pipeline state
-    if (extensions.extended_dynamic_state3 &&
+    if (extensions.extended_dynamic_state3 && 
         (driver_id == VK_DRIVER_ID_SAMSUNG_PROPRIETARY)) {
         LOG_WARNING(Render_Vulkan,
                     "Samsung: Disabling broken extendedDynamicState3ColorBlendEquation");
         features.extended_dynamic_state3.extendedDynamicState3ColorBlendEnable = false;
         features.extended_dynamic_state3.extendedDynamicState3ColorBlendEquation = false;
     }
-
+    
     // Intel Windows < 27.20.100.0: Broken VertexInputDynamicState
     // Same for NVIDIA Proprietary < 580.119.02, unknown when VIDS was first NOT broken
     // Disable VertexInputDynamicState on old Intel Windows drivers
@@ -1143,7 +1289,7 @@ bool Device::GetSuitability(bool requires_swapchain) {
             RemoveExtensionFeature(extensions.vertex_input_dynamic_state, features.vertex_input_dynamic_state, VK_EXT_VERTEX_INPUT_DYNAMIC_STATE_EXTENSION_NAME);
         }
     }
-
+    
     if (u32(Settings::values.dyna_state.GetValue()) == 0) {
         LOG_INFO(Render_Vulkan, "Extended Dynamic State disabled by user setting, clearing all EDS features");
         features.extended_dynamic_state.extendedDynamicState = false;
@@ -1350,6 +1496,15 @@ 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={}",
+                 properties.maintenance5.polygonModePointSize,
+                 properties.maintenance5.depthStencilSwizzleOneSupport,
+                 properties.maintenance5.earlyFragmentMultisampleCoverageAfterSampleCounting &&
+                 properties.maintenance5.earlyFragmentSampleMaskTestBeforeSampleCounting,
+                 properties.maintenance5.nonStrictWideLinesUseParallelogram);
+    }
     RemoveExtensionFeatureIfUnsuitable(extensions.maintenance5, features.maintenance5,
                                        VK_KHR_MAINTENANCE_5_EXTENSION_NAME);
 
@@ -1358,13 +1513,41 @@ void Device::RemoveUnsuitableExtensions() {
     RemoveExtensionFeatureIfUnsuitable(extensions.maintenance6, features.maintenance6,
                                        VK_KHR_MAINTENANCE_6_EXTENSION_NAME);
 
-    // VK_KHR_maintenance7
+    // VK_KHR_maintenance7 (proposed for Vulkan 1.4, no features)
     extensions.maintenance7 = loaded_extensions.contains(VK_KHR_MAINTENANCE_7_EXTENSION_NAME);
     RemoveExtensionIfUnsuitable(extensions.maintenance7, VK_KHR_MAINTENANCE_7_EXTENSION_NAME);
 
-    // VK_KHR_maintenance8
+    // VK_KHR_maintenance8 (proposed for Vulkan 1.4, no features)
     extensions.maintenance8 = loaded_extensions.contains(VK_KHR_MAINTENANCE_8_EXTENSION_NAME);
     RemoveExtensionIfUnsuitable(extensions.maintenance8, VK_KHR_MAINTENANCE_8_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) {

src/video_core/vulkan_common/vulkan_device.h

@@ -6,6 +6,7 @@
 
 #pragma once
 
+#include <algorithm>
 #include <optional>
 #include <set>
 #include <span>
@@ -302,12 +303,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.
@@ -340,6 +343,7 @@ public:
         return properties.float_controls;
     }
 
+
     /// Returns true if ASTC is natively supported.
     bool IsOptimalAstcSupported() const {
         return features.features.textureCompressionASTC_LDR;
@@ -355,6 +359,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;
@@ -385,11 +394,26 @@ public:
         return properties.subgroup_size_control.requiredSubgroupSizeStages & stage;
     }
 
-    /// Returns true if the device supports the provided subgroup feature.
-    bool IsSubgroupFeatureSupported(VkSubgroupFeatureFlagBits feature) const {
-        return properties.subgroup_properties.supportedOperations & feature;
+    /// Returns true if the device supports the provided subgroup feature for the given stages.
+    bool IsSubgroupFeatureSupported(VkSubgroupFeatureFlagBits 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;
@@ -462,7 +486,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.
@@ -634,6 +658,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;
     }
@@ -662,9 +696,15 @@ public:
         return dynamic_state3_alpha_to_one;
     }
 
-    /// Returns true if the device supports VK_EXT_vertex_input_dynamic_state.
+    /// Returns true when the user enabled extended core dynamic states (level > 0).
+    bool UsesAdvancedCoreDynamicState() const {
+        return u32(Settings::values.dyna_state.GetValue()) > 0;
+    }
+
+    /// 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
@@ -711,10 +751,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;
@@ -849,7 +890,7 @@ public:
 
     /// Returns true if early fragment tests optimizations are available.
     bool SupportsEarlyFragmentTests() const {
-        return extensions.maintenance5 &&
+        return extensions.maintenance5 && 
                properties.maintenance5.earlyFragmentMultisampleCoverageAfterSampleCounting &&
                properties.maintenance5.earlyFragmentSampleMaskTestBeforeSampleCounting;
     }
@@ -869,12 +910,6 @@ public:
         return extensions.maintenance8;
     }
 
-    /// Returns true if the device supports UINT8 index buffer conversion via compute shader.
-    bool SupportsUint8Indices() const {
-        return features.bit8_storage.storageBuffer8BitAccess &&
-               features.bit16_storage.storageBuffer16BitAccess;
-    }
-
     [[nodiscard]] static constexpr bool CheckBrokenCompute(VkDriverId driver_id,
                                                            u32 driver_version) {
         if (driver_id == VK_DRIVER_ID_INTEL_PROPRIETARY_WINDOWS) {
@@ -906,6 +941,10 @@ public:
     /// GPU logging integration
     void InitializeGPULogging();
     void ShutdownGPULogging();
+    
+    static constexpr size_t SampleLocationTableSize() {
+        return sample_location_table_size;
+    }
 
 private:
     /// Checks if the physical device is suitable and configures the object state
@@ -916,13 +955,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.
@@ -1024,6 +1063,8 @@ private:
     bool supports_d24_depth{};                 ///< Supports D24 depth buffers.
     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{};
@@ -1035,6 +1076,8 @@ private:
     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};