diff --git a/src/compiler/ir/Builder.zig b/src/compiler/ir/Builder.zig index 8183b64..748dde2 100644 --- a/src/compiler/ir/Builder.zig +++ b/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); +} diff --git a/src/compiler/ir/Rewriter.zig b/src/compiler/ir/Rewriter.zig index deffd00..9284733 100644 --- a/src/compiler/ir/Rewriter.zig +++ b/src/compiler/ir/Rewriter.zig @@ -2,6 +2,7 @@ const std = @import("std"); const ids = @import("id.zig"); const module_ir = @import("module.zig"); const Builder = @import("Builder.zig"); +const validator = @import("validator/validator.zig"); const Self = @This(); @@ -52,11 +53,11 @@ pub fn countUses(self: *const Self, value: ids.ValueId) usize { } pub fn replaceAllUses(self: *Self, old: ids.ValueId, replacement: ids.ValueId) Error!usize { - const old_value = self.module.values.get(old) orelse return error.InvalidValue; - const replacement_value = self.module.values.get(replacement) orelse return error.InvalidValue; + const old_value = self.module.values.get(old) orelse return Error.InvalidValue; + const replacement_value = self.module.values.get(replacement) orelse return Error.InvalidValue; if (old_value.type != replacement_value.type) - return error.TypeMismatch; + return Error.TypeMismatch; if (old == replacement) return 0; @@ -77,17 +78,17 @@ pub fn replaceAllUses(self: *Self, old: ids.ValueId, replacement: ids.ValueId) E } pub fn eraseInstruction(self: *Self, instruction_id: ids.InstructionId) Error!void { - const instruction = self.module.instructions.get(instruction_id) orelse return error.InvalidInstruction; + const instruction = self.module.instructions.get(instruction_id) orelse return Error.InvalidInstruction; if (instruction.operation.hasSideEffects()) - return error.SideEffectingInstruction; + return Error.SideEffectingInstruction; if (instruction.result) |result| { if (self.countUses(result) != 0) - return error.ResultStillUsed; + return Error.ResultStillUsed; } - const block = self.module.blocks.getMut(instruction.parent_block) orelse return error.InvalidBlock; + const block = self.module.blocks.getMut(instruction.parent_block) orelse return Error.InvalidBlock; var owned_index: ?usize = null; for (block.instructions.items, 0..) |candidate, index| { @@ -97,29 +98,23 @@ pub fn eraseInstruction(self: *Self, instruction_id: ids.InstructionId) Error!vo } } - _ = block.instructions.orderedRemove(owned_index orelse return error.InstructionNotOwnedByBlock); + _ = block.instructions.orderedRemove(owned_index orelse return Error.InstructionNotOwnedByBlock); if (instruction.result) |result| _ = self.module.values.remove(result); _ = self.module.instructions.remove(instruction_id); } -pub fn redirectEdges( - self: *Self, - source: ids.BlockId, - old_target: ids.BlockId, - new_target: ids.BlockId, - new_arguments: []const ids.ValueId, -) Error!usize { - const source_block = self.module.blocks.get(source) orelse return error.InvalidBlock; - const target_block = self.module.blocks.get(new_target) orelse return error.InvalidBlock; +pub fn redirectEdges(self: *Self, source: ids.BlockId, old_target: ids.BlockId, new_target: ids.BlockId, new_arguments: []const ids.ValueId) Error!usize { + const source_block = self.module.blocks.get(source) orelse return Error.InvalidBlock; + const target_block = self.module.blocks.get(new_target) orelse return Error.InvalidBlock; if (source_block.parent_function != target_block.parent_function) - return error.InvalidFunction; + return Error.InvalidFunction; try self.validateArguments(target_block, new_arguments); const mutable_source = self.module.blocks.getMut(source).?; - const terminator = if (mutable_source.terminator) |*value| value else return error.InvalidBlock; + const terminator = if (mutable_source.terminator) |*value| value else return Error.InvalidBlock; var count: usize = 0; @@ -141,29 +136,23 @@ pub fn redirectEdges( return count; } -pub fn addBlockParameter( - self: *Self, - block_id: ids.BlockId, - ty: ids.TypeId, - name: ?[]const u8, - incoming: []const IncomingValue, -) Error!ids.ValueId { - const block = self.module.blocks.get(block_id) orelse return error.InvalidBlock; - const function = self.module.functions.get(block.parent_function) orelse return error.InvalidFunction; +pub fn addBlockParameter(self: *Self, block_id: ids.BlockId, ty: ids.TypeId, name: ?[]const u8, incoming: []const IncomingValue) Error!ids.ValueId { + const block = self.module.blocks.get(block_id) orelse return Error.InvalidBlock; + const function = self.module.functions.get(block.parent_function) orelse return Error.InvalidFunction; for (incoming) |item| { - const value = self.module.values.get(item.value) orelse return error.InvalidValue; + const value = self.module.values.get(item.value) orelse return Error.InvalidValue; if (value.type != ty) - return error.TypeMismatch; + return Error.TypeMismatch; if (!functionHasEdgeTo(self.module, function, item.predecessor, block_id)) - return error.UnexpectedIncomingValue; + return Error.UnexpectedIncomingValue; } for (function.blocks.items) |predecessor| { const edge_count = countEdgesTo(self.module.blocks.get(predecessor).?, block_id); if (edge_count != 0 and findIncoming(incoming, predecessor) == null) - return error.MissingIncomingValue; + return Error.MissingIncomingValue; } var builder = Builder.init(self.module); @@ -177,19 +166,14 @@ pub fn addBlockParameter( return parameter; } -pub fn removeBlockParameter( - self: *Self, - block_id: ids.BlockId, - parameter_index: usize, - replacement: ids.ValueId, -) Error!void { - const block = self.module.blocks.get(block_id) orelse return error.InvalidBlock; - if (parameter_index >= block.parameters.items.len) return error.InvalidParameterIndex; +pub fn removeBlockParameter(self: *Self, block_id: ids.BlockId, parameter_index: usize, replacement: ids.ValueId) Error!void { + const block = self.module.blocks.get(block_id) orelse return Error.InvalidBlock; + if (parameter_index >= block.parameters.items.len) return Error.InvalidParameterIndex; const parameter = block.parameters.items[parameter_index]; - if (parameter == replacement) return error.InvalidValue; + if (parameter == replacement) return Error.InvalidValue; _ = try self.replaceAllUses(parameter, replacement); - const function = self.module.functions.get(block.parent_function) orelse return error.InvalidFunction; + const function = self.module.functions.get(block.parent_function) orelse return Error.InvalidFunction; for (function.blocks.items) |predecessor| { try self.removeArgumentFromEdges(predecessor, block_id, parameter_index); } @@ -198,7 +182,7 @@ pub fn removeBlockParameter( _ = mutable_block.parameters.orderedRemove(parameter_index); for (mutable_block.parameters.items[parameter_index..], parameter_index..) |value_id, index| { - const value = self.module.values.getMut(value_id) orelse return error.InvalidValue; + const value = self.module.values.getMut(value_id) orelse return Error.InvalidValue; value.definition.block_parameter.index = @intCast(index); } @@ -206,21 +190,15 @@ pub fn removeBlockParameter( } fn validateArguments(self: *const Self, target: *const module_ir.Block, arguments: []const ids.ValueId) Error!void { - if (arguments.len != target.parameters.items.len) return error.TypeMismatch; + if (arguments.len != target.parameters.items.len) return Error.TypeMismatch; for (arguments, target.parameters.items) |argument, parameter| { - const argument_value = self.module.values.get(argument) orelse return error.InvalidValue; - const parameter_value = self.module.values.get(parameter) orelse return error.InvalidValue; - if (argument_value.type != parameter_value.type) return error.TypeMismatch; + const argument_value = self.module.values.get(argument) orelse return Error.InvalidValue; + const parameter_value = self.module.values.get(parameter) orelse return Error.InvalidValue; + if (argument_value.type != parameter_value.type) return Error.TypeMismatch; } } -fn redirectOne( - self: *Self, - edge: *module_ir.Edge, - old_target: ids.BlockId, - new_target: ids.BlockId, - arguments: []const ids.ValueId, -) !bool { +fn redirectOne(self: *Self, edge: *module_ir.Edge, old_target: ids.BlockId, new_target: ids.BlockId, arguments: []const ids.ValueId) !bool { if (edge.target != old_target) return false; @@ -230,8 +208,8 @@ fn redirectOne( } fn appendArgumentToEdges(self: *Self, predecessor: ids.BlockId, target: ids.BlockId, value: ids.ValueId) !void { - const block = self.module.blocks.getMut(predecessor) orelse return error.InvalidBlock; - const terminator = if (block.terminator) |*item| item else return error.InvalidBlock; + const block = self.module.blocks.getMut(predecessor) orelse return Error.InvalidBlock; + const terminator = if (block.terminator) |*item| item else return Error.InvalidBlock; switch (terminator.*) { .branch => |*edge| { @@ -257,8 +235,8 @@ fn appendEdgeArgument(self: *Self, edge: *module_ir.Edge, value: ids.ValueId) !v } fn removeArgumentFromEdges(self: *Self, predecessor: ids.BlockId, target: ids.BlockId, index: usize) !void { - const block = self.module.blocks.getMut(predecessor) orelse return error.InvalidBlock; - const terminator = if (block.terminator) |*item| item else return error.InvalidBlock; + const block = self.module.blocks.getMut(predecessor) orelse return Error.InvalidBlock; + const terminator = if (block.terminator) |*item| item else return Error.InvalidBlock; switch (terminator.*) { .branch => |*edge| { @@ -278,7 +256,7 @@ fn removeArgumentFromEdges(self: *Self, predecessor: ids.BlockId, target: ids.Bl fn removeEdgeArgument(self: *Self, edge: *module_ir.Edge, index: usize) !void { if (index >= edge.arguments.len) - return error.InvalidParameterIndex; + return Error.InvalidParameterIndex; const arguments = try self.module.allocator().alloc(ids.ValueId, edge.arguments.len - 1); @memcpy(arguments[0..index], edge.arguments[0..index]); @@ -299,12 +277,7 @@ fn findIncoming(incoming: []const IncomingValue, predecessor: ids.BlockId) ?ids. return null; } -fn functionHasEdgeTo( - module: *const module_ir.Module, - function: *const module_ir.Function, - predecessor: ids.BlockId, - target: ids.BlockId, -) bool { +fn functionHasEdgeTo(module: *const module_ir.Module, function: *const module_ir.Function, predecessor: ids.BlockId, target: ids.BlockId) bool { for (function.blocks.items) |block_id| { if (block_id != predecessor) continue; @@ -321,3 +294,135 @@ fn countEdgesTo(block: *const module_ir.Block, target: ids.BlockId) usize { else => 0, }; } + +test "Rewriter: replace all ID uses, safely erase dead instruction" { + // shader compute @main + // { + // %0: constant u32 = bits(0x1) + // %1: constant u32 = bits(0x2) + // + // fn @main() -> void + // { + // .entry(): + // %2: u32 = integer_add %0, %1 + // %3: u32 = integer_multiply %2, %1 + // return + // } + // } + + var module = module_ir.Module.init(std.testing.allocator, .compute); + defer module.deinit(); + + var builder = Builder.init(&module); + + const void_type = try builder.internType(.void); + const u32_type = try builder.internType(.{ .integer = .{ .bits = 32, .signedness = .unsigned } }); + + const one = try builder.internConstant(u32_type, .{ .integer_bits = 1 }); + const two = try builder.internConstant(u32_type, .{ .integer_bits = 2 }); + + const main = try builder.addFunction(void_type, "main"); + builder.setEntryPoint(main); + + const entry = try builder.addBlock(main, "entry"); + const sum = (try builder.appendInstruction(entry, u32_type, .{ + .binary = .{ + .opcode = .integer_add, + .lhs = one, + .rhs = two, + }, + }, null)).?; + _ = try builder.appendInstruction(entry, u32_type, .{ + .binary = .{ + .opcode = .integer_multiply, + .lhs = sum, + .rhs = two, + }, + }, null); + try builder.setTerminator(entry, .return_void); + + try validator.validate(&module); + + const sum_instruction = module.values.get(sum).?.definition.instruction; + var rewriter = Self.init(&module); + + try std.testing.expectEqual(@as(usize, 1), try rewriter.replaceAllUses(sum, one)); + try rewriter.eraseInstruction(sum_instruction); + + try std.testing.expect(module.values.get(sum) == null); + try std.testing.expect(module.instructions.get(sum_instruction) == null); + + try validator.validate(&module); +} + +test "Rewriter: add block parameter and sync branch calls" { + // shader compute @main + // { + // %0: constant u32 = bits(0x1) + // + // fn @main() -> void + // { + // .entry(): + // branch .merge() + // + // .merge(): + // return + // + // .alternate(): + // return + // } + // } + + var module = module_ir.Module.init(std.testing.allocator, .compute); + defer module.deinit(); + + var builder = Builder.init(&module); + + const void_type = try builder.internType(.void); + const u32_type = try builder.internType(.{ .integer = .{ .bits = 32, .signedness = .unsigned } }); + + const one = try builder.internConstant(u32_type, .{ .integer_bits = 1 }); + + const main = try builder.addFunction(void_type, "main"); + builder.setEntryPoint(main); + + const entry = try builder.addBlock(main, "entry"); + const merge = try builder.addBlock(main, "merge"); + const alternate = try builder.addBlock(main, "alternate"); + + try builder.setTerminator(entry, .{ .branch = try builder.edge(merge, &.{}) }); + try builder.setTerminator(merge, .return_void); + try builder.setTerminator(alternate, .return_void); + + var rewriter = Self.init(&module); + + const parameter = try rewriter.addBlockParameter(merge, u32_type, "incoming", &.{ + .{ + .predecessor = entry, + .value = one, + }, + }); + const merge_edge = module.blocks.get(entry).?.terminator.?.branch; + try std.testing.expectEqualSlices(ids.ValueId, &.{one}, merge_edge.arguments); + + _ = try builder.appendInstruction(merge, u32_type, .{ + .binary = .{ + .opcode = .integer_add, + .lhs = parameter, + .rhs = one, + }, + }, null); + try validator.validate(&module); + + try rewriter.removeBlockParameter(merge, 0, one); + + try std.testing.expectEqual(@as(usize, 0), module.blocks.get(merge).?.parameters.items.len); + try std.testing.expectEqual(@as(usize, 0), module.blocks.get(entry).?.terminator.?.branch.arguments.len); + + try validator.validate(&module); + + try std.testing.expectEqual(@as(usize, 1), try rewriter.redirectEdges(entry, merge, alternate, &.{})); + try std.testing.expectEqual(alternate, module.blocks.get(entry).?.terminator.?.branch.target); + + try validator.validate(&module); +} diff --git a/src/compiler/ir/cfg.zig b/src/compiler/ir/cfg.zig index cd1aadf..d3609c4 100644 --- a/src/compiler/ir/cfg.zig +++ b/src/compiler/ir/cfg.zig @@ -19,8 +19,8 @@ reachable: []bool, dominators: []bool, pub fn init(allocator: std.mem.Allocator, module: *const module_ir.Module, function_id: ids.FunctionId) Error!Self { - const function = module.functions.get(function_id) orelse return error.InvalidFunction; - const entry = function.entry_block orelse return error.MissingEntryBlock; + const function = module.functions.get(function_id) orelse return Error.InvalidFunction; + const entry = function.entry_block orelse return Error.MissingEntryBlock; const blocks = try allocator.dupe(ids.BlockId, function.blocks.items); errdefer allocator.free(blocks); @@ -84,8 +84,8 @@ pub fn dominates(self: *const Self, dominator: ids.BlockId, block: ids.BlockId) fn buildPredecessors(self: *Self, module: *const module_ir.Module) Error!void { for (self.blocks) |source| { - const block = module.blocks.get(source) orelse return error.InvalidBlock; - const terminator = block.terminator orelse return error.MissingTerminator; + const block = module.blocks.get(source) orelse return Error.InvalidBlock; + const terminator = block.terminator orelse return Error.MissingTerminator; switch (terminator) { .branch => |edge| try self.addPredecessor(edge.target, source), .conditional_branch => |branch| { @@ -101,12 +101,12 @@ fn buildReachability(self: *Self, module: *const module_ir.Module, entry: ids.Bl var queue: std.ArrayList(ids.BlockId) = .empty; defer queue.deinit(self.allocator); try queue.append(self.allocator, entry); - self.reachable[self.indexOf(entry) orelse return error.InvalidBlock] = true; + self.reachable[self.indexOf(entry) orelse return Error.InvalidBlock] = true; var cursor: usize = 0; while (cursor < queue.items.len) : (cursor += 1) { - const block = module.blocks.get(queue.items[cursor]) orelse return error.InvalidBlock; - const terminator = block.terminator orelse return error.MissingTerminator; + const block = module.blocks.get(queue.items[cursor]) orelse return Error.InvalidBlock; + const terminator = block.terminator orelse return Error.MissingTerminator; switch (terminator) { .branch => |edge| try self.markReachable(&queue, edge.target), .conditional_branch => |branch| { @@ -168,12 +168,12 @@ fn buildDominators(self: *Self, entry: ids.BlockId) void { } fn addPredecessor(self: *Self, target: ids.BlockId, source: ids.BlockId) Error!void { - const target_index = self.indexOf(target) orelse return error.CrossFunctionEdge; + const target_index = self.indexOf(target) orelse return Error.CrossFunctionEdge; try self.predecessors_by_block[target_index].append(self.allocator, source); } fn markReachable(self: *Self, queue: *std.ArrayList(ids.BlockId), target: ids.BlockId) Error!void { - const target_index = self.indexOf(target) orelse return error.CrossFunctionEdge; + const target_index = self.indexOf(target) orelse return Error.CrossFunctionEdge; if (self.reachable[target_index]) return; diff --git a/src/compiler/ir/ir.zig b/src/compiler/ir/ir.zig index b6f1755..aede46f 100644 --- a/src/compiler/ir/ir.zig +++ b/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; +} diff --git a/src/compiler/ir/module.zig b/src/compiler/ir/module.zig index c77cca3..e384b27 100644 --- a/src/compiler/ir/module.zig +++ b/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); +} diff --git a/src/compiler/ir/parser/parser.zig b/src/compiler/ir/parser/parser.zig index 81b344b..a5aad3c 100644 --- a/src/compiler/ir/parser/parser.zig +++ b/src/compiler/ir/parser/parser.zig @@ -10,23 +10,23 @@ const ast = @import("ast.zig"); const lowerer = @import("lower.zig"); pub const Error = error{ - UnexpectedToken, + DuplicateName, + DuplicateValue, + InvalidCompositeIndex, InvalidNumber, + InvalidOpcode, + InvalidResult, + InvalidSemantic, InvalidStage, InvalidType, - InvalidOpcode, - InvalidSemantic, - DuplicateValue, - DuplicateName, - UnknownValue, - UnknownConstant, - UnknownInterface, - UnknownFunction, - UnknownBlock, - MissingTerminator, MissingResultType, - InvalidResult, - InvalidCompositeIndex, + MissingTerminator, + UnexpectedToken, + UnknownBlock, + UnknownConstant, + UnknownFunction, + UnknownInterface, + UnknownValue, }; pub const max_file_size = 64 * 1024 * 1024; @@ -59,7 +59,7 @@ const Parser = struct { try self.expectDiscard(.equal); const direction_token = try self.expect(.identifier); - const direction = std.meta.stringToEnum(module_ir.InterfaceDirection, direction_token.text) orelse return error.InvalidSemantic; + const direction = std.meta.stringToEnum(module_ir.InterfaceDirection, direction_token.text) orelse return Error.InvalidSemantic; try self.expectDiscard(.left_square); const semantic_name = (try self.expect(.identifier)).text; @@ -92,9 +92,9 @@ const Parser = struct { const builtin_name = (try self.expect(.identifier)).text; try self.expectDiscard(.right_paren); - const builtin = std.meta.stringToEnum(module_ir.Builtin, builtin_name) orelse return error.InvalidSemantic; + const builtin = std.meta.stringToEnum(module_ir.Builtin, builtin_name) orelse return Error.InvalidSemantic; break :blk .{ .builtin = builtin }; - } else return error.InvalidSemantic; + } else return Error.InvalidSemantic; try self.expectDiscard(.right_square); @@ -137,7 +137,7 @@ const Parser = struct { const bits = try self.parseUnsigned(u64, .number); try self.expectDiscard(.right_paren); - const ir_type = self.module.?.types.get(ty) orelse return error.InvalidType; + const ir_type = self.module.?.types.get(ty) orelse return Error.InvalidType; break :blk switch (ir_type.*) { .integer => .{ @@ -146,16 +146,16 @@ const Parser = struct { .floating => .{ .float_bits = bits, }, - else => return error.InvalidType, + else => return Error.InvalidType, }; } - return error.UnexpectedToken; + return Error.UnexpectedToken; }, .left_square => .{ .composite = try self.parseConstantList(), }, .number => try self.parseDirectConstant(ty), - else => return error.UnexpectedToken, + else => return Error.UnexpectedToken, }; return .{ @@ -167,12 +167,12 @@ const Parser = struct { fn parseDirectConstant(self: *Parser, ty: ids.TypeId) !ParsedConstantValue { const text = (try self.expect(.number)).text; - const ir_type = self.module.?.types.get(ty) orelse return error.InvalidType; + const ir_type = self.module.?.types.get(ty) orelse return Error.InvalidType; return switch (ir_type.*) { .integer => |integer| .{ .integer_bits = try parseIntegerLiteral(integer, text) }, .floating => |float| .{ .float_bits = try parseFloatLiteral(float.bits, text) }, - else => error.InvalidType, + else => Error.InvalidType, }; } @@ -210,7 +210,7 @@ const Parser = struct { while ((try self.peek()).tag != .right_brace) { if ((try self.peek()).tag != .dot_name) - return error.UnexpectedToken; + return Error.UnexpectedToken; try function.blocks.append(self.allocator, try self.parseBlock()); } @@ -246,7 +246,7 @@ const Parser = struct { const token = try self.peek(); if (token.tag == .right_brace or token.tag == .dot_name) - return error.MissingTerminator; + return Error.MissingTerminator; if (token.tag == .value_ref) { try block.instructions.append(self.allocator, try self.parseInstruction(true)); @@ -254,7 +254,7 @@ const Parser = struct { } if (token.tag != .identifier) - return error.UnexpectedToken; + return Error.UnexpectedToken; if (isTerminatorName(token.text)) { block.terminator = try self.parseTerminator(); @@ -285,7 +285,7 @@ const Parser = struct { if (std.mem.startsWith(u8, name, "cmp_")) { const opcode_name = name["cmp_".len..]; - const opcode = std.meta.stringToEnum(inst_ir.CompareOpcode, opcode_name) orelse return error.InvalidOpcode; + const opcode = std.meta.stringToEnum(inst_ir.CompareOpcode, opcode_name) orelse return Error.InvalidOpcode; const lhs = try self.parseValueRef(); try self.expectDiscard(.comma); @@ -337,7 +337,7 @@ const Parser = struct { } if (indices.items.len == 0) - return error.InvalidCompositeIndex; + return Error.InvalidCompositeIndex; return .{ .composite_extract = .{ .composite = composite, @@ -377,7 +377,7 @@ const Parser = struct { }; } - return error.InvalidOpcode; + return Error.InvalidOpcode; } fn parseTerminator(self: *Parser) !ParsedTerminator { @@ -414,7 +414,7 @@ const Parser = struct { if (std.mem.eql(u8, name, "unreachable")) return .unreachable_value; - return error.UnexpectedToken; + return Error.UnexpectedToken; } fn parseEdge(self: *Parser) !ParsedEdge { @@ -441,7 +441,7 @@ const Parser = struct { return module.internType(.boolean); if (std.mem.startsWith(u8, token.text, "vec")) { - const length = parseTextUnsigned(u8, token.text[3..]) catch return error.InvalidType; + const length = parseTextUnsigned(u8, token.text[3..]) catch return Error.InvalidType; try self.expectDiscard(.left_square); const element_type = try self.parseType(); @@ -497,7 +497,7 @@ const Parser = struct { try self.expectDiscard(.left_square); const address_name = (try self.expect(.identifier)).text; - const address_space = std.meta.stringToEnum(type_ir.AddressSpace, address_name) orelse return error.InvalidType; + const address_space = std.meta.stringToEnum(type_ir.AddressSpace, address_name) orelse return Error.InvalidType; try self.expectDiscard(.comma); const pointee_type = try self.parseType(); @@ -515,7 +515,7 @@ const Parser = struct { try self.expectDiscard(.left_square); const kind_name = (try self.expect(.identifier)).text; - const kind = std.meta.stringToEnum(type_ir.ResourceKind, kind_name) orelse return error.InvalidType; + const kind = std.meta.stringToEnum(type_ir.ResourceKind, kind_name) orelse return Error.InvalidType; try self.expectDiscard(.right_square); @@ -527,7 +527,7 @@ const Parser = struct { } if (token.text.len > 1 and (token.text[0] == 'i' or token.text[0] == 'u')) { - const bits = parseTextUnsigned(u16, token.text[1..]) catch return error.InvalidType; + const bits = parseTextUnsigned(u16, token.text[1..]) catch return Error.InvalidType; return module.internType(.{ .integer = .{ .bits = bits, .signedness = if (token.text[0] == 'i') .signed else .unsigned, @@ -535,14 +535,14 @@ const Parser = struct { } if (token.text.len > 1 and token.text[0] == 'f') { - const bits = parseTextUnsigned(u16, token.text[1..]) catch return error.InvalidType; + const bits = parseTextUnsigned(u16, token.text[1..]) catch return Error.InvalidType; return module.internType(.{ .floating = .{ .bits = bits, }, }); } - return error.InvalidType; + return Error.InvalidType; } fn parseConstantList(self: *Parser) ![]const u32 { @@ -595,13 +595,13 @@ const Parser = struct { fn parseUnsigned(self: *Parser, comptime T: type, tag: TokenTag) !T { const token = try self.expect(tag); - return parseTextUnsigned(T, token.text) catch error.InvalidNumber; + return parseTextUnsigned(T, token.text) catch Error.InvalidNumber; } fn expectIdentifier(self: *Parser, expected: []const u8) !void { const token = try self.expect(.identifier); if (!std.mem.eql(u8, token.text, expected)) - return error.UnexpectedToken; + return Error.UnexpectedToken; } fn expectDiscard(self: *Parser, tag: TokenTag) !void { @@ -611,7 +611,7 @@ const Parser = struct { fn expect(self: *Parser, tag: TokenTag) !Token { const token = try self.take(); if (token.tag != tag) - return error.UnexpectedToken; + return Error.UnexpectedToken; return token; } @@ -629,7 +629,7 @@ const Parser = struct { fn take(self: *Parser) !Token { const token = self.lexer.take(); if (token.tag == .invalid) - return error.UnexpectedToken; + return Error.UnexpectedToken; return token; } }; @@ -644,25 +644,25 @@ fn isTerminatorName(name: []const u8) bool { fn parseIntegerLiteral(integer: type_ir.IntegerType, text: []const u8) !u64 { if (integer.bits == 0 or integer.bits > 64) - return error.InvalidType; + return Error.InvalidType; if (integer.signedness == .unsigned) { - const value = std.fmt.parseInt(u64, text, 10) catch return error.InvalidNumber; + const value = std.fmt.parseInt(u64, text, 10) catch return Error.InvalidNumber; if (integer.bits < 64) { const shift: u6 = @intCast(integer.bits); const maximum = (@as(u64, 1) << shift) - 1; if (value > maximum) - return error.InvalidNumber; + return Error.InvalidNumber; } return value; } - const value = std.fmt.parseInt(i64, text, 10) catch return error.InvalidNumber; + const value = std.fmt.parseInt(i64, text, 10) catch return Error.InvalidNumber; if (integer.bits < 64) { const sign_shift: u6 = @intCast(integer.bits - 1); const magnitude = @as(i64, 1) << sign_shift; if (value < -magnitude or value > magnitude - 1) - return error.InvalidNumber; + return Error.InvalidNumber; const width: u6 = @intCast(integer.bits); const mask = (@as(u64, 1) << width) - 1; @@ -675,18 +675,18 @@ fn parseIntegerLiteral(integer: type_ir.IntegerType, text: []const u8) !u64 { fn parseFloatLiteral(bits: u16, text: []const u8) !u64 { return switch (bits) { 16 => blk: { - const value = std.fmt.parseFloat(f16, text) catch return error.InvalidNumber; + const value = std.fmt.parseFloat(f16, text) catch return Error.InvalidNumber; break :blk @as(u16, @bitCast(value)); }, 32 => blk: { - const value = std.fmt.parseFloat(f32, text) catch return error.InvalidNumber; + const value = std.fmt.parseFloat(f32, text) catch return Error.InvalidNumber; break :blk @as(u32, @bitCast(value)); }, 64 => blk: { - const value = std.fmt.parseFloat(f64, text) catch return error.InvalidNumber; + const value = std.fmt.parseFloat(f64, text) catch return Error.InvalidNumber; break :blk @as(u64, @bitCast(value)); }, - else => error.InvalidType, + else => Error.InvalidType, }; } @@ -695,7 +695,7 @@ fn parseTextUnsigned(comptime T: type, text: []const u8) !T { const digits = if (base == 16) text[2..] else text; if (digits.len == 0) - return error.InvalidNumber; + return Error.InvalidNumber; return std.fmt.parseInt(T, digits, base); } @@ -712,7 +712,7 @@ pub fn parseString(backing_allocator: std.mem.Allocator, source: []const u8) !mo try parser.expectIdentifier("shader"); const stage_token = try parser.expect(.identifier); - const stage = std.meta.stringToEnum(module_ir.Stage, stage_token.text) orelse return error.InvalidStage; + const stage = std.meta.stringToEnum(module_ir.Stage, stage_token.text) orelse return Error.InvalidStage; var module = module_ir.Module.init(backing_allocator, stage); errdefer module.deinit(); @@ -735,10 +735,10 @@ pub fn parseString(backing_allocator: std.mem.Allocator, source: []const u8) !mo if (std.mem.eql(u8, token.text, "fn")) { try parsed.functions.append(temporary_allocator, try parser.parseFunction()); } else { - return error.UnexpectedToken; + return Error.UnexpectedToken; } }, - else => return error.UnexpectedToken, + else => return Error.UnexpectedToken, } } try parser.expectDiscard(.right_brace); @@ -765,3 +765,194 @@ pub fn parseFileInDir(backing_allocator: std.mem.Allocator, io: std.Io, director return parseString(backing_allocator, source); } + +test "Parser: interface" { + const printer = @import("../printer.zig"); + + const source = + \\ shader vertex @main + \\ { + \\ @in_color: vec4[f32] = input[location(0), component(0), index(0)] + \\ @out_color: vec4[f32] = output[location(0), component(0), index(0)] + \\ @position: vec4[f32] = output[builtin(position)] + \\ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\ } + ; + + var module = try parseString(std.testing.allocator, source); + defer module.deinit(); + const printed = try printer.allocPrint(std.testing.allocator, &module); + defer std.testing.allocator.free(printed); + + try std.testing.expect(std.mem.indexOf(u8, printed, "@in_color: vec4[f32] = input[location(0), component(0), index(0)]") != null); + try std.testing.expect(std.mem.indexOf(u8, printed, "@out_color: vec4[f32] = output[location(0), component(0), index(0)]") != null); + try std.testing.expect(std.mem.indexOf(u8, printed, "@position: vec4[f32] = output[builtin(position)]") != null); +} + +test "Parser: types, operations, calls, terminators" { + const printer = @import("../printer.zig"); + + const source = + \\ shader fragment @main + \\ { + \\ %0: constant bool = true + \\ %1: constant u32 = bits(0x1) + \\ %2: constant u32 = bits(0x2) + \\ %3: constant f32 = bits(0x3f800000) + \\ %4: constant array[u32, 2] = [#1, #2] + \\ %5: constant struct[u32, u32] = [#1, #2] + \\ %6: constant ptr[private, u32] = null + \\ %7: constant resourceHandle[sampler] = null + \\ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %9: u32 = bitwise_not %1 + \\ %10: u32 = integer_add %9, %2 + \\ %11: bool = cmp_equal %1, %2 + \\ %12: u32 = select %11, %1, %2 + \\ %13: u32 = bitcast %12 + \\ %14: vec2[u32] = composite_construct %1, %2 + \\ %15: u32 = composite_extract %14[0] + \\ %16: f32 = negate %3 + \\ %17: f32 = float_add %3, %16 + \\ %18: u32 = call @helper(%15) + \\ return + \\ } + \\ + \\ fn @helper(%8: u32) -> u32 + \\ { + \\ .entry(): + \\ return %8 + \\ } + \\ + \\ fn @discarder() -> void + \\ { + \\ .entry(): + \\ discard + \\ } + \\ + \\ fn @dead() -> void + \\ { + \\ .entry(): + \\ unreachable + \\ } + \\ } + ; + + var module = try parseString(std.testing.allocator, source); + defer module.deinit(); + const printed = try printer.allocPrint(std.testing.allocator, &module); + defer std.testing.allocator.free(printed); + + var reparsed = try parseString(std.testing.allocator, printed); + defer reparsed.deinit(); + const printed_again = try printer.allocPrint(std.testing.allocator, &reparsed); + defer std.testing.allocator.free(printed_again); + try std.testing.expectEqualStrings(printed, printed_again); + try std.testing.expect(std.mem.indexOf(u8, printed, "cmp_equal %1, %2") != null); + try std.testing.expect(std.mem.indexOf(u8, printed, "cmp.") == null); +} + +test "Parser: named value IDs" { + const printer = @import("../printer.zig"); + + const source = + \\ shader compute @main + \\ { + \\ %one_value: constant u32 = bits(0x1) + \\ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %sum_value: u32 = integer_add %one_value, %one_value + \\ branch .merge(%sum_value) + \\ + \\ .merge(%merged_value: u32): + \\ %product_value: u32 = integer_multiply %merged_value, %one_value + \\ return + \\ } + \\ } + ; + + var module = try parseString(std.testing.allocator, source); + defer module.deinit(); + const printed = try printer.allocPrint(std.testing.allocator, &module); + defer std.testing.allocator.free(printed); + + try std.testing.expect(std.mem.indexOf(u8, printed, "%one_value: constant u32 = bits(0x1)") != null); + try std.testing.expect(std.mem.indexOf(u8, printed, "%sum_value: u32 = integer_add %one_value, %one_value") != null); + try std.testing.expect(std.mem.indexOf(u8, printed, "branch .merge(%sum_value)") != null); + try std.testing.expect(std.mem.indexOf(u8, printed, ".merge(%merged_value: u32)") != null); + try std.testing.expect(std.mem.indexOf(u8, printed, "%product_value: u32 = integer_multiply %merged_value, %one_value") != null); + + var reparsed = try parseString(std.testing.allocator, printed); + defer reparsed.deinit(); + const printed_again = try printer.allocPrint(std.testing.allocator, &reparsed); + defer std.testing.allocator.free(printed_again); + try std.testing.expectEqualStrings(printed, printed_again); +} + +test "Parser: numeric constants" { + const printer = @import("../printer.zig"); + + const source = + \\ shader compute @main + \\ { + \\ %0: constant u8 = 255 + \\ %1: constant i8 = -1 + \\ %2: constant f16 = 1.5 + \\ %3: constant f32 = -0.0 + \\ %4: constant f64 = 2.5e0 + \\ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\ } + ; + + var module = try parseString(std.testing.allocator, source); + defer module.deinit(); + const printed = try printer.allocPrint(std.testing.allocator, &module); + defer std.testing.allocator.free(printed); + + try std.testing.expect(std.mem.indexOf(u8, printed, "%0: constant u8 = bits(0xff)") != null); + try std.testing.expect(std.mem.indexOf(u8, printed, "%1: constant i8 = bits(0xff)") != null); + try std.testing.expect(std.mem.indexOf(u8, printed, "%2: constant f16 = bits(0x3e00)") != null); + try std.testing.expect(std.mem.indexOf(u8, printed, "%3: constant f32 = bits(0x80000000)") != null); + try std.testing.expect(std.mem.indexOf(u8, printed, "%4: constant f64 = bits(0x4004000000000000)") != null); + + const out_of_range = + \\ shader compute @main + \\ { + \\ %0: constant u8 = 256 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\ } + ; + try std.testing.expectError(Error.InvalidNumber, parseString(std.testing.allocator, out_of_range)); +} + +test "Parser: Error unknown value" { + const source = + \\ shader compute @main + \\ { + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ return %99 + \\ } + \\ } + ; + try std.testing.expectError(Error.UnknownValue, parseString(std.testing.allocator, source)); +} diff --git a/src/compiler/ir/pass_manager.zig b/src/compiler/ir/pass_manager.zig deleted file mode 100644 index 2a0eda1..0000000 --- a/src/compiler/ir/pass_manager.zig +++ /dev/null @@ -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", -}; diff --git a/src/compiler/ir/transformer_manager.zig b/src/compiler/ir/transformer_manager.zig new file mode 100644 index 0000000..8a376d1 --- /dev/null +++ b/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); +} diff --git a/src/compiler/ir/transformers/inline_all_functions.zig b/src/compiler/ir/transformers/inline_all_functions.zig new file mode 100644 index 0000000..c8b1086 --- /dev/null +++ b/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)); +} diff --git a/src/compiler/ir/validator/dominance.zig b/src/compiler/ir/validator/dominance.zig index 4eba3e5..1d94207 100644 --- a/src/compiler/ir/validator/dominance.zig +++ b/src/compiler/ir/validator/dominance.zig @@ -19,8 +19,8 @@ const DominanceUseContext = struct { pub fn validate(module: *const module_ir.Module, function_id: ids.FunctionId) Error!void { var analysis = cfg.init(module.backingAllocator(), module, function_id) catch |err| switch (err) { - error.OutOfMemory => return error.OutOfMemory, - else => return error.InvalidBlock, + std.mem.Allocator.Error.OutOfMemory => return Error.OutOfMemory, + else => return Error.InvalidBlock, }; defer analysis.deinit(); @@ -41,7 +41,7 @@ pub fn validate(module: *const module_ir.Module, function_id: ids.FunctionId) Er instruction.operation.visitValueUses(&context, checkDominanceUse); if (!context.valid) - return error.DefinitionDoesNotDominateUse; + return Error.DefinitionDoesNotDominateUse; } var context: DominanceUseContext = .{ @@ -55,7 +55,7 @@ pub fn validate(module: *const module_ir.Module, function_id: ids.FunctionId) Er module_ir.visitTerminatorValueUses(block.terminator.?, &context, checkDominanceUse); if (!context.valid) - return error.DefinitionDoesNotDominateUse; + return Error.DefinitionDoesNotDominateUse; } } diff --git a/src/compiler/ir/validator/validator.zig b/src/compiler/ir/validator/validator.zig index 4c4d0a0..8501f0e 100644 --- a/src/compiler/ir/validator/validator.zig +++ b/src/compiler/ir/validator/validator.zig @@ -3,44 +3,42 @@ const ids = @import("../id.zig"); const type_ir = @import("../type.zig"); const inst_ir = @import("../instruction.zig"); const module_ir = @import("../module.zig"); +const Builder = @import("../Builder.zig"); const dominance = @import("dominance.zig"); pub const ValidationError = error{ - MissingEntryPoint, - InvalidEntryPoint, - InvalidType, - InvalidConstant, - InvalidValue, - InvalidFunction, + CrossFunctionReference, + DefinitionDoesNotDominateUse, + EntryBlockHasPredecessor, InvalidBlock, + InvalidConstant, + InvalidEntryPoint, + InvalidFunction, InvalidInstruction, + InvalidStructuredControl, + InvalidType, + InvalidValue, + MissingEntryPoint, MissingFunctionEntryBlock, MissingTerminator, - EntryBlockHasPredecessor, - WrongParent, - WrongDefinition, - WrongParameterIndex, - WrongResultPresence, - WrongOperandType, - WrongResultType, WrongBranchArgumentCount, WrongBranchArgumentType, - CrossFunctionReference, - WrongReturnType, + WrongDefinition, WrongInterfaceDirection, - InvalidStructuredControl, - DefinitionDoesNotDominateUse, + WrongOperandType, + WrongParameterIndex, + WrongParent, + WrongResultPresence, + WrongResultType, + WrongReturnType, }; pub const Error = ValidationError || std.mem.Allocator.Error; -/// Early validator for the foundational IR. It covers object ownership, CFG -/// edges, single definitions, function boundaries, and the currently modeled -/// operation types, and SSA dominance. pub fn validate(module: *const module_ir.Module) Error!void { - const entry_point = module.entry_point orelse return error.MissingEntryPoint; + const entry_point = module.entry_point orelse return Error.MissingEntryPoint; if (!module.functions.isLive(entry_point)) - return error.InvalidEntryPoint; + return Error.InvalidEntryPoint; for (module.types.entries.items) |entry| { const ty = entry orelse continue; @@ -51,12 +49,12 @@ pub fn validate(module: *const module_ir.Module) Error!void { const constant = entry orelse continue; if (!module.types.isLive(constant.type)) - return error.InvalidType; + return Error.InvalidType; if (constant.value == .composite) { for (constant.value.composite) |element| { if (!module.constants.isLive(element)) - return error.InvalidConstant; + return Error.InvalidConstant; } } } @@ -65,30 +63,30 @@ pub fn validate(module: *const module_ir.Module) Error!void { const value = entry orelse continue; if (!module.types.isLive(value.type)) - return error.InvalidType; + return Error.InvalidType; const value_id = ids.ValueId.fromIndex(value_index); switch (value.definition) { .constant => |id| { - const constant = module.constants.get(id) orelse return error.InvalidConstant; + const constant = module.constants.get(id) orelse return Error.InvalidConstant; if (constant.type != value.type) - return error.WrongResultType; + return Error.WrongResultType; }, .function_parameter => |definition| { - const function = module.functions.get(definition.function) orelse return error.InvalidFunction; + const function = module.functions.get(definition.function) orelse return Error.InvalidFunction; if (definition.index >= function.parameters.items.len or function.parameters.items[definition.index] != value_id) - return error.WrongParameterIndex; + return Error.WrongParameterIndex; }, .block_parameter => |definition| { - const block = module.blocks.get(definition.block) orelse return error.InvalidBlock; + const block = module.blocks.get(definition.block) orelse return Error.InvalidBlock; if (definition.index >= block.parameters.items.len or block.parameters.items[definition.index] != value_id) - return error.WrongParameterIndex; + return Error.WrongParameterIndex; }, .instruction => |instruction_id| { - const instruction = module.instructions.get(instruction_id) orelse return error.InvalidInstruction; + const instruction = module.instructions.get(instruction_id) orelse return Error.InvalidInstruction; if (instruction.result != value_id) - return error.WrongDefinition; + return Error.WrongDefinition; }, .undef => {}, } @@ -97,13 +95,13 @@ pub fn validate(module: *const module_ir.Module) Error!void { for (module.interface_variables.entries.items) |entry| { const variable = entry orelse continue; if (!module.types.isLive(variable.type)) - return error.InvalidType; + return Error.InvalidType; } for (module.resources.entries.items) |entry| { const resource = entry orelse continue; if (!module.types.isLive(resource.type)) - return error.InvalidType; + return Error.InvalidType; } for (module.functions.entries.items, 0..) |entry, function_index| { @@ -111,32 +109,32 @@ pub fn validate(module: *const module_ir.Module) Error!void { const function_id = ids.FunctionId.fromIndex(function_index); if (!module.types.isLive(function.return_type)) - return error.InvalidType; + return Error.InvalidType; if (function.parameter_types.items.len != function.parameters.items.len) - return error.WrongParameterIndex; + return Error.WrongParameterIndex; for (function.parameter_types.items, function.parameters.items, 0..) |parameter_type, parameter_id, index| { - const parameter = module.values.get(parameter_id) orelse return error.InvalidValue; + const parameter = module.values.get(parameter_id) orelse return Error.InvalidValue; if (parameter.type != parameter_type) - return error.WrongResultType; + return Error.WrongResultType; if (parameter.definition != .function_parameter or parameter.definition.function_parameter.function != function_id or parameter.definition.function_parameter.index != index) - return error.WrongDefinition; + return Error.WrongDefinition; } - const entry_block = function.entry_block orelse return error.MissingFunctionEntryBlock; - const entry_block_value = module.blocks.get(entry_block) orelse return error.InvalidBlock; - if (entry_block_value.parent_function != function_id) return error.WrongParent; + const entry_block = function.entry_block orelse return Error.MissingFunctionEntryBlock; + const entry_block_value = module.blocks.get(entry_block) orelse return Error.InvalidBlock; + if (entry_block_value.parent_function != function_id) return Error.WrongParent; for (function.blocks.items) |block_id| { - const block = module.blocks.get(block_id) orelse return error.InvalidBlock; + const block = module.blocks.get(block_id) orelse return Error.InvalidBlock; if (block.parent_function != function_id) - return error.WrongParent; + return Error.WrongParent; try validateBlock(module, function_id, block_id, block); } @@ -146,7 +144,7 @@ pub fn validate(module: *const module_ir.Module) Error!void { if (block.terminator) |terminator| { if (targetsBlock(terminator, entry_block)) - return error.EntryBlockHasPredecessor; + return Error.EntryBlockHasPredecessor; } } @@ -158,23 +156,23 @@ fn validateType(module: *const module_ir.Module, ty: type_ir.Type) ValidationErr switch (ty) { .vector => |vector| { if (!module.types.isLive(vector.element_type) or vector.length < 2) - return error.InvalidType; + return ValidationError.InvalidType; }, .array => |array| { if (!module.types.isLive(array.element_type) or array.length == 0) - return error.InvalidType; + return ValidationError.InvalidType; }, .structure => |structure| for (structure.members) |member| { if (!module.types.isLive(member)) - return error.InvalidType; + return ValidationError.InvalidType; }, .pointer => |pointer| { if (!module.types.isLive(pointer.pointee_type)) - return error.InvalidType; + return ValidationError.InvalidType; }, .resource_handle => |handle| if (handle.data_type) |data_type| { if (!module.types.isLive(data_type)) - return error.InvalidType; + return ValidationError.InvalidType; }, else => {}, } @@ -187,11 +185,11 @@ fn validateBlock( block: *const module_ir.Block, ) ValidationError!void { for (block.parameters.items, 0..) |parameter_id, index| { - const parameter = module.values.get(parameter_id) orelse return error.InvalidValue; + const parameter = module.values.get(parameter_id) orelse return ValidationError.InvalidValue; if (parameter.definition != .block_parameter or parameter.definition.block_parameter.block != block_id or parameter.definition.block_parameter.index != index) - return error.WrongDefinition; + return ValidationError.WrongDefinition; } switch (block.structured_control) { @@ -204,97 +202,97 @@ fn validateBlock( } for (block.instructions.items) |instruction_id| { - const instruction = module.instructions.get(instruction_id) orelse return error.InvalidInstruction; + const instruction = module.instructions.get(instruction_id) orelse return ValidationError.InvalidInstruction; if (instruction.parent_block != block_id) - return error.WrongParent; + return ValidationError.WrongParent; if (instruction.result) |result_id| { - const result = module.values.get(result_id) orelse return error.InvalidValue; + const result = module.values.get(result_id) orelse return ValidationError.InvalidValue; if (result.definition != .instruction or result.definition.instruction != instruction_id) - return error.WrongDefinition; + return ValidationError.WrongDefinition; } try validateOperation(module, function_id, instruction); } - const terminator = block.terminator orelse return error.MissingTerminator; + const terminator = block.terminator orelse return ValidationError.MissingTerminator; try validateTerminator(module, function_id, terminator); } fn validateOperation(module: *const module_ir.Module, function_id: ids.FunctionId, instruction: *const inst_ir.Instruction) ValidationError!void { - const result_type = if (instruction.result) |result| module.typeOf(result) orelse return error.InvalidValue else null; + const result_type = if (instruction.result) |result| module.typeOf(result) orelse return ValidationError.InvalidValue else null; switch (instruction.operation) { .unary => |op| { const operand_type = try operandType(module, function_id, op.operand); if (result_type == null) - return error.WrongResultPresence; + return ValidationError.WrongResultPresence; if (result_type.? != operand_type) - return error.WrongResultType; + return ValidationError.WrongResultType; }, .binary => |op| { const lhs_type = try operandType(module, function_id, op.lhs); const rhs_type = try operandType(module, function_id, op.rhs); if (lhs_type != rhs_type) - return error.WrongOperandType; + return ValidationError.WrongOperandType; if (result_type == null or result_type.? != lhs_type) - return error.WrongResultType; + return ValidationError.WrongResultType; }, .compare => |op| { const lhs_type = try operandType(module, function_id, op.lhs); if (try operandType(module, function_id, op.rhs) != lhs_type) - return error.WrongOperandType; + return ValidationError.WrongOperandType; - const result = result_type orelse return error.WrongResultPresence; + const result = result_type orelse return ValidationError.WrongResultPresence; if (!isBoolean(module, result)) - return error.WrongResultType; + return ValidationError.WrongResultType; }, .select => |op| { if (!isBoolean(module, try operandType(module, function_id, op.condition))) - return error.WrongOperandType; + return ValidationError.WrongOperandType; const true_type = try operandType(module, function_id, op.true_value); if (try operandType(module, function_id, op.false_value) != true_type) - return error.WrongOperandType; + return ValidationError.WrongOperandType; if (result_type == null or result_type.? != true_type) - return error.WrongResultType; + return ValidationError.WrongResultType; }, .bitcast => |operand| { _ = try operandType(module, function_id, operand); if (result_type == null) - return error.WrongResultPresence; + return ValidationError.WrongResultPresence; }, .composite_construct => |op| { - const result = result_type orelse return error.WrongResultPresence; - const ty = module.types.get(result) orelse return error.InvalidType; + const result = result_type orelse return ValidationError.WrongResultPresence; + const ty = module.types.get(result) orelse return ValidationError.InvalidType; switch (ty.*) { .vector => |vector| { if (op.elements.len != vector.length) - return error.WrongOperandType; + return ValidationError.WrongOperandType; for (op.elements) |element| { if (try operandType(module, function_id, element) != vector.element_type) - return error.WrongOperandType; + return ValidationError.WrongOperandType; } }, .structure => |structure| { if (op.elements.len != structure.members.len) - return error.WrongOperandType; + return ValidationError.WrongOperandType; for (op.elements, structure.members) |element, member_type| { if (try operandType(module, function_id, element) != member_type) - return error.WrongOperandType; + return ValidationError.WrongOperandType; } }, - else => return error.WrongResultType, + else => return ValidationError.WrongResultType, } }, .composite_extract => |op| { @@ -302,127 +300,127 @@ fn validateOperation(module: *const module_ir.Module, function_id: ids.FunctionI const extracted_type = try indexedType(module, composite_type, op.indices); if (result_type == null or result_type.? != extracted_type) - return error.WrongResultType; + return ValidationError.WrongResultType; }, .load_interface => |op| { - const variable = module.interface_variables.get(op.variable) orelse return error.InvalidValue; + const variable = module.interface_variables.get(op.variable) orelse return ValidationError.InvalidValue; if (variable.direction != .input) - return error.WrongInterfaceDirection; + return ValidationError.WrongInterfaceDirection; if (op.element_index) |index| _ = try operandType(module, function_id, index); if (result_type == null or result_type.? != variable.type) - return error.WrongResultType; + return ValidationError.WrongResultType; }, .store_interface => |op| { if (result_type != null) - return error.WrongResultPresence; + return ValidationError.WrongResultPresence; - const variable = module.interface_variables.get(op.variable) orelse return error.InvalidValue; + const variable = module.interface_variables.get(op.variable) orelse return ValidationError.InvalidValue; if (variable.direction != .output) - return error.WrongInterfaceDirection; + return ValidationError.WrongInterfaceDirection; if (try operandType(module, function_id, op.value) != variable.type) - return error.WrongOperandType; + return ValidationError.WrongOperandType; if (op.element_index) |index| _ = try operandType(module, function_id, index); }, .call => |op| { - const callee = module.functions.get(op.function) orelse return error.InvalidFunction; + const callee = module.functions.get(op.function) orelse return ValidationError.InvalidFunction; if (op.arguments.len != callee.parameter_types.items.len) - return error.WrongOperandType; + return ValidationError.WrongOperandType; for (op.arguments, callee.parameter_types.items) |argument, parameter_type| { if (try operandType(module, function_id, argument) != parameter_type) - return error.WrongOperandType; + return ValidationError.WrongOperandType; } - const return_type = module.types.get(callee.return_type) orelse return error.InvalidType; + const return_type = module.types.get(callee.return_type) orelse return ValidationError.InvalidType; if (return_type.* == .void) { if (result_type != null) - return error.WrongResultPresence; + return ValidationError.WrongResultPresence; } else if (result_type == null or result_type.? != callee.return_type) - return error.WrongResultType; + return ValidationError.WrongResultType; }, } } fn validateTerminator(module: *const module_ir.Module, function_id: ids.FunctionId, terminator: module_ir.Terminator) ValidationError!void { - const function = module.functions.get(function_id) orelse return error.InvalidFunction; + const function = module.functions.get(function_id) orelse return ValidationError.InvalidFunction; switch (terminator) { .branch => |edge| try validateEdge(module, function_id, edge), .conditional_branch => |branch| { if (!isBoolean(module, try operandType(module, function_id, branch.condition))) - return error.WrongOperandType; + return ValidationError.WrongOperandType; try validateEdge(module, function_id, branch.true_edge); try validateEdge(module, function_id, branch.false_edge); }, .return_void => { if (module.types.get(function.return_type).?.* != .void) - return error.WrongReturnType; + return ValidationError.WrongReturnType; }, .return_value => |value| { if (try operandType(module, function_id, value) != function.return_type) - return error.WrongReturnType; + return ValidationError.WrongReturnType; }, .discard => { if (module.stage != .fragment) - return error.WrongReturnType; + return ValidationError.WrongReturnType; }, .@"unreachable" => {}, } } fn validateEdge(module: *const module_ir.Module, function_id: ids.FunctionId, edge: module_ir.Edge) ValidationError!void { - const target = module.blocks.get(edge.target) orelse return error.InvalidBlock; + const target = module.blocks.get(edge.target) orelse return ValidationError.InvalidBlock; if (target.parent_function != function_id) - return error.CrossFunctionReference; + return ValidationError.CrossFunctionReference; if (edge.arguments.len != target.parameters.items.len) - return error.WrongBranchArgumentCount; + return ValidationError.WrongBranchArgumentCount; for (edge.arguments, target.parameters.items) |argument, parameter| { if (try operandType(module, function_id, argument) != module.typeOf(parameter).?) - return error.WrongBranchArgumentType; + return ValidationError.WrongBranchArgumentType; } } fn validateTarget(module: *const module_ir.Module, function_id: ids.FunctionId, target_id: ids.BlockId) ValidationError!void { - const target = module.blocks.get(target_id) orelse return error.InvalidStructuredControl; + const target = module.blocks.get(target_id) orelse return ValidationError.InvalidStructuredControl; if (target.parent_function != function_id) - return error.InvalidStructuredControl; + return ValidationError.InvalidStructuredControl; } fn operandType(module: *const module_ir.Module, function_id: ids.FunctionId, value_id: ids.ValueId) ValidationError!ids.TypeId { - const value = module.values.get(value_id) orelse return error.InvalidValue; - const owner = valueFunction(module, value_id) catch return error.InvalidValue; + const value = module.values.get(value_id) orelse return ValidationError.InvalidValue; + const owner = valueFunction(module, value_id) catch return ValidationError.InvalidValue; if (owner) |actual| { if (actual != function_id) - return error.CrossFunctionReference; + return ValidationError.CrossFunctionReference; } return value.type; } fn valueFunction(module: *const module_ir.Module, value_id: ids.ValueId) ValidationError!?ids.FunctionId { - const value = module.values.get(value_id) orelse return error.InvalidValue; + const value = module.values.get(value_id) orelse return ValidationError.InvalidValue; return switch (value.definition) { .constant, .undef => null, .function_parameter => |definition| definition.function, - .block_parameter => |definition| (module.blocks.get(definition.block) orelse return error.InvalidBlock).parent_function, + .block_parameter => |definition| (module.blocks.get(definition.block) orelse return ValidationError.InvalidBlock).parent_function, .instruction => |instruction_id| blk: { - const instruction = module.instructions.get(instruction_id) orelse return error.InvalidInstruction; - const block = module.blocks.get(instruction.parent_block) orelse return error.InvalidBlock; + const instruction = module.instructions.get(instruction_id) orelse return ValidationError.InvalidInstruction; + const block = module.blocks.get(instruction.parent_block) orelse return ValidationError.InvalidBlock; break :blk block.parent_function; }, }; @@ -430,27 +428,27 @@ fn valueFunction(module: *const module_ir.Module, value_id: ids.ValueId) Validat fn indexedType(module: *const module_ir.Module, root: ids.TypeId, indices: []const u32) ValidationError!ids.TypeId { if (indices.len == 0) - return error.WrongOperandType; + return ValidationError.WrongOperandType; var current = root; for (indices) |index| { - const ty = module.types.get(current) orelse return error.InvalidType; + const ty = module.types.get(current) orelse return ValidationError.InvalidType; current = switch (ty.*) { .vector => |vector| if (index < vector.length) vector.element_type else - return error.WrongOperandType, + return ValidationError.WrongOperandType, .array => |array| if (index < array.length) array.element_type else - return error.WrongOperandType, + return ValidationError.WrongOperandType, .structure => |structure| if (index < structure.members.len) structure.members[index] else - return error.WrongOperandType, + return ValidationError.WrongOperandType, - else => return error.WrongOperandType, + else => return ValidationError.WrongOperandType, }; } return current; @@ -468,3 +466,108 @@ fn targetsBlock(terminator: module_ir.Terminator, target: ids.BlockId) bool { else => false, }; } + +test "Validator: Error wrong block argument count" { + // shader compute @main + // { + // fn @main() -> void + // { + // .entry(): + // branch .merge() + // + // .merge(%0: u32): + // return + // } + // } + + var module = module_ir.Module.init(std.testing.allocator, .compute); + defer module.deinit(); + var builder = Builder.init(&module); + + const void_type = try builder.internType(.void); + const u32_type = try builder.internType(.{ .integer = .{ .bits = 32, .signedness = .unsigned } }); + const main = try builder.addFunction(void_type, "main"); + builder.setEntryPoint(main); + const entry = try builder.addBlock(main, "entry"); + const merge = try builder.addBlock(main, "merge"); + _ = try builder.addBlockParameter(merge, u32_type, null); + try builder.setTerminator(entry, .{ .branch = try builder.edge(merge, &.{}) }); + try builder.setTerminator(merge, .return_void); + + try std.testing.expectError(Error.WrongBranchArgumentCount, validate(&module)); +} + +test "Validator: Error SSA definition does not dominate its use" { + // shader compute @main + // { + // %0: constant bool = true + // %1: constant u32 = bits(0x1) + // + // fn @main() -> void + // { + // .entry(): + // conditional_branch %0, .left(), .right() + // + // .left(): + // %2: u32 = integer_add %1, %1 + // branch .merge() + // + // .right(): + // branch .merge() + // + // .merge(): + // %3: u32 = integer_multiply %2, %1 + // return + // } + // } + + var module = module_ir.Module.init(std.testing.allocator, .compute); + defer module.deinit(); + var builder = Builder.init(&module); + + const void_type = try builder.internType(.void); + const bool_type = try builder.internType(.boolean); + const u32_type = try builder.internType(.{ .integer = .{ .bits = 32, .signedness = .unsigned } }); + const condition = try builder.internConstant(bool_type, .{ .boolean = true }); + const one = try builder.internConstant(u32_type, .{ .integer_bits = 1 }); + const main = try builder.addFunction(void_type, "main"); + builder.setEntryPoint(main); + const entry = try builder.addBlock(main, "entry"); + const left = try builder.addBlock(main, "left"); + const right = try builder.addBlock(main, "right"); + const merge = try builder.addBlock(main, "merge"); + + try builder.setTerminator( + entry, + .{ + .conditional_branch = .{ + .condition = condition, + .true_edge = try builder.edge(left, &.{}), + .false_edge = try builder.edge(right, &.{}), + }, + }, + ); + + const left_value = (try builder.appendInstruction(left, u32_type, .{ + .binary = .{ + .opcode = .integer_add, + .lhs = one, + .rhs = one, + }, + }, null)).?; + + try builder.setTerminator(left, .{ .branch = try builder.edge(merge, &.{}) }); + try builder.setTerminator(right, .{ .branch = try builder.edge(merge, &.{}) }); + + _ = try builder.appendInstruction(merge, u32_type, .{ + .binary = .{ + .opcode = .integer_multiply, + .lhs = left_value, + .rhs = one, + }, + }, null); + + try builder.setTerminator(merge, .return_void); + + try std.testing.expectError(Error.DefinitionDoesNotDominateUse, validate(&module)); +} diff --git a/src/compiler/root.zig b/src/compiler/root.zig index c8b1c6e..e496b6c 100644 --- a/src/compiler/root.zig +++ b/src/compiler/root.zig @@ -3,824 +3,15 @@ //! This module exposes the project-specific intermediate representation in //! `ir` and the SPIR-V frontend in `spirv`. //! -//! Together they form -//! the first stage of the compiler pipeline: SPIR-V binary modules are decoded, -//! translated into a smaller and easier-to-transform IR, validated, and then made -//! available to later optimization or code-generation passes. - -const std = @import("std"); +//! Together they form the first stage of the compiler pipeline: SPIR-V binary +//! modules are decoded, translated into a smaller and easier-to-transform IR, +//! validated, and then made available to later optimization or code-generation +//! transformers. pub const ir = @import("ir/ir.zig"); pub const spirv = @import("spirv/root.zig"); -const VisitorStatistics = struct { - functions: usize = 0, - blocks: usize = 0, -}; - -test "IR builder generation" { - // shader vertex @main - // { - // @color: vec4[f32] = input[location(0), component(0), index(0)] - // @out_color: vec4[f32] = output[location(0), component(0), index(0)] - // %0: constant bool = true - // %1: constant f32 = bits(0x3f800000) - // - // fn @main() -> void - // { - // .entry(): - // %3: vec4[f32] = load_interface @color - // conditional_branch %0, .pass(), .merge(%3) - // - // .pass(): - // %4: vec4[f32] = composite_construct %1, %1, %1, %1 - // branch .merge(%4) - // - // .merge(%2: vec4[f32]): - // store_interface @out_color, %2 - // return - // } - // } - - var module = ir.module.Module.init(std.testing.allocator, .vertex); - defer module.deinit(); - var builder = ir.Builder.init(&module); - - const void_type = try builder.internType(.void); - const bool_type = try builder.internType(.boolean); - const f32_type = try builder.internType(.{ .floating = .{ .bits = 32 } }); - const duplicate_f32 = try builder.internType(.{ .floating = .{ .bits = 32 } }); - try std.testing.expectEqual(f32_type, duplicate_f32); - const vec4_type = try builder.internType(.{ .vector = .{ .element_type = f32_type, .length = 4 } }); - - const true_value = try builder.internConstant(bool_type, .{ .boolean = true }); - const one = try builder.internConstant(f32_type, .{ .float_bits = @as(u32, @bitCast(@as(f32, 1.0))) }); - - const input = try builder.addInterfaceVariable(vec4_type, .input, .{ .location = .{ .location = 0 } }, "color"); - const output = try builder.addInterfaceVariable(vec4_type, .output, .{ .location = .{ .location = 0 } }, "out_color"); - const main = try builder.addFunction(void_type, "main"); - builder.setEntryPoint(main); - const entry = try builder.addBlock(main, "entry"); - const pass = try builder.addBlock(main, "pass"); - const merge = try builder.addBlock(main, "merge"); - const merged = try builder.addBlockParameter(merge, vec4_type, "merged"); - - const loaded = (try builder.appendInstruction(entry, vec4_type, .{ - .load_interface = .{ .variable = input }, - }, "loaded")).?; - try builder.setTerminator(entry, .{ .conditional_branch = .{ - .condition = true_value, - .true_edge = try builder.edge(pass, &.{}), - .false_edge = try builder.edge(merge, &.{loaded}), - } }); - - const splat = (try builder.appendInstruction(pass, vec4_type, .{ - .composite_construct = .{ .elements = &.{ one, one, one, one } }, - }, "white")).?; - try builder.setTerminator(pass, .{ .branch = try builder.edge(merge, &.{splat}) }); - _ = try builder.appendInstruction(merge, null, .{ - .store_interface = .{ .variable = output, .value = merged }, - }, null); - try builder.setTerminator(merge, .return_void); - - try ir.validator.validate(&module); - - var control_flow = try ir.cfg.init(std.testing.allocator, &module, main); - defer control_flow.deinit(); - try std.testing.expectEqual(@as(usize, 2), control_flow.predecessors(merge).?.len); - try std.testing.expect(control_flow.dominates(entry, merge)); - try std.testing.expect(!control_flow.dominates(pass, merge)); - - const text = try ir.printer.allocPrint(std.testing.allocator, &module); - - defer std.testing.allocator.free(text); - try std.testing.expect(std.mem.indexOf(u8, text, "shader vertex @main") != null); - try std.testing.expect(std.mem.indexOf(u8, text, "@color: vec4[f32] = input[location(0), component(0), index(0)]") != null); - try std.testing.expect(std.mem.indexOf(u8, text, "@out_color: vec4[f32] = output[location(0), component(0), index(0)]") != null); - try std.testing.expect(std.mem.indexOf(u8, text, "conditional_branch %0, .pass(), .merge(%loaded)") != null); - try std.testing.expect(std.mem.indexOf(u8, text, ".merge(%merged: vec4[f32])") != null); - try std.testing.expect(std.mem.indexOf(u8, text, "store_interface @out_color, %merged") != null); - - var parsed = try ir.parser.parseString(std.testing.allocator, text); - defer parsed.deinit(); - const round_trip = try ir.printer.allocPrint(std.testing.allocator, &parsed); - defer std.testing.allocator.free(round_trip); - try std.testing.expectEqualStrings(text, round_trip); - - const io = std.Options.debug_io; - const path = ".zig-cache/ir-parser-round-trip.ir"; - const file = try std.Io.Dir.cwd().createFile(io, path, .{ .truncate = true }); - { - defer file.close(io); - var file_buffer: [4096]u8 = @splat(0); - var file_writer = file.writer(io, &file_buffer); - try file_writer.interface.writeAll(text); - try file_writer.interface.flush(); - } - defer std.Io.Dir.cwd().deleteFile(io, path) catch @panic("Caught an error while handling an error"); - - var parsed_file = try ir.parser.parseFile(std.testing.allocator, io, path); - defer parsed_file.deinit(); - const file_round_trip = try ir.printer.allocPrint(std.testing.allocator, &parsed_file); - defer std.testing.allocator.free(file_round_trip); - try std.testing.expectEqualStrings(text, file_round_trip); -} - -test "IR parse interface" { - const source = - \\ shader vertex @main - \\ { - \\ @in_color: vec4[f32] = input[location(0), component(0), index(0)] - \\ @out_color: vec4[f32] = output[location(0), component(0), index(0)] - \\ @position: vec4[f32] = output[builtin(position)] - \\ - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ return - \\ } - \\ } - ; - - var module = try ir.parser.parseString(std.testing.allocator, source); - defer module.deinit(); - const printed = try ir.printer.allocPrint(std.testing.allocator, &module); - defer std.testing.allocator.free(printed); - - try std.testing.expect(std.mem.indexOf(u8, printed, "@in_color: vec4[f32] = input[location(0), component(0), index(0)]") != null); - try std.testing.expect(std.mem.indexOf(u8, printed, "@out_color: vec4[f32] = output[location(0), component(0), index(0)]") != null); - try std.testing.expect(std.mem.indexOf(u8, printed, "@position: vec4[f32] = output[builtin(position)]") != null); -} - -test "IR parse types, operations, calls, terminators" { - const source = - \\ shader fragment @main - \\ { - \\ %0: constant bool = true - \\ %1: constant u32 = bits(0x1) - \\ %2: constant u32 = bits(0x2) - \\ %3: constant f32 = bits(0x3f800000) - \\ %4: constant array[u32, 2] = [#1, #2] - \\ %5: constant struct[u32, u32] = [#1, #2] - \\ %6: constant ptr[private, u32] = null - \\ %7: constant resourceHandle[sampler] = null - \\ - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ %9: u32 = bitwise_not %1 - \\ %10: u32 = integer_add %9, %2 - \\ %11: bool = cmp_equal %1, %2 - \\ %12: u32 = select %11, %1, %2 - \\ %13: u32 = bitcast %12 - \\ %14: vec2[u32] = composite_construct %1, %2 - \\ %15: u32 = composite_extract %14[0] - \\ %16: f32 = negate %3 - \\ %17: f32 = float_add %3, %16 - \\ %18: u32 = call @helper(%15) - \\ return - \\ } - \\ - \\ fn @helper(%8: u32) -> u32 - \\ { - \\ .entry(): - \\ return %8 - \\ } - \\ - \\ fn @discarder() -> void - \\ { - \\ .entry(): - \\ discard - \\ } - \\ - \\ fn @dead() -> void - \\ { - \\ .entry(): - \\ unreachable - \\ } - \\ } - ; - - var module = try ir.parser.parseString(std.testing.allocator, source); - defer module.deinit(); - const printed = try ir.printer.allocPrint(std.testing.allocator, &module); - defer std.testing.allocator.free(printed); - - var reparsed = try ir.parser.parseString(std.testing.allocator, printed); - defer reparsed.deinit(); - const printed_again = try ir.printer.allocPrint(std.testing.allocator, &reparsed); - defer std.testing.allocator.free(printed_again); - try std.testing.expectEqualStrings(printed, printed_again); - try std.testing.expect(std.mem.indexOf(u8, printed, "cmp_equal %1, %2") != null); - try std.testing.expect(std.mem.indexOf(u8, printed, "cmp.") == null); -} - -test "IR parse named value IDs" { - const source = - \\ shader compute @main - \\ { - \\ %one_value: constant u32 = bits(0x1) - \\ - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ %sum_value: u32 = integer_add %one_value, %one_value - \\ branch .merge(%sum_value) - \\ - \\ .merge(%merged_value: u32): - \\ %product_value: u32 = integer_multiply %merged_value, %one_value - \\ return - \\ } - \\ } - ; - - var module = try ir.parser.parseString(std.testing.allocator, source); - defer module.deinit(); - const printed = try ir.printer.allocPrint(std.testing.allocator, &module); - defer std.testing.allocator.free(printed); - - try std.testing.expect(std.mem.indexOf(u8, printed, "%one_value: constant u32 = bits(0x1)") != null); - try std.testing.expect(std.mem.indexOf(u8, printed, "%sum_value: u32 = integer_add %one_value, %one_value") != null); - try std.testing.expect(std.mem.indexOf(u8, printed, "branch .merge(%sum_value)") != null); - try std.testing.expect(std.mem.indexOf(u8, printed, ".merge(%merged_value: u32)") != null); - try std.testing.expect(std.mem.indexOf(u8, printed, "%product_value: u32 = integer_multiply %merged_value, %one_value") != null); - - var reparsed = try ir.parser.parseString(std.testing.allocator, printed); - defer reparsed.deinit(); - const printed_again = try ir.printer.allocPrint(std.testing.allocator, &reparsed); - defer std.testing.allocator.free(printed_again); - try std.testing.expectEqualStrings(printed, printed_again); -} - -test "IR parse numeric constants" { - const source = - \\ shader compute @main - \\ { - \\ %0: constant u8 = 255 - \\ %1: constant i8 = -1 - \\ %2: constant f16 = 1.5 - \\ %3: constant f32 = -0.0 - \\ %4: constant f64 = 2.5e0 - \\ - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ return - \\ } - \\ } - ; - - var module = try ir.parser.parseString(std.testing.allocator, source); - defer module.deinit(); - const printed = try ir.printer.allocPrint(std.testing.allocator, &module); - defer std.testing.allocator.free(printed); - - try std.testing.expect(std.mem.indexOf(u8, printed, "%0: constant u8 = bits(0xff)") != null); - try std.testing.expect(std.mem.indexOf(u8, printed, "%1: constant i8 = bits(0xff)") != null); - try std.testing.expect(std.mem.indexOf(u8, printed, "%2: constant f16 = bits(0x3e00)") != null); - try std.testing.expect(std.mem.indexOf(u8, printed, "%3: constant f32 = bits(0x80000000)") != null); - try std.testing.expect(std.mem.indexOf(u8, printed, "%4: constant f64 = bits(0x4004000000000000)") != null); - - const out_of_range = - \\ shader compute @main - \\ { - \\ %0: constant u8 = 256 - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ return - \\ } - \\ } - ; - try std.testing.expectError(error.InvalidNumber, ir.parser.parseString(std.testing.allocator, out_of_range)); -} - -test "IR parser error: unknown value" { - const source = - \\ shader compute @main - \\ { - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ return %99 - \\ } - \\ } - ; - try std.testing.expectError(error.UnknownValue, ir.parser.parseString(std.testing.allocator, source)); -} - -test "Validator error: wrong block argument count" { - // shader compute @main - // { - // fn @main() -> void - // { - // .entry(): - // branch .merge() - // - // .merge(%0: u32): - // return - // } - // } - - var module = ir.module.Module.init(std.testing.allocator, .compute); - defer module.deinit(); - var builder = ir.Builder.init(&module); - - const void_type = try builder.internType(.void); - const u32_type = try builder.internType(.{ .integer = .{ .bits = 32, .signedness = .unsigned } }); - const main = try builder.addFunction(void_type, "main"); - builder.setEntryPoint(main); - const entry = try builder.addBlock(main, "entry"); - const merge = try builder.addBlock(main, "merge"); - _ = try builder.addBlockParameter(merge, u32_type, null); - try builder.setTerminator(entry, .{ .branch = try builder.edge(merge, &.{}) }); - try builder.setTerminator(merge, .return_void); - - try std.testing.expectError(error.WrongBranchArgumentCount, ir.validator.validate(&module)); -} - -test "Central store IDs disposal" { - var module = ir.module.Module.init(std.testing.allocator, .fragment); - defer module.deinit(); - const first = try module.internType(.boolean); - try std.testing.expect(module.types.remove(first)); - const second = try module.internType(.boolean); - try std.testing.expect(first.index() != second.index()); - try std.testing.expect(module.types.get(first) == null); -} - -test "Validator error: SSA definition does not dominate its use" { - // shader compute @main - // { - // %0: constant bool = true - // %1: constant u32 = bits(0x1) - // - // fn @main() -> void - // { - // .entry(): - // conditional_branch %0, .left(), .right() - // - // .left(): - // %2: u32 = integer_add %1, %1 - // branch .merge() - // - // .right(): - // branch .merge() - // - // .merge(): - // %3: u32 = integer_multiply %2, %1 - // return - // } - // } - - var module = ir.module.Module.init(std.testing.allocator, .compute); - defer module.deinit(); - var builder = ir.Builder.init(&module); - - const void_type = try builder.internType(.void); - const bool_type = try builder.internType(.boolean); - const u32_type = try builder.internType(.{ .integer = .{ .bits = 32, .signedness = .unsigned } }); - const condition = try builder.internConstant(bool_type, .{ .boolean = true }); - const one = try builder.internConstant(u32_type, .{ .integer_bits = 1 }); - const main = try builder.addFunction(void_type, "main"); - builder.setEntryPoint(main); - const entry = try builder.addBlock(main, "entry"); - const left = try builder.addBlock(main, "left"); - const right = try builder.addBlock(main, "right"); - const merge = try builder.addBlock(main, "merge"); - - try builder.setTerminator( - entry, - .{ - .conditional_branch = .{ - .condition = condition, - .true_edge = try builder.edge(left, &.{}), - .false_edge = try builder.edge(right, &.{}), - }, - }, - ); - - const left_value = (try builder.appendInstruction(left, u32_type, .{ - .binary = .{ - .opcode = .integer_add, - .lhs = one, - .rhs = one, - }, - }, null)).?; - - try builder.setTerminator(left, .{ .branch = try builder.edge(merge, &.{}) }); - try builder.setTerminator(right, .{ .branch = try builder.edge(merge, &.{}) }); - - _ = try builder.appendInstruction(merge, u32_type, .{ - .binary = .{ - .opcode = .integer_multiply, - .lhs = left_value, - .rhs = one, - }, - }, null); - - try builder.setTerminator(merge, .return_void); - - try std.testing.expectError(error.DefinitionDoesNotDominateUse, ir.validator.validate(&module)); -} - -test "Rewriter replace all ID uses, safely erase dead instruction" { - // shader compute @main - // { - // %0: constant u32 = bits(0x1) - // %1: constant u32 = bits(0x2) - // - // fn @main() -> void - // { - // .entry(): - // %2: u32 = integer_add %0, %1 - // %3: u32 = integer_multiply %2, %1 - // return - // } - // } - - var module = ir.module.Module.init(std.testing.allocator, .compute); - defer module.deinit(); - - var builder = ir.Builder.init(&module); - - const void_type = try builder.internType(.void); - const u32_type = try builder.internType(.{ .integer = .{ .bits = 32, .signedness = .unsigned } }); - - const one = try builder.internConstant(u32_type, .{ .integer_bits = 1 }); - const two = try builder.internConstant(u32_type, .{ .integer_bits = 2 }); - - const main = try builder.addFunction(void_type, "main"); - builder.setEntryPoint(main); - - const entry = try builder.addBlock(main, "entry"); - const sum = (try builder.appendInstruction(entry, u32_type, .{ - .binary = .{ - .opcode = .integer_add, - .lhs = one, - .rhs = two, - }, - }, null)).?; - _ = try builder.appendInstruction(entry, u32_type, .{ - .binary = .{ - .opcode = .integer_multiply, - .lhs = sum, - .rhs = two, - }, - }, null); - try builder.setTerminator(entry, .return_void); - - try ir.validator.validate(&module); - - const sum_instruction = module.values.get(sum).?.definition.instruction; - var rewriter = ir.Rewriter.init(&module); - - try std.testing.expectEqual(@as(usize, 1), try rewriter.replaceAllUses(sum, one)); - try rewriter.eraseInstruction(sum_instruction); - - try std.testing.expect(module.values.get(sum) == null); - try std.testing.expect(module.instructions.get(sum_instruction) == null); - - try ir.validator.validate(&module); -} - -test "Rewriter add block parameter and sync branch calls" { - // shader compute @main - // { - // %0: constant u32 = bits(0x1) - // - // fn @main() -> void - // { - // .entry(): - // branch .merge() - // - // .merge(): - // return - // - // .alternate(): - // return - // } - // } - - var module = ir.module.Module.init(std.testing.allocator, .compute); - defer module.deinit(); - - var builder = ir.Builder.init(&module); - - const void_type = try builder.internType(.void); - const u32_type = try builder.internType(.{ .integer = .{ .bits = 32, .signedness = .unsigned } }); - - const one = try builder.internConstant(u32_type, .{ .integer_bits = 1 }); - - const main = try builder.addFunction(void_type, "main"); - builder.setEntryPoint(main); - - const entry = try builder.addBlock(main, "entry"); - const merge = try builder.addBlock(main, "merge"); - const alternate = try builder.addBlock(main, "alternate"); - - try builder.setTerminator(entry, .{ .branch = try builder.edge(merge, &.{}) }); - try builder.setTerminator(merge, .return_void); - try builder.setTerminator(alternate, .return_void); - - var rewriter = ir.Rewriter.init(&module); - - const parameter = try rewriter.addBlockParameter(merge, u32_type, "incoming", &.{ - .{ - .predecessor = entry, - .value = one, - }, - }); - const merge_edge = module.blocks.get(entry).?.terminator.?.branch; - try std.testing.expectEqualSlices(ir.id.ValueId, &.{one}, merge_edge.arguments); - - _ = try builder.appendInstruction(merge, u32_type, .{ - .binary = .{ - .opcode = .integer_add, - .lhs = parameter, - .rhs = one, - }, - }, null); - try ir.validator.validate(&module); - - try rewriter.removeBlockParameter(merge, 0, one); - - try std.testing.expectEqual(@as(usize, 0), module.blocks.get(merge).?.parameters.items.len); - try std.testing.expectEqual(@as(usize, 0), module.blocks.get(entry).?.terminator.?.branch.arguments.len); - - try ir.validator.validate(&module); - - try std.testing.expectEqual(@as(usize, 1), try rewriter.redirectEdges(entry, merge, alternate, &.{})); - try std.testing.expectEqual(alternate, module.blocks.get(entry).?.terminator.?.branch.target); - - try ir.validator.validate(&module); -} - -fn establishNoCalls(_: *ir.module.Module, _: *ir.pass_manager.Context) !bool { - return false; -} - -fn countVisitedFunction(context: ?*anyopaque, _: ir.id.FunctionId, _: *const ir.module.Function) !void { - const statistics: *VisitorStatistics = @ptrCast(@alignCast(context.?)); - statistics.functions += 1; -} - -fn countVisitedBlock(context: ?*anyopaque, _: ir.id.BlockId, _: *const ir.module.Block) !void { - const statistics: *VisitorStatistics = @ptrCast(@alignCast(context.?)); - statistics.blocks += 1; -} - -test "Pass manager track independent IR properties" { - // shader compute @main - // { - // fn @main() -> void - // { - // .entry(): - // return - // } - // } - - var module = ir.module.Module.init(std.testing.allocator, .compute); - defer module.deinit(); - - var builder = ir.Builder.init(&module); - - const void_type = try builder.internType(.void); - - const main = try builder.addFunction(void_type, "main"); - builder.setEntryPoint(main); - - const entry = try builder.addBlock(main, "entry"); - try builder.setTerminator(entry, .return_void); - - module.properties.valid_cfg = true; - - var manager = ir.pass_manager.Manager.init(std.testing.allocator); - defer manager.deinit(); - - try manager.add(.{ - .name = "establish-no-calls", - .required = .{ .valid_cfg = true }, - .produced = .{ .no_function_calls = true }, - .run = establishNoCalls, - }); - - var context: ir.pass_manager.Context = .{ .allocator = std.testing.allocator }; - - try std.testing.expect(!try manager.run(&module, &context)); - try std.testing.expect(module.properties.no_function_calls); - - var statistics: VisitorStatistics = .{}; - - try ir.visitor.walk(&module, .{ - .context = &statistics, - .visitFunction = countVisitedFunction, - .visitBlock = countVisitedBlock, - }); - - try std.testing.expectEqual(@as(usize, 1), statistics.functions); - try std.testing.expectEqual(@as(usize, 1), statistics.blocks); -} - -test "SPIR-V parser error: zero-word instruction" { - const words = [_]u32{ - spirv.spec.magic_number, - 0x0001_0000, - 0, - 2, - 0, - instructionWord(.nop, 0), - }; - try std.testing.expectError(error.ZeroWordInstruction, spirv.Parser.init(&words)); - - const truncated = [_]u32{ - spirv.spec.magic_number, - 0x0001_0000, - 0, - 2, - 0, - instructionWord(.i_add, 5), - 1, - }; - try std.testing.expectError(error.TruncatedInstruction, spirv.Parser.init(&truncated)); -} - -test "SPIR-V structured branches and OpPhi to block parameters" { - const assembly = - \\ OpCapability Shader - \\ OpMemoryModel Logical GLSL450 - \\ OpEntryPoint GLCompute %main "main" - \\ OpExecutionMode %main LocalSize 1 1 1 - \\ OpName %main "main" - \\ OpName %entry "entry" - \\ OpName %true "true" - \\ OpName %one "one" - \\ OpName %then "then" - \\ OpName %then_value "then_value" - \\ OpName %else "else" - \\ OpName %else_value "else_value" - \\ OpName %merge "merge" - \\ OpName %merged "merged" - \\ OpName %product "product" - \\ - \\ %void = OpTypeVoid - \\ %bool = OpTypeBool - \\ %uint = OpTypeInt 32 0 - \\ %fn_void = OpTypeFunction %void - \\ %true = OpConstantTrue %bool - \\ %one = OpConstant %uint 1 - \\ - \\ %main = OpFunction %void None %fn_void - \\ %entry = OpLabel - \\ OpSelectionMerge %merge None - \\ OpBranchConditional %true %then %else - \\ %then = OpLabel - \\ %then_value = OpIAdd %uint %one %one - \\ OpBranch %merge - \\ %else = OpLabel - \\ %else_value = OpISub %uint %one %one - \\ OpBranch %merge - \\ %merge = OpLabel - \\ %merged = OpPhi %uint %then_value %then %else_value %else - \\ %product = OpIMul %uint %merged %one - \\ OpReturn - \\ OpFunctionEnd - ; - const words = try assembleSpirv(std.testing.allocator, assembly); - defer std.testing.allocator.free(words); - - var module = try spirv.translator.translate(std.testing.allocator, words, .{ .entry_point = "main" }); - defer module.deinit(); - - try std.testing.expectEqual(ir.module.Stage.compute, module.stage); - try std.testing.expectEqual([3]u32{ 1, 1, 1 }, module.execution_modes.workgroup_size.?); - try std.testing.expect(module.properties.valid_cfg); - try std.testing.expect(module.properties.valid_ssa); - - const function = module.functions.get(module.entry_point.?).?; - try std.testing.expectEqual(@as(usize, 4), function.blocks.items.len); - const entry = module.blocks.get(function.blocks.items[0]).?; - try std.testing.expect(entry.structured_control == .selection); - const merge = module.blocks.get(function.blocks.items[3]).?; - try std.testing.expectEqual(@as(usize, 1), merge.parameters.items.len); - try std.testing.expectEqual(@as(usize, 1), merge.instructions.items.len); - const multiply = module.instructions.get(merge.instructions.items[0]).?; - try std.testing.expectEqual(ir.instruction.BinaryOpcode.integer_multiply, multiply.operation.binary.opcode); - - const text = try ir.printer.allocPrint(std.testing.allocator, &module); - defer std.testing.allocator.free(text); - - try std.testing.expect(std.mem.indexOf(u8, text, "%one: constant u32 = bits(0x1)") != null); - try std.testing.expect(std.mem.indexOf(u8, text, "%true: constant bool = true") != null); - try std.testing.expect(std.mem.indexOf(u8, text, "conditional_branch %true, .then(), .else()") != null); - try std.testing.expect(std.mem.indexOf(u8, text, "%then_value: u32 = integer_add %one, %one") != null); - try std.testing.expect(std.mem.indexOf(u8, text, "branch .merge(%then_value)") != null); - try std.testing.expect(std.mem.indexOf(u8, text, "%else_value: u32 = integer_subtract %one, %one") != null); - try std.testing.expect(std.mem.indexOf(u8, text, ".merge(%merged: u32)") != null); - try std.testing.expect(std.mem.indexOf(u8, text, "%product: u32 = integer_multiply %merged, %one") != null); - try std.testing.expect(std.mem.indexOf(u8, text, "integerMultiply") == null); - - var parsed = try ir.parser.parseString(std.testing.allocator, text); - defer parsed.deinit(); - const round_trip = try ir.printer.allocPrint(std.testing.allocator, &parsed); - defer std.testing.allocator.free(round_trip); - try std.testing.expectEqualStrings(text, round_trip); -} - -test "SPIR-V decorated vertex interfaces and load-store operations" { - const assembly = - \\ OpCapability Shader - \\ OpMemoryModel Logical GLSL450 - \\ OpEntryPoint Vertex %main "main" %in_color %out_color - \\ OpName %in_color "in_color" - \\ OpName %out_color "out_color" - \\ OpDecorate %in_color Location 0 - \\ OpDecorate %out_color Location 0 - \\ - \\ %void = OpTypeVoid - \\ %float = OpTypeFloat 32 - \\ %vec4 = OpTypeVector %float 4 - \\ %input_vec4 = OpTypePointer Input %vec4 - \\ %output_vec4 = OpTypePointer Output %vec4 - \\ %fn_void = OpTypeFunction %void - \\ %in_color = OpVariable %input_vec4 Input - \\ %out_color = OpVariable %output_vec4 Output - \\ - \\ %main = OpFunction %void None %fn_void - \\ %entry = OpLabel - \\ %color = OpLoad %vec4 %in_color - \\ OpStore %out_color %color - \\ OpReturn - \\ OpFunctionEnd - ; - const words = try assembleSpirv(std.testing.allocator, assembly); - defer std.testing.allocator.free(words); - - var module = try spirv.translator.translate(std.testing.allocator, words, .{ .entry_point = "main" }); - defer module.deinit(); - try std.testing.expectEqual(ir.module.Stage.vertex, module.stage); - try std.testing.expectEqual(@as(usize, 2), module.interface_variables.entries.items.len); - - const text = try ir.printer.allocPrint(std.testing.allocator, &module); - defer std.testing.allocator.free(text); - try std.testing.expect(std.mem.indexOf(u8, text, "load_interface @in_color") != null); - try std.testing.expect(std.mem.indexOf(u8, text, "store_interface @out_color") != null); -} - -fn instructionWord(opcode: spirv.spec.Opcode, word_count: u16) u32 { - return (@as(u32, word_count) << 16) | @intFromEnum(opcode); -} - -fn assembleSpirv(allocator: std.mem.Allocator, assembly: []const u8) ![]u32 { - var io_backend: std.Io.Threaded = .init(allocator, .{}); - defer io_backend.deinit(); - const io = io_backend.io(); - - var child = try std.process.spawn(io, .{ - .argv = &.{ "spirv-as", "--target-env", "spv1.0", "-o", "-", "-" }, - .stdin = .pipe, - .stdout = .pipe, - .stderr = .pipe, - }); - defer child.kill(io); - - { - const stdin = child.stdin.?; - var stdin_writer = stdin.writer(io, &.{}); - try stdin_writer.interface.writeAll(assembly); - try stdin_writer.interface.flush(); - stdin.close(io); - child.stdin = null; - } - - var stdout_buffer: [4096]u8 = undefined; - var stdout_reader = child.stdout.?.reader(io, &stdout_buffer); - const binary = try stdout_reader.interface.allocRemaining(allocator, .limited(1024 * 1024)); - defer allocator.free(binary); - - var stderr_buffer: [4096]u8 = undefined; - var stderr_reader = child.stderr.?.reader(io, &stderr_buffer); - const stderr = try stderr_reader.interface.allocRemaining(allocator, .limited(64 * 1024)); - defer allocator.free(stderr); - - const term = try child.wait(io); - switch (term) { - .exited => |code| if (code != 0) { - std.log.err("spirv-as failed:\n{s}", .{stderr}); - return error.SpirvAssemblyFailed; - }, - else => { - std.log.err("spirv-as terminated unexpectedly:\n{s}", .{stderr}); - return error.SpirvAssemblyFailed; - }, - } - - if (binary.len % @sizeOf(u32) != 0) return error.InvalidSpirvBinaryLength; - const words = try allocator.alloc(u32, binary.len / @sizeOf(u32)); - errdefer allocator.free(words); - for (words, 0..) |*word, index| { - const offset = index * @sizeOf(u32); - word.* = std.mem.readInt(u32, binary[offset..][0..4], .little); - } - return words; +test { + _ = ir; + _ = spirv; } diff --git a/src/compiler/spirv/Parser.zig b/src/compiler/spirv/Parser.zig index 48be5df..88056f9 100644 --- a/src/compiler/spirv/Parser.zig +++ b/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); +} diff --git a/src/compiler/spirv/root.zig b/src/compiler/spirv/root.zig index c4b3186..49ae358 100644 --- a/src/compiler/spirv/root.zig +++ b/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; +} diff --git a/src/compiler/spirv/translator.zig b/src/compiler/spirv/translator.zig index 78b1856..463c01a 100644 --- a/src/compiler/spirv/translator.zig +++ b/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; +} diff --git a/src/intel/compiler/compiler.zig b/src/intel/compiler/compiler.zig new file mode 100644 index 0000000..2a48e91 --- /dev/null +++ b/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; +} diff --git a/src/intel/compiler/id.zig b/src/intel/compiler/ir/id.zig similarity index 85% rename from src/intel/compiler/id.zig rename to src/intel/compiler/ir/id.zig index 6b86474..c1c1061 100644 --- a/src/intel/compiler/id.zig +++ b/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); diff --git a/src/intel/compiler/instruction.zig b/src/intel/compiler/ir/instruction.zig similarity index 90% rename from src/intel/compiler/instruction.zig rename to src/intel/compiler/ir/instruction.zig index 03845ec..8e9658c 100644 --- a/src/intel/compiler/instruction.zig +++ b/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, diff --git a/src/intel/compiler/ir/ir.zig b/src/intel/compiler/ir/ir.zig new file mode 100644 index 0000000..154b7af --- /dev/null +++ b/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; diff --git a/src/intel/compiler/operand.zig b/src/intel/compiler/ir/operand.zig similarity index 98% rename from src/intel/compiler/operand.zig rename to src/intel/compiler/ir/operand.zig index 21bc0e6..8219115 100644 --- a/src/intel/compiler/operand.zig +++ b/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 { diff --git a/src/intel/compiler/printer.zig b/src/intel/compiler/ir/printer.zig similarity index 56% rename from src/intel/compiler/printer.zig rename to src/intel/compiler/ir/printer.zig index 28c9738..e81606c 100644 --- a/src/intel/compiler/printer.zig +++ b/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}); - if (program.entry_function) |entry| { - try writer.writeAll("; .entry: "); - try writeFunctionRef(program, writer, entry); - try writer.writeByte('\n'); - } - - try writer.writeByte('\n'); + try writer.print("; .dispatch_width: {t}\n\n", .{program.dispatch_width}); for (program.virtual_registers.entries.items, 0..) |entry, index| { const register = entry orelse continue; try writeVirtualRegisterRef(program, writer, ids.VirtualRegisterId.fromIndex(index)); - try writer.print(": vgrf {t}[{d}] = class({t}), size({d}), alignment({d}), 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| { - const function = entry orelse continue; + try writer.writeByte('\n'); - try writer.writeAll("\nfn "); - try writeFunctionRef(program, writer, ids.FunctionId.fromIndex(function_index)); - try writer.writeByte('('); - 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("", .{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 (program.blocks.entries.items, 0..) |entry, block_index| { + const block = entry orelse continue; + const block_id = ids.BlockId.fromIndex(block_index); - for (function.blocks.items) |block_id| { - const block = program.blocks.get(block_id) orelse continue; + try writeBlockRef(program, writer, block_id); + try writer.writeAll(":\n"); - try writer.writeAll(indent); - try writeBlockRef(program, writer, block_id); - try writer.writeAll(":\n"); - - switch (block.structured_control) { - .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.*); + switch (block.structured_control) { + .none => {}, + .selection => |selection| { + try writer.writeAll(indent ++ "structured_selection "); + try writeBlockRef(program, writer, selection.merge_block); try writer.writeByte('\n'); - } - - if (block.terminator) |terminator| { - try writer.writeAll(indent ** 2); - try writeTerminator(program, writer, terminator); - try writer.writeAll("\n\n"); - } else { - try writer.writeAll(indent ** 2 ++ "\n\n"); - } + }, + .loop => |loop| { + try writer.writeAll(indent ++ "structured_loop merge("); + try writeBlockRef(program, writer, loop.merge_block); + try writer.writeAll("), continue("); + try writeBlockRef(program, writer, loop.continue_block); + try writer.writeAll(")\n"); + }, } - 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 ++ "\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(" = move "); - try writeSource(program, writer, op.source); + try writer.writeAll("mov "); + try writeDestination(program, writer, execution_size, op.destination); + try writer.writeAll(", "); + try writeSource(program, writer, execution_size, op.source); }, .binary => |op| { - try writeDestination(program, writer, op.destination); - try writer.print(" = {t} ", .{op.opcode}); - try writeSource(program, writer, op.lhs); + try writer.print("{t} ", .{op.opcode}); + try writeDestination(program, writer, execution_size, op.destination); 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); - }, - .call => |op| { - 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(')'); + try writeSource(program, writer, execution_size, op.lhs); + try writer.writeAll(", "); + try writeSource(program, writer, execution_size, op.rhs); }, .send => |op| { + try writer.writeAll("send "); if (op.response) |response| { try writeRegisterSpan(program, writer, response); - try writer.writeAll(" = "); + try 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) { - .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 writeRegister(program, writer, source.register); try writer.print(":{t}", .{source.type}); + if (source.register != .immediate) + try writeSourceRegion(writer, execution_size, source.register, source.region); if (source.absolute) try writer.writeByte(')'); } -fn writeDestination(program: *const program_ir.Program, writer: *std.Io.Writer, destination: operand.Destination) !void { - try writeRegisterAtOffset(program, writer, destination.register, destination.region.byte_offset); - try writer.print("<{d}>:{t}", .{ destination.region.horizontal_stride, destination.type }); +fn writeDestination(program: *const program_ir.Program, writer: *std.Io.Writer, execution_size: device.ExecutionSize, destination: operand.Destination) !void { + _ = execution_size; + try writeRegister(program, writer, destination.register); + try writer.print(":{t}", .{destination.type}); + try writeDestinationRegion(writer, destination.register, destination.region); } -fn writeRegisterAtOffset( - program: *const program_ir.Program, - writer: *std.Io.Writer, - register: operand.RegisterRef, - byte_offset: u16, -) !void { +fn writeSourceRegion(writer: *std.Io.Writer, execution_size: device.ExecutionSize, register: operand.RegisterRef, region: operand.Region) !void { + const byte_offset = registerByteOffset(register) + region.byte_offset; + const execution_width: u8 = @intFromEnum(execution_size); + const is_default = region.vertical_stride == execution_width and + region.width == execution_width and + region.horizontal_stride == 1; + const is_broadcast = region.vertical_stride == 0 and + region.width == 1 and + region.horizontal_stride == 0; + + if (byte_offset == 0 and is_default) + return; + + try writer.writeByte('['); + if (byte_offset != 0) + try writer.print("byte={d}", .{byte_offset}); + + if (is_broadcast) { + if (byte_offset != 0) + try writer.writeAll(", "); + try writer.writeAll("broadcast"); + } else if (!is_default) { + if (byte_offset != 0) + try writer.writeAll(", "); + try writer.print("vstride={d}, width={d}, hstride={d}", .{ + region.vertical_stride, + region.width, + region.horizontal_stride, + }); + } + try writer.writeByte(']'); +} + +fn writeDestinationRegion(writer: *std.Io.Writer, register: operand.RegisterRef, region: operand.DestinationRegion) !void { + const byte_offset = registerByteOffset(register) + region.byte_offset; + if (byte_offset == 0 and region.horizontal_stride == 1) + return; + + try writer.writeByte('['); + if (byte_offset != 0) + try writer.print("byte={d}", .{byte_offset}); + if (region.horizontal_stride != 1) { + if (byte_offset != 0) + try writer.writeAll(", "); + try writer.print("hstride={d}", .{region.horizontal_stride}); + } + try writer.writeByte(']'); +} + +fn registerByteOffset(register: operand.RegisterRef) u16 { + return switch (register) { + .physical_grf => |physical| physical.byte_offset, + else => 0, + }; +} + +fn writeRegister(program: *const program_ir.Program, writer: *std.Io.Writer, register: operand.RegisterRef) !void { switch (register) { - .virtual => |virtual| { - try writeVirtualRegisterRef(program, writer, virtual); - try writer.print(".{d}", .{byte_offset}); - }, - .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}); - }, + .virtual => |virtual| try writeVirtualRegisterRef(program, writer, virtual), + .physical_grf => |physical| try writer.print("r{d}", .{physical.number}), + .architecture => |architecture| try writeArchitectureRegister(writer, architecture), .immediate => |immediate| try writeImmediate(writer, immediate), .null => try writer.writeAll("null"), } @@ -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| { diff --git a/src/intel/compiler/program.zig b/src/intel/compiler/ir/program.zig similarity index 60% rename from src/intel/compiler/program.zig rename to src/intel/compiler/ir/program.zig index 147115d..f0db2b0 100644 --- a/src/intel/compiler/program.zig +++ b/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( - backing_allocator: std.mem.Allocator, - stage: Stage, - device_info: device.DeviceInfo, - dispatch_width: device.DispatchWidth, - ) Program { + pub fn init(backing_allocator: std.mem.Allocator, stage: Stage, device_info: device.DeviceInfo, dispatch_width: device.DispatchWidth) Program { return .{ .arena = std.heap.ArenaAllocator.init(backing_allocator), .stage = stage, @@ -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 { - const function = self.functions.getMut(function_id) orelse return error.InvalidFunction; + pub fn addBlock(self: *Program, name: ?[]const u8) !ids.BlockId { const owned_name = if (name) |value| try self.allocator().dupe(u8, value) else null; const block_id = try self.blocks.add(self.allocator(), .{ - .parent_function = function_id, .name = owned_name, }); - errdefer _ = self.blocks.remove(block_id); - - try function.blocks.append(self.allocator(), block_id); - if (function.entry_block == null) - function.entry_block = block_id; + if (self.entry_block == null) + self.entry_block = block_id; return block_id; } - pub fn appendInstruction( - self: *Program, - block_id: ids.BlockId, - execution_size: device.ExecutionSize, - predicate: ?operand.Predicate, - operation: instructions.Operation, - ) !ids.InstructionId { + pub fn setEntryBlock(self: *Program, block_id: ids.BlockId) !void { + if (!self.blocks.isLive(block_id)) + return error.InvalidBlock; + self.entry_block = block_id; + } + + pub fn appendInstruction(self: *Program, block_id: ids.BlockId, execution_size: device.ExecutionSize, predicate: ?operand.Predicate, operation: instructions.Operation) !ids.InstructionId { const block = self.blocks.getMut(block_id) orelse return error.InvalidBlock; - 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; diff --git a/src/intel/compiler/validator.zig b/src/intel/compiler/ir/validator.zig similarity index 54% rename from src/intel/compiler/validator.zig rename to src/intel/compiler/ir/validator.zig index c7230c7..201735a 100644 --- a/src/intel/compiler/validator.zig +++ b/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_function = program.functions.get(entry_function_id) orelse return error.InvalidFunction; - const entry_block = entry_function.entry_block orelse return error.MissingEntryBlock; - 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; - } - } + const entry_block = program.entry_block orelse return error.MissingEntryBlock; + if (!program.blocks.isLive(entry_block)) + return error.InvalidBlock; for (program.virtual_registers.entries.items) |entry| { const register = entry orelse continue; @@ -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 validateTerminator(program, block.parent_function, block.terminator.?); + try validateStructuredControl(program, block.structured_control); + 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( - 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; +fn validateTerminator(program: *const program_ir.Program, terminator: instruction.Terminator) Error!void { 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, function_id, branch.false_block); + try validateBlockTarget(program, branch.true_block); + try validateBlockTarget(program, branch.false_block); }, - .return_void => if (function.return_type != null) - 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" => {}, + .end_thread, .@"unreachable" => {}, } } -fn validateStructuredControl( - program: *const program_ir.Program, - function_id: ids.FunctionId, - control: instruction.StructuredControl, -) Error!void { +fn validateStructuredControl(program: *const program_ir.Program, 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, function_id, loop.continue_block); + try validateBlockTarget(program, loop.merge_block); + try validateBlockTarget(program, loop.continue_block); }, } } -fn validateBlockTarget(program: *const program_ir.Program, function_id: ids.FunctionId, block_id: ids.BlockId) Error!void { - const block = program.blocks.get(block_id) orelse return error.InvalidBlock; - if (block.parent_function != function_id) - 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; +fn validateBlockTarget(program: *const program_ir.Program, block_id: ids.BlockId) Error!void { + if (!program.blocks.isLive(block_id)) + return error.InvalidBlock; } diff --git a/src/intel/compiler/lower.zig b/src/intel/compiler/lower.zig deleted file mode 100644 index c8f1dfb..0000000 --- a/src/intel/compiler/lower.zig +++ /dev/null @@ -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(); -} diff --git a/src/intel/compiler/lower/lower.zig b/src/intel/compiler/lower/lower.zig new file mode 100644 index 0000000..bc75bf5 --- /dev/null +++ b/src/intel/compiler/lower/lower.zig @@ -0,0 +1,1186 @@ +const std = @import("std"); +const shader_ir = @import("shader_ir").ir; +const device = @import("../device.zig"); +const ids = @import("../ir/id.zig"); +const instruction = @import("../ir/instruction.zig"); +const operand = @import("../ir/operand.zig"); +const printer = @import("../ir/printer.zig"); +const program_ir = @import("../ir/program.zig"); +const validator = @import("../ir/validator.zig"); + +pub const Options = struct { + dispatch_width: device.DispatchWidth = .simd8, +}; + +pub const Error = std.mem.Allocator.Error || error{ + MissingEntryPoint, + InvalidEntryPoint, + InvalidModule, + InvalidLoweredProgram, + SanitizationFailed, + UnsanitizedModule, + UnsupportedGeneration, + UnsupportedStage, + UnsupportedDispatchWidth, + UnsupportedType, + UnsupportedOperation, + UnsupportedTerminator, +}; + +const PredicateValue = union(enum) { + constant: bool, + dynamic: operand.Predicate, +}; + +const ValueLocation = union(enum) { + source: operand.Source, + predicate: PredicateValue, +}; + +const LoweringState = struct { + lowerer: *Lowerer, + program: *program_ir.Program, + block_map: []?ids.BlockId, + value_locations: []?ValueLocation, + + fn lowerType(self: *const LoweringState, type_id: shader_ir.id.TypeId) Error!operand.DataType { + const ty = self.lowerer.module.types.get(type_id) orelse return Error.InvalidModule; + return switch (ty.*) { + .void => Error.UnsupportedType, + .integer => |integer| if (integer.bits == 32) + switch (integer.signedness) { + .unsigned => .u32, + .signed => .i32, + } + else + Error.UnsupportedType, + .floating => |floating| if (floating.bits == 32) .f32 else Error.UnsupportedType, + else => Error.UnsupportedType, + }; + } + + fn isBoolean(self: *const LoweringState, type_id: shader_ir.id.TypeId) Error!bool { + const ty = self.lowerer.module.types.get(type_id) orelse return Error.InvalidModule; + return ty.* == .boolean; + } + + fn mappedBlock(self: *const LoweringState, source_id: shader_ir.id.BlockId) Error!ids.BlockId { + if (source_id.index() >= self.block_map.len) + return Error.InvalidModule; + return self.block_map[source_id.index()] orelse Error.InvalidModule; + } + + fn putLocation(self: *LoweringState, value_id: shader_ir.id.ValueId, new_location: ValueLocation) Error!void { + if (value_id.index() >= self.value_locations.len or self.value_locations[value_id.index()] != null) + return Error.InvalidModule; + self.value_locations[value_id.index()] = new_location; + } + + fn addRegister(self: *LoweringState, data_type: operand.DataType, class: operand.RegisterClass, name: ?[]const u8) Error!ids.VirtualRegisterId { + return self.program.addVirtualRegister(.{ + .size_bytes = @as(u32, data_type.sizeBytes()) * @intFromEnum(self.lowerer.options.dispatch_width), + .alignment_bytes = self.lowerer.device_info.grf_size_bytes, + .element_type = data_type, + .lane_count = @intFromEnum(self.lowerer.options.dispatch_width), + .class = class, + .name = name, + }) catch |err| return mapProgramError(err); + } + + fn registerSource(self: *const LoweringState, register_id: ids.VirtualRegisterId, data_type: operand.DataType) operand.Source { + _ = self; + return .{ + .register = .{ .virtual = register_id }, + .type = data_type, + .region = operand.Region.contiguous(.simd8), + }; + } + + fn addRegisterLocation(self: *LoweringState, value_id: shader_ir.id.ValueId, class: operand.RegisterClass) Error!operand.Source { + const value = self.lowerer.module.values.get(value_id) orelse return Error.InvalidModule; + const data_type = try self.lowerType(value.type); + const register_id = try self.addRegister(data_type, class, value.name); + const register_source = self.registerSource(register_id, data_type); + try self.putLocation(value_id, .{ .source = register_source }); + return register_source; + } + + fn location(self: *LoweringState, value_id: shader_ir.id.ValueId) Error!ValueLocation { + if (value_id.index() >= self.value_locations.len) + return Error.InvalidModule; + + if (self.value_locations[value_id.index()]) |existing| + return existing; + + const value = self.lowerer.module.values.get(value_id) orelse return Error.InvalidModule; + switch (value.definition) { + .constant => |constant_id| { + const constant = self.lowerer.module.constants.get(constant_id) orelse return Error.InvalidModule; + + if (constant.type != value.type) + return Error.InvalidModule; + + const result: ValueLocation = if (try self.isBoolean(value.type)) switch (constant.value) { + .boolean => |boolean| .{ .predicate = .{ .constant = boolean } }, + else => return Error.UnsupportedType, + } else .{ + .source = try self.constantSource(value.type, constant.value), + }; + self.value_locations[value_id.index()] = result; + return result; + }, + .undef => { + if (try self.isBoolean(value.type)) + return Error.UnsupportedType; + _ = try self.addRegisterLocation(value_id, .temporary); + return self.value_locations[value_id.index()].?; + }, + else => return Error.InvalidModule, + } + } + + fn source(self: *LoweringState, value_id: shader_ir.id.ValueId) Error!operand.Source { + return switch (try self.location(value_id)) { + .source => |value| value, + .predicate => Error.UnsupportedType, + }; + } + + fn predicate(self: *LoweringState, value_id: shader_ir.id.ValueId) Error!PredicateValue { + return switch (try self.location(value_id)) { + .source => Error.UnsupportedType, + .predicate => |value| value, + }; + } + + fn destination(self: *LoweringState, value_id: shader_ir.id.ValueId) Error!operand.Destination { + const source_value = try self.source(value_id); + + if (source_value.negate or source_value.absolute) + return Error.InvalidLoweredProgram; + + return switch (source_value.register) { + .virtual => .{ + .register = source_value.register, + .type = source_value.type, + .region = .{ .byte_offset = source_value.region.byte_offset }, + }, + else => Error.InvalidLoweredProgram, + }; + } + + fn constantSource(self: *const LoweringState, type_id: shader_ir.id.TypeId, value: shader_ir.constant.ConstantValue) Error!operand.Source { + const data_type = try self.lowerType(type_id); + const immediate: operand.Immediate = switch (data_type) { + .u32 => switch (value) { + .integer_bits => |bits| .{ .u32 = @truncate(bits) }, + else => return Error.UnsupportedType, + }, + .i32 => switch (value) { + .integer_bits => |bits| .{ .i32 = @bitCast(@as(u32, @truncate(bits))) }, + else => return Error.UnsupportedType, + }, + .f32 => switch (value) { + .float_bits => |bits| .{ .f32 = @bitCast(@as(u32, @truncate(bits))) }, + else => return Error.UnsupportedType, + }, + else => unreachable, + }; + + return .{ + .register = .{ .immediate = immediate }, + .type = data_type, + .region = operand.Region.broadcast(), + }; + } + + fn appendInstruction(self: *LoweringState, block_id: ids.BlockId, predicate_value: ?operand.Predicate, operation: instruction.Operation) Error!void { + _ = self.program.appendInstruction(block_id, .simd8, predicate_value, operation) catch |err| + return mapProgramError(err); + } + + fn appendMove(self: *LoweringState, block_id: ids.BlockId, predicate_value: ?operand.Predicate, destination_value: operand.Destination, source_value: operand.Source) Error!void { + try self.appendInstruction(block_id, predicate_value, .{ + .move = .{ + .destination = destination_value, + .source = source_value, + }, + }); + } + + fn sourceEntryFunction(self: *const LoweringState) Error!struct { shader_ir.id.FunctionId, *const shader_ir.module.Function } { + const source_entry = self.lowerer.module.entry_point orelse return Error.MissingEntryPoint; + const function = self.lowerer.module.functions.get(source_entry) orelse return Error.InvalidEntryPoint; + const return_type = self.lowerer.module.types.get(function.return_type) orelse return Error.InvalidModule; + if (return_type.* != .void or function.parameters.items.len != 0) + return Error.InvalidEntryPoint; + return .{ source_entry, function }; + } + + fn lowerBlocks(self: *LoweringState) Error!void { + const source_function_id, const function = try self.sourceEntryFunction(); + + for (function.blocks.items) |source_block_id| { + const source_block = self.lowerer.module.blocks.get(source_block_id) orelse return Error.InvalidModule; + if (source_block.parent_function != source_function_id) + return Error.InvalidModule; + const target_block_id = self.program.addBlock(source_block.name) catch |err| + return mapProgramError(err); + if (source_block_id.index() >= self.block_map.len or self.block_map[source_block_id.index()] != null) + return Error.InvalidModule; + self.block_map[source_block_id.index()] = target_block_id; + } + + const source_entry = function.entry_block orelse return Error.InvalidModule; + self.program.setEntryBlock(try self.mappedBlock(source_entry)) catch |err| return mapProgramError(err); + } + + fn lowerParameters(self: *LoweringState) Error!void { + for (self.lowerer.module.blocks.entries.items) |entry| { + const block = entry orelse continue; + for (block.parameters.items) |parameter_id| + _ = try self.addRegisterLocation(parameter_id, .temporary); + } + } + + fn lowerInstructions(self: *LoweringState, allocator: std.mem.Allocator) Error!void { + const visited = try allocator.alloc(bool, self.lowerer.module.blocks.entries.items.len); + defer allocator.free(visited); + @memset(visited, false); + + const source_entry = try self.sourceEntryFunction(); + const function = source_entry[1]; + try self.lowerBlockInstructions(function.entry_block orelse return Error.InvalidModule, visited); + + for (function.blocks.items) |source_block_id| { + if (!visited[source_block_id.index()]) + try self.lowerBlockInstructions(source_block_id, visited); + } + } + + fn lowerBlockInstructions(self: *LoweringState, source_block_id: shader_ir.id.BlockId, visited: []bool) Error!void { + if (source_block_id.index() >= visited.len) + return Error.InvalidModule; + + if (visited[source_block_id.index()]) + return; + + visited[source_block_id.index()] = true; + + const block = self.lowerer.module.blocks.get(source_block_id) orelse return Error.InvalidModule; + const target_block_id = try self.mappedBlock(source_block_id); + for (block.instructions.items) |instruction_id| { + const source_instruction = self.lowerer.module.instructions.get(instruction_id) orelse return Error.InvalidModule; + if (source_instruction.parent_block != source_block_id) + return Error.InvalidModule; + try self.lowerInstruction(target_block_id, source_instruction.*); + } + + switch (block.terminator orelse return Error.InvalidModule) { + .branch => |edge| try self.lowerBlockInstructions(edge.target, visited), + .conditional_branch => |branch| { + try self.lowerBlockInstructions(branch.true_edge.target, visited); + try self.lowerBlockInstructions(branch.false_edge.target, visited); + }, + else => {}, + } + } + + fn lowerInstruction(self: *LoweringState, block_id: ids.BlockId, source_instruction: shader_ir.instruction.Instruction) Error!void { + switch (source_instruction.operation) { + .unary => |operation| try self.lowerUnary(block_id, source_instruction.result, operation), + .binary => |operation| try self.lowerBinary(block_id, source_instruction.result, operation), + .compare => |operation| try self.lowerCompare(block_id, source_instruction.result, operation), + .select => |operation| try self.lowerSelect(block_id, source_instruction.result, operation), + .bitcast => |value_id| try self.lowerBitcast(source_instruction.result, value_id), + .load_interface => |operation| try self.lowerLoadInterface(block_id, source_instruction.result, operation), + .store_interface => |operation| try self.lowerStoreInterface(block_id, source_instruction.result, operation), + .composite_construct, .composite_extract => return Error.UnsupportedOperation, + .call => return Error.UnsanitizedModule, + } + } + + fn requireResult(result: ?shader_ir.id.ValueId) Error!shader_ir.id.ValueId { + return result orelse Error.InvalidModule; + } + + fn requireNoResult(result: ?shader_ir.id.ValueId) Error!void { + if (result != null) + return Error.InvalidModule; + } + + fn lowerUnary(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.Unary) Error!void { + const result_id = try requireResult(result); + switch (operation.opcode) { + .logical_not => { + const source_predicate = try self.predicate(operation.operand); + const inverted: PredicateValue = switch (source_predicate) { + .constant => |value| .{ .constant = !value }, + .dynamic => |value| .{ .dynamic = .{ + .flag = value.flag, + .inverse = !value.inverse, + } }, + }; + try self.putLocation(result_id, .{ .predicate = inverted }); + }, + .negate => { + const source_value = try self.source(operation.operand); + if (source_value.type != .i32 and source_value.type != .f32) + return Error.UnsupportedOperation; + _ = try self.addRegisterLocation(result_id, .temporary); + var negated = source_value; + negated.negate = !negated.negate; + try self.appendMove(block_id, null, try self.destination(result_id), negated); + }, + .bitwise_not => { + const source_value = try self.source(operation.operand); + if (source_value.type != .u32 and source_value.type != .i32) + return Error.UnsupportedOperation; + _ = try self.addRegisterLocation(result_id, .temporary); + const all_ones: operand.Immediate = switch (source_value.type) { + .u32 => .{ .u32 = std.math.maxInt(u32) }, + .i32 => .{ .i32 = -1 }, + else => unreachable, + }; + try self.appendInstruction(block_id, null, .{ + .binary = .{ + .opcode = .bitwise_xor, + .destination = try self.destination(result_id), + .lhs = source_value, + .rhs = .{ + .register = .{ .immediate = all_ones }, + .type = source_value.type, + .region = operand.Region.broadcast(), + }, + }, + }); + }, + } + } + + fn lowerBinary(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.Binary) Error!void { + const result_id = try requireResult(result); + const lhs = try self.source(operation.lhs); + var rhs = try self.source(operation.rhs); + _ = try self.addRegisterLocation(result_id, .temporary); + const destination_value = try self.destination(result_id); + + if (lhs.type != destination_value.type or rhs.type != destination_value.type) + return Error.InvalidModule; + + const opcode: instruction.BinaryOpcode = switch (operation.opcode) { + .integer_add => if (lhs.type == .u32 or lhs.type == .i32) .add else return Error.UnsupportedOperation, + .float_add => if (lhs.type == .f32) .add else return Error.UnsupportedOperation, + + .integer_subtract => if (lhs.type == .u32 or lhs.type == .i32) subtract: { + rhs.negate = !rhs.negate; + break :subtract .add; + } else return Error.UnsupportedOperation, + + .float_subtract => if (lhs.type == .f32) subtract: { + rhs.negate = !rhs.negate; + break :subtract .add; + } else return Error.UnsupportedOperation, + + .integer_multiply => if (lhs.type == .u32 or lhs.type == .i32) .multiply else return Error.UnsupportedOperation, + .float_multiply => if (lhs.type == .f32) .multiply else return Error.UnsupportedOperation, + .shift_left => if (lhs.type == .u32 or lhs.type == .i32) .shift_left else return Error.UnsupportedOperation, + .logical_shift_right => if (lhs.type == .u32) .shift_right else return Error.UnsupportedOperation, + .arithmetic_shift_right => if (lhs.type == .i32) .shift_right else return Error.UnsupportedOperation, + .bitwise_and => if (lhs.type == .u32 or lhs.type == .i32) .bitwise_and else return Error.UnsupportedOperation, + .bitwise_or => if (lhs.type == .u32 or lhs.type == .i32) .bitwise_or else return Error.UnsupportedOperation, + .bitwise_xor => if (lhs.type == .u32 or lhs.type == .i32) .bitwise_xor else return Error.UnsupportedOperation, + .unsigned_divide, + .signed_divide, + .unsigned_modulo, + .signed_modulo, + .float_divide, + .float_modulo, + .logical_and, + .logical_or, + => return Error.UnsupportedOperation, + }; + + try self.appendInstruction(block_id, null, .{ + .binary = .{ + .opcode = opcode, + .destination = destination_value, + .lhs = lhs, + .rhs = rhs, + }, + }); + } + + fn lowerCompare(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.Compare) Error!void { + const result_id = try requireResult(result); + const result_value = self.lowerer.module.values.get(result_id) orelse return Error.InvalidModule; + + if (!try self.isBoolean(result_value.type)) + return Error.InvalidModule; + + const lhs = try self.source(operation.lhs); + const rhs = try self.source(operation.rhs); + if (lhs.type != rhs.type) + return Error.InvalidModule; + + const opcode: instruction.CompareOpcode = switch (operation.opcode) { + .equal => if (lhs.type == .u32 or lhs.type == .i32) .equal else return Error.UnsupportedOperation, + .not_equal => if (lhs.type == .u32 or lhs.type == .i32) .not_equal else return Error.UnsupportedOperation, + .unsigned_less => if (lhs.type == .u32) .less_than else return Error.UnsupportedOperation, + .signed_less => if (lhs.type == .i32) .less_than else return Error.UnsupportedOperation, + .ordered_float_equal, + .unordered_float_equal, + .ordered_float_not_equal, + .unordered_float_not_equal, + .ordered_float_less, + .unordered_float_less, + => return Error.UnsupportedOperation, + }; + + const flag_id = self.program.addVirtualFlag(.{ .name = result_value.name }) catch |err| + return mapProgramError(err); + + const predicate_value: operand.Predicate = .{ .flag = .{ .virtual = flag_id } }; + try self.putLocation(result_id, .{ .predicate = .{ .dynamic = predicate_value } }); + try self.appendInstruction(block_id, null, .{ + .compare = .{ + .opcode = opcode, + .destination = predicate_value.flag, + .lhs = lhs, + .rhs = rhs, + }, + }); + } + + fn lowerSelect(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.Select) Error!void { + const result_id = try requireResult(result); + const true_value = try self.source(operation.true_value); + const false_value = try self.source(operation.false_value); + _ = try self.addRegisterLocation(result_id, .temporary); + const destination_value = try self.destination(result_id); + + if (true_value.type != destination_value.type or false_value.type != destination_value.type) + return Error.InvalidModule; + + switch (try self.predicate(operation.condition)) { + .constant => |condition| try self.appendMove( + block_id, + null, + destination_value, + if (condition) true_value else false_value, + ), + .dynamic => |condition| { + try self.appendMove(block_id, .{ + .flag = condition.flag, + .inverse = !condition.inverse, + }, destination_value, false_value); + try self.appendMove(block_id, condition, destination_value, true_value); + }, + } + } + + fn lowerBitcast(self: *LoweringState, result: ?shader_ir.id.ValueId, source_id: shader_ir.id.ValueId) Error!void { + const result_id = try requireResult(result); + const result_value = self.lowerer.module.values.get(result_id) orelse return Error.InvalidModule; + const target_type = try self.lowerType(result_value.type); + var source_value = try self.source(source_id); + + source_value.register = switch (source_value.register) { + .immediate => |immediate| .{ .immediate = bitcastImmediate(immediate, target_type) }, + else => source_value.register, + }; + + source_value.type = target_type; + try self.putLocation(result_id, .{ .source = source_value }); + } + + fn lowerLoadInterface(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.LoadInterface) Error!void { + const result_id = try requireResult(result); + + if (operation.element_index != null) + return Error.UnsupportedOperation; + + const variable = self.lowerer.module.interface_variables.get(operation.variable) orelse return Error.InvalidModule; + + if (variable.direction != .input) + return Error.InvalidModule; + + const result_value = self.lowerer.module.values.get(result_id) orelse return Error.InvalidModule; + + if (result_value.type != variable.type) + return Error.InvalidModule; + + _ = try self.addRegisterLocation(result_id, .varying); + try self.appendInstruction(block_id, null, .{ + .load_input = .{ + .destination = try self.destination(result_id), + .semantic = try lowerInterfaceSemantic(variable.semantic), + }, + }); + } + + fn lowerStoreInterface(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.StoreInterface) Error!void { + try requireNoResult(result); + + if (operation.element_index != null) + return Error.UnsupportedOperation; + + const variable = self.lowerer.module.interface_variables.get(operation.variable) orelse return Error.InvalidModule; + + if (variable.direction != .output) + return Error.InvalidModule; + + const source_value = try self.source(operation.value); + const value = self.lowerer.module.values.get(operation.value) orelse return Error.InvalidModule; + + if (value.type != variable.type) + return Error.InvalidModule; + + try self.appendInstruction(block_id, null, .{ + .store_output = .{ + .semantic = try lowerInterfaceSemantic(variable.semantic), + .source = source_value, + }, + }); + } + + fn lowerControlAndTerminators(self: *LoweringState, allocator: std.mem.Allocator) Error!void { + const source_entry = try self.sourceEntryFunction(); + const function = source_entry[1]; + + for (function.blocks.items) |source_block_id| { + const source_block = self.lowerer.module.blocks.get(source_block_id) orelse return Error.InvalidModule; + const target_block_id = try self.mappedBlock(source_block_id); + const target_block = self.program.blocks.getMut(target_block_id) orelse return Error.InvalidLoweredProgram; + target_block.structured_control = switch (source_block.structured_control) { + .none => .none, + .selection => |selection| .{ .selection = .{ + .merge_block = try self.mappedBlock(selection.merge_block), + } }, + .loop => |loop| .{ .loop = .{ + .merge_block = try self.mappedBlock(loop.merge_block), + .continue_block = try self.mappedBlock(loop.continue_block), + } }, + }; + + const source_terminator = source_block.terminator orelse return Error.InvalidModule; + const target_terminator: instruction.Terminator = switch (source_terminator) { + .branch => |edge| branch: { + try self.lowerEdgeCopies(allocator, target_block_id, edge, null); + break :branch .{ .jump = try self.mappedBlock(edge.target) }; + }, + .conditional_branch => |branch| conditional: { + switch (try self.predicate(branch.condition)) { + .constant => |condition| { + const edge = if (condition) branch.true_edge else branch.false_edge; + try self.lowerEdgeCopies(allocator, target_block_id, edge, null); + break :conditional .{ .jump = try self.mappedBlock(edge.target) }; + }, + .dynamic => |condition| { + try self.lowerEdgeCopies(allocator, target_block_id, branch.true_edge, condition); + try self.lowerEdgeCopies(allocator, target_block_id, branch.false_edge, .{ + .flag = condition.flag, + .inverse = !condition.inverse, + }); + break :conditional .{ .conditional_branch = .{ + .predicate = condition, + .true_block = try self.mappedBlock(branch.true_edge.target), + .false_block = try self.mappedBlock(branch.false_edge.target), + } }; + }, + } + }, + .return_void => .end_thread, + .return_value => return Error.InvalidEntryPoint, + .discard => return Error.UnsupportedTerminator, + .@"unreachable" => .@"unreachable", + }; + self.program.setTerminator(target_block_id, target_terminator) catch |err| + return mapProgramError(err); + } + } + + fn lowerEdgeCopies( + self: *LoweringState, + allocator: std.mem.Allocator, + source_block: ids.BlockId, + edge: shader_ir.module.Edge, + predicate_value: ?operand.Predicate, + ) Error!void { + const target_source_block = self.lowerer.module.blocks.get(edge.target) orelse return Error.InvalidModule; + + if (edge.arguments.len != target_source_block.parameters.items.len) + return Error.InvalidModule; + + if (edge.arguments.len == 0) + return; + + const temporaries = try allocator.alloc(ids.VirtualRegisterId, edge.arguments.len); + defer allocator.free(temporaries); + + // Capture every source before writing any destination so loop backedges and + // swaps retain the parallel-copy semantics of shared-IR block arguments. + for (edge.arguments, 0..) |argument_id, index| { + const argument = try self.source(argument_id); + const temporary = try self.addRegister(argument.type, .temporary, null); + temporaries[index] = temporary; + try self.appendMove(source_block, predicate_value, .{ + .register = .{ .virtual = temporary }, + .type = argument.type, + }, argument); + } + + for (target_source_block.parameters.items, temporaries) |parameter_id, temporary| { + const destination_value = try self.destination(parameter_id); + try self.appendMove(source_block, predicate_value, destination_value, self.registerSource(temporary, destination_value.type)); + } + } +}; + +pub const Lowerer = struct { + module: *shader_ir.module.Module, + device_info: device.DeviceInfo, + options: Options, + + pub fn init(module: *shader_ir.module.Module, device_info: device.DeviceInfo, options: Options) Lowerer { + return .{ + .module = module, + .device_info = device_info, + .options = options, + }; + } + + pub fn lower(self: *Lowerer, allocator: std.mem.Allocator) Error!program_ir.Program { + // Only supports gen9 for now as it is the only gen I have access to + if (self.device_info.generation != .gen9) + return Error.UnsupportedGeneration; + + if (self.options.dispatch_width != .simd8 or !self.device_info.supportsDispatch(self.options.dispatch_width)) + return Error.UnsupportedDispatchWidth; + + shader_ir.validator.validate(self.module) catch |err| return switch (err) { + error.OutOfMemory => Error.OutOfMemory, + error.MissingEntryPoint => Error.MissingEntryPoint, + error.InvalidEntryPoint => Error.InvalidEntryPoint, + else => Error.InvalidModule, + }; + + var transformer_manager = shader_ir.transformer_manager.Manager.init(allocator); + defer transformer_manager.deinit(); + transformer_manager.add(shader_ir.inline_all_functions.transformer) catch return Error.OutOfMemory; + + var transformer_context: shader_ir.transformer_manager.Context = .{ .allocator = allocator }; + _ = transformer_manager.run(self.module, &transformer_context) catch |err| return switch (err) { + error.OutOfMemory => Error.OutOfMemory, + else => Error.SanitizationFailed, + }; + if (!self.module.properties.no_function_calls) + return Error.UnsanitizedModule; + + var program = program_ir.Program.init(allocator, self.module.stage, self.device_info, self.options.dispatch_width); + errdefer program.deinit(); + + const block_map = try allocator.alloc(?ids.BlockId, self.module.blocks.entries.items.len); + defer allocator.free(block_map); + @memset(block_map, null); + + const value_locations = try allocator.alloc(?ValueLocation, self.module.values.entries.items.len); + defer allocator.free(value_locations); + @memset(value_locations, null); + + var state: LoweringState = .{ + .lowerer = self, + .program = &program, + .block_map = block_map, + .value_locations = value_locations, + }; + + try state.lowerBlocks(); + try state.lowerParameters(); + try state.lowerInstructions(allocator); + try state.lowerControlAndTerminators(allocator); + + program.properties.instructions_selected = true; + program.properties.block_parameters_lowered = true; + validator.validate(&program) catch return Error.InvalidLoweredProgram; + return program; + } +}; + +fn lowerInterfaceSemantic(semantic: shader_ir.module.InterfaceSemantic) Error!instruction.InterfaceSemantic { + return switch (semantic) { + .location => |location| if (location.index == 0) + .{ + .location = .{ + .location = location.location, + .component = location.component, + }, + } + else + Error.UnsupportedOperation, + .builtin => |builtin| .{ + .builtin = .{ + .builtin = switch (builtin) { + .position => .position, + .vertex_index => .vertex_index, + .instance_index => .instance_index, + .frag_coord, .frag_depth, .global_invocation_id => return Error.UnsupportedOperation, + }, + }, + }, + }; +} + +fn bitcastImmediate(immediate: operand.Immediate, target_type: operand.DataType) operand.Immediate { + const bits: u32 = switch (immediate) { + .u32 => |value| value, + .i32 => |value| @bitCast(value), + .f32 => |value| @bitCast(value), + }; + return switch (target_type) { + .u32 => .{ .u32 = bits }, + .i32 => .{ .i32 = @bitCast(bits) }, + .f32 => .{ .f32 = @bitCast(bits) }, + else => unreachable, + }; +} + +fn mapProgramError(err: anyerror) Error { + return switch (err) { + Error.OutOfMemory => Error.OutOfMemory, + else => Error.InvalidLoweredProgram, + }; +} + +/// Convenience entry point for callers that do not need to retain a lowerer. +pub inline fn lower(allocator: std.mem.Allocator, module: *shader_ir.module.Module, device_info: device.DeviceInfo, options: Options) Error!program_ir.Program { + var lowerer = Lowerer.init(module, device_info, options); + return lowerer.lower(allocator); +} + +const test_device: device.DeviceInfo = .{ + .generation = .gen9, + .platform = .skylake, + .pci_device_id = 0x1912, + .grf_count = 128, +}; + +fn expectLowered(source: []const u8, expected: []const u8) !void { + var module = try shader_ir.parser.parseString(std.testing.allocator, source); + defer module.deinit(); + + var program = try lower(std.testing.allocator, &module, test_device, .{}); + defer program.deinit(); + + const actual = try printer.allocPrint(std.testing.allocator, &program); + defer std.testing.allocator.free(actual); + try std.testing.expectEqualStrings(expected, actual); +} + +fn expectLoweredFragments(source: []const u8, expected: []const []const u8, unexpected: []const []const u8) !void { + var module = try shader_ir.parser.parseString(std.testing.allocator, source); + defer module.deinit(); + + var program = try lower(std.testing.allocator, &module, test_device, .{}); + defer program.deinit(); + + const actual = try printer.allocPrint(std.testing.allocator, &program); + defer std.testing.allocator.free(actual); + + for (expected) |fragment| + try std.testing.expect(std.mem.indexOf(u8, actual, fragment) != null); + for (unexpected) |fragment| + try std.testing.expect(std.mem.indexOf(u8, actual, fragment) == null); +} + +fn expectLoweringError(source: []const u8, expected: Error) !void { + var module = try shader_ir.parser.parseString(std.testing.allocator, source); + defer module.deinit(); + + var program = lower(std.testing.allocator, &module, test_device, .{}) catch |actual| { + try std.testing.expectEqual(expected, actual); + return; + }; + defer program.deinit(); + return error.TestExpectedError; +} + +test "Lower: basic shader" { + const source = + \\shader vertex @main + \\{ + \\ @out_value: u32 = output[location(0), component(0), index(0)] + \\ %one: constant u32 = bits(0x1) + \\ %two: constant u32 = bits(0x2) + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %sum: u32 = integer_add %one, %two + \\ %condition: bool = cmp_unsigned_less %one, %two + \\ conditional_branch %condition, .left(), .right() + \\ .left(): + \\ branch .merge(%sum) + \\ .right(): + \\ branch .merge(%two) + \\ .merge(%value: u32): + \\ store_interface @out_value, %value + \\ return + \\ } + \\} + ; + + const expected = + \\; Flint program: + \\; .stage: vertex + \\; .generation: gen9 + \\; .platform: skylake + \\; .dispatch_width: simd8 + \\ + \\%value: vgrf u32[8], class(temporary), size(32), alignment(32), spillable + \\%sum: vgrf u32[8], class(temporary), size(32), alignment(32), spillable + \\%v2: vgrf u32[8], class(temporary), size(32), alignment(32), spillable + \\%v3: vgrf u32[8], class(temporary), size(32), alignment(32), spillable + \\%condition: vflag + \\ + \\.entry: + \\ [simd8] add %sum:u32, 1:u32, 2:u32 + \\ [simd8] cmp_less_than %condition, 1:u32, 2:u32 + \\ conditional_branch (+%condition), .left, .right + \\ + \\.left: + \\ [simd8] mov %v2:u32, %sum:u32 + \\ [simd8] mov %value:u32, %v2:u32 + \\ jump .merge + \\ + \\.right: + \\ [simd8] mov %v3:u32, 2:u32 + \\ [simd8] mov %value:u32, %v3:u32 + \\ jump .merge + \\ + \\.merge: + \\ [simd8] store_output location(0), component(0), %value:u32 + \\ end_thread + \\ + \\ + ; + + try expectLowered(source, expected); +} + +test "Lower: control flow" { + const source = + \\shader vertex @main + \\{ + \\ %one: constant u32 = bits(0x1) + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ branch .producer() + \\ .producer(): + \\ %sum: u32 = integer_add %one, %one + \\ branch .merge() + \\ .merge(): + \\ %doubled: u32 = integer_add %sum, %one + \\ return + \\ } + \\} + ; + + const expected = + \\; Flint program: + \\; .stage: vertex + \\; .generation: gen9 + \\; .platform: skylake + \\; .dispatch_width: simd8 + \\ + \\%sum: vgrf u32[8], class(temporary), size(32), alignment(32), spillable + \\%doubled: vgrf u32[8], class(temporary), size(32), alignment(32), spillable + \\ + \\.entry: + \\ jump .producer + \\ + \\.producer: + \\ [simd8] add %sum:u32, 1:u32, 1:u32 + \\ jump .merge + \\ + \\.merge: + \\ [simd8] add %doubled:u32, %sum:u32, 1:u32 + \\ end_thread + \\ + \\ + ; + + try expectLowered(source, expected); +} + +test "Lower: function call" { + const source = + \\shader vertex @main + \\{ + \\ %one: constant u32 = bits(0x1) + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %result: u32 = call @identity(%one) + \\ return + \\ } + \\ fn @identity(%value: u32) -> u32 + \\ { + \\ .entry(): + \\ return %value + \\ } + \\} + ; + + const expected = + \\; Flint program: + \\; .stage: vertex + \\; .generation: gen9 + \\; .platform: skylake + \\; .dispatch_width: simd8 + \\ + \\%result: vgrf u32[8], class(temporary), size(32), alignment(32), spillable + \\%v1: vgrf u32[8], class(temporary), size(32), alignment(32), spillable + \\ + \\.entry: + \\ jump .b2 + \\ + \\.b1: + \\ end_thread + \\ + \\.b2: + \\ [simd8] mov %v1:u32, 1:u32 + \\ [simd8] mov %result:u32, %v1:u32 + \\ jump .b1 + \\ + \\ + ; + + try expectLowered(source, expected); +} + +test "Lower: unary/binary operations" { + const source = + \\shader vertex @main + \\{ + \\ %u_one: constant u32 = bits(0x1) + \\ %u_two: constant u32 = bits(0x2) + \\ %i_one: constant i32 = bits(0x1) + \\ %i_two: constant i32 = bits(0x2) + \\ %f_one: constant f32 = bits(0x3f800000) + \\ %f_two: constant f32 = bits(0x40000000) + \\ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %integer_negated: i32 = negate %i_one + \\ %float_negated: f32 = negate %f_one + \\ %inverted: u32 = bitwise_not %u_one + \\ %integer_difference: i32 = integer_subtract %i_one, %i_two + \\ %float_difference: f32 = float_subtract %f_one, %f_two + \\ %integer_product: u32 = integer_multiply %u_one, %u_two + \\ %float_product: f32 = float_multiply %f_one, %f_two + \\ %shifted_left: u32 = shift_left %u_one, %u_two + \\ %logical_right: u32 = logical_shift_right %u_two, %u_one + \\ %arithmetic_right: i32 = arithmetic_shift_right %i_two, %i_one + \\ %masked: u32 = bitwise_and %u_one, %u_two + \\ %combined: u32 = bitwise_or %u_one, %u_two + \\ %toggled: u32 = bitwise_xor %u_one, %u_two + \\ return + \\ } + \\} + ; + + try expectLoweredFragments(source, &.{ + "[simd8] mov %integer_negated:i32, -1:i32", + "[simd8] mov %float_negated:f32, -1:f32", + "[simd8] bitwise_xor %inverted:u32, 1:u32, 4294967295:u32", + "[simd8] add %integer_difference:i32, 1:i32, -2:i32", + "[simd8] add %float_difference:f32, 1:f32, -2:f32", + "[simd8] multiply %integer_product:u32, 1:u32, 2:u32", + "[simd8] multiply %float_product:f32, 1:f32, 2:f32", + "[simd8] shift_left %shifted_left:u32, 1:u32, 2:u32", + "[simd8] shift_right %logical_right:u32, 2:u32, 1:u32", + "[simd8] shift_right %arithmetic_right:i32, 2:i32, 1:i32", + "[simd8] bitwise_and %masked:u32, 1:u32, 2:u32", + "[simd8] bitwise_or %combined:u32, 1:u32, 2:u32", + "[simd8] bitwise_xor %toggled:u32, 1:u32, 2:u32", + }, &.{}); +} + +test "Lower: selects and bitcasts" { + const source = + \\shader vertex @main + \\{ + \\ %always: constant bool = true + \\ %one: constant u32 = bits(0x1) + \\ %two: constant u32 = bits(0x2) + \\ %float_one: constant f32 = bits(0x3f800000) + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %condition: bool = cmp_unsigned_less %one, %two + \\ %dynamic_choice: u32 = select %condition, %one, %two + \\ %inverted_condition: bool = logical_not %condition + \\ %inverted_choice: u32 = select %inverted_condition, %one, %two + \\ %constant_choice: u32 = select %always, %one, %two + \\ %one_bits: u32 = bitcast %float_one + \\ %constant_sum: u32 = integer_add %one_bits, %one + \\ %negative: f32 = negate %float_one + \\ %negative_bits: u32 = bitcast %negative + \\ %register_sum: u32 = integer_add %negative_bits, %one + \\ return + \\ } + \\} + ; + + try expectLoweredFragments(source, &.{ + "[simd8] cmp_less_than %condition, 1:u32, 2:u32", + "[simd8] (-%condition) mov %dynamic_choice:u32, 2:u32", + "[simd8] (+%condition) mov %dynamic_choice:u32, 1:u32", + "[simd8] (+%condition) mov %inverted_choice:u32, 2:u32", + "[simd8] (-%condition) mov %inverted_choice:u32, 1:u32", + "[simd8] mov %constant_choice:u32, 1:u32", + "[simd8] add %constant_sum:u32, 1065353216:u32, 1:u32", + "[simd8] mov %negative:f32, -1:f32", + "[simd8] add %register_sum:u32, %negative:u32, 1:u32", + }, &.{ + "%one_bits: vgrf", + "%negative_bits: vgrf", + }); +} + +test "Lower: vertex interfaces" { + const source = + \\shader vertex @main + \\{ + \\ @attribute_in: u32 = input[location(2), component(1), index(0)] + \\ @vertex_id_in: u32 = input[builtin(vertex_index)] + \\ @value_out: u32 = output[location(3), component(2), index(0)] + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %attribute: u32 = load_interface @attribute_in + \\ %vertex_id: u32 = load_interface @vertex_id_in + \\ %value: u32 = integer_add %attribute, %vertex_id + \\ store_interface @value_out, %value + \\ return + \\ } + \\} + ; + + try expectLoweredFragments(source, &.{ + "%attribute: vgrf u32[8], class(varying)", + "%vertex_id: vgrf u32[8], class(varying)", + "[simd8] load_input %attribute:u32, location(2), component(1)", + "[simd8] load_input %vertex_id:u32, builtin(vertex_index), component(0)", + "[simd8] add %value:u32, %attribute:u32, %vertex_id:u32", + "[simd8] store_output location(3), component(2), %value:u32", + }, &.{}); +} + +test "Lower: constant conditional branch" { + const source = + \\shader vertex @main + \\{ + \\ %always: constant bool = true + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ conditional_branch %always, .taken(), .untaken() + \\ .taken(): + \\ return + \\ .untaken(): + \\ return + \\ } + \\} + ; + + try expectLoweredFragments(source, &.{ + ".entry:\n jump .taken", + ".taken:\n end_thread", + ".untaken:\n end_thread", + }, &.{ + "conditional_branch", + "vflag", + }); +} + +test "Lower: unsupported operations" { + try expectLoweringError( + \\shader vertex @main + \\{ + \\ %one: constant u32 = bits(0x1) + \\ %two: constant u32 = bits(0x2) + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %quotient: u32 = unsigned_divide %one, %two + \\ return + \\ } + \\} + , Error.UnsupportedOperation); + + try expectLoweringError( + \\shader vertex @main + \\{ + \\ %one: constant u32 = bits(0x1) + \\ %two: constant u32 = bits(0x2) + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %pair: vec2[u32] = composite_construct %one, %two + \\ return + \\ } + \\} + , Error.UnsupportedOperation); + + try expectLoweringError( + \\shader vertex @main + \\{ + \\ %one: constant u16 = bits(0x1) + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %sum: u16 = integer_add %one, %one + \\ return + \\ } + \\} + , Error.UnsupportedType); +} + +test "Lower: unreachable terminator" { + const source = + \\shader vertex @main + \\{ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ unreachable + \\ } + \\} + ; + + try expectLoweredFragments(source, &.{ + ".entry:\n unreachable", + }, &.{}); +} + +test "Lower: unsupported target configuration" { + var module = try shader_ir.parser.parseString(std.testing.allocator, + \\shader vertex @main + \\{ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\} + ); + defer module.deinit(); + + var gen10 = test_device; + gen10.generation = .gen10; + try std.testing.expectError(Error.UnsupportedGeneration, lower(std.testing.allocator, &module, gen10, .{})); + try std.testing.expectError(Error.UnsupportedDispatchWidth, lower(std.testing.allocator, &module, test_device, .{ .dispatch_width = .simd16 })); +} diff --git a/src/intel/compiler/root.zig b/src/intel/compiler/root.zig deleted file mode 100644 index c8db6c1..0000000 --- a/src/intel/compiler/root.zig +++ /dev/null @@ -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.?); -} diff --git a/src/intel/lib.zig b/src/intel/lib.zig index a6a6010..c376216 100644 --- a/src/intel/lib.zig +++ b/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;