[Flint] adding math and compare eu encoding

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 |
+| vkCreateSemaphore                                | ⚙️ WIP         |
-| vkCreateShaderModule                             | ⚙️ WIP         |
+| vkCreateShaderModule                             | ✅ Implemented |
-| vkCreateSwapchainKHR                             | ⚙️ WIP         |
+| vkCreateSwapchainKHR                             | ✅ Implemented |
-| vkCreateWaylandSurfaceKHR                        | ⚙️ WIP         |
+| vkCreateWaylandSurfaceKHR                        | ✅ Implemented |
 | vkCreateWin32SurfaceKHR                          | ⚙️ WIP         |
-| vkCreateXcbSurfaceKHR                            | ⚙️ WIP         |
+| vkCreateXcbSurfaceKHR                            | ✅ Implemented |
-| vkCreateXlibSurfaceKHR                           | ⚙️ WIP         |
+| 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         |
+| vkCmdDispatch                                    | ✅ Implemented |
-| vkCmdDispatchBaseKHR                             | ⚙️ WIP         |
+| 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         |
+| vkCreateComputePipelines                         | ✅ Implemented |
-| vkCreateDescriptorPool                           | ⚙️ WIP         |
+| vkCreateDescriptorPool                           | ✅ Implemented |
-| vkCreateDescriptorSetLayout                      | ⚙️ WIP         |
+| 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         |
+| vkCreateShaderModule                             | ✅ Implemented |
-| vkCreateSwapchainKHR                             | ⚙️ WIP         |
+| vkCreateSwapchainKHR                             | ✅ Implemented |
-| vkCreateWaylandSurfaceKHR                        | ⚙️ WIP         |
+| vkCreateWaylandSurfaceKHR                        | ✅ Implemented |
 | vkCreateWin32SurfaceKHR                          | ⚙️ WIP         |
-| vkCreateXcbSurfaceKHR                            | ⚙️ WIP         |
+| vkCreateXcbSurfaceKHR                            | ✅ Implemented |
-| vkCreateXlibSurfaceKHR                           | ⚙️ WIP         |
+| vkCreateXlibSurfaceKHR                           | ✅ Implemented |
 | vkDestroyBuffer                                  | ✅ Implemented |
 | vkDestroyBufferView                              | ⚙️ WIP         |
 | vkDestroyCommandPool                             | ✅ Implemented |
-| vkDestroyDescriptorPool                          | ⚙️ WIP         |
+| vkDestroyDescriptorPool                          | ✅ Implemented |
-| vkDestroyDescriptorSetLayout                     | ⚙️ WIP         |
+| 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         |
+| vkDestroyShaderModule                            | ✅ Implemented |
-| vkDestroySurfaceKHR                              | ⚙️ WIP         |
+| vkDestroySurfaceKHR                              | ✅ Implemented |
-| vkDestroySwapchainKHR                            | ⚙️ WIP         |
+| 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         |

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);
 }

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};

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);
+    }
+}

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,
+};

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(&copy, .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) {
+    const opcode: eu.Opcode = switch (binary.opcode) {
-        .add => 64,
+        .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) {
+    const function: eu.MathFunction = switch (math.opcode) {
-        .integer_quotient => 12,
+        .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))))));
+}

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;
+    if (!program.blocks.isLive(entry_id))
-    var live_block_count: usize = 0;
+        return Error.InvalidProgram;
-    for (program.blocks.entries.items) |block| {
+
-        if (block != null)
+    const block_offsets = try allocator.alloc(?usize, program.blocks.entries.items.len);
-            live_block_count += 1;
+    defer allocator.free(block_offsets);
-    }
+    @memset(block_offsets, null);
-
+
-    if (live_block_count != 1) {
+    var block_order: std.ArrayList(ids.BlockId) = .empty;
-        std.log.scoped(.FlintEuEncoder).err("cannot encode control flow: program has {d} live blocks; only one is currently supported", .{live_block_count});
+    defer block_order.deinit(allocator);
-        return Error.UnsupportedControlFlow;
+    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| {
+    for (block_order.items) |block_id| {
-        const instruction = program.instructions.get(instruction_id) orelse return Error.InvalidProgram;
+        const block = program.blocks.get(block_id) orelse return Error.InvalidProgram;
-        if (instruction.predicate != null) {
+        block_offsets[block_id.index()] = kernel.items.len;
-            std.log.scoped(.FlintEuEncoder).err("cannot encode instruction {d} ({t}): predication is not supported", .{ instruction_id.index(), std.meta.activeTag(instruction.operation) });
-            return Error.UnsupportedPredication;
-        }
 
-        const encoded = switch (instruction.operation) {
+        for (block.instructions.items) |instruction_id|
-            .move => |move| eu.encodeMove(instruction.execution_size, move),
+            try encodeInstruction(allocator, &kernel, program, instruction_id);
-            .surface_message => |message| eu.encodeSurfaceMessage(instruction.execution_size, message),
-            .binary => |binary| eu.encodeBinary(instruction.execution_size, binary),
-            .math => |math| eu.encodeMath(instruction.execution_size, math),
-            else => {
-                std.log.scoped(.FlintEuEncoder).err("cannot encode instruction {d}: unsupported operation {t}", .{ instruction_id.index(), std.meta.activeTag(instruction.operation) });
-                return Error.UnsupportedOperation;
-            },
-        } catch |err| {
-            std.log.scoped(.FlintEuEncoder).err("failed to encode instruction {d} ({t}): {s}", .{ instruction_id.index(), std.meta.activeTag(instruction.operation), @errorName(err) });
-            return err;
-        };
-        try appendInstruction(allocator, &kernel, encoded);
-    }
 
-    const terminator = entry.terminator orelse return Error.InvalidProgram;
+        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,
+                });
+            },
+            .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) |instruction|
+                for (instructions) |encoded|
-                try appendInstruction(allocator, &kernel, instruction);
+                    try appendInstruction(allocator, &kernel, encoded);
                 program.program_data.total_grf_count = eu.eot_payload_grf + 1;
             },
-        else => {
+            .@"unreachable" => return Error.UnsupportedControlFlow,
-            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);
+}

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,

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;
+}

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);
+}

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;