[Flint] centralize compute lowering, adding surface messages with GRF

allocation

src/intel/FlintPipeline.zig

@@ -15,16 +15,7 @@ const PipelineKind = enum {
     compute,
 };
 
-pub const ComputeArtifact = struct {
+pub const ComputeArtifact = compiler.targets.ComputeArtifact;
-    program: compiler.Program,
-    resources: compiler.targets.ComputeResourceLayout,
-
-    fn deinit(self: *ComputeArtifact, allocator: std.mem.Allocator) void {
-        self.resources.deinit(allocator);
-        self.program.deinit();
-        self.* = undefined;
-    }
-};
 
 const CommonStage = struct {
     stage: shader_ir.ir.module.Stage,
@@ -144,7 +135,7 @@ fn compileStage(allocator: std.mem.Allocator, info: *const vk.PipelineShaderStag
 
 fn lowerToFlint(allocator: std.mem.Allocator, module: *base.ShaderModule.IrModule, device_info: ?compiler.device.DeviceInfo) VkError!?ComputeArtifact {
     const target = device_info orelse return null;
-    var program = compiler.targets.lower(allocator, module, target, .{}) catch |err| switch (err) {
+    return compiler.targets.compileCompute(allocator, module, target, .{}) catch |err| switch (err) {
         error.OutOfMemory => return VkError.OutOfHostMemory,
 
         error.UnsupportedGeneration,
@@ -156,44 +147,14 @@ fn lowerToFlint(allocator: std.mem.Allocator, module: *base.ShaderModule.IrModul
         error.UnsupportedType,
         error.UnsupportedOperation,
         error.UnsupportedTerminator,
+        error.TooManyStorageBuffers,
         => return null,
 
         else => {
-            std.log.scoped(.FlintPipeline).err("shader lowering failed: {s}", .{@errorName(err)});
+            std.log.scoped(.FlintPipeline).err("compute compilation failed: {s}", .{@errorName(err)});
             return VkError.ValidationFailed;
         },
     };
-    errdefer program.deinit();
-
-    var resources = compiler.targets.layoutComputeResources(allocator, &program) catch |err| switch (err) {
-        error.OutOfMemory => return VkError.OutOfHostMemory,
-        error.UnsupportedGeneration,
-        error.TooManyStorageBuffers,
-        => {
-            program.deinit();
-            return null;
-        },
-    };
-    errdefer resources.deinit(allocator);
-
-    compiler.targets.lowerComputeResources(&program, &resources) catch |err| switch (err) {
-        error.UnsupportedGeneration => {
-            resources.deinit(allocator);
-            program.deinit();
-            return null;
-        },
-        error.InvalidProgram,
-        error.InvalidResourceLayout,
-        => {
-            std.log.scoped(.FlintPipeline).err("Flint compute resource lowering failed: {s}", .{@errorName(err)});
-            return VkError.ValidationFailed;
-        },
-    };
-
-    return .{
-        .program = program,
-        .resources = resources,
-    };
 }
 
 fn validateComputePipelineLayout(layout: *const base.PipelineLayout, resources: *const compiler.targets.ComputeResourceLayout) VkError!void {
@@ -321,6 +282,8 @@ test "Flint pipeline: lower common compute IR" {
     try std.testing.expect(program.properties.block_parameters_lowered);
     try std.testing.expect(!program.properties.system_values_lowered);
     try std.testing.expect(program.properties.resources_lowered);
+    try std.testing.expect(program.properties.messages_lowered);
+    try std.testing.expect(program.properties.registers_allocated);
     try std.testing.expect(!program.properties.instructions_selected);
     try std.testing.expectEqual([3]u32{ 1, 1, 1 }, program.workgroup_size);
     try std.testing.expectEqual(@as(usize, 1), program.storage_buffers.entries.items.len);
@@ -330,6 +293,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 %id_x:u32, component(0)") != null);
+    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, "store_buffer bti(0), 0:u32, %id_x:u32") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "surface_write bti(0), 0:u32, r0:u32") != null);
 }

src/intel/compiler/ir/instruction.zig

@@ -28,6 +28,20 @@ pub const StoreBuffer = struct {
     source: operand.Source,
 };
 
+pub const SurfaceRead = struct {
+    destination: operand.Destination,
+    binding_table: u8,
+    address: operand.Source,
+    immediate_offset: u32 = 0,
+};
+
+pub const SurfaceWrite = struct {
+    binding_table: u8,
+    address: operand.Source,
+    immediate_offset: u32 = 0,
+    data: operand.Source,
+};
+
 pub const Move = struct {
     destination: operand.Destination,
     source: operand.Source,
@@ -70,6 +84,8 @@ pub const Operation = union(enum) {
     load_global_invocation_id: LoadGlobalInvocationId,
     load_buffer: LoadBuffer,
     store_buffer: StoreBuffer,
+    surface_read: SurfaceRead,
+    surface_write: SurfaceWrite,
     move: Move,
     binary: Binary,
     compare: Compare,

src/intel/compiler/ir/printer.zig

@@ -136,6 +136,22 @@ fn writeOperation(program: *const program_ir.Program, writer: *std.Io.Writer, ex
             try writer.writeAll(", ");
             try writeSource(program, writer, execution_size, op.source);
         },
+        .surface_read => |op| {
+            try writer.writeAll("surface_read ");
+            try writeDestination(program, writer, execution_size, op.destination);
+            try writer.print(", bti({d}), ", .{op.binding_table});
+            try writeSource(program, writer, execution_size, op.address);
+            if (op.immediate_offset != 0)
+                try writer.print(", offset({d})", .{op.immediate_offset});
+        },
+        .surface_write => |op| {
+            try writer.print("surface_write bti({d}), ", .{op.binding_table});
+            try writeSource(program, writer, execution_size, op.address);
+            if (op.immediate_offset != 0)
+                try writer.print(", offset({d})", .{op.immediate_offset});
+            try writer.writeAll(", ");
+            try writeSource(program, writer, execution_size, op.data);
+        },
         .move => |op| {
             try writer.writeAll("mov ");
             try writeDestination(program, writer, execution_size, op.destination);

src/intel/compiler/ir/validator.zig

@@ -32,6 +32,7 @@ pub const Error = error{
     UnloweredParallelCopy,
     UnloweredSystemValue,
     UnloweredResource,
+    UnloweredMessage,
     InvalidPayloadLayout,
     EntryBlockHasParameters,
     DuplicateBlockParameter,
@@ -138,6 +139,8 @@ fn validateInstruction(program: *const program_ir.Program, inst: instruction.Ins
                 return Error.InvalidGlobalInvocationId;
         },
         .load_buffer => |op| {
+            if (program.properties.messages_lowered)
+                return Error.UnloweredMessage;
             try validateBufferReference(program, op.buffer);
             try validateDestination(program, op.destination);
             try validateBufferOffset(program, op.byte_offset);
@@ -145,12 +148,26 @@ fn validateInstruction(program: *const program_ir.Program, inst: instruction.Ins
                 return Error.InvalidBufferAccess;
         },
         .store_buffer => |op| {
+            if (program.properties.messages_lowered)
+                return Error.UnloweredMessage;
             try validateBufferReference(program, op.buffer);
             try validateBufferOffset(program, op.byte_offset);
             try validateSource(program, op.source);
             if (!op.source.type.isInitialTargetType())
                 return Error.InvalidBufferAccess;
         },
+        .surface_read => |op| {
+            try validateDestination(program, op.destination);
+            try validateBufferOffset(program, op.address);
+            if (!op.destination.type.isInitialTargetType())
+                return Error.InvalidBufferAccess;
+        },
+        .surface_write => |op| {
+            try validateBufferOffset(program, op.address);
+            try validateSource(program, op.data);
+            if (!op.data.type.isInitialTargetType())
+                return Error.InvalidBufferAccess;
+        },
         .move => |op| {
             try validateDestination(program, op.destination);
             try validateSource(program, op.source);

src/intel/compiler/lower/common_ir.zig

@@ -913,12 +913,6 @@ pub const Lowerer = struct {
 
         program.properties.common_ir_lowered = true;
         validator.validate(&program) catch return Error.InvalidLoweredProgram;
-
-        block_arguments.run(allocator, &program) catch |err| return switch (err) {
-            error.OutOfMemory => Error.OutOfMemory,
-            else => Error.InvalidLoweredProgram,
-        };
-        validator.validate(&program) catch return Error.InvalidLoweredProgram;
         return program;
     }
 };
@@ -964,6 +958,7 @@ fn expectLowered(source: []const u8, expected: []const u8) !void {
 
     var program = try lower(std.testing.allocator, &module, test_device, .{});
     defer program.deinit();
+    try block_arguments.run(std.testing.allocator, &program);
     try std.testing.expect(program.properties.common_ir_lowered);
     try std.testing.expect(!program.properties.instructions_selected);
 
@@ -979,6 +974,7 @@ fn expectLoweredFragments(source: []const u8, expected: []const []const u8, unex
 
     var program = try lower(std.testing.allocator, &module, test_device, .{});
     defer program.deinit();
+    try block_arguments.run(std.testing.allocator, &program);
     try std.testing.expect(program.properties.common_ir_lowered);
     try std.testing.expect(!program.properties.instructions_selected);
 

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

@@ -1,5 +1,6 @@
 const std = @import("std");
 
+pub const message_lowering = @import("message_lowering.zig");
 pub const resource_layout = @import("resource_layout.zig");
 pub const resource_lowering = @import("resource_lowering.zig");
 pub const ResourceLayout = resource_layout.Layout;

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

@@ -0,0 +1,110 @@
+const instruction = @import("../../../ir/instruction.zig");
+const operand = @import("../../../ir/operand.zig");
+const program_ir = @import("../../../ir/program.zig");
+
+pub const Error = error{
+    ResourcesNotLowered,
+    InvalidProgram,
+};
+
+pub fn run(program: *program_ir.Program) Error!void {
+    if (!program.properties.resources_lowered)
+        return error.ResourcesNotLowered;
+    if (program.properties.messages_lowered)
+        return;
+
+    for (program.instructions.entries.items) |*entry| {
+        const inst = if (entry.*) |*value| value else continue;
+        inst.operation = switch (inst.operation) {
+            .load_buffer => |op| .{ .surface_read = .{
+                .destination = op.destination,
+                .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,
+                .address = op.byte_offset,
+                .immediate_offset = op.immediate_offset,
+                .data = op.source,
+            } },
+            else => inst.operation,
+        };
+    }
+
+    program.properties.messages_lowered = true;
+}
+
+fn bindingTableIndex(reference: instruction.BufferReference) ?u8 {
+    return switch (reference) {
+        .binding_table => |index| index,
+        .logical => null,
+    };
+}
+
+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,
+};
+
+fn immediate(value: u32) operand.Source {
+    return .{
+        .register = .{ .immediate = .{ .u32 = value } },
+        .type = .u32,
+        .region = operand.Region.broadcast(),
+    };
+}
+
+test "[gen9] compute message lowering: select surface messages" {
+    var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, test_device, .simd8);
+    defer program.deinit();
+
+    const value = try program.addVirtualRegister(.{
+        .size_bytes = 32,
+        .alignment_bytes = 32,
+        .element_type = .u32,
+        .lane_count = 8,
+        .class = .temporary,
+    });
+    const entry = try program.addBlock("entry");
+    const load = try program.appendInstruction(entry, .simd8, null, .{ .load_buffer = .{
+        .destination = .{ .register = .{ .virtual = value }, .type = .u32 },
+        .buffer = .{ .binding_table = 2 },
+        .byte_offset = immediate(16),
+        .immediate_offset = 4,
+    } });
+    const store = try program.appendInstruction(entry, .simd8, null, .{ .store_buffer = .{
+        .buffer = .{ .binding_table = 3 },
+        .byte_offset = immediate(32),
+        .immediate_offset = 8,
+        .source = .{
+            .register = .{ .virtual = value },
+            .type = .u32,
+            .region = operand.Region.contiguous(.simd8),
+        },
+    } });
+    try program.setTerminator(entry, .end_thread);
+    program.properties.resources_lowered = true;
+
+    try run(&program);
+
+    const read = program.instructions.get(load).?.operation.surface_read;
+    try std.testing.expectEqual(@as(u8, 2), read.binding_table);
+    try std.testing.expectEqual(@as(u32, 4), read.immediate_offset);
+    const write = program.instructions.get(store).?.operation.surface_write;
+    try std.testing.expectEqual(@as(u8, 3), write.binding_table);
+    try std.testing.expectEqual(@as(u32, 8), write.immediate_offset);
+    try std.testing.expect(program.properties.messages_lowered);
+}
+
+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));
+}

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

@@ -0,0 +1,74 @@
+const std = @import("std");
+const shader_ir = @import("shader_ir").ir;
+
+const device = @import("../../../device.zig");
+const program_ir = @import("../../../ir/program.zig");
+const common_ir = @import("../../../lower/common_ir.zig");
+const block_arguments = @import("../../../lower/block_arguments.zig");
+const parallel_copies = @import("../../../lower/parallel_copies.zig");
+const flag_allocation = @import("../flag_allocation.zig");
+const register_allocation = @import("../register_allocation.zig");
+
+const compute = @import("compute.zig");
+const message_lowering = @import("message_lowering.zig");
+const resource_layout = @import("resource_layout.zig");
+const resource_lowering = @import("resource_lowering.zig");
+
+pub const Error = common_ir.Error || block_arguments.Error || parallel_copies.Error ||
+    message_lowering.Error || resource_layout.Error || resource_lowering.Error || flag_allocation.Error || register_allocation.Error || compute.Error || error{
+    UnsupportedGeneration,
+    UnsupportedStage,
+    UnsupportedDispatchWidth,
+    UnsupportedGrfSize,
+};
+
+pub const Artifact = struct {
+    program: program_ir.Program,
+    resources: resource_layout.Layout,
+
+    pub fn deinit(self: *Artifact, allocator: std.mem.Allocator) void {
+        self.resources.deinit(allocator);
+        self.program.deinit();
+        self.* = undefined;
+    }
+};
+
+pub fn compile(allocator: std.mem.Allocator, module: *shader_ir.module.Module, device_info: device.DeviceInfo, options: common_ir.Options) Error!Artifact {
+    if (device_info.generation != .gen9)
+        return Error.UnsupportedGeneration;
+    if (module.stage != .compute)
+        return Error.UnsupportedStage;
+    if (options.dispatch_width != .simd8 or !device_info.supportsDispatch(.simd8))
+        return Error.UnsupportedDispatchWidth;
+    if (device_info.grf_size_bytes != 32)
+        return Error.UnsupportedGrfSize;
+    if (module.execution_modes.workgroup_size) |workgroup_size|
+        try compute.validateWorkgroupSize(workgroup_size);
+
+    var program = try common_ir.lower(
+        allocator,
+        module,
+        device_info,
+        options,
+    );
+    errdefer program.deinit();
+
+    try block_arguments.run(allocator, &program);
+    try parallel_copies.run(allocator, &program);
+
+    var resources = try resource_layout.Layout.init(
+        allocator,
+        &program,
+    );
+    errdefer resources.deinit(allocator);
+
+    try resource_lowering.run(&program, &resources);
+    try message_lowering.run(&program);
+    try flag_allocation.run(allocator, &program);
+    try register_allocation.run(allocator, &program);
+
+    return .{
+        .program = program,
+        .resources = resources,
+    };
+}

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

@@ -1,50 +1,20 @@
 const std = @import("std");
 const shader_ir = @import("shader_ir").ir;
 const device = @import("../../device.zig");
-const program_ir = @import("../../ir/program.zig");
 const common_ir = @import("../../lower/common_ir.zig");
-const parallel_copies = @import("../../lower/parallel_copies.zig");
 
 pub const compute = @import("compute/compute.zig");
+pub const compute_pipeline = @import("compute/pipeline.zig");
 pub const flag_allocation = @import("flag_allocation.zig");
+pub const register_allocation = @import("register_allocation.zig");
 pub const validator = @import("validator.zig");
 
 pub const Options = common_ir.Options;
-pub const ResourceLoweringError = compute.resource_lowering.Error;
+pub const ComputeArtifact = compute_pipeline.Artifact;
+pub const Error = compute_pipeline.Error;
 
-pub const Error = common_ir.Error || compute.Error || flag_allocation.Error || error{
+pub fn compileCompute(allocator: std.mem.Allocator, module: *shader_ir.module.Module, device_info: device.DeviceInfo, options: Options) Error!ComputeArtifact {
-    UnsupportedGeneration,
+    return compute_pipeline.compile(allocator, module, device_info, options);
-    UnsupportedStage,
-    UnsupportedDispatchWidth,
-    UnsupportedGrfSize,
-};
-
-pub fn lower(allocator: std.mem.Allocator, module: *shader_ir.module.Module, device_info: device.DeviceInfo, options: Options) Error!program_ir.Program {
-    if (device_info.generation != .gen9)
-        return Error.UnsupportedGeneration;
-    if (module.stage != .compute)
-        return Error.UnsupportedStage;
-    if (options.dispatch_width != .simd8 or !device_info.supportsDispatch(.simd8))
-        return Error.UnsupportedDispatchWidth;
-    if (device_info.grf_size_bytes != 32)
-        return Error.UnsupportedGrfSize;
-    if (module.execution_modes.workgroup_size) |workgroup_size|
-        try compute.validateWorkgroupSize(workgroup_size);
-
-    var program = try common_ir.lower(allocator, module, device_info, options);
-    errdefer program.deinit();
-    parallel_copies.run(allocator, &program) catch |err| return switch (err) {
-        error.OutOfMemory => Error.OutOfMemory,
-        error.InvalidProgram => Error.InvalidLoweredProgram,
-    };
-    try flag_allocation.run(allocator, &program);
-    validator.validate(&program) catch return Error.InvalidLoweredProgram;
-    return program;
-}
-
-pub fn lowerComputeResources(program: *program_ir.Program, layout: *const compute.ResourceLayout) ResourceLoweringError!void {
-    try compute.resource_lowering.run(program, layout);
-    validator.validate(program) catch return ResourceLoweringError.InvalidProgram;
 }
 
 test "[gen9] target: reject unsupported target configurations" {
@@ -68,17 +38,17 @@ test "[gen9] target: reject unsupported target configurations" {
     };
     var other_generation = gen9_device;
     other_generation.generation = .gen11;
-    try std.testing.expectError(Error.UnsupportedGeneration, lower(std.testing.allocator, &module, other_generation, .{}));
+    try std.testing.expectError(Error.UnsupportedGeneration, compileCompute(std.testing.allocator, &module, other_generation, .{}));
 
     module.stage = .fragment;
-    try std.testing.expectError(Error.UnsupportedStage, lower(std.testing.allocator, &module, gen9_device, .{}));
+    try std.testing.expectError(Error.UnsupportedStage, compileCompute(std.testing.allocator, &module, gen9_device, .{}));
     module.stage = .compute;
 
-    try std.testing.expectError(Error.UnsupportedDispatchWidth, lower(std.testing.allocator, &module, gen9_device, .{ .dispatch_width = .simd16 }));
+    try std.testing.expectError(Error.UnsupportedDispatchWidth, compileCompute(std.testing.allocator, &module, gen9_device, .{ .dispatch_width = .simd16 }));
 
     var wide_grf = gen9_device;
     wide_grf.grf_size_bytes = 64;
-    try std.testing.expectError(Error.UnsupportedGrfSize, lower(std.testing.allocator, &module, wide_grf, .{}));
+    try std.testing.expectError(Error.UnsupportedGrfSize, compileCompute(std.testing.allocator, &module, wide_grf, .{}));
 }
 
 test "[gen9] target: lower 256 KiB SSBO copy loop" {
@@ -123,16 +93,19 @@ test "[gen9] target: lower 256 KiB SSBO copy loop" {
         .pci_device_id = 0x1912,
         .grf_count = 128,
     };
-    var program = try lower(std.testing.allocator, &module, gen9_device, .{});
+    var artifact = try compileCompute(std.testing.allocator, &module, gen9_device, .{});
-    defer program.deinit();
+    defer artifact.deinit(std.testing.allocator);
+    const program = &artifact.program;
+    const resources = &artifact.resources;
 
     try std.testing.expect(program.properties.common_ir_lowered);
     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);
+    try std.testing.expect(program.properties.registers_allocated);
 
-    var resources = try compute.ResourceLayout.init(std.testing.allocator, &program);
+    try std.testing.expect(program.properties.resources_lowered);
-    defer resources.deinit(std.testing.allocator);
+    try std.testing.expect(program.properties.messages_lowered);
     try std.testing.expectEqual(@as(usize, 2), resources.bindings.len);
     try std.testing.expectEqual(compute.resource_layout.Binding{
         .set = 0,
@@ -145,9 +118,6 @@ test "[gen9] target: lower 256 KiB SSBO copy loop" {
         .binding_table_index = 1,
     }, resources.bindings[1]);
 
-    try lowerComputeResources(&program, &resources);
-    try validator.validate(&program);
-
     var load_offsets: [4]bool = @splat(false);
     var store_offsets: [4]bool = @splat(false);
     var load_count: usize = 0;
@@ -155,16 +125,16 @@ test "[gen9] target: lower 256 KiB SSBO copy loop" {
     for (program.instructions.entries.items) |instruction_entry| {
         const inst = instruction_entry orelse continue;
         switch (inst.operation) {
-            .load_buffer => |operation| {
+            .surface_read => |operation| {
-                try std.testing.expectEqual(@as(u8, 0), operation.buffer.binding_table);
+                try std.testing.expectEqual(@as(u8, 0), operation.binding_table);
                 try std.testing.expect(operation.immediate_offset % @sizeOf(u32) == 0);
                 const component = operation.immediate_offset / @sizeOf(u32);
                 try std.testing.expect(component < load_offsets.len);
                 load_offsets[component] = true;
                 load_count += 1;
             },
-            .store_buffer => |operation| {
+            .surface_write => |operation| {
-                try std.testing.expectEqual(@as(u8, 1), operation.buffer.binding_table);
+                try std.testing.expectEqual(@as(u8, 1), operation.binding_table);
                 try std.testing.expect(operation.immediate_offset % @sizeOf(u32) == 0);
                 const component = operation.immediate_offset / @sizeOf(u32);
                 try std.testing.expect(component < store_offsets.len);

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

@@ -0,0 +1,257 @@
+const std = @import("std");
+
+const ids = @import("../../ir/id.zig");
+const instruction = @import("../../ir/instruction.zig");
+const operand = @import("../../ir/operand.zig");
+const program_ir = @import("../../ir/program.zig");
+
+pub const Error = std.mem.Allocator.Error || error{
+    BlockParametersNotLowered,
+    ParallelCopiesNotLowered,
+    InvalidProgram,
+    OutOfRegisters,
+};
+
+pub fn run(allocator: std.mem.Allocator, program: *program_ir.Program) Error!void {
+    if (!program.properties.block_parameters_lowered)
+        return error.BlockParametersNotLowered;
+    if (!program.properties.parallel_copies_lowered)
+        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;
+
+    const allocations = try allocator.alloc(?operand.PhysicalGrf, program.virtual_registers.entries.items.len);
+    defer allocator.free(allocations);
+    @memset(allocations, null);
+
+    var next_byte: usize = @as(usize, program.program_data.payload_grf_count) * grf_size;
+    next_byte = try reserveExistingPhysicalRegisters(program, next_byte, grf_size);
+    const capacity: usize = @as(usize, program.device_info.grf_count) * grf_size;
+
+    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;
+        if (end > capacity)
+            return error.OutOfRegisters;
+
+        allocations[index] = .{
+            .number = @intCast(start / grf_size),
+            .byte_offset = @intCast(start % grf_size),
+        };
+        next_byte = end;
+    }
+
+    try rewriteProgram(program, allocations);
+    program.program_data.total_grf_count = @intCast(std.math.divCeil(usize, next_byte, grf_size) catch return error.InvalidProgram);
+    program.properties.registers_allocated = true;
+}
+
+fn reserveExistingPhysicalRegisters(program: *const program_ir.Program, initial: usize, grf_size: usize) Error!usize {
+    var next_byte = initial;
+    if (program.payload.header_grf) |header|
+        reservePhysical(&next_byte, header, grf_size);
+
+    for (program.instructions.entries.items) |entry| {
+        const inst = entry orelse continue;
+        switch (inst.operation) {
+            .load_global_invocation_id => |op| reserveRegister(&next_byte, op.destination.register, grf_size),
+            .load_buffer => |op| {
+                reserveRegister(&next_byte, op.destination.register, grf_size);
+                reserveRegister(&next_byte, op.byte_offset.register, grf_size);
+            },
+            .store_buffer => |op| {
+                reserveRegister(&next_byte, op.byte_offset.register, grf_size);
+                reserveRegister(&next_byte, op.source.register, grf_size);
+            },
+            .surface_read => |op| {
+                reserveRegister(&next_byte, op.destination.register, grf_size);
+                reserveRegister(&next_byte, op.address.register, grf_size);
+            },
+            .surface_write => |op| {
+                reserveRegister(&next_byte, op.address.register, grf_size);
+                reserveRegister(&next_byte, op.data.register, grf_size);
+            },
+            .move => |op| {
+                reserveRegister(&next_byte, op.destination.register, grf_size);
+                reserveRegister(&next_byte, op.source.register, grf_size);
+            },
+            .binary => |op| {
+                reserveRegister(&next_byte, op.destination.register, grf_size);
+                reserveRegister(&next_byte, op.lhs.register, grf_size);
+                reserveRegister(&next_byte, op.rhs.register, grf_size);
+            },
+            .compare => |op| {
+                reserveRegister(&next_byte, op.lhs.register, grf_size);
+                reserveRegister(&next_byte, op.rhs.register, grf_size);
+            },
+            .parallel_copy => return error.ParallelCopiesNotLowered,
+        }
+    }
+    return next_byte;
+}
+
+fn reserveRegister(next_byte: *usize, register: operand.RegisterRef, grf_size: usize) void {
+    switch (register) {
+        .physical_grf => |physical| reservePhysical(next_byte, physical, grf_size),
+        else => {},
+    }
+}
+
+fn reservePhysical(next_byte: *usize, physical: operand.PhysicalGrf, grf_size: usize) void {
+    const end = (@as(usize, physical.number) + 1) * grf_size;
+    next_byte.* = @max(next_byte.*, end);
+}
+
+fn rewriteProgram(program: *program_ir.Program, allocations: []const ?operand.PhysicalGrf) Error!void {
+    for (program.instructions.entries.items) |*entry| {
+        const inst = if (entry.*) |*value| value else continue;
+        switch (inst.operation) {
+            .load_global_invocation_id => |*op| try rewriteDestination(program, &op.destination, allocations),
+            .load_buffer => |*op| {
+                try rewriteDestination(program, &op.destination, allocations);
+                try rewriteSource(program, &op.byte_offset, allocations);
+            },
+            .store_buffer => |*op| {
+                try rewriteSource(program, &op.byte_offset, allocations);
+                try rewriteSource(program, &op.source, allocations);
+            },
+            .surface_read => |*op| {
+                try rewriteDestination(program, &op.destination, allocations);
+                try rewriteSource(program, &op.address, allocations);
+            },
+            .surface_write => |*op| {
+                try rewriteSource(program, &op.address, allocations);
+                try rewriteSource(program, &op.data, allocations);
+            },
+            .move => |*op| {
+                try rewriteDestination(program, &op.destination, allocations);
+                try rewriteSource(program, &op.source, allocations);
+            },
+            .binary => |*op| {
+                try rewriteDestination(program, &op.destination, allocations);
+                try rewriteSource(program, &op.lhs, allocations);
+                try rewriteSource(program, &op.rhs, allocations);
+            },
+            .compare => |*op| {
+                try rewriteSource(program, &op.lhs, allocations);
+                try rewriteSource(program, &op.rhs, allocations);
+            },
+            .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;
+        switch (terminator.*) {
+            .jump => |*edge| try rewriteEdge(program, edge, allocations),
+            .conditional_branch => |*branch| {
+                try rewriteEdge(program, &branch.true_edge, allocations);
+                try rewriteEdge(program, &branch.false_edge, allocations);
+            },
+            .end_thread, .@"unreachable" => {},
+        }
+    }
+}
+
+fn rewriteEdge(program: *const program_ir.Program, edge: *instruction.Edge, allocations: []const ?operand.PhysicalGrf) Error!void {
+    for (@constCast(edge.arguments)) |*argument| switch (argument.*) {
+        .source => |*edge_source| try rewriteSource(program, edge_source, allocations),
+        .predicate => {},
+    };
+}
+
+fn rewriteSource(program: *const program_ir.Program, value: *operand.Source, allocations: []const ?operand.PhysicalGrf) Error!void {
+    try rewriteRegister(program, &value.register, allocations);
+}
+
+fn rewriteDestination(program: *const program_ir.Program, destination: *operand.Destination, allocations: []const ?operand.PhysicalGrf) Error!void {
+    try rewriteRegister(program, &destination.register, allocations);
+}
+
+fn rewriteRegister(program: *const program_ir.Program, register: *operand.RegisterRef, allocations: []const ?operand.PhysicalGrf) Error!void {
+    const virtual = switch (register.*) {
+        .virtual => |value| value,
+        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;
+    register.* = .{ .physical_grf = physical };
+}
+
+const test_device = @import("../../device.zig").DeviceInfo{
+    .generation = .gen9,
+    .platform = .skylake,
+    .pci_device_id = 0x1912,
+    .grf_count = 128,
+};
+
+fn addRegister(program: *program_ir.Program, size: u32, alignment: u16) !ids.VirtualRegisterId {
+    return program.addVirtualRegister(.{
+        .size_bytes = size,
+        .alignment_bytes = alignment,
+        .element_type = .u32,
+        .lane_count = 8,
+        .class = .temporary,
+    });
+}
+
+fn source(register: ids.VirtualRegisterId) operand.Source {
+    return .{
+        .register = .{ .virtual = register },
+        .type = .u32,
+        .region = operand.Region.contiguous(.simd8),
+    };
+}
+
+fn markPrerequisites(program: *program_ir.Program) void {
+    program.properties.block_parameters_lowered = true;
+    program.properties.parallel_copies_lowered = true;
+}
+
+test "[gen9] register allocation: assign non-overlapping physical GRFs" {
+    var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, test_device, .simd8);
+    defer program.deinit();
+    program.program_data.payload_grf_count = 1;
+
+    const first = try addRegister(&program, 32, 32);
+    const second = try addRegister(&program, 64, 32);
+    const entry = try program.addBlock("entry");
+    const move = try program.appendInstruction(entry, .simd8, null, .{ .move = .{
+        .destination = .{ .register = .{ .virtual = second }, .type = .u32 },
+        .source = source(first),
+    } });
+    try program.setTerminator(entry, .end_thread);
+    markPrerequisites(&program);
+
+    try run(std.testing.allocator, &program);
+
+    const operation = program.instructions.get(move).?.operation.move;
+    try std.testing.expectEqual(operand.PhysicalGrf{ .number = 1 }, operation.source.register.physical_grf);
+    try std.testing.expectEqual(operand.PhysicalGrf{ .number = 2 }, operation.destination.register.physical_grf);
+    try std.testing.expectEqual(@as(u16, 4), program.program_data.total_grf_count);
+    try std.testing.expect(program.properties.registers_allocated);
+}
+
+test "[gen9] register allocation: report GRF exhaustion" {
+    var limited_device = test_device;
+    limited_device.grf_count = 2;
+    var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, limited_device, .simd8);
+    defer program.deinit();
+
+    _ = try addRegister(&program, 96, 32);
+    const entry = try program.addBlock("entry");
+    try program.setTerminator(entry, .end_thread);
+    markPrerequisites(&program);
+
+    try std.testing.expectError(error.OutOfRegisters, run(std.testing.allocator, &program));
+    try std.testing.expect(!program.properties.registers_allocated);
+}

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

@@ -64,6 +64,16 @@ fn validateInstruction(inst: instruction.Instruction) Error!void {
             try validateSource(op.byte_offset);
             try validateSource(op.source);
         },
+        .surface_read => |op| {
+            try validateBindingTableIndex(op.binding_table);
+            try validateDestination(op.destination);
+            try validateSource(op.address);
+        },
+        .surface_write => |op| {
+            try validateBindingTableIndex(op.binding_table);
+            try validateSource(op.address);
+            try validateSource(op.data);
+        },
         .move => |op| {
             try validateDestination(op.destination);
             try validateSource(op.source);
@@ -94,11 +104,15 @@ fn validateInstruction(inst: instruction.Instruction) Error!void {
 fn validateBufferReference(reference: instruction.BufferReference) Error!void {
     switch (reference) {
         .logical => {},
-        .binding_table => |index| if (index >= compute.resource_layout.max_storage_buffers)
+        .binding_table => |index| try validateBindingTableIndex(index),
-            return Error.InvalidBindingTableIndex,
     }
 }
 
+fn validateBindingTableIndex(index: u8) Error!void {
+    if (index >= compute.resource_layout.max_storage_buffers)
+        return Error.InvalidBindingTableIndex;
+}
+
 fn validateSource(source: operand.Source) Error!void {
     try validateType(source.type);
     switch (source.register) {

src/intel/compiler/targets/targets.zig

@@ -6,39 +6,23 @@ const common_ir = @import("../lower/common_ir.zig");
 
 pub const gen9 = @import("gen9/gen9.zig");
 
+pub const ComputeArtifact = gen9.ComputeArtifact;
 pub const ComputeResourceLayout = gen9.compute.ResourceLayout;
-pub const ResourceLayoutError = gen9.compute.resource_layout.Error || error{UnsupportedGeneration};
-pub const ResourceLoweringError = gen9.ResourceLoweringError || error{UnsupportedGeneration};
-
 pub const Error = gen9.Error || error{UnsupportedGeneration};
 pub const ValidationError = gen9.validator.Error || error{UnsupportedGeneration};
 
-pub fn lower(
+pub fn compileCompute(
     allocator: std.mem.Allocator,
     module: *shader_ir.module.Module,
     device_info: device.DeviceInfo,
     options: common_ir.Options,
-) Error!program_ir.Program {
+) Error!ComputeArtifact {
     return switch (device_info.generation) {
-        .gen9 => gen9.lower(allocator, module, device_info, options),
+        .gen9 => gen9.compileCompute(allocator, module, device_info, options),
         .gen10, .gen11 => Error.UnsupportedGeneration,
     };
 }
 
-pub fn layoutComputeResources(allocator: std.mem.Allocator, program: *const program_ir.Program) ResourceLayoutError!ComputeResourceLayout {
-    return switch (program.device_info.generation) {
-        .gen9 => ComputeResourceLayout.init(allocator, program),
-        .gen10, .gen11 => ResourceLayoutError.UnsupportedGeneration,
-    };
-}
-
-pub fn lowerComputeResources(program: *program_ir.Program, layout: *const ComputeResourceLayout) ResourceLoweringError!void {
-    return switch (program.device_info.generation) {
-        .gen9 => gen9.lowerComputeResources(program, layout),
-        .gen10, .gen11 => ResourceLoweringError.UnsupportedGeneration,
-    };
-}
-
 pub fn validate(program: *const program_ir.Program) ValidationError!void {
     return switch (program.device_info.generation) {
         .gen9 => gen9.validator.validate(program),