vk_shader_decompiler: Misc changes

Update Sirit and its usage in vk_shader_decompiler. Highlights:
- Implement tessellation shaders
- Implement geometry shaders
- Implement some missing features
- Use native half float instructions when available.

3 files changed, 1648 insertions, 697 deletions

diff --git a/externals/sirit b/externals/sirit
-Subproject f7c4b07a7e14edb1dcd93bc9879c823423705c2
+Subproject e1a6729df7f11e33f6dc0939b18995a57c8bf3d
diff --git a/src/video_core/renderer_vulkan/vk_shader_decompiler.cpp b/src/video_core/renderer_vulkan/vk_shader_decompiler.cpp
index 76894275b..8f517bdc1 100644
--- a/src/video_core/renderer_vulkan/vk_shader_decompiler.cpp
+++ b/src/video_core/renderer_vulkan/vk_shader_decompiler.cpp
@@ -3,8 +3,10 @@
 // Refer to the license.txt file included.
 
 #include <functional>
+#include <limits>
 #include <map>
-#include <set>
+#include <type_traits>
+#include <utility>
 
 #include <fmt/format.h>
 
@@ -23,7 +25,9 @@
 #include "video_core/shader/node.h"
 #include "video_core/shader/shader_ir.h"
 
-namespace Vulkan::VKShader {
+namespace Vulkan {
+
+namespace {
 
 using Sirit::Id;
 using Tegra::Engines::ShaderType;
@@ -35,22 +39,60 @@ using namespace VideoCommon::Shader;
 using Maxwell = Tegra::Engines::Maxwell3D::Regs;
 using Operation = const OperationNode&;
 
+class ASTDecompiler;
+class ExprDecompiler;
+
 // TODO(Rodrigo): Use rasterizer's value
-constexpr u32 MAX_CONSTBUFFER_FLOATS = 0x4000;
-constexpr u32 MAX_CONSTBUFFER_ELEMENTS = MAX_CONSTBUFFER_FLOATS / 4;
-constexpr u32 STAGE_BINDING_STRIDE = 0x100;
+constexpr u32 MaxConstBufferFloats = 0x4000;
+constexpr u32 MaxConstBufferElements = MaxConstBufferFloats / 4;
+
+constexpr u32 NumInputPatches = 32; // This value seems to be the standard
+
+enum class Type { Void, Bool, Bool2, Float, Int, Uint, HalfFloat };
+
+class Expression final {
+public:
+    Expression(Id id, Type type) : id{id}, type{type} {
+        ASSERT(type != Type::Void);
+    }
+    Expression() : type{Type::Void} {}
 
-enum class Type { Bool, Bool2, Float, Int, Uint, HalfFloat };
+    Id id{};
+    Type type{};
+};
+static_assert(std::is_standard_layout_v<Expression>);
 
-struct SamplerImage {
-    Id image_type;
-    Id sampled_image_type;
-    Id sampler;
+struct TexelBuffer {
+    Id image_type{};
+    Id image{};
 };
 
-namespace {
+struct SampledImage {
+    Id image_type{};
+    Id sampled_image_type{};
+    Id sampler{};
+};
+
+struct StorageImage {
+    Id image_type{};
+    Id image{};
+};
+
+struct AttributeType {
+    Type type;
+    Id scalar;
+    Id vector;
+};
+
+struct VertexIndices {
+    std::optional<u32> position;
+    std::optional<u32> viewport;
+    std::optional<u32> point_size;
+    std::optional<u32> clip_distances;
+};
 
 spv::Dim GetSamplerDim(const Sampler& sampler) {
+    ASSERT(!sampler.IsBuffer());
     switch (sampler.GetType()) {
     case Tegra::Shader::TextureType::Texture1D:
         return spv::Dim::Dim1D;
@@ -66,6 +108,138 @@ spv::Dim GetSamplerDim(const Sampler& sampler) {
     }
 }
 
+std::pair<spv::Dim, bool> GetImageDim(const Image& image) {
+    switch (image.GetType()) {
+    case Tegra::Shader::ImageType::Texture1D:
+        return {spv::Dim::Dim1D, false};
+    case Tegra::Shader::ImageType::TextureBuffer:
+        return {spv::Dim::Buffer, false};
+    case Tegra::Shader::ImageType::Texture1DArray:
+        return {spv::Dim::Dim1D, true};
+    case Tegra::Shader::ImageType::Texture2D:
+        return {spv::Dim::Dim2D, false};
+    case Tegra::Shader::ImageType::Texture2DArray:
+        return {spv::Dim::Dim2D, true};
+    case Tegra::Shader::ImageType::Texture3D:
+        return {spv::Dim::Dim3D, false};
+    default:
+        UNIMPLEMENTED_MSG("Unimplemented image type={}", static_cast<u32>(image.GetType()));
+        return {spv::Dim::Dim2D, false};
+    }
+}
+
+/// Returns the number of vertices present in a primitive topology.
+u32 GetNumPrimitiveTopologyVertices(Maxwell::PrimitiveTopology primitive_topology) {
+    switch (primitive_topology) {
+    case Maxwell::PrimitiveTopology::Points:
+        return 1;
+    case Maxwell::PrimitiveTopology::Lines:
+    case Maxwell::PrimitiveTopology::LineLoop:
+    case Maxwell::PrimitiveTopology::LineStrip:
+        return 2;
+    case Maxwell::PrimitiveTopology::Triangles:
+    case Maxwell::PrimitiveTopology::TriangleStrip:
+    case Maxwell::PrimitiveTopology::TriangleFan:
+        return 3;
+    case Maxwell::PrimitiveTopology::LinesAdjacency:
+    case Maxwell::PrimitiveTopology::LineStripAdjacency:
+        return 4;
+    case Maxwell::PrimitiveTopology::TrianglesAdjacency:
+    case Maxwell::PrimitiveTopology::TriangleStripAdjacency:
+        return 6;
+    case Maxwell::PrimitiveTopology::Quads:
+        UNIMPLEMENTED_MSG("Quads");
+        return 3;
+    case Maxwell::PrimitiveTopology::QuadStrip:
+        UNIMPLEMENTED_MSG("QuadStrip");
+        return 3;
+    case Maxwell::PrimitiveTopology::Polygon:
+        UNIMPLEMENTED_MSG("Polygon");
+        return 3;
+    case Maxwell::PrimitiveTopology::Patches:
+        UNIMPLEMENTED_MSG("Patches");
+        return 3;
+    default:
+        UNREACHABLE();
+        return 3;
+    }
+}
+
+spv::ExecutionMode GetExecutionMode(Maxwell::TessellationPrimitive primitive) {
+    switch (primitive) {
+    case Maxwell::TessellationPrimitive::Isolines:
+        return spv::ExecutionMode::Isolines;
+    case Maxwell::TessellationPrimitive::Triangles:
+        return spv::ExecutionMode::Triangles;
+    case Maxwell::TessellationPrimitive::Quads:
+        return spv::ExecutionMode::Quads;
+    }
+    UNREACHABLE();
+    return spv::ExecutionMode::Triangles;
+}
+
+spv::ExecutionMode GetExecutionMode(Maxwell::TessellationSpacing spacing) {
+    switch (spacing) {
+    case Maxwell::TessellationSpacing::Equal:
+        return spv::ExecutionMode::SpacingEqual;
+    case Maxwell::TessellationSpacing::FractionalOdd:
+        return spv::ExecutionMode::SpacingFractionalOdd;
+    case Maxwell::TessellationSpacing::FractionalEven:
+        return spv::ExecutionMode::SpacingFractionalEven;
+    }
+    UNREACHABLE();
+    return spv::ExecutionMode::SpacingEqual;
+}
+
+spv::ExecutionMode GetExecutionMode(Maxwell::PrimitiveTopology input_topology) {
+    switch (input_topology) {
+    case Maxwell::PrimitiveTopology::Points:
+        return spv::ExecutionMode::InputPoints;
+    case Maxwell::PrimitiveTopology::Lines:
+    case Maxwell::PrimitiveTopology::LineLoop:
+    case Maxwell::PrimitiveTopology::LineStrip:
+        return spv::ExecutionMode::InputLines;
+    case Maxwell::PrimitiveTopology::Triangles:
+    case Maxwell::PrimitiveTopology::TriangleStrip:
+    case Maxwell::PrimitiveTopology::TriangleFan:
+        return spv::ExecutionMode::Triangles;
+    case Maxwell::PrimitiveTopology::LinesAdjacency:
+    case Maxwell::PrimitiveTopology::LineStripAdjacency:
+        return spv::ExecutionMode::InputLinesAdjacency;
+    case Maxwell::PrimitiveTopology::TrianglesAdjacency:
+    case Maxwell::PrimitiveTopology::TriangleStripAdjacency:
+        return spv::ExecutionMode::InputTrianglesAdjacency;
+    case Maxwell::PrimitiveTopology::Quads:
+        UNIMPLEMENTED_MSG("Quads");
+        return spv::ExecutionMode::Triangles;
+    case Maxwell::PrimitiveTopology::QuadStrip:
+        UNIMPLEMENTED_MSG("QuadStrip");
+        return spv::ExecutionMode::Triangles;
+    case Maxwell::PrimitiveTopology::Polygon:
+        UNIMPLEMENTED_MSG("Polygon");
+        return spv::ExecutionMode::Triangles;
+    case Maxwell::PrimitiveTopology::Patches:
+        UNIMPLEMENTED_MSG("Patches");
+        return spv::ExecutionMode::Triangles;
+    }
+    UNREACHABLE();
+    return spv::ExecutionMode::Triangles;
+}
+
+spv::ExecutionMode GetExecutionMode(Tegra::Shader::OutputTopology output_topology) {
+    switch (output_topology) {
+    case Tegra::Shader::OutputTopology::PointList:
+        return spv::ExecutionMode::OutputPoints;
+    case Tegra::Shader::OutputTopology::LineStrip:
+        return spv::ExecutionMode::OutputLineStrip;
+    case Tegra::Shader::OutputTopology::TriangleStrip:
+        return spv::ExecutionMode::OutputTriangleStrip;
+    default:
+        UNREACHABLE();
+        return spv::ExecutionMode::OutputPoints;
+    }
+}
+
 /// Returns true if an attribute index is one of the 32 generic attributes
 constexpr bool IsGenericAttribute(Attribute::Index attribute) {
     return attribute >= Attribute::Index::Attribute_0 &&
@@ -73,7 +247,7 @@ constexpr bool IsGenericAttribute(Attribute::Index attribute) {
 }
 
 /// Returns the location of a generic attribute
-constexpr u32 GetGenericAttributeLocation(Attribute::Index attribute) {
+u32 GetGenericAttributeLocation(Attribute::Index attribute) {
     ASSERT(IsGenericAttribute(attribute));
     return static_cast<u32>(attribute) - static_cast<u32>(Attribute::Index::Attribute_0);
 }
@@ -87,20 +261,146 @@ bool IsPrecise(Operation operand) {
     return false;
 }
 
-} // namespace
-
-class ASTDecompiler;
-class ExprDecompiler;
-
-class SPIRVDecompiler : public Sirit::Module {
+class SPIRVDecompiler final : public Sirit::Module {
 public:
-    explicit SPIRVDecompiler(const VKDevice& device, const ShaderIR& ir, ShaderType stage)
-        : Module(0x00010300), device{device}, ir{ir}, stage{stage}, header{ir.GetHeader()} {
+    explicit SPIRVDecompiler(const VKDevice& device, const ShaderIR& ir, ShaderType stage,
+                             const Specialization& specialization)
+        : Module(0x00010300), device{device}, ir{ir}, stage{stage}, header{ir.GetHeader()},
+          specialization{specialization} {
         AddCapability(spv::Capability::Shader);
+        AddCapability(spv::Capability::UniformAndStorageBuffer16BitAccess);
+        AddCapability(spv::Capability::ImageQuery);
+        AddCapability(spv::Capability::Image1D);
+        AddCapability(spv::Capability::ImageBuffer);
+        AddCapability(spv::Capability::ImageGatherExtended);
+        AddCapability(spv::Capability::SampledBuffer);
+        AddCapability(spv::Capability::StorageImageWriteWithoutFormat);
+        AddCapability(spv::Capability::SubgroupBallotKHR);
+        AddCapability(spv::Capability::SubgroupVoteKHR);
+        AddExtension("SPV_KHR_shader_ballot");
+        AddExtension("SPV_KHR_subgroup_vote");
         AddExtension("SPV_KHR_storage_buffer_storage_class");
         AddExtension("SPV_KHR_variable_pointers");
+
+        if (ir.UsesViewportIndex()) {
+            AddCapability(spv::Capability::MultiViewport);
+            if (device.IsExtShaderViewportIndexLayerSupported()) {
+                AddExtension("SPV_EXT_shader_viewport_index_layer");
+                AddCapability(spv::Capability::ShaderViewportIndexLayerEXT);
+            }
+        }
+
+        if (device.IsFloat16Supported()) {
+            AddCapability(spv::Capability::Float16);
+        }
+        t_scalar_half = Name(TypeFloat(device.IsFloat16Supported() ? 16 : 32), "scalar_half");
+        t_half = Name(TypeVector(t_scalar_half, 2), "half");
+
+        const Id main = Decompile();
+
+        switch (stage) {
+        case ShaderType::Vertex:
+            AddEntryPoint(spv::ExecutionModel::Vertex, main, "main", interfaces);
+            break;
+        case ShaderType::TesselationControl:
+            AddCapability(spv::Capability::Tessellation);
+            AddEntryPoint(spv::ExecutionModel::TessellationControl, main, "main", interfaces);
+            AddExecutionMode(main, spv::ExecutionMode::OutputVertices,
+                             header.common2.threads_per_input_primitive);
+            break;
+        case ShaderType::TesselationEval:
+            AddCapability(spv::Capability::Tessellation);
+            AddEntryPoint(spv::ExecutionModel::TessellationEvaluation, main, "main", interfaces);
+            AddExecutionMode(main, GetExecutionMode(specialization.tessellation.primitive));
+            AddExecutionMode(main, GetExecutionMode(specialization.tessellation.spacing));
+            AddExecutionMode(main, specialization.tessellation.clockwise
+                                       ? spv::ExecutionMode::VertexOrderCw
+                                       : spv::ExecutionMode::VertexOrderCcw);
+            break;
+        case ShaderType::Geometry:
+            AddCapability(spv::Capability::Geometry);
+            AddEntryPoint(spv::ExecutionModel::Geometry, main, "main", interfaces);
+            AddExecutionMode(main, GetExecutionMode(specialization.primitive_topology));
+            AddExecutionMode(main, GetExecutionMode(header.common3.output_topology));
+            AddExecutionMode(main, spv::ExecutionMode::OutputVertices,
+                             header.common4.max_output_vertices);
+            // TODO(Rodrigo): Where can we get this info from?
+            AddExecutionMode(main, spv::ExecutionMode::Invocations, 1U);
+            break;
+        case ShaderType::Fragment:
+            AddEntryPoint(spv::ExecutionModel::Fragment, main, "main", interfaces);
+            AddExecutionMode(main, spv::ExecutionMode::OriginUpperLeft);
+            if (header.ps.omap.depth) {
+                AddExecutionMode(main, spv::ExecutionMode::DepthReplacing);
+            }
+            break;
+        case ShaderType::Compute:
+            const auto workgroup_size = specialization.workgroup_size;
+            AddExecutionMode(main, spv::ExecutionMode::LocalSize, workgroup_size[0],
+                             workgroup_size[1], workgroup_size[2]);
+            AddEntryPoint(spv::ExecutionModel::GLCompute, main, "main", interfaces);
+            break;
+        }
     }
 
+private:
+    Id Decompile() {
+        DeclareCommon();
+        DeclareVertex();
+        DeclareTessControl();
+        DeclareTessEval();
+        DeclareGeometry();
+        DeclareFragment();
+        DeclareCompute();
+        DeclareRegisters();
+        DeclarePredicates();
+        DeclareLocalMemory();
+        DeclareSharedMemory();
+        DeclareInternalFlags();
+        DeclareInputAttributes();
+        DeclareOutputAttributes();
+
+        u32 binding = specialization.base_binding;
+        binding = DeclareConstantBuffers(binding);
+        binding = DeclareGlobalBuffers(binding);
+        binding = DeclareTexelBuffers(binding);
+        binding = DeclareSamplers(binding);
+        binding = DeclareImages(binding);
+
+        const Id main = OpFunction(t_void, {}, TypeFunction(t_void));
+        AddLabel();
+
+        if (ir.IsDecompiled()) {
+            DeclareFlowVariables();
+            DecompileAST();
+        } else {
+            AllocateLabels();
+            DecompileBranchMode();
+        }
+
+        OpReturn();
+        OpFunctionEnd();
+
+        return main;
+    }
+
+    void DefinePrologue() {
+        if (stage == ShaderType::Vertex) {
+            // Clear Position to avoid reading trash on the Z conversion.
+            const auto position_index = out_indices.position.value();
+            const Id position = AccessElement(t_out_float4, out_vertex, position_index);
+            OpStore(position, v_varying_default);
+
+            if (specialization.point_size) {
+                const u32 point_size_index = out_indices.point_size.value();
+                const Id out_point_size = AccessElement(t_out_float, out_vertex, point_size_index);
+                OpStore(out_point_size, Constant(t_float, *specialization.point_size));
+            }
+        }
+    }
+
+    void DecompileAST();
+
     void DecompileBranchMode() {
         const u32 first_address = ir.GetBasicBlocks().begin()->first;
         const Id loop_label = OpLabel("loop");
@@ -111,14 +411,15 @@ public:
 
         std::vector<Sirit::Literal> literals;
         std::vector<Id> branch_labels;
-        for (const auto& pair : labels) {
-            const auto [literal, label] = pair;
+        for (const auto& [literal, label] : labels) {
             literals.push_back(literal);
             branch_labels.push_back(label);
         }
 
-        jmp_to = Emit(OpVariable(TypePointer(spv::StorageClass::Function, t_uint),
-                                 spv::StorageClass::Function, Constant(t_uint, first_address)));
+        jmp_to = OpVariable(TypePointer(spv::StorageClass::Function, t_uint),
+                            spv::StorageClass::Function, Constant(t_uint, first_address));
+        AddLocalVariable(jmp_to);
+
         std::tie(ssy_flow_stack, ssy_flow_stack_top) = CreateFlowStack();
         std::tie(pbk_flow_stack, pbk_flow_stack_top) = CreateFlowStack();
 
@@ -128,151 +429,118 @@ public:
         Name(pbk_flow_stack, "pbk_flow_stack");
         Name(pbk_flow_stack_top, "pbk_flow_stack_top");
 
-        Emit(OpBranch(loop_label));
-        Emit(loop_label);
-        Emit(OpLoopMerge(merge_label, continue_label, spv::LoopControlMask::Unroll));
-        Emit(OpBranch(dummy_label));
+        DefinePrologue();
+
+        OpBranch(loop_label);
+        AddLabel(loop_label);
+        OpLoopMerge(merge_label, continue_label, spv::LoopControlMask::MaskNone);
+        OpBranch(dummy_label);
 
-        Emit(dummy_label);
+        AddLabel(dummy_label);
         const Id default_branch = OpLabel();
-        const Id jmp_to_load = Emit(OpLoad(t_uint, jmp_to));
-        Emit(OpSelectionMerge(jump_label, spv::SelectionControlMask::MaskNone));
-        Emit(OpSwitch(jmp_to_load, default_branch, literals, branch_labels));
+        const Id jmp_to_load = OpLoad(t_uint, jmp_to);
+        OpSelectionMerge(jump_label, spv::SelectionControlMask::MaskNone);
+        OpSwitch(jmp_to_load, default_branch, literals, branch_labels);
 
-        Emit(default_branch);
-        Emit(OpReturn());
+        AddLabel(default_branch);
+        OpReturn();
 
-        for (const auto& pair : ir.GetBasicBlocks()) {
-            const auto& [address, bb] = pair;
-            Emit(labels.at(address));
+        for (const auto& [address, bb] : ir.GetBasicBlocks()) {
+            AddLabel(labels.at(address));
 
             VisitBasicBlock(bb);
 
             const auto next_it = labels.lower_bound(address + 1);
             const Id next_label = next_it != labels.end() ? next_it->second : default_branch;
-            Emit(OpBranch(next_label));
+            OpBranch(next_label);
         }
 
-        Emit(jump_label);
-        Emit(OpBranch(continue_label));
-        Emit(continue_label);
-        Emit(OpBranch(loop_label));
-        Emit(merge_label);
+        AddLabel(jump_label);
+        OpBranch(continue_label);
+        AddLabel(continue_label);
+        OpBranch(loop_label);
+        AddLabel(merge_label);
     }
 
-    void DecompileAST();
+private:
+    friend class ASTDecompiler;
+    friend class ExprDecompiler;
 
-    void Decompile() {
-        const bool is_fully_decompiled = ir.IsDecompiled();
-        AllocateBindings();
-        if (!is_fully_decompiled) {
-            AllocateLabels();
-        }
+    static constexpr auto INTERNAL_FLAGS_COUNT = static_cast<std::size_t>(InternalFlag::Amount);
 
-        DeclareVertex();
-        DeclareGeometry();
-        DeclareFragment();
-        DeclareRegisters();
-        DeclarePredicates();
-        if (is_fully_decompiled) {
-            DeclareFlowVariables();
+    void AllocateLabels() {
+        for (const auto& pair : ir.GetBasicBlocks()) {
+            const u32 address = pair.first;
+            labels.emplace(address, OpLabel(fmt::format("label_0x{:x}", address)));
         }
-        DeclareLocalMemory();
-        DeclareInternalFlags();
-        DeclareInputAttributes();
-        DeclareOutputAttributes();
-        DeclareConstantBuffers();
-        DeclareGlobalBuffers();
-        DeclareSamplers();
+    }
 
-        execute_function =
-            Emit(OpFunction(t_void, spv::FunctionControlMask::Inline, TypeFunction(t_void)));
-        Emit(OpLabel());
+    void DeclareCommon() {
+        thread_id =
+            DeclareInputBuiltIn(spv::BuiltIn::SubgroupLocalInvocationId, t_in_uint, "thread_id");
+    }
 
-        if (is_fully_decompiled) {
-            DecompileAST();
-        } else {
-            DecompileBranchMode();
+    void DeclareVertex() {
+        if (stage != ShaderType::Vertex) {
+            return;
         }
+        Id out_vertex_struct;
+        std::tie(out_vertex_struct, out_indices) = DeclareVertexStruct();
+        const Id vertex_ptr = TypePointer(spv::StorageClass::Output, out_vertex_struct);
+        out_vertex = OpVariable(vertex_ptr, spv::StorageClass::Output);
+        interfaces.push_back(AddGlobalVariable(Name(out_vertex, "out_vertex")));
 
-        Emit(OpReturn());
-        Emit(OpFunctionEnd());
+        // Declare input attributes
+        vertex_index = DeclareInputBuiltIn(spv::BuiltIn::VertexIndex, t_in_uint, "vertex_index");
+        instance_index =
+            DeclareInputBuiltIn(spv::BuiltIn::InstanceIndex, t_in_uint, "instance_index");
     }
 
-    ShaderEntries GetShaderEntries() const {
-        ShaderEntries entries;
-        entries.const_buffers_base_binding = const_buffers_base_binding;
-        entries.global_buffers_base_binding = global_buffers_base_binding;
-        entries.samplers_base_binding = samplers_base_binding;
-        for (const auto& cbuf : ir.GetConstantBuffers()) {
-            entries.const_buffers.emplace_back(cbuf.second, cbuf.first);
-        }
-        for (const auto& gmem_pair : ir.GetGlobalMemory()) {
-            const auto& [base, usage] = gmem_pair;
-            entries.global_buffers.emplace_back(base.cbuf_index, base.cbuf_offset);
-        }
-        for (const auto& sampler : ir.GetSamplers()) {
-            entries.samplers.emplace_back(sampler);
-        }
-        for (const auto& attribute : ir.GetInputAttributes()) {
-            if (IsGenericAttribute(attribute)) {
-                entries.attributes.insert(GetGenericAttributeLocation(attribute));
-            }
+    void DeclareTessControl() {
+        if (stage != ShaderType::TesselationControl) {
+            return;
         }
-        entries.clip_distances = ir.GetClipDistances();
-        entries.shader_length = ir.GetLength();
-        entries.entry_function = execute_function;
-        entries.interfaces = interfaces;
-        return entries;
-    }
-
-private:
-    friend class ASTDecompiler;
-    friend class ExprDecompiler;
-
-    static constexpr auto INTERNAL_FLAGS_COUNT = static_cast<std::size_t>(InternalFlag::Amount);
+        DeclareInputVertexArray(NumInputPatches);
+        DeclareOutputVertexArray(header.common2.threads_per_input_primitive);
 
-    void AllocateBindings() {
-        const u32 binding_base = static_cast<u32>(stage) * STAGE_BINDING_STRIDE;
-        u32 binding_iterator = binding_base;
+        tess_level_outer = DeclareBuiltIn(
+            spv::BuiltIn::TessLevelOuter, spv::StorageClass::Output,
+            TypePointer(spv::StorageClass::Output, TypeArray(t_float, Constant(t_uint, 4U))),
+            "tess_level_outer");
+        Decorate(tess_level_outer, spv::Decoration::Patch);
 
-        const auto Allocate = [&binding_iterator](std::size_t count) {
-            const u32 current_binding = binding_iterator;
-            binding_iterator += static_cast<u32>(count);
-            return current_binding;
-        };
-        const_buffers_base_binding = Allocate(ir.GetConstantBuffers().size());
-        global_buffers_base_binding = Allocate(ir.GetGlobalMemory().size());
-        samplers_base_binding = Allocate(ir.GetSamplers().size());
+        tess_level_inner = DeclareBuiltIn(
+            spv::BuiltIn::TessLevelInner, spv::StorageClass::Output,
+            TypePointer(spv::StorageClass::Output, TypeArray(t_float, Constant(t_uint, 2U))),
+            "tess_level_inner");
+        Decorate(tess_level_inner, spv::Decoration::Patch);
 
-        ASSERT_MSG(binding_iterator - binding_base < STAGE_BINDING_STRIDE,
-                   "Stage binding stride is too small");
+        invocation_id = DeclareInputBuiltIn(spv::BuiltIn::InvocationId, t_in_int, "invocation_id");
     }
 
-    void AllocateLabels() {
-        for (const auto& pair : ir.GetBasicBlocks()) {
-            const u32 address = pair.first;
-            labels.emplace(address, OpLabel(fmt::format("label_0x{:x}", address)));
-        }
-    }
-
-    void DeclareVertex() {
-        if (stage != ShaderType::Vertex)
+    void DeclareTessEval() {
+        if (stage != ShaderType::TesselationEval) {
             return;
+        }
+        DeclareInputVertexArray(NumInputPatches);
+        DeclareOutputVertex();
 
-        DeclareVertexRedeclarations();
+        tess_coord = DeclareInputBuiltIn(spv::BuiltIn::TessCoord, t_in_float3, "tess_coord");
     }
 
     void DeclareGeometry() {
-        if (stage != ShaderType::Geometry)
+        if (stage != ShaderType::Geometry) {
             return;
-
-        UNIMPLEMENTED();
+        }
+        const u32 num_input = GetNumPrimitiveTopologyVertices(specialization.primitive_topology);
+        DeclareInputVertexArray(num_input);
+        DeclareOutputVertex();
     }
 
     void DeclareFragment() {
-        if (stage != ShaderType::Fragment)
+        if (stage != ShaderType::Fragment) {
             return;
+        }
 
         for (u32 rt = 0; rt < static_cast<u32>(frag_colors.size()); ++rt) {
             if (!IsRenderTargetUsed(rt)) {
@@ -296,10 +564,19 @@ private:
             interfaces.push_back(frag_depth);
         }
 
-        frag_coord = DeclareBuiltIn(spv::BuiltIn::FragCoord, spv::StorageClass::Input, t_in_float4,
-                                    "frag_coord");
-        front_facing = DeclareBuiltIn(spv::BuiltIn::FrontFacing, spv::StorageClass::Input,
-                                      t_in_bool, "front_facing");
+        frag_coord = DeclareInputBuiltIn(spv::BuiltIn::FragCoord, t_in_float4, "frag_coord");
+        front_facing = DeclareInputBuiltIn(spv::BuiltIn::FrontFacing, t_in_bool, "front_facing");
+        point_coord = DeclareInputBuiltIn(spv::BuiltIn::PointCoord, t_in_float2, "point_coord");
+    }
+
+    void DeclareCompute() {
+        if (stage != ShaderType::Compute) {
+            return;
+        }
+
+        workgroup_id = DeclareInputBuiltIn(spv::BuiltIn::WorkgroupId, t_in_uint3, "workgroup_id");
+        local_invocation_id =
+            DeclareInputBuiltIn(spv::BuiltIn::LocalInvocationId, t_in_uint3, "local_invocation_id");
     }
 
     void DeclareRegisters() {
@@ -327,21 +604,44 @@ private:
     }
 
     void DeclareLocalMemory() {
-        if (const u64 local_memory_size = header.GetLocalMemorySize(); local_memory_size > 0) {
-            const auto element_count = static_cast<u32>(Common::AlignUp(local_memory_size, 4) / 4);
-            const Id type_array = TypeArray(t_float, Constant(t_uint, element_count));
-            const Id type_pointer = TypePointer(spv::StorageClass::Private, type_array);
-            Name(type_pointer, "LocalMemory");
+        // TODO(Rodrigo): Unstub kernel local memory size and pass it from a register at
+        // specialization time.
+        const u64 lmem_size = stage == ShaderType::Compute ? 0x400 : header.GetLocalMemorySize();
+        if (lmem_size == 0) {
+            return;
+        }
+        const auto element_count = static_cast<u32>(Common::AlignUp(lmem_size, 4) / 4);
+        const Id type_array = TypeArray(t_float, Constant(t_uint, element_count));
+        const Id type_pointer = TypePointer(spv::StorageClass::Private, type_array);
+        Name(type_pointer, "LocalMemory");
+
+        local_memory =
+            OpVariable(type_pointer, spv::StorageClass::Private, ConstantNull(type_array));
+        AddGlobalVariable(Name(local_memory, "local_memory"));
+    }
+
+    void DeclareSharedMemory() {
+        if (stage != ShaderType::Compute) {
+            return;
+        }
+        t_smem_uint = TypePointer(spv::StorageClass::Workgroup, t_uint);
 
-            local_memory =
-                OpVariable(type_pointer, spv::StorageClass::Private, ConstantNull(type_array));
-            AddGlobalVariable(Name(local_memory, "local_memory"));
+        const u32 smem_size = specialization.shared_memory_size;
+        if (smem_size == 0) {
+            // Avoid declaring an empty array.
+            return;
         }
+        const auto element_count = static_cast<u32>(Common::AlignUp(smem_size, 4) / 4);
+        const Id type_array = TypeArray(t_uint, Constant(t_uint, element_count));
+        const Id type_pointer = TypePointer(spv::StorageClass::Workgroup, type_array);
+        Name(type_pointer, "SharedMemory");
+
+        shared_memory = OpVariable(type_pointer, spv::StorageClass::Workgroup);
+        AddGlobalVariable(Name(shared_memory, "shared_memory"));
     }
 
     void DeclareInternalFlags() {
-        constexpr std::array<const char*, INTERNAL_FLAGS_COUNT> names = {"zero", "sign", "carry",
-                                                                         "overflow"};
+        constexpr std::array names = {"zero", "sign", "carry", "overflow"};
         for (std::size_t flag = 0; flag < INTERNAL_FLAGS_COUNT; ++flag) {
             const auto flag_code = static_cast<InternalFlag>(flag);
             const Id id = OpVariable(t_prv_bool, spv::StorageClass::Private, v_false);
@@ -349,17 +649,53 @@ private:
         }
     }
 
+    void DeclareInputVertexArray(u32 length) {
+        constexpr auto storage = spv::StorageClass::Input;
+        std::tie(in_indices, in_vertex) = DeclareVertexArray(storage, "in_indices", length);
+    }
+
+    void DeclareOutputVertexArray(u32 length) {
+        constexpr auto storage = spv::StorageClass::Output;
+        std::tie(out_indices, out_vertex) = DeclareVertexArray(storage, "out_indices", length);
+    }
+
+    std::tuple<VertexIndices, Id> DeclareVertexArray(spv::StorageClass storage_class,
+                                                     std::string name, u32 length) {
+        const auto [struct_id, indices] = DeclareVertexStruct();
+        const Id vertex_array = TypeArray(struct_id, Constant(t_uint, length));
+        const Id vertex_ptr = TypePointer(storage_class, vertex_array);
+        const Id vertex = OpVariable(vertex_ptr, storage_class);
+        AddGlobalVariable(Name(vertex, std::move(name)));
+        interfaces.push_back(vertex);
+        return {indices, vertex};
+    }
+
+    void DeclareOutputVertex() {
+        Id out_vertex_struct;
+        std::tie(out_vertex_struct, out_indices) = DeclareVertexStruct();
+        const Id out_vertex_ptr = TypePointer(spv::StorageClass::Output, out_vertex_struct);
+        out_vertex = OpVariable(out_vertex_ptr, spv::StorageClass::Output);
+        interfaces.push_back(AddGlobalVariable(Name(out_vertex, "out_vertex")));
+    }
+
     void DeclareInputAttributes() {
         for (const auto index : ir.GetInputAttributes()) {
             if (!IsGenericAttribute(index)) {
                 continue;
             }
 
-            UNIMPLEMENTED_IF(stage == ShaderType::Geometry);
-
             const u32 location = GetGenericAttributeLocation(index);
-            const Id id = OpVariable(t_in_float4, spv::StorageClass::Input);
-            Name(AddGlobalVariable(id), fmt::format("in_attr{}", location));
+            const auto type_descriptor = GetAttributeType(location);
+            Id type;
+            if (IsInputAttributeArray()) {
+                type = GetTypeVectorDefinitionLut(type_descriptor.type).at(3);
+                type = TypeArray(type, Constant(t_uint, GetNumInputVertices()));
+                type = TypePointer(spv::StorageClass::Input, type);
+            } else {
+                type = type_descriptor.vector;
+            }
+            const Id id = OpVariable(type, spv::StorageClass::Input);
+            AddGlobalVariable(Name(id, fmt::format("in_attr{}", location)));
             input_attributes.emplace(index, id);
             interfaces.push_back(id);
 
@@ -389,8 +725,21 @@ private:
             if (!IsGenericAttribute(index)) {
                 continue;
             }
-            const auto location = GetGenericAttributeLocation(index);
-            const Id id = OpVariable(t_out_float4, spv::StorageClass::Output);
+            const u32 location = GetGenericAttributeLocation(index);
+            Id type = t_float4;
+            Id varying_default = v_varying_default;
+            if (IsOutputAttributeArray()) {
+                const u32 num = GetNumOutputVertices();
+                type = TypeArray(type, Constant(t_uint, num));
+                if (device.GetDriverID() != vk::DriverIdKHR::eIntelProprietaryWindows) {
+                    // Intel's proprietary driver fails to setup defaults for arrayed output
+                    // attributes.
+                    varying_default = ConstantComposite(type, std::vector(num, varying_default));
+                }
+            }
+            type = TypePointer(spv::StorageClass::Output, type);
+
+            const Id id = OpVariable(type, spv::StorageClass::Output, varying_default);
             Name(AddGlobalVariable(id), fmt::format("out_attr{}", location));
             output_attributes.emplace(index, id);
             interfaces.push_back(id);
@@ -399,10 +748,8 @@ private:
         }
     }
 
-    void DeclareConstantBuffers() {
-        u32 binding = const_buffers_base_binding;
-        for (const auto& entry : ir.GetConstantBuffers()) {
-            const auto [index, size] = entry;
+    u32 DeclareConstantBuffers(u32 binding) {
+        for (const auto& [index, size] : ir.GetConstantBuffers()) {
             const Id type = device.IsKhrUniformBufferStandardLayoutSupported() ? t_cbuf_scalar_ubo
                                                                                : t_cbuf_std140_ubo;
             const Id id = OpVariable(type, spv::StorageClass::Uniform);
@@ -412,12 +759,11 @@ private:
             Decorate(id, spv::Decoration::DescriptorSet, DESCRIPTOR_SET);
             constant_buffers.emplace(index, id);
         }
+        return binding;
     }
 
-    void DeclareGlobalBuffers() {
-        u32 binding = global_buffers_base_binding;
-        for (const auto& entry : ir.GetGlobalMemory()) {
-            const auto [base, usage] = entry;
+    u32 DeclareGlobalBuffers(u32 binding) {
+        for (const auto& [base, usage] : ir.GetGlobalMemory()) {
             const Id id = OpVariable(t_gmem_ssbo, spv::StorageClass::StorageBuffer);
             AddGlobalVariable(
                 Name(id, fmt::format("gmem_{}_{}", base.cbuf_index, base.cbuf_offset)));
@@ -426,89 +772,187 @@ private:
             Decorate(id, spv::Decoration::DescriptorSet, DESCRIPTOR_SET);
             global_buffers.emplace(base, id);
         }
+        return binding;
+    }
+
+    u32 DeclareTexelBuffers(u32 binding) {
+        for (const auto& sampler : ir.GetSamplers()) {
+            if (!sampler.IsBuffer()) {
+                continue;
+            }
+            ASSERT(!sampler.IsArray());
+            ASSERT(!sampler.IsShadow());
+
+            constexpr auto dim = spv::Dim::Buffer;
+            constexpr int depth = 0;
+            constexpr int arrayed = 0;
+            constexpr bool ms = false;
+            constexpr int sampled = 1;
+            constexpr auto format = spv::ImageFormat::Unknown;
+            const Id image_type = TypeImage(t_float, dim, depth, arrayed, ms, sampled, format);
+            const Id pointer_type = TypePointer(spv::StorageClass::UniformConstant, image_type);
+            const Id id = OpVariable(pointer_type, spv::StorageClass::UniformConstant);
+            AddGlobalVariable(Name(id, fmt::format("sampler_{}", sampler.GetIndex())));
+            Decorate(id, spv::Decoration::Binding, binding++);
+            Decorate(id, spv::Decoration::DescriptorSet, DESCRIPTOR_SET);
+
+            texel_buffers.emplace(sampler.GetIndex(), TexelBuffer{image_type, id});
+        }
+        return binding;
     }
 
-    void DeclareSamplers() {
-        u32 binding = samplers_base_binding;
+    u32 DeclareSamplers(u32 binding) {
         for (const auto& sampler : ir.GetSamplers()) {
+            if (sampler.IsBuffer()) {
+                continue;
+            }
             const auto dim = GetSamplerDim(sampler);
             const int depth = sampler.IsShadow() ? 1 : 0;
             const int arrayed = sampler.IsArray() ? 1 : 0;
-            // TODO(Rodrigo): Sampled 1 indicates that the image will be used with a sampler. When
-            // SULD and SUST instructions are implemented, replace this value.
-            const int sampled = 1;
-            const Id image_type =
-                TypeImage(t_float, dim, depth, arrayed, false, sampled, spv::ImageFormat::Unknown);
+            constexpr bool ms = false;
+            constexpr int sampled = 1;
+            constexpr auto format = spv::ImageFormat::Unknown;
+            const Id image_type = TypeImage(t_float, dim, depth, arrayed, ms, sampled, format);
             const Id sampled_image_type = TypeSampledImage(image_type);
             const Id pointer_type =
                 TypePointer(spv::StorageClass::UniformConstant, sampled_image_type);
             const Id id = OpVariable(pointer_type, spv::StorageClass::UniformConstant);
             AddGlobalVariable(Name(id, fmt::format("sampler_{}", sampler.GetIndex())));
+            Decorate(id, spv::Decoration::Binding, binding++);
+            Decorate(id, spv::Decoration::DescriptorSet, DESCRIPTOR_SET);
 
-            sampler_images.insert(
-                {static_cast<u32>(sampler.GetIndex()), {image_type, sampled_image_type, id}});
+            sampled_images.emplace(sampler.GetIndex(),
+                                   SampledImage{image_type, sampled_image_type, id});
+        }
+        return binding;
+    }
+
+    u32 DeclareImages(u32 binding) {
+        for (const auto& image : ir.GetImages()) {
+            const auto [dim, arrayed] = GetImageDim(image);
+            constexpr int depth = 0;
+            constexpr bool ms = false;
+            constexpr int sampled = 2; // This won't be accessed with a sampler
+            constexpr auto format = spv::ImageFormat::Unknown;
+            const Id image_type = TypeImage(t_uint, dim, depth, arrayed, ms, sampled, format, {});
+            const Id pointer_type = TypePointer(spv::StorageClass::UniformConstant, image_type);
+            const Id id = OpVariable(pointer_type, spv::StorageClass::UniformConstant);
+            AddGlobalVariable(Name(id, fmt::format("image_{}", image.GetIndex())));
 
             Decorate(id, spv::Decoration::Binding, binding++);
             Decorate(id, spv::Decoration::DescriptorSet, DESCRIPTOR_SET);
+            if (image.IsRead() && !image.IsWritten()) {
+                Decorate(id, spv::Decoration::NonWritable);
+            } else if (image.IsWritten() && !image.IsRead()) {
+                Decorate(id, spv::Decoration::NonReadable);
+            }
+
+            images.emplace(static_cast<u32>(image.GetIndex()), StorageImage{image_type, id});
         }
+        return binding;
+    }
+
+    bool IsInputAttributeArray() const {
+        return stage == ShaderType::TesselationControl || stage == ShaderType::TesselationEval ||
+               stage == ShaderType::Geometry;
     }
 
-    void DeclareVertexRedeclarations() {
-        vertex_index = DeclareBuiltIn(spv::BuiltIn::VertexIndex, spv::StorageClass::Input,
-                                      t_in_uint, "vertex_index");
-        instance_index = DeclareBuiltIn(spv::BuiltIn::InstanceIndex, spv::StorageClass::Input,
-                                        t_in_uint, "instance_index");
+    bool IsOutputAttributeArray() const {
+        return stage == ShaderType::TesselationControl;
+    }
 
-        bool is_clip_distances_declared = false;
-        for (const auto index : ir.GetOutputAttributes()) {
-            if (index == Attribute::Index::ClipDistances0123 ||
-                index == Attribute::Index::ClipDistances4567) {
-                is_clip_distances_declared = true;
-            }
+    u32 GetNumInputVertices() const {
+        switch (stage) {
+        case ShaderType::Geometry:
+            return GetNumPrimitiveTopologyVertices(specialization.primitive_topology);
+        case ShaderType::TesselationControl:
+        case ShaderType::TesselationEval:
+            return NumInputPatches;
+        default:
+            UNREACHABLE();
+            return 1;
         }
+    }
 
-        std::vector<Id> members;
-        members.push_back(t_float4);
-        if (ir.UsesPointSize()) {
-            members.push_back(t_float);
-        }
-        if (is_clip_distances_declared) {
-            members.push_back(TypeArray(t_float, Constant(t_uint, 8)));
-        }
-
-        const Id gl_per_vertex_struct = Name(TypeStruct(members), "PerVertex");
-        Decorate(gl_per_vertex_struct, spv::Decoration::Block);
-
-        u32 declaration_index = 0;
-        const auto MemberDecorateBuiltIn = [&](spv::BuiltIn builtin, std::string name,
-                                               bool condition) {
-            if (!condition)
-                return u32{};
-            MemberName(gl_per_vertex_struct, declaration_index, name);
-            MemberDecorate(gl_per_vertex_struct, declaration_index, spv::Decoration::BuiltIn,
-                           static_cast<u32>(builtin));
-            return declaration_index++;
+    u32 GetNumOutputVertices() const {
+        switch (stage) {
+        case ShaderType::TesselationControl:
+            return header.common2.threads_per_input_primitive;
+        default:
+            UNREACHABLE();
+            return 1;
+        }
+    }
+
+    std::tuple<Id, VertexIndices> DeclareVertexStruct() {
+        struct BuiltIn {
+            Id type;
+            spv::BuiltIn builtin;
+            const char* name;
+        };
+        std::vector<BuiltIn> members;
+        members.reserve(4);
+
+        const auto AddBuiltIn = [&](Id type, spv::BuiltIn builtin, const char* name) {
+            const auto index = static_cast<u32>(members.size());
+            members.push_back(BuiltIn{type, builtin, name});
+            return index;
         };
 
-        position_index = MemberDecorateBuiltIn(spv::BuiltIn::Position, "position", true);
-        point_size_index =
-            MemberDecorateBuiltIn(spv::BuiltIn::PointSize, "point_size", ir.UsesPointSize());
-        clip_distances_index = MemberDecorateBuiltIn(spv::BuiltIn::ClipDistance, "clip_distances",
-                                                     is_clip_distances_declared);
+        VertexIndices indices;
+        indices.position = AddBuiltIn(t_float4, spv::BuiltIn::Position, "position");
+
+        if (ir.UsesViewportIndex()) {
+            if (stage != ShaderType::Vertex || device.IsExtShaderViewportIndexLayerSupported()) {
+                indices.viewport = AddBuiltIn(t_int, spv::BuiltIn::ViewportIndex, "viewport_index");
+            } else {
+                LOG_ERROR(Render_Vulkan,
+                          "Shader requires ViewportIndex but it's not supported on this "
+                          "stage with this device.");
+            }
+        }
+
+        if (ir.UsesPointSize() || specialization.point_size) {
+            indices.point_size = AddBuiltIn(t_float, spv::BuiltIn::PointSize, "point_size");
+        }
+
+        const auto& output_attributes = ir.GetOutputAttributes();
+        const bool declare_clip_distances =
+            std::any_of(output_attributes.begin(), output_attributes.end(), [](const auto& index) {
+                return index == Attribute::Index::ClipDistances0123 ||
+                       index == Attribute::Index::ClipDistances4567;
+            });
+        if (declare_clip_distances) {
+            indices.clip_distances = AddBuiltIn(TypeArray(t_float, Constant(t_uint, 8)),
+                                                spv::BuiltIn::ClipDistance, "clip_distances");
+        }
+
+        std::vector<Id> member_types;
+        member_types.reserve(members.size());
+        for (std::size_t i = 0; i < members.size(); ++i) {
+            member_types.push_back(members[i].type);
+        }
+        const Id per_vertex_struct = Name(TypeStruct(member_types), "PerVertex");
+        Decorate(per_vertex_struct, spv::Decoration::Block);
+
+        for (std::size_t index = 0; index < members.size(); ++index) {
+            const auto& member = members[index];
+            MemberName(per_vertex_struct, static_cast<u32>(index), member.name);
+            MemberDecorate(per_vertex_struct, static_cast<u32>(index), spv::Decoration::BuiltIn,
+                           static_cast<u32>(member.builtin));
+        }
 
-        const Id type_pointer = TypePointer(spv::StorageClass::Output, gl_per_vertex_struct);
-        per_vertex = OpVariable(type_pointer, spv::StorageClass::Output);
-        AddGlobalVariable(Name(per_vertex, "per_vertex"));
-        interfaces.push_back(per_vertex);
+        return {per_vertex_struct, indices};
     }
 
     void VisitBasicBlock(const NodeBlock& bb) {
         for (const auto& node : bb) {
-            static_cast<void>(Visit(node));
+            [[maybe_unused]] const Type type = Visit(node).type;
+            ASSERT(type == Type::Void);
         }
     }
 
-    Id Visit(const Node& node) {
+    Expression Visit(const Node& node) {
         if (const auto operation = std::get_if<OperationNode>(&*node)) {
             const auto operation_index = static_cast<std::size_t>(operation->GetCode());
             const auto decompiler = operation_decompilers[operation_index];
@@ -516,18 +960,21 @@ private:
                 UNREACHABLE_MSG("Operation decompiler {} not defined", operation_index);
             }
             return (this->*decompiler)(*operation);
+        }
 
-        } else if (const auto gpr = std::get_if<GprNode>(&*node)) {
+        if (const auto gpr = std::get_if<GprNode>(&*node)) {
             const u32 index = gpr->GetIndex();
             if (index == Register::ZeroIndex) {
-                return Constant(t_float, 0.0f);
+                return {v_float_zero, Type::Float};
             }
-            return Emit(OpLoad(t_float, registers.at(index)));
+            return {OpLoad(t_float, registers.at(index)), Type::Float};
+        }
 
-        } else if (const auto immediate = std::get_if<ImmediateNode>(&*node)) {
-            return BitcastTo<Type::Float>(Constant(t_uint, immediate->GetValue()));
+        if (const auto immediate = std::get_if<ImmediateNode>(&*node)) {
+            return {Constant(t_uint, immediate->GetValue()), Type::Uint};
+        }
 
-        } else if (const auto predicate = std::get_if<PredicateNode>(&*node)) {
+        if (const auto predicate = std::get_if<PredicateNode>(&*node)) {
             const auto value = [&]() -> Id {
                 switch (const auto index = predicate->GetIndex(); index) {
                 case Tegra::Shader::Pred::UnusedIndex:
@@ -535,74 +982,107 @@ private:
                 case Tegra::Shader::Pred::NeverExecute:
                     return v_false;
                 default:
-                    return Emit(OpLoad(t_bool, predicates.at(index)));
+                    return OpLoad(t_bool, predicates.at(index));
                 }
             }();
             if (predicate->IsNegated()) {
-                return Emit(OpLogicalNot(t_bool, value));
+                return {OpLogicalNot(t_bool, value), Type::Bool};
             }
-            return value;
+            return {value, Type::Bool};
+        }
 
-        } else if (const auto abuf = std::get_if<AbufNode>(&*node)) {
+        if (const auto abuf = std::get_if<AbufNode>(&*node)) {
             const auto attribute = abuf->GetIndex();
-            const auto element = abuf->GetElement();
+            const u32 element = abuf->GetElement();
+            const auto& buffer = abuf->GetBuffer();
+
+            const auto ArrayPass = [&](Id pointer_type, Id composite, std::vector<u32> indices) {
+                std::vector<Id> members;
+                members.reserve(std::size(indices) + 1);
+
+                if (buffer && IsInputAttributeArray()) {
+                    members.push_back(AsUint(Visit(buffer)));
+                }
+                for (const u32 index : indices) {
+                    members.push_back(Constant(t_uint, index));
+                }
+                return OpAccessChain(pointer_type, composite, members);
+            };
 
             switch (attribute) {
-            case Attribute::Index::Position:
-                if (stage != ShaderType::Fragment) {
-                    UNIMPLEMENTED();
-                    break;
-                } else {
+            case Attribute::Index::Position: {
+                if (stage == ShaderType::Fragment) {
                     if (element == 3) {
-                        return Constant(t_float, 1.0f);
+                        return {Constant(t_float, 1.0f), Type::Float};
                     }
-                    return Emit(OpLoad(t_float, AccessElement(t_in_float, frag_coord, element)));
+                    return {OpLoad(t_float, AccessElement(t_in_float, frag_coord, element)),
+                            Type::Float};
+                }
+                const auto elements = {in_indices.position.value(), element};
+                return {OpLoad(t_float, ArrayPass(t_in_float, in_vertex, elements)), Type::Float};
+            }
+            case Attribute::Index::PointCoord: {
+                switch (element) {
+                case 0:
+                case 1:
+                    return {OpCompositeExtract(t_float, OpLoad(t_float2, point_coord), element),
+                            Type::Float};
                 }
+                UNIMPLEMENTED_MSG("Unimplemented point coord element={}", element);
+                return {v_float_zero, Type::Float};
+            }
             case Attribute::Index::TessCoordInstanceIDVertexID:
                 // TODO(Subv): Find out what the values are for the first two elements when inside a
                 // vertex shader, and what's the value of the fourth element when inside a Tess Eval
                 // shader.
-                ASSERT(stage == ShaderType::Vertex);
                 switch (element) {
+                case 0:
+                case 1:
+                    return {OpLoad(t_float, AccessElement(t_in_float, tess_coord, element)),
+                            Type::Float};
                 case 2:
-                    return BitcastFrom<Type::Uint>(Emit(OpLoad(t_uint, instance_index)));
+                    return {OpLoad(t_uint, instance_index), Type::Uint};
                 case 3:
-                    return BitcastFrom<Type::Uint>(Emit(OpLoad(t_uint, vertex_index)));
+                    return {OpLoad(t_uint, vertex_index), Type::Uint};
                 }
                 UNIMPLEMENTED_MSG("Unmanaged TessCoordInstanceIDVertexID element={}", element);
-                return Constant(t_float, 0);
+                return {Constant(t_uint, 0U), Type::Uint};
             case Attribute::Index::FrontFacing:
                 // TODO(Subv): Find out what the values are for the other elements.
                 ASSERT(stage == ShaderType::Fragment);
                 if (element == 3) {
-                    const Id is_front_facing = Emit(OpLoad(t_bool, front_facing));
-                    const Id true_value =
-                        BitcastTo<Type::Float>(Constant(t_int, static_cast<s32>(-1)));
-                    const Id false_value = BitcastTo<Type::Float>(Constant(t_int, 0));
-                    return Emit(OpSelect(t_float, is_front_facing, true_value, false_value));
+                    const Id is_front_facing = OpLoad(t_bool, front_facing);
+                    const Id true_value = Constant(t_int, static_cast<s32>(-1));
+                    const Id false_value = Constant(t_int, 0);
+                    return {OpSelect(t_int, is_front_facing, true_value, false_value), Type::Int};
                 }
                 UNIMPLEMENTED_MSG("Unmanaged FrontFacing element={}", element);
-                return Constant(t_float, 0.0f);
+                return {v_float_zero, Type::Float};
             default:
                 if (IsGenericAttribute(attribute)) {
-                    const Id pointer =
-                        AccessElement(t_in_float, input_attributes.at(attribute), element);
-                    return Emit(OpLoad(t_float, pointer));
+                    const u32 location = GetGenericAttributeLocation(attribute);
+                    const auto type_descriptor = GetAttributeType(location);
+                    const Type type = type_descriptor.type;
+                    const Id attribute_id = input_attributes.at(attribute);
+                    const Id pointer = ArrayPass(type_descriptor.scalar, attribute_id, {element});
+                    return {OpLoad(GetTypeDefinition(type), pointer), type};
                 }
                 break;
             }
             UNIMPLEMENTED_MSG("Unhandled input attribute: {}", static_cast<u32>(attribute));
+            return {v_float_zero, Type::Float};
+        }
 
-        } else if (const auto cbuf = std::get_if<CbufNode>(&*node)) {
+        if (const auto cbuf = std::get_if<CbufNode>(&*node)) {
             const Node& offset = cbuf->GetOffset();
             const Id buffer_id = constant_buffers.at(cbuf->GetIndex());
 
             Id pointer{};
             if (device.IsKhrUniformBufferStandardLayoutSupported()) {
-                const Id buffer_offset = Emit(OpShiftRightLogical(
-                    t_uint, BitcastTo<Type::Uint>(Visit(offset)), Constant(t_uint, 2u)));
-                pointer = Emit(
-                    OpAccessChain(t_cbuf_float, buffer_id, Constant(t_uint, 0u), buffer_offset));
+                const Id buffer_offset =
+                    OpShiftRightLogical(t_uint, AsUint(Visit(offset)), Constant(t_uint, 2U));
+                pointer =
+                    OpAccessChain(t_cbuf_float, buffer_id, Constant(t_uint, 0U), buffer_offset);
             } else {
                 Id buffer_index{};
                 Id buffer_element{};
@@ -614,53 +1094,76 @@ private:
                     buffer_element = Constant(t_uint, (offset_imm / 4) % 4);
                 } else if (std::holds_alternative<OperationNode>(*offset)) {
                     // Indirect access
-                    const Id offset_id = BitcastTo<Type::Uint>(Visit(offset));
-                    const Id unsafe_offset = Emit(OpUDiv(t_uint, offset_id, Constant(t_uint, 4)));
-                    const Id final_offset = Emit(OpUMod(
-                        t_uint, unsafe_offset, Constant(t_uint, MAX_CONSTBUFFER_ELEMENTS - 1)));
-                    buffer_index = Emit(OpUDiv(t_uint, final_offset, Constant(t_uint, 4)));
-                    buffer_element = Emit(OpUMod(t_uint, final_offset, Constant(t_uint, 4)));
+                    const Id offset_id = AsUint(Visit(offset));
+                    const Id unsafe_offset = OpUDiv(t_uint, offset_id, Constant(t_uint, 4));
+                    const Id final_offset =
+                        OpUMod(t_uint, unsafe_offset, Constant(t_uint, MaxConstBufferElements - 1));
+                    buffer_index = OpUDiv(t_uint, final_offset, Constant(t_uint, 4));
+                    buffer_element = OpUMod(t_uint, final_offset, Constant(t_uint, 4));
                 } else {
                     UNREACHABLE_MSG("Unmanaged offset node type");
                 }
-                pointer = Emit(OpAccessChain(t_cbuf_float, buffer_id, Constant(t_uint, 0),
-                                             buffer_index, buffer_element));
+                pointer = OpAccessChain(t_cbuf_float, buffer_id, Constant(t_uint, 0), buffer_index,
+                                        buffer_element);
             }
-            return Emit(OpLoad(t_float, pointer));
+            return {OpLoad(t_float, pointer), Type::Float};
+        }
 
-        } else if (const auto gmem = std::get_if<GmemNode>(&*node)) {
+        if (const auto gmem = std::get_if<GmemNode>(&*node)) {
             const Id gmem_buffer = global_buffers.at(gmem->GetDescriptor());
-            const Id real = BitcastTo<Type::Uint>(Visit(gmem->GetRealAddress()));
-            const Id base = BitcastTo<Type::Uint>(Visit(gmem->GetBaseAddress()));
+            const Id real = AsUint(Visit(gmem->GetRealAddress()));
+            const Id base = AsUint(Visit(gmem->GetBaseAddress()));
+
+            Id offset = OpISub(t_uint, real, base);
+            offset = OpUDiv(t_uint, offset, Constant(t_uint, 4U));
+            return {OpLoad(t_float,
+                           OpAccessChain(t_gmem_float, gmem_buffer, Constant(t_uint, 0U), offset)),
+                    Type::Float};
+        }
 
-            Id offset = Emit(OpISub(t_uint, real, base));
-            offset = Emit(OpUDiv(t_uint, offset, Constant(t_uint, 4u)));
-            return Emit(OpLoad(t_float, Emit(OpAccessChain(t_gmem_float, gmem_buffer,
-                                                           Constant(t_uint, 0u), offset))));
+        if (const auto lmem = std::get_if<LmemNode>(&*node)) {
+            Id address = AsUint(Visit(lmem->GetAddress()));
+            address = OpShiftRightLogical(t_uint, address, Constant(t_uint, 2U));
+            const Id pointer = OpAccessChain(t_prv_float, local_memory, address);
+            return {OpLoad(t_float, pointer), Type::Float};
+        }
+
+        if (const auto smem = std::get_if<SmemNode>(&*node)) {
+            Id address = AsUint(Visit(smem->GetAddress()));
+            address = OpShiftRightLogical(t_uint, address, Constant(t_uint, 2U));
+            const Id pointer = OpAccessChain(t_smem_uint, shared_memory, address);
+            return {OpLoad(t_uint, pointer), Type::Uint};
+        }
 
-        } else if (const auto conditional = std::get_if<ConditionalNode>(&*node)) {
+        if (const auto internal_flag = std::get_if<InternalFlagNode>(&*node)) {
+            const Id flag = internal_flags.at(static_cast<std::size_t>(internal_flag->GetFlag()));
+            return {OpLoad(t_bool, flag), Type::Bool};
+        }
+
+        if (const auto conditional = std::get_if<ConditionalNode>(&*node)) {
             // It's invalid to call conditional on nested nodes, use an operation instead
             const Id true_label = OpLabel();
             const Id skip_label = OpLabel();
-            const Id condition = Visit(conditional->GetCondition());
-            Emit(OpSelectionMerge(skip_label, spv::SelectionControlMask::MaskNone));
-            Emit(OpBranchConditional(condition, true_label, skip_label));
-            Emit(true_label);
+            const Id condition = AsBool(Visit(conditional->GetCondition()));
+            OpSelectionMerge(skip_label, spv::SelectionControlMask::MaskNone);
+            OpBranchConditional(condition, true_label, skip_label);
+            AddLabel(true_label);
 
-            ++conditional_nest_count;
+            conditional_branch_set = true;
+            inside_branch = false;
             VisitBasicBlock(conditional->GetCode());
-            --conditional_nest_count;
-
-            if (inside_branch == 0) {
-                Emit(OpBranch(skip_label));
+            conditional_branch_set = false;
+            if (!inside_branch) {
+                OpBranch(skip_label);
             } else {
-                inside_branch--;
+                inside_branch = false;
             }
-            Emit(skip_label);
+            AddLabel(skip_label);
             return {};
+        }
 
-        } else if (const auto comment = std::get_if<CommentNode>(&*node)) {
-            Name(Emit(OpUndef(t_void)), comment->GetText());
+        if (const auto comment = std::get_if<CommentNode>(&*node)) {
+            Name(OpUndef(t_void), comment->GetText());
             return {};
         }
 
@@ -669,94 +1172,126 @@ private:
     }
 
     template <Id (Module::*func)(Id, Id), Type result_type, Type type_a = result_type>
-    Id Unary(Operation operation) {
+    Expression Unary(Operation operation) {
         const Id type_def = GetTypeDefinition(result_type);
-        const Id op_a = VisitOperand<type_a>(operation, 0);
+        const Id op_a = As(Visit(operation[0]), type_a);
 
-        const Id value = BitcastFrom<result_type>(Emit((this->*func)(type_def, op_a)));
+        const Id value = (this->*func)(type_def, op_a);
         if (IsPrecise(operation)) {
             Decorate(value, spv::Decoration::NoContraction);
         }
-        return value;
+        return {value, result_type};
     }
 
     template <Id (Module::*func)(Id, Id, Id), Type result_type, Type type_a = result_type,
               Type type_b = type_a>
-    Id Binary(Operation operation) {
+    Expression Binary(Operation operation) {
         const Id type_def = GetTypeDefinition(result_type);
-        const Id op_a = VisitOperand<type_a>(operation, 0);
-        const Id op_b = VisitOperand<type_b>(operation, 1);
+        const Id op_a = As(Visit(operation[0]), type_a);
+        const Id op_b = As(Visit(operation[1]), type_b);
 
-        const Id value = BitcastFrom<result_type>(Emit((this->*func)(type_def, op_a, op_b)));
+        const Id value = (this->*func)(type_def, op_a, op_b);
         if (IsPrecise(operation)) {
             Decorate(value, spv::Decoration::NoContraction);
         }
-        return value;
+        return {value, result_type};
     }
 
     template <Id (Module::*func)(Id, Id, Id, Id), Type result_type, Type type_a = result_type,
               Type type_b = type_a, Type type_c = type_b>
-    Id Ternary(Operation operation) {
+    Expression Ternary(Operation operation) {
         const Id type_def = GetTypeDefinition(result_type);
-        const Id op_a = VisitOperand<type_a>(operation, 0);
-        const Id op_b = VisitOperand<type_b>(operation, 1);
-        const Id op_c = VisitOperand<type_c>(operation, 2);
+        const Id op_a = As(Visit(operation[0]), type_a);
+        const Id op_b = As(Visit(operation[1]), type_b);
+        const Id op_c = As(Visit(operation[2]), type_c);
 
-        const Id value = BitcastFrom<result_type>(Emit((this->*func)(type_def, op_a, op_b, op_c)));
+        const Id value = (this->*func)(type_def, op_a, op_b, op_c);
         if (IsPrecise(operation)) {
             Decorate(value, spv::Decoration::NoContraction);
         }
-        return value;
+        return {value, result_type};
     }
 
     template <Id (Module::*func)(Id, Id, Id, Id, Id), Type result_type, Type type_a = result_type,
               Type type_b = type_a, Type type_c = type_b, Type type_d = type_c>
-    Id Quaternary(Operation operation) {
+    Expression Quaternary(Operation operation) {
         const Id type_def = GetTypeDefinition(result_type);
-        const Id op_a = VisitOperand<type_a>(operation, 0);
-        const Id op_b = VisitOperand<type_b>(operation, 1);
-        const Id op_c = VisitOperand<type_c>(operation, 2);
-        const Id op_d = VisitOperand<type_d>(operation, 3);
+        const Id op_a = As(Visit(operation[0]), type_a);
+        const Id op_b = As(Visit(operation[1]), type_b);
+        const Id op_c = As(Visit(operation[2]), type_c);
+        const Id op_d = As(Visit(operation[3]), type_d);
 
-        const Id value =
-            BitcastFrom<result_type>(Emit((this->*func)(type_def, op_a, op_b, op_c, op_d)));
+        const Id value = (this->*func)(type_def, op_a, op_b, op_c, op_d);
         if (IsPrecise(operation)) {
             Decorate(value, spv::Decoration::NoContraction);
         }
-        return value;
+        return {value, result_type};
     }
 
-    Id Assign(Operation operation) {
+    Expression Assign(Operation operation) {
         const Node& dest = operation[0];
         const Node& src = operation[1];
 
-        Id target{};
+        Expression target{};
         if (const auto gpr = std::get_if<GprNode>(&*dest)) {
             if (gpr->GetIndex() == Register::ZeroIndex) {
                 // Writing to Register::ZeroIndex is a no op
                 return {};
             }
-            target = registers.at(gpr->GetIndex());
+            target = {registers.at(gpr->GetIndex()), Type::Float};
 
         } else if (const auto abuf = std::get_if<AbufNode>(&*dest)) {
-            target = [&]() -> Id {
+            const auto& buffer = abuf->GetBuffer();
+            const auto ArrayPass = [&](Id pointer_type, Id composite, std::vector<u32> indices) {
+                std::vector<Id> members;
+                members.reserve(std::size(indices) + 1);
+
+                if (buffer && IsOutputAttributeArray()) {
+                    members.push_back(AsUint(Visit(buffer)));
+                }
+                for (const u32 index : indices) {
+                    members.push_back(Constant(t_uint, index));
+                }
+                return OpAccessChain(pointer_type, composite, members);
+            };
+
+            target = [&]() -> Expression {
+                const u32 element = abuf->GetElement();
                 switch (const auto attribute = abuf->GetIndex(); attribute) {
-                case Attribute::Index::Position:
-                    return AccessElement(t_out_float, per_vertex, position_index,
-                                         abuf->GetElement());
+                case Attribute::Index::Position: {
+                    const u32 index = out_indices.position.value();
+                    return {ArrayPass(t_out_float, out_vertex, {index, element}), Type::Float};
+                }
                 case Attribute::Index::LayerViewportPointSize:
-                    UNIMPLEMENTED_IF(abuf->GetElement() != 3);
-                    return AccessElement(t_out_float, per_vertex, point_size_index);
-                case Attribute::Index::ClipDistances0123:
-                    return AccessElement(t_out_float, per_vertex, clip_distances_index,
-                                         abuf->GetElement());
-                case Attribute::Index::ClipDistances4567:
-                    return AccessElement(t_out_float, per_vertex, clip_distances_index,
-                                         abuf->GetElement() + 4);
+                    switch (element) {
+                    case 2: {
+                        if (!out_indices.viewport) {
+                            return {};
+                        }
+                        const u32 index = out_indices.viewport.value();
+                        return {AccessElement(t_out_int, out_vertex, index), Type::Int};
+                    }
+                    case 3: {
+                        const auto index = out_indices.point_size.value();
+                        return {AccessElement(t_out_float, out_vertex, index), Type::Float};
+                    }
+                    default:
+                        UNIMPLEMENTED_MSG("LayerViewportPoint element={}", abuf->GetElement());
+                        return {};
+                    }
+                case Attribute::Index::ClipDistances0123: {
+                    const u32 index = out_indices.clip_distances.value();
+                    return {AccessElement(t_out_float, out_vertex, index, element), Type::Float};
+                }
+                case Attribute::Index::ClipDistances4567: {
+                    const u32 index = out_indices.clip_distances.value();
+                    return {AccessElement(t_out_float, out_vertex, index, element + 4),
+                            Type::Float};
+                }
                 default:
                     if (IsGenericAttribute(attribute)) {
-                        return AccessElement(t_out_float, output_attributes.at(attribute),
-                                             abuf->GetElement());
+                        const Id composite = output_attributes.at(attribute);
+                        return {ArrayPass(t_out_float, composite, {element}), Type::Float};
                     }
                     UNIMPLEMENTED_MSG("Unhandled output attribute: {}",
                                       static_cast<u32>(attribute));
@@ -764,72 +1299,154 @@ private:
                 }
             }();
 
+        } else if (const auto patch = std::get_if<PatchNode>(&*dest)) {
+            target = [&]() -> Expression {
+                const u32 offset = patch->GetOffset();
+                switch (offset) {
+                case 0:
+                case 1:
+                case 2:
+                case 3:
+                    return {AccessElement(t_out_float, tess_level_outer, offset % 4), Type::Float};
+                case 4:
+                case 5:
+                    return {AccessElement(t_out_float, tess_level_inner, offset % 4), Type::Float};
+                }
+                UNIMPLEMENTED_MSG("Unhandled patch output offset: {}", offset);
+                return {};
+            }();
+
         } else if (const auto lmem = std::get_if<LmemNode>(&*dest)) {
-            Id address = BitcastTo<Type::Uint>(Visit(lmem->GetAddress()));
-            address = Emit(OpUDiv(t_uint, address, Constant(t_uint, 4)));
-            target = Emit(OpAccessChain(t_prv_float, local_memory, {address}));
+            Id address = AsUint(Visit(lmem->GetAddress()));
+            address = OpUDiv(t_uint, address, Constant(t_uint, 4));
+            target = {OpAccessChain(t_prv_float, local_memory, address), Type::Float};
+
+        } else if (const auto smem = std::get_if<SmemNode>(&*dest)) {
+            ASSERT(stage == ShaderType::Compute);
+            Id address = AsUint(Visit(smem->GetAddress()));
+            address = OpShiftRightLogical(t_uint, address, Constant(t_uint, 2U));
+            target = {OpAccessChain(t_smem_uint, shared_memory, address), Type::Uint};
+
+        } else if (const auto gmem = std::get_if<GmemNode>(&*dest)) {
+            const Id real = AsUint(Visit(gmem->GetRealAddress()));
+            const Id base = AsUint(Visit(gmem->GetBaseAddress()));
+            const Id diff = OpISub(t_uint, real, base);
+            const Id offset = OpShiftRightLogical(t_uint, diff, Constant(t_uint, 2));
+
+            const Id gmem_buffer = global_buffers.at(gmem->GetDescriptor());
+            target = {OpAccessChain(t_gmem_float, gmem_buffer, Constant(t_uint, 0), offset),
+                      Type::Float};
+
+        } else {
+            UNIMPLEMENTED();
         }
 
-        Emit(OpStore(target, Visit(src)));
+        OpStore(target.id, As(Visit(src), target.type));
         return {};
     }
 
-    Id FCastHalf0(Operation operation) {
-        UNIMPLEMENTED();
-        return {};
+    template <u32 offset>
+    Expression FCastHalf(Operation operation) {
+        const Id value = AsHalfFloat(Visit(operation[0]));
+        return {GetFloatFromHalfScalar(OpCompositeExtract(t_scalar_half, value, offset)),
+                Type::Float};
     }
 
-    Id FCastHalf1(Operation operation) {
-        UNIMPLEMENTED();
-        return {};
-    }
+    Expression FSwizzleAdd(Operation operation) {
+        const Id minus = Constant(t_float, -1.0f);
+        const Id plus = v_float_one;
+        const Id zero = v_float_zero;
+        const Id lut_a = ConstantComposite(t_float4, minus, plus, minus, zero);
+        const Id lut_b = ConstantComposite(t_float4, minus, minus, plus, minus);
 
-    Id FSwizzleAdd(Operation operation) {
-        UNIMPLEMENTED();
-        return {};
-    }
+        Id mask = OpLoad(t_uint, thread_id);
+        mask = OpBitwiseAnd(t_uint, mask, Constant(t_uint, 3));
+        mask = OpShiftLeftLogical(t_uint, mask, Constant(t_uint, 1));
+        mask = OpShiftRightLogical(t_uint, AsUint(Visit(operation[2])), mask);
+        mask = OpBitwiseAnd(t_uint, mask, Constant(t_uint, 3));
 
-    Id HNegate(Operation operation) {
-        UNIMPLEMENTED();
-        return {};
+        const Id modifier_a = OpVectorExtractDynamic(t_float, lut_a, mask);
+        const Id modifier_b = OpVectorExtractDynamic(t_float, lut_b, mask);
+
+        const Id op_a = OpFMul(t_float, AsFloat(Visit(operation[0])), modifier_a);
+        const Id op_b = OpFMul(t_float, AsFloat(Visit(operation[1])), modifier_b);
+        return {OpFAdd(t_float, op_a, op_b), Type::Float};
     }
 
-    Id HClamp(Operation operation) {
-        UNIMPLEMENTED();
-        return {};
+    Expression HNegate(Operation operation) {
+        const bool is_f16 = device.IsFloat16Supported();
+        const Id minus_one = Constant(t_scalar_half, is_f16 ? 0xbc00 : 0xbf800000);
+        const Id one = Constant(t_scalar_half, is_f16 ? 0x3c00 : 0x3f800000);
+        const auto GetNegate = [&](std::size_t index) {
+            return OpSelect(t_scalar_half, AsBool(Visit(operation[index])), minus_one, one);
+        };
+        const Id negation = OpCompositeConstruct(t_half, GetNegate(1), GetNegate(2));
+        return {OpFMul(t_half, AsHalfFloat(Visit(operation[0])), negation), Type::HalfFloat};
     }
 
-    Id HCastFloat(Operation operation) {
-        UNIMPLEMENTED();
-        return {};
+    Expression HClamp(Operation operation) {
+        const auto Pack = [&](std::size_t index) {
+            const Id scalar = GetHalfScalarFromFloat(AsFloat(Visit(operation[index])));
+            return OpCompositeConstruct(t_half, scalar, scalar);
+        };
+        const Id value = AsHalfFloat(Visit(operation[0]));
+        const Id min = Pack(1);
+        const Id max = Pack(2);
+
+        const Id clamped = OpFClamp(t_half, value, min, max);
+        if (IsPrecise(operation)) {
+            Decorate(clamped, spv::Decoration::NoContraction);
+        }
+        return {clamped, Type::HalfFloat};
     }
 
-    Id HUnpack(Operation operation) {
-        UNIMPLEMENTED();
-        return {};
+    Expression HCastFloat(Operation operation) {
+        const Id value = GetHalfScalarFromFloat(AsFloat(Visit(operation[0])));
+        return {OpCompositeConstruct(t_half, value, Constant(t_scalar_half, 0)), Type::HalfFloat};
     }
 
-    Id HMergeF32(Operation operation) {
-        UNIMPLEMENTED();
-        return {};
+    Expression HUnpack(Operation operation) {
+        Expression operand = Visit(operation[0]);
+        const auto type = std::get<Tegra::Shader::HalfType>(operation.GetMeta());
+        if (type == Tegra::Shader::HalfType::H0_H1) {
+            return operand;
+        }
+        const auto value = [&] {
+            switch (std::get<Tegra::Shader::HalfType>(operation.GetMeta())) {
+            case Tegra::Shader::HalfType::F32:
+                return GetHalfScalarFromFloat(AsFloat(operand));
+            case Tegra::Shader::HalfType::H0_H0:
+                return OpCompositeExtract(t_scalar_half, AsHalfFloat(operand), 0);
+            case Tegra::Shader::HalfType::H1_H1:
+                return OpCompositeExtract(t_scalar_half, AsHalfFloat(operand), 1);
+            default:
+                UNREACHABLE();
+                return ConstantNull(t_half);
+            }
+        }();
+        return {OpCompositeConstruct(t_half, value, value), Type::HalfFloat};
     }
 
-    Id HMergeH0(Operation operation) {
-        UNIMPLEMENTED();
-        return {};
+    Expression HMergeF32(Operation operation) {
+        const Id value = AsHalfFloat(Visit(operation[0]));
+        return {GetFloatFromHalfScalar(OpCompositeExtract(t_scalar_half, value, 0)), Type::Float};
     }
 
-    Id HMergeH1(Operation operation) {
-        UNIMPLEMENTED();
-        return {};
+    template <u32 offset>
+    Expression HMergeHN(Operation operation) {
+        const Id target = AsHalfFloat(Visit(operation[0]));
+        const Id source = AsHalfFloat(Visit(operation[1]));
+        const Id object = OpCompositeExtract(t_scalar_half, source, offset);
+        return {OpCompositeInsert(t_half, object, target, offset), Type::HalfFloat};
     }
 
-    Id HPack2(Operation operation) {
-        UNIMPLEMENTED();
-        return {};
+    Expression HPack2(Operation operation) {
+        const Id low = GetHalfScalarFromFloat(AsFloat(Visit(operation[0])));
+        const Id high = GetHalfScalarFromFloat(AsFloat(Visit(operation[1])));
+        return {OpCompositeConstruct(t_half, low, high), Type::HalfFloat};
     }
 
-    Id LogicalAssign(Operation operation) {
+    Expression LogicalAssign(Operation operation) {
         const Node& dest = operation[0];
         const Node& src = operation[1];
 
@@ -850,106 +1467,190 @@ private:
             target = internal_flags.at(static_cast<u32>(flag->GetFlag()));
         }
 
-        Emit(OpStore(target, Visit(src)));
+        OpStore(target, AsBool(Visit(src)));
         return {};
     }
 
-    Id LogicalPick2(Operation operation) {
-        UNIMPLEMENTED();
-        return {};
+    Id GetTextureSampler(Operation operation) {
+        const auto& meta = std::get<MetaTexture>(operation.GetMeta());
+        ASSERT(!meta.sampler.IsBuffer());
+
+        const auto& entry = sampled_images.at(meta.sampler.GetIndex());
+        return OpLoad(entry.sampled_image_type, entry.sampler);
     }
 
-    Id LogicalAnd2(Operation operation) {
-        UNIMPLEMENTED();
-        return {};
+    Id GetTextureImage(Operation operation) {
+        const auto& meta = std::get<MetaTexture>(operation.GetMeta());
+        const u32 index = meta.sampler.GetIndex();
+        if (meta.sampler.IsBuffer()) {
+            const auto& entry = texel_buffers.at(index);
+            return OpLoad(entry.image_type, entry.image);
+        } else {
+            const auto& entry = sampled_images.at(index);
+            return OpImage(entry.image_type, GetTextureSampler(operation));
+        }
     }
 
-    Id GetTextureSampler(Operation operation) {
-        const auto meta = std::get_if<MetaTexture>(&operation.GetMeta());
-        const auto entry = sampler_images.at(static_cast<u32>(meta->sampler.GetIndex()));
-        return Emit(OpLoad(entry.sampled_image_type, entry.sampler));
+    Id GetImage(Operation operation) {
+        const auto& meta = std::get<MetaImage>(operation.GetMeta());
+        const auto entry = images.at(meta.image.GetIndex());
+        return OpLoad(entry.image_type, entry.image);
     }
 
-    Id GetTextureImage(Operation operation) {
-        const auto meta = std::get_if<MetaTexture>(&operation.GetMeta());
-        const auto entry = sampler_images.at(static_cast<u32>(meta->sampler.GetIndex()));
-        return Emit(OpImage(entry.image_type, GetTextureSampler(operation)));
+    Id AssembleVector(const std::vector<Id>& coords, Type type) {
+        const Id coords_type = GetTypeVectorDefinitionLut(type).at(coords.size() - 1);
+        return coords.size() == 1 ? coords[0] : OpCompositeConstruct(coords_type, coords);
     }
 
-    Id GetTextureCoordinates(Operation operation) {
-        const auto meta = std::get_if<MetaTexture>(&operation.GetMeta());
+    Id GetCoordinates(Operation operation, Type type) {
         std::vector<Id> coords;
         for (std::size_t i = 0; i < operation.GetOperandsCount(); ++i) {
-            coords.push_back(Visit(operation[i]));
+            coords.push_back(As(Visit(operation[i]), type));
         }
-        if (meta->sampler.IsArray()) {
-            const Id array_integer = BitcastTo<Type::Int>(Visit(meta->array));
-            coords.push_back(Emit(OpConvertSToF(t_float, array_integer)));
+        if (const auto meta = std::get_if<MetaTexture>(&operation.GetMeta())) {
+            // Add array coordinate for textures
+            if (meta->sampler.IsArray()) {
+                Id array = AsInt(Visit(meta->array));
+                if (type == Type::Float) {
+                    array = OpConvertSToF(t_float, array);
+                }
+                coords.push_back(array);
+            }
         }
-        if (meta->sampler.IsShadow()) {
-            coords.push_back(Visit(meta->depth_compare));
+        return AssembleVector(coords, type);
+    }
+
+    Id GetOffsetCoordinates(Operation operation) {
+        const auto& meta = std::get<MetaTexture>(operation.GetMeta());
+        std::vector<Id> coords;
+        coords.reserve(meta.aoffi.size());
+        for (const auto& coord : meta.aoffi) {
+            coords.push_back(AsInt(Visit(coord)));
         }
+        return AssembleVector(coords, Type::Int);
+    }
 
-        const std::array<Id, 4> t_float_lut = {nullptr, t_float2, t_float3, t_float4};
-        return coords.size() == 1
-                   ? coords[0]
-                   : Emit(OpCompositeConstruct(t_float_lut.at(coords.size() - 1), coords));
+    std::pair<Id, Id> GetDerivatives(Operation operation) {
+        const auto& meta = std::get<MetaTexture>(operation.GetMeta());
+        const auto& derivatives = meta.derivates;
+        ASSERT(derivatives.size() % 2 == 0);
+
+        const std::size_t components = derivatives.size() / 2;
+        std::vector<Id> dx, dy;
+        dx.reserve(components);
+        dy.reserve(components);
+        for (std::size_t index = 0; index < components; ++index) {
+            dx.push_back(AsFloat(Visit(derivatives.at(index * 2 + 0))));
+            dy.push_back(AsFloat(Visit(derivatives.at(index * 2 + 1))));
+        }
+        return {AssembleVector(dx, Type::Float), AssembleVector(dy, Type::Float)};
     }
 
-    Id GetTextureElement(Operation operation, Id sample_value) {
-        const auto meta = std::get_if<MetaTexture>(&operation.GetMeta());
-        ASSERT(meta);
-        return Emit(OpCompositeExtract(t_float, sample_value, meta->element));
+    Expression GetTextureElement(Operation operation, Id sample_value, Type type) {
+        const auto& meta = std::get<MetaTexture>(operation.GetMeta());
+        const auto type_def = GetTypeDefinition(type);
+        return {OpCompositeExtract(type_def, sample_value, meta.element), type};
     }
 
-    Id Texture(Operation operation) {
-        const Id texture = Emit(OpImageSampleImplicitLod(t_float4, GetTextureSampler(operation),
-                                                         GetTextureCoordinates(operation)));
-        return GetTextureElement(operation, texture);
+    Expression Texture(Operation operation) {
+        const auto& meta = std::get<MetaTexture>(operation.GetMeta());
+        UNIMPLEMENTED_IF(!meta.aoffi.empty());
+
+        const bool can_implicit = stage == ShaderType::Fragment;
+        const Id sampler = GetTextureSampler(operation);
+        const Id coords = GetCoordinates(operation, Type::Float);
+
+        if (meta.depth_compare) {
+            // Depth sampling
+            UNIMPLEMENTED_IF(meta.bias);
+            const Id dref = AsFloat(Visit(meta.depth_compare));
+            if (can_implicit) {
+                return {OpImageSampleDrefImplicitLod(t_float, sampler, coords, dref, {}),
+                        Type::Float};
+            } else {
+                return {OpImageSampleDrefExplicitLod(t_float, sampler, coords, dref,
+                                                     spv::ImageOperandsMask::Lod, v_float_zero),
+                        Type::Float};
+            }
+        }
+
+        std::vector<Id> operands;
+        spv::ImageOperandsMask mask{};
+        if (meta.bias) {
+            mask = mask | spv::ImageOperandsMask::Bias;
+            operands.push_back(AsFloat(Visit(meta.bias)));
+        }
+
+        Id texture;
+        if (can_implicit) {
+            texture = OpImageSampleImplicitLod(t_float4, sampler, coords, mask, operands);
+        } else {
+            texture = OpImageSampleExplicitLod(t_float4, sampler, coords,
+                                               mask | spv::ImageOperandsMask::Lod, v_float_zero,
+                                               operands);
+        }
+        return GetTextureElement(operation, texture, Type::Float);
     }
 
-    Id TextureLod(Operation operation) {
-        const auto meta = std::get_if<MetaTexture>(&operation.GetMeta());
-        const Id texture = Emit(OpImageSampleExplicitLod(
-            t_float4, GetTextureSampler(operation), GetTextureCoordinates(operation),
-            spv::ImageOperandsMask::Lod, Visit(meta->lod)));
-        return GetTextureElement(operation, texture);
+    Expression TextureLod(Operation operation) {
+        const auto& meta = std::get<MetaTexture>(operation.GetMeta());
+
+        const Id sampler = GetTextureSampler(operation);
+        const Id coords = GetCoordinates(operation, Type::Float);
+        const Id lod = AsFloat(Visit(meta.lod));
+
+        spv::ImageOperandsMask mask = spv::ImageOperandsMask::Lod;
+        std::vector<Id> operands;
+        if (!meta.aoffi.empty()) {
+            mask = mask | spv::ImageOperandsMask::Offset;
+            operands.push_back(GetOffsetCoordinates(operation));
+        }
+
+        if (meta.sampler.IsShadow()) {
+            const Id dref = AsFloat(Visit(meta.depth_compare));
+            return {
+                OpImageSampleDrefExplicitLod(t_float, sampler, coords, dref, mask, lod, operands),
+                Type::Float};
+        }
+        const Id texture = OpImageSampleExplicitLod(t_float4, sampler, coords, mask, lod, operands);
+        return GetTextureElement(operation, texture, Type::Float);
     }
 
-    Id TextureGather(Operation operation) {
-        const auto meta = std::get_if<MetaTexture>(&operation.GetMeta());
-        const auto coords = GetTextureCoordinates(operation);
+    Expression TextureGather(Operation operation) {
+        const auto& meta = std::get<MetaTexture>(operation.GetMeta());
+        UNIMPLEMENTED_IF(!meta.aoffi.empty());
 
-        Id texture;
-        if (meta->sampler.IsShadow()) {
-            texture = Emit(OpImageDrefGather(t_float4, GetTextureSampler(operation), coords,
-                                             Visit(meta->component)));
+        const Id coords = GetCoordinates(operation, Type::Float);
+        Id texture{};
+        if (meta.sampler.IsShadow()) {
+            texture = OpImageDrefGather(t_float4, GetTextureSampler(operation), coords,
+                                        AsFloat(Visit(meta.depth_compare)));
         } else {
             u32 component_value = 0;
-            if (meta->component) {
-                const auto component = std::get_if<ImmediateNode>(&*meta->component);
+            if (meta.component) {
+                const auto component = std::get_if<ImmediateNode>(&*meta.component);
                 ASSERT_MSG(component, "Component is not an immediate value");
                 component_value = component->GetValue();
             }
-            texture = Emit(OpImageGather(t_float4, GetTextureSampler(operation), coords,
-                                         Constant(t_uint, component_value)));
+            texture = OpImageGather(t_float4, GetTextureSampler(operation), coords,
+                                    Constant(t_uint, component_value));
         }
-
-        return GetTextureElement(operation, texture);
+        return GetTextureElement(operation, texture, Type::Float);
     }
 
-    Id TextureQueryDimensions(Operation operation) {
-        const auto meta = std::get_if<MetaTexture>(&operation.GetMeta());
-        const auto image_id = GetTextureImage(operation);
-        AddCapability(spv::Capability::ImageQuery);
+    Expression TextureQueryDimensions(Operation operation) {
+        const auto& meta = std::get<MetaTexture>(operation.GetMeta());
+        UNIMPLEMENTED_IF(!meta.aoffi.empty());
+        UNIMPLEMENTED_IF(meta.depth_compare);
 
-        if (meta->element == 3) {
-            return BitcastTo<Type::Float>(Emit(OpImageQueryLevels(t_int, image_id)));
+        const auto image_id = GetTextureImage(operation);
+        if (meta.element == 3) {
+            return {OpImageQueryLevels(t_int, image_id), Type::Int};
         }
 
-        const Id lod = VisitOperand<Type::Uint>(operation, 0);
+        const Id lod = AsUint(Visit(operation[0]));
         const std::size_t coords_count = [&]() {
-            switch (const auto type = meta->sampler.GetType(); type) {
+            switch (const auto type = meta.sampler.GetType(); type) {
             case Tegra::Shader::TextureType::Texture1D:
                 return 1;
             case Tegra::Shader::TextureType::Texture2D:
@@ -963,141 +1664,190 @@ private:
             }
         }();
 
-        if (meta->element >= coords_count) {
-            return Constant(t_float, 0.0f);
+        if (meta.element >= coords_count) {
+            return {v_float_zero, Type::Float};
         }
 
         const std::array<Id, 3> types = {t_int, t_int2, t_int3};
-        const Id sizes = Emit(OpImageQuerySizeLod(types.at(coords_count - 1), image_id, lod));
-        const Id size = Emit(OpCompositeExtract(t_int, sizes, meta->element));
-        return BitcastTo<Type::Float>(size);
+        const Id sizes = OpImageQuerySizeLod(types.at(coords_count - 1), image_id, lod);
+        const Id size = OpCompositeExtract(t_int, sizes, meta.element);
+        return {size, Type::Int};
     }
 
-    Id TextureQueryLod(Operation operation) {
-        UNIMPLEMENTED();
-        return {};
+    Expression TextureQueryLod(Operation operation) {
+        const auto& meta = std::get<MetaTexture>(operation.GetMeta());
+        UNIMPLEMENTED_IF(!meta.aoffi.empty());
+        UNIMPLEMENTED_IF(meta.depth_compare);
+
+        if (meta.element >= 2) {
+            UNREACHABLE_MSG("Invalid element");
+            return {v_float_zero, Type::Float};
+        }
+        const auto sampler_id = GetTextureSampler(operation);
+
+        const Id multiplier = Constant(t_float, 256.0f);
+        const Id multipliers = ConstantComposite(t_float2, multiplier, multiplier);
+
+        const Id coords = GetCoordinates(operation, Type::Float);
+        Id size = OpImageQueryLod(t_float2, sampler_id, coords);
+        size = OpFMul(t_float2, size, multipliers);
+        size = OpConvertFToS(t_int2, size);
+        return GetTextureElement(operation, size, Type::Int);
     }
 
-    Id TexelFetch(Operation operation) {
+    Expression TexelFetch(Operation operation) {
+        const auto& meta = std::get<MetaTexture>(operation.GetMeta());
+        UNIMPLEMENTED_IF(meta.depth_compare);
+
+        const Id image = GetTextureImage(operation);
+        const Id coords = GetCoordinates(operation, Type::Int);
+        Id fetch;
+        if (meta.lod && !meta.sampler.IsBuffer()) {
+            fetch = OpImageFetch(t_float4, image, coords, spv::ImageOperandsMask::Lod,
+                                 AsInt(Visit(meta.lod)));
+        } else {
+            fetch = OpImageFetch(t_float4, image, coords);
+        }
+        return GetTextureElement(operation, fetch, Type::Float);
+    }
+
+    Expression TextureGradient(Operation operation) {
+        const auto& meta = std::get<MetaTexture>(operation.GetMeta());
+        UNIMPLEMENTED_IF(!meta.aoffi.empty());
+
+        const Id sampler = GetTextureSampler(operation);
+        const Id coords = GetCoordinates(operation, Type::Float);
+        const auto [dx, dy] = GetDerivatives(operation);
+        const std::vector grad = {dx, dy};
+
+        static constexpr auto mask = spv::ImageOperandsMask::Grad;
+        const Id texture = OpImageSampleImplicitLod(t_float4, sampler, coords, mask, grad);
+        return GetTextureElement(operation, texture, Type::Float);
+    }
+
+    Expression ImageLoad(Operation operation) {
         UNIMPLEMENTED();
         return {};
     }
 
-    Id TextureGradient(Operation operation) {
-        UNIMPLEMENTED();
+    Expression ImageStore(Operation operation) {
+        const auto meta{std::get<MetaImage>(operation.GetMeta())};
+        std::vector<Id> colors;
+        for (const auto& value : meta.values) {
+            colors.push_back(AsUint(Visit(value)));
+        }
+
+        const Id coords = GetCoordinates(operation, Type::Int);
+        const Id texel = OpCompositeConstruct(t_uint4, colors);
+
+        OpImageWrite(GetImage(operation), coords, texel, {});
         return {};
     }
 
-    Id ImageLoad(Operation operation) {
+    Expression AtomicImageAdd(Operation operation) {
         UNIMPLEMENTED();
         return {};
     }
 
-    Id ImageStore(Operation operation) {
+    Expression AtomicImageMin(Operation operation) {
         UNIMPLEMENTED();
         return {};
     }
 
-    Id AtomicImageAdd(Operation operation) {
+    Expression AtomicImageMax(Operation operation) {
         UNIMPLEMENTED();
         return {};
     }
 
-    Id AtomicImageAnd(Operation operation) {
+    Expression AtomicImageAnd(Operation operation) {
         UNIMPLEMENTED();
         return {};
     }
 
-    Id AtomicImageOr(Operation operation) {
+    Expression AtomicImageOr(Operation operation) {
         UNIMPLEMENTED();
         return {};
     }
 
-    Id AtomicImageXor(Operation operation) {
+    Expression AtomicImageXor(Operation operation) {
         UNIMPLEMENTED();
         return {};
     }
 
-    Id AtomicImageExchange(Operation operation) {
+    Expression AtomicImageExchange(Operation operation) {
         UNIMPLEMENTED();
         return {};
     }
 
-    Id Branch(Operation operation) {
-        const auto target = std::get_if<ImmediateNode>(&*operation[0]);
-        UNIMPLEMENTED_IF(!target);
-
-        Emit(OpStore(jmp_to, Constant(t_uint, target->GetValue())));
-        Emit(OpBranch(continue_label));
-        inside_branch = conditional_nest_count;
-        if (conditional_nest_count == 0) {
-            Emit(OpLabel());
+    Expression Branch(Operation operation) {
+        const auto& target = std::get<ImmediateNode>(*operation[0]);
+        OpStore(jmp_to, Constant(t_uint, target.GetValue()));
+        OpBranch(continue_label);
+        inside_branch = true;
+        if (!conditional_branch_set) {
+            AddLabel();
         }
         return {};
     }
 
-    Id BranchIndirect(Operation operation) {
-        const Id op_a = VisitOperand<Type::Uint>(operation, 0);
+    Expression BranchIndirect(Operation operation) {
+        const Id op_a = AsUint(Visit(operation[0]));
 
-        Emit(OpStore(jmp_to, op_a));
-        Emit(OpBranch(continue_label));
-        inside_branch = conditional_nest_count;
-        if (conditional_nest_count == 0) {
-            Emit(OpLabel());
+        OpStore(jmp_to, op_a);
+        OpBranch(continue_label);
+        inside_branch = true;
+        if (!conditional_branch_set) {
+            AddLabel();
         }
         return {};
     }
 
-    Id PushFlowStack(Operation operation) {
-        const auto target = std::get_if<ImmediateNode>(&*operation[0]);
-        ASSERT(target);
-
+    Expression PushFlowStack(Operation operation) {
+        const auto& target = std::get<ImmediateNode>(*operation[0]);
         const auto [flow_stack, flow_stack_top] = GetFlowStack(operation);
-        const Id current = Emit(OpLoad(t_uint, flow_stack_top));
-        const Id next = Emit(OpIAdd(t_uint, current, Constant(t_uint, 1)));
-        const Id access = Emit(OpAccessChain(t_func_uint, flow_stack, current));
+        const Id current = OpLoad(t_uint, flow_stack_top);
+        const Id next = OpIAdd(t_uint, current, Constant(t_uint, 1));
+        const Id access = OpAccessChain(t_func_uint, flow_stack, current);
 
-        Emit(OpStore(access, Constant(t_uint, target->GetValue())));
-        Emit(OpStore(flow_stack_top, next));
+        OpStore(access, Constant(t_uint, target.GetValue()));
+        OpStore(flow_stack_top, next);
         return {};
     }
 
-    Id PopFlowStack(Operation operation) {
+    Expression PopFlowStack(Operation operation) {
         const auto [flow_stack, flow_stack_top] = GetFlowStack(operation);
-        const Id current = Emit(OpLoad(t_uint, flow_stack_top));
-        const Id previous = Emit(OpISub(t_uint, current, Constant(t_uint, 1)));
-        const Id access = Emit(OpAccessChain(t_func_uint, flow_stack, previous));
-        const Id target = Emit(OpLoad(t_uint, access));
-
-        Emit(OpStore(flow_stack_top, previous));
-        Emit(OpStore(jmp_to, target));
-        Emit(OpBranch(continue_label));
-        inside_branch = conditional_nest_count;
-        if (conditional_nest_count == 0) {
-            Emit(OpLabel());
+        const Id current = OpLoad(t_uint, flow_stack_top);
+        const Id previous = OpISub(t_uint, current, Constant(t_uint, 1));
+        const Id access = OpAccessChain(t_func_uint, flow_stack, previous);
+        const Id target = OpLoad(t_uint, access);
+
+        OpStore(flow_stack_top, previous);
+        OpStore(jmp_to, target);
+        OpBranch(continue_label);
+        inside_branch = true;
+        if (!conditional_branch_set) {
+            AddLabel();
         }
         return {};
     }
 
-    Id PreExit() {
-        switch (stage) {
-        case ShaderType::Vertex: {
-            // TODO(Rodrigo): We should use VK_EXT_depth_range_unrestricted instead, but it doesn't
-            // seem to be working on Nvidia's drivers and Intel (mesa and blob) doesn't support it.
-            const Id z_pointer = AccessElement(t_out_float, per_vertex, position_index, 2u);
-            Id depth = Emit(OpLoad(t_float, z_pointer));
-            depth = Emit(OpFAdd(t_float, depth, Constant(t_float, 1.0f)));
-            depth = Emit(OpFMul(t_float, depth, Constant(t_float, 0.5f)));
-            Emit(OpStore(z_pointer, depth));
-            break;
+    void PreExit() {
+        if (stage == ShaderType::Vertex) {
+            const u32 position_index = out_indices.position.value();
+            const Id z_pointer = AccessElement(t_out_float, out_vertex, position_index, 2U);
+            const Id w_pointer = AccessElement(t_out_float, out_vertex, position_index, 3U);
+            Id depth = OpLoad(t_float, z_pointer);
+            depth = OpFAdd(t_float, depth, OpLoad(t_float, w_pointer));
+            depth = OpFMul(t_float, depth, Constant(t_float, 0.5f));
+            OpStore(z_pointer, depth);
         }
-        case ShaderType::Fragment: {
+        if (stage == ShaderType::Fragment) {
             const auto SafeGetRegister = [&](u32 reg) {
                 // TODO(Rodrigo): Replace with contains once C++20 releases
                 if (const auto it = registers.find(reg); it != registers.end()) {
-                    return Emit(OpLoad(t_float, it->second));
+                    return OpLoad(t_float, it->second);
                 }
-                return Constant(t_float, 0.0f);
+                return v_float_zero;
             };
 
             UNIMPLEMENTED_IF_MSG(header.ps.omap.sample_mask != 0,
@@ -1112,8 +1862,8 @@ private:
                 // TODO(Subv): Figure out how dual-source blending is configured in the Switch.
                 for (u32 component = 0; component < 4; ++component) {
                     if (header.ps.IsColorComponentOutputEnabled(rt, component)) {
-                        Emit(OpStore(AccessElement(t_out_float, frag_colors.at(rt), component),
-                                     SafeGetRegister(current_reg)));
+                        OpStore(AccessElement(t_out_float, frag_colors.at(rt), component),
+                                SafeGetRegister(current_reg));
                         ++current_reg;
                     }
                 }
@@ -1121,110 +1871,117 @@ private:
             if (header.ps.omap.depth) {
                 // The depth output is always 2 registers after the last color output, and
                 // current_reg already contains one past the last color register.
-                Emit(OpStore(frag_depth, SafeGetRegister(current_reg + 1)));
+                OpStore(frag_depth, SafeGetRegister(current_reg + 1));
             }
-            break;
-        }
         }
-
-        return {};
     }
 
-    Id Exit(Operation operation) {
+    Expression Exit(Operation operation) {
         PreExit();
-        inside_branch = conditional_nest_count;
-        if (conditional_nest_count > 0) {
-            Emit(OpReturn());
+        inside_branch = true;
+        if (conditional_branch_set) {
+            OpReturn();
         } else {
             const Id dummy = OpLabel();
-            Emit(OpBranch(dummy));
-            Emit(dummy);
-            Emit(OpReturn());
-            Emit(OpLabel());
+            OpBranch(dummy);
+            AddLabel(dummy);
+            OpReturn();
+            AddLabel();
         }
         return {};
     }
 
-    Id Discard(Operation operation) {
-        inside_branch = conditional_nest_count;
-        if (conditional_nest_count > 0) {
-            Emit(OpKill());
+    Expression Discard(Operation operation) {
+        inside_branch = true;
+        if (conditional_branch_set) {
+            OpKill();
         } else {
             const Id dummy = OpLabel();
-            Emit(OpBranch(dummy));
-            Emit(dummy);
-            Emit(OpKill());
-            Emit(OpLabel());
+            OpBranch(dummy);
+            AddLabel(dummy);
+            OpKill();
+            AddLabel();
         }
         return {};
     }
 
-    Id EmitVertex(Operation operation) {
-        UNIMPLEMENTED();
+    Expression EmitVertex(Operation) {
+        OpEmitVertex();
         return {};
     }
 
-    Id EndPrimitive(Operation operation) {
-        UNIMPLEMENTED();
+    Expression EndPrimitive(Operation operation) {
+        OpEndPrimitive();
         return {};
     }
 
-    Id YNegate(Operation operation) {
-        UNIMPLEMENTED();
-        return {};
+    Expression InvocationId(Operation) {
+        return {OpLoad(t_int, invocation_id), Type::Int};
     }
 
-    template <u32 element>
-    Id LocalInvocationId(Operation) {
+    Expression YNegate(Operation) {
         UNIMPLEMENTED();
-        return {};
+        return {Constant(t_float, 1.0f), Type::Float};
     }
 
     template <u32 element>
-    Id WorkGroupId(Operation) {
-        UNIMPLEMENTED();
-        return {};
+    Expression LocalInvocationId(Operation) {
+        const Id id = OpLoad(t_uint3, local_invocation_id);
+        return {OpCompositeExtract(t_uint, id, element), Type::Uint};
     }
 
-    Id BallotThread(Operation) {
-        UNIMPLEMENTED();
-        return {};
+    template <u32 element>
+    Expression WorkGroupId(Operation operation) {
+        const Id id = OpLoad(t_uint3, workgroup_id);
+        return {OpCompositeExtract(t_uint, id, element), Type::Uint};
     }
 
-    Id VoteAll(Operation) {
-        UNIMPLEMENTED();
-        return {};
-    }
+    Expression BallotThread(Operation operation) {
+        const Id predicate = AsBool(Visit(operation[0]));
+        const Id ballot = OpSubgroupBallotKHR(t_uint4, predicate);
 
-    Id VoteAny(Operation) {
-        UNIMPLEMENTED();
-        return {};
+        if (!device.IsWarpSizePotentiallyBiggerThanGuest()) {
+            // Guest-like devices can just return the first index.
+            return {OpCompositeExtract(t_uint, ballot, 0U), Type::Uint};
+        }
+
+        // The others will have to return what is local to the current thread.
+        // For instance a device with a warp size of 64 will return the upper uint when the current
+        // thread is 38.
+        const Id tid = OpLoad(t_uint, thread_id);
+        const Id thread_index = OpShiftRightLogical(t_uint, tid, Constant(t_uint, 5));
+        return {OpVectorExtractDynamic(t_uint, ballot, thread_index), Type::Uint};
     }
 
-    Id VoteEqual(Operation) {
-        UNIMPLEMENTED();
-        return {};
+    template <Id (Module::*func)(Id, Id)>
+    Expression Vote(Operation operation) {
+        // TODO(Rodrigo): Handle devices with different warp sizes
+        const Id predicate = AsBool(Visit(operation[0]));
+        return {(this->*func)(t_bool, predicate), Type::Bool};
     }
 
-    Id ThreadId(Operation) {
-        UNIMPLEMENTED();
-        return {};
+    Expression ThreadId(Operation) {
+        return {OpLoad(t_uint, thread_id), Type::Uint};
     }
 
-    Id ShuffleIndexed(Operation) {
-        UNIMPLEMENTED();
-        return {};
+    Expression ShuffleIndexed(Operation operation) {
+        const Id value = AsFloat(Visit(operation[0]));
+        const Id index = AsUint(Visit(operation[1]));
+        return {OpSubgroupReadInvocationKHR(t_float, value, index), Type::Float};
     }
 
-    Id DeclareBuiltIn(spv::BuiltIn builtin, spv::StorageClass storage, Id type,
-                      const std::string& name) {
+    Id DeclareBuiltIn(spv::BuiltIn builtin, spv::StorageClass storage, Id type, std::string name) {
         const Id id = OpVariable(type, storage);
         Decorate(id, spv::Decoration::BuiltIn, static_cast<u32>(builtin));
-        AddGlobalVariable(Name(id, name));
+        AddGlobalVariable(Name(id, std::move(name)));
         interfaces.push_back(id);
         return id;
     }
 
+    Id DeclareInputBuiltIn(spv::BuiltIn builtin, Id type, std::string name) {
+        return DeclareBuiltIn(builtin, spv::StorageClass::Input, type, std::move(name));
+    }
+
     bool IsRenderTargetUsed(u32 rt) const {
         for (u32 component = 0; component < 4; ++component) {
             if (header.ps.IsColorComponentOutputEnabled(rt, component)) {
@@ -1242,66 +1999,148 @@ private:
             members.push_back(Constant(t_uint, element));
         }
 
-        return Emit(OpAccessChain(pointer_type, composite, members));
+        return OpAccessChain(pointer_type, composite, members);
     }
 
-    template <Type type>
-    Id VisitOperand(Operation operation, std::size_t operand_index) {
-        const Id value = Visit(operation[operand_index]);
-
-        switch (type) {
+    Id As(Expression expr, Type wanted_type) {
+        switch (wanted_type) {
         case Type::Bool:
+            return AsBool(expr);
         case Type::Bool2:
+            return AsBool2(expr);
         case Type::Float:
-            return value;
+            return AsFloat(expr);
         case Type::Int:
-            return Emit(OpBitcast(t_int, value));
+            return AsInt(expr);
         case Type::Uint:
-            return Emit(OpBitcast(t_uint, value));
+            return AsUint(expr);
         case Type::HalfFloat:
-            UNIMPLEMENTED();
+            return AsHalfFloat(expr);
+        default:
+            UNREACHABLE();
+            return expr.id;
         }
-        UNREACHABLE();
-        return value;
     }
 
-    template <Type type>
-    Id BitcastFrom(Id value) {
-        switch (type) {
-        case Type::Bool:
-        case Type::Bool2:
+    Id AsBool(Expression expr) {
+        ASSERT(expr.type == Type::Bool);
+        return expr.id;
+    }
+
+    Id AsBool2(Expression expr) {
+        ASSERT(expr.type == Type::Bool2);
+        return expr.id;
+    }
+
+    Id AsFloat(Expression expr) {
+        switch (expr.type) {
         case Type::Float:
-            return value;
+            return expr.id;
         case Type::Int:
         case Type::Uint:
-            return Emit(OpBitcast(t_float, value));
+            return OpBitcast(t_float, expr.id);
         case Type::HalfFloat:
-            UNIMPLEMENTED();
+            if (device.IsFloat16Supported()) {
+                return OpBitcast(t_float, expr.id);
+            }
+            return OpBitcast(t_float, OpPackHalf2x16(t_uint, expr.id));
+        default:
+            UNREACHABLE();
+            return expr.id;
         }
-        UNREACHABLE();
-        return value;
     }
 
-    template <Type type>
-    Id BitcastTo(Id value) {
-        switch (type) {
-        case Type::Bool:
-        case Type::Bool2:
+    Id AsInt(Expression expr) {
+        switch (expr.type) {
+        case Type::Int:
+            return expr.id;
+        case Type::Float:
+        case Type::Uint:
+            return OpBitcast(t_int, expr.id);
+        case Type::HalfFloat:
+            if (device.IsFloat16Supported()) {
+                return OpBitcast(t_int, expr.id);
+            }
+            return OpPackHalf2x16(t_int, expr.id);
+        default:
             UNREACHABLE();
+            return expr.id;
+        }
+    }
+
+    Id AsUint(Expression expr) {
+        switch (expr.type) {
+        case Type::Uint:
+            return expr.id;
         case Type::Float:
-            return Emit(OpBitcast(t_float, value));
         case Type::Int:
-            return Emit(OpBitcast(t_int, value));
-        case Type::Uint:
-            return Emit(OpBitcast(t_uint, value));
+            return OpBitcast(t_uint, expr.id);
         case Type::HalfFloat:
-            UNIMPLEMENTED();
+            if (device.IsFloat16Supported()) {
+                return OpBitcast(t_uint, expr.id);
+            }
+            return OpPackHalf2x16(t_uint, expr.id);
+        default:
+            UNREACHABLE();
+            return expr.id;
+        }
+    }
+
+    Id AsHalfFloat(Expression expr) {
+        switch (expr.type) {
+        case Type::HalfFloat:
+            return expr.id;
+        case Type::Float:
+        case Type::Int:
+        case Type::Uint:
+            if (device.IsFloat16Supported()) {
+                return OpBitcast(t_half, expr.id);
+            }
+            return OpUnpackHalf2x16(t_half, AsUint(expr));
+        default:
+            UNREACHABLE();
+            return expr.id;
+        }
+    }
+
+    Id GetHalfScalarFromFloat(Id value) {
+        if (device.IsFloat16Supported()) {
+            return OpFConvert(t_scalar_half, value);
         }
-        UNREACHABLE();
         return value;
     }
 
-    Id GetTypeDefinition(Type type) {
+    Id GetFloatFromHalfScalar(Id value) {
+        if (device.IsFloat16Supported()) {
+            return OpFConvert(t_float, value);
+        }
+        return value;
+    }
+
+    AttributeType GetAttributeType(u32 location) const {
+        if (stage != ShaderType::Vertex) {
+            return {Type::Float, t_in_float, t_in_float4};
+        }
+        switch (specialization.attribute_types.at(location)) {
+        case Maxwell::VertexAttribute::Type::SignedNorm:
+        case Maxwell::VertexAttribute::Type::UnsignedNorm:
+        case Maxwell::VertexAttribute::Type::Float:
+            return {Type::Float, t_in_float, t_in_float4};
+        case Maxwell::VertexAttribute::Type::SignedInt:
+            return {Type::Int, t_in_int, t_in_int4};
+        case Maxwell::VertexAttribute::Type::UnsignedInt:
+            return {Type::Uint, t_in_uint, t_in_uint4};
+        case Maxwell::VertexAttribute::Type::UnsignedScaled:
+        case Maxwell::VertexAttribute::Type::SignedScaled:
+            UNIMPLEMENTED();
+            return {Type::Float, t_in_float, t_in_float4};
+        default:
+            UNREACHABLE();
+            return {Type::Float, t_in_float, t_in_float4};
+        }
+    }
+
+    Id GetTypeDefinition(Type type) const {
         switch (type) {
         case Type::Bool:
             return t_bool;
@@ -1314,10 +2153,25 @@ private:
         case Type::Uint:
             return t_uint;
         case Type::HalfFloat:
+            return t_half;
+        default:
+            UNREACHABLE();
+            return {};
+        }
+    }
+
+    std::array<Id, 4> GetTypeVectorDefinitionLut(Type type) const {
+        switch (type) {
+        case Type::Float:
+            return {nullptr, t_float2, t_float3, t_float4};
+        case Type::Int:
+            return {nullptr, t_int2, t_int3, t_int4};
+        case Type::Uint:
+            return {nullptr, t_uint2, t_uint3, t_uint4};
+        default:
             UNIMPLEMENTED();
+            return {};
         }
-        UNREACHABLE();
-        return {};
     }
 
     std::tuple<Id, Id> CreateFlowStack() {
@@ -1327,9 +2181,11 @@ private:
         constexpr auto storage_class = spv::StorageClass::Function;
 
         const Id flow_stack_type = TypeArray(t_uint, Constant(t_uint, FLOW_STACK_SIZE));
-        const Id stack = Emit(OpVariable(TypePointer(storage_class, flow_stack_type), storage_class,
-                                         ConstantNull(flow_stack_type)));
-        const Id top = Emit(OpVariable(t_func_uint, storage_class, Constant(t_uint, 0)));
+        const Id stack = OpVariable(TypePointer(storage_class, flow_stack_type), storage_class,
+                                    ConstantNull(flow_stack_type));
+        const Id top = OpVariable(t_func_uint, storage_class, Constant(t_uint, 0));
+        AddLocalVariable(stack);
+        AddLocalVariable(top);
         return std::tie(stack, top);
     }
 
@@ -1358,8 +2214,8 @@ private:
         &SPIRVDecompiler::Unary<&Module::OpFNegate, Type::Float>,
         &SPIRVDecompiler::Unary<&Module::OpFAbs, Type::Float>,
         &SPIRVDecompiler::Ternary<&Module::OpFClamp, Type::Float>,
-        &SPIRVDecompiler::FCastHalf0,
-        &SPIRVDecompiler::FCastHalf1,
+        &SPIRVDecompiler::FCastHalf<0>,
+        &SPIRVDecompiler::FCastHalf<1>,
         &SPIRVDecompiler::Binary<&Module::OpFMin, Type::Float>,
         &SPIRVDecompiler::Binary<&Module::OpFMax, Type::Float>,
         &SPIRVDecompiler::Unary<&Module::OpCos, Type::Float>,
@@ -1407,7 +2263,7 @@ private:
         &SPIRVDecompiler::Unary<&Module::OpBitcast, Type::Uint, Type::Int>,
         &SPIRVDecompiler::Binary<&Module::OpShiftLeftLogical, Type::Uint>,
         &SPIRVDecompiler::Binary<&Module::OpShiftRightLogical, Type::Uint>,
-        &SPIRVDecompiler::Binary<&Module::OpShiftRightArithmetic, Type::Uint>,
+        &SPIRVDecompiler::Binary<&Module::OpShiftRightLogical, Type::Uint>,
         &SPIRVDecompiler::Binary<&Module::OpBitwiseAnd, Type::Uint>,
         &SPIRVDecompiler::Binary<&Module::OpBitwiseOr, Type::Uint>,
         &SPIRVDecompiler::Binary<&Module::OpBitwiseXor, Type::Uint>,
@@ -1426,8 +2282,8 @@ private:
         &SPIRVDecompiler::HCastFloat,
         &SPIRVDecompiler::HUnpack,
         &SPIRVDecompiler::HMergeF32,
-        &SPIRVDecompiler::HMergeH0,
-        &SPIRVDecompiler::HMergeH1,
+        &SPIRVDecompiler::HMergeHN<0>,
+        &SPIRVDecompiler::HMergeHN<1>,
         &SPIRVDecompiler::HPack2,
 
         &SPIRVDecompiler::LogicalAssign,
@@ -1435,8 +2291,9 @@ private:
         &SPIRVDecompiler::Binary<&Module::OpLogicalOr, Type::Bool>,
         &SPIRVDecompiler::Binary<&Module::OpLogicalNotEqual, Type::Bool>,
         &SPIRVDecompiler::Unary<&Module::OpLogicalNot, Type::Bool>,
-        &SPIRVDecompiler::LogicalPick2,
-        &SPIRVDecompiler::LogicalAnd2,
+        &SPIRVDecompiler::Binary<&Module::OpVectorExtractDynamic, Type::Bool, Type::Bool2,
+                                 Type::Uint>,
+        &SPIRVDecompiler::Unary<&Module::OpAll, Type::Bool, Type::Bool2>,
 
         &SPIRVDecompiler::Binary<&Module::OpFOrdLessThan, Type::Bool, Type::Float>,
         &SPIRVDecompiler::Binary<&Module::OpFOrdEqual, Type::Bool, Type::Float>,
@@ -1444,7 +2301,7 @@ private:
         &SPIRVDecompiler::Binary<&Module::OpFOrdGreaterThan, Type::Bool, Type::Float>,
         &SPIRVDecompiler::Binary<&Module::OpFOrdNotEqual, Type::Bool, Type::Float>,
         &SPIRVDecompiler::Binary<&Module::OpFOrdGreaterThanEqual, Type::Bool, Type::Float>,
-        &SPIRVDecompiler::Unary<&Module::OpIsNan, Type::Bool>,
+        &SPIRVDecompiler::Unary<&Module::OpIsNan, Type::Bool, Type::Float>,
 
         &SPIRVDecompiler::Binary<&Module::OpSLessThan, Type::Bool, Type::Int>,
         &SPIRVDecompiler::Binary<&Module::OpIEqual, Type::Bool, Type::Int>,
@@ -1460,19 +2317,19 @@ private:
         &SPIRVDecompiler::Binary<&Module::OpINotEqual, Type::Bool, Type::Uint>,
         &SPIRVDecompiler::Binary<&Module::OpUGreaterThanEqual, Type::Bool, Type::Uint>,
 
-        &SPIRVDecompiler::Binary<&Module::OpFOrdLessThan, Type::Bool, Type::HalfFloat>,
-        &SPIRVDecompiler::Binary<&Module::OpFOrdEqual, Type::Bool, Type::HalfFloat>,
-        &SPIRVDecompiler::Binary<&Module::OpFOrdLessThanEqual, Type::Bool, Type::HalfFloat>,
-        &SPIRVDecompiler::Binary<&Module::OpFOrdGreaterThan, Type::Bool, Type::HalfFloat>,
-        &SPIRVDecompiler::Binary<&Module::OpFOrdNotEqual, Type::Bool, Type::HalfFloat>,
-        &SPIRVDecompiler::Binary<&Module::OpFOrdGreaterThanEqual, Type::Bool, Type::HalfFloat>,
+        &SPIRVDecompiler::Binary<&Module::OpFOrdLessThan, Type::Bool2, Type::HalfFloat>,
+        &SPIRVDecompiler::Binary<&Module::OpFOrdEqual, Type::Bool2, Type::HalfFloat>,
+        &SPIRVDecompiler::Binary<&Module::OpFOrdLessThanEqual, Type::Bool2, Type::HalfFloat>,
+        &SPIRVDecompiler::Binary<&Module::OpFOrdGreaterThan, Type::Bool2, Type::HalfFloat>,
+        &SPIRVDecompiler::Binary<&Module::OpFOrdNotEqual, Type::Bool2, Type::HalfFloat>,
+        &SPIRVDecompiler::Binary<&Module::OpFOrdGreaterThanEqual, Type::Bool2, Type::HalfFloat>,
         // TODO(Rodrigo): Should these use the OpFUnord* variants?
-        &SPIRVDecompiler::Binary<&Module::OpFOrdLessThan, Type::Bool, Type::HalfFloat>,
-        &SPIRVDecompiler::Binary<&Module::OpFOrdEqual, Type::Bool, Type::HalfFloat>,
-        &SPIRVDecompiler::Binary<&Module::OpFOrdLessThanEqual, Type::Bool, Type::HalfFloat>,
-        &SPIRVDecompiler::Binary<&Module::OpFOrdGreaterThan, Type::Bool, Type::HalfFloat>,
-        &SPIRVDecompiler::Binary<&Module::OpFOrdNotEqual, Type::Bool, Type::HalfFloat>,
-        &SPIRVDecompiler::Binary<&Module::OpFOrdGreaterThanEqual, Type::Bool, Type::HalfFloat>,
+        &SPIRVDecompiler::Binary<&Module::OpFOrdLessThan, Type::Bool2, Type::HalfFloat>,
+        &SPIRVDecompiler::Binary<&Module::OpFOrdEqual, Type::Bool2, Type::HalfFloat>,
+        &SPIRVDecompiler::Binary<&Module::OpFOrdLessThanEqual, Type::Bool2, Type::HalfFloat>,
+        &SPIRVDecompiler::Binary<&Module::OpFOrdGreaterThan, Type::Bool2, Type::HalfFloat>,
+        &SPIRVDecompiler::Binary<&Module::OpFOrdNotEqual, Type::Bool2, Type::HalfFloat>,
+        &SPIRVDecompiler::Binary<&Module::OpFOrdGreaterThanEqual, Type::Bool2, Type::HalfFloat>,
 
         &SPIRVDecompiler::Texture,
         &SPIRVDecompiler::TextureLod,
@@ -1509,9 +2366,9 @@ private:
         &SPIRVDecompiler::WorkGroupId<2>,
 
         &SPIRVDecompiler::BallotThread,
-        &SPIRVDecompiler::VoteAll,
-        &SPIRVDecompiler::VoteAny,
-        &SPIRVDecompiler::VoteEqual,
+        &SPIRVDecompiler::Vote<&Module::OpSubgroupAllKHR>,
+        &SPIRVDecompiler::Vote<&Module::OpSubgroupAnyKHR>,
+        &SPIRVDecompiler::Vote<&Module::OpSubgroupAllEqualKHR>,
 
         &SPIRVDecompiler::ThreadId,
         &SPIRVDecompiler::ShuffleIndexed,
@@ -1522,8 +2379,7 @@ private:
     const ShaderIR& ir;
     const ShaderType stage;
     const Tegra::Shader::Header header;
-    u64 conditional_nest_count{};
-    u64 inside_branch{};
+    const Specialization& specialization;
 
     const Id t_void = Name(TypeVoid(), "void");
 
@@ -1551,20 +2407,28 @@ private:
     const Id t_func_uint = Name(TypePointer(spv::StorageClass::Function, t_uint), "func_uint");
 
     const Id t_in_bool = Name(TypePointer(spv::StorageClass::Input, t_bool), "in_bool");
+    const Id t_in_int = Name(TypePointer(spv::StorageClass::Input, t_int), "in_int");
+    const Id t_in_int4 = Name(TypePointer(spv::StorageClass::Input, t_int4), "in_int4");
     const Id t_in_uint = Name(TypePointer(spv::StorageClass::Input, t_uint), "in_uint");
+    const Id t_in_uint3 = Name(TypePointer(spv::StorageClass::Input, t_uint3), "in_uint3");
+    const Id t_in_uint4 = Name(TypePointer(spv::StorageClass::Input, t_uint4), "in_uint4");
     const Id t_in_float = Name(TypePointer(spv::StorageClass::Input, t_float), "in_float");
+    const Id t_in_float2 = Name(TypePointer(spv::StorageClass::Input, t_float2), "in_float2");
+    const Id t_in_float3 = Name(TypePointer(spv::StorageClass::Input, t_float3), "in_float3");
     const Id t_in_float4 = Name(TypePointer(spv::StorageClass::Input, t_float4), "in_float4");
 
+    const Id t_out_int = Name(TypePointer(spv::StorageClass::Output, t_int), "out_int");
+
     const Id t_out_float = Name(TypePointer(spv::StorageClass::Output, t_float), "out_float");
     const Id t_out_float4 = Name(TypePointer(spv::StorageClass::Output, t_float4), "out_float4");
 
     const Id t_cbuf_float = TypePointer(spv::StorageClass::Uniform, t_float);
     const Id t_cbuf_std140 = Decorate(
-        Name(TypeArray(t_float4, Constant(t_uint, MAX_CONSTBUFFER_ELEMENTS)), "CbufStd140Array"),
-        spv::Decoration::ArrayStride, 16u);
+        Name(TypeArray(t_float4, Constant(t_uint, MaxConstBufferElements)), "CbufStd140Array"),
+        spv::Decoration::ArrayStride, 16U);
     const Id t_cbuf_scalar = Decorate(
-        Name(TypeArray(t_float, Constant(t_uint, MAX_CONSTBUFFER_FLOATS)), "CbufScalarArray"),
-        spv::Decoration::ArrayStride, 4u);
+        Name(TypeArray(t_float, Constant(t_uint, MaxConstBufferFloats)), "CbufScalarArray"),
+        spv::Decoration::ArrayStride, 4U);
     const Id t_cbuf_std140_struct = MemberDecorate(
         Decorate(TypeStruct(t_cbuf_std140), spv::Decoration::Block), 0, spv::Decoration::Offset, 0);
     const Id t_cbuf_scalar_struct = MemberDecorate(
@@ -1572,28 +2436,43 @@ private:
     const Id t_cbuf_std140_ubo = TypePointer(spv::StorageClass::Uniform, t_cbuf_std140_struct);
     const Id t_cbuf_scalar_ubo = TypePointer(spv::StorageClass::Uniform, t_cbuf_scalar_struct);
 
+    Id t_smem_uint{};
+
     const Id t_gmem_float = TypePointer(spv::StorageClass::StorageBuffer, t_float);
     const Id t_gmem_array =
-        Name(Decorate(TypeRuntimeArray(t_float), spv::Decoration::ArrayStride, 4u), "GmemArray");
+        Name(Decorate(TypeRuntimeArray(t_float), spv::Decoration::ArrayStride, 4U), "GmemArray");
     const Id t_gmem_struct = MemberDecorate(
         Decorate(TypeStruct(t_gmem_array), spv::Decoration::Block), 0, spv::Decoration::Offset, 0);
     const Id t_gmem_ssbo = TypePointer(spv::StorageClass::StorageBuffer, t_gmem_struct);
 
     const Id v_float_zero = Constant(t_float, 0.0f);
+    const Id v_float_one = Constant(t_float, 1.0f);
+
+    // Nvidia uses these defaults for varyings (e.g. position and generic attributes)
+    const Id v_varying_default =
+        ConstantComposite(t_float4, v_float_zero, v_float_zero, v_float_zero, v_float_one);
+
     const Id v_true = ConstantTrue(t_bool);
     const Id v_false = ConstantFalse(t_bool);
 
-    Id per_vertex{};
+    Id t_scalar_half{};
+    Id t_half{};
+
+    Id out_vertex{};
+    Id in_vertex{};
     std::map<u32, Id> registers;
     std::map<Tegra::Shader::Pred, Id> predicates;
     std::map<u32, Id> flow_variables;
     Id local_memory{};
+    Id shared_memory{};
     std::array<Id, INTERNAL_FLAGS_COUNT> internal_flags{};
     std::map<Attribute::Index, Id> input_attributes;
     std::map<Attribute::Index, Id> output_attributes;
     std::map<u32, Id> constant_buffers;
     std::map<GlobalMemoryBase, Id> global_buffers;
-    std::map<u32, SamplerImage> sampler_images;
+    std::map<u32, TexelBuffer> texel_buffers;
+    std::map<u32, SampledImage> sampled_images;
+    std::map<u32, StorageImage> images;
 
     Id instance_index{};
     Id vertex_index{};
@@ -1601,18 +2480,20 @@ private:
     Id frag_depth{};
     Id frag_coord{};
     Id front_facing{};
-
-    u32 position_index{};
-    u32 point_size_index{};
-    u32 clip_distances_index{};
+    Id point_coord{};
+    Id tess_level_outer{};
+    Id tess_level_inner{};
+    Id tess_coord{};
+    Id invocation_id{};
+    Id workgroup_id{};
+    Id local_invocation_id{};
+    Id thread_id{};
+
+    VertexIndices in_indices;
+    VertexIndices out_indices;
 
     std::vector<Id> interfaces;
 
-    u32 const_buffers_base_binding{};
-    u32 global_buffers_base_binding{};
-    u32 samplers_base_binding{};
-
-    Id execute_function{};
     Id jmp_to{};
     Id ssy_flow_stack_top{};
     Id pbk_flow_stack_top{};
@@ -1620,6 +2501,9 @@ private:
     Id pbk_flow_stack{};
     Id continue_label{};
     std::map<u32, Id> labels;
+
+    bool conditional_branch_set{};
+    bool inside_branch{};
 };
 
 class ExprDecompiler {
@@ -1630,25 +2514,25 @@ public:
         const Id type_def = decomp.GetTypeDefinition(Type::Bool);
         const Id op1 = Visit(expr.operand1);
         const Id op2 = Visit(expr.operand2);
-        return decomp.Emit(decomp.OpLogicalAnd(type_def, op1, op2));
+        return decomp.OpLogicalAnd(type_def, op1, op2);
     }
 
     Id operator()(const ExprOr& expr) {
         const Id type_def = decomp.GetTypeDefinition(Type::Bool);
         const Id op1 = Visit(expr.operand1);
         const Id op2 = Visit(expr.operand2);
-        return decomp.Emit(decomp.OpLogicalOr(type_def, op1, op2));
+        return decomp.OpLogicalOr(type_def, op1, op2);
     }
 
     Id operator()(const ExprNot& expr) {
         const Id type_def = decomp.GetTypeDefinition(Type::Bool);
         const Id op1 = Visit(expr.operand1);
-        return decomp.Emit(decomp.OpLogicalNot(type_def, op1));
+        return decomp.OpLogicalNot(type_def, op1);
     }
 
     Id operator()(const ExprPredicate& expr) {
         const auto pred = static_cast<Tegra::Shader::Pred>(expr.predicate);
-        return decomp.Emit(decomp.OpLoad(decomp.t_bool, decomp.predicates.at(pred)));
+        return decomp.OpLoad(decomp.t_bool, decomp.predicates.at(pred));
     }
 
     Id operator()(const ExprCondCode& expr) {
@@ -1670,12 +2554,15 @@ public:
             }
         } else if (const auto flag = std::get_if<InternalFlagNode>(&*cc)) {
             target = decomp.internal_flags.at(static_cast<u32>(flag->GetFlag()));
+        } else {
+            UNREACHABLE();
         }
-        return decomp.Emit(decomp.OpLoad(decomp.t_bool, target));
+
+        return decomp.OpLoad(decomp.t_bool, target);
     }
 
     Id operator()(const ExprVar& expr) {
-        return decomp.Emit(decomp.OpLoad(decomp.t_bool, decomp.flow_variables.at(expr.var_index)));
+        return decomp.OpLoad(decomp.t_bool, decomp.flow_variables.at(expr.var_index));
     }
 
     Id operator()(const ExprBoolean& expr) {
@@ -1684,9 +2571,9 @@ public:
 
     Id operator()(const ExprGprEqual& expr) {
         const Id target = decomp.Constant(decomp.t_uint, expr.value);
-        const Id gpr = decomp.BitcastTo<Type::Uint>(
-            decomp.Emit(decomp.OpLoad(decomp.t_float, decomp.registers.at(expr.gpr))));
-        return decomp.Emit(decomp.OpLogicalEqual(decomp.t_uint, gpr, target));
+        Id gpr = decomp.OpLoad(decomp.t_float, decomp.registers.at(expr.gpr));
+        gpr = decomp.OpBitcast(decomp.t_uint, gpr);
+        return decomp.OpLogicalEqual(decomp.t_uint, gpr, target);
     }
 
     Id Visit(const Expr& node) {
@@ -1714,16 +2601,16 @@ public:
         const Id condition = expr_parser.Visit(ast.condition);
         const Id then_label = decomp.OpLabel();
         const Id endif_label = decomp.OpLabel();
-        decomp.Emit(decomp.OpSelectionMerge(endif_label, spv::SelectionControlMask::MaskNone));
-        decomp.Emit(decomp.OpBranchConditional(condition, then_label, endif_label));
-        decomp.Emit(then_label);
+        decomp.OpSelectionMerge(endif_label, spv::SelectionControlMask::MaskNone);
+        decomp.OpBranchConditional(condition, then_label, endif_label);
+        decomp.AddLabel(then_label);
         ASTNode current = ast.nodes.GetFirst();
         while (current) {
             Visit(current);
             current = current->GetNext();
         }
-        decomp.Emit(decomp.OpBranch(endif_label));
-        decomp.Emit(endif_label);
+        decomp.OpBranch(endif_label);
+        decomp.AddLabel(endif_label);
     }
 
     void operator()([[maybe_unused]] const ASTIfElse& ast) {
@@ -1741,7 +2628,7 @@ public:
     void operator()(const ASTVarSet& ast) {
         ExprDecompiler expr_parser{decomp};
         const Id condition = expr_parser.Visit(ast.condition);
-        decomp.Emit(decomp.OpStore(decomp.flow_variables.at(ast.index), condition));
+        decomp.OpStore(decomp.flow_variables.at(ast.index), condition);
     }
 
     void operator()([[maybe_unused]] const ASTLabel& ast) {
@@ -1758,12 +2645,11 @@ public:
         const Id loop_start_block = decomp.OpLabel();
         const Id loop_end_block = decomp.OpLabel();
         current_loop_exit = endloop_label;
-        decomp.Emit(decomp.OpBranch(loop_label));
-        decomp.Emit(loop_label);
-        decomp.Emit(
-            decomp.OpLoopMerge(endloop_label, loop_end_block, spv::LoopControlMask::MaskNone));
-        decomp.Emit(decomp.OpBranch(loop_start_block));
-        decomp.Emit(loop_start_block);
+        decomp.OpBranch(loop_label);
+        decomp.AddLabel(loop_label);
+        decomp.OpLoopMerge(endloop_label, loop_end_block, spv::LoopControlMask::MaskNone);
+        decomp.OpBranch(loop_start_block);
+        decomp.AddLabel(loop_start_block);
         ASTNode current = ast.nodes.GetFirst();
         while (current) {
             Visit(current);
@@ -1771,8 +2657,8 @@ public:
         }
         ExprDecompiler expr_parser{decomp};
         const Id condition = expr_parser.Visit(ast.condition);
-        decomp.Emit(decomp.OpBranchConditional(condition, loop_label, endloop_label));
-        decomp.Emit(endloop_label);
+        decomp.OpBranchConditional(condition, loop_label, endloop_label);
+        decomp.AddLabel(endloop_label);
     }
 
     void operator()(const ASTReturn& ast) {
@@ -1781,27 +2667,27 @@ public:
             const Id condition = expr_parser.Visit(ast.condition);
             const Id then_label = decomp.OpLabel();
             const Id endif_label = decomp.OpLabel();
-            decomp.Emit(decomp.OpSelectionMerge(endif_label, spv::SelectionControlMask::MaskNone));
-            decomp.Emit(decomp.OpBranchConditional(condition, then_label, endif_label));
-            decomp.Emit(then_label);
+            decomp.OpSelectionMerge(endif_label, spv::SelectionControlMask::MaskNone);
+            decomp.OpBranchConditional(condition, then_label, endif_label);
+            decomp.AddLabel(then_label);
             if (ast.kills) {
-                decomp.Emit(decomp.OpKill());
+                decomp.OpKill();
             } else {
                 decomp.PreExit();
-                decomp.Emit(decomp.OpReturn());
+                decomp.OpReturn();
             }
-            decomp.Emit(endif_label);
+            decomp.AddLabel(endif_label);
         } else {
             const Id next_block = decomp.OpLabel();
-            decomp.Emit(decomp.OpBranch(next_block));
-            decomp.Emit(next_block);
+            decomp.OpBranch(next_block);
+            decomp.AddLabel(next_block);
             if (ast.kills) {
-                decomp.Emit(decomp.OpKill());
+                decomp.OpKill();
             } else {
                 decomp.PreExit();
-                decomp.Emit(decomp.OpReturn());
+                decomp.OpReturn();
             }
-            decomp.Emit(decomp.OpLabel());
+            decomp.AddLabel(decomp.OpLabel());
         }
     }
 
@@ -1811,17 +2697,17 @@ public:
             const Id condition = expr_parser.Visit(ast.condition);
             const Id then_label = decomp.OpLabel();
             const Id endif_label = decomp.OpLabel();
-            decomp.Emit(decomp.OpSelectionMerge(endif_label, spv::SelectionControlMask::MaskNone));
-            decomp.Emit(decomp.OpBranchConditional(condition, then_label, endif_label));
-            decomp.Emit(then_label);
-            decomp.Emit(decomp.OpBranch(current_loop_exit));
-            decomp.Emit(endif_label);
+            decomp.OpSelectionMerge(endif_label, spv::SelectionControlMask::MaskNone);
+            decomp.OpBranchConditional(condition, then_label, endif_label);
+            decomp.AddLabel(then_label);
+            decomp.OpBranch(current_loop_exit);
+            decomp.AddLabel(endif_label);
         } else {
             const Id next_block = decomp.OpLabel();
-            decomp.Emit(decomp.OpBranch(next_block));
-            decomp.Emit(next_block);
-            decomp.Emit(decomp.OpBranch(current_loop_exit));
-            decomp.Emit(decomp.OpLabel());
+            decomp.OpBranch(next_block);
+            decomp.AddLabel(next_block);
+            decomp.OpBranch(current_loop_exit);
+            decomp.AddLabel(decomp.OpLabel());
         }
     }
 
@@ -1842,20 +2728,51 @@ void SPIRVDecompiler::DecompileAST() {
         flow_variables.emplace(i, AddGlobalVariable(id));
     }
 
+    DefinePrologue();
+
     const ASTNode program = ir.GetASTProgram();
     ASTDecompiler decompiler{*this};
     decompiler.Visit(program);
 
     const Id next_block = OpLabel();
-    Emit(OpBranch(next_block));
-    Emit(next_block);
+    OpBranch(next_block);
+    AddLabel(next_block);
+}
+
+} // Anonymous namespace
+
+ShaderEntries GenerateShaderEntries(const VideoCommon::Shader::ShaderIR& ir) {
+    ShaderEntries entries;
+    for (const auto& cbuf : ir.GetConstantBuffers()) {
+        entries.const_buffers.emplace_back(cbuf.second, cbuf.first);
+    }
+    for (const auto& [base, usage] : ir.GetGlobalMemory()) {
+        entries.global_buffers.emplace_back(base.cbuf_index, base.cbuf_offset, usage.is_written);
+    }
+    for (const auto& sampler : ir.GetSamplers()) {
+        if (sampler.IsBuffer()) {
+            entries.texel_buffers.emplace_back(sampler);
+        } else {
+            entries.samplers.emplace_back(sampler);
+        }
+    }
+    for (const auto& image : ir.GetImages()) {
+        entries.images.emplace_back(image);
+    }
+    for (const auto& attribute : ir.GetInputAttributes()) {
+        if (IsGenericAttribute(attribute)) {
+            entries.attributes.insert(GetGenericAttributeLocation(attribute));
+        }
+    }
+    entries.clip_distances = ir.GetClipDistances();
+    entries.shader_length = ir.GetLength();
+    entries.uses_warps = ir.UsesWarps();
+    return entries;
 }
 
-DecompilerResult Decompile(const VKDevice& device, const VideoCommon::Shader::ShaderIR& ir,
-                           ShaderType stage) {
-    auto decompiler = std::make_unique<SPIRVDecompiler>(device, ir, stage);
-    decompiler->Decompile();
-    return {std::move(decompiler), decompiler->GetShaderEntries()};
+std::vector<u32> Decompile(const VKDevice& device, const VideoCommon::Shader::ShaderIR& ir,
+                           ShaderType stage, const Specialization& specialization) {
+    return SPIRVDecompiler(device, ir, stage, specialization).Assemble();
 }
 
-} // namespace Vulkan::VKShader
+} // namespace Vulkan
diff --git a/src/video_core/renderer_vulkan/vk_shader_decompiler.h b/src/video_core/renderer_vulkan/vk_shader_decompiler.h
index 203fc00d0..2b01321b6 100644
--- a/src/video_core/renderer_vulkan/vk_shader_decompiler.h
+++ b/src/video_core/renderer_vulkan/vk_shader_decompiler.h
@@ -5,29 +5,28 @@
 #pragma once
 
 #include <array>
+#include <bitset>
 #include <memory>
 #include <set>
+#include <type_traits>
 #include <utility>
 #include <vector>
 
-#include <sirit/sirit.h>
-
 #include "common/common_types.h"
 #include "video_core/engines/maxwell_3d.h"
+#include "video_core/engines/shader_type.h"
 #include "video_core/shader/shader_ir.h"
 
-namespace VideoCommon::Shader {
-class ShaderIR;
-}
-
 namespace Vulkan {
 class VKDevice;
 }
 
-namespace Vulkan::VKShader {
+namespace Vulkan {
 
 using Maxwell = Tegra::Engines::Maxwell3D::Regs;
+using TexelBufferEntry = VideoCommon::Shader::Sampler;
 using SamplerEntry = VideoCommon::Shader::Sampler;
+using ImageEntry = VideoCommon::Shader::Image;
 
 constexpr u32 DESCRIPTOR_SET = 0;
 
@@ -46,39 +45,74 @@ private:
 
 class GlobalBufferEntry {
 public:
-    explicit GlobalBufferEntry(u32 cbuf_index, u32 cbuf_offset)
-        : cbuf_index{cbuf_index}, cbuf_offset{cbuf_offset} {}
+    constexpr explicit GlobalBufferEntry(u32 cbuf_index, u32 cbuf_offset, bool is_written)
+        : cbuf_index{cbuf_index}, cbuf_offset{cbuf_offset}, is_written{is_written} {}
 
-    u32 GetCbufIndex() const {
+    constexpr u32 GetCbufIndex() const {
         return cbuf_index;
     }
 
-    u32 GetCbufOffset() const {
+    constexpr u32 GetCbufOffset() const {
         return cbuf_offset;
     }
 
+    constexpr bool IsWritten() const {
+        return is_written;
+    }
+
 private:
     u32 cbuf_index{};
     u32 cbuf_offset{};
+    bool is_written{};
 };
 
 struct ShaderEntries {
-    u32 const_buffers_base_binding{};
-    u32 global_buffers_base_binding{};
-    u32 samplers_base_binding{};
+    u32 NumBindings() const {
+        return static_cast<u32>(const_buffers.size() + global_buffers.size() +
+                                texel_buffers.size() + samplers.size() + images.size());
+    }
+
     std::vector<ConstBufferEntry> const_buffers;
     std::vector<GlobalBufferEntry> global_buffers;
+    std::vector<TexelBufferEntry> texel_buffers;
     std::vector<SamplerEntry> samplers;
+    std::vector<ImageEntry> images;
     std::set<u32> attributes;
     std::array<bool, Maxwell::NumClipDistances> clip_distances{};
     std::size_t shader_length{};
-    Sirit::Id entry_function{};
-    std::vector<Sirit::Id> interfaces;
+    bool uses_warps{};
+};
+
+struct Specialization final {
+    u32 base_binding{};
+
+    // Compute specific
+    std::array<u32, 3> workgroup_size{};
+    u32 shared_memory_size{};
+
+    // Graphics specific
+    Maxwell::PrimitiveTopology primitive_topology{};
+    std::optional<float> point_size{};
+    std::array<Maxwell::VertexAttribute::Type, Maxwell::NumVertexAttributes> attribute_types{};
+
+    // Tessellation specific
+    struct {
+        Maxwell::TessellationPrimitive primitive{};
+        Maxwell::TessellationSpacing spacing{};
+        bool clockwise{};
+    } tessellation;
+};
+// Old gcc versions don't consider this trivially copyable.
+// static_assert(std::is_trivially_copyable_v<Specialization>);
+
+struct SPIRVShader {
+    std::vector<u32> code;
+    ShaderEntries entries;
 };
 
-using DecompilerResult = std::pair<std::unique_ptr<Sirit::Module>, ShaderEntries>;
+ShaderEntries GenerateShaderEntries(const VideoCommon::Shader::ShaderIR& ir);
 
-DecompilerResult Decompile(const VKDevice& device, const VideoCommon::Shader::ShaderIR& ir,
-                           Tegra::Engines::ShaderType stage);
+std::vector<u32> Decompile(const VKDevice& device, const VideoCommon::Shader::ShaderIR& ir,
+                           Tegra::Engines::ShaderType stage, const Specialization& specialization);
 
-} // namespace Vulkan::VKShader
+} // namespace Vulkan