[IR/Flint] adding shared SPIR-V IR and Gen9 vertex lowering

2026-08-05 17:20:49 +02:00
parent 85ae10d69a
commit 937b84cbc3
19 changed files with 1672 additions and 250 deletions
@@ -17,6 +17,7 @@ const ImplementationDesc = struct {
        *std.Build.Module,
        *std.Build.Module,
        *std.Build.Module,
+        *std.Build.Module,
        std.Build.ResolvedTarget,
        std.builtin.OptimizeMode,
        bool,
@@ -125,6 +126,7 @@ pub fn build(b: *std.Build) !void {
    base_mod.addImport("vulkan", vulkan);
    base_mod.addImport("zmath", zmath);
    base_mod.addImport("drm", drm);
+    base_mod.addImport("shader_ir", ir_mod);

    const base_c_includes = b.addTranslateC(.{
        .root_source_file = b.path("src/vulkan/c_includes.h"),
@@ -174,7 +176,7 @@ pub fn build(b: *std.Build) !void {
            for (implementations[0..impl_index], implementation_modules[0..impl_index]) |child_impl, child_mod|
                lib_mod.addImport(child_impl.name, child_mod);
        } else if (impl.custom) |func| {
-            func(b, options, lib, lib_mod, base_mod, vulkan, base_c_mod, target, optimize, use_llvm) catch continue;
+            func(b, options, lib, lib_mod, base_mod, vulkan, base_c_mod, ir_mod, target, optimize, use_llvm) catch continue;
        }

        const icd_file = b.addWriteFile(
@@ -433,6 +435,7 @@ fn customSoft(
    _: *std.Build.Module,
    _: *std.Build.Module,
    base_c_mod: *std.Build.Module,
+    _: *std.Build.Module,
    target: std.Build.ResolvedTarget,
    optimize: std.builtin.OptimizeMode,
    use_llvm: bool,
@@ -462,23 +465,20 @@ fn customSoft(
 // Flint specialized functions

 fn customFlint(
-    b: *std.Build,
+    _: *std.Build,
    _: *Step.Options,
    _: *Step.Compile,
    lib_mod: *std.Build.Module,
    _: *std.Build.Module,
    _: *std.Build.Module,
    base_c_mod: *std.Build.Module,
-    target: std.Build.ResolvedTarget,
-    optimize: std.builtin.OptimizeMode,
+    shader_ir_mod: *std.Build.Module,
+    _: std.Build.ResolvedTarget,
+    _: std.builtin.OptimizeMode,
    _: bool,
 ) !void {
    lib_mod.addImport("intel_c", base_c_mod);
-    lib_mod.addImport("shader_ir", b.createModule(.{
-        .root_source_file = b.path("src/compiler/root.zig"),
-        .target = target,
-        .optimize = optimize,
-    }));
+    lib_mod.addImport("shader_ir", shader_ir_mod);
 }

 // Phi specialized functions
@@ -491,6 +491,7 @@ fn customPhi(
    _: *std.Build.Module,
    _: *std.Build.Module,
    base_c_mod: *std.Build.Module,
+    _: *std.Build.Module,
    target: std.Build.ResolvedTarget,
    optimize: std.builtin.OptimizeMode,
    use_llvm: bool,
@@ -573,14 +573,39 @@ final large-shader implementation.

 ## SPIR-V frontend

-The compiler currently provides a word parser and an initial translator in
-`spirv/`. The parser validates the header, word counts, truncation, and literal
-strings. The translator selects one entry point and lowers a defined subset:
+The compiler currently provides a word parser, an owned `SourceModule`, and an
+initial translator in `spirv/`. The parser validates the header, word counts,
+truncation, and literal strings. `SourceModule` copies and retains validated
+SPIR-V so API objects can instantiate multiple entry points without borrowing
+application memory.
+
+Use `translator.instantiate` when retaining a source module:
+
+```zig
+var source = try ir.spirv.SourceModule.init(allocator, words);
+defer source.deinit(allocator);
+
+var module = try ir.spirv.translator.instantiate(allocator, &source, .{
+    .entry_point = "main",
+    .stage = .compute,
+    .specializations = &.{.{
+        .constant_id = 7, // SPIR-V SpecId
+        .data = std.mem.asBytes(&workgroup_width),
+    }},
+});
+defer module.deinit();
+```
+
+`translator.translate` remains a convenience wrapper for borrowed words. Each
+translation selects one entry point and returns an independent mutable IR
+module. The translator lowers a defined subset:

 - Vertex, fragment, and compute stages.
 - Basic scalar, vector, array, structure, pointer, and function types.
- Ordinary and composite constants; unapplied specialization constants are
-  refused.
+- Ordinary constants plus scalar boolean, integer, and floating-point
+  specialization constants selected through `SpecId`. Missing overrides use the
+  SPIR-V defaults, and specialization composites are rebuilt from their
+  specialized elements. `OpSpecConstantOp` is not evaluated yet.
 - Functions, blocks, branches, structured merge marks, and returns.
 - `OpPhi` into block parameters and edge arguments.
 - The arithmetic, comparison, select, bitcast, and composite operations named
@@ -0,0 +1,54 @@
+const std = @import("std");
+const Parser = @import("Parser.zig");
+
+const Self = @This();
+
+words: []u32,
+parsed: Parser,
+
+pub const Error = std.mem.Allocator.Error || Parser.Error;
+
+pub fn init(allocator: std.mem.Allocator, words: []const u32) Error!Self {
+    const owned_words = try allocator.dupe(u32, words);
+    errdefer allocator.free(owned_words);
+
+    return .{
+        .words = owned_words,
+        .parsed = try Parser.init(owned_words),
+    };
+}
+
+pub fn deinit(self: *Self, allocator: std.mem.Allocator) void {
+    allocator.free(self.words);
+    self.* = undefined;
+}
+
+pub fn code(self: *const Self) []const u32 {
+    return self.words;
+}
+
+pub fn parser(self: *const Self) Parser {
+    return self.parsed;
+}
+
+test "SPIR-V: source module owns and validates its words" {
+    var words = [_]u32{
+        0x07230203,
+        0x00010000,
+        0,
+        1,
+        0,
+    };
+
+    var source = try Self.init(std.testing.allocator, &words);
+    defer source.deinit(std.testing.allocator);
+
+    words[0] = 0;
+    try std.testing.expectEqual(@as(u32, 0x07230203), source.code()[0]);
+    try std.testing.expectEqual(@as(u8, 1), source.parser().header.major());
+}
+
+test "SPIR-V: source module rejects malformed input" {
+    const malformed = [_]u32{ 0, 0, 0, 0, 0 };
+    try std.testing.expectError(error.InvalidMagic, Self.init(std.testing.allocator, &malformed));
+}
@@ -3,19 +3,21 @@
 //! This namespace contains the SPIR-V parser and translator used to import shader
 //! modules into the compiler IR.
 //!
-//! `Parser` validates the SPIR-V header and iterates over binary instructions.
-//! `spec` exposes a minimalistic SPIR-V header translation.
+//! `Parser` validates borrowed SPIR-V words, while `SourceModule` owns and
+//! structurally validates words that need to outlive an API call. `spec` exposes a
+//! minimalistic SPIR-V header translation.
 //!
-//! The main entry point is `translator.translate`, which finds the requested entry
-//! point, maps its execution model to an IR shader stage, lowers supported types,
-//! constants, interfaces, instructions, and structured control flow, then validates
-//! the generated IR module.
+//! Use `translator.instantiate` to lower one entry point from a retained source
+//! module. `translator.translate` remains as a convenience wrapper for borrowed
+//! words.

 pub const Parser = @import("Parser.zig");
+pub const SourceModule = @import("SourceModule.zig");
 pub const translator = @import("translator.zig");
 pub const spec = @import("spirv.zig");

 test {
    _ = Parser;
+    _ = SourceModule;
    _ = translator;
 }
@@ -161,6 +161,7 @@ pub const ExecutionMode = enum(u32) {
 };

 pub const Decoration = enum(u32) {
+    spec_id = 1,
    built_in = 11,
    location = 30,
    component = 31,
@@ -1,10 +1,19 @@
 const std = @import("std");
+const builtin_info = @import("builtin");
 const Parser = @import("Parser.zig");
+const SourceModule = @import("SourceModule.zig");
 const spirv = @import("spirv.zig");
 const ir = @import("../ir/ir.zig");

+pub const SpecializationValue = struct {
+    constant_id: u32,
+    data: []const u8,
+};
+
 pub const Options = struct {
    entry_point: []const u8,
+    stage: ?ir.module.Stage = null,
+    specializations: []const SpecializationValue = &.{},
 };

 pub const TranslationError = error{
@@ -24,6 +33,8 @@ pub const TranslationError = error{
    UnsupportedType,
    UnsupportedConstant,
    SpecializationConstantsNotApplied,
+    InvalidSpecialization,
+    DuplicateSpecializationConstant,
    UnsupportedOpcode,
 };

@@ -34,6 +45,7 @@ const EntryPoint = struct {
 };

 const Decorations = struct {
+    spec_id: ?u32 = null,
    location: ?u32 = null,
    component: u8 = 0,
    index: u8 = 0,
@@ -57,6 +69,7 @@ const Context = struct {
    variable_defs: []?Parser.Instruction,
    names: []?[]const u8,
    decorations: []Decorations,
+    specializations: []const SpecializationValue,

    types: []?ir.id.TypeId,
    values: []?ir.id.ValueId,
@@ -82,6 +95,17 @@ const Context = struct {
        return self.names[index];
    }

+    fn specializationData(self: *const Context, result_id: u32) TranslationError!?[]const u8 {
+        const index = try self.idIndex(result_id);
+        const spec_id = self.decorations[index].spec_id orelse return null;
+
+        for (self.specializations) |specialization| {
+            if (specialization.constant_id == spec_id)
+                return specialization.data;
+        }
+        return null;
+    }
+
    fn translateType(self: *Context, spv_id: u32) anyerror!ir.id.TypeId {
        const index = try self.idIndex(spv_id);
        if (self.types[index]) |translated|
@@ -232,7 +256,7 @@ const Context = struct {
                try expectOperandCount(operands, 2);
                break :blk try self.builder.internConstant(try self.translateType(operands[0]), .null);
            },
-            .constant_composite => blk: {
+            .constant_composite, .spec_constant_composite => blk: {
                if (operands.len < 2)
                    return error.InvalidInstruction;

@@ -252,12 +276,45 @@ const Context = struct {
                    .{ .composite = elements },
                );
            },
-            .spec_constant_true,
-            .spec_constant_false,
-            .spec_constant,
-            .spec_constant_composite,
-            .spec_constant_op,
-            => return error.SpecializationConstantsNotApplied,
+            .spec_constant_true, .spec_constant_false => blk: {
+                try expectOperandCount(operands, 2);
+                const ty = try self.translateType(operands[0]);
+                const type_data = self.module.types.get(ty) orelse return error.InvalidId;
+                if (type_data.* != .boolean)
+                    return error.UnsupportedConstant;
+
+                const value = if (try self.specializationData(spv_id)) |data|
+                    try specializationBoolean(data)
+                else
+                    instruction.opcode == .spec_constant_true;
+                break :blk try self.builder.internConstant(ty, .{ .boolean = value });
+            },
+            .spec_constant => blk: {
+                if (operands.len < 3 or operands.len > 4)
+                    return error.InvalidInstruction;
+
+                const ty = try self.translateType(operands[0]);
+                const type_data = self.module.types.get(ty) orelse return error.InvalidId;
+                const default_bits = try literalBits(operands[2..]);
+                const override = try self.specializationData(spv_id);
+
+                break :blk switch (type_data.*) {
+                    .integer => |integer| try self.builder.internConstant(ty, .{
+                        .integer_bits = if (override) |data|
+                            try specializationBits(data, integer.bits)
+                        else
+                            default_bits,
+                    }),
+                    .floating => |floating| try self.builder.internConstant(ty, .{
+                        .float_bits = if (override) |data|
+                            try specializationBits(data, floating.bits)
+                        else
+                            default_bits,
+                    }),
+                    else => return error.UnsupportedConstant,
+                };
+            },
+            .spec_constant_op => return error.SpecializationConstantsNotApplied,

            else => return error.MissingDefinition,
        };
@@ -290,9 +347,12 @@ const Context = struct {
    }
 };

-pub fn translate(allocator: std.mem.Allocator, words: []const u32, options: Options) !ir.module.Module {
-    const parser = try Parser.init(words);
-    const entry_point = try findEntryPoint(parser, options.entry_point);
+/// Translates one entry point from a retained SPIR-V source into an independent
+/// common IR module. The returned module does not borrow from `source`.
+pub fn instantiate(allocator: std.mem.Allocator, source: *const SourceModule, options: Options) !ir.module.Module {
+    try validateSpecializations(options.specializations);
+    const parser = source.parser();
+    const entry_point = try findEntryPoint(parser, options.entry_point, options.stage);
    const stage = try translateStage(entry_point.model);

    var module = ir.module.Module.init(allocator, stage);
@@ -314,6 +374,7 @@ pub fn translate(allocator: std.mem.Allocator, words: []const u32, options: Opti
        .variable_defs = try allocOptional(Parser.Instruction, scratch, bound),
        .names = try allocOptional([]const u8, scratch, bound),
        .decorations = try scratch.alloc(Decorations, bound),
+        .specializations = options.specializations,
        .types = try allocOptional(ir.id.TypeId, scratch, bound),
        .values = try allocOptional(ir.id.ValueId, scratch, bound),
        .blocks = try allocOptional(ir.id.BlockId, scratch, bound),
@@ -334,6 +395,13 @@ pub fn translate(allocator: std.mem.Allocator, words: []const u32, options: Opti
    return module;
 }

+/// Convenience wrapper for callers that do not retain a source module.
+pub fn translate(allocator: std.mem.Allocator, words: []const u32, options: Options) !ir.module.Module {
+    var source = try SourceModule.init(allocator, words);
+    defer source.deinit(allocator);
+    return instantiate(allocator, &source, options);
+}
+
 fn collectDeclarations(context: *Context) !void {
    var iterator = context.parser.iterator();
    while (try iterator.next()) |instruction| {
@@ -378,6 +446,12 @@ fn collectDecoration(context: *Context, operands: []const u32) !void {
    const index = try context.idIndex(operands[0]);
    const decoration: spirv.Decoration = @enumFromInt(operands[1]);
    switch (decoration) {
+        .spec_id => {
+            try expectOperandCount(operands, 3);
+            if (context.decorations[index].spec_id != null)
+                return error.InvalidInstruction;
+            context.decorations[index].spec_id = operands[2];
+        },
        .built_in => {
            try expectOperandCount(operands, 3);
            context.decorations[index].builtin = operands[2];
@@ -455,7 +529,7 @@ fn translateInterfaces(context: *Context, interface_ids: []const u32) !void {
    }
 }

-fn findEntryPoint(parser: Parser, requested_name: []const u8) !EntryPoint {
+fn findEntryPoint(parser: Parser, requested_name: []const u8, requested_stage: ?ir.module.Stage) !EntryPoint {
    var found: ?EntryPoint = null;
    var iterator = parser.iterator();
    while (try iterator.next()) |instruction| {
@@ -473,11 +547,21 @@ fn findEntryPoint(parser: Parser, requested_name: []const u8) !EntryPoint {
        if (!try Parser.literalStringEquals(instruction.operands[2 .. 2 + string_words], requested_name))
            continue;

+        const model: spirv.ExecutionModel = @enumFromInt(instruction.operands[0]);
+        if (requested_stage) |stage| {
+            const candidate_stage = translateStage(model) catch |err| switch (err) {
+                error.UnsupportedExecutionModel => continue,
+                else => return err,
+            };
+            if (candidate_stage != stage)
+                continue;
+        }
+
        if (found != null)
            return error.AmbiguousEntryPoint;

        found = .{
-            .model = @enumFromInt(instruction.operands[0]),
+            .model = model,
            .function_id = instruction.operands[1],
            .interface_ids = instruction.operands[2 + string_words ..],
        };
@@ -1012,6 +1096,35 @@ fn translateCompareOpcode(opcode: spirv.Opcode) ir.instruction.CompareOpcode {
    };
 }

+fn validateSpecializations(specializations: []const SpecializationValue) TranslationError!void {
+    for (specializations, 0..) |specialization, index| {
+        for (specializations[0..index]) |previous| {
+            if (previous.constant_id == specialization.constant_id)
+                return error.DuplicateSpecializationConstant;
+        }
+    }
+}
+
+fn specializationBoolean(data: []const u8) TranslationError!bool {
+    if (data.len != @sizeOf(u32))
+        return error.InvalidSpecialization;
+    return std.mem.readInt(u32, data[0..4], builtin_info.target.cpu.arch.endian()) != 0;
+}
+
+fn specializationBits(data: []const u8, bit_width: u16) TranslationError!u64 {
+    const expected_size: usize = (@as(usize, bit_width) + 7) / 8;
+    if (data.len != expected_size)
+        return error.InvalidSpecialization;
+
+    return switch (expected_size) {
+        1 => data[0],
+        2 => std.mem.readInt(u16, data[0..2], builtin_info.target.cpu.arch.endian()),
+        4 => std.mem.readInt(u32, data[0..4], builtin_info.target.cpu.arch.endian()),
+        8 => std.mem.readInt(u64, data[0..8], builtin_info.target.cpu.arch.endian()),
+        else => error.InvalidSpecialization,
+    };
+}
+
 fn literalBits(words: []const u32) TranslationError!u64 {
    return switch (words.len) {
        1 => words[0],
@@ -1230,6 +1343,120 @@ test "SPIR-V: fragment execution modes and translated properties" {
    try std.testing.expect(entry.terminator.? == .return_void);
 }

+test "SPIR-V: retained source instantiates independent entry points" {
+    const assembly =
+        \\OpCapability Shader
+        \\OpMemoryModel Logical GLSL450
+        \\OpEntryPoint Vertex %vertex_main "main"
+        \\OpEntryPoint GLCompute %compute_main "main"
+        \\OpExecutionMode %compute_main LocalSize 2 1 1
+        \\%void = OpTypeVoid
+        \\%fn_void = OpTypeFunction %void
+        \\%vertex_main = OpFunction %void None %fn_void
+        \\    %vertex_entry = OpLabel
+        \\    OpReturn
+        \\OpFunctionEnd
+        \\%compute_main = OpFunction %void None %fn_void
+        \\    %compute_entry = OpLabel
+        \\    OpReturn
+        \\OpFunctionEnd
+    ;
+    const words = try assembleSpirv(std.testing.allocator, assembly);
+    defer std.testing.allocator.free(words);
+
+    var source = try SourceModule.init(std.testing.allocator, words);
+    defer source.deinit(std.testing.allocator);
+
+    var vertex_module = try instantiate(std.testing.allocator, &source, .{
+        .entry_point = "main",
+        .stage = .vertex,
+    });
+    defer vertex_module.deinit();
+    var compute_module = try instantiate(std.testing.allocator, &source, .{
+        .entry_point = "main",
+        .stage = .compute,
+    });
+    defer compute_module.deinit();
+
+    try std.testing.expectEqual(ir.module.Stage.vertex, vertex_module.stage);
+    try std.testing.expectEqual(ir.module.Stage.compute, compute_module.stage);
+    try std.testing.expectEqual(@as(?[3]u32, .{ 2, 1, 1 }), compute_module.execution_modes.workgroup_size);
+    try std.testing.expect(vertex_module.entry_point != null);
+    try std.testing.expect(compute_module.entry_point != null);
+}
+
+test "SPIR-V: scalar specialization constants and defaults" {
+    const assembly =
+        \\OpCapability Shader
+        \\OpMemoryModel Logical GLSL450
+        \\OpEntryPoint GLCompute %main "main"
+        \\OpExecutionMode %main LocalSize 1 1 1
+        \\OpName %number "number"
+        \\OpName %enabled "enabled"
+        \\OpName %pair "pair"
+        \\OpDecorate %number SpecId 7
+        \\OpDecorate %enabled SpecId 8
+        \\%void = OpTypeVoid
+        \\%bool = OpTypeBool
+        \\%u32 = OpTypeInt 32 0
+        \\%vec2_u32 = OpTypeVector %u32 2
+        \\%fn_void = OpTypeFunction %void
+        \\%number = OpSpecConstant %u32 3
+        \\%enabled = OpSpecConstantFalse %bool
+        \\%pair = OpSpecConstantComposite %vec2_u32 %number %number
+        \\%main = OpFunction %void None %fn_void
+        \\    %entry = OpLabel
+        \\    %sum = OpIAdd %u32 %number %number
+        \\    %selected = OpSelect %u32 %enabled %sum %number
+        \\    %first = OpCompositeExtract %u32 %pair 0
+        \\    OpReturn
+        \\OpFunctionEnd
+    ;
+    const words = try assembleSpirv(std.testing.allocator, assembly);
+    defer std.testing.allocator.free(words);
+
+    var source = try SourceModule.init(std.testing.allocator, words);
+    defer source.deinit(std.testing.allocator);
+
+    var defaults = try instantiate(std.testing.allocator, &source, .{
+        .entry_point = "main",
+        .stage = .compute,
+    });
+    defer defaults.deinit();
+    try expectNamedIntegerConstant(&defaults, "number", 3);
+    try expectNamedBooleanConstant(&defaults, "enabled", false);
+
+    const number_override: u32 = 42;
+    const enabled_override: u32 = 1;
+    const specializations = [_]SpecializationValue{
+        .{ .constant_id = 7, .data = std.mem.asBytes(&number_override) },
+        .{ .constant_id = 8, .data = std.mem.asBytes(&enabled_override) },
+    };
+    var specialized = try instantiate(std.testing.allocator, &source, .{
+        .entry_point = "main",
+        .stage = .compute,
+        .specializations = &specializations,
+    });
+    defer specialized.deinit();
+    try expectNamedIntegerConstant(&specialized, "number", 42);
+    try expectNamedBooleanConstant(&specialized, "enabled", true);
+
+    const invalid_size: u16 = 9;
+    try std.testing.expectError(error.InvalidSpecialization, instantiate(std.testing.allocator, &source, .{
+        .entry_point = "main",
+        .stage = .compute,
+        .specializations = &.{.{ .constant_id = 7, .data = std.mem.asBytes(&invalid_size) }},
+    }));
+    try std.testing.expectError(error.DuplicateSpecializationConstant, instantiate(std.testing.allocator, &source, .{
+        .entry_point = "main",
+        .stage = .compute,
+        .specializations = &.{
+            .{ .constant_id = 7, .data = std.mem.asBytes(&number_override) },
+            .{ .constant_id = 7, .data = std.mem.asBytes(&number_override) },
+        },
+    }));
+}
+
 test "SPIR-V: entry point lookup errors" {
    const single_entry_assembly =
        \\OpCapability Shader
@@ -1475,6 +1702,29 @@ test "SPIR-V: preserves location components and builtin interfaces" {
    try std.testing.expectEqual(ir.module.Builtin.position, position.semantic.builtin);
 }

+fn expectNamedIntegerConstant(module: *const ir.module.Module, name: []const u8, expected: u64) !void {
+    const value = findNamedConstant(module, name) orelse return error.MissingNamedConstant;
+    try std.testing.expect(value == .integer_bits);
+    try std.testing.expectEqual(expected, value.integer_bits);
+}
+
+fn expectNamedBooleanConstant(module: *const ir.module.Module, name: []const u8, expected: bool) !void {
+    const value = findNamedConstant(module, name) orelse return error.MissingNamedConstant;
+    try std.testing.expect(value == .boolean);
+    try std.testing.expectEqual(expected, value.boolean);
+}
+
+fn findNamedConstant(module: *const ir.module.Module, name: []const u8) ?ir.constant.ConstantValue {
+    for (module.values.entries.items) |entry| {
+        const value = entry orelse continue;
+        const value_name = value.name orelse continue;
+        if (!std.mem.eql(u8, value_name, name) or value.definition != .constant)
+            continue;
+        return module.constants.get(value.definition.constant).?.value;
+    }
+    return null;
+}
+
 fn assembleSpirv(allocator: std.mem.Allocator, assembly: []const u8) ![]u32 {
    var io_backend: std.Io.Threaded = .init(allocator, .{});
    defer io_backend.deinit();
@@ -7,6 +7,7 @@ const lib = @import("lib.zig");
 const pci_ids = @import("pci_ids.zig").map;

 const FlintDevice = @import("FlintDevice.zig");
+const compiler_device = @import("compiler/device.zig");

 const VkError = base.VkError;
 const SurfaceKHR = base.SurfaceKHR;
@@ -29,6 +30,7 @@ pub const extensions = [_]vk.ExtensionProperties{

 interface: Interface,
 kmd_type: lib.KmdType,
+compiler_info: ?compiler_device.DeviceInfo,
 node_path: [base.drm.max_node_name:0]u8,

 pub fn create(allocator: std.mem.Allocator, instance: *base.Instance, drm_device: *const base.drm.Device, kmd_type: lib.KmdType) VkError!*Self {
@@ -224,6 +226,7 @@ pub fn create(allocator: std.mem.Allocator, instance: *base.Instance, drm_device
    self.* = .{
        .interface = interface,
        .kmd_type = kmd_type,
+        .compiler_info = compiler_device.DeviceInfo.fromPciDeviceId(interface.props.device_id),
        .node_path = @splat(0),
    };
    const node_path = drm_device.nodePath();
@@ -1,47 +1,210 @@
 const std = @import("std");
 const vk = @import("vulkan");
 const base = @import("base");
+const shader_ir = @import("shader_ir");
+const compiler = @import("compiler/compiler.zig");
+const FlintPhysicalDevice = @import("FlintPhysicalDevice.zig");

 const VkError = base.VkError;

 const Self = @This();
 pub const Interface = base.Pipeline;

+const PipelineKind = enum {
+    graphics,
+    compute,
+};
+
+const CommonStage = struct {
+    stage: shader_ir.ir.module.Stage,
+    module: base.ShaderModule.IrModule,
+    program: ?compiler.Program,
+
+    fn deinit(self: *CommonStage) void {
+        if (self.program) |*program|
+            program.deinit();
+        self.module.deinit();
+        self.* = undefined;
+    }
+};
+
 interface: Interface,
-host_allocator: base.VulkanAllocator,
+artifact_allocator: base.VulkanAllocator,
+stages: []CommonStage,

 pub fn createCompute(device: *base.Device, allocator: std.mem.Allocator, cache: ?*base.PipelineCache, info: *const vk.ComputePipelineCreateInfo) VkError!*Self {
    const self = allocator.create(Self) catch return VkError.OutOfHostMemory;
-    errdefer allocator.destroy(self);
+    var initialized = false;
+    errdefer if (initialized) self.interface.destroy(allocator) else allocator.destroy(self);

    var interface = try Interface.initCompute(device, allocator, cache, info);
    interface.vtable = &.{ .destroy = destroy };
+
    self.* = .{
        .interface = interface,
-        .host_allocator = base.VulkanAllocator.from(allocator).clone(),
+        .artifact_allocator = base.VulkanAllocator.from(allocator).clone(),
+        .stages = &.{},
    };
-    errdefer self.interface.layout.unref(allocator);
+    initialized = true;

+    self.stages = try compileStages(self.artifact_allocator.allocator(), &.{info.stage}, .compute, compilerDeviceInfo(device));
    return self;
 }

 pub fn createGraphics(device: *base.Device, allocator: std.mem.Allocator, cache: ?*base.PipelineCache, info: *const vk.GraphicsPipelineCreateInfo) VkError!*Self {
    const self = allocator.create(Self) catch return VkError.OutOfHostMemory;
-    errdefer allocator.destroy(self);
+    var initialized = false;
+    errdefer if (initialized) self.interface.destroy(allocator) else allocator.destroy(self);

    var interface = try Interface.initGraphics(device, allocator, cache, info);
    interface.vtable = &.{ .destroy = destroy };

    self.* = .{
        .interface = interface,
-        .host_allocator = base.VulkanAllocator.from(allocator).clone(),
+        .artifact_allocator = base.VulkanAllocator.from(allocator).clone(),
+        .stages = &.{},
    };
-    errdefer self.interface.layout.unref(allocator);
+    initialized = true;

+    const stage_infos = if (info.p_stages) |stages|
+        stages[0..info.stage_count]
+    else
+        return VkError.ValidationFailed;
+    self.stages = try compileStages(self.artifact_allocator.allocator(), stage_infos, .graphics, compilerDeviceInfo(device));
    return self;
 }

+fn compileStages(allocator: std.mem.Allocator, infos: []const vk.PipelineShaderStageCreateInfo, pipeline_kind: PipelineKind, device_info: ?compiler.device.DeviceInfo) VkError![]CommonStage {
+    if (infos.len == 0)
+        return VkError.ValidationFailed;
+
+    const stages = allocator.alloc(CommonStage, infos.len) catch return VkError.OutOfHostMemory;
+    var initialized: usize = 0;
+    errdefer {
+        for (stages[0..initialized]) |*stage|
+            stage.deinit();
+        allocator.free(stages);
+    }
+
+    for (infos, stages) |*info, *stage| {
+        stage.* = try compileStage(allocator, info, pipeline_kind, device_info);
+        initialized += 1;
+    }
+    return stages;
+}
+
+fn compileStage(allocator: std.mem.Allocator, info: *const vk.PipelineShaderStageCreateInfo, pipeline_kind: PipelineKind, device_info: ?compiler.device.DeviceInfo) VkError!CommonStage {
+    const specializations = try specializationValues(allocator, info.p_specialization_info);
+    defer if (specializations.len != 0) allocator.free(specializations);
+
+    const expected_stage = commonStage(info.stage) orelse return VkError.ValidationFailed;
+    switch (pipeline_kind) {
+        .compute => if (expected_stage != .compute) return VkError.ValidationFailed,
+        .graphics => if (expected_stage == .compute) return VkError.ValidationFailed,
+    }
+
+    const shader_module = base.NonDispatchable(base.ShaderModule).fromHandleObject(info.module) catch |err| return err;
+    var module = shader_module.instantiateIr(allocator, .{
+        .entry_point = std.mem.span(info.p_name),
+        .stage = expected_stage,
+        .specializations = specializations,
+    }) catch |err| {
+        std.log.scoped(.FlintPipeline).err("common shader translation failed: {s}", .{@errorName(err)});
+        return switch (err) {
+            error.OutOfMemory => VkError.OutOfHostMemory,
+            else => VkError.ValidationFailed,
+        };
+    };
+    errdefer module.deinit();
+
+    std.debug.assert(module.stage == expected_stage);
+
+    var program = try lowerToFlint(allocator, &module, device_info);
+    errdefer if (program) |*value| value.deinit();
+
+    return .{
+        .stage = expected_stage,
+        .module = module,
+        .program = program,
+    };
+}
+
+fn lowerToFlint(allocator: std.mem.Allocator, module: *base.ShaderModule.IrModule, device_info: ?compiler.device.DeviceInfo) VkError!?compiler.Program {
+    const target = device_info orelse return null;
+    return compiler.lower.lower(allocator, module, target, .{}) catch |err| switch (err) {
+        error.OutOfMemory => VkError.OutOfHostMemory,
+        error.UnsupportedGeneration,
+        error.UnsupportedStage,
+        error.UnsupportedDispatchWidth,
+        error.UnsupportedType,
+        error.UnsupportedOperation,
+        error.UnsupportedTerminator,
+        => null,
+        else => {
+            std.log.scoped(.FlintPipeline).err("Flint shader lowering failed: {s}", .{@errorName(err)});
+            return VkError.ValidationFailed;
+        },
+    };
+}
+
+fn compilerDeviceInfo(device: *const base.Device) ?compiler.device.DeviceInfo {
+    const physical_device: *const FlintPhysicalDevice = @alignCast(@fieldParentPtr("interface", device.physical_device));
+    return physical_device.compiler_info;
+}
+
+fn specializationValues(allocator: std.mem.Allocator, info: ?*const vk.SpecializationInfo) VkError![]shader_ir.spirv.translator.SpecializationValue {
+    const specialization = info orelse return &.{};
+    if (specialization.map_entry_count == 0)
+        return &.{};
+
+    const entries = specialization.p_map_entries orelse return VkError.ValidationFailed;
+    const data: []const u8 = if (specialization.data_size == 0)
+        &.{}
+    else
+        @as([*]const u8, @ptrCast(@alignCast(specialization.p_data)))[0..specialization.data_size];
+
+    const values = allocator.alloc(shader_ir.spirv.translator.SpecializationValue, specialization.map_entry_count) catch
+        return VkError.OutOfHostMemory;
+    errdefer allocator.free(values);
+
+    for (entries[0..specialization.map_entry_count], values) |entry, *value| {
+        const offset: usize = entry.offset;
+        const end = std.math.add(usize, offset, entry.size) catch return VkError.ValidationFailed;
+        if (end > data.len)
+            return VkError.ValidationFailed;
+        value.* = .{
+            .constant_id = entry.constant_id,
+            .data = data[offset..end],
+        };
+    }
+    return values;
+}
+
+fn commonStage(stage: vk.ShaderStageFlags) ?shader_ir.ir.module.Stage {
+    const bits: u32 = @bitCast(stage);
+    const vertex_bits: u32 = @bitCast(vk.ShaderStageFlags{ .vertex_bit = true });
+    const fragment_bits: u32 = @bitCast(vk.ShaderStageFlags{ .fragment_bit = true });
+    const compute_bits: u32 = @bitCast(vk.ShaderStageFlags{ .compute_bit = true });
+
+    return if (bits == vertex_bits)
+        .vertex
+    else if (bits == fragment_bits)
+        .fragment
+    else if (bits == compute_bits)
+        .compute
+    else
+        null;
+}
+
+fn deinitStages(allocator: std.mem.Allocator, stages: []CommonStage) void {
+    for (stages) |*stage|
+        stage.deinit();
+    if (stages.len != 0)
+        allocator.free(stages);
+}
+
 pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void {
    const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
+    deinitStages(self.artifact_allocator.allocator(), self.stages);
    allocator.destroy(self);
 }
@@ -8,7 +8,6 @@ const Self = @This();
 pub const Interface = base.ShaderModule;

 interface: Interface,
-code: []u32,
 ref_count: std.atomic.Value(usize),

 pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const vk.ShaderModuleCreateInfo) VkError!*Self {
@@ -16,14 +15,11 @@ pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const v
    errdefer allocator.destroy(self);

    var interface = try Interface.init(device, allocator, info);
+    errdefer interface.deinit();
    interface.vtable = &.{ .destroy = destroy };
-    if (info.code_size % @sizeOf(u32) != 0) return VkError.ValidationFailed;
-    const code = allocator.dupe(u32, info.p_code[0 .. info.code_size / @sizeOf(u32)]) catch return VkError.OutOfHostMemory;
-    errdefer allocator.free(code);

    self.* = .{
        .interface = interface,
-        .code = code,
        .ref_count = std.atomic.Value(usize).init(1),
    };
    return self;
@@ -35,7 +31,7 @@ pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void {
 }

 pub fn drop(self: *Self, allocator: std.mem.Allocator) void {
-    allocator.free(self.code);
+    self.interface.deinit();
    allocator.destroy(self);
 }

@@ -44,7 +40,6 @@ pub fn ref(self: *Self) void {
 }

 pub fn unref(self: *Self, allocator: std.mem.Allocator) void {
-    if (self.ref_count.fetchSub(1, .release) == 1) {
+    if (self.ref_count.fetchSub(1, .acq_rel) == 1)
        self.drop(allocator);
-    }
 }
@@ -183,4 +183,5 @@ test "[ir] ID stability after removal" {

 test {
    _ = lower;
+    _ = lower.vertex_abi;
 }
@@ -38,6 +38,50 @@ pub const DeviceInfo = struct {
            .simd32 => self.supports_simd32,
        };
    }
+
+    pub fn fromPciDeviceId(raw_pci_device_id: u32) ?DeviceInfo {
+        if (raw_pci_device_id > 0xffff)
+            return null;
+        const pci_device_id: u16 = @intCast(raw_pci_device_id);
+
+        const platform: Platform = switch (pci_device_id & 0xff00) {
+            0x1900 => .skylake,
+            0x5900 => .kabylake,
+            0x3e00 => switch (pci_device_id) {
+                0x3ea0, 0x3ea1, 0x3ea2, 0x3ea3, 0x3ea4 => .whiskey_lake,
+                else => .coffee_lake,
+            },
+            0x9b00 => .comet_lake,
+            0x8a00 => .ice_lake,
+            0x4500 => .elkhart_lake,
+            0x4e00 => .jasper_lake,
+            else => switch (pci_device_id) {
+                0x0a84, 0x1a84, 0x1a85, 0x5a84, 0x5a85 => .broxton,
+                0x3184, 0x3185 => .gemini_lake,
+                0x87c0, 0x87ca => .kabylake,
+                else => return null,
+            },
+        };
+
+        const generation: Generation = switch (platform) {
+            .skylake,
+            .broxton,
+            .kabylake,
+            .gemini_lake,
+            .coffee_lake,
+            .whiskey_lake,
+            .comet_lake,
+            => .gen9,
+            .ice_lake, .elkhart_lake, .jasper_lake => .gen11,
+        };
+
+        return .{
+            .generation = generation,
+            .platform = platform,
+            .pci_device_id = pci_device_id,
+            .grf_count = 128,
+        };
+    }
 };

 pub const DispatchWidth = enum(u8) {
@@ -54,3 +98,16 @@ pub const ExecutionSize = enum(u8) {
    simd16 = 16,
    simd32 = 32,
 };
+
+test "compiler device: classify supported Intel PCI IDs" {
+    const std = @import("std");
+
+    try std.testing.expectEqual(Platform.skylake, DeviceInfo.fromPciDeviceId(0x1912).?.platform);
+    try std.testing.expectEqual(Platform.broxton, DeviceInfo.fromPciDeviceId(0x5a84).?.platform);
+    try std.testing.expectEqual(Platform.kabylake, DeviceInfo.fromPciDeviceId(0x5916).?.platform);
+    try std.testing.expectEqual(Platform.whiskey_lake, DeviceInfo.fromPciDeviceId(0x3ea0).?.platform);
+    try std.testing.expectEqual(Platform.comet_lake, DeviceInfo.fromPciDeviceId(0x9bc5).?.platform);
+    try std.testing.expectEqual(Generation.gen11, DeviceInfo.fromPciDeviceId(0x8a52).?.generation);
+    try std.testing.expectEqual(Generation.gen11, DeviceInfo.fromPciDeviceId(0x4e55).?.generation);
+    try std.testing.expectEqual(@as(?DeviceInfo, null), DeviceInfo.fromPciDeviceId(0x46a6));
+}
@@ -10,6 +10,7 @@ const Self = @This();

 pub const Error = std.mem.Allocator.Error || error{
    InvalidBlock,
+    InvalidInstruction,
    InvalidInsertionIndex,
    TerminatorAlreadySet,
 };
@@ -55,13 +56,7 @@ pub fn edge(self: *Self, target: ids.BlockId, arguments: []const pseudo.EdgeArgu
    };
 }

-pub fn appendInstruction(
-    self: *Self,
-    block_id: ids.BlockId,
-    execution_size: device.ExecutionSize,
-    predicate: ?operand.Predicate,
-    operation: instruction.Operation,
-) Error!ids.InstructionId {
+pub fn appendInstruction(self: *Self, block_id: ids.BlockId, execution_size: device.ExecutionSize, predicate: ?operand.Predicate, operation: instruction.Operation) Error!ids.InstructionId {
    const block = self.program.blocks.get(block_id) orelse return Error.InvalidBlock;
    return self.insertInstruction(block_id, block.instructions.items.len, execution_size, predicate, operation);
 }
@@ -91,6 +86,12 @@ pub fn insertInstruction(
    return instruction_id;
 }

+pub fn replaceOperation(self: *Self, instruction_id: ids.InstructionId, operation: instruction.Operation) Error!void {
+    const inst = self.program.instructions.getMut(instruction_id) orelse return Error.InvalidInstruction;
+    const owned_operation = try instruction.cloneOperation(self.program.allocator(), operation);
+    inst.operation = owned_operation;
+}
+
 pub fn setStructuredControl(self: *Self, block_id: ids.BlockId, control: instruction.StructuredControl) Error!void {
    const block = self.program.blocks.getMut(block_id) orelse return Error.InvalidBlock;
    block.structured_control = control;
@@ -162,6 +163,13 @@ test "[ir] Builder: construction and ordered insertion" {
    try std.testing.expectEqual(entry, program.instructions.get(first).?.parent_block);
    try std.testing.expectEqual(entry, program.instructions.get(second).?.parent_block);

+    try builder.replaceOperation(first, moveImmediate(register_id, 3));
+    const replaced = program.instructions.get(first).?;
+    try std.testing.expectEqual(entry, replaced.parent_block);
+    try std.testing.expectEqual(device.ExecutionSize.simd8, replaced.execution_size);
+    try std.testing.expectEqual(@as(u32, 3), replaced.operation.move.source.register.immediate.u32);
+    try std.testing.expectError(Error.InvalidInstruction, builder.replaceOperation(ids.InstructionId.fromIndex(999), moveImmediate(register_id, 4)));
+
    try builder.setStructuredControl(entry, .{ .selection = .{ .merge_block = exit } });
    try builder.setTerminator(entry, .{ .jump = try builder.edge(exit, &.{}) });
    try builder.setTerminator(exit, .end_thread);
@@ -8,6 +8,7 @@ const operand = @import("operand.zig");
 pub const Stage = shared_ir.Stage;

 pub const Properties = packed struct {
+    common_ir_lowered: bool = false,
    instructions_selected: bool = false,
    block_parameters_lowered: bool = false,
    parallel_copies_lowered: bool = false,
@@ -24,7 +25,7 @@ pub const Properties = packed struct {
    flags_allocated: bool = false,
    branches_resolved: bool = false,

-    _padding: u20 = 0,
+    _padding: u19 = 0,
 };

 pub const VertexPayload = struct {
@@ -31,6 +31,10 @@ pub const Error = error{
    DuplicateParallelCopyDestination,
    PredicatedParallelCopy,
    UnloweredParallelCopy,
+    UnloweredStageIo,
+    UnloweredMessage,
+    InvalidInterfaceSemantic,
+    InvalidMessage,
    EntryBlockHasParameters,
    DuplicateBlockParameter,
    EdgeArgumentCountMismatch,
@@ -125,8 +129,18 @@ fn validateInstruction(program: *const program_ir.Program, inst: instruction.Ins
        try validateFlag(program, predicate.flag);

    switch (inst.operation) {
-        .load_input => |op| try validateDestination(program, op.destination),
-        .store_output => |op| try validateSource(program, op.source),
+        .load_input => |op| {
+            if (program.properties.stage_io_lowered)
+                return Error.UnloweredStageIo;
+            try validateDestination(program, op.destination);
+            try validateInterfaceSemantic(op.semantic, .input);
+        },
+        .store_output => |op| {
+            if (program.properties.stage_io_lowered)
+                return Error.UnloweredStageIo;
+            try validateSource(program, op.source);
+            try validateInterfaceSemantic(op.semantic, .output);
+        },
        .move => |op| {
            try validateDestination(program, op.destination);
            try validateSource(program, op.source);
@@ -142,9 +156,17 @@ fn validateInstruction(program: *const program_ir.Program, inst: instruction.Ins
            try validateSource(program, op.rhs);
        },
        .send => |op| {
+            if (program.properties.messages_lowered)
+                return Error.UnloweredMessage;
            try validateSpan(program, op.payload);
            if (op.response) |response|
                try validateSpan(program, response);
+            switch (op.message) {
+                .urb_write => |urb_write| {
+                    if (op.response != null or (!urb_write.channels.x and !urb_write.channels.y and !urb_write.channels.z and !urb_write.channels.w))
+                        return Error.InvalidMessage;
+                },
+            }
        },
        .parallel_copy => |op| {
            if (program.properties.parallel_copies_lowered)
@@ -156,6 +178,29 @@ fn validateInstruction(program: *const program_ir.Program, inst: instruction.Ins
    }
 }

+const InterfaceDirection = enum { input, output };
+
+fn validateInterfaceSemantic(semantic: instruction.InterfaceSemantic, direction: InterfaceDirection) Error!void {
+    switch (semantic) {
+        .location => |location| {
+            if (location.component > 3)
+                return Error.InvalidInterfaceSemantic;
+        },
+        .builtin => |builtin| switch (direction) {
+            .input => switch (builtin.builtin) {
+                .vertex_index, .instance_index => if (builtin.component != 0)
+                    return Error.InvalidInterfaceSemantic,
+                .position => return Error.InvalidInterfaceSemantic,
+            },
+            .output => switch (builtin.builtin) {
+                .position => if (builtin.component > 3)
+                    return Error.InvalidInterfaceSemantic,
+                .vertex_index, .instance_index => return Error.InvalidInterfaceSemantic,
+            },
+        },
+    }
+}
+
 fn validateParallelCopy(program: *const program_ir.Program, copy: pseudo.ParallelCopy) Error!void {
    if (copy.register_copies.len == 0 and copy.flag_copies.len == 0)
        return Error.EmptyParallelCopy;
@@ -280,7 +325,18 @@ fn validateSpan(program: *const program_ir.Program, span: operand.RegisterSpan)
    if (span.register_count == 0)
        return Error.InvalidRegisterSpan;
    switch (span.base) {
-        .virtual, .physical_grf => try validateRegisterRef(program, span.base),
+        .virtual => |register_id| {
+            try validateRegisterRef(program, span.base);
+            const register = program.virtual_registers.get(register_id) orelse return Error.InvalidVirtualRegister;
+            const required_size = @as(u32, span.register_count) * program.device_info.grf_size_bytes;
+            if (register.size_bytes < required_size)
+                return Error.InvalidRegisterSpan;
+        },
+        .physical_grf => |physical| {
+            try validateRegisterRef(program, span.base);
+            if (physical.byte_offset != 0 or @as(u32, physical.number) + span.register_count > program.device_info.grf_count)
+                return Error.InvalidRegisterSpan;
+        },
        else => return Error.InvalidRegisterSpan,
    }
 }
@@ -1,5 +1,6 @@
 const std = @import("std");
-const shader_ir = @import("shader_ir").ir;
+const shader_compiler = @import("shader_ir");
+const shader_ir = shader_compiler.ir;
 const device = @import("../device.zig");
 const Builder = @import("../ir/Builder.zig");
 const ids = @import("../ir/id.zig");
@@ -11,6 +12,7 @@ const program_ir = @import("../ir/program.zig");
 const validator = @import("../ir/validator.zig");

 pub const block_arguments = @import("block_arguments.zig");
+pub const vertex_abi = @import("vertex_abi.zig");

 pub const Options = struct {
    dispatch_width: device.DispatchWidth = .simd8,
@@ -33,21 +35,26 @@ pub const Error = std.mem.Allocator.Error || error{

 const PredicateValue = pseudo.PredicateValue;

+const LoweredType = struct {
+    element_type: operand.DataType,
+    component_count: usize,
+};
+
 const ValueLocation = union(enum) {
-    source: operand.Source,
+    components: []const operand.Source,
    predicate: PredicateValue,
 };

 const LoweringState = struct {
    lowerer: *Lowerer,
    builder: Builder,
+    storage: std.mem.Allocator,
    block_map: []?ids.BlockId,
    value_locations: []?ValueLocation,

-    fn lowerType(self: *const LoweringState, type_id: shader_ir.id.TypeId) Error!operand.DataType {
+    fn lowerScalarType(self: *const LoweringState, type_id: shader_ir.id.TypeId) Error!operand.DataType {
        const ty = self.lowerer.module.types.get(type_id) orelse return Error.InvalidModule;
        return switch (ty.*) {
-            .void => Error.UnsupportedType,
            .integer => |integer| if (integer.bits == 32)
                switch (integer.signedness) {
                    .unsigned => .u32,
@@ -60,6 +67,24 @@ const LoweringState = struct {
        };
    }

+    fn lowerType(self: *const LoweringState, type_id: shader_ir.id.TypeId) Error!LoweredType {
+        const ty = self.lowerer.module.types.get(type_id) orelse return Error.InvalidModule;
+        return switch (ty.*) {
+            .integer, .floating => .{
+                .element_type = try self.lowerScalarType(type_id),
+                .component_count = 1,
+            },
+            .vector => |vector| if (vector.length >= 2 and vector.length <= 4)
+                .{
+                    .element_type = try self.lowerScalarType(vector.element_type),
+                    .component_count = vector.length,
+                }
+            else
+                Error.UnsupportedType,
+            else => Error.UnsupportedType,
+        };
+    }
+
    fn isBoolean(self: *const LoweringState, type_id: shader_ir.id.TypeId) Error!bool {
        const ty = self.lowerer.module.types.get(type_id) orelse return Error.InvalidModule;
        return ty.* == .boolean;
@@ -97,13 +122,28 @@ const LoweringState = struct {
        };
    }

-    fn addRegisterLocation(self: *LoweringState, value_id: shader_ir.id.ValueId, class: operand.RegisterClass) Error!operand.Source {
+    fn componentName(self: *LoweringState, name: ?[]const u8, component_index: usize, component_count: usize) Error!?[]const u8 {
+        if (name == null or component_count == 1)
+            return name;
+        const suffixes = "xyzw";
+        const formatted = try std.fmt.allocPrint(self.storage, "{s}_{c}", .{ name.?, suffixes[component_index] });
+        return @as([]const u8, formatted);
+    }
+
+    fn addRegisterLocation(self: *LoweringState, value_id: shader_ir.id.ValueId, class: operand.RegisterClass) Error![]const operand.Source {
        const value = self.lowerer.module.values.get(value_id) orelse return Error.InvalidModule;
-        const data_type = try self.lowerType(value.type);
-        const register_id = try self.addRegister(data_type, class, value.name);
-        const register_source = self.registerSource(register_id, data_type);
-        try self.putLocation(value_id, .{ .source = register_source });
-        return register_source;
+        const lowered_type = try self.lowerType(value.type);
+        const result = try self.storage.alloc(operand.Source, lowered_type.component_count);
+        for (result, 0..) |*component, component_index| {
+            const register_id = try self.addRegister(
+                lowered_type.element_type,
+                class,
+                try self.componentName(value.name, component_index, lowered_type.component_count),
+            );
+            component.* = self.registerSource(register_id, lowered_type.element_type);
+        }
+        try self.putLocation(value_id, .{ .components = result });
+        return result;
    }

    fn location(self: *LoweringState, value_id: shader_ir.id.ValueId) Error!ValueLocation {
@@ -117,7 +157,6 @@ const LoweringState = struct {
        switch (value.definition) {
            .constant => |constant_id| {
                const constant = self.lowerer.module.constants.get(constant_id) orelse return Error.InvalidModule;
-
                if (constant.type != value.type)
                    return Error.InvalidModule;

@@ -125,7 +164,7 @@ const LoweringState = struct {
                    .boolean => |boolean| .{ .predicate = .{ .constant = boolean } },
                    else => return Error.UnsupportedType,
                } else .{
-                    .source = try self.constantSource(value.type, constant.value),
+                    .components = try self.constantComponents(value.type, constant.value),
                };
                self.value_locations[value_id.index()] = result;
                return result;
@@ -140,23 +179,28 @@ const LoweringState = struct {
        }
    }

-    fn source(self: *LoweringState, value_id: shader_ir.id.ValueId) Error!operand.Source {
+    fn components(self: *LoweringState, value_id: shader_ir.id.ValueId) Error![]const operand.Source {
        return switch (try self.location(value_id)) {
-            .source => |value| value,
+            .components => |values| values,
            .predicate => Error.UnsupportedType,
        };
    }

+    fn source(self: *LoweringState, value_id: shader_ir.id.ValueId) Error!operand.Source {
+        const values = try self.components(value_id);
+        if (values.len != 1)
+            return Error.UnsupportedType;
+        return values[0];
+    }
+
    fn predicate(self: *LoweringState, value_id: shader_ir.id.ValueId) Error!PredicateValue {
        return switch (try self.location(value_id)) {
-            .source => Error.UnsupportedType,
+            .components => Error.UnsupportedType,
            .predicate => |value| value,
        };
    }

-    fn destination(self: *LoweringState, value_id: shader_ir.id.ValueId) Error!operand.Destination {
-        const source_value = try self.source(value_id);
-
+    fn destinationFromSource(source_value: operand.Source) Error!operand.Destination {
        if (source_value.negate or source_value.absolute)
            return Error.InvalidLoweredProgram;

@@ -170,8 +214,50 @@ const LoweringState = struct {
        };
    }

-    fn constantSource(self: *const LoweringState, type_id: shader_ir.id.TypeId, value: shader_ir.constant.ConstantValue) Error!operand.Source {
-        const data_type = try self.lowerType(type_id);
+    fn destination(self: *LoweringState, value_id: shader_ir.id.ValueId) Error!operand.Destination {
+        return destinationFromSource(try self.source(value_id));
+    }
+
+    fn constantComponents(self: *LoweringState, type_id: shader_ir.id.TypeId, value: shader_ir.constant.ConstantValue) Error![]const operand.Source {
+        const lowered_type = try self.lowerType(type_id);
+        const result = try self.storage.alloc(operand.Source, lowered_type.component_count);
+        if (lowered_type.component_count == 1) {
+            result[0] = try self.constantScalarSource(type_id, value);
+            return result;
+        }
+
+        const ty = self.lowerer.module.types.get(type_id) orelse return Error.InvalidModule;
+        const vector = switch (ty.*) {
+            .vector => |vector| vector,
+            else => return Error.InvalidModule,
+        };
+        switch (value) {
+            .composite => |elements| {
+                if (elements.len != lowered_type.component_count)
+                    return Error.InvalidModule;
+                for (elements, result) |constant_id, *component| {
+                    const element = self.lowerer.module.constants.get(constant_id) orelse return Error.InvalidModule;
+                    if (element.type != vector.element_type)
+                        return Error.InvalidModule;
+                    component.* = try self.constantScalarSource(element.type, element.value);
+                }
+            },
+            .null => {
+                const zero: shader_ir.constant.ConstantValue = switch (lowered_type.element_type) {
+                    .u32, .i32 => .{ .integer_bits = 0 },
+                    .f32 => .{ .float_bits = 0 },
+                    else => unreachable,
+                };
+                for (result) |*component|
+                    component.* = try self.constantScalarSource(vector.element_type, zero);
+            },
+            else => return Error.UnsupportedType,
+        }
+        return result;
+    }
+
+    fn constantScalarSource(self: *const LoweringState, type_id: shader_ir.id.TypeId, value: shader_ir.constant.ConstantValue) Error!operand.Source {
+        const data_type = try self.lowerScalarType(type_id);
        const immediate: operand.Immediate = switch (data_type) {
            .u32 => switch (value) {
                .integer_bits => |bits| .{ .u32 = @truncate(bits) },
@@ -254,7 +340,8 @@ const LoweringState = struct {
                    self.builder.addBlockParameter(target_block_id, .{ .flag = flag_id }) catch |err|
                        return mapProgramError(err);
                } else {
-                    const parameter_source = try self.addRegisterLocation(parameter_id, .temporary);
+                    const parameter_components = try self.addRegisterLocation(parameter_id, .temporary);
+                    for (parameter_components) |parameter_source| {
                        const register_id = switch (parameter_source.register) {
                            .virtual => |id| id,
                            else => return Error.InvalidLoweredProgram,
@@ -265,6 +352,7 @@ const LoweringState = struct {
                }
            }
        }
+    }

    fn lowerInstructions(self: *LoweringState, allocator: std.mem.Allocator) Error!void {
        const visited = try allocator.alloc(bool, self.lowerer.module.blocks.entries.items.len);
@@ -315,10 +403,11 @@ const LoweringState = struct {
            .binary => |operation| try self.lowerBinary(block_id, source_instruction.result, operation),
            .compare => |operation| try self.lowerCompare(block_id, source_instruction.result, operation),
            .select => |operation| try self.lowerSelect(block_id, source_instruction.result, operation),
-            .bitcast => |value_id| try self.lowerBitcast(source_instruction.result, value_id),
+            .bitcast => |value_id| try self.lowerBitcast(block_id, source_instruction.result, value_id),
            .load_interface => |operation| try self.lowerLoadInterface(block_id, source_instruction.result, operation),
            .store_interface => |operation| try self.lowerStoreInterface(block_id, source_instruction.result, operation),
-            .composite_construct, .composite_extract => return Error.UnsupportedOperation,
+            .composite_construct => |operation| try self.lowerCompositeConstruct(source_instruction.result, operation),
+            .composite_extract => |operation| try self.lowerCompositeExtract(source_instruction.result, operation),
            .call => return Error.UnsanitizedModule,
        }
    }
@@ -334,8 +423,7 @@ const LoweringState = struct {

    fn lowerUnary(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.Unary) Error!void {
        const result_id = try requireResult(result);
-        switch (operation.opcode) {
-            .logical_not => {
+        if (operation.opcode == .logical_not) {
            const source_predicate = try self.predicate(operation.operand);
            const inverted: PredicateValue = switch (source_predicate) {
                .constant => |value| .{ .constant = !value },
@@ -345,74 +433,71 @@ const LoweringState = struct {
                } },
            };
            try self.putLocation(result_id, .{ .predicate = inverted });
-            },
+            return;
+        }
+
+        const source_components = try self.components(operation.operand);
+        const result_components = try self.addRegisterLocation(result_id, .temporary);
+        if (source_components.len != result_components.len)
+            return Error.InvalidModule;
+
+        for (source_components, result_components) |source_component, result_component| {
+            if (source_component.type != result_component.type)
+                return Error.InvalidModule;
+            switch (operation.opcode) {
                .negate => {
-                const source_value = try self.source(operation.operand);
-                if (source_value.type != .i32 and source_value.type != .f32)
+                    if (source_component.type != .i32 and source_component.type != .f32)
                        return Error.UnsupportedOperation;
-                _ = try self.addRegisterLocation(result_id, .temporary);
-                var negated = source_value;
+                    var negated = source_component;
                    negated.negate = !negated.negate;
-                try self.appendMove(block_id, null, try self.destination(result_id), negated);
+                    try self.appendMove(block_id, null, try destinationFromSource(result_component), negated);
                },
                .bitwise_not => {
-                const source_value = try self.source(operation.operand);
-                if (source_value.type != .u32 and source_value.type != .i32)
-                    return Error.UnsupportedOperation;
-                _ = try self.addRegisterLocation(result_id, .temporary);
-                const all_ones: operand.Immediate = switch (source_value.type) {
+                    const all_ones: operand.Immediate = switch (source_component.type) {
                        .u32 => .{ .u32 = std.math.maxInt(u32) },
                        .i32 => .{ .i32 = -1 },
-                    else => unreachable,
+                        else => return Error.UnsupportedOperation,
                    };
                    try self.appendInstruction(block_id, null, .{
                        .binary = .{
                            .opcode = .bitwise_xor,
-                        .destination = try self.destination(result_id),
-                        .lhs = source_value,
+                            .destination = try destinationFromSource(result_component),
+                            .lhs = source_component,
                            .rhs = .{
                                .register = .{ .immediate = all_ones },
-                            .type = source_value.type,
+                                .type = source_component.type,
                                .region = operand.Region.broadcast(),
                            },
                        },
                    });
                },
+                .logical_not => unreachable,
+            }
        }
    }

    fn lowerBinary(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.Binary) Error!void {
        const result_id = try requireResult(result);
-        const lhs = try self.source(operation.lhs);
-        var rhs = try self.source(operation.rhs);
-        _ = try self.addRegisterLocation(result_id, .temporary);
-        const destination_value = try self.destination(result_id);
-
-        if (lhs.type != destination_value.type or rhs.type != destination_value.type)
+        const lhs_components = try self.components(operation.lhs);
+        const rhs_components = try self.components(operation.rhs);
+        const result_components = try self.addRegisterLocation(result_id, .temporary);
+        if (lhs_components.len == 0 or lhs_components.len != rhs_components.len or lhs_components.len != result_components.len)
            return Error.InvalidModule;

+        const data_type = lhs_components[0].type;
        const opcode: instruction.BinaryOpcode = switch (operation.opcode) {
-            .integer_add => if (lhs.type == .u32 or lhs.type == .i32) .add else return Error.UnsupportedOperation,
-            .float_add => if (lhs.type == .f32) .add else return Error.UnsupportedOperation,
-
-            .integer_subtract => if (lhs.type == .u32 or lhs.type == .i32) subtract: {
-                rhs.negate = !rhs.negate;
-                break :subtract .add;
-            } else return Error.UnsupportedOperation,
-
-            .float_subtract => if (lhs.type == .f32) subtract: {
-                rhs.negate = !rhs.negate;
-                break :subtract .add;
-            } else return Error.UnsupportedOperation,
-
-            .integer_multiply => if (lhs.type == .u32 or lhs.type == .i32) .multiply else return Error.UnsupportedOperation,
-            .float_multiply => if (lhs.type == .f32) .multiply else return Error.UnsupportedOperation,
-            .shift_left => if (lhs.type == .u32 or lhs.type == .i32) .shift_left else return Error.UnsupportedOperation,
-            .logical_shift_right => if (lhs.type == .u32) .shift_right else return Error.UnsupportedOperation,
-            .arithmetic_shift_right => if (lhs.type == .i32) .shift_right else return Error.UnsupportedOperation,
-            .bitwise_and => if (lhs.type == .u32 or lhs.type == .i32) .bitwise_and else return Error.UnsupportedOperation,
-            .bitwise_or => if (lhs.type == .u32 or lhs.type == .i32) .bitwise_or else return Error.UnsupportedOperation,
-            .bitwise_xor => if (lhs.type == .u32 or lhs.type == .i32) .bitwise_xor else return Error.UnsupportedOperation,
+            .integer_add => if (data_type == .u32 or data_type == .i32) .add else return Error.UnsupportedOperation,
+            .float_add => if (data_type == .f32) .add else return Error.UnsupportedOperation,
+            .integer_subtract => if (data_type == .u32 or data_type == .i32) .add else return Error.UnsupportedOperation,
+            .float_subtract => if (data_type == .f32) .add else return Error.UnsupportedOperation,
+            .integer_multiply => if (data_type == .u32 or data_type == .i32) .multiply else return Error.UnsupportedOperation,
+            .float_multiply => if (data_type == .f32) .multiply else return Error.UnsupportedOperation,
+            .shift_left => if (data_type == .u32 or data_type == .i32) .shift_left else return Error.UnsupportedOperation,
+            .logical_shift_right => if (data_type == .u32) .shift_right else return Error.UnsupportedOperation,
+            .arithmetic_shift_right => if (data_type == .i32) .shift_right else return Error.UnsupportedOperation,
+            .bitwise_and => if (data_type == .u32 or data_type == .i32) .bitwise_and else return Error.UnsupportedOperation,
+            .bitwise_or => if (data_type == .u32 or data_type == .i32) .bitwise_or else return Error.UnsupportedOperation,
+            .bitwise_xor => if (data_type == .u32 or data_type == .i32) .bitwise_xor else return Error.UnsupportedOperation,
            .unsigned_divide,
            .signed_divide,
            .unsigned_modulo,
@@ -424,15 +509,22 @@ const LoweringState = struct {
            => return Error.UnsupportedOperation,
        };

+        for (lhs_components, rhs_components, result_components) |lhs, rhs_value, result_component| {
+            if (lhs.type != data_type or rhs_value.type != data_type or result_component.type != data_type)
+                return Error.InvalidModule;
+            var rhs = rhs_value;
+            if (operation.opcode == .integer_subtract or operation.opcode == .float_subtract)
+                rhs.negate = !rhs.negate;
            try self.appendInstruction(block_id, null, .{
                .binary = .{
                    .opcode = opcode,
-                .destination = destination_value,
+                    .destination = try destinationFromSource(result_component),
                    .lhs = lhs,
                    .rhs = rhs,
                },
            });
        }
+    }

    fn lowerCompare(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.Compare) Error!void {
        const result_id = try requireResult(result);
@@ -441,8 +533,12 @@ const LoweringState = struct {
        if (!try self.isBoolean(result_value.type))
            return Error.InvalidModule;

-        const lhs = try self.source(operation.lhs);
-        const rhs = try self.source(operation.rhs);
+        const lhs_components = try self.components(operation.lhs);
+        const rhs_components = try self.components(operation.rhs);
+        if (lhs_components.len != 1 or rhs_components.len != 1)
+            return Error.UnsupportedOperation;
+        const lhs = lhs_components[0];
+        const rhs = rhs_components[0];
        if (lhs.type != rhs.type)
            return Error.InvalidModule;

@@ -477,44 +573,89 @@ const LoweringState = struct {

    fn lowerSelect(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.Select) Error!void {
        const result_id = try requireResult(result);
-        const true_value = try self.source(operation.true_value);
-        const false_value = try self.source(operation.false_value);
-        _ = try self.addRegisterLocation(result_id, .temporary);
-        const destination_value = try self.destination(result_id);
+        const true_components = try self.components(operation.true_value);
+        const false_components = try self.components(operation.false_value);
+        const result_components = try self.addRegisterLocation(result_id, .temporary);
+        if (true_components.len != false_components.len or true_components.len != result_components.len)
+            return Error.InvalidModule;

+        const condition = try self.predicate(operation.condition);
+        for (true_components, false_components, result_components) |true_value, false_value, result_component| {
+            const destination_value = try destinationFromSource(result_component);
            if (true_value.type != destination_value.type or false_value.type != destination_value.type)
                return Error.InvalidModule;

-        switch (try self.predicate(operation.condition)) {
-            .constant => |condition| try self.appendMove(
+            switch (condition) {
+                .constant => |constant| try self.appendMove(
                    block_id,
                    null,
                    destination_value,
-                if (condition) true_value else false_value,
+                    if (constant) true_value else false_value,
                ),
-            .dynamic => |condition| {
+                .dynamic => |dynamic| {
                    try self.appendMove(block_id, .{
-                    .flag = condition.flag,
-                    .inverse = !condition.inverse,
+                        .flag = dynamic.flag,
+                        .inverse = !dynamic.inverse,
                    }, destination_value, false_value);
-                try self.appendMove(block_id, condition, destination_value, true_value);
+                    try self.appendMove(block_id, dynamic, destination_value, true_value);
                },
            }
        }
+    }

-    fn lowerBitcast(self: *LoweringState, result: ?shader_ir.id.ValueId, source_id: shader_ir.id.ValueId) Error!void {
+    fn lowerBitcast(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, source_id: shader_ir.id.ValueId) Error!void {
        const result_id = try requireResult(result);
        const result_value = self.lowerer.module.values.get(result_id) orelse return Error.InvalidModule;
        const target_type = try self.lowerType(result_value.type);
-        var source_value = try self.source(source_id);
+        const source_components = try self.components(source_id);
+        const result_components = try self.addRegisterLocation(result_id, .temporary);
+        if (source_components.len != target_type.component_count or source_components.len != result_components.len)
+            return Error.UnsupportedOperation;

-        source_value.register = switch (source_value.register) {
-            .immediate => |immediate| .{ .immediate = bitcastImmediate(immediate, target_type) },
-            else => source_value.register,
+        for (source_components, result_components) |source_component, result_component| {
+            // The source operand type selects the reinterpretation used by the
+            // move; the target-typed register materializes it before any CFG edge.
+            var cast_source = source_component;
+            cast_source.register = switch (cast_source.register) {
+                .immediate => |immediate| .{ .immediate = bitcastImmediate(immediate, target_type.element_type) },
+                else => cast_source.register,
            };
+            cast_source.type = target_type.element_type;
+            try self.appendMove(block_id, null, try destinationFromSource(result_component), cast_source);
+        }
+    }

-        source_value.type = target_type;
-        try self.putLocation(result_id, .{ .source = source_value });
+    fn lowerCompositeConstruct(self: *LoweringState, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.CompositeConstruct) Error!void {
+        const result_id = try requireResult(result);
+        const result_value = self.lowerer.module.values.get(result_id) orelse return Error.InvalidModule;
+        const result_type = try self.lowerType(result_value.type);
+        if (result_type.component_count < 2 or operation.elements.len != result_type.component_count)
+            return Error.UnsupportedOperation;
+
+        const result_components = try self.storage.alloc(operand.Source, result_type.component_count);
+        for (operation.elements, result_components) |element_id, *component| {
+            const element_components = try self.components(element_id);
+            if (element_components.len != 1 or element_components[0].type != result_type.element_type)
+                return Error.InvalidModule;
+            component.* = element_components[0];
+        }
+        try self.putLocation(result_id, .{ .components = result_components });
+    }
+
+    fn lowerCompositeExtract(self: *LoweringState, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.CompositeExtract) Error!void {
+        const result_id = try requireResult(result);
+        if (operation.indices.len != 1)
+            return Error.UnsupportedOperation;
+        const source_components = try self.components(operation.composite);
+        const component_index: usize = operation.indices[0];
+        if (component_index >= source_components.len)
+            return Error.InvalidModule;
+
+        const result_value = self.lowerer.module.values.get(result_id) orelse return Error.InvalidModule;
+        const result_type = try self.lowerType(result_value.type);
+        if (result_type.component_count != 1 or result_type.element_type != source_components[component_index].type)
+            return Error.InvalidModule;
+        try self.putLocation(result_id, .{ .components = source_components[component_index .. component_index + 1] });
    }

    fn lowerLoadInterface(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.LoadInterface) Error!void {
@@ -533,14 +674,16 @@ const LoweringState = struct {
        if (result_value.type != variable.type)
            return Error.InvalidModule;

-        _ = try self.addRegisterLocation(result_id, .varying);
+        const result_components = try self.addRegisterLocation(result_id, .varying);
+        for (result_components, 0..) |result_component, component_index| {
            try self.appendInstruction(block_id, null, .{
                .load_input = .{
-                .destination = try self.destination(result_id),
-                .semantic = try lowerInterfaceSemantic(variable.semantic),
+                    .destination = try destinationFromSource(result_component),
+                    .semantic = try lowerInterfaceSemantic(variable.semantic, @intCast(component_index)),
                },
            });
        }
+    }

    fn lowerStoreInterface(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.StoreInterface) Error!void {
        try requireNoResult(result);
@@ -553,19 +696,21 @@ const LoweringState = struct {
        if (variable.direction != .output)
            return Error.InvalidModule;

-        const source_value = try self.source(operation.value);
+        const source_components = try self.components(operation.value);
        const value = self.lowerer.module.values.get(operation.value) orelse return Error.InvalidModule;

        if (value.type != variable.type)
            return Error.InvalidModule;

+        for (source_components, 0..) |source_component, component_index| {
            try self.appendInstruction(block_id, null, .{
                .store_output = .{
-                .semantic = try lowerInterfaceSemantic(variable.semantic),
-                .source = source_value,
+                    .semantic = try lowerInterfaceSemantic(variable.semantic, @intCast(component_index)),
+                    .source = source_component,
                },
            });
        }
+    }

    fn lowerControlAndTerminators(self: *LoweringState, allocator: std.mem.Allocator) Error!void {
        const source_entry = try self.sourceEntryFunction();
@@ -624,18 +769,19 @@ const LoweringState = struct {
        if (edge.arguments.len != target_source_block.parameters.items.len)
            return Error.InvalidModule;

-        const arguments = try allocator.alloc(pseudo.EdgeArgument, edge.arguments.len);
-        errdefer allocator.free(arguments);
-        for (edge.arguments, arguments) |argument_id, *argument| {
-            argument.* = switch (try self.location(argument_id)) {
-                .source => |source_value| .{ .source = source_value },
-                .predicate => |predicate_value| .{ .predicate = predicate_value },
-            };
+        var arguments: std.ArrayList(pseudo.EdgeArgument) = .empty;
+        defer arguments.deinit(allocator);
+        for (edge.arguments) |argument_id| {
+            switch (try self.location(argument_id)) {
+                .components => |bundle| for (bundle) |component|
+                    try arguments.append(allocator, .{ .source = component }),
+                .predicate => |predicate_value| try arguments.append(allocator, .{ .predicate = predicate_value }),
+            }
        }

        return .{
            .target = try self.mappedBlock(edge.target),
-            .arguments = arguments,
+            .arguments = try arguments.toOwnedSlice(allocator),
        };
    }
 };
@@ -668,6 +814,8 @@ pub const Lowerer = struct {
        // Only supports gen9 for now as it is the only gen I have access to
        if (self.device_info.generation != .gen9)
            return Error.UnsupportedGeneration;
+        if (self.module.stage != .vertex)
+            return Error.UnsupportedStage;

        if (self.options.dispatch_width != .simd8 or !self.device_info.supportsDispatch(self.options.dispatch_width))
            return Error.UnsupportedDispatchWidth;
@@ -705,6 +853,7 @@ pub const Lowerer = struct {
        var state: LoweringState = .{
            .lowerer = self,
            .builder = Builder.init(&program),
+            .storage = program.allocator(),
            .block_map = block_map,
            .value_locations = value_locations,
        };
@@ -714,7 +863,7 @@ pub const Lowerer = struct {
        try state.lowerInstructions(allocator);
        try state.lowerControlAndTerminators(allocator);

-        program.properties.instructions_selected = true;
+        program.properties.common_ir_lowered = true;
        validator.validate(&program) catch return Error.InvalidLoweredProgram;

        block_arguments.run(allocator, &program) catch |err| return switch (err) {
@@ -726,25 +875,30 @@ pub const Lowerer = struct {
    }
 };

-fn lowerInterfaceSemantic(semantic: shader_ir.module.InterfaceSemantic) Error!instruction.InterfaceSemantic {
+fn lowerInterfaceSemantic(semantic: shader_ir.module.InterfaceSemantic, component_offset: u8) Error!instruction.InterfaceSemantic {
    return switch (semantic) {
-        .location => |location| if (location.index == 0)
-            .{
+        .location => |location| location: {
+            if (location.index != 0)
+                return Error.UnsupportedOperation;
+            const component = std.math.add(u8, location.component, component_offset) catch return Error.UnsupportedOperation;
+            if (component > 3)
+                return Error.UnsupportedOperation;
+            break :location .{
                .location = .{
                    .location = location.location,
-                    .component = location.component,
+                    .component = component,
+                },
+            };
        },
-            }
-        else
-            Error.UnsupportedOperation,
        .builtin => |builtin| .{
            .builtin = .{
                .builtin = switch (builtin) {
                    .position => .position,
-                    .vertex_index => .vertex_index,
-                    .instance_index => .instance_index,
+                    .vertex_index => if (component_offset == 0) .vertex_index else return Error.UnsupportedOperation,
+                    .instance_index => if (component_offset == 0) .instance_index else return Error.UnsupportedOperation,
                    .frag_coord, .frag_depth, .global_invocation_id => return Error.UnsupportedOperation,
                },
+                .component = component_offset,
            },
        },
    };
@@ -790,6 +944,8 @@ fn expectLowered(source: []const u8, expected: []const u8) !void {

    var program = try lower(std.testing.allocator, &module, test_device, .{});
    defer program.deinit();
+    try std.testing.expect(program.properties.common_ir_lowered);
+    try std.testing.expect(!program.properties.instructions_selected);

    const actual = try printer.allocPrint(std.testing.allocator, &program);
    defer std.testing.allocator.free(actual);
@@ -802,6 +958,8 @@ fn expectLoweredFragments(source: []const u8, expected: []const []const u8, unex

    var program = try lower(std.testing.allocator, &module, test_device, .{});
    defer program.deinit();
+    try std.testing.expect(program.properties.common_ir_lowered);
+    try std.testing.expect(!program.properties.instructions_selected);

    const actual = try printer.allocPrint(std.testing.allocator, &program);
    defer std.testing.allocator.free(actual);
@@ -1063,13 +1221,12 @@ test "[ir] Lower: selects and bitcasts" {
        "[simd8] (+%condition) mov %inverted_choice:u32, 2:u32",
        "[simd8] (-%condition) mov %inverted_choice:u32, 1:u32",
        "[simd8] mov %constant_choice:u32, 1:u32",
-        "[simd8] add %constant_sum:u32, 1065353216:u32, 1:u32",
+        "[simd8] mov %one_bits:u32, 1065353216:u32",
+        "[simd8] add %constant_sum:u32, %one_bits:u32, 1:u32",
        "[simd8] mov %negative:f32, -1:f32",
-        "[simd8] add %register_sum:u32, %negative:u32, 1:u32",
-    }, &.{
-        "%one_bits: vgrf",
-        "%negative_bits: vgrf",
-    });
+        "[simd8] mov %negative_bits:u32, %negative:u32",
+        "[simd8] add %register_sum:u32, %negative_bits:u32, 1:u32",
+    }, &.{});
 }

 test "[ir] Lower: vertex interfaces" {
@@ -1101,6 +1258,139 @@ test "[ir] Lower: vertex interfaces" {
    }, &.{});
 }

+test "[ir] Lower: vector operations, composites, and interfaces" {
+    const source =
+        \\shader vertex @main
+        \\{
+        \\    @attribute_in: vec4[f32] = input[location(0), component(0), index(0)]
+        \\    @position_out: vec4[f32] = output[builtin(position)]
+        \\    %one_u32: constant u32 = bits(0x1)
+        \\    %two_u32: constant u32 = bits(0x2)
+        \\    %two_f32: constant f32 = bits(0x40000000)
+        \\    %scale_constant: constant vec4[f32] = [#2, #2, #2, #2]
+        \\    %zero_constant: constant vec4[f32] = null
+        \\    fn @main() -> void
+        \\    {
+        \\        .entry():
+        \\            %attribute: vec4[f32] = load_interface @attribute_in
+        \\            %scaled: vec4[f32] = float_multiply %attribute, %scale_constant
+        \\            %with_zero: vec4[f32] = float_add %scaled, %zero_constant
+        \\            %first: f32 = composite_extract %with_zero[0]
+        \\            %rebuilt: vec4[f32] = composite_construct %first, %first, %first, %first
+        \\            %condition: bool = cmp_unsigned_less %one_u32, %two_u32
+        \\            %selected: vec4[f32] = select %condition, %with_zero, %rebuilt
+        \\            %selected_bits: vec4[u32] = bitcast %selected
+        \\            %restored: vec4[f32] = bitcast %selected_bits
+        \\            store_interface @position_out, %restored
+        \\            return
+        \\    }
+        \\}
+    ;
+
+    try expectLoweredFragments(source, &.{
+        "%attribute_x: vgrf f32[8], class(varying)",
+        "%attribute_w: vgrf f32[8], class(varying)",
+        "[simd8] load_input %attribute_x:f32, location(0), component(0)",
+        "[simd8] load_input %attribute_w:f32, location(0), component(3)",
+        "[simd8] multiply %scaled_x:f32, %attribute_x:f32, 2:f32",
+        "[simd8] multiply %scaled_w:f32, %attribute_w:f32, 2:f32",
+        "[simd8] add %with_zero_x:f32, %scaled_x:f32, 0:f32",
+        "[simd8] (+%condition) mov %selected_x:f32, %with_zero_x:f32",
+        "[simd8] mov %selected_bits_x:u32, %selected_x:u32",
+        "[simd8] mov %restored_w:f32, %selected_bits_w:f32",
+        "[simd8] store_output builtin(position), component(0), %restored_x:f32",
+        "[simd8] store_output builtin(position), component(3), %restored_w:f32",
+    }, &.{
+        "%scale_constant_",
+        "%zero_constant_",
+        "%rebuilt_",
+    });
+}
+
+test "[ir] Lower: SPIR-V vec4 end-to-end" {
+    // Assembled from a vertex shader that loads a vec4 input, multiplies it by
+    // vec4(2.0), and stores the result to Position.
+    const words = [_]u32{
+        119734787,  65536,   458752,     15,         0,          131089,     1,          196622,
+        0,          1,       458767,     0,          1,          1852399981, 0,          2,
+        3,          262149,  1,          1852399981, 0,          327685,     2,          1885302377,
+        1953067887, 7237481, 393221,     3,          1601467759, 1769172848, 1852795252, 0,
+        327685,     4,       1769172848, 1852795252, 0,          262149,     5,          1818321779,
+        25701,      262215,  2,          30,         0,          262215,     3,          11,
+        0,          131091,  6,          196630,     7,          32,         262167,     8,
+        7,          4,       262176,     9,          1,          8,          262176,     10,
+        3,          8,       196641,     11,         6,          262187,     7,          12,
+        1073741824, 458796,  8,          13,         12,         12,         12,         12,
+        262203,     9,       2,          1,          262203,     10,         3,          3,
+        327734,     6,       1,          0,          11,         131320,     14,         262205,
+        8,          4,       2,          327813,     8,          5,          4,          13,
+        196670,     3,       5,          65789,      65592,
+    };
+
+    var module = try shader_compiler.spirv.translator.translate(std.testing.allocator, &words, .{
+        .entry_point = "main",
+        .stage = .vertex,
+    });
+    defer module.deinit();
+
+    var program = try lower(std.testing.allocator, &module, test_device, .{});
+    defer program.deinit();
+    const text = try printer.allocPrint(std.testing.allocator, &program);
+    defer std.testing.allocator.free(text);
+
+    for ([_][]const u8{
+        "[simd8] load_input %position_x:f32, location(0), component(0)",
+        "[simd8] load_input %position_w:f32, location(0), component(3)",
+        "[simd8] multiply %scaled_x:f32, %position_x:f32, 2:f32",
+        "[simd8] multiply %scaled_w:f32, %position_w:f32, 2:f32",
+        "[simd8] store_output builtin(position), component(0), %scaled_x:f32",
+        "[simd8] store_output builtin(position), component(3), %scaled_w:f32",
+    }) |fragment|
+        try std.testing.expect(std.mem.indexOf(u8, text, fragment) != null);
+}
+
+test "[ir] Lower: vector block parameter" {
+    const source =
+        \\shader vertex @main
+        \\{
+        \\    %one: constant u32 = bits(0x1)
+        \\    %two: constant u32 = bits(0x2)
+        \\    fn @main() -> void
+        \\    {
+        \\        .entry():
+        \\            %pair: vec2[u32] = composite_construct %one, %two
+        \\            branch .merge(%pair)
+        \\        .merge(%merged: vec2[u32]):
+        \\            %first: u32 = composite_extract %merged[0]
+        \\            return
+        \\    }
+        \\}
+    ;
+
+    try expectLoweredFragments(source, &.{
+        "%merged_x: vgrf u32[8]",
+        "%merged_y: vgrf u32[8]",
+        "parallel_copy [%merged_x:u32 <- 1:u32, %merged_y:u32 <- 2:u32]",
+    }, &.{
+        ".merge(",
+    });
+}
+
+test "[ir] Lower: reject vector interface component overflow" {
+    try expectLoweringError(
+        \\shader vertex @main
+        \\{
+        \\    @attribute_in: vec2[f32] = input[location(0), component(3), index(0)]
+        \\    fn @main() -> void
+        \\    {
+        \\        .entry():
+        \\            %attribute: vec2[f32] = load_interface @attribute_in
+        \\            return
+        \\    }
+        \\}
+    , Error.UnsupportedOperation);
+}
+
 test "[ir] Lower: constant conditional branch" {
    const source =
        \\shader vertex @main
@@ -1186,11 +1476,11 @@ test "[ir] Lower: unsupported operations" {
        \\    fn @main() -> void
        \\    {
        \\        .entry():
-        \\            %pair: vec2[u32] = composite_construct %one, %two
+        \\            %wide: vec5[u32] = composite_construct %one, %two, %one, %two, %one
        \\            return
        \\    }
        \\}
-    , Error.UnsupportedOperation);
+    , Error.UnsupportedType);

    try expectLoweringError(
        \\shader vertex @main
@@ -1239,5 +1529,10 @@ test "[ir] Lower: unsupported target configuration" {
    var gen10 = test_device;
    gen10.generation = .gen10;
    try std.testing.expectError(Error.UnsupportedGeneration, lower(std.testing.allocator, &module, gen10, .{}));
+
+    module.stage = .fragment;
+    try std.testing.expectError(Error.UnsupportedStage, lower(std.testing.allocator, &module, test_device, .{}));
+    module.stage = .vertex;
+
    try std.testing.expectError(Error.UnsupportedDispatchWidth, lower(std.testing.allocator, &module, test_device, .{ .dispatch_width = .simd16 }));
 }
@@ -0,0 +1,485 @@
+const std = @import("std");
+const Builder = @import("../ir/Builder.zig");
+const ids = @import("../ir/id.zig");
+const instruction = @import("../ir/instruction.zig");
+const operand = @import("../ir/operand.zig");
+const program_ir = @import("../ir/program.zig");
+const validator = @import("../ir/validator.zig");
+
+pub const InputComponent = struct {
+    location: u32,
+    component: u8,
+    payload_grf_offset: u16,
+};
+
+pub const Layout = struct {
+    input_components: []const InputComponent,
+    position_urb_offset: u16,
+};
+
+pub const Error = std.mem.Allocator.Error || error{
+    InvalidProgram,
+    UnsupportedTarget,
+    MissingVertexPayload,
+    InvalidLayout,
+    UnsupportedStageIo,
+    MissingPosition,
+    ExistingUrbWrite,
+};
+
+pub fn run(allocator: std.mem.Allocator, program: *program_ir.Program, layout: Layout) Error!void {
+    validator.validate(program) catch return Error.InvalidProgram;
+    if (program.properties.stage_io_lowered)
+        return;
+
+    if (!program.properties.common_ir_lowered or !program.properties.block_parameters_lowered or
+        program.properties.registers_allocated or program.properties.messages_lowered)
+        return Error.InvalidProgram;
+    if (program.device_info.generation != .gen9 or program.stage != .vertex or
+        program.dispatch_width != .simd8 or program.device_info.grf_size_bytes != 32)
+        return Error.UnsupportedTarget;
+
+    const vertex_payload = program.payload.vertex orelse return Error.MissingVertexPayload;
+    try validateLayout(program, vertex_payload, layout);
+
+    var position_components = instruction.ChannelMask{ .x = false, .y = false, .z = false, .w = false };
+    var end_thread_count: usize = 0;
+    try preflight(program, layout, &position_components, &end_thread_count);
+    if (!position_components.x or !position_components.y or !position_components.z or !position_components.w or end_thread_count == 0)
+        return Error.MissingPosition;
+
+    var builder = Builder.init(program);
+    const position_payload = builder.addVirtualRegister(.{
+        .size_bytes = 4 * program.device_info.grf_size_bytes,
+        .alignment_bytes = program.device_info.grf_size_bytes,
+        .element_type = .f32,
+        .lane_count = 4 * @intFromEnum(program.dispatch_width),
+        .class = .payload,
+        .spillable = false,
+        .name = "position_urb_payload",
+    }) catch |err| return mapBuilderError(err);
+
+    for (program.blocks.entries.items) |entry| {
+        const block = entry orelse continue;
+        for (block.instructions.items) |instruction_id| {
+            const inst = program.instructions.get(instruction_id) orelse return Error.InvalidProgram;
+            const replacement: ?instruction.Operation = switch (inst.operation) {
+                .load_input => |load| .{ .move = .{
+                    .destination = load.destination,
+                    .source = .{
+                        .register = .{ .physical_grf = try inputPhysicalGrf(program, vertex_payload, layout, load.semantic) },
+                        .type = load.destination.type,
+                        .region = operand.Region.contiguous(.simd8),
+                    },
+                } },
+                .store_output => |store| blk: {
+                    const component = try positionComponent(store.semantic);
+                    break :blk .{ .move = .{
+                        .destination = .{
+                            .register = .{ .virtual = position_payload },
+                            .type = .f32,
+                            .region = .{ .byte_offset = @as(u16, component) * program.device_info.grf_size_bytes },
+                        },
+                        .source = store.source,
+                    } };
+                },
+                else => null,
+            };
+            if (replacement) |operation|
+                builder.replaceOperation(instruction_id, operation) catch |err| return mapBuilderError(err);
+        }
+    }
+
+    for (program.blocks.entries.items, 0..) |entry, block_index| {
+        const block = entry orelse continue;
+        if (block.terminator.? != .end_thread)
+            continue;
+        _ = builder.appendInstruction(ids.BlockId.fromIndex(block_index), .simd8, null, .{
+            .send = .{
+                .message = .{ .urb_write = .{
+                    .offset = layout.position_urb_offset,
+                    .channels = .{},
+                    .end_of_thread = true,
+                } },
+                .payload = .{
+                    .base = .{ .virtual = position_payload },
+                    .register_count = 4,
+                },
+            },
+        }) catch |err| return mapBuilderError(err);
+    }
+
+    program.properties.stage_io_lowered = true;
+    validator.validate(program) catch return Error.InvalidProgram;
+    _ = allocator;
+}
+
+fn validateLayout(program: *const program_ir.Program, vertex_payload: program_ir.VertexPayload, layout: Layout) Error!void {
+    if (vertex_payload.first_attribute_grf.byte_offset != 0 or vertex_payload.attribute_grf_count == 0)
+        return Error.InvalidLayout;
+    if (@as(u32, vertex_payload.first_attribute_grf.number) + vertex_payload.attribute_grf_count > program.device_info.grf_count)
+        return Error.InvalidLayout;
+
+    for (layout.input_components, 0..) |mapping, index| {
+        if (mapping.component > 3 or mapping.payload_grf_offset >= vertex_payload.attribute_grf_count)
+            return Error.InvalidLayout;
+        for (layout.input_components[0..index]) |previous| {
+            if (previous.location == mapping.location and previous.component == mapping.component)
+                return Error.InvalidLayout;
+        }
+    }
+}
+
+fn preflight(program: *const program_ir.Program, layout: Layout, position_components: *instruction.ChannelMask, end_thread_count: *usize) Error!void {
+    for (program.blocks.entries.items) |entry| {
+        const block = entry orelse continue;
+        for (block.instructions.items) |instruction_id| {
+            const inst = program.instructions.get(instruction_id) orelse return Error.InvalidProgram;
+            switch (inst.operation) {
+                .load_input => |load| {
+                    if (inst.execution_size != .simd8 or findInput(layout, load.semantic) == null)
+                        return Error.UnsupportedStageIo;
+                },
+                .store_output => |store| {
+                    if (inst.execution_size != .simd8 or store.source.type != .f32)
+                        return Error.UnsupportedStageIo;
+                    switch (try positionComponent(store.semantic)) {
+                        0 => position_components.x = true,
+                        1 => position_components.y = true,
+                        2 => position_components.z = true,
+                        3 => position_components.w = true,
+                        else => unreachable,
+                    }
+                },
+                .send => |send| switch (send.message) {
+                    .urb_write => return Error.ExistingUrbWrite,
+                },
+                else => {},
+            }
+        }
+        switch (block.terminator orelse return Error.InvalidProgram) {
+            .end_thread => end_thread_count.* += 1,
+            else => {},
+        }
+    }
+}
+
+fn findInput(layout: Layout, semantic: instruction.InterfaceSemantic) ?InputComponent {
+    const location = switch (semantic) {
+        .location => |location| location,
+        .builtin => return null,
+    };
+    for (layout.input_components) |mapping| {
+        if (mapping.location == location.location and mapping.component == location.component)
+            return mapping;
+    }
+    return null;
+}
+
+fn inputPhysicalGrf(program: *const program_ir.Program, vertex_payload: program_ir.VertexPayload, layout: Layout, semantic: instruction.InterfaceSemantic) Error!operand.PhysicalGrf {
+    const mapping = findInput(layout, semantic) orelse return Error.UnsupportedStageIo;
+    const number = @as(u32, vertex_payload.first_attribute_grf.number) + mapping.payload_grf_offset;
+    if (number >= program.device_info.grf_count)
+        return Error.InvalidLayout;
+    return .{ .number = @intCast(number) };
+}
+
+fn positionComponent(semantic: instruction.InterfaceSemantic) Error!u8 {
+    return switch (semantic) {
+        .builtin => |builtin| if (builtin.builtin == .position and builtin.component <= 3)
+            builtin.component
+        else
+            Error.UnsupportedStageIo,
+        .location => Error.UnsupportedStageIo,
+    };
+}
+
+fn mapBuilderError(err: anyerror) Error {
+    return switch (err) {
+        error.OutOfMemory => Error.OutOfMemory,
+        else => Error.InvalidProgram,
+    };
+}
+
+fn appendTestShaderBody(program: *program_ir.Program, position_component_count: u8) !ids.BlockId {
+    var builder = Builder.init(program);
+    program.properties.common_ir_lowered = true;
+    program.properties.block_parameters_lowered = true;
+    program.payload.vertex = .{
+        .first_attribute_grf = .{ .number = 4 },
+        .attribute_grf_count = 4,
+    };
+
+    const attribute = try builder.addVirtualRegister(.{
+        .size_bytes = 32,
+        .alignment_bytes = 32,
+        .element_type = .f32,
+        .lane_count = 8,
+        .class = .varying,
+        .name = "attribute",
+    });
+    const entry = try builder.addBlock("entry");
+    _ = try builder.appendInstruction(entry, .simd8, null, .{
+        .load_input = .{
+            .destination = .{ .register = .{ .virtual = attribute }, .type = .f32 },
+            .semantic = .{ .location = .{ .location = 2, .component = 1 } },
+        },
+    });
+
+    for (0..position_component_count) |component| {
+        const position = try builder.addVirtualRegister(.{
+            .size_bytes = 32,
+            .alignment_bytes = 32,
+            .element_type = .f32,
+            .lane_count = 8,
+            .class = .temporary,
+            .name = "position",
+        });
+        _ = try builder.appendInstruction(entry, .simd8, null, .{
+            .store_output = .{
+                .semantic = .{ .builtin = .{ .builtin = .position, .component = @intCast(component) } },
+                .source = .{
+                    .register = .{ .virtual = position },
+                    .type = .f32,
+                    .region = operand.Region.contiguous(.simd8),
+                },
+            },
+        });
+    }
+    return entry;
+}
+
+const test_input_layout = [_]InputComponent{.{
+    .location = 2,
+    .component = 1,
+    .payload_grf_offset = 3,
+}};
+
+const test_layout: Layout = .{
+    .input_components = &test_input_layout,
+    .position_urb_offset = 7,
+};
+
+test "vertex ABI: lower explicit input payload and position URB output" {
+    const device = @import("../device.zig");
+    const printer = @import("../ir/printer.zig");
+
+    const device_info: device.DeviceInfo = .{
+        .generation = .gen9,
+        .platform = .skylake,
+        .pci_device_id = 0x1912,
+        .grf_count = 128,
+    };
+    var program = program_ir.Program.init(std.testing.allocator, .vertex, device_info, .simd8);
+    defer program.deinit();
+    var builder = Builder.init(&program);
+
+    program.properties.common_ir_lowered = true;
+    program.properties.block_parameters_lowered = true;
+    program.payload.vertex = .{
+        .first_attribute_grf = .{ .number = 4 },
+        .attribute_grf_count = 4,
+    };
+
+    const attribute = try builder.addVirtualRegister(.{
+        .size_bytes = 32,
+        .alignment_bytes = 32,
+        .element_type = .f32,
+        .lane_count = 8,
+        .class = .varying,
+        .name = "attribute",
+    });
+    const position_names = [_][]const u8{ "position_x", "position_y", "position_z", "position_w" };
+    var position: [4]ids.VirtualRegisterId = undefined;
+    for (&position, position_names) |*register_id, name| {
+        register_id.* = try builder.addVirtualRegister(.{
+            .size_bytes = 32,
+            .alignment_bytes = 32,
+            .element_type = .f32,
+            .lane_count = 8,
+            .class = .temporary,
+            .name = name,
+        });
+    }
+
+    const entry = try builder.addBlock("entry");
+    _ = try builder.appendInstruction(entry, .simd8, null, .{
+        .load_input = .{
+            .destination = .{ .register = .{ .virtual = attribute }, .type = .f32 },
+            .semantic = .{ .location = .{ .location = 2, .component = 1 } },
+        },
+    });
+    for (position, 0..) |register_id, component| {
+        _ = try builder.appendInstruction(entry, .simd8, null, .{
+            .store_output = .{
+                .semantic = .{ .builtin = .{ .builtin = .position, .component = @intCast(component) } },
+                .source = .{
+                    .register = .{ .virtual = register_id },
+                    .type = .f32,
+                    .region = operand.Region.contiguous(.simd8),
+                },
+            },
+        });
+    }
+    try builder.setTerminator(entry, .end_thread);
+    try validator.validate(&program);
+
+    const input_layout = [_]InputComponent{.{
+        .location = 2,
+        .component = 1,
+        .payload_grf_offset = 3,
+    }};
+    try run(std.testing.allocator, &program, .{
+        .input_components = &input_layout,
+        .position_urb_offset = 7,
+    });
+    try run(std.testing.allocator, &program, .{
+        .input_components = &input_layout,
+        .position_urb_offset = 7,
+    });
+
+    try std.testing.expect(program.properties.stage_io_lowered);
+    try std.testing.expect(!program.properties.messages_lowered);
+    try std.testing.expect(!program.properties.instructions_selected);
+
+    const text = try printer.allocPrint(std.testing.allocator, &program);
+    defer std.testing.allocator.free(text);
+    for ([_][]const u8{
+        "[simd8] mov %attribute:f32, r7:f32",
+        "[simd8] mov %position_urb_payload:f32, %position_x:f32",
+        "[simd8] mov %position_urb_payload:f32[byte=32], %position_y:f32",
+        "[simd8] mov %position_urb_payload:f32[byte=64], %position_z:f32",
+        "[simd8] mov %position_urb_payload:f32[byte=96], %position_w:f32",
+        "send urb_write[offset(7), channels(xyzw), end_of_thread], payload(%position_urb_payload[4])",
+    }) |fragment|
+        try std.testing.expect(std.mem.indexOf(u8, text, fragment) != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "load_input") == null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "store_output") == null);
+}
+
+test "vertex ABI: reject invalid layout and incomplete position" {
+    const device = @import("../device.zig");
+    const device_info: device.DeviceInfo = .{
+        .generation = .gen9,
+        .platform = .skylake,
+        .pci_device_id = 0x1912,
+        .grf_count = 128,
+    };
+
+    var invalid_layout_program = program_ir.Program.init(std.testing.allocator, .vertex, device_info, .simd8);
+    defer invalid_layout_program.deinit();
+    var invalid_layout_builder = Builder.init(&invalid_layout_program);
+    const invalid_layout_entry = try appendTestShaderBody(&invalid_layout_program, 4);
+    try invalid_layout_builder.setTerminator(invalid_layout_entry, .end_thread);
+
+    const out_of_range_input = [_]InputComponent{.{
+        .location = 2,
+        .component = 1,
+        .payload_grf_offset = 4,
+    }};
+    try std.testing.expectError(Error.InvalidLayout, run(std.testing.allocator, &invalid_layout_program, .{
+        .input_components = &out_of_range_input,
+        .position_urb_offset = 7,
+    }));
+
+    const duplicate_inputs = [_]InputComponent{
+        test_input_layout[0],
+        test_input_layout[0],
+    };
+    try std.testing.expectError(Error.InvalidLayout, run(std.testing.allocator, &invalid_layout_program, .{
+        .input_components = &duplicate_inputs,
+        .position_urb_offset = 7,
+    }));
+    try std.testing.expect(!invalid_layout_program.properties.stage_io_lowered);
+
+    var incomplete_program = program_ir.Program.init(std.testing.allocator, .vertex, device_info, .simd8);
+    defer incomplete_program.deinit();
+    var incomplete_builder = Builder.init(&incomplete_program);
+    const incomplete_entry = try appendTestShaderBody(&incomplete_program, 3);
+    try incomplete_builder.setTerminator(incomplete_entry, .end_thread);
+
+    try std.testing.expectError(Error.MissingPosition, run(std.testing.allocator, &incomplete_program, test_layout));
+    try std.testing.expect(!incomplete_program.properties.stage_io_lowered);
+}
+
+test "vertex ABI: reject an existing logical URB write" {
+    const device = @import("../device.zig");
+    const device_info: device.DeviceInfo = .{
+        .generation = .gen9,
+        .platform = .skylake,
+        .pci_device_id = 0x1912,
+        .grf_count = 128,
+    };
+    var program = program_ir.Program.init(std.testing.allocator, .vertex, device_info, .simd8);
+    defer program.deinit();
+    var builder = Builder.init(&program);
+    const entry = try appendTestShaderBody(&program, 4);
+    const payload = try builder.addVirtualRegister(.{
+        .size_bytes = 32,
+        .alignment_bytes = 32,
+        .element_type = .u32,
+        .lane_count = 8,
+        .class = .payload,
+        .spillable = false,
+        .name = "existing_payload",
+    });
+    _ = try builder.appendInstruction(entry, .simd8, null, .{
+        .send = .{
+            .message = .{ .urb_write = .{ .offset = 0 } },
+            .payload = .{
+                .base = .{ .virtual = payload },
+                .register_count = 1,
+            },
+        },
+    });
+    try builder.setTerminator(entry, .end_thread);
+
+    try std.testing.expectError(Error.ExistingUrbWrite, run(std.testing.allocator, &program, test_layout));
+    try std.testing.expect(!program.properties.stage_io_lowered);
+}
+
+test "vertex ABI: append an EOT URB write to every shader exit" {
+    const device = @import("../device.zig");
+    const device_info: device.DeviceInfo = .{
+        .generation = .gen9,
+        .platform = .skylake,
+        .pci_device_id = 0x1912,
+        .grf_count = 128,
+    };
+    var program = program_ir.Program.init(std.testing.allocator, .vertex, device_info, .simd8);
+    defer program.deinit();
+    var builder = Builder.init(&program);
+    const entry = try appendTestShaderBody(&program, 4);
+    const first_exit = try builder.addBlock("first_exit");
+    const second_exit = try builder.addBlock("second_exit");
+    const condition = try builder.addVirtualFlag(.{ .name = "condition" });
+    try builder.setTerminator(entry, .{ .conditional_branch = .{
+        .predicate = .{ .flag = .{ .virtual = condition } },
+        .true_edge = try builder.edge(first_exit, &.{}),
+        .false_edge = try builder.edge(second_exit, &.{}),
+    } });
+    try builder.setTerminator(first_exit, .end_thread);
+    try builder.setTerminator(second_exit, .end_thread);
+
+    try run(std.testing.allocator, &program, test_layout);
+
+    var urb_write_count: usize = 0;
+    for (program.blocks.entries.items) |block_entry| {
+        const block = block_entry orelse continue;
+        for (block.instructions.items) |instruction_id| {
+            const inst = program.instructions.get(instruction_id).?;
+            switch (inst.operation) {
+                .send => |send| switch (send.message) {
+                    .urb_write => |urb_write| {
+                        try std.testing.expect(urb_write.end_of_thread);
+                        try std.testing.expectEqual(@as(u16, 7), urb_write.offset);
+                        urb_write_count += 1;
+                    },
+                },
+                else => {},
+            }
+        }
+    }
+    try std.testing.expectEqual(@as(usize, 2), urb_write_count);
+    try validator.validate(&program);
+}
@@ -15,6 +15,7 @@ pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const v
    errdefer allocator.destroy(self);

    var interface = try Interface.init(device, allocator, info);
+    errdefer interface.deinit();
    interface.vtable = &.{ .destroy = destroy };

    self.* = .{
@@ -30,6 +31,7 @@ pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void {
 }

 pub fn drop(self: *Self, allocator: std.mem.Allocator) void {
+    self.interface.deinit();
    allocator.destroy(self);
 }

@@ -38,7 +40,6 @@ pub fn ref(self: *Self) void {
 }

 pub fn unref(self: *Self, allocator: std.mem.Allocator) void {
-    if (self.ref_count.fetchSub(1, .release) == 1) {
+    if (self.ref_count.fetchSub(1, .acq_rel) == 1)
        self.drop(allocator);
-    }
 }
@@ -5,15 +5,14 @@ const spv = @import("spv");

 const VkError = base.VkError;

-
 const Self = @This();
 pub const Interface = base.ShaderModule;

 interface: Interface,
 module: spv.Module,

-/// Pipelines need SPIR-V module reference so shader module may not
-/// be destroy on call to `vkDestroyShaderModule`
+/// Pipelines need a SPIR-V module reference so the shader module may outlive
+/// the application's `vkDestroyShaderModule` call.
 ref_count: std.atomic.Value(usize),

 pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const vk.ShaderModuleCreateInfo) VkError!*Self {
@@ -21,6 +20,7 @@ pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const v
    errdefer allocator.destroy(self);

    var interface = try Interface.init(device, allocator, info);
+    errdefer interface.deinit();

    const device_allocator = device.device_allocator.allocator();

@@ -28,11 +28,9 @@ pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const v
        .destroy = destroy,
    };

-    const code = info.p_code[0..@divExact(info.code_size, 4)];
-
    self.* = .{
        .interface = interface,
-        .module = spv.Module.init(device_allocator, code, .{
+        .module = spv.Module.init(device_allocator, interface.code(), .{
            .use_simd_vectors_specializations = base.config.soft_shaders_simd,
        }) catch |err| switch (err) {
            spv.Module.ModuleError.OutOfMemory => return VkError.OutOfHostMemory,
@@ -59,9 +57,8 @@ pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void {

 pub fn drop(self: *Self, allocator: std.mem.Allocator) void {
    const device_allocator = self.interface.owner.device_allocator.allocator();
-
    self.module.deinit(device_allocator);
-
+    self.interface.deinit();
    allocator.destroy(self);
 }

@@ -70,7 +67,6 @@ pub fn ref(self: *Self) void {
 }

 pub fn unref(self: *Self, allocator: std.mem.Allocator) void {
-    if (self.ref_count.fetchSub(1, .release) == 1) {
+    if (self.ref_count.fetchSub(1, .acq_rel) == 1)
        self.drop(allocator);
-    }
 }
@@ -1,5 +1,6 @@
 const std = @import("std");
 const vk = @import("vulkan");
+const shader_ir = @import("shader_ir");

 const VkError = @import("error_set.zig").VkError;

@@ -7,8 +8,11 @@ const Device = @import("Device.zig");

 const Self = @This();
 pub const ObjectType: vk.ObjectType = .shader_module;
+pub const IrModule = shader_ir.ir.module.Module;
+pub const InstantiateOptions = shader_ir.spirv.translator.Options;

 owner: *Device,
+source: shader_ir.spirv.SourceModule,

 vtable: *const VTable,

@@ -16,16 +20,40 @@ pub const VTable = struct {
    destroy: *const fn (*Self, std.mem.Allocator) void,
 };

-pub fn init(device: *Device, allocator: std.mem.Allocator, info: *const vk.ShaderModuleCreateInfo) VkError!Self {
-    _ = allocator;
-    _ = info;
+pub fn init(device: *Device, _: std.mem.Allocator, info: *const vk.ShaderModuleCreateInfo) VkError!Self {
+    if (info.code_size % @sizeOf(u32) != 0)
+        return VkError.ValidationFailed;
+
+    const source_allocator = device.device_allocator.allocator();
+    const words = info.p_code[0 .. info.code_size / @sizeOf(u32)];
+    var source = shader_ir.spirv.SourceModule.init(source_allocator, words) catch |err| return switch (err) {
+        error.OutOfMemory => VkError.OutOfHostMemory,
+        else => VkError.ValidationFailed,
+    };
+    errdefer source.deinit(source_allocator);
+
    return .{
        .owner = device,
+        .source = source,
        // SAFETY: the backend assigns the vtable before returning the shader module.
        .vtable = undefined,
    };
 }

+pub fn deinit(self: *Self) void {
+    self.source.deinit(self.owner.device_allocator.allocator());
+}
+
+pub fn code(self: *const Self) []const u32 {
+    return self.source.code();
+}
+
+/// Instantiates one entry point as a fresh, backend-agnostic IR module.
+/// The returned module does not borrow from this shader module.
+pub fn instantiateIr(self: *const Self, allocator: std.mem.Allocator, options: InstantiateOptions) !IrModule {
+    return shader_ir.spirv.translator.instantiate(allocator, &self.source, options);
+}
+
 pub inline fn destroy(self: *Self, allocator: std.mem.Allocator) void {
    self.vtable.destroy(self, allocator);
 }