From d357b722ba759d7328ffc0af2f49b19d9ce32840 Mon Sep 17 00:00:00 2001 From: Kbz-8 Date: Sun, 30 Aug 2026 19:57:09 +0200 Subject: [PATCH] [Flint] adding array_length and some math operations encoding --- src/compiler/ir/parser/ast.zig | 1 + src/compiler/ir/parser/lower.zig | 13 ++ src/compiler/ir/parser/parser.zig | 20 +++ src/intel/FlintCommandBuffer.zig | 15 +- src/intel/FlintPipeline.zig | 13 -- src/intel/compiler/ir/instruction.zig | 20 +++ src/intel/compiler/ir/printer.zig | 17 ++ src/intel/compiler/ir/validator.zig | 86 +++++---- src/intel/compiler/lower/common_ir.zig | 170 ++++++++++++++++-- .../gen9/compute/array_length_lowering.zig | 92 ++++++++++ .../targets/gen9/compute/dispatch.zig | 28 ++- .../targets/gen9/compute/eu_encoder.zig | 131 +++++++++++--- .../targets/gen9/compute/kernel_encoder.zig | 27 ++- .../targets/gen9/compute/pipeline.zig | 31 ++-- .../gen9/compute/resource_lowering.zig | 2 + src/intel/compiler/targets/gen9/gen9.zig | 37 ++++ .../targets/gen9/register_allocation.zig | 18 ++ src/intel/compiler/targets/gen9/validator.zig | 10 ++ 18 files changed, 622 insertions(+), 109 deletions(-) create mode 100644 src/intel/compiler/targets/gen9/compute/array_length_lowering.zig diff --git a/src/compiler/ir/parser/ast.zig b/src/compiler/ir/parser/ast.zig index f18446f..e3665b7 100644 --- a/src/compiler/ir/parser/ast.zig +++ b/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 }, }; diff --git a/src/compiler/ir/parser/lower.zig b/src/compiler/ir/parser/lower.zig index 537af13..9820af1 100644 --- a/src/compiler/ir/parser/lower.zig +++ b/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; diff --git a/src/compiler/ir/parser/parser.zig b/src/compiler/ir/parser/parser.zig index 8eff048..854d31a 100644 --- a/src/compiler/ir/parser/parser.zig +++ b/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; diff --git a/src/intel/FlintCommandBuffer.zig b/src/intel/FlintCommandBuffer.zig index 3ff346a..82f3984 100644 --- a/src/intel/FlintCommandBuffer.zig +++ b/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 | diff --git a/src/intel/FlintPipeline.zig b/src/intel/FlintPipeline.zig index a615eeb..b5c3dbb 100644 --- a/src/intel/FlintPipeline.zig +++ b/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; diff --git a/src/intel/compiler/ir/instruction.zig b/src/intel/compiler/ir/instruction.zig index c8d4e9f..06f23a0 100644 --- a/src/intel/compiler/ir/instruction.zig +++ b/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, }; diff --git a/src/intel/compiler/ir/printer.zig b/src/intel/compiler/ir/printer.zig index e33bb02..e9561b8 100644 --- a/src/intel/compiler/ir/printer.zig +++ b/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), } } diff --git a/src/intel/compiler/ir/validator.zig b/src/intel/compiler/ir/validator.zig index fa4d992..4ff486e 100644 --- a/src/intel/compiler/ir/validator.zig +++ b/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) diff --git a/src/intel/compiler/lower/common_ir.zig b/src/intel/compiler/lower/common_ir.zig index c8598da..d59bcc0 100644 --- a/src/intel/compiler/lower/common_ir.zig +++ b/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) - return Error.UnsupportedOperation, + .builtin => |builtin| switch (builtin) { + .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 \\{ diff --git a/src/intel/compiler/targets/gen9/compute/array_length_lowering.zig b/src/intel/compiler/targets/gen9/compute/array_length_lowering.zig new file mode 100644 index 0000000..2f2455b --- /dev/null +++ b/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, + }; +} diff --git a/src/intel/compiler/targets/gen9/compute/dispatch.zig b/src/intel/compiler/targets/gen9/compute/dispatch.zig index 1cdd839..012d3f7 100644 --- a/src/intel/compiler/targets/gen9/compute/dispatch.zig +++ b/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); diff --git a/src/intel/compiler/targets/gen9/compute/eu_encoder.zig b/src/intel/compiler/targets/gen9/compute/eu_encoder.zig index e7a10e2..4359619 100644 --- a/src/intel/compiler/targets/gen9/compute/eu_encoder.zig +++ b/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), - ); - - 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, - } - + setDestination(&encoded, .grf, try hardwareType(move.destination.type), destination, try horizontalStride(move.destination.region.horizontal_stride)); + try setSource0(&encoded, move.source); 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, diff --git a/src/intel/compiler/targets/gen9/compute/kernel_encoder.zig b/src/intel/compiler/targets/gen9/compute/kernel_encoder.zig index fcd7cb4..8e7829a 100644 --- a/src/intel/compiler/targets/gen9/compute/kernel_encoder.zig +++ b/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), - .surface_message => |message| try eu.encodeSurfaceMessage(instruction.execution_size, message), - else => return Error.UnsupportedOperation, + .move => |move| eu.encodeMove(instruction.execution_size, move), + .surface_message => |message| eu.encodeSurfaceMessage(instruction.execution_size, message), + .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); diff --git a/src/intel/compiler/targets/gen9/compute/pipeline.zig b/src/intel/compiler/targets/gen9/compute/pipeline.zig index 81a9cfa..03a6dcf 100644 --- a/src/intel/compiler/targets/gen9/compute/pipeline.zig +++ b/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( - allocator, - &program, - ); + var resources = try resource_layout.Layout.init(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) { - error.UnsupportedControlFlow, - error.UnsupportedOperation, - error.UnsupportedPredication, - error.UnsupportedExecutionSize, - error.UnsupportedDataType, - error.UnsupportedOperand, - error.EotRegisterUnavailable, - => null, - else => return err, + const kernel = kernel_encoder.encode(allocator, &program) catch |err| encoding_error: { + std.log.scoped(.FlintCompiler).err("Gen9 EU kernel encoding failed: {s}", .{@errorName(err)}); + break :encoding_error switch (err) { + error.UnsupportedControlFlow, + error.UnsupportedOperation, + error.UnsupportedPredication, + error.UnsupportedExecutionSize, + error.UnsupportedDataType, + error.UnsupportedOperand, + error.EotRegisterUnavailable, + => null, + else => return err, + }; }; errdefer if (kernel) |bytes| allocator.free(bytes); diff --git a/src/intel/compiler/targets/gen9/compute/resource_lowering.zig b/src/intel/compiler/targets/gen9/compute/resource_lowering.zig index 7c91e90..ea87a05 100644 --- a/src/intel/compiler/targets/gen9/compute/resource_lowering.zig +++ b/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, }; } diff --git a/src/intel/compiler/targets/gen9/gen9.zig b/src/intel/compiler/targets/gen9/gen9.zig index 2dde287..58975eb 100644 --- a/src/intel/compiler/targets/gen9/gen9.zig +++ b/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); + } +} diff --git a/src/intel/compiler/targets/gen9/register_allocation.zig b/src/intel/compiler/targets/gen9/register_allocation.zig index 3f55ce2..32a4f26 100644 --- a/src/intel/compiler/targets/gen9/register_allocation.zig +++ b/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, } } diff --git a/src/intel/compiler/targets/gen9/validator.zig b/src/intel/compiler/targets/gen9/validator.zig index 2f49708..161bf38 100644 --- a/src/intel/compiler/targets/gen9/validator.zig +++ b/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);