[Flint] adding array_length and some math operations encoding

src/compiler/ir/parser/ast.zig

@@ -101,5 +101,6 @@ pub const ParsedOperation = union(enum) {
     store_interface: struct { interface_name: []const u8, value: ValueRef },
     load_buffer: struct { resource_name: []const u8, byte_offset: ValueRef },
     store_buffer: struct { resource_name: []const u8, byte_offset: ValueRef, value: ValueRef },
+    array_length: struct { resource_name: []const u8, byte_offset: ValueRef, stride: u32 },
     call: struct { function_name: []const u8, arguments: []const ValueRef },
 };

src/compiler/ir/parser/lower.zig

@@ -285,6 +285,19 @@ fn lowerOperation(
                 .inferred_type = null,
             };
         },
+        .array_length => |op| blk: {
+            const resource_id = resources.get(op.resource_name) orelse return error.UnknownResource;
+            const byte_offset = resolveValue(values, op.byte_offset) orelse return error.UnknownValue;
+
+            break :blk .{
+                .operation = .{ .array_length = .{
+                    .resource = resource_id,
+                    .byte_offset = byte_offset,
+                    .stride = op.stride,
+                } },
+                .inferred_type = null,
+            };
+        },
         .call => |op| blk: {
             const function_id = functions.get(op.function_name) orelse return error.UnknownFunction;
             var arguments: std.ArrayList(ids.ValueId) = .empty;

src/compiler/ir/parser/parser.zig

@@ -412,6 +412,19 @@ const Parser = struct {
             } };
         }
 
+        if (std.mem.eql(u8, name, "array_length")) {
+            const resource_name = (try self.expect(.at_name)).text;
+            try self.expectDiscard(.comma);
+            const byte_offset = try self.parseValueRef();
+            try self.expectDiscard(.comma);
+            try self.expectIdentifier("stride");
+            return .{ .array_length = .{
+                .resource_name = resource_name,
+                .byte_offset = byte_offset,
+                .stride = try self.parseUnsigned(u32, .number),
+            } };
+        }
+
         if (std.mem.eql(u8, name, "call")) {
             const function_name = (try self.expect(.at_name)).text;
             try self.expectDiscard(.left_paren);
@@ -522,6 +535,13 @@ const Parser = struct {
             });
         }
 
+        if (std.mem.eql(u8, token.text, "runtime_array")) {
+            try self.expectDiscard(.left_square);
+            const element_type = try self.parseType();
+            try self.expectDiscard(.right_square);
+            return module.internType(.{ .runtime_array = .{ .element_type = element_type } });
+        }
+
         if (std.mem.eql(u8, token.text, "struct")) {
             try self.expectDiscard(.left_square);
             var members: std.ArrayList(ids.TypeId) = .empty;

src/intel/FlintCommandBuffer.zig

@@ -346,7 +346,7 @@ pub fn dispatchBase(interface: *Interface, base_group_x: u32, base_group_y: u32,
     var ranges: [gen9_dispatch.max_surfaces]?MemoryRange = @splat(null);
     var sizes: [gen9_dispatch.max_surfaces]u64 = @splat(0);
     for (artifact.resources.bindings) |resource| {
-        if (resource.set >= base.vulkan_max_descriptor_sets or @as(usize, resource.binding_table_index) >= gen9_dispatch.max_surfaces)
+        if (resource.set >= base.vulkan_max_descriptor_sets or @as(usize, resource.binding_table_index) >= gen9_dispatch.max_storage_surfaces)
             return VkError.ValidationFailed;
 
         const descriptor_set = self.bound_compute_descriptor_sets[resource.set] orelse return VkError.ValidationFailed;
@@ -400,7 +400,7 @@ pub fn dispatchBase(interface: *Interface, base_group_x: u32, base_group_y: u32,
     self.gpu_allocations.append(self.interface.host_allocator.allocator(), state) catch return VkError.OutOfHostMemory;
     state_owned = false;
 
-    for (0..@as(usize, state_layout.surface_count)) |index| {
+    for (0..@as(usize, state_layout.storage_surface_count)) |index| {
         const range = ranges[index] orelse return VkError.ValidationFailed;
         if (range.offset > std.math.maxInt(u32))
             return VkError.FeatureNotPresent;
@@ -415,6 +415,17 @@ pub fn dispatchBase(interface: *Interface, base_group_x: u32, base_group_y: u32,
         }) catch return VkError.OutOfHostMemory;
     }
 
+    const size_table_surface = @as(usize, state_layout.storage_surface_count);
+    self.relocations.append(self.interface.host_allocator.allocator(), .{
+        .source_handle = state_handle,
+        .target_handle = state_handle,
+        .offset = state_layout.surface_address_offsets[size_table_surface],
+        .delta = state_layout.size_table_offset,
+        .read = true,
+        .write = false,
+        .domain = .render,
+    }) catch return VkError.OutOfHostMemory;
+
     try self.emitSlice(&gen9_dispatch.pipeControl(gen9_dispatch.pipe_control.cs_stall |
         gen9_dispatch.pipe_control.dc_flush |
         gen9_dispatch.pipe_control.render_target_flush |

src/intel/FlintPipeline.zig

@@ -171,19 +171,6 @@ fn lowerToFlint(allocator: std.mem.Allocator, module: *base.ShaderModule.IrModul
     const target = device_info orelse return null;
     return compiler.targets.compileCompute(allocator, module, target, .{}) catch |err| switch (err) {
         error.OutOfMemory => return VkError.OutOfHostMemory,
-
-        error.UnsupportedGeneration,
-        error.UnsupportedStage,
-        error.UnsupportedDispatchWidth,
-        error.UnsupportedGrfSize,
-        error.UnsupportedWorkgroupSize,
-        error.MissingWorkgroupSize,
-        error.UnsupportedType,
-        error.UnsupportedOperation,
-        error.UnsupportedTerminator,
-        error.TooManyStorageBuffers,
-        => return null,
-
         else => {
             std.log.scoped(.FlintPipeline).err("compute compilation failed: {s}", .{@errorName(err)});
             return VkError.ValidationFailed;

src/intel/compiler/ir/instruction.zig

@@ -28,6 +28,13 @@ pub const StoreBuffer = struct {
     source: operand.Source,
 };
 
+pub const ArrayLength = struct {
+    destination: operand.Destination,
+    buffer: BufferReference,
+    byte_offset: operand.Source,
+    stride: u32,
+};
+
 pub const SurfaceRead = struct {
     destination: operand.Destination,
     binding_table: u8,
@@ -93,15 +100,28 @@ pub const Compare = struct {
     rhs: operand.Source,
 };
 
+pub const MathOpcode = enum {
+    integer_quotient,
+};
+
+pub const Math = struct {
+    opcode: MathOpcode,
+    destination: operand.Destination,
+    lhs: operand.Source,
+    rhs: operand.Source,
+};
+
 pub const Operation = union(enum) {
     load_global_invocation_id: LoadGlobalInvocationId,
     load_buffer: LoadBuffer,
     store_buffer: StoreBuffer,
+    array_length: ArrayLength,
     surface_read: SurfaceRead,
     surface_write: SurfaceWrite,
     surface_message: SurfaceMessage,
     move: Move,
     binary: Binary,
+    math: Math,
     compare: Compare,
     parallel_copy: pseudo.ParallelCopy,
 };

src/intel/compiler/ir/printer.zig

@@ -136,6 +136,15 @@ fn writeOperation(program: *const program_ir.Program, writer: *std.Io.Writer, ex
             try writer.writeAll(", ");
             try writeSource(program, writer, execution_size, op.source);
         },
+        .array_length => |op| {
+            try writer.writeAll("array_length ");
+            try writeDestination(program, writer, execution_size, op.destination);
+            try writer.writeAll(", ");
+            try writeBufferReference(program, writer, op.buffer);
+            try writer.writeAll(", ");
+            try writeSource(program, writer, execution_size, op.byte_offset);
+            try writer.print(", stride({d})", .{op.stride});
+        },
         .surface_read => |op| {
             try writer.writeAll("surface_read ");
             try writeDestination(program, writer, execution_size, op.destination);
@@ -185,6 +194,14 @@ fn writeOperation(program: *const program_ir.Program, writer: *std.Io.Writer, ex
             try writer.writeAll(", ");
             try writeSource(program, writer, execution_size, op.rhs);
         },
+        .math => |op| {
+            try writer.print("{t} ", .{op.opcode});
+            try writeDestination(program, writer, execution_size, op.destination);
+            try writer.writeAll(", ");
+            try writeSource(program, writer, execution_size, op.lhs);
+            try writer.writeAll(", ");
+            try writeSource(program, writer, execution_size, op.rhs);
+        },
         .parallel_copy => |op| try writeParallelCopy(program, writer, execution_size, op),
     }
 }

src/intel/compiler/ir/validator.zig

@@ -5,42 +5,43 @@ const program_ir = @import("program.zig");
 const pseudo = @import("pseudo.zig");
 
 pub const Error = error{
-    MissingEntryBlock,
-    InvalidBlock,
-    MissingTerminator,
-    InvalidInstruction,
-    InvalidVirtualRegister,
-    InvalidVirtualFlag,
-    UnallocatedVirtualFlag,
-    InvalidPhysicalRegister,
-    InvalidRegisterSize,
-    InvalidRegisterAlignment,
-    InvalidLaneCount,
-    InvalidRegion,
-    InvalidDestination,
-    InvalidImmediateType,
-    InvalidStorageBuffer,
-    InvalidBufferReference,
-    InvalidGlobalInvocationId,
-    InvalidBufferAccess,
-    InvalidWorkgroupSize,
-    EmptyParallelCopy,
-    InvalidParallelCopyDestination,
-    ParallelCopyTypeMismatch,
-    DuplicateParallelCopyDestination,
-    PredicatedParallelCopy,
-    UnloweredParallelCopy,
-    UnloweredSystemValue,
-    UnloweredResource,
-    UnloweredMessage,
-    InvalidMessage,
-    InvalidPayloadLayout,
-    EntryBlockHasParameters,
     DuplicateBlockParameter,
+    DuplicateParallelCopyDestination,
     EdgeArgumentCountMismatch,
     EdgeArgumentKindMismatch,
     EdgeArgumentTypeMismatch,
+    EmptyParallelCopy,
+    EntryBlockHasParameters,
+    InvalidBlock,
+    InvalidBufferAccess,
+    InvalidBufferReference,
+    InvalidDestination,
+    InvalidGlobalInvocationId,
+    InvalidImmediateType,
+    InvalidInstruction,
+    InvalidLaneCount,
+    InvalidMath,
+    InvalidMessage,
+    InvalidParallelCopyDestination,
+    InvalidPayloadLayout,
+    InvalidPhysicalRegister,
+    InvalidRegion,
+    InvalidRegisterAlignment,
+    InvalidRegisterSize,
+    InvalidStorageBuffer,
+    InvalidVirtualFlag,
+    InvalidVirtualRegister,
+    InvalidWorkgroupSize,
+    MissingEntryBlock,
+    MissingTerminator,
+    ParallelCopyTypeMismatch,
+    PredicatedParallelCopy,
+    UnallocatedVirtualFlag,
     UnloweredBlockParameter,
+    UnloweredMessage,
+    UnloweredParallelCopy,
+    UnloweredResource,
+    UnloweredSystemValue,
 };
 
 pub fn validate(program: *const program_ir.Program) Error!void {
@@ -157,6 +158,13 @@ fn validateInstruction(program: *const program_ir.Program, inst: instruction.Ins
             if (!op.source.type.isInitialTargetType())
                 return Error.InvalidBufferAccess;
         },
+        .array_length => |op| {
+            try validateBufferReference(program, op.buffer);
+            try validateDestination(program, op.destination);
+            try validateBufferOffset(program, op.byte_offset);
+            if (op.destination.type != .u32 or op.stride == 0)
+                return Error.InvalidBufferAccess;
+        },
         .surface_read => |op| {
             try validateDestination(program, op.destination);
             try validateBufferOffset(program, op.address);
@@ -199,6 +207,22 @@ fn validateInstruction(program: *const program_ir.Program, inst: instruction.Ins
             try validateSource(program, op.lhs);
             try validateSource(program, op.rhs);
         },
+        .math => |op| {
+            try validateDestination(program, op.destination);
+            try validateSource(program, op.lhs);
+            try validateSource(program, op.rhs);
+
+            switch (op.opcode) {
+                .integer_quotient => {
+                    if (inst.execution_size != .simd8)
+                        return Error.InvalidMath;
+                    if (op.destination.type != .u32 and op.destination.type != .i32)
+                        return Error.InvalidMath;
+                    if (op.lhs.type != op.destination.type or op.rhs.type != op.destination.type)
+                        return Error.InvalidMath;
+                },
+            }
+        },
 
         .parallel_copy => |op| {
             if (program.properties.parallel_copies_lowered)

src/intel/compiler/lower/common_ir.zig

@@ -1,4 +1,5 @@
 const std = @import("std");
+const base = @import("base");
 const shader_compiler = @import("shader_ir");
 const shader_ir = shader_compiler.ir;
 const device = @import("../device.zig");
@@ -431,7 +432,7 @@ const LoweringState = struct {
             .load_buffer => |operation| try self.lowerLoadBuffer(block_id, source_instruction.result, operation),
             .store_buffer => |operation| try self.lowerStoreBuffer(block_id, source_instruction.result, operation),
             .call => return Error.UnsanitizedModule,
-            .array_length => return Error.UnsupportedOperation,
+            .array_length => |operation| try self.lowerArrayLength(block_id, source_instruction.result, operation),
         }
     }
 
@@ -507,6 +508,32 @@ const LoweringState = struct {
         if (lhs_components.len == 0 or lhs_components.len != rhs_components.len or lhs_components.len != result_components.len)
             return Error.InvalidModule;
 
+        switch (operation.opcode) {
+            .unsigned_divide,
+            .signed_divide,
+            => {
+                const signed = operation.opcode == .signed_divide;
+                const target_type: operand.DataType = if (signed) .i32 else .u32;
+
+                for (lhs_components, rhs_components, result_components) |lhs_value, rhs_value, result_component| {
+                    var dst = try destinationFromSource(result_component);
+                    dst.type = target_type;
+
+                    try self.appendInstruction(block_id, null, .{
+                        .math = .{
+                            .opcode = .integer_quotient,
+                            .destination = dst,
+                            .lhs = try retypeIntegerSource(lhs_value, target_type),
+                            .rhs = try retypeIntegerSource(rhs_value, target_type),
+                        },
+                    });
+                }
+                return;
+            },
+
+            else => {},
+        }
+
         const data_type = lhs_components[0].type;
         const opcode: instruction.BinaryOpcode = switch (operation.opcode) {
             .integer_add => if (data_type == .u32 or data_type == .i32) .add else return Error.UnsupportedOperation,
@@ -521,8 +548,11 @@ const LoweringState = struct {
             .bitwise_and => if (data_type == .u32 or data_type == .i32) .bitwise_and else return Error.UnsupportedOperation,
             .bitwise_or => if (data_type == .u32 or data_type == .i32) .bitwise_or else return Error.UnsupportedOperation,
             .bitwise_xor => if (data_type == .u32 or data_type == .i32) .bitwise_xor else return Error.UnsupportedOperation,
+
             .unsigned_divide,
             .signed_divide,
+            => unreachable,
+
             .unsigned_modulo,
             .signed_modulo,
             .float_divide,
@@ -549,6 +579,47 @@ const LoweringState = struct {
         }
     }
 
+    fn retypeIntegerSource(source_value: operand.Source, target_type: operand.DataType) Error!operand.Source {
+        if (target_type != .u32 and target_type != .i32)
+            return Error.UnsupportedType;
+
+        var result = source_value;
+
+        switch (source_value.type) {
+            .u32, .i32 => {},
+            else => return Error.UnsupportedType,
+        }
+
+        result.register = switch (source_value.register) {
+            .immediate => |immediate| .{
+                .immediate = switch (target_type) {
+                    .u32 => .{
+                        .u32 = switch (immediate) {
+                            .u32 => |v| v,
+                            .i32 => |v| @bitCast(v),
+                            else => return Error.UnsupportedType,
+                        },
+                    },
+
+                    .i32 => .{
+                        .i32 = switch (immediate) {
+                            .u32 => |v| @bitCast(v),
+                            .i32 => |v| v,
+                            else => return Error.UnsupportedType,
+                        },
+                    },
+
+                    else => unreachable,
+                },
+            },
+
+            else => source_value.register,
+        };
+
+        result.type = target_type;
+        return result;
+    }
+
     fn lowerCompare(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.Compare) Error!void {
         const result_id = try requireResult(result);
         const result_value = self.lowerer.module.values.get(result_id) orelse return Error.InvalidModule;
@@ -689,21 +760,36 @@ const LoweringState = struct {
         const variable = self.lowerer.module.interface_variables.get(operation.variable) orelse return Error.InvalidModule;
         if (variable.direction != .input)
             return Error.InvalidModule;
+
         switch (variable.semantic) {
-            .builtin => |builtin| if (builtin != .global_invocation_id)
+            .builtin => |builtin| switch (builtin) {
-                return Error.UnsupportedOperation,
+                .global_invocation_id => try self.lowerGlobalInvocationId(block_id, result_id, variable),
+                .num_workgroups => try self.lowerNumWorkgroups(result_id),
+                .workgroup_size => base.unsupported("workgroup size builtin is not yet supported in Flint", .{}),
+                else => return Error.UnsupportedOperation,
+            },
             .location => return Error.UnsupportedOperation,
         }
+    }
 
+    fn lowerGlobalInvocationId(
+        self: *LoweringState,
+        block_id: ids.BlockId,
+        result_id: shader_ir.id.ValueId,
+        variable: *const shader_ir.module.InterfaceVariable,
+    ) Error!void {
         const result_value = self.lowerer.module.values.get(result_id) orelse return Error.InvalidModule;
         if (result_value.type != variable.type)
             return Error.InvalidModule;
+
         const result_components = try self.addRegisterLocation(result_id, .temporary);
         if (result_components.len != 3)
             return Error.UnsupportedOperation;
+
         for (result_components, 0..) |result_component, component_index| {
             if (result_component.type != .u32)
                 return Error.UnsupportedOperation;
+
             try self.appendInstruction(block_id, null, .{
                 .load_global_invocation_id = .{
                     .destination = try destinationFromSource(result_component),
@@ -713,6 +799,28 @@ const LoweringState = struct {
         }
     }
 
+    fn lowerNumWorkgroups(self: *LoweringState, result_id: shader_ir.id.ValueId) Error!void {
+        const result_value = self.lowerer.module.values.get(result_id) orelse return Error.InvalidModule;
+        const lowered_type = try self.lowerType(result_value.type);
+
+        if (lowered_type.element_type != .u32 or lowered_type.component_count != 3)
+            return Error.UnsupportedType;
+
+        const vec = try self.storage.alloc(operand.Source, 3);
+
+        for (self.lowerer.module.execution_modes.workgroup_size.?, vec) |value, *component| {
+            component.* = .{
+                .register = .{
+                    .immediate = .{ .u32 = value },
+                },
+                .type = .u32,
+                .region = operand.Region.broadcast(),
+            };
+        }
+
+        try self.putLocation(result_id, .{ .components = vec });
+    }
+
     fn lowerStoreInterface(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.StoreInterface) Error!void {
         _ = self;
         _ = block_id;
@@ -759,6 +867,26 @@ const LoweringState = struct {
         }
     }
 
+    fn lowerArrayLength(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.ArrayLength) Error!void {
+        const result_id = try requireResult(result);
+        const byte_offset = try self.source(operation.byte_offset);
+        if (byte_offset.type != .u32)
+            return Error.UnsupportedType;
+
+        const result_components = try self.addRegisterLocation(result_id, .temporary);
+        if (result_components.len != 1 or result_components[0].type != .u32)
+            return Error.UnsupportedType;
+
+        try self.appendInstruction(block_id, null, .{
+            .array_length = .{
+                .destination = try destinationFromSource(result_components[0]),
+                .buffer = .{ .logical = try self.storageBuffer(operation.resource) },
+                .byte_offset = byte_offset,
+                .stride = operation.stride,
+            },
+        });
+    }
+
     fn lowerControlAndTerminators(self: *LoweringState, allocator: std.mem.Allocator) Error!void {
         const source_entry = try self.sourceEntryFunction();
         const function = source_entry[1];
@@ -1297,6 +1425,28 @@ test "[ir] Lower: vector storage-buffer operations" {
     }, &.{});
 }
 
+test "[ir] Lower: runtime array length" {
+    const source =
+        \\shader compute @main
+        \\{
+        \\    @storage: runtime_array[u32] = storage_buffer[set(0), binding(3)]
+        \\    %offset: constant u32 = 16
+        \\    fn @main() -> void
+        \\    {
+        \\        .entry():
+        \\            %length: u32 = array_length @storage, %offset, stride 4
+        \\            return
+        \\    }
+        \\}
+    ;
+
+    try expectLoweredFragments(source, &.{
+        "@storage = storage_buffer[set(0), binding(3)]",
+        "%length: vgrf u32[8], class(temporary)",
+        "[simd8] array_length %length:u32, @storage, 16:u32, stride(4)",
+    }, &.{});
+}
+
 test "[ir] Lower: vector block parameter" {
     const source =
         \\shader compute @main
@@ -1402,20 +1552,6 @@ test "[ir] Lower: boolean block parameter" {
 }
 
 test "[ir] Lower: unsupported operations" {
-    try expectLoweringError(
-        \\shader compute @main
-        \\{
-        \\    %one: constant u32 = bits(0x1)
-        \\    %two: constant u32 = bits(0x2)
-        \\    fn @main() -> void
-        \\    {
-        \\        .entry():
-        \\            %quotient: u32 = unsigned_divide %one, %two
-        \\            return
-        \\    }
-        \\}
-    , Error.UnsupportedOperation);
-
     try expectLoweringError(
         \\shader compute @main
         \\{

src/intel/compiler/targets/gen9/compute/array_length_lowering.zig

@@ -0,0 +1,92 @@
+const std = @import("std");
+
+const Builder = @import("../../../ir/Builder.zig");
+const ids = @import("../../../ir/id.zig");
+const instruction = @import("../../../ir/instruction.zig");
+const operand = @import("../../../ir/operand.zig");
+const program_ir = @import("../../../ir/program.zig");
+const resource_layout = @import("resource_layout.zig");
+
+pub const Error = std.mem.Allocator.Error || error{InvalidProgram};
+
+pub fn run(program: *program_ir.Program, layout: *const resource_layout.Layout) Error!void {
+    if (!program.properties.resources_lowered)
+        return Error.InvalidProgram;
+    if (layout.bindings.len >= std.math.maxInt(u8))
+        return Error.InvalidProgram;
+
+    var builder = Builder.init(program);
+    for (program.blocks.entries.items, 0..) |entry, block_index| {
+        _ = entry orelse continue;
+        const block_id = ids.BlockId.fromIndex(block_index);
+        var instruction_index: usize = 0;
+
+        while (true) {
+            const block = program.blocks.get(block_id) orelse return Error.InvalidProgram;
+            if (instruction_index >= block.instructions.items.len)
+                break;
+
+            const instruction_id = block.instructions.items[instruction_index];
+            const inst = program.instructions.get(instruction_id) orelse return Error.InvalidProgram;
+            const op = switch (inst.operation) {
+                .array_length => |value| value,
+                else => {
+                    instruction_index += 1;
+                    continue;
+                },
+            };
+            const resource_index = switch (op.buffer) {
+                .binding_table => |value| value,
+                .logical => return Error.InvalidProgram,
+            };
+            if (resource_index >= layout.bindings.len or op.stride == 0)
+                return Error.InvalidProgram;
+
+            const execution_size = inst.execution_size;
+            const predicate = inst.predicate;
+            const result_source: operand.Source = .{
+                .register = op.destination.register,
+                .type = .u32,
+                .region = operand.Region.contiguous(execution_size),
+            };
+            var negated_offset = op.byte_offset;
+            negated_offset.negate = !negated_offset.negate;
+
+            const mutable = program.instructions.getMut(instruction_id) orelse return Error.InvalidProgram;
+            mutable.operation = .{ .load_buffer = .{
+                .destination = op.destination,
+                .buffer = .{ .binding_table = @intCast(layout.bindings.len) },
+                .byte_offset = immediate(@as(u32, resource_index) * @sizeOf(u32)),
+            } };
+
+            _ = builder.insertInstruction(block_id, instruction_index + 1, execution_size, predicate, .{ .binary = .{
+                .opcode = .add,
+                .destination = op.destination,
+                .lhs = result_source,
+                .rhs = negated_offset,
+            } }) catch |err| return mapBuilderError(err);
+            _ = builder.insertInstruction(block_id, instruction_index + 2, execution_size, predicate, .{ .math = .{
+                .opcode = .integer_quotient,
+                .destination = op.destination,
+                .lhs = result_source,
+                .rhs = immediate(op.stride),
+            } }) catch |err| return mapBuilderError(err);
+            instruction_index += 3;
+        }
+    }
+}
+
+fn immediate(value: u32) operand.Source {
+    return .{
+        .register = .{ .immediate = .{ .u32 = value } },
+        .type = .u32,
+        .region = operand.Region.broadcast(),
+    };
+}
+
+fn mapBuilderError(err: Builder.Error) Error {
+    return switch (err) {
+        error.OutOfMemory => Error.OutOfMemory,
+        else => Error.InvalidProgram,
+    };
+}

src/intel/compiler/targets/gen9/compute/dispatch.zig

@@ -1,6 +1,7 @@
 const std = @import("std");
 
-pub const max_surfaces: usize = 4;
+pub const max_storage_surfaces: usize = 4;
+pub const max_surfaces: usize = max_storage_surfaces + 1;
 pub const page_size: usize = 4096;
 pub const surface_state_size: usize = 64;
 pub const interface_descriptor_size: usize = 32;
@@ -22,12 +23,14 @@ pub const StateLayout = struct {
     surface_offsets: [max_surfaces]u32,
     surface_address_offsets: [max_surfaces]u32,
     surface_count: u8,
+    storage_surface_count: u8,
+    size_table_offset: u32,
     binding_table_offset: u32,
     interface_descriptor_offset: u32,
 };
 
 pub fn writeState(destination: []u8, kernel: []const u8, buffer_sizes: []const u64) Error!StateLayout {
-    if (buffer_sizes.len > max_surfaces)
+    if (buffer_sizes.len > max_storage_surfaces)
         return Error.TooManySurfaces;
 
     var layout: StateLayout = .{
@@ -35,7 +38,9 @@ pub fn writeState(destination: []u8, kernel: []const u8, buffer_sizes: []const u
         .kernel_offset = 0,
         .surface_offsets = @splat(0),
         .surface_address_offsets = @splat(0),
-        .surface_count = @intCast(buffer_sizes.len),
+        .surface_count = @intCast(buffer_sizes.len + 1),
+        .storage_surface_count = @intCast(buffer_sizes.len),
+        .size_table_offset = 0,
         .binding_table_offset = 0,
         .interface_descriptor_offset = 0,
     };
@@ -48,11 +53,23 @@ pub fn writeState(destination: []u8, kernel: []const u8, buffer_sizes: []const u
         cursor += surface_state_size;
         if (size == 0)
             return Error.EmptyBuffer;
+        if (size > std.math.maxInt(u32))
+            return Error.UnsupportedBufferSize;
     }
 
+    const size_table_surface = buffer_sizes.len;
+    cursor = alignForward(cursor, surface_state_size);
+    layout.surface_offsets[size_table_surface] = @intCast(cursor);
+    layout.surface_address_offsets[size_table_surface] = @intCast(cursor + 8 * @sizeOf(u32));
+    cursor += surface_state_size;
+
     cursor = alignForward(cursor, 32);
     layout.binding_table_offset = @intCast(cursor);
-    cursor += buffer_sizes.len * @sizeOf(u32);
+    cursor += layout.surface_count * @sizeOf(u32);
+
+    cursor = alignForward(cursor, @alignOf(u32));
+    layout.size_table_offset = @intCast(cursor);
+    cursor += @max(buffer_sizes.len, 1) * @sizeOf(u32);
 
     cursor = alignForward(cursor, 64);
     layout.interface_descriptor_offset = @intCast(cursor);
@@ -67,7 +84,10 @@ pub fn writeState(destination: []u8, kernel: []const u8, buffer_sizes: []const u
     for (buffer_sizes, 0..) |size, index| {
         _ = try encodeRawBufferSurface(destination, layout.surface_offsets[index], size);
         putU32(destination, layout.binding_table_offset + @as(u32, @intCast(index * @sizeOf(u32))), layout.surface_offsets[index]);
+        putU32(destination, layout.size_table_offset + @as(u32, @intCast(index * @sizeOf(u32))), @intCast(size));
     }
+    _ = try encodeRawBufferSurface(destination, layout.surface_offsets[size_table_surface], @max(buffer_sizes.len, 1) * @sizeOf(u32));
+    putU32(destination, layout.binding_table_offset + @as(u32, @intCast(size_table_surface * @sizeOf(u32))), layout.surface_offsets[size_table_surface]);
 
     const idd = layout.interface_descriptor_offset;
     putU32(destination, idd + 0, layout.kernel_offset);

src/intel/compiler/targets/gen9/compute/eu_encoder.zig

@@ -54,28 +54,8 @@ const Grf = struct {
 pub fn encodeMove(execution_size: device.ExecutionSize, move: ir_instruction.Move) Error!EncodedInstruction {
     var encoded = try instructionHeader(1, execution_size);
     const destination = try resolveGrf(move.destination.register, move.destination.region.byte_offset);
-
+    setDestination(&encoded, .grf, try hardwareType(move.destination.type), destination, try horizontalStride(move.destination.region.horizontal_stride));
-    setDestination(
+    try setSource0(&encoded, move.source);
-        &encoded,
-        .grf,
-        try hardwareType(move.destination.type),
-        destination,
-        try horizontalStride(move.destination.region.horizontal_stride),
-    );
-
-    switch (move.source.register) {
-        .physical_grf => {
-            const source = try resolveGrf(move.source.register, move.source.region.byte_offset);
-            try setSource0Register(&encoded, move.source, source);
-        },
-        .immediate => |immediate| {
-            if (move.source.negate or move.source.absolute)
-                return Error.UnsupportedOperand;
-            setSource0Immediate(&encoded, try hardwareType(move.source.type), immediate);
-        },
-        else => return Error.UnsupportedOperand,
-    }
-
     return encoded;
 }
 
@@ -141,6 +121,38 @@ pub fn encodeSurfaceMessage(execution_size: device.ExecutionSize, message: ir_in
     return encoded;
 }
 
+pub fn encodeBinary(execution_size: device.ExecutionSize, binary: ir_instruction.Binary) Error!EncodedInstruction {
+    const opcode: u7 = switch (binary.opcode) {
+        .add => 64,
+        else => return Error.UnsupportedOperand,
+    };
+
+    var encoded = try instructionHeader(opcode, execution_size);
+    const destination = try resolveGrf(binary.destination.register, binary.destination.region.byte_offset);
+    setDestination(&encoded, .grf, try hardwareType(binary.destination.type), destination, try horizontalStride(binary.destination.region.horizontal_stride));
+    try setSource0(&encoded, binary.lhs);
+    try setSource1(&encoded, binary.rhs);
+    return encoded;
+}
+
+pub fn encodeMath(execution_size: device.ExecutionSize, math: ir_instruction.Math) Error!EncodedInstruction {
+    if (execution_size != .simd8)
+        return Error.UnsupportedExecutionSize;
+
+    var encoded = try instructionHeader(56, execution_size);
+    const destination = try resolveGrf(math.destination.register, math.destination.region.byte_offset);
+    setDestination(&encoded, .grf, try hardwareType(math.destination.type), destination, try horizontalStride(math.destination.region.horizontal_stride));
+
+    try setSource0(&encoded, math.lhs);
+    try setSource1(&encoded, math.rhs);
+
+    encoded.setBits(27, 24, switch (math.opcode) {
+        .integer_quotient => 12,
+    });
+
+    return encoded;
+}
+
 fn instructionHeader(opcode: u7, execution_size: device.ExecutionSize) Error!EncodedInstruction {
     var encoded: EncodedInstruction = .{};
     encoded.setBits(6, 0, opcode);
@@ -180,6 +192,81 @@ fn setSource0Immediate(encoded: *EncodedInstruction, data_type: HardwareType, im
     });
 }
 
+fn setSource0(encoded: *EncodedInstruction, source: operand.Source) Error!void {
+    switch (source.register) {
+        .physical_grf => {
+            const register = try resolveGrf(source.register, source.region.byte_offset);
+            try setSource0Register(encoded, source, register);
+        },
+        .immediate => |immediate| {
+            setSource0Immediate(encoded, try hardwareType(source.type), try applyImmediateModifiers(immediate, source.negate, source.absolute));
+        },
+
+        else => return Error.UnsupportedOperand,
+    }
+}
+
+fn setSource1Register(encoded: *EncodedInstruction, source: operand.Source, register: Grf) Error!void {
+    encoded.setBits(90, 89, @intFromEnum(RegisterFile.grf));
+    encoded.setBits(94, 91, @intFromEnum(try hardwareType(source.type)));
+
+    // Direct addressing.
+    encoded.setBits(100, 96, register.byte_offset);
+    encoded.setBits(108, 101, register.number);
+
+    // Source modifiers.
+    encoded.setBits(109, 109, @intFromBool(source.absolute));
+    encoded.setBits(110, 110, @intFromBool(source.negate));
+
+    // AddressMode = direct.
+    encoded.setBits(111, 111, 0);
+
+    // Align1 region.
+    encoded.setBits(113, 112, try horizontalStride(source.region.horizontal_stride));
+    encoded.setBits(116, 114, try regionWidth(source.region.width));
+    encoded.setBits(120, 117, try verticalStride(source.region.vertical_stride));
+}
+
+fn setSource1Immediate(encoded: *EncodedInstruction, data_type: HardwareType, immediate: operand.Immediate) void {
+    encoded.setBits(90, 89, @intFromEnum(RegisterFile.immediate));
+
+    encoded.setBits(94, 91, @intFromEnum(data_type));
+
+    encoded.setBits(127, 96, switch (immediate) {
+        .u32 => |value| value,
+        .i32 => |value| @as(u32, @bitCast(value)),
+        .f32 => |value| @as(u32, @bitCast(value)),
+    });
+}
+
+fn setSource1(encoded: *EncodedInstruction, source: operand.Source) Error!void {
+    switch (source.register) {
+        .physical_grf => {
+            const register = try resolveGrf(source.register, source.region.byte_offset);
+            try setSource1Register(encoded, source, register);
+        },
+
+        .immediate => |immediate| {
+            setSource1Immediate(encoded, try hardwareType(source.type), try applyImmediateModifiers(immediate, source.negate, source.absolute));
+        },
+
+        else => return Error.UnsupportedOperand,
+    }
+}
+
+fn applyImmediateModifiers(immediate: operand.Immediate, negate: bool, absolute: bool) Error!operand.Immediate {
+    if (absolute)
+        return Error.UnsupportedOperand;
+    if (!negate)
+        return immediate;
+
+    return switch (immediate) {
+        .u32 => |value| .{ .u32 = 0 -% value },
+        .i32 => |value| .{ .i32 = 0 -% value },
+        .f32 => |value| .{ .f32 = -value },
+    };
+}
+
 fn resolveGrf(register: operand.RegisterRef, region_byte_offset: u16) Error!Grf {
     const physical = switch (register) {
         .physical_grf => |value| value,

src/intel/compiler/targets/gen9/compute/kernel_encoder.zig

@@ -25,21 +25,33 @@ pub fn encode(allocator: std.mem.Allocator, program: *program_ir.Program) Error!
             live_block_count += 1;
     }
 
-    if (live_block_count != 1)
+    if (live_block_count != 1) {
+        std.log.scoped(.FlintEuEncoder).err("cannot encode control flow: program has {d} live blocks; only one is currently supported", .{live_block_count});
         return Error.UnsupportedControlFlow;
+    }
 
     var kernel: std.ArrayList(u8) = .empty;
     errdefer kernel.deinit(allocator);
 
     for (entry.instructions.items) |instruction_id| {
         const instruction = program.instructions.get(instruction_id) orelse return Error.InvalidProgram;
-        if (instruction.predicate != null)
+        if (instruction.predicate != null) {
+            std.log.scoped(.FlintEuEncoder).err("cannot encode instruction {d} ({t}): predication is not supported", .{ instruction_id.index(), std.meta.activeTag(instruction.operation) });
             return Error.UnsupportedPredication;
+        }
 
         const encoded = switch (instruction.operation) {
-            .move => |move| try eu.encodeMove(instruction.execution_size, move),
+            .move => |move| eu.encodeMove(instruction.execution_size, move),
-            .surface_message => |message| try eu.encodeSurfaceMessage(instruction.execution_size, message),
+            .surface_message => |message| eu.encodeSurfaceMessage(instruction.execution_size, message),
-            else => return Error.UnsupportedOperation,
+            .binary => |binary| eu.encodeBinary(instruction.execution_size, binary),
+            .math => |math| eu.encodeMath(instruction.execution_size, math),
+            else => {
+                std.log.scoped(.FlintEuEncoder).err("cannot encode instruction {d}: unsupported operation {t}", .{ instruction_id.index(), std.meta.activeTag(instruction.operation) });
+                return Error.UnsupportedOperation;
+            },
+        } catch |err| {
+            std.log.scoped(.FlintEuEncoder).err("failed to encode instruction {d} ({t}): {s}", .{ instruction_id.index(), std.meta.activeTag(instruction.operation), @errorName(err) });
+            return err;
         };
         try appendInstruction(allocator, &kernel, encoded);
     }
@@ -53,7 +65,10 @@ pub fn encode(allocator: std.mem.Allocator, program: *program_ir.Program) Error!
                 try appendInstruction(allocator, &kernel, instruction);
             program.program_data.total_grf_count = eu.eot_payload_grf + 1;
         },
-        else => return Error.UnsupportedControlFlow,
+        else => {
+            std.log.scoped(.FlintEuEncoder).err("cannot encode entry-block terminator {t}", .{std.meta.activeTag(terminator)});
+            return Error.UnsupportedControlFlow;
+        },
     }
 
     return kernel.toOwnedSlice(allocator);

src/intel/compiler/targets/gen9/compute/pipeline.zig

@@ -11,6 +11,7 @@ const register_allocation = @import("../register_allocation.zig");
 
 const compute = @import("compute.zig");
 const abi = @import("abi.zig");
+const array_length_lowering = @import("array_length_lowering.zig");
 const kernel_encoder = @import("kernel_encoder.zig");
 const message_addresses = @import("message_addresses.zig");
 const message_lowering = @import("message_lowering.zig");
@@ -22,6 +23,7 @@ pub const Error = common_ir.Error ||
     block_arguments.Error ||
     parallel_copies.Error ||
     abi.Error ||
+    array_length_lowering.Error ||
     kernel_encoder.Error ||
     message_addresses.Error ||
     message_lowering.Error ||
@@ -76,29 +78,30 @@ pub fn compile(allocator: std.mem.Allocator, module: *shader_ir.module.Module, d
     try block_arguments.run(allocator, &program);
     try parallel_copies.run(allocator, &program);
 
-    var resources = try resource_layout.Layout.init(
+    var resources = try resource_layout.Layout.init(allocator, &program);
-        allocator,
-        &program,
-    );
     errdefer resources.deinit(allocator);
 
     try resource_lowering.run(&program, &resources);
+    try array_length_lowering.run(&program, &resources);
     try message_lowering.run(&program);
     try message_addresses.run(&program);
     try message_payloads.run(&program);
     try flag_allocation.run(allocator, &program);
     try register_allocation.run(allocator, &program);
 
-    const kernel = kernel_encoder.encode(allocator, &program) catch |err| switch (err) {
+    const kernel = kernel_encoder.encode(allocator, &program) catch |err| encoding_error: {
-        error.UnsupportedControlFlow,
+        std.log.scoped(.FlintCompiler).err("Gen9 EU kernel encoding failed: {s}", .{@errorName(err)});
-        error.UnsupportedOperation,
+        break :encoding_error switch (err) {
-        error.UnsupportedPredication,
+            error.UnsupportedControlFlow,
-        error.UnsupportedExecutionSize,
+            error.UnsupportedOperation,
-        error.UnsupportedDataType,
+            error.UnsupportedPredication,
-        error.UnsupportedOperand,
+            error.UnsupportedExecutionSize,
-        error.EotRegisterUnavailable,
+            error.UnsupportedDataType,
-        => null,
+            error.UnsupportedOperand,
-        else => return err,
+            error.EotRegisterUnavailable,
+            => null,
+            else => return err,
+        };
     };
     errdefer if (kernel) |bytes| allocator.free(bytes);
 

src/intel/compiler/targets/gen9/compute/resource_lowering.zig

@@ -53,6 +53,7 @@ fn bufferReference(operation: instruction.Operation) ?instruction.BufferReferenc
     return switch (operation) {
         .load_buffer => |op| op.buffer,
         .store_buffer => |op| op.buffer,
+        .array_length => |op| op.buffer,
         else => null,
     };
 }
@@ -61,6 +62,7 @@ fn bufferReferenceMut(operation: *instruction.Operation) ?*instruction.BufferRef
     return switch (operation.*) {
         .load_buffer => |*op| &op.buffer,
         .store_buffer => |*op| &op.buffer,
+        .array_length => |*op| &op.buffer,
         else => null,
     };
 }

src/intel/compiler/targets/gen9/gen9.zig

@@ -150,3 +150,40 @@ test "[gen9] target: lower 256 KiB SSBO copy loop" {
     try std.testing.expectEqual(@as(usize, 4), load_count);
     try std.testing.expectEqual(@as(usize, 4), store_count);
 }
+
+test "[gen9] target: encode runtime array length" {
+    var module = try shader_ir.parser.parseString(std.testing.allocator,
+        \\shader compute @main
+        \\{
+        \\    @storage: runtime_array[u32] = storage_buffer[set(0), binding(0)]
+        \\    %offset: constant u32 = 16
+        \\    fn @main() -> void
+        \\    {
+        \\        .entry():
+        \\            %length: u32 = array_length @storage, %offset, stride 4
+        \\            return
+        \\    }
+        \\}
+    );
+    defer module.deinit();
+    module.execution_modes.workgroup_size = .{ 1, 1, 1 };
+
+    const gen9_device: device.DeviceInfo = .{
+        .generation = .gen9,
+        .platform = .skylake,
+        .pci_device_id = 0x1912,
+        .grf_count = 128,
+    };
+    var artifact = try compileCompute(std.testing.allocator, &module, gen9_device, .{});
+    defer artifact.deinit(std.testing.allocator);
+
+    if (artifact.kernel == null) {
+        const encoded = try compute.kernel_encoder.encode(std.testing.allocator, &artifact.program);
+        std.testing.allocator.free(encoded);
+        return error.TestExpectedEncodedKernel;
+    }
+    for (artifact.program.instructions.entries.items) |entry| {
+        const inst = entry orelse continue;
+        try std.testing.expect(inst.operation != .array_length);
+    }
+}

src/intel/compiler/targets/gen9/register_allocation.zig

@@ -68,6 +68,10 @@ fn reserveExistingPhysicalRegisters(program: *const program_ir.Program, initial:
                 reserveRegister(&next_byte, op.byte_offset.register, grf_size);
                 reserveRegister(&next_byte, op.source.register, grf_size);
             },
+            .array_length => |op| {
+                reserveRegister(&next_byte, op.destination.register, grf_size);
+                reserveRegister(&next_byte, op.byte_offset.register, grf_size);
+            },
             .surface_read => |op| {
                 reserveRegister(&next_byte, op.destination.register, grf_size);
                 reserveRegister(&next_byte, op.address.register, grf_size);
@@ -94,6 +98,11 @@ fn reserveExistingPhysicalRegisters(program: *const program_ir.Program, initial:
                 reserveRegister(&next_byte, op.lhs.register, grf_size);
                 reserveRegister(&next_byte, op.rhs.register, grf_size);
             },
+            .math => |op| {
+                reserveRegister(&next_byte, op.destination.register, grf_size);
+                reserveRegister(&next_byte, op.lhs.register, grf_size);
+                reserveRegister(&next_byte, op.rhs.register, grf_size);
+            },
             .parallel_copy => return Error.ParallelCopiesNotLowered,
         }
     }
@@ -125,6 +134,10 @@ fn rewriteProgram(program: *program_ir.Program, allocations: []const ?operand.Ph
                 try rewriteSource(program, &op.byte_offset, allocations);
                 try rewriteSource(program, &op.source, allocations);
             },
+            .array_length => |*op| {
+                try rewriteDestination(program, &op.destination, allocations);
+                try rewriteSource(program, &op.byte_offset, allocations);
+            },
             .surface_read => |*op| {
                 try rewriteDestination(program, &op.destination, allocations);
                 try rewriteSource(program, &op.address, allocations);
@@ -151,6 +164,11 @@ fn rewriteProgram(program: *program_ir.Program, allocations: []const ?operand.Ph
                 try rewriteSource(program, &op.lhs, allocations);
                 try rewriteSource(program, &op.rhs, allocations);
             },
+            .math => |*op| {
+                try rewriteDestination(program, &op.destination, allocations);
+                try rewriteSource(program, &op.lhs, allocations);
+                try rewriteSource(program, &op.rhs, allocations);
+            },
             .parallel_copy => return Error.ParallelCopiesNotLowered,
         }
     }

src/intel/compiler/targets/gen9/validator.zig

@@ -64,6 +64,11 @@ fn validateInstruction(inst: instruction.Instruction) Error!void {
             try validateSource(op.byte_offset);
             try validateSource(op.source);
         },
+        .array_length => |op| {
+            try validateBufferReference(op.buffer);
+            try validateDestination(op.destination);
+            try validateSource(op.byte_offset);
+        },
         .surface_read => |op| {
             try validateBindingTableIndex(op.binding_table);
             try validateDestination(op.destination);
@@ -89,6 +94,11 @@ fn validateInstruction(inst: instruction.Instruction) Error!void {
             try validateSource(op.lhs);
             try validateSource(op.rhs);
         },
+        .math => |op| {
+            try validateDestination(op.destination);
+            try validateSource(op.lhs);
+            try validateSource(op.rhs);
+        },
         .parallel_copy => |copy| {
             for (copy.register_copies) |item| {
                 try validateDestination(item.destination);