[Flint] adding compute resource layout lowering

src/intel/FlintCommandBuffer.zig

@@ -174,16 +174,24 @@ pub fn beginRenderPass(interface: *Interface, render_pass: *base.RenderPass, fra
 
 pub fn bindDescriptorSets(interface: *Interface, bind_point: vk.PipelineBindPoint, first_set: u32, sets: [base.vulkan_max_descriptor_sets]?*base.DescriptorSet, dynamic_offsets: []const u32) VkError!void {
     const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
-    if (bind_point != .compute) return;
-    if (first_set >= base.vulkan_max_descriptor_sets) return VkError.ValidationFailed;
+    if (bind_point != .compute)
+        return;
+
+    if (dynamic_offsets.len != 0)
+        return VkError.FeatureNotPresent;
+
+    if (first_set >= base.vulkan_max_descriptor_sets)
+        return VkError.ValidationFailed;
 
     for (sets, 0..) |set, index| {
         const base_set = set orelse break;
         const destination = first_set + index;
-        if (destination >= base.vulkan_max_descriptor_sets) return VkError.ValidationFailed;
+
+        if (destination >= base.vulkan_max_descriptor_sets)
+            return VkError.ValidationFailed;
+
         self.bound_compute_descriptor_sets[destination] = @alignCast(@fieldParentPtr("interface", base_set));
     }
-    _ = dynamic_offsets;
 }
 
 pub fn bindPipeline(interface: *Interface, bind_point: vk.PipelineBindPoint, pipeline: *base.Pipeline) VkError!void {
@@ -289,13 +297,38 @@ pub fn dispatch(interface: *Interface, group_count_x: u32, group_count_y: u32, g
 }
 
 pub fn dispatchBase(interface: *Interface, base_group_x: u32, base_group_y: u32, base_group_z: u32, group_count_x: u32, group_count_y: u32, group_count_z: u32) VkError!void {
-    _ = interface;
-    _ = base_group_x;
-    _ = base_group_y;
-    _ = base_group_z;
-    _ = group_count_x;
-    _ = group_count_y;
-    _ = group_count_z;
+    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;
+    }
+
+    const pipeline = self.bound_compute_pipeline orelse return VkError.ValidationFailed;
+    const artifact = pipeline.computeArtifact() orelse return VkError.FeatureNotPresent;
+    for (artifact.resources.bindings) |resource| {
+        if (resource.set >= base.vulkan_max_descriptor_sets)
+            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;
+    }
 }
 
 pub fn setDeviceMask(interface: *Interface, device_mask: u32) VkError!void {

src/intel/FlintPipeline.zig

@@ -15,14 +15,25 @@ const PipelineKind = enum {
     compute,
 };
 
+pub const ComputeArtifact = struct {
+    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,
     module: base.ShaderModule.IrModule,
-    program: ?compiler.Program,
+    artifact: ?ComputeArtifact,
 
-    fn deinit(self: *CommonStage) void {
-        if (self.program) |*program|
-            program.deinit();
+    fn deinit(self: *CommonStage, allocator: std.mem.Allocator) void {
+        if (self.artifact) |*artifact|
+            artifact.deinit(allocator);
         self.module.deinit();
         self.* = undefined;
     }
@@ -48,6 +59,8 @@ 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));
+    if (self.computeArtifact()) |artifact|
+        try validateComputePipelineLayout(self.interface.layout, &artifact.resources);
     return self;
 }
 
@@ -74,12 +87,7 @@ pub fn createGraphics(device: *base.Device, allocator: std.mem.Allocator, cache:
     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, infos: []const vk.PipelineShaderStageCreateInfo, pipeline_kind: PipelineKind, device_info: ?compiler.device.DeviceInfo) VkError![]CommonStage {
     if (infos.len == 0)
         return VkError.ValidationFailed;
 
@@ -87,7 +95,7 @@ fn compileStages(
     var initialized: usize = 0;
     errdefer {
         for (stages[0..initialized]) |*stage|
-            stage.deinit();
+            stage.deinit(allocator);
         allocator.free(stages);
     }
 
@@ -98,12 +106,7 @@ fn compileStages(
     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, 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);
 
@@ -129,24 +132,21 @@ fn compileStage(
 
     std.debug.assert(module.stage == expected_stage);
 
-    var program = try lowerToFlint(allocator, &module, device_info);
-    errdefer if (program) |*value| value.deinit();
+    var artifact = try lowerToFlint(allocator, &module, device_info);
+    errdefer if (artifact) |*value| value.deinit(allocator);
 
     return .{
         .stage = expected_stage,
         .module = module,
-        .program = program,
+        .artifact = artifact,
     };
 }
 
-fn lowerToFlint(
-    allocator: std.mem.Allocator,
-    module: *base.ShaderModule.IrModule,
-    device_info: ?compiler.device.DeviceInfo,
-) VkError!?compiler.Program {
+fn lowerToFlint(allocator: std.mem.Allocator, module: *base.ShaderModule.IrModule, device_info: ?compiler.device.DeviceInfo) VkError!?ComputeArtifact {
     const target = device_info orelse return null;
-    const program = compiler.targets.lower(allocator, module, target, .{}) catch |err| switch (err) {
+    var program = compiler.targets.lower(allocator, module, target, .{}) catch |err| switch (err) {
         error.OutOfMemory => return VkError.OutOfHostMemory,
+
         error.UnsupportedGeneration,
         error.UnsupportedStage,
         error.UnsupportedDispatchWidth,
@@ -157,12 +157,60 @@ fn lowerToFlint(
         error.UnsupportedOperation,
         error.UnsupportedTerminator,
         => return null,
+
         else => {
-            std.log.scoped(.FlintPipeline).err("Flint shader lowering failed: {s}", .{@errorName(err)});
+            std.log.scoped(.FlintPipeline).err("shader lowering failed: {s}", .{@errorName(err)});
             return VkError.ValidationFailed;
         },
     };
-    return program;
+    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 {
+    for (resources.bindings) |resource| {
+        if (resource.set >= layout.set_count)
+            return VkError.ValidationFailed;
+
+        const set_layout = layout.set_layouts[resource.set] orelse return VkError.ValidationFailed;
+
+        if (resource.binding >= set_layout.bindings.len)
+            return VkError.ValidationFailed;
+
+        const binding = set_layout.bindings[resource.binding];
+
+        if (binding.descriptor_type != .storage_buffer or binding.array_size == 0)
+            return VkError.ValidationFailed;
+    }
 }
 
 fn compilerDeviceInfo(device: *const base.Device) ?compiler.device.DeviceInfo {
@@ -216,11 +264,17 @@ fn commonStage(stage: vk.ShaderStageFlags) ?shader_ir.ir.module.Stage {
 
 fn deinitStages(allocator: std.mem.Allocator, stages: []CommonStage) void {
     for (stages) |*stage|
-        stage.deinit();
+        stage.deinit(allocator);
     if (stages.len != 0)
         allocator.free(stages);
 }
 
+pub fn computeArtifact(self: *const Self) ?*const ComputeArtifact {
+    if (self.interface.bind_point != .compute or self.stages.len != 1 or self.stages[0].stage != .compute)
+        return null;
+    return if (self.stages[0].artifact) |*artifact| artifact else null;
+}
+
 pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void {
     const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
     deinitStages(self.artifact_allocator.allocator(), self.stages);
@@ -259,20 +313,23 @@ test "Flint pipeline: lower common compute IR" {
     }, null);
     try builder.setTerminator(entry, .return_void);
 
-    var program = (try lowerToFlint(std.testing.allocator, &module, device_info)).?;
-    defer program.deinit();
+    var artifact = (try lowerToFlint(std.testing.allocator, &module, device_info)).?;
+    defer artifact.deinit(std.testing.allocator);
+    const program = &artifact.program;
 
     try std.testing.expect(program.properties.common_ir_lowered);
     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.resources_lowered);
     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);
-    try compiler.targets.validate(&program);
+    try std.testing.expectEqual(@as(usize, 1), artifact.resources.bindings.len);
+    try std.testing.expectEqual(@as(u8, 0), artifact.resources.bindings[0].binding_table_index);
+    try compiler.targets.validate(program);
 
-    const text = try compiler.printer.allocPrint(std.testing.allocator, &program);
+    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, "store_buffer @storage, 0:u32, %id_x:u32") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "store_buffer bti(0), 0:u32, %id_x:u32") != null);
 }

src/intel/compiler/compiler.zig

@@ -51,7 +51,7 @@ test "[ir] basic compute shader" {
     });
     _ = try builder.appendInstruction(entry, .simd8, null, .{
         .store_buffer = .{
-            .buffer = storage,
+            .buffer = .{ .logical = storage },
             .byte_offset = .{
                 .register = .{ .immediate = .{ .u32 = 0 } },
                 .type = .u32,

src/intel/compiler/ir/instruction.zig

@@ -9,15 +9,20 @@ pub const LoadGlobalInvocationId = struct {
     component: u8,
 };
 
+pub const BufferReference = union(enum) {
+    logical: ids.StorageBufferId,
+    binding_table: u8,
+};
+
 pub const LoadBuffer = struct {
     destination: operand.Destination,
-    buffer: ids.StorageBufferId,
+    buffer: BufferReference,
     byte_offset: operand.Source,
     immediate_offset: u32 = 0,
 };
 
 pub const StoreBuffer = struct {
-    buffer: ids.StorageBufferId,
+    buffer: BufferReference,
     byte_offset: operand.Source,
     immediate_offset: u32 = 0,
     source: operand.Source,

src/intel/compiler/ir/printer.zig

@@ -120,7 +120,7 @@ fn writeOperation(program: *const program_ir.Program, writer: *std.Io.Writer, ex
             try writer.writeAll("load_buffer ");
             try writeDestination(program, writer, execution_size, op.destination);
             try writer.writeAll(", ");
-            try writeStorageBufferRef(program, writer, op.buffer);
+            try writeBufferReference(program, writer, op.buffer);
             try writer.writeAll(", ");
             try writeSource(program, writer, execution_size, op.byte_offset);
             if (op.immediate_offset != 0)
@@ -128,7 +128,7 @@ fn writeOperation(program: *const program_ir.Program, writer: *std.Io.Writer, ex
         },
         .store_buffer => |op| {
             try writer.writeAll("store_buffer ");
-            try writeStorageBufferRef(program, writer, op.buffer);
+            try writeBufferReference(program, writer, op.buffer);
             try writer.writeAll(", ");
             try writeSource(program, writer, execution_size, op.byte_offset);
             if (op.immediate_offset != 0)
@@ -356,6 +356,13 @@ fn writeFlagRef(program: *const program_ir.Program, writer: *std.Io.Writer, flag
     }
 }
 
+fn writeBufferReference(program: *const program_ir.Program, writer: *std.Io.Writer, reference: inst_ir.BufferReference) !void {
+    switch (reference) {
+        .logical => |buffer| try writeStorageBufferRef(program, writer, buffer),
+        .binding_table => |index| try writer.print("bti({d})", .{index}),
+    }
+}
+
 fn writeStorageBufferRef(program: *const program_ir.Program, writer: *std.Io.Writer, buffer_id: ids.StorageBufferId) !void {
     const buffer = program.storage_buffers.get(buffer_id);
     try writeNamedRef(writer, if (buffer) |value| value.name else null, "buffer", buffer_id.index(), '@');

src/intel/compiler/ir/validator.zig

@@ -19,6 +19,7 @@ pub const Error = error{
     InvalidDestination,
     InvalidImmediateType,
     InvalidStorageBuffer,
+    InvalidBufferReference,
     InvalidGlobalInvocationId,
     InvalidBufferAccess,
     InvalidWorkgroupSize,
@@ -136,20 +137,14 @@ fn validateInstruction(program: *const program_ir.Program, inst: instruction.Ins
                 return Error.InvalidGlobalInvocationId;
         },
         .load_buffer => |op| {
-            if (program.properties.resources_lowered)
-                return Error.UnloweredResource;
-            if (!program.storage_buffers.isLive(op.buffer))
-                return Error.InvalidStorageBuffer;
+            try validateBufferReference(program, op.buffer);
             try validateDestination(program, op.destination);
             try validateBufferOffset(program, op.byte_offset);
             if (!op.destination.type.isInitialTargetType())
                 return Error.InvalidBufferAccess;
         },
         .store_buffer => |op| {
-            if (program.properties.resources_lowered)
-                return Error.UnloweredResource;
-            if (!program.storage_buffers.isLive(op.buffer))
-                return Error.InvalidStorageBuffer;
+            try validateBufferReference(program, op.buffer);
             try validateBufferOffset(program, op.byte_offset);
             try validateSource(program, op.source);
             if (!op.source.type.isInitialTargetType())
@@ -180,6 +175,19 @@ fn validateInstruction(program: *const program_ir.Program, inst: instruction.Ins
     }
 }
 
+fn validateBufferReference(program: *const program_ir.Program, reference: instruction.BufferReference) Error!void {
+    switch (reference) {
+        .logical => |buffer| {
+            if (program.properties.resources_lowered)
+                return Error.UnloweredResource;
+            if (!program.storage_buffers.isLive(buffer))
+                return Error.InvalidStorageBuffer;
+        },
+        .binding_table => if (!program.properties.resources_lowered)
+            return Error.InvalidBufferReference,
+    }
+}
+
 fn validateBufferOffset(program: *const program_ir.Program, source: operand.Source) Error!void {
     try validateSource(program, source);
     if (source.type != .u32)
@@ -408,7 +416,7 @@ test "[ir] validator checks compute system values and resources" {
     const buffer_load_id = try builder.appendInstruction(entry, .simd8, null, .{
         .load_buffer = .{
             .destination = .{ .register = .{ .virtual = register }, .type = .u32 },
-            .buffer = buffer,
+            .buffer = .{ .logical = buffer },
             .byte_offset = .{
                 .register = .{ .immediate = .{ .u32 = 0 } },
                 .type = .u32,
@@ -427,10 +435,12 @@ test "[ir] validator checks compute system values and resources" {
     try std.testing.expectError(Error.UnloweredSystemValue, validate(&program));
     program.properties.system_values_lowered = false;
 
-    program.instructions.getMut(buffer_load_id).?.operation.load_buffer.buffer = ids.StorageBufferId.fromIndex(99);
+    program.instructions.getMut(buffer_load_id).?.operation.load_buffer.buffer = .{ .logical = ids.StorageBufferId.fromIndex(99) };
     try std.testing.expectError(Error.InvalidStorageBuffer, validate(&program));
-    program.instructions.getMut(buffer_load_id).?.operation.load_buffer.buffer = buffer;
+    program.instructions.getMut(buffer_load_id).?.operation.load_buffer.buffer = .{ .logical = buffer };
 
     program.properties.resources_lowered = true;
     try std.testing.expectError(Error.UnloweredResource, validate(&program));
+    program.instructions.getMut(buffer_load_id).?.operation.load_buffer.buffer = .{ .binding_table = 0 };
+    try validate(&program);
 }

src/intel/compiler/lower/common_ir.zig

@@ -731,7 +731,7 @@ const LoweringState = struct {
             try self.appendInstruction(block_id, null, .{
                 .load_buffer = .{
                     .destination = try destinationFromSource(result_component),
-                    .buffer = buffer,
+                    .buffer = .{ .logical = buffer },
                     .byte_offset = byte_offset,
                     .immediate_offset = @intCast(component_index * result_component.type.sizeBytes()),
                 },
@@ -749,7 +749,7 @@ const LoweringState = struct {
         for (source_components, 0..) |source_component, component_index| {
             try self.appendInstruction(block_id, null, .{
                 .store_buffer = .{
-                    .buffer = buffer,
+                    .buffer = .{ .logical = buffer },
                     .byte_offset = byte_offset,
                     .immediate_offset = @intCast(component_index * source_component.type.sizeBytes()),
                     .source = source_component,

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

@@ -1,5 +1,9 @@
 const std = @import("std");
 
+pub const resource_layout = @import("resource_layout.zig");
+pub const resource_lowering = @import("resource_lowering.zig");
+pub const ResourceLayout = resource_layout.Layout;
+
 pub const Error = error{UnsupportedWorkgroupSize};
 
 pub fn validateWorkgroupSize(size: [3]u32) Error!void {

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

@@ -0,0 +1,141 @@
+const std = @import("std");
+const ids = @import("../../../ir/id.zig");
+const program_ir = @import("../../../ir/program.zig");
+
+pub const max_storage_buffers: usize = 4;
+
+pub const Error = std.mem.Allocator.Error || error{
+    TooManyStorageBuffers,
+};
+
+pub const Binding = struct {
+    set: u32,
+    binding: u32,
+    binding_table_index: u8,
+};
+
+const Candidate = struct {
+    resource: ids.StorageBufferId,
+    set: u32,
+    binding: u32,
+};
+
+pub const Layout = struct {
+    bindings: []Binding,
+    resource_indices: []?u8,
+
+    pub fn init(allocator: std.mem.Allocator, program: *const program_ir.Program) Error!Layout {
+        var candidates: std.ArrayList(Candidate) = .empty;
+        defer candidates.deinit(allocator);
+
+        for (program.storage_buffers.entries.items, 0..) |entry, index| {
+            const buffer = entry orelse continue;
+            try candidates.append(allocator, .{
+                .resource = ids.StorageBufferId.fromIndex(index),
+                .set = buffer.set,
+                .binding = buffer.binding,
+            });
+        }
+        std.mem.sort(Candidate, candidates.items, {}, lessThan);
+
+        var unique_count: usize = 0;
+        for (candidates.items, 0..) |candidate, index| {
+            if (index == 0 or candidate.set != candidates.items[index - 1].set or candidate.binding != candidates.items[index - 1].binding)
+                unique_count += 1;
+        }
+        if (unique_count > max_storage_buffers)
+            return Error.TooManyStorageBuffers;
+
+        const bindings = try allocator.alloc(Binding, unique_count);
+        errdefer allocator.free(bindings);
+        const resource_indices = try allocator.alloc(?u8, program.storage_buffers.entries.items.len);
+        errdefer allocator.free(resource_indices);
+        @memset(resource_indices, null);
+
+        var binding_index: usize = 0;
+        for (candidates.items, 0..) |candidate, index| {
+            if (index == 0 or candidate.set != candidates.items[index - 1].set or candidate.binding != candidates.items[index - 1].binding) {
+                bindings[binding_index] = .{
+                    .set = candidate.set,
+                    .binding = candidate.binding,
+                    .binding_table_index = @intCast(binding_index),
+                };
+                binding_index += 1;
+            }
+            resource_indices[candidate.resource.index()] = @intCast(binding_index - 1);
+        }
+        std.debug.assert(binding_index == bindings.len);
+
+        return .{
+            .bindings = bindings,
+            .resource_indices = resource_indices,
+        };
+    }
+
+    pub fn deinit(self: *Layout, allocator: std.mem.Allocator) void {
+        allocator.free(self.bindings);
+        allocator.free(self.resource_indices);
+        self.* = undefined;
+    }
+
+    pub fn bindingTableIndex(self: *const Layout, resource: ids.StorageBufferId) ?u8 {
+        if (resource.index() >= self.resource_indices.len)
+            return null;
+        return self.resource_indices[resource.index()];
+    }
+};
+
+fn lessThan(_: void, lhs: Candidate, rhs: Candidate) bool {
+    if (lhs.set != rhs.set)
+        return lhs.set < rhs.set;
+    if (lhs.binding != rhs.binding)
+        return lhs.binding < rhs.binding;
+    return lhs.resource.index() < rhs.resource.index();
+}
+
+test "[gen9] compute resource layout: assign stable binding-table indices" {
+    const device = @import("../../../device.zig");
+
+    const device_info: device.DeviceInfo = .{
+        .generation = .gen9,
+        .platform = .skylake,
+        .pci_device_id = 0x1912,
+        .grf_count = 128,
+    };
+    var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, device_info, .simd8);
+    defer program.deinit();
+
+    const third = try program.addStorageBuffer(.{ .set = 2, .binding = 7 });
+    const first = try program.addStorageBuffer(.{ .set = 0, .binding = 3 });
+    const alias = try program.addStorageBuffer(.{ .set = 0, .binding = 3 });
+    const second = try program.addStorageBuffer(.{ .set = 1, .binding = 0 });
+
+    var layout = try Layout.init(std.testing.allocator, &program);
+    defer layout.deinit(std.testing.allocator);
+
+    try std.testing.expectEqual(@as(usize, 3), layout.bindings.len);
+    try std.testing.expectEqual(Binding{ .set = 0, .binding = 3, .binding_table_index = 0 }, layout.bindings[0]);
+    try std.testing.expectEqual(Binding{ .set = 1, .binding = 0, .binding_table_index = 1 }, layout.bindings[1]);
+    try std.testing.expectEqual(Binding{ .set = 2, .binding = 7, .binding_table_index = 2 }, layout.bindings[2]);
+    try std.testing.expectEqual(@as(?u8, 0), layout.bindingTableIndex(first));
+    try std.testing.expectEqual(@as(?u8, 0), layout.bindingTableIndex(alias));
+    try std.testing.expectEqual(@as(?u8, 1), layout.bindingTableIndex(second));
+    try std.testing.expectEqual(@as(?u8, 2), layout.bindingTableIndex(third));
+}
+
+test "[gen9] compute resource layout: enforce advertised storage-buffer limit" {
+    const device = @import("../../../device.zig");
+
+    const device_info: device.DeviceInfo = .{
+        .generation = .gen9,
+        .platform = .skylake,
+        .pci_device_id = 0x1912,
+        .grf_count = 128,
+    };
+    var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, device_info, .simd8);
+    defer program.deinit();
+    for (0..max_storage_buffers + 1) |binding|
+        _ = try program.addStorageBuffer(.{ .set = 0, .binding = @intCast(binding) });
+
+    try std.testing.expectError(Error.TooManyStorageBuffers, Layout.init(std.testing.allocator, &program));
+}

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

@@ -0,0 +1,127 @@
+const instruction = @import("../../../ir/instruction.zig");
+const program_ir = @import("../../../ir/program.zig");
+const validator = @import("../../../ir/validator.zig");
+const resource_layout = @import("resource_layout.zig");
+
+pub const Error = error{
+    InvalidProgram,
+    InvalidResourceLayout,
+};
+
+pub fn run(program: *program_ir.Program, layout: *const resource_layout.Layout) Error!void {
+    validator.validate(program) catch return Error.InvalidProgram;
+    if (program.properties.resources_lowered)
+        return;
+    if (layout.resource_indices.len != program.storage_buffers.entries.items.len)
+        return Error.InvalidResourceLayout;
+
+    for (program.blocks.entries.items) |block_entry| {
+        const block = block_entry orelse continue;
+        for (block.instructions.items) |instruction_id| {
+            const inst = program.instructions.get(instruction_id) orelse return Error.InvalidProgram;
+            const reference = bufferReference(inst.operation) orelse continue;
+            const resource = switch (reference) {
+                .logical => |value| value,
+                .binding_table => return Error.InvalidProgram,
+            };
+            const binding_table_index = layout.bindingTableIndex(resource) orelse return Error.InvalidResourceLayout;
+            if (binding_table_index >= layout.bindings.len)
+                return Error.InvalidResourceLayout;
+            const buffer = program.storage_buffers.get(resource) orelse return Error.InvalidProgram;
+            const binding = layout.bindings[binding_table_index];
+            if (binding.binding_table_index != binding_table_index or binding.set != buffer.set or binding.binding != buffer.binding)
+                return Error.InvalidResourceLayout;
+        }
+    }
+
+    for (program.blocks.entries.items) |block_entry| {
+        const block = block_entry orelse continue;
+        for (block.instructions.items) |instruction_id| {
+            const inst = program.instructions.getMut(instruction_id) orelse unreachable;
+            const reference = bufferReferenceMut(&inst.operation) orelse continue;
+            const resource = reference.logical;
+            const binding_table_index = layout.bindingTableIndex(resource).?;
+            reference.* = .{ .binding_table = binding_table_index };
+        }
+    }
+
+    program.properties.resources_lowered = true;
+    validator.validate(program) catch return Error.InvalidProgram;
+}
+
+fn bufferReference(operation: instruction.Operation) ?instruction.BufferReference {
+    return switch (operation) {
+        .load_buffer => |op| op.buffer,
+        .store_buffer => |op| op.buffer,
+        else => null,
+    };
+}
+
+fn bufferReferenceMut(operation: *instruction.Operation) ?*instruction.BufferReference {
+    return switch (operation.*) {
+        .load_buffer => |*op| &op.buffer,
+        .store_buffer => |*op| &op.buffer,
+        else => null,
+    };
+}
+
+test "[gen9] compute resource lowering: resolve logical buffers" {
+    const std = @import("std");
+    const Builder = @import("../../../ir/Builder.zig");
+    const device = @import("../../../device.zig");
+    const operand = @import("../../../ir/operand.zig");
+    const printer = @import("../../../ir/printer.zig");
+
+    const device_info: device.DeviceInfo = .{
+        .generation = .gen9,
+        .platform = .skylake,
+        .pci_device_id = 0x1912,
+        .grf_count = 128,
+    };
+    var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, device_info, .simd8);
+    defer program.deinit();
+    var builder = Builder.init(&program);
+
+    const value = try builder.addVirtualRegister(.{
+        .size_bytes = 32,
+        .alignment_bytes = 32,
+        .element_type = .u32,
+        .lane_count = 8,
+        .class = .temporary,
+    });
+    const buffer = try builder.addStorageBuffer(.{ .set = 1, .binding = 3, .name = "storage" });
+    const entry = try builder.addBlock("entry");
+    const store_id = try builder.appendInstruction(entry, .simd8, null, .{
+        .store_buffer = .{
+            .buffer = .{ .logical = buffer },
+            .byte_offset = .{
+                .register = .{ .immediate = .{ .u32 = 0 } },
+                .type = .u32,
+                .region = operand.Region.broadcast(),
+            },
+            .source = .{
+                .register = .{ .virtual = value },
+                .type = .u32,
+                .region = operand.Region.contiguous(.simd8),
+            },
+        },
+    });
+    try builder.setTerminator(entry, .end_thread);
+
+    var layout = try resource_layout.Layout.init(std.testing.allocator, &program);
+    defer layout.deinit(std.testing.allocator);
+    layout.bindings[0].binding = 4;
+    try std.testing.expectError(Error.InvalidResourceLayout, run(&program, &layout));
+    try std.testing.expect(!program.properties.resources_lowered);
+    try std.testing.expect(program.instructions.get(store_id).?.operation.store_buffer.buffer == .logical);
+    layout.bindings[0].binding = 3;
+
+    try run(&program, &layout);
+    try validator.validate(&program);
+
+    try std.testing.expect(program.properties.resources_lowered);
+    try std.testing.expectEqual(@as(u8, 0), program.instructions.get(store_id).?.operation.store_buffer.buffer.binding_table);
+    const text = try printer.allocPrint(std.testing.allocator, &program);
+    defer std.testing.allocator.free(text);
+    try std.testing.expect(std.mem.indexOf(u8, text, "store_buffer bti(0), 0:u32") != null);
+}

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

@@ -8,6 +8,8 @@ pub const compute = @import("compute/compute.zig");
 pub const validator = @import("validator.zig");
 
 pub const Options = common_ir.Options;
+pub const ResourceLoweringError = compute.resource_lowering.Error;
+
 pub const Error = common_ir.Error || compute.Error || error{
     UnsupportedGeneration,
     UnsupportedStage,
@@ -38,6 +40,11 @@ pub fn lower(
     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" {
     var module = try shader_ir.parser.parseString(std.testing.allocator,
         \\shader compute @main

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

@@ -12,6 +12,7 @@ pub const Error = shared.Error || compute.Error || error{
     UnsupportedExecutionSize,
     UnsupportedDataType,
     InvalidPhysicalFlag,
+    InvalidBindingTableIndex,
     InvalidPayloadLayout,
 };
 
@@ -54,10 +55,12 @@ fn validateInstruction(inst: instruction.Instruction) Error!void {
     switch (inst.operation) {
         .load_global_invocation_id => |op| try validateDestination(op.destination),
         .load_buffer => |op| {
+            try validateBufferReference(op.buffer);
             try validateDestination(op.destination);
             try validateSource(op.byte_offset);
         },
         .store_buffer => |op| {
+            try validateBufferReference(op.buffer);
             try validateSource(op.byte_offset);
             try validateSource(op.source);
         },
@@ -88,6 +91,14 @@ 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)
+            return Error.InvalidBindingTableIndex,
+    }
+}
+
 fn validateSource(source: operand.Source) Error!void {
     try validateType(source.type);
     switch (source.register) {

src/intel/compiler/targets/targets.zig

@@ -6,6 +6,10 @@ const common_ir = @import("../lower/common_ir.zig");
 
 pub const gen9 = @import("gen9/gen9.zig");
 
+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};
 
@@ -21,6 +25,20 @@ pub fn lower(
     };
 }
 
+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),