diff --git a/src/software/interpreter/Program.zig b/src/software/interpreter/Program.zig index 0d3fe6c..11db471 100644 --- a/src/software/interpreter/Program.zig +++ b/src/software/interpreter/Program.zig @@ -43,6 +43,7 @@ stage: module_ir.Stage, entry_pc: u32, register_count: usize, scratch_count: usize, +array_lengths: []const bc.ArrayLength, code: []const bc.Instruction, edges: []const bc.Edge, copies: []const bc.Copy, @@ -66,6 +67,7 @@ pub fn compile(backing_allocator: std.mem.Allocator, module: *const module_ir.Mo .entry_pc = lowerer.entry_pc, .register_count = lowerer.register_count, .scratch_count = lowerer.scratch_count, + .array_lengths = lowerer.array_lengths.items, .code = lowerer.code.items, .edges = lowerer.edges.items, .copies = lowerer.copies.items, @@ -107,6 +109,7 @@ const Lowerer = struct { register_count: usize = 0, scratch_count: usize = 0, entry_pc: u32 = 0, + array_lengths: std.ArrayList(bc.ArrayLength) = .empty, code: std.ArrayList(bc.Instruction) = .empty, edges: std.ArrayList(bc.Edge) = .empty, copies: std.ArrayList(bc.Copy) = .empty, @@ -387,7 +390,22 @@ const Lowerer = struct { try self.emit(.store_buffer, src.components, src.base, byte_offset, bc.invalid_register, bc.invalid_register, @intFromEnum(op.resource)); }, .call => return CompileError.UnsupportedOperation, - .array_length => return CompileError.UnsupportedOperation, + .array_length => |op| { + const dst = result orelse return CompileError.InvalidOperation; + const byte_offset = try self.bufferOffset(op.byte_offset); + _ = try self.storageBuffer(op.resource); + + if (dst.components != 1 or dst.kind != .unsigned_integer) + return CompileError.InvalidOperation; + + const metadata_index = try u32Index(self.array_lengths.items.len); + try self.array_lengths.append(self.allocator, .{ + .resource = @intFromEnum(op.resource), + .stride = op.stride, + }); + + try self.emit(.array_length, 1, dst.base, byte_offset, bc.invalid_register, bc.invalid_register, metadata_index); + }, } } diff --git a/src/software/interpreter/Runtime.zig b/src/software/interpreter/Runtime.zig index e108964..430a778 100644 --- a/src/software/interpreter/Runtime.zig +++ b/src/software/interpreter/Runtime.zig @@ -98,45 +98,13 @@ pub fn run(self: *Self, program: *const Program, options: RunOptions) RuntimeErr pc += 1; switch (instruction.opcode) { - .copy => self.copy(instruction), - .negate_i32 => self.unaryInt(instruction, .negate), - .negate_f32 => self.unaryFloat(instruction), - .logical_not => self.unaryInt(instruction, .logical_not), - .bitwise_not => self.unaryInt(instruction, .bitwise_not), - .integer_add => try self.binaryInt(instruction, .add), - .integer_subtract => try self.binaryInt(instruction, .subtract), - .integer_multiply => try self.binaryInt(instruction, .multiply), - .unsigned_divide => try self.binaryInt(instruction, .unsigned_divide), - .signed_divide => try self.binaryInt(instruction, .signed_divide), - .unsigned_modulo => try self.binaryInt(instruction, .unsigned_modulo), - .signed_modulo => try self.binaryInt(instruction, .signed_modulo), - .shift_left => try self.binaryInt(instruction, .shift_left), - .logical_shift_right => try self.binaryInt(instruction, .logical_shift_right), + .@"unreachable" => return RuntimeError.UnreachableExecuted, + .array_length => try self.arrayLength(program, options.resource_buffers, instruction), .arithmetic_shift_right => try self.binaryInt(instruction, .arithmetic_shift_right), .bitwise_and => try self.binaryInt(instruction, .bitwise_and), + .bitwise_not => self.unaryInt(instruction, .bitwise_not), .bitwise_or => try self.binaryInt(instruction, .bitwise_or), .bitwise_xor => try self.binaryInt(instruction, .bitwise_xor), - .logical_and => try self.binaryInt(instruction, .logical_and), - .logical_or => try self.binaryInt(instruction, .logical_or), - .float_add => self.binaryFloat(instruction, .add), - .float_subtract => self.binaryFloat(instruction, .subtract), - .float_multiply => self.binaryFloat(instruction, .multiply), - .float_divide => self.binaryFloat(instruction, .divide), - .float_modulo => self.binaryFloat(instruction, .modulo), - .compare_equal => self.compareInt(instruction, .equal), - .compare_not_equal => self.compareInt(instruction, .not_equal), - .compare_unsigned_less => self.compareInt(instruction, .unsigned_less), - .compare_signed_less => self.compareInt(instruction, .signed_less), - .compare_ordered_float_equal => self.compareFloat(instruction, .ordered_equal), - .compare_unordered_float_equal => self.compareFloat(instruction, .unordered_equal), - .compare_ordered_float_not_equal => self.compareFloat(instruction, .ordered_not_equal), - .compare_unordered_float_not_equal => self.compareFloat(instruction, .unordered_not_equal), - .compare_ordered_float_less => self.compareFloat(instruction, .ordered_less), - .compare_unordered_float_less => self.compareFloat(instruction, .unordered_less), - .select => self.select(instruction), - .load_buffer => try self.loadBuffer(program, options.resource_buffers, instruction), - .store_buffer => try self.storeBuffer(program, options.resource_buffers, instruction), - .jump_edge => pc = try self.applyEdge(program, instruction.immediate), .branch => { if (instruction.immediate >= program.branches.len) return RuntimeError.InvalidBytecode; @@ -144,9 +112,42 @@ pub fn run(self: *Self, program: *const Program, options: RunOptions) RuntimeErr const branch = program.branches[instruction.immediate]; pc = try self.applyEdge(program, if (self.registers[instruction.a] != 0) branch.true_edge else branch.false_edge); }, - .return_void => return .returned, + .compare_equal => self.compareInt(instruction, .equal), + .compare_not_equal => self.compareInt(instruction, .not_equal), + .compare_ordered_float_equal => self.compareFloat(instruction, .ordered_equal), + .compare_ordered_float_less => self.compareFloat(instruction, .ordered_less), + .compare_ordered_float_not_equal => self.compareFloat(instruction, .ordered_not_equal), + .compare_signed_less => self.compareInt(instruction, .signed_less), + .compare_unordered_float_equal => self.compareFloat(instruction, .unordered_equal), + .compare_unordered_float_less => self.compareFloat(instruction, .unordered_less), + .compare_unordered_float_not_equal => self.compareFloat(instruction, .unordered_not_equal), + .compare_unsigned_less => self.compareInt(instruction, .unsigned_less), + .copy => self.copy(instruction), .discard => return .discarded, - .@"unreachable" => return RuntimeError.UnreachableExecuted, + .float_add => self.binaryFloat(instruction, .add), + .float_divide => self.binaryFloat(instruction, .divide), + .float_modulo => self.binaryFloat(instruction, .modulo), + .float_multiply => self.binaryFloat(instruction, .multiply), + .float_subtract => self.binaryFloat(instruction, .subtract), + .integer_add => try self.binaryInt(instruction, .add), + .integer_multiply => try self.binaryInt(instruction, .multiply), + .integer_subtract => try self.binaryInt(instruction, .subtract), + .jump_edge => pc = try self.applyEdge(program, instruction.immediate), + .load_buffer => try self.loadBuffer(program, options.resource_buffers, instruction), + .logical_and => try self.binaryInt(instruction, .logical_and), + .logical_not => self.unaryInt(instruction, .logical_not), + .logical_or => try self.binaryInt(instruction, .logical_or), + .logical_shift_right => try self.binaryInt(instruction, .logical_shift_right), + .negate_f32 => self.unaryFloat(instruction), + .negate_i32 => self.unaryInt(instruction, .negate), + .return_void => return .returned, + .select => self.select(instruction), + .shift_left => try self.binaryInt(instruction, .shift_left), + .signed_divide => try self.binaryInt(instruction, .signed_divide), + .signed_modulo => try self.binaryInt(instruction, .signed_modulo), + .store_buffer => try self.storeBuffer(program, options.resource_buffers, instruction), + .unsigned_divide => try self.binaryInt(instruction, .unsigned_divide), + .unsigned_modulo => try self.binaryInt(instruction, .unsigned_modulo), } } } @@ -173,6 +174,33 @@ const BinaryFloat = enum { add, subtract, multiply, divide, modulo }; const CompareInt = enum { equal, not_equal, unsigned_less, signed_less }; const CompareFloat = enum { ordered_equal, unordered_equal, ordered_not_equal, unordered_not_equal, ordered_less, unordered_less }; +fn arrayLength(self: *Self, program: *const Program, resource_buffers: []const ?[]u8, instruction: bc.Instruction) RuntimeError!void { + if (instruction.components != 1) + return RuntimeError.InvalidBytecode; + + if (instruction.a >= self.registers.len or instruction.b >= self.registers.len) + return RuntimeError.InvalidBytecode; + + if (instruction.immediate >= program.array_lengths.len) + return RuntimeError.InvalidBytecode; + + const metadata = program.array_lengths[instruction.immediate]; + + if (metadata.stride == 0) + return RuntimeError.InvalidBytecode; + + const buffer = try resourceBuffer(program, resource_buffers, metadata.resource); + const byte_offset: usize = self.registers[instruction.b]; + + if (byte_offset > buffer.len) + return RuntimeError.BufferOutOfBounds; + + const byte_length = buffer.len - byte_offset; + const element_count = byte_length / metadata.stride; + + self.registers[instruction.a] = std.math.cast(u32, element_count) orelse return RuntimeError.IntegerOverflow; +} + fn copy(self: *Self, instruction: bc.Instruction) void { for (0..instruction.components) |component| self.registers[@as(usize, instruction.a) + component] = self.registers[@as(usize, instruction.b) + component]; diff --git a/src/software/interpreter/Shader.zig b/src/software/interpreter/Shader.zig index 33b1c0f..74f7efc 100644 --- a/src/software/interpreter/Shader.zig +++ b/src/software/interpreter/Shader.zig @@ -54,9 +54,7 @@ pub fn compile(allocator: std.mem.Allocator, module: *SoftShaderModule, stage: * }; errdefer program.deinit(); - if (!hasCompatibleInterface(&program, expected_stage) or - (expected_stage == .compute and module_ir.execution_modes.workgroup_size == null)) - { + if (!hasCompatibleInterface(&program, expected_stage) or (expected_stage == .compute and module_ir.execution_modes.workgroup_size == null)) { std.log.scoped(.IrInterpreter).err("unsupported stage interface or execution modes", .{}); return VkError.ValidationFailed; } diff --git a/src/software/interpreter/bytecode.zig b/src/software/interpreter/bytecode.zig index 8d59754..09cc4e2 100644 --- a/src/software/interpreter/bytecode.zig +++ b/src/software/interpreter/bytecode.zig @@ -36,49 +36,50 @@ comptime { } pub const Opcode = enum(u16) { - copy, - negate_i32, - negate_f32, - logical_not, - bitwise_not, - integer_add, - integer_subtract, - integer_multiply, - unsigned_divide, - signed_divide, - unsigned_modulo, - signed_modulo, - float_add, - float_subtract, - float_multiply, - float_divide, - float_modulo, - shift_left, - logical_shift_right, + @"unreachable", + array_length, arithmetic_shift_right, bitwise_and, + bitwise_not, bitwise_or, bitwise_xor, - logical_and, - logical_or, + branch, compare_equal, compare_not_equal, - compare_unsigned_less, - compare_signed_less, compare_ordered_float_equal, - compare_unordered_float_equal, - compare_ordered_float_not_equal, - compare_unordered_float_not_equal, compare_ordered_float_less, + compare_ordered_float_not_equal, + compare_signed_less, + compare_unordered_float_equal, compare_unordered_float_less, - select, - load_buffer, - store_buffer, - jump_edge, - branch, - return_void, + compare_unordered_float_not_equal, + compare_unsigned_less, + copy, discard, - @"unreachable", + float_add, + float_divide, + float_modulo, + float_multiply, + float_subtract, + integer_add, + integer_multiply, + integer_subtract, + jump_edge, + load_buffer, + logical_and, + logical_not, + logical_or, + logical_shift_right, + negate_f32, + negate_i32, + return_void, + select, + shift_left, + signed_divide, + signed_modulo, + store_buffer, + unsigned_divide, + unsigned_modulo, }; pub const Copy = struct { @@ -99,3 +100,8 @@ pub const Branch = struct { true_edge: u32, false_edge: u32, }; + +pub const ArrayLength = struct { + resource: u32, + stride: u32, +}; diff --git a/src/software/interpreter/test/arithmetic.zig b/src/software/interpreter/test/arithmetic.zig index 3e64dc5..5e2006a 100644 --- a/src/software/interpreter/test/arithmetic.zig +++ b/src/software/interpreter/test/arithmetic.zig @@ -14,7 +14,7 @@ fn bitsF32(value: u32) f32 { return @bitCast(value); } -test "[interpreter] vector floating-point arithmetic" { +test "[interpreter] vector float arithmetic" { var module = try ir.parser.parseString(std.testing.allocator, \\ shader vertex @main \\ { diff --git a/src/software/interpreter/test/loops.zig b/src/software/interpreter/test/loops.zig index f6cc385..e324298 100644 --- a/src/software/interpreter/test/loops.zig +++ b/src/software/interpreter/test/loops.zig @@ -6,7 +6,7 @@ const Runtime = @import("../Runtime.zig"); const ir = shader_ir.ir; -test "[interpreter] loop back edges copy block arguments in parallel" { +test "[interpreter] loop copy block arguments in parallel" { var module = try ir.parser.parseString(std.testing.allocator, \\ shader compute @main \\ { diff --git a/src/software/interpreter/test/storage_buffers.zig b/src/software/interpreter/test/storage_buffers.zig index f10a5b0..52b2ca8 100644 --- a/src/software/interpreter/test/storage_buffers.zig +++ b/src/software/interpreter/test/storage_buffers.zig @@ -64,7 +64,26 @@ const bounds_shader = \\ } ; -test "[interpreter] storage-buffer vector load and store use portable little-endian words" { +const array_length_shader = + \\ shader compute @main + \\ { + \\ @source: runtime_array[u32] = storage_buffer[set(0), binding(0)] + \\ @destination: u32 = storage_buffer[set(0), binding(1)] + \\ + \\ %zero: constant u32 = 0 + \\ %offset: constant u32 = 8 + \\ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %length: u32 = array_length @source, %offset, stride 4 + \\ store_buffer @destination, %zero, %length + \\ return + \\ } + \\ } +; + +test "[interpreter] ssbo vector load/store use portable little-endian words" { var module = try ir.parser.parseString(std.testing.allocator, copy_shader); defer module.deinit(); @@ -92,7 +111,7 @@ test "[interpreter] storage-buffer vector load and store use portable little-end try std.testing.expectEqual(@as(u8, 0xcc), destination[19]); } -test "[interpreter] storage-buffer scalar load and store interpret little-endian words" { +test "[interpreter] ssbo scalar load/store interpret little-endian words" { var module = try ir.parser.parseString(std.testing.allocator, scalar_shader); defer module.deinit(); @@ -108,7 +127,7 @@ test "[interpreter] storage-buffer scalar load and store interpret little-endian try std.testing.expectEqualSlices(u8, &[_]u8{ 0x79, 0x56, 0x34, 0x12 }, &destination); } -test "[interpreter] storage-buffer accesses report unbound and out-of-bounds resources" { +test "[interpreter] ssbo access oob resources" { var module = try ir.parser.parseString(std.testing.allocator, bounds_shader); defer module.deinit(); @@ -125,3 +144,22 @@ test "[interpreter] storage-buffer accesses report unbound and out-of-bounds res const unchanged = [_]u8{0xa5} ** 8; try std.testing.expectEqualSlices(u8, &unchanged, &buffer); } + +test "[interpreter] ssbo array length" { + var module = try ir.parser.parseString(std.testing.allocator, array_length_shader); + defer module.deinit(); + + var program = try Program.compile(std.testing.allocator, &module); + defer program.deinit(); + var runtime = try Runtime.init(std.testing.allocator, &program); + defer runtime.deinit(); + + var source = [_]u8{ 0xff, 0x78, 0x56, 0x34, 0x12, 0xef, 0xcd, 0xab, 0x90, 0x04, 0x03, 0x02, 0x01, 0xdd, 0xcc, 0xbb, 0xaa }; + var destination = [_]u8{0xcc} ** 20; + const resources = [_]?[]u8{ source[0..], destination[0..] }; + + try std.testing.expectEqual(Runtime.Outcome.returned, try runtime.run(&program, .{ .resource_buffers = &resources })); + try std.testing.expectEqualSlices(u8, &[_]u8{ 2, 0, 0, 0 }, destination[0..4]); + const unchanged = [_]u8{0xcc} ** 16; + try std.testing.expectEqualSlices(u8, &unchanged, destination[4..]); +} diff --git a/src/software/interpreter/test/termination.zig b/src/software/interpreter/test/termination.zig index 3649e5d..5b9de96 100644 --- a/src/software/interpreter/test/termination.zig +++ b/src/software/interpreter/test/termination.zig @@ -6,7 +6,7 @@ const Runtime = @import("../Runtime.zig"); const ir = shader_ir.ir; -test "[interpreter] fragment discard is an execution outcome" { +test "[interpreter] fragment discard" { var module = try ir.parser.parseString(std.testing.allocator, \\ shader fragment @main \\ { @@ -27,7 +27,7 @@ test "[interpreter] fragment discard is an execution outcome" { try std.testing.expectEqual(Runtime.Outcome.discarded, try runtime.run(&program, .{})); } -test "[interpreter] execution budget stops an infinite loop" { +test "[interpreter] execution budget stops infinite loop" { var module = try ir.parser.parseString(std.testing.allocator, \\ shader compute @main \\ {