[Soft] adding new experimental shader interpreter based on new IR

README.md

@@ -235,29 +235,35 @@ To bring forth the driver:
 zig build phi --release=[fast|safe|small]
 ```
 
-To build Phi without a Xeon Phi card, enable host emulation:
+Shouldst thou endeavor to forge Phi in the absence of thy Xeon Phi card, thou must permit the emulation of the host:
 
 ```
 zig build phi -Dphi-host-emulation=true --release=[fast|safe|small]
 ```
 
-In this mode the daemon is compiled natively through Zig's `std.Build` C compilation support and installed as `zig-out/bin/phi_device-host`. The driver exposes one synthetic Phi physical device and communicates with the daemon over an IPv4 TCP socket bound only to `127.0.0.1`. The client socket is implemented directly with Zig's `std.Io.net` API; it does not reuse daemon C transport functions or load `libmicmgmt` or `libscif`.
+Take heed, for this path is forged solely for development; thou shalt find that even the execution of shaders may fail to manifest when thou requirest them most.
 
-The driver first connects to an already-running daemon. If none is listening, it launches its embedded daemon locally and retries the connection. The temporary embedded executable removes its own filesystem entry after launch. The daemon may also be started manually:
+Behold how this creation is wrought: In this state, the daemon is compiled natively by the power of Zig's `std.Build`, and set within `zig-out/bin/phi_device-host`.
+The driver revealeth but one synthetic Phi device, its voice reaching the daemon through an IPv4 TCP socket tethered only to `127.0.0.1`.
+For the client socket, thou shalt find it fashioned directly from the `std.Io.net` API; it escheweth the use of C transport functions and requireth no summoning of `libmicmgmt` or `libscif`.
+
+The driver first endeavoreth to bind itself to a daemon already manifest. If the void remaineth silent, it shall launch its embedded daemon from within, retrying the bond.
+This transient essence shall vanish from the filesystem once its task is begun. Yet, thou art not bound to this way; thou mayest also start the daemon by thine own hand:
 
 ```
 ./zig-out/bin/phi_device-host
 ```
 
-The loopback port defaults to `43616` and may be selected at build time for both the driver and daemon:
+The loopback port resideth by default at `43616`, yet thou mayest decree another during the hour of forging, for both driver and daemon alike:
 
 ```
 zig build phi -Dphi-host-emulation=true -Dphi-emulation-port=43617
 ```
 
-Builds without `-Dphi-host-emulation=true` retain the real Xeon Phi path: the daemon is cross-compiled with the configured `k1om-mpss-linux-gcc`, deployed over SSH/SCP, and communicates through SCIF.
+Shouldst thou forgo the `-Dphi-host-emulation=true` decree, the true path of the Xeon Phi remaineth.
+In this manner, the daemon is cross-compiled by way of `k1om-mpss-linux-gcc`, delivered unto its destination via SSH/SCP, and speaketh through the SCIF.\
 
-And for those who seek another manner of building:
+For those who seek a different method of creation:
 
 ```
 zig build --help

build.zig

@@ -435,7 +435,7 @@ fn customSoft(
     _: *std.Build.Module,
     _: *std.Build.Module,
     base_c_mod: *std.Build.Module,
-    _: *std.Build.Module,
+    shader_ir_mod: *std.Build.Module,
     target: std.Build.ResolvedTarget,
     optimize: std.builtin.OptimizeMode,
     use_llvm: bool,
@@ -448,14 +448,17 @@ fn customSoft(
 
     lib_mod.addImport("soft_c", base_c_mod);
     lib_mod.addImport("spv", spv.module("spv"));
+    lib_mod.addImport("shader_ir", shader_ir_mod);
 
     const single_threaded_option = b.option(bool, "soft-single-threaded", "Single threaded runtime mode") orelse false;
+    const ir_interpreter_option = b.option(bool, "soft-ir-interpreter", "Enable the experimental backend-agnostic IR interpreter") orelse false;
     const shaders_simd_option = b.option(bool, "soft-shader-simd", "Shaders SIMD acceleration") orelse true;
     const compute_dump_early_results_table_option = b.option(u32, "soft-compute-dump-early-results-table", "Dump compute shaders results table before invocation");
     const compute_dump_final_results_table_option = b.option(u32, "soft-compute-dump-final-results-table", "Dump compute shaders results table after invocation");
     const approxiamte_rgb_option = b.option(bool, "soft-approximates-rgb", "Approximate sRGB <-> RGB conversions") orelse true;
 
     options.addOption(bool, "soft_single_threaded", single_threaded_option);
+    options.addOption(bool, "soft_ir_interpreter", ir_interpreter_option);
     options.addOption(bool, "soft_shaders_simd", shaders_simd_option);
     options.addOption(?u32, "soft_compute_dump_early_results_table", compute_dump_early_results_table_option);
     options.addOption(?u32, "soft_compute_dump_final_results_table", compute_dump_final_results_table_option);

src/intel/compiler/lower/vertex_abi.zig

@@ -64,24 +64,28 @@ pub fn run(allocator: std.mem.Allocator, program: *program_ir.Program, layout: L
         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),
+                .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 },
+                    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,
                         },
-                        .source = store.source,
-                    } };
+                    };
                 },
                 else => null,
             };
@@ -96,11 +100,13 @@ pub fn run(allocator: std.mem.Allocator, program: *program_ir.Program, layout: L
             continue;
         _ = builder.appendInstruction(ids.BlockId.fromIndex(block_index), .simd8, null, .{
             .send = .{
-                .message = .{ .urb_write = .{
-                    .offset = layout.position_urb_offset,
-                    .channels = .{},
-                    .end_of_thread = true,
-                } },
+                .message = .{
+                    .urb_write = .{
+                        .offset = layout.position_urb_offset,
+                        .channels = .{},
+                        .end_of_thread = true,
+                    },
+                },
                 .payload = .{
                     .base = .{ .virtual = position_payload },
                     .register_count = 4,

src/software/SoftPipeline.zig

@@ -38,6 +38,7 @@ const SoftInstance = @import("SoftInstance.zig");
 const SoftSampler = @import("SoftSampler.zig");
 const SoftShaderModule = @import("SoftShaderModule.zig");
 const SoftPipelineCache = @import("SoftPipelineCache.zig");
+const InterpreterShader = @import("interpreter/Shader.zig");
 
 const Self = @This();
 pub const Interface = base.Pipeline;
@@ -51,6 +52,7 @@ const Shader = struct {
     module: *SoftShaderModule,
     runtimes: []Runtime,
     entry: []const u8,
+    interpreter: ?InterpreterShader,
 };
 
 const Stages = enum {
@@ -185,6 +187,8 @@ pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void {
 
     var it = self.stages.iterator();
     while (it.next()) |entry| {
+        if (entry.value.interpreter) |*interpreter|
+            interpreter.deinit();
         entry.value.module.unref(allocator);
         for (entry.value.runtimes) |*runtime| {
             runtime.rt.function_stack.clearAndFree(device_allocator); // Hacky to avoid leaks
@@ -255,11 +259,15 @@ fn createShader(
         }
     }
 
-    return .{
+    var shader: Shader = .{
         .module = module,
         .runtimes = runtimes,
         .entry = runtimes_allocator.dupe(u8, entry) catch return VkError.OutOfDeviceMemory,
+        .interpreter = null,
     };
+    if (comptime base.config.soft_ir_interpreter)
+        shader.interpreter = try InterpreterShader.compile(runtimes_allocator, module, stage, runtimes_count);
+    return shader;
 }
 
 fn initRuntime(allocator: std.mem.Allocator, module: *SoftShaderModule, stage: *const vk.PipelineShaderStageCreateInfo, image_api: spv.Runtime.ImageAPI) VkError!spv.Runtime {

src/software/device/ComputeDispatcher.zig

@@ -7,6 +7,7 @@ const PipelineState = ExecutionDevice.PipelineState;
 
 const SoftDevice = @import("../SoftDevice.zig");
 const SoftPipeline = @import("../SoftPipeline.zig");
+const ir_compute = @import("../interpreter/compute.zig");
 
 const VkError = base.VkError;
 const SpvRuntimeError = spv.Runtime.RuntimeError;
@@ -70,6 +71,10 @@ pub fn dispatchBase(self: *Self, base_group_x: u32, base_group_y: u32, base_grou
 
     const pipeline = self.state.pipeline orelse return VkError.InvalidPipelineDrv;
     const shader = pipeline.stages.getPtr(.compute) orelse return VkError.InvalidPipelineDrv;
+    if (comptime base.config.soft_ir_interpreter) {
+        if (shader.interpreter) |*interpreter_shader|
+            return ir_compute.dispatch(interpreter_shader, base_group_x, base_group_y, base_group_z, group_count_x, group_count_y, group_count_z);
+    }
     const spv_module = &shader.module.module;
     self.batch_size = if (spv_module.reflection_infos.has_atomics) 1 else shader.runtimes.len;
 

src/software/device/vertex_dispatcher.zig

@@ -10,6 +10,7 @@ const SpvRuntimeError = spv.Runtime.RuntimeError;
 const Renderer = @import("Renderer.zig");
 const SoftPipeline = @import("../SoftPipeline.zig");
 const blitter = @import("blitter.zig");
+const ir_vertex = @import("../interpreter/vertex.zig");
 
 const VkError = base.VkError;
 const interface_blob_padding = @sizeOf(F32x4);
@@ -41,6 +42,28 @@ pub fn runWrapper(data: RunData) void {
 
 inline fn run(data: RunData) !void {
     const shader = data.pipeline.stages.getPtrAssertContains(.vertex);
+    if (comptime base.config.soft_ir_interpreter) {
+        // Interpolation decorations are not represented in the common IR yet,
+        // so fragment-linked graphics pipelines retain the SPIR-V path.
+        if (data.pipeline.stages.getPtr(.fragment) == null) {
+            if (shader.interpreter) |*interpreter_shader| {
+                return ir_vertex.run(
+                    data.allocator,
+                    data.pipeline,
+                    interpreter_shader,
+                    data.batch_id,
+                    data.batch_size,
+                    data.vertex_count,
+                    data.first_vertex,
+                    data.first_instance,
+                    data.indices,
+                    data.primitive_restart,
+                    data.instance_index,
+                    data.draw_call,
+                );
+            }
+        }
+    }
     const runtime = &shader.runtimes[data.batch_id];
     const mutex = &runtime.mutex;
     const rt = &runtime.rt;

src/software/interpreter/Program.zig

@@ -0,0 +1,498 @@
+const std = @import("std");
+const shader_ir = @import("shader_ir");
+const bc = @import("bytecode.zig");
+
+const ir = shader_ir.ir;
+const ids = ir.id;
+const inst_ir = ir.instruction;
+const module_ir = ir.module;
+
+pub const CompileError = error{
+    InvalidConstant,
+    InvalidControlFlow,
+    InvalidInterface,
+    InvalidOperation,
+    InvalidValue,
+    TooManyInstructions,
+    TooManyRegisters,
+    UnsupportedOperation,
+    UnsupportedType,
+};
+
+pub const InterfaceBinding = struct {
+    direction: module_ir.InterfaceDirection,
+    semantic: module_ir.InterfaceSemantic,
+    span: bc.Span,
+};
+
+pub const RegisterInit = struct {
+    register: bc.Register,
+    value: u32,
+};
+
+const Self = @This();
+
+arena: std.heap.ArenaAllocator,
+stage: module_ir.Stage,
+entry_pc: u32,
+register_count: usize,
+scratch_count: usize,
+code: []const bc.Instruction,
+edges: []const bc.Edge,
+copies: []const bc.Copy,
+branches: []const bc.Branch,
+initializers: []const RegisterInit,
+interfaces: []const ?InterfaceBinding,
+
+pub fn compile(backing_allocator: std.mem.Allocator, module: *const module_ir.Module) !Self {
+    try ir.validator.validate(module);
+
+    var result: Self = undefined;
+    result.arena = std.heap.ArenaAllocator.init(backing_allocator);
+    errdefer result.arena.deinit();
+
+    var lowerer = try Lowerer.init(result.arena.allocator(), module);
+    try lowerer.lower();
+
+    result.stage = module.stage;
+    result.entry_pc = lowerer.entry_pc;
+    result.register_count = lowerer.register_count;
+    result.scratch_count = lowerer.scratch_count;
+    result.code = lowerer.code.items;
+    result.edges = lowerer.edges.items;
+    result.copies = lowerer.copies.items;
+    result.branches = lowerer.branches.items;
+    result.initializers = lowerer.initializers.items;
+    result.interfaces = lowerer.interfaces;
+    return result;
+}
+
+pub fn deinit(self: *Self) void {
+    self.arena.deinit();
+    self.* = undefined;
+}
+
+pub fn interfaceBinding(self: *const Self, variable: ids.InterfaceVariableId) ?InterfaceBinding {
+    if (variable.index() >= self.interfaces.len)
+        return null;
+
+    return self.interfaces[variable.index()];
+}
+
+const Lowerer = struct {
+    allocator: std.mem.Allocator,
+    module: *const module_ir.Module,
+    function: *const module_ir.Function,
+    entry_block: ids.BlockId,
+    values: []?bc.Span,
+    interfaces: []?InterfaceBinding,
+    block_pcs: []?u32,
+    register_count: usize = 0,
+    scratch_count: usize = 0,
+    entry_pc: u32 = 0,
+    code: std.ArrayList(bc.Instruction) = .empty,
+    edges: std.ArrayList(bc.Edge) = .empty,
+    copies: std.ArrayList(bc.Copy) = .empty,
+    branches: std.ArrayList(bc.Branch) = .empty,
+    initializers: std.ArrayList(RegisterInit) = .empty,
+
+    fn init(allocator: std.mem.Allocator, module: *const module_ir.Module) !Lowerer {
+        const entry_id = module.entry_point orelse return CompileError.InvalidControlFlow;
+        const function = module.functions.get(entry_id) orelse return CompileError.InvalidControlFlow;
+        const entry_block = function.entry_block orelse return CompileError.InvalidControlFlow;
+
+        if (function.parameters.items.len != 0)
+            return CompileError.UnsupportedOperation;
+
+        const values = try allocator.alloc(?bc.Span, module.values.entries.items.len);
+        @memset(values, null);
+        const interfaces = try allocator.alloc(?InterfaceBinding, module.interface_variables.entries.items.len);
+        @memset(interfaces, null);
+        const block_pcs = try allocator.alloc(?u32, module.blocks.entries.items.len);
+        @memset(block_pcs, null);
+
+        return .{
+            .allocator = allocator,
+            .module = module,
+            .function = function,
+            .entry_block = entry_block,
+            .values = values,
+            .interfaces = interfaces,
+            .block_pcs = block_pcs,
+        };
+    }
+
+    fn lower(self: *Lowerer) !void {
+        for (self.module.values.entries.items, 0..) |entry, index| {
+            const value = entry orelse continue;
+            self.values[index] = try self.allocate(value.type);
+        }
+        for (self.module.interface_variables.entries.items, 0..) |entry, index| {
+            const variable = entry orelse continue;
+            self.interfaces[index] = .{
+                .direction = variable.direction,
+                .semantic = variable.semantic,
+                .span = try self.allocate(variable.type),
+            };
+        }
+        try self.initializeConstants();
+
+        for (self.function.blocks.items) |block_id| {
+            const block = self.module.blocks.get(block_id) orelse return CompileError.InvalidControlFlow;
+            self.block_pcs[block_id.index()] = try u32Index(self.code.items.len);
+            for (block.instructions.items) |instruction_id| {
+                const instruction = self.module.instructions.get(instruction_id) orelse return CompileError.InvalidOperation;
+                try self.lowerInstruction(instruction);
+            }
+            try self.lowerTerminator(block.terminator orelse return CompileError.InvalidControlFlow);
+        }
+
+        for (self.edges.items) |*edge| {
+            if (edge.target_block >= self.block_pcs.len)
+                return CompileError.InvalidControlFlow;
+
+            edge.target_pc = self.block_pcs[edge.target_block] orelse return CompileError.InvalidControlFlow;
+        }
+        self.entry_pc = self.block_pcs[self.entry_block.index()] orelse return CompileError.InvalidControlFlow;
+    }
+
+    fn allocate(self: *Lowerer, type_id: ids.TypeId) !bc.Span {
+        const ty = self.module.types.get(type_id) orelse return CompileError.UnsupportedType;
+        var kind: bc.ValueKind = undefined;
+        var components: u8 = 1;
+
+        switch (ty.*) {
+            .boolean => kind = .boolean,
+            .integer => |integer| {
+                if (integer.bits != 32)
+                    return CompileError.UnsupportedType;
+
+                kind = if (integer.signedness == .signed) .signed_integer else .unsigned_integer;
+            },
+            .floating => |floating| {
+                if (floating.bits != 32)
+                    return CompileError.UnsupportedType;
+
+                kind = .floating;
+            },
+            .vector => |vector| {
+                const element = self.module.types.get(vector.element_type) orelse return CompileError.UnsupportedType;
+                components = vector.length;
+
+                kind = switch (element.*) {
+                    .boolean => .boolean,
+                    .integer => |integer| blk: {
+                        if (integer.bits != 32)
+                            return CompileError.UnsupportedType;
+
+                        break :blk if (integer.signedness == .signed) .signed_integer else .unsigned_integer;
+                    },
+                    .floating => |floating| if (floating.bits == 32) .floating else return CompileError.UnsupportedType,
+                    else => return CompileError.UnsupportedType,
+                };
+            },
+            else => return CompileError.UnsupportedType,
+        }
+        const end = std.math.add(usize, self.register_count, components) catch return CompileError.TooManyRegisters;
+
+        if (end > @as(usize, std.math.maxInt(bc.Register)) + 1)
+            return CompileError.TooManyRegisters;
+
+        const allocated: bc.Span = .{ .base = @intCast(self.register_count), .components = components, .kind = kind };
+        self.register_count = end;
+        return allocated;
+    }
+
+    fn initializeConstants(self: *Lowerer) !void {
+        for (self.module.values.entries.items, self.values) |entry, destination| {
+            const value = entry orelse continue;
+            if (value.definition == .constant)
+                try self.initializeConstant(destination orelse return CompileError.InvalidValue, value.definition.constant);
+        }
+    }
+
+    fn initializeConstant(self: *Lowerer, destination: bc.Span, constant_id: ids.ConstantId) !void {
+        const constant = self.module.constants.get(constant_id) orelse return CompileError.InvalidConstant;
+        switch (constant.value) {
+            .boolean => |value| {
+                if (destination.components != 1 or destination.kind != .boolean)
+                    return CompileError.InvalidConstant;
+
+                try self.initializers.append(self.allocator, .{ .register = destination.base, .value = @intFromBool(value) });
+            },
+            .integer_bits, .float_bits => |value| {
+                if (destination.components != 1)
+                    return CompileError.InvalidConstant;
+
+                try self.initializers.append(self.allocator, .{ .register = destination.base, .value = @truncate(value) });
+            },
+            .null, .undef => for (0..destination.components) |component|
+                try self.initializers.append(self.allocator, .{ .register = try offset(destination.base, component), .value = 0 }),
+            .composite => |elements| {
+                if (elements.len != destination.components)
+                    return CompileError.InvalidConstant;
+
+                for (elements, 0..) |element, component| {
+                    try self.initializeConstant(.{
+                        .base = try offset(destination.base, component),
+                        .components = 1,
+                        .kind = destination.kind,
+                    }, element);
+                }
+            },
+        }
+    }
+
+    fn lowerInstruction(self: *Lowerer, instruction: *const inst_ir.Instruction) !void {
+        const result = if (instruction.result) |id| try self.span(id) else null;
+        switch (instruction.operation) {
+            .unary => |op| {
+                const dst = result orelse return CompileError.InvalidOperation;
+                const src = try self.span(op.operand);
+
+                if (!dst.sameShape(src))
+                    return CompileError.InvalidOperation;
+
+                const opcode: bc.Opcode = switch (op.opcode) {
+                    .negate => switch (dst.kind) {
+                        .signed_integer => .negate_i32,
+                        .floating => .negate_f32,
+                        else => return CompileError.InvalidOperation,
+                    },
+                    .logical_not => if (dst.kind == .boolean) .logical_not else return CompileError.InvalidOperation,
+                    .bitwise_not => switch (dst.kind) {
+                        .signed_integer, .unsigned_integer => .bitwise_not,
+                        else => return CompileError.InvalidOperation,
+                    },
+                };
+                try self.emit(opcode, dst.components, dst.base, src.base, bc.invalid_register, bc.invalid_register, 0);
+            },
+            .binary => |op| {
+                const dst = result orelse return CompileError.InvalidOperation;
+                const lhs = try self.span(op.lhs);
+                const rhs = try self.span(op.rhs);
+
+                if (!dst.sameShape(lhs) or !dst.sameShape(rhs))
+                    return CompileError.InvalidOperation;
+
+                try self.emit(try binaryOpcode(op.opcode, dst.kind), dst.components, dst.base, lhs.base, rhs.base, bc.invalid_register, 0);
+            },
+            .compare => |op| {
+                const dst = result orelse return CompileError.InvalidOperation;
+                const lhs = try self.span(op.lhs);
+                const rhs = try self.span(op.rhs);
+
+                if (dst.kind != .boolean or dst.components != 1 or lhs.components != 1 or !lhs.sameShape(rhs))
+                    return CompileError.UnsupportedOperation;
+
+                try self.emit(try compareOpcode(op.opcode, lhs.kind), 1, dst.base, lhs.base, rhs.base, bc.invalid_register, 0);
+            },
+            .select => |op| {
+                const dst = result orelse return CompileError.InvalidOperation;
+                const condition = try self.span(op.condition);
+                const yes = try self.span(op.true_value);
+                const no = try self.span(op.false_value);
+
+                if (condition.kind != .boolean or condition.components != 1 or !dst.sameShape(yes) or !dst.sameShape(no))
+                    return CompileError.InvalidOperation;
+
+                try self.emit(.select, dst.components, dst.base, condition.base, yes.base, no.base, 0);
+            },
+            .bitcast => |id| {
+                const dst = result orelse return CompileError.InvalidOperation;
+                const src = try self.span(id);
+
+                if (dst.components != src.components)
+                    return CompileError.UnsupportedOperation;
+
+                try self.emitCopy(dst, src);
+            },
+            .composite_construct => |op| {
+                const dst = result orelse return CompileError.InvalidOperation;
+
+                if (dst.components != op.elements.len)
+                    return CompileError.UnsupportedOperation;
+
+                for (op.elements, 0..) |id, component| {
+                    const src = try self.span(id);
+
+                    if (src.components != 1)
+                        return CompileError.UnsupportedOperation;
+
+                    try self.emit(.copy, 1, try offset(dst.base, component), src.base, bc.invalid_register, bc.invalid_register, 0);
+                }
+            },
+            .composite_extract => |op| {
+                const dst = result orelse return CompileError.InvalidOperation;
+                const src = try self.span(op.composite);
+
+                if (dst.components != 1 or op.indices.len != 1 or op.indices[0] >= src.components)
+                    return CompileError.UnsupportedOperation;
+
+                try self.emit(.copy, 1, dst.base, try offset(src.base, op.indices[0]), bc.invalid_register, bc.invalid_register, 0);
+            },
+            .load_interface => |op| {
+                if (op.element_index != null)
+                    return CompileError.UnsupportedOperation;
+
+                const dst = result orelse return CompileError.InvalidOperation;
+                const binding = self.interfaceFor(op.variable) orelse return CompileError.InvalidInterface;
+
+                if (binding.direction != .input or !dst.sameShape(binding.span))
+                    return CompileError.InvalidInterface;
+
+                try self.emitCopy(dst, binding.span);
+            },
+            .store_interface => |op| {
+                if (op.element_index != null)
+                    return CompileError.UnsupportedOperation;
+
+                const binding = self.interfaceFor(op.variable) orelse return CompileError.InvalidInterface;
+                const src = try self.span(op.value);
+
+                if (binding.direction != .output or !binding.span.sameShape(src))
+                    return CompileError.InvalidInterface;
+
+                try self.emitCopy(binding.span, src);
+            },
+            .call => return CompileError.UnsupportedOperation,
+        }
+    }
+
+    fn lowerTerminator(self: *Lowerer, terminator: module_ir.Terminator) !void {
+        switch (terminator) {
+            .branch => |edge| try self.emit(.jump_edge, 1, bc.invalid_register, bc.invalid_register, bc.invalid_register, bc.invalid_register, try self.addEdge(edge)),
+            .conditional_branch => |branch| {
+                const condition = try self.span(branch.condition);
+                if (condition.kind != .boolean or condition.components != 1) return CompileError.InvalidControlFlow;
+                const index = try u32Index(self.branches.items.len);
+                try self.branches.append(self.allocator, .{
+                    .true_edge = try self.addEdge(branch.true_edge),
+                    .false_edge = try self.addEdge(branch.false_edge),
+                });
+                try self.emit(.branch, 1, condition.base, bc.invalid_register, bc.invalid_register, bc.invalid_register, index);
+            },
+            .return_void => try self.emit(.return_void, 1, bc.invalid_register, bc.invalid_register, bc.invalid_register, bc.invalid_register, 0),
+            .return_value => return CompileError.UnsupportedOperation,
+            .discard => try self.emit(.discard, 1, bc.invalid_register, bc.invalid_register, bc.invalid_register, bc.invalid_register, 0),
+            .@"unreachable" => try self.emit(.@"unreachable", 1, bc.invalid_register, bc.invalid_register, bc.invalid_register, bc.invalid_register, 0),
+        }
+    }
+
+    fn addEdge(self: *Lowerer, edge: module_ir.Edge) !u32 {
+        const target = self.module.blocks.get(edge.target) orelse return CompileError.InvalidControlFlow;
+
+        if (target.parameters.items.len != edge.arguments.len)
+            return CompileError.InvalidControlFlow;
+
+        const first = try u32Index(self.copies.items.len);
+        var scratch: usize = 0;
+        for (edge.arguments, target.parameters.items) |source_id, destination_id| {
+            const source = try self.span(source_id);
+            const destination = try self.span(destination_id);
+            if (!source.sameShape(destination)) return CompileError.InvalidControlFlow;
+            if (source.base == destination.base) continue;
+            try self.copies.append(self.allocator, .{
+                .destination = destination.base,
+                .source = source.base,
+                .components = source.components,
+                .scratch_base = @intCast(scratch),
+            });
+            scratch += source.components;
+        }
+
+        if (scratch > std.math.maxInt(bc.Register))
+            return CompileError.TooManyRegisters;
+
+        self.scratch_count = @max(self.scratch_count, scratch);
+        const copy_count = self.copies.items.len - first;
+
+        if (copy_count > std.math.maxInt(u16))
+            return CompileError.TooManyInstructions;
+
+        const index = try u32Index(self.edges.items.len);
+        try self.edges.append(self.allocator, .{
+            .target_block = @intFromEnum(edge.target),
+            .first_copy = first,
+            .copy_count = @intCast(copy_count),
+        });
+        return index;
+    }
+
+    fn emitCopy(self: *Lowerer, dst: bc.Span, src: bc.Span) !void {
+        if (dst.components != src.components)
+            return CompileError.InvalidOperation;
+        try self.emit(.copy, dst.components, dst.base, src.base, bc.invalid_register, bc.invalid_register, 0);
+    }
+
+    fn emit(self: *Lowerer, opcode: bc.Opcode, components: u16, a: bc.Register, b: bc.Register, c: bc.Register, d: bc.Register, immediate: u32) !void {
+        if (self.code.items.len >= std.math.maxInt(u32))
+            return CompileError.TooManyInstructions;
+        try self.code.append(self.allocator, .{ .opcode = opcode, .components = components, .a = a, .b = b, .c = c, .d = d, .immediate = immediate });
+    }
+
+    fn span(self: *const Lowerer, id: ids.ValueId) !bc.Span {
+        if (id.index() >= self.values.len)
+            return CompileError.InvalidValue;
+        return self.values[id.index()] orelse CompileError.InvalidValue;
+    }
+
+    fn interfaceFor(self: *const Lowerer, id: ids.InterfaceVariableId) ?InterfaceBinding {
+        if (id.index() >= self.interfaces.len)
+            return null;
+        return self.interfaces[id.index()];
+    }
+};
+
+fn binaryOpcode(op: inst_ir.BinaryOpcode, kind: bc.ValueKind) !bc.Opcode {
+    return switch (op) {
+        .integer_add => if (kind == .signed_integer or kind == .unsigned_integer) .integer_add else CompileError.InvalidOperation,
+        .integer_subtract => if (kind == .signed_integer or kind == .unsigned_integer) .integer_subtract else CompileError.InvalidOperation,
+        .integer_multiply => if (kind == .signed_integer or kind == .unsigned_integer) .integer_multiply else CompileError.InvalidOperation,
+        .unsigned_divide => if (kind == .unsigned_integer) .unsigned_divide else CompileError.InvalidOperation,
+        .signed_divide => if (kind == .signed_integer) .signed_divide else CompileError.InvalidOperation,
+        .unsigned_modulo => if (kind == .unsigned_integer) .unsigned_modulo else CompileError.InvalidOperation,
+        .signed_modulo => if (kind == .signed_integer) .signed_modulo else CompileError.InvalidOperation,
+        .float_add => if (kind == .floating) .float_add else CompileError.InvalidOperation,
+        .float_subtract => if (kind == .floating) .float_subtract else CompileError.InvalidOperation,
+        .float_multiply => if (kind == .floating) .float_multiply else CompileError.InvalidOperation,
+        .float_divide => if (kind == .floating) .float_divide else CompileError.InvalidOperation,
+        .float_modulo => if (kind == .floating) .float_modulo else CompileError.InvalidOperation,
+        .shift_left => if (kind == .signed_integer or kind == .unsigned_integer) .shift_left else CompileError.InvalidOperation,
+        .logical_shift_right => if (kind == .signed_integer or kind == .unsigned_integer) .logical_shift_right else CompileError.InvalidOperation,
+        .arithmetic_shift_right => if (kind == .signed_integer) .arithmetic_shift_right else CompileError.InvalidOperation,
+        .bitwise_and => if (kind == .signed_integer or kind == .unsigned_integer) .bitwise_and else CompileError.InvalidOperation,
+        .bitwise_or => if (kind == .signed_integer or kind == .unsigned_integer) .bitwise_or else CompileError.InvalidOperation,
+        .bitwise_xor => if (kind == .signed_integer or kind == .unsigned_integer) .bitwise_xor else CompileError.InvalidOperation,
+        .logical_and => if (kind == .boolean) .logical_and else CompileError.InvalidOperation,
+        .logical_or => if (kind == .boolean) .logical_or else CompileError.InvalidOperation,
+    };
+}
+
+fn compareOpcode(op: inst_ir.CompareOpcode, kind: bc.ValueKind) !bc.Opcode {
+    return switch (op) {
+        .equal => if (kind == .signed_integer or kind == .unsigned_integer or kind == .boolean) .compare_equal else CompileError.InvalidOperation,
+        .not_equal => if (kind == .signed_integer or kind == .unsigned_integer or kind == .boolean) .compare_not_equal else CompileError.InvalidOperation,
+        .unsigned_less => if (kind == .unsigned_integer) .compare_unsigned_less else CompileError.InvalidOperation,
+        .signed_less => if (kind == .signed_integer) .compare_signed_less else CompileError.InvalidOperation,
+        .ordered_float_equal => if (kind == .floating) .compare_ordered_float_equal else CompileError.InvalidOperation,
+        .unordered_float_equal => if (kind == .floating) .compare_unordered_float_equal else CompileError.InvalidOperation,
+        .ordered_float_not_equal => if (kind == .floating) .compare_ordered_float_not_equal else CompileError.InvalidOperation,
+        .unordered_float_not_equal => if (kind == .floating) .compare_unordered_float_not_equal else CompileError.InvalidOperation,
+        .ordered_float_less => if (kind == .floating) .compare_ordered_float_less else CompileError.InvalidOperation,
+        .unordered_float_less => if (kind == .floating) .compare_unordered_float_less else CompileError.InvalidOperation,
+    };
+}
+
+fn offset(base: bc.Register, component: usize) !bc.Register {
+    const value = @as(usize, base) + component;
+    if (value > std.math.maxInt(bc.Register))
+        return CompileError.TooManyRegisters;
+    return @intCast(value);
+}
+
+fn u32Index(value: usize) !u32 {
+    if (value > std.math.maxInt(u32))
+        return CompileError.TooManyInstructions;
+    return @intCast(value);
+}

src/software/interpreter/Runtime.zig

@@ -0,0 +1,307 @@
+const std = @import("std");
+const shader_ir = @import("shader_ir");
+const bc = @import("bytecode.zig");
+const Program = @import("Program.zig");
+
+const ids = shader_ir.ir.id;
+
+pub const RuntimeError = error{
+    DivisionByZero,
+    IntegerOverflow,
+    InvalidBytecode,
+    InvalidInterface,
+    ShiftOutOfRange,
+    StepLimitExceeded,
+    UnreachableExecuted,
+    WrongInterfaceDirection,
+    WrongInterfaceType,
+};
+
+pub const Outcome = enum {
+    returned,
+    discarded,
+};
+
+pub const RunOptions = struct {
+    max_steps: usize = 1_000_000,
+};
+
+const Self = @This();
+
+allocator: std.mem.Allocator,
+registers: []u32,
+scratch: []u32,
+
+pub fn init(allocator: std.mem.Allocator, program: *const Program) !Self {
+    const registers = try allocator.alloc(u32, program.register_count);
+    errdefer allocator.free(registers);
+    const scratch = try allocator.alloc(u32, program.scratch_count);
+    errdefer allocator.free(scratch);
+    @memset(registers, 0);
+    @memset(scratch, 0);
+    for (program.initializers) |initializer|
+        registers[initializer.register] = initializer.value;
+    return .{ .allocator = allocator, .registers = registers, .scratch = scratch };
+}
+
+pub fn deinit(self: *Self) void {
+    self.allocator.free(self.registers);
+    self.allocator.free(self.scratch);
+    self.* = undefined;
+}
+
+pub fn writeInput(self: *Self, program: *const Program, variable: ids.InterfaceVariableId, values: []const u32) RuntimeError!void {
+    const binding = program.interfaceBinding(variable) orelse return RuntimeError.InvalidInterface;
+
+    if (binding.direction != .input)
+        return RuntimeError.WrongInterfaceDirection;
+
+    if (binding.span.components != values.len)
+        return RuntimeError.WrongInterfaceType;
+
+    @memcpy(self.registers[binding.span.base..][0..values.len], values);
+}
+
+pub fn readOutput(self: *const Self, program: *const Program, variable: ids.InterfaceVariableId, values: []u32) RuntimeError!void {
+    const binding = program.interfaceBinding(variable) orelse return RuntimeError.InvalidInterface;
+
+    if (binding.direction != .output)
+        return RuntimeError.WrongInterfaceDirection;
+
+    if (binding.span.components != values.len)
+        return RuntimeError.WrongInterfaceType;
+
+    @memcpy(values, self.registers[binding.span.base..][0..values.len]);
+}
+
+pub fn run(self: *Self, program: *const Program, options: RunOptions) RuntimeError!Outcome {
+    var pc = program.entry_pc;
+    var steps: usize = 0;
+    while (true) {
+        if (steps >= options.max_steps)
+            return RuntimeError.StepLimitExceeded;
+
+        steps += 1;
+
+        if (pc >= program.code.len)
+            return RuntimeError.InvalidBytecode;
+
+        const instruction = program.code[pc];
+        pc += 1;
+
+        switch (instruction.opcode) {
+            .copy => self.copy(instruction),
+            .negate_i32 => self.unaryInt(instruction, .negate),
+            .negate_f32 => self.unaryFloat(instruction),
+            .logical_not => self.unaryInt(instruction, .logical_not),
+            .bitwise_not => self.unaryInt(instruction, .bitwise_not),
+            .integer_add => try self.binaryInt(instruction, .add),
+            .integer_subtract => try self.binaryInt(instruction, .subtract),
+            .integer_multiply => try self.binaryInt(instruction, .multiply),
+            .unsigned_divide => try self.binaryInt(instruction, .unsigned_divide),
+            .signed_divide => try self.binaryInt(instruction, .signed_divide),
+            .unsigned_modulo => try self.binaryInt(instruction, .unsigned_modulo),
+            .signed_modulo => try self.binaryInt(instruction, .signed_modulo),
+            .shift_left => try self.binaryInt(instruction, .shift_left),
+            .logical_shift_right => try self.binaryInt(instruction, .logical_shift_right),
+            .arithmetic_shift_right => try self.binaryInt(instruction, .arithmetic_shift_right),
+            .bitwise_and => try self.binaryInt(instruction, .bitwise_and),
+            .bitwise_or => try self.binaryInt(instruction, .bitwise_or),
+            .bitwise_xor => try self.binaryInt(instruction, .bitwise_xor),
+            .logical_and => try self.binaryInt(instruction, .logical_and),
+            .logical_or => try self.binaryInt(instruction, .logical_or),
+            .float_add => self.binaryFloat(instruction, .add),
+            .float_subtract => self.binaryFloat(instruction, .subtract),
+            .float_multiply => self.binaryFloat(instruction, .multiply),
+            .float_divide => self.binaryFloat(instruction, .divide),
+            .float_modulo => self.binaryFloat(instruction, .modulo),
+            .compare_equal => self.compareInt(instruction, .equal),
+            .compare_not_equal => self.compareInt(instruction, .not_equal),
+            .compare_unsigned_less => self.compareInt(instruction, .unsigned_less),
+            .compare_signed_less => self.compareInt(instruction, .signed_less),
+            .compare_ordered_float_equal => self.compareFloat(instruction, .ordered_equal),
+            .compare_unordered_float_equal => self.compareFloat(instruction, .unordered_equal),
+            .compare_ordered_float_not_equal => self.compareFloat(instruction, .ordered_not_equal),
+            .compare_unordered_float_not_equal => self.compareFloat(instruction, .unordered_not_equal),
+            .compare_ordered_float_less => self.compareFloat(instruction, .ordered_less),
+            .compare_unordered_float_less => self.compareFloat(instruction, .unordered_less),
+            .select => self.select(instruction),
+            .jump_edge => pc = try self.applyEdge(program, instruction.immediate),
+            .branch => {
+                if (instruction.immediate >= program.branches.len)
+                    return RuntimeError.InvalidBytecode;
+
+                const branch = program.branches[instruction.immediate];
+                pc = try self.applyEdge(program, if (self.registers[instruction.a] != 0) branch.true_edge else branch.false_edge);
+            },
+            .return_void => return .returned,
+            .discard => return .discarded,
+            .@"unreachable" => return RuntimeError.UnreachableExecuted,
+        }
+    }
+}
+
+const UnaryInt = enum { negate, logical_not, bitwise_not };
+const BinaryInt = enum {
+    add,
+    subtract,
+    multiply,
+    unsigned_divide,
+    signed_divide,
+    unsigned_modulo,
+    signed_modulo,
+    shift_left,
+    logical_shift_right,
+    arithmetic_shift_right,
+    bitwise_and,
+    bitwise_or,
+    bitwise_xor,
+    logical_and,
+    logical_or,
+};
+const BinaryFloat = enum { add, subtract, multiply, divide, modulo };
+const CompareInt = enum { equal, not_equal, unsigned_less, signed_less };
+const CompareFloat = enum { ordered_equal, unordered_equal, ordered_not_equal, unordered_not_equal, ordered_less, unordered_less };
+
+fn copy(self: *Self, instruction: bc.Instruction) void {
+    for (0..instruction.components) |component|
+        self.registers[@as(usize, instruction.a) + component] = self.registers[@as(usize, instruction.b) + component];
+}
+
+fn unaryInt(self: *Self, instruction: bc.Instruction, comptime operation: UnaryInt) void {
+    for (0..instruction.components) |component| {
+        const value = self.registers[@as(usize, instruction.b) + component];
+        self.registers[@as(usize, instruction.a) + component] = switch (operation) {
+            .negate => 0 -% value,
+            .logical_not => @intFromBool(value == 0),
+            .bitwise_not => ~value,
+        };
+    }
+}
+
+fn unaryFloat(self: *Self, instruction: bc.Instruction) void {
+    for (0..instruction.components) |component| {
+        const value: f32 = @bitCast(self.registers[@as(usize, instruction.b) + component]);
+        self.registers[@as(usize, instruction.a) + component] = @bitCast(-value);
+    }
+}
+
+fn binaryInt(self: *Self, instruction: bc.Instruction, comptime operation: BinaryInt) RuntimeError!void {
+    for (0..instruction.components) |component| {
+        const lhs = self.registers[@as(usize, instruction.b) + component];
+        const rhs = self.registers[@as(usize, instruction.c) + component];
+        self.registers[@as(usize, instruction.a) + component] = switch (operation) {
+            .add => lhs +% rhs,
+            .subtract => lhs -% rhs,
+            .multiply => lhs *% rhs,
+            .unsigned_divide => if (rhs == 0) return RuntimeError.DivisionByZero else lhs / rhs,
+            .unsigned_modulo => if (rhs == 0) return RuntimeError.DivisionByZero else lhs % rhs,
+            .signed_divide => blk: {
+                const signed_lhs: i32 = @bitCast(lhs);
+                const signed_rhs: i32 = @bitCast(rhs);
+
+                if (signed_rhs == 0)
+                    return RuntimeError.DivisionByZero;
+
+                if (signed_lhs == std.math.minInt(i32) and signed_rhs == -1)
+                    return RuntimeError.IntegerOverflow;
+
+                break :blk @bitCast(@divTrunc(signed_lhs, signed_rhs));
+            },
+            .signed_modulo => blk: {
+                const signed_lhs: i32 = @bitCast(lhs);
+                const signed_rhs: i32 = @bitCast(rhs);
+
+                if (signed_rhs == 0)
+                    return RuntimeError.DivisionByZero;
+
+                if (signed_lhs == std.math.minInt(i32) and signed_rhs == -1)
+                    break :blk 0;
+
+                break :blk @bitCast(@mod(signed_lhs, signed_rhs));
+            },
+            .shift_left => if (rhs >= 32) return RuntimeError.ShiftOutOfRange else lhs << @intCast(rhs),
+            .logical_shift_right => if (rhs >= 32) return RuntimeError.ShiftOutOfRange else lhs >> @intCast(rhs),
+            .arithmetic_shift_right => if (rhs >= 32) return RuntimeError.ShiftOutOfRange else @bitCast(@as(i32, @bitCast(lhs)) >> @intCast(rhs)),
+            .bitwise_and => lhs & rhs,
+            .bitwise_or => lhs | rhs,
+            .bitwise_xor => lhs ^ rhs,
+            .logical_and => @intFromBool(lhs != 0 and rhs != 0),
+            .logical_or => @intFromBool(lhs != 0 or rhs != 0),
+        };
+    }
+}
+
+fn binaryFloat(self: *Self, instruction: bc.Instruction, comptime operation: BinaryFloat) void {
+    for (0..instruction.components) |component| {
+        const lhs: f32 = @bitCast(self.registers[@as(usize, instruction.b) + component]);
+        const rhs: f32 = @bitCast(self.registers[@as(usize, instruction.c) + component]);
+        const result = switch (operation) {
+            .add => lhs + rhs,
+            .subtract => lhs - rhs,
+            .multiply => lhs * rhs,
+            .divide => lhs / rhs,
+            .modulo => lhs - rhs * @floor(lhs / rhs),
+        };
+        self.registers[@as(usize, instruction.a) + component] = @bitCast(result);
+    }
+}
+
+fn compareInt(self: *Self, instruction: bc.Instruction, comptime operation: CompareInt) void {
+    const lhs = self.registers[instruction.b];
+    const rhs = self.registers[instruction.c];
+    self.registers[instruction.a] = @intFromBool(switch (operation) {
+        .equal => lhs == rhs,
+        .not_equal => lhs != rhs,
+        .unsigned_less => lhs < rhs,
+        .signed_less => @as(i32, @bitCast(lhs)) < @as(i32, @bitCast(rhs)),
+    });
+}
+
+fn compareFloat(self: *Self, instruction: bc.Instruction, comptime operation: CompareFloat) void {
+    const lhs: f32 = @bitCast(self.registers[instruction.b]);
+    const rhs: f32 = @bitCast(self.registers[instruction.c]);
+    const unordered = std.math.isNan(lhs) or std.math.isNan(rhs);
+    self.registers[instruction.a] = @intFromBool(switch (operation) {
+        .ordered_equal => !unordered and lhs == rhs,
+        .unordered_equal => unordered or lhs == rhs,
+        .ordered_not_equal => !unordered and lhs != rhs,
+        .unordered_not_equal => unordered or lhs != rhs,
+        .ordered_less => !unordered and lhs < rhs,
+        .unordered_less => unordered or lhs < rhs,
+    });
+}
+
+fn select(self: *Self, instruction: bc.Instruction) void {
+    const selected = if (self.registers[instruction.b] != 0) instruction.c else instruction.d;
+    for (0..instruction.components) |component|
+        self.registers[@as(usize, instruction.a) + component] = self.registers[@as(usize, selected) + component];
+}
+
+fn applyEdge(self: *Self, program: *const Program, edge_index: u32) RuntimeError!u32 {
+    if (edge_index >= program.edges.len)
+        return RuntimeError.InvalidBytecode;
+
+    const edge = program.edges[edge_index];
+    const end = @as(usize, edge.first_copy) + edge.copy_count;
+
+    if (end > program.copies.len)
+        return RuntimeError.InvalidBytecode;
+
+    const copies = program.copies[edge.first_copy..end];
+    for (copies) |item| {
+        for (0..item.components) |component| {
+            self.scratch[@as(usize, item.scratch_base) + component] = self.registers[@as(usize, item.source) + component];
+        }
+    }
+    for (copies) |item| {
+        for (0..item.components) |component| {
+            self.registers[@as(usize, item.destination) + component] = self.scratch[@as(usize, item.scratch_base) + component];
+        }
+    }
+
+    if (edge.target_pc >= program.code.len)
+        return RuntimeError.InvalidBytecode;
+
+    return edge.target_pc;
+}

src/software/interpreter/Shader.zig

@@ -0,0 +1,162 @@
+const std = @import("std");
+const vk = @import("vulkan");
+const base = @import("base");
+const shader_ir = @import("shader_ir");
+
+const Program = @import("Program.zig");
+const Runtime = @import("Runtime.zig");
+const SoftShaderModule = @import("../SoftShaderModule.zig");
+
+const VkError = base.VkError;
+const ir = shader_ir.ir;
+
+pub const RuntimeSlot = struct {
+    mutex: std.Io.Mutex = .init,
+    runtime: Runtime,
+};
+
+const Self = @This();
+
+program: Program,
+runtimes: []RuntimeSlot,
+workgroup_size: ?[3]u32,
+
+/// Compiles a stage when the current interpreter can execute its complete
+/// interface. `null` deliberately selects the existing SPIR-V runtime.
+pub fn compile(
+    allocator: std.mem.Allocator,
+    module: *SoftShaderModule,
+    stage: *const vk.PipelineShaderStageCreateInfo,
+    runtime_count: usize,
+) VkError!?Self {
+    const expected_stage = commonStage(stage.stage) orelse return null;
+    if (expected_stage == .fragment)
+        return null;
+
+    const specializations = try specializationValues(allocator, stage.p_specialization_info);
+    defer if (specializations.len != 0) allocator.free(specializations);
+
+    var module_ir = module.interface.instantiateIr(allocator, .{
+        .entry_point = std.mem.span(stage.p_name),
+        .stage = expected_stage,
+        .specializations = specializations,
+    }) catch |err| {
+        if (err == error.OutOfMemory)
+            return VkError.OutOfDeviceMemory;
+        std.log.scoped(.SoftIrInterpreter).debug("IR translation fallback: {s}", .{@errorName(err)});
+        return null;
+    };
+    defer module_ir.deinit();
+
+    var program = Program.compile(allocator, &module_ir) catch |err| {
+        if (err == error.OutOfMemory)
+            return VkError.OutOfDeviceMemory;
+        std.log.scoped(.SoftIrInterpreter).debug("bytecode lowering fallback: {s}", .{@errorName(err)});
+        return null;
+    };
+    errdefer program.deinit();
+
+    if (!hasCompatibleInterface(&program, expected_stage) or
+        (expected_stage == .compute and module_ir.execution_modes.workgroup_size == null))
+    {
+        std.log.scoped(.SoftIrInterpreter).debug("stage interface or execution modes require the SPIR-V runtime", .{});
+        program.deinit();
+        return null;
+    }
+
+    const runtimes = allocator.alloc(RuntimeSlot, runtime_count) catch return VkError.OutOfDeviceMemory;
+    var initialized: usize = 0;
+    errdefer {
+        for (runtimes[0..initialized]) |*slot|
+            slot.runtime.deinit();
+        allocator.free(runtimes);
+    }
+    for (runtimes) |*slot| {
+        slot.* = .{ .runtime = Runtime.init(allocator, &program) catch return VkError.OutOfDeviceMemory };
+        initialized += 1;
+    }
+
+    std.log.scoped(.SoftIrInterpreter).debug("compiled {s} stage to {d} bytecode instructions", .{
+        @tagName(expected_stage),
+        program.code.len,
+    });
+    return .{
+        .program = program,
+        .runtimes = runtimes,
+        .workgroup_size = module_ir.execution_modes.workgroup_size,
+    };
+}
+
+pub fn deinit(self: *Self) void {
+    for (self.runtimes) |*slot|
+        slot.runtime.deinit();
+    self.program.deinit();
+    self.* = undefined;
+}
+
+fn hasCompatibleInterface(program: *const Program, stage: ir.module.Stage) bool {
+    var has_position = false;
+    for (program.interfaces) |optional_binding| {
+        const binding = optional_binding orelse continue;
+        switch (binding.semantic) {
+            .location => |location| {
+                if (stage == .compute or location.index != 0 or
+                    @as(u16, location.component) + binding.span.components > 4)
+                    return false;
+            },
+            .builtin => |builtin| switch (stage) {
+                .vertex => switch (builtin) {
+                    .vertex_index, .instance_index => if (binding.direction != .input) return false,
+                    .position => {
+                        if (binding.direction != .output or binding.span.kind != .floating or binding.span.components != 4)
+                            return false;
+                        has_position = true;
+                    },
+                    else => return false,
+                },
+                .compute => if (builtin != .global_invocation_id or binding.direction != .input or binding.span.components != 3)
+                    return false,
+                .fragment => return false,
+            },
+        }
+    }
+    return stage != .vertex or has_position;
+}
+
+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.OutOfDeviceMemory;
+    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) ?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;
+}

src/software/interpreter/bytecode.zig

@@ -0,0 +1,99 @@
+const std = @import("std");
+
+pub const Register = u16;
+pub const invalid_register = std.math.maxInt(Register);
+
+pub const ValueKind = enum(u8) {
+    boolean,
+    signed_integer,
+    unsigned_integer,
+    floating,
+};
+
+pub const Span = struct {
+    base: Register,
+    components: u8,
+    kind: ValueKind,
+
+    pub fn sameShape(a: Span, b: Span) bool {
+        return a.components == b.components and a.kind == b.kind;
+    }
+};
+
+/// Native-endian internal bytecode. It is not a serialized or stable ABI.
+pub const Instruction = extern struct {
+    opcode: Opcode,
+    components: u16 = 1,
+    a: Register = invalid_register,
+    b: Register = invalid_register,
+    c: Register = invalid_register,
+    d: Register = invalid_register,
+    immediate: u32 = 0,
+};
+
+comptime {
+    std.debug.assert(@sizeOf(Instruction) == 16);
+}
+
+pub const Opcode = enum(u16) {
+    copy,
+    negate_i32,
+    negate_f32,
+    logical_not,
+    bitwise_not,
+    integer_add,
+    integer_subtract,
+    integer_multiply,
+    unsigned_divide,
+    signed_divide,
+    unsigned_modulo,
+    signed_modulo,
+    float_add,
+    float_subtract,
+    float_multiply,
+    float_divide,
+    float_modulo,
+    shift_left,
+    logical_shift_right,
+    arithmetic_shift_right,
+    bitwise_and,
+    bitwise_or,
+    bitwise_xor,
+    logical_and,
+    logical_or,
+    compare_equal,
+    compare_not_equal,
+    compare_unsigned_less,
+    compare_signed_less,
+    compare_ordered_float_equal,
+    compare_unordered_float_equal,
+    compare_ordered_float_not_equal,
+    compare_unordered_float_not_equal,
+    compare_ordered_float_less,
+    compare_unordered_float_less,
+    select,
+    jump_edge,
+    branch,
+    return_void,
+    discard,
+    @"unreachable",
+};
+
+pub const Copy = struct {
+    destination: Register,
+    source: Register,
+    components: u8,
+    scratch_base: Register,
+};
+
+pub const Edge = struct {
+    target_block: u32,
+    target_pc: u32 = 0,
+    first_copy: u32,
+    copy_count: u16,
+};
+
+pub const Branch = struct {
+    true_edge: u32,
+    false_edge: u32,
+};

src/software/interpreter/compute.zig

@@ -0,0 +1,48 @@
+const std = @import("std");
+const base = @import("base");
+const shader_ir = @import("shader_ir");
+
+const Shader = @import("Shader.zig");
+
+const VkError = base.VkError;
+const ir = shader_ir.ir;
+
+pub fn dispatch(shader: *Shader, base_group_x: u32, base_group_y: u32, base_group_z: u32, group_count_x: u32, group_count_y: u32, group_count_z: u32) VkError!void {
+    const local_size = shader.workgroup_size orelse return VkError.ValidationFailed;
+    const local_xy = std.math.mul(usize, local_size[0], local_size[1]) catch return VkError.ValidationFailed;
+    const local_count = std.math.mul(usize, local_xy, local_size[2]) catch return VkError.ValidationFailed;
+    if (shader.runtimes.len == 0)
+        return VkError.InvalidPipelineDrv;
+
+    const global_id = findGlobalInvocationId(&shader.program);
+    var runtime = &shader.runtimes[0].runtime;
+    for (0..group_count_z) |group_z| {
+        for (0..group_count_y) |group_y| {
+            for (0..group_count_x) |group_x| {
+                for (0..local_count) |local_index| {
+                    if (global_id) |variable| {
+                        const local_z = local_index / local_xy;
+                        const local_remainder = local_index - local_z * local_xy;
+                        const local_y = local_remainder / local_size[0];
+                        const local_x = local_remainder - local_y * local_size[0];
+                        runtime.writeInput(&shader.program, variable, &.{
+                            (base_group_x + @as(u32, @intCast(group_x))) * local_size[0] + @as(u32, @intCast(local_x)),
+                            (base_group_y + @as(u32, @intCast(group_y))) * local_size[1] + @as(u32, @intCast(local_y)),
+                            (base_group_z + @as(u32, @intCast(group_z))) * local_size[2] + @as(u32, @intCast(local_z)),
+                        }) catch return VkError.Unknown;
+                    }
+                    _ = runtime.run(&shader.program, .{}) catch return VkError.Unknown;
+                }
+            }
+        }
+    }
+}
+
+fn findGlobalInvocationId(program: *const @import("Program.zig")) ?ir.id.InterfaceVariableId {
+    for (program.interfaces, 0..) |optional_binding, index| {
+        const binding = optional_binding orelse continue;
+        if (binding.direction == .input and binding.semantic == .builtin and binding.semantic.builtin == .global_invocation_id)
+            return ir.id.InterfaceVariableId.fromIndex(index);
+    }
+    return null;
+}

src/software/interpreter/root.zig

@@ -0,0 +1,16 @@
+//! Software bytecode interpreter for the backend-agnostic shader IR.
+//!
+//! This first slice supports allocation-free scalar execution of 32-bit scalar
+//! and vector arithmetic, interface I/O, control flow, and block parameters.
+
+pub const bytecode = @import("bytecode.zig");
+pub const Program = @import("Program.zig");
+pub const Runtime = @import("Runtime.zig");
+
+pub const Outcome = Runtime.Outcome;
+pub const RunOptions = Runtime.RunOptions;
+pub const RuntimeError = Runtime.RuntimeError;
+
+comptime {
+    _ = @import("test/test.zig");
+}

src/software/interpreter/test/arithmetic.zig

@@ -0,0 +1,56 @@
+const std = @import("std");
+const shader_ir = @import("shader_ir");
+
+const Program = @import("../Program.zig");
+const Runtime = @import("../Runtime.zig");
+
+const ir = shader_ir.ir;
+
+fn f32Bits(value: f32) u32 {
+    return @bitCast(value);
+}
+
+fn bitsF32(value: u32) f32 {
+    return @bitCast(value);
+}
+
+test "[interpreter] vector floating-point arithmetic" {
+    var module = try ir.parser.parseString(std.testing.allocator,
+        \\ shader vertex @main
+        \\ {
+        \\     @lhs: vec4[f32] = input[location(0), component(0), index(0)]
+        \\     @rhs: vec4[f32] = input[location(1), component(0), index(0)]
+        \\     @output: vec4[f32] = output[location(0), component(0), index(0)]
+        \\
+        \\     fn @main() -> void
+        \\     {
+        \\         .entry():
+        \\             %lhs_value: vec4[f32] = load_interface @lhs
+        \\             %rhs_value: vec4[f32] = load_interface @rhs
+        \\             %product: vec4[f32] = float_multiply %lhs_value, %rhs_value
+        \\             store_interface @output, %product
+        \\             return
+        \\     }
+        \\ }
+    );
+    defer module.deinit();
+
+    const lhs = ir.id.InterfaceVariableId.fromIndex(0);
+    const rhs = ir.id.InterfaceVariableId.fromIndex(1);
+    const output = ir.id.InterfaceVariableId.fromIndex(2);
+
+    var program = try Program.compile(std.testing.allocator, &module);
+    defer program.deinit();
+    var runtime = try Runtime.init(std.testing.allocator, &program);
+    defer runtime.deinit();
+
+    try runtime.writeInput(&program, lhs, &.{ f32Bits(2), f32Bits(-3), f32Bits(0.5), f32Bits(8) });
+    try runtime.writeInput(&program, rhs, &.{ f32Bits(4), f32Bits(2), f32Bits(6), f32Bits(0.25) });
+    try std.testing.expectEqual(Runtime.Outcome.returned, try runtime.run(&program, .{}));
+
+    var result: [4]u32 = undefined;
+    try runtime.readOutput(&program, output, &result);
+    const expected = [_]f32{ 8, -6, 3, 2 };
+    for (result, expected) |actual, wanted|
+        try std.testing.expectEqual(wanted, bitsF32(actual));
+}

src/software/interpreter/test/branching.zig

@@ -0,0 +1,65 @@
+const std = @import("std");
+const shader_ir = @import("shader_ir");
+
+const Program = @import("../Program.zig");
+const Runtime = @import("../Runtime.zig");
+
+const ir = shader_ir.ir;
+
+fn i32Bits(value: i32) u32 {
+    return @bitCast(value);
+}
+
+test "[interpreter] branches with block parameters" {
+    var module = try ir.parser.parseString(std.testing.allocator,
+        \\ shader vertex @main
+        \\ {
+        \\     @input: i32 = input[location(0), component(0), index(0)]
+        \\     @output: i32 = output[location(0), component(0), index(0)]
+        \\
+        \\     %one: constant i32 = 1
+        \\     %ten: constant i32 = 10
+        \\
+        \\     fn @main() -> void
+        \\     {
+        \\         .entry():
+        \\             %value: i32 = load_interface @input
+        \\             %condition: bool = cmp_signed_less %value, %ten
+        \\             conditional_branch %condition, .less(), .greater_equal()
+        \\
+        \\         .less():
+        \\             %incremented: i32 = integer_add %value, %one
+        \\             branch .merge(%incremented)
+        \\
+        \\         .greater_equal():
+        \\             %decremented: i32 = integer_subtract %value, %one
+        \\             branch .merge(%decremented)
+        \\
+        \\         .merge(%merged: i32):
+        \\             store_interface @output, %merged
+        \\             return
+        \\     }
+        \\ }
+    );
+
+    const input = ir.id.InterfaceVariableId.fromIndex(0);
+    const output = ir.id.InterfaceVariableId.fromIndex(1);
+
+    var program = try Program.compile(std.testing.allocator, &module);
+    defer program.deinit();
+    module.deinit();
+
+    var runtime = try Runtime.init(std.testing.allocator, &program);
+    defer runtime.deinit();
+
+    try runtime.writeInput(&program, input, &.{i32Bits(5)});
+    try std.testing.expectEqual(Runtime.Outcome.returned, try runtime.run(&program, .{}));
+    var result: [1]u32 = undefined;
+    try runtime.readOutput(&program, output, &result);
+    try std.testing.expectEqual(@as(i32, 6), @as(i32, @bitCast(result[0])));
+
+    try runtime.writeInput(&program, input, &.{i32Bits(20)});
+    try std.testing.expectEqual(Runtime.Outcome.returned, try runtime.run(&program, .{}));
+    try runtime.readOutput(&program, output, &result);
+    try std.testing.expectEqual(@as(i32, 19), @as(i32, @bitCast(result[0])));
+}

src/software/interpreter/test/loops.zig

@@ -0,0 +1,58 @@
+const std = @import("std");
+const shader_ir = @import("shader_ir");
+
+const Program = @import("../Program.zig");
+const Runtime = @import("../Runtime.zig");
+
+const ir = shader_ir.ir;
+
+test "[interpreter] loop back edges copy block arguments in parallel" {
+    var module = try ir.parser.parseString(std.testing.allocator,
+        \\ shader compute @main
+        \\ {
+        \\     @output_a: u32 = output[location(0), component(0), index(0)]
+        \\     @output_b: u32 = output[location(1), component(0), index(0)]
+        \\
+        \\     %zero: constant u32 = 0
+        \\     %one: constant u32 = 1
+        \\     %two: constant u32 = 2
+        \\     %three: constant u32 = 3
+        \\
+        \\     fn @main() -> void
+        \\     {
+        \\         .entry():
+        \\             branch .loop(%one, %two, %zero)
+        \\
+        \\         .loop(%a: u32, %b: u32, %index: u32):
+        \\             %condition: bool = cmp_unsigned_less %index, %three
+        \\             conditional_branch %condition, .body(), .exit(%a, %b)
+        \\
+        \\         .body():
+        \\             %next_index: u32 = integer_add %index, %one
+        \\             branch .loop(%b, %a, %next_index)
+        \\
+        \\         .exit(%final_a: u32, %final_b: u32):
+        \\             store_interface @output_a, %final_a
+        \\             store_interface @output_b, %final_b
+        \\             return
+        \\     }
+        \\ }
+    );
+    defer module.deinit();
+
+    const output_a = ir.id.InterfaceVariableId.fromIndex(0);
+    const output_b = ir.id.InterfaceVariableId.fromIndex(1);
+
+    var program = try Program.compile(std.testing.allocator, &module);
+    defer program.deinit();
+    var runtime = try Runtime.init(std.testing.allocator, &program);
+    defer runtime.deinit();
+
+    try std.testing.expectEqual(Runtime.Outcome.returned, try runtime.run(&program, .{}));
+    var a: [1]u32 = undefined;
+    var b: [1]u32 = undefined;
+    try runtime.readOutput(&program, output_a, &a);
+    try runtime.readOutput(&program, output_b, &b);
+    try std.testing.expectEqual(@as(u32, 2), a[0]);
+    try std.testing.expectEqual(@as(u32, 1), b[0]);
+}

src/software/interpreter/test/termination.zig

@@ -0,0 +1,51 @@
+const std = @import("std");
+const shader_ir = @import("shader_ir");
+
+const Program = @import("../Program.zig");
+const Runtime = @import("../Runtime.zig");
+
+const ir = shader_ir.ir;
+
+test "[interpreter] fragment discard is an execution outcome" {
+    var module = try ir.parser.parseString(std.testing.allocator,
+        \\ shader fragment @main
+        \\ {
+        \\     fn @main() -> void
+        \\     {
+        \\         .entry():
+        \\             discard
+        \\     }
+        \\ }
+    );
+    defer module.deinit();
+
+    var program = try Program.compile(std.testing.allocator, &module);
+    defer program.deinit();
+    var runtime = try Runtime.init(std.testing.allocator, &program);
+    defer runtime.deinit();
+
+    try std.testing.expectEqual(Runtime.Outcome.discarded, try runtime.run(&program, .{}));
+}
+
+test "[interpreter] execution budget stops an infinite loop" {
+    var module = try ir.parser.parseString(std.testing.allocator,
+        \\ shader compute @main
+        \\ {
+        \\     fn @main() -> void
+        \\     {
+        \\         .entry():
+        \\             branch .loop()
+        \\         .loop():
+        \\             branch .loop()
+        \\     }
+        \\ }
+    );
+    defer module.deinit();
+
+    var program = try Program.compile(std.testing.allocator, &module);
+    defer program.deinit();
+    var runtime = try Runtime.init(std.testing.allocator, &program);
+    defer runtime.deinit();
+
+    try std.testing.expectError(Runtime.RuntimeError.StepLimitExceeded, runtime.run(&program, .{ .max_steps = 8 }));
+}

src/software/interpreter/test/test.zig

@@ -0,0 +1,14 @@
+const std = @import("std");
+
+const bytecode = @import("../bytecode.zig");
+
+test "[interpreter] bytecode instruction size" {
+    try std.testing.expectEqual(@as(usize, 16), @sizeOf(bytecode.Instruction));
+}
+
+comptime {
+    _ = @import("arithmetic.zig");
+    _ = @import("branching.zig");
+    _ = @import("loops.zig");
+    _ = @import("termination.zig");
+}

src/software/interpreter/vertex.zig

@@ -0,0 +1,173 @@
+const std = @import("std");
+const vk = @import("vulkan");
+const base = @import("base");
+const shader_ir = @import("shader_ir");
+
+const bc = @import("bytecode.zig");
+const Shader = @import("Shader.zig");
+const SoftPipeline = @import("../SoftPipeline.zig");
+const Renderer = @import("../device/Renderer.zig");
+const blitter = @import("../device/blitter.zig");
+
+const VkError = base.VkError;
+const ir = shader_ir.ir;
+const interface_blob_padding = @sizeOf(base.zm.F32x4);
+
+pub fn run(
+    allocator: std.mem.Allocator,
+    pipeline: *SoftPipeline,
+    shader: *Shader,
+    batch_id: usize,
+    batch_size: usize,
+    vertex_count: usize,
+    first_vertex: usize,
+    first_instance: usize,
+    indices: ?[]const u32,
+    primitive_restart: ?[]const bool,
+    instance_index: usize,
+    draw_call: *Renderer.DrawCall,
+) VkError!void {
+    const slot = &shader.runtimes[batch_id];
+    const io = draw_call.renderer.device.interface.io();
+    slot.mutex.lock(io) catch return VkError.DeviceLost;
+    defer slot.mutex.unlock(io);
+
+    var invocation_index = batch_id;
+    while (invocation_index < vertex_count) : (invocation_index += batch_size) {
+        const output = &draw_call.vertices[(instance_index * vertex_count) + invocation_index];
+        if (primitive_restart) |restart| {
+            if (restart[invocation_index]) {
+                output.primitive_restart = true;
+                continue;
+            }
+        }
+
+        const vertex_index: u32 = if (indices) |draw_indices| draw_indices[invocation_index] else @intCast(first_vertex + invocation_index);
+        try populateInputs(
+            &slot.runtime,
+            &shader.program,
+            pipeline,
+            draw_call,
+            vertex_index,
+            @intCast(first_instance + instance_index),
+        );
+        const outcome = slot.runtime.run(&shader.program, .{}) catch return VkError.Unknown;
+        if (outcome == .discarded)
+            continue;
+        try collectOutputs(allocator, &slot.runtime, &shader.program, output);
+    }
+}
+
+fn populateInputs(
+    runtime: anytype,
+    program: *const @import("Program.zig"),
+    pipeline: *SoftPipeline,
+    draw_call: *Renderer.DrawCall,
+    vertex_index: u32,
+    instance_index: u32,
+) VkError!void {
+    for (program.interfaces, 0..) |optional_binding, index| {
+        const binding = optional_binding orelse continue;
+        if (binding.direction != .input)
+            continue;
+        const variable = ir.id.InterfaceVariableId.fromIndex(index);
+        var values: [4]u32 = @splat(0);
+        switch (binding.semantic) {
+            .builtin => |builtin| values[0] = switch (builtin) {
+                .vertex_index => vertex_index,
+                .instance_index => instance_index,
+                else => return VkError.InvalidPipelineDrv,
+            },
+            .location => |location| {
+                const attribute = findAttribute(
+                    pipeline.interface.mode.graphics.input_assembly.attribute_description orelse &.{},
+                    location.location,
+                ) orelse {
+                    runtime.writeInput(program, variable, values[0..binding.span.components]) catch return VkError.Unknown;
+                    continue;
+                };
+                const binding_description = findBinding(
+                    pipeline.interface.mode.graphics.input_assembly.binding_description orelse return VkError.ValidationFailed,
+                    attribute.binding,
+                ) orelse return VkError.ValidationFailed;
+                const vertex_buffer = draw_call.renderer.state.data.graphics.vertex_buffers[attribute.binding];
+                const buffer = vertex_buffer.buffer;
+                const memory = buffer.interface.memory orelse return VkError.InvalidDeviceMemoryDrv;
+                const input_index = switch (binding_description.input_rate) {
+                    .vertex => @as(usize, vertex_index),
+                    .instance => @as(usize, instance_index),
+                    else => return VkError.ValidationFailed,
+                };
+                const offset = buffer.interface.offset + vertex_buffer.offset + binding_description.stride * input_index + attribute.offset;
+                const input_size = base.format.texelSize(attribute.format);
+                var robust_bytes: [64]u8 = @splat(0);
+                if (input_size > robust_bytes.len)
+                    return VkError.Unknown;
+                if (offset < memory.size) {
+                    const available = @min(input_size, @as(usize, @intCast(memory.size - offset)));
+                    const mapped = memory.map(offset, available) catch &.{};
+                    @memcpy(robust_bytes[0..mapped.len], mapped);
+                }
+                values = if (base.format.isUnnormalizedInteger(attribute.format))
+                    blitter.readInt4(robust_bytes[0..input_size], attribute.format)
+                else
+                    @bitCast(blitter.readFloat4(robust_bytes[0..input_size], attribute.format));
+                const first_component: usize = location.component;
+                const end_component = first_component + binding.span.components;
+                runtime.writeInput(program, variable, values[first_component..end_component]) catch return VkError.Unknown;
+                continue;
+            },
+        }
+        runtime.writeInput(program, variable, values[0..binding.span.components]) catch return VkError.Unknown;
+    }
+}
+
+fn collectOutputs(allocator: std.mem.Allocator, runtime: anytype, program: *const @import("Program.zig"), output: *Renderer.Vertex) VkError!void {
+    for (program.interfaces, 0..) |optional_binding, index| {
+        const binding = optional_binding orelse continue;
+        if (binding.direction != .output)
+            continue;
+        const variable = ir.id.InterfaceVariableId.fromIndex(index);
+        var values: [4]u32 = @splat(0);
+        runtime.readOutput(program, variable, values[0..binding.span.components]) catch return VkError.Unknown;
+        switch (binding.semantic) {
+            .builtin => |builtin| switch (builtin) {
+                .position => @memcpy(std.mem.asBytes(&output.position), std.mem.asBytes(&values)),
+                else => return VkError.InvalidPipelineDrv,
+            },
+            .location => |location| {
+                if (location.location >= output.outputs.len or location.component >= output.outputs[0].len)
+                    return VkError.ValidationFailed;
+                const size = @as(usize, binding.span.components) * @sizeOf(u32);
+                const blob = allocator.alloc(u8, size + interface_blob_padding) catch return VkError.OutOfDeviceMemory;
+                @memset(blob, 0);
+                @memcpy(blob[0..size], std.mem.asBytes(&values)[0..size]);
+                output.outputs[location.location][location.component] = .{
+                    .interpolation_type = switch (binding.span.kind) {
+                        .signed_integer, .unsigned_integer, .boolean => .flat,
+                        .floating => .smooth,
+                    },
+                    .centroid = false,
+                    .blob = blob,
+                    .size = size,
+                };
+            },
+        }
+    }
+}
+
+fn findAttribute(attributes: []const vk.VertexInputAttributeDescription, location: u32) ?vk.VertexInputAttributeDescription {
+    for (attributes) |attribute|
+        if (attribute.location == location) return attribute;
+    return null;
+}
+
+fn findBinding(bindings: []const vk.VertexInputBindingDescription, binding: u32) ?vk.VertexInputBindingDescription {
+    for (bindings) |description|
+        if (description.binding == binding) return description;
+    return null;
+}
+
+comptime {
+    _ = bc;
+}

src/software/lib.zig

@@ -6,6 +6,7 @@ pub const c = @import("soft_c");
 pub const config = base.config;
 
 pub const Device = @import("device/Device.zig");
+pub const interpreter = @import("interpreter/root.zig");
 
 pub const SoftInstance = @import("SoftInstance.zig");
 pub const SoftDevice = @import("SoftDevice.zig");
@@ -84,6 +85,7 @@ comptime {
 
 test {
     std.testing.refAllDecls(Device);
+    std.testing.refAllDecls(interpreter);
     std.testing.refAllDecls(SoftBinarySemaphore);
     std.testing.refAllDecls(SoftBuffer);
     std.testing.refAllDecls(SoftBufferView);

test/test_runner.zig

@@ -43,17 +43,7 @@ pub fn main() !void {
         var status = Status.pass;
         slowest.startTiming();
 
-        const friendly_name = blk: {
-            const name = t.name;
-            var it = std.mem.splitScalar(u8, name, '.');
-            while (it.next()) |value| {
-                if (std.mem.eql(u8, value, "test")) {
-                    const rest = it.rest();
-                    break :blk if (rest.len > 0) rest else name;
-                }
-            }
-            break :blk name;
-        };
+        const friendly_name = friendlyName(t.name);
 
         current_test = friendly_name;
         std.testing.allocator_instance = .{};
@@ -220,6 +210,13 @@ pub const panic = std.debug.FullPanic(struct {
     }
 }.panicFn);
 
+fn friendlyName(name: []const u8) []const u8 {
+    const marker = ".test.";
+    const index = std.mem.lastIndexOf(u8, name, marker) orelse return name;
+    const suffix = name[index + marker.len ..];
+    return if (suffix.len > 0) suffix else name;
+}
+
 fn isUnnamed(t: std.builtin.TestFn) bool {
     const marker = ".test_";
     const test_name = t.name;