[IR] adding more unit tests
This commit is contained in:
@@ -943,7 +943,7 @@ test "Parser: numeric constants" {
|
|||||||
try std.testing.expectError(Error.InvalidNumber, parseString(std.testing.allocator, out_of_range));
|
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 =
|
const source =
|
||||||
\\ shader compute @main
|
\\ shader compute @main
|
||||||
\\ {
|
\\ {
|
||||||
@@ -956,3 +956,372 @@ test "Parser: Error unknown value" {
|
|||||||
;
|
;
|
||||||
try std.testing.expectError(Error.UnknownValue, parseString(std.testing.allocator, source));
|
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
|
||||||
|
\\ }
|
||||||
|
\\}
|
||||||
|
);
|
||||||
|
}
|
||||||
|
|||||||
@@ -3,7 +3,7 @@ const ids = @import("../id.zig");
|
|||||||
const type_ir = @import("../type.zig");
|
const type_ir = @import("../type.zig");
|
||||||
const inst_ir = @import("../instruction.zig");
|
const inst_ir = @import("../instruction.zig");
|
||||||
const module_ir = @import("../module.zig");
|
const module_ir = @import("../module.zig");
|
||||||
const Builder = @import("../Builder.zig");
|
|
||||||
const dominance = @import("dominance.zig");
|
const dominance = @import("dominance.zig");
|
||||||
|
|
||||||
pub const ValidationError = error{
|
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" {
|
fn expectValidationError(expected: Error, source: []const u8) !void {
|
||||||
// shader compute @main
|
const parser = @import("../parser/parser.zig");
|
||||||
// {
|
try std.testing.expectError(expected, parser.parseString(std.testing.allocator, source));
|
||||||
// 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" {
|
test "Validator: well-typed control flow and operations" {
|
||||||
// shader compute @main
|
const parser = @import("../parser/parser.zig");
|
||||||
// {
|
|
||||||
// %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);
|
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();
|
defer module.deinit();
|
||||||
var builder = Builder.init(&module);
|
|
||||||
|
|
||||||
const void_type = try builder.internType(.void);
|
try validate(&module);
|
||||||
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(
|
test "Validator: required entry point" {
|
||||||
entry,
|
try expectValidationError(Error.MissingEntryPoint,
|
||||||
.{
|
\\shader compute
|
||||||
.conditional_branch = .{
|
\\{
|
||||||
.condition = condition,
|
\\ fn @helper() -> void
|
||||||
.true_edge = try builder.edge(left, &.{}),
|
\\ {
|
||||||
.false_edge = try builder.edge(right, &.{}),
|
\\ .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, .{
|
try expectValidationError(Error.InvalidType,
|
||||||
.binary = .{
|
\\shader compute @main
|
||||||
.opcode = .integer_add,
|
\\{
|
||||||
.lhs = one,
|
\\ fn @main(%value: array[u32, 0]) -> void
|
||||||
.rhs = one,
|
\\ {
|
||||||
},
|
\\ .entry():
|
||||||
}, null)).?;
|
\\ return
|
||||||
|
\\ }
|
||||||
|
\\}
|
||||||
|
);
|
||||||
|
}
|
||||||
|
|
||||||
try builder.setTerminator(left, .{ .branch = try builder.edge(merge, &.{}) });
|
test "Validator: reject predecessor of the entry block" {
|
||||||
try builder.setTerminator(right, .{ .branch = try builder.edge(merge, &.{}) });
|
try expectValidationError(Error.EntryBlockHasPredecessor,
|
||||||
|
\\shader compute @main
|
||||||
|
\\{
|
||||||
|
\\ fn @main() -> void
|
||||||
|
\\ {
|
||||||
|
\\ .entry():
|
||||||
|
\\ return
|
||||||
|
\\ .back_edge():
|
||||||
|
\\ branch .entry()
|
||||||
|
\\ }
|
||||||
|
\\}
|
||||||
|
);
|
||||||
|
}
|
||||||
|
|
||||||
_ = try builder.appendInstruction(merge, u32_type, .{
|
test "Validator: check branch arguments" {
|
||||||
.binary = .{
|
try expectValidationError(Error.WrongBranchArgumentCount,
|
||||||
.opcode = .integer_multiply,
|
\\shader compute @main
|
||||||
.lhs = left_value,
|
\\{
|
||||||
.rhs = one,
|
\\ fn @main() -> void
|
||||||
},
|
\\ {
|
||||||
}, null);
|
\\ .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));
|
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));
|
||||||
|
}
|
||||||
|
|||||||
@@ -155,7 +155,108 @@ pub fn copyLiteralString(allocator: anytype, words: []const u32) ![]u8 {
|
|||||||
return result;
|
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{
|
const words = [_]u32{
|
||||||
spirv.magic_number,
|
spirv.magic_number,
|
||||||
0x0001_0000,
|
0x0001_0000,
|
||||||
@@ -178,6 +279,10 @@ test "SPIR-V: parser error zero-word instruction" {
|
|||||||
try std.testing.expectError(error.TruncatedInstruction, Self.init(&truncated));
|
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 {
|
fn instructionWord(opcode: spirv.Opcode, word_count: u16) u32 {
|
||||||
return (@as(u32, word_count) << 16) | @intFromEnum(opcode);
|
return (@as(u32, word_count) << 16) | @intFromEnum(opcode);
|
||||||
}
|
}
|
||||||
|
|||||||
@@ -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);
|
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 {
|
fn assembleSpirv(allocator: std.mem.Allocator, assembly: []const u8) ![]u32 {
|
||||||
var io_backend: std.Io.Threaded = .init(allocator, .{});
|
var io_backend: std.Io.Threaded = .init(allocator, .{});
|
||||||
defer io_backend.deinit();
|
defer io_backend.deinit();
|
||||||
|
|||||||
Reference in New Issue
Block a user