From 04ff263b40f76d427488485c006fda55c6e5b504 Mon Sep 17 00:00:00 2001 From: Kbz-8 Date: Fri, 14 Aug 2026 20:41:31 +0200 Subject: [PATCH] [Flint] adding compute resource layout lowering --- src/intel/FlintCommandBuffer.zig | 55 +++++-- src/intel/FlintPipeline.zig | 127 +++++++++++----- src/intel/compiler/compiler.zig | 2 +- src/intel/compiler/ir/instruction.zig | 9 +- src/intel/compiler/ir/printer.zig | 11 +- src/intel/compiler/ir/validator.zig | 32 ++-- src/intel/compiler/lower/common_ir.zig | 4 +- .../compiler/targets/gen9/compute/compute.zig | 4 + .../targets/gen9/compute/resource_layout.zig | 141 ++++++++++++++++++ .../gen9/compute/resource_lowering.zig | 127 ++++++++++++++++ src/intel/compiler/targets/gen9/gen9.zig | 7 + src/intel/compiler/targets/gen9/validator.zig | 11 ++ src/intel/compiler/targets/targets.zig | 18 +++ 13 files changed, 484 insertions(+), 64 deletions(-) create mode 100644 src/intel/compiler/targets/gen9/compute/resource_layout.zig create mode 100644 src/intel/compiler/targets/gen9/compute/resource_lowering.zig diff --git a/src/intel/FlintCommandBuffer.zig b/src/intel/FlintCommandBuffer.zig index 8bc86a4..e6fdc23 100644 --- a/src/intel/FlintCommandBuffer.zig +++ b/src/intel/FlintCommandBuffer.zig @@ -174,16 +174,24 @@ pub fn beginRenderPass(interface: *Interface, render_pass: *base.RenderPass, fra pub fn bindDescriptorSets(interface: *Interface, bind_point: vk.PipelineBindPoint, first_set: u32, sets: [base.vulkan_max_descriptor_sets]?*base.DescriptorSet, dynamic_offsets: []const u32) VkError!void { const self: *Self = @alignCast(@fieldParentPtr("interface", interface)); - if (bind_point != .compute) return; - if (first_set >= base.vulkan_max_descriptor_sets) return VkError.ValidationFailed; + if (bind_point != .compute) + return; + + if (dynamic_offsets.len != 0) + return VkError.FeatureNotPresent; + + if (first_set >= base.vulkan_max_descriptor_sets) + return VkError.ValidationFailed; for (sets, 0..) |set, index| { const base_set = set orelse break; const destination = first_set + index; - if (destination >= base.vulkan_max_descriptor_sets) return VkError.ValidationFailed; + + if (destination >= base.vulkan_max_descriptor_sets) + return VkError.ValidationFailed; + self.bound_compute_descriptor_sets[destination] = @alignCast(@fieldParentPtr("interface", base_set)); } - _ = dynamic_offsets; } pub fn bindPipeline(interface: *Interface, bind_point: vk.PipelineBindPoint, pipeline: *base.Pipeline) VkError!void { @@ -289,13 +297,38 @@ pub fn dispatch(interface: *Interface, group_count_x: u32, group_count_y: u32, g } pub fn dispatchBase(interface: *Interface, base_group_x: u32, base_group_y: u32, base_group_z: u32, group_count_x: u32, group_count_y: u32, group_count_z: u32) VkError!void { - _ = interface; - _ = base_group_x; - _ = base_group_y; - _ = base_group_z; - _ = group_count_x; - _ = group_count_y; - _ = group_count_z; + const self: *Self = @alignCast(@fieldParentPtr("interface", interface)); + if (group_count_x == 0 or group_count_y == 0 or group_count_z == 0) + return; + + inline for ([_]struct { u32, u32 }{ + .{ base_group_x, group_count_x }, + .{ base_group_y, group_count_y }, + .{ base_group_z, group_count_z }, + }) |dimension| { + const group_end = std.math.add(u32, dimension[0], dimension[1]) catch return VkError.ValidationFailed; + if (group_end > 65535) + return VkError.ValidationFailed; + } + + const pipeline = self.bound_compute_pipeline orelse return VkError.ValidationFailed; + const artifact = pipeline.computeArtifact() orelse return VkError.FeatureNotPresent; + for (artifact.resources.bindings) |resource| { + if (resource.set >= base.vulkan_max_descriptor_sets) + return VkError.ValidationFailed; + + const descriptor_set = self.bound_compute_descriptor_sets[resource.set] orelse return VkError.ValidationFailed; + const expected_layout = pipeline.interface.layout.set_layouts[resource.set] orelse return VkError.ValidationFailed; + + if (descriptor_set.interface.layout != expected_layout) + return VkError.ValidationFailed; + + const descriptor = try descriptor_set.getBuffer(resource.binding, 0); + const buffer = descriptor.buffer orelse return VkError.ValidationFailed; + + if (!buffer.usage.storage_buffer_bit or buffer.memory == null) + return VkError.ValidationFailed; + } } pub fn setDeviceMask(interface: *Interface, device_mask: u32) VkError!void { diff --git a/src/intel/FlintPipeline.zig b/src/intel/FlintPipeline.zig index bc51b45..0b0275c 100644 --- a/src/intel/FlintPipeline.zig +++ b/src/intel/FlintPipeline.zig @@ -15,14 +15,25 @@ const PipelineKind = enum { compute, }; +pub const ComputeArtifact = struct { + program: compiler.Program, + resources: compiler.targets.ComputeResourceLayout, + + fn deinit(self: *ComputeArtifact, allocator: std.mem.Allocator) void { + self.resources.deinit(allocator); + self.program.deinit(); + self.* = undefined; + } +}; + const CommonStage = struct { stage: shader_ir.ir.module.Stage, module: base.ShaderModule.IrModule, - program: ?compiler.Program, + artifact: ?ComputeArtifact, - fn deinit(self: *CommonStage) void { - if (self.program) |*program| - program.deinit(); + fn deinit(self: *CommonStage, allocator: std.mem.Allocator) void { + if (self.artifact) |*artifact| + artifact.deinit(allocator); self.module.deinit(); self.* = undefined; } @@ -48,6 +59,8 @@ pub fn createCompute(device: *base.Device, allocator: std.mem.Allocator, cache: initialized = true; self.stages = try compileStages(self.artifact_allocator.allocator(), &.{info.stage}, .compute, compilerDeviceInfo(device)); + if (self.computeArtifact()) |artifact| + try validateComputePipelineLayout(self.interface.layout, &artifact.resources); return self; } @@ -74,12 +87,7 @@ pub fn createGraphics(device: *base.Device, allocator: std.mem.Allocator, cache: return self; } -fn compileStages( - allocator: std.mem.Allocator, - infos: []const vk.PipelineShaderStageCreateInfo, - pipeline_kind: PipelineKind, - device_info: ?compiler.device.DeviceInfo, -) VkError![]CommonStage { +fn compileStages(allocator: std.mem.Allocator, infos: []const vk.PipelineShaderStageCreateInfo, pipeline_kind: PipelineKind, device_info: ?compiler.device.DeviceInfo) VkError![]CommonStage { if (infos.len == 0) return VkError.ValidationFailed; @@ -87,7 +95,7 @@ fn compileStages( var initialized: usize = 0; errdefer { for (stages[0..initialized]) |*stage| - stage.deinit(); + stage.deinit(allocator); allocator.free(stages); } @@ -98,12 +106,7 @@ fn compileStages( return stages; } -fn compileStage( - allocator: std.mem.Allocator, - info: *const vk.PipelineShaderStageCreateInfo, - pipeline_kind: PipelineKind, - device_info: ?compiler.device.DeviceInfo, -) VkError!CommonStage { +fn compileStage(allocator: std.mem.Allocator, info: *const vk.PipelineShaderStageCreateInfo, pipeline_kind: PipelineKind, device_info: ?compiler.device.DeviceInfo) VkError!CommonStage { const specializations = try specializationValues(allocator, info.p_specialization_info); defer if (specializations.len != 0) allocator.free(specializations); @@ -129,24 +132,21 @@ fn compileStage( std.debug.assert(module.stage == expected_stage); - var program = try lowerToFlint(allocator, &module, device_info); - errdefer if (program) |*value| value.deinit(); + var artifact = try lowerToFlint(allocator, &module, device_info); + errdefer if (artifact) |*value| value.deinit(allocator); return .{ .stage = expected_stage, .module = module, - .program = program, + .artifact = artifact, }; } -fn lowerToFlint( - allocator: std.mem.Allocator, - module: *base.ShaderModule.IrModule, - device_info: ?compiler.device.DeviceInfo, -) VkError!?compiler.Program { +fn lowerToFlint(allocator: std.mem.Allocator, module: *base.ShaderModule.IrModule, device_info: ?compiler.device.DeviceInfo) VkError!?ComputeArtifact { const target = device_info orelse return null; - const program = compiler.targets.lower(allocator, module, target, .{}) catch |err| switch (err) { + var program = compiler.targets.lower(allocator, module, target, .{}) catch |err| switch (err) { error.OutOfMemory => return VkError.OutOfHostMemory, + error.UnsupportedGeneration, error.UnsupportedStage, error.UnsupportedDispatchWidth, @@ -157,12 +157,60 @@ fn lowerToFlint( error.UnsupportedOperation, error.UnsupportedTerminator, => return null, + else => { - std.log.scoped(.FlintPipeline).err("Flint shader lowering failed: {s}", .{@errorName(err)}); + std.log.scoped(.FlintPipeline).err("shader lowering failed: {s}", .{@errorName(err)}); return VkError.ValidationFailed; }, }; - return program; + errdefer program.deinit(); + + var resources = compiler.targets.layoutComputeResources(allocator, &program) catch |err| switch (err) { + error.OutOfMemory => return VkError.OutOfHostMemory, + error.UnsupportedGeneration, + error.TooManyStorageBuffers, + => { + program.deinit(); + return null; + }, + }; + errdefer resources.deinit(allocator); + + compiler.targets.lowerComputeResources(&program, &resources) catch |err| switch (err) { + error.UnsupportedGeneration => { + resources.deinit(allocator); + program.deinit(); + return null; + }, + error.InvalidProgram, + error.InvalidResourceLayout, + => { + std.log.scoped(.FlintPipeline).err("Flint compute resource lowering failed: {s}", .{@errorName(err)}); + return VkError.ValidationFailed; + }, + }; + + return .{ + .program = program, + .resources = resources, + }; +} + +fn validateComputePipelineLayout(layout: *const base.PipelineLayout, resources: *const compiler.targets.ComputeResourceLayout) VkError!void { + for (resources.bindings) |resource| { + if (resource.set >= layout.set_count) + return VkError.ValidationFailed; + + const set_layout = layout.set_layouts[resource.set] orelse return VkError.ValidationFailed; + + if (resource.binding >= set_layout.bindings.len) + return VkError.ValidationFailed; + + const binding = set_layout.bindings[resource.binding]; + + if (binding.descriptor_type != .storage_buffer or binding.array_size == 0) + return VkError.ValidationFailed; + } } fn compilerDeviceInfo(device: *const base.Device) ?compiler.device.DeviceInfo { @@ -216,11 +264,17 @@ fn commonStage(stage: vk.ShaderStageFlags) ?shader_ir.ir.module.Stage { fn deinitStages(allocator: std.mem.Allocator, stages: []CommonStage) void { for (stages) |*stage| - stage.deinit(); + stage.deinit(allocator); if (stages.len != 0) allocator.free(stages); } +pub fn computeArtifact(self: *const Self) ?*const ComputeArtifact { + if (self.interface.bind_point != .compute or self.stages.len != 1 or self.stages[0].stage != .compute) + return null; + return if (self.stages[0].artifact) |*artifact| artifact else null; +} + pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void { const self: *Self = @alignCast(@fieldParentPtr("interface", interface)); deinitStages(self.artifact_allocator.allocator(), self.stages); @@ -259,20 +313,23 @@ test "Flint pipeline: lower common compute IR" { }, null); try builder.setTerminator(entry, .return_void); - var program = (try lowerToFlint(std.testing.allocator, &module, device_info)).?; - defer program.deinit(); + var artifact = (try lowerToFlint(std.testing.allocator, &module, device_info)).?; + defer artifact.deinit(std.testing.allocator); + const program = &artifact.program; try std.testing.expect(program.properties.common_ir_lowered); try std.testing.expect(program.properties.block_parameters_lowered); try std.testing.expect(!program.properties.system_values_lowered); - try std.testing.expect(!program.properties.resources_lowered); + try std.testing.expect(program.properties.resources_lowered); try std.testing.expect(!program.properties.instructions_selected); try std.testing.expectEqual([3]u32{ 1, 1, 1 }, program.workgroup_size); try std.testing.expectEqual(@as(usize, 1), program.storage_buffers.entries.items.len); - try compiler.targets.validate(&program); + try std.testing.expectEqual(@as(usize, 1), artifact.resources.bindings.len); + try std.testing.expectEqual(@as(u8, 0), artifact.resources.bindings[0].binding_table_index); + try compiler.targets.validate(program); - const text = try compiler.printer.allocPrint(std.testing.allocator, &program); + const text = try compiler.printer.allocPrint(std.testing.allocator, program); defer std.testing.allocator.free(text); try std.testing.expect(std.mem.indexOf(u8, text, "load_global_invocation_id %id_x:u32, component(0)") != null); - try std.testing.expect(std.mem.indexOf(u8, text, "store_buffer @storage, 0:u32, %id_x:u32") != null); + try std.testing.expect(std.mem.indexOf(u8, text, "store_buffer bti(0), 0:u32, %id_x:u32") != null); } diff --git a/src/intel/compiler/compiler.zig b/src/intel/compiler/compiler.zig index 7fbf221..b2991cb 100644 --- a/src/intel/compiler/compiler.zig +++ b/src/intel/compiler/compiler.zig @@ -51,7 +51,7 @@ test "[ir] basic compute shader" { }); _ = try builder.appendInstruction(entry, .simd8, null, .{ .store_buffer = .{ - .buffer = storage, + .buffer = .{ .logical = storage }, .byte_offset = .{ .register = .{ .immediate = .{ .u32 = 0 } }, .type = .u32, diff --git a/src/intel/compiler/ir/instruction.zig b/src/intel/compiler/ir/instruction.zig index 2c2b6e7..e55cc63 100644 --- a/src/intel/compiler/ir/instruction.zig +++ b/src/intel/compiler/ir/instruction.zig @@ -9,15 +9,20 @@ pub const LoadGlobalInvocationId = struct { component: u8, }; +pub const BufferReference = union(enum) { + logical: ids.StorageBufferId, + binding_table: u8, +}; + pub const LoadBuffer = struct { destination: operand.Destination, - buffer: ids.StorageBufferId, + buffer: BufferReference, byte_offset: operand.Source, immediate_offset: u32 = 0, }; pub const StoreBuffer = struct { - buffer: ids.StorageBufferId, + buffer: BufferReference, byte_offset: operand.Source, immediate_offset: u32 = 0, source: operand.Source, diff --git a/src/intel/compiler/ir/printer.zig b/src/intel/compiler/ir/printer.zig index a3e9a97..15af158 100644 --- a/src/intel/compiler/ir/printer.zig +++ b/src/intel/compiler/ir/printer.zig @@ -120,7 +120,7 @@ fn writeOperation(program: *const program_ir.Program, writer: *std.Io.Writer, ex try writer.writeAll("load_buffer "); try writeDestination(program, writer, execution_size, op.destination); try writer.writeAll(", "); - try writeStorageBufferRef(program, writer, op.buffer); + try writeBufferReference(program, writer, op.buffer); try writer.writeAll(", "); try writeSource(program, writer, execution_size, op.byte_offset); if (op.immediate_offset != 0) @@ -128,7 +128,7 @@ fn writeOperation(program: *const program_ir.Program, writer: *std.Io.Writer, ex }, .store_buffer => |op| { try writer.writeAll("store_buffer "); - try writeStorageBufferRef(program, writer, op.buffer); + try writeBufferReference(program, writer, op.buffer); try writer.writeAll(", "); try writeSource(program, writer, execution_size, op.byte_offset); if (op.immediate_offset != 0) @@ -356,6 +356,13 @@ fn writeFlagRef(program: *const program_ir.Program, writer: *std.Io.Writer, flag } } +fn writeBufferReference(program: *const program_ir.Program, writer: *std.Io.Writer, reference: inst_ir.BufferReference) !void { + switch (reference) { + .logical => |buffer| try writeStorageBufferRef(program, writer, buffer), + .binding_table => |index| try writer.print("bti({d})", .{index}), + } +} + fn writeStorageBufferRef(program: *const program_ir.Program, writer: *std.Io.Writer, buffer_id: ids.StorageBufferId) !void { const buffer = program.storage_buffers.get(buffer_id); try writeNamedRef(writer, if (buffer) |value| value.name else null, "buffer", buffer_id.index(), '@'); diff --git a/src/intel/compiler/ir/validator.zig b/src/intel/compiler/ir/validator.zig index eaf97a2..e7b34a9 100644 --- a/src/intel/compiler/ir/validator.zig +++ b/src/intel/compiler/ir/validator.zig @@ -19,6 +19,7 @@ pub const Error = error{ InvalidDestination, InvalidImmediateType, InvalidStorageBuffer, + InvalidBufferReference, InvalidGlobalInvocationId, InvalidBufferAccess, InvalidWorkgroupSize, @@ -136,20 +137,14 @@ fn validateInstruction(program: *const program_ir.Program, inst: instruction.Ins return Error.InvalidGlobalInvocationId; }, .load_buffer => |op| { - if (program.properties.resources_lowered) - return Error.UnloweredResource; - if (!program.storage_buffers.isLive(op.buffer)) - return Error.InvalidStorageBuffer; + try validateBufferReference(program, op.buffer); try validateDestination(program, op.destination); try validateBufferOffset(program, op.byte_offset); if (!op.destination.type.isInitialTargetType()) return Error.InvalidBufferAccess; }, .store_buffer => |op| { - if (program.properties.resources_lowered) - return Error.UnloweredResource; - if (!program.storage_buffers.isLive(op.buffer)) - return Error.InvalidStorageBuffer; + try validateBufferReference(program, op.buffer); try validateBufferOffset(program, op.byte_offset); try validateSource(program, op.source); if (!op.source.type.isInitialTargetType()) @@ -180,6 +175,19 @@ fn validateInstruction(program: *const program_ir.Program, inst: instruction.Ins } } +fn validateBufferReference(program: *const program_ir.Program, reference: instruction.BufferReference) Error!void { + switch (reference) { + .logical => |buffer| { + if (program.properties.resources_lowered) + return Error.UnloweredResource; + if (!program.storage_buffers.isLive(buffer)) + return Error.InvalidStorageBuffer; + }, + .binding_table => if (!program.properties.resources_lowered) + return Error.InvalidBufferReference, + } +} + fn validateBufferOffset(program: *const program_ir.Program, source: operand.Source) Error!void { try validateSource(program, source); if (source.type != .u32) @@ -408,7 +416,7 @@ test "[ir] validator checks compute system values and resources" { const buffer_load_id = try builder.appendInstruction(entry, .simd8, null, .{ .load_buffer = .{ .destination = .{ .register = .{ .virtual = register }, .type = .u32 }, - .buffer = buffer, + .buffer = .{ .logical = buffer }, .byte_offset = .{ .register = .{ .immediate = .{ .u32 = 0 } }, .type = .u32, @@ -427,10 +435,12 @@ test "[ir] validator checks compute system values and resources" { try std.testing.expectError(Error.UnloweredSystemValue, validate(&program)); program.properties.system_values_lowered = false; - program.instructions.getMut(buffer_load_id).?.operation.load_buffer.buffer = ids.StorageBufferId.fromIndex(99); + program.instructions.getMut(buffer_load_id).?.operation.load_buffer.buffer = .{ .logical = ids.StorageBufferId.fromIndex(99) }; try std.testing.expectError(Error.InvalidStorageBuffer, validate(&program)); - program.instructions.getMut(buffer_load_id).?.operation.load_buffer.buffer = buffer; + program.instructions.getMut(buffer_load_id).?.operation.load_buffer.buffer = .{ .logical = buffer }; program.properties.resources_lowered = true; try std.testing.expectError(Error.UnloweredResource, validate(&program)); + program.instructions.getMut(buffer_load_id).?.operation.load_buffer.buffer = .{ .binding_table = 0 }; + try validate(&program); } diff --git a/src/intel/compiler/lower/common_ir.zig b/src/intel/compiler/lower/common_ir.zig index 5b2c6a1..39ce6aa 100644 --- a/src/intel/compiler/lower/common_ir.zig +++ b/src/intel/compiler/lower/common_ir.zig @@ -731,7 +731,7 @@ const LoweringState = struct { try self.appendInstruction(block_id, null, .{ .load_buffer = .{ .destination = try destinationFromSource(result_component), - .buffer = buffer, + .buffer = .{ .logical = buffer }, .byte_offset = byte_offset, .immediate_offset = @intCast(component_index * result_component.type.sizeBytes()), }, @@ -749,7 +749,7 @@ const LoweringState = struct { for (source_components, 0..) |source_component, component_index| { try self.appendInstruction(block_id, null, .{ .store_buffer = .{ - .buffer = buffer, + .buffer = .{ .logical = buffer }, .byte_offset = byte_offset, .immediate_offset = @intCast(component_index * source_component.type.sizeBytes()), .source = source_component, diff --git a/src/intel/compiler/targets/gen9/compute/compute.zig b/src/intel/compiler/targets/gen9/compute/compute.zig index e9c7640..cedc13f 100644 --- a/src/intel/compiler/targets/gen9/compute/compute.zig +++ b/src/intel/compiler/targets/gen9/compute/compute.zig @@ -1,5 +1,9 @@ const std = @import("std"); +pub const resource_layout = @import("resource_layout.zig"); +pub const resource_lowering = @import("resource_lowering.zig"); +pub const ResourceLayout = resource_layout.Layout; + pub const Error = error{UnsupportedWorkgroupSize}; pub fn validateWorkgroupSize(size: [3]u32) Error!void { diff --git a/src/intel/compiler/targets/gen9/compute/resource_layout.zig b/src/intel/compiler/targets/gen9/compute/resource_layout.zig new file mode 100644 index 0000000..3769e35 --- /dev/null +++ b/src/intel/compiler/targets/gen9/compute/resource_layout.zig @@ -0,0 +1,141 @@ +const std = @import("std"); +const ids = @import("../../../ir/id.zig"); +const program_ir = @import("../../../ir/program.zig"); + +pub const max_storage_buffers: usize = 4; + +pub const Error = std.mem.Allocator.Error || error{ + TooManyStorageBuffers, +}; + +pub const Binding = struct { + set: u32, + binding: u32, + binding_table_index: u8, +}; + +const Candidate = struct { + resource: ids.StorageBufferId, + set: u32, + binding: u32, +}; + +pub const Layout = struct { + bindings: []Binding, + resource_indices: []?u8, + + pub fn init(allocator: std.mem.Allocator, program: *const program_ir.Program) Error!Layout { + var candidates: std.ArrayList(Candidate) = .empty; + defer candidates.deinit(allocator); + + for (program.storage_buffers.entries.items, 0..) |entry, index| { + const buffer = entry orelse continue; + try candidates.append(allocator, .{ + .resource = ids.StorageBufferId.fromIndex(index), + .set = buffer.set, + .binding = buffer.binding, + }); + } + std.mem.sort(Candidate, candidates.items, {}, lessThan); + + var unique_count: usize = 0; + for (candidates.items, 0..) |candidate, index| { + if (index == 0 or candidate.set != candidates.items[index - 1].set or candidate.binding != candidates.items[index - 1].binding) + unique_count += 1; + } + if (unique_count > max_storage_buffers) + return Error.TooManyStorageBuffers; + + const bindings = try allocator.alloc(Binding, unique_count); + errdefer allocator.free(bindings); + const resource_indices = try allocator.alloc(?u8, program.storage_buffers.entries.items.len); + errdefer allocator.free(resource_indices); + @memset(resource_indices, null); + + var binding_index: usize = 0; + for (candidates.items, 0..) |candidate, index| { + if (index == 0 or candidate.set != candidates.items[index - 1].set or candidate.binding != candidates.items[index - 1].binding) { + bindings[binding_index] = .{ + .set = candidate.set, + .binding = candidate.binding, + .binding_table_index = @intCast(binding_index), + }; + binding_index += 1; + } + resource_indices[candidate.resource.index()] = @intCast(binding_index - 1); + } + std.debug.assert(binding_index == bindings.len); + + return .{ + .bindings = bindings, + .resource_indices = resource_indices, + }; + } + + pub fn deinit(self: *Layout, allocator: std.mem.Allocator) void { + allocator.free(self.bindings); + allocator.free(self.resource_indices); + self.* = undefined; + } + + pub fn bindingTableIndex(self: *const Layout, resource: ids.StorageBufferId) ?u8 { + if (resource.index() >= self.resource_indices.len) + return null; + return self.resource_indices[resource.index()]; + } +}; + +fn lessThan(_: void, lhs: Candidate, rhs: Candidate) bool { + if (lhs.set != rhs.set) + return lhs.set < rhs.set; + if (lhs.binding != rhs.binding) + return lhs.binding < rhs.binding; + return lhs.resource.index() < rhs.resource.index(); +} + +test "[gen9] compute resource layout: assign stable binding-table indices" { + const device = @import("../../../device.zig"); + + const device_info: device.DeviceInfo = .{ + .generation = .gen9, + .platform = .skylake, + .pci_device_id = 0x1912, + .grf_count = 128, + }; + var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, device_info, .simd8); + defer program.deinit(); + + const third = try program.addStorageBuffer(.{ .set = 2, .binding = 7 }); + const first = try program.addStorageBuffer(.{ .set = 0, .binding = 3 }); + const alias = try program.addStorageBuffer(.{ .set = 0, .binding = 3 }); + const second = try program.addStorageBuffer(.{ .set = 1, .binding = 0 }); + + var layout = try Layout.init(std.testing.allocator, &program); + defer layout.deinit(std.testing.allocator); + + try std.testing.expectEqual(@as(usize, 3), layout.bindings.len); + try std.testing.expectEqual(Binding{ .set = 0, .binding = 3, .binding_table_index = 0 }, layout.bindings[0]); + try std.testing.expectEqual(Binding{ .set = 1, .binding = 0, .binding_table_index = 1 }, layout.bindings[1]); + try std.testing.expectEqual(Binding{ .set = 2, .binding = 7, .binding_table_index = 2 }, layout.bindings[2]); + try std.testing.expectEqual(@as(?u8, 0), layout.bindingTableIndex(first)); + try std.testing.expectEqual(@as(?u8, 0), layout.bindingTableIndex(alias)); + try std.testing.expectEqual(@as(?u8, 1), layout.bindingTableIndex(second)); + try std.testing.expectEqual(@as(?u8, 2), layout.bindingTableIndex(third)); +} + +test "[gen9] compute resource layout: enforce advertised storage-buffer limit" { + const device = @import("../../../device.zig"); + + const device_info: device.DeviceInfo = .{ + .generation = .gen9, + .platform = .skylake, + .pci_device_id = 0x1912, + .grf_count = 128, + }; + var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, device_info, .simd8); + defer program.deinit(); + for (0..max_storage_buffers + 1) |binding| + _ = try program.addStorageBuffer(.{ .set = 0, .binding = @intCast(binding) }); + + try std.testing.expectError(Error.TooManyStorageBuffers, Layout.init(std.testing.allocator, &program)); +} diff --git a/src/intel/compiler/targets/gen9/compute/resource_lowering.zig b/src/intel/compiler/targets/gen9/compute/resource_lowering.zig new file mode 100644 index 0000000..7c91e90 --- /dev/null +++ b/src/intel/compiler/targets/gen9/compute/resource_lowering.zig @@ -0,0 +1,127 @@ +const instruction = @import("../../../ir/instruction.zig"); +const program_ir = @import("../../../ir/program.zig"); +const validator = @import("../../../ir/validator.zig"); +const resource_layout = @import("resource_layout.zig"); + +pub const Error = error{ + InvalidProgram, + InvalidResourceLayout, +}; + +pub fn run(program: *program_ir.Program, layout: *const resource_layout.Layout) Error!void { + validator.validate(program) catch return Error.InvalidProgram; + if (program.properties.resources_lowered) + return; + if (layout.resource_indices.len != program.storage_buffers.entries.items.len) + return Error.InvalidResourceLayout; + + for (program.blocks.entries.items) |block_entry| { + const block = block_entry orelse continue; + for (block.instructions.items) |instruction_id| { + const inst = program.instructions.get(instruction_id) orelse return Error.InvalidProgram; + const reference = bufferReference(inst.operation) orelse continue; + const resource = switch (reference) { + .logical => |value| value, + .binding_table => return Error.InvalidProgram, + }; + const binding_table_index = layout.bindingTableIndex(resource) orelse return Error.InvalidResourceLayout; + if (binding_table_index >= layout.bindings.len) + return Error.InvalidResourceLayout; + const buffer = program.storage_buffers.get(resource) orelse return Error.InvalidProgram; + const binding = layout.bindings[binding_table_index]; + if (binding.binding_table_index != binding_table_index or binding.set != buffer.set or binding.binding != buffer.binding) + return Error.InvalidResourceLayout; + } + } + + for (program.blocks.entries.items) |block_entry| { + const block = block_entry orelse continue; + for (block.instructions.items) |instruction_id| { + const inst = program.instructions.getMut(instruction_id) orelse unreachable; + const reference = bufferReferenceMut(&inst.operation) orelse continue; + const resource = reference.logical; + const binding_table_index = layout.bindingTableIndex(resource).?; + reference.* = .{ .binding_table = binding_table_index }; + } + } + + program.properties.resources_lowered = true; + validator.validate(program) catch return Error.InvalidProgram; +} + +fn bufferReference(operation: instruction.Operation) ?instruction.BufferReference { + return switch (operation) { + .load_buffer => |op| op.buffer, + .store_buffer => |op| op.buffer, + else => null, + }; +} + +fn bufferReferenceMut(operation: *instruction.Operation) ?*instruction.BufferReference { + return switch (operation.*) { + .load_buffer => |*op| &op.buffer, + .store_buffer => |*op| &op.buffer, + else => null, + }; +} + +test "[gen9] compute resource lowering: resolve logical buffers" { + const std = @import("std"); + const Builder = @import("../../../ir/Builder.zig"); + const device = @import("../../../device.zig"); + const operand = @import("../../../ir/operand.zig"); + const printer = @import("../../../ir/printer.zig"); + + const device_info: device.DeviceInfo = .{ + .generation = .gen9, + .platform = .skylake, + .pci_device_id = 0x1912, + .grf_count = 128, + }; + var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, device_info, .simd8); + defer program.deinit(); + var builder = Builder.init(&program); + + const value = try builder.addVirtualRegister(.{ + .size_bytes = 32, + .alignment_bytes = 32, + .element_type = .u32, + .lane_count = 8, + .class = .temporary, + }); + const buffer = try builder.addStorageBuffer(.{ .set = 1, .binding = 3, .name = "storage" }); + const entry = try builder.addBlock("entry"); + const store_id = try builder.appendInstruction(entry, .simd8, null, .{ + .store_buffer = .{ + .buffer = .{ .logical = buffer }, + .byte_offset = .{ + .register = .{ .immediate = .{ .u32 = 0 } }, + .type = .u32, + .region = operand.Region.broadcast(), + }, + .source = .{ + .register = .{ .virtual = value }, + .type = .u32, + .region = operand.Region.contiguous(.simd8), + }, + }, + }); + try builder.setTerminator(entry, .end_thread); + + var layout = try resource_layout.Layout.init(std.testing.allocator, &program); + defer layout.deinit(std.testing.allocator); + layout.bindings[0].binding = 4; + try std.testing.expectError(Error.InvalidResourceLayout, run(&program, &layout)); + try std.testing.expect(!program.properties.resources_lowered); + try std.testing.expect(program.instructions.get(store_id).?.operation.store_buffer.buffer == .logical); + layout.bindings[0].binding = 3; + + try run(&program, &layout); + try validator.validate(&program); + + try std.testing.expect(program.properties.resources_lowered); + try std.testing.expectEqual(@as(u8, 0), program.instructions.get(store_id).?.operation.store_buffer.buffer.binding_table); + const text = try printer.allocPrint(std.testing.allocator, &program); + defer std.testing.allocator.free(text); + try std.testing.expect(std.mem.indexOf(u8, text, "store_buffer bti(0), 0:u32") != null); +} diff --git a/src/intel/compiler/targets/gen9/gen9.zig b/src/intel/compiler/targets/gen9/gen9.zig index 4a19780..ce8d553 100644 --- a/src/intel/compiler/targets/gen9/gen9.zig +++ b/src/intel/compiler/targets/gen9/gen9.zig @@ -8,6 +8,8 @@ pub const compute = @import("compute/compute.zig"); pub const validator = @import("validator.zig"); pub const Options = common_ir.Options; +pub const ResourceLoweringError = compute.resource_lowering.Error; + pub const Error = common_ir.Error || compute.Error || error{ UnsupportedGeneration, UnsupportedStage, @@ -38,6 +40,11 @@ pub fn lower( return program; } +pub fn lowerComputeResources(program: *program_ir.Program, layout: *const compute.ResourceLayout) ResourceLoweringError!void { + try compute.resource_lowering.run(program, layout); + validator.validate(program) catch return ResourceLoweringError.InvalidProgram; +} + test "[gen9] target: reject unsupported target configurations" { var module = try shader_ir.parser.parseString(std.testing.allocator, \\shader compute @main diff --git a/src/intel/compiler/targets/gen9/validator.zig b/src/intel/compiler/targets/gen9/validator.zig index 17b7ddd..581ec9a 100644 --- a/src/intel/compiler/targets/gen9/validator.zig +++ b/src/intel/compiler/targets/gen9/validator.zig @@ -12,6 +12,7 @@ pub const Error = shared.Error || compute.Error || error{ UnsupportedExecutionSize, UnsupportedDataType, InvalidPhysicalFlag, + InvalidBindingTableIndex, InvalidPayloadLayout, }; @@ -54,10 +55,12 @@ fn validateInstruction(inst: instruction.Instruction) Error!void { switch (inst.operation) { .load_global_invocation_id => |op| try validateDestination(op.destination), .load_buffer => |op| { + try validateBufferReference(op.buffer); try validateDestination(op.destination); try validateSource(op.byte_offset); }, .store_buffer => |op| { + try validateBufferReference(op.buffer); try validateSource(op.byte_offset); try validateSource(op.source); }, @@ -88,6 +91,14 @@ fn validateInstruction(inst: instruction.Instruction) Error!void { } } +fn validateBufferReference(reference: instruction.BufferReference) Error!void { + switch (reference) { + .logical => {}, + .binding_table => |index| if (index >= compute.resource_layout.max_storage_buffers) + return Error.InvalidBindingTableIndex, + } +} + fn validateSource(source: operand.Source) Error!void { try validateType(source.type); switch (source.register) { diff --git a/src/intel/compiler/targets/targets.zig b/src/intel/compiler/targets/targets.zig index bf35fce..a3e67d0 100644 --- a/src/intel/compiler/targets/targets.zig +++ b/src/intel/compiler/targets/targets.zig @@ -6,6 +6,10 @@ const common_ir = @import("../lower/common_ir.zig"); pub const gen9 = @import("gen9/gen9.zig"); +pub const ComputeResourceLayout = gen9.compute.ResourceLayout; +pub const ResourceLayoutError = gen9.compute.resource_layout.Error || error{UnsupportedGeneration}; +pub const ResourceLoweringError = gen9.ResourceLoweringError || error{UnsupportedGeneration}; + pub const Error = gen9.Error || error{UnsupportedGeneration}; pub const ValidationError = gen9.validator.Error || error{UnsupportedGeneration}; @@ -21,6 +25,20 @@ pub fn lower( }; } +pub fn layoutComputeResources(allocator: std.mem.Allocator, program: *const program_ir.Program) ResourceLayoutError!ComputeResourceLayout { + return switch (program.device_info.generation) { + .gen9 => ComputeResourceLayout.init(allocator, program), + .gen10, .gen11 => ResourceLayoutError.UnsupportedGeneration, + }; +} + +pub fn lowerComputeResources(program: *program_ir.Program, layout: *const ComputeResourceLayout) ResourceLoweringError!void { + return switch (program.device_info.generation) { + .gen9 => gen9.lowerComputeResources(program, layout), + .gen10, .gen11 => ResourceLoweringError.UnsupportedGeneration, + }; +} + pub fn validate(program: *const program_ir.Program) ValidationError!void { return switch (program.device_info.generation) { .gen9 => gen9.validator.validate(program),