[Flint] adding IR lowering and printer

[IR] switching from "passes" to "transformers" for clarity

mirror-test.txt

@@ -1 +0,0 @@
-tests

src/compiler/ir/Builder.zig

@@ -225,3 +225,116 @@ fn constantEql(a: constant_ir.ConstantValue, b: constant_ir.ConstantValue) bool
         .composite => |value| b == .composite and std.mem.eql(ids.ConstantId, value, b.composite),
     };
 }
+
+test "Builder: generation" {
+    const cfg = @import("cfg.zig");
+    const parser = @import("parser/parser.zig");
+    const printer = @import("printer.zig");
+    const validator = @import("validator/validator.zig");
+
+    // shader vertex @main
+    // {
+    //     @color: vec4[f32] = input[location(0), component(0), index(0)]
+    //     @out_color: vec4[f32] = output[location(0), component(0), index(0)]
+    //     %0: constant bool = true
+    //     %1: constant f32 = bits(0x3f800000)
+    //
+    //     fn @main() -> void
+    //     {
+    //         .entry():
+    //             %3: vec4[f32] = load_interface @color
+    //             conditional_branch %0, .pass(), .merge(%3)
+    //
+    //         .pass():
+    //             %4: vec4[f32] = composite_construct %1, %1, %1, %1
+    //             branch .merge(%4)
+    //
+    //         .merge(%2: vec4[f32]):
+    //             store_interface @out_color, %2
+    //             return
+    //     }
+    // }
+
+    var module = module_ir.Module.init(std.testing.allocator, .vertex);
+    defer module.deinit();
+    var builder = Self.init(&module);
+
+    const void_type = try builder.internType(.void);
+    const bool_type = try builder.internType(.boolean);
+    const f32_type = try builder.internType(.{ .floating = .{ .bits = 32 } });
+    const duplicate_f32 = try builder.internType(.{ .floating = .{ .bits = 32 } });
+    try std.testing.expectEqual(f32_type, duplicate_f32);
+    const vec4_type = try builder.internType(.{ .vector = .{ .element_type = f32_type, .length = 4 } });
+
+    const true_value = try builder.internConstant(bool_type, .{ .boolean = true });
+    const one = try builder.internConstant(f32_type, .{ .float_bits = @as(u32, @bitCast(@as(f32, 1.0))) });
+
+    const input = try builder.addInterfaceVariable(vec4_type, .input, .{ .location = .{ .location = 0 } }, "color");
+    const output = try builder.addInterfaceVariable(vec4_type, .output, .{ .location = .{ .location = 0 } }, "out_color");
+    const main = try builder.addFunction(void_type, "main");
+    builder.setEntryPoint(main);
+    const entry = try builder.addBlock(main, "entry");
+    const pass = try builder.addBlock(main, "pass");
+    const merge = try builder.addBlock(main, "merge");
+    const merged = try builder.addBlockParameter(merge, vec4_type, "merged");
+
+    const loaded = (try builder.appendInstruction(entry, vec4_type, .{
+        .load_interface = .{ .variable = input },
+    }, "loaded")).?;
+    try builder.setTerminator(entry, .{ .conditional_branch = .{
+        .condition = true_value,
+        .true_edge = try builder.edge(pass, &.{}),
+        .false_edge = try builder.edge(merge, &.{loaded}),
+    } });
+
+    const splat = (try builder.appendInstruction(pass, vec4_type, .{
+        .composite_construct = .{ .elements = &.{ one, one, one, one } },
+    }, "white")).?;
+    try builder.setTerminator(pass, .{ .branch = try builder.edge(merge, &.{splat}) });
+    _ = try builder.appendInstruction(merge, null, .{
+        .store_interface = .{ .variable = output, .value = merged },
+    }, null);
+    try builder.setTerminator(merge, .return_void);
+
+    try validator.validate(&module);
+
+    var control_flow = try cfg.init(std.testing.allocator, &module, main);
+    defer control_flow.deinit();
+    try std.testing.expectEqual(@as(usize, 2), control_flow.predecessors(merge).?.len);
+    try std.testing.expect(control_flow.dominates(entry, merge));
+    try std.testing.expect(!control_flow.dominates(pass, merge));
+
+    const text = try printer.allocPrint(std.testing.allocator, &module);
+
+    defer std.testing.allocator.free(text);
+    try std.testing.expect(std.mem.indexOf(u8, text, "shader vertex @main") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "@color: vec4[f32] = input[location(0), component(0), index(0)]") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "@out_color: vec4[f32] = output[location(0), component(0), index(0)]") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "conditional_branch %0, .pass(), .merge(%loaded)") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, ".merge(%merged: vec4[f32])") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "store_interface @out_color, %merged") != null);
+
+    var parsed = try parser.parseString(std.testing.allocator, text);
+    defer parsed.deinit();
+    const round_trip = try printer.allocPrint(std.testing.allocator, &parsed);
+    defer std.testing.allocator.free(round_trip);
+    try std.testing.expectEqualStrings(text, round_trip);
+
+    const io = std.Options.debug_io;
+    const path = ".zig-cache/ir-parser-round-trip.ir";
+    const file = try std.Io.Dir.cwd().createFile(io, path, .{ .truncate = true });
+    {
+        defer file.close(io);
+        var file_buffer: [4096]u8 = @splat(0);
+        var file_writer = file.writer(io, &file_buffer);
+        try file_writer.interface.writeAll(text);
+        try file_writer.interface.flush();
+    }
+    defer std.Io.Dir.cwd().deleteFile(io, path) catch @panic("Caught an error while handling an error");
+
+    var parsed_file = try parser.parseFile(std.testing.allocator, io, path);
+    defer parsed_file.deinit();
+    const file_round_trip = try printer.allocPrint(std.testing.allocator, &parsed_file);
+    defer std.testing.allocator.free(file_round_trip);
+    try std.testing.expectEqualStrings(text, file_round_trip);
+}

src/compiler/ir/Rewriter.zig

@@ -2,6 +2,7 @@ const std = @import("std");
 const ids = @import("id.zig");
 const module_ir = @import("module.zig");
 const Builder = @import("Builder.zig");
+const validator = @import("validator/validator.zig");
 
 const Self = @This();
 
@@ -52,11 +53,11 @@ pub fn countUses(self: *const Self, value: ids.ValueId) usize {
 }
 
 pub fn replaceAllUses(self: *Self, old: ids.ValueId, replacement: ids.ValueId) Error!usize {
-    const old_value = self.module.values.get(old) orelse return error.InvalidValue;
-    const replacement_value = self.module.values.get(replacement) orelse return error.InvalidValue;
+    const old_value = self.module.values.get(old) orelse return Error.InvalidValue;
+    const replacement_value = self.module.values.get(replacement) orelse return Error.InvalidValue;
 
     if (old_value.type != replacement_value.type)
-        return error.TypeMismatch;
+        return Error.TypeMismatch;
 
     if (old == replacement)
         return 0;
@@ -77,17 +78,17 @@ pub fn replaceAllUses(self: *Self, old: ids.ValueId, replacement: ids.ValueId) E
 }
 
 pub fn eraseInstruction(self: *Self, instruction_id: ids.InstructionId) Error!void {
-    const instruction = self.module.instructions.get(instruction_id) orelse return error.InvalidInstruction;
+    const instruction = self.module.instructions.get(instruction_id) orelse return Error.InvalidInstruction;
 
     if (instruction.operation.hasSideEffects())
-        return error.SideEffectingInstruction;
+        return Error.SideEffectingInstruction;
 
     if (instruction.result) |result| {
         if (self.countUses(result) != 0)
-            return error.ResultStillUsed;
+            return Error.ResultStillUsed;
     }
 
-    const block = self.module.blocks.getMut(instruction.parent_block) orelse return error.InvalidBlock;
+    const block = self.module.blocks.getMut(instruction.parent_block) orelse return Error.InvalidBlock;
     var owned_index: ?usize = null;
 
     for (block.instructions.items, 0..) |candidate, index| {
@@ -97,29 +98,23 @@ pub fn eraseInstruction(self: *Self, instruction_id: ids.InstructionId) Error!vo
         }
     }
 
-    _ = block.instructions.orderedRemove(owned_index orelse return error.InstructionNotOwnedByBlock);
+    _ = block.instructions.orderedRemove(owned_index orelse return Error.InstructionNotOwnedByBlock);
     if (instruction.result) |result|
         _ = self.module.values.remove(result);
     _ = self.module.instructions.remove(instruction_id);
 }
 
-pub fn redirectEdges(
-    self: *Self,
-    source: ids.BlockId,
-    old_target: ids.BlockId,
-    new_target: ids.BlockId,
-    new_arguments: []const ids.ValueId,
-) Error!usize {
-    const source_block = self.module.blocks.get(source) orelse return error.InvalidBlock;
-    const target_block = self.module.blocks.get(new_target) orelse return error.InvalidBlock;
+pub fn redirectEdges(self: *Self, source: ids.BlockId, old_target: ids.BlockId, new_target: ids.BlockId, new_arguments: []const ids.ValueId) Error!usize {
+    const source_block = self.module.blocks.get(source) orelse return Error.InvalidBlock;
+    const target_block = self.module.blocks.get(new_target) orelse return Error.InvalidBlock;
 
     if (source_block.parent_function != target_block.parent_function)
-        return error.InvalidFunction;
+        return Error.InvalidFunction;
 
     try self.validateArguments(target_block, new_arguments);
 
     const mutable_source = self.module.blocks.getMut(source).?;
-    const terminator = if (mutable_source.terminator) |*value| value else return error.InvalidBlock;
+    const terminator = if (mutable_source.terminator) |*value| value else return Error.InvalidBlock;
 
     var count: usize = 0;
 
@@ -141,29 +136,23 @@ pub fn redirectEdges(
     return count;
 }
 
-pub fn addBlockParameter(
-    self: *Self,
-    block_id: ids.BlockId,
-    ty: ids.TypeId,
-    name: ?[]const u8,
-    incoming: []const IncomingValue,
-) Error!ids.ValueId {
-    const block = self.module.blocks.get(block_id) orelse return error.InvalidBlock;
-    const function = self.module.functions.get(block.parent_function) orelse return error.InvalidFunction;
+pub fn addBlockParameter(self: *Self, block_id: ids.BlockId, ty: ids.TypeId, name: ?[]const u8, incoming: []const IncomingValue) Error!ids.ValueId {
+    const block = self.module.blocks.get(block_id) orelse return Error.InvalidBlock;
+    const function = self.module.functions.get(block.parent_function) orelse return Error.InvalidFunction;
 
     for (incoming) |item| {
-        const value = self.module.values.get(item.value) orelse return error.InvalidValue;
+        const value = self.module.values.get(item.value) orelse return Error.InvalidValue;
 
         if (value.type != ty)
-            return error.TypeMismatch;
+            return Error.TypeMismatch;
 
         if (!functionHasEdgeTo(self.module, function, item.predecessor, block_id))
-            return error.UnexpectedIncomingValue;
+            return Error.UnexpectedIncomingValue;
     }
     for (function.blocks.items) |predecessor| {
         const edge_count = countEdgesTo(self.module.blocks.get(predecessor).?, block_id);
         if (edge_count != 0 and findIncoming(incoming, predecessor) == null)
-            return error.MissingIncomingValue;
+            return Error.MissingIncomingValue;
     }
 
     var builder = Builder.init(self.module);
@@ -177,19 +166,14 @@ pub fn addBlockParameter(
     return parameter;
 }
 
-pub fn removeBlockParameter(
-    self: *Self,
-    block_id: ids.BlockId,
-    parameter_index: usize,
-    replacement: ids.ValueId,
-) Error!void {
-    const block = self.module.blocks.get(block_id) orelse return error.InvalidBlock;
-    if (parameter_index >= block.parameters.items.len) return error.InvalidParameterIndex;
+pub fn removeBlockParameter(self: *Self, block_id: ids.BlockId, parameter_index: usize, replacement: ids.ValueId) Error!void {
+    const block = self.module.blocks.get(block_id) orelse return Error.InvalidBlock;
+    if (parameter_index >= block.parameters.items.len) return Error.InvalidParameterIndex;
     const parameter = block.parameters.items[parameter_index];
-    if (parameter == replacement) return error.InvalidValue;
+    if (parameter == replacement) return Error.InvalidValue;
     _ = try self.replaceAllUses(parameter, replacement);
 
-    const function = self.module.functions.get(block.parent_function) orelse return error.InvalidFunction;
+    const function = self.module.functions.get(block.parent_function) orelse return Error.InvalidFunction;
     for (function.blocks.items) |predecessor| {
         try self.removeArgumentFromEdges(predecessor, block_id, parameter_index);
     }
@@ -198,7 +182,7 @@ pub fn removeBlockParameter(
     _ = mutable_block.parameters.orderedRemove(parameter_index);
 
     for (mutable_block.parameters.items[parameter_index..], parameter_index..) |value_id, index| {
-        const value = self.module.values.getMut(value_id) orelse return error.InvalidValue;
+        const value = self.module.values.getMut(value_id) orelse return Error.InvalidValue;
         value.definition.block_parameter.index = @intCast(index);
     }
 
@@ -206,21 +190,15 @@ pub fn removeBlockParameter(
 }
 
 fn validateArguments(self: *const Self, target: *const module_ir.Block, arguments: []const ids.ValueId) Error!void {
-    if (arguments.len != target.parameters.items.len) return error.TypeMismatch;
+    if (arguments.len != target.parameters.items.len) return Error.TypeMismatch;
     for (arguments, target.parameters.items) |argument, parameter| {
-        const argument_value = self.module.values.get(argument) orelse return error.InvalidValue;
-        const parameter_value = self.module.values.get(parameter) orelse return error.InvalidValue;
-        if (argument_value.type != parameter_value.type) return error.TypeMismatch;
+        const argument_value = self.module.values.get(argument) orelse return Error.InvalidValue;
+        const parameter_value = self.module.values.get(parameter) orelse return Error.InvalidValue;
+        if (argument_value.type != parameter_value.type) return Error.TypeMismatch;
     }
 }
 
-fn redirectOne(
-    self: *Self,
-    edge: *module_ir.Edge,
-    old_target: ids.BlockId,
-    new_target: ids.BlockId,
-    arguments: []const ids.ValueId,
-) !bool {
+fn redirectOne(self: *Self, edge: *module_ir.Edge, old_target: ids.BlockId, new_target: ids.BlockId, arguments: []const ids.ValueId) !bool {
     if (edge.target != old_target)
         return false;
 
@@ -230,8 +208,8 @@ fn redirectOne(
 }
 
 fn appendArgumentToEdges(self: *Self, predecessor: ids.BlockId, target: ids.BlockId, value: ids.ValueId) !void {
-    const block = self.module.blocks.getMut(predecessor) orelse return error.InvalidBlock;
-    const terminator = if (block.terminator) |*item| item else return error.InvalidBlock;
+    const block = self.module.blocks.getMut(predecessor) orelse return Error.InvalidBlock;
+    const terminator = if (block.terminator) |*item| item else return Error.InvalidBlock;
 
     switch (terminator.*) {
         .branch => |*edge| {
@@ -257,8 +235,8 @@ fn appendEdgeArgument(self: *Self, edge: *module_ir.Edge, value: ids.ValueId) !v
 }
 
 fn removeArgumentFromEdges(self: *Self, predecessor: ids.BlockId, target: ids.BlockId, index: usize) !void {
-    const block = self.module.blocks.getMut(predecessor) orelse return error.InvalidBlock;
-    const terminator = if (block.terminator) |*item| item else return error.InvalidBlock;
+    const block = self.module.blocks.getMut(predecessor) orelse return Error.InvalidBlock;
+    const terminator = if (block.terminator) |*item| item else return Error.InvalidBlock;
 
     switch (terminator.*) {
         .branch => |*edge| {
@@ -278,7 +256,7 @@ fn removeArgumentFromEdges(self: *Self, predecessor: ids.BlockId, target: ids.Bl
 
 fn removeEdgeArgument(self: *Self, edge: *module_ir.Edge, index: usize) !void {
     if (index >= edge.arguments.len)
-        return error.InvalidParameterIndex;
+        return Error.InvalidParameterIndex;
 
     const arguments = try self.module.allocator().alloc(ids.ValueId, edge.arguments.len - 1);
     @memcpy(arguments[0..index], edge.arguments[0..index]);
@@ -299,12 +277,7 @@ fn findIncoming(incoming: []const IncomingValue, predecessor: ids.BlockId) ?ids.
     return null;
 }
 
-fn functionHasEdgeTo(
-    module: *const module_ir.Module,
-    function: *const module_ir.Function,
-    predecessor: ids.BlockId,
-    target: ids.BlockId,
-) bool {
+fn functionHasEdgeTo(module: *const module_ir.Module, function: *const module_ir.Function, predecessor: ids.BlockId, target: ids.BlockId) bool {
     for (function.blocks.items) |block_id| {
         if (block_id != predecessor)
             continue;
@@ -321,3 +294,135 @@ fn countEdgesTo(block: *const module_ir.Block, target: ids.BlockId) usize {
         else => 0,
     };
 }
+
+test "Rewriter: replace all ID uses, safely erase dead instruction" {
+    // shader compute @main
+    // {
+    //     %0: constant u32 = bits(0x1)
+    //     %1: constant u32 = bits(0x2)
+    //
+    //     fn @main() -> void
+    //     {
+    //         .entry():
+    //             %2: u32 = integer_add %0, %1
+    //             %3: u32 = integer_multiply %2, %1
+    //             return
+    //     }
+    // }
+
+    var module = module_ir.Module.init(std.testing.allocator, .compute);
+    defer module.deinit();
+
+    var builder = Builder.init(&module);
+
+    const void_type = try builder.internType(.void);
+    const u32_type = try builder.internType(.{ .integer = .{ .bits = 32, .signedness = .unsigned } });
+
+    const one = try builder.internConstant(u32_type, .{ .integer_bits = 1 });
+    const two = try builder.internConstant(u32_type, .{ .integer_bits = 2 });
+
+    const main = try builder.addFunction(void_type, "main");
+    builder.setEntryPoint(main);
+
+    const entry = try builder.addBlock(main, "entry");
+    const sum = (try builder.appendInstruction(entry, u32_type, .{
+        .binary = .{
+            .opcode = .integer_add,
+            .lhs = one,
+            .rhs = two,
+        },
+    }, null)).?;
+    _ = try builder.appendInstruction(entry, u32_type, .{
+        .binary = .{
+            .opcode = .integer_multiply,
+            .lhs = sum,
+            .rhs = two,
+        },
+    }, null);
+    try builder.setTerminator(entry, .return_void);
+
+    try validator.validate(&module);
+
+    const sum_instruction = module.values.get(sum).?.definition.instruction;
+    var rewriter = Self.init(&module);
+
+    try std.testing.expectEqual(@as(usize, 1), try rewriter.replaceAllUses(sum, one));
+    try rewriter.eraseInstruction(sum_instruction);
+
+    try std.testing.expect(module.values.get(sum) == null);
+    try std.testing.expect(module.instructions.get(sum_instruction) == null);
+
+    try validator.validate(&module);
+}
+
+test "Rewriter: add block parameter and sync branch calls" {
+    // shader compute @main
+    // {
+    //     %0: constant u32 = bits(0x1)
+    //
+    //     fn @main() -> void
+    //     {
+    //         .entry():
+    //             branch .merge()
+    //
+    //         .merge():
+    //             return
+    //
+    //         .alternate():
+    //             return
+    //     }
+    // }
+
+    var module = module_ir.Module.init(std.testing.allocator, .compute);
+    defer module.deinit();
+
+    var builder = Builder.init(&module);
+
+    const void_type = try builder.internType(.void);
+    const u32_type = try builder.internType(.{ .integer = .{ .bits = 32, .signedness = .unsigned } });
+
+    const one = try builder.internConstant(u32_type, .{ .integer_bits = 1 });
+
+    const main = try builder.addFunction(void_type, "main");
+    builder.setEntryPoint(main);
+
+    const entry = try builder.addBlock(main, "entry");
+    const merge = try builder.addBlock(main, "merge");
+    const alternate = try builder.addBlock(main, "alternate");
+
+    try builder.setTerminator(entry, .{ .branch = try builder.edge(merge, &.{}) });
+    try builder.setTerminator(merge, .return_void);
+    try builder.setTerminator(alternate, .return_void);
+
+    var rewriter = Self.init(&module);
+
+    const parameter = try rewriter.addBlockParameter(merge, u32_type, "incoming", &.{
+        .{
+            .predecessor = entry,
+            .value = one,
+        },
+    });
+    const merge_edge = module.blocks.get(entry).?.terminator.?.branch;
+    try std.testing.expectEqualSlices(ids.ValueId, &.{one}, merge_edge.arguments);
+
+    _ = try builder.appendInstruction(merge, u32_type, .{
+        .binary = .{
+            .opcode = .integer_add,
+            .lhs = parameter,
+            .rhs = one,
+        },
+    }, null);
+    try validator.validate(&module);
+
+    try rewriter.removeBlockParameter(merge, 0, one);
+
+    try std.testing.expectEqual(@as(usize, 0), module.blocks.get(merge).?.parameters.items.len);
+    try std.testing.expectEqual(@as(usize, 0), module.blocks.get(entry).?.terminator.?.branch.arguments.len);
+
+    try validator.validate(&module);
+
+    try std.testing.expectEqual(@as(usize, 1), try rewriter.redirectEdges(entry, merge, alternate, &.{}));
+    try std.testing.expectEqual(alternate, module.blocks.get(entry).?.terminator.?.branch.target);
+
+    try validator.validate(&module);
+}

src/compiler/ir/cfg.zig

@@ -19,8 +19,8 @@ reachable: []bool,
 dominators: []bool,
 
 pub fn init(allocator: std.mem.Allocator, module: *const module_ir.Module, function_id: ids.FunctionId) Error!Self {
-    const function = module.functions.get(function_id) orelse return error.InvalidFunction;
-    const entry = function.entry_block orelse return error.MissingEntryBlock;
+    const function = module.functions.get(function_id) orelse return Error.InvalidFunction;
+    const entry = function.entry_block orelse return Error.MissingEntryBlock;
     const blocks = try allocator.dupe(ids.BlockId, function.blocks.items);
     errdefer allocator.free(blocks);
 
@@ -84,8 +84,8 @@ pub fn dominates(self: *const Self, dominator: ids.BlockId, block: ids.BlockId)
 
 fn buildPredecessors(self: *Self, module: *const module_ir.Module) Error!void {
     for (self.blocks) |source| {
-        const block = module.blocks.get(source) orelse return error.InvalidBlock;
-        const terminator = block.terminator orelse return error.MissingTerminator;
+        const block = module.blocks.get(source) orelse return Error.InvalidBlock;
+        const terminator = block.terminator orelse return Error.MissingTerminator;
         switch (terminator) {
             .branch => |edge| try self.addPredecessor(edge.target, source),
             .conditional_branch => |branch| {
@@ -101,12 +101,12 @@ fn buildReachability(self: *Self, module: *const module_ir.Module, entry: ids.Bl
     var queue: std.ArrayList(ids.BlockId) = .empty;
     defer queue.deinit(self.allocator);
     try queue.append(self.allocator, entry);
-    self.reachable[self.indexOf(entry) orelse return error.InvalidBlock] = true;
+    self.reachable[self.indexOf(entry) orelse return Error.InvalidBlock] = true;
 
     var cursor: usize = 0;
     while (cursor < queue.items.len) : (cursor += 1) {
-        const block = module.blocks.get(queue.items[cursor]) orelse return error.InvalidBlock;
-        const terminator = block.terminator orelse return error.MissingTerminator;
+        const block = module.blocks.get(queue.items[cursor]) orelse return Error.InvalidBlock;
+        const terminator = block.terminator orelse return Error.MissingTerminator;
         switch (terminator) {
             .branch => |edge| try self.markReachable(&queue, edge.target),
             .conditional_branch => |branch| {
@@ -168,12 +168,12 @@ fn buildDominators(self: *Self, entry: ids.BlockId) void {
 }
 
 fn addPredecessor(self: *Self, target: ids.BlockId, source: ids.BlockId) Error!void {
-    const target_index = self.indexOf(target) orelse return error.CrossFunctionEdge;
+    const target_index = self.indexOf(target) orelse return Error.CrossFunctionEdge;
     try self.predecessors_by_block[target_index].append(self.allocator, source);
 }
 
 fn markReachable(self: *Self, queue: *std.ArrayList(ids.BlockId), target: ids.BlockId) Error!void {
-    const target_index = self.indexOf(target) orelse return error.CrossFunctionEdge;
+    const target_index = self.indexOf(target) orelse return Error.CrossFunctionEdge;
 
     if (self.reachable[target_index])
         return;

src/compiler/ir/ir.zig

@@ -48,11 +48,22 @@ pub const cfg = @import("cfg.zig");
 pub const constant = @import("constant.zig");
 pub const id = @import("id.zig");
 pub const instruction = @import("instruction.zig");
+pub const inline_all_functions = @import("transformers/inline_all_functions.zig");
 pub const module = @import("module.zig");
 pub const parser = @import("parser/parser.zig");
-pub const pass_manager = @import("pass_manager.zig");
+pub const transformer_manager = @import("transformer_manager.zig");
 pub const printer = @import("printer.zig");
 pub const types = @import("type.zig");
 pub const validator = @import("validator/validator.zig");
 pub const value = @import("value.zig");
 pub const visitor = @import("visitor.zig");
+
+test {
+    _ = Builder;
+    _ = Rewriter;
+    _ = inline_all_functions;
+    _ = module;
+    _ = parser;
+    _ = transformer_manager;
+    _ = validator;
+}

src/compiler/ir/module.zig

@@ -233,3 +233,13 @@ fn replaceOne(operand: *ids.ValueId, old: ids.ValueId, replacement: ids.ValueId,
     operand.* = replacement;
     count.* += 1;
 }
+
+test "Module: central store IDs graveyard" {
+    var module = Module.init(std.testing.allocator, .fragment);
+    defer module.deinit();
+    const first = try module.internType(.boolean);
+    try std.testing.expect(module.types.remove(first));
+    const second = try module.internType(.boolean);
+    try std.testing.expect(first.index() != second.index());
+    try std.testing.expect(module.types.get(first) == null);
+}

src/compiler/ir/parser/parser.zig

@@ -10,23 +10,23 @@ const ast = @import("ast.zig");
 const lowerer = @import("lower.zig");
 
 pub const Error = error{
-    UnexpectedToken,
+    DuplicateName,
+    DuplicateValue,
+    InvalidCompositeIndex,
     InvalidNumber,
+    InvalidOpcode,
+    InvalidResult,
+    InvalidSemantic,
     InvalidStage,
     InvalidType,
-    InvalidOpcode,
-    InvalidSemantic,
-    DuplicateValue,
-    DuplicateName,
-    UnknownValue,
-    UnknownConstant,
-    UnknownInterface,
-    UnknownFunction,
-    UnknownBlock,
-    MissingTerminator,
     MissingResultType,
-    InvalidResult,
-    InvalidCompositeIndex,
+    MissingTerminator,
+    UnexpectedToken,
+    UnknownBlock,
+    UnknownConstant,
+    UnknownFunction,
+    UnknownInterface,
+    UnknownValue,
 };
 
 pub const max_file_size = 64 * 1024 * 1024;
@@ -59,7 +59,7 @@ const Parser = struct {
         try self.expectDiscard(.equal);
 
         const direction_token = try self.expect(.identifier);
-        const direction = std.meta.stringToEnum(module_ir.InterfaceDirection, direction_token.text) orelse return error.InvalidSemantic;
+        const direction = std.meta.stringToEnum(module_ir.InterfaceDirection, direction_token.text) orelse return Error.InvalidSemantic;
         try self.expectDiscard(.left_square);
 
         const semantic_name = (try self.expect(.identifier)).text;
@@ -92,9 +92,9 @@ const Parser = struct {
             const builtin_name = (try self.expect(.identifier)).text;
             try self.expectDiscard(.right_paren);
 
-            const builtin = std.meta.stringToEnum(module_ir.Builtin, builtin_name) orelse return error.InvalidSemantic;
+            const builtin = std.meta.stringToEnum(module_ir.Builtin, builtin_name) orelse return Error.InvalidSemantic;
             break :blk .{ .builtin = builtin };
-        } else return error.InvalidSemantic;
+        } else return Error.InvalidSemantic;
 
         try self.expectDiscard(.right_square);
 
@@ -137,7 +137,7 @@ const Parser = struct {
                     const bits = try self.parseUnsigned(u64, .number);
                     try self.expectDiscard(.right_paren);
 
-                    const ir_type = self.module.?.types.get(ty) orelse return error.InvalidType;
+                    const ir_type = self.module.?.types.get(ty) orelse return Error.InvalidType;
 
                     break :blk switch (ir_type.*) {
                         .integer => .{
@@ -146,16 +146,16 @@ const Parser = struct {
                         .floating => .{
                             .float_bits = bits,
                         },
-                        else => return error.InvalidType,
+                        else => return Error.InvalidType,
                     };
                 }
-                return error.UnexpectedToken;
+                return Error.UnexpectedToken;
             },
             .left_square => .{
                 .composite = try self.parseConstantList(),
             },
             .number => try self.parseDirectConstant(ty),
-            else => return error.UnexpectedToken,
+            else => return Error.UnexpectedToken,
         };
 
         return .{
@@ -167,12 +167,12 @@ const Parser = struct {
 
     fn parseDirectConstant(self: *Parser, ty: ids.TypeId) !ParsedConstantValue {
         const text = (try self.expect(.number)).text;
-        const ir_type = self.module.?.types.get(ty) orelse return error.InvalidType;
+        const ir_type = self.module.?.types.get(ty) orelse return Error.InvalidType;
 
         return switch (ir_type.*) {
             .integer => |integer| .{ .integer_bits = try parseIntegerLiteral(integer, text) },
             .floating => |float| .{ .float_bits = try parseFloatLiteral(float.bits, text) },
-            else => error.InvalidType,
+            else => Error.InvalidType,
         };
     }
 
@@ -210,7 +210,7 @@ const Parser = struct {
 
         while ((try self.peek()).tag != .right_brace) {
             if ((try self.peek()).tag != .dot_name)
-                return error.UnexpectedToken;
+                return Error.UnexpectedToken;
             try function.blocks.append(self.allocator, try self.parseBlock());
         }
 
@@ -246,7 +246,7 @@ const Parser = struct {
             const token = try self.peek();
 
             if (token.tag == .right_brace or token.tag == .dot_name)
-                return error.MissingTerminator;
+                return Error.MissingTerminator;
 
             if (token.tag == .value_ref) {
                 try block.instructions.append(self.allocator, try self.parseInstruction(true));
@@ -254,7 +254,7 @@ const Parser = struct {
             }
 
             if (token.tag != .identifier)
-                return error.UnexpectedToken;
+                return Error.UnexpectedToken;
 
             if (isTerminatorName(token.text)) {
                 block.terminator = try self.parseTerminator();
@@ -285,7 +285,7 @@ const Parser = struct {
 
         if (std.mem.startsWith(u8, name, "cmp_")) {
             const opcode_name = name["cmp_".len..];
-            const opcode = std.meta.stringToEnum(inst_ir.CompareOpcode, opcode_name) orelse return error.InvalidOpcode;
+            const opcode = std.meta.stringToEnum(inst_ir.CompareOpcode, opcode_name) orelse return Error.InvalidOpcode;
             const lhs = try self.parseValueRef();
 
             try self.expectDiscard(.comma);
@@ -337,7 +337,7 @@ const Parser = struct {
             }
 
             if (indices.items.len == 0)
-                return error.InvalidCompositeIndex;
+                return Error.InvalidCompositeIndex;
             return .{
                 .composite_extract = .{
                     .composite = composite,
@@ -377,7 +377,7 @@ const Parser = struct {
             };
         }
 
-        return error.InvalidOpcode;
+        return Error.InvalidOpcode;
     }
 
     fn parseTerminator(self: *Parser) !ParsedTerminator {
@@ -414,7 +414,7 @@ const Parser = struct {
         if (std.mem.eql(u8, name, "unreachable"))
             return .unreachable_value;
 
-        return error.UnexpectedToken;
+        return Error.UnexpectedToken;
     }
 
     fn parseEdge(self: *Parser) !ParsedEdge {
@@ -441,7 +441,7 @@ const Parser = struct {
             return module.internType(.boolean);
 
         if (std.mem.startsWith(u8, token.text, "vec")) {
-            const length = parseTextUnsigned(u8, token.text[3..]) catch return error.InvalidType;
+            const length = parseTextUnsigned(u8, token.text[3..]) catch return Error.InvalidType;
             try self.expectDiscard(.left_square);
 
             const element_type = try self.parseType();
@@ -497,7 +497,7 @@ const Parser = struct {
             try self.expectDiscard(.left_square);
 
             const address_name = (try self.expect(.identifier)).text;
-            const address_space = std.meta.stringToEnum(type_ir.AddressSpace, address_name) orelse return error.InvalidType;
+            const address_space = std.meta.stringToEnum(type_ir.AddressSpace, address_name) orelse return Error.InvalidType;
             try self.expectDiscard(.comma);
 
             const pointee_type = try self.parseType();
@@ -515,7 +515,7 @@ const Parser = struct {
             try self.expectDiscard(.left_square);
 
             const kind_name = (try self.expect(.identifier)).text;
-            const kind = std.meta.stringToEnum(type_ir.ResourceKind, kind_name) orelse return error.InvalidType;
+            const kind = std.meta.stringToEnum(type_ir.ResourceKind, kind_name) orelse return Error.InvalidType;
 
             try self.expectDiscard(.right_square);
 
@@ -527,7 +527,7 @@ const Parser = struct {
         }
 
         if (token.text.len > 1 and (token.text[0] == 'i' or token.text[0] == 'u')) {
-            const bits = parseTextUnsigned(u16, token.text[1..]) catch return error.InvalidType;
+            const bits = parseTextUnsigned(u16, token.text[1..]) catch return Error.InvalidType;
             return module.internType(.{ .integer = .{
                 .bits = bits,
                 .signedness = if (token.text[0] == 'i') .signed else .unsigned,
@@ -535,14 +535,14 @@ const Parser = struct {
         }
 
         if (token.text.len > 1 and token.text[0] == 'f') {
-            const bits = parseTextUnsigned(u16, token.text[1..]) catch return error.InvalidType;
+            const bits = parseTextUnsigned(u16, token.text[1..]) catch return Error.InvalidType;
             return module.internType(.{
                 .floating = .{
                     .bits = bits,
                 },
             });
         }
-        return error.InvalidType;
+        return Error.InvalidType;
     }
 
     fn parseConstantList(self: *Parser) ![]const u32 {
@@ -595,13 +595,13 @@ const Parser = struct {
 
     fn parseUnsigned(self: *Parser, comptime T: type, tag: TokenTag) !T {
         const token = try self.expect(tag);
-        return parseTextUnsigned(T, token.text) catch error.InvalidNumber;
+        return parseTextUnsigned(T, token.text) catch Error.InvalidNumber;
     }
 
     fn expectIdentifier(self: *Parser, expected: []const u8) !void {
         const token = try self.expect(.identifier);
         if (!std.mem.eql(u8, token.text, expected))
-            return error.UnexpectedToken;
+            return Error.UnexpectedToken;
     }
 
     fn expectDiscard(self: *Parser, tag: TokenTag) !void {
@@ -611,7 +611,7 @@ const Parser = struct {
     fn expect(self: *Parser, tag: TokenTag) !Token {
         const token = try self.take();
         if (token.tag != tag)
-            return error.UnexpectedToken;
+            return Error.UnexpectedToken;
         return token;
     }
 
@@ -629,7 +629,7 @@ const Parser = struct {
     fn take(self: *Parser) !Token {
         const token = self.lexer.take();
         if (token.tag == .invalid)
-            return error.UnexpectedToken;
+            return Error.UnexpectedToken;
         return token;
     }
 };
@@ -644,25 +644,25 @@ fn isTerminatorName(name: []const u8) bool {
 
 fn parseIntegerLiteral(integer: type_ir.IntegerType, text: []const u8) !u64 {
     if (integer.bits == 0 or integer.bits > 64)
-        return error.InvalidType;
+        return Error.InvalidType;
 
     if (integer.signedness == .unsigned) {
-        const value = std.fmt.parseInt(u64, text, 10) catch return error.InvalidNumber;
+        const value = std.fmt.parseInt(u64, text, 10) catch return Error.InvalidNumber;
         if (integer.bits < 64) {
             const shift: u6 = @intCast(integer.bits);
             const maximum = (@as(u64, 1) << shift) - 1;
             if (value > maximum)
-                return error.InvalidNumber;
+                return Error.InvalidNumber;
         }
         return value;
     }
 
-    const value = std.fmt.parseInt(i64, text, 10) catch return error.InvalidNumber;
+    const value = std.fmt.parseInt(i64, text, 10) catch return Error.InvalidNumber;
     if (integer.bits < 64) {
         const sign_shift: u6 = @intCast(integer.bits - 1);
         const magnitude = @as(i64, 1) << sign_shift;
         if (value < -magnitude or value > magnitude - 1)
-            return error.InvalidNumber;
+            return Error.InvalidNumber;
 
         const width: u6 = @intCast(integer.bits);
         const mask = (@as(u64, 1) << width) - 1;
@@ -675,18 +675,18 @@ fn parseIntegerLiteral(integer: type_ir.IntegerType, text: []const u8) !u64 {
 fn parseFloatLiteral(bits: u16, text: []const u8) !u64 {
     return switch (bits) {
         16 => blk: {
-            const value = std.fmt.parseFloat(f16, text) catch return error.InvalidNumber;
+            const value = std.fmt.parseFloat(f16, text) catch return Error.InvalidNumber;
             break :blk @as(u16, @bitCast(value));
         },
         32 => blk: {
-            const value = std.fmt.parseFloat(f32, text) catch return error.InvalidNumber;
+            const value = std.fmt.parseFloat(f32, text) catch return Error.InvalidNumber;
             break :blk @as(u32, @bitCast(value));
         },
         64 => blk: {
-            const value = std.fmt.parseFloat(f64, text) catch return error.InvalidNumber;
+            const value = std.fmt.parseFloat(f64, text) catch return Error.InvalidNumber;
             break :blk @as(u64, @bitCast(value));
         },
-        else => error.InvalidType,
+        else => Error.InvalidType,
     };
 }
 
@@ -695,7 +695,7 @@ fn parseTextUnsigned(comptime T: type, text: []const u8) !T {
     const digits = if (base == 16) text[2..] else text;
 
     if (digits.len == 0)
-        return error.InvalidNumber;
+        return Error.InvalidNumber;
 
     return std.fmt.parseInt(T, digits, base);
 }
@@ -712,7 +712,7 @@ pub fn parseString(backing_allocator: std.mem.Allocator, source: []const u8) !mo
 
     try parser.expectIdentifier("shader");
     const stage_token = try parser.expect(.identifier);
-    const stage = std.meta.stringToEnum(module_ir.Stage, stage_token.text) orelse return error.InvalidStage;
+    const stage = std.meta.stringToEnum(module_ir.Stage, stage_token.text) orelse return Error.InvalidStage;
 
     var module = module_ir.Module.init(backing_allocator, stage);
     errdefer module.deinit();
@@ -735,10 +735,10 @@ pub fn parseString(backing_allocator: std.mem.Allocator, source: []const u8) !mo
                 if (std.mem.eql(u8, token.text, "fn")) {
                     try parsed.functions.append(temporary_allocator, try parser.parseFunction());
                 } else {
-                    return error.UnexpectedToken;
+                    return Error.UnexpectedToken;
                 }
             },
-            else => return error.UnexpectedToken,
+            else => return Error.UnexpectedToken,
         }
     }
     try parser.expectDiscard(.right_brace);
@@ -765,3 +765,194 @@ pub fn parseFileInDir(backing_allocator: std.mem.Allocator, io: std.Io, director
 
     return parseString(backing_allocator, source);
 }
+
+test "Parser: interface" {
+    const printer = @import("../printer.zig");
+
+    const source =
+        \\ shader vertex @main
+        \\ {
+        \\     @in_color: vec4[f32] = input[location(0), component(0), index(0)]
+        \\     @out_color: vec4[f32] = output[location(0), component(0), index(0)]
+        \\     @position: vec4[f32] = output[builtin(position)]
+        \\
+        \\     fn @main() -> void
+        \\     {
+        \\         .entry():
+        \\             return
+        \\     }
+        \\ }
+    ;
+
+    var module = try parseString(std.testing.allocator, source);
+    defer module.deinit();
+    const printed = try printer.allocPrint(std.testing.allocator, &module);
+    defer std.testing.allocator.free(printed);
+
+    try std.testing.expect(std.mem.indexOf(u8, printed, "@in_color: vec4[f32] = input[location(0), component(0), index(0)]") != null);
+    try std.testing.expect(std.mem.indexOf(u8, printed, "@out_color: vec4[f32] = output[location(0), component(0), index(0)]") != null);
+    try std.testing.expect(std.mem.indexOf(u8, printed, "@position: vec4[f32] = output[builtin(position)]") != null);
+}
+
+test "Parser: types, operations, calls, terminators" {
+    const printer = @import("../printer.zig");
+
+    const source =
+        \\ shader fragment @main
+        \\ {
+        \\     %0: constant bool = true
+        \\     %1: constant u32 = bits(0x1)
+        \\     %2: constant u32 = bits(0x2)
+        \\     %3: constant f32 = bits(0x3f800000)
+        \\     %4: constant array[u32, 2] = [#1, #2]
+        \\     %5: constant struct[u32, u32] = [#1, #2]
+        \\     %6: constant ptr[private, u32] = null
+        \\     %7: constant resourceHandle[sampler] = null
+        \\
+        \\     fn @main() -> void
+        \\     {
+        \\         .entry():
+        \\             %9: u32 = bitwise_not %1
+        \\             %10: u32 = integer_add %9, %2
+        \\             %11: bool = cmp_equal %1, %2
+        \\             %12: u32 = select %11, %1, %2
+        \\             %13: u32 = bitcast %12
+        \\             %14: vec2[u32] = composite_construct %1, %2
+        \\             %15: u32 = composite_extract %14[0]
+        \\             %16: f32 = negate %3
+        \\             %17: f32 = float_add %3, %16
+        \\             %18: u32 = call @helper(%15)
+        \\             return
+        \\     }
+        \\
+        \\     fn @helper(%8: u32) -> u32
+        \\     {
+        \\         .entry():
+        \\             return %8
+        \\     }
+        \\
+        \\     fn @discarder() -> void
+        \\     {
+        \\         .entry():
+        \\             discard
+        \\     }
+        \\
+        \\     fn @dead() -> void
+        \\     {
+        \\         .entry():
+        \\             unreachable
+        \\     }
+        \\ }
+    ;
+
+    var module = try parseString(std.testing.allocator, source);
+    defer module.deinit();
+    const printed = try printer.allocPrint(std.testing.allocator, &module);
+    defer std.testing.allocator.free(printed);
+
+    var reparsed = try parseString(std.testing.allocator, printed);
+    defer reparsed.deinit();
+    const printed_again = try printer.allocPrint(std.testing.allocator, &reparsed);
+    defer std.testing.allocator.free(printed_again);
+    try std.testing.expectEqualStrings(printed, printed_again);
+    try std.testing.expect(std.mem.indexOf(u8, printed, "cmp_equal %1, %2") != null);
+    try std.testing.expect(std.mem.indexOf(u8, printed, "cmp.") == null);
+}
+
+test "Parser: named value IDs" {
+    const printer = @import("../printer.zig");
+
+    const source =
+        \\ shader compute @main
+        \\ {
+        \\     %one_value: constant u32 = bits(0x1)
+        \\
+        \\     fn @main() -> void
+        \\     {
+        \\         .entry():
+        \\             %sum_value: u32 = integer_add %one_value, %one_value
+        \\             branch .merge(%sum_value)
+        \\
+        \\         .merge(%merged_value: u32):
+        \\             %product_value: u32 = integer_multiply %merged_value, %one_value
+        \\             return
+        \\     }
+        \\ }
+    ;
+
+    var module = try parseString(std.testing.allocator, source);
+    defer module.deinit();
+    const printed = try printer.allocPrint(std.testing.allocator, &module);
+    defer std.testing.allocator.free(printed);
+
+    try std.testing.expect(std.mem.indexOf(u8, printed, "%one_value: constant u32 = bits(0x1)") != null);
+    try std.testing.expect(std.mem.indexOf(u8, printed, "%sum_value: u32 = integer_add %one_value, %one_value") != null);
+    try std.testing.expect(std.mem.indexOf(u8, printed, "branch .merge(%sum_value)") != null);
+    try std.testing.expect(std.mem.indexOf(u8, printed, ".merge(%merged_value: u32)") != null);
+    try std.testing.expect(std.mem.indexOf(u8, printed, "%product_value: u32 = integer_multiply %merged_value, %one_value") != null);
+
+    var reparsed = try parseString(std.testing.allocator, printed);
+    defer reparsed.deinit();
+    const printed_again = try printer.allocPrint(std.testing.allocator, &reparsed);
+    defer std.testing.allocator.free(printed_again);
+    try std.testing.expectEqualStrings(printed, printed_again);
+}
+
+test "Parser: numeric constants" {
+    const printer = @import("../printer.zig");
+
+    const source =
+        \\ shader compute @main
+        \\ {
+        \\     %0: constant u8 = 255
+        \\     %1: constant i8 = -1
+        \\     %2: constant f16 = 1.5
+        \\     %3: constant f32 = -0.0
+        \\     %4: constant f64 = 2.5e0
+        \\
+        \\     fn @main() -> void
+        \\     {
+        \\         .entry():
+        \\             return
+        \\     }
+        \\ }
+    ;
+
+    var module = try parseString(std.testing.allocator, source);
+    defer module.deinit();
+    const printed = try printer.allocPrint(std.testing.allocator, &module);
+    defer std.testing.allocator.free(printed);
+
+    try std.testing.expect(std.mem.indexOf(u8, printed, "%0: constant u8 = bits(0xff)") != null);
+    try std.testing.expect(std.mem.indexOf(u8, printed, "%1: constant i8 = bits(0xff)") != null);
+    try std.testing.expect(std.mem.indexOf(u8, printed, "%2: constant f16 = bits(0x3e00)") != null);
+    try std.testing.expect(std.mem.indexOf(u8, printed, "%3: constant f32 = bits(0x80000000)") != null);
+    try std.testing.expect(std.mem.indexOf(u8, printed, "%4: constant f64 = bits(0x4004000000000000)") != null);
+
+    const out_of_range =
+        \\ shader compute @main
+        \\ {
+        \\     %0: constant u8 = 256
+        \\     fn @main() -> void
+        \\     {
+        \\         .entry():
+        \\             return
+        \\     }
+        \\ }
+    ;
+    try std.testing.expectError(Error.InvalidNumber, parseString(std.testing.allocator, out_of_range));
+}
+
+test "Parser: Error unknown value" {
+    const source =
+        \\ shader compute @main
+        \\ {
+        \\     fn @main() -> void
+        \\     {
+        \\         .entry():
+        \\             return %99
+        \\     }
+        \\ }
+    ;
+    try std.testing.expectError(Error.UnknownValue, parseString(std.testing.allocator, source));
+}

src/compiler/ir/pass_manager.zig

@@ -1,83 +0,0 @@
-const std = @import("std");
-const module_ir = @import("module.zig");
-const validator = @import("validator/validator.zig");
-
-pub const Context = struct {
-    allocator: std.mem.Allocator,
-    validate_after_each_pass: bool = true,
-};
-
-pub const Pass = struct {
-    name: []const u8,
-    required: module_ir.Properties = .{},
-    produced: module_ir.Properties = .{},
-    invalidated: module_ir.Properties = .{},
-    run: *const fn (module: *module_ir.Module, context: *Context) anyerror!bool,
-};
-
-pub const Manager = struct {
-    allocator: std.mem.Allocator,
-    passes: std.ArrayList(Pass) = .empty,
-
-    pub fn init(allocator: std.mem.Allocator) Manager {
-        return .{ .allocator = allocator };
-    }
-
-    pub fn deinit(self: *Manager) void {
-        self.passes.deinit(self.allocator);
-        self.* = undefined;
-    }
-
-    pub fn add(self: *Manager, pass: Pass) !void {
-        try self.passes.append(self.allocator, pass);
-    }
-
-    pub fn run(self: *Manager, module: *module_ir.Module, context: *Context) !bool {
-        var changed = false;
-        for (self.passes.items) |pass| {
-            if (!satisfies(module.properties, pass.required))
-                return error.RequiredPropertyMissing;
-
-            changed = (try pass.run(module, context)) or changed;
-            applyInvalidated(&module.properties, pass.invalidated);
-            applyProduced(&module.properties, pass.produced);
-
-            if (context.validate_after_each_pass)
-                try validator.validate(module);
-        }
-        return changed;
-    }
-};
-
-fn satisfies(actual: module_ir.Properties, required: module_ir.Properties) bool {
-    inline for (property_names) |name| {
-        if (@field(required, name) and !@field(actual, name))
-            return false;
-    }
-    return true;
-}
-
-fn applyProduced(properties: *module_ir.Properties, produced: module_ir.Properties) void {
-    inline for (property_names) |name| {
-        if (@field(produced, name))
-            @field(properties, name) = true;
-    }
-}
-
-fn applyInvalidated(properties: *module_ir.Properties, invalidated: module_ir.Properties) void {
-    inline for (property_names) |name| {
-        if (@field(invalidated, name))
-            @field(properties, name) = false;
-    }
-}
-
-const property_names = .{
-    "valid_cfg",
-    "valid_ssa",
-    "structured_control_flow",
-    "no_function_calls",
-    "no_local_memory",
-    "no_matrix_types",
-    "no_large_composites",
-    "explicit_resource_offsets",
-};

src/compiler/ir/transformer_manager.zig

@@ -0,0 +1,157 @@
+const std = @import("std");
+const ids = @import("id.zig");
+const module_ir = @import("module.zig");
+const Builder = @import("Builder.zig");
+const validator = @import("validator/validator.zig");
+const visitor = @import("visitor.zig");
+
+pub const Context = struct {
+    allocator: std.mem.Allocator,
+    validate_after_each_transformer: bool = true,
+};
+
+pub const Transformer = struct {
+    name: []const u8,
+    required: module_ir.Properties = .{},
+    produced: module_ir.Properties = .{},
+    invalidated: module_ir.Properties = .{},
+    run: *const fn (module: *module_ir.Module, context: *Context) anyerror!bool,
+};
+
+pub const Manager = struct {
+    allocator: std.mem.Allocator,
+    transformers: std.ArrayList(Transformer) = .empty,
+
+    pub fn init(allocator: std.mem.Allocator) Manager {
+        return .{ .allocator = allocator };
+    }
+
+    pub fn deinit(self: *Manager) void {
+        self.transformers.deinit(self.allocator);
+        self.* = undefined;
+    }
+
+    pub fn add(self: *Manager, transformer: Transformer) !void {
+        try self.transformers.append(self.allocator, transformer);
+    }
+
+    pub fn run(self: *Manager, module: *module_ir.Module, context: *Context) !bool {
+        var changed = false;
+        for (self.transformers.items) |transformer| {
+            if (!satisfies(module.properties, transformer.required))
+                return error.RequiredPropertyMissing;
+
+            changed = (try transformer.run(module, context)) or changed;
+            applyInvalidated(&module.properties, transformer.invalidated);
+            applyProduced(&module.properties, transformer.produced);
+
+            if (context.validate_after_each_transformer)
+                try validator.validate(module);
+        }
+        return changed;
+    }
+};
+
+fn satisfies(actual: module_ir.Properties, required: module_ir.Properties) bool {
+    inline for (property_names) |name| {
+        if (@field(required, name) and !@field(actual, name))
+            return false;
+    }
+    return true;
+}
+
+fn applyProduced(properties: *module_ir.Properties, produced: module_ir.Properties) void {
+    inline for (property_names) |name| {
+        if (@field(produced, name))
+            @field(properties, name) = true;
+    }
+}
+
+fn applyInvalidated(properties: *module_ir.Properties, invalidated: module_ir.Properties) void {
+    inline for (property_names) |name| {
+        if (@field(invalidated, name))
+            @field(properties, name) = false;
+    }
+}
+
+const property_names = .{
+    "valid_cfg",
+    "valid_ssa",
+    "structured_control_flow",
+    "no_function_calls",
+    "no_local_memory",
+    "no_matrix_types",
+    "no_large_composites",
+    "explicit_resource_offsets",
+};
+
+const VisitorStatistics = struct {
+    functions: usize = 0,
+    blocks: usize = 0,
+};
+
+fn establishNoCalls(_: *module_ir.Module, _: *Context) !bool {
+    return false;
+}
+
+fn countVisitedFunction(context: ?*anyopaque, _: ids.FunctionId, _: *const module_ir.Function) !void {
+    const statistics: *VisitorStatistics = @ptrCast(@alignCast(context.?));
+    statistics.functions += 1;
+}
+
+fn countVisitedBlock(context: ?*anyopaque, _: ids.BlockId, _: *const module_ir.Block) !void {
+    const statistics: *VisitorStatistics = @ptrCast(@alignCast(context.?));
+    statistics.blocks += 1;
+}
+
+test "Transformers manager: tracks independent IR properties" {
+    // shader compute @main
+    // {
+    //     fn @main() -> void
+    //     {
+    //         .entry():
+    //             return
+    //     }
+    // }
+
+    var module = module_ir.Module.init(std.testing.allocator, .compute);
+    defer module.deinit();
+
+    var builder = Builder.init(&module);
+
+    const void_type = try builder.internType(.void);
+
+    const main = try builder.addFunction(void_type, "main");
+    builder.setEntryPoint(main);
+
+    const entry = try builder.addBlock(main, "entry");
+    try builder.setTerminator(entry, .return_void);
+
+    module.properties.valid_cfg = true;
+
+    var manager = Manager.init(std.testing.allocator);
+    defer manager.deinit();
+
+    try manager.add(.{
+        .name = "establish-no-calls",
+        .required = .{ .valid_cfg = true },
+        .produced = .{ .no_function_calls = true },
+        .run = establishNoCalls,
+    });
+
+    var context: Context = .{ .allocator = std.testing.allocator };
+
+    try std.testing.expect(!try manager.run(&module, &context));
+    try std.testing.expect(module.properties.no_function_calls);
+
+    var statistics: VisitorStatistics = .{};
+
+    try visitor.walk(&module, .{
+        .context = &statistics,
+        .visitFunction = countVisitedFunction,
+        .visitBlock = countVisitedBlock,
+    });
+
+    try std.testing.expectEqual(@as(usize, 1), statistics.functions);
+    try std.testing.expectEqual(@as(usize, 1), statistics.blocks);
+}

src/compiler/ir/transformers/inline_all_functions.zig

@@ -0,0 +1,751 @@
+const std = @import("std");
+const Builder = @import("../Builder.zig");
+const Rewriter = @import("../Rewriter.zig");
+const ids = @import("../id.zig");
+const instruction_ir = @import("../instruction.zig");
+const module_ir = @import("../module.zig");
+const transformer_manager = @import("../transformer_manager.zig");
+
+pub const Error = Rewriter.Error || error{
+    InvalidModule,
+    RecursiveCall,
+    UnsupportedEntryBlockParameters,
+};
+
+const VisitState = enum {
+    unvisited,
+    visiting,
+    complete,
+};
+
+const CallGraph = struct {
+    module: *const module_ir.Module,
+    states: []VisitState,
+    postorder: *std.ArrayList(ids.FunctionId),
+    allocator: std.mem.Allocator,
+
+    fn visit(self: *CallGraph, function_id: ids.FunctionId) Error!void {
+        if (function_id.index() >= self.states.len)
+            return Error.InvalidModule;
+
+        switch (self.states[function_id.index()]) {
+            .visiting => return Error.RecursiveCall,
+            .complete => return,
+            .unvisited => {},
+        }
+        self.states[function_id.index()] = .visiting;
+
+        const function = self.module.functions.get(function_id) orelse return Error.InvalidModule;
+        for (function.blocks.items) |block_id| {
+            const block = self.module.blocks.get(block_id) orelse return Error.InvalidModule;
+            for (block.instructions.items) |instruction_id| {
+                const inst = self.module.instructions.get(instruction_id) orelse return Error.InvalidModule;
+                switch (inst.operation) {
+                    .call => |call| try self.visit(call.function),
+                    else => {},
+                }
+            }
+        }
+
+        self.states[function_id.index()] = .complete;
+        try self.postorder.append(self.allocator, function_id);
+    }
+};
+
+pub const transformer: transformer_manager.Transformer = .{
+    .name = "inline-all-functions",
+    .produced = .{ .no_function_calls = true },
+    .invalidated = .{ .structured_control_flow = true },
+    .run = transform,
+};
+
+fn transform(module: *module_ir.Module, context: *transformer_manager.Context) !bool {
+    const scratch_allocator = context.allocator;
+    const entry_point = module.entry_point orelse return Error.InvalidModule;
+    if (!module.functions.isLive(entry_point))
+        return Error.InvalidModule;
+
+    var changed = try removeUnreachableBlocks(module, scratch_allocator);
+
+    const states = try scratch_allocator.alloc(VisitState, module.functions.entries.items.len);
+    defer scratch_allocator.free(states);
+    @memset(states, .unvisited);
+
+    var postorder: std.ArrayList(ids.FunctionId) = .empty;
+    defer postorder.deinit(scratch_allocator);
+
+    var call_graph: CallGraph = .{
+        .module = module,
+        .states = states,
+        .postorder = &postorder,
+        .allocator = scratch_allocator,
+    };
+    try call_graph.visit(entry_point);
+
+    for (postorder.items) |function_id|
+        changed = (try inlineCallsInFunction(module, scratch_allocator, function_id)) or changed;
+
+    changed = removeNonEntryFunctions(module, entry_point) or changed;
+
+    for (module.instructions.entries.items) |entry| {
+        const inst = entry orelse continue;
+        if (inst.operation == .call)
+            return Error.InvalidModule;
+    }
+
+    return changed;
+}
+
+fn inlineCallsInFunction(module: *module_ir.Module, scratch_allocator: std.mem.Allocator, function_id: ids.FunctionId) Error!bool {
+    const function = module.functions.get(function_id) orelse return Error.InvalidModule;
+    var calls: std.ArrayList(ids.InstructionId) = .empty;
+    defer calls.deinit(scratch_allocator);
+
+    for (function.blocks.items) |block_id| {
+        const block = module.blocks.get(block_id) orelse return Error.InvalidModule;
+        for (block.instructions.items) |instruction_id| {
+            const inst = module.instructions.get(instruction_id) orelse return Error.InvalidModule;
+            if (inst.operation == .call)
+                try calls.append(scratch_allocator, instruction_id);
+        }
+    }
+
+    for (calls.items) |call_id|
+        try inlineCall(module, scratch_allocator, function_id, call_id);
+
+    return calls.items.len != 0;
+}
+
+fn inlineCall(module: *module_ir.Module, scratch_allocator: std.mem.Allocator, caller_id: ids.FunctionId, call_id: ids.InstructionId) Error!void {
+    const call_inst = module.instructions.get(call_id) orelse return Error.InvalidModule;
+    if (call_inst.operation != .call)
+        return Error.InvalidModule;
+
+    const call = call_inst.operation.call;
+    if (call.function == caller_id)
+        return Error.RecursiveCall;
+
+    const callee = module.functions.get(call.function) orelse return Error.InvalidModule;
+    const callee_entry = module.blocks.get(callee.entry_block orelse return Error.InvalidModule) orelse return Error.InvalidModule;
+    if (callee_entry.parameters.items.len != 0)
+        return Error.UnsupportedEntryBlockParameters;
+
+    const caller_block_id = call_inst.parent_block;
+    const caller_block = module.blocks.get(caller_block_id) orelse return Error.InvalidModule;
+    if (caller_block.parent_function != caller_id)
+        return Error.InvalidModule;
+
+    var call_index: ?usize = null;
+    for (caller_block.instructions.items, 0..) |instruction_id, index| {
+        if (instruction_id == call_id) {
+            call_index = index;
+            break;
+        }
+    }
+    const index = call_index orelse return Error.InvalidModule;
+
+    const arguments = try scratch_allocator.dupe(ids.ValueId, call.arguments);
+    defer scratch_allocator.free(arguments);
+
+    const suffix = try scratch_allocator.dupe(ids.InstructionId, caller_block.instructions.items[index + 1 ..]);
+    defer scratch_allocator.free(suffix);
+
+    const old_terminator = caller_block.terminator orelse return Error.InvalidModule;
+    const old_structured_control = caller_block.structured_control;
+    const call_result = call_inst.result;
+    const result_type = if (call_result) |result_id|
+        (module.values.get(result_id) orelse return Error.InvalidModule).type
+    else
+        null;
+    const result_name = if (call_result) |result_id|
+        (module.values.get(result_id) orelse return Error.InvalidModule).name
+    else
+        null;
+
+    var builder = Builder.init(module);
+    const continuation_id = try builder.addBlock(caller_id, null);
+    const continuation_result = if (result_type) |ty|
+        try builder.addBlockParameter(continuation_id, ty, result_name)
+    else
+        null;
+
+    {
+        const continuation = module.blocks.getMut(continuation_id) orelse return Error.InvalidModule;
+        try continuation.instructions.appendSlice(module.allocator(), suffix);
+        continuation.terminator = old_terminator;
+        continuation.structured_control = switch (old_structured_control) {
+            .loop => .none,
+            else => old_structured_control,
+        };
+    }
+    for (suffix) |instruction_id| {
+        const moved = module.instructions.getMut(instruction_id) orelse return Error.InvalidModule;
+        moved.parent_block = continuation_id;
+    }
+
+    {
+        const mutable_caller_block = module.blocks.getMut(caller_block_id) orelse return Error.InvalidModule;
+        mutable_caller_block.instructions.shrinkRetainingCapacity(index);
+        mutable_caller_block.terminator = null;
+        mutable_caller_block.structured_control = switch (old_structured_control) {
+            .loop => old_structured_control,
+            else => .none,
+        };
+    }
+
+    if (call_result) |old_result| {
+        const replacement = continuation_result orelse return Error.InvalidModule;
+        var rewriter = Rewriter.init(module);
+        _ = try rewriter.replaceAllUses(old_result, replacement);
+    }
+
+    const cloned_entry = try cloneCallee(
+        module,
+        scratch_allocator,
+        caller_id,
+        call.function,
+        arguments,
+        continuation_id,
+        continuation_result != null,
+    );
+
+    const cloned_entry_block = module.blocks.get(cloned_entry) orelse return Error.InvalidModule;
+    if (cloned_entry_block.parameters.items.len != 0)
+        return Error.UnsupportedEntryBlockParameters;
+
+    const branch_arguments = try module.allocator().alloc(ids.ValueId, 0);
+    const mutable_caller_block = module.blocks.getMut(caller_block_id) orelse return Error.InvalidModule;
+    mutable_caller_block.terminator = .{ .branch = .{
+        .target = cloned_entry,
+        .arguments = branch_arguments,
+    } };
+
+    if (call_result) |result_id|
+        _ = module.values.remove(result_id);
+    _ = module.instructions.remove(call_id);
+}
+
+fn cloneCallee(
+    module: *module_ir.Module,
+    scratch_allocator: std.mem.Allocator,
+    caller_id: ids.FunctionId,
+    callee_id: ids.FunctionId,
+    arguments: []const ids.ValueId,
+    continuation_id: ids.BlockId,
+    returns_value: bool,
+) Error!ids.BlockId {
+    const callee = module.functions.get(callee_id) orelse return Error.InvalidModule;
+    if (callee.parameters.items.len != arguments.len)
+        return Error.InvalidModule;
+
+    const callee_blocks = try scratch_allocator.dupe(ids.BlockId, callee.blocks.items);
+    defer scratch_allocator.free(callee_blocks);
+    const callee_entry = callee.entry_block orelse return Error.InvalidModule;
+
+    const block_map = try scratch_allocator.alloc(?ids.BlockId, module.blocks.entries.items.len);
+    defer scratch_allocator.free(block_map);
+    @memset(block_map, null);
+
+    const value_map = try scratch_allocator.alloc(?ids.ValueId, module.values.entries.items.len);
+    defer scratch_allocator.free(value_map);
+    @memset(value_map, null);
+
+    const instruction_map = try scratch_allocator.alloc(?ids.InstructionId, module.instructions.entries.items.len);
+    defer scratch_allocator.free(instruction_map);
+    @memset(instruction_map, null);
+
+    for (callee.parameters.items, arguments) |parameter, argument| {
+        if (parameter.index() >= value_map.len)
+            return Error.InvalidModule;
+        value_map[parameter.index()] = argument;
+    }
+
+    var builder = Builder.init(module);
+    for (callee_blocks) |old_block_id| {
+        const new_block_id = try builder.addBlock(caller_id, null);
+        if (old_block_id.index() >= block_map.len)
+            return Error.InvalidModule;
+        block_map[old_block_id.index()] = new_block_id;
+    }
+
+    for (callee_blocks) |old_block_id| {
+        const old_block = module.blocks.get(old_block_id) orelse return Error.InvalidModule;
+        const new_block_id = mappedBlock(block_map, old_block_id) orelse return Error.InvalidModule;
+        for (old_block.parameters.items) |old_parameter| {
+            const old_value = module.values.get(old_parameter) orelse return Error.InvalidModule;
+            const new_parameter = try builder.addBlockParameter(new_block_id, old_value.type, null);
+            if (old_parameter.index() >= value_map.len)
+                return Error.InvalidModule;
+            value_map[old_parameter.index()] = new_parameter;
+        }
+    }
+
+    for (callee_blocks) |old_block_id| {
+        const old_block = module.blocks.get(old_block_id) orelse return Error.InvalidModule;
+        const new_block_id = mappedBlock(block_map, old_block_id) orelse return Error.InvalidModule;
+
+        for (old_block.instructions.items) |old_instruction_id| {
+            const old_instruction = (module.instructions.get(old_instruction_id) orelse return Error.InvalidModule).*;
+            const new_instruction_id = try module.instructions.add(module.allocator(), .{
+                .parent_block = new_block_id,
+                .result = null,
+                .operation = old_instruction.operation,
+                .source = old_instruction.source,
+            });
+            if (old_instruction_id.index() >= instruction_map.len)
+                return Error.InvalidModule;
+            instruction_map[old_instruction_id.index()] = new_instruction_id;
+
+            if (old_instruction.result) |old_result| {
+                const old_value = module.values.get(old_result) orelse return Error.InvalidModule;
+                const new_result = try module.values.add(module.allocator(), .{
+                    .type = old_value.type,
+                    .definition = .{ .instruction = new_instruction_id },
+                    .name = null,
+                });
+                module.instructions.getMut(new_instruction_id).?.result = new_result;
+                if (old_result.index() >= value_map.len)
+                    return Error.InvalidModule;
+                value_map[old_result.index()] = new_result;
+            }
+
+            const new_block = module.blocks.getMut(new_block_id) orelse return Error.InvalidModule;
+            try new_block.instructions.append(module.allocator(), new_instruction_id);
+        }
+    }
+
+    for (callee_blocks) |old_block_id| {
+        const old_block = module.blocks.get(old_block_id) orelse return Error.InvalidModule;
+        const new_block_id = mappedBlock(block_map, old_block_id) orelse return Error.InvalidModule;
+
+        for (old_block.instructions.items) |old_instruction_id| {
+            const new_instruction_id = mappedInstruction(instruction_map, old_instruction_id) orelse return Error.InvalidModule;
+            const new_instruction = module.instructions.getMut(new_instruction_id) orelse return Error.InvalidModule;
+            new_instruction.operation = try remapOperation(module, value_map, new_instruction.operation);
+        }
+
+        const new_block = module.blocks.getMut(new_block_id) orelse return Error.InvalidModule;
+        new_block.structured_control = try remapStructuredControl(block_map, old_block.structured_control);
+        new_block.terminator = try remapTerminator(
+            module,
+            value_map,
+            block_map,
+            old_block.terminator orelse return Error.InvalidModule,
+            continuation_id,
+            returns_value,
+        );
+    }
+
+    return mappedBlock(block_map, callee_entry) orelse return Error.InvalidModule;
+}
+
+fn remapOperation(
+    module: *module_ir.Module,
+    value_map: []const ?ids.ValueId,
+    operation: instruction_ir.Operation,
+) Error!instruction_ir.Operation {
+    return switch (operation) {
+        .unary => |op| .{ .unary = .{
+            .opcode = op.opcode,
+            .operand = try mappedValue(module, value_map, op.operand),
+        } },
+        .binary => |op| .{ .binary = .{
+            .opcode = op.opcode,
+            .lhs = try mappedValue(module, value_map, op.lhs),
+            .rhs = try mappedValue(module, value_map, op.rhs),
+        } },
+        .compare => |op| .{ .compare = .{
+            .opcode = op.opcode,
+            .lhs = try mappedValue(module, value_map, op.lhs),
+            .rhs = try mappedValue(module, value_map, op.rhs),
+        } },
+        .select => |op| .{ .select = .{
+            .condition = try mappedValue(module, value_map, op.condition),
+            .true_value = try mappedValue(module, value_map, op.true_value),
+            .false_value = try mappedValue(module, value_map, op.false_value),
+        } },
+        .bitcast => |operand| .{ .bitcast = try mappedValue(module, value_map, operand) },
+        .composite_construct => |op| .{ .composite_construct = .{
+            .elements = try remapValues(module, value_map, op.elements),
+        } },
+        .composite_extract => |op| .{ .composite_extract = .{
+            .composite = try mappedValue(module, value_map, op.composite),
+            .indices = op.indices,
+        } },
+        .load_interface => |op| .{ .load_interface = .{
+            .variable = op.variable,
+            .element_index = if (op.element_index) |index| try mappedValue(module, value_map, index) else null,
+        } },
+        .store_interface => |op| .{ .store_interface = .{
+            .variable = op.variable,
+            .value = try mappedValue(module, value_map, op.value),
+            .element_index = if (op.element_index) |index| try mappedValue(module, value_map, index) else null,
+        } },
+        .call => Error.InvalidModule,
+    };
+}
+
+fn remapTerminator(
+    module: *module_ir.Module,
+    value_map: []const ?ids.ValueId,
+    block_map: []const ?ids.BlockId,
+    terminator: module_ir.Terminator,
+    continuation_id: ids.BlockId,
+    returns_value: bool,
+) Error!module_ir.Terminator {
+    return switch (terminator) {
+        .branch => |edge| .{ .branch = try remapEdge(module, value_map, block_map, edge) },
+        .conditional_branch => |branch| .{ .conditional_branch = .{
+            .condition = try mappedValue(module, value_map, branch.condition),
+            .true_edge = try remapEdge(module, value_map, block_map, branch.true_edge),
+            .false_edge = try remapEdge(module, value_map, block_map, branch.false_edge),
+        } },
+        .return_void => if (returns_value)
+            Error.InvalidModule
+        else
+            .{ .branch = .{
+                .target = continuation_id,
+                .arguments = try module.allocator().alloc(ids.ValueId, 0),
+            } },
+        .return_value => |value| if (!returns_value)
+            Error.InvalidModule
+        else
+            .{ .branch = .{
+                .target = continuation_id,
+                .arguments = try remapValues(module, value_map, &.{value}),
+            } },
+        .discard => .discard,
+        .@"unreachable" => .@"unreachable",
+    };
+}
+
+fn remapEdge(module: *module_ir.Module, value_map: []const ?ids.ValueId, block_map: []const ?ids.BlockId, edge: module_ir.Edge) Error!module_ir.Edge {
+    return .{
+        .target = mappedBlock(block_map, edge.target) orelse return Error.InvalidModule,
+        .arguments = try remapValues(module, value_map, edge.arguments),
+    };
+}
+
+fn remapStructuredControl(block_map: []const ?ids.BlockId, control: module_ir.StructuredControl) Error!module_ir.StructuredControl {
+    return switch (control) {
+        .none => .none,
+        .selection => |selection| .{ .selection = .{
+            .merge_block = mappedBlock(block_map, selection.merge_block) orelse return Error.InvalidModule,
+        } },
+        .loop => |loop| .{ .loop = .{
+            .merge_block = mappedBlock(block_map, loop.merge_block) orelse return Error.InvalidModule,
+            .continue_block = mappedBlock(block_map, loop.continue_block) orelse return Error.InvalidModule,
+        } },
+    };
+}
+
+fn remapValues(module: *module_ir.Module, value_map: []const ?ids.ValueId, values: []const ids.ValueId) Error![]const ids.ValueId {
+    const result = try module.allocator().alloc(ids.ValueId, values.len);
+    for (values, result) |value, *mapped|
+        mapped.* = try mappedValue(module, value_map, value);
+    return result;
+}
+
+fn mappedValue(module: *const module_ir.Module, value_map: []const ?ids.ValueId, value_id: ids.ValueId) Error!ids.ValueId {
+    if (value_id.index() < value_map.len) {
+        if (value_map[value_id.index()]) |mapped|
+            return mapped;
+    }
+
+    const value = module.values.get(value_id) orelse return Error.InvalidModule;
+    return switch (value.definition) {
+        .constant, .undef => value_id,
+        else => Error.InvalidModule,
+    };
+}
+
+fn removeUnreachableBlocks(module: *module_ir.Module, scratch_allocator: std.mem.Allocator) Error!bool {
+    const reachable = try scratch_allocator.alloc(bool, module.blocks.entries.items.len);
+    defer scratch_allocator.free(reachable);
+
+    var queue: std.ArrayList(ids.BlockId) = .empty;
+    defer queue.deinit(scratch_allocator);
+
+    var changed = false;
+    for (module.functions.entries.items, 0..) |entry, function_index| {
+        const function = entry orelse continue;
+        const function_id = ids.FunctionId.fromIndex(function_index);
+        const entry_block = function.entry_block orelse return Error.InvalidModule;
+
+        @memset(reachable, false);
+        queue.clearRetainingCapacity();
+        try enqueueReachable(module, scratch_allocator, function_id, reachable, &queue, entry_block);
+
+        var cursor: usize = 0;
+        while (cursor < queue.items.len) : (cursor += 1) {
+            const block = module.blocks.get(queue.items[cursor]) orelse return Error.InvalidModule;
+            switch (block.terminator orelse return Error.InvalidModule) {
+                .branch => |edge| try enqueueReachable(module, scratch_allocator, function_id, reachable, &queue, edge.target),
+                .conditional_branch => |branch| {
+                    try enqueueReachable(module, scratch_allocator, function_id, reachable, &queue, branch.true_edge.target);
+                    try enqueueReachable(module, scratch_allocator, function_id, reachable, &queue, branch.false_edge.target);
+                },
+                else => {},
+            }
+        }
+
+        for (queue.items) |block_id| {
+            const block = module.blocks.get(block_id) orelse return Error.InvalidModule;
+            switch (block.structured_control) {
+                .none => {},
+                .selection => |selection| if (!isReachable(reachable, selection.merge_block))
+                    return Error.InvalidModule,
+                .loop => |loop| if (!isReachable(reachable, loop.merge_block) or !isReachable(reachable, loop.continue_block))
+                    return Error.InvalidModule,
+            }
+        }
+
+        const mutable_function = module.functions.getMut(function_id) orelse return Error.InvalidModule;
+        var block_index: usize = 0;
+        while (block_index < mutable_function.blocks.items.len) {
+            const block_id = mutable_function.blocks.items[block_index];
+            if (isReachable(reachable, block_id)) {
+                block_index += 1;
+                continue;
+            }
+
+            removeBlock(module, block_id);
+            _ = mutable_function.blocks.orderedRemove(block_index);
+            changed = true;
+        }
+    }
+
+    return changed;
+}
+
+fn enqueueReachable(
+    module: *const module_ir.Module,
+    scratch_allocator: std.mem.Allocator,
+    function_id: ids.FunctionId,
+    reachable: []bool,
+    queue: *std.ArrayList(ids.BlockId),
+    block_id: ids.BlockId,
+) Error!void {
+    if (block_id.index() >= reachable.len)
+        return Error.InvalidModule;
+    if (reachable[block_id.index()])
+        return;
+
+    const block = module.blocks.get(block_id) orelse return Error.InvalidModule;
+    if (block.parent_function != function_id)
+        return Error.InvalidModule;
+
+    reachable[block_id.index()] = true;
+    try queue.append(scratch_allocator, block_id);
+}
+
+fn isReachable(reachable: []const bool, block_id: ids.BlockId) bool {
+    return block_id.index() < reachable.len and reachable[block_id.index()];
+}
+
+fn removeBlock(module: *module_ir.Module, block_id: ids.BlockId) void {
+    const block = module.blocks.get(block_id) orelse return;
+    for (block.parameters.items) |parameter_id|
+        _ = module.values.remove(parameter_id);
+    for (block.instructions.items) |instruction_id| {
+        const inst = module.instructions.get(instruction_id) orelse continue;
+        if (inst.result) |result_id|
+            _ = module.values.remove(result_id);
+        _ = module.instructions.remove(instruction_id);
+    }
+    _ = module.blocks.remove(block_id);
+}
+
+fn mappedBlock(block_map: []const ?ids.BlockId, block_id: ids.BlockId) ?ids.BlockId {
+    if (block_id.index() >= block_map.len)
+        return null;
+    return block_map[block_id.index()];
+}
+
+fn mappedInstruction(instruction_map: []const ?ids.InstructionId, instruction_id: ids.InstructionId) ?ids.InstructionId {
+    if (instruction_id.index() >= instruction_map.len)
+        return null;
+    return instruction_map[instruction_id.index()];
+}
+
+fn removeNonEntryFunctions(module: *module_ir.Module, entry_point: ids.FunctionId) bool {
+    var changed = false;
+    for (module.functions.entries.items, 0..) |entry, function_index| {
+        const function = entry orelse continue;
+        const function_id = ids.FunctionId.fromIndex(function_index);
+        if (function_id == entry_point)
+            continue;
+
+        for (function.parameters.items) |parameter_id|
+            _ = module.values.remove(parameter_id);
+
+        for (function.blocks.items) |block_id| {
+            const block = module.blocks.get(block_id) orelse continue;
+            for (block.parameters.items) |parameter_id|
+                _ = module.values.remove(parameter_id);
+            for (block.instructions.items) |instruction_id| {
+                const inst = module.instructions.get(instruction_id) orelse continue;
+                if (inst.result) |result_id|
+                    _ = module.values.remove(result_id);
+                _ = module.instructions.remove(instruction_id);
+            }
+            _ = module.blocks.remove(block_id);
+        }
+
+        _ = module.functions.remove(function_id);
+        changed = true;
+    }
+    return changed;
+}
+
+fn runForTest(module: *module_ir.Module) !bool {
+    var manager = transformer_manager.Manager.init(std.testing.allocator);
+    defer manager.deinit();
+    try manager.add(transformer);
+
+    var context: transformer_manager.Context = .{ .allocator = std.testing.allocator };
+    return manager.run(module, &context);
+}
+
+fn liveFunctionCount(module: *const module_ir.Module) usize {
+    var count: usize = 0;
+    for (module.functions.entries.items) |entry| {
+        if (entry != null)
+            count += 1;
+    }
+    return count;
+}
+
+test "Inline All Functions: nested calls and multiple returns" {
+    const parser = @import("../parser/parser.zig");
+    const validator = @import("../validator/validator.zig");
+
+    var module = try parser.parseString(std.testing.allocator,
+        \\shader vertex @main
+        \\{
+        \\    %condition: constant bool = true
+        \\    %one: constant u32 = bits(0x1)
+        \\    %two: constant u32 = bits(0x2)
+        \\    fn @main() -> void
+        \\    {
+        \\        .entry():
+        \\            %result: u32 = call @outer(%one)
+        \\            %sum: u32 = integer_add %result, %two
+        \\            return
+        \\    }
+        \\    fn @outer(%outer_value: u32) -> u32
+        \\    {
+        \\        .entry():
+        \\            %nested: u32 = call @identity(%outer_value)
+        \\            return %nested
+        \\    }
+        \\    fn @identity(%identity_value: u32) -> u32
+        \\    {
+        \\        .entry():
+        \\            conditional_branch %condition, .left(), .right()
+        \\        .left():
+        \\            return %identity_value
+        \\        .right():
+        \\            return %two
+        \\    }
+        \\}
+    );
+    defer module.deinit();
+
+    try std.testing.expect(try runForTest(&module));
+    try validator.validate(&module);
+    try std.testing.expect(module.properties.no_function_calls);
+    try std.testing.expectEqual(@as(usize, 1), liveFunctionCount(&module));
+
+    for (module.instructions.entries.items) |entry| {
+        const inst = entry orelse continue;
+        try std.testing.expect(inst.operation != .call);
+    }
+}
+
+test "Inline All Functions: reject reachable recursion" {
+    const parser = @import("../parser/parser.zig");
+
+    var module = try parser.parseString(std.testing.allocator,
+        \\shader vertex @main
+        \\{
+        \\    %one: constant u32 = bits(0x1)
+        \\    fn @main() -> void
+        \\    {
+        \\        .entry():
+        \\            %result: u32 = call @recurse(%one)
+        \\            return
+        \\    }
+        \\    fn @recurse(%value: u32) -> u32
+        \\    {
+        \\        .entry():
+        \\            %nested: u32 = call @recurse(%value)
+        \\            return %nested
+        \\    }
+        \\}
+    );
+    defer module.deinit();
+
+    try std.testing.expectError(Error.RecursiveCall, runForTest(&module));
+    try std.testing.expectEqual(@as(usize, 2), liveFunctionCount(&module));
+}
+
+test "Inline All Functions: remove calls in unreachable blocks" {
+    const parser = @import("../parser/parser.zig");
+    const validator = @import("../validator/validator.zig");
+
+    var module = try parser.parseString(std.testing.allocator,
+        \\shader vertex @main
+        \\{
+        \\    %one: constant u32 = bits(0x1)
+        \\    fn @main() -> void
+        \\    {
+        \\        .entry():
+        \\            return
+        \\        .dead():
+        \\            %local: u32 = integer_add %one, %one
+        \\            %result: u32 = call @identity(%local)
+        \\            return
+        \\    }
+        \\    fn @identity(%identity_value: u32) -> u32
+        \\    {
+        \\        .entry():
+        \\            return %identity_value
+        \\    }
+        \\}
+    );
+    defer module.deinit();
+
+    try std.testing.expect(try runForTest(&module));
+    try validator.validate(&module);
+    try std.testing.expectEqual(@as(usize, 1), liveFunctionCount(&module));
+    const entry_function = module.functions.get(module.entry_point.?).?;
+    try std.testing.expectEqual(@as(usize, 1), entry_function.blocks.items.len);
+}
+
+test "Inline All Functions: remove unreachable recursive helpers" {
+    const parser = @import("../parser/parser.zig");
+    const validator = @import("../validator/validator.zig");
+
+    var module = try parser.parseString(std.testing.allocator,
+        \\shader vertex @main
+        \\{
+        \\    fn @main() -> void
+        \\    {
+        \\        .entry():
+        \\            return
+        \\    }
+        \\    fn @dead() -> void
+        \\    {
+        \\        .entry():
+        \\            call @dead()
+        \\            return
+        \\    }
+        \\}
+    );
+    defer module.deinit();
+
+    try std.testing.expect(try runForTest(&module));
+    try validator.validate(&module);
+    try std.testing.expectEqual(@as(usize, 1), liveFunctionCount(&module));
+}

src/compiler/ir/validator/dominance.zig

@@ -19,8 +19,8 @@ const DominanceUseContext = struct {
 
 pub fn validate(module: *const module_ir.Module, function_id: ids.FunctionId) Error!void {
     var analysis = cfg.init(module.backingAllocator(), module, function_id) catch |err| switch (err) {
-        error.OutOfMemory => return error.OutOfMemory,
-        else => return error.InvalidBlock,
+        std.mem.Allocator.Error.OutOfMemory => return Error.OutOfMemory,
+        else => return Error.InvalidBlock,
     };
     defer analysis.deinit();
 
@@ -41,7 +41,7 @@ pub fn validate(module: *const module_ir.Module, function_id: ids.FunctionId) Er
             instruction.operation.visitValueUses(&context, checkDominanceUse);
 
             if (!context.valid)
-                return error.DefinitionDoesNotDominateUse;
+                return Error.DefinitionDoesNotDominateUse;
         }
 
         var context: DominanceUseContext = .{
@@ -55,7 +55,7 @@ pub fn validate(module: *const module_ir.Module, function_id: ids.FunctionId) Er
         module_ir.visitTerminatorValueUses(block.terminator.?, &context, checkDominanceUse);
 
         if (!context.valid)
-            return error.DefinitionDoesNotDominateUse;
+            return Error.DefinitionDoesNotDominateUse;
     }
 }
 

src/compiler/ir/validator/validator.zig

@@ -3,44 +3,42 @@ const ids = @import("../id.zig");
 const type_ir = @import("../type.zig");
 const inst_ir = @import("../instruction.zig");
 const module_ir = @import("../module.zig");
+const Builder = @import("../Builder.zig");
 const dominance = @import("dominance.zig");
 
 pub const ValidationError = error{
-    MissingEntryPoint,
-    InvalidEntryPoint,
-    InvalidType,
-    InvalidConstant,
-    InvalidValue,
-    InvalidFunction,
+    CrossFunctionReference,
+    DefinitionDoesNotDominateUse,
+    EntryBlockHasPredecessor,
     InvalidBlock,
+    InvalidConstant,
+    InvalidEntryPoint,
+    InvalidFunction,
     InvalidInstruction,
+    InvalidStructuredControl,
+    InvalidType,
+    InvalidValue,
+    MissingEntryPoint,
     MissingFunctionEntryBlock,
     MissingTerminator,
-    EntryBlockHasPredecessor,
-    WrongParent,
-    WrongDefinition,
-    WrongParameterIndex,
-    WrongResultPresence,
-    WrongOperandType,
-    WrongResultType,
     WrongBranchArgumentCount,
     WrongBranchArgumentType,
-    CrossFunctionReference,
-    WrongReturnType,
+    WrongDefinition,
     WrongInterfaceDirection,
-    InvalidStructuredControl,
-    DefinitionDoesNotDominateUse,
+    WrongOperandType,
+    WrongParameterIndex,
+    WrongParent,
+    WrongResultPresence,
+    WrongResultType,
+    WrongReturnType,
 };
 
 pub const Error = ValidationError || std.mem.Allocator.Error;
 
-/// Early validator for the foundational IR. It covers object ownership, CFG
-/// edges, single definitions, function boundaries, and the currently modeled
-/// operation types, and SSA dominance.
 pub fn validate(module: *const module_ir.Module) Error!void {
-    const entry_point = module.entry_point orelse return error.MissingEntryPoint;
+    const entry_point = module.entry_point orelse return Error.MissingEntryPoint;
     if (!module.functions.isLive(entry_point))
-        return error.InvalidEntryPoint;
+        return Error.InvalidEntryPoint;
 
     for (module.types.entries.items) |entry| {
         const ty = entry orelse continue;
@@ -51,12 +49,12 @@ pub fn validate(module: *const module_ir.Module) Error!void {
         const constant = entry orelse continue;
 
         if (!module.types.isLive(constant.type))
-            return error.InvalidType;
+            return Error.InvalidType;
 
         if (constant.value == .composite) {
             for (constant.value.composite) |element| {
                 if (!module.constants.isLive(element))
-                    return error.InvalidConstant;
+                    return Error.InvalidConstant;
             }
         }
     }
@@ -65,30 +63,30 @@ pub fn validate(module: *const module_ir.Module) Error!void {
         const value = entry orelse continue;
 
         if (!module.types.isLive(value.type))
-            return error.InvalidType;
+            return Error.InvalidType;
 
         const value_id = ids.ValueId.fromIndex(value_index);
 
         switch (value.definition) {
             .constant => |id| {
-                const constant = module.constants.get(id) orelse return error.InvalidConstant;
+                const constant = module.constants.get(id) orelse return Error.InvalidConstant;
                 if (constant.type != value.type)
-                    return error.WrongResultType;
+                    return Error.WrongResultType;
             },
             .function_parameter => |definition| {
-                const function = module.functions.get(definition.function) orelse return error.InvalidFunction;
+                const function = module.functions.get(definition.function) orelse return Error.InvalidFunction;
                 if (definition.index >= function.parameters.items.len or function.parameters.items[definition.index] != value_id)
-                    return error.WrongParameterIndex;
+                    return Error.WrongParameterIndex;
             },
             .block_parameter => |definition| {
-                const block = module.blocks.get(definition.block) orelse return error.InvalidBlock;
+                const block = module.blocks.get(definition.block) orelse return Error.InvalidBlock;
                 if (definition.index >= block.parameters.items.len or block.parameters.items[definition.index] != value_id)
-                    return error.WrongParameterIndex;
+                    return Error.WrongParameterIndex;
             },
             .instruction => |instruction_id| {
-                const instruction = module.instructions.get(instruction_id) orelse return error.InvalidInstruction;
+                const instruction = module.instructions.get(instruction_id) orelse return Error.InvalidInstruction;
                 if (instruction.result != value_id)
-                    return error.WrongDefinition;
+                    return Error.WrongDefinition;
             },
             .undef => {},
         }
@@ -97,13 +95,13 @@ pub fn validate(module: *const module_ir.Module) Error!void {
     for (module.interface_variables.entries.items) |entry| {
         const variable = entry orelse continue;
         if (!module.types.isLive(variable.type))
-            return error.InvalidType;
+            return Error.InvalidType;
     }
 
     for (module.resources.entries.items) |entry| {
         const resource = entry orelse continue;
         if (!module.types.isLive(resource.type))
-            return error.InvalidType;
+            return Error.InvalidType;
     }
 
     for (module.functions.entries.items, 0..) |entry, function_index| {
@@ -111,32 +109,32 @@ pub fn validate(module: *const module_ir.Module) Error!void {
         const function_id = ids.FunctionId.fromIndex(function_index);
 
         if (!module.types.isLive(function.return_type))
-            return error.InvalidType;
+            return Error.InvalidType;
 
         if (function.parameter_types.items.len != function.parameters.items.len)
-            return error.WrongParameterIndex;
+            return Error.WrongParameterIndex;
 
         for (function.parameter_types.items, function.parameters.items, 0..) |parameter_type, parameter_id, index| {
-            const parameter = module.values.get(parameter_id) orelse return error.InvalidValue;
+            const parameter = module.values.get(parameter_id) orelse return Error.InvalidValue;
 
             if (parameter.type != parameter_type)
-                return error.WrongResultType;
+                return Error.WrongResultType;
 
             if (parameter.definition != .function_parameter or
                 parameter.definition.function_parameter.function != function_id or
                 parameter.definition.function_parameter.index != index)
-                return error.WrongDefinition;
+                return Error.WrongDefinition;
         }
 
-        const entry_block = function.entry_block orelse return error.MissingFunctionEntryBlock;
-        const entry_block_value = module.blocks.get(entry_block) orelse return error.InvalidBlock;
-        if (entry_block_value.parent_function != function_id) return error.WrongParent;
+        const entry_block = function.entry_block orelse return Error.MissingFunctionEntryBlock;
+        const entry_block_value = module.blocks.get(entry_block) orelse return Error.InvalidBlock;
+        if (entry_block_value.parent_function != function_id) return Error.WrongParent;
 
         for (function.blocks.items) |block_id| {
-            const block = module.blocks.get(block_id) orelse return error.InvalidBlock;
+            const block = module.blocks.get(block_id) orelse return Error.InvalidBlock;
 
             if (block.parent_function != function_id)
-                return error.WrongParent;
+                return Error.WrongParent;
 
             try validateBlock(module, function_id, block_id, block);
         }
@@ -146,7 +144,7 @@ pub fn validate(module: *const module_ir.Module) Error!void {
 
             if (block.terminator) |terminator| {
                 if (targetsBlock(terminator, entry_block))
-                    return error.EntryBlockHasPredecessor;
+                    return Error.EntryBlockHasPredecessor;
             }
         }
 
@@ -158,23 +156,23 @@ fn validateType(module: *const module_ir.Module, ty: type_ir.Type) ValidationErr
     switch (ty) {
         .vector => |vector| {
             if (!module.types.isLive(vector.element_type) or vector.length < 2)
-                return error.InvalidType;
+                return ValidationError.InvalidType;
         },
         .array => |array| {
             if (!module.types.isLive(array.element_type) or array.length == 0)
-                return error.InvalidType;
+                return ValidationError.InvalidType;
         },
         .structure => |structure| for (structure.members) |member| {
             if (!module.types.isLive(member))
-                return error.InvalidType;
+                return ValidationError.InvalidType;
         },
         .pointer => |pointer| {
             if (!module.types.isLive(pointer.pointee_type))
-                return error.InvalidType;
+                return ValidationError.InvalidType;
         },
         .resource_handle => |handle| if (handle.data_type) |data_type| {
             if (!module.types.isLive(data_type))
-                return error.InvalidType;
+                return ValidationError.InvalidType;
         },
         else => {},
     }
@@ -187,11 +185,11 @@ fn validateBlock(
     block: *const module_ir.Block,
 ) ValidationError!void {
     for (block.parameters.items, 0..) |parameter_id, index| {
-        const parameter = module.values.get(parameter_id) orelse return error.InvalidValue;
+        const parameter = module.values.get(parameter_id) orelse return ValidationError.InvalidValue;
         if (parameter.definition != .block_parameter or
             parameter.definition.block_parameter.block != block_id or
             parameter.definition.block_parameter.index != index)
-            return error.WrongDefinition;
+            return ValidationError.WrongDefinition;
     }
 
     switch (block.structured_control) {
@@ -204,97 +202,97 @@ fn validateBlock(
     }
 
     for (block.instructions.items) |instruction_id| {
-        const instruction = module.instructions.get(instruction_id) orelse return error.InvalidInstruction;
+        const instruction = module.instructions.get(instruction_id) orelse return ValidationError.InvalidInstruction;
 
         if (instruction.parent_block != block_id)
-            return error.WrongParent;
+            return ValidationError.WrongParent;
 
         if (instruction.result) |result_id| {
-            const result = module.values.get(result_id) orelse return error.InvalidValue;
+            const result = module.values.get(result_id) orelse return ValidationError.InvalidValue;
 
             if (result.definition != .instruction or result.definition.instruction != instruction_id)
-                return error.WrongDefinition;
+                return ValidationError.WrongDefinition;
         }
 
         try validateOperation(module, function_id, instruction);
     }
 
-    const terminator = block.terminator orelse return error.MissingTerminator;
+    const terminator = block.terminator orelse return ValidationError.MissingTerminator;
     try validateTerminator(module, function_id, terminator);
 }
 
 fn validateOperation(module: *const module_ir.Module, function_id: ids.FunctionId, instruction: *const inst_ir.Instruction) ValidationError!void {
-    const result_type = if (instruction.result) |result| module.typeOf(result) orelse return error.InvalidValue else null;
+    const result_type = if (instruction.result) |result| module.typeOf(result) orelse return ValidationError.InvalidValue else null;
 
     switch (instruction.operation) {
         .unary => |op| {
             const operand_type = try operandType(module, function_id, op.operand);
 
             if (result_type == null)
-                return error.WrongResultPresence;
+                return ValidationError.WrongResultPresence;
 
             if (result_type.? != operand_type)
-                return error.WrongResultType;
+                return ValidationError.WrongResultType;
         },
         .binary => |op| {
             const lhs_type = try operandType(module, function_id, op.lhs);
             const rhs_type = try operandType(module, function_id, op.rhs);
 
             if (lhs_type != rhs_type)
-                return error.WrongOperandType;
+                return ValidationError.WrongOperandType;
 
             if (result_type == null or result_type.? != lhs_type)
-                return error.WrongResultType;
+                return ValidationError.WrongResultType;
         },
         .compare => |op| {
             const lhs_type = try operandType(module, function_id, op.lhs);
             if (try operandType(module, function_id, op.rhs) != lhs_type)
-                return error.WrongOperandType;
+                return ValidationError.WrongOperandType;
 
-            const result = result_type orelse return error.WrongResultPresence;
+            const result = result_type orelse return ValidationError.WrongResultPresence;
             if (!isBoolean(module, result))
-                return error.WrongResultType;
+                return ValidationError.WrongResultType;
         },
         .select => |op| {
             if (!isBoolean(module, try operandType(module, function_id, op.condition)))
-                return error.WrongOperandType;
+                return ValidationError.WrongOperandType;
 
             const true_type = try operandType(module, function_id, op.true_value);
             if (try operandType(module, function_id, op.false_value) != true_type)
-                return error.WrongOperandType;
+                return ValidationError.WrongOperandType;
 
             if (result_type == null or result_type.? != true_type)
-                return error.WrongResultType;
+                return ValidationError.WrongResultType;
         },
         .bitcast => |operand| {
             _ = try operandType(module, function_id, operand);
             if (result_type == null)
-                return error.WrongResultPresence;
+                return ValidationError.WrongResultPresence;
         },
         .composite_construct => |op| {
-            const result = result_type orelse return error.WrongResultPresence;
-            const ty = module.types.get(result) orelse return error.InvalidType;
+            const result = result_type orelse return ValidationError.WrongResultPresence;
+            const ty = module.types.get(result) orelse return ValidationError.InvalidType;
 
             switch (ty.*) {
                 .vector => |vector| {
                     if (op.elements.len != vector.length)
-                        return error.WrongOperandType;
+                        return ValidationError.WrongOperandType;
 
                     for (op.elements) |element| {
                         if (try operandType(module, function_id, element) != vector.element_type)
-                            return error.WrongOperandType;
+                            return ValidationError.WrongOperandType;
                     }
                 },
                 .structure => |structure| {
                     if (op.elements.len != structure.members.len)
-                        return error.WrongOperandType;
+                        return ValidationError.WrongOperandType;
 
                     for (op.elements, structure.members) |element, member_type| {
                         if (try operandType(module, function_id, element) != member_type)
-                            return error.WrongOperandType;
+                            return ValidationError.WrongOperandType;
                     }
                 },
-                else => return error.WrongResultType,
+                else => return ValidationError.WrongResultType,
             }
         },
         .composite_extract => |op| {
@@ -302,127 +300,127 @@ fn validateOperation(module: *const module_ir.Module, function_id: ids.FunctionI
             const extracted_type = try indexedType(module, composite_type, op.indices);
 
             if (result_type == null or result_type.? != extracted_type)
-                return error.WrongResultType;
+                return ValidationError.WrongResultType;
         },
         .load_interface => |op| {
-            const variable = module.interface_variables.get(op.variable) orelse return error.InvalidValue;
+            const variable = module.interface_variables.get(op.variable) orelse return ValidationError.InvalidValue;
 
             if (variable.direction != .input)
-                return error.WrongInterfaceDirection;
+                return ValidationError.WrongInterfaceDirection;
 
             if (op.element_index) |index|
                 _ = try operandType(module, function_id, index);
 
             if (result_type == null or result_type.? != variable.type)
-                return error.WrongResultType;
+                return ValidationError.WrongResultType;
         },
         .store_interface => |op| {
             if (result_type != null)
-                return error.WrongResultPresence;
+                return ValidationError.WrongResultPresence;
 
-            const variable = module.interface_variables.get(op.variable) orelse return error.InvalidValue;
+            const variable = module.interface_variables.get(op.variable) orelse return ValidationError.InvalidValue;
 
             if (variable.direction != .output)
-                return error.WrongInterfaceDirection;
+                return ValidationError.WrongInterfaceDirection;
 
             if (try operandType(module, function_id, op.value) != variable.type)
-                return error.WrongOperandType;
+                return ValidationError.WrongOperandType;
 
             if (op.element_index) |index|
                 _ = try operandType(module, function_id, index);
         },
         .call => |op| {
-            const callee = module.functions.get(op.function) orelse return error.InvalidFunction;
+            const callee = module.functions.get(op.function) orelse return ValidationError.InvalidFunction;
 
             if (op.arguments.len != callee.parameter_types.items.len)
-                return error.WrongOperandType;
+                return ValidationError.WrongOperandType;
 
             for (op.arguments, callee.parameter_types.items) |argument, parameter_type| {
                 if (try operandType(module, function_id, argument) != parameter_type)
-                    return error.WrongOperandType;
+                    return ValidationError.WrongOperandType;
             }
 
-            const return_type = module.types.get(callee.return_type) orelse return error.InvalidType;
+            const return_type = module.types.get(callee.return_type) orelse return ValidationError.InvalidType;
 
             if (return_type.* == .void) {
                 if (result_type != null)
-                    return error.WrongResultPresence;
+                    return ValidationError.WrongResultPresence;
             } else if (result_type == null or result_type.? != callee.return_type)
-                return error.WrongResultType;
+                return ValidationError.WrongResultType;
         },
     }
 }
 
 fn validateTerminator(module: *const module_ir.Module, function_id: ids.FunctionId, terminator: module_ir.Terminator) ValidationError!void {
-    const function = module.functions.get(function_id) orelse return error.InvalidFunction;
+    const function = module.functions.get(function_id) orelse return ValidationError.InvalidFunction;
 
     switch (terminator) {
         .branch => |edge| try validateEdge(module, function_id, edge),
         .conditional_branch => |branch| {
             if (!isBoolean(module, try operandType(module, function_id, branch.condition)))
-                return error.WrongOperandType;
+                return ValidationError.WrongOperandType;
 
             try validateEdge(module, function_id, branch.true_edge);
             try validateEdge(module, function_id, branch.false_edge);
         },
         .return_void => {
             if (module.types.get(function.return_type).?.* != .void)
-                return error.WrongReturnType;
+                return ValidationError.WrongReturnType;
         },
         .return_value => |value| {
             if (try operandType(module, function_id, value) != function.return_type)
-                return error.WrongReturnType;
+                return ValidationError.WrongReturnType;
         },
         .discard => {
             if (module.stage != .fragment)
-                return error.WrongReturnType;
+                return ValidationError.WrongReturnType;
         },
         .@"unreachable" => {},
     }
 }
 
 fn validateEdge(module: *const module_ir.Module, function_id: ids.FunctionId, edge: module_ir.Edge) ValidationError!void {
-    const target = module.blocks.get(edge.target) orelse return error.InvalidBlock;
+    const target = module.blocks.get(edge.target) orelse return ValidationError.InvalidBlock;
 
     if (target.parent_function != function_id)
-        return error.CrossFunctionReference;
+        return ValidationError.CrossFunctionReference;
 
     if (edge.arguments.len != target.parameters.items.len)
-        return error.WrongBranchArgumentCount;
+        return ValidationError.WrongBranchArgumentCount;
 
     for (edge.arguments, target.parameters.items) |argument, parameter| {
         if (try operandType(module, function_id, argument) != module.typeOf(parameter).?)
-            return error.WrongBranchArgumentType;
+            return ValidationError.WrongBranchArgumentType;
     }
 }
 
 fn validateTarget(module: *const module_ir.Module, function_id: ids.FunctionId, target_id: ids.BlockId) ValidationError!void {
-    const target = module.blocks.get(target_id) orelse return error.InvalidStructuredControl;
+    const target = module.blocks.get(target_id) orelse return ValidationError.InvalidStructuredControl;
     if (target.parent_function != function_id)
-        return error.InvalidStructuredControl;
+        return ValidationError.InvalidStructuredControl;
 }
 
 fn operandType(module: *const module_ir.Module, function_id: ids.FunctionId, value_id: ids.ValueId) ValidationError!ids.TypeId {
-    const value = module.values.get(value_id) orelse return error.InvalidValue;
-    const owner = valueFunction(module, value_id) catch return error.InvalidValue;
+    const value = module.values.get(value_id) orelse return ValidationError.InvalidValue;
+    const owner = valueFunction(module, value_id) catch return ValidationError.InvalidValue;
 
     if (owner) |actual| {
         if (actual != function_id)
-            return error.CrossFunctionReference;
+            return ValidationError.CrossFunctionReference;
     }
     return value.type;
 }
 
 fn valueFunction(module: *const module_ir.Module, value_id: ids.ValueId) ValidationError!?ids.FunctionId {
-    const value = module.values.get(value_id) orelse return error.InvalidValue;
+    const value = module.values.get(value_id) orelse return ValidationError.InvalidValue;
 
     return switch (value.definition) {
         .constant, .undef => null,
         .function_parameter => |definition| definition.function,
-        .block_parameter => |definition| (module.blocks.get(definition.block) orelse return error.InvalidBlock).parent_function,
+        .block_parameter => |definition| (module.blocks.get(definition.block) orelse return ValidationError.InvalidBlock).parent_function,
         .instruction => |instruction_id| blk: {
-            const instruction = module.instructions.get(instruction_id) orelse return error.InvalidInstruction;
-            const block = module.blocks.get(instruction.parent_block) orelse return error.InvalidBlock;
+            const instruction = module.instructions.get(instruction_id) orelse return ValidationError.InvalidInstruction;
+            const block = module.blocks.get(instruction.parent_block) orelse return ValidationError.InvalidBlock;
             break :blk block.parent_function;
         },
     };
@@ -430,27 +428,27 @@ fn valueFunction(module: *const module_ir.Module, value_id: ids.ValueId) Validat
 
 fn indexedType(module: *const module_ir.Module, root: ids.TypeId, indices: []const u32) ValidationError!ids.TypeId {
     if (indices.len == 0)
-        return error.WrongOperandType;
+        return ValidationError.WrongOperandType;
 
     var current = root;
 
     for (indices) |index| {
-        const ty = module.types.get(current) orelse return error.InvalidType;
+        const ty = module.types.get(current) orelse return ValidationError.InvalidType;
         current = switch (ty.*) {
             .vector => |vector| if (index < vector.length)
                 vector.element_type
             else
-                return error.WrongOperandType,
+                return ValidationError.WrongOperandType,
             .array => |array| if (index < array.length)
                 array.element_type
             else
-                return error.WrongOperandType,
+                return ValidationError.WrongOperandType,
             .structure => |structure| if (index < structure.members.len)
                 structure.members[index]
             else
-                return error.WrongOperandType,
+                return ValidationError.WrongOperandType,
 
-            else => return error.WrongOperandType,
+            else => return ValidationError.WrongOperandType,
         };
     }
     return current;
@@ -468,3 +466,108 @@ fn targetsBlock(terminator: module_ir.Terminator, target: ids.BlockId) bool {
         else => false,
     };
 }
+
+test "Validator: Error wrong block argument count" {
+    // shader compute @main
+    // {
+    //     fn @main() -> void
+    //     {
+    //         .entry():
+    //             branch .merge()
+    //
+    //         .merge(%0: u32):
+    //             return
+    //     }
+    // }
+
+    var module = module_ir.Module.init(std.testing.allocator, .compute);
+    defer module.deinit();
+    var builder = Builder.init(&module);
+
+    const void_type = try builder.internType(.void);
+    const u32_type = try builder.internType(.{ .integer = .{ .bits = 32, .signedness = .unsigned } });
+    const main = try builder.addFunction(void_type, "main");
+    builder.setEntryPoint(main);
+    const entry = try builder.addBlock(main, "entry");
+    const merge = try builder.addBlock(main, "merge");
+    _ = try builder.addBlockParameter(merge, u32_type, null);
+    try builder.setTerminator(entry, .{ .branch = try builder.edge(merge, &.{}) });
+    try builder.setTerminator(merge, .return_void);
+
+    try std.testing.expectError(Error.WrongBranchArgumentCount, validate(&module));
+}
+
+test "Validator: Error SSA definition does not dominate its use" {
+    // shader compute @main
+    // {
+    //     %0: constant bool = true
+    //     %1: constant u32 = bits(0x1)
+    //
+    //     fn @main() -> void
+    //     {
+    //         .entry():
+    //             conditional_branch %0, .left(), .right()
+    //
+    //         .left():
+    //             %2: u32 = integer_add %1, %1
+    //             branch .merge()
+    //
+    //         .right():
+    //             branch .merge()
+    //
+    //         .merge():
+    //             %3: u32 = integer_multiply %2, %1
+    //             return
+    //     }
+    // }
+
+    var module = module_ir.Module.init(std.testing.allocator, .compute);
+    defer module.deinit();
+    var builder = Builder.init(&module);
+
+    const void_type = try builder.internType(.void);
+    const bool_type = try builder.internType(.boolean);
+    const u32_type = try builder.internType(.{ .integer = .{ .bits = 32, .signedness = .unsigned } });
+    const condition = try builder.internConstant(bool_type, .{ .boolean = true });
+    const one = try builder.internConstant(u32_type, .{ .integer_bits = 1 });
+    const main = try builder.addFunction(void_type, "main");
+    builder.setEntryPoint(main);
+    const entry = try builder.addBlock(main, "entry");
+    const left = try builder.addBlock(main, "left");
+    const right = try builder.addBlock(main, "right");
+    const merge = try builder.addBlock(main, "merge");
+
+    try builder.setTerminator(
+        entry,
+        .{
+            .conditional_branch = .{
+                .condition = condition,
+                .true_edge = try builder.edge(left, &.{}),
+                .false_edge = try builder.edge(right, &.{}),
+            },
+        },
+    );
+
+    const left_value = (try builder.appendInstruction(left, u32_type, .{
+        .binary = .{
+            .opcode = .integer_add,
+            .lhs = one,
+            .rhs = one,
+        },
+    }, null)).?;
+
+    try builder.setTerminator(left, .{ .branch = try builder.edge(merge, &.{}) });
+    try builder.setTerminator(right, .{ .branch = try builder.edge(merge, &.{}) });
+
+    _ = try builder.appendInstruction(merge, u32_type, .{
+        .binary = .{
+            .opcode = .integer_multiply,
+            .lhs = left_value,
+            .rhs = one,
+        },
+    }, null);
+
+    try builder.setTerminator(merge, .return_void);
+
+    try std.testing.expectError(Error.DefinitionDoesNotDominateUse, validate(&module));
+}

src/compiler/root.zig

@@ -3,824 +3,15 @@
 //! This module exposes the project-specific intermediate representation in
 //! `ir` and the SPIR-V frontend in `spirv`.
 //!
-//! Together they form
-//! the first stage of the compiler pipeline: SPIR-V binary modules are decoded,
-//! translated into a smaller and easier-to-transform IR, validated, and then made
-//! available to later optimization or code-generation passes.
-
-const std = @import("std");
+//! Together they form the first stage of the compiler pipeline: SPIR-V binary
+//! modules are decoded, translated into a smaller and easier-to-transform IR,
+//! validated, and then made available to later optimization or code-generation
+//! transformers.
 
 pub const ir = @import("ir/ir.zig");
 pub const spirv = @import("spirv/root.zig");
 
-const VisitorStatistics = struct {
-    functions: usize = 0,
-    blocks: usize = 0,
-};
-
-test "IR builder generation" {
-    // shader vertex @main
-    // {
-    //     @color: vec4[f32] = input[location(0), component(0), index(0)]
-    //     @out_color: vec4[f32] = output[location(0), component(0), index(0)]
-    //     %0: constant bool = true
-    //     %1: constant f32 = bits(0x3f800000)
-    //
-    //     fn @main() -> void
-    //     {
-    //         .entry():
-    //             %3: vec4[f32] = load_interface @color
-    //             conditional_branch %0, .pass(), .merge(%3)
-    //
-    //         .pass():
-    //             %4: vec4[f32] = composite_construct %1, %1, %1, %1
-    //             branch .merge(%4)
-    //
-    //         .merge(%2: vec4[f32]):
-    //             store_interface @out_color, %2
-    //             return
-    //     }
-    // }
-
-    var module = ir.module.Module.init(std.testing.allocator, .vertex);
-    defer module.deinit();
-    var builder = ir.Builder.init(&module);
-
-    const void_type = try builder.internType(.void);
-    const bool_type = try builder.internType(.boolean);
-    const f32_type = try builder.internType(.{ .floating = .{ .bits = 32 } });
-    const duplicate_f32 = try builder.internType(.{ .floating = .{ .bits = 32 } });
-    try std.testing.expectEqual(f32_type, duplicate_f32);
-    const vec4_type = try builder.internType(.{ .vector = .{ .element_type = f32_type, .length = 4 } });
-
-    const true_value = try builder.internConstant(bool_type, .{ .boolean = true });
-    const one = try builder.internConstant(f32_type, .{ .float_bits = @as(u32, @bitCast(@as(f32, 1.0))) });
-
-    const input = try builder.addInterfaceVariable(vec4_type, .input, .{ .location = .{ .location = 0 } }, "color");
-    const output = try builder.addInterfaceVariable(vec4_type, .output, .{ .location = .{ .location = 0 } }, "out_color");
-    const main = try builder.addFunction(void_type, "main");
-    builder.setEntryPoint(main);
-    const entry = try builder.addBlock(main, "entry");
-    const pass = try builder.addBlock(main, "pass");
-    const merge = try builder.addBlock(main, "merge");
-    const merged = try builder.addBlockParameter(merge, vec4_type, "merged");
-
-    const loaded = (try builder.appendInstruction(entry, vec4_type, .{
-        .load_interface = .{ .variable = input },
-    }, "loaded")).?;
-    try builder.setTerminator(entry, .{ .conditional_branch = .{
-        .condition = true_value,
-        .true_edge = try builder.edge(pass, &.{}),
-        .false_edge = try builder.edge(merge, &.{loaded}),
-    } });
-
-    const splat = (try builder.appendInstruction(pass, vec4_type, .{
-        .composite_construct = .{ .elements = &.{ one, one, one, one } },
-    }, "white")).?;
-    try builder.setTerminator(pass, .{ .branch = try builder.edge(merge, &.{splat}) });
-    _ = try builder.appendInstruction(merge, null, .{
-        .store_interface = .{ .variable = output, .value = merged },
-    }, null);
-    try builder.setTerminator(merge, .return_void);
-
-    try ir.validator.validate(&module);
-
-    var control_flow = try ir.cfg.init(std.testing.allocator, &module, main);
-    defer control_flow.deinit();
-    try std.testing.expectEqual(@as(usize, 2), control_flow.predecessors(merge).?.len);
-    try std.testing.expect(control_flow.dominates(entry, merge));
-    try std.testing.expect(!control_flow.dominates(pass, merge));
-
-    const text = try ir.printer.allocPrint(std.testing.allocator, &module);
-
-    defer std.testing.allocator.free(text);
-    try std.testing.expect(std.mem.indexOf(u8, text, "shader vertex @main") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "@color: vec4[f32] = input[location(0), component(0), index(0)]") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "@out_color: vec4[f32] = output[location(0), component(0), index(0)]") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "conditional_branch %0, .pass(), .merge(%loaded)") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, ".merge(%merged: vec4[f32])") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "store_interface @out_color, %merged") != null);
-
-    var parsed = try ir.parser.parseString(std.testing.allocator, text);
-    defer parsed.deinit();
-    const round_trip = try ir.printer.allocPrint(std.testing.allocator, &parsed);
-    defer std.testing.allocator.free(round_trip);
-    try std.testing.expectEqualStrings(text, round_trip);
-
-    const io = std.Options.debug_io;
-    const path = ".zig-cache/ir-parser-round-trip.ir";
-    const file = try std.Io.Dir.cwd().createFile(io, path, .{ .truncate = true });
-    {
-        defer file.close(io);
-        var file_buffer: [4096]u8 = @splat(0);
-        var file_writer = file.writer(io, &file_buffer);
-        try file_writer.interface.writeAll(text);
-        try file_writer.interface.flush();
-    }
-    defer std.Io.Dir.cwd().deleteFile(io, path) catch @panic("Caught an error while handling an error");
-
-    var parsed_file = try ir.parser.parseFile(std.testing.allocator, io, path);
-    defer parsed_file.deinit();
-    const file_round_trip = try ir.printer.allocPrint(std.testing.allocator, &parsed_file);
-    defer std.testing.allocator.free(file_round_trip);
-    try std.testing.expectEqualStrings(text, file_round_trip);
-}
-
-test "IR parse interface" {
-    const source =
-        \\ shader vertex @main
-        \\ {
-        \\     @in_color: vec4[f32] = input[location(0), component(0), index(0)]
-        \\     @out_color: vec4[f32] = output[location(0), component(0), index(0)]
-        \\     @position: vec4[f32] = output[builtin(position)]
-        \\
-        \\     fn @main() -> void
-        \\     {
-        \\         .entry():
-        \\             return
-        \\     }
-        \\ }
-    ;
-
-    var module = try ir.parser.parseString(std.testing.allocator, source);
-    defer module.deinit();
-    const printed = try ir.printer.allocPrint(std.testing.allocator, &module);
-    defer std.testing.allocator.free(printed);
-
-    try std.testing.expect(std.mem.indexOf(u8, printed, "@in_color: vec4[f32] = input[location(0), component(0), index(0)]") != null);
-    try std.testing.expect(std.mem.indexOf(u8, printed, "@out_color: vec4[f32] = output[location(0), component(0), index(0)]") != null);
-    try std.testing.expect(std.mem.indexOf(u8, printed, "@position: vec4[f32] = output[builtin(position)]") != null);
-}
-
-test "IR parse types, operations, calls, terminators" {
-    const source =
-        \\ shader fragment @main
-        \\ {
-        \\     %0: constant bool = true
-        \\     %1: constant u32 = bits(0x1)
-        \\     %2: constant u32 = bits(0x2)
-        \\     %3: constant f32 = bits(0x3f800000)
-        \\     %4: constant array[u32, 2] = [#1, #2]
-        \\     %5: constant struct[u32, u32] = [#1, #2]
-        \\     %6: constant ptr[private, u32] = null
-        \\     %7: constant resourceHandle[sampler] = null
-        \\
-        \\     fn @main() -> void
-        \\     {
-        \\         .entry():
-        \\             %9: u32 = bitwise_not %1
-        \\             %10: u32 = integer_add %9, %2
-        \\             %11: bool = cmp_equal %1, %2
-        \\             %12: u32 = select %11, %1, %2
-        \\             %13: u32 = bitcast %12
-        \\             %14: vec2[u32] = composite_construct %1, %2
-        \\             %15: u32 = composite_extract %14[0]
-        \\             %16: f32 = negate %3
-        \\             %17: f32 = float_add %3, %16
-        \\             %18: u32 = call @helper(%15)
-        \\             return
-        \\     }
-        \\
-        \\     fn @helper(%8: u32) -> u32
-        \\     {
-        \\         .entry():
-        \\             return %8
-        \\     }
-        \\
-        \\     fn @discarder() -> void
-        \\     {
-        \\         .entry():
-        \\             discard
-        \\     }
-        \\
-        \\     fn @dead() -> void
-        \\     {
-        \\         .entry():
-        \\             unreachable
-        \\     }
-        \\ }
-    ;
-
-    var module = try ir.parser.parseString(std.testing.allocator, source);
-    defer module.deinit();
-    const printed = try ir.printer.allocPrint(std.testing.allocator, &module);
-    defer std.testing.allocator.free(printed);
-
-    var reparsed = try ir.parser.parseString(std.testing.allocator, printed);
-    defer reparsed.deinit();
-    const printed_again = try ir.printer.allocPrint(std.testing.allocator, &reparsed);
-    defer std.testing.allocator.free(printed_again);
-    try std.testing.expectEqualStrings(printed, printed_again);
-    try std.testing.expect(std.mem.indexOf(u8, printed, "cmp_equal %1, %2") != null);
-    try std.testing.expect(std.mem.indexOf(u8, printed, "cmp.") == null);
-}
-
-test "IR parse named value IDs" {
-    const source =
-        \\ shader compute @main
-        \\ {
-        \\     %one_value: constant u32 = bits(0x1)
-        \\
-        \\     fn @main() -> void
-        \\     {
-        \\         .entry():
-        \\             %sum_value: u32 = integer_add %one_value, %one_value
-        \\             branch .merge(%sum_value)
-        \\
-        \\         .merge(%merged_value: u32):
-        \\             %product_value: u32 = integer_multiply %merged_value, %one_value
-        \\             return
-        \\     }
-        \\ }
-    ;
-
-    var module = try ir.parser.parseString(std.testing.allocator, source);
-    defer module.deinit();
-    const printed = try ir.printer.allocPrint(std.testing.allocator, &module);
-    defer std.testing.allocator.free(printed);
-
-    try std.testing.expect(std.mem.indexOf(u8, printed, "%one_value: constant u32 = bits(0x1)") != null);
-    try std.testing.expect(std.mem.indexOf(u8, printed, "%sum_value: u32 = integer_add %one_value, %one_value") != null);
-    try std.testing.expect(std.mem.indexOf(u8, printed, "branch .merge(%sum_value)") != null);
-    try std.testing.expect(std.mem.indexOf(u8, printed, ".merge(%merged_value: u32)") != null);
-    try std.testing.expect(std.mem.indexOf(u8, printed, "%product_value: u32 = integer_multiply %merged_value, %one_value") != null);
-
-    var reparsed = try ir.parser.parseString(std.testing.allocator, printed);
-    defer reparsed.deinit();
-    const printed_again = try ir.printer.allocPrint(std.testing.allocator, &reparsed);
-    defer std.testing.allocator.free(printed_again);
-    try std.testing.expectEqualStrings(printed, printed_again);
-}
-
-test "IR parse numeric constants" {
-    const source =
-        \\ shader compute @main
-        \\ {
-        \\     %0: constant u8 = 255
-        \\     %1: constant i8 = -1
-        \\     %2: constant f16 = 1.5
-        \\     %3: constant f32 = -0.0
-        \\     %4: constant f64 = 2.5e0
-        \\
-        \\     fn @main() -> void
-        \\     {
-        \\         .entry():
-        \\             return
-        \\     }
-        \\ }
-    ;
-
-    var module = try ir.parser.parseString(std.testing.allocator, source);
-    defer module.deinit();
-    const printed = try ir.printer.allocPrint(std.testing.allocator, &module);
-    defer std.testing.allocator.free(printed);
-
-    try std.testing.expect(std.mem.indexOf(u8, printed, "%0: constant u8 = bits(0xff)") != null);
-    try std.testing.expect(std.mem.indexOf(u8, printed, "%1: constant i8 = bits(0xff)") != null);
-    try std.testing.expect(std.mem.indexOf(u8, printed, "%2: constant f16 = bits(0x3e00)") != null);
-    try std.testing.expect(std.mem.indexOf(u8, printed, "%3: constant f32 = bits(0x80000000)") != null);
-    try std.testing.expect(std.mem.indexOf(u8, printed, "%4: constant f64 = bits(0x4004000000000000)") != null);
-
-    const out_of_range =
-        \\ shader compute @main
-        \\ {
-        \\     %0: constant u8 = 256
-        \\     fn @main() -> void
-        \\     {
-        \\         .entry():
-        \\             return
-        \\     }
-        \\ }
-    ;
-    try std.testing.expectError(error.InvalidNumber, ir.parser.parseString(std.testing.allocator, out_of_range));
-}
-
-test "IR parser error: unknown value" {
-    const source =
-        \\ shader compute @main
-        \\ {
-        \\     fn @main() -> void
-        \\     {
-        \\         .entry():
-        \\             return %99
-        \\     }
-        \\ }
-    ;
-    try std.testing.expectError(error.UnknownValue, ir.parser.parseString(std.testing.allocator, source));
-}
-
-test "Validator error: wrong block argument count" {
-    // shader compute @main
-    // {
-    //     fn @main() -> void
-    //     {
-    //         .entry():
-    //             branch .merge()
-    //
-    //         .merge(%0: u32):
-    //             return
-    //     }
-    // }
-
-    var module = ir.module.Module.init(std.testing.allocator, .compute);
-    defer module.deinit();
-    var builder = ir.Builder.init(&module);
-
-    const void_type = try builder.internType(.void);
-    const u32_type = try builder.internType(.{ .integer = .{ .bits = 32, .signedness = .unsigned } });
-    const main = try builder.addFunction(void_type, "main");
-    builder.setEntryPoint(main);
-    const entry = try builder.addBlock(main, "entry");
-    const merge = try builder.addBlock(main, "merge");
-    _ = try builder.addBlockParameter(merge, u32_type, null);
-    try builder.setTerminator(entry, .{ .branch = try builder.edge(merge, &.{}) });
-    try builder.setTerminator(merge, .return_void);
-
-    try std.testing.expectError(error.WrongBranchArgumentCount, ir.validator.validate(&module));
-}
-
-test "Central store IDs disposal" {
-    var module = ir.module.Module.init(std.testing.allocator, .fragment);
-    defer module.deinit();
-    const first = try module.internType(.boolean);
-    try std.testing.expect(module.types.remove(first));
-    const second = try module.internType(.boolean);
-    try std.testing.expect(first.index() != second.index());
-    try std.testing.expect(module.types.get(first) == null);
-}
-
-test "Validator error: SSA definition does not dominate its use" {
-    // shader compute @main
-    // {
-    //     %0: constant bool = true
-    //     %1: constant u32 = bits(0x1)
-    //
-    //     fn @main() -> void
-    //     {
-    //         .entry():
-    //             conditional_branch %0, .left(), .right()
-    //
-    //         .left():
-    //             %2: u32 = integer_add %1, %1
-    //             branch .merge()
-    //
-    //         .right():
-    //             branch .merge()
-    //
-    //         .merge():
-    //             %3: u32 = integer_multiply %2, %1
-    //             return
-    //     }
-    // }
-
-    var module = ir.module.Module.init(std.testing.allocator, .compute);
-    defer module.deinit();
-    var builder = ir.Builder.init(&module);
-
-    const void_type = try builder.internType(.void);
-    const bool_type = try builder.internType(.boolean);
-    const u32_type = try builder.internType(.{ .integer = .{ .bits = 32, .signedness = .unsigned } });
-    const condition = try builder.internConstant(bool_type, .{ .boolean = true });
-    const one = try builder.internConstant(u32_type, .{ .integer_bits = 1 });
-    const main = try builder.addFunction(void_type, "main");
-    builder.setEntryPoint(main);
-    const entry = try builder.addBlock(main, "entry");
-    const left = try builder.addBlock(main, "left");
-    const right = try builder.addBlock(main, "right");
-    const merge = try builder.addBlock(main, "merge");
-
-    try builder.setTerminator(
-        entry,
-        .{
-            .conditional_branch = .{
-                .condition = condition,
-                .true_edge = try builder.edge(left, &.{}),
-                .false_edge = try builder.edge(right, &.{}),
-            },
-        },
-    );
-
-    const left_value = (try builder.appendInstruction(left, u32_type, .{
-        .binary = .{
-            .opcode = .integer_add,
-            .lhs = one,
-            .rhs = one,
-        },
-    }, null)).?;
-
-    try builder.setTerminator(left, .{ .branch = try builder.edge(merge, &.{}) });
-    try builder.setTerminator(right, .{ .branch = try builder.edge(merge, &.{}) });
-
-    _ = try builder.appendInstruction(merge, u32_type, .{
-        .binary = .{
-            .opcode = .integer_multiply,
-            .lhs = left_value,
-            .rhs = one,
-        },
-    }, null);
-
-    try builder.setTerminator(merge, .return_void);
-
-    try std.testing.expectError(error.DefinitionDoesNotDominateUse, ir.validator.validate(&module));
-}
-
-test "Rewriter replace all ID uses, safely erase dead instruction" {
-    // shader compute @main
-    // {
-    //     %0: constant u32 = bits(0x1)
-    //     %1: constant u32 = bits(0x2)
-    //
-    //     fn @main() -> void
-    //     {
-    //         .entry():
-    //             %2: u32 = integer_add %0, %1
-    //             %3: u32 = integer_multiply %2, %1
-    //             return
-    //     }
-    // }
-
-    var module = ir.module.Module.init(std.testing.allocator, .compute);
-    defer module.deinit();
-
-    var builder = ir.Builder.init(&module);
-
-    const void_type = try builder.internType(.void);
-    const u32_type = try builder.internType(.{ .integer = .{ .bits = 32, .signedness = .unsigned } });
-
-    const one = try builder.internConstant(u32_type, .{ .integer_bits = 1 });
-    const two = try builder.internConstant(u32_type, .{ .integer_bits = 2 });
-
-    const main = try builder.addFunction(void_type, "main");
-    builder.setEntryPoint(main);
-
-    const entry = try builder.addBlock(main, "entry");
-    const sum = (try builder.appendInstruction(entry, u32_type, .{
-        .binary = .{
-            .opcode = .integer_add,
-            .lhs = one,
-            .rhs = two,
-        },
-    }, null)).?;
-    _ = try builder.appendInstruction(entry, u32_type, .{
-        .binary = .{
-            .opcode = .integer_multiply,
-            .lhs = sum,
-            .rhs = two,
-        },
-    }, null);
-    try builder.setTerminator(entry, .return_void);
-
-    try ir.validator.validate(&module);
-
-    const sum_instruction = module.values.get(sum).?.definition.instruction;
-    var rewriter = ir.Rewriter.init(&module);
-
-    try std.testing.expectEqual(@as(usize, 1), try rewriter.replaceAllUses(sum, one));
-    try rewriter.eraseInstruction(sum_instruction);
-
-    try std.testing.expect(module.values.get(sum) == null);
-    try std.testing.expect(module.instructions.get(sum_instruction) == null);
-
-    try ir.validator.validate(&module);
-}
-
-test "Rewriter add block parameter and sync branch calls" {
-    // shader compute @main
-    // {
-    //     %0: constant u32 = bits(0x1)
-    //
-    //     fn @main() -> void
-    //     {
-    //         .entry():
-    //             branch .merge()
-    //
-    //         .merge():
-    //             return
-    //
-    //         .alternate():
-    //             return
-    //     }
-    // }
-
-    var module = ir.module.Module.init(std.testing.allocator, .compute);
-    defer module.deinit();
-
-    var builder = ir.Builder.init(&module);
-
-    const void_type = try builder.internType(.void);
-    const u32_type = try builder.internType(.{ .integer = .{ .bits = 32, .signedness = .unsigned } });
-
-    const one = try builder.internConstant(u32_type, .{ .integer_bits = 1 });
-
-    const main = try builder.addFunction(void_type, "main");
-    builder.setEntryPoint(main);
-
-    const entry = try builder.addBlock(main, "entry");
-    const merge = try builder.addBlock(main, "merge");
-    const alternate = try builder.addBlock(main, "alternate");
-
-    try builder.setTerminator(entry, .{ .branch = try builder.edge(merge, &.{}) });
-    try builder.setTerminator(merge, .return_void);
-    try builder.setTerminator(alternate, .return_void);
-
-    var rewriter = ir.Rewriter.init(&module);
-
-    const parameter = try rewriter.addBlockParameter(merge, u32_type, "incoming", &.{
-        .{
-            .predecessor = entry,
-            .value = one,
-        },
-    });
-    const merge_edge = module.blocks.get(entry).?.terminator.?.branch;
-    try std.testing.expectEqualSlices(ir.id.ValueId, &.{one}, merge_edge.arguments);
-
-    _ = try builder.appendInstruction(merge, u32_type, .{
-        .binary = .{
-            .opcode = .integer_add,
-            .lhs = parameter,
-            .rhs = one,
-        },
-    }, null);
-    try ir.validator.validate(&module);
-
-    try rewriter.removeBlockParameter(merge, 0, one);
-
-    try std.testing.expectEqual(@as(usize, 0), module.blocks.get(merge).?.parameters.items.len);
-    try std.testing.expectEqual(@as(usize, 0), module.blocks.get(entry).?.terminator.?.branch.arguments.len);
-
-    try ir.validator.validate(&module);
-
-    try std.testing.expectEqual(@as(usize, 1), try rewriter.redirectEdges(entry, merge, alternate, &.{}));
-    try std.testing.expectEqual(alternate, module.blocks.get(entry).?.terminator.?.branch.target);
-
-    try ir.validator.validate(&module);
-}
-
-fn establishNoCalls(_: *ir.module.Module, _: *ir.pass_manager.Context) !bool {
-    return false;
-}
-
-fn countVisitedFunction(context: ?*anyopaque, _: ir.id.FunctionId, _: *const ir.module.Function) !void {
-    const statistics: *VisitorStatistics = @ptrCast(@alignCast(context.?));
-    statistics.functions += 1;
-}
-
-fn countVisitedBlock(context: ?*anyopaque, _: ir.id.BlockId, _: *const ir.module.Block) !void {
-    const statistics: *VisitorStatistics = @ptrCast(@alignCast(context.?));
-    statistics.blocks += 1;
-}
-
-test "Pass manager track independent IR properties" {
-    // shader compute @main
-    // {
-    //     fn @main() -> void
-    //     {
-    //         .entry():
-    //             return
-    //     }
-    // }
-
-    var module = ir.module.Module.init(std.testing.allocator, .compute);
-    defer module.deinit();
-
-    var builder = ir.Builder.init(&module);
-
-    const void_type = try builder.internType(.void);
-
-    const main = try builder.addFunction(void_type, "main");
-    builder.setEntryPoint(main);
-
-    const entry = try builder.addBlock(main, "entry");
-    try builder.setTerminator(entry, .return_void);
-
-    module.properties.valid_cfg = true;
-
-    var manager = ir.pass_manager.Manager.init(std.testing.allocator);
-    defer manager.deinit();
-
-    try manager.add(.{
-        .name = "establish-no-calls",
-        .required = .{ .valid_cfg = true },
-        .produced = .{ .no_function_calls = true },
-        .run = establishNoCalls,
-    });
-
-    var context: ir.pass_manager.Context = .{ .allocator = std.testing.allocator };
-
-    try std.testing.expect(!try manager.run(&module, &context));
-    try std.testing.expect(module.properties.no_function_calls);
-
-    var statistics: VisitorStatistics = .{};
-
-    try ir.visitor.walk(&module, .{
-        .context = &statistics,
-        .visitFunction = countVisitedFunction,
-        .visitBlock = countVisitedBlock,
-    });
-
-    try std.testing.expectEqual(@as(usize, 1), statistics.functions);
-    try std.testing.expectEqual(@as(usize, 1), statistics.blocks);
-}
-
-test "SPIR-V parser error: zero-word instruction" {
-    const words = [_]u32{
-        spirv.spec.magic_number,
-        0x0001_0000,
-        0,
-        2,
-        0,
-        instructionWord(.nop, 0),
-    };
-    try std.testing.expectError(error.ZeroWordInstruction, spirv.Parser.init(&words));
-
-    const truncated = [_]u32{
-        spirv.spec.magic_number,
-        0x0001_0000,
-        0,
-        2,
-        0,
-        instructionWord(.i_add, 5),
-        1,
-    };
-    try std.testing.expectError(error.TruncatedInstruction, spirv.Parser.init(&truncated));
-}
-
-test "SPIR-V structured branches and OpPhi to block parameters" {
-    const assembly =
-        \\ OpCapability Shader
-        \\ OpMemoryModel Logical GLSL450
-        \\ OpEntryPoint GLCompute %main "main"
-        \\ OpExecutionMode %main LocalSize 1 1 1
-        \\ OpName %main "main"
-        \\ OpName %entry "entry"
-        \\ OpName %true "true"
-        \\ OpName %one "one"
-        \\ OpName %then "then"
-        \\ OpName %then_value "then_value"
-        \\ OpName %else "else"
-        \\ OpName %else_value "else_value"
-        \\ OpName %merge "merge"
-        \\ OpName %merged "merged"
-        \\ OpName %product "product"
-        \\
-        \\ %void = OpTypeVoid
-        \\ %bool = OpTypeBool
-        \\ %uint = OpTypeInt 32 0
-        \\ %fn_void = OpTypeFunction %void
-        \\ %true = OpConstantTrue %bool
-        \\ %one = OpConstant %uint 1
-        \\
-        \\ %main = OpFunction %void None %fn_void
-        \\     %entry = OpLabel
-        \\     OpSelectionMerge %merge None
-        \\     OpBranchConditional %true %then %else
-        \\     %then = OpLabel
-        \\     %then_value = OpIAdd %uint %one %one
-        \\     OpBranch %merge
-        \\     %else = OpLabel
-        \\     %else_value = OpISub %uint %one %one
-        \\     OpBranch %merge
-        \\     %merge = OpLabel
-        \\     %merged = OpPhi %uint %then_value %then %else_value %else
-        \\     %product = OpIMul %uint %merged %one
-        \\     OpReturn
-        \\ OpFunctionEnd
-    ;
-    const words = try assembleSpirv(std.testing.allocator, assembly);
-    defer std.testing.allocator.free(words);
-
-    var module = try spirv.translator.translate(std.testing.allocator, words, .{ .entry_point = "main" });
-    defer module.deinit();
-
-    try std.testing.expectEqual(ir.module.Stage.compute, module.stage);
-    try std.testing.expectEqual([3]u32{ 1, 1, 1 }, module.execution_modes.workgroup_size.?);
-    try std.testing.expect(module.properties.valid_cfg);
-    try std.testing.expect(module.properties.valid_ssa);
-
-    const function = module.functions.get(module.entry_point.?).?;
-    try std.testing.expectEqual(@as(usize, 4), function.blocks.items.len);
-    const entry = module.blocks.get(function.blocks.items[0]).?;
-    try std.testing.expect(entry.structured_control == .selection);
-    const merge = module.blocks.get(function.blocks.items[3]).?;
-    try std.testing.expectEqual(@as(usize, 1), merge.parameters.items.len);
-    try std.testing.expectEqual(@as(usize, 1), merge.instructions.items.len);
-    const multiply = module.instructions.get(merge.instructions.items[0]).?;
-    try std.testing.expectEqual(ir.instruction.BinaryOpcode.integer_multiply, multiply.operation.binary.opcode);
-
-    const text = try ir.printer.allocPrint(std.testing.allocator, &module);
-    defer std.testing.allocator.free(text);
-
-    try std.testing.expect(std.mem.indexOf(u8, text, "%one: constant u32 = bits(0x1)") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "%true: constant bool = true") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "conditional_branch %true, .then(), .else()") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "%then_value: u32 = integer_add %one, %one") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "branch .merge(%then_value)") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "%else_value: u32 = integer_subtract %one, %one") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, ".merge(%merged: u32)") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "%product: u32 = integer_multiply %merged, %one") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "integerMultiply") == null);
-
-    var parsed = try ir.parser.parseString(std.testing.allocator, text);
-    defer parsed.deinit();
-    const round_trip = try ir.printer.allocPrint(std.testing.allocator, &parsed);
-    defer std.testing.allocator.free(round_trip);
-    try std.testing.expectEqualStrings(text, round_trip);
-}
-
-test "SPIR-V decorated vertex interfaces and load-store operations" {
-    const assembly =
-        \\ OpCapability Shader
-        \\ OpMemoryModel Logical GLSL450
-        \\ OpEntryPoint Vertex %main "main" %in_color %out_color
-        \\ OpName %in_color "in_color"
-        \\ OpName %out_color "out_color"
-        \\ OpDecorate %in_color Location 0
-        \\ OpDecorate %out_color Location 0
-        \\
-        \\ %void = OpTypeVoid
-        \\ %float = OpTypeFloat 32
-        \\ %vec4 = OpTypeVector %float 4
-        \\ %input_vec4 = OpTypePointer Input %vec4
-        \\ %output_vec4 = OpTypePointer Output %vec4
-        \\ %fn_void = OpTypeFunction %void
-        \\ %in_color = OpVariable %input_vec4 Input
-        \\ %out_color = OpVariable %output_vec4 Output
-        \\
-        \\ %main = OpFunction %void None %fn_void
-        \\     %entry = OpLabel
-        \\     %color = OpLoad %vec4 %in_color
-        \\     OpStore %out_color %color
-        \\     OpReturn
-        \\ OpFunctionEnd
-    ;
-    const words = try assembleSpirv(std.testing.allocator, assembly);
-    defer std.testing.allocator.free(words);
-
-    var module = try spirv.translator.translate(std.testing.allocator, words, .{ .entry_point = "main" });
-    defer module.deinit();
-    try std.testing.expectEqual(ir.module.Stage.vertex, module.stage);
-    try std.testing.expectEqual(@as(usize, 2), module.interface_variables.entries.items.len);
-
-    const text = try ir.printer.allocPrint(std.testing.allocator, &module);
-    defer std.testing.allocator.free(text);
-    try std.testing.expect(std.mem.indexOf(u8, text, "load_interface @in_color") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "store_interface @out_color") != null);
-}
-
-fn instructionWord(opcode: spirv.spec.Opcode, word_count: u16) u32 {
-    return (@as(u32, word_count) << 16) | @intFromEnum(opcode);
-}
-
-fn assembleSpirv(allocator: std.mem.Allocator, assembly: []const u8) ![]u32 {
-    var io_backend: std.Io.Threaded = .init(allocator, .{});
-    defer io_backend.deinit();
-    const io = io_backend.io();
-
-    var child = try std.process.spawn(io, .{
-        .argv = &.{ "spirv-as", "--target-env", "spv1.0", "-o", "-", "-" },
-        .stdin = .pipe,
-        .stdout = .pipe,
-        .stderr = .pipe,
-    });
-    defer child.kill(io);
-
-    {
-        const stdin = child.stdin.?;
-        var stdin_writer = stdin.writer(io, &.{});
-        try stdin_writer.interface.writeAll(assembly);
-        try stdin_writer.interface.flush();
-        stdin.close(io);
-        child.stdin = null;
-    }
-
-    var stdout_buffer: [4096]u8 = undefined;
-    var stdout_reader = child.stdout.?.reader(io, &stdout_buffer);
-    const binary = try stdout_reader.interface.allocRemaining(allocator, .limited(1024 * 1024));
-    defer allocator.free(binary);
-
-    var stderr_buffer: [4096]u8 = undefined;
-    var stderr_reader = child.stderr.?.reader(io, &stderr_buffer);
-    const stderr = try stderr_reader.interface.allocRemaining(allocator, .limited(64 * 1024));
-    defer allocator.free(stderr);
-
-    const term = try child.wait(io);
-    switch (term) {
-        .exited => |code| if (code != 0) {
-            std.log.err("spirv-as failed:\n{s}", .{stderr});
-            return error.SpirvAssemblyFailed;
-        },
-        else => {
-            std.log.err("spirv-as terminated unexpectedly:\n{s}", .{stderr});
-            return error.SpirvAssemblyFailed;
-        },
-    }
-
-    if (binary.len % @sizeOf(u32) != 0) return error.InvalidSpirvBinaryLength;
-    const words = try allocator.alloc(u32, binary.len / @sizeOf(u32));
-    errdefer allocator.free(words);
-    for (words, 0..) |*word, index| {
-        const offset = index * @sizeOf(u32);
-        word.* = std.mem.readInt(u32, binary[offset..][0..4], .little);
-    }
-    return words;
+test {
+    _ = ir;
+    _ = spirv;
 }

src/compiler/spirv/Parser.zig

@@ -1,3 +1,4 @@
+const std = @import("std");
 const spirv = @import("spirv.zig");
 
 const Self = @This();
@@ -153,3 +154,30 @@ pub fn copyLiteralString(allocator: anytype, words: []const u32) ![]u8 {
     }
     return result;
 }
+
+test "SPIR-V: parser error zero-word instruction" {
+    const words = [_]u32{
+        spirv.magic_number,
+        0x0001_0000,
+        0,
+        2,
+        0,
+        instructionWord(.nop, 0),
+    };
+    try std.testing.expectError(error.ZeroWordInstruction, Self.init(&words));
+
+    const truncated = [_]u32{
+        spirv.magic_number,
+        0x0001_0000,
+        0,
+        2,
+        0,
+        instructionWord(.i_add, 5),
+        1,
+    };
+    try std.testing.expectError(error.TruncatedInstruction, Self.init(&truncated));
+}
+
+fn instructionWord(opcode: spirv.Opcode, word_count: u16) u32 {
+    return (@as(u32, word_count) << 16) | @intFromEnum(opcode);
+}

src/compiler/spirv/root.zig

@@ -14,3 +14,8 @@
 pub const Parser = @import("Parser.zig");
 pub const translator = @import("translator.zig");
 pub const spec = @import("spirv.zig");
+
+test {
+    _ = Parser;
+    _ = translator;
+}

src/compiler/spirv/translator.zig

@@ -1072,3 +1072,179 @@ fn allocOptional(comptime T: type, allocator: std.mem.Allocator, count: usize) !
     @memset(values, null);
     return values;
 }
+
+test "SPIR-V: structured branches and OpPhi to block parameters" {
+    const assembly =
+        \\ OpCapability Shader
+        \\ OpMemoryModel Logical GLSL450
+        \\ OpEntryPoint GLCompute %main "main"
+        \\ OpExecutionMode %main LocalSize 1 1 1
+        \\ OpName %main "main"
+        \\ OpName %entry "entry"
+        \\ OpName %true "true"
+        \\ OpName %one "one"
+        \\ OpName %then "then"
+        \\ OpName %then_value "then_value"
+        \\ OpName %else "else"
+        \\ OpName %else_value "else_value"
+        \\ OpName %merge "merge"
+        \\ OpName %merged "merged"
+        \\ OpName %product "product"
+        \\
+        \\ %void = OpTypeVoid
+        \\ %bool = OpTypeBool
+        \\ %uint = OpTypeInt 32 0
+        \\ %fn_void = OpTypeFunction %void
+        \\ %true = OpConstantTrue %bool
+        \\ %one = OpConstant %uint 1
+        \\
+        \\ %main = OpFunction %void None %fn_void
+        \\     %entry = OpLabel
+        \\     OpSelectionMerge %merge None
+        \\     OpBranchConditional %true %then %else
+        \\     %then = OpLabel
+        \\     %then_value = OpIAdd %uint %one %one
+        \\     OpBranch %merge
+        \\     %else = OpLabel
+        \\     %else_value = OpISub %uint %one %one
+        \\     OpBranch %merge
+        \\     %merge = OpLabel
+        \\     %merged = OpPhi %uint %then_value %then %else_value %else
+        \\     %product = OpIMul %uint %merged %one
+        \\     OpReturn
+        \\ OpFunctionEnd
+    ;
+    const words = try assembleSpirv(std.testing.allocator, assembly);
+    defer std.testing.allocator.free(words);
+
+    var module = try translate(std.testing.allocator, words, .{ .entry_point = "main" });
+    defer module.deinit();
+
+    try std.testing.expectEqual(ir.module.Stage.compute, module.stage);
+    try std.testing.expectEqual([3]u32{ 1, 1, 1 }, module.execution_modes.workgroup_size.?);
+    try std.testing.expect(module.properties.valid_cfg);
+    try std.testing.expect(module.properties.valid_ssa);
+
+    const function = module.functions.get(module.entry_point.?).?;
+    try std.testing.expectEqual(@as(usize, 4), function.blocks.items.len);
+    const entry = module.blocks.get(function.blocks.items[0]).?;
+    try std.testing.expect(entry.structured_control == .selection);
+    const merge = module.blocks.get(function.blocks.items[3]).?;
+    try std.testing.expectEqual(@as(usize, 1), merge.parameters.items.len);
+    try std.testing.expectEqual(@as(usize, 1), merge.instructions.items.len);
+    const multiply = module.instructions.get(merge.instructions.items[0]).?;
+    try std.testing.expectEqual(ir.instruction.BinaryOpcode.integer_multiply, multiply.operation.binary.opcode);
+
+    const text = try ir.printer.allocPrint(std.testing.allocator, &module);
+    defer std.testing.allocator.free(text);
+
+    try std.testing.expect(std.mem.indexOf(u8, text, "%one: constant u32 = bits(0x1)") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "%true: constant bool = true") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "conditional_branch %true, .then(), .else()") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "%then_value: u32 = integer_add %one, %one") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "branch .merge(%then_value)") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "%else_value: u32 = integer_subtract %one, %one") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, ".merge(%merged: u32)") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "%product: u32 = integer_multiply %merged, %one") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "integerMultiply") == null);
+
+    var parsed = try ir.parser.parseString(std.testing.allocator, text);
+    defer parsed.deinit();
+    const round_trip = try ir.printer.allocPrint(std.testing.allocator, &parsed);
+    defer std.testing.allocator.free(round_trip);
+    try std.testing.expectEqualStrings(text, round_trip);
+}
+
+test "SPIR-V: decorated vertex interfaces and load-store operations" {
+    const assembly =
+        \\ OpCapability Shader
+        \\ OpMemoryModel Logical GLSL450
+        \\ OpEntryPoint Vertex %main "main" %in_color %out_color
+        \\ OpName %in_color "in_color"
+        \\ OpName %out_color "out_color"
+        \\ OpDecorate %in_color Location 0
+        \\ OpDecorate %out_color Location 0
+        \\
+        \\ %void = OpTypeVoid
+        \\ %float = OpTypeFloat 32
+        \\ %vec4 = OpTypeVector %float 4
+        \\ %input_vec4 = OpTypePointer Input %vec4
+        \\ %output_vec4 = OpTypePointer Output %vec4
+        \\ %fn_void = OpTypeFunction %void
+        \\ %in_color = OpVariable %input_vec4 Input
+        \\ %out_color = OpVariable %output_vec4 Output
+        \\
+        \\ %main = OpFunction %void None %fn_void
+        \\     %entry = OpLabel
+        \\     %color = OpLoad %vec4 %in_color
+        \\     OpStore %out_color %color
+        \\     OpReturn
+        \\ OpFunctionEnd
+    ;
+    const words = try assembleSpirv(std.testing.allocator, assembly);
+    defer std.testing.allocator.free(words);
+
+    var module = try translate(std.testing.allocator, words, .{ .entry_point = "main" });
+    defer module.deinit();
+    try std.testing.expectEqual(ir.module.Stage.vertex, module.stage);
+    try std.testing.expectEqual(@as(usize, 2), module.interface_variables.entries.items.len);
+
+    const text = try ir.printer.allocPrint(std.testing.allocator, &module);
+    defer std.testing.allocator.free(text);
+    try std.testing.expect(std.mem.indexOf(u8, text, "load_interface @in_color") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "store_interface @out_color") != null);
+}
+
+fn assembleSpirv(allocator: std.mem.Allocator, assembly: []const u8) ![]u32 {
+    var io_backend: std.Io.Threaded = .init(allocator, .{});
+    defer io_backend.deinit();
+    const io = io_backend.io();
+
+    var child = try std.process.spawn(io, .{
+        .argv = &.{ "spirv-as", "--target-env", "spv1.0", "-o", "-", "-" },
+        .stdin = .pipe,
+        .stdout = .pipe,
+        .stderr = .pipe,
+    });
+    defer child.kill(io);
+
+    {
+        const stdin = child.stdin.?;
+        var stdin_writer = stdin.writer(io, &.{});
+        try stdin_writer.interface.writeAll(assembly);
+        try stdin_writer.interface.flush();
+        stdin.close(io);
+        child.stdin = null;
+    }
+
+    var stdout_buffer: [4096]u8 = undefined;
+    var stdout_reader = child.stdout.?.reader(io, &stdout_buffer);
+    const binary = try stdout_reader.interface.allocRemaining(allocator, .limited(1024 * 1024));
+    defer allocator.free(binary);
+
+    var stderr_buffer: [4096]u8 = undefined;
+    var stderr_reader = child.stderr.?.reader(io, &stderr_buffer);
+    const stderr = try stderr_reader.interface.allocRemaining(allocator, .limited(64 * 1024));
+    defer allocator.free(stderr);
+
+    const term = try child.wait(io);
+    switch (term) {
+        .exited => |code| if (code != 0) {
+            std.log.err("spirv-as failed:\n{s}", .{stderr});
+            return error.SpirvAssemblyFailed;
+        },
+        else => {
+            std.log.err("spirv-as terminated unexpectedly:\n{s}", .{stderr});
+            return error.SpirvAssemblyFailed;
+        },
+    }
+
+    if (binary.len % @sizeOf(u32) != 0) return error.InvalidSpirvBinaryLength;
+    const words = try allocator.alloc(u32, binary.len / @sizeOf(u32));
+    errdefer allocator.free(words);
+    for (words, 0..) |*word, index| {
+        const offset = index * @sizeOf(u32);
+        word.* = std.mem.readInt(u32, binary[offset..][0..4], .little);
+    }
+    return words;
+}

src/intel/compiler/compiler.zig

@@ -0,0 +1,183 @@
+//! Flint-specific shader IR for Intel Gen hardware.
+//! This is the mutable, non-SSA layer between the common shader IR and machine code.
+
+pub const device = @import("device.zig");
+pub const ir = @import("ir/ir.zig");
+pub const lower = @import("lower/lower.zig");
+
+pub const id = ir.id;
+pub const instruction = ir.instruction;
+pub const operand = ir.operand;
+pub const printer = ir.printer;
+pub const program = ir.program;
+pub const validator = ir.validator;
+
+pub const Program = ir.Program;
+pub const Stage = ir.Stage;
+
+const std = @import("std");
+
+test "Flint IR foundation" {
+    // ; Flint program:
+    // ;   .stage: vertex
+    // ;   .generation: gen9
+    // ;   .platform: skylake
+    // ;   .dispatch_width: simd8
+    //
+    // %position: vgrf f32[8] = class(varying), size(32), alignment(32), spillable
+    // %urb_payload: vgrf u32[16] = class(payload), size(64), alignment(32)
+    //
+    // .entry:
+    //     [simd8] load_input %position:f32, location(0), component(0)
+    //     [simd8] multiply %position:f32, %position:f32, 1:f32
+    //     [simd8] mov %position:f32, %position:f32[byte=4, broadcast]
+    //     [simd8] store_output builtin(position), component(0), %position:f32
+    //     [simd8] send urb_write[offset(0), channels(xyzw), end_of_thread], payload(%urb_payload[2])
+    //     end_thread
+
+    const device_info: device.DeviceInfo = .{
+        .generation = .gen9,
+        .platform = .skylake,
+        .pci_device_id = 0x1912,
+        .grf_count = 128,
+    };
+
+    var shader = Program.init(std.testing.allocator, .vertex, device_info, .simd8);
+    defer shader.deinit();
+
+    const position = try shader.addVirtualRegister(.{
+        .size_bytes = 32,
+        .alignment_bytes = 32,
+        .element_type = .f32,
+        .lane_count = 8,
+        .class = .varying,
+        .name = "position",
+    });
+    const urb_payload = try shader.addVirtualRegister(.{
+        .size_bytes = 64,
+        .alignment_bytes = 32,
+        .element_type = .u32,
+        .lane_count = 16,
+        .class = .payload,
+        .spillable = false,
+        .name = "urb_payload",
+    });
+    const entry = try shader.addBlock("entry");
+    try shader.setEntryBlock(entry);
+
+    _ = try shader.appendInstruction(entry, .simd8, null, .{
+        .load_input = .{
+            .destination = .{
+                .register = .{ .virtual = position },
+                .type = .f32,
+            },
+            .semantic = .{
+                .location = .{
+                    .location = 0,
+                },
+            },
+        },
+    });
+    _ = try shader.appendInstruction(entry, .simd8, null, .{
+        .binary = .{
+            .opcode = .multiply,
+            .destination = .{
+                .register = .{ .virtual = position },
+                .type = .f32,
+            },
+            .lhs = .{
+                .register = .{ .virtual = position },
+                .type = .f32,
+                .region = operand.Region.contiguous(.simd8),
+            },
+            .rhs = .{
+                .register = .{
+                    .immediate = .{ .f32 = 1.0 },
+                },
+                .type = .f32,
+                .region = operand.Region.broadcast(),
+            },
+        },
+    });
+    _ = try shader.appendInstruction(entry, .simd8, null, .{
+        .move = .{
+            .destination = .{
+                .register = .{ .virtual = position },
+                .type = .f32,
+            },
+            .source = .{
+                .register = .{ .virtual = position },
+                .type = .f32,
+                .region = .{
+                    .byte_offset = 4,
+                    .vertical_stride = 0,
+                    .width = 1,
+                    .horizontal_stride = 0,
+                },
+            },
+        },
+    });
+    _ = try shader.appendInstruction(entry, .simd8, null, .{
+        .store_output = .{
+            .semantic = .{
+                .builtin = .{ .builtin = .position },
+            },
+            .source = .{
+                .register = .{ .virtual = position },
+                .type = .f32,
+                .region = operand.Region.contiguous(.simd8),
+            },
+        },
+    });
+    _ = try shader.appendInstruction(entry, .simd8, null, .{
+        .send = .{
+            .message = .{
+                .urb_write = .{
+                    .offset = 0,
+                    .end_of_thread = true,
+                },
+            },
+            .payload = .{
+                .base = .{ .virtual = urb_payload },
+                .register_count = 2,
+            },
+        },
+    });
+    try shader.setTerminator(entry, .end_thread);
+
+    shader.properties.instructions_selected = true;
+    try validator.validate(&shader);
+
+    try std.testing.expectEqual(entry, shader.entry_block.?);
+    try std.testing.expect(shader.properties.instructions_selected);
+
+    const text = try printer.allocPrint(std.testing.allocator, &shader);
+    defer std.testing.allocator.free(text);
+
+    try std.testing.expect(std.mem.indexOf(u8, text, "Flint program") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "vertex") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "gen9") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "skylake") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "simd8") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "%position: vgrf f32[8]") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "[simd8] multiply %position:f32, %position:f32, 1:f32") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "[simd8] mov %position:f32, %position:f32[byte=4, broadcast]") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "send urb_write[offset(0), channels(xyzw), end_of_thread], payload(%urb_payload[2])") != null);
+}
+
+test "ID stability after removal" {
+    var store: id.Store(id.VirtualFlagId, operand.VirtualFlag) = .{};
+    defer store.entries.deinit(std.testing.allocator);
+
+    const first = try store.add(std.testing.allocator, .{ .name = "first" });
+    try std.testing.expect(store.remove(first));
+    const second = try store.add(std.testing.allocator, .{ .name = "second" });
+
+    try std.testing.expect(first != second);
+    try std.testing.expect(store.get(first) == null);
+    try std.testing.expectEqualStrings("second", store.get(second).?.name.?);
+}
+
+test {
+    _ = lower;
+}

src/intel/compiler/device.zig

@@ -0,0 +1,56 @@
+pub const Generation = enum {
+    gen9,
+    gen10,
+    gen11,
+};
+
+pub const Platform = enum {
+    skylake,
+    broxton,
+    kabylake,
+    gemini_lake,
+    coffee_lake,
+    whiskey_lake,
+    comet_lake,
+    ice_lake,
+    elkhart_lake,
+    jasper_lake,
+};
+
+pub const DeviceInfo = struct {
+    generation: Generation,
+    platform: Platform,
+    pci_device_id: u16,
+
+    grf_count: u16,
+    grf_size_bytes: u16 = 32,
+
+    supports_int64: bool = false,
+    supports_float64: bool = false,
+    supports_half_float: bool = false,
+    supports_simd16: bool = false,
+    supports_simd32: bool = false,
+
+    pub fn supportsDispatch(self: DeviceInfo, width: DispatchWidth) bool {
+        return switch (width) {
+            .simd8 => true,
+            .simd16 => self.supports_simd16,
+            .simd32 => self.supports_simd32,
+        };
+    }
+};
+
+pub const DispatchWidth = enum(u8) {
+    simd8 = 8,
+    simd16 = 16,
+    simd32 = 32,
+};
+
+pub const ExecutionSize = enum(u8) {
+    simd1 = 1,
+    simd2 = 2,
+    simd4 = 4,
+    simd8 = 8,
+    simd16 = 16,
+    simd32 = 32,
+};

src/intel/compiler/ir/id.zig

@@ -0,0 +1,14 @@
+const shared_ids = @import("shader_ir").ir.id;
+
+pub const BlockTag = opaque {};
+pub const InstructionTag = opaque {};
+pub const VirtualRegisterTag = opaque {};
+pub const VirtualFlagTag = opaque {};
+
+pub const BlockId = shared_ids.Id(BlockTag);
+pub const InstructionId = shared_ids.Id(InstructionTag);
+pub const VirtualRegisterId = shared_ids.Id(VirtualRegisterTag);
+pub const VirtualFlagId = shared_ids.Id(VirtualFlagTag);
+
+pub const Id = shared_ids.Id;
+pub const Store = shared_ids.Store;

src/intel/compiler/ir/instruction.zig

@@ -0,0 +1,137 @@
+const std = @import("std");
+const device = @import("../device.zig");
+const ids = @import("id.zig");
+const operand = @import("operand.zig");
+
+pub const Builtin = enum {
+    position,
+    vertex_index,
+    instance_index,
+};
+
+pub const InterfaceSemantic = union(enum) {
+    location: struct {
+        location: u32,
+        component: u8 = 0,
+    },
+    builtin: struct {
+        builtin: Builtin,
+        component: u8 = 0,
+    },
+};
+
+pub const LoadInput = struct {
+    destination: operand.Destination,
+    semantic: InterfaceSemantic,
+};
+
+pub const StoreOutput = struct {
+    semantic: InterfaceSemantic,
+    source: operand.Source,
+};
+
+pub const Move = struct {
+    destination: operand.Destination,
+    source: operand.Source,
+};
+
+pub const BinaryOpcode = enum {
+    add,
+    multiply,
+    bitwise_and,
+    bitwise_or,
+    bitwise_xor,
+    shift_left,
+    shift_right,
+};
+
+pub const Binary = struct {
+    opcode: BinaryOpcode,
+    destination: operand.Destination,
+    lhs: operand.Source,
+    rhs: operand.Source,
+};
+
+pub const CompareOpcode = enum {
+    equal,
+    not_equal,
+    less_than,
+    less_or_equal,
+    greater_than,
+    greater_or_equal,
+};
+
+pub const Compare = struct {
+    opcode: CompareOpcode,
+    destination: operand.FlagRef,
+    lhs: operand.Source,
+    rhs: operand.Source,
+};
+
+pub const ChannelMask = packed struct(u4) {
+    x: bool = true,
+    y: bool = true,
+    z: bool = true,
+    w: bool = true,
+};
+
+pub const UrbWrite = struct {
+    offset: u16,
+    channels: ChannelMask = .{},
+    end_of_thread: bool = false,
+};
+
+pub const Message = union(enum) {
+    urb_write: UrbWrite,
+};
+
+pub const Send = struct {
+    message: Message,
+    payload: operand.RegisterSpan,
+    response: ?operand.RegisterSpan = null,
+};
+
+pub const Operation = union(enum) {
+    load_input: LoadInput,
+    store_output: StoreOutput,
+    move: Move,
+    binary: Binary,
+    compare: Compare,
+    send: Send,
+};
+
+pub const Instruction = struct {
+    parent_block: ids.BlockId,
+    execution_size: device.ExecutionSize,
+    predicate: ?operand.Predicate = null,
+    operation: Operation,
+};
+
+pub const Terminator = union(enum) {
+    jump: ids.BlockId,
+    conditional_branch: struct {
+        predicate: operand.Predicate,
+        true_block: ids.BlockId,
+        false_block: ids.BlockId,
+    },
+    end_thread,
+    @"unreachable",
+};
+
+pub const StructuredControl = union(enum) {
+    none,
+    selection: struct {
+        merge_block: ids.BlockId,
+    },
+    loop: struct {
+        merge_block: ids.BlockId,
+        continue_block: ids.BlockId,
+    },
+};
+
+pub const Block = struct {
+    instructions: std.ArrayList(ids.InstructionId) = .empty,
+    terminator: ?Terminator = null,
+    structured_control: StructuredControl = .none,
+    name: ?[]const u8 = null,
+};

src/intel/compiler/ir/ir.zig

@@ -0,0 +1,9 @@
+pub const id = @import("id.zig");
+pub const instruction = @import("instruction.zig");
+pub const operand = @import("operand.zig");
+pub const printer = @import("printer.zig");
+pub const program = @import("program.zig");
+pub const validator = @import("validator.zig");
+
+pub const Program = program.Program;
+pub const Stage = program.Stage;

src/intel/compiler/ir/operand.zig

@@ -0,0 +1,148 @@
+const device = @import("../device.zig");
+const ids = @import("id.zig");
+
+pub const DataType = enum {
+    u8,
+    i8,
+    u16,
+    i16,
+    f16,
+    u32,
+    i32,
+    f32,
+    u64,
+    i64,
+    f64,
+
+    pub fn sizeBytes(self: DataType) u8 {
+        return switch (self) {
+            .u8, .i8 => 1,
+            .u16, .i16, .f16 => 2,
+            .u32, .i32, .f32 => 4,
+            .u64, .i64, .f64 => 8,
+        };
+    }
+
+    pub fn isInitialTargetType(self: DataType) bool {
+        return switch (self) {
+            .u32, .i32, .f32 => true,
+            else => false,
+        };
+    }
+};
+
+pub const RegisterClass = enum {
+    uniform,
+    varying,
+    payload,
+    response,
+    temporary,
+};
+
+pub const VirtualRegister = struct {
+    size_bytes: u32,
+    alignment_bytes: u16,
+    element_type: DataType,
+    lane_count: u8,
+    class: RegisterClass,
+    spillable: bool = true,
+    name: ?[]const u8 = null,
+};
+
+pub const VirtualFlag = struct {
+    name: ?[]const u8 = null,
+};
+
+pub const PhysicalGrf = struct {
+    number: u16,
+    byte_offset: u8 = 0,
+};
+
+pub const PhysicalFlag = struct {
+    register: u8 = 0,
+    subregister: u8 = 0,
+};
+
+pub const ArchitectureRegister = union(enum) {
+    flag: u8,
+    address: u8,
+    accumulator: u8,
+    notification: u8,
+    instruction_pointer,
+};
+
+pub const Immediate = union(enum) {
+    u32: u32,
+    i32: i32,
+    f32: f32,
+};
+
+pub const RegisterRef = union(enum) {
+    virtual: ids.VirtualRegisterId,
+    physical_grf: PhysicalGrf,
+    architecture: ArchitectureRegister,
+    immediate: Immediate,
+    null,
+};
+
+pub const FlagRef = union(enum) {
+    virtual: ids.VirtualFlagId,
+    physical: PhysicalFlag,
+};
+
+pub const Predicate = struct {
+    flag: FlagRef,
+    inverse: bool = false,
+};
+
+pub const Region = struct {
+    byte_offset: u16 = 0,
+    vertical_stride: u8,
+    width: u8,
+    horizontal_stride: u8,
+
+    pub fn scalar() Region {
+        return .{
+            .vertical_stride = 0,
+            .width = 1,
+            .horizontal_stride = 0,
+        };
+    }
+
+    pub fn contiguous(execution_size: device.ExecutionSize) Region {
+        const width: u8 = @intFromEnum(execution_size);
+        return .{
+            .vertical_stride = width,
+            .width = width,
+            .horizontal_stride = 1,
+        };
+    }
+
+    pub fn broadcast() Region {
+        return scalar();
+    }
+};
+
+pub const DestinationRegion = struct {
+    byte_offset: u16 = 0,
+    horizontal_stride: u8 = 1,
+};
+
+pub const Source = struct {
+    register: RegisterRef,
+    type: DataType,
+    region: Region,
+    negate: bool = false,
+    absolute: bool = false,
+};
+
+pub const Destination = struct {
+    register: RegisterRef,
+    type: DataType,
+    region: DestinationRegion = .{},
+};
+
+pub const RegisterSpan = struct {
+    base: RegisterRef,
+    register_count: u8,
+};

src/intel/compiler/ir/printer.zig

@@ -0,0 +1,352 @@
+const std = @import("std");
+const device = @import("../device.zig");
+const ids = @import("id.zig");
+const inst_ir = @import("instruction.zig");
+const operand = @import("operand.zig");
+const program_ir = @import("program.zig");
+
+const indent = "    ";
+
+pub fn write(program: *const program_ir.Program, writer: *std.Io.Writer) std.Io.Writer.Error!void {
+    try writer.writeAll("; Flint program:\n");
+    try writer.print(";   .stage: {t}\n", .{program.stage});
+    try writer.print(";   .generation: {t}\n", .{program.device_info.generation});
+    try writer.print(";   .platform: {t}\n", .{program.device_info.platform});
+    try writer.print(";   .dispatch_width: {t}\n\n", .{program.dispatch_width});
+
+    for (program.virtual_registers.entries.items, 0..) |entry, index| {
+        const register = entry orelse continue;
+        try writeVirtualRegisterRef(program, writer, ids.VirtualRegisterId.fromIndex(index));
+        try writer.print(": vgrf {t}[{d}], class({t}), size({d}), alignment({d}){s}\n", .{
+            register.element_type,
+            register.lane_count,
+            register.class,
+            register.size_bytes,
+            register.alignment_bytes,
+            if (register.spillable) ", spillable" else "",
+        });
+    }
+
+    for (program.virtual_flags.entries.items, 0..) |entry, index| {
+        _ = entry orelse continue;
+        try writeVirtualFlagRef(program, writer, ids.VirtualFlagId.fromIndex(index));
+        try writer.writeAll(": vflag\n");
+    }
+
+    try writer.writeByte('\n');
+
+    for (program.blocks.entries.items, 0..) |entry, block_index| {
+        const block = entry orelse continue;
+        const block_id = ids.BlockId.fromIndex(block_index);
+
+        try writeBlockRef(program, writer, block_id);
+        try writer.writeAll(":\n");
+
+        switch (block.structured_control) {
+            .none => {},
+            .selection => |selection| {
+                try writer.writeAll(indent ++ "structured_selection ");
+                try writeBlockRef(program, writer, selection.merge_block);
+                try writer.writeByte('\n');
+            },
+            .loop => |loop| {
+                try writer.writeAll(indent ++ "structured_loop merge(");
+                try writeBlockRef(program, writer, loop.merge_block);
+                try writer.writeAll("), continue(");
+                try writeBlockRef(program, writer, loop.continue_block);
+                try writer.writeAll(")\n");
+            },
+        }
+
+        for (block.instructions.items) |instruction_id| {
+            const instruction = program.instructions.get(instruction_id) orelse continue;
+            try writer.writeAll(indent);
+            try writeInstruction(program, writer, instruction.*);
+            try writer.writeByte('\n');
+        }
+
+        if (block.terminator) |terminator| {
+            try writer.writeAll(indent);
+            try writeTerminator(program, writer, terminator);
+            try writer.writeAll("\n\n");
+        } else {
+            try writer.writeAll(indent ++ "<missing terminator>\n\n");
+        }
+    }
+}
+
+pub fn allocPrint(allocator: std.mem.Allocator, program: *const program_ir.Program) ![]u8 {
+    var output: std.Io.Writer.Allocating = .init(allocator);
+    defer output.deinit();
+    try write(program, &output.writer);
+    return output.toOwnedSlice();
+}
+
+fn writeInstruction(program: *const program_ir.Program, writer: *std.Io.Writer, instruction: inst_ir.Instruction) !void {
+    try writer.print("[simd{d}] ", .{@intFromEnum(instruction.execution_size)});
+    if (instruction.predicate) |predicate| {
+        try writePredicate(program, writer, predicate);
+        try writer.writeByte(' ');
+    }
+    try writeOperation(program, writer, instruction.execution_size, instruction.operation);
+}
+
+fn writeOperation(program: *const program_ir.Program, writer: *std.Io.Writer, execution_size: device.ExecutionSize, operation: inst_ir.Operation) !void {
+    switch (operation) {
+        .load_input => |op| {
+            try writer.writeAll("load_input ");
+            try writeDestination(program, writer, execution_size, op.destination);
+            try writer.writeAll(", ");
+            try writeInterfaceSemantic(writer, op.semantic);
+        },
+        .store_output => |op| {
+            try writer.writeAll("store_output ");
+            try writeInterfaceSemantic(writer, op.semantic);
+            try writer.writeAll(", ");
+            try writeSource(program, writer, execution_size, op.source);
+        },
+        .move => |op| {
+            try writer.writeAll("mov ");
+            try writeDestination(program, writer, execution_size, op.destination);
+            try writer.writeAll(", ");
+            try writeSource(program, writer, execution_size, op.source);
+        },
+        .binary => |op| {
+            try writer.print("{t} ", .{op.opcode});
+            try writeDestination(program, writer, execution_size, op.destination);
+            try writer.writeAll(", ");
+            try writeSource(program, writer, execution_size, op.lhs);
+            try writer.writeAll(", ");
+            try writeSource(program, writer, execution_size, op.rhs);
+        },
+        .compare => |op| {
+            try writer.print("cmp_{t} ", .{op.opcode});
+            try writeFlagRef(program, writer, op.destination);
+            try writer.writeAll(", ");
+            try writeSource(program, writer, execution_size, op.lhs);
+            try writer.writeAll(", ");
+            try writeSource(program, writer, execution_size, op.rhs);
+        },
+        .send => |op| {
+            try writer.writeAll("send ");
+            if (op.response) |response| {
+                try writeRegisterSpan(program, writer, response);
+                try writer.writeAll(", ");
+            }
+            try writeMessage(writer, op.message);
+            try writer.writeAll(", payload(");
+            try writeRegisterSpan(program, writer, op.payload);
+            try writer.writeByte(')');
+        },
+    }
+}
+
+fn writeTerminator(program: *const program_ir.Program, writer: *std.Io.Writer, terminator: inst_ir.Terminator) !void {
+    switch (terminator) {
+        .jump => |target| {
+            try writer.writeAll("jump ");
+            try writeBlockRef(program, writer, target);
+        },
+        .conditional_branch => |branch| {
+            try writer.writeAll("conditional_branch ");
+            try writePredicate(program, writer, branch.predicate);
+            try writer.writeAll(", ");
+            try writeBlockRef(program, writer, branch.true_block);
+            try writer.writeAll(", ");
+            try writeBlockRef(program, writer, branch.false_block);
+        },
+        .end_thread => try writer.writeAll("end_thread"),
+        .@"unreachable" => try writer.writeAll("unreachable"),
+    }
+}
+
+fn writeSource(program: *const program_ir.Program, writer: *std.Io.Writer, execution_size: device.ExecutionSize, source: operand.Source) !void {
+    if (source.negate)
+        try writer.writeByte('-');
+    if (source.absolute)
+        try writer.writeAll("abs(");
+
+    try writeRegister(program, writer, source.register);
+    try writer.print(":{t}", .{source.type});
+    if (source.register != .immediate)
+        try writeSourceRegion(writer, execution_size, source.register, source.region);
+
+    if (source.absolute)
+        try writer.writeByte(')');
+}
+
+fn writeDestination(program: *const program_ir.Program, writer: *std.Io.Writer, execution_size: device.ExecutionSize, destination: operand.Destination) !void {
+    _ = execution_size;
+    try writeRegister(program, writer, destination.register);
+    try writer.print(":{t}", .{destination.type});
+    try writeDestinationRegion(writer, destination.register, destination.region);
+}
+
+fn writeSourceRegion(writer: *std.Io.Writer, execution_size: device.ExecutionSize, register: operand.RegisterRef, region: operand.Region) !void {
+    const byte_offset = registerByteOffset(register) + region.byte_offset;
+    const execution_width: u8 = @intFromEnum(execution_size);
+    const is_default = region.vertical_stride == execution_width and
+        region.width == execution_width and
+        region.horizontal_stride == 1;
+    const is_broadcast = region.vertical_stride == 0 and
+        region.width == 1 and
+        region.horizontal_stride == 0;
+
+    if (byte_offset == 0 and is_default)
+        return;
+
+    try writer.writeByte('[');
+    if (byte_offset != 0)
+        try writer.print("byte={d}", .{byte_offset});
+
+    if (is_broadcast) {
+        if (byte_offset != 0)
+            try writer.writeAll(", ");
+        try writer.writeAll("broadcast");
+    } else if (!is_default) {
+        if (byte_offset != 0)
+            try writer.writeAll(", ");
+        try writer.print("vstride={d}, width={d}, hstride={d}", .{
+            region.vertical_stride,
+            region.width,
+            region.horizontal_stride,
+        });
+    }
+    try writer.writeByte(']');
+}
+
+fn writeDestinationRegion(writer: *std.Io.Writer, register: operand.RegisterRef, region: operand.DestinationRegion) !void {
+    const byte_offset = registerByteOffset(register) + region.byte_offset;
+    if (byte_offset == 0 and region.horizontal_stride == 1)
+        return;
+
+    try writer.writeByte('[');
+    if (byte_offset != 0)
+        try writer.print("byte={d}", .{byte_offset});
+    if (region.horizontal_stride != 1) {
+        if (byte_offset != 0)
+            try writer.writeAll(", ");
+        try writer.print("hstride={d}", .{region.horizontal_stride});
+    }
+    try writer.writeByte(']');
+}
+
+fn registerByteOffset(register: operand.RegisterRef) u16 {
+    return switch (register) {
+        .physical_grf => |physical| physical.byte_offset,
+        else => 0,
+    };
+}
+
+fn writeRegister(program: *const program_ir.Program, writer: *std.Io.Writer, register: operand.RegisterRef) !void {
+    switch (register) {
+        .virtual => |virtual| try writeVirtualRegisterRef(program, writer, virtual),
+        .physical_grf => |physical| try writer.print("r{d}", .{physical.number}),
+        .architecture => |architecture| try writeArchitectureRegister(writer, architecture),
+        .immediate => |immediate| try writeImmediate(writer, immediate),
+        .null => try writer.writeAll("null"),
+    }
+}
+
+fn writeArchitectureRegister(writer: *std.Io.Writer, register: operand.ArchitectureRegister) !void {
+    switch (register) {
+        .flag => |index| try writer.print("f{d}", .{index}),
+        .address => |index| try writer.print("a{d}", .{index}),
+        .accumulator => |index| try writer.print("acc{d}", .{index}),
+        .notification => |index| try writer.print("n{d}", .{index}),
+        .instruction_pointer => try writer.writeAll("ip"),
+    }
+}
+
+fn writeImmediate(writer: *std.Io.Writer, immediate: operand.Immediate) !void {
+    switch (immediate) {
+        .u32 => |value| try writer.print("{d}", .{value}),
+        .i32 => |value| try writer.print("{d}", .{value}),
+        .f32 => |value| try writer.print("{d}", .{value}),
+    }
+}
+
+fn writePredicate(program: *const program_ir.Program, writer: *std.Io.Writer, predicate: operand.Predicate) !void {
+    try writer.writeAll(if (predicate.inverse) "(-" else "(+");
+    try writeFlagRef(program, writer, predicate.flag);
+    try writer.writeByte(')');
+}
+
+fn writeFlagRef(program: *const program_ir.Program, writer: *std.Io.Writer, flag: operand.FlagRef) !void {
+    switch (flag) {
+        .virtual => |virtual| try writeVirtualFlagRef(program, writer, virtual),
+        .physical => |physical| try writer.print("f{d}.{d}", .{ physical.register, physical.subregister }),
+    }
+}
+
+fn writeRegisterSpan(program: *const program_ir.Program, writer: *std.Io.Writer, span: operand.RegisterSpan) !void {
+    try writeRegister(program, writer, span.base);
+    const byte_offset = registerByteOffset(span.base);
+    if (byte_offset != 0)
+        try writer.print("[byte={d}]", .{byte_offset});
+    try writer.print("[{d}]", .{span.register_count});
+}
+
+fn writeInterfaceSemantic(writer: *std.Io.Writer, semantic: inst_ir.InterfaceSemantic) !void {
+    switch (semantic) {
+        .location => |location| try writer.print("location({d}), component({d})", .{ location.location, location.component }),
+        .builtin => |builtin| try writer.print("builtin({t}), component({d})", .{ builtin.builtin, builtin.component }),
+    }
+}
+
+fn writeMessage(writer: *std.Io.Writer, message: inst_ir.Message) !void {
+    switch (message) {
+        .urb_write => |urb| {
+            try writer.print("urb_write[offset({d}), channels(", .{urb.offset});
+            try writeChannelMask(writer, urb.channels);
+            try writer.writeByte(')');
+            if (urb.end_of_thread)
+                try writer.writeAll(", end_of_thread");
+            try writer.writeByte(']');
+        },
+    }
+}
+
+fn writeChannelMask(writer: *std.Io.Writer, mask: inst_ir.ChannelMask) !void {
+    if (mask.x) try writer.writeByte('x');
+    if (mask.y) try writer.writeByte('y');
+    if (mask.z) try writer.writeByte('z');
+    if (mask.w) try writer.writeByte('w');
+}
+
+fn writeVirtualRegisterRef(program: *const program_ir.Program, writer: *std.Io.Writer, register_id: ids.VirtualRegisterId) !void {
+    const register = program.virtual_registers.get(register_id);
+    try writeNamedRef(writer, if (register) |value| value.name else null, "v", register_id.index(), '%');
+}
+
+fn writeVirtualFlagRef(program: *const program_ir.Program, writer: *std.Io.Writer, flag_id: ids.VirtualFlagId) !void {
+    const flag = program.virtual_flags.get(flag_id);
+    try writeNamedRef(writer, if (flag) |value| value.name else null, "f", flag_id.index(), '%');
+}
+
+fn writeBlockRef(program: *const program_ir.Program, writer: *std.Io.Writer, block_id: ids.BlockId) !void {
+    const block = program.blocks.get(block_id);
+    try writeNamedRef(writer, if (block) |value| value.name else null, "b", block_id.index(), '.');
+}
+
+fn writeNamedRef(writer: *std.Io.Writer, name: ?[]const u8, fallback: []const u8, index: usize, prefix: u8) !void {
+    try writer.writeByte(prefix);
+    if (name) |text| {
+        if (isValidName(text)) {
+            try writer.writeAll(text);
+            return;
+        }
+    }
+    try writer.print("{s}{d}", .{ fallback, index });
+}
+
+fn isValidName(name: []const u8) bool {
+    if (name.len == 0 or (!std.ascii.isAlphabetic(name[0]) and name[0] != '_'))
+        return false;
+
+    for (name[1..]) |byte| {
+        if (!std.ascii.isAlphanumeric(byte) and byte != '_')
+            return false;
+    }
+    return true;
+}

src/intel/compiler/ir/program.zig

@@ -0,0 +1,135 @@
+const std = @import("std");
+const shared_ir = @import("shader_ir").ir.module;
+const device = @import("../device.zig");
+const ids = @import("id.zig");
+const instructions = @import("instruction.zig");
+const operand = @import("operand.zig");
+
+pub const Stage = shared_ir.Stage;
+
+pub const Properties = packed struct {
+    instructions_selected: bool = false,
+    block_parameters_lowered: bool = false,
+
+    stage_io_lowered: bool = false,
+    resources_lowered: bool = false,
+    messages_lowered: bool = false,
+    control_flow_lowered: bool = false,
+
+    regions_legalized: bool = false,
+    types_legalized: bool = false,
+
+    registers_allocated: bool = false,
+    flags_allocated: bool = false,
+    branches_resolved: bool = false,
+
+    _padding: u21 = 0,
+};
+
+pub const VertexPayload = struct {
+    first_attribute_grf: operand.PhysicalGrf,
+    attribute_grf_count: u16,
+};
+
+pub const PayloadLayout = struct {
+    header_grf: ?operand.PhysicalGrf = null,
+    vertex: ?VertexPayload = null,
+};
+
+pub const ProgramData = struct {
+    payload_grf_count: u16 = 0,
+    total_grf_count: u16 = 0,
+    scratch_size_bytes: u32 = 0,
+};
+
+pub const BlockStore = ids.Store(ids.BlockId, instructions.Block);
+pub const InstructionStore = ids.Store(ids.InstructionId, instructions.Instruction);
+pub const VirtualRegisterStore = ids.Store(ids.VirtualRegisterId, operand.VirtualRegister);
+pub const VirtualFlagStore = ids.Store(ids.VirtualFlagId, operand.VirtualFlag);
+
+pub const Program = struct {
+    arena: std.heap.ArenaAllocator,
+
+    stage: Stage,
+    device_info: device.DeviceInfo,
+    dispatch_width: device.DispatchWidth,
+
+    entry_block: ?ids.BlockId = null,
+
+    blocks: BlockStore = .{},
+    instructions: InstructionStore = .{},
+    virtual_registers: VirtualRegisterStore = .{},
+    virtual_flags: VirtualFlagStore = .{},
+
+    payload: PayloadLayout = .{},
+    program_data: ProgramData = .{},
+    properties: Properties = .{},
+
+    pub fn init(backing_allocator: std.mem.Allocator, stage: Stage, device_info: device.DeviceInfo, dispatch_width: device.DispatchWidth) Program {
+        return .{
+            .arena = std.heap.ArenaAllocator.init(backing_allocator),
+            .stage = stage,
+            .device_info = device_info,
+            .dispatch_width = dispatch_width,
+        };
+    }
+
+    pub fn deinit(self: *Program) void {
+        self.arena.deinit();
+        self.* = undefined;
+    }
+
+    pub fn allocator(self: *Program) std.mem.Allocator {
+        return self.arena.allocator();
+    }
+
+    pub fn addVirtualRegister(self: *Program, register: operand.VirtualRegister) !ids.VirtualRegisterId {
+        var owned = register;
+        if (register.name) |name|
+            owned.name = try self.allocator().dupe(u8, name);
+        return self.virtual_registers.add(self.allocator(), owned);
+    }
+
+    pub fn addVirtualFlag(self: *Program, flag: operand.VirtualFlag) !ids.VirtualFlagId {
+        var owned = flag;
+        if (flag.name) |name|
+            owned.name = try self.allocator().dupe(u8, name);
+        return self.virtual_flags.add(self.allocator(), owned);
+    }
+
+    pub fn addBlock(self: *Program, name: ?[]const u8) !ids.BlockId {
+        const owned_name = if (name) |value| try self.allocator().dupe(u8, value) else null;
+        const block_id = try self.blocks.add(self.allocator(), .{
+            .name = owned_name,
+        });
+        if (self.entry_block == null)
+            self.entry_block = block_id;
+        return block_id;
+    }
+
+    pub fn setEntryBlock(self: *Program, block_id: ids.BlockId) !void {
+        if (!self.blocks.isLive(block_id))
+            return error.InvalidBlock;
+        self.entry_block = block_id;
+    }
+
+    pub fn appendInstruction(self: *Program, block_id: ids.BlockId, execution_size: device.ExecutionSize, predicate: ?operand.Predicate, operation: instructions.Operation) !ids.InstructionId {
+        const block = self.blocks.getMut(block_id) orelse return error.InvalidBlock;
+
+        const instruction_id = try self.instructions.add(self.allocator(), .{
+            .parent_block = block_id,
+            .execution_size = execution_size,
+            .predicate = predicate,
+            .operation = operation,
+        });
+        try block.instructions.append(self.allocator(), instruction_id);
+        return instruction_id;
+    }
+
+    pub fn setTerminator(self: *Program, block_id: ids.BlockId, terminator: instructions.Terminator) !void {
+        const block = self.blocks.getMut(block_id) orelse return error.InvalidBlock;
+        if (block.terminator != null)
+            return error.TerminatorAlreadySet;
+        block.terminator = terminator;
+    }
+};

src/intel/compiler/ir/validator.zig

@@ -0,0 +1,194 @@
+const ids = @import("id.zig");
+const instruction = @import("instruction.zig");
+const operand = @import("operand.zig");
+const program_ir = @import("program.zig");
+
+pub const Error = error{
+    UnsupportedGeneration,
+    UnsupportedStage,
+    UnsupportedDispatchWidth,
+    UnsupportedExecutionSize,
+    UnsupportedDataType,
+    MissingEntryBlock,
+    InvalidBlock,
+    MissingTerminator,
+    InvalidInstruction,
+    InvalidVirtualRegister,
+    InvalidVirtualFlag,
+    InvalidPhysicalRegister,
+    InvalidPhysicalFlag,
+    InvalidRegisterSize,
+    InvalidRegisterAlignment,
+    InvalidLaneCount,
+    InvalidRegion,
+    InvalidDestination,
+    InvalidImmediateType,
+    InvalidRegisterSpan,
+};
+
+pub fn validate(program: *const program_ir.Program) Error!void {
+    if (program.device_info.generation != .gen9)
+        return error.UnsupportedGeneration;
+    if (program.stage != .vertex)
+        return error.UnsupportedStage;
+    if (program.dispatch_width != .simd8 or !program.device_info.supportsDispatch(.simd8))
+        return error.UnsupportedDispatchWidth;
+
+    const entry_block = program.entry_block orelse return error.MissingEntryBlock;
+    if (!program.blocks.isLive(entry_block))
+        return error.InvalidBlock;
+
+    for (program.virtual_registers.entries.items) |entry| {
+        const register = entry orelse continue;
+        if (register.size_bytes == 0)
+            return error.InvalidRegisterSize;
+        if (register.alignment_bytes == 0 or
+            (register.alignment_bytes & (register.alignment_bytes - 1)) != 0)
+            return error.InvalidRegisterAlignment;
+        if (register.lane_count == 0)
+            return error.InvalidLaneCount;
+        try validateType(register.element_type);
+    }
+
+    for (program.blocks.entries.items, 0..) |entry, block_index| {
+        const block = entry orelse continue;
+        if (block.terminator == null)
+            return error.MissingTerminator;
+
+        const block_id = ids.BlockId.fromIndex(block_index);
+        for (block.instructions.items) |instruction_id| {
+            const inst = program.instructions.get(instruction_id) orelse return error.InvalidInstruction;
+            if (inst.parent_block != block_id)
+                return error.InvalidInstruction;
+            try validateInstruction(program, inst.*);
+        }
+
+        try validateStructuredControl(program, block.structured_control);
+        try validateTerminator(program, block.terminator.?);
+    }
+}
+
+fn validateInstruction(program: *const program_ir.Program, inst: instruction.Instruction) Error!void {
+    switch (inst.execution_size) {
+        .simd1, .simd8 => {},
+        else => return error.UnsupportedExecutionSize,
+    }
+
+    if (inst.predicate) |predicate|
+        try validateFlag(program, predicate.flag);
+
+    switch (inst.operation) {
+        .load_input => |op| try validateDestination(program, op.destination),
+        .store_output => |op| try validateSource(program, op.source),
+        .move => |op| {
+            try validateDestination(program, op.destination);
+            try validateSource(program, op.source);
+        },
+        .binary => |op| {
+            try validateDestination(program, op.destination);
+            try validateSource(program, op.lhs);
+            try validateSource(program, op.rhs);
+        },
+        .compare => |op| {
+            try validateFlag(program, op.destination);
+            try validateSource(program, op.lhs);
+            try validateSource(program, op.rhs);
+        },
+        .send => |op| {
+            try validateSpan(program, op.payload);
+            if (op.response) |response|
+                try validateSpan(program, response);
+        },
+    }
+}
+
+fn validateType(data_type: operand.DataType) Error!void {
+    if (!data_type.isInitialTargetType())
+        return error.UnsupportedDataType;
+}
+
+fn validateSource(program: *const program_ir.Program, source: operand.Source) Error!void {
+    try validateType(source.type);
+    if (source.region.width == 0)
+        return error.InvalidRegion;
+    try validateRegisterRef(program, source.register);
+
+    if (source.register == .immediate) {
+        const matches = switch (source.register.immediate) {
+            .u32 => source.type == .u32,
+            .i32 => source.type == .i32,
+            .f32 => source.type == .f32,
+        };
+        if (!matches)
+            return error.InvalidImmediateType;
+    }
+}
+
+fn validateDestination(program: *const program_ir.Program, destination: operand.Destination) Error!void {
+    try validateType(destination.type);
+    if (destination.region.horizontal_stride == 0)
+        return error.InvalidRegion;
+    switch (destination.register) {
+        .immediate, .null => return error.InvalidDestination,
+        else => try validateRegisterRef(program, destination.register),
+    }
+}
+
+fn validateRegisterRef(program: *const program_ir.Program, register: operand.RegisterRef) Error!void {
+    switch (register) {
+        .virtual => |id| if (!program.virtual_registers.isLive(id))
+            return error.InvalidVirtualRegister,
+        .physical_grf => |physical| {
+            if (physical.number >= program.device_info.grf_count or
+                physical.byte_offset >= program.device_info.grf_size_bytes)
+                return error.InvalidPhysicalRegister;
+        },
+        .architecture, .immediate, .null => {},
+    }
+}
+
+fn validateFlag(program: *const program_ir.Program, flag: operand.FlagRef) Error!void {
+    switch (flag) {
+        .virtual => |id| if (!program.virtual_flags.isLive(id))
+            return error.InvalidVirtualFlag,
+        .physical => |physical| if (physical.register != 0 or physical.subregister > 1)
+            return error.InvalidPhysicalFlag,
+    }
+}
+
+fn validateSpan(program: *const program_ir.Program, span: operand.RegisterSpan) Error!void {
+    if (span.register_count == 0)
+        return error.InvalidRegisterSpan;
+    switch (span.base) {
+        .virtual, .physical_grf => try validateRegisterRef(program, span.base),
+        else => return error.InvalidRegisterSpan,
+    }
+}
+
+fn validateTerminator(program: *const program_ir.Program, terminator: instruction.Terminator) Error!void {
+    switch (terminator) {
+        .jump => |target| try validateBlockTarget(program, target),
+        .conditional_branch => |branch| {
+            try validateFlag(program, branch.predicate.flag);
+            try validateBlockTarget(program, branch.true_block);
+            try validateBlockTarget(program, branch.false_block);
+        },
+        .end_thread, .@"unreachable" => {},
+    }
+}
+
+fn validateStructuredControl(program: *const program_ir.Program, control: instruction.StructuredControl) Error!void {
+    switch (control) {
+        .none => {},
+        .selection => |selection| try validateBlockTarget(program, selection.merge_block),
+        .loop => |loop| {
+            try validateBlockTarget(program, loop.merge_block);
+            try validateBlockTarget(program, loop.continue_block);
+        },
+    }
+}
+
+fn validateBlockTarget(program: *const program_ir.Program, block_id: ids.BlockId) Error!void {
+    if (!program.blocks.isLive(block_id))
+        return error.InvalidBlock;
+}

src/intel/lib.zig

@@ -2,15 +2,12 @@ const std = @import("std");
 const vk = @import("vulkan");
 pub const base = @import("base");
 
+pub const kmd = @import("kmd.zig");
+pub const compiler = @import("compiler/compiler.zig");
+
 pub const c = @import("intel_c");
 pub const config = base.config;
 
-pub const FlintInstance = @import("FlintInstance.zig");
-pub const FlintDevice = @import("FlintDevice.zig");
-pub const FlintPhysicalDevice = @import("FlintPhysicalDevice.zig");
-pub const FlintQueue = @import("FlintQueue.zig");
-pub const kmd = @import("kmd.zig");
-
 pub const FlintBinarySemaphore = @import("FlintBinarySemaphore.zig");
 pub const FlintBuffer = @import("FlintBuffer.zig");
 pub const FlintBufferView = @import("FlintBufferView.zig");
@@ -19,16 +16,20 @@ pub const FlintCommandPool = @import("FlintCommandPool.zig");
 pub const FlintDescriptorPool = @import("FlintDescriptorPool.zig");
 pub const FlintDescriptorSet = @import("FlintDescriptorSet.zig");
 pub const FlintDescriptorSetLayout = @import("FlintDescriptorSetLayout.zig");
+pub const FlintDevice = @import("FlintDevice.zig");
 pub const FlintDeviceMemory = @import("FlintDeviceMemory.zig");
 pub const FlintEvent = @import("FlintEvent.zig");
 pub const FlintFence = @import("FlintFence.zig");
 pub const FlintFramebuffer = @import("FlintFramebuffer.zig");
 pub const FlintImage = @import("FlintImage.zig");
 pub const FlintImageView = @import("FlintImageView.zig");
+pub const FlintInstance = @import("FlintInstance.zig");
+pub const FlintPhysicalDevice = @import("FlintPhysicalDevice.zig");
 pub const FlintPipeline = @import("FlintPipeline.zig");
 pub const FlintPipelineCache = @import("FlintPipelineCache.zig");
 pub const FlintPipelineLayout = @import("FlintPipelineLayout.zig");
 pub const FlintQueryPool = @import("FlintQueryPool.zig");
+pub const FlintQueue = @import("FlintQueue.zig");
 pub const FlintRenderPass = @import("FlintRenderPass.zig");
 pub const FlintSampler = @import("FlintSampler.zig");
 pub const FlintShaderModule = @import("FlintShaderModule.zig");
@@ -89,6 +90,7 @@ test {
     std.testing.refAllDecls(FlintRenderPass);
     std.testing.refAllDecls(FlintSampler);
     std.testing.refAllDecls(FlintShaderModule);
+    std.testing.refAllDecls(compiler);
     std.testing.refAllDecls(kmd);
     std.testing.refAllDecls(base);
 }

src/software/SoftCommandBuffer.zig

@@ -128,9 +128,10 @@ pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const v
 
     self.* = .{
         .interface = interface,
-        .command_allocator = .init(interface.host_allocator.allocator()),
+        .command_allocator = undefined,
         .commands = .empty,
     };
+    self.command_allocator = .init(interface.host_allocator.allocator());
     return self;
 }