[Flint] adding gen9 eu encoding

build.zig

@@ -470,8 +470,8 @@ fn customSoft(
 // Flint specialized functions
 
 fn customFlint(
-    _: *std.Build,
-    _: *Step.Options,
+    b: *std.Build,
+    options: *Step.Options,
     _: *Step.Compile,
     lib_mod: *std.Build.Module,
     _: *std.Build.Module,
@@ -484,6 +484,12 @@ fn customFlint(
 ) !void {
     lib_mod.addImport("intel_c", base_c_mod);
     lib_mod.addImport("shader_ir", shader_ir_mod);
+
+    const dump_common_ir = b.option(bool, "flint-dump-common-ir", "Print backend-agnostic shader IR after translation") orelse false;
+    const dump_ir = b.option(bool, "flint-dump-ir", "Print final Flint IR after backend lowering") orelse false;
+
+    options.addOption(bool, "flint_dump_common_ir", dump_common_ir);
+    options.addOption(bool, "flint_dump_ir", dump_ir);
 }
 
 // Phi specialized functions

src/intel/FlintCommandBuffer.zig

@@ -12,6 +12,7 @@ const MemoryRange = @import("MemoryRange.zig");
 
 const copy = @import("copy_commands.zig");
 const blitter = @import("blitter.zig");
+const gen9_dispatch = @import("compiler/targets/gen9/compute/dispatch.zig");
 
 const Self = @This();
 pub const Interface = base.CommandBuffer;
@@ -19,6 +20,8 @@ pub const Interface = base.CommandBuffer;
 interface: Interface,
 batch: std.ArrayList(u32),
 relocations: std.ArrayList(kmd.Relocation),
+gpu_allocations: std.ArrayList(kmd.Memory),
+engine: ?kmd.Engine,
 bound_compute_pipeline: ?*FlintPipeline,
 bound_compute_descriptor_sets: [base.vulkan_max_descriptor_sets]?*FlintDescriptorSet,
 
@@ -84,6 +87,8 @@ pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const v
         .interface = interface,
         .batch = .empty,
         .relocations = .empty,
+        .gpu_allocations = .empty,
+        .engine = null,
         .bound_compute_pipeline = null,
         .bound_compute_descriptor_sets = @splat(null),
     };
@@ -93,8 +98,10 @@ pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const v
 pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void {
     const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
     const command_allocator = self.interface.host_allocator.allocator();
+    self.releaseGpuAllocations();
     self.batch.deinit(command_allocator);
     self.relocations.deinit(command_allocator);
+    self.gpu_allocations.deinit(command_allocator);
     allocator.destroy(self);
 }
 
@@ -105,7 +112,7 @@ pub fn submitGpuBatch(self: *Self, syncs: []const kmd.SyncDependency) VkError!vo
     // Empty command buffers still need a no-op submission to carry queue synchronization.
     const device: *FlintDevice = @alignCast(@fieldParentPtr("interface", self.interface.owner));
     const allocator = self.interface.host_allocator.allocator();
-    try device.kmd.submitBatch(self.interface.owner.io(), allocator, self.batch.items, self.relocations.items, syncs);
+    try device.kmd.submitBatch(self.interface.owner.io(), allocator, self.engine orelse .blitter, self.batch.items, self.relocations.items, syncs);
 }
 
 pub fn begin(interface: *Interface, info: *const vk.CommandBufferBeginInfo) VkError!void {
@@ -119,22 +126,46 @@ pub fn end(interface: *Interface) VkError!void {
 
 pub fn reset(interface: *Interface, flags: vk.CommandBufferResetFlags) VkError!void {
     const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
+    self.releaseGpuAllocations();
     if (flags.release_resources_bit) {
         const command_allocator = self.interface.host_allocator.allocator();
         self.batch.clearAndFree(command_allocator);
         self.relocations.clearAndFree(command_allocator);
+        self.gpu_allocations.clearAndFree(command_allocator);
     } else {
         self.batch.clearRetainingCapacity();
         self.relocations.clearRetainingCapacity();
+        self.gpu_allocations.clearRetainingCapacity();
     }
+    self.engine = null;
     self.bound_compute_pipeline = null;
     self.bound_compute_descriptor_sets = @splat(null);
 }
 
+fn releaseGpuAllocations(self: *Self) void {
+    const device: *FlintDevice = @alignCast(@fieldParentPtr("interface", self.interface.owner));
+    for (self.gpu_allocations.items) |*allocation|
+        allocation.deinit(&device.kmd, self.interface.owner.io());
+    self.gpu_allocations.clearRetainingCapacity();
+}
+
+pub fn requireEngine(self: *Self, engine: kmd.Engine) VkError!void {
+    if (self.engine) |current| {
+        if (current != engine)
+            return VkError.FeatureNotPresent;
+    } else {
+        self.engine = engine;
+    }
+}
+
 pub fn emit(self: *Self, dword: u32) VkError!void {
     self.batch.append(self.interface.host_allocator.allocator(), dword) catch return VkError.OutOfHostMemory;
 }
 
+fn emitSlice(self: *Self, words: []const u32) VkError!void {
+    self.batch.appendSlice(self.interface.host_allocator.allocator(), words) catch return VkError.OutOfHostMemory;
+}
+
 pub fn emitRelocatedAddress(self: *Self, range: MemoryRange, read: bool, write: bool) VkError!void {
     const address_offset = self.batch.items.len * @sizeOf(u32);
     try self.emit(@intCast(range.offset));
@@ -145,6 +176,7 @@ pub fn emitRelocatedAddress(self: *Self, range: MemoryRange, read: bool, write:
         .delta = @intCast(range.offset),
         .read = read,
         .write = write,
+        .domain = if ((self.engine orelse .blitter) == .render) .render else .none,
     }) catch return VkError.OutOfHostMemory;
 }
 
@@ -300,35 +332,132 @@ pub fn dispatchBase(interface: *Interface, base_group_x: u32, base_group_y: u32,
     const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
     if (group_count_x == 0 or group_count_y == 0 or group_count_z == 0)
         return;
-
-    inline for ([_]struct { u32, u32 }{
-        .{ base_group_x, group_count_x },
-        .{ base_group_y, group_count_y },
-        .{ base_group_z, group_count_z },
-    }) |dimension| {
-        const group_end = std.math.add(u32, dimension[0], dimension[1]) catch return VkError.ValidationFailed;
-        if (group_end > 65535)
-            return VkError.ValidationFailed;
-    }
+    if (base_group_x != 0 or base_group_y != 0 or base_group_z != 0 or
+        group_count_x != 1 or group_count_y != 1 or group_count_z != 1)
+        return VkError.FeatureNotPresent;
 
     const pipeline = self.bound_compute_pipeline orelse return VkError.ValidationFailed;
     const artifact = pipeline.computeArtifact() orelse return VkError.FeatureNotPresent;
+    const kernel = artifact.kernel orelse return VkError.FeatureNotPresent;
+    if (!std.mem.eql(u32, &artifact.program.workgroup_size, &.{ 1, 1, 1 }) or
+        artifact.program.program_data.scratch_size_bytes != 0)
+        return VkError.FeatureNotPresent;
+
+    var ranges: [gen9_dispatch.max_surfaces]?MemoryRange = @splat(null);
+    var sizes: [gen9_dispatch.max_surfaces]u64 = @splat(0);
     for (artifact.resources.bindings) |resource| {
-        if (resource.set >= base.vulkan_max_descriptor_sets)
+        if (resource.set >= base.vulkan_max_descriptor_sets or @as(usize, resource.binding_table_index) >= gen9_dispatch.max_surfaces)
             return VkError.ValidationFailed;
 
         const descriptor_set = self.bound_compute_descriptor_sets[resource.set] orelse return VkError.ValidationFailed;
         const expected_layout = pipeline.interface.layout.set_layouts[resource.set] orelse return VkError.ValidationFailed;
-
         if (descriptor_set.interface.layout != expected_layout)
             return VkError.ValidationFailed;
 
         const descriptor = try descriptor_set.getBuffer(resource.binding, 0);
         const buffer = descriptor.buffer orelse return VkError.ValidationFailed;
-
         if (!buffer.usage.storage_buffer_bit or buffer.memory == null)
             return VkError.ValidationFailed;
+
+        const range = try MemoryRange.fromBuffer(buffer, descriptor.offset, descriptor.size);
+        ranges[resource.binding_table_index] = range;
+        sizes[resource.binding_table_index] = range.size;
     }
+
+    const old_engine = self.engine;
+    try self.requireEngine(.render);
+    const old_batch_len = self.batch.items.len;
+    const old_relocation_len = self.relocations.items.len;
+    const old_allocation_len = self.gpu_allocations.items.len;
+    errdefer {
+        self.engine = old_engine;
+        self.batch.items.len = old_batch_len;
+        self.relocations.items.len = old_relocation_len;
+        while (self.gpu_allocations.items.len > old_allocation_len) {
+            const device: *FlintDevice = @alignCast(@fieldParentPtr("interface", self.interface.owner));
+            self.gpu_allocations.items[self.gpu_allocations.items.len - 1].deinit(&device.kmd, self.interface.owner.io());
+            self.gpu_allocations.items.len -= 1;
+        }
+    }
+
+    const device: *FlintDevice = @alignCast(@fieldParentPtr("interface", self.interface.owner));
+    var state = try device.kmd.allocateMemory(self.interface.owner.io(), gen9_dispatch.page_size);
+    var state_owned = true;
+    errdefer if (state_owned) state.deinit(&device.kmd, self.interface.owner.io());
+
+    const mapped = try state.map(&device.kmd, self.interface.owner.io(), 0, gen9_dispatch.page_size);
+    const state_layout = gen9_dispatch.writeState(mapped, kernel, sizes[0..artifact.resources.bindings.len]) catch |err| switch (err) {
+        error.StateTooLarge,
+        error.UnsupportedBufferSize,
+        error.EmptyBuffer,
+        error.TooManySurfaces,
+        => return VkError.FeatureNotPresent,
+    };
+    state.unmap();
+    try state.flushRange(&device.kmd, self.interface.owner.io(), 0, state_layout.size);
+    const state_handle = try state.handle();
+
+    self.gpu_allocations.append(self.interface.host_allocator.allocator(), state) catch return VkError.OutOfHostMemory;
+    state_owned = false;
+
+    for (0..@as(usize, state_layout.surface_count)) |index| {
+        const range = ranges[index] orelse return VkError.ValidationFailed;
+        if (range.offset > std.math.maxInt(u32))
+            return VkError.FeatureNotPresent;
+        self.relocations.append(self.interface.host_allocator.allocator(), .{
+            .source_handle = state_handle,
+            .target_handle = try range.memory.allocation.handle(),
+            .offset = state_layout.surface_address_offsets[index],
+            .delta = @intCast(range.offset),
+            .read = true,
+            .write = true,
+            .domain = .render,
+        }) catch return VkError.OutOfHostMemory;
+    }
+
+    try self.emitSlice(&gen9_dispatch.pipeControl(gen9_dispatch.pipe_control.cs_stall |
+        gen9_dispatch.pipe_control.dc_flush |
+        gen9_dispatch.pipe_control.render_target_flush |
+        gen9_dispatch.pipe_control.depth_flush));
+    try self.emitSlice(&gen9_dispatch.pipeControl(gen9_dispatch.pipe_control.cs_stall |
+        gen9_dispatch.pipe_control.texture_invalidate |
+        gen9_dispatch.pipe_control.constant_invalidate |
+        gen9_dispatch.pipe_control.state_invalidate |
+        gen9_dispatch.pipe_control.instruction_invalidate));
+    try self.emitSlice(&gen9_dispatch.ccStatePointers);
+    try self.emitSlice(&gen9_dispatch.pipelineSelectGpgpu);
+    try self.emitSlice(&gen9_dispatch.pipeControl(gen9_dispatch.pipe_control.cs_stall |
+        gen9_dispatch.pipe_control.dc_flush |
+        gen9_dispatch.pipe_control.render_target_flush));
+
+    const sba_start = self.batch.items.len * @sizeOf(u32);
+    try self.emitSlice(&gen9_dispatch.stateBaseAddress());
+    inline for (.{
+        .{ 4, kmd.Domain.render },
+        .{ 6, kmd.Domain.render },
+        .{ 10, kmd.Domain.instruction },
+    }) |base_address| {
+        self.relocations.append(self.interface.host_allocator.allocator(), .{
+            .target_handle = state_handle,
+            .offset = sba_start + base_address[0] * @sizeOf(u32),
+            .delta = gen9_dispatch.base_address_delta,
+            .read = true,
+            .domain = base_address[1],
+        }) catch return VkError.OutOfHostMemory;
+    }
+
+    try self.emitSlice(&gen9_dispatch.pipeControl(gen9_dispatch.pipe_control.cs_stall |
+        gen9_dispatch.pipe_control.texture_invalidate |
+        gen9_dispatch.pipe_control.constant_invalidate |
+        gen9_dispatch.pipe_control.state_invalidate |
+        gen9_dispatch.pipe_control.instruction_invalidate));
+    try self.emitSlice(&gen9_dispatch.pipeControl(gen9_dispatch.pipe_control.cs_stall));
+    try self.emitSlice(&gen9_dispatch.mediaVfeState());
+    try self.emitSlice(&gen9_dispatch.interfaceDescriptorLoad(state_layout.interface_descriptor_offset));
+    try self.emitSlice(&gen9_dispatch.gpgpuWalker(.{ 1, 1, 1 }, 1));
+    try self.emitSlice(&gen9_dispatch.mediaStateFlush);
+    try self.emitSlice(&gen9_dispatch.pipeControl(gen9_dispatch.pipe_control.cs_stall |
+        gen9_dispatch.pipe_control.dc_flush));
 }
 
 pub fn setDeviceMask(interface: *Interface, device_mask: u32) VkError!void {
@@ -382,23 +511,30 @@ pub fn endRenderPass(interface: *Interface) VkError!void {
 pub fn executeCommands(interface: *Interface, commands: *Interface) VkError!void {
     const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
     const secondary: *Self = @alignCast(@fieldParentPtr("interface", commands));
+    if (secondary.gpu_allocations.items.len != 0)
+        return VkError.FeatureNotPresent;
+    if (secondary.engine) |engine|
+        try self.requireEngine(engine);
+
     const allocator = self.interface.host_allocator.allocator();
     const relocation_offset = self.batch.items.len * @sizeOf(u32);
-
     self.batch.appendSlice(allocator, secondary.batch.items) catch return VkError.OutOfHostMemory;
     for (secondary.relocations.items) |relocation| {
         self.relocations.append(allocator, .{
+            .source_handle = relocation.source_handle,
             .target_handle = relocation.target_handle,
-            .offset = relocation.offset + relocation_offset,
+            .offset = relocation.offset + if (relocation.source_handle == null) relocation_offset else 0,
             .delta = relocation.delta,
             .read = relocation.read,
             .write = relocation.write,
+            .domain = relocation.domain,
         }) catch return VkError.OutOfHostMemory;
     }
 }
 
 pub fn fillBuffer(interface: *Interface, buffer: *base.Buffer, offset: vk.DeviceSize, size: vk.DeviceSize, data: u32) VkError!void {
     const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
+    try self.requireEngine(.blitter);
     const dst_range = try copy.fillRange(buffer, offset, size);
 
     var filled: vk.DeviceSize = 0;

src/intel/FlintPipeline.zig

@@ -49,7 +49,7 @@ pub fn createCompute(device: *base.Device, allocator: std.mem.Allocator, cache:
     };
     initialized = true;
 
-    self.stages = try compileStages(self.artifact_allocator.allocator(), &.{info.stage}, .compute, compilerDeviceInfo(device));
+    self.stages = try compileStages(self.artifact_allocator.allocator(), device.io(), &.{info.stage}, .compute, compilerDeviceInfo(device));
     if (self.computeArtifact()) |artifact|
         try validateComputePipelineLayout(self.interface.layout, &artifact.resources);
     return self;
@@ -74,11 +74,11 @@ pub fn createGraphics(device: *base.Device, allocator: std.mem.Allocator, cache:
         stages[0..info.stage_count]
     else
         return VkError.ValidationFailed;
-    self.stages = try compileStages(self.artifact_allocator.allocator(), stage_infos, .graphics, compilerDeviceInfo(device));
+    self.stages = try compileStages(self.artifact_allocator.allocator(), device.io(), stage_infos, .graphics, compilerDeviceInfo(device));
     return self;
 }
 
-fn compileStages(allocator: std.mem.Allocator, infos: []const vk.PipelineShaderStageCreateInfo, pipeline_kind: PipelineKind, device_info: ?compiler.device.DeviceInfo) VkError![]CommonStage {
+fn compileStages(allocator: std.mem.Allocator, io: std.Io, infos: []const vk.PipelineShaderStageCreateInfo, pipeline_kind: PipelineKind, device_info: ?compiler.device.DeviceInfo) VkError![]CommonStage {
     if (infos.len == 0)
         return VkError.ValidationFailed;
 
@@ -91,13 +91,13 @@ fn compileStages(allocator: std.mem.Allocator, infos: []const vk.PipelineShaderS
     }
 
     for (infos, stages) |*info, *stage| {
-        stage.* = try compileStage(allocator, info, pipeline_kind, device_info);
+        stage.* = try compileStage(allocator, io, info, pipeline_kind, device_info);
         initialized += 1;
     }
     return stages;
 }
 
-fn compileStage(allocator: std.mem.Allocator, info: *const vk.PipelineShaderStageCreateInfo, pipeline_kind: PipelineKind, device_info: ?compiler.device.DeviceInfo) VkError!CommonStage {
+fn compileStage(allocator: std.mem.Allocator, io: std.Io, info: *const vk.PipelineShaderStageCreateInfo, pipeline_kind: PipelineKind, device_info: ?compiler.device.DeviceInfo) VkError!CommonStage {
     const specializations = try specializationValues(allocator, info.p_specialization_info);
     defer if (specializations.len != 0) allocator.free(specializations);
 
@@ -122,9 +122,15 @@ fn compileStage(allocator: std.mem.Allocator, info: *const vk.PipelineShaderStag
     errdefer module.deinit();
 
     std.debug.assert(module.stage == expected_stage);
+    if (base.config.flint_dump_common_ir)
+        dumpCommonIr(allocator, io, std.mem.span(info.p_name), &module);
 
     var artifact = try lowerToFlint(allocator, &module, device_info);
     errdefer if (artifact) |*value| value.deinit(allocator);
+    if (base.config.flint_dump_ir) {
+        if (artifact) |*value|
+            dumpFlintIr(allocator, io, std.mem.span(info.p_name), &value.program);
+    }
 
     return .{
         .stage = expected_stage,
@@ -133,6 +139,34 @@ fn compileStage(allocator: std.mem.Allocator, info: *const vk.PipelineShaderStag
     };
 }
 
+fn dumpCommonIr(allocator: std.mem.Allocator, io: std.Io, entry_point: []const u8, module: *const base.ShaderModule.IrModule) void {
+    const text = shader_ir.ir.printer.allocPrint(allocator, module) catch |err| {
+        std.log.scoped(.FlintPipeline).err("could not print backend-agnostic IR: {s}", .{@errorName(err)});
+        return;
+    };
+    defer allocator.free(text);
+
+    var stdout_buffer: [1024]u8 = undefined;
+    var stdout_file_writer: std.Io.File.Writer = .init(.stdout(), io, &stdout_buffer);
+    const stdout_writer = &stdout_file_writer.interface;
+    stdout_writer.print("\n=== backend-agnostic IR: {s} ===\n{s}\n", .{ entry_point, text }) catch @panic("Debug printing failed");
+    stdout_writer.flush() catch @panic("Debug printing failed");
+}
+
+fn dumpFlintIr(allocator: std.mem.Allocator, io: std.Io, entry_point: []const u8, program: *const compiler.program.Program) void {
+    const text = compiler.printer.allocPrint(allocator, program) catch |err| {
+        std.log.scoped(.FlintPipeline).err("could not print Flint IR: {s}", .{@errorName(err)});
+        return;
+    };
+    defer allocator.free(text);
+
+    var stdout_buffer: [1024]u8 = undefined;
+    var stdout_file_writer: std.Io.File.Writer = .init(.stdout(), io, &stdout_buffer);
+    const stdout_writer = &stdout_file_writer.interface;
+    stdout_writer.print("\n=== Flint IR: {s} ===\n{s}\n", .{ entry_point, text }) catch @panic("Debug printing failed");
+    stdout_writer.flush() catch @panic("Debug printing failed");
+}
+
 fn lowerToFlint(allocator: std.mem.Allocator, module: *base.ShaderModule.IrModule, device_info: ?compiler.device.DeviceInfo) VkError!?ComputeArtifact {
     const target = device_info orelse return null;
     return compiler.targets.compileCompute(allocator, module, target, .{}) catch |err| switch (err) {
@@ -279,6 +313,9 @@ test "Flint pipeline: lower common compute IR" {
     const program = &artifact.program;
 
     try std.testing.expect(program.properties.common_ir_lowered);
+    try std.testing.expect(program.properties.compute_abi_lowered);
+    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.resources_lowered);
@@ -295,6 +332,6 @@ 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, "load_global_invocation_id r1:u32, component(0)") != null);
     try std.testing.expect(std.mem.indexOf(u8, text, "surface_message write bti(0)") != null);
 }

src/intel/FlintQueue.zig

@@ -97,6 +97,7 @@ pub fn submit(interface: *Interface, infos: []Interface.SubmitInfo, fence: ?*bas
                 try device.kmd.submitBatch(
                     interface.owner.io(),
                     allocator,
+                    .blitter,
                     &.{},
                     &.{},
                     syncs.items,
@@ -129,6 +130,7 @@ pub fn submit(interface: *Interface, infos: []Interface.SubmitInfo, fence: ?*bas
         try device.kmd.submitBatch(
             interface.owner.io(),
             allocator,
+            .blitter,
             &.{},
             &.{},
             syncs[0..sync_count],

src/intel/compiler/device.zig

@@ -47,18 +47,39 @@ pub const DeviceInfo = struct {
         const platform: Platform = switch (pci_device_id & 0xff00) {
             0x1900 => .skylake,
             0x5900 => .kabylake,
+
             0x3e00 => switch (pci_device_id) {
-                0x3ea0, 0x3ea1, 0x3ea2, 0x3ea3, 0x3ea4 => .whiskey_lake,
+                0x3ea0,
+                0x3ea1,
+                0x3ea2,
+                0x3ea3,
+                0x3ea4,
+                => .whiskey_lake,
+
                 else => .coffee_lake,
             },
+
             0x9b00 => .comet_lake,
             0x8a00 => .ice_lake,
             0x4500 => .elkhart_lake,
             0x4e00 => .jasper_lake,
+
             else => switch (pci_device_id) {
-                0x0a84, 0x1a84, 0x1a85, 0x5a84, 0x5a85 => .broxton,
-                0x3184, 0x3185 => .gemini_lake,
-                0x87c0, 0x87ca => .kabylake,
+                0x0a84,
+                0x1a84,
+                0x1a85,
+                0x5a84,
+                0x5a85,
+                => .broxton,
+
+                0x3184,
+                0x3185,
+                => .gemini_lake,
+
+                0x87c0,
+                0x87ca,
+                => .kabylake,
+
                 else => return null,
             },
         };
@@ -72,7 +93,11 @@ pub const DeviceInfo = struct {
             .whiskey_lake,
             .comet_lake,
             => .gen9,
-            .ice_lake, .elkhart_lake, .jasper_lake => .gen11,
+
+            .ice_lake,
+            .elkhart_lake,
+            .jasper_lake,
+            => .gen11,
         };
 
         return .{

src/intel/compiler/ir/program.zig

@@ -11,6 +11,7 @@ pub const Properties = packed struct {
     block_parameters_lowered: bool = false,
     parallel_copies_lowered: bool = false,
 
+    compute_abi_lowered: bool = false,
     system_values_lowered: bool = false,
     resources_lowered: bool = false,
     messages_lowered: bool = false,
@@ -25,7 +26,7 @@ pub const Properties = packed struct {
     flags_allocated: bool = false,
     branches_resolved: bool = false,
 
-    _padding: u17 = 0,
+    _padding: u16 = 0,
 };
 
 pub const StorageBuffer = struct {

src/intel/compiler/lower/block_arguments.zig

@@ -13,7 +13,7 @@ pub const Error = std.mem.Allocator.Error || error{
 };
 
 pub fn run(allocator: std.mem.Allocator, program: *program_ir.Program) Error!void {
-    validator.validate(program) catch return error.InvalidProgram;
+    validator.validate(program) catch return Error.InvalidProgram;
     if (program.properties.block_parameters_lowered)
         return;
 
@@ -28,8 +28,8 @@ pub fn run(allocator: std.mem.Allocator, program: *program_ir.Program) Error!voi
 
     var emitted_parallel_copy = false;
     for (original_blocks.items) |block_id| {
-        const block = program.blocks.get(block_id) orelse return error.InvalidProgram;
-        const terminator = block.terminator orelse return error.InvalidProgram;
+        const block = program.blocks.get(block_id) orelse return Error.InvalidProgram;
+        const terminator = block.terminator orelse return Error.InvalidProgram;
         const rewritten: instruction.Terminator = switch (terminator) {
             .jump => |edge| .{ .jump = try rewriteEdge(
                 allocator,
@@ -74,9 +74,9 @@ fn rewriteEdge(
     if (edge.arguments.len == 0)
         return .{ .target = edge.target, .arguments = &.{} };
 
-    const target = builder.program.blocks.get(edge.target) orelse return error.InvalidProgram;
+    const target = builder.program.blocks.get(edge.target) orelse return Error.InvalidProgram;
     if (target.parameters.items.len != edge.arguments.len)
-        return error.InvalidProgram;
+        return Error.InvalidProgram;
 
     var register_copies: std.ArrayList(pseudo.RegisterCopy) = .empty;
     defer register_copies.deinit(allocator);
@@ -88,10 +88,10 @@ fn rewriteEdge(
             .register => |destination_id| {
                 const source = switch (argument) {
                     .source => |value| value,
-                    .predicate => return error.InvalidProgram,
+                    .predicate => return Error.InvalidProgram,
                 };
                 const destination = builder.program.virtual_registers.get(destination_id) orelse
-                    return error.InvalidProgram;
+                    return Error.InvalidProgram;
                 try register_copies.append(allocator, .{
                     .destination = .{
                         .register = .{ .virtual = destination_id },
@@ -102,7 +102,7 @@ fn rewriteEdge(
             },
             .flag => |destination_id| {
                 const source = switch (argument) {
-                    .source => return error.InvalidProgram,
+                    .source => return Error.InvalidProgram,
                     .predicate => |value| value,
                 };
                 try flag_copies.append(allocator, .{
@@ -135,8 +135,8 @@ fn executionSize(dispatch_width: device.DispatchWidth) device.ExecutionSize {
 
 fn mapBuilderError(err: anyerror) Error {
     return switch (err) {
-        error.OutOfMemory => error.OutOfMemory,
-        else => error.InvalidProgram,
+        Error.OutOfMemory => Error.OutOfMemory,
+        else => Error.InvalidProgram,
     };
 }
 

src/intel/compiler/lower/common_ir.zig

@@ -870,7 +870,7 @@ pub const Lowerer = struct {
 
         var transformer_context: shader_ir.transformer_manager.Context = .{ .allocator = allocator };
         _ = transformer_manager.run(self.module, &transformer_context) catch |err| return switch (err) {
-            error.OutOfMemory => Error.OutOfMemory,
+            Error.OutOfMemory => Error.OutOfMemory,
             else => Error.SanitizationFailed,
         };
         if (!self.module.properties.no_function_calls)
@@ -997,7 +997,7 @@ fn expectLoweringError(source: []const u8, expected: Error) !void {
         return;
     };
     defer program.deinit();
-    return error.TestExpectedError;
+    return Error.TestExpectedError;
 }
 
 test "[ir] Lower: basic shader" {

src/intel/compiler/lower/parallel_copies.zig

@@ -28,7 +28,7 @@ const FlagWrite = struct {
 };
 
 pub fn run(allocator: std.mem.Allocator, program: *program_ir.Program) Error!void {
-    validator.validate(program) catch return error.InvalidProgram;
+    validator.validate(program) catch return Error.InvalidProgram;
     if (program.properties.parallel_copies_lowered)
         return;
 
@@ -39,12 +39,12 @@ pub fn run(allocator: std.mem.Allocator, program: *program_ir.Program) Error!voi
         var instruction_index: usize = 0;
 
         while (true) {
-            const block = program.blocks.get(block_id) orelse return error.InvalidProgram;
+            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 inst = program.instructions.get(instruction_id) orelse return Error.InvalidProgram;
             const parallel_copy = switch (inst.operation) {
                 .parallel_copy => |copy| copy,
                 else => {
@@ -53,7 +53,7 @@ pub fn run(allocator: std.mem.Allocator, program: *program_ir.Program) Error!voi
                 },
             };
             if (inst.predicate != null)
-                return error.InvalidProgram;
+                return Error.InvalidProgram;
             const execution_size = inst.execution_size;
 
             var emitted: std.ArrayList(EmittedInstruction) = .empty;
@@ -61,16 +61,16 @@ pub fn run(allocator: std.mem.Allocator, program: *program_ir.Program) Error!voi
             try lowerParallelCopy(allocator, &builder, execution_size, parallel_copy, &emitted);
 
             if (emitted.items.len == 0) {
-                const mutable_block = program.blocks.getMut(block_id) orelse return error.InvalidProgram;
+                const mutable_block = program.blocks.getMut(block_id) orelse return Error.InvalidProgram;
                 const removed_id = mutable_block.instructions.orderedRemove(instruction_index);
                 if (removed_id != instruction_id or !program.instructions.remove(instruction_id))
-                    return error.InvalidProgram;
+                    return Error.InvalidProgram;
                 continue;
             }
 
             builder.replaceOperation(instruction_id, emitted.items[0].operation) catch |err|
                 return mapBuilderError(err);
-            const replacement = program.instructions.getMut(instruction_id) orelse return error.InvalidProgram;
+            const replacement = program.instructions.getMut(instruction_id) orelse return Error.InvalidProgram;
             replacement.predicate = emitted.items[0].predicate;
 
             for (emitted.items[1..], 1..) |item, offset| {
@@ -87,7 +87,7 @@ pub fn run(allocator: std.mem.Allocator, program: *program_ir.Program) Error!voi
     }
 
     program.properties.parallel_copies_lowered = true;
-    validator.validate(program) catch return error.InvalidProgram;
+    validator.validate(program) catch return Error.InvalidProgram;
 }
 
 fn lowerParallelCopy(
@@ -130,9 +130,9 @@ fn scheduleRegisterCopies(
 
         const cycle_copy = &pending.items[0];
         const destination_id = destinationVirtualRegister(cycle_copy.destination) orelse
-            return error.InvalidProgram;
+            return Error.InvalidProgram;
         const destination_register = builder.program.virtual_registers.get(destination_id) orelse
-            return error.InvalidProgram;
+            return Error.InvalidProgram;
         const temporary = builder.addVirtualRegister(.{
             .size_bytes = destination_register.size_bytes,
             .alignment_bytes = destination_register.alignment_bytes,
@@ -295,8 +295,8 @@ fn immediateU32(value: u32) operand.Source {
 
 fn mapBuilderError(err: anyerror) Error {
     return switch (err) {
-        error.OutOfMemory => error.OutOfMemory,
-        else => error.InvalidProgram,
+        Error.OutOfMemory => Error.OutOfMemory,
+        else => Error.InvalidProgram,
     };
 }
 

src/intel/compiler/targets/gen9/compute/abi.zig

@@ -0,0 +1,60 @@
+const operand = @import("../../../ir/operand.zig");
+const program_ir = @import("../../../ir/program.zig");
+
+pub const Error = error{
+    InvalidPayloadLayout,
+};
+
+const thread_header: operand.PhysicalGrf = .{
+    .number = 0,
+    .byte_offset = 0,
+};
+
+pub fn run(program: *program_ir.Program) Error!void {
+    if (program.properties.compute_abi_lowered)
+        return;
+
+    if (program.payload.header_grf) |header| {
+        if (header.number != thread_header.number or header.byte_offset != thread_header.byte_offset)
+            return Error.InvalidPayloadLayout;
+    }
+    if (program.program_data.payload_grf_count > 1)
+        return Error.InvalidPayloadLayout;
+
+    program.payload.header_grf = thread_header;
+    program.program_data.payload_grf_count = 1;
+    program.properties.compute_abi_lowered = true;
+}
+
+const std = @import("std");
+const device = @import("../../../device.zig");
+
+const test_device: device.DeviceInfo = .{
+    .generation = .gen9,
+    .platform = .skylake,
+    .pci_device_id = 0x1912,
+    .grf_count = 128,
+};
+
+test "[gen9] compute ABI: reserve thread header" {
+    var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, test_device, .simd8);
+    defer program.deinit();
+
+    try run(&program);
+
+    try std.testing.expectEqual(thread_header, program.payload.header_grf.?);
+    try std.testing.expectEqual(@as(u16, 1), program.program_data.payload_grf_count);
+    try std.testing.expect(program.properties.compute_abi_lowered);
+
+    try run(&program);
+    try std.testing.expectEqual(@as(u16, 1), program.program_data.payload_grf_count);
+}
+
+test "[gen9] compute ABI: reject conflicting payload" {
+    var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, test_device, .simd8);
+    defer program.deinit();
+    program.payload.header_grf = .{ .number = 1 };
+
+    try std.testing.expectError(Error.InvalidPayloadLayout, run(&program));
+    try std.testing.expect(!program.properties.compute_abi_lowered);
+}

src/intel/compiler/targets/gen9/compute/compute.zig

@@ -1,6 +1,11 @@
 const std = @import("std");
 
+pub const abi = @import("abi.zig");
+pub const dispatch = @import("dispatch.zig");
+pub const eu_encoder = @import("eu_encoder.zig");
+pub const kernel_encoder = @import("kernel_encoder.zig");
 pub const message_addresses = @import("message_addresses.zig");
+pub const message_descriptor = @import("message_descriptor.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");

src/intel/compiler/targets/gen9/compute/dispatch.zig

@@ -0,0 +1,166 @@
+const std = @import("std");
+
+pub const max_surfaces: usize = 4;
+pub const page_size: usize = 4096;
+pub const surface_state_size: usize = 64;
+pub const interface_descriptor_size: usize = 32;
+
+const mocs: u32 = 0x78;
+pub const base_address_delta: u32 = 1 | (mocs << 4);
+const raw_surface_format: u32 = 0x1ff;
+
+pub const Error = error{
+    EmptyBuffer,
+    StateTooLarge,
+    UnsupportedBufferSize,
+    TooManySurfaces,
+};
+
+pub const StateLayout = struct {
+    size: usize,
+    kernel_offset: u32,
+    surface_offsets: [max_surfaces]u32,
+    surface_address_offsets: [max_surfaces]u32,
+    surface_count: u8,
+    binding_table_offset: u32,
+    interface_descriptor_offset: u32,
+};
+
+pub fn writeState(destination: []u8, kernel: []const u8, buffer_sizes: []const u64) Error!StateLayout {
+    if (buffer_sizes.len > max_surfaces)
+        return Error.TooManySurfaces;
+
+    var layout: StateLayout = .{
+        .size = 0,
+        .kernel_offset = 0,
+        .surface_offsets = @splat(0),
+        .surface_address_offsets = @splat(0),
+        .surface_count = @intCast(buffer_sizes.len),
+        .binding_table_offset = 0,
+        .interface_descriptor_offset = 0,
+    };
+
+    var cursor = alignForward(kernel.len, 64);
+    for (buffer_sizes, 0..) |size, index| {
+        cursor = alignForward(cursor, surface_state_size);
+        layout.surface_offsets[index] = @intCast(cursor);
+        layout.surface_address_offsets[index] = @intCast(cursor + 8 * @sizeOf(u32));
+        cursor += surface_state_size;
+        if (size == 0)
+            return Error.EmptyBuffer;
+    }
+
+    cursor = alignForward(cursor, 32);
+    layout.binding_table_offset = @intCast(cursor);
+    cursor += buffer_sizes.len * @sizeOf(u32);
+
+    cursor = alignForward(cursor, 64);
+    layout.interface_descriptor_offset = @intCast(cursor);
+    cursor += interface_descriptor_size;
+    layout.size = alignForward(cursor, page_size);
+    if (layout.size > destination.len or layout.size > page_size)
+        return Error.StateTooLarge;
+
+    @memset(destination[0..layout.size], 0);
+    @memcpy(destination[layout.kernel_offset .. layout.kernel_offset + kernel.len], kernel);
+
+    for (buffer_sizes, 0..) |size, index| {
+        _ = try encodeRawBufferSurface(destination, layout.surface_offsets[index], size);
+        putU32(destination, layout.binding_table_offset + @as(u32, @intCast(index * @sizeOf(u32))), layout.surface_offsets[index]);
+    }
+
+    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 + 6 * @sizeOf(u32), 1);
+
+    return layout;
+}
+
+fn encodeRawBufferSurface(destination: []u8, offset: u32, byte_size: u64) Error!void {
+    if (byte_size == 0)
+        return Error.EmptyBuffer;
+
+    const aligned_size = std.mem.alignForward(u64, byte_size, 4);
+    const padded_size = aligned_size + (aligned_size - byte_size);
+    if (padded_size == 0 or padded_size > (@as(u64, 1) << 32))
+        return Error.UnsupportedBufferSize;
+    const length_minus_one: u32 = @intCast(padded_size - 1);
+
+    putU32(destination, offset + 0, (4 << 29) |
+        (raw_surface_format << 18) |
+        (1 << 16) |
+        (1 << 14));
+    putU32(destination, offset + 1 * @sizeOf(u32), mocs << 24);
+    putU32(destination, offset + 2 * @sizeOf(u32), (length_minus_one & 0x7f) |
+        (((length_minus_one >> 7) & 0x3fff) << 16));
+    putU32(destination, offset + 3 * @sizeOf(u32), ((length_minus_one >> 21) & 0x7ff) << 21);
+}
+
+pub const ccStatePointers = [_]u32{
+    0x780e0000,
+    0,
+};
+
+pub const pipelineSelectGpgpu = [_]u32{0x69040302};
+
+pub fn pipeControl(bits: u32) [6]u32 {
+    return .{ 0x7a000004, bits, 0, 0, 0, 0 };
+}
+
+pub const pipe_control = struct {
+    pub const state_invalidate: u32 = 1 << 2;
+    pub const constant_invalidate: u32 = 1 << 3;
+    pub const dc_flush: u32 = 1 << 5;
+    pub const texture_invalidate: u32 = 1 << 10;
+    pub const instruction_invalidate: u32 = 1 << 11;
+    pub const render_target_flush: u32 = 1 << 12;
+    pub const depth_flush: u32 = 1 << 0;
+    pub const cs_stall: u32 = 1 << 20;
+};
+
+pub fn stateBaseAddress() [19]u32 {
+    var words: [19]u32 = @splat(0);
+    words[0] = 0x61010011;
+    words[3] = mocs << 16;
+    words[4] = base_address_delta;
+    words[6] = base_address_delta;
+    words[10] = base_address_delta;
+    words[13] = (1 << 12) | 1;
+    words[15] = (1 << 12) | 1;
+    return words;
+}
+
+pub fn mediaVfeState() [9]u32 {
+    var words: [9]u32 = @splat(0);
+    words[0] = 0x70000007;
+    words[3] = (1 << 16) | (2 << 8);
+    words[5] = 2 << 16;
+    return words;
+}
+
+pub fn interfaceDescriptorLoad(offset: u32) [4]u32 {
+    return .{ 0x70020002, 0, interface_descriptor_size, offset };
+}
+
+pub fn gpgpuWalker(group_count: [3]u32, right_mask: u32) [15]u32 {
+    var words: [15]u32 = @splat(0);
+    words[0] = 0x7105000d;
+    words[7] = group_count[0];
+    words[10] = group_count[1];
+    words[12] = group_count[2];
+    words[13] = right_mask;
+    words[14] = 0xffffffff;
+    return words;
+}
+
+pub const mediaStateFlush = [_]u32{ 0x70040000, 0 };
+
+fn alignForward(value: usize, alignment: usize) usize {
+    return std.mem.alignForward(usize, value, alignment);
+}
+
+fn putU32(destination: []u8, offset: u32, value: u32) void {
+    std.mem.writeInt(u32, destination[offset..][0..@sizeOf(u32)], value, .little);
+}

src/intel/compiler/targets/gen9/compute/eu_encoder.zig

@@ -0,0 +1,250 @@
+const std = @import("std");
+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");
+
+pub const Error = error{
+    UnsupportedExecutionSize,
+    UnsupportedDataType,
+    UnsupportedOperand,
+    InvalidRegister,
+    InvalidRegion,
+};
+
+pub const eot_payload_grf: u8 = 112;
+
+pub const EncodedInstruction = struct {
+    words: [2]u64 = .{ 0, 0 },
+
+    pub fn setBits(self: *EncodedInstruction, high: u7, low: u7, value: u64) void {
+        const width = @as(u8, high) - @as(u8, low) + 1;
+        const word = @as(usize, high) / 64;
+        const word_low: u6 = @intCast(@as(u8, low) % 64);
+        const mask = (@as(u64, std.math.maxInt(u64)) >> @intCast(64 - width)) << word_low;
+        self.words[word] = (self.words[word] & ~mask) | ((value << word_low) & mask);
+    }
+
+    pub fn bits(self: EncodedInstruction, high: u7, low: u7) u64 {
+        const width = @as(u8, high) - @as(u8, low) + 1;
+        const word = @as(usize, high) / 64;
+        const word_low: u6 = @intCast(@as(u8, low) % 64);
+        return (self.words[word] >> word_low) & (@as(u64, std.math.maxInt(u64)) >> @intCast(64 - width));
+    }
+};
+
+const RegisterFile = enum(u2) {
+    architecture = 0,
+    grf = 1,
+    immediate = 3,
+};
+
+const HardwareType = enum(u4) {
+    unsigned_dword = 0,
+    signed_dword = 1,
+    unsigned_word = 2,
+    float = 7,
+};
+
+const Grf = struct {
+    number: u8,
+    byte_offset: u5,
+};
+
+pub fn encodeMove(execution_size: device.ExecutionSize, move: ir_instruction.Move) Error!EncodedInstruction {
+    var encoded = try instructionHeader(1, execution_size);
+    const destination = try resolveGrf(move.destination.register, move.destination.region.byte_offset);
+
+    setDestination(
+        &encoded,
+        .grf,
+        try hardwareType(move.destination.type),
+        destination,
+        try horizontalStride(move.destination.region.horizontal_stride),
+    );
+
+    switch (move.source.register) {
+        .physical_grf => {
+            const source = try resolveGrf(move.source.register, move.source.region.byte_offset);
+            try setSource0Register(&encoded, move.source, source);
+        },
+        .immediate => |immediate| {
+            if (move.source.negate or move.source.absolute)
+                return Error.UnsupportedOperand;
+            setSource0Immediate(&encoded, try hardwareType(move.source.type), immediate);
+        },
+        else => return Error.UnsupportedOperand,
+    }
+
+    return encoded;
+}
+
+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);
+    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));
+    copy.setBits(46, 43, @intFromEnum(HardwareType.unsigned_dword));
+    copy.setBits(76, 69, header.number);
+    copy.setBits(81, 80, 1);
+    copy.setBits(84, 82, 3);
+    copy.setBits(88, 85, 4);
+
+    var send = try instructionHeader(49, .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));
+    send.setBits(46, 43, @intFromEnum(HardwareType.unsigned_word));
+    send.setBits(76, 69, eot_payload_grf);
+    send.setBits(81, 80, 1);
+    send.setBits(84, 82, 3);
+    send.setBits(88, 85, 4);
+    send.setBits(90, 89, @intFromEnum(RegisterFile.immediate));
+    send.setBits(94, 91, @intFromEnum(HardwareType.unsigned_dword));
+    send.setBits(124, 96, 0x02000010); // mlen=1, no response, do not dereference URB
+    send.setBits(27, 24, 7); // Thread Spawner
+    send.setBits(127, 127, 1);
+
+    return .{ copy, send };
+}
+
+pub fn encodeSurfaceMessage(execution_size: device.ExecutionSize, message: ir_instruction.SurfaceMessage) Error!EncodedInstruction {
+    var encoded = try instructionHeader(49, execution_size);
+    const descriptor = message_descriptor.encode(message);
+    const payload = try resolveGrf(message.payload.base, 0);
+    if (payload.byte_offset != 0)
+        return Error.InvalidRegister;
+
+    if (message.response) |response| {
+        const destination = try resolveGrf(response.base, 0);
+        if (destination.byte_offset != 0)
+            return Error.InvalidRegister;
+        setDestination(&encoded, .grf, .unsigned_word, destination, 1);
+    } else {
+        setDestination(&encoded, .architecture, .unsigned_word, .{ .number = 0, .byte_offset = 0 }, 1);
+    }
+
+    encoded.setBits(42, 41, @intFromEnum(RegisterFile.grf));
+    encoded.setBits(46, 43, @intFromEnum(HardwareType.unsigned_dword));
+    encoded.setBits(76, 69, payload.number);
+    encoded.setBits(68, 64, payload.byte_offset);
+    encoded.setBits(81, 80, 1); // horizontal stride 1
+    encoded.setBits(84, 82, 3); // width 8
+    encoded.setBits(88, 85, 4); // vertical stride 8
+    encoded.setBits(90, 89, @intFromEnum(RegisterFile.immediate));
+    encoded.setBits(94, 91, @intFromEnum(HardwareType.unsigned_dword));
+    encoded.setBits(124, 96, descriptor.value);
+    encoded.setBits(27, 24, descriptor.sfid);
+    return encoded;
+}
+
+fn instructionHeader(opcode: u7, execution_size: device.ExecutionSize) Error!EncodedInstruction {
+    var encoded: EncodedInstruction = .{};
+    encoded.setBits(6, 0, opcode);
+    encoded.setBits(23, 21, try executionSize(execution_size));
+    return encoded;
+}
+
+fn setDestination(encoded: *EncodedInstruction, file: RegisterFile, data_type: HardwareType, register: Grf, horizontal_stride: u2) void {
+    encoded.setBits(36, 35, @intFromEnum(file));
+    encoded.setBits(40, 37, @intFromEnum(data_type));
+    encoded.setBits(52, 48, register.byte_offset);
+    encoded.setBits(60, 53, register.number);
+    encoded.setBits(62, 61, horizontal_stride);
+}
+
+fn setSource0Register(encoded: *EncodedInstruction, source: operand.Source, register: Grf) Error!void {
+    encoded.setBits(42, 41, @intFromEnum(RegisterFile.grf));
+    encoded.setBits(46, 43, @intFromEnum(try hardwareType(source.type)));
+    encoded.setBits(68, 64, register.byte_offset);
+    encoded.setBits(76, 69, register.number);
+    encoded.setBits(77, 77, @intFromBool(source.absolute));
+    encoded.setBits(78, 78, @intFromBool(source.negate));
+    encoded.setBits(81, 80, try horizontalStride(source.region.horizontal_stride));
+    encoded.setBits(84, 82, try regionWidth(source.region.width));
+    encoded.setBits(88, 85, try verticalStride(source.region.vertical_stride));
+}
+
+fn setSource0Immediate(encoded: *EncodedInstruction, data_type: HardwareType, immediate: operand.Immediate) void {
+    encoded.setBits(42, 41, @intFromEnum(RegisterFile.immediate));
+    encoded.setBits(46, 43, @intFromEnum(data_type));
+    encoded.setBits(90, 89, @intFromEnum(RegisterFile.architecture));
+    encoded.setBits(94, 91, @intFromEnum(data_type));
+    encoded.setBits(127, 96, switch (immediate) {
+        .u32 => |value| value,
+        .i32 => |value| @as(u32, @bitCast(value)),
+        .f32 => |value| @as(u32, @bitCast(value)),
+    });
+}
+
+fn resolveGrf(register: operand.RegisterRef, region_byte_offset: u16) Error!Grf {
+    const physical = switch (register) {
+        .physical_grf => |value| value,
+        else => return Error.UnsupportedOperand,
+    };
+
+    const byte_address = @as(u32, physical.number) * 32 + physical.byte_offset + region_byte_offset;
+    const number = byte_address / 32;
+
+    if (number >= 128)
+        return Error.InvalidRegister;
+
+    return .{
+        .number = @intCast(number),
+        .byte_offset = @intCast(byte_address % 32),
+    };
+}
+
+fn hardwareType(data_type: operand.DataType) Error!HardwareType {
+    return switch (data_type) {
+        .u32 => .unsigned_dword,
+        .i32 => .signed_dword,
+        .f32 => .float,
+        else => Error.UnsupportedDataType,
+    };
+}
+
+fn executionSize(size: device.ExecutionSize) Error!u3 {
+    return switch (size) {
+        .simd1 => 0,
+        .simd8 => 3,
+        else => Error.UnsupportedExecutionSize,
+    };
+}
+
+fn horizontalStride(stride: u8) Error!u2 {
+    return switch (stride) {
+        0 => 0,
+        1 => 1,
+        2 => 2,
+        4 => 3,
+        else => Error.InvalidRegion,
+    };
+}
+
+fn regionWidth(width: u8) Error!u3 {
+    return switch (width) {
+        1 => 0,
+        2 => 1,
+        4 => 2,
+        8 => 3,
+        16 => 4,
+        else => Error.InvalidRegion,
+    };
+}
+
+fn verticalStride(stride: u8) Error!u4 {
+    return switch (stride) {
+        0 => 0,
+        1 => 1,
+        2 => 2,
+        4 => 3,
+        8 => 4,
+        16 => 5,
+        32 => 6,
+        else => Error.InvalidRegion,
+    };
+}

src/intel/compiler/targets/gen9/compute/kernel_encoder.zig

@@ -0,0 +1,67 @@
+const std = @import("std");
+
+const eu = @import("eu_encoder.zig");
+const program_ir = @import("../../../ir/program.zig");
+
+pub const Error = std.mem.Allocator.Error || eu.Error || error{
+    InvalidProgram,
+    UnsupportedControlFlow,
+    UnsupportedOperation,
+    UnsupportedPredication,
+    EotRegisterUnavailable,
+};
+
+pub fn encode(allocator: std.mem.Allocator, program: *program_ir.Program) Error![]u8 {
+    if (!program.properties.registers_allocated)
+        return Error.InvalidProgram;
+    if (program.program_data.total_grf_count > eu.eot_payload_grf)
+        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)
+        return Error.UnsupportedControlFlow;
+
+    var kernel: std.ArrayList(u8) = .empty;
+    errdefer kernel.deinit(allocator);
+
+    for (entry.instructions.items) |instruction_id| {
+        const instruction = program.instructions.get(instruction_id) orelse return Error.InvalidProgram;
+        if (instruction.predicate != null)
+            return Error.UnsupportedPredication;
+
+        const encoded = switch (instruction.operation) {
+            .move => |move| try eu.encodeMove(instruction.execution_size, move),
+            .surface_message => |message| try eu.encodeSurfaceMessage(instruction.execution_size, message),
+            else => return Error.UnsupportedOperation,
+        };
+        try appendInstruction(allocator, &kernel, encoded);
+    }
+
+    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 => return Error.UnsupportedControlFlow,
+    }
+
+    return kernel.toOwnedSlice(allocator);
+}
+
+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);
+}

src/intel/compiler/targets/gen9/compute/message_addresses.zig

@@ -18,7 +18,7 @@ const AddressAdjustment = struct {
 
 pub fn run(program: *program_ir.Program) Error!void {
     if (!program.properties.messages_lowered)
-        return error.MessagesNotLowered;
+        return Error.MessagesNotLowered;
     if (program.properties.message_addresses_lowered)
         return;
 
@@ -29,18 +29,18 @@ pub fn run(program: *program_ir.Program) Error!void {
         var instruction_index: usize = 0;
 
         while (true) {
-            const block = program.blocks.get(block_id) orelse return error.InvalidProgram;
+            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 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;
+                return Error.InvalidProgram;
 
             if (adjustment.immediate_offset == 0) {
                 instruction_index += 1;
@@ -51,15 +51,18 @@ pub fn run(program: *program_ir.Program) Error!void {
                 .immediate => |immediate| {
                     const base = switch (immediate) {
                         .u32 => |value| value,
-                        else => return error.InvalidProgram,
+                        else => return Error.InvalidProgram,
                     };
-                    const mutable = program.instructions.getMut(instruction_id) orelse return error.InvalidProgram;
-                    const address = messageAddressMut(&mutable.operation) orelse 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 => {
+                .virtual,
+                .physical_grf,
+                .architecture,
+                => {
                     const execution_width: u32 = @intFromEnum(inst.execution_size);
                     const size_bytes = execution_width * @sizeOf(u32);
                     const address_register = builder.addVirtualRegister(.{
@@ -82,8 +85,8 @@ pub fn run(program: *program_ir.Program) Error!void {
                         },
                     }) 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;
+                    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,
@@ -92,7 +95,7 @@ pub fn run(program: *program_ir.Program) Error!void {
                     address.immediate_offset.* = 0;
                     instruction_index += 2;
                 },
-                .null => return error.InvalidProgram,
+                .null => return Error.InvalidProgram,
             }
         }
     }
@@ -131,8 +134,8 @@ fn immediateSource(value: u32) operand.Source {
 
 fn mapBuilderError(err: Builder.Error) Error {
     return switch (err) {
-        error.OutOfMemory => error.OutOfMemory,
-        else => error.InvalidProgram,
+        error.OutOfMemory => Error.OutOfMemory,
+        else => Error.InvalidProgram,
     };
 }
 

src/intel/compiler/targets/gen9/compute/message_descriptor.zig

@@ -0,0 +1,90 @@
+const instruction = @import("../../../ir/instruction.zig");
+
+pub const Descriptor = struct {
+    sfid: u8,
+    value: u32,
+    message_length: u8,
+    response_length: u8,
+};
+
+const dc1_sfid: u8 = 12;
+const simd8_one_channel_control: u8 = 0x2e;
+
+const MessageType = enum(u8) {
+    untyped_surface_read = 1,
+    untyped_surface_write = 9,
+};
+
+pub fn encode(message: instruction.SurfaceMessage) Descriptor {
+    const lengths: struct { message: u8, response: u8 } = switch (message.kind) {
+        .read => .{ .message = 1, .response = 1 },
+        .write => .{ .message = 2, .response = 0 },
+    };
+    const message_type: MessageType = switch (message.kind) {
+        .read => .untyped_surface_read,
+        .write => .untyped_surface_write,
+    };
+
+    return .{
+        .sfid = dc1_sfid,
+        .value = makeDescriptor(
+            message.binding_table,
+            simd8_one_channel_control,
+            message_type,
+            lengths.message,
+            lengths.response,
+        ),
+        .message_length = lengths.message,
+        .response_length = lengths.response,
+    };
+}
+
+fn makeDescriptor(binding_table: u8, message_control: u8, message_type: MessageType, message_length: u8, response_length: u8) u32 {
+    return @as(u32, binding_table) |
+        (@as(u32, message_control) << 8) |
+        (@as(u32, @intFromEnum(message_type)) << 14) |
+        (@as(u32, response_length) << 20) |
+        (@as(u32, message_length) << 25);
+}
+
+test "[gen9] message descriptor: encode SIMD8 one-channel surface read" {
+    const std = @import("std");
+    const descriptor = encode(.{
+        .kind = .read,
+        .binding_table = 3,
+        .payload = .{ .base = .{ .physical_grf = .{ .number = 1 } }, .register_count = 1 },
+        .response = .{ .base = .{ .physical_grf = .{ .number = 2 } }, .register_count = 1 },
+        .data_type = .u32,
+    });
+
+    try std.testing.expectEqual(@as(u8, 12), descriptor.sfid);
+    try std.testing.expectEqual(@as(u8, 1), descriptor.message_length);
+    try std.testing.expectEqual(@as(u8, 1), descriptor.response_length);
+    try std.testing.expectEqual(@as(u8, 3), @as(u8, @truncate(descriptor.value)));
+    try std.testing.expectEqual(@as(u8, 0x2e), @as(u8, @truncate(descriptor.value >> 8)) & 0x3f);
+    try std.testing.expectEqual(@as(u8, 1), @as(u8, @truncate(descriptor.value >> 14)) & 0x1f);
+    try std.testing.expectEqual(@as(u8, 1), @as(u8, @truncate(descriptor.value >> 20)) & 0x1f);
+    try std.testing.expectEqual(@as(u8, 1), @as(u8, @truncate(descriptor.value >> 25)) & 0x0f);
+    try std.testing.expectEqual(@as(u32, 0x02106e03), descriptor.value);
+}
+
+test "[gen9] message descriptor: encode SIMD8 one-channel surface write" {
+    const std = @import("std");
+    const descriptor = encode(.{
+        .kind = .write,
+        .binding_table = 7,
+        .payload = .{ .base = .{ .physical_grf = .{ .number = 1 } }, .register_count = 2 },
+        .response = null,
+        .data_type = .u32,
+    });
+
+    try std.testing.expectEqual(@as(u8, 12), descriptor.sfid);
+    try std.testing.expectEqual(@as(u8, 2), descriptor.message_length);
+    try std.testing.expectEqual(@as(u8, 0), descriptor.response_length);
+    try std.testing.expectEqual(@as(u8, 7), @as(u8, @truncate(descriptor.value)));
+    try std.testing.expectEqual(@as(u8, 0x2e), @as(u8, @truncate(descriptor.value >> 8)) & 0x3f);
+    try std.testing.expectEqual(@as(u8, 9), @as(u8, @truncate(descriptor.value >> 14)) & 0x1f);
+    try std.testing.expectEqual(@as(u8, 0), @as(u8, @truncate(descriptor.value >> 20)) & 0x1f);
+    try std.testing.expectEqual(@as(u8, 2), @as(u8, @truncate(descriptor.value >> 25)) & 0x0f);
+    try std.testing.expectEqual(@as(u32, 0x04026e07), descriptor.value);
+}

src/intel/compiler/targets/gen9/compute/message_lowering.zig

@@ -9,7 +9,7 @@ pub const Error = error{
 
 pub fn run(program: *program_ir.Program) Error!void {
     if (!program.properties.resources_lowered)
-        return error.ResourcesNotLowered;
+        return Error.ResourcesNotLowered;
     if (program.properties.messages_lowered)
         return;
 
@@ -18,12 +18,12 @@ pub fn run(program: *program_ir.Program) Error!void {
         inst.operation = switch (inst.operation) {
             .load_buffer => |op| .{ .surface_read = .{
                 .destination = op.destination,
-                .binding_table = bindingTableIndex(op.buffer) orelse return error.InvalidProgram,
+                .binding_table = bindingTableIndex(op.buffer) orelse return Error.InvalidProgram,
                 .address = op.byte_offset,
                 .immediate_offset = op.immediate_offset,
             } },
             .store_buffer => |op| .{ .surface_write = .{
-                .binding_table = bindingTableIndex(op.buffer) orelse return error.InvalidProgram,
+                .binding_table = bindingTableIndex(op.buffer) orelse return Error.InvalidProgram,
                 .address = op.byte_offset,
                 .immediate_offset = op.immediate_offset,
                 .data = op.source,
@@ -106,5 +106,5 @@ test "[gen9] compute message lowering: reject unresolved resources" {
     var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, test_device, .simd8);
     defer program.deinit();
 
-    try std.testing.expectError(error.ResourcesNotLowered, run(&program));
+    try std.testing.expectError(Error.ResourcesNotLowered, run(&program));
 }

src/intel/compiler/targets/gen9/compute/message_payloads.zig

@@ -14,11 +14,11 @@ pub const Error = std.mem.Allocator.Error || error{
 
 pub fn run(program: *program_ir.Program) Error!void {
     if (!program.properties.message_addresses_lowered)
-        return error.MessageAddressesNotLowered;
+        return Error.MessageAddressesNotLowered;
     if (program.properties.message_payloads_lowered)
         return;
     if (program.device_info.grf_size_bytes != 32)
-        return error.InvalidProgram;
+        return Error.InvalidProgram;
 
     var builder = Builder.init(program);
     for (program.blocks.entries.items, 0..) |entry, block_index| {
@@ -27,17 +27,17 @@ pub fn run(program: *program_ir.Program) Error!void {
         var instruction_index: usize = 0;
 
         while (true) {
-            const block = program.blocks.get(block_id) orelse return error.InvalidProgram;
+            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 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;
+                        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 = .{
@@ -45,7 +45,7 @@ pub fn run(program: *program_ir.Program) Error!void {
                         .source = op.address,
                     } }) catch |err| return mapBuilderError(err);
 
-                    const mutable = program.instructions.getMut(instruction_id) orelse return error.InvalidProgram;
+                    const mutable = program.instructions.getMut(instruction_id) orelse return Error.InvalidProgram;
                     mutable.operation = .{ .surface_message = .{
                         .kind = .read,
                         .binding_table = op.binding_table,
@@ -57,7 +57,7 @@ pub fn run(program: *program_ir.Program) Error!void {
                 },
                 .surface_write => |op| {
                     if (op.immediate_offset != 0 or op.address.type != .u32)
-                        return error.InvalidProgram;
+                        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),
@@ -68,7 +68,7 @@ pub fn run(program: *program_ir.Program) Error!void {
                         .source = op.data,
                     } }) catch |err| return mapBuilderError(err);
 
-                    const mutable = program.instructions.getMut(instruction_id) orelse return error.InvalidProgram;
+                    const mutable = program.instructions.getMut(instruction_id) orelse return Error.InvalidProgram;
                     mutable.operation = .{ .surface_message = .{
                         .kind = .write,
                         .binding_table = op.binding_table,
@@ -107,20 +107,20 @@ fn payloadDestination(register: ids.VirtualRegisterId, byte_offset: u16, data_ty
 
 fn responseSpan(destination: operand.Destination) Error!operand.RegisterSpan {
     if (destination.region.byte_offset != 0 or destination.region.horizontal_stride != 1)
-        return error.InvalidProgram;
+        return Error.InvalidProgram;
     return switch (destination.register) {
         .virtual, .physical_grf => .{
             .base = destination.register,
             .register_count = 1,
         },
-        else => error.InvalidProgram,
+        else => Error.InvalidProgram,
     };
 }
 
 fn mapBuilderError(err: Builder.Error) Error {
     return switch (err) {
-        error.OutOfMemory => error.OutOfMemory,
-        else => error.InvalidProgram,
+        error.OutOfMemory => Error.OutOfMemory,
+        else => Error.InvalidProgram,
     };
 }
 

src/intel/compiler/targets/gen9/compute/pipeline.zig

@@ -10,25 +10,42 @@ const flag_allocation = @import("../flag_allocation.zig");
 const register_allocation = @import("../register_allocation.zig");
 
 const compute = @import("compute.zig");
+const abi = @import("abi.zig");
+const kernel_encoder = @import("kernel_encoder.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_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,
-    UnsupportedGrfSize,
-};
+pub const Error = common_ir.Error ||
+    block_arguments.Error ||
+    parallel_copies.Error ||
+    abi.Error ||
+    kernel_encoder.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,
+        UnsupportedGrfSize,
+    };
 
 pub const Artifact = struct {
     program: program_ir.Program,
     resources: resource_layout.Layout,
+    kernel: ?[]u8,
 
     pub fn deinit(self: *Artifact, allocator: std.mem.Allocator) void {
+        if (self.kernel) |kernel|
+            allocator.free(kernel);
         self.resources.deinit(allocator);
         self.program.deinit();
         self.* = undefined;
@@ -55,6 +72,7 @@ pub fn compile(allocator: std.mem.Allocator, module: *shader_ir.module.Module, d
     );
     errdefer program.deinit();
 
+    try abi.run(&program);
     try block_arguments.run(allocator, &program);
     try parallel_copies.run(allocator, &program);
 
@@ -71,8 +89,22 @@ pub fn compile(allocator: std.mem.Allocator, module: *shader_ir.module.Module, d
     try flag_allocation.run(allocator, &program);
     try register_allocation.run(allocator, &program);
 
+    const kernel = kernel_encoder.encode(allocator, &program) catch |err| switch (err) {
+        error.UnsupportedControlFlow,
+        error.UnsupportedOperation,
+        error.UnsupportedPredication,
+        error.UnsupportedExecutionSize,
+        error.UnsupportedDataType,
+        error.UnsupportedOperand,
+        error.EotRegisterUnavailable,
+        => null,
+        else => return err,
+    };
+    errdefer if (kernel) |bytes| allocator.free(bytes);
+
     return .{
         .program = program,
         .resources = resources,
+        .kernel = kernel,
     };
 }

src/intel/compiler/targets/gen9/flag_allocation.zig

@@ -18,15 +18,15 @@ const physical_flag_count = 2;
 
 pub fn run(allocator: std.mem.Allocator, program: *program_ir.Program) Error!void {
     if (!program.properties.block_parameters_lowered)
-        return error.BlockParametersNotLowered;
+        return Error.BlockParametersNotLowered;
 
     if (!program.properties.parallel_copies_lowered)
-        return error.ParallelCopiesNotLowered;
+        return Error.ParallelCopiesNotLowered;
 
     if (program.properties.flags_allocated)
         return;
 
-    validator.validate(program) catch return error.InvalidProgram;
+    validator.validate(program) catch return Error.InvalidProgram;
 
     const allocations = try allocator.alloc(?operand.PhysicalFlag, program.virtual_flags.entries.items.len);
     defer allocator.free(allocations);
@@ -38,9 +38,9 @@ pub fn run(allocator: std.mem.Allocator, program: *program_ir.Program) Error!voi
     for (allocations) |*allocation| {
         const marker = allocation.* orelse continue;
         if (marker.subregister != std.math.maxInt(u8))
-            return error.InvalidProgram;
+            return Error.InvalidProgram;
 
-        const subregister = std.mem.indexOfScalar(bool, &occupied, false) orelse return error.OutOfFlagRegisters;
+        const subregister = std.mem.indexOfScalar(bool, &occupied, false) orelse return Error.OutOfFlagRegisters;
 
         allocation.* = .{
             .register = 0,
@@ -51,7 +51,7 @@ pub fn run(allocator: std.mem.Allocator, program: *program_ir.Program) Error!voi
 
     try visitProgramFlags(program, allocations, &occupied, true);
     program.properties.flags_allocated = true;
-    validator.validate(program) catch return error.InvalidProgram;
+    validator.validate(program) catch return Error.InvalidProgram;
 }
 
 fn visitProgramFlags(
@@ -63,14 +63,14 @@ fn visitProgramFlags(
     for (program.instructions.entries.items, 0..) |entry, instruction_index| {
         _ = entry orelse continue;
         const inst = program.instructions.getMut(ids.InstructionId.fromIndex(instruction_index)) orelse
-            return error.InvalidProgram;
+            return Error.InvalidProgram;
 
         if (inst.predicate) |*predicate|
             try visitFlagRef(program, &predicate.flag, allocations, occupied, rewrite);
 
         switch (inst.operation) {
             .compare => |*compare| try visitFlagRef(program, &compare.destination, allocations, occupied, rewrite),
-            .parallel_copy => return error.ParallelCopiesNotLowered,
+            .parallel_copy => return Error.ParallelCopiesNotLowered,
             else => {},
         }
     }
@@ -78,8 +78,8 @@ fn visitProgramFlags(
     for (program.blocks.entries.items, 0..) |entry, block_index| {
         _ = entry orelse continue;
         const block = program.blocks.getMut(ids.BlockId.fromIndex(block_index)) orelse
-            return error.InvalidProgram;
-        const terminator = if (block.terminator) |*value| value else return error.InvalidProgram;
+            return Error.InvalidProgram;
+        const terminator = if (block.terminator) |*value| value else return Error.InvalidProgram;
 
         switch (terminator.*) {
             .jump => |*edge| try visitEdge(program, edge, allocations, occupied, rewrite),
@@ -129,7 +129,7 @@ fn visitFlagRef(
     switch (flag.*) {
         .virtual => |virtual| {
             if (!program.virtual_flags.isLive(virtual) or virtual.index() >= allocations.len)
-                return error.InvalidProgram;
+                return Error.InvalidProgram;
 
             if (!rewrite) {
                 // Mark this virtual flag as referenced without assigning a physical
@@ -139,14 +139,14 @@ fn visitFlagRef(
                 return;
             }
 
-            const physical = allocations[virtual.index()] orelse return error.InvalidProgram;
+            const physical = allocations[virtual.index()] orelse return Error.InvalidProgram;
             if (physical.subregister >= physical_flag_count)
-                return error.InvalidProgram;
+                return Error.InvalidProgram;
             flag.* = .{ .physical = physical };
         },
         .physical => |physical| {
             if (physical.register != 0 or physical.subregister >= physical_flag_count)
-                return error.InvalidProgram;
+                return Error.InvalidProgram;
             occupied[physical.subregister] = true;
         },
     }
@@ -251,7 +251,7 @@ test "[gen9] flag allocation: report exhaustion without rewriting" {
     try program.setTerminator(entry, .end_thread);
     markPrerequisites(&program);
 
-    try std.testing.expectError(error.OutOfFlagRegisters, run(std.testing.allocator, &program));
+    try std.testing.expectError(Error.OutOfFlagRegisters, run(std.testing.allocator, &program));
     try std.testing.expect(!program.properties.flags_allocated);
     try std.testing.expectEqual(first, program.instructions.get(first_compare).?.operation.compare.destination.virtual);
 }

src/intel/compiler/targets/gen9/gen9.zig

@@ -99,6 +99,9 @@ test "[gen9] target: lower 256 KiB SSBO copy loop" {
     const resources = &artifact.resources;
 
     try std.testing.expect(program.properties.common_ir_lowered);
+    try std.testing.expect(program.properties.compute_abi_lowered);
+    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.parallel_copies_lowered);
     try std.testing.expect(program.properties.flags_allocated);

src/intel/compiler/targets/gen9/register_allocation.zig

@@ -14,15 +14,15 @@ pub const Error = std.mem.Allocator.Error || error{
 
 pub fn run(allocator: std.mem.Allocator, program: *program_ir.Program) Error!void {
     if (!program.properties.block_parameters_lowered)
-        return error.BlockParametersNotLowered;
+        return Error.BlockParametersNotLowered;
     if (!program.properties.parallel_copies_lowered)
-        return error.ParallelCopiesNotLowered;
+        return Error.ParallelCopiesNotLowered;
     if (program.properties.registers_allocated)
         return;
 
     const grf_size = program.device_info.grf_size_bytes;
     if (grf_size == 0)
-        return error.InvalidProgram;
+        return Error.InvalidProgram;
 
     const allocations = try allocator.alloc(?operand.PhysicalGrf, program.virtual_registers.entries.items.len);
     defer allocator.free(allocations);
@@ -35,9 +35,9 @@ pub fn run(allocator: std.mem.Allocator, program: *program_ir.Program) Error!voi
     for (program.virtual_registers.entries.items, 0..) |entry, index| {
         const register = entry orelse continue;
         const start = std.mem.alignForward(usize, next_byte, register.alignment_bytes);
-        const end = std.math.add(usize, start, register.size_bytes) catch return error.OutOfRegisters;
+        const end = std.math.add(usize, start, register.size_bytes) catch return Error.OutOfRegisters;
         if (end > capacity)
-            return error.OutOfRegisters;
+            return Error.OutOfRegisters;
 
         allocations[index] = .{
             .number = @intCast(start / grf_size),
@@ -47,7 +47,7 @@ pub fn run(allocator: std.mem.Allocator, program: *program_ir.Program) Error!voi
     }
 
     try rewriteProgram(program, allocations);
-    program.program_data.total_grf_count = @intCast(std.math.divCeil(usize, next_byte, grf_size) catch return error.InvalidProgram);
+    program.program_data.total_grf_count = @intCast(std.math.divCeil(usize, next_byte, grf_size) catch return Error.InvalidProgram);
     program.properties.registers_allocated = true;
 }
 
@@ -94,7 +94,7 @@ fn reserveExistingPhysicalRegisters(program: *const program_ir.Program, initial:
                 reserveRegister(&next_byte, op.lhs.register, grf_size);
                 reserveRegister(&next_byte, op.rhs.register, grf_size);
             },
-            .parallel_copy => return error.ParallelCopiesNotLowered,
+            .parallel_copy => return Error.ParallelCopiesNotLowered,
         }
     }
     return next_byte;
@@ -151,15 +151,15 @@ fn rewriteProgram(program: *program_ir.Program, allocations: []const ?operand.Ph
                 try rewriteSource(program, &op.lhs, allocations);
                 try rewriteSource(program, &op.rhs, allocations);
             },
-            .parallel_copy => return error.ParallelCopiesNotLowered,
+            .parallel_copy => return Error.ParallelCopiesNotLowered,
         }
     }
 
     for (program.blocks.entries.items) |*entry| {
         const block = if (entry.*) |*value| value else continue;
         if (block.parameters.items.len != 0)
-            return error.BlockParametersNotLowered;
-        const terminator = if (block.terminator) |*value| value else return error.InvalidProgram;
+            return Error.BlockParametersNotLowered;
+        const terminator = if (block.terminator) |*value| value else return Error.InvalidProgram;
         switch (terminator.*) {
             .jump => |*edge| try rewriteEdge(program, edge, allocations),
             .conditional_branch => |*branch| {
@@ -192,8 +192,8 @@ fn rewriteRegister(program: *const program_ir.Program, register: *operand.Regist
         else => return,
     };
     if (!program.virtual_registers.isLive(virtual) or virtual.index() >= allocations.len)
-        return error.InvalidProgram;
-    const physical = allocations[virtual.index()] orelse return error.InvalidProgram;
+        return Error.InvalidProgram;
+    const physical = allocations[virtual.index()] orelse return Error.InvalidProgram;
     register.* = .{ .physical_grf = physical };
 }
 
@@ -262,6 +262,6 @@ test "[gen9] register allocation: report GRF exhaustion" {
     try program.setTerminator(entry, .end_thread);
     markPrerequisites(&program);
 
-    try std.testing.expectError(error.OutOfRegisters, run(std.testing.allocator, &program));
+    try std.testing.expectError(Error.OutOfRegisters, run(std.testing.allocator, &program));
     try std.testing.expect(!program.properties.registers_allocated);
 }

src/intel/copy_commands.zig

@@ -9,6 +9,7 @@ const FlintCommandBuffer = @import("FlintCommandBuffer.zig");
 const MemoryRange = @import("MemoryRange.zig");
 
 pub fn emitLinearCopy(cmd: *FlintCommandBuffer, src: MemoryRange, dst: MemoryRange) VkError!void {
+    try cmd.requireEngine(.blitter);
     if (src.size != dst.size) return VkError.ValidationFailed;
 
     var copied: vk.DeviceSize = 0;

src/intel/i915/i915.zig

@@ -7,13 +7,17 @@ pub const gem_close = 0x09;
 
 pub const mmap_offset_wb = 2;
 pub const gem_domain_cpu = 0x00000001;
+pub const gem_domain_render = 0x00000004;
+pub const gem_domain_instruction = 0x00000010;
 pub const gem_domain_gtt = 0x00000040;
+pub const exec_render = 1 << 0;
 pub const exec_blt = 3 << 0;
 pub const exec_fence_array: u64 = 1 << 19;
 pub const exec_fence_wait: u32 = 1 << 0;
 pub const exec_fence_signal: u32 = 1 << 1;
 pub const exec_object_write = 1 << 2;
 pub const mi_flush_dw: u32 = (0x26 << 23) | 3;
+pub const mi_batch_buffer_end: u32 = 0x05000000;
 
 pub const GemCreate = extern struct {
     size: u64,

src/intel/i915/kmd.zig

@@ -7,6 +7,11 @@ const common_kmd = @import("../kmd.zig");
 
 const VkError = base.VkError;
 
+const RelocationGroup = struct {
+    source_handle: u32,
+    entries: std.ArrayList(_i915.RelocationEntry) = .empty,
+};
+
 const Mapping = struct {
     bytes: []align(std.heap.page_size_min) u8,
 
@@ -54,8 +59,19 @@ pub const Device = struct {
         return memory;
     }
 
-    pub fn submitBatch(self: *Device, io: std.Io, allocator: std.mem.Allocator, commands: []const u32, relocations: []const common_kmd.Relocation, syncs: []const common_kmd.SyncDependency) VkError!void {
-        const trailer_words = 6;
+    pub fn submitBatch(
+        self: *Device,
+        io: std.Io,
+        allocator: std.mem.Allocator,
+        engine: common_kmd.Engine,
+        commands: []const u32,
+        relocations: []const common_kmd.Relocation,
+        syncs: []const common_kmd.SyncDependency,
+    ) VkError!void {
+        const trailer_words: usize = switch (engine) {
+            .blitter => 6,
+            .render => if (commands.len % 2 == 0) 2 else 1,
+        };
         const batch_size = (commands.len + trailer_words) * @sizeOf(u32);
         var batch = try self.allocateMemory(io, batch_size);
         defer batch.deinit(self, io);
@@ -64,65 +80,93 @@ pub const Device = struct {
             const batch_map = try batch.map(self, io, 0, batch_size);
             const batch_words = std.mem.bytesAsSlice(u32, batch_map);
             @memcpy(batch_words[0..commands.len], commands);
-            batch_words[commands.len + 0] = _i915.mi_flush_dw;
-            batch_words[commands.len + 1] = 0;
-            batch_words[commands.len + 2] = 0;
-            batch_words[commands.len + 3] = 0;
-            batch_words[commands.len + 4] = 0;
-            batch_words[commands.len + 5] = 0x05000000;
+            @memset(batch_words[commands.len..], 0);
+            switch (engine) {
+                .blitter => {
+                    batch_words[commands.len] = _i915.mi_flush_dw;
+                    batch_words[commands.len + 5] = _i915.mi_batch_buffer_end;
+                },
+                .render => batch_words[commands.len] = _i915.mi_batch_buffer_end,
+            }
             batch.unmap();
         }
         try batch.flushRange(self, io, 0, batch_size);
 
-        var objects = std.ArrayList(_i915.ExecObject2).empty;
-        defer objects.deinit(allocator);
-
         var object_handles = std.ArrayList(u32).empty;
         defer object_handles.deinit(allocator);
-
         for (relocations) |relocation| {
-            if (std.mem.indexOfScalar(u32, object_handles.items, relocation.target_handle) == null) {
-                object_handles.append(allocator, relocation.target_handle) catch return VkError.OutOfHostMemory;
-                objects.append(allocator, .{
-                    .handle = relocation.target_handle,
-                    .relocation_count = 0,
-                    .relocs_ptr = 0,
-                    .alignment = 0,
-                    .offset = 0,
-                    .flags = if (relocation.write) _i915.exec_object_write else 0,
-                    .rsvd1 = 0,
-                    .rsvd2 = 0,
-                }) catch return VkError.OutOfHostMemory;
-            } else if (relocation.write) {
-                const index = std.mem.indexOfScalar(u32, object_handles.items, relocation.target_handle).?;
-                objects.items[index].flags |= _i915.exec_object_write;
+            if (relocation.source_handle) |source| {
+                if (std.mem.indexOfScalar(u32, object_handles.items, source) == null)
+                    object_handles.append(allocator, source) catch return VkError.OutOfHostMemory;
             }
+            if (std.mem.indexOfScalar(u32, object_handles.items, relocation.target_handle) == null)
+                object_handles.append(allocator, relocation.target_handle) catch return VkError.OutOfHostMemory;
         }
+        if (std.mem.indexOfScalar(u32, object_handles.items, batch.handle) == null)
+            object_handles.append(allocator, batch.handle) catch return VkError.OutOfHostMemory;
 
-        var i915_relocations = std.ArrayList(_i915.RelocationEntry).empty;
-        defer i915_relocations.deinit(allocator);
-
+        var groups = std.ArrayList(RelocationGroup).empty;
+        defer {
+            for (groups.items) |*group| group.entries.deinit(allocator);
+            groups.deinit(allocator);
+        }
         for (relocations) |relocation| {
-            i915_relocations.append(allocator, .{
+            const source = relocation.source_handle orelse batch.handle;
+            var group_index = std.mem.indexOfScalar(u32, object_handles.items, source) orelse return VkError.DeviceLost;
+            for (groups.items, 0..) |group, index| {
+                if (group.source_handle == source) {
+                    group_index = index;
+                    break;
+                }
+            } else {
+                groups.append(allocator, .{ .source_handle = source }) catch return VkError.OutOfHostMemory;
+                group_index = groups.items.len - 1;
+            }
+
+            const domain: u32 = switch (relocation.domain) {
+                .none => 0,
+                .render => _i915.gem_domain_render,
+                .instruction => _i915.gem_domain_instruction,
+            };
+            groups.items[group_index].entries.append(allocator, .{
                 .target_handle = relocation.target_handle,
                 .delta = relocation.delta,
                 .offset = relocation.offset,
                 .presumed_offset = 0,
-                .read_domains = 0,
-                .write_domain = 0,
+                .read_domains = if (relocation.read) domain else 0,
+                .write_domain = if (relocation.write) domain else 0,
             }) catch return VkError.OutOfHostMemory;
         }
 
-        objects.append(allocator, .{
-            .handle = batch.handle,
-            .relocation_count = @intCast(i915_relocations.items.len),
-            .relocs_ptr = @intFromPtr(i915_relocations.items.ptr),
-            .alignment = 0,
-            .offset = 0,
-            .flags = 0,
-            .rsvd1 = 0,
-            .rsvd2 = 0,
-        }) catch return VkError.OutOfHostMemory;
+        var objects = std.ArrayList(_i915.ExecObject2).empty;
+        defer objects.deinit(allocator);
+        for (object_handles.items) |handle| {
+            var flags: u64 = 0;
+            for (relocations) |relocation| {
+                if (relocation.target_handle == handle and relocation.write)
+                    flags |= _i915.exec_object_write;
+            }
+
+            var relocation_count: u32 = 0;
+            var relocs_ptr: u64 = 0;
+            for (groups.items) |group| {
+                if (group.source_handle == handle) {
+                    relocation_count = @intCast(group.entries.items.len);
+                    relocs_ptr = @intFromPtr(group.entries.items.ptr);
+                    break;
+                }
+            }
+            objects.append(allocator, .{
+                .handle = handle,
+                .relocation_count = relocation_count,
+                .relocs_ptr = relocs_ptr,
+                .alignment = 0,
+                .offset = 0,
+                .flags = flags,
+                .rsvd1 = 0,
+                .rsvd2 = 0,
+            }) catch return VkError.OutOfHostMemory;
+        }
 
         var exec_fences = std.ArrayList(_i915.ExecFence).empty;
         defer exec_fences.deinit(allocator);
@@ -142,7 +186,10 @@ pub const Device = struct {
             .DR4 = 0,
             .num_cliprects = @intCast(exec_fences.items.len),
             .cliprects_ptr = if (exec_fences.items.len == 0) 0 else @intFromPtr(exec_fences.items.ptr),
-            .flags = _i915.exec_blt | (if (exec_fences.items.len == 0) 0 else _i915.exec_fence_array),
+            .flags = @as(u64, switch (engine) {
+                .blitter => _i915.exec_blt,
+                .render => _i915.exec_render,
+            }) | (if (exec_fences.items.len == 0) 0 else _i915.exec_fence_array),
             .rsvd1 = 0,
             .rsvd2 = 0,
         };

src/intel/kmd.zig

@@ -19,12 +19,25 @@ pub const blt_depth_8: u32 = 0 << 24;
 pub const rop_source_copy: u32 = 0xcc << 16;
 pub const max_blt_span: vk.DeviceSize = 32 * 1024 - 1;
 
+pub const Engine = enum {
+    blitter,
+    render,
+};
+
+pub const Domain = enum {
+    none,
+    render,
+    instruction,
+};
+
 pub const Relocation = struct {
+    source_handle: ?u32 = null,
     target_handle: u32,
     offset: u64,
     delta: u32,
     read: bool = false,
     write: bool = false,
+    domain: Domain = .none,
 };
 
 pub const SyncDependency = struct {
@@ -63,10 +76,10 @@ pub const Device = union(KmdType) {
         };
     }
 
-    pub fn submitBatch(self: *Device, io: std.Io, allocator: std.mem.Allocator, commands: []const u32, relocations: []const Relocation, syncs: []const SyncDependency) VkError!void {
+    pub fn submitBatch(self: *Device, io: std.Io, allocator: std.mem.Allocator, engine: Engine, commands: []const u32, relocations: []const Relocation, syncs: []const SyncDependency) VkError!void {
         return switch (self.*) {
-            .i915 => |*device| device.submitBatch(io, allocator, commands, relocations, syncs),
-            .xe => |*device| device.submitBatch(io, allocator, commands, relocations, syncs),
+            .i915 => |*device| device.submitBatch(io, allocator, engine, commands, relocations, syncs),
+            .xe => |*device| device.submitBatch(io, allocator, engine, commands, relocations, syncs),
             .invalid => VkError.DeviceLost,
         };
     }

src/intel/xe/kmd.zig

@@ -23,7 +23,15 @@ pub const Device = struct {
         return VkError.OutOfDeviceMemory;
     }
 
-    pub fn submitBatch(_: *Device, _: std.Io, _: std.mem.Allocator, _: []const u32, _: []const common_kmd.Relocation, _: []const common_kmd.SyncDependency) VkError!void {
+    pub fn submitBatch(
+        _: *Device,
+        _: std.Io,
+        _: std.mem.Allocator,
+        _: common_kmd.Engine,
+        _: []const u32,
+        _: []const common_kmd.Relocation,
+        _: []const common_kmd.SyncDependency,
+    ) VkError!void {
         return VkError.FeatureNotPresent;
     }
 };