From 1042b4a4226b1a0b6d0860c8be27bb193d2ae79a Mon Sep 17 00:00:00 2001 From: Kbz-8 Date: Wed, 29 Jul 2026 19:57:30 +0200 Subject: [PATCH] [IR] adding more unit tests --- src/compiler/ir/parser/parser.zig | 371 ++++++++++++++- src/compiler/ir/validator/validator.zig | 592 ++++++++++++++++++++---- src/compiler/spirv/Parser.zig | 107 ++++- src/compiler/spirv/translator.zig | 280 +++++++++++ 4 files changed, 1257 insertions(+), 93 deletions(-) diff --git a/src/compiler/ir/parser/parser.zig b/src/compiler/ir/parser/parser.zig index a5aad3c..742fe07 100644 --- a/src/compiler/ir/parser/parser.zig +++ b/src/compiler/ir/parser/parser.zig @@ -943,7 +943,7 @@ test "Parser: numeric constants" { try std.testing.expectError(Error.InvalidNumber, parseString(std.testing.allocator, out_of_range)); } -test "Parser: Error unknown value" { +test "Parser: error unknown value" { const source = \\ shader compute @main \\ { @@ -956,3 +956,372 @@ test "Parser: Error unknown value" { ; try std.testing.expectError(Error.UnknownValue, parseString(std.testing.allocator, source)); } + +fn expectParseError(expected: Error, source: []const u8) !void { + try std.testing.expectError(expected, parseString(std.testing.allocator, source)); +} + +test "Parser: infer instruction result types and reorders module declarations" { + const printer = @import("../printer.zig"); + + var module = try parseString(std.testing.allocator, + \\shader compute @main + \\{ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %sum = integer_add %late, %late + \\ %equal = cmp_equal %sum, %late + \\ return + \\ } + \\ %late: constant u32 = 1 + \\} + ); + defer module.deinit(); + + const text = try printer.allocPrint(std.testing.allocator, &module); + defer std.testing.allocator.free(text); + try std.testing.expect(std.mem.indexOf(u8, text, "%sum: u32 = integer_add %late, %late") != null); + try std.testing.expect(std.mem.indexOf(u8, text, "%equal: bool = cmp_equal %sum, %late") != null); +} + +test "Parser: check instruction result presence during lowering" { + try expectParseError(Error.InvalidResult, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ integer_add %one, %one + \\ return + \\ } + \\} + ); + + try expectParseError(Error.MissingResultType, + \\shader compute @main + \\{ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %bad = call @sink() + \\ return + \\ } + \\ fn @sink() -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\} + ); + + try expectParseError(Error.InvalidResult, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ call @identity(%one) + \\ return + \\ } + \\ fn @identity(%value: u32) -> u32 + \\ { + \\ .entry(): + \\ return %value + \\ } + \\} + ); +} + +test "Parser: report unresolved reference namespaces" { + try expectParseError(Error.UnknownFunction, + \\shader compute @missing + \\{ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\} + ); + + try expectParseError(Error.UnknownFunction, + \\shader compute @main + \\{ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ call @missing() + \\ return + \\ } + \\} + ); + + try expectParseError(Error.UnknownInterface, + \\shader compute @main + \\{ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %value = load_interface @missing + \\ return + \\ } + \\} + ); + + try expectParseError(Error.UnknownBlock, + \\shader compute @main + \\{ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ branch .missing() + \\ } + \\} + ); + + try expectParseError(Error.UnknownConstant, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ %pair: constant vec2[u32] = [#0, #9] + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\} + ); + + try expectParseError(Error.UnknownValue, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %first = integer_add %later, %one + \\ %later = integer_add %one, %one + \\ return + \\ } + \\} + ); +} + +test "Parser: reject duplicate names and values" { + try expectParseError(Error.DuplicateName, + \\shader vertex @main + \\{ + \\ @color: u32 = input[location(0), component(0), index(0)] + \\ @color: u32 = output[location(0), component(0), index(0)] + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\} + ); + + try expectParseError(Error.DuplicateName, + \\shader compute @main + \\{ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\} + ); + + try expectParseError(Error.DuplicateName, + \\shader compute @main + \\{ + \\ fn @main() -> void + \\ { + \\ .same(): + \\ return + \\ .same(): + \\ return + \\ } + \\} + ); + + try expectParseError(Error.DuplicateValue, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %result = integer_add %one, %one + \\ %result = integer_multiply %one, %one + \\ return + \\ } + \\} + ); +} + +test "Parser: reject malformed block structure and opcodes" { + try expectParseError(Error.MissingTerminator, + \\shader compute @main + \\{ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ } + \\} + ); + + try expectParseError(Error.UnexpectedToken, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ return + \\ %late = integer_add %one, %one + \\ } + \\} + ); + + try expectParseError(Error.UnexpectedToken, + \\shader compute @main + \\{ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\} + \\trailing + ); + + try expectParseError(Error.InvalidOpcode, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %bad = cmp_unknown %one, %one + \\ return + \\ } + \\} + ); +} + +test "Parser: nested composite extraction and index errors" { + const printer = @import("../printer.zig"); + + var module = try parseString(std.testing.allocator, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %inner = composite_construct %one, %one + \\ %outer = composite_construct %inner, %inner + \\ %element = composite_extract %outer[1][0] + \\ return + \\ } + \\} + ); + defer module.deinit(); + + const text = try printer.allocPrint(std.testing.allocator, &module); + defer std.testing.allocator.free(text); + try std.testing.expect(std.mem.indexOf(u8, text, "%element: u32 = composite_extract %outer[1][0]") != null); + + try expectParseError(Error.InvalidCompositeIndex, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %vector = composite_construct %one, %one + \\ %element = composite_extract %vector + \\ return + \\ } + \\} + ); + + try expectParseError(Error.InvalidCompositeIndex, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %vector = composite_construct %one, %one + \\ %element = composite_extract %vector[2] + \\ return + \\ } + \\} + ); + + try expectParseError(Error.InvalidCompositeIndex, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %element = composite_extract %one[0] + \\ return + \\ } + \\} + ); +} + +test "Parser: report invalid stage, semantics, and types" { + try expectParseError(Error.InvalidStage, + \\shader geometry + \\{ + \\} + ); + + try expectParseError(Error.InvalidSemantic, + \\shader vertex @main + \\{ + \\ @value: u32 = varying[location(0), component(0), index(0)] + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\} + ); + + try expectParseError(Error.InvalidSemantic, + \\shader vertex @main + \\{ + \\ @value: u32 = input[builtin(not_a_builtin)] + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\} + ); + + try expectParseError(Error.InvalidType, + \\shader compute @main + \\{ + \\ fn @main(%value: ptr[unknown, u32]) -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\} + ); +} diff --git a/src/compiler/ir/validator/validator.zig b/src/compiler/ir/validator/validator.zig index 8501f0e..91b5e89 100644 --- a/src/compiler/ir/validator/validator.zig +++ b/src/compiler/ir/validator/validator.zig @@ -3,7 +3,7 @@ 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{ @@ -467,107 +467,517 @@ fn targetsBlock(terminator: module_ir.Terminator, target: ids.BlockId) bool { }; } -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)); +fn expectValidationError(expected: Error, source: []const u8) !void { + const parser = @import("../parser/parser.zig"); + try std.testing.expectError(expected, parser.parseString(std.testing.allocator, source)); } -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 - // } - // } +test "Validator: well-typed control flow and operations" { + const parser = @import("../parser/parser.zig"); - var module = module_ir.Module.init(std.testing.allocator, .compute); + var module = try parser.parseString(std.testing.allocator, + \\shader vertex @main + \\{ + \\ @input: u32 = input[location(0), component(0), index(0)] + \\ @output: u32 = output[location(0), component(0), index(0)] + \\ %condition: constant bool = true + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %loaded: u32 = load_interface @input + \\ conditional_branch %condition, .left(%loaded), .right(%one) + \\ .left(%left_value: u32): + \\ %left_result: u32 = integer_add %left_value, %one + \\ branch .merge(%left_result) + \\ .right(%right_value: u32): + \\ %right_result: u32 = call @identity(%right_value) + \\ branch .merge(%right_result) + \\ .merge(%result: u32): + \\ store_interface @output, %result + \\ return + \\ } + \\ fn @identity(%value: u32) -> u32 + \\ { + \\ .entry(): + \\ return %value + \\ } + \\} + ); 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 validate(&module); +} - try builder.setTerminator( - entry, - .{ - .conditional_branch = .{ - .condition = condition, - .true_edge = try builder.edge(left, &.{}), - .false_edge = try builder.edge(right, &.{}), - }, - }, +test "Validator: required entry point" { + try expectValidationError(Error.MissingEntryPoint, + \\shader compute + \\{ + \\ fn @helper() -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\} + ); +} + +test "Validator: required function entry block" { + try expectValidationError(Error.MissingFunctionEntryBlock, + \\shader compute @main + \\{ + \\ fn @main() -> void + \\ { + \\ } + \\} + ); +} + +test "Validator: reject invalid aggregate types" { + try expectValidationError(Error.InvalidType, + \\shader compute @main + \\{ + \\ fn @main(%value: vec1[u32]) -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\} ); - const left_value = (try builder.appendInstruction(left, u32_type, .{ - .binary = .{ - .opcode = .integer_add, - .lhs = one, - .rhs = one, - }, - }, null)).?; + try expectValidationError(Error.InvalidType, + \\shader compute @main + \\{ + \\ fn @main(%value: array[u32, 0]) -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\} + ); +} - try builder.setTerminator(left, .{ .branch = try builder.edge(merge, &.{}) }); - try builder.setTerminator(right, .{ .branch = try builder.edge(merge, &.{}) }); +test "Validator: reject predecessor of the entry block" { + try expectValidationError(Error.EntryBlockHasPredecessor, + \\shader compute @main + \\{ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ return + \\ .back_edge(): + \\ branch .entry() + \\ } + \\} + ); +} - _ = try builder.appendInstruction(merge, u32_type, .{ - .binary = .{ - .opcode = .integer_multiply, - .lhs = left_value, - .rhs = one, - }, - }, null); +test "Validator: check branch arguments" { + try expectValidationError(Error.WrongBranchArgumentCount, + \\shader compute @main + \\{ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ branch .merge() + \\ .merge(%value: u32): + \\ return + \\ } + \\} + ); - try builder.setTerminator(merge, .return_void); + try expectValidationError(Error.WrongBranchArgumentType, + \\shader compute @main + \\{ + \\ %condition: constant bool = true + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ branch .merge(%condition) + \\ .merge(%value: u32): + \\ return + \\ } + \\} + ); +} + +test "Validator: check terminator operand and return types" { + try expectValidationError(Error.WrongOperandType, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ conditional_branch %one, .left(), .right() + \\ .left(): + \\ return + \\ .right(): + \\ return + \\ } + \\} + ); + + try expectValidationError(Error.WrongReturnType, + \\shader compute @main + \\{ + \\ fn @main() -> u32 + \\ { + \\ .entry(): + \\ return + \\ } + \\} + ); + + try expectValidationError(Error.WrongReturnType, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ return %one + \\ } + \\} + ); + + try expectValidationError(Error.WrongReturnType, + \\shader compute @main + \\{ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ discard + \\ } + \\} + ); +} + +test "Validator: check unary, binary, compare, and select types" { + try expectValidationError(Error.WrongResultType, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %result: bool = bitwise_not %one + \\ return + \\ } + \\} + ); + + try expectValidationError(Error.WrongOperandType, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ %float: constant f32 = 1.0 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %result: u32 = integer_add %one, %float + \\ return + \\ } + \\} + ); + + try expectValidationError(Error.WrongResultType, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %result: u32 = cmp_equal %one, %one + \\ return + \\ } + \\} + ); + + try expectValidationError(Error.WrongOperandType, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %result: u32 = select %one, %one, %one + \\ return + \\ } + \\} + ); + + try expectValidationError(Error.WrongOperandType, + \\shader compute @main + \\{ + \\ %condition: constant bool = true + \\ %one: constant u32 = 1 + \\ %float: constant f32 = 1.0 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %result: u32 = select %condition, %one, %float + \\ return + \\ } + \\} + ); +} + +test "Validator: check composite operations" { + try expectValidationError(Error.WrongOperandType, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %result: vec2[u32] = composite_construct %one + \\ return + \\ } + \\} + ); + + try expectValidationError(Error.WrongResultType, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %vector: vec2[u32] = composite_construct %one, %one + \\ %result: bool = composite_extract %vector[0] + \\ return + \\ } + \\} + ); +} + +test "Validator: check interface direction and value types" { + try expectValidationError(Error.WrongInterfaceDirection, + \\shader vertex @main + \\{ + \\ @output: u32 = output[location(0), component(0), index(0)] + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %value: u32 = load_interface @output + \\ return + \\ } + \\} + ); + + try expectValidationError(Error.WrongInterfaceDirection, + \\shader vertex @main + \\{ + \\ @input: u32 = input[location(0), component(0), index(0)] + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ store_interface @input, %one + \\ return + \\ } + \\} + ); + + try expectValidationError(Error.WrongOperandType, + \\shader vertex @main + \\{ + \\ @output: u32 = output[location(0), component(0), index(0)] + \\ %value: constant f32 = 1.0 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ store_interface @output, %value + \\ return + \\ } + \\} + ); + + try expectValidationError(Error.WrongResultPresence, + \\shader vertex @main + \\{ + \\ @output: u32 = output[location(0), component(0), index(0)] + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %result: bool = store_interface @output, %one + \\ return + \\ } + \\} + ); +} + +test "Validator: check function calls" { + try expectValidationError(Error.WrongOperandType, + \\shader compute @main + \\{ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %result: u32 = call @identity() + \\ return + \\ } + \\ fn @identity(%value: u32) -> u32 + \\ { + \\ .entry(): + \\ return %value + \\ } + \\} + ); + + try expectValidationError(Error.WrongOperandType, + \\shader compute @main + \\{ + \\ %condition: constant bool = true + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %result: u32 = call @identity(%condition) + \\ return + \\ } + \\ fn @identity(%value: u32) -> u32 + \\ { + \\ .entry(): + \\ return %value + \\ } + \\} + ); + + try expectValidationError(Error.WrongResultType, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %result: f32 = call @identity(%one) + \\ return + \\ } + \\ fn @identity(%value: u32) -> u32 + \\ { + \\ .entry(): + \\ return %value + \\ } + \\} + ); + + try expectValidationError(Error.WrongResultPresence, + \\shader compute @main + \\{ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %result: u32 = call @helper() + \\ return + \\ } + \\ fn @helper() -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\} + ); +} + +test "Validator: reject cross-function value references" { + try expectValidationError(Error.CrossFunctionReference, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %main_value: u32 = integer_add %one, %one + \\ return + \\ } + \\ fn @helper() -> void + \\ { + \\ .entry(): + \\ %result: u32 = integer_add %main_value, %one + \\ return + \\ } + \\} + ); +} + +test "Validator: reject an SSA definition that does not dominate its use" { + try expectValidationError(Error.DefinitionDoesNotDominateUse, + \\shader compute @main + \\{ + \\ %condition: constant bool = true + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ conditional_branch %condition, .left(), .right() + \\ .left(): + \\ %left_value: u32 = integer_add %one, %one + \\ branch .merge() + \\ .right(): + \\ branch .merge() + \\ .merge(): + \\ %result: u32 = integer_multiply %left_value, %one + \\ return + \\ } + \\} + ); +} + +test "Validator: reject a same-block use before its definition" { + const parser = @import("../parser/parser.zig"); + + var module = try parser.parseString(std.testing.allocator, + \\shader compute @main + \\{ + \\ %one: constant u32 = 1 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %first: u32 = integer_add %one, %one + \\ %second: u32 = integer_multiply %first, %one + \\ return + \\ } + \\} + ); + defer module.deinit(); + + const function = module.functions.get(module.entry_point.?).?; + const block = module.blocks.get(function.entry_block.?).?; + const first_instruction = block.instructions.items[0]; + const second_instruction = block.instructions.items[1]; + const later_result = module.instructions.get(second_instruction).?.result.?; + module.instructions.getMut(first_instruction).?.operation.binary.lhs = later_result; try std.testing.expectError(Error.DefinitionDoesNotDominateUse, validate(&module)); } + +test "Validator: reject a structured-control target in another function" { + const parser = @import("../parser/parser.zig"); + + var module = try parser.parseString(std.testing.allocator, + \\shader compute @main + \\{ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\ fn @helper() -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\} + ); + defer module.deinit(); + + const main = module.functions.get(module.entry_point.?).?; + const helper = module.functions.get(ids.FunctionId.fromIndex(1)).?; + module.blocks.getMut(main.entry_block.?).?.structured_control = .{ + .selection = .{ .merge_block = helper.entry_block.? }, + }; + + try std.testing.expectError(Error.InvalidStructuredControl, validate(&module)); +} diff --git a/src/compiler/spirv/Parser.zig b/src/compiler/spirv/Parser.zig index 88056f9..a338758 100644 --- a/src/compiler/spirv/Parser.zig +++ b/src/compiler/spirv/Parser.zig @@ -155,7 +155,108 @@ pub fn copyLiteralString(allocator: anytype, words: []const u32) ![]u8 { return result; } -test "SPIR-V: parser error zero-word instruction" { +test "SPIR-V: parser validates module headers" { + const short = [_]u32{ spirv.magic_number, 0x0001_0000, 0, 1 }; + try std.testing.expectError(error.HeaderTooShort, Self.init(&short)); + + var invalid_magic = validHeader(0x0001_0000); + invalid_magic[0] = 0x1234_5678; + try std.testing.expectError(error.InvalidMagic, Self.init(&invalid_magic)); + + var byte_swapped = validHeader(0x0001_0000); + byte_swapped[0] = spirv.byte_swapped_magic_number; + try std.testing.expectError(error.ByteSwappedModule, Self.init(&byte_swapped)); + + var invalid_major = validHeader(0x0002_0000); + try std.testing.expectError(error.InvalidVersion, Self.init(&invalid_major)); + + var invalid_minor = validHeader(0x0001_0700); + try std.testing.expectError(error.InvalidVersion, Self.init(&invalid_minor)); + + var invalid_reserved_bits = validHeader(0x0101_0001); + try std.testing.expectError(error.InvalidVersion, Self.init(&invalid_reserved_bits)); + + var zero_bound = validHeader(0x0001_0000); + zero_bound[3] = 0; + try std.testing.expectError(error.InvalidIdBound, Self.init(&zero_bound)); + + var nonzero_schema = validHeader(0x0001_0000); + nonzero_schema[4] = 1; + try std.testing.expectError(error.InvalidSchema, Self.init(&nonzero_schema)); + + var version_1_6 = validHeader(0x0001_0600); + version_1_6[2] = 0xfeed_beef; + version_1_6[3] = 42; + const parser = try Self.init(&version_1_6); + try std.testing.expectEqual(@as(u8, 1), parser.header.major()); + try std.testing.expectEqual(@as(u8, 6), parser.header.minor()); + try std.testing.expectEqual(@as(u32, 0xfeed_beef), parser.header.generator); + try std.testing.expectEqual(@as(u32, 42), parser.header.bound); + + var instruction_iterator = parser.iterator(); + try std.testing.expectEqual(@as(?Instruction, null), try instruction_iterator.next()); +} + +test "SPIR-V: parser iterates instructions and operands" { + const words = [_]u32{ + spirv.magic_number, + 0x0001_0000, + 0, + 8, + 0, + instructionWord(.nop, 1), + instructionWord(.i_add, 5), + 1, + 2, + 3, + 4, + }; + const parser = try Self.init(&words); + var instruction_iterator = parser.iterator(); + + const nop = (try instruction_iterator.next()).?; + try std.testing.expectEqual(spirv.Opcode.nop, nop.opcode); + try std.testing.expectEqual(@as(usize, spirv.header_word_count), nop.word_offset); + try std.testing.expectEqual(@as(usize, 0), nop.operands.len); + try std.testing.expectEqual(@as(?u32, null), nop.operand(0)); + + const add = (try instruction_iterator.next()).?; + try std.testing.expectEqual(spirv.Opcode.i_add, add.opcode); + try std.testing.expectEqual(@as(usize, spirv.header_word_count + 1), add.word_offset); + try std.testing.expectEqualSlices(u32, &.{ 1, 2, 3, 4 }, add.operands); + try std.testing.expectEqual(@as(?u32, 1), add.operand(0)); + try std.testing.expectEqual(@as(?u32, 4), add.operand(3)); + try std.testing.expectEqual(@as(?u32, null), add.operand(4)); + try std.testing.expectEqual(@as(?Instruction, null), try instruction_iterator.next()); +} + +test "SPIR-V: parser literal string helpers" { + const empty = [_]u32{0}; + try std.testing.expectEqual(@as(usize, 1), try literalStringWordCount(&empty)); + try std.testing.expect(try literalStringEquals(&empty, "")); + + const abc = [_]u32{0x0063_6261}; + try std.testing.expectEqual(@as(usize, 1), try literalStringWordCount(&abc)); + try std.testing.expect(try literalStringEquals(&abc, "abc")); + try std.testing.expect(!try literalStringEquals(&abc, "ab")); + try std.testing.expect(!try literalStringEquals(&abc, "abcd")); + + const main = [_]u32{ 0x6e69_616d, 0 }; + try std.testing.expectEqual(@as(usize, 2), try literalStringWordCount(&main)); + try std.testing.expect(try literalStringEquals(&main, "main")); + try std.testing.expect(!try literalStringEquals(&main, "Main")); + + const copy = try copyLiteralString(std.testing.allocator, &main); + defer std.testing.allocator.free(copy); + try std.testing.expectEqualStrings("main", copy); + + const unterminated = [_]u32{0x6463_6261}; + try std.testing.expectError(error.UnterminatedString, literalStringWordCount(&unterminated)); + try std.testing.expectError(error.UnterminatedString, literalStringEquals(&unterminated, "abcd")); + try std.testing.expectError(error.UnterminatedString, copyLiteralString(std.testing.allocator, &unterminated)); +} + +test "SPIR-V: parser rejects malformed instruction framing" { const words = [_]u32{ spirv.magic_number, 0x0001_0000, @@ -178,6 +279,10 @@ test "SPIR-V: parser error zero-word instruction" { try std.testing.expectError(error.TruncatedInstruction, Self.init(&truncated)); } +fn validHeader(version: u32) [spirv.header_word_count]u32 { + return .{ spirv.magic_number, version, 0, 1, 0 }; +} + fn instructionWord(opcode: spirv.Opcode, word_count: u16) u32 { return (@as(u32, word_count) << 16) | @intFromEnum(opcode); } diff --git a/src/compiler/spirv/translator.zig b/src/compiler/spirv/translator.zig index 463c01a..77441b9 100644 --- a/src/compiler/spirv/translator.zig +++ b/src/compiler/spirv/translator.zig @@ -1195,6 +1195,286 @@ test "SPIR-V: decorated vertex interfaces and load-store operations" { try std.testing.expect(std.mem.indexOf(u8, text, "store_interface @out_color") != null); } +test "SPIR-V: fragment execution modes and translated properties" { + const assembly = + \\OpCapability Shader + \\OpMemoryModel Logical GLSL450 + \\OpEntryPoint Fragment %main "main" + \\OpExecutionMode %main OriginUpperLeft + \\OpExecutionMode %main EarlyFragmentTests + \\%void = OpTypeVoid + \\%fn_void = OpTypeFunction %void + \\%main = OpFunction %void None %fn_void + \\ %entry = OpLabel + \\ 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.fragment, module.stage); + try std.testing.expect(module.execution_modes.early_fragment_tests); + try std.testing.expectEqual(@as(?[3]u32, null), module.execution_modes.workgroup_size); + try std.testing.expect(module.properties.valid_cfg); + try std.testing.expect(module.properties.valid_ssa); + try std.testing.expect(module.properties.structured_control_flow); + try std.testing.expect(module.properties.no_function_calls); + + const function = module.functions.get(module.entry_point.?).?; + try std.testing.expectEqualStrings("main", function.name.?); + try std.testing.expectEqual(@as(usize, 1), function.blocks.items.len); + const entry = module.blocks.get(function.entry_block.?).?; + try std.testing.expect(entry.terminator.? == .return_void); +} + +test "SPIR-V: entry point lookup errors" { + const single_entry_assembly = + \\OpCapability Shader + \\OpMemoryModel Logical GLSL450 + \\OpEntryPoint GLCompute %main "main" + \\OpExecutionMode %main LocalSize 1 1 1 + \\%void = OpTypeVoid + \\%fn_void = OpTypeFunction %void + \\%main = OpFunction %void None %fn_void + \\ %entry = OpLabel + \\ OpReturn + \\OpFunctionEnd + ; + const single_entry_words = try assembleSpirv(std.testing.allocator, single_entry_assembly); + defer std.testing.allocator.free(single_entry_words); + try std.testing.expectError(error.EntryPointNotFound, translate(std.testing.allocator, single_entry_words, .{ .entry_point = "missing" })); + + const ambiguous_assembly = + \\OpCapability Shader + \\OpMemoryModel Logical GLSL450 + \\OpEntryPoint GLCompute %first "main" + \\OpEntryPoint GLCompute %second "main" + \\%void = OpTypeVoid + \\%fn_void = OpTypeFunction %void + \\%first = OpFunction %void None %fn_void + \\ %first_entry = OpLabel + \\ OpReturn + \\OpFunctionEnd + \\%second = OpFunction %void None %fn_void + \\ %second_entry = OpLabel + \\ OpReturn + \\OpFunctionEnd + ; + const ambiguous_words = try assembleSpirv(std.testing.allocator, ambiguous_assembly); + defer std.testing.allocator.free(ambiguous_words); + try std.testing.expectError(error.AmbiguousEntryPoint, translate(std.testing.allocator, ambiguous_words, .{ .entry_point = "main" })); + + const unsupported_assembly = + \\OpCapability Shader + \\OpCapability Geometry + \\OpMemoryModel Logical GLSL450 + \\OpEntryPoint Geometry %main "main" + \\%void = OpTypeVoid + \\%fn_void = OpTypeFunction %void + \\%main = OpFunction %void None %fn_void + \\ %entry = OpLabel + \\ OpReturn + \\OpFunctionEnd + ; + const unsupported_words = try assembleSpirv(std.testing.allocator, unsupported_assembly); + defer std.testing.allocator.free(unsupported_words); + try std.testing.expectError(error.UnsupportedExecutionModel, translate(std.testing.allocator, unsupported_words, .{ .entry_point = "main" })); +} + +test "SPIR-V: operation mappings to backend-agnostic IR" { + const assembly = + \\OpCapability Shader + \\OpMemoryModel Logical GLSL450 + \\OpEntryPoint GLCompute %main "main" + \\OpExecutionMode %main LocalSize 1 1 1 + \\%void = OpTypeVoid + \\%bool = OpTypeBool + \\%uint = OpTypeInt 32 0 + \\%float = OpTypeFloat 32 + \\%vec2 = OpTypeVector %uint 2 + \\%fn_void = OpTypeFunction %void + \\%true = OpConstantTrue %bool + \\%one = OpConstant %uint 1 + \\%two = OpConstant %uint 2 + \\%main = OpFunction %void None %fn_void + \\ %entry = OpLabel + \\ %not = OpLogicalNot %bool %true + \\ %sum = OpIAdd %uint %one %two + \\ %less = OpULessThan %bool %one %two + \\ %selected = OpSelect %uint %less %one %two + \\ %cast = OpBitcast %float %one + \\ %vector = OpCompositeConstruct %vec2 %one %two + \\ %element = OpCompositeExtract %uint %vector 1 + \\ 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(); + + const function = module.functions.get(module.entry_point.?).?; + const block = module.blocks.get(function.entry_block.?).?; + try std.testing.expectEqual(@as(usize, 7), block.instructions.items.len); + + const logical_not = module.instructions.get(block.instructions.items[0]).?; + try std.testing.expectEqual(ir.instruction.UnaryOpcode.logical_not, logical_not.operation.unary.opcode); + + const add = module.instructions.get(block.instructions.items[1]).?; + try std.testing.expectEqual(ir.instruction.BinaryOpcode.integer_add, add.operation.binary.opcode); + + const less = module.instructions.get(block.instructions.items[2]).?; + try std.testing.expectEqual(ir.instruction.CompareOpcode.unsigned_less, less.operation.compare.opcode); + + const select = module.instructions.get(block.instructions.items[3]).?; + try std.testing.expect(select.operation == .select); + + const bitcast = module.instructions.get(block.instructions.items[4]).?; + try std.testing.expect(bitcast.operation == .bitcast); + + const construct = module.instructions.get(block.instructions.items[5]).?; + try std.testing.expectEqual(@as(usize, 2), construct.operation.composite_construct.elements.len); + + const extract = module.instructions.get(block.instructions.items[6]).?; + try std.testing.expectEqualSlices(u32, &.{1}, extract.operation.composite_extract.indices); +} + +test "SPIR-V: structured loop and OpPhi back edge" { + const assembly = + \\OpCapability Shader + \\OpMemoryModel Logical GLSL450 + \\OpEntryPoint GLCompute %main "main" + \\OpExecutionMode %main LocalSize 1 1 1 + \\OpName %entry "entry" + \\OpName %header "header" + \\OpName %body "body" + \\OpName %continue "continue" + \\OpName %merge "merge" + \\%void = OpTypeVoid + \\%bool = OpTypeBool + \\%uint = OpTypeInt 32 0 + \\%fn_void = OpTypeFunction %void + \\%true = OpConstantTrue %bool + \\%zero = OpConstant %uint 0 + \\%one = OpConstant %uint 1 + \\%main = OpFunction %void None %fn_void + \\ %entry = OpLabel + \\ OpBranch %header + \\ %header = OpLabel + \\ %index = OpPhi %uint %zero %entry %next %continue + \\ OpLoopMerge %merge %continue None + \\ OpBranchConditional %true %body %merge + \\ %body = OpLabel + \\ OpBranch %continue + \\ %continue = OpLabel + \\ %next = OpIAdd %uint %index %one + \\ OpBranch %header + \\ %merge = OpLabel + \\ 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(); + + const function = module.functions.get(module.entry_point.?).?; + try std.testing.expectEqual(@as(usize, 5), function.blocks.items.len); + const entry_id = function.blocks.items[0]; + const header_id = function.blocks.items[1]; + const continue_id = function.blocks.items[3]; + const merge_id = function.blocks.items[4]; + + const entry = module.blocks.get(entry_id).?; + try std.testing.expectEqual(@as(usize, 1), entry.terminator.?.branch.arguments.len); + + const header = module.blocks.get(header_id).?; + try std.testing.expectEqual(@as(usize, 1), header.parameters.items.len); + try std.testing.expect(header.structured_control == .loop); + try std.testing.expectEqual(merge_id, header.structured_control.loop.merge_block); + try std.testing.expectEqual(continue_id, header.structured_control.loop.continue_block); + + const continue_block = module.blocks.get(continue_id).?; + try std.testing.expectEqual(header_id, continue_block.terminator.?.branch.target); + try std.testing.expectEqual(@as(usize, 1), continue_block.terminator.?.branch.arguments.len); +} + +test "SPIR-V: rejects a missing OpPhi incoming value" { + const assembly = + \\OpCapability Shader + \\OpMemoryModel Logical GLSL450 + \\OpEntryPoint GLCompute %main "main" + \\OpExecutionMode %main LocalSize 1 1 1 + \\%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 + \\ OpBranchConditional %true %left %right + \\ %left = OpLabel + \\ OpBranch %merge + \\ %right = OpLabel + \\ OpBranch %merge + \\ %merge = OpLabel + \\ %value = OpPhi %uint %one %left + \\ OpReturn + \\OpFunctionEnd + ; + const words = try assembleSpirv(std.testing.allocator, assembly); + defer std.testing.allocator.free(words); + + try std.testing.expectError(error.MissingPhiIncomingValue, translate(std.testing.allocator, words, .{ .entry_point = "main" })); +} + +test "SPIR-V: preserves location components and builtin interfaces" { + const assembly = + \\OpCapability Shader + \\OpMemoryModel Logical GLSL450 + \\OpEntryPoint Vertex %main "main" %input_value %position + \\OpDecorate %input_value Location 3 + \\OpDecorate %input_value Component 2 + \\OpDecorate %input_value Index 1 + \\OpDecorate %position BuiltIn Position + \\%void = OpTypeVoid + \\%float = OpTypeFloat 32 + \\%vec4 = OpTypeVector %float 4 + \\%input_vec4 = OpTypePointer Input %vec4 + \\%output_vec4 = OpTypePointer Output %vec4 + \\%fn_void = OpTypeFunction %void + \\%input_value = OpVariable %input_vec4 Input + \\%position = OpVariable %output_vec4 Output + \\%main = OpFunction %void None %fn_void + \\ %entry = OpLabel + \\ 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(); + + const input = module.interface_variables.get(ir.id.InterfaceVariableId.fromIndex(0)).?; + try std.testing.expectEqual(ir.module.InterfaceDirection.input, input.direction); + try std.testing.expect(input.semantic == .location); + try std.testing.expectEqual(@as(u32, 3), input.semantic.location.location); + try std.testing.expectEqual(@as(u8, 2), input.semantic.location.component); + try std.testing.expectEqual(@as(u8, 1), input.semantic.location.index); + + const position = module.interface_variables.get(ir.id.InterfaceVariableId.fromIndex(1)).?; + try std.testing.expectEqual(ir.module.InterfaceDirection.output, position.direction); + try std.testing.expect(position.semantic == .builtin); + try std.testing.expectEqual(ir.module.Builtin.position, position.semantic.builtin); +} + fn assembleSpirv(allocator: std.mem.Allocator, assembly: []const u8) ![]u32 { var io_backend: std.Io.Threaded = .init(allocator, .{}); defer io_backend.deinit();