From 2343cf4a586c33611818dcf05847172b630b66ef Mon Sep 17 00:00:00 2001 From: Kbz-8 Date: Mon, 31 Aug 2026 11:27:21 +0200 Subject: [PATCH] [Flint] adding math and compare eu encoding --- README.md | 56 ++--- src/intel/FlintPipeline.zig | 5 +- .../compiler/targets/gen9/compute/compute.zig | 2 + .../targets/gen9/compute/dispatch.zig | 25 ++- .../compiler/targets/gen9/compute/eu.zig | 96 ++++++++ .../targets/gen9/compute/eu_encoder.zig | 178 ++++++++++++++- .../targets/gen9/compute/kernel_encoder.zig | 211 ++++++++++++++---- .../targets/gen9/compute/pipeline.zig | 7 + .../compiler/targets/gen9/compute/regions.zig | 49 ++++ .../targets/gen9/compute/system_values.zig | 96 ++++++++ src/intel/compiler/targets/gen9/gen9.zig | 1 + 11 files changed, 642 insertions(+), 84 deletions(-) create mode 100644 src/intel/compiler/targets/gen9/compute/eu.zig create mode 100644 src/intel/compiler/targets/gen9/compute/regions.zig create mode 100644 src/intel/compiler/targets/gen9/compute/system_values.zig diff --git a/README.md b/README.md index aa006d2..e09902d 100644 --- a/README.md +++ b/README.md @@ -331,13 +331,13 @@ Assume thou that functions lacking in this array are, for now, not intended to b | vkCreateQueryPool | ⚙️ WIP | | vkCreateRenderPass | ⚙️ WIP | | vkCreateSampler | ⚙️ WIP | -| vkCreateSemaphore | ✅ Implemented | -| vkCreateShaderModule | ⚙️ WIP | -| vkCreateSwapchainKHR | ⚙️ WIP | -| vkCreateWaylandSurfaceKHR | ⚙️ WIP | +| vkCreateSemaphore | ⚙️ WIP | +| vkCreateShaderModule | ✅ Implemented | +| vkCreateSwapchainKHR | ✅ Implemented | +| vkCreateWaylandSurfaceKHR | ✅ Implemented | | vkCreateWin32SurfaceKHR | ⚙️ WIP | -| vkCreateXcbSurfaceKHR | ⚙️ WIP | -| vkCreateXlibSurfaceKHR | ⚙️ WIP | +| vkCreateXcbSurfaceKHR | ✅ Implemented | +| vkCreateXlibSurfaceKHR | ✅ Implemented | | vkDestroyBuffer | ✅ Implemented | | vkDestroyBufferView | ⚙️ WIP | | vkDestroyCommandPool | ✅ Implemented | @@ -359,7 +359,7 @@ Assume thou that functions lacking in this array are, for now, not intended to b | vkDestroySemaphore | ✅ Implemented | | vkDestroyShaderModule | ⚙️ WIP | | vkDestroySurfaceKHR | ⚙️ WIP | -| vkDestroySwapchainKHR | ⚙️ WIP | +| vkDestroySwapchainKHR | ✅ Implemented | | vkDeviceWaitIdle | ✅ Implemented | | vkEndCommandBuffer | ✅ Implemented | | vkEnumerateDeviceExtensionProperties | ⚙️ WIP | @@ -473,9 +473,9 @@ Assume thou that functions lacking in this array are, for now, not intended to b | vkBindImageMemory | ✅ Implemented | | vkCmdBeginQuery | ⚙️ WIP | | vkCmdBeginRenderPass | ⚙️ WIP | -| vkCmdBindDescriptorSets | ⚙️ WIP | +| vkCmdBindDescriptorSets | ✅ Implemented | | vkCmdBindIndexBuffer | ⚙️ WIP | -| vkCmdBindPipeline | ⚙️ WIP | +| vkCmdBindPipeline | ✅ Implemented | | vkCmdBindVertexBuffers | ⚙️ WIP | | vkCmdBlitImage | ⚙️ WIP | | vkCmdClearAttachments | ⚙️ WIP | @@ -486,8 +486,8 @@ Assume thou that functions lacking in this array are, for now, not intended to b | vkCmdCopyImage | ✅ Implemented | | vkCmdCopyImageToBuffer | ✅ Implemented | | vkCmdCopyQueryPoolResults | ⚙️ WIP | -| vkCmdDispatch | ⚙️ WIP | -| vkCmdDispatchBaseKHR | ⚙️ WIP | +| vkCmdDispatch | ✅ Implemented | +| vkCmdDispatchBaseKHR | ✅ Implemented | | vkCmdDispatchIndirect | ⚙️ WIP | | vkCmdDraw | ⚙️ WIP | | vkCmdDrawIndexed | ⚙️ WIP | @@ -520,9 +520,9 @@ Assume thou that functions lacking in this array are, for now, not intended to b | vkCreateBuffer | ✅ Implemented | | vkCreateBufferView | ⚙️ WIP | | vkCreateCommandPool | ✅ Implemented | -| vkCreateComputePipelines | ⚙️ WIP | -| vkCreateDescriptorPool | ⚙️ WIP | -| vkCreateDescriptorSetLayout | ⚙️ WIP | +| vkCreateComputePipelines | ✅ Implemented | +| vkCreateDescriptorPool | ✅ Implemented | +| vkCreateDescriptorSetLayout | ✅ Implemented | | vkCreateDevice | ✅ Implemented | | vkCreateEvent | ⚙️ WIP | | vkCreateFence | ✅ Implemented | @@ -537,17 +537,17 @@ Assume thou that functions lacking in this array are, for now, not intended to b | vkCreateRenderPass | ⚙️ WIP | | vkCreateSampler | ⚙️ WIP | | vkCreateSemaphore | ✅ Implemented | -| vkCreateShaderModule | ⚙️ WIP | -| vkCreateSwapchainKHR | ⚙️ WIP | -| vkCreateWaylandSurfaceKHR | ⚙️ WIP | +| vkCreateShaderModule | ✅ Implemented | +| vkCreateSwapchainKHR | ✅ Implemented | +| vkCreateWaylandSurfaceKHR | ✅ Implemented | | vkCreateWin32SurfaceKHR | ⚙️ WIP | -| vkCreateXcbSurfaceKHR | ⚙️ WIP | -| vkCreateXlibSurfaceKHR | ⚙️ WIP | +| vkCreateXcbSurfaceKHR | ✅ Implemented | +| vkCreateXlibSurfaceKHR | ✅ Implemented | | vkDestroyBuffer | ✅ Implemented | | vkDestroyBufferView | ⚙️ WIP | | vkDestroyCommandPool | ✅ Implemented | -| vkDestroyDescriptorPool | ⚙️ WIP | -| vkDestroyDescriptorSetLayout | ⚙️ WIP | +| vkDestroyDescriptorPool | ✅ Implemented | +| vkDestroyDescriptorSetLayout | ✅ Implemented | | vkDestroyDevice | ✅ Implemented | | vkDestroyEvent | ⚙️ WIP | | vkDestroyFence | ✅ Implemented | @@ -555,16 +555,16 @@ Assume thou that functions lacking in this array are, for now, not intended to b | vkDestroyImage | ✅ Implemented | | vkDestroyImageView | ⚙️ WIP | | vkDestroyInstance | ✅ Implemented | -| vkDestroyPipeline | ⚙️ WIP | +| vkDestroyPipeline | ✅ Implemented | | vkDestroyPipelineCache | ⚙️ WIP | | vkDestroyPipelineLayout | ⚙️ WIP | | vkDestroyQueryPool | ⚙️ WIP | | vkDestroyRenderPass | ⚙️ WIP | | vkDestroySampler | ⚙️ WIP | | vkDestroySemaphore | ✅ Implemented | -| vkDestroyShaderModule | ⚙️ WIP | -| vkDestroySurfaceKHR | ⚙️ WIP | -| vkDestroySwapchainKHR | ⚙️ WIP | +| vkDestroyShaderModule | ✅ Implemented | +| vkDestroySurfaceKHR | ✅ Implemented | +| vkDestroySwapchainKHR | ✅ Implemented | | vkDeviceWaitIdle | ✅ Implemented | | vkEndCommandBuffer | ✅ Implemented | | vkEnumerateDeviceExtensionProperties | ✅ Implemented | @@ -575,7 +575,7 @@ Assume thou that functions lacking in this array are, for now, not intended to b | vkEnumeratePhysicalDevices | ✅ Implemented | | vkFlushMappedMemoryRanges | ✅ Implemented | | vkFreeCommandBuffers | ✅ Implemented | -| vkFreeDescriptorSets | ⚙️ WIP | +| vkFreeDescriptorSets | ✅ Implemented | | vkFreeMemory | ✅ Implemented | | vkGetBufferDeviceAddress | ⚙️ WIP | | vkGetBufferDeviceAddressEXT | ⚙️ WIP | @@ -592,7 +592,7 @@ Assume thou that functions lacking in this array are, for now, not intended to b | vkGetImageMemoryRequirements | ⚙️ WIP | | vkGetImageSparseMemoryRequirements | ⚙️ WIP | | vkGetImageSubresourceLayout | ⚙️ WIP | -| vkGetInstanceProcAddr | ⚙️ WIP | +| vkGetInstanceProcAddr | ✅ Implemented | | vkGetPhysicalDeviceFeatures | ✅ Implemented | | vkGetPhysicalDeviceFormatProperties | ⚙️ WIP | | vkGetPhysicalDeviceImageFormatProperties | ⚙️ WIP | @@ -621,7 +621,7 @@ Assume thou that functions lacking in this array are, for now, not intended to b | vkQueueWaitIdle | ✅ Implemented | | vkResetCommandBuffer | ✅ Implemented | | vkResetCommandPool | ✅ Implemented | -| vkResetDescriptorPool | ⚙️ WIP | +| vkResetDescriptorPool | ✅ Implemented | | vkResetEvent | ⚙️ WIP | | vkResetFences | ✅ Implemented | | vkResetQueryPool | ⚙️ WIP | diff --git a/src/intel/FlintPipeline.zig b/src/intel/FlintPipeline.zig index b5c3dbb..f90bbf6 100644 --- a/src/intel/FlintPipeline.zig +++ b/src/intel/FlintPipeline.zig @@ -304,7 +304,7 @@ test "Flint pipeline: lower common compute IR" { try std.testing.expectEqual(@as(u16, 1), program.program_data.payload_grf_count); try std.testing.expectEqual(@as(u16, 0), program.payload.header_grf.?.number); try std.testing.expect(program.properties.block_parameters_lowered); - try std.testing.expect(!program.properties.system_values_lowered); + 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); @@ -319,6 +319,7 @@ 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 r1:u32, component(0)") != null); + try std.testing.expect(std.mem.indexOf(u8, text, "load_global_invocation_id") == null); + try std.testing.expect(std.mem.indexOf(u8, text, "mov r1:u32, 0:u32") != null); try std.testing.expect(std.mem.indexOf(u8, text, "surface_message write bti(0)") != null); } diff --git a/src/intel/compiler/targets/gen9/compute/compute.zig b/src/intel/compiler/targets/gen9/compute/compute.zig index 031f6a1..8ebed89 100644 --- a/src/intel/compiler/targets/gen9/compute/compute.zig +++ b/src/intel/compiler/targets/gen9/compute/compute.zig @@ -10,6 +10,8 @@ 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 regions = @import("regions.zig"); +pub const system_values = @import("system_values.zig"); pub const ResourceLayout = resource_layout.Layout; pub const Error = error{UnsupportedWorkgroupSize}; diff --git a/src/intel/compiler/targets/gen9/compute/dispatch.zig b/src/intel/compiler/targets/gen9/compute/dispatch.zig index 012d3f7..9f1b46d 100644 --- a/src/intel/compiler/targets/gen9/compute/dispatch.zig +++ b/src/intel/compiler/targets/gen9/compute/dispatch.zig @@ -92,7 +92,7 @@ pub fn writeState(destination: []u8, kernel: []const u8, buffer_sizes: []const u const idd = layout.interface_descriptor_offset; putU32(destination, idd + 0, layout.kernel_offset); putU32(destination, idd + 4, 0); - putU32(destination, idd + 4 * @sizeOf(u32), @as(u32, @intCast(buffer_sizes.len)) | layout.binding_table_offset); + putU32(destination, idd + 4 * @sizeOf(u32), @as(u32, layout.surface_count) | layout.binding_table_offset); putU32(destination, idd + 6 * @sizeOf(u32), 1); return layout; @@ -184,3 +184,26 @@ fn alignForward(value: usize, alignment: usize) usize { fn putU32(destination: []u8, offset: u32, value: u32) void { std.mem.writeInt(u32, destination[offset..][0..@sizeOf(u32)], value, .little); } + +test "[gen9] dispatch: interface descriptor exposes internal size-table surface" { + var state: [page_size]u8 = undefined; + const layout = try writeState(&state, &.{ 0xaa, 0xbb }, &.{ 4096, 8192 }); + + try std.testing.expectEqual(@as(u8, 3), layout.surface_count); + + const descriptor_binding_table = std.mem.readInt( + u32, + state[layout.interface_descriptor_offset + 4 * @sizeOf(u32) ..][0..@sizeOf(u32)], + .little, + ); + try std.testing.expectEqual(layout.binding_table_offset | @as(u32, layout.surface_count), descriptor_binding_table); + + for (0..layout.surface_count) |index| { + const entry = std.mem.readInt( + u32, + state[layout.binding_table_offset + index * @sizeOf(u32) ..][0..@sizeOf(u32)], + .little, + ); + try std.testing.expectEqual(layout.surface_offsets[index], entry); + } +} diff --git a/src/intel/compiler/targets/gen9/compute/eu.zig b/src/intel/compiler/targets/gen9/compute/eu.zig new file mode 100644 index 0000000..87fc909 --- /dev/null +++ b/src/intel/compiler/targets/gen9/compute/eu.zig @@ -0,0 +1,96 @@ +pub const Opcode = enum(u7) { + add = 0x40, + addc = 0x4E, + @"and" = 0x05, + asr = 0x0C, + avg = 0x42, + bfe = 0x18, + bfi1 = 0x19, + bfi2 = 0x1A, + bfrev = 0x17, + brc = 0x23, + brd = 0x21, + @"break" = 0x28, + call = 0x2C, + calla = 0x2B, + cbit = 0x4D, + cmp = 0x10, + cmpn = 0x11, + cont = 0x29, + csel = 0x12, + dp2 = 0x57, + dp3 = 0x56, + dp4 = 0x54, + dph = 0x55, + @"else" = 0x24, + endif = 0x25, + fbh = 0x4B, + fbl = 0x4C, + frc = 0x43, + goto = 0x2E, + halt = 0x2A, + @"if" = 0x22, + illegal = 0x00, + jmpi = 0x20, + join = 0x2F, + line = 0x59, + lrp = 0x5C, + lzd = 0x4A, + mac = 0x48, + mach = 0x49, + mad = 0x5B, + madm = 0x5D, + math = 0x38, + mov = 0x01, + movi = 0x03, + mul = 0x41, + nop = 0x7E, + not = 0x04, + @"or" = 0x06, + pln = 0x5A, + ret = 0x2D, + rndd = 0x45, + rnde = 0x46, + rndu = 0x44, + rndz = 0x47, + sad2 = 0x50, + sada2 = 0x51, + sel = 0x02, + send = 0x31, + sendc = 0x32, + sends = 0x33, + sendsc = 0x34, + shl = 0x09, + shr = 0x08, + smov = 0x0A, + subb = 0x4F, + wait = 0x30, + @"while" = 0x27, + xor = 0x07, +}; + +pub const CompareCondition = enum(u3) { + none = 0, + zero = 1, + not_zero = 2, + greater = 3, + greater_or_equal = 4, + less = 5, + less_or_equal = 6, +}; + +pub const MathFunction = enum(u4) { + inv = 1, + log = 2, + exp = 3, + sqrt = 4, + rsq = 5, + sin = 6, + cos = 7, + sincos = 8, + fdiv = 9, + pow = 10, + idiv_and_rem = 11, + idiv = 12, + irem = 13, +}; diff --git a/src/intel/compiler/targets/gen9/compute/eu_encoder.zig b/src/intel/compiler/targets/gen9/compute/eu_encoder.zig index 4359619..59277fa 100644 --- a/src/intel/compiler/targets/gen9/compute/eu_encoder.zig +++ b/src/intel/compiler/targets/gen9/compute/eu_encoder.zig @@ -3,6 +3,7 @@ const device = @import("../../../device.zig"); const ir_instruction = @import("../../../ir/instruction.zig"); const operand = @import("../../../ir/operand.zig"); const message_descriptor = @import("message_descriptor.zig"); +const eu = @import("eu.zig"); pub const Error = error{ UnsupportedExecutionSize, @@ -52,7 +53,7 @@ const Grf = struct { }; pub fn encodeMove(execution_size: device.ExecutionSize, move: ir_instruction.Move) Error!EncodedInstruction { - var encoded = try instructionHeader(1, execution_size); + var encoded = try instructionHeader(.mov, 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)); try setSource0(&encoded, move.source); @@ -63,7 +64,7 @@ pub fn encodeEndThread(header: operand.PhysicalGrf) Error![2]EncodedInstruction if (header.number != 0 or header.byte_offset != 0) return Error.InvalidRegister; - var copy = try instructionHeader(1, .simd8); + var copy = try instructionHeader(.mov, .simd8); copy.setBits(34, 34, 1); // NoMask setDestination(©, .grf, .unsigned_dword, .{ .number = eot_payload_grf, .byte_offset = 0 }, 1); copy.setBits(42, 41, @intFromEnum(RegisterFile.grf)); @@ -73,7 +74,7 @@ pub fn encodeEndThread(header: operand.PhysicalGrf) Error![2]EncodedInstruction copy.setBits(84, 82, 3); copy.setBits(88, 85, 4); - var send = try instructionHeader(49, .simd8); + var send = try instructionHeader(.send, .simd8); send.setBits(34, 34, 1); // NoMask setDestination(&send, .architecture, .unsigned_word, .{ .number = 0, .byte_offset = 0 }, 1); send.setBits(42, 41, @intFromEnum(RegisterFile.grf)); @@ -92,7 +93,7 @@ pub fn encodeEndThread(header: operand.PhysicalGrf) Error![2]EncodedInstruction } pub fn encodeSurfaceMessage(execution_size: device.ExecutionSize, message: ir_instruction.SurfaceMessage) Error!EncodedInstruction { - var encoded = try instructionHeader(49, execution_size); + var encoded = try instructionHeader(.send, execution_size); const descriptor = message_descriptor.encode(message); const payload = try resolveGrf(message.payload.base, 0); if (payload.byte_offset != 0) @@ -121,9 +122,65 @@ pub fn encodeSurfaceMessage(execution_size: device.ExecutionSize, message: ir_in return encoded; } +pub fn encodeJump(displacement_bytes: i32) Error!EncodedInstruction { + return encodeJumpWithPredicate(displacement_bytes, null); +} + +pub fn encodePredicatedJump(displacement_bytes: i32, predicate: operand.Predicate) Error!EncodedInstruction { + return encodeJumpWithPredicate(displacement_bytes, predicate); +} + +fn encodeJumpWithPredicate(displacement_bytes: i32, predicate: ?operand.Predicate) Error!EncodedInstruction { + var encoded = try instructionHeader(.jmpi, .simd1); + encoded.setBits(34, 34, 1); // NoMask + + // JMPI updates the instruction pointer: IP = IP + displacement. + setDestination(&encoded, .architecture, .signed_dword, .{ .number = 0xa0, .byte_offset = 0 }, 1); + encoded.setBits(42, 41, @intFromEnum(RegisterFile.architecture)); + encoded.setBits(46, 43, @intFromEnum(HardwareType.signed_dword)); + encoded.setBits(76, 69, 0xa0); + encoded.setBits(81, 80, 0); + encoded.setBits(84, 82, 0); + encoded.setBits(88, 85, 0); + setSource1Immediate(&encoded, .signed_dword, .{ .i32 = displacement_bytes }); + + if (predicate) |value| { + const flag = switch (value.flag) { + .physical => |physical| physical, + .virtual => return Error.UnsupportedOperand, + }; + if (flag.register != 0 or flag.subregister > 1) + return Error.InvalidRegister; + + encoded.setBits(19, 16, 1); // Normal predicate control. + encoded.setBits(20, 20, @intFromBool(value.inverse)); + encoded.setBits(33, 33, flag.register); + encoded.setBits(32, 32, flag.subregister); + } + + return encoded; +} + +pub fn patchJump(encoded_bytes: []u8, displacement_bytes: i32) Error!void { + if (encoded_bytes.len < 16) + return Error.InvalidRegister; + + var encoded: EncodedInstruction = .{ .words = .{ + std.mem.readInt(u64, encoded_bytes[0..8], .little), + std.mem.readInt(u64, encoded_bytes[8..16], .little), + } }; + if (encoded.bits(6, 0) != @intFromEnum(eu.Opcode.jmpi)) + return Error.UnsupportedOperand; + encoded.setBits(127, 96, @as(u32, @bitCast(displacement_bytes))); + std.mem.writeInt(u64, encoded_bytes[0..8], encoded.words[0], .little); + std.mem.writeInt(u64, encoded_bytes[8..16], encoded.words[1], .little); +} + pub fn encodeBinary(execution_size: device.ExecutionSize, binary: ir_instruction.Binary) Error!EncodedInstruction { - const opcode: u7 = switch (binary.opcode) { - .add => 64, + const opcode: eu.Opcode = switch (binary.opcode) { + .bitwise_xor => .xor, + .add => .add, + .multiply => .mul, else => return Error.UnsupportedOperand, }; @@ -135,27 +192,57 @@ pub fn encodeBinary(execution_size: device.ExecutionSize, binary: ir_instruction return encoded; } +pub fn encodeCompare(execution_size: device.ExecutionSize, compare: ir_instruction.Compare) Error!EncodedInstruction { + const flag = switch (compare.destination) { + .physical => |value| value, + .virtual => return Error.UnsupportedOperand, + }; + if (flag.register != 0 or flag.subregister > 1) + return Error.InvalidRegister; + + var encoded = try instructionHeader(.cmp, execution_size); + setDestination(&encoded, .architecture, try hardwareType(compare.lhs.type), .{ .number = 0, .byte_offset = 0 }, 1); + try setSource0(&encoded, compare.lhs); + try setSource1(&encoded, compare.rhs); + + const condition: eu.CompareCondition = switch (compare.opcode) { + .equal => .zero, + .not_equal => .not_zero, + .greater_than => .greater, + .greater_or_equal => .greater_or_equal, + .less_than => .less, + .less_or_equal => .less_or_equal, + }; + + encoded.setBits(27, 24, @intFromEnum(condition)); + encoded.setBits(33, 33, flag.register); + encoded.setBits(32, 32, flag.subregister); + 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); + var encoded = try instructionHeader(.math, 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, - }); + const function: eu.MathFunction = switch (math.opcode) { + .integer_quotient => .idiv, + }; + + encoded.setBits(27, 24, @intFromEnum(function)); return encoded; } -fn instructionHeader(opcode: u7, execution_size: device.ExecutionSize) Error!EncodedInstruction { +fn instructionHeader(opcode: eu.Opcode, execution_size: device.ExecutionSize) Error!EncodedInstruction { var encoded: EncodedInstruction = .{}; - encoded.setBits(6, 0, opcode); + encoded.setBits(6, 0, @intFromEnum(opcode)); encoded.setBits(23, 21, try executionSize(execution_size)); return encoded; } @@ -335,3 +422,70 @@ fn verticalStride(stride: u8) Error!u4 { else => Error.InvalidRegion, }; } + +fn testBinary(opcode: ir_instruction.BinaryOpcode) ir_instruction.Binary { + return .{ + .opcode = opcode, + .destination = .{ + .register = .{ .physical_grf = .{ .number = 3 } }, + .type = .u32, + }, + .lhs = .{ + .register = .{ .physical_grf = .{ .number = 1 } }, + .type = .u32, + .region = operand.Region.contiguous(.simd8), + }, + .rhs = .{ + .register = .{ .immediate = .{ .u32 = 16 } }, + .type = .u32, + .region = operand.Region.broadcast(), + }, + }; +} + +test "[gen9] EU encoder: encode integer multiply" { + const encoded = try encodeBinary(.simd8, testBinary(.multiply)); + try std.testing.expectEqual(@as(u64, 65), encoded.bits(6, 0)); + try std.testing.expectEqual(@as(u64, 3), encoded.bits(23, 21)); +} + +test "[gen9] EU encoder: encode bitwise XOR" { + const encoded = try encodeBinary(.simd8, testBinary(.bitwise_xor)); + try std.testing.expectEqual(@as(u64, 7), encoded.bits(6, 0)); + try std.testing.expectEqual(@as(u64, 16), encoded.bits(127, 96)); +} + +test "[gen9] EU encoder: encode unsigned less-than comparison" { + const encoded = try encodeCompare(.simd8, .{ + .opcode = .less_than, + .destination = .{ .physical = .{ .register = 0, .subregister = 1 } }, + .lhs = .{ + .register = .{ .physical_grf = .{ .number = 1 } }, + .type = .u32, + .region = operand.Region.contiguous(.simd8), + }, + .rhs = .{ + .register = .{ .physical_grf = .{ .number = 2 } }, + .type = .u32, + .region = operand.Region.contiguous(.simd8), + }, + }); + + try std.testing.expectEqual(@as(u64, 16), encoded.bits(6, 0)); + try std.testing.expectEqual(@as(u64, 5), encoded.bits(27, 24)); + try std.testing.expectEqual(@as(u64, 0), encoded.bits(33, 33)); + try std.testing.expectEqual(@as(u64, 1), encoded.bits(32, 32)); +} + +test "[gen9] EU encoder: encode predicated jump" { + const encoded = try encodePredicatedJump(-32, .{ + .flag = .{ .physical = .{ .register = 0, .subregister = 1 } }, + .inverse = true, + }); + + try std.testing.expectEqual(@as(u64, @intFromEnum(eu.Opcode.jmpi)), encoded.bits(6, 0)); + try std.testing.expectEqual(@as(u64, 1), encoded.bits(19, 16)); + try std.testing.expectEqual(@as(u64, 1), encoded.bits(20, 20)); + try std.testing.expectEqual(@as(u64, 1), encoded.bits(32, 32)); + try std.testing.expectEqual(@as(i32, -32), @as(i32, @bitCast(@as(u32, @truncate(encoded.bits(127, 96)))))); +} diff --git a/src/intel/compiler/targets/gen9/compute/kernel_encoder.zig b/src/intel/compiler/targets/gen9/compute/kernel_encoder.zig index 8e7829a..ecacf57 100644 --- a/src/intel/compiler/targets/gen9/compute/kernel_encoder.zig +++ b/src/intel/compiler/targets/gen9/compute/kernel_encoder.zig @@ -1,6 +1,7 @@ const std = @import("std"); const eu = @import("eu_encoder.zig"); +const ids = @import("../../../ir/id.zig"); const program_ir = @import("../../../ir/program.zig"); pub const Error = std.mem.Allocator.Error || eu.Error || error{ @@ -11,6 +12,11 @@ pub const Error = std.mem.Allocator.Error || eu.Error || error{ EotRegisterUnavailable, }; +const JumpFixup = struct { + instruction_offset: usize, + target: ids.BlockId, +}; + pub fn encode(allocator: std.mem.Allocator, program: *program_ir.Program) Error![]u8 { if (!program.properties.registers_allocated) return Error.InvalidProgram; @@ -18,65 +24,188 @@ pub fn encode(allocator: std.mem.Allocator, program: *program_ir.Program) Error! return Error.EotRegisterUnavailable; const entry_id = program.entry_block orelse return Error.InvalidProgram; - const entry = program.blocks.get(entry_id) orelse return Error.InvalidProgram; - var live_block_count: usize = 0; - for (program.blocks.entries.items) |block| { - if (block != null) - 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; + if (!program.blocks.isLive(entry_id)) + return Error.InvalidProgram; + + const block_offsets = try allocator.alloc(?usize, program.blocks.entries.items.len); + defer allocator.free(block_offsets); + @memset(block_offsets, null); + + var block_order: std.ArrayList(ids.BlockId) = .empty; + defer block_order.deinit(allocator); + try block_order.append(allocator, entry_id); + for (program.blocks.entries.items, 0..) |block, block_index| { + if (block != null and block_index != entry_id.index()) + try block_order.append(allocator, ids.BlockId.fromIndex(block_index)); } + var fixups: std.ArrayList(JumpFixup) = .empty; + defer fixups.deinit(allocator); 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) { - 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; - } + for (block_order.items) |block_id| { + const block = program.blocks.get(block_id) orelse return Error.InvalidProgram; + block_offsets[block_id.index()] = kernel.items.len; - const encoded = switch (instruction.operation) { - .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; + for (block.instructions.items) |instruction_id| + try encodeInstruction(allocator, &kernel, program, instruction_id); + + const terminator = block.terminator orelse return Error.InvalidProgram; + switch (terminator) { + .jump => |edge| { + const instruction_offset = kernel.items.len; + try appendInstruction(allocator, &kernel, try eu.encodeJump(0)); + try fixups.append(allocator, .{ + .instruction_offset = instruction_offset, + .target = edge.target, + }); }, - } 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); + .conditional_branch => |branch| { + const true_instruction_offset = kernel.items.len; + try appendInstruction(allocator, &kernel, try eu.encodePredicatedJump(0, branch.predicate)); + try fixups.append(allocator, .{ + .instruction_offset = true_instruction_offset, + .target = branch.true_edge.target, + }); + + const false_instruction_offset = kernel.items.len; + try appendInstruction(allocator, &kernel, try eu.encodeJump(0)); + try fixups.append(allocator, .{ + .instruction_offset = false_instruction_offset, + .target = branch.false_edge.target, + }); + }, + .end_thread => { + const header = program.payload.header_grf orelse return Error.InvalidProgram; + const instructions = try eu.encodeEndThread(header); + for (instructions) |encoded| + try appendInstruction(allocator, &kernel, encoded); + program.program_data.total_grf_count = eu.eot_payload_grf + 1; + }, + .@"unreachable" => return Error.UnsupportedControlFlow, + } } - const terminator = entry.terminator orelse return Error.InvalidProgram; - switch (terminator) { - .end_thread => { - const header = program.payload.header_grf orelse return Error.InvalidProgram; - const instructions = try eu.encodeEndThread(header); - for (instructions) |instruction| - try appendInstruction(allocator, &kernel, instruction); - program.program_data.total_grf_count = eu.eot_payload_grf + 1; - }, - else => { - std.log.scoped(.FlintEuEncoder).err("cannot encode entry-block terminator {t}", .{std.meta.activeTag(terminator)}); + for (fixups.items) |fixup| { + if (fixup.target.index() >= block_offsets.len) + return Error.InvalidProgram; + const target_offset = block_offsets[fixup.target.index()] orelse return Error.InvalidProgram; + const next_instruction_offset = fixup.instruction_offset + 16; + const displacement = std.math.cast(i32, @as(i64, @intCast(target_offset)) - @as(i64, @intCast(next_instruction_offset))) orelse return Error.UnsupportedControlFlow; - }, + try eu.patchJump(kernel.items[fixup.instruction_offset..][0..16], displacement); } return kernel.toOwnedSlice(allocator); } +fn encodeInstruction(allocator: std.mem.Allocator, kernel: *std.ArrayList(u8), program: *const program_ir.Program, instruction_id: ids.InstructionId) Error!void { + const inst = program.instructions.get(instruction_id) orelse return Error.InvalidProgram; + if (inst.predicate != null) { + std.log.scoped(.FlintEuEncoder).err("cannot encode instruction {d} ({t}): predication is not supported", .{ instruction_id.index(), std.meta.activeTag(inst.operation) }); + return Error.UnsupportedPredication; + } + + const encoded = switch (inst.operation) { + .move => |move| eu.encodeMove(inst.execution_size, move), + .surface_message => |message| eu.encodeSurfaceMessage(inst.execution_size, message), + .binary => |binary| eu.encodeBinary(inst.execution_size, binary), + .compare => |compare| eu.encodeCompare(inst.execution_size, compare), + .math => |math| eu.encodeMath(inst.execution_size, math), + else => { + std.log.scoped(.FlintEuEncoder).err("cannot encode instruction {d}: unsupported operation {t}", .{ instruction_id.index(), std.meta.activeTag(inst.operation) }); + return Error.UnsupportedOperation; + }, + } catch |err| { + std.log.scoped(.FlintEuEncoder).err("failed to encode instruction {d} ({t}): {s}", .{ instruction_id.index(), std.meta.activeTag(inst.operation), @errorName(err) }); + if (err == error.InvalidRegion) switch (inst.operation) { + .move => |move| std.log.scoped(.FlintEuEncoder).err("move in block {d}: destination {t} byte={d} hstride={d}; source {t} byte={d} vstride={d} width={d} hstride={d}", .{ + inst.parent_block.index(), + move.destination.register, + move.destination.region.byte_offset, + move.destination.region.horizontal_stride, + move.source.register, + move.source.region.byte_offset, + move.source.region.vertical_stride, + move.source.region.width, + move.source.region.horizontal_stride, + }), + else => {}, + }; + return err; + }; + try appendInstruction(allocator, kernel, encoded); +} + fn appendInstruction(allocator: std.mem.Allocator, kernel: *std.ArrayList(u8), instruction: eu.EncodedInstruction) std.mem.Allocator.Error!void { var bytes: [16]u8 = undefined; std.mem.writeInt(u64, bytes[0..8], instruction.words[0], .little); std.mem.writeInt(u64, bytes[8..16], instruction.words[1], .little); try kernel.appendSlice(allocator, &bytes); } + +test "[gen9] kernel encoder: patch unconditional jump between blocks" { + 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 entry = try program.addBlock("entry"); + const exit = try program.addBlock("exit"); + try program.setEntryBlock(entry); + try program.setTerminator(entry, .{ .jump = .{ .target = exit, .arguments = &.{} } }); + try program.setTerminator(exit, .end_thread); + program.payload.header_grf = .{ .number = 0 }; + program.properties.registers_allocated = true; + + const kernel = try encode(std.testing.allocator, &program); + defer std.testing.allocator.free(kernel); + + try std.testing.expectEqual(@as(usize, 48), kernel.len); + try std.testing.expectEqual(@as(u7, 32), @as(u7, @truncate(std.mem.readInt(u64, kernel[0..8], .little)))); + try std.testing.expectEqual(@as(i32, 0), @as(i32, @bitCast(std.mem.readInt(u32, kernel[12..16], .little)))); +} + +test "[gen9] kernel encoder: patch conditional branch targets" { + 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 entry = try program.addBlock("entry"); + const true_block = try program.addBlock("true"); + const false_block = try program.addBlock("false"); + try program.setEntryBlock(entry); + try program.setTerminator(entry, .{ .conditional_branch = .{ + .predicate = .{ .flag = .{ .physical = .{ .register = 0, .subregister = 1 } } }, + .true_edge = .{ .target = true_block, .arguments = &.{} }, + .false_edge = .{ .target = false_block, .arguments = &.{} }, + } }); + try program.setTerminator(true_block, .end_thread); + try program.setTerminator(false_block, .end_thread); + program.payload.header_grf = .{ .number = 0 }; + program.properties.registers_allocated = true; + + const kernel = try encode(std.testing.allocator, &program); + defer std.testing.allocator.free(kernel); + + try std.testing.expectEqual(@as(usize, 96), kernel.len); + try std.testing.expectEqual(@as(i32, 16), @as(i32, @bitCast(std.mem.readInt(u32, kernel[12..16], .little)))); + try std.testing.expectEqual(@as(i32, 32), @as(i32, @bitCast(std.mem.readInt(u32, kernel[28..32], .little)))); + const first_word = std.mem.readInt(u64, kernel[0..8], .little); + try std.testing.expectEqual(@as(u64, 1), (first_word >> 16) & 0xf); + try std.testing.expectEqual(@as(u64, 1), (first_word >> 32) & 0x1); +} diff --git a/src/intel/compiler/targets/gen9/compute/pipeline.zig b/src/intel/compiler/targets/gen9/compute/pipeline.zig index 03a6dcf..17082fb 100644 --- a/src/intel/compiler/targets/gen9/compute/pipeline.zig +++ b/src/intel/compiler/targets/gen9/compute/pipeline.zig @@ -18,6 +18,8 @@ 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"); +const regions = @import("regions.zig"); +const system_values = @import("system_values.zig"); pub const Error = common_ir.Error || block_arguments.Error || @@ -30,6 +32,7 @@ pub const Error = common_ir.Error || message_payloads.Error || resource_layout.Error || resource_lowering.Error || + system_values.Error || flag_allocation.Error || register_allocation.Error || compute.Error || @@ -75,6 +78,7 @@ pub fn compile(allocator: std.mem.Allocator, module: *shader_ir.module.Module, d errdefer program.deinit(); try abi.run(&program); + try system_values.run(&program); try block_arguments.run(allocator, &program); try parallel_copies.run(allocator, &program); @@ -86,6 +90,7 @@ pub fn compile(allocator: std.mem.Allocator, module: *shader_ir.module.Module, d try message_lowering.run(&program); try message_addresses.run(&program); try message_payloads.run(&program); + regions.run(&program); try flag_allocation.run(allocator, &program); try register_allocation.run(allocator, &program); @@ -98,6 +103,8 @@ pub fn compile(allocator: std.mem.Allocator, module: *shader_ir.module.Module, d error.UnsupportedExecutionSize, error.UnsupportedDataType, error.UnsupportedOperand, + error.InvalidRegister, + error.InvalidRegion, error.EotRegisterUnavailable, => null, else => return err, diff --git a/src/intel/compiler/targets/gen9/compute/regions.zig b/src/intel/compiler/targets/gen9/compute/regions.zig new file mode 100644 index 0000000..0705fb5 --- /dev/null +++ b/src/intel/compiler/targets/gen9/compute/regions.zig @@ -0,0 +1,49 @@ +const operand = @import("../../../ir/operand.zig"); +const program_ir = @import("../../../ir/program.zig"); + +pub fn run(program: *program_ir.Program) void { + if (program.properties.regions_legalized) + return; + + for (program.instructions.entries.items) |*entry| { + const inst = if (entry.*) |*value| value else continue; + switch (inst.operation) { + .load_buffer => |*op| legalizeSource(&op.byte_offset, inst.execution_size), + .store_buffer => |*op| { + legalizeSource(&op.byte_offset, inst.execution_size); + legalizeSource(&op.source, inst.execution_size); + }, + .array_length => |*op| legalizeSource(&op.byte_offset, inst.execution_size), + .surface_read => |*op| legalizeSource(&op.address, inst.execution_size), + .surface_write => |*op| { + legalizeSource(&op.address, inst.execution_size); + legalizeSource(&op.data, inst.execution_size); + }, + .move => |*op| legalizeSource(&op.source, inst.execution_size), + .binary => |*op| { + legalizeSource(&op.lhs, inst.execution_size); + legalizeSource(&op.rhs, inst.execution_size); + }, + .math => |*op| { + legalizeSource(&op.lhs, inst.execution_size); + legalizeSource(&op.rhs, inst.execution_size); + }, + .compare => |*op| { + legalizeSource(&op.lhs, inst.execution_size); + legalizeSource(&op.rhs, inst.execution_size); + }, + else => {}, + } + } + + program.properties.regions_legalized = true; +} + +fn legalizeSource(source: *operand.Source, execution_size: @import("../../../device.zig").ExecutionSize) void { + const byte_offset = source.region.byte_offset; + source.region = switch (source.register) { + .immediate => operand.Region.broadcast(), + else => operand.Region.contiguous(execution_size), + }; + source.region.byte_offset = byte_offset; +} diff --git a/src/intel/compiler/targets/gen9/compute/system_values.zig b/src/intel/compiler/targets/gen9/compute/system_values.zig new file mode 100644 index 0000000..a56d42d --- /dev/null +++ b/src/intel/compiler/targets/gen9/compute/system_values.zig @@ -0,0 +1,96 @@ +const std = @import("std"); + +const instruction = @import("../../../ir/instruction.zig"); +const operand = @import("../../../ir/operand.zig"); +const program_ir = @import("../../../ir/program.zig"); +const validator = @import("../../../ir/validator.zig"); + +pub const Error = error{InvalidProgram}; + +pub fn run(program: *program_ir.Program) Error!void { + validator.validate(program) catch return Error.InvalidProgram; + if (program.properties.system_values_lowered) + return; + + // Flint dispatch currently accepts only one invocation in one workgroup at + // base group zero, so every component of GlobalInvocationId is zero. + if (!std.mem.eql(u32, &program.workgroup_size, &.{ 1, 1, 1 })) + return; + + for (program.instructions.entries.items) |*entry| { + const inst = if (entry.*) |*value| value else continue; + inst.operation = switch (inst.operation) { + .load_global_invocation_id => |op| .{ .move = .{ + .destination = op.destination, + .source = zero(), + } }, + else => inst.operation, + }; + } + + program.properties.system_values_lowered = true; + validator.validate(program) catch return Error.InvalidProgram; +} + +fn zero() operand.Source { + return .{ + .register = .{ .immediate = .{ .u32 = 0 } }, + .type = .u32, + .region = operand.Region.broadcast(), + }; +} + +test "[gen9] system values: lower global invocation ID for single invocation" { + const Builder = @import("../../../ir/Builder.zig"); + 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(); + var builder = Builder.init(&program); + + const destination = try builder.addVirtualRegister(.{ + .size_bytes = 32, + .alignment_bytes = 32, + .element_type = .u32, + .lane_count = 8, + .class = .temporary, + }); + const entry = try builder.addBlock("entry"); + const load = try builder.appendInstruction(entry, .simd8, null, .{ .load_global_invocation_id = .{ + .destination = .{ .register = .{ .virtual = destination }, .type = .u32 }, + .component = 2, + } }); + try builder.setTerminator(entry, .end_thread); + try builder.setEntryBlock(entry); + + try run(&program); + + try std.testing.expect(program.properties.system_values_lowered); + const move = program.instructions.get(load).?.operation.move; + try std.testing.expectEqual(@as(u32, 0), move.source.register.immediate.u32); +} + +test "[gen9] system values: preserve IDs for unsupported workgroup sizes" { + 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, .{ 2, 1, 1 }, device_info, .simd8); + defer program.deinit(); + const entry = try program.addBlock("entry"); + try program.setTerminator(entry, .end_thread); + try program.setEntryBlock(entry); + + try run(&program); + try std.testing.expect(!program.properties.system_values_lowered); +} diff --git a/src/intel/compiler/targets/gen9/gen9.zig b/src/intel/compiler/targets/gen9/gen9.zig index 58975eb..c47db3c 100644 --- a/src/intel/compiler/targets/gen9/gen9.zig +++ b/src/intel/compiler/targets/gen9/gen9.zig @@ -95,6 +95,7 @@ test "[gen9] target: lower 256 KiB SSBO copy loop" { }; var artifact = try compileCompute(std.testing.allocator, &module, gen9_device, .{}); defer artifact.deinit(std.testing.allocator); + try std.testing.expect(artifact.kernel != null); const program = &artifact.program; const resources = &artifact.resources;