[Flint] adding IR lowering and printer

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

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 old_value = self.module.values.get(old) orelse return Error.InvalidValue;
-    const replacement_value = self.module.values.get(replacement) 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(
+pub fn redirectEdges(self: *Self, source: ids.BlockId, old_target: ids.BlockId, new_target: ids.BlockId, new_arguments: []const ids.ValueId) Error!usize {
-    self: *Self,
+    const source_block = self.module.blocks.get(source) orelse return Error.InvalidBlock;
-    source: ids.BlockId,
+    const target_block = self.module.blocks.get(new_target) orelse return Error.InvalidBlock;
-    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(
+pub fn addBlockParameter(self: *Self, block_id: ids.BlockId, ty: ids.TypeId, name: ?[]const u8, incoming: []const IncomingValue) Error!ids.ValueId {
-    self: *Self,
+    const block = self.module.blocks.get(block_id) orelse return Error.InvalidBlock;
-    block_id: ids.BlockId,
+    const function = self.module.functions.get(block.parent_function) orelse return Error.InvalidFunction;
-    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(
+pub fn removeBlockParameter(self: *Self, block_id: ids.BlockId, parameter_index: usize, replacement: ids.ValueId) Error!void {
-    self: *Self,
+    const block = self.module.blocks.get(block_id) orelse return Error.InvalidBlock;
-    block_id: ids.BlockId,
+    if (parameter_index >= block.parameters.items.len) return Error.InvalidParameterIndex;
-    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 argument_value = self.module.values.get(argument) orelse return Error.InvalidValue;
-        const parameter_value = self.module.values.get(parameter) 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;
+        if (argument_value.type != parameter_value.type) return Error.TypeMismatch;
     }
 }
 
-fn redirectOne(
+fn redirectOne(self: *Self, edge: *module_ir.Edge, old_target: ids.BlockId, new_target: ids.BlockId, arguments: []const ids.ValueId) !bool {
-    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 block = self.module.blocks.getMut(predecessor) orelse return Error.InvalidBlock;
-    const terminator = if (block.terminator) |*item| item else 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 block = self.module.blocks.getMut(predecessor) orelse return Error.InvalidBlock;
-    const terminator = if (block.terminator) |*item| item else 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(
+fn functionHasEdgeTo(module: *const module_ir.Module, function: *const module_ir.Function, predecessor: ids.BlockId, target: ids.BlockId) bool {
-    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 function = module.functions.get(function_id) orelse return Error.InvalidFunction;
-    const entry = function.entry_block orelse return error.MissingEntryBlock;
+    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 block = module.blocks.get(source) orelse return Error.InvalidBlock;
-        const terminator = block.terminator orelse return error.MissingTerminator;
+        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 block = module.blocks.get(queue.items[cursor]) orelse return Error.InvalidBlock;
-        const terminator = block.terminator orelse return error.MissingTerminator;
+        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,
+    MissingTerminator,
-    InvalidCompositeIndex,
+    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,
+        std.mem.Allocator.Error.OutOfMemory => return Error.OutOfMemory,
-        else => return error.InvalidBlock,
+        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,
+    CrossFunctionReference,
-    InvalidEntryPoint,
+    DefinitionDoesNotDominateUse,
-    InvalidType,
+    EntryBlockHasPredecessor,
-    InvalidConstant,
-    InvalidValue,
-    InvalidFunction,
     InvalidBlock,
+    InvalidConstant,
+    InvalidEntryPoint,
+    InvalidFunction,
     InvalidInstruction,
+    InvalidStructuredControl,
+    InvalidType,
+    InvalidValue,
+    MissingEntryPoint,
     MissingFunctionEntryBlock,
     MissingTerminator,
-    EntryBlockHasPredecessor,
-    WrongParent,
-    WrongDefinition,
-    WrongParameterIndex,
-    WrongResultPresence,
-    WrongOperandType,
-    WrongResultType,
     WrongBranchArgumentCount,
     WrongBranchArgumentType,
-    CrossFunctionReference,
+    WrongDefinition,
-    WrongReturnType,
     WrongInterfaceDirection,
-    InvalidStructuredControl,
+    WrongOperandType,
-    DefinitionDoesNotDominateUse,
+    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 = function.entry_block orelse return Error.MissingFunctionEntryBlock;
-        const entry_block_value = module.blocks.get(entry_block) orelse return error.InvalidBlock;
+        const entry_block_value = module.blocks.get(entry_block) orelse return Error.InvalidBlock;
-        if (entry_block_value.parent_function != function_id) return error.WrongParent;
+        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 result = result_type orelse return ValidationError.WrongResultPresence;
-            const ty = module.types.get(result) orelse return error.InvalidType;
+            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 value = module.values.get(value_id) orelse return ValidationError.InvalidValue;
-    const owner = valueFunction(module, value_id) catch return error.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 instruction = module.instructions.get(instruction_id) orelse return ValidationError.InvalidInstruction;
-            const block = module.blocks.get(instruction.parent_block) orelse return error.InvalidBlock;
+            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
+//! Together they form the first stage of the compiler pipeline: SPIR-V binary
-//! the first stage of the compiler pipeline: SPIR-V binary modules are decoded,
+//! modules are decoded, translated into a smaller and easier-to-transform IR,
-//! translated into a smaller and easier-to-transform IR, validated, and then made
+//! validated, and then made available to later optimization or code-generation
-//! available to later optimization or code-generation passes.
+//! transformers.
-
-const std = @import("std");
 
 pub const ir = @import("ir/ir.zig");
 pub const spirv = @import("spirv/root.zig");
 
-const VisitorStatistics = struct {
+test {
-    functions: usize = 0,
+    _ = ir;
-    blocks: usize = 0,
+    _ = spirv;
-};
-
-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;
 }

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/ir/id.zig

@@ -1,12 +1,10 @@
 const shared_ids = @import("shader_ir").ir.id;
 
-pub const FunctionTag = opaque {};
 pub const BlockTag = opaque {};
 pub const InstructionTag = opaque {};
 pub const VirtualRegisterTag = opaque {};
 pub const VirtualFlagTag = opaque {};
 
-pub const FunctionId = shared_ids.Id(FunctionTag);
 pub const BlockId = shared_ids.Id(BlockTag);
 pub const InstructionId = shared_ids.Id(InstructionTag);
 pub const VirtualRegisterId = shared_ids.Id(VirtualRegisterTag);

src/intel/compiler/ir/instruction.zig

@@ -1,5 +1,5 @@
 const std = @import("std");
-const device = @import("device.zig");
+const device = @import("../device.zig");
 const ids = @import("id.zig");
 const operand = @import("operand.zig");
 
@@ -91,19 +91,12 @@ pub const Send = struct {
     response: ?operand.RegisterSpan = null,
 };
 
-pub const Call = struct {
-    function: ids.FunctionId,
-    destination: ?operand.Destination = null,
-    arguments: []const operand.Source,
-};
-
 pub const Operation = union(enum) {
     load_input: LoadInput,
     store_output: StoreOutput,
     move: Move,
     binary: Binary,
     compare: Compare,
-    call: Call,
     send: Send,
 };
 
@@ -121,8 +114,6 @@ pub const Terminator = union(enum) {
         true_block: ids.BlockId,
         false_block: ids.BlockId,
     },
-    return_void,
-    return_value: operand.Source,
     end_thread,
     @"unreachable",
 };
@@ -139,7 +130,6 @@ pub const StructuredControl = union(enum) {
 };
 
 pub const Block = struct {
-    parent_function: ids.FunctionId,
     instructions: std.ArrayList(ids.InstructionId) = .empty,
     terminator: ?Terminator = null,
     structured_control: StructuredControl = .none,

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

@@ -1,4 +1,4 @@
-const device = @import("device.zig");
+const device = @import("../device.zig");
 const ids = @import("id.zig");
 
 pub const DataType = enum {

src/intel/compiler/ir/printer.zig

@@ -1,4 +1,5 @@
 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");
@@ -11,25 +12,18 @@ pub fn write(program: *const program_ir.Program, writer: *std.Io.Writer) std.Io.
     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", .{program.dispatch_width});
+    try writer.print(";   .dispatch_width: {t}\n\n", .{program.dispatch_width});
-    if (program.entry_function) |entry| {
-        try writer.writeAll(";   .entry: ");
-        try writeFunctionRef(program, writer, entry);
-        try writer.writeByte('\n');
-    }
-
-    try writer.writeByte('\n');
 
     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}), spillable({})\n", .{
+        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,
-            register.spillable,
+            if (register.spillable) ", spillable" else "",
         });
     }
 
@@ -39,69 +33,45 @@ pub fn write(program: *const program_ir.Program, writer: *std.Io.Writer) std.Io.
         try writer.writeAll(": vflag\n");
     }
 
-    for (program.functions.entries.items, 0..) |entry, function_index| {
+    try writer.writeByte('\n');
-        const function = entry orelse continue;
 
-        try writer.writeAll("\nfn ");
+    for (program.blocks.entries.items, 0..) |entry, block_index| {
-        try writeFunctionRef(program, writer, ids.FunctionId.fromIndex(function_index));
+        const block = entry orelse continue;
-        try writer.writeByte('(');
+        const block_id = ids.BlockId.fromIndex(block_index);
-        for (function.parameters.items, 0..) |parameter, index| {
-            if (index != 0)
-                try writer.writeAll(", ");
-            try writeVirtualRegisterRef(program, writer, parameter);
-            try writer.writeAll(": ");
-            if (program.virtual_registers.get(parameter)) |register|
-                try writer.print("{t}", .{register.element_type})
-            else
-                try writer.print("<invalid-vgrf-{d}>", .{parameter.index()});
-        }
-        try writer.writeAll(") -> ");
-        if (function.return_type) |return_type|
-            try writer.print("{t}", .{return_type})
-        else
-            try writer.writeAll("void");
-        try writer.writeAll("\n{\n");
 
-        for (function.blocks.items) |block_id| {
+        try writeBlockRef(program, writer, block_id);
-            const block = program.blocks.get(block_id) orelse continue;
+        try writer.writeAll(":\n");
 
-            try writer.writeAll(indent);
+        switch (block.structured_control) {
-            try writeBlockRef(program, writer, block_id);
+            .none => {},
-            try writer.writeAll(":\n");
+            .selection => |selection| {
-
+                try writer.writeAll(indent ++ "structured_selection ");
-            switch (block.structured_control) {
+                try writeBlockRef(program, writer, selection.merge_block);
-                .none => {},
-                .selection => |selection| {
-                    try writer.writeAll(indent ** 2 ++ "structured_selection ");
-                    try writeBlockRef(program, writer, selection.merge_block);
-                    try writer.writeByte('\n');
-                },
-                .loop => |loop| {
-                    try writer.writeAll(indent ** 2 ++ "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 ** 2);
-                try writeInstruction(program, writer, instruction.*);
                 try writer.writeByte('\n');
-            }
+            },
-
+            .loop => |loop| {
-            if (block.terminator) |terminator| {
+                try writer.writeAll(indent ++ "structured_loop merge(");
-                try writer.writeAll(indent ** 2);
+                try writeBlockRef(program, writer, loop.merge_block);
-                try writeTerminator(program, writer, terminator);
+                try writer.writeAll("), continue(");
-                try writer.writeAll("\n\n");
+                try writeBlockRef(program, writer, loop.continue_block);
-            } else {
+                try writer.writeAll(")\n");
-                try writer.writeAll(indent ** 2 ++ "<missing terminator>\n\n");
+            },
-            }
         }
 
-        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");
+        }
     }
 }
 
@@ -118,62 +88,51 @@ fn writeInstruction(program: *const program_ir.Program, writer: *std.Io.Writer,
         try writePredicate(program, writer, predicate);
         try writer.writeByte(' ');
     }
-    try writeOperation(program, writer, instruction.operation);
+    try writeOperation(program, writer, instruction.execution_size, instruction.operation);
 }
 
-fn writeOperation(program: *const program_ir.Program, writer: *std.Io.Writer, operation: inst_ir.Operation) !void {
+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 writeDestination(program, writer, op.destination);
+            try writer.writeAll("load_input ");
-            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, op.source);
+            try writeSource(program, writer, execution_size, op.source);
         },
         .move => |op| {
-            try writeDestination(program, writer, op.destination);
+            try writer.writeAll("mov ");
-            try writer.writeAll(" = move ");
+            try writeDestination(program, writer, execution_size, op.destination);
-            try writeSource(program, writer, op.source);
+            try writer.writeAll(", ");
+            try writeSource(program, writer, execution_size, op.source);
         },
         .binary => |op| {
-            try writeDestination(program, writer, op.destination);
+            try writer.print("{t} ", .{op.opcode});
-            try writer.print(" = {t} ", .{op.opcode});
+            try writeDestination(program, writer, execution_size, op.destination);
-            try writeSource(program, writer, op.lhs);
             try writer.writeAll(", ");
-            try writeSource(program, writer, op.rhs);
+            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.print(" = cmp_{t} ", .{op.opcode});
-            try writeSource(program, writer, op.lhs);
             try writer.writeAll(", ");
-            try writeSource(program, writer, op.rhs);
+            try writeSource(program, writer, execution_size, op.lhs);
-        },
+            try writer.writeAll(", ");
-        .call => |op| {
+            try writeSource(program, writer, execution_size, op.rhs);
-            if (op.destination) |destination| {
-                try writeDestination(program, writer, destination);
-                try writer.writeAll(" = ");
-            }
-            try writer.writeAll("call ");
-            try writeFunctionRef(program, writer, op.function);
-            try writer.writeByte('(');
-            for (op.arguments, 0..) |argument, index| {
-                if (index != 0)
-                    try writer.writeAll(", ");
-                try writeSource(program, writer, argument);
-            }
-            try writer.writeByte(')');
         },
         .send => |op| {
+            try writer.writeAll("send ");
             if (op.response) |response| {
                 try writeRegisterSpan(program, writer, response);
-                try writer.writeAll(" = ");
+                try writer.writeAll(", ");
             }
-            try writer.writeAll("send ");
             try writeMessage(writer, op.message);
             try writer.writeAll(", payload(");
             try writeRegisterSpan(program, writer, op.payload);
@@ -196,63 +155,94 @@ fn writeTerminator(program: *const program_ir.Program, writer: *std.Io.Writer, t
             try writer.writeAll(", ");
             try writeBlockRef(program, writer, branch.false_block);
         },
-        .return_void => try writer.writeAll("return"),
-        .return_value => |value| {
-            try writer.writeAll("return ");
-            try writeSource(program, writer, value);
-        },
         .end_thread => try writer.writeAll("end_thread"),
         .@"unreachable" => try writer.writeAll("unreachable"),
     }
 }
 
-fn writeSource(program: *const program_ir.Program, writer: *std.Io.Writer, source: operand.Source) !void {
+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(");
 
-    switch (source.register) {
+    try writeRegister(program, writer, source.register);
-        .immediate => |immediate| try writeImmediate(writer, immediate),
-        else => {
-            try writeRegisterAtOffset(program, writer, source.register, source.region.byte_offset);
-            try writer.print("<{d};{d},{d}>", .{
-                source.region.vertical_stride,
-                source.region.width,
-                source.region.horizontal_stride,
-            });
-        },
-    }
     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, destination: operand.Destination) !void {
+fn writeDestination(program: *const program_ir.Program, writer: *std.Io.Writer, execution_size: device.ExecutionSize, destination: operand.Destination) !void {
-    try writeRegisterAtOffset(program, writer, destination.register, destination.region.byte_offset);
+    _ = execution_size;
-    try writer.print("<{d}>:{t}", .{ destination.region.horizontal_stride, destination.type });
+    try writeRegister(program, writer, destination.register);
+    try writer.print(":{t}", .{destination.type});
+    try writeDestinationRegion(writer, destination.register, destination.region);
 }
 
-fn writeRegisterAtOffset(
+fn writeSourceRegion(writer: *std.Io.Writer, execution_size: device.ExecutionSize, register: operand.RegisterRef, region: operand.Region) !void {
-    program: *const program_ir.Program,
+    const byte_offset = registerByteOffset(register) + region.byte_offset;
-    writer: *std.Io.Writer,
+    const execution_width: u8 = @intFromEnum(execution_size);
-    register: operand.RegisterRef,
+    const is_default = region.vertical_stride == execution_width and
-    byte_offset: u16,
+        region.width == execution_width and
-) !void {
+        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| {
+        .virtual => |virtual| try writeVirtualRegisterRef(program, writer, virtual),
-            try writeVirtualRegisterRef(program, writer, virtual);
+        .physical_grf => |physical| try writer.print("r{d}", .{physical.number}),
-            try writer.print(".{d}", .{byte_offset});
+        .architecture => |architecture| try writeArchitectureRegister(writer, architecture),
-        },
-        .physical_grf => |physical| try writer.print("r{d}.{d}", .{
-            physical.number,
-            @as(u16, physical.byte_offset) + byte_offset,
-        }),
-        .architecture => |architecture| {
-            try writeArchitectureRegister(writer, architecture);
-            try writer.print(".{d}", .{byte_offset});
-        },
         .immediate => |immediate| try writeImmediate(writer, immediate),
         .null => try writer.writeAll("null"),
     }
@@ -290,7 +280,10 @@ fn writeFlagRef(program: *const program_ir.Program, writer: *std.Io.Writer, flag
 }
 
 fn writeRegisterSpan(program: *const program_ir.Program, writer: *std.Io.Writer, span: operand.RegisterSpan) !void {
-    try writeRegisterAtOffset(program, writer, span.base, 0);
+    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});
 }
 
@@ -336,11 +329,6 @@ fn writeBlockRef(program: *const program_ir.Program, writer: *std.Io.Writer, blo
     try writeNamedRef(writer, if (block) |value| value.name else null, "b", block_id.index(), '.');
 }
 
-fn writeFunctionRef(program: *const program_ir.Program, writer: *std.Io.Writer, function_id: ids.FunctionId) !void {
-    const function = program.functions.get(function_id);
-    try writeNamedRef(writer, if (function) |value| value.name else null, "fn", function_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| {

src/intel/compiler/ir/program.zig

@@ -1,6 +1,6 @@
 const std = @import("std");
 const shared_ir = @import("shader_ir").ir.module;
-const device = @import("device.zig");
+const device = @import("../device.zig");
 const ids = @import("id.zig");
 const instructions = @import("instruction.zig");
 const operand = @import("operand.zig");
@@ -10,7 +10,6 @@ pub const Stage = shared_ir.Stage;
 pub const Properties = packed struct {
     instructions_selected: bool = false,
     block_parameters_lowered: bool = false,
-    calls_lowered: bool = false,
 
     stage_io_lowered: bool = false,
     resources_lowered: bool = false,
@@ -24,7 +23,7 @@ pub const Properties = packed struct {
     flags_allocated: bool = false,
     branches_resolved: bool = false,
 
-    _padding: u20 = 0,
+    _padding: u21 = 0,
 };
 
 pub const VertexPayload = struct {
@@ -43,15 +42,6 @@ pub const ProgramData = struct {
     scratch_size_bytes: u32 = 0,
 };
 
-pub const Function = struct {
-    return_type: ?operand.DataType,
-    parameters: std.ArrayList(ids.VirtualRegisterId) = .empty,
-    blocks: std.ArrayList(ids.BlockId) = .empty,
-    entry_block: ?ids.BlockId = null,
-    name: ?[]const u8 = null,
-};
-
-pub const FunctionStore = ids.Store(ids.FunctionId, Function);
 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);
@@ -64,9 +54,8 @@ pub const Program = struct {
     device_info: device.DeviceInfo,
     dispatch_width: device.DispatchWidth,
 
-    entry_function: ?ids.FunctionId = null,
+    entry_block: ?ids.BlockId = null,
 
-    functions: FunctionStore = .{},
     blocks: BlockStore = .{},
     instructions: InstructionStore = .{},
     virtual_registers: VirtualRegisterStore = .{},
@@ -76,12 +65,7 @@ pub const Program = struct {
     program_data: ProgramData = .{},
     properties: Properties = .{},
 
-    pub fn init(
+    pub fn init(backing_allocator: std.mem.Allocator, stage: Stage, device_info: device.DeviceInfo, dispatch_width: device.DispatchWidth) Program {
-        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,
@@ -99,27 +83,6 @@ pub const Program = struct {
         return self.arena.allocator();
     }
 
-    pub fn addFunction(self: *Program, return_type: ?operand.DataType, name: ?[]const u8) !ids.FunctionId {
-        const owned_name = if (name) |value| try self.allocator().dupe(u8, value) else null;
-        return self.functions.add(self.allocator(), .{
-            .return_type = return_type,
-            .name = owned_name,
-        });
-    }
-
-    pub fn setEntryFunction(self: *Program, function_id: ids.FunctionId) !void {
-        if (!self.functions.isLive(function_id))
-            return error.InvalidFunction;
-        self.entry_function = function_id;
-    }
-
-    pub fn addFunctionParameter(self: *Program, function_id: ids.FunctionId, register_id: ids.VirtualRegisterId) !void {
-        const function = self.functions.getMut(function_id) orelse return error.InvalidFunction;
-        if (!self.virtual_registers.isLive(register_id))
-            return error.InvalidVirtualRegister;
-        try function.parameters.append(self.allocator(), register_id);
-    }
-
     pub fn addVirtualRegister(self: *Program, register: operand.VirtualRegister) !ids.VirtualRegisterId {
         var owned = register;
         if (register.name) |name|
@@ -134,40 +97,30 @@ pub const Program = struct {
         return self.virtual_flags.add(self.allocator(), owned);
     }
 
-    pub fn addBlock(self: *Program, function_id: ids.FunctionId, name: ?[]const u8) !ids.BlockId {
+    pub fn addBlock(self: *Program, name: ?[]const u8) !ids.BlockId {
-        const function = self.functions.getMut(function_id) orelse return error.InvalidFunction;
         const owned_name = if (name) |value| try self.allocator().dupe(u8, value) else null;
         const block_id = try self.blocks.add(self.allocator(), .{
-            .parent_function = function_id,
             .name = owned_name,
         });
-        errdefer _ = self.blocks.remove(block_id);
+        if (self.entry_block == null)
-
+            self.entry_block = block_id;
-        try function.blocks.append(self.allocator(), block_id);
-        if (function.entry_block == null)
-            function.entry_block = block_id;
         return block_id;
     }
 
-    pub fn appendInstruction(
+    pub fn setEntryBlock(self: *Program, block_id: ids.BlockId) !void {
-        self: *Program,
+        if (!self.blocks.isLive(block_id))
-        block_id: ids.BlockId,
+            return error.InvalidBlock;
-        execution_size: device.ExecutionSize,
+        self.entry_block = block_id;
-        predicate: ?operand.Predicate,
+    }
-        operation: instructions.Operation,
+
-    ) !ids.InstructionId {
+    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;
-        var owned_operation = operation;
-        switch (owned_operation) {
-            .call => |*call| call.arguments = try self.allocator().dupe(operand.Source, call.arguments),
-            else => {},
-        }
 
         const instruction_id = try self.instructions.add(self.allocator(), .{
             .parent_block = block_id,
             .execution_size = execution_size,
             .predicate = predicate,
-            .operation = owned_operation,
+            .operation = operation,
         });
         try block.instructions.append(self.allocator(), instruction_id);
         return instruction_id;

src/intel/compiler/ir/validator.zig

@@ -9,12 +9,8 @@ pub const Error = error{
     UnsupportedDispatchWidth,
     UnsupportedExecutionSize,
     UnsupportedDataType,
-    MissingEntryFunction,
-    InvalidFunction,
     MissingEntryBlock,
     InvalidBlock,
-    WrongParentFunction,
-    CrossFunctionBranch,
     MissingTerminator,
     InvalidInstruction,
     InvalidVirtualRegister,
@@ -28,10 +24,6 @@ pub const Error = error{
     InvalidDestination,
     InvalidImmediateType,
     InvalidRegisterSpan,
-    WrongArgumentCount,
-    WrongArgumentType,
-    WrongReturnType,
-    InvalidEntryTerminator,
 };
 
 pub fn validate(program: *const program_ir.Program) Error!void {
@@ -42,35 +34,9 @@ pub fn validate(program: *const program_ir.Program) Error!void {
     if (program.dispatch_width != .simd8 or !program.device_info.supportsDispatch(.simd8))
         return error.UnsupportedDispatchWidth;
 
-    const entry_function_id = program.entry_function orelse return error.MissingEntryFunction;
+    const entry_block = program.entry_block orelse return error.MissingEntryBlock;
-    const entry_function = program.functions.get(entry_function_id) orelse return error.InvalidFunction;
+    if (!program.blocks.isLive(entry_block))
-    const entry_block = entry_function.entry_block orelse return error.MissingEntryBlock;
+        return error.InvalidBlock;
-    const entry_block_data = program.blocks.get(entry_block) orelse return error.InvalidBlock;
-    if (entry_block_data.parent_function != entry_function_id)
-        return error.WrongParentFunction;
-
-    for (program.functions.entries.items, 0..) |entry, function_index| {
-        const function = entry orelse continue;
-        const function_id = ids.FunctionId.fromIndex(function_index);
-
-        if (function.return_type) |return_type|
-            try validateType(return_type);
-        const function_entry = function.entry_block orelse return error.MissingEntryBlock;
-        const function_entry_data = program.blocks.get(function_entry) orelse return error.InvalidBlock;
-        if (function_entry_data.parent_function != function_id)
-            return error.WrongParentFunction;
-
-        for (function.parameters.items) |parameter| {
-            if (!program.virtual_registers.isLive(parameter))
-                return error.InvalidVirtualRegister;
-        }
-
-        for (function.blocks.items) |block_id| {
-            const block = program.blocks.get(block_id) orelse return error.InvalidBlock;
-            if (block.parent_function != function_id)
-                return error.WrongParentFunction;
-        }
-    }
 
     for (program.virtual_registers.entries.items) |entry| {
         const register = entry orelse continue;
@@ -86,8 +52,6 @@ pub fn validate(program: *const program_ir.Program) Error!void {
 
     for (program.blocks.entries.items, 0..) |entry, block_index| {
         const block = entry orelse continue;
-        if (!program.functions.isLive(block.parent_function) or !functionContainsBlock(program, block.parent_function, ids.BlockId.fromIndex(block_index)))
-            return error.WrongParentFunction;
         if (block.terminator == null)
             return error.MissingTerminator;
 
@@ -99,8 +63,8 @@ pub fn validate(program: *const program_ir.Program) Error!void {
             try validateInstruction(program, inst.*);
         }
 
-        try validateStructuredControl(program, block.parent_function, block.structured_control);
+        try validateStructuredControl(program, block.structured_control);
-        try validateTerminator(program, block.parent_function, block.terminator.?);
+        try validateTerminator(program, block.terminator.?);
     }
 }
 
@@ -130,7 +94,6 @@ fn validateInstruction(program: *const program_ir.Program, inst: instruction.Ins
             try validateSource(program, op.lhs);
             try validateSource(program, op.rhs);
         },
-        .call => |op| try validateCall(program, op),
         .send => |op| {
             try validateSpan(program, op.payload);
             if (op.response) |response|
@@ -139,28 +102,6 @@ fn validateInstruction(program: *const program_ir.Program, inst: instruction.Ins
     }
 }
 
-fn validateCall(program: *const program_ir.Program, call: instruction.Call) Error!void {
-    const function = program.functions.get(call.function) orelse return error.InvalidFunction;
-    if (call.arguments.len != function.parameters.items.len)
-        return error.WrongArgumentCount;
-
-    for (call.arguments, function.parameters.items) |argument, parameter_id| {
-        try validateSource(program, argument);
-        const parameter = program.virtual_registers.get(parameter_id) orelse return error.InvalidVirtualRegister;
-        if (argument.type != parameter.element_type)
-            return error.WrongArgumentType;
-    }
-
-    if (function.return_type) |return_type| {
-        const destination = call.destination orelse return error.WrongReturnType;
-        try validateDestination(program, destination);
-        if (destination.type != return_type)
-            return error.WrongReturnType;
-    } else if (call.destination != null) {
-        return error.WrongReturnType;
-    }
-}
-
 fn validateType(data_type: operand.DataType) Error!void {
     if (!data_type.isInitialTargetType())
         return error.UnsupportedDataType;
@@ -224,59 +165,30 @@ fn validateSpan(program: *const program_ir.Program, span: operand.RegisterSpan)
     }
 }
 
-fn validateTerminator(
+fn validateTerminator(program: *const program_ir.Program, terminator: instruction.Terminator) Error!void {
-    program: *const program_ir.Program,
-    function_id: ids.FunctionId,
-    terminator: instruction.Terminator,
-) Error!void {
-    const function = program.functions.get(function_id) orelse return error.InvalidFunction;
     switch (terminator) {
-        .jump => |target| try validateBlockTarget(program, function_id, target),
+        .jump => |target| try validateBlockTarget(program, target),
         .conditional_branch => |branch| {
             try validateFlag(program, branch.predicate.flag);
-            try validateBlockTarget(program, function_id, branch.true_block);
+            try validateBlockTarget(program, branch.true_block);
-            try validateBlockTarget(program, function_id, branch.false_block);
+            try validateBlockTarget(program, branch.false_block);
         },
-        .return_void => if (function.return_type != null)
+        .end_thread, .@"unreachable" => {},
-            return error.WrongReturnType,
-        .return_value => |value| {
-            const return_type = function.return_type orelse return error.WrongReturnType;
-            try validateSource(program, value);
-            if (value.type != return_type)
-                return error.WrongReturnType;
-        },
-        .end_thread => if (program.entry_function != function_id)
-            return error.InvalidEntryTerminator,
-        .@"unreachable" => {},
     }
 }
 
-fn validateStructuredControl(
+fn validateStructuredControl(program: *const program_ir.Program, control: instruction.StructuredControl) Error!void {
-    program: *const program_ir.Program,
-    function_id: ids.FunctionId,
-    control: instruction.StructuredControl,
-) Error!void {
     switch (control) {
         .none => {},
-        .selection => |selection| try validateBlockTarget(program, function_id, selection.merge_block),
+        .selection => |selection| try validateBlockTarget(program, selection.merge_block),
         .loop => |loop| {
-            try validateBlockTarget(program, function_id, loop.merge_block);
+            try validateBlockTarget(program, loop.merge_block);
-            try validateBlockTarget(program, function_id, loop.continue_block);
+            try validateBlockTarget(program, loop.continue_block);
         },
     }
 }
 
-fn validateBlockTarget(program: *const program_ir.Program, function_id: ids.FunctionId, block_id: ids.BlockId) Error!void {
+fn validateBlockTarget(program: *const program_ir.Program, block_id: ids.BlockId) Error!void {
-    const block = program.blocks.get(block_id) orelse return error.InvalidBlock;
+    if (!program.blocks.isLive(block_id))
-    if (block.parent_function != function_id)
+        return error.InvalidBlock;
-        return error.CrossFunctionBranch;
-}
-
-fn functionContainsBlock(program: *const program_ir.Program, function_id: ids.FunctionId, block_id: ids.BlockId) bool {
-    const function = program.functions.get(function_id) orelse return false;
-    for (function.blocks.items) |candidate| {
-        if (candidate == block_id)
-            return true;
-    }
-    return false;
 }

src/intel/compiler/lower.zig

@@ -1,74 +0,0 @@
-const std = @import("std");
-const shader_ir = @import("shader_ir").ir;
-const device = @import("device.zig");
-const program_ir = @import("program.zig");
-
-pub const Options = struct {
-    dispatch_width: device.DispatchWidth = .simd8,
-};
-
-pub const Error = std.mem.Allocator.Error || error{
-    MissingEntryPoint,
-    InvalidEntryPoint,
-    UnsupportedGeneration,
-    UnsupportedStage,
-    UnsupportedDispatchWidth,
-    UnsupportedType,
-    UnsupportedOperation,
-    UnsupportedTerminator,
-    LoweringNotImplemented,
-};
-
-pub const Lowerer = struct {
-    allocator: std.mem.Allocator,
-    module: *const shader_ir.module.Module,
-    device_info: device.DeviceInfo,
-    options: Options,
-
-    pub fn init(
-        allocator: std.mem.Allocator,
-        module: *const shader_ir.module.Module,
-        device_info: device.DeviceInfo,
-        options: Options,
-    ) Lowerer {
-        return .{
-            .allocator = allocator,
-            .module = module,
-            .device_info = device_info,
-            .options = options,
-        };
-    }
-
-    pub fn lower(self: *Lowerer) Error!program_ir.Program {
-        var program = program_ir.Program.init(self.allocator, self.module.stage, self.device_info, self.options.dispatch_width);
-        errdefer program.deinit();
-
-        const entry_point = self.source.entryPoint(self.module.stage);
-        if (entry_point) |ep| {
-            program.entry_block = try program.addBlock(ep.name);
-            _ = program.appendInstruction(program.entry_block orelse return error.InvalidBlock, self.options.dispatch_width, null, .{ .name = ep.name });
-        }
-
-        const blocks = self.source.blocks();
-        for (blocks) |block| {
-            const block_id = try program.addBlock(block.name);
-            for (block.instructions) |inst| {
-                _ = program.appendInstruction(block_id, self.options.dispatch_width, null, .{ .name = inst.name });
-            }
-            _ = program.setTerminator(block_id, .{ .name = block.terminator.name });
-        }
-
-        return program;
-    }
-};
-
-/// Convenience entry point for callers that do not need to retain a lowerer.
-pub inline fn lower(
-    allocator: std.mem.Allocator,
-    module: *const shader_ir.module.Module,
-    device_info: device.DeviceInfo,
-    options: Options,
-) Error!program_ir.Program {
-    var lowerer = Lowerer.init(allocator, module, device_info, options);
-    return lowerer.lower();
-}

src/intel/compiler/root.zig

@@ -1,266 +0,0 @@
-//! Backend-specific shader IR for Intel Gen hardware.
-//! This is the mutable, non-SSA layer between the shared shader IR and machine  encoding.
-
-pub const device = @import("device.zig");
-pub const id = @import("id.zig");
-pub const instruction = @import("instruction.zig");
-pub const lower = @import("lower.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 Function = program.Function;
-pub const FunctionId = id.FunctionId;
-pub const Program = program.Program;
-pub const Stage = program.Stage;
-
-const std = @import("std");
-
-test "Flint IR foundation" {
-    // ; Flint program:
-    // ;   .stage: vertex
-    // ;   .generation: gen9
-    // ;   .platform: skylake
-    // ;   .dispatch_width: simd8
-    // ;   .entry: @main
-    //
-    // %position: vgrf f32[8] = class(varying), size(32), alignment(32), spillable(true)
-    // %urb_payload: vgrf u32[16] = class(payload), size(64), alignment(32), spillable(false)
-    //
-    // fn @main() -> void
-    // {
-    //     .entry:
-    //         [simd8] %position.0<1>:f32 = load_input location(0), component(0)
-    //         [simd8] %position.0<1>:f32 = multiply %position.0<8;8,1>:f32, 1:f32
-    //         [simd8] store_output builtin(position), component(0), %position.0<8;8,1>:f32
-    //         [simd8] send urb_write[offset(0), channels(xyzw), end_of_thread], payload(%urb_payload.0[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 main = try shader.addFunction(null, "main");
-    try shader.setEntryFunction(main);
-
-    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(main, "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, .{
-        .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(main, shader.entry_function.?);
-    try std.testing.expectEqual(entry, shader.functions.get(main).?.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, "@main") != 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] %position.0<1>:f32 = multiply %position.0<8;8,1>:f32, 1:f32") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "send urb_write[offset(0), channels(xyzw), end_of_thread], payload(%urb_payload.0[2])") != null);
-}
-
-test "Flint IR function calls" {
-    // ; Flint program:
-    // ;   .stage: vertex
-    // ;   .generation: gen9
-    // ;   .platform: skylake
-    // ;   .dispatch_width: simd8
-    // ;   .entry: @main
-    //
-    // %value: vgrf u32[8] = class(temporary), size(32), alignment(32), spillable(true)
-    // %result: vgrf u32[8] = class(temporary), size(32), alignment(32), spillable(true)
-    //
-    // fn @main() -> void
-    // {
-    //     .entry:
-    //         [simd8] %result.0<1>:u32 = call @identity(1:u32)
-    //         end_thread
-    //
-    // }
-    //
-    // fn @identity(%value: u32) -> u32
-    // {
-    //     .entry:
-    //         return %value.0<8;8,1>:u32
-    //
-    // }
-
-    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 main = try shader.addFunction(null, "main");
-    const identity = try shader.addFunction(.u32, "identity");
-    try shader.setEntryFunction(main);
-
-    const parameter = try shader.addVirtualRegister(.{
-        .size_bytes = 32,
-        .alignment_bytes = 32,
-        .element_type = .u32,
-        .lane_count = 8,
-        .class = .temporary,
-        .name = "value",
-    });
-    const result = try shader.addVirtualRegister(.{
-        .size_bytes = 32,
-        .alignment_bytes = 32,
-        .element_type = .u32,
-        .lane_count = 8,
-        .class = .temporary,
-        .name = "result",
-    });
-    try shader.addFunctionParameter(identity, parameter);
-
-    const main_entry = try shader.addBlock(main, "entry");
-    const identity_entry = try shader.addBlock(identity, "entry");
-
-    const call_id = try shader.appendInstruction(main_entry, .simd8, null, .{
-        .call = .{
-            .function = identity,
-            .destination = .{
-                .register = .{ .virtual = result },
-                .type = .u32,
-            },
-            .arguments = &.{
-                .{
-                    .register = .{ .immediate = .{ .u32 = 1 } },
-                    .type = .u32,
-                    .region = operand.Region.broadcast(),
-                },
-            },
-        },
-    });
-    try shader.setTerminator(main_entry, .end_thread);
-    try shader.setTerminator(identity_entry, .{ .return_value = .{
-        .register = .{ .virtual = parameter },
-        .type = .u32,
-        .region = operand.Region.contiguous(.simd8),
-    } });
-
-    try validator.validate(&shader);
-
-    const call = shader.instructions.get(call_id).?.operation.call;
-    try std.testing.expectEqual(identity, call.function);
-    try std.testing.expectEqual(@as(usize, 1), call.arguments.len);
-
-    const text = try printer.allocPrint(std.testing.allocator, &shader);
-    defer std.testing.allocator.free(text);
-
-    try std.testing.expect(std.mem.indexOf(u8, text, "fn @identity(%value: u32) -> u32") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "[simd8] %result.0<1>:u32 = call @identity(1:u32)") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "return %value.0<8;8,1>:u32") != 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.?);
-}

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,20 +16,23 @@ 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");
-pub const compiler = @import("compiler/root.zig");
 
 pub const Instance = FlintInstance;