[Flint] moving ir gen specific code to separate file

.gitea/workflows/Build.yml

@@ -55,7 +55,7 @@ jobs:
       - name: ZLint pass
         run: |
           curl -fsSL https://raw.githubusercontent.com/DonIsaac/zlint/refs/heads/main/tasks/install.sh | bash
-          zlint
+          zlint --deny-warnings
 
       - name: Building Ape
         run: zig build ape --release=safe

src/compiler/ir/ir.zig

@@ -57,13 +57,3 @@ pub const types = @import("type.zig");
 pub const validator = @import("validator/validator.zig");
 pub const value = @import("value.zig");
 pub const visitor = @import("visitor.zig");
-
-test {
-    _ = Builder;
-    _ = Rewriter;
-    _ = inline_all_functions;
-    _ = module;
-    _ = parser;
-    _ = transformer_manager;
-    _ = validator;
-}

src/intel/FlintPipeline.zig

@@ -74,7 +74,12 @@ 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;
 
@@ -93,7 +98,12 @@ fn compileStages(allocator: std.mem.Allocator, infos: []const vk.PipelineShaderS
     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,22 +139,30 @@ 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!?compiler.Program {
+fn lowerToFlint(
+    allocator: std.mem.Allocator,
+    module: *base.ShaderModule.IrModule,
+    device_info: ?compiler.device.DeviceInfo,
+) VkError!?compiler.Program {
     const target = device_info orelse return null;
-    return compiler.lower.lower(allocator, module, target, .{}) catch |err| switch (err) {
+    const program = compiler.targets.lower(allocator, module, target, .{}) catch |err| switch (err) {
-        error.OutOfMemory => VkError.OutOfHostMemory,
+        error.OutOfMemory => return VkError.OutOfHostMemory,
         error.UnsupportedGeneration,
         error.UnsupportedStage,
         error.UnsupportedDispatchWidth,
+        error.UnsupportedGrfSize,
+        error.UnsupportedWorkgroupSize,
+        error.MissingWorkgroupSize,
         error.UnsupportedType,
         error.UnsupportedOperation,
         error.UnsupportedTerminator,
-        => null,
+        => return null,
         else => {
             std.log.scoped(.FlintPipeline).err("Flint shader lowering failed: {s}", .{@errorName(err)});
             return VkError.ValidationFailed;
         },
     };
+    return program;
 }
 
 fn compilerDeviceInfo(device: *const base.Device) ?compiler.device.DeviceInfo {
@@ -208,3 +226,53 @@ pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void {
     deinitStages(self.artifact_allocator.allocator(), self.stages);
     allocator.destroy(self);
 }
+
+test "Flint pipeline: lower common compute IR" {
+    const device_info: compiler.device.DeviceInfo = .{
+        .generation = .gen9,
+        .platform = .skylake,
+        .pci_device_id = 0x1912,
+        .grf_count = 128,
+    };
+    var module = shader_ir.ir.module.Module.init(std.testing.allocator, .compute);
+    defer module.deinit();
+    module.execution_modes.workgroup_size = .{ 1, 1, 1 };
+    var builder = shader_ir.ir.Builder.init(&module);
+
+    const void_type = try builder.internType(.void);
+    const u32_type = try builder.internType(.{ .integer = .{ .bits = 32, .signedness = .unsigned } });
+    const vec3_type = try builder.internType(.{ .vector = .{ .element_type = u32_type, .length = 3 } });
+    const global_id = try builder.addInterfaceVariable(vec3_type, .input, .{ .builtin = .global_invocation_id }, "global_id");
+    const storage = try builder.addResource(u32_type, .storage_buffer, 0, 2, "storage");
+    const zero = try builder.internConstant(u32_type, .{ .integer_bits = 0 });
+    const main = try builder.addFunction(void_type, "main");
+    builder.setEntryPoint(main);
+    const entry = try builder.addBlock(main, "entry");
+    const id = (try builder.appendInstruction(entry, vec3_type, .{
+        .load_interface = .{ .variable = global_id },
+    }, "id")).?;
+    const x = (try builder.appendInstruction(entry, u32_type, .{
+        .composite_extract = .{ .composite = id, .indices = &.{0} },
+    }, "x")).?;
+    _ = try builder.appendInstruction(entry, null, .{
+        .store_buffer = .{ .resource = storage, .byte_offset = zero, .value = x },
+    }, null);
+    try builder.setTerminator(entry, .return_void);
+
+    var program = (try lowerToFlint(std.testing.allocator, &module, device_info)).?;
+    defer program.deinit();
+
+    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.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);
+
+    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);
+}

src/intel/compiler/compiler.zig

@@ -4,6 +4,7 @@
 pub const device = @import("device.zig");
 pub const ir = @import("ir/ir.zig");
 pub const lower = @import("lower/lower.zig");
+pub const targets = @import("targets/targets.zig");
 
 pub const Builder = ir.Builder;
 pub const id = ir.id;
@@ -15,28 +16,10 @@ pub const pseudo = ir.pseudo;
 pub const validator = ir.validator;
 
 pub const Program = ir.Program;
-pub const Stage = ir.Stage;
 
 const std = @import("std");
 
-test "[ir] basic shader" {
+test "[ir] basic compute shader" {
-    // ; Flint program:
-    // ;   .stage: vertex
-    // ;   .generation: gen9
-    // ;   .platform: skylake
-    // ;   .dispatch_width: simd8
-    //
-    // %position: vgrf f32[8] = class(varying), size(32), alignment(32), spillable
-    // %urb_payload: vgrf u32[16] = class(payload), size(64), alignment(32)
-    //
-    // .entry:
-    //     [simd8] load_input %position:f32, location(0), component(0)
-    //     [simd8] multiply %position:f32, %position:f32, 1:f32
-    //     [simd8] mov %position:f32, %position:f32[byte=4, broadcast]
-    //     [simd8] store_output builtin(position), component(0), %position:f32
-    //     [simd8] send urb_write[offset(0), channels(xyzw), end_of_thread], payload(%urb_payload[2])
-    //     end_thread
-
     const device_info: device.DeviceInfo = .{
         .generation = .gen9,
         .platform = .skylake,
@@ -44,128 +27,52 @@ test "[ir] basic shader" {
         .grf_count = 128,
     };
 
-    var shader = Program.init(std.testing.allocator, .vertex, device_info, .simd8);
+    var shader = Program.init(std.testing.allocator, .{ 1, 1, 1 }, device_info, .simd8);
     defer shader.deinit();
     var builder = Builder.init(&shader);
 
-    const position = try builder.addVirtualRegister(.{
+    const value = try builder.addVirtualRegister(.{
         .size_bytes = 32,
         .alignment_bytes = 32,
-        .element_type = .f32,
-        .lane_count = 8,
-        .class = .varying,
-        .name = "position",
-    });
-    const urb_payload = try builder.addVirtualRegister(.{
-        .size_bytes = 64,
-        .alignment_bytes = 32,
         .element_type = .u32,
-        .lane_count = 16,
+        .lane_count = 8,
-        .class = .payload,
+        .class = .temporary,
-        .spillable = false,
+        .name = "value",
-        .name = "urb_payload",
     });
+    const storage = try builder.addStorageBuffer(.{ .set = 0, .binding = 1, .name = "storage" });
     const entry = try builder.addBlock("entry");
     try builder.setEntryBlock(entry);
 
     _ = try builder.appendInstruction(entry, .simd8, null, .{
-        .load_input = .{
+        .load_global_invocation_id = .{
-            .destination = .{
+            .destination = .{ .register = .{ .virtual = value }, .type = .u32 },
-                .register = .{ .virtual = position },
+            .component = 0,
-                .type = .f32,
-            },
-            .semantic = .{
-                .location = .{
-                    .location = 0,
-                },
-            },
         },
     });
     _ = try builder.appendInstruction(entry, .simd8, null, .{
-        .binary = .{
+        .store_buffer = .{
-            .opcode = .multiply,
+            .buffer = storage,
-            .destination = .{
+            .byte_offset = .{
-                .register = .{ .virtual = position },
+                .register = .{ .immediate = .{ .u32 = 0 } },
-                .type = .f32,
+                .type = .u32,
-            },
-            .lhs = .{
-                .register = .{ .virtual = position },
-                .type = .f32,
-                .region = operand.Region.contiguous(.simd8),
-            },
-            .rhs = .{
-                .register = .{
-                    .immediate = .{ .f32 = 1.0 },
-                },
-                .type = .f32,
                 .region = operand.Region.broadcast(),
             },
-        },
-    });
-    _ = try builder.appendInstruction(entry, .simd8, null, .{
-        .move = .{
-            .destination = .{
-                .register = .{ .virtual = position },
-                .type = .f32,
-            },
             .source = .{
-                .register = .{ .virtual = position },
+                .register = .{ .virtual = value },
-                .type = .f32,
+                .type = .u32,
-                .region = .{
-                    .byte_offset = 4,
-                    .vertical_stride = 0,
-                    .width = 1,
-                    .horizontal_stride = 0,
-                },
-            },
-        },
-    });
-    _ = try builder.appendInstruction(entry, .simd8, null, .{
-        .store_output = .{
-            .semantic = .{
-                .builtin = .{ .builtin = .position },
-            },
-            .source = .{
-                .register = .{ .virtual = position },
-                .type = .f32,
                 .region = operand.Region.contiguous(.simd8),
             },
         },
     });
-    _ = try builder.appendInstruction(entry, .simd8, null, .{
-        .send = .{
-            .message = .{
-                .urb_write = .{
-                    .offset = 0,
-                    .end_of_thread = true,
-                },
-            },
-            .payload = .{
-                .base = .{ .virtual = urb_payload },
-                .register_count = 2,
-            },
-        },
-    });
     try builder.setTerminator(entry, .end_thread);
-
-    shader.properties.instructions_selected = true;
     try validator.validate(&shader);
 
-    try std.testing.expectEqual(entry, shader.entry_block.?);
-    try std.testing.expect(shader.properties.instructions_selected);
-
     const text = try printer.allocPrint(std.testing.allocator, &shader);
     defer std.testing.allocator.free(text);
-
+    try std.testing.expect(std.mem.indexOf(u8, text, "Flint compute program") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "Flint program") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "@storage = storage_buffer[set(0), binding(1)]") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "vertex") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "load_global_invocation_id %value:u32, component(0)") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "gen9") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "store_buffer @storage, 0:u32, %value:u32") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "skylake") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "simd8") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "%position: vgrf f32[8]") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "[simd8] multiply %position:f32, %position:f32, 1:f32") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "[simd8] mov %position:f32, %position:f32[byte=4, broadcast]") != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "send urb_write[offset(0), channels(xyzw), end_of_thread], payload(%urb_payload[2])") != null);
 }
 
 test "[ir] ID stability after removal" {
@@ -180,8 +87,3 @@ test "[ir] ID stability after removal" {
     try std.testing.expect(store.get(first) == null);
     try std.testing.expectEqualStrings("second", store.get(second).?.name.?);
 }
-
-test {
-    _ = lower;
-    _ = lower.vertex_abi;
-}

src/intel/compiler/ir/Builder.zig

@@ -29,6 +29,10 @@ pub fn addVirtualFlag(self: *Self, flag: operand.VirtualFlag) Error!ids.VirtualF
     return self.program.addVirtualFlag(flag);
 }
 
+pub fn addStorageBuffer(self: *Self, buffer: program_ir.StorageBuffer) Error!ids.StorageBufferId {
+    return self.program.addStorageBuffer(buffer);
+}
+
 pub fn addBlock(self: *Self, name: ?[]const u8) Error!ids.BlockId {
     return self.program.addBlock(name);
 }
@@ -132,7 +136,7 @@ test "[ir] Builder: construction and ordered insertion" {
         .grf_count = 128,
     };
 
-    var program = program_ir.Program.init(std.testing.allocator, .vertex, device_info, .simd8);
+    var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, device_info, .simd8);
     defer program.deinit();
     var builder = Self.init(&program);
 

src/intel/compiler/ir/id.zig

@@ -4,11 +4,13 @@ pub const BlockTag = opaque {};
 pub const InstructionTag = opaque {};
 pub const VirtualRegisterTag = opaque {};
 pub const VirtualFlagTag = opaque {};
+pub const StorageBufferTag = opaque {};
 
 pub const BlockId = shared_ids.Id(BlockTag);
 pub const InstructionId = shared_ids.Id(InstructionTag);
 pub const VirtualRegisterId = shared_ids.Id(VirtualRegisterTag);
 pub const VirtualFlagId = shared_ids.Id(VirtualFlagTag);
+pub const StorageBufferId = shared_ids.Id(StorageBufferTag);
 
 pub const Id = shared_ids.Id;
 pub const Store = shared_ids.Store;

src/intel/compiler/ir/instruction.zig

@@ -4,30 +4,22 @@ const ids = @import("id.zig");
 const operand = @import("operand.zig");
 const pseudo = @import("pseudo.zig");
 
-pub const Builtin = enum {
+pub const LoadGlobalInvocationId = struct {
-    position,
-    vertex_index,
-    instance_index,
-};
-
-pub const InterfaceSemantic = union(enum) {
-    location: struct {
-        location: u32,
-        component: u8 = 0,
-    },
-    builtin: struct {
-        builtin: Builtin,
-        component: u8 = 0,
-    },
-};
-
-pub const LoadInput = struct {
     destination: operand.Destination,
-    semantic: InterfaceSemantic,
+    component: u8,
 };
 
-pub const StoreOutput = struct {
+pub const LoadBuffer = struct {
-    semantic: InterfaceSemantic,
+    destination: operand.Destination,
+    buffer: ids.StorageBufferId,
+    byte_offset: operand.Source,
+    immediate_offset: u32 = 0,
+};
+
+pub const StoreBuffer = struct {
+    buffer: ids.StorageBufferId,
+    byte_offset: operand.Source,
+    immediate_offset: u32 = 0,
     source: operand.Source,
 };
 
@@ -69,36 +61,13 @@ pub const Compare = struct {
     rhs: operand.Source,
 };
 
-pub const ChannelMask = packed struct(u4) {
-    x: bool = true,
-    y: bool = true,
-    z: bool = true,
-    w: bool = true,
-};
-
-pub const UrbWrite = struct {
-    offset: u16,
-    channels: ChannelMask = .{},
-    end_of_thread: bool = false,
-};
-
-pub const Message = union(enum) {
-    urb_write: UrbWrite,
-};
-
-pub const Send = struct {
-    message: Message,
-    payload: operand.RegisterSpan,
-    response: ?operand.RegisterSpan = null,
-};
-
 pub const Operation = union(enum) {
-    load_input: LoadInput,
+    load_global_invocation_id: LoadGlobalInvocationId,
-    store_output: StoreOutput,
+    load_buffer: LoadBuffer,
+    store_buffer: StoreBuffer,
     move: Move,
     binary: Binary,
     compare: Compare,
-    send: Send,
     parallel_copy: pseudo.ParallelCopy,
 };
 

src/intel/compiler/ir/ir.zig

@@ -8,4 +8,3 @@ pub const pseudo = @import("pseudo.zig");
 pub const validator = @import("validator.zig");
 
 pub const Program = program.Program;
-pub const Stage = program.Stage;

src/intel/compiler/ir/operand.zig

@@ -33,7 +33,6 @@ pub const DataType = enum {
 
 pub const RegisterClass = enum {
     uniform,
-    varying,
     payload,
     response,
     temporary,

src/intel/compiler/ir/printer.zig

@@ -9,12 +9,20 @@ const pseudo = @import("pseudo.zig");
 const indent = "    ";
 
 pub fn write(program: *const program_ir.Program, writer: *std.Io.Writer) std.Io.Writer.Error!void {
-    try writer.writeAll("; Flint program:\n");
+    try writer.writeAll("; Flint compute program:\n");
-    try writer.print(";   .stage: {t}\n", .{program.stage});
+    try writer.print(";   .workgroup_size: [{d}, {d}, {d}]\n", .{ program.workgroup_size[0], program.workgroup_size[1], program.workgroup_size[2] });
     try writer.print(";   .generation: {t}\n", .{program.device_info.generation});
     try writer.print(";   .platform: {t}\n", .{program.device_info.platform});
     try writer.print(";   .dispatch_width: {t}\n\n", .{program.dispatch_width});
 
+    for (program.storage_buffers.entries.items, 0..) |entry, index| {
+        const buffer = entry orelse continue;
+        try writeStorageBufferRef(program, writer, ids.StorageBufferId.fromIndex(index));
+        try writer.print(" = storage_buffer[set({d}), binding({d})]\n", .{ buffer.set, buffer.binding });
+    }
+    if (program.storage_buffers.entries.items.len != 0)
+        try writer.writeByte('\n');
+
     for (program.virtual_registers.entries.items, 0..) |entry, index| {
         const register = entry orelse continue;
         try writeVirtualRegisterRef(program, writer, ids.VirtualRegisterId.fromIndex(index));
@@ -103,15 +111,28 @@ fn writeInstruction(program: *const program_ir.Program, writer: *std.Io.Writer,
 
 fn writeOperation(program: *const program_ir.Program, writer: *std.Io.Writer, execution_size: device.ExecutionSize, operation: inst_ir.Operation) !void {
     switch (operation) {
-        .load_input => |op| {
+        .load_global_invocation_id => |op| {
-            try writer.writeAll("load_input ");
+            try writer.writeAll("load_global_invocation_id ");
+            try writeDestination(program, writer, execution_size, op.destination);
+            try writer.print(", component({d})", .{op.component});
+        },
+        .load_buffer => |op| {
+            try writer.writeAll("load_buffer ");
             try writeDestination(program, writer, execution_size, op.destination);
             try writer.writeAll(", ");
-            try writeInterfaceSemantic(writer, op.semantic);
+            try writeStorageBufferRef(program, writer, op.buffer);
+            try writer.writeAll(", ");
+            try writeSource(program, writer, execution_size, op.byte_offset);
+            if (op.immediate_offset != 0)
+                try writer.print(", offset({d})", .{op.immediate_offset});
         },
-        .store_output => |op| {
+        .store_buffer => |op| {
-            try writer.writeAll("store_output ");
+            try writer.writeAll("store_buffer ");
-            try writeInterfaceSemantic(writer, op.semantic);
+            try writeStorageBufferRef(program, writer, op.buffer);
+            try writer.writeAll(", ");
+            try writeSource(program, writer, execution_size, op.byte_offset);
+            if (op.immediate_offset != 0)
+                try writer.print(", offset({d})", .{op.immediate_offset});
             try writer.writeAll(", ");
             try writeSource(program, writer, execution_size, op.source);
         },
@@ -138,17 +159,6 @@ fn writeOperation(program: *const program_ir.Program, writer: *std.Io.Writer, ex
             try writeSource(program, writer, execution_size, op.rhs);
         },
         .parallel_copy => |op| try writeParallelCopy(program, writer, execution_size, op),
-        .send => |op| {
-            try writer.writeAll("send ");
-            if (op.response) |response| {
-                try writeRegisterSpan(program, writer, response);
-                try writer.writeAll(", ");
-            }
-            try writeMessage(writer, op.message);
-            try writer.writeAll(", payload(");
-            try writeRegisterSpan(program, writer, op.payload);
-            try writer.writeByte(')');
-        },
     }
 }
 
@@ -346,39 +356,9 @@ fn writeFlagRef(program: *const program_ir.Program, writer: *std.Io.Writer, flag
     }
 }
 
-fn writeRegisterSpan(program: *const program_ir.Program, writer: *std.Io.Writer, span: operand.RegisterSpan) !void {
+fn writeStorageBufferRef(program: *const program_ir.Program, writer: *std.Io.Writer, buffer_id: ids.StorageBufferId) !void {
-    try writeRegister(program, writer, span.base);
+    const buffer = program.storage_buffers.get(buffer_id);
-    const byte_offset = registerByteOffset(span.base);
+    try writeNamedRef(writer, if (buffer) |value| value.name else null, "buffer", buffer_id.index(), '@');
-    if (byte_offset != 0)
-        try writer.print("[byte={d}]", .{byte_offset});
-    try writer.print("[{d}]", .{span.register_count});
-}
-
-fn writeInterfaceSemantic(writer: *std.Io.Writer, semantic: inst_ir.InterfaceSemantic) !void {
-    switch (semantic) {
-        .location => |location| try writer.print("location({d}), component({d})", .{ location.location, location.component }),
-        .builtin => |builtin| try writer.print("builtin({t}), component({d})", .{ builtin.builtin, builtin.component }),
-    }
-}
-
-fn writeMessage(writer: *std.Io.Writer, message: inst_ir.Message) !void {
-    switch (message) {
-        .urb_write => |urb| {
-            try writer.print("urb_write[offset({d}), channels(", .{urb.offset});
-            try writeChannelMask(writer, urb.channels);
-            try writer.writeByte(')');
-            if (urb.end_of_thread)
-                try writer.writeAll(", end_of_thread");
-            try writer.writeByte(']');
-        },
-    }
-}
-
-fn writeChannelMask(writer: *std.Io.Writer, mask: inst_ir.ChannelMask) !void {
-    if (mask.x) try writer.writeByte('x');
-    if (mask.y) try writer.writeByte('y');
-    if (mask.z) try writer.writeByte('z');
-    if (mask.w) try writer.writeByte('w');
 }
 
 fn writeVirtualRegisterRef(program: *const program_ir.Program, writer: *std.Io.Writer, register_id: ids.VirtualRegisterId) !void {

src/intel/compiler/ir/program.zig

@@ -1,19 +1,17 @@
 const std = @import("std");
-const shared_ir = @import("shader_ir").ir.module;
+
 const device = @import("../device.zig");
 const ids = @import("id.zig");
 const instructions = @import("instruction.zig");
 const operand = @import("operand.zig");
 
-pub const Stage = shared_ir.Stage;
-
 pub const Properties = packed struct {
     common_ir_lowered: bool = false,
     instructions_selected: bool = false,
     block_parameters_lowered: bool = false,
     parallel_copies_lowered: bool = false,
 
-    stage_io_lowered: bool = false,
+    system_values_lowered: bool = false,
     resources_lowered: bool = false,
     messages_lowered: bool = false,
     control_flow_lowered: bool = false,
@@ -28,14 +26,14 @@ pub const Properties = packed struct {
     _padding: u19 = 0,
 };
 
-pub const VertexPayload = struct {
+pub const StorageBuffer = struct {
-    first_attribute_grf: operand.PhysicalGrf,
+    set: u32,
-    attribute_grf_count: u16,
+    binding: u32,
+    name: ?[]const u8 = null,
 };
 
 pub const PayloadLayout = struct {
     header_grf: ?operand.PhysicalGrf = null,
-    vertex: ?VertexPayload = null,
 };
 
 pub const ProgramData = struct {
@@ -48,11 +46,12 @@ pub const BlockStore = ids.Store(ids.BlockId, instructions.Block);
 pub const InstructionStore = ids.Store(ids.InstructionId, instructions.Instruction);
 pub const VirtualRegisterStore = ids.Store(ids.VirtualRegisterId, operand.VirtualRegister);
 pub const VirtualFlagStore = ids.Store(ids.VirtualFlagId, operand.VirtualFlag);
+pub const StorageBufferStore = ids.Store(ids.StorageBufferId, StorageBuffer);
 
 pub const Program = struct {
     arena: std.heap.ArenaAllocator,
 
-    stage: Stage,
+    workgroup_size: [3]u32,
     device_info: device.DeviceInfo,
     dispatch_width: device.DispatchWidth,
 
@@ -62,15 +61,16 @@ pub const Program = struct {
     instructions: InstructionStore = .{},
     virtual_registers: VirtualRegisterStore = .{},
     virtual_flags: VirtualFlagStore = .{},
+    storage_buffers: StorageBufferStore = .{},
 
     payload: PayloadLayout = .{},
     program_data: ProgramData = .{},
     properties: Properties = .{},
 
-    pub fn init(backing_allocator: std.mem.Allocator, stage: Stage, device_info: device.DeviceInfo, dispatch_width: device.DispatchWidth) Program {
+    pub fn init(backing_allocator: std.mem.Allocator, workgroup_size: [3]u32, device_info: device.DeviceInfo, dispatch_width: device.DispatchWidth) Program {
         return .{
             .arena = std.heap.ArenaAllocator.init(backing_allocator),
-            .stage = stage,
+            .workgroup_size = workgroup_size,
             .device_info = device_info,
             .dispatch_width = dispatch_width,
         };
@@ -99,6 +99,13 @@ pub const Program = struct {
         return self.virtual_flags.add(self.allocator(), owned);
     }
 
+    pub fn addStorageBuffer(self: *Program, buffer: StorageBuffer) !ids.StorageBufferId {
+        var owned = buffer;
+        if (buffer.name) |name|
+            owned.name = try self.allocator().dupe(u8, name);
+        return self.storage_buffers.add(self.allocator(), owned);
+    }
+
     pub fn addBlock(self: *Program, name: ?[]const u8) !ids.BlockId {
         const owned_name = if (name) |value| try self.allocator().dupe(u8, value) else null;
         const block_id = try self.blocks.add(self.allocator(), .{

src/intel/compiler/ir/pseudo.zig

@@ -48,7 +48,7 @@ test "[ir] pseudo: parallel copy ownership and printing" {
         .grf_count = 128,
     };
 
-    var program = program_ir.Program.init(std.testing.allocator, .vertex, device_info, .simd8);
+    var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, device_info, .simd8);
     defer program.deinit();
     var builder = Builder.init(&program);
 
@@ -134,7 +134,7 @@ test "[ir] pseudo: validator rejects invalid parallel copies" {
         .grf_count = 128,
     };
 
-    var program = program_ir.Program.init(std.testing.allocator, .vertex, device_info, .simd8);
+    var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, device_info, .simd8);
     defer program.deinit();
     var builder = Builder.init(&program);
 

src/intel/compiler/ir/validator.zig

@@ -5,11 +5,6 @@ const program_ir = @import("program.zig");
 const pseudo = @import("pseudo.zig");
 
 pub const Error = error{
-    UnsupportedGeneration,
-    UnsupportedStage,
-    UnsupportedDispatchWidth,
-    UnsupportedExecutionSize,
-    UnsupportedDataType,
     MissingEntryBlock,
     InvalidBlock,
     MissingTerminator,
@@ -17,24 +12,25 @@ pub const Error = error{
     InvalidVirtualRegister,
     InvalidVirtualFlag,
     InvalidPhysicalRegister,
-    InvalidPhysicalFlag,
     InvalidRegisterSize,
     InvalidRegisterAlignment,
     InvalidLaneCount,
     InvalidRegion,
     InvalidDestination,
     InvalidImmediateType,
-    InvalidRegisterSpan,
+    InvalidStorageBuffer,
+    InvalidGlobalInvocationId,
+    InvalidBufferAccess,
+    InvalidWorkgroupSize,
     EmptyParallelCopy,
     InvalidParallelCopyDestination,
     ParallelCopyTypeMismatch,
     DuplicateParallelCopyDestination,
     PredicatedParallelCopy,
     UnloweredParallelCopy,
-    UnloweredStageIo,
+    UnloweredSystemValue,
-    UnloweredMessage,
+    UnloweredResource,
-    InvalidInterfaceSemantic,
+    InvalidPayloadLayout,
-    InvalidMessage,
     EntryBlockHasParameters,
     DuplicateBlockParameter,
     EdgeArgumentCountMismatch,
@@ -44,17 +40,15 @@ pub const Error = error{
 };
 
 pub fn validate(program: *const program_ir.Program) Error!void {
-    if (program.device_info.generation != .gen9)
+    if (program.workgroup_size[0] == 0 or program.workgroup_size[1] == 0 or program.workgroup_size[2] == 0)
-        return Error.UnsupportedGeneration;
+        return Error.InvalidWorkgroupSize;
-    if (program.stage != .vertex)
-        return Error.UnsupportedStage;
-    if (program.dispatch_width != .simd8 or !program.device_info.supportsDispatch(.simd8))
-        return Error.UnsupportedDispatchWidth;
 
     const entry_block = program.entry_block orelse return Error.MissingEntryBlock;
     if (!program.blocks.isLive(entry_block))
         return Error.InvalidBlock;
 
+    try validatePayload(program);
+
     for (program.virtual_registers.entries.items) |entry| {
         const register = entry orelse continue;
         if (register.size_bytes == 0)
@@ -64,7 +58,6 @@ pub fn validate(program: *const program_ir.Program) Error!void {
             return Error.InvalidRegisterAlignment;
         if (register.lane_count == 0)
             return Error.InvalidLaneCount;
-        try validateType(register.element_type);
     }
 
     for (program.blocks.entries.items, 0..) |entry, block_index| {
@@ -92,6 +85,17 @@ pub fn validate(program: *const program_ir.Program) Error!void {
     }
 }
 
+fn validatePayload(program: *const program_ir.Program) Error!void {
+    if (program.program_data.payload_grf_count > program.device_info.grf_count)
+        return Error.InvalidPayloadLayout;
+
+    if (program.payload.header_grf) |header| {
+        try validateRegisterRef(program, .{ .physical_grf = header });
+        if (header.byte_offset != 0)
+            return Error.InvalidPayloadLayout;
+    }
+}
+
 fn validateBlockParameter(program: *const program_ir.Program, block_index: usize, parameter_index: usize, parameter: pseudo.BlockParameter) Error!void {
     switch (parameter) {
         .register => |register_id| if (!program.virtual_registers.isLive(register_id))
@@ -120,26 +124,36 @@ fn blockParametersEqual(a: pseudo.BlockParameter, b: pseudo.BlockParameter) bool
 }
 
 fn validateInstruction(program: *const program_ir.Program, inst: instruction.Instruction) Error!void {
-    switch (inst.execution_size) {
-        .simd1, .simd8 => {},
-        else => return Error.UnsupportedExecutionSize,
-    }
-
     if (inst.predicate) |predicate|
         try validateFlag(program, predicate.flag);
 
     switch (inst.operation) {
-        .load_input => |op| {
+        .load_global_invocation_id => |op| {
-            if (program.properties.stage_io_lowered)
+            if (program.properties.system_values_lowered)
-                return Error.UnloweredStageIo;
+                return Error.UnloweredSystemValue;
             try validateDestination(program, op.destination);
-            try validateInterfaceSemantic(op.semantic, .input);
+            if (op.component >= 3 or op.destination.type != .u32)
+                return Error.InvalidGlobalInvocationId;
         },
-        .store_output => |op| {
+        .load_buffer => |op| {
-            if (program.properties.stage_io_lowered)
+            if (program.properties.resources_lowered)
-                return Error.UnloweredStageIo;
+                return Error.UnloweredResource;
+            if (!program.storage_buffers.isLive(op.buffer))
+                return Error.InvalidStorageBuffer;
+            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 validateBufferOffset(program, op.byte_offset);
             try validateSource(program, op.source);
-            try validateInterfaceSemantic(op.semantic, .output);
+            if (!op.source.type.isInitialTargetType())
+                return Error.InvalidBufferAccess;
         },
         .move => |op| {
             try validateDestination(program, op.destination);
@@ -155,19 +169,7 @@ fn validateInstruction(program: *const program_ir.Program, inst: instruction.Ins
             try validateSource(program, op.lhs);
             try validateSource(program, op.rhs);
         },
-        .send => |op| {
+
-            if (program.properties.messages_lowered)
-                return Error.UnloweredMessage;
-            try validateSpan(program, op.payload);
-            if (op.response) |response|
-                try validateSpan(program, response);
-            switch (op.message) {
-                .urb_write => |urb_write| {
-                    if (op.response != null or (!urb_write.channels.x and !urb_write.channels.y and !urb_write.channels.z and !urb_write.channels.w))
-                        return Error.InvalidMessage;
-                },
-            }
-        },
         .parallel_copy => |op| {
             if (program.properties.parallel_copies_lowered)
                 return Error.UnloweredParallelCopy;
@@ -178,27 +180,10 @@ fn validateInstruction(program: *const program_ir.Program, inst: instruction.Ins
     }
 }
 
-const InterfaceDirection = enum { input, output };
+fn validateBufferOffset(program: *const program_ir.Program, source: operand.Source) Error!void {
-
+    try validateSource(program, source);
-fn validateInterfaceSemantic(semantic: instruction.InterfaceSemantic, direction: InterfaceDirection) Error!void {
+    if (source.type != .u32)
-    switch (semantic) {
+        return Error.InvalidBufferAccess;
-        .location => |location| {
-            if (location.component > 3)
-                return Error.InvalidInterfaceSemantic;
-        },
-        .builtin => |builtin| switch (direction) {
-            .input => switch (builtin.builtin) {
-                .vertex_index, .instance_index => if (builtin.component != 0)
-                    return Error.InvalidInterfaceSemantic,
-                .position => return Error.InvalidInterfaceSemantic,
-            },
-            .output => switch (builtin.builtin) {
-                .position => if (builtin.component > 3)
-                    return Error.InvalidInterfaceSemantic,
-                .vertex_index, .instance_index => return Error.InvalidInterfaceSemantic,
-            },
-        },
-    }
 }
 
 fn validateParallelCopy(program: *const program_ir.Program, copy: pseudo.ParallelCopy) Error!void {
@@ -267,13 +252,7 @@ fn isBroadcast(region: operand.Region) bool {
     return region.vertical_stride == 0 and region.width == 1 and region.horizontal_stride == 0;
 }
 
-fn validateType(data_type: operand.DataType) Error!void {
-    if (!data_type.isInitialTargetType())
-        return Error.UnsupportedDataType;
-}
-
 fn validateSource(program: *const program_ir.Program, source: operand.Source) Error!void {
-    try validateType(source.type);
     if (source.region.width == 0)
         return Error.InvalidRegion;
     try validateRegisterRef(program, source.register);
@@ -290,7 +269,6 @@ fn validateSource(program: *const program_ir.Program, source: operand.Source) Er
 }
 
 fn validateDestination(program: *const program_ir.Program, destination: operand.Destination) Error!void {
-    try validateType(destination.type);
     if (destination.region.horizontal_stride == 0)
         return Error.InvalidRegion;
     switch (destination.register) {
@@ -316,28 +294,7 @@ fn validateFlag(program: *const program_ir.Program, flag: operand.FlagRef) Error
     switch (flag) {
         .virtual => |id| if (!program.virtual_flags.isLive(id))
             return Error.InvalidVirtualFlag,
-        .physical => |physical| if (physical.register != 0 or physical.subregister > 1)
+        .physical => {},
-            return Error.InvalidPhysicalFlag,
-    }
-}
-
-fn validateSpan(program: *const program_ir.Program, span: operand.RegisterSpan) Error!void {
-    if (span.register_count == 0)
-        return Error.InvalidRegisterSpan;
-    switch (span.base) {
-        .virtual => |register_id| {
-            try validateRegisterRef(program, span.base);
-            const register = program.virtual_registers.get(register_id) orelse return Error.InvalidVirtualRegister;
-            const required_size = @as(u32, span.register_count) * program.device_info.grf_size_bytes;
-            if (register.size_bytes < required_size)
-                return Error.InvalidRegisterSpan;
-        },
-        .physical_grf => |physical| {
-            try validateRegisterRef(program, span.base);
-            if (physical.byte_offset != 0 or @as(u32, physical.number) + span.register_count > program.device_info.grf_count)
-                return Error.InvalidRegisterSpan;
-        },
-        else => return Error.InvalidRegisterSpan,
     }
 }
 
@@ -417,3 +374,63 @@ fn validateBlockTarget(program: *const program_ir.Program, block_id: ids.BlockId
     if (!program.blocks.isLive(block_id))
         return Error.InvalidBlock;
 }
+
+test "[ir] validator checks compute system values and resources" {
+    const std = @import("std");
+    const Builder = @import("Builder.zig");
+    const device = @import("../device.zig");
+
+    const device_info: device.DeviceInfo = .{
+        .generation = .gen9,
+        .platform = .skylake,
+        .pci_device_id = 0x1912,
+        .grf_count = 128,
+    };
+    var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, device_info, .simd8);
+    defer program.deinit();
+    var builder = Builder.init(&program);
+
+    const register = try builder.addVirtualRegister(.{
+        .size_bytes = 32,
+        .alignment_bytes = 32,
+        .element_type = .u32,
+        .lane_count = 8,
+        .class = .temporary,
+    });
+    const buffer = try builder.addStorageBuffer(.{ .set = 0, .binding = 0 });
+    const entry = try builder.addBlock("entry");
+    const system_value_id = try builder.appendInstruction(entry, .simd8, null, .{
+        .load_global_invocation_id = .{
+            .destination = .{ .register = .{ .virtual = register }, .type = .u32 },
+            .component = 0,
+        },
+    });
+    const buffer_load_id = try builder.appendInstruction(entry, .simd8, null, .{
+        .load_buffer = .{
+            .destination = .{ .register = .{ .virtual = register }, .type = .u32 },
+            .buffer = buffer,
+            .byte_offset = .{
+                .register = .{ .immediate = .{ .u32 = 0 } },
+                .type = .u32,
+                .region = operand.Region.broadcast(),
+            },
+        },
+    });
+    try builder.setTerminator(entry, .end_thread);
+    try validate(&program);
+
+    program.instructions.getMut(system_value_id).?.operation.load_global_invocation_id.component = 3;
+    try std.testing.expectError(Error.InvalidGlobalInvocationId, validate(&program));
+    program.instructions.getMut(system_value_id).?.operation.load_global_invocation_id.component = 0;
+
+    program.properties.system_values_lowered = true;
+    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);
+    try std.testing.expectError(Error.InvalidStorageBuffer, validate(&program));
+    program.instructions.getMut(buffer_load_id).?.operation.load_buffer.buffer = buffer;
+
+    program.properties.resources_lowered = true;
+    try std.testing.expectError(Error.UnloweredResource, validate(&program));
+}

src/intel/compiler/lower/block_arguments.zig

@@ -150,7 +150,7 @@ test "[ir] block arguments: lower register and flag parameters" {
         .grf_count = 128,
     };
 
-    var program = program_ir.Program.init(std.testing.allocator, .vertex, device_info, .simd8);
+    var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, device_info, .simd8);
     defer program.deinit();
     var builder = Builder.init(&program);
 
@@ -224,7 +224,7 @@ test "[ir] block arguments: split same-target conditional edges" {
         .grf_count = 128,
     };
 
-    var program = program_ir.Program.init(std.testing.allocator, .vertex, device_info, .simd8);
+    var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, device_info, .simd8);
     defer program.deinit();
     var builder = Builder.init(&program);
 

src/intel/compiler/lower/vertex_abi.zig

@@ -1,491 +0,0 @@
-const std = @import("std");
-const Builder = @import("../ir/Builder.zig");
-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");
-const validator = @import("../ir/validator.zig");
-
-pub const InputComponent = struct {
-    location: u32,
-    component: u8,
-    payload_grf_offset: u16,
-};
-
-pub const Layout = struct {
-    input_components: []const InputComponent,
-    position_urb_offset: u16,
-};
-
-pub const Error = std.mem.Allocator.Error || error{
-    InvalidProgram,
-    UnsupportedTarget,
-    MissingVertexPayload,
-    InvalidLayout,
-    UnsupportedStageIo,
-    MissingPosition,
-    ExistingUrbWrite,
-};
-
-pub fn run(allocator: std.mem.Allocator, program: *program_ir.Program, layout: Layout) Error!void {
-    validator.validate(program) catch return Error.InvalidProgram;
-    if (program.properties.stage_io_lowered)
-        return;
-
-    if (!program.properties.common_ir_lowered or !program.properties.block_parameters_lowered or
-        program.properties.registers_allocated or program.properties.messages_lowered)
-        return Error.InvalidProgram;
-    if (program.device_info.generation != .gen9 or program.stage != .vertex or
-        program.dispatch_width != .simd8 or program.device_info.grf_size_bytes != 32)
-        return Error.UnsupportedTarget;
-
-    const vertex_payload = program.payload.vertex orelse return Error.MissingVertexPayload;
-    try validateLayout(program, vertex_payload, layout);
-
-    var position_components = instruction.ChannelMask{ .x = false, .y = false, .z = false, .w = false };
-    var end_thread_count: usize = 0;
-    try preflight(program, layout, &position_components, &end_thread_count);
-    if (!position_components.x or !position_components.y or !position_components.z or !position_components.w or end_thread_count == 0)
-        return Error.MissingPosition;
-
-    var builder = Builder.init(program);
-    const position_payload = builder.addVirtualRegister(.{
-        .size_bytes = 4 * program.device_info.grf_size_bytes,
-        .alignment_bytes = program.device_info.grf_size_bytes,
-        .element_type = .f32,
-        .lane_count = 4 * @intFromEnum(program.dispatch_width),
-        .class = .payload,
-        .spillable = false,
-        .name = "position_urb_payload",
-    }) catch |err| return mapBuilderError(err);
-
-    for (program.blocks.entries.items) |entry| {
-        const block = entry orelse continue;
-        for (block.instructions.items) |instruction_id| {
-            const inst = program.instructions.get(instruction_id) orelse return Error.InvalidProgram;
-            const replacement: ?instruction.Operation = switch (inst.operation) {
-                .load_input => |load| .{
-                    .move = .{
-                        .destination = load.destination,
-                        .source = .{
-                            .register = .{ .physical_grf = try inputPhysicalGrf(program, vertex_payload, layout, load.semantic) },
-                            .type = load.destination.type,
-                            .region = operand.Region.contiguous(.simd8),
-                        },
-                    },
-                },
-                .store_output => |store| blk: {
-                    const component = try positionComponent(store.semantic);
-                    break :blk .{
-                        .move = .{
-                            .destination = .{
-                                .register = .{ .virtual = position_payload },
-                                .type = .f32,
-                                .region = .{ .byte_offset = @as(u16, component) * program.device_info.grf_size_bytes },
-                            },
-                            .source = store.source,
-                        },
-                    };
-                },
-                else => null,
-            };
-            if (replacement) |operation|
-                builder.replaceOperation(instruction_id, operation) catch |err| return mapBuilderError(err);
-        }
-    }
-
-    for (program.blocks.entries.items, 0..) |entry, block_index| {
-        const block = entry orelse continue;
-        if (block.terminator.? != .end_thread)
-            continue;
-        _ = builder.appendInstruction(ids.BlockId.fromIndex(block_index), .simd8, null, .{
-            .send = .{
-                .message = .{
-                    .urb_write = .{
-                        .offset = layout.position_urb_offset,
-                        .channels = .{},
-                        .end_of_thread = true,
-                    },
-                },
-                .payload = .{
-                    .base = .{ .virtual = position_payload },
-                    .register_count = 4,
-                },
-            },
-        }) catch |err| return mapBuilderError(err);
-    }
-
-    program.properties.stage_io_lowered = true;
-    validator.validate(program) catch return Error.InvalidProgram;
-    _ = allocator;
-}
-
-fn validateLayout(program: *const program_ir.Program, vertex_payload: program_ir.VertexPayload, layout: Layout) Error!void {
-    if (vertex_payload.first_attribute_grf.byte_offset != 0 or vertex_payload.attribute_grf_count == 0)
-        return Error.InvalidLayout;
-    if (@as(u32, vertex_payload.first_attribute_grf.number) + vertex_payload.attribute_grf_count > program.device_info.grf_count)
-        return Error.InvalidLayout;
-
-    for (layout.input_components, 0..) |mapping, index| {
-        if (mapping.component > 3 or mapping.payload_grf_offset >= vertex_payload.attribute_grf_count)
-            return Error.InvalidLayout;
-        for (layout.input_components[0..index]) |previous| {
-            if (previous.location == mapping.location and previous.component == mapping.component)
-                return Error.InvalidLayout;
-        }
-    }
-}
-
-fn preflight(program: *const program_ir.Program, layout: Layout, position_components: *instruction.ChannelMask, end_thread_count: *usize) Error!void {
-    for (program.blocks.entries.items) |entry| {
-        const block = entry orelse continue;
-        for (block.instructions.items) |instruction_id| {
-            const inst = program.instructions.get(instruction_id) orelse return Error.InvalidProgram;
-            switch (inst.operation) {
-                .load_input => |load| {
-                    if (inst.execution_size != .simd8 or findInput(layout, load.semantic) == null)
-                        return Error.UnsupportedStageIo;
-                },
-                .store_output => |store| {
-                    if (inst.execution_size != .simd8 or store.source.type != .f32)
-                        return Error.UnsupportedStageIo;
-                    switch (try positionComponent(store.semantic)) {
-                        0 => position_components.x = true,
-                        1 => position_components.y = true,
-                        2 => position_components.z = true,
-                        3 => position_components.w = true,
-                        else => unreachable,
-                    }
-                },
-                .send => |send| switch (send.message) {
-                    .urb_write => return Error.ExistingUrbWrite,
-                },
-                else => {},
-            }
-        }
-        switch (block.terminator orelse return Error.InvalidProgram) {
-            .end_thread => end_thread_count.* += 1,
-            else => {},
-        }
-    }
-}
-
-fn findInput(layout: Layout, semantic: instruction.InterfaceSemantic) ?InputComponent {
-    const location = switch (semantic) {
-        .location => |location| location,
-        .builtin => return null,
-    };
-    for (layout.input_components) |mapping| {
-        if (mapping.location == location.location and mapping.component == location.component)
-            return mapping;
-    }
-    return null;
-}
-
-fn inputPhysicalGrf(program: *const program_ir.Program, vertex_payload: program_ir.VertexPayload, layout: Layout, semantic: instruction.InterfaceSemantic) Error!operand.PhysicalGrf {
-    const mapping = findInput(layout, semantic) orelse return Error.UnsupportedStageIo;
-    const number = @as(u32, vertex_payload.first_attribute_grf.number) + mapping.payload_grf_offset;
-    if (number >= program.device_info.grf_count)
-        return Error.InvalidLayout;
-    return .{ .number = @intCast(number) };
-}
-
-fn positionComponent(semantic: instruction.InterfaceSemantic) Error!u8 {
-    return switch (semantic) {
-        .builtin => |builtin| if (builtin.builtin == .position and builtin.component <= 3)
-            builtin.component
-        else
-            Error.UnsupportedStageIo,
-        .location => Error.UnsupportedStageIo,
-    };
-}
-
-fn mapBuilderError(err: anyerror) Error {
-    return switch (err) {
-        error.OutOfMemory => Error.OutOfMemory,
-        else => Error.InvalidProgram,
-    };
-}
-
-fn appendTestShaderBody(program: *program_ir.Program, position_component_count: u8) !ids.BlockId {
-    var builder = Builder.init(program);
-    program.properties.common_ir_lowered = true;
-    program.properties.block_parameters_lowered = true;
-    program.payload.vertex = .{
-        .first_attribute_grf = .{ .number = 4 },
-        .attribute_grf_count = 4,
-    };
-
-    const attribute = try builder.addVirtualRegister(.{
-        .size_bytes = 32,
-        .alignment_bytes = 32,
-        .element_type = .f32,
-        .lane_count = 8,
-        .class = .varying,
-        .name = "attribute",
-    });
-    const entry = try builder.addBlock("entry");
-    _ = try builder.appendInstruction(entry, .simd8, null, .{
-        .load_input = .{
-            .destination = .{ .register = .{ .virtual = attribute }, .type = .f32 },
-            .semantic = .{ .location = .{ .location = 2, .component = 1 } },
-        },
-    });
-
-    for (0..position_component_count) |component| {
-        const position = try builder.addVirtualRegister(.{
-            .size_bytes = 32,
-            .alignment_bytes = 32,
-            .element_type = .f32,
-            .lane_count = 8,
-            .class = .temporary,
-            .name = "position",
-        });
-        _ = try builder.appendInstruction(entry, .simd8, null, .{
-            .store_output = .{
-                .semantic = .{ .builtin = .{ .builtin = .position, .component = @intCast(component) } },
-                .source = .{
-                    .register = .{ .virtual = position },
-                    .type = .f32,
-                    .region = operand.Region.contiguous(.simd8),
-                },
-            },
-        });
-    }
-    return entry;
-}
-
-const test_input_layout = [_]InputComponent{.{
-    .location = 2,
-    .component = 1,
-    .payload_grf_offset = 3,
-}};
-
-const test_layout: Layout = .{
-    .input_components = &test_input_layout,
-    .position_urb_offset = 7,
-};
-
-test "vertex ABI: lower explicit input payload and position URB output" {
-    const device = @import("../device.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, .vertex, device_info, .simd8);
-    defer program.deinit();
-    var builder = Builder.init(&program);
-
-    program.properties.common_ir_lowered = true;
-    program.properties.block_parameters_lowered = true;
-    program.payload.vertex = .{
-        .first_attribute_grf = .{ .number = 4 },
-        .attribute_grf_count = 4,
-    };
-
-    const attribute = try builder.addVirtualRegister(.{
-        .size_bytes = 32,
-        .alignment_bytes = 32,
-        .element_type = .f32,
-        .lane_count = 8,
-        .class = .varying,
-        .name = "attribute",
-    });
-    const position_names = [_][]const u8{ "position_x", "position_y", "position_z", "position_w" };
-    var position: [4]ids.VirtualRegisterId = undefined;
-    for (&position, position_names) |*register_id, name| {
-        register_id.* = try builder.addVirtualRegister(.{
-            .size_bytes = 32,
-            .alignment_bytes = 32,
-            .element_type = .f32,
-            .lane_count = 8,
-            .class = .temporary,
-            .name = name,
-        });
-    }
-
-    const entry = try builder.addBlock("entry");
-    _ = try builder.appendInstruction(entry, .simd8, null, .{
-        .load_input = .{
-            .destination = .{ .register = .{ .virtual = attribute }, .type = .f32 },
-            .semantic = .{ .location = .{ .location = 2, .component = 1 } },
-        },
-    });
-    for (position, 0..) |register_id, component| {
-        _ = try builder.appendInstruction(entry, .simd8, null, .{
-            .store_output = .{
-                .semantic = .{ .builtin = .{ .builtin = .position, .component = @intCast(component) } },
-                .source = .{
-                    .register = .{ .virtual = register_id },
-                    .type = .f32,
-                    .region = operand.Region.contiguous(.simd8),
-                },
-            },
-        });
-    }
-    try builder.setTerminator(entry, .end_thread);
-    try validator.validate(&program);
-
-    const input_layout = [_]InputComponent{.{
-        .location = 2,
-        .component = 1,
-        .payload_grf_offset = 3,
-    }};
-    try run(std.testing.allocator, &program, .{
-        .input_components = &input_layout,
-        .position_urb_offset = 7,
-    });
-    try run(std.testing.allocator, &program, .{
-        .input_components = &input_layout,
-        .position_urb_offset = 7,
-    });
-
-    try std.testing.expect(program.properties.stage_io_lowered);
-    try std.testing.expect(!program.properties.messages_lowered);
-    try std.testing.expect(!program.properties.instructions_selected);
-
-    const text = try printer.allocPrint(std.testing.allocator, &program);
-    defer std.testing.allocator.free(text);
-    for ([_][]const u8{
-        "[simd8] mov %attribute:f32, r7:f32",
-        "[simd8] mov %position_urb_payload:f32, %position_x:f32",
-        "[simd8] mov %position_urb_payload:f32[byte=32], %position_y:f32",
-        "[simd8] mov %position_urb_payload:f32[byte=64], %position_z:f32",
-        "[simd8] mov %position_urb_payload:f32[byte=96], %position_w:f32",
-        "send urb_write[offset(7), channels(xyzw), end_of_thread], payload(%position_urb_payload[4])",
-    }) |fragment|
-        try std.testing.expect(std.mem.indexOf(u8, text, fragment) != null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "load_input") == null);
-    try std.testing.expect(std.mem.indexOf(u8, text, "store_output") == null);
-}
-
-test "vertex ABI: reject invalid layout and incomplete position" {
-    const device = @import("../device.zig");
-    const device_info: device.DeviceInfo = .{
-        .generation = .gen9,
-        .platform = .skylake,
-        .pci_device_id = 0x1912,
-        .grf_count = 128,
-    };
-
-    var invalid_layout_program = program_ir.Program.init(std.testing.allocator, .vertex, device_info, .simd8);
-    defer invalid_layout_program.deinit();
-    var invalid_layout_builder = Builder.init(&invalid_layout_program);
-    const invalid_layout_entry = try appendTestShaderBody(&invalid_layout_program, 4);
-    try invalid_layout_builder.setTerminator(invalid_layout_entry, .end_thread);
-
-    const out_of_range_input = [_]InputComponent{.{
-        .location = 2,
-        .component = 1,
-        .payload_grf_offset = 4,
-    }};
-    try std.testing.expectError(Error.InvalidLayout, run(std.testing.allocator, &invalid_layout_program, .{
-        .input_components = &out_of_range_input,
-        .position_urb_offset = 7,
-    }));
-
-    const duplicate_inputs = [_]InputComponent{
-        test_input_layout[0],
-        test_input_layout[0],
-    };
-    try std.testing.expectError(Error.InvalidLayout, run(std.testing.allocator, &invalid_layout_program, .{
-        .input_components = &duplicate_inputs,
-        .position_urb_offset = 7,
-    }));
-    try std.testing.expect(!invalid_layout_program.properties.stage_io_lowered);
-
-    var incomplete_program = program_ir.Program.init(std.testing.allocator, .vertex, device_info, .simd8);
-    defer incomplete_program.deinit();
-    var incomplete_builder = Builder.init(&incomplete_program);
-    const incomplete_entry = try appendTestShaderBody(&incomplete_program, 3);
-    try incomplete_builder.setTerminator(incomplete_entry, .end_thread);
-
-    try std.testing.expectError(Error.MissingPosition, run(std.testing.allocator, &incomplete_program, test_layout));
-    try std.testing.expect(!incomplete_program.properties.stage_io_lowered);
-}
-
-test "vertex ABI: reject an existing logical URB write" {
-    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, .vertex, device_info, .simd8);
-    defer program.deinit();
-    var builder = Builder.init(&program);
-    const entry = try appendTestShaderBody(&program, 4);
-    const payload = try builder.addVirtualRegister(.{
-        .size_bytes = 32,
-        .alignment_bytes = 32,
-        .element_type = .u32,
-        .lane_count = 8,
-        .class = .payload,
-        .spillable = false,
-        .name = "existing_payload",
-    });
-    _ = try builder.appendInstruction(entry, .simd8, null, .{
-        .send = .{
-            .message = .{ .urb_write = .{ .offset = 0 } },
-            .payload = .{
-                .base = .{ .virtual = payload },
-                .register_count = 1,
-            },
-        },
-    });
-    try builder.setTerminator(entry, .end_thread);
-
-    try std.testing.expectError(Error.ExistingUrbWrite, run(std.testing.allocator, &program, test_layout));
-    try std.testing.expect(!program.properties.stage_io_lowered);
-}
-
-test "vertex ABI: append an EOT URB write to every shader exit" {
-    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, .vertex, device_info, .simd8);
-    defer program.deinit();
-    var builder = Builder.init(&program);
-    const entry = try appendTestShaderBody(&program, 4);
-    const first_exit = try builder.addBlock("first_exit");
-    const second_exit = try builder.addBlock("second_exit");
-    const condition = try builder.addVirtualFlag(.{ .name = "condition" });
-    try builder.setTerminator(entry, .{ .conditional_branch = .{
-        .predicate = .{ .flag = .{ .virtual = condition } },
-        .true_edge = try builder.edge(first_exit, &.{}),
-        .false_edge = try builder.edge(second_exit, &.{}),
-    } });
-    try builder.setTerminator(first_exit, .end_thread);
-    try builder.setTerminator(second_exit, .end_thread);
-
-    try run(std.testing.allocator, &program, test_layout);
-
-    var urb_write_count: usize = 0;
-    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).?;
-            switch (inst.operation) {
-                .send => |send| switch (send.message) {
-                    .urb_write => |urb_write| {
-                        try std.testing.expect(urb_write.end_of_thread);
-                        try std.testing.expectEqual(@as(u16, 7), urb_write.offset);
-                        urb_write_count += 1;
-                    },
-                },
-                else => {},
-            }
-        }
-    }
-    try std.testing.expectEqual(@as(usize, 2), urb_write_count);
-    try validator.validate(&program);
-}

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

@@ -0,0 +1,20 @@
+const std = @import("std");
+
+pub const Error = error{UnsupportedWorkgroupSize};
+
+pub fn validateWorkgroupSize(size: [3]u32) Error!void {
+    if (size[0] == 0 or size[1] == 0 or size[2] == 0 or size[0] > 128 or size[1] > 128 or size[2] > 64)
+        return Error.UnsupportedWorkgroupSize;
+    const xy = std.math.mul(u32, size[0], size[1]) catch return Error.UnsupportedWorkgroupSize;
+    const invocations = std.math.mul(u32, xy, size[2]) catch return Error.UnsupportedWorkgroupSize;
+    if (invocations > 128)
+        return Error.UnsupportedWorkgroupSize;
+}
+
+test "[gen9] compute: validate workgroup limits" {
+    try validateWorkgroupSize(.{ 1, 1, 1 });
+    try validateWorkgroupSize(.{ 128, 1, 1 });
+    try std.testing.expectError(Error.UnsupportedWorkgroupSize, validateWorkgroupSize(.{ 0, 1, 1 }));
+    try std.testing.expectError(Error.UnsupportedWorkgroupSize, validateWorkgroupSize(.{ 129, 1, 1 }));
+    try std.testing.expectError(Error.UnsupportedWorkgroupSize, validateWorkgroupSize(.{ 64, 3, 1 }));
+}

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

@@ -0,0 +1,73 @@
+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");
+
+pub const compute = @import("compute/compute.zig");
+pub const validator = @import("validator.zig");
+
+pub const Options = common_ir.Options;
+pub const Error = common_ir.Error || compute.Error || error{
+    UnsupportedGeneration,
+    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();
+    validator.validate(&program) catch return Error.InvalidLoweredProgram;
+    return program;
+}
+
+test "[gen9] target: reject unsupported target configurations" {
+    var module = try shader_ir.parser.parseString(std.testing.allocator,
+        \\shader compute @main
+        \\{
+        \\    fn @main() -> void
+        \\    {
+        \\        .entry():
+        \\            return
+        \\    }
+        \\}
+    );
+    defer module.deinit();
+
+    const gen9_device: device.DeviceInfo = .{
+        .generation = .gen9,
+        .platform = .skylake,
+        .pci_device_id = 0x1912,
+        .grf_count = 128,
+    };
+    var other_generation = gen9_device;
+    other_generation.generation = .gen11;
+    try std.testing.expectError(Error.UnsupportedGeneration, lower(std.testing.allocator, &module, other_generation, .{}));
+
+    module.stage = .fragment;
+    try std.testing.expectError(Error.UnsupportedStage, lower(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 }));
+
+    var wide_grf = gen9_device;
+    wide_grf.grf_size_bytes = 64;
+    try std.testing.expectError(Error.UnsupportedGrfSize, lower(std.testing.allocator, &module, wide_grf, .{}));
+}

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

@@ -0,0 +1,173 @@
+const std = @import("std");
+const compute = @import("compute/compute.zig");
+const shared = @import("../../ir/validator.zig");
+const instruction = @import("../../ir/instruction.zig");
+const operand = @import("../../ir/operand.zig");
+const program_ir = @import("../../ir/program.zig");
+
+pub const Error = shared.Error || compute.Error || error{
+    UnsupportedGeneration,
+    UnsupportedDispatchWidth,
+    UnsupportedGrfSize,
+    UnsupportedExecutionSize,
+    UnsupportedDataType,
+    InvalidPhysicalFlag,
+    InvalidPayloadLayout,
+};
+
+pub fn validate(program: *const program_ir.Program) Error!void {
+    try shared.validate(program);
+
+    if (program.device_info.generation != .gen9)
+        return Error.UnsupportedGeneration;
+    try compute.validateWorkgroupSize(program.workgroup_size);
+    if (program.dispatch_width != .simd8 or !program.device_info.supportsDispatch(.simd8))
+        return Error.UnsupportedDispatchWidth;
+    if (program.device_info.grf_size_bytes != 32)
+        return Error.UnsupportedGrfSize;
+
+    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.InvalidInstruction;
+            switch (inst.execution_size) {
+                .simd1, .simd8 => {},
+                else => return Error.UnsupportedExecutionSize,
+            }
+            try validateInstruction(inst.*);
+        }
+        try validateTerminator(block.terminator.?);
+    }
+
+    for (program.virtual_registers.entries.items) |entry| {
+        const register = entry orelse continue;
+        if (!register.element_type.isInitialTargetType())
+            return Error.UnsupportedDataType;
+    }
+
+    try validatePayload(program);
+}
+
+fn validateInstruction(inst: instruction.Instruction) Error!void {
+    if (inst.predicate) |predicate|
+        try validateFlag(predicate.flag);
+    switch (inst.operation) {
+        .load_global_invocation_id => |op| try validateDestination(op.destination),
+        .load_buffer => |op| {
+            try validateDestination(op.destination);
+            try validateSource(op.byte_offset);
+        },
+        .store_buffer => |op| {
+            try validateSource(op.byte_offset);
+            try validateSource(op.source);
+        },
+        .move => |op| {
+            try validateDestination(op.destination);
+            try validateSource(op.source);
+        },
+        .binary => |op| {
+            try validateDestination(op.destination);
+            try validateSource(op.lhs);
+            try validateSource(op.rhs);
+        },
+        .compare => |op| {
+            try validateFlag(op.destination);
+            try validateSource(op.lhs);
+            try validateSource(op.rhs);
+        },
+        .parallel_copy => |copy| {
+            for (copy.register_copies) |item| {
+                try validateDestination(item.destination);
+                try validateSource(item.source);
+            }
+            for (copy.flag_copies) |item| switch (item.source) {
+                .constant => {},
+                .dynamic => |predicate| try validateFlag(predicate.flag),
+            };
+        },
+    }
+}
+
+fn validateSource(source: operand.Source) Error!void {
+    try validateType(source.type);
+    switch (source.register) {
+        .immediate => |immediate| try validateImmediate(immediate),
+        else => {},
+    }
+}
+
+fn validateDestination(destination: operand.Destination) Error!void {
+    try validateType(destination.type);
+}
+
+fn validateType(data_type: operand.DataType) Error!void {
+    if (!data_type.isInitialTargetType())
+        return Error.UnsupportedDataType;
+}
+
+fn validateImmediate(immediate: operand.Immediate) Error!void {
+    switch (immediate) {
+        .u32, .i32, .f32 => {},
+    }
+}
+
+fn validateTerminator(terminator: instruction.Terminator) Error!void {
+    switch (terminator) {
+        .conditional_branch => |branch| {
+            try validateFlag(branch.predicate.flag);
+            try validateEdge(branch.true_edge);
+            try validateEdge(branch.false_edge);
+        },
+        .jump => |edge| try validateEdge(edge),
+        else => {},
+    }
+}
+
+fn validateEdge(edge: instruction.Edge) Error!void {
+    for (edge.arguments) |argument| switch (argument) {
+        .source => {},
+        .predicate => |predicate_value| switch (predicate_value) {
+            .constant => {},
+            .dynamic => |predicate| try validateFlag(predicate.flag),
+        },
+    };
+}
+
+fn validateFlag(flag: operand.FlagRef) Error!void {
+    switch (flag) {
+        .virtual => {},
+        .physical => |physical| if (physical.register != 0 or physical.subregister > 1)
+            return Error.InvalidPhysicalFlag,
+    }
+}
+
+fn validatePayload(program: *const program_ir.Program) Error!void {
+    if (program.payload.header_grf) |header| {
+        if (header.number != 0 or header.byte_offset != 0)
+            return Error.InvalidPayloadLayout;
+    }
+}
+
+test "[gen9] validator: layer target legality over shared structural validation" {
+    const Builder = @import("../../ir/Builder.zig");
+    const device = @import("../../device.zig");
+
+    const gen11_device: device.DeviceInfo = .{
+        .generation = .gen11,
+        .platform = .ice_lake,
+        .pci_device_id = 0x8a52,
+        .grf_count = 128,
+        .supports_simd16 = true,
+    };
+    var program = program_ir.Program.init(std.testing.allocator, .{ 1, 1, 1 }, gen11_device, .simd16);
+    defer program.deinit();
+    var builder = Builder.init(&program);
+    const entry = try builder.addBlock("entry");
+    try builder.setTerminator(entry, .end_thread);
+
+    try shared.validate(&program);
+    try std.testing.expectError(Error.UnsupportedGeneration, validate(&program));
+
+    program.device_info.generation = .gen9;
+    try std.testing.expectError(Error.UnsupportedDispatchWidth, validate(&program));
+}

src/intel/compiler/targets/targets.zig

@@ -0,0 +1,29 @@
+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");
+
+pub const gen9 = @import("gen9/gen9.zig");
+
+pub const Error = gen9.Error || error{UnsupportedGeneration};
+pub const ValidationError = gen9.validator.Error || error{UnsupportedGeneration};
+
+pub fn lower(
+    allocator: std.mem.Allocator,
+    module: *shader_ir.module.Module,
+    device_info: device.DeviceInfo,
+    options: common_ir.Options,
+) Error!program_ir.Program {
+    return switch (device_info.generation) {
+        .gen9 => gen9.lower(allocator, module, device_info, options),
+        .gen10, .gen11 => Error.UnsupportedGeneration,
+    };
+}
+
+pub fn validate(program: *const program_ir.Program) ValidationError!void {
+    return switch (program.device_info.generation) {
+        .gen9 => gen9.validator.validate(program),
+        .gen10, .gen11 => ValidationError.UnsupportedGeneration,
+    };
+}

src/software/interpreter/Runtime.zig

@@ -39,12 +39,16 @@ scratch: []u32,
 pub fn init(allocator: std.mem.Allocator, program: *const Program) !Self {
     const registers = try allocator.alloc(u32, program.register_count);
     errdefer allocator.free(registers);
+
     const scratch = try allocator.alloc(u32, program.scratch_count);
     errdefer allocator.free(scratch);
+
     @memset(registers, 0);
     @memset(scratch, 0);
+
     for (program.initializers) |initializer|
         registers[initializer.register] = initializer.value;
+
     return .{ .allocator = allocator, .registers = registers, .scratch = scratch };
 }
 

src/software/interpreter/root.zig

@@ -1,7 +1,4 @@
 //! Software bytecode interpreter for the backend-agnostic shader IR.
-//!
-//! This first slice supports allocation-free scalar execution of 32-bit scalar
-//! and vector arithmetic, interface I/O, control flow, and block parameters.
 
 pub const bytecode = @import("bytecode.zig");
 pub const Program = @import("Program.zig");