9 files changed, 243 insertions, 144 deletions

diff --git a/src/shader_recompiler/backend/spirv/emit_context.cpp b/src/shader_recompiler/backend/spirv/emit_context.cpp
index 6c8f16562..4a4de3676 100644
--- a/src/shader_recompiler/backend/spirv/emit_context.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_context.cpp
@@ -48,6 +48,25 @@ Id ImageType(EmitContext& ctx, const TextureDescriptor& desc) {
     }
     throw InvalidArgument("Invalid texture type {}", desc.type);
 }
+
+Id DefineVariable(EmitContext& ctx, Id type, std::optional<spv::BuiltIn> builtin,
+                  spv::StorageClass storage_class) {
+    const Id pointer_type{ctx.TypePointer(storage_class, type)};
+    const Id id{ctx.AddGlobalVariable(pointer_type, storage_class)};
+    if (builtin) {
+        ctx.Decorate(id, spv::Decoration::BuiltIn, *builtin);
+    }
+    ctx.interfaces.push_back(id);
+    return id;
+}
+
+Id DefineInput(EmitContext& ctx, Id type, std::optional<spv::BuiltIn> builtin = std::nullopt) {
+    return DefineVariable(ctx, type, builtin, spv::StorageClass::Input);
+}
+
+Id DefineOutput(EmitContext& ctx, Id type, std::optional<spv::BuiltIn> builtin = std::nullopt) {
+    return DefineVariable(ctx, type, builtin, spv::StorageClass::Output);
+}
 } // Anonymous namespace
 
 void VectorTypes::Define(Sirit::Module& sirit_ctx, Id base_type, std::string_view name) {
@@ -144,59 +163,8 @@ void EmitContext::DefineCommonConstants() {
 }
 
 void EmitContext::DefineInterfaces(const Info& info, Stage stage) {
-    const auto define{
-        [this](Id type, std::optional<spv::BuiltIn> builtin, spv::StorageClass storage_class) {
-            const Id pointer_type{TypePointer(storage_class, type)};
-            const Id id{AddGlobalVariable(pointer_type, storage_class)};
-            if (builtin) {
-                Decorate(id, spv::Decoration::BuiltIn, *builtin);
-            }
-            interfaces.push_back(id);
-            return id;
-        }};
-    using namespace std::placeholders;
-    const auto define_input{std::bind(define, _1, _2, spv::StorageClass::Input)};
-    const auto define_output{std::bind(define, _1, _2, spv::StorageClass::Output)};
-
-    if (info.uses_workgroup_id) {
-        workgroup_id = define_input(U32[3], spv::BuiltIn::WorkgroupId);
-    }
-    if (info.uses_local_invocation_id) {
-        local_invocation_id = define_input(U32[3], spv::BuiltIn::LocalInvocationId);
-    }
-    if (info.loads_position) {
-        const bool is_fragment{stage != Stage::Fragment};
-        const spv::BuiltIn built_in{is_fragment ? spv::BuiltIn::Position : spv::BuiltIn::FragCoord};
-        input_position = define_input(F32[4], built_in);
-    }
-    for (size_t i = 0; i < info.loads_generics.size(); ++i) {
-        if (info.loads_generics[i]) {
-            // FIXME: Declare size from input
-            input_generics[i] = define_input(F32[4], std::nullopt);
-            Decorate(input_generics[i], spv::Decoration::Location, static_cast<u32>(i));
-            Name(input_generics[i], fmt::format("in_attr{}", i));
-        }
-    }
-    if (info.stores_position) {
-        output_position = define_output(F32[4], spv::BuiltIn::Position);
-    }
-    for (size_t i = 0; i < info.stores_generics.size(); ++i) {
-        if (info.stores_generics[i]) {
-            output_generics[i] = define_output(F32[4], std::nullopt);
-            Decorate(output_generics[i], spv::Decoration::Location, static_cast<u32>(i));
-            Name(output_generics[i], fmt::format("out_attr{}", i));
-        }
-    }
-    if (stage == Stage::Fragment) {
-        for (size_t i = 0; i < 8; ++i) {
-            if (!info.stores_frag_color[i]) {
-                continue;
-            }
-            frag_color[i] = define_output(F32[4], std::nullopt);
-            Decorate(frag_color[i], spv::Decoration::Location, static_cast<u32>(i));
-            Name(frag_color[i], fmt::format("frag_color{}", i));
-        }
-    }
+    DefineInputs(info, stage);
+    DefineOutputs(info, stage);
 }
 
 void EmitContext::DefineConstantBuffers(const Info& info, u32& binding) {
@@ -225,33 +193,6 @@ void EmitContext::DefineConstantBuffers(const Info& info, u32& binding) {
     }
 }
 
-void EmitContext::DefineConstantBuffers(const Info& info, Id UniformDefinitions::*member_type,
-                                        u32 binding, Id type, char type_char, u32 element_size) {
-    const Id array_type{TypeArray(type, Constant(U32[1], 65536U / element_size))};
-    Decorate(array_type, spv::Decoration::ArrayStride, element_size);
-
-    const Id struct_type{TypeStruct(array_type)};
-    Name(struct_type, fmt::format("cbuf_block_{}{}", type_char, element_size * CHAR_BIT));
-    Decorate(struct_type, spv::Decoration::Block);
-    MemberName(struct_type, 0, "data");
-    MemberDecorate(struct_type, 0, spv::Decoration::Offset, 0U);
-
-    const Id struct_pointer_type{TypePointer(spv::StorageClass::Uniform, struct_type)};
-    const Id uniform_type{TypePointer(spv::StorageClass::Uniform, type)};
-    uniform_types.*member_type = uniform_type;
-
-    for (const ConstantBufferDescriptor& desc : info.constant_buffer_descriptors) {
-        const Id id{AddGlobalVariable(struct_pointer_type, spv::StorageClass::Uniform)};
-        Decorate(id, spv::Decoration::Binding, binding);
-        Decorate(id, spv::Decoration::DescriptorSet, 0U);
-        Name(id, fmt::format("c{}", desc.index));
-        for (size_t i = 0; i < desc.count; ++i) {
-            cbufs[desc.index + i].*member_type = id;
-        }
-        binding += desc.count;
-    }
-}
-
 void EmitContext::DefineStorageBuffers(const Info& info, u32& binding) {
     if (info.storage_buffers_descriptors.empty()) {
         return;
@@ -311,4 +252,94 @@ void EmitContext::DefineLabels(IR::Program& program) {
     }
 }
 
+void EmitContext::DefineInputs(const Info& info, Stage stage) {
+    if (info.uses_workgroup_id) {
+        workgroup_id = DefineInput(*this, U32[3], spv::BuiltIn::WorkgroupId);
+    }
+    if (info.uses_local_invocation_id) {
+        local_invocation_id = DefineInput(*this, U32[3], spv::BuiltIn::LocalInvocationId);
+    }
+    if (info.loads_position) {
+        const bool is_fragment{stage != Stage::Fragment};
+        const spv::BuiltIn built_in{is_fragment ? spv::BuiltIn::Position : spv::BuiltIn::FragCoord};
+        input_position = DefineInput(*this, F32[4], built_in);
+    }
+    if (info.loads_instance_id) {
+        if (profile.support_vertex_instance_id) {
+            instance_id = DefineInput(*this, U32[1], spv::BuiltIn::InstanceId);
+        } else {
+            instance_index = DefineInput(*this, U32[1], spv::BuiltIn::InstanceIndex);
+            base_instance = DefineInput(*this, U32[1], spv::BuiltIn::BaseInstance);
+        }
+    }
+    if (info.loads_vertex_id) {
+        if (profile.support_vertex_instance_id) {
+            vertex_id = DefineInput(*this, U32[1], spv::BuiltIn::VertexId);
+        } else {
+            vertex_index = DefineInput(*this, U32[1], spv::BuiltIn::VertexIndex);
+            base_vertex = DefineInput(*this, U32[1], spv::BuiltIn::BaseVertex);
+        }
+    }
+    for (size_t index = 0; index < info.loads_generics.size(); ++index) {
+        if (!info.loads_generics[index]) {
+            continue;
+        }
+        // FIXME: Declare size from input
+        const Id id{DefineInput(*this, F32[4])};
+        Decorate(id, spv::Decoration::Location, static_cast<u32>(index));
+        Name(id, fmt::format("in_attr{}", index));
+        input_generics[index] = id;
+    }
+}
+
+void EmitContext::DefineConstantBuffers(const Info& info, Id UniformDefinitions::*member_type,
+                                        u32 binding, Id type, char type_char, u32 element_size) {
+    const Id array_type{TypeArray(type, Constant(U32[1], 65536U / element_size))};
+    Decorate(array_type, spv::Decoration::ArrayStride, element_size);
+
+    const Id struct_type{TypeStruct(array_type)};
+    Name(struct_type, fmt::format("cbuf_block_{}{}", type_char, element_size * CHAR_BIT));
+    Decorate(struct_type, spv::Decoration::Block);
+    MemberName(struct_type, 0, "data");
+    MemberDecorate(struct_type, 0, spv::Decoration::Offset, 0U);
+
+    const Id struct_pointer_type{TypePointer(spv::StorageClass::Uniform, struct_type)};
+    const Id uniform_type{TypePointer(spv::StorageClass::Uniform, type)};
+    uniform_types.*member_type = uniform_type;
+
+    for (const ConstantBufferDescriptor& desc : info.constant_buffer_descriptors) {
+        const Id id{AddGlobalVariable(struct_pointer_type, spv::StorageClass::Uniform)};
+        Decorate(id, spv::Decoration::Binding, binding);
+        Decorate(id, spv::Decoration::DescriptorSet, 0U);
+        Name(id, fmt::format("c{}", desc.index));
+        for (size_t i = 0; i < desc.count; ++i) {
+            cbufs[desc.index + i].*member_type = id;
+        }
+        binding += desc.count;
+    }
+}
+
+void EmitContext::DefineOutputs(const Info& info, Stage stage) {
+    if (info.stores_position) {
+        output_position = DefineOutput(*this, F32[4], spv::BuiltIn::Position);
+    }
+    for (size_t i = 0; i < info.stores_generics.size(); ++i) {
+        if (info.stores_generics[i]) {
+            output_generics[i] = DefineOutput(*this, F32[4]);
+            Decorate(output_generics[i], spv::Decoration::Location, static_cast<u32>(i));
+            Name(output_generics[i], fmt::format("out_attr{}", i));
+        }
+    }
+    if (stage == Stage::Fragment) {
+        for (size_t i = 0; i < 8; ++i) {
+            if (!info.stores_frag_color[i]) {
+                continue;
+            }
+            frag_color[i] = DefineOutput(*this, F32[4]);
+            Decorate(frag_color[i], spv::Decoration::Location, static_cast<u32>(i));
+            Name(frag_color[i], fmt::format("frag_color{}", i));
+        }
+    }
+}
+
 } // namespace Shader::Backend::SPIRV
diff --git a/src/shader_recompiler/backend/spirv/emit_context.h b/src/shader_recompiler/backend/spirv/emit_context.h
index 2d7961ac3..9b9e0d6b1 100644
--- a/src/shader_recompiler/backend/spirv/emit_context.h
+++ b/src/shader_recompiler/backend/spirv/emit_context.h
@@ -82,6 +82,12 @@ public:
 
     Id workgroup_id{};
     Id local_invocation_id{};
+    Id instance_id{};
+    Id instance_index{};
+    Id base_instance{};
+    Id vertex_id{};
+    Id vertex_index{};
+    Id base_vertex{};
 
     Id input_position{};
     std::array<Id, 32> input_generics{};
@@ -99,11 +105,15 @@ private:
     void DefineCommonConstants();
     void DefineInterfaces(const Info& info, Stage stage);
     void DefineConstantBuffers(const Info& info, u32& binding);
-    void DefineConstantBuffers(const Info& info, Id UniformDefinitions::*member_type, u32 binding,
-                               Id type, char type_char, u32 element_size);
     void DefineStorageBuffers(const Info& info, u32& binding);
     void DefineTextures(const Info& info, u32& binding);
     void DefineLabels(IR::Program& program);
+
+    void DefineConstantBuffers(const Info& info, Id UniformDefinitions::*member_type, u32 binding,
+                               Id type, char type_char, u32 element_size);
+
+    void DefineInputs(const Info& info, Stage stage);
+    void DefineOutputs(const Info& info, Stage stage);
 };
 
 } // namespace Shader::Backend::SPIRV
diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp
index b8978b94a..efd0b70b7 100644
--- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp
@@ -113,6 +113,43 @@ Id TypeId(const EmitContext& ctx, IR::Type type) {
     }
 }
 
+Id DefineMain(EmitContext& ctx, IR::Program& program) {
+    const Id void_function{ctx.TypeFunction(ctx.void_id)};
+    const Id main{ctx.OpFunction(ctx.void_id, spv::FunctionControlMask::MaskNone, void_function)};
+    for (IR::Block* const block : program.blocks) {
+        ctx.AddLabel(block->Definition<Id>());
+        for (IR::Inst& inst : block->Instructions()) {
+            EmitInst(ctx, &inst);
+        }
+    }
+    ctx.OpFunctionEnd();
+    return main;
+}
+
+void DefineEntryPoint(Environment& env, EmitContext& ctx, Id main) {
+    const std::span interfaces(ctx.interfaces.data(), ctx.interfaces.size());
+    spv::ExecutionModel execution_model{};
+    switch (env.ShaderStage()) {
+    case Shader::Stage::Compute: {
+        const std::array<u32, 3> workgroup_size{env.WorkgroupSize()};
+        execution_model = spv::ExecutionModel::GLCompute;
+        ctx.AddExecutionMode(main, spv::ExecutionMode::LocalSize, workgroup_size[0],
+                             workgroup_size[1], workgroup_size[2]);
+        break;
+    }
+    case Shader::Stage::VertexB:
+        execution_model = spv::ExecutionModel::Vertex;
+        break;
+    case Shader::Stage::Fragment:
+        execution_model = spv::ExecutionModel::Fragment;
+        ctx.AddExecutionMode(main, spv::ExecutionMode::OriginUpperLeft);
+        break;
+    default:
+        throw NotImplementedException("Stage {}", env.ShaderStage());
+    }
+    ctx.AddEntryPoint(execution_model, main, "main", interfaces);
+}
+
 void SetupDenormControl(const Profile& profile, const IR::Program& program, EmitContext& ctx,
                         Id main_func) {
     if (!profile.support_float_controls) {
@@ -173,6 +210,25 @@ void SetupDenormControl(const Profile& profile, const IR::Program& program, Emit
     }
 }
 
+void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ctx) {
+    if (info.uses_sampled_1d) {
+        ctx.AddCapability(spv::Capability::Sampled1D);
+    }
+    if (info.uses_sparse_residency) {
+        ctx.AddCapability(spv::Capability::SparseResidency);
+    }
+    if (info.uses_demote_to_helper_invocation) {
+        ctx.AddExtension("SPV_EXT_demote_to_helper_invocation");
+        ctx.AddCapability(spv::Capability::DemoteToHelperInvocationEXT);
+    }
+    if (!profile.support_vertex_instance_id && (info.loads_instance_id || info.loads_vertex_id)) {
+        ctx.AddExtension("SPV_KHR_shader_draw_parameters");
+        ctx.AddCapability(spv::Capability::DrawParameters);
+    }
+    // TODO: Track this usage
+    ctx.AddCapability(spv::Capability::ImageGatherExtended);
+}
+
 Id PhiArgDef(EmitContext& ctx, IR::Inst* inst, size_t index) {
     // Phi nodes can have forward declarations, if an argument is not defined provide a forward
     // declaration of it. Invoke will take care of giving it the right definition when it's
@@ -202,53 +258,10 @@ Id PhiArgDef(EmitContext& ctx, IR::Inst* inst, size_t index) {
 std::vector<u32> EmitSPIRV(const Profile& profile, Environment& env, IR::Program& program,
                            u32& binding) {
     EmitContext ctx{profile, program, binding};
-    const Id void_function{ctx.TypeFunction(ctx.void_id)};
-    const Id func{ctx.OpFunction(ctx.void_id, spv::FunctionControlMask::MaskNone, void_function)};
-    for (IR::Block* const block : program.blocks) {
-        ctx.AddLabel(block->Definition<Id>());
-        for (IR::Inst& inst : block->Instructions()) {
-            EmitInst(ctx, &inst);
-        }
-    }
-    ctx.OpFunctionEnd();
-
-    const std::span interfaces(ctx.interfaces.data(), ctx.interfaces.size());
-    spv::ExecutionModel execution_model{};
-    switch (env.ShaderStage()) {
-    case Shader::Stage::Compute: {
-        const std::array<u32, 3> workgroup_size{env.WorkgroupSize()};
-        execution_model = spv::ExecutionModel::GLCompute;
-        ctx.AddExecutionMode(func, spv::ExecutionMode::LocalSize, workgroup_size[0],
-                             workgroup_size[1], workgroup_size[2]);
-        break;
-    }
-    case Shader::Stage::VertexB:
-        execution_model = spv::ExecutionModel::Vertex;
-        break;
-    case Shader::Stage::Fragment:
-        execution_model = spv::ExecutionModel::Fragment;
-        ctx.AddExecutionMode(func, spv::ExecutionMode::OriginUpperLeft);
-        break;
-    default:
-        throw NotImplementedException("Stage {}", env.ShaderStage());
-    }
-    ctx.AddEntryPoint(execution_model, func, "main", interfaces);
-
-    SetupDenormControl(profile, program, ctx, func);
-    const Info& info{program.info};
-    if (info.uses_sampled_1d) {
-        ctx.AddCapability(spv::Capability::Sampled1D);
-    }
-    if (info.uses_sparse_residency) {
-        ctx.AddCapability(spv::Capability::SparseResidency);
-    }
-    if (info.uses_demote_to_helper_invocation) {
-        ctx.AddExtension("SPV_EXT_demote_to_helper_invocation");
-        ctx.AddCapability(spv::Capability::DemoteToHelperInvocationEXT);
-    }
-    // TODO: Track this usage
-    ctx.AddCapability(spv::Capability::ImageGatherExtended);
-
+    const Id main{DefineMain(ctx, program)};
+    DefineEntryPoint(env, ctx, main);
+    SetupDenormControl(profile, program, ctx, main);
+    SetupCapabilities(profile, program.info, ctx);
     return ctx.Assemble();
 }
 
diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.h b/src/shader_recompiler/backend/spirv/emit_spirv.h
index 1fe65f8a9..e297a0e20 100644
--- a/src/shader_recompiler/backend/spirv/emit_spirv.h
+++ b/src/shader_recompiler/backend/spirv/emit_spirv.h
@@ -81,8 +81,8 @@ void EmitLoadStorageS8(EmitContext& ctx);
 void EmitLoadStorageU16(EmitContext& ctx);
 void EmitLoadStorageS16(EmitContext& ctx);
 Id EmitLoadStorage32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
-void EmitLoadStorage64(EmitContext& ctx);
-void EmitLoadStorage128(EmitContext& ctx);
+Id EmitLoadStorage64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
+Id EmitLoadStorage128(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
 void EmitWriteStorageU8(EmitContext& ctx);
 void EmitWriteStorageS8(EmitContext& ctx);
 void EmitWriteStorageU16(EmitContext& ctx);
diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp
index 02d115740..052b84151 100644
--- a/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp
@@ -19,6 +19,10 @@ Id InputAttrPointer(EmitContext& ctx, IR::Attribute attr) {
     case IR::Attribute::PositionZ:
     case IR::Attribute::PositionW:
         return ctx.OpAccessChain(ctx.input_f32, ctx.input_position, element_id());
+    case IR::Attribute::InstanceId:
+        return ctx.OpLoad(ctx.U32[1], ctx.instance_id);
+    case IR::Attribute::VertexId:
+        return ctx.OpLoad(ctx.U32[1], ctx.vertex_id);
     default:
         throw NotImplementedException("Read attribute {}", attr);
     }
@@ -125,6 +129,18 @@ Id EmitGetCbufU64(EmitContext& ctx, const IR::Value& binding, const IR::Value& o
 }
 
 Id EmitGetAttribute(EmitContext& ctx, IR::Attribute attr) {
+    if (!ctx.profile.support_vertex_instance_id) {
+        switch (attr) {
+        case IR::Attribute::InstanceId:
+            return ctx.OpISub(ctx.U32[1], ctx.OpLoad(ctx.U32[1], ctx.instance_index),
+                              ctx.OpLoad(ctx.U32[1], ctx.base_instance));
+        case IR::Attribute::VertexId:
+            return ctx.OpISub(ctx.U32[1], ctx.OpLoad(ctx.U32[1], ctx.vertex_index),
+                              ctx.OpLoad(ctx.U32[1], ctx.base_vertex));
+        default:
+            break;
+        }
+    }
     return ctx.OpLoad(ctx.F32[1], InputAttrPointer(ctx, attr));
 }
 
diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_memory.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_memory.cpp
index 7d3efc741..088bd3059 100644
--- a/src/shader_recompiler/backend/spirv/emit_spirv_memory.cpp
+++ b/src/shader_recompiler/backend/spirv/emit_spirv_memory.cpp
@@ -7,8 +7,8 @@
 #include "shader_recompiler/backend/spirv/emit_spirv.h"
 
 namespace Shader::Backend::SPIRV {
-
-static Id StorageIndex(EmitContext& ctx, const IR::Value& offset, size_t element_size) {
+namespace {
+Id StorageIndex(EmitContext& ctx, const IR::Value& offset, size_t element_size) {
     if (offset.IsImmediate()) {
         const u32 imm_offset{static_cast<u32>(offset.U32() / element_size)};
         return ctx.Constant(ctx.U32[1], imm_offset);
@@ -22,6 +22,32 @@ static Id StorageIndex(EmitContext& ctx, const IR::Value& offset, size_t element
     return ctx.OpShiftRightLogical(ctx.U32[1], index, shift_id);
 }
 
+Id EmitLoadStorage(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
+                   u32 num_components) {
+    // TODO: Support reinterpreting bindings, guaranteed to be aligned
+    if (!binding.IsImmediate()) {
+        throw NotImplementedException("Dynamic storage buffer indexing");
+    }
+    const Id ssbo{ctx.ssbos[binding.U32()]};
+    const Id base_index{StorageIndex(ctx, offset, sizeof(u32))};
+    std::array<Id, 4> components;
+    for (u32 element = 0; element < num_components; ++element) {
+        Id index{base_index};
+        if (element > 0) {
+            index = ctx.OpIAdd(ctx.U32[1], base_index, ctx.Constant(ctx.U32[1], element));
+        }
+        const Id pointer{ctx.OpAccessChain(ctx.storage_u32, ssbo, ctx.u32_zero_value, index)};
+        components[element] = ctx.OpLoad(ctx.U32[1], pointer);
+    }
+    if (num_components == 1) {
+        return components[0];
+    } else {
+        const std::span components_span(components.data(), num_components);
+        return ctx.OpCompositeConstruct(ctx.U32[num_components], components_span);
+    }
+}
+} // Anonymous namespace
+
 void EmitLoadGlobalU8(EmitContext&) {
     throw NotImplementedException("SPIR-V Instruction");
 }
@@ -95,21 +121,15 @@ void EmitLoadStorageS16(EmitContext&) {
 }
 
 Id EmitLoadStorage32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) {
-    if (!binding.IsImmediate()) {
-        throw NotImplementedException("Dynamic storage buffer indexing");
-    }
-    const Id ssbo{ctx.ssbos[binding.U32()]};
-    const Id index{StorageIndex(ctx, offset, sizeof(u32))};
-    const Id pointer{ctx.OpAccessChain(ctx.storage_u32, ssbo, ctx.u32_zero_value, index)};
-    return ctx.OpLoad(ctx.U32[1], pointer);
+    return EmitLoadStorage(ctx, binding, offset, 1);
 }
 
-void EmitLoadStorage64(EmitContext&) {
-    throw NotImplementedException("SPIR-V Instruction");
+Id EmitLoadStorage64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) {
+    return EmitLoadStorage(ctx, binding, offset, 2);
 }
 
-void EmitLoadStorage128(EmitContext&) {
-    throw NotImplementedException("SPIR-V Instruction");
+Id EmitLoadStorage128(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) {
+    return EmitLoadStorage(ctx, binding, offset, 4);
 }
 
 void EmitWriteStorageU8(EmitContext&) {
diff --git a/src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp b/src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp
index e72505d61..e7fa3fce0 100644
--- a/src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp
+++ b/src/shader_recompiler/ir_opt/collect_shader_info_pass.cpp
@@ -38,6 +38,12 @@ void GetAttribute(Info& info, IR::Attribute attribute) {
     case IR::Attribute::PositionW:
         info.loads_position = true;
         break;
+    case IR::Attribute::InstanceId:
+        info.loads_instance_id = true;
+        break;
+    case IR::Attribute::VertexId:
+        info.loads_vertex_id = true;
+        break;
     default:
         throw NotImplementedException("Get attribute {}", attribute);
     }
diff --git a/src/shader_recompiler/profile.h b/src/shader_recompiler/profile.h
index c6a143598..770299524 100644
--- a/src/shader_recompiler/profile.h
+++ b/src/shader_recompiler/profile.h
@@ -8,6 +8,7 @@ namespace Shader {
 
 struct Profile {
     bool unified_descriptor_binding{};
+    bool support_vertex_instance_id{};
     bool support_float_controls{};
     bool support_separate_denorm_behavior{};
     bool support_separate_rounding_mode{};
diff --git a/src/shader_recompiler/shader_info.h b/src/shader_recompiler/shader_info.h
index 6eff762e2..f97730b34 100644
--- a/src/shader_recompiler/shader_info.h
+++ b/src/shader_recompiler/shader_info.h
@@ -59,6 +59,8 @@ struct Info {
 
     std::array<bool, 32> loads_generics{};
     bool loads_position{};
+    bool loads_instance_id{};
+    bool loads_vertex_id{};
 
     std::array<bool, 8> stores_frag_color{};
     bool stores_frag_depth{};