[Soft] adding array length support to ir interpreter

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);
+            },
         }
     }
 

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),
+            .@"unreachable" => return RuntimeError.UnreachableExecuted,
-            .negate_i32 => self.unaryInt(instruction, .negate),
+            .array_length => try self.arrayLength(program, options.resource_buffers, instruction),
-            .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),
             .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];

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
+    if (!hasCompatibleInterface(&program, expected_stage) or (expected_stage == .compute and module_ir.execution_modes.workgroup_size == null)) {
-        (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;
     }

src/software/interpreter/bytecode.zig

@@ -36,49 +36,50 @@ comptime {
 }
 
 pub const Opcode = enum(u16) {
-    copy,
+    @"unreachable",
-    negate_i32,
+    array_length,
-    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,
     arithmetic_shift_right,
     bitwise_and,
+    bitwise_not,
     bitwise_or,
     bitwise_xor,
-    logical_and,
+    branch,
-    logical_or,
     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,
+    compare_unordered_float_not_equal,
-    load_buffer,
+    compare_unsigned_less,
-    store_buffer,
+    copy,
-    jump_edge,
-    branch,
-    return_void,
     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,
+};

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
         \\ {

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
         \\ {

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..]);
+}

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
         \\ {