From 0788470ee53c308632cb31ce7ddb0d2fcfd6e85c Mon Sep 17 00:00:00 2001 From: Kbz-8 Date: Fri, 28 Aug 2026 19:35:16 +0200 Subject: [PATCH] [Flint] normalize surface addresses and pack Gen9 message payloads --- src/intel/FlintPipeline.zig | 4 +- src/intel/compiler/ir/instruction.zig | 14 ++ src/intel/compiler/ir/printer.zig | 11 + src/intel/compiler/ir/program.zig | 4 +- src/intel/compiler/ir/validator.zig | 26 +++ src/intel/compiler/lower/block_arguments.zig | 6 +- .../compiler/targets/gen9/compute/compute.zig | 2 + .../gen9/compute/message_addresses.zig | 208 ++++++++++++++++++ .../targets/gen9/compute/message_payloads.zig | 200 +++++++++++++++++ .../targets/gen9/compute/pipeline.zig | 6 +- src/intel/compiler/targets/gen9/gen9.zig | 34 ++- .../targets/gen9/register_allocation.zig | 10 + src/intel/compiler/targets/gen9/validator.zig | 1 + src/software/SoftCommandBuffer.zig | 5 +- 14 files changed, 503 insertions(+), 28 deletions(-) create mode 100644 src/intel/compiler/targets/gen9/compute/message_addresses.zig create mode 100644 src/intel/compiler/targets/gen9/compute/message_payloads.zig diff --git a/src/intel/FlintPipeline.zig b/src/intel/FlintPipeline.zig index 989dd78..1c2edd2 100644 --- a/src/intel/FlintPipeline.zig +++ b/src/intel/FlintPipeline.zig @@ -283,6 +283,8 @@ test "Flint pipeline: lower common compute IR" { try std.testing.expect(!program.properties.system_values_lowered); try std.testing.expect(program.properties.resources_lowered); try std.testing.expect(program.properties.messages_lowered); + try std.testing.expect(program.properties.message_addresses_lowered); + try std.testing.expect(program.properties.message_payloads_lowered); try std.testing.expect(program.properties.registers_allocated); try std.testing.expect(!program.properties.instructions_selected); try std.testing.expectEqual([3]u32{ 1, 1, 1 }, program.workgroup_size); @@ -294,5 +296,5 @@ test "Flint pipeline: lower common compute IR" { 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 r0:u32, component(0)") != null); - try std.testing.expect(std.mem.indexOf(u8, text, "surface_write bti(0), 0:u32, r0:u32") != null); + try std.testing.expect(std.mem.indexOf(u8, text, "surface_message write bti(0)") != null); } diff --git a/src/intel/compiler/ir/instruction.zig b/src/intel/compiler/ir/instruction.zig index dd80d99..c8d4e9f 100644 --- a/src/intel/compiler/ir/instruction.zig +++ b/src/intel/compiler/ir/instruction.zig @@ -42,6 +42,19 @@ pub const SurfaceWrite = struct { data: operand.Source, }; +pub const SurfaceMessageKind = enum { + read, + write, +}; + +pub const SurfaceMessage = struct { + kind: SurfaceMessageKind, + binding_table: u8, + payload: operand.RegisterSpan, + response: ?operand.RegisterSpan, + data_type: operand.DataType, +}; + pub const Move = struct { destination: operand.Destination, source: operand.Source, @@ -86,6 +99,7 @@ pub const Operation = union(enum) { store_buffer: StoreBuffer, surface_read: SurfaceRead, surface_write: SurfaceWrite, + surface_message: SurfaceMessage, move: Move, binary: Binary, compare: Compare, diff --git a/src/intel/compiler/ir/printer.zig b/src/intel/compiler/ir/printer.zig index 0b22c64..e33bb02 100644 --- a/src/intel/compiler/ir/printer.zig +++ b/src/intel/compiler/ir/printer.zig @@ -152,6 +152,17 @@ fn writeOperation(program: *const program_ir.Program, writer: *std.Io.Writer, ex try writer.writeAll(", "); try writeSource(program, writer, execution_size, op.data); }, + .surface_message => |op| { + try writer.print("surface_message {t} bti({d}), payload(", .{ op.kind, op.binding_table }); + try writeRegister(program, writer, op.payload.base); + try writer.print(", {d})", .{op.payload.register_count}); + if (op.response) |response| { + try writer.writeAll(", response("); + try writeRegister(program, writer, response.base); + try writer.print(", {d})", .{response.register_count}); + } + try writer.print(", type({t})", .{op.data_type}); + }, .move => |op| { try writer.writeAll("mov "); try writeDestination(program, writer, execution_size, op.destination); diff --git a/src/intel/compiler/ir/program.zig b/src/intel/compiler/ir/program.zig index 834ed75..1be541f 100644 --- a/src/intel/compiler/ir/program.zig +++ b/src/intel/compiler/ir/program.zig @@ -14,6 +14,8 @@ pub const Properties = packed struct { system_values_lowered: bool = false, resources_lowered: bool = false, messages_lowered: bool = false, + message_addresses_lowered: bool = false, + message_payloads_lowered: bool = false, control_flow_lowered: bool = false, regions_legalized: bool = false, @@ -23,7 +25,7 @@ pub const Properties = packed struct { flags_allocated: bool = false, branches_resolved: bool = false, - _padding: u19 = 0, + _padding: u17 = 0, }; pub const StorageBuffer = struct { diff --git a/src/intel/compiler/ir/validator.zig b/src/intel/compiler/ir/validator.zig index 012f263..fa4d992 100644 --- a/src/intel/compiler/ir/validator.zig +++ b/src/intel/compiler/ir/validator.zig @@ -33,6 +33,7 @@ pub const Error = error{ UnloweredSystemValue, UnloweredResource, UnloweredMessage, + InvalidMessage, InvalidPayloadLayout, EntryBlockHasParameters, DuplicateBlockParameter, @@ -168,6 +169,22 @@ fn validateInstruction(program: *const program_ir.Program, inst: instruction.Ins if (!op.data.type.isInitialTargetType()) return Error.InvalidBufferAccess; }, + .surface_message => |op| { + try validateRegisterSpan(program, op.payload); + if (!op.data_type.isInitialTargetType()) + return Error.InvalidMessage; + switch (op.kind) { + .read => { + if (op.payload.register_count != 1 or op.response == null) + return Error.InvalidMessage; + try validateRegisterSpan(program, op.response.?); + if (op.response.?.register_count != 1) + return Error.InvalidMessage; + }, + .write => if (op.payload.register_count != 2 or op.response != null) + return Error.InvalidMessage, + } + }, .move => |op| { try validateDestination(program, op.destination); try validateSource(program, op.source); @@ -303,6 +320,15 @@ fn validateDestination(program: *const program_ir.Program, destination: operand. } } +fn validateRegisterSpan(program: *const program_ir.Program, span: operand.RegisterSpan) Error!void { + if (span.register_count == 0) + return Error.InvalidMessage; + switch (span.base) { + .virtual, .physical_grf => try validateRegisterRef(program, span.base), + else => return Error.InvalidMessage, + } +} + fn validateRegisterRef(program: *const program_ir.Program, register: operand.RegisterRef) Error!void { switch (register) { .virtual => |id| if (!program.virtual_registers.isLive(id)) diff --git a/src/intel/compiler/lower/block_arguments.zig b/src/intel/compiler/lower/block_arguments.zig index 935a7ab..3caaf61 100644 --- a/src/intel/compiler/lower/block_arguments.zig +++ b/src/intel/compiler/lower/block_arguments.zig @@ -6,6 +6,7 @@ const operand = @import("../ir/operand.zig"); const program_ir = @import("../ir/program.zig"); const pseudo = @import("../ir/pseudo.zig"); const validator = @import("../ir/validator.zig"); +const device = @import("../device.zig"); pub const Error = std.mem.Allocator.Error || error{ InvalidProgram, @@ -128,7 +129,7 @@ fn rewriteEdge( return .{ .target = edge_block, .arguments = &.{} }; } -fn executionSize(dispatch_width: @import("../device.zig").DispatchWidth) @import("../device.zig").ExecutionSize { +fn executionSize(dispatch_width: device.DispatchWidth) device.ExecutionSize { return @enumFromInt(@intFromEnum(dispatch_width)); } @@ -140,7 +141,6 @@ fn mapBuilderError(err: anyerror) Error { } test "[ir] block arguments: lower register and flag parameters" { - const device = @import("../device.zig"); const printer = @import("../ir/printer.zig"); const device_info: device.DeviceInfo = .{ @@ -215,8 +215,6 @@ test "[ir] block arguments: lower register and flag parameters" { } test "[ir] block arguments: split same-target conditional edges" { - const device = @import("../device.zig"); - const device_info: device.DeviceInfo = .{ .generation = .gen9, .platform = .skylake, diff --git a/src/intel/compiler/targets/gen9/compute/compute.zig b/src/intel/compiler/targets/gen9/compute/compute.zig index e629c9f..08e7c73 100644 --- a/src/intel/compiler/targets/gen9/compute/compute.zig +++ b/src/intel/compiler/targets/gen9/compute/compute.zig @@ -1,6 +1,8 @@ const std = @import("std"); +pub const message_addresses = @import("message_addresses.zig"); pub const message_lowering = @import("message_lowering.zig"); +pub const message_payloads = @import("message_payloads.zig"); pub const resource_layout = @import("resource_layout.zig"); pub const resource_lowering = @import("resource_lowering.zig"); pub const ResourceLayout = resource_layout.Layout; diff --git a/src/intel/compiler/targets/gen9/compute/message_addresses.zig b/src/intel/compiler/targets/gen9/compute/message_addresses.zig new file mode 100644 index 0000000..93fb0eb --- /dev/null +++ b/src/intel/compiler/targets/gen9/compute/message_addresses.zig @@ -0,0 +1,208 @@ +const std = @import("std"); + +const Builder = @import("../../../ir/Builder.zig"); +const ids = @import("../../../ir/id.zig"); +const operand = @import("../../../ir/operand.zig"); +const program_ir = @import("../../../ir/program.zig"); +const instruction = @import("../../../ir/instruction.zig"); + +pub const Error = std.mem.Allocator.Error || error{ + MessagesNotLowered, + InvalidProgram, +}; + +const AddressAdjustment = struct { + address: operand.Source, + immediate_offset: u32, +}; + +pub fn run(program: *program_ir.Program) Error!void { + if (!program.properties.messages_lowered) + return error.MessagesNotLowered; + if (program.properties.message_addresses_lowered) + return; + + 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 adjustment = addressAdjustment(inst.operation) orelse { + instruction_index += 1; + continue; + }; + if (adjustment.address.type != .u32) + return error.InvalidProgram; + + if (adjustment.immediate_offset == 0) { + instruction_index += 1; + continue; + } + + switch (adjustment.address.register) { + .immediate => |immediate| { + const base = switch (immediate) { + .u32 => |value| value, + else => return error.InvalidProgram, + }; + const mutable = program.instructions.getMut(instruction_id) orelse return error.InvalidProgram; + const address = messageAddressMut(&mutable.operation) orelse return error.InvalidProgram; + address.source.register = .{ .immediate = .{ .u32 = base +% adjustment.immediate_offset } }; + address.immediate_offset.* = 0; + instruction_index += 1; + }, + .virtual, .physical_grf, .architecture => { + const execution_width: u32 = @intFromEnum(inst.execution_size); + const size_bytes = execution_width * @sizeOf(u32); + const address_register = builder.addVirtualRegister(.{ + .size_bytes = size_bytes, + .alignment_bytes = @intCast(@min(size_bytes, program.device_info.grf_size_bytes)), + .element_type = .u32, + .lane_count = @intCast(execution_width), + .class = .temporary, + }) catch |err| return mapBuilderError(err); + + _ = builder.insertInstruction(block_id, instruction_index, inst.execution_size, inst.predicate, .{ + .binary = .{ + .opcode = .add, + .destination = .{ + .register = .{ .virtual = address_register }, + .type = .u32, + }, + .lhs = adjustment.address, + .rhs = immediateSource(adjustment.immediate_offset), + }, + }) catch |err| return mapBuilderError(err); + + const mutable = program.instructions.getMut(instruction_id) orelse return error.InvalidProgram; + const address = messageAddressMut(&mutable.operation) orelse return error.InvalidProgram; + address.source.* = .{ + .register = .{ .virtual = address_register }, + .type = .u32, + .region = operand.Region.contiguous(inst.execution_size), + }; + address.immediate_offset.* = 0; + instruction_index += 2; + }, + .null => return error.InvalidProgram, + } + } + } + + program.properties.message_addresses_lowered = true; +} + +fn addressAdjustment(operation: instruction.Operation) ?AddressAdjustment { + return switch (operation) { + .surface_read => |op| .{ .address = op.address, .immediate_offset = op.immediate_offset }, + .surface_write => |op| .{ .address = op.address, .immediate_offset = op.immediate_offset }, + else => null, + }; +} + +const MutableAddress = struct { + source: *operand.Source, + immediate_offset: *u32, +}; + +fn messageAddressMut(operation: *instruction.Operation) ?MutableAddress { + return switch (operation.*) { + .surface_read => |*op| .{ .source = &op.address, .immediate_offset = &op.immediate_offset }, + .surface_write => |*op| .{ .source = &op.address, .immediate_offset = &op.immediate_offset }, + else => null, + }; +} + +fn immediateSource(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, + }; +} + +const device = @import("../../../device.zig"); + +const test_device: device.DeviceInfo = .{ + .generation = .gen9, + .platform = .skylake, + .pci_device_id = 0x1912, + .grf_count = 128, +}; + +fn markPrerequisite(program: *program_ir.Program) void { + program.properties.messages_lowered = true; +} + +test "[gen9] message addresses: fold immediate offsets" { + var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, test_device, .simd8); + defer program.deinit(); + + const entry = try program.addBlock("entry"); + const message = try program.appendInstruction(entry, .simd8, null, .{ .surface_write = .{ + .binding_table = 0, + .address = immediateSource(12), + .immediate_offset = 4, + .data = immediateSource(7), + } }); + try program.setTerminator(entry, .end_thread); + markPrerequisite(&program); + + try run(&program); + + const write = program.instructions.get(message).?.operation.surface_write; + try std.testing.expectEqual(@as(u32, 16), write.address.register.immediate.u32); + try std.testing.expectEqual(@as(u32, 0), write.immediate_offset); + try std.testing.expect(program.properties.message_addresses_lowered); +} + +test "[gen9] message addresses: materialize dynamic offsets" { + var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, test_device, .simd8); + defer program.deinit(); + + const base = try program.addVirtualRegister(.{ + .size_bytes = 32, + .alignment_bytes = 32, + .element_type = .u32, + .lane_count = 8, + .class = .temporary, + }); + const entry = try program.addBlock("entry"); + const message = try program.appendInstruction(entry, .simd8, null, .{ .surface_read = .{ + .destination = .{ .register = .{ .virtual = base }, .type = .u32 }, + .binding_table = 0, + .address = .{ + .register = .{ .virtual = base }, + .type = .u32, + .region = operand.Region.contiguous(.simd8), + }, + .immediate_offset = 8, + } }); + try program.setTerminator(entry, .end_thread); + markPrerequisite(&program); + + try run(&program); + + const block = program.blocks.get(entry).?; + try std.testing.expectEqual(@as(usize, 2), block.instructions.items.len); + try std.testing.expect(program.instructions.get(block.instructions.items[0]).?.operation == .binary); + const read = program.instructions.get(message).?.operation.surface_read; + try std.testing.expect(read.address.register == .virtual); + try std.testing.expect(read.address.register.virtual != base); + try std.testing.expectEqual(@as(u32, 0), read.immediate_offset); +} diff --git a/src/intel/compiler/targets/gen9/compute/message_payloads.zig b/src/intel/compiler/targets/gen9/compute/message_payloads.zig new file mode 100644 index 0000000..3904f4b --- /dev/null +++ b/src/intel/compiler/targets/gen9/compute/message_payloads.zig @@ -0,0 +1,200 @@ +const std = @import("std"); + +const Builder = @import("../../../ir/Builder.zig"); +const device = @import("../../../device.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"); + +pub const Error = std.mem.Allocator.Error || error{ + MessageAddressesNotLowered, + InvalidProgram, +}; + +pub fn run(program: *program_ir.Program) Error!void { + if (!program.properties.message_addresses_lowered) + return error.MessageAddressesNotLowered; + if (program.properties.message_payloads_lowered) + return; + if (program.device_info.grf_size_bytes != 32) + 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 execution_size = inst.execution_size; + switch (inst.operation) { + .surface_read => |op| { + if (op.immediate_offset != 0 or op.address.type != .u32) + return error.InvalidProgram; + const response = try responseSpan(op.destination); + const payload = try addPayloadRegister(&builder, execution_size, 1); + _ = builder.insertInstruction(block_id, instruction_index, execution_size, null, .{ .move = .{ + .destination = payloadDestination(payload, 0, .u32), + .source = op.address, + } }) catch |err| return mapBuilderError(err); + + const mutable = program.instructions.getMut(instruction_id) orelse return error.InvalidProgram; + mutable.operation = .{ .surface_message = .{ + .kind = .read, + .binding_table = op.binding_table, + .payload = .{ .base = .{ .virtual = payload }, .register_count = 1 }, + .response = response, + .data_type = op.destination.type, + } }; + instruction_index += 2; + }, + .surface_write => |op| { + if (op.immediate_offset != 0 or op.address.type != .u32) + return error.InvalidProgram; + const payload = try addPayloadRegister(&builder, execution_size, 2); + _ = builder.insertInstruction(block_id, instruction_index, execution_size, null, .{ .move = .{ + .destination = payloadDestination(payload, 0, .u32), + .source = op.address, + } }) catch |err| return mapBuilderError(err); + _ = builder.insertInstruction(block_id, instruction_index + 1, execution_size, null, .{ .move = .{ + .destination = payloadDestination(payload, 32, op.data.type), + .source = op.data, + } }) catch |err| return mapBuilderError(err); + + const mutable = program.instructions.getMut(instruction_id) orelse return error.InvalidProgram; + mutable.operation = .{ .surface_message = .{ + .kind = .write, + .binding_table = op.binding_table, + .payload = .{ .base = .{ .virtual = payload }, .register_count = 2 }, + .response = null, + .data_type = op.data.type, + } }; + instruction_index += 3; + }, + else => instruction_index += 1, + } + } + } + + program.properties.message_payloads_lowered = true; +} + +fn addPayloadRegister(builder: *Builder, execution_size: device.ExecutionSize, register_count: u8) Error!ids.VirtualRegisterId { + return builder.addVirtualRegister(.{ + .size_bytes = @as(u32, register_count) * 32, + .alignment_bytes = 32, + .element_type = .u32, + .lane_count = @intFromEnum(execution_size), + .class = .temporary, + .spillable = false, + }) catch |err| return mapBuilderError(err); +} + +fn payloadDestination(register: ids.VirtualRegisterId, byte_offset: u16, data_type: operand.DataType) operand.Destination { + return .{ + .register = .{ .virtual = register }, + .type = data_type, + .region = .{ .byte_offset = byte_offset }, + }; +} + +fn responseSpan(destination: operand.Destination) Error!operand.RegisterSpan { + if (destination.region.byte_offset != 0 or destination.region.horizontal_stride != 1) + return error.InvalidProgram; + return switch (destination.register) { + .virtual, .physical_grf => .{ + .base = destination.register, + .register_count = 1, + }, + else => error.InvalidProgram, + }; +} + +fn mapBuilderError(err: Builder.Error) Error { + return switch (err) { + error.OutOfMemory => error.OutOfMemory, + else => error.InvalidProgram, + }; +} + +const test_device: device.DeviceInfo = .{ + .generation = .gen9, + .platform = .skylake, + .pci_device_id = 0x1912, + .grf_count = 128, +}; + +fn immediate(value: u32) operand.Source { + return .{ + .register = .{ .immediate = .{ .u32 = value } }, + .type = .u32, + .region = operand.Region.broadcast(), + }; +} + +test "[gen9] message payloads: pack surface write address and data" { + var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, test_device, .simd8); + defer program.deinit(); + + const entry = try program.addBlock("entry"); + const message = try program.appendInstruction(entry, .simd8, null, .{ .surface_write = .{ + .binding_table = 2, + .address = immediate(16), + .data = immediate(42), + } }); + try program.setTerminator(entry, .end_thread); + program.properties.message_addresses_lowered = true; + + try run(&program); + + const block = program.blocks.get(entry).?; + try std.testing.expectEqual(@as(usize, 3), block.instructions.items.len); + const address_move = program.instructions.get(block.instructions.items[0]).?.operation.move; + const data_move = program.instructions.get(block.instructions.items[1]).?.operation.move; + try std.testing.expectEqual(@as(u16, 0), address_move.destination.region.byte_offset); + try std.testing.expectEqual(@as(u16, 32), data_move.destination.region.byte_offset); + try std.testing.expectEqual(address_move.destination.register.virtual, data_move.destination.register.virtual); + + const send = program.instructions.get(message).?.operation.surface_message; + try std.testing.expectEqual(instruction.SurfaceMessageKind.write, send.kind); + try std.testing.expectEqual(@as(u8, 2), send.binding_table); + try std.testing.expectEqual(@as(u8, 2), send.payload.register_count); + try std.testing.expect(send.response == null); + try std.testing.expect(program.properties.message_payloads_lowered); +} + +test "[gen9] message payloads: prepare surface read response" { + var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, test_device, .simd8); + defer program.deinit(); + + const result = try program.addVirtualRegister(.{ + .size_bytes = 32, + .alignment_bytes = 32, + .element_type = .u32, + .lane_count = 8, + .class = .response, + }); + const entry = try program.addBlock("entry"); + const message = try program.appendInstruction(entry, .simd8, null, .{ .surface_read = .{ + .destination = .{ .register = .{ .virtual = result }, .type = .u32 }, + .binding_table = 1, + .address = immediate(0), + } }); + try program.setTerminator(entry, .end_thread); + program.properties.message_addresses_lowered = true; + + try run(&program); + + const send = program.instructions.get(message).?.operation.surface_message; + try std.testing.expectEqual(instruction.SurfaceMessageKind.read, send.kind); + try std.testing.expectEqual(@as(u8, 1), send.payload.register_count); + try std.testing.expectEqual(result, send.response.?.base.virtual); + try std.testing.expectEqual(@as(u8, 1), send.response.?.register_count); +} diff --git a/src/intel/compiler/targets/gen9/compute/pipeline.zig b/src/intel/compiler/targets/gen9/compute/pipeline.zig index fde9220..146b160 100644 --- a/src/intel/compiler/targets/gen9/compute/pipeline.zig +++ b/src/intel/compiler/targets/gen9/compute/pipeline.zig @@ -10,12 +10,14 @@ const flag_allocation = @import("../flag_allocation.zig"); const register_allocation = @import("../register_allocation.zig"); const compute = @import("compute.zig"); +const message_addresses = @import("message_addresses.zig"); const message_lowering = @import("message_lowering.zig"); +const message_payloads = @import("message_payloads.zig"); const resource_layout = @import("resource_layout.zig"); const resource_lowering = @import("resource_lowering.zig"); pub const Error = common_ir.Error || block_arguments.Error || parallel_copies.Error || - message_lowering.Error || resource_layout.Error || resource_lowering.Error || flag_allocation.Error || register_allocation.Error || compute.Error || error{ + message_addresses.Error || message_lowering.Error || message_payloads.Error || resource_layout.Error || resource_lowering.Error || flag_allocation.Error || register_allocation.Error || compute.Error || error{ UnsupportedGeneration, UnsupportedStage, UnsupportedDispatchWidth, @@ -64,6 +66,8 @@ pub fn compile(allocator: std.mem.Allocator, module: *shader_ir.module.Module, d try resource_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); diff --git a/src/intel/compiler/targets/gen9/gen9.zig b/src/intel/compiler/targets/gen9/gen9.zig index 8b2af61..bd99325 100644 --- a/src/intel/compiler/targets/gen9/gen9.zig +++ b/src/intel/compiler/targets/gen9/gen9.zig @@ -106,6 +106,8 @@ test "[gen9] target: lower 256 KiB SSBO copy loop" { try std.testing.expect(program.properties.resources_lowered); try std.testing.expect(program.properties.messages_lowered); + try std.testing.expect(program.properties.message_addresses_lowered); + try std.testing.expect(program.properties.message_payloads_lowered); try std.testing.expectEqual(@as(usize, 2), resources.bindings.len); try std.testing.expectEqual(compute.resource_layout.Binding{ .set = 0, @@ -118,28 +120,24 @@ test "[gen9] target: lower 256 KiB SSBO copy loop" { .binding_table_index = 1, }, resources.bindings[1]); - var load_offsets: [4]bool = @splat(false); - var store_offsets: [4]bool = @splat(false); var load_count: usize = 0; var store_count: usize = 0; for (program.instructions.entries.items) |instruction_entry| { const inst = instruction_entry orelse continue; switch (inst.operation) { - .surface_read => |operation| { - try std.testing.expectEqual(@as(u8, 0), operation.binding_table); - try std.testing.expect(operation.immediate_offset % @sizeOf(u32) == 0); - const component = operation.immediate_offset / @sizeOf(u32); - try std.testing.expect(component < load_offsets.len); - load_offsets[component] = true; - load_count += 1; - }, - .surface_write => |operation| { - try std.testing.expectEqual(@as(u8, 1), operation.binding_table); - try std.testing.expect(operation.immediate_offset % @sizeOf(u32) == 0); - const component = operation.immediate_offset / @sizeOf(u32); - try std.testing.expect(component < store_offsets.len); - store_offsets[component] = true; - store_count += 1; + .surface_message => |operation| switch (operation.kind) { + .read => { + try std.testing.expectEqual(@as(u8, 0), operation.binding_table); + try std.testing.expectEqual(@as(u8, 1), operation.payload.register_count); + try std.testing.expect(operation.response != null); + load_count += 1; + }, + .write => { + try std.testing.expectEqual(@as(u8, 1), operation.binding_table); + try std.testing.expectEqual(@as(u8, 2), operation.payload.register_count); + try std.testing.expect(operation.response == null); + store_count += 1; + }, }, .parallel_copy => return error.UnloweredParallelCopy, else => {}, @@ -148,6 +146,4 @@ 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); - try std.testing.expectEqual([4]bool{ true, true, true, true }, load_offsets); - try std.testing.expectEqual([4]bool{ true, true, true, true }, store_offsets); } diff --git a/src/intel/compiler/targets/gen9/register_allocation.zig b/src/intel/compiler/targets/gen9/register_allocation.zig index d6b2097..797997a 100644 --- a/src/intel/compiler/targets/gen9/register_allocation.zig +++ b/src/intel/compiler/targets/gen9/register_allocation.zig @@ -76,6 +76,11 @@ fn reserveExistingPhysicalRegisters(program: *const program_ir.Program, initial: reserveRegister(&next_byte, op.address.register, grf_size); reserveRegister(&next_byte, op.data.register, grf_size); }, + .surface_message => |op| { + reserveRegister(&next_byte, op.payload.base, grf_size); + if (op.response) |response| + reserveRegister(&next_byte, response.base, grf_size); + }, .move => |op| { reserveRegister(&next_byte, op.destination.register, grf_size); reserveRegister(&next_byte, op.source.register, grf_size); @@ -128,6 +133,11 @@ fn rewriteProgram(program: *program_ir.Program, allocations: []const ?operand.Ph try rewriteSource(program, &op.address, allocations); try rewriteSource(program, &op.data, allocations); }, + .surface_message => |*op| { + try rewriteRegister(program, &op.payload.base, allocations); + if (op.response) |*response| + try rewriteRegister(program, &response.base, allocations); + }, .move => |*op| { try rewriteDestination(program, &op.destination, allocations); try rewriteSource(program, &op.source, allocations); diff --git a/src/intel/compiler/targets/gen9/validator.zig b/src/intel/compiler/targets/gen9/validator.zig index f15c0e6..2f49708 100644 --- a/src/intel/compiler/targets/gen9/validator.zig +++ b/src/intel/compiler/targets/gen9/validator.zig @@ -74,6 +74,7 @@ fn validateInstruction(inst: instruction.Instruction) Error!void { try validateSource(op.address); try validateSource(op.data); }, + .surface_message => |op| try validateBindingTableIndex(op.binding_table), .move => |op| { try validateDestination(op.destination); try validateSource(op.source); diff --git a/src/software/SoftCommandBuffer.zig b/src/software/SoftCommandBuffer.zig index 7dd95f3..1774c90 100644 --- a/src/software/SoftCommandBuffer.zig +++ b/src/software/SoftCommandBuffer.zig @@ -13,6 +13,7 @@ const SoftPipeline = @import("SoftPipeline.zig"); const SoftRenderPass = @import("SoftRenderPass.zig"); const ExecutionDevice = @import("device/Device.zig"); +const Renderer = @import("device/Renderer.zig"); const blitter = @import("device/blitter.zig"); const Self = @This(); @@ -1375,7 +1376,7 @@ pub fn setDepthBias(interface: *Interface, constant_factor: f32, clamp: f32, slo const CommandImpl = struct { const Impl = @This(); - depth_bias: @import("device/Renderer.zig").DepthBias, + depth_bias: Renderer.DepthBias, pub fn execute(context: *anyopaque, device: *ExecutionDevice) VkError!void { const impl: *Impl = @ptrCast(@alignCast(context)); @@ -1402,7 +1403,7 @@ pub fn setDepthBounds(interface: *Interface, min: f32, max: f32) VkError!void { const CommandImpl = struct { const Impl = @This(); - depth_bounds: @import("device/Renderer.zig").DepthBounds, + depth_bounds: Renderer.DepthBounds, pub fn execute(context: *anyopaque, device: *ExecutionDevice) VkError!void { const impl: *Impl = @ptrCast(@alignCast(context));