diff --git a/.gitea/workflows/Build.yml b/.gitea/workflows/Build.yml index 3352c15..bef7ba0 100644 --- a/.gitea/workflows/Build.yml +++ b/.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 diff --git a/src/compiler/ir/ir.zig b/src/compiler/ir/ir.zig index aede46f..29af086 100644 --- a/src/compiler/ir/ir.zig +++ b/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; -} diff --git a/src/intel/FlintPipeline.zig b/src/intel/FlintPipeline.zig index 47cb99f..bc51b45 100644 --- a/src/intel/FlintPipeline.zig +++ b/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) { - error.OutOfMemory => VkError.OutOfHostMemory, + const program = compiler.targets.lower(allocator, module, target, .{}) catch |err| switch (err) { + 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); +} diff --git a/src/intel/compiler/compiler.zig b/src/intel/compiler/compiler.zig index 0c34a25..7fbf221 100644 --- a/src/intel/compiler/compiler.zig +++ b/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" { - // ; 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 - +test "[ir] basic compute shader" { 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, - .class = .payload, - .spillable = false, - .name = "urb_payload", + .lane_count = 8, + .class = .temporary, + .name = "value", }); + 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 = .{ - .destination = .{ - .register = .{ .virtual = position }, - .type = .f32, - }, - .semantic = .{ - .location = .{ - .location = 0, - }, - }, + .load_global_invocation_id = .{ + .destination = .{ .register = .{ .virtual = value }, .type = .u32 }, + .component = 0, }, }); _ = try builder.appendInstruction(entry, .simd8, null, .{ - .binary = .{ - .opcode = .multiply, - .destination = .{ - .register = .{ .virtual = position }, - .type = .f32, - }, - .lhs = .{ - .register = .{ .virtual = position }, - .type = .f32, - .region = operand.Region.contiguous(.simd8), - }, - .rhs = .{ - .register = .{ - .immediate = .{ .f32 = 1.0 }, - }, - .type = .f32, + .store_buffer = .{ + .buffer = storage, + .byte_offset = .{ + .register = .{ .immediate = .{ .u32 = 0 } }, + .type = .u32, .region = operand.Region.broadcast(), }, - }, - }); - _ = try builder.appendInstruction(entry, .simd8, null, .{ - .move = .{ - .destination = .{ - .register = .{ .virtual = position }, - .type = .f32, - }, .source = .{ - .register = .{ .virtual = position }, - .type = .f32, - .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, + .register = .{ .virtual = value }, + .type = .u32, .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 program") != null); - try std.testing.expect(std.mem.indexOf(u8, text, "vertex") != null); - try std.testing.expect(std.mem.indexOf(u8, text, "gen9") != 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); + try std.testing.expect(std.mem.indexOf(u8, text, "Flint compute 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, "load_global_invocation_id %value:u32, component(0)") != null); + try std.testing.expect(std.mem.indexOf(u8, text, "store_buffer @storage, 0:u32, %value:u32") != 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; -} diff --git a/src/intel/compiler/ir/Builder.zig b/src/intel/compiler/ir/Builder.zig index 87b394f..7ee8d4a 100644 --- a/src/intel/compiler/ir/Builder.zig +++ b/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); diff --git a/src/intel/compiler/ir/id.zig b/src/intel/compiler/ir/id.zig index c1c1061..1aae900 100644 --- a/src/intel/compiler/ir/id.zig +++ b/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; diff --git a/src/intel/compiler/ir/instruction.zig b/src/intel/compiler/ir/instruction.zig index c4d4afa..2c2b6e7 100644 --- a/src/intel/compiler/ir/instruction.zig +++ b/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 { - 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 { +pub const LoadGlobalInvocationId = struct { destination: operand.Destination, - semantic: InterfaceSemantic, + component: u8, }; -pub const StoreOutput = struct { - semantic: InterfaceSemantic, +pub const LoadBuffer = struct { + 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, - store_output: StoreOutput, + load_global_invocation_id: LoadGlobalInvocationId, + load_buffer: LoadBuffer, + store_buffer: StoreBuffer, move: Move, binary: Binary, compare: Compare, - send: Send, parallel_copy: pseudo.ParallelCopy, }; diff --git a/src/intel/compiler/ir/ir.zig b/src/intel/compiler/ir/ir.zig index 7770dee..2497dd7 100644 --- a/src/intel/compiler/ir/ir.zig +++ b/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; diff --git a/src/intel/compiler/ir/operand.zig b/src/intel/compiler/ir/operand.zig index 8219115..12e84bd 100644 --- a/src/intel/compiler/ir/operand.zig +++ b/src/intel/compiler/ir/operand.zig @@ -33,7 +33,6 @@ pub const DataType = enum { pub const RegisterClass = enum { uniform, - varying, payload, response, temporary, diff --git a/src/intel/compiler/ir/printer.zig b/src/intel/compiler/ir/printer.zig index e673abf..a3e9a97 100644 --- a/src/intel/compiler/ir/printer.zig +++ b/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.print("; .stage: {t}\n", .{program.stage}); + try writer.writeAll("; Flint compute program:\n"); + 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| { - try writer.writeAll("load_input "); + .load_global_invocation_id => |op| { + 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| { - try writer.writeAll("store_output "); - try writeInterfaceSemantic(writer, op.semantic); + .store_buffer => |op| { + try writer.writeAll("store_buffer "); + 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 { - try writeRegister(program, writer, span.base); - const byte_offset = registerByteOffset(span.base); - 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 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(), '@'); } fn writeVirtualRegisterRef(program: *const program_ir.Program, writer: *std.Io.Writer, register_id: ids.VirtualRegisterId) !void { diff --git a/src/intel/compiler/ir/program.zig b/src/intel/compiler/ir/program.zig index f5f69c7..834ed75 100644 --- a/src/intel/compiler/ir/program.zig +++ b/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 { - first_attribute_grf: operand.PhysicalGrf, - attribute_grf_count: u16, +pub const StorageBuffer = struct { + set: u32, + 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(), .{ diff --git a/src/intel/compiler/ir/pseudo.zig b/src/intel/compiler/ir/pseudo.zig index 54a5b80..0bc06cb 100644 --- a/src/intel/compiler/ir/pseudo.zig +++ b/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); diff --git a/src/intel/compiler/ir/validator.zig b/src/intel/compiler/ir/validator.zig index 8cd3e19..eaf97a2 100644 --- a/src/intel/compiler/ir/validator.zig +++ b/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, - UnloweredMessage, - InvalidInterfaceSemantic, - InvalidMessage, + UnloweredSystemValue, + UnloweredResource, + InvalidPayloadLayout, 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) - return Error.UnsupportedGeneration; - if (program.stage != .vertex) - return Error.UnsupportedStage; - if (program.dispatch_width != .simd8 or !program.device_info.supportsDispatch(.simd8)) - return Error.UnsupportedDispatchWidth; + if (program.workgroup_size[0] == 0 or program.workgroup_size[1] == 0 or program.workgroup_size[2] == 0) + return Error.InvalidWorkgroupSize; 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| { - if (program.properties.stage_io_lowered) - return Error.UnloweredStageIo; + .load_global_invocation_id => |op| { + if (program.properties.system_values_lowered) + 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| { - if (program.properties.stage_io_lowered) - return Error.UnloweredStageIo; + .load_buffer => |op| { + if (program.properties.resources_lowered) + 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 validateInterfaceSemantic(semantic: instruction.InterfaceSemantic, direction: InterfaceDirection) Error!void { - switch (semantic) { - .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 validateBufferOffset(program: *const program_ir.Program, source: operand.Source) Error!void { + try validateSource(program, source); + if (source.type != .u32) + return Error.InvalidBufferAccess; } 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) - 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, + .physical => {}, } } @@ -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)); +} diff --git a/src/intel/compiler/lower/block_arguments.zig b/src/intel/compiler/lower/block_arguments.zig index 39592b1..935a7ab 100644 --- a/src/intel/compiler/lower/block_arguments.zig +++ b/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); diff --git a/src/intel/compiler/lower/common_ir.zig b/src/intel/compiler/lower/common_ir.zig new file mode 100644 index 0000000..5b2c6a1 --- /dev/null +++ b/src/intel/compiler/lower/common_ir.zig @@ -0,0 +1,1491 @@ +const std = @import("std"); +const shader_compiler = @import("shader_ir"); +const shader_ir = shader_compiler.ir; +const device = @import("../device.zig"); +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 printer = @import("../ir/printer.zig"); +const pseudo = @import("../ir/pseudo.zig"); +const program_ir = @import("../ir/program.zig"); +const validator = @import("../ir/validator.zig"); + +pub const block_arguments = @import("block_arguments.zig"); + +pub const Options = struct { + dispatch_width: device.DispatchWidth = .simd8, +}; + +pub const Error = std.mem.Allocator.Error || error{ + MissingEntryPoint, + InvalidEntryPoint, + InvalidModule, + InvalidLoweredProgram, + SanitizationFailed, + UnsanitizedModule, + UnsupportedStage, + MissingWorkgroupSize, + UnsupportedType, + UnsupportedOperation, + UnsupportedTerminator, +}; + +const PredicateValue = pseudo.PredicateValue; + +const LoweredType = struct { + element_type: operand.DataType, + component_count: usize, +}; + +const ValueLocation = union(enum) { + components: []const operand.Source, + predicate: PredicateValue, +}; + +const LoweringState = struct { + lowerer: *Lowerer, + builder: Builder, + storage: std.mem.Allocator, + block_map: []?ids.BlockId, + value_locations: []?ValueLocation, + storage_buffer_map: []?ids.StorageBufferId, + + fn lowerScalarType(self: *const LoweringState, type_id: shader_ir.id.TypeId) Error!operand.DataType { + const ty = self.lowerer.module.types.get(type_id) orelse return Error.InvalidModule; + return switch (ty.*) { + .integer => |integer| if (integer.bits == 32) + switch (integer.signedness) { + .unsigned => .u32, + .signed => .i32, + } + else + Error.UnsupportedType, + .floating => |floating| if (floating.bits == 32) .f32 else Error.UnsupportedType, + else => Error.UnsupportedType, + }; + } + + fn lowerType(self: *const LoweringState, type_id: shader_ir.id.TypeId) Error!LoweredType { + const ty = self.lowerer.module.types.get(type_id) orelse return Error.InvalidModule; + return switch (ty.*) { + .integer, .floating => .{ + .element_type = try self.lowerScalarType(type_id), + .component_count = 1, + }, + .vector => |vector| if (vector.length >= 2 and vector.length <= 4) + .{ + .element_type = try self.lowerScalarType(vector.element_type), + .component_count = vector.length, + } + else + Error.UnsupportedType, + else => Error.UnsupportedType, + }; + } + + fn isBoolean(self: *const LoweringState, type_id: shader_ir.id.TypeId) Error!bool { + const ty = self.lowerer.module.types.get(type_id) orelse return Error.InvalidModule; + return ty.* == .boolean; + } + + fn mappedBlock(self: *const LoweringState, source_id: shader_ir.id.BlockId) Error!ids.BlockId { + if (source_id.index() >= self.block_map.len) + return Error.InvalidModule; + return self.block_map[source_id.index()] orelse Error.InvalidModule; + } + + fn storageBuffer(self: *LoweringState, source_id: shader_ir.id.ResourceId) Error!ids.StorageBufferId { + if (source_id.index() >= self.storage_buffer_map.len) + return Error.InvalidModule; + if (self.storage_buffer_map[source_id.index()]) |existing| + return existing; + + const resource = self.lowerer.module.resources.get(source_id) orelse return Error.InvalidModule; + if (resource.kind != .storage_buffer) + return Error.UnsupportedOperation; + const buffer_id = self.builder.addStorageBuffer(.{ + .set = resource.set, + .binding = resource.binding, + .name = resource.name, + }) catch |err| return mapProgramError(err); + self.storage_buffer_map[source_id.index()] = buffer_id; + return buffer_id; + } + + fn putLocation(self: *LoweringState, value_id: shader_ir.id.ValueId, new_location: ValueLocation) Error!void { + if (value_id.index() >= self.value_locations.len or self.value_locations[value_id.index()] != null) + return Error.InvalidModule; + self.value_locations[value_id.index()] = new_location; + } + + fn addRegister(self: *LoweringState, data_type: operand.DataType, class: operand.RegisterClass, name: ?[]const u8) Error!ids.VirtualRegisterId { + return self.builder.addVirtualRegister(.{ + .size_bytes = @as(u32, data_type.sizeBytes()) * @intFromEnum(self.lowerer.options.dispatch_width), + .alignment_bytes = self.lowerer.device_info.grf_size_bytes, + .element_type = data_type, + .lane_count = @intFromEnum(self.lowerer.options.dispatch_width), + .class = class, + .name = name, + }) catch |err| return mapProgramError(err); + } + + fn executionSize(self: *const LoweringState) device.ExecutionSize { + return @enumFromInt(@intFromEnum(self.lowerer.options.dispatch_width)); + } + + fn registerSource(self: *const LoweringState, register_id: ids.VirtualRegisterId, data_type: operand.DataType) operand.Source { + return .{ + .register = .{ .virtual = register_id }, + .type = data_type, + .region = operand.Region.contiguous(self.executionSize()), + }; + } + + fn componentName(self: *LoweringState, name: ?[]const u8, component_index: usize, component_count: usize) Error!?[]const u8 { + if (name == null or component_count == 1) + return name; + const suffixes = "xyzw"; + const formatted = try std.fmt.allocPrint(self.storage, "{s}_{c}", .{ name.?, suffixes[component_index] }); + return @as([]const u8, formatted); + } + + fn addRegisterLocation(self: *LoweringState, value_id: shader_ir.id.ValueId, class: operand.RegisterClass) Error![]const operand.Source { + const value = self.lowerer.module.values.get(value_id) orelse return Error.InvalidModule; + const lowered_type = try self.lowerType(value.type); + const result = try self.storage.alloc(operand.Source, lowered_type.component_count); + for (result, 0..) |*component, component_index| { + const register_id = try self.addRegister( + lowered_type.element_type, + class, + try self.componentName(value.name, component_index, lowered_type.component_count), + ); + component.* = self.registerSource(register_id, lowered_type.element_type); + } + try self.putLocation(value_id, .{ .components = result }); + return result; + } + + fn location(self: *LoweringState, value_id: shader_ir.id.ValueId) Error!ValueLocation { + if (value_id.index() >= self.value_locations.len) + return Error.InvalidModule; + + if (self.value_locations[value_id.index()]) |existing| + return existing; + + const value = self.lowerer.module.values.get(value_id) orelse return Error.InvalidModule; + switch (value.definition) { + .constant => |constant_id| { + const constant = self.lowerer.module.constants.get(constant_id) orelse return Error.InvalidModule; + if (constant.type != value.type) + return Error.InvalidModule; + + const result: ValueLocation = if (try self.isBoolean(value.type)) switch (constant.value) { + .boolean => |boolean| .{ .predicate = .{ .constant = boolean } }, + else => return Error.UnsupportedType, + } else .{ + .components = try self.constantComponents(value.type, constant.value), + }; + self.value_locations[value_id.index()] = result; + return result; + }, + .undef => { + if (try self.isBoolean(value.type)) + return Error.UnsupportedType; + _ = try self.addRegisterLocation(value_id, .temporary); + return self.value_locations[value_id.index()].?; + }, + else => return Error.InvalidModule, + } + } + + fn components(self: *LoweringState, value_id: shader_ir.id.ValueId) Error![]const operand.Source { + return switch (try self.location(value_id)) { + .components => |values| values, + .predicate => Error.UnsupportedType, + }; + } + + fn source(self: *LoweringState, value_id: shader_ir.id.ValueId) Error!operand.Source { + const values = try self.components(value_id); + if (values.len != 1) + return Error.UnsupportedType; + return values[0]; + } + + fn predicate(self: *LoweringState, value_id: shader_ir.id.ValueId) Error!PredicateValue { + return switch (try self.location(value_id)) { + .components => Error.UnsupportedType, + .predicate => |value| value, + }; + } + + fn destinationFromSource(source_value: operand.Source) Error!operand.Destination { + if (source_value.negate or source_value.absolute) + return Error.InvalidLoweredProgram; + + return switch (source_value.register) { + .virtual => .{ + .register = source_value.register, + .type = source_value.type, + .region = .{ .byte_offset = source_value.region.byte_offset }, + }, + else => Error.InvalidLoweredProgram, + }; + } + + fn destination(self: *LoweringState, value_id: shader_ir.id.ValueId) Error!operand.Destination { + return destinationFromSource(try self.source(value_id)); + } + + fn constantComponents(self: *LoweringState, type_id: shader_ir.id.TypeId, value: shader_ir.constant.ConstantValue) Error![]const operand.Source { + const lowered_type = try self.lowerType(type_id); + const result = try self.storage.alloc(operand.Source, lowered_type.component_count); + if (lowered_type.component_count == 1) { + result[0] = try self.constantScalarSource(type_id, value); + return result; + } + + const ty = self.lowerer.module.types.get(type_id) orelse return Error.InvalidModule; + const vector = switch (ty.*) { + .vector => |vector| vector, + else => return Error.InvalidModule, + }; + switch (value) { + .composite => |elements| { + if (elements.len != lowered_type.component_count) + return Error.InvalidModule; + for (elements, result) |constant_id, *component| { + const element = self.lowerer.module.constants.get(constant_id) orelse return Error.InvalidModule; + if (element.type != vector.element_type) + return Error.InvalidModule; + component.* = try self.constantScalarSource(element.type, element.value); + } + }, + .null => { + const zero: shader_ir.constant.ConstantValue = switch (lowered_type.element_type) { + .u32, .i32 => .{ .integer_bits = 0 }, + .f32 => .{ .float_bits = 0 }, + else => unreachable, + }; + for (result) |*component| + component.* = try self.constantScalarSource(vector.element_type, zero); + }, + else => return Error.UnsupportedType, + } + return result; + } + + fn constantScalarSource(self: *const LoweringState, type_id: shader_ir.id.TypeId, value: shader_ir.constant.ConstantValue) Error!operand.Source { + const data_type = try self.lowerScalarType(type_id); + const immediate: operand.Immediate = switch (data_type) { + .u32 => switch (value) { + .integer_bits => |bits| .{ .u32 = @truncate(bits) }, + else => return Error.UnsupportedType, + }, + .i32 => switch (value) { + .integer_bits => |bits| .{ .i32 = @bitCast(@as(u32, @truncate(bits))) }, + else => return Error.UnsupportedType, + }, + .f32 => switch (value) { + .float_bits => |bits| .{ .f32 = @bitCast(@as(u32, @truncate(bits))) }, + else => return Error.UnsupportedType, + }, + else => unreachable, + }; + + return .{ + .register = .{ .immediate = immediate }, + .type = data_type, + .region = operand.Region.broadcast(), + }; + } + + fn appendInstruction(self: *LoweringState, block_id: ids.BlockId, predicate_value: ?operand.Predicate, operation: instruction.Operation) Error!void { + _ = self.builder.appendInstruction(block_id, self.executionSize(), predicate_value, operation) catch |err| + return mapProgramError(err); + } + + fn appendMove(self: *LoweringState, block_id: ids.BlockId, predicate_value: ?operand.Predicate, destination_value: operand.Destination, source_value: operand.Source) Error!void { + try self.appendInstruction(block_id, predicate_value, .{ + .move = .{ + .destination = destination_value, + .source = source_value, + }, + }); + } + + fn sourceEntryFunction(self: *const LoweringState) Error!struct { shader_ir.id.FunctionId, *const shader_ir.module.Function } { + const source_entry = self.lowerer.module.entry_point orelse return Error.MissingEntryPoint; + const function = self.lowerer.module.functions.get(source_entry) orelse return Error.InvalidEntryPoint; + const return_type = self.lowerer.module.types.get(function.return_type) orelse return Error.InvalidModule; + if (return_type.* != .void or function.parameters.items.len != 0) + return Error.InvalidEntryPoint; + return .{ source_entry, function }; + } + + fn lowerBlocks(self: *LoweringState) Error!void { + const source_function_id, const function = try self.sourceEntryFunction(); + + for (function.blocks.items) |source_block_id| { + const source_block = self.lowerer.module.blocks.get(source_block_id) orelse return Error.InvalidModule; + if (source_block.parent_function != source_function_id) + return Error.InvalidModule; + const target_block_id = self.builder.addBlock(source_block.name) catch |err| + return mapProgramError(err); + if (source_block_id.index() >= self.block_map.len or self.block_map[source_block_id.index()] != null) + return Error.InvalidModule; + self.block_map[source_block_id.index()] = target_block_id; + } + + const source_entry = function.entry_block orelse return Error.InvalidModule; + self.builder.setEntryBlock(try self.mappedBlock(source_entry)) catch |err| return mapProgramError(err); + } + + fn lowerParameters(self: *LoweringState) Error!void { + const source_entry = try self.sourceEntryFunction(); + const function = source_entry[1]; + + for (function.blocks.items) |source_block_id| { + const source_block = self.lowerer.module.blocks.get(source_block_id) orelse return Error.InvalidModule; + const target_block_id = try self.mappedBlock(source_block_id); + + for (source_block.parameters.items) |parameter_id| { + const value = self.lowerer.module.values.get(parameter_id) orelse return Error.InvalidModule; + if (try self.isBoolean(value.type)) { + const flag_id = self.builder.addVirtualFlag(.{ .name = value.name }) catch |err| + return mapProgramError(err); + const predicate_value: operand.Predicate = .{ .flag = .{ .virtual = flag_id } }; + try self.putLocation(parameter_id, .{ .predicate = .{ .dynamic = predicate_value } }); + self.builder.addBlockParameter(target_block_id, .{ .flag = flag_id }) catch |err| + return mapProgramError(err); + } else { + const parameter_components = try self.addRegisterLocation(parameter_id, .temporary); + for (parameter_components) |parameter_source| { + const register_id = switch (parameter_source.register) { + .virtual => |id| id, + else => return Error.InvalidLoweredProgram, + }; + self.builder.addBlockParameter(target_block_id, .{ .register = register_id }) catch |err| + return mapProgramError(err); + } + } + } + } + } + + fn lowerInstructions(self: *LoweringState, allocator: std.mem.Allocator) Error!void { + const visited = try allocator.alloc(bool, self.lowerer.module.blocks.entries.items.len); + defer allocator.free(visited); + @memset(visited, false); + + const source_entry = try self.sourceEntryFunction(); + const function = source_entry[1]; + try self.lowerBlockInstructions(function.entry_block orelse return Error.InvalidModule, visited); + + for (function.blocks.items) |source_block_id| { + if (!visited[source_block_id.index()]) + try self.lowerBlockInstructions(source_block_id, visited); + } + } + + fn lowerBlockInstructions(self: *LoweringState, source_block_id: shader_ir.id.BlockId, visited: []bool) Error!void { + if (source_block_id.index() >= visited.len) + return Error.InvalidModule; + + if (visited[source_block_id.index()]) + return; + + visited[source_block_id.index()] = true; + + const block = self.lowerer.module.blocks.get(source_block_id) orelse return Error.InvalidModule; + const target_block_id = try self.mappedBlock(source_block_id); + for (block.instructions.items) |instruction_id| { + const source_instruction = self.lowerer.module.instructions.get(instruction_id) orelse return Error.InvalidModule; + if (source_instruction.parent_block != source_block_id) + return Error.InvalidModule; + try self.lowerInstruction(target_block_id, source_instruction.*); + } + + switch (block.terminator orelse return Error.InvalidModule) { + .branch => |edge| try self.lowerBlockInstructions(edge.target, visited), + .conditional_branch => |branch| { + try self.lowerBlockInstructions(branch.true_edge.target, visited); + try self.lowerBlockInstructions(branch.false_edge.target, visited); + }, + else => {}, + } + } + + fn lowerInstruction(self: *LoweringState, block_id: ids.BlockId, source_instruction: shader_ir.instruction.Instruction) Error!void { + switch (source_instruction.operation) { + .unary => |operation| try self.lowerUnary(block_id, source_instruction.result, operation), + .binary => |operation| try self.lowerBinary(block_id, source_instruction.result, operation), + .compare => |operation| try self.lowerCompare(block_id, source_instruction.result, operation), + .select => |operation| try self.lowerSelect(block_id, source_instruction.result, operation), + .bitcast => |value_id| try self.lowerBitcast(block_id, source_instruction.result, value_id), + .load_interface => |operation| try self.lowerLoadInterface(block_id, source_instruction.result, operation), + .store_interface => |operation| try self.lowerStoreInterface(block_id, source_instruction.result, operation), + .composite_construct => |operation| try self.lowerCompositeConstruct(source_instruction.result, operation), + .composite_extract => |operation| try self.lowerCompositeExtract(source_instruction.result, operation), + .load_buffer => |operation| try self.lowerLoadBuffer(block_id, source_instruction.result, operation), + .store_buffer => |operation| try self.lowerStoreBuffer(block_id, source_instruction.result, operation), + .call => return Error.UnsanitizedModule, + } + } + + fn requireResult(result: ?shader_ir.id.ValueId) Error!shader_ir.id.ValueId { + return result orelse Error.InvalidModule; + } + + fn requireNoResult(result: ?shader_ir.id.ValueId) Error!void { + if (result != null) + return Error.InvalidModule; + } + + fn lowerUnary(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.Unary) Error!void { + const result_id = try requireResult(result); + if (operation.opcode == .logical_not) { + const source_predicate = try self.predicate(operation.operand); + const inverted: PredicateValue = switch (source_predicate) { + .constant => |value| .{ .constant = !value }, + .dynamic => |value| .{ .dynamic = .{ + .flag = value.flag, + .inverse = !value.inverse, + } }, + }; + try self.putLocation(result_id, .{ .predicate = inverted }); + return; + } + + const source_components = try self.components(operation.operand); + const result_components = try self.addRegisterLocation(result_id, .temporary); + if (source_components.len != result_components.len) + return Error.InvalidModule; + + for (source_components, result_components) |source_component, result_component| { + if (source_component.type != result_component.type) + return Error.InvalidModule; + switch (operation.opcode) { + .negate => { + if (source_component.type != .i32 and source_component.type != .f32) + return Error.UnsupportedOperation; + var negated = source_component; + negated.negate = !negated.negate; + try self.appendMove(block_id, null, try destinationFromSource(result_component), negated); + }, + .bitwise_not => { + const all_ones: operand.Immediate = switch (source_component.type) { + .u32 => .{ .u32 = std.math.maxInt(u32) }, + .i32 => .{ .i32 = -1 }, + else => return Error.UnsupportedOperation, + }; + try self.appendInstruction(block_id, null, .{ + .binary = .{ + .opcode = .bitwise_xor, + .destination = try destinationFromSource(result_component), + .lhs = source_component, + .rhs = .{ + .register = .{ .immediate = all_ones }, + .type = source_component.type, + .region = operand.Region.broadcast(), + }, + }, + }); + }, + .logical_not => unreachable, + } + } + } + + fn lowerBinary(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.Binary) Error!void { + const result_id = try requireResult(result); + const lhs_components = try self.components(operation.lhs); + const rhs_components = try self.components(operation.rhs); + const result_components = try self.addRegisterLocation(result_id, .temporary); + if (lhs_components.len == 0 or lhs_components.len != rhs_components.len or lhs_components.len != result_components.len) + return Error.InvalidModule; + + const data_type = lhs_components[0].type; + const opcode: instruction.BinaryOpcode = switch (operation.opcode) { + .integer_add => if (data_type == .u32 or data_type == .i32) .add else return Error.UnsupportedOperation, + .float_add => if (data_type == .f32) .add else return Error.UnsupportedOperation, + .integer_subtract => if (data_type == .u32 or data_type == .i32) .add else return Error.UnsupportedOperation, + .float_subtract => if (data_type == .f32) .add else return Error.UnsupportedOperation, + .integer_multiply => if (data_type == .u32 or data_type == .i32) .multiply else return Error.UnsupportedOperation, + .float_multiply => if (data_type == .f32) .multiply else return Error.UnsupportedOperation, + .shift_left => if (data_type == .u32 or data_type == .i32) .shift_left else return Error.UnsupportedOperation, + .logical_shift_right => if (data_type == .u32) .shift_right else return Error.UnsupportedOperation, + .arithmetic_shift_right => if (data_type == .i32) .shift_right else return Error.UnsupportedOperation, + .bitwise_and => if (data_type == .u32 or data_type == .i32) .bitwise_and else return Error.UnsupportedOperation, + .bitwise_or => if (data_type == .u32 or data_type == .i32) .bitwise_or else return Error.UnsupportedOperation, + .bitwise_xor => if (data_type == .u32 or data_type == .i32) .bitwise_xor else return Error.UnsupportedOperation, + .unsigned_divide, + .signed_divide, + .unsigned_modulo, + .signed_modulo, + .float_divide, + .float_modulo, + .logical_and, + .logical_or, + => return Error.UnsupportedOperation, + }; + + for (lhs_components, rhs_components, result_components) |lhs, rhs_value, result_component| { + if (lhs.type != data_type or rhs_value.type != data_type or result_component.type != data_type) + return Error.InvalidModule; + var rhs = rhs_value; + if (operation.opcode == .integer_subtract or operation.opcode == .float_subtract) + rhs.negate = !rhs.negate; + try self.appendInstruction(block_id, null, .{ + .binary = .{ + .opcode = opcode, + .destination = try destinationFromSource(result_component), + .lhs = lhs, + .rhs = rhs, + }, + }); + } + } + + fn lowerCompare(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.Compare) Error!void { + const result_id = try requireResult(result); + const result_value = self.lowerer.module.values.get(result_id) orelse return Error.InvalidModule; + + if (!try self.isBoolean(result_value.type)) + return Error.InvalidModule; + + const lhs_components = try self.components(operation.lhs); + const rhs_components = try self.components(operation.rhs); + if (lhs_components.len != 1 or rhs_components.len != 1) + return Error.UnsupportedOperation; + const lhs = lhs_components[0]; + const rhs = rhs_components[0]; + if (lhs.type != rhs.type) + return Error.InvalidModule; + + const opcode: instruction.CompareOpcode = switch (operation.opcode) { + .equal => if (lhs.type == .u32 or lhs.type == .i32) .equal else return Error.UnsupportedOperation, + .not_equal => if (lhs.type == .u32 or lhs.type == .i32) .not_equal else return Error.UnsupportedOperation, + .unsigned_less => if (lhs.type == .u32) .less_than else return Error.UnsupportedOperation, + .signed_less => if (lhs.type == .i32) .less_than else return Error.UnsupportedOperation, + .ordered_float_equal, + .unordered_float_equal, + .ordered_float_not_equal, + .unordered_float_not_equal, + .ordered_float_less, + .unordered_float_less, + => return Error.UnsupportedOperation, + }; + + const flag_id = self.builder.addVirtualFlag(.{ .name = result_value.name }) catch |err| + return mapProgramError(err); + + const predicate_value: operand.Predicate = .{ .flag = .{ .virtual = flag_id } }; + try self.putLocation(result_id, .{ .predicate = .{ .dynamic = predicate_value } }); + try self.appendInstruction(block_id, null, .{ + .compare = .{ + .opcode = opcode, + .destination = predicate_value.flag, + .lhs = lhs, + .rhs = rhs, + }, + }); + } + + fn lowerSelect(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.Select) Error!void { + const result_id = try requireResult(result); + const true_components = try self.components(operation.true_value); + const false_components = try self.components(operation.false_value); + const result_components = try self.addRegisterLocation(result_id, .temporary); + if (true_components.len != false_components.len or true_components.len != result_components.len) + return Error.InvalidModule; + + const condition = try self.predicate(operation.condition); + for (true_components, false_components, result_components) |true_value, false_value, result_component| { + const destination_value = try destinationFromSource(result_component); + if (true_value.type != destination_value.type or false_value.type != destination_value.type) + return Error.InvalidModule; + + switch (condition) { + .constant => |constant| try self.appendMove( + block_id, + null, + destination_value, + if (constant) true_value else false_value, + ), + .dynamic => |dynamic| { + try self.appendMove(block_id, .{ + .flag = dynamic.flag, + .inverse = !dynamic.inverse, + }, destination_value, false_value); + try self.appendMove(block_id, dynamic, destination_value, true_value); + }, + } + } + } + + fn lowerBitcast(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, source_id: shader_ir.id.ValueId) Error!void { + const result_id = try requireResult(result); + const result_value = self.lowerer.module.values.get(result_id) orelse return Error.InvalidModule; + const target_type = try self.lowerType(result_value.type); + const source_components = try self.components(source_id); + const result_components = try self.addRegisterLocation(result_id, .temporary); + if (source_components.len != target_type.component_count or source_components.len != result_components.len) + return Error.UnsupportedOperation; + + for (source_components, result_components) |source_component, result_component| { + // The source operand type selects the reinterpretation used by the + // move; the target-typed register materializes it before any CFG edge. + var cast_source = source_component; + cast_source.register = switch (cast_source.register) { + .immediate => |immediate| .{ .immediate = bitcastImmediate(immediate, target_type.element_type) }, + else => cast_source.register, + }; + cast_source.type = target_type.element_type; + try self.appendMove(block_id, null, try destinationFromSource(result_component), cast_source); + } + } + + fn lowerCompositeConstruct(self: *LoweringState, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.CompositeConstruct) Error!void { + const result_id = try requireResult(result); + const result_value = self.lowerer.module.values.get(result_id) orelse return Error.InvalidModule; + const result_type = try self.lowerType(result_value.type); + if (result_type.component_count < 2 or operation.elements.len != result_type.component_count) + return Error.UnsupportedOperation; + + const result_components = try self.storage.alloc(operand.Source, result_type.component_count); + for (operation.elements, result_components) |element_id, *component| { + const element_components = try self.components(element_id); + if (element_components.len != 1 or element_components[0].type != result_type.element_type) + return Error.InvalidModule; + component.* = element_components[0]; + } + try self.putLocation(result_id, .{ .components = result_components }); + } + + fn lowerCompositeExtract(self: *LoweringState, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.CompositeExtract) Error!void { + const result_id = try requireResult(result); + if (operation.indices.len != 1) + return Error.UnsupportedOperation; + const source_components = try self.components(operation.composite); + const component_index: usize = operation.indices[0]; + if (component_index >= source_components.len) + return Error.InvalidModule; + + const result_value = self.lowerer.module.values.get(result_id) orelse return Error.InvalidModule; + const result_type = try self.lowerType(result_value.type); + if (result_type.component_count != 1 or result_type.element_type != source_components[component_index].type) + return Error.InvalidModule; + try self.putLocation(result_id, .{ .components = source_components[component_index .. component_index + 1] }); + } + + fn lowerLoadInterface(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.LoadInterface) Error!void { + const result_id = try requireResult(result); + if (operation.element_index != null) + return Error.UnsupportedOperation; + + const variable = self.lowerer.module.interface_variables.get(operation.variable) orelse return Error.InvalidModule; + if (variable.direction != .input) + return Error.InvalidModule; + switch (variable.semantic) { + .builtin => |builtin| if (builtin != .global_invocation_id) + return Error.UnsupportedOperation, + .location => return Error.UnsupportedOperation, + } + + const result_value = self.lowerer.module.values.get(result_id) orelse return Error.InvalidModule; + if (result_value.type != variable.type) + return Error.InvalidModule; + const result_components = try self.addRegisterLocation(result_id, .temporary); + if (result_components.len != 3) + return Error.UnsupportedOperation; + for (result_components, 0..) |result_component, component_index| { + if (result_component.type != .u32) + return Error.UnsupportedOperation; + try self.appendInstruction(block_id, null, .{ + .load_global_invocation_id = .{ + .destination = try destinationFromSource(result_component), + .component = @intCast(component_index), + }, + }); + } + } + + fn lowerStoreInterface(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.StoreInterface) Error!void { + _ = self; + _ = block_id; + _ = operation; + try requireNoResult(result); + return Error.UnsupportedOperation; + } + + fn lowerLoadBuffer(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.LoadBuffer) Error!void { + const result_id = try requireResult(result); + const byte_offset = try self.source(operation.byte_offset); + if (byte_offset.type != .u32) + return Error.UnsupportedType; + const buffer = try self.storageBuffer(operation.resource); + const result_components = try self.addRegisterLocation(result_id, .temporary); + for (result_components, 0..) |result_component, component_index| { + try self.appendInstruction(block_id, null, .{ + .load_buffer = .{ + .destination = try destinationFromSource(result_component), + .buffer = buffer, + .byte_offset = byte_offset, + .immediate_offset = @intCast(component_index * result_component.type.sizeBytes()), + }, + }); + } + } + + fn lowerStoreBuffer(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.StoreBuffer) Error!void { + try requireNoResult(result); + const byte_offset = try self.source(operation.byte_offset); + if (byte_offset.type != .u32) + return Error.UnsupportedType; + const buffer = try self.storageBuffer(operation.resource); + const source_components = try self.components(operation.value); + for (source_components, 0..) |source_component, component_index| { + try self.appendInstruction(block_id, null, .{ + .store_buffer = .{ + .buffer = buffer, + .byte_offset = byte_offset, + .immediate_offset = @intCast(component_index * source_component.type.sizeBytes()), + .source = source_component, + }, + }); + } + } + + fn lowerControlAndTerminators(self: *LoweringState, allocator: std.mem.Allocator) Error!void { + const source_entry = try self.sourceEntryFunction(); + const function = source_entry[1]; + + for (function.blocks.items) |source_block_id| { + const source_block = self.lowerer.module.blocks.get(source_block_id) orelse return Error.InvalidModule; + const target_block_id = try self.mappedBlock(source_block_id); + const structured_control: instruction.StructuredControl = switch (source_block.structured_control) { + .none => .none, + .selection => |selection| .{ .selection = .{ + .merge_block = try self.mappedBlock(selection.merge_block), + } }, + .loop => |loop| .{ .loop = .{ + .merge_block = try self.mappedBlock(loop.merge_block), + .continue_block = try self.mappedBlock(loop.continue_block), + } }, + }; + self.builder.setStructuredControl(target_block_id, structured_control) catch |err| + return mapProgramError(err); + + const source_terminator = source_block.terminator orelse return Error.InvalidModule; + const target_terminator: instruction.Terminator = switch (source_terminator) { + .branch => |edge| .{ .jump = try self.lowerEdge(allocator, edge) }, + .conditional_branch => |branch| conditional: { + switch (try self.predicate(branch.condition)) { + .constant => |condition| { + const edge = if (condition) branch.true_edge else branch.false_edge; + break :conditional .{ .jump = try self.lowerEdge(allocator, edge) }; + }, + .dynamic => |condition| { + const true_edge = try self.lowerEdge(allocator, branch.true_edge); + errdefer allocator.free(true_edge.arguments); + const false_edge = try self.lowerEdge(allocator, branch.false_edge); + break :conditional .{ .conditional_branch = .{ + .predicate = condition, + .true_edge = true_edge, + .false_edge = false_edge, + } }; + }, + } + }, + .return_void => .end_thread, + .return_value => return Error.InvalidEntryPoint, + .discard => return Error.UnsupportedTerminator, + .@"unreachable" => .@"unreachable", + }; + defer freeTerminatorArguments(allocator, target_terminator); + self.builder.setTerminator(target_block_id, target_terminator) catch |err| + return mapProgramError(err); + } + } + + fn lowerEdge(self: *LoweringState, allocator: std.mem.Allocator, edge: shader_ir.module.Edge) Error!instruction.Edge { + const target_source_block = self.lowerer.module.blocks.get(edge.target) orelse return Error.InvalidModule; + if (edge.arguments.len != target_source_block.parameters.items.len) + return Error.InvalidModule; + + var arguments: std.ArrayList(pseudo.EdgeArgument) = .empty; + defer arguments.deinit(allocator); + for (edge.arguments) |argument_id| { + switch (try self.location(argument_id)) { + .components => |bundle| for (bundle) |component| + try arguments.append(allocator, .{ .source = component }), + .predicate => |predicate_value| try arguments.append(allocator, .{ .predicate = predicate_value }), + } + } + + return .{ + .target = try self.mappedBlock(edge.target), + .arguments = try arguments.toOwnedSlice(allocator), + }; + } +}; + +fn freeTerminatorArguments(allocator: std.mem.Allocator, terminator: instruction.Terminator) void { + switch (terminator) { + .jump => |edge| allocator.free(edge.arguments), + .conditional_branch => |branch| { + allocator.free(branch.true_edge.arguments); + allocator.free(branch.false_edge.arguments); + }, + else => {}, + } +} + +pub const Lowerer = struct { + module: *shader_ir.module.Module, + device_info: device.DeviceInfo, + options: Options, + + pub fn init(module: *shader_ir.module.Module, device_info: device.DeviceInfo, options: Options) Lowerer { + return .{ + .module = module, + .device_info = device_info, + .options = options, + }; + } + + pub fn lower(self: *Lowerer, allocator: std.mem.Allocator) Error!program_ir.Program { + shader_ir.validator.validate(self.module) catch |err| return switch (err) { + error.OutOfMemory => Error.OutOfMemory, + error.MissingEntryPoint => Error.MissingEntryPoint, + error.InvalidEntryPoint => Error.InvalidEntryPoint, + else => Error.InvalidModule, + }; + + var transformer_manager = shader_ir.transformer_manager.Manager.init(allocator); + defer transformer_manager.deinit(); + transformer_manager.add(shader_ir.inline_all_functions.transformer) catch return Error.OutOfMemory; + + var transformer_context: shader_ir.transformer_manager.Context = .{ .allocator = allocator }; + _ = transformer_manager.run(self.module, &transformer_context) catch |err| return switch (err) { + error.OutOfMemory => Error.OutOfMemory, + else => Error.SanitizationFailed, + }; + if (!self.module.properties.no_function_calls) + return Error.UnsanitizedModule; + + if (self.module.stage != .compute) + return Error.UnsupportedStage; + const workgroup_size = self.module.execution_modes.workgroup_size orelse return Error.MissingWorkgroupSize; + if (workgroup_size[0] == 0 or workgroup_size[1] == 0 or workgroup_size[2] == 0) + return Error.InvalidModule; + + var program = program_ir.Program.init(allocator, workgroup_size, self.device_info, self.options.dispatch_width); + errdefer program.deinit(); + + const block_map = try allocator.alloc(?ids.BlockId, self.module.blocks.entries.items.len); + defer allocator.free(block_map); + @memset(block_map, null); + + const value_locations = try allocator.alloc(?ValueLocation, self.module.values.entries.items.len); + defer allocator.free(value_locations); + @memset(value_locations, null); + + const storage_buffer_map = try allocator.alloc(?ids.StorageBufferId, self.module.resources.entries.items.len); + defer allocator.free(storage_buffer_map); + @memset(storage_buffer_map, null); + + var state: LoweringState = .{ + .lowerer = self, + .builder = Builder.init(&program), + .storage = program.allocator(), + .block_map = block_map, + .value_locations = value_locations, + .storage_buffer_map = storage_buffer_map, + }; + + try state.lowerBlocks(); + try state.lowerParameters(); + try state.lowerInstructions(allocator); + try state.lowerControlAndTerminators(allocator); + + program.properties.common_ir_lowered = true; + validator.validate(&program) catch return Error.InvalidLoweredProgram; + + block_arguments.run(allocator, &program) catch |err| return switch (err) { + error.OutOfMemory => Error.OutOfMemory, + else => Error.InvalidLoweredProgram, + }; + validator.validate(&program) catch return Error.InvalidLoweredProgram; + return program; + } +}; + +fn bitcastImmediate(immediate: operand.Immediate, target_type: operand.DataType) operand.Immediate { + const bits: u32 = switch (immediate) { + .u32 => |value| value, + .i32 => |value| @bitCast(value), + .f32 => |value| @bitCast(value), + }; + return switch (target_type) { + .u32 => .{ .u32 = bits }, + .i32 => .{ .i32 = @bitCast(bits) }, + .f32 => .{ .f32 = @bitCast(bits) }, + else => unreachable, + }; +} + +fn mapProgramError(err: anyerror) Error { + return switch (err) { + Error.OutOfMemory => Error.OutOfMemory, + else => Error.InvalidLoweredProgram, + }; +} + +/// Convenience entry point for callers that do not need to retain a lowerer. +pub inline fn lower(allocator: std.mem.Allocator, module: *shader_ir.module.Module, device_info: device.DeviceInfo, options: Options) Error!program_ir.Program { + var lowerer = Lowerer.init(module, device_info, options); + return lowerer.lower(allocator); +} + +const test_device: device.DeviceInfo = .{ + .generation = .gen9, + .platform = .skylake, + .pci_device_id = 0x1912, + .grf_count = 128, +}; + +fn expectLowered(source: []const u8, expected: []const u8) !void { + var module = try shader_ir.parser.parseString(std.testing.allocator, source); + defer module.deinit(); + module.execution_modes.workgroup_size = .{ 1, 1, 1 }; + + var program = try lower(std.testing.allocator, &module, test_device, .{}); + defer program.deinit(); + try std.testing.expect(program.properties.common_ir_lowered); + try std.testing.expect(!program.properties.instructions_selected); + + const actual = try printer.allocPrint(std.testing.allocator, &program); + defer std.testing.allocator.free(actual); + try std.testing.expectEqualStrings(expected, actual); +} + +fn expectLoweredFragments(source: []const u8, expected: []const []const u8, unexpected: []const []const u8) !void { + var module = try shader_ir.parser.parseString(std.testing.allocator, source); + defer module.deinit(); + module.execution_modes.workgroup_size = .{ 1, 1, 1 }; + + var program = try lower(std.testing.allocator, &module, test_device, .{}); + defer program.deinit(); + try std.testing.expect(program.properties.common_ir_lowered); + try std.testing.expect(!program.properties.instructions_selected); + + const actual = try printer.allocPrint(std.testing.allocator, &program); + defer std.testing.allocator.free(actual); + + for (expected) |fragment| + try std.testing.expect(std.mem.indexOf(u8, actual, fragment) != null); + for (unexpected) |fragment| + try std.testing.expect(std.mem.indexOf(u8, actual, fragment) == null); +} + +fn expectLoweringError(source: []const u8, expected: Error) !void { + var module = try shader_ir.parser.parseString(std.testing.allocator, source); + defer module.deinit(); + module.execution_modes.workgroup_size = .{ 1, 1, 1 }; + + var program = lower(std.testing.allocator, &module, test_device, .{}) catch |actual| { + try std.testing.expectEqual(expected, actual); + return; + }; + defer program.deinit(); + return error.TestExpectedError; +} + +test "[ir] Lower: basic shader" { + const source = + \\shader compute @main + \\{ + \\ %one: constant u32 = bits(0x1) + \\ %two: constant u32 = bits(0x2) + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %sum: u32 = integer_add %one, %two + \\ %condition: bool = cmp_unsigned_less %one, %two + \\ conditional_branch %condition, .left(), .right() + \\ .left(): + \\ branch .merge(%sum) + \\ .right(): + \\ branch .merge(%two) + \\ .merge(%value: u32): + \\ return + \\ } + \\} + ; + + const expected = + \\; Flint compute program: + \\; .workgroup_size: [1, 1, 1] + \\; .generation: gen9 + \\; .platform: skylake + \\; .dispatch_width: simd8 + \\ + \\%value: vgrf u32[8], class(temporary), size(32), alignment(32), spillable + \\%sum: vgrf u32[8], class(temporary), size(32), alignment(32), spillable + \\%condition: vflag + \\ + \\.entry: + \\ [simd8] add %sum:u32, 1:u32, 2:u32 + \\ [simd8] cmp_less_than %condition, 1:u32, 2:u32 + \\ conditional_branch (+%condition), .left, .right + \\ + \\.left: + \\ jump .b4 + \\ + \\.right: + \\ jump .b5 + \\ + \\.merge: + \\ end_thread + \\ + \\.b4: + \\ [simd8] parallel_copy [%value:u32 <- %sum:u32] + \\ jump .merge + \\ + \\.b5: + \\ [simd8] parallel_copy [%value:u32 <- 2:u32] + \\ jump .merge + \\ + \\ + ; + + try expectLowered(source, expected); +} + +test "[ir] Lower: control flow" { + const source = + \\shader compute @main + \\{ + \\ %one: constant u32 = bits(0x1) + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ branch .producer() + \\ .producer(): + \\ %sum: u32 = integer_add %one, %one + \\ branch .merge() + \\ .merge(): + \\ %doubled: u32 = integer_add %sum, %one + \\ return + \\ } + \\} + ; + + const expected = + \\; Flint compute program: + \\; .workgroup_size: [1, 1, 1] + \\; .generation: gen9 + \\; .platform: skylake + \\; .dispatch_width: simd8 + \\ + \\%sum: vgrf u32[8], class(temporary), size(32), alignment(32), spillable + \\%doubled: vgrf u32[8], class(temporary), size(32), alignment(32), spillable + \\ + \\.entry: + \\ jump .producer + \\ + \\.producer: + \\ [simd8] add %sum:u32, 1:u32, 1:u32 + \\ jump .merge + \\ + \\.merge: + \\ [simd8] add %doubled:u32, %sum:u32, 1:u32 + \\ end_thread + \\ + \\ + ; + + try expectLowered(source, expected); +} + +test "[ir] Lower: function call" { + const source = + \\shader compute @main + \\{ + \\ %one: constant u32 = bits(0x1) + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %result: u32 = call @identity(%one) + \\ return + \\ } + \\ fn @identity(%value: u32) -> u32 + \\ { + \\ .entry(): + \\ return %value + \\ } + \\} + ; + + const expected = + \\; Flint compute program: + \\; .workgroup_size: [1, 1, 1] + \\; .generation: gen9 + \\; .platform: skylake + \\; .dispatch_width: simd8 + \\ + \\%result: vgrf u32[8], class(temporary), size(32), alignment(32), spillable + \\ + \\.entry: + \\ jump .b2 + \\ + \\.b1: + \\ end_thread + \\ + \\.b2: + \\ jump .b3 + \\ + \\.b3: + \\ [simd8] parallel_copy [%result:u32 <- 1:u32] + \\ jump .b1 + \\ + \\ + ; + + try expectLowered(source, expected); +} + +test "[ir] Lower: unary/binary operations" { + const source = + \\shader compute @main + \\{ + \\ %u_one: constant u32 = bits(0x1) + \\ %u_two: constant u32 = bits(0x2) + \\ %i_one: constant i32 = bits(0x1) + \\ %i_two: constant i32 = bits(0x2) + \\ %f_one: constant f32 = bits(0x3f800000) + \\ %f_two: constant f32 = bits(0x40000000) + \\ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %integer_negated: i32 = negate %i_one + \\ %float_negated: f32 = negate %f_one + \\ %inverted: u32 = bitwise_not %u_one + \\ %integer_difference: i32 = integer_subtract %i_one, %i_two + \\ %float_difference: f32 = float_subtract %f_one, %f_two + \\ %integer_product: u32 = integer_multiply %u_one, %u_two + \\ %float_product: f32 = float_multiply %f_one, %f_two + \\ %shifted_left: u32 = shift_left %u_one, %u_two + \\ %logical_right: u32 = logical_shift_right %u_two, %u_one + \\ %arithmetic_right: i32 = arithmetic_shift_right %i_two, %i_one + \\ %masked: u32 = bitwise_and %u_one, %u_two + \\ %combined: u32 = bitwise_or %u_one, %u_two + \\ %toggled: u32 = bitwise_xor %u_one, %u_two + \\ return + \\ } + \\} + ; + + try expectLoweredFragments(source, &.{ + "[simd8] mov %integer_negated:i32, -1:i32", + "[simd8] mov %float_negated:f32, -1:f32", + "[simd8] bitwise_xor %inverted:u32, 1:u32, 4294967295:u32", + "[simd8] add %integer_difference:i32, 1:i32, -2:i32", + "[simd8] add %float_difference:f32, 1:f32, -2:f32", + "[simd8] multiply %integer_product:u32, 1:u32, 2:u32", + "[simd8] multiply %float_product:f32, 1:f32, 2:f32", + "[simd8] shift_left %shifted_left:u32, 1:u32, 2:u32", + "[simd8] shift_right %logical_right:u32, 2:u32, 1:u32", + "[simd8] shift_right %arithmetic_right:i32, 2:i32, 1:i32", + "[simd8] bitwise_and %masked:u32, 1:u32, 2:u32", + "[simd8] bitwise_or %combined:u32, 1:u32, 2:u32", + "[simd8] bitwise_xor %toggled:u32, 1:u32, 2:u32", + }, &.{}); +} + +test "[ir] Lower: selects and bitcasts" { + const source = + \\shader compute @main + \\{ + \\ %always: constant bool = true + \\ %one: constant u32 = bits(0x1) + \\ %two: constant u32 = bits(0x2) + \\ %float_one: constant f32 = bits(0x3f800000) + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %condition: bool = cmp_unsigned_less %one, %two + \\ %dynamic_choice: u32 = select %condition, %one, %two + \\ %inverted_condition: bool = logical_not %condition + \\ %inverted_choice: u32 = select %inverted_condition, %one, %two + \\ %constant_choice: u32 = select %always, %one, %two + \\ %one_bits: u32 = bitcast %float_one + \\ %constant_sum: u32 = integer_add %one_bits, %one + \\ %negative: f32 = negate %float_one + \\ %negative_bits: u32 = bitcast %negative + \\ %register_sum: u32 = integer_add %negative_bits, %one + \\ return + \\ } + \\} + ; + + try expectLoweredFragments(source, &.{ + "[simd8] cmp_less_than %condition, 1:u32, 2:u32", + "[simd8] (-%condition) mov %dynamic_choice:u32, 2:u32", + "[simd8] (+%condition) mov %dynamic_choice:u32, 1:u32", + "[simd8] (+%condition) mov %inverted_choice:u32, 2:u32", + "[simd8] (-%condition) mov %inverted_choice:u32, 1:u32", + "[simd8] mov %constant_choice:u32, 1:u32", + "[simd8] mov %one_bits:u32, 1065353216:u32", + "[simd8] add %constant_sum:u32, %one_bits:u32, 1:u32", + "[simd8] mov %negative:f32, -1:f32", + "[simd8] mov %negative_bits:u32, %negative:u32", + "[simd8] add %register_sum:u32, %negative_bits:u32, 1:u32", + }, &.{}); +} + +test "[ir] Lower: global invocation ID" { + const source = + \\shader compute @main + \\{ + \\ @global_id: vec3[u32] = input[builtin(global_invocation_id)] + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %id: vec3[u32] = load_interface @global_id + \\ %x: u32 = composite_extract %id[0] + \\ return + \\ } + \\} + ; + + try expectLoweredFragments(source, &.{ + "%id_x: vgrf u32[8], class(temporary)", + "%id_z: vgrf u32[8], class(temporary)", + "[simd8] load_global_invocation_id %id_x:u32, component(0)", + "[simd8] load_global_invocation_id %id_y:u32, component(1)", + "[simd8] load_global_invocation_id %id_z:u32, component(2)", + }, &.{}); +} + +test "[ir] Lower: vector storage-buffer operations" { + const source = + \\shader compute @main + \\{ + \\ @source: vec4[u32] = storage_buffer[set(0), binding(1)] + \\ @destination: vec4[u32] = storage_buffer[set(0), binding(2)] + \\ %offset: constant u32 = 16 + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %value: vec4[u32] = load_buffer @source, %offset + \\ store_buffer @destination, %offset, %value + \\ return + \\ } + \\} + ; + + try expectLoweredFragments(source, &.{ + "@source = storage_buffer[set(0), binding(1)]", + "@destination = storage_buffer[set(0), binding(2)]", + "[simd8] load_buffer %value_x:u32, @source, 16:u32", + "[simd8] load_buffer %value_y:u32, @source, 16:u32, offset(4)", + "[simd8] load_buffer %value_w:u32, @source, 16:u32, offset(12)", + "[simd8] store_buffer @destination, 16:u32, %value_x:u32", + "[simd8] store_buffer @destination, 16:u32, offset(12), %value_w:u32", + }, &.{}); +} + +test "[ir] Lower: vector block parameter" { + const source = + \\shader compute @main + \\{ + \\ %one: constant u32 = bits(0x1) + \\ %two: constant u32 = bits(0x2) + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %pair: vec2[u32] = composite_construct %one, %two + \\ branch .merge(%pair) + \\ .merge(%merged: vec2[u32]): + \\ %first: u32 = composite_extract %merged[0] + \\ return + \\ } + \\} + ; + + try expectLoweredFragments(source, &.{ + "%merged_x: vgrf u32[8]", + "%merged_y: vgrf u32[8]", + "parallel_copy [%merged_x:u32 <- 1:u32, %merged_y:u32 <- 2:u32]", + }, &.{ + ".merge(", + }); +} + +test "[ir] Lower: reject non-compute interfaces" { + try expectLoweringError( + \\shader compute @main + \\{ + \\ @input: u32 = input[location(0), component(0), index(0)] + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %value: u32 = load_interface @input + \\ return + \\ } + \\} + , Error.UnsupportedOperation); +} + +test "[ir] Lower: constant conditional branch" { + const source = + \\shader compute @main + \\{ + \\ %always: constant bool = true + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ conditional_branch %always, .taken(), .untaken() + \\ .taken(): + \\ return + \\ .untaken(): + \\ return + \\ } + \\} + ; + + try expectLoweredFragments(source, &.{ + ".entry:\n jump .taken", + ".taken:\n end_thread", + ".untaken:\n end_thread", + }, &.{ + "conditional_branch", + "vflag", + }); +} + +test "[ir] Lower: boolean block parameter" { + const source = + \\shader compute @main + \\{ + \\ %one: constant u32 = bits(0x1) + \\ %two: constant u32 = bits(0x2) + \\ %never: constant bool = false + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %condition: bool = cmp_unsigned_less %one, %two + \\ conditional_branch %condition, .left(), .right() + \\ .left(): + \\ branch .merge(%condition) + \\ .right(): + \\ branch .merge(%never) + \\ .merge(%merged: bool): + \\ conditional_branch %merged, .taken(), .not_taken() + \\ .taken(): + \\ return + \\ .not_taken(): + \\ return + \\ } + \\} + ; + + try expectLoweredFragments(source, &.{ + "%condition: vflag", + "%merged: vflag", + "parallel_copy [%merged <- (+%condition)]", + "parallel_copy [%merged <- false]", + ".merge:\n conditional_branch (+%merged), .taken, .not_taken", + }, &.{}); +} + +test "[ir] Lower: unsupported operations" { + try expectLoweringError( + \\shader compute @main + \\{ + \\ %one: constant u32 = bits(0x1) + \\ %two: constant u32 = bits(0x2) + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %quotient: u32 = unsigned_divide %one, %two + \\ return + \\ } + \\} + , Error.UnsupportedOperation); + + try expectLoweringError( + \\shader compute @main + \\{ + \\ %one: constant u32 = bits(0x1) + \\ %two: constant u32 = bits(0x2) + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %wide: vec5[u32] = composite_construct %one, %two, %one, %two, %one + \\ return + \\ } + \\} + , Error.UnsupportedType); + + try expectLoweringError( + \\shader compute @main + \\{ + \\ %one: constant u16 = bits(0x1) + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ %sum: u16 = integer_add %one, %one + \\ return + \\ } + \\} + , Error.UnsupportedType); +} + +test "[ir] Lower: unreachable terminator" { + const source = + \\shader compute @main + \\{ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ unreachable + \\ } + \\} + ; + + try expectLoweredFragments(source, &.{ + ".entry:\n unreachable", + }, &.{}); +} + +test "[ir] Lower: common lowering is generation and dispatch-width agnostic" { + var module = try shader_ir.parser.parseString(std.testing.allocator, + \\shader compute @main + \\{ + \\ fn @main() -> void + \\ { + \\ .entry(): + \\ return + \\ } + \\} + ); + defer module.deinit(); + module.execution_modes.workgroup_size = .{ 1, 1, 1 }; + + var other_target = test_device; + other_target.generation = .gen11; + other_target.grf_size_bytes = 64; + other_target.supports_simd16 = true; + var program = try lower(std.testing.allocator, &module, other_target, .{ .dispatch_width = .simd16 }); + defer program.deinit(); + + try std.testing.expectEqual(device.Generation.gen11, program.device_info.generation); + try std.testing.expectEqual(device.DispatchWidth.simd16, program.dispatch_width); + try validator.validate(&program); +} diff --git a/src/intel/compiler/lower/lower.zig b/src/intel/compiler/lower/lower.zig index 23d23d9..744d425 100644 --- a/src/intel/compiler/lower/lower.zig +++ b/src/intel/compiler/lower/lower.zig @@ -1,1553 +1,2 @@ -const std = @import("std"); -const shader_compiler = @import("shader_ir"); -const shader_ir = shader_compiler.ir; -const device = @import("../device.zig"); -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 printer = @import("../ir/printer.zig"); -const pseudo = @import("../ir/pseudo.zig"); -const program_ir = @import("../ir/program.zig"); -const validator = @import("../ir/validator.zig"); - pub const block_arguments = @import("block_arguments.zig"); -pub const vertex_abi = @import("vertex_abi.zig"); - -pub const Options = struct { - dispatch_width: device.DispatchWidth = .simd8, -}; - -pub const Error = std.mem.Allocator.Error || error{ - MissingEntryPoint, - InvalidEntryPoint, - InvalidModule, - InvalidLoweredProgram, - SanitizationFailed, - UnsanitizedModule, - UnsupportedGeneration, - UnsupportedStage, - UnsupportedDispatchWidth, - UnsupportedType, - UnsupportedOperation, - UnsupportedTerminator, -}; - -const PredicateValue = pseudo.PredicateValue; - -const LoweredType = struct { - element_type: operand.DataType, - component_count: usize, -}; - -const ValueLocation = union(enum) { - components: []const operand.Source, - predicate: PredicateValue, -}; - -const LoweringState = struct { - lowerer: *Lowerer, - builder: Builder, - storage: std.mem.Allocator, - block_map: []?ids.BlockId, - value_locations: []?ValueLocation, - - fn lowerScalarType(self: *const LoweringState, type_id: shader_ir.id.TypeId) Error!operand.DataType { - const ty = self.lowerer.module.types.get(type_id) orelse return Error.InvalidModule; - return switch (ty.*) { - .integer => |integer| if (integer.bits == 32) - switch (integer.signedness) { - .unsigned => .u32, - .signed => .i32, - } - else - Error.UnsupportedType, - .floating => |floating| if (floating.bits == 32) .f32 else Error.UnsupportedType, - else => Error.UnsupportedType, - }; - } - - fn lowerType(self: *const LoweringState, type_id: shader_ir.id.TypeId) Error!LoweredType { - const ty = self.lowerer.module.types.get(type_id) orelse return Error.InvalidModule; - return switch (ty.*) { - .integer, .floating => .{ - .element_type = try self.lowerScalarType(type_id), - .component_count = 1, - }, - .vector => |vector| if (vector.length >= 2 and vector.length <= 4) - .{ - .element_type = try self.lowerScalarType(vector.element_type), - .component_count = vector.length, - } - else - Error.UnsupportedType, - else => Error.UnsupportedType, - }; - } - - fn isBoolean(self: *const LoweringState, type_id: shader_ir.id.TypeId) Error!bool { - const ty = self.lowerer.module.types.get(type_id) orelse return Error.InvalidModule; - return ty.* == .boolean; - } - - fn mappedBlock(self: *const LoweringState, source_id: shader_ir.id.BlockId) Error!ids.BlockId { - if (source_id.index() >= self.block_map.len) - return Error.InvalidModule; - return self.block_map[source_id.index()] orelse Error.InvalidModule; - } - - fn putLocation(self: *LoweringState, value_id: shader_ir.id.ValueId, new_location: ValueLocation) Error!void { - if (value_id.index() >= self.value_locations.len or self.value_locations[value_id.index()] != null) - return Error.InvalidModule; - self.value_locations[value_id.index()] = new_location; - } - - fn addRegister(self: *LoweringState, data_type: operand.DataType, class: operand.RegisterClass, name: ?[]const u8) Error!ids.VirtualRegisterId { - return self.builder.addVirtualRegister(.{ - .size_bytes = @as(u32, data_type.sizeBytes()) * @intFromEnum(self.lowerer.options.dispatch_width), - .alignment_bytes = self.lowerer.device_info.grf_size_bytes, - .element_type = data_type, - .lane_count = @intFromEnum(self.lowerer.options.dispatch_width), - .class = class, - .name = name, - }) catch |err| return mapProgramError(err); - } - - fn registerSource(self: *const LoweringState, register_id: ids.VirtualRegisterId, data_type: operand.DataType) operand.Source { - _ = self; - return .{ - .register = .{ .virtual = register_id }, - .type = data_type, - .region = operand.Region.contiguous(.simd8), - }; - } - - fn componentName(self: *LoweringState, name: ?[]const u8, component_index: usize, component_count: usize) Error!?[]const u8 { - if (name == null or component_count == 1) - return name; - const suffixes = "xyzw"; - const formatted = try std.fmt.allocPrint(self.storage, "{s}_{c}", .{ name.?, suffixes[component_index] }); - return @as([]const u8, formatted); - } - - fn addRegisterLocation(self: *LoweringState, value_id: shader_ir.id.ValueId, class: operand.RegisterClass) Error![]const operand.Source { - const value = self.lowerer.module.values.get(value_id) orelse return Error.InvalidModule; - const lowered_type = try self.lowerType(value.type); - const result = try self.storage.alloc(operand.Source, lowered_type.component_count); - for (result, 0..) |*component, component_index| { - const register_id = try self.addRegister( - lowered_type.element_type, - class, - try self.componentName(value.name, component_index, lowered_type.component_count), - ); - component.* = self.registerSource(register_id, lowered_type.element_type); - } - try self.putLocation(value_id, .{ .components = result }); - return result; - } - - fn location(self: *LoweringState, value_id: shader_ir.id.ValueId) Error!ValueLocation { - if (value_id.index() >= self.value_locations.len) - return Error.InvalidModule; - - if (self.value_locations[value_id.index()]) |existing| - return existing; - - const value = self.lowerer.module.values.get(value_id) orelse return Error.InvalidModule; - switch (value.definition) { - .constant => |constant_id| { - const constant = self.lowerer.module.constants.get(constant_id) orelse return Error.InvalidModule; - if (constant.type != value.type) - return Error.InvalidModule; - - const result: ValueLocation = if (try self.isBoolean(value.type)) switch (constant.value) { - .boolean => |boolean| .{ .predicate = .{ .constant = boolean } }, - else => return Error.UnsupportedType, - } else .{ - .components = try self.constantComponents(value.type, constant.value), - }; - self.value_locations[value_id.index()] = result; - return result; - }, - .undef => { - if (try self.isBoolean(value.type)) - return Error.UnsupportedType; - _ = try self.addRegisterLocation(value_id, .temporary); - return self.value_locations[value_id.index()].?; - }, - else => return Error.InvalidModule, - } - } - - fn components(self: *LoweringState, value_id: shader_ir.id.ValueId) Error![]const operand.Source { - return switch (try self.location(value_id)) { - .components => |values| values, - .predicate => Error.UnsupportedType, - }; - } - - fn source(self: *LoweringState, value_id: shader_ir.id.ValueId) Error!operand.Source { - const values = try self.components(value_id); - if (values.len != 1) - return Error.UnsupportedType; - return values[0]; - } - - fn predicate(self: *LoweringState, value_id: shader_ir.id.ValueId) Error!PredicateValue { - return switch (try self.location(value_id)) { - .components => Error.UnsupportedType, - .predicate => |value| value, - }; - } - - fn destinationFromSource(source_value: operand.Source) Error!operand.Destination { - if (source_value.negate or source_value.absolute) - return Error.InvalidLoweredProgram; - - return switch (source_value.register) { - .virtual => .{ - .register = source_value.register, - .type = source_value.type, - .region = .{ .byte_offset = source_value.region.byte_offset }, - }, - else => Error.InvalidLoweredProgram, - }; - } - - fn destination(self: *LoweringState, value_id: shader_ir.id.ValueId) Error!operand.Destination { - return destinationFromSource(try self.source(value_id)); - } - - fn constantComponents(self: *LoweringState, type_id: shader_ir.id.TypeId, value: shader_ir.constant.ConstantValue) Error![]const operand.Source { - const lowered_type = try self.lowerType(type_id); - const result = try self.storage.alloc(operand.Source, lowered_type.component_count); - if (lowered_type.component_count == 1) { - result[0] = try self.constantScalarSource(type_id, value); - return result; - } - - const ty = self.lowerer.module.types.get(type_id) orelse return Error.InvalidModule; - const vector = switch (ty.*) { - .vector => |vector| vector, - else => return Error.InvalidModule, - }; - switch (value) { - .composite => |elements| { - if (elements.len != lowered_type.component_count) - return Error.InvalidModule; - for (elements, result) |constant_id, *component| { - const element = self.lowerer.module.constants.get(constant_id) orelse return Error.InvalidModule; - if (element.type != vector.element_type) - return Error.InvalidModule; - component.* = try self.constantScalarSource(element.type, element.value); - } - }, - .null => { - const zero: shader_ir.constant.ConstantValue = switch (lowered_type.element_type) { - .u32, .i32 => .{ .integer_bits = 0 }, - .f32 => .{ .float_bits = 0 }, - else => unreachable, - }; - for (result) |*component| - component.* = try self.constantScalarSource(vector.element_type, zero); - }, - else => return Error.UnsupportedType, - } - return result; - } - - fn constantScalarSource(self: *const LoweringState, type_id: shader_ir.id.TypeId, value: shader_ir.constant.ConstantValue) Error!operand.Source { - const data_type = try self.lowerScalarType(type_id); - const immediate: operand.Immediate = switch (data_type) { - .u32 => switch (value) { - .integer_bits => |bits| .{ .u32 = @truncate(bits) }, - else => return Error.UnsupportedType, - }, - .i32 => switch (value) { - .integer_bits => |bits| .{ .i32 = @bitCast(@as(u32, @truncate(bits))) }, - else => return Error.UnsupportedType, - }, - .f32 => switch (value) { - .float_bits => |bits| .{ .f32 = @bitCast(@as(u32, @truncate(bits))) }, - else => return Error.UnsupportedType, - }, - else => unreachable, - }; - - return .{ - .register = .{ .immediate = immediate }, - .type = data_type, - .region = operand.Region.broadcast(), - }; - } - - fn appendInstruction(self: *LoweringState, block_id: ids.BlockId, predicate_value: ?operand.Predicate, operation: instruction.Operation) Error!void { - _ = self.builder.appendInstruction(block_id, .simd8, predicate_value, operation) catch |err| - return mapProgramError(err); - } - - fn appendMove(self: *LoweringState, block_id: ids.BlockId, predicate_value: ?operand.Predicate, destination_value: operand.Destination, source_value: operand.Source) Error!void { - try self.appendInstruction(block_id, predicate_value, .{ - .move = .{ - .destination = destination_value, - .source = source_value, - }, - }); - } - - fn sourceEntryFunction(self: *const LoweringState) Error!struct { shader_ir.id.FunctionId, *const shader_ir.module.Function } { - const source_entry = self.lowerer.module.entry_point orelse return Error.MissingEntryPoint; - const function = self.lowerer.module.functions.get(source_entry) orelse return Error.InvalidEntryPoint; - const return_type = self.lowerer.module.types.get(function.return_type) orelse return Error.InvalidModule; - if (return_type.* != .void or function.parameters.items.len != 0) - return Error.InvalidEntryPoint; - return .{ source_entry, function }; - } - - fn lowerBlocks(self: *LoweringState) Error!void { - const source_function_id, const function = try self.sourceEntryFunction(); - - for (function.blocks.items) |source_block_id| { - const source_block = self.lowerer.module.blocks.get(source_block_id) orelse return Error.InvalidModule; - if (source_block.parent_function != source_function_id) - return Error.InvalidModule; - const target_block_id = self.builder.addBlock(source_block.name) catch |err| - return mapProgramError(err); - if (source_block_id.index() >= self.block_map.len or self.block_map[source_block_id.index()] != null) - return Error.InvalidModule; - self.block_map[source_block_id.index()] = target_block_id; - } - - const source_entry = function.entry_block orelse return Error.InvalidModule; - self.builder.setEntryBlock(try self.mappedBlock(source_entry)) catch |err| return mapProgramError(err); - } - - fn lowerParameters(self: *LoweringState) Error!void { - const source_entry = try self.sourceEntryFunction(); - const function = source_entry[1]; - - for (function.blocks.items) |source_block_id| { - const source_block = self.lowerer.module.blocks.get(source_block_id) orelse return Error.InvalidModule; - const target_block_id = try self.mappedBlock(source_block_id); - - for (source_block.parameters.items) |parameter_id| { - const value = self.lowerer.module.values.get(parameter_id) orelse return Error.InvalidModule; - if (try self.isBoolean(value.type)) { - const flag_id = self.builder.addVirtualFlag(.{ .name = value.name }) catch |err| - return mapProgramError(err); - const predicate_value: operand.Predicate = .{ .flag = .{ .virtual = flag_id } }; - try self.putLocation(parameter_id, .{ .predicate = .{ .dynamic = predicate_value } }); - self.builder.addBlockParameter(target_block_id, .{ .flag = flag_id }) catch |err| - return mapProgramError(err); - } else { - const parameter_components = try self.addRegisterLocation(parameter_id, .temporary); - for (parameter_components) |parameter_source| { - const register_id = switch (parameter_source.register) { - .virtual => |id| id, - else => return Error.InvalidLoweredProgram, - }; - self.builder.addBlockParameter(target_block_id, .{ .register = register_id }) catch |err| - return mapProgramError(err); - } - } - } - } - } - - fn lowerInstructions(self: *LoweringState, allocator: std.mem.Allocator) Error!void { - const visited = try allocator.alloc(bool, self.lowerer.module.blocks.entries.items.len); - defer allocator.free(visited); - @memset(visited, false); - - const source_entry = try self.sourceEntryFunction(); - const function = source_entry[1]; - try self.lowerBlockInstructions(function.entry_block orelse return Error.InvalidModule, visited); - - for (function.blocks.items) |source_block_id| { - if (!visited[source_block_id.index()]) - try self.lowerBlockInstructions(source_block_id, visited); - } - } - - fn lowerBlockInstructions(self: *LoweringState, source_block_id: shader_ir.id.BlockId, visited: []bool) Error!void { - if (source_block_id.index() >= visited.len) - return Error.InvalidModule; - - if (visited[source_block_id.index()]) - return; - - visited[source_block_id.index()] = true; - - const block = self.lowerer.module.blocks.get(source_block_id) orelse return Error.InvalidModule; - const target_block_id = try self.mappedBlock(source_block_id); - for (block.instructions.items) |instruction_id| { - const source_instruction = self.lowerer.module.instructions.get(instruction_id) orelse return Error.InvalidModule; - if (source_instruction.parent_block != source_block_id) - return Error.InvalidModule; - try self.lowerInstruction(target_block_id, source_instruction.*); - } - - switch (block.terminator orelse return Error.InvalidModule) { - .branch => |edge| try self.lowerBlockInstructions(edge.target, visited), - .conditional_branch => |branch| { - try self.lowerBlockInstructions(branch.true_edge.target, visited); - try self.lowerBlockInstructions(branch.false_edge.target, visited); - }, - else => {}, - } - } - - fn lowerInstruction(self: *LoweringState, block_id: ids.BlockId, source_instruction: shader_ir.instruction.Instruction) Error!void { - switch (source_instruction.operation) { - .unary => |operation| try self.lowerUnary(block_id, source_instruction.result, operation), - .binary => |operation| try self.lowerBinary(block_id, source_instruction.result, operation), - .compare => |operation| try self.lowerCompare(block_id, source_instruction.result, operation), - .select => |operation| try self.lowerSelect(block_id, source_instruction.result, operation), - .bitcast => |value_id| try self.lowerBitcast(block_id, source_instruction.result, value_id), - .load_interface => |operation| try self.lowerLoadInterface(block_id, source_instruction.result, operation), - .store_interface => |operation| try self.lowerStoreInterface(block_id, source_instruction.result, operation), - .composite_construct => |operation| try self.lowerCompositeConstruct(source_instruction.result, operation), - .composite_extract => |operation| try self.lowerCompositeExtract(source_instruction.result, operation), - .load_buffer, .store_buffer => return Error.UnsupportedOperation, - .call => return Error.UnsanitizedModule, - } - } - - fn requireResult(result: ?shader_ir.id.ValueId) Error!shader_ir.id.ValueId { - return result orelse Error.InvalidModule; - } - - fn requireNoResult(result: ?shader_ir.id.ValueId) Error!void { - if (result != null) - return Error.InvalidModule; - } - - fn lowerUnary(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.Unary) Error!void { - const result_id = try requireResult(result); - if (operation.opcode == .logical_not) { - const source_predicate = try self.predicate(operation.operand); - const inverted: PredicateValue = switch (source_predicate) { - .constant => |value| .{ .constant = !value }, - .dynamic => |value| .{ .dynamic = .{ - .flag = value.flag, - .inverse = !value.inverse, - } }, - }; - try self.putLocation(result_id, .{ .predicate = inverted }); - return; - } - - const source_components = try self.components(operation.operand); - const result_components = try self.addRegisterLocation(result_id, .temporary); - if (source_components.len != result_components.len) - return Error.InvalidModule; - - for (source_components, result_components) |source_component, result_component| { - if (source_component.type != result_component.type) - return Error.InvalidModule; - switch (operation.opcode) { - .negate => { - if (source_component.type != .i32 and source_component.type != .f32) - return Error.UnsupportedOperation; - var negated = source_component; - negated.negate = !negated.negate; - try self.appendMove(block_id, null, try destinationFromSource(result_component), negated); - }, - .bitwise_not => { - const all_ones: operand.Immediate = switch (source_component.type) { - .u32 => .{ .u32 = std.math.maxInt(u32) }, - .i32 => .{ .i32 = -1 }, - else => return Error.UnsupportedOperation, - }; - try self.appendInstruction(block_id, null, .{ - .binary = .{ - .opcode = .bitwise_xor, - .destination = try destinationFromSource(result_component), - .lhs = source_component, - .rhs = .{ - .register = .{ .immediate = all_ones }, - .type = source_component.type, - .region = operand.Region.broadcast(), - }, - }, - }); - }, - .logical_not => unreachable, - } - } - } - - fn lowerBinary(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.Binary) Error!void { - const result_id = try requireResult(result); - const lhs_components = try self.components(operation.lhs); - const rhs_components = try self.components(operation.rhs); - const result_components = try self.addRegisterLocation(result_id, .temporary); - if (lhs_components.len == 0 or lhs_components.len != rhs_components.len or lhs_components.len != result_components.len) - return Error.InvalidModule; - - const data_type = lhs_components[0].type; - const opcode: instruction.BinaryOpcode = switch (operation.opcode) { - .integer_add => if (data_type == .u32 or data_type == .i32) .add else return Error.UnsupportedOperation, - .float_add => if (data_type == .f32) .add else return Error.UnsupportedOperation, - .integer_subtract => if (data_type == .u32 or data_type == .i32) .add else return Error.UnsupportedOperation, - .float_subtract => if (data_type == .f32) .add else return Error.UnsupportedOperation, - .integer_multiply => if (data_type == .u32 or data_type == .i32) .multiply else return Error.UnsupportedOperation, - .float_multiply => if (data_type == .f32) .multiply else return Error.UnsupportedOperation, - .shift_left => if (data_type == .u32 or data_type == .i32) .shift_left else return Error.UnsupportedOperation, - .logical_shift_right => if (data_type == .u32) .shift_right else return Error.UnsupportedOperation, - .arithmetic_shift_right => if (data_type == .i32) .shift_right else return Error.UnsupportedOperation, - .bitwise_and => if (data_type == .u32 or data_type == .i32) .bitwise_and else return Error.UnsupportedOperation, - .bitwise_or => if (data_type == .u32 or data_type == .i32) .bitwise_or else return Error.UnsupportedOperation, - .bitwise_xor => if (data_type == .u32 or data_type == .i32) .bitwise_xor else return Error.UnsupportedOperation, - .unsigned_divide, - .signed_divide, - .unsigned_modulo, - .signed_modulo, - .float_divide, - .float_modulo, - .logical_and, - .logical_or, - => return Error.UnsupportedOperation, - }; - - for (lhs_components, rhs_components, result_components) |lhs, rhs_value, result_component| { - if (lhs.type != data_type or rhs_value.type != data_type or result_component.type != data_type) - return Error.InvalidModule; - var rhs = rhs_value; - if (operation.opcode == .integer_subtract or operation.opcode == .float_subtract) - rhs.negate = !rhs.negate; - try self.appendInstruction(block_id, null, .{ - .binary = .{ - .opcode = opcode, - .destination = try destinationFromSource(result_component), - .lhs = lhs, - .rhs = rhs, - }, - }); - } - } - - fn lowerCompare(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.Compare) Error!void { - const result_id = try requireResult(result); - const result_value = self.lowerer.module.values.get(result_id) orelse return Error.InvalidModule; - - if (!try self.isBoolean(result_value.type)) - return Error.InvalidModule; - - const lhs_components = try self.components(operation.lhs); - const rhs_components = try self.components(operation.rhs); - if (lhs_components.len != 1 or rhs_components.len != 1) - return Error.UnsupportedOperation; - const lhs = lhs_components[0]; - const rhs = rhs_components[0]; - if (lhs.type != rhs.type) - return Error.InvalidModule; - - const opcode: instruction.CompareOpcode = switch (operation.opcode) { - .equal => if (lhs.type == .u32 or lhs.type == .i32) .equal else return Error.UnsupportedOperation, - .not_equal => if (lhs.type == .u32 or lhs.type == .i32) .not_equal else return Error.UnsupportedOperation, - .unsigned_less => if (lhs.type == .u32) .less_than else return Error.UnsupportedOperation, - .signed_less => if (lhs.type == .i32) .less_than else return Error.UnsupportedOperation, - .ordered_float_equal, - .unordered_float_equal, - .ordered_float_not_equal, - .unordered_float_not_equal, - .ordered_float_less, - .unordered_float_less, - => return Error.UnsupportedOperation, - }; - - const flag_id = self.builder.addVirtualFlag(.{ .name = result_value.name }) catch |err| - return mapProgramError(err); - - const predicate_value: operand.Predicate = .{ .flag = .{ .virtual = flag_id } }; - try self.putLocation(result_id, .{ .predicate = .{ .dynamic = predicate_value } }); - try self.appendInstruction(block_id, null, .{ - .compare = .{ - .opcode = opcode, - .destination = predicate_value.flag, - .lhs = lhs, - .rhs = rhs, - }, - }); - } - - fn lowerSelect(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.Select) Error!void { - const result_id = try requireResult(result); - const true_components = try self.components(operation.true_value); - const false_components = try self.components(operation.false_value); - const result_components = try self.addRegisterLocation(result_id, .temporary); - if (true_components.len != false_components.len or true_components.len != result_components.len) - return Error.InvalidModule; - - const condition = try self.predicate(operation.condition); - for (true_components, false_components, result_components) |true_value, false_value, result_component| { - const destination_value = try destinationFromSource(result_component); - if (true_value.type != destination_value.type or false_value.type != destination_value.type) - return Error.InvalidModule; - - switch (condition) { - .constant => |constant| try self.appendMove( - block_id, - null, - destination_value, - if (constant) true_value else false_value, - ), - .dynamic => |dynamic| { - try self.appendMove(block_id, .{ - .flag = dynamic.flag, - .inverse = !dynamic.inverse, - }, destination_value, false_value); - try self.appendMove(block_id, dynamic, destination_value, true_value); - }, - } - } - } - - fn lowerBitcast(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, source_id: shader_ir.id.ValueId) Error!void { - const result_id = try requireResult(result); - const result_value = self.lowerer.module.values.get(result_id) orelse return Error.InvalidModule; - const target_type = try self.lowerType(result_value.type); - const source_components = try self.components(source_id); - const result_components = try self.addRegisterLocation(result_id, .temporary); - if (source_components.len != target_type.component_count or source_components.len != result_components.len) - return Error.UnsupportedOperation; - - for (source_components, result_components) |source_component, result_component| { - // The source operand type selects the reinterpretation used by the - // move; the target-typed register materializes it before any CFG edge. - var cast_source = source_component; - cast_source.register = switch (cast_source.register) { - .immediate => |immediate| .{ .immediate = bitcastImmediate(immediate, target_type.element_type) }, - else => cast_source.register, - }; - cast_source.type = target_type.element_type; - try self.appendMove(block_id, null, try destinationFromSource(result_component), cast_source); - } - } - - fn lowerCompositeConstruct(self: *LoweringState, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.CompositeConstruct) Error!void { - const result_id = try requireResult(result); - const result_value = self.lowerer.module.values.get(result_id) orelse return Error.InvalidModule; - const result_type = try self.lowerType(result_value.type); - if (result_type.component_count < 2 or operation.elements.len != result_type.component_count) - return Error.UnsupportedOperation; - - const result_components = try self.storage.alloc(operand.Source, result_type.component_count); - for (operation.elements, result_components) |element_id, *component| { - const element_components = try self.components(element_id); - if (element_components.len != 1 or element_components[0].type != result_type.element_type) - return Error.InvalidModule; - component.* = element_components[0]; - } - try self.putLocation(result_id, .{ .components = result_components }); - } - - fn lowerCompositeExtract(self: *LoweringState, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.CompositeExtract) Error!void { - const result_id = try requireResult(result); - if (operation.indices.len != 1) - return Error.UnsupportedOperation; - const source_components = try self.components(operation.composite); - const component_index: usize = operation.indices[0]; - if (component_index >= source_components.len) - return Error.InvalidModule; - - const result_value = self.lowerer.module.values.get(result_id) orelse return Error.InvalidModule; - const result_type = try self.lowerType(result_value.type); - if (result_type.component_count != 1 or result_type.element_type != source_components[component_index].type) - return Error.InvalidModule; - try self.putLocation(result_id, .{ .components = source_components[component_index .. component_index + 1] }); - } - - fn lowerLoadInterface(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.LoadInterface) Error!void { - const result_id = try requireResult(result); - - if (operation.element_index != null) - return Error.UnsupportedOperation; - - const variable = self.lowerer.module.interface_variables.get(operation.variable) orelse return Error.InvalidModule; - - if (variable.direction != .input) - return Error.InvalidModule; - - const result_value = self.lowerer.module.values.get(result_id) orelse return Error.InvalidModule; - - if (result_value.type != variable.type) - return Error.InvalidModule; - - const result_components = try self.addRegisterLocation(result_id, .varying); - for (result_components, 0..) |result_component, component_index| { - try self.appendInstruction(block_id, null, .{ - .load_input = .{ - .destination = try destinationFromSource(result_component), - .semantic = try lowerInterfaceSemantic(variable.semantic, @intCast(component_index)), - }, - }); - } - } - - fn lowerStoreInterface(self: *LoweringState, block_id: ids.BlockId, result: ?shader_ir.id.ValueId, operation: shader_ir.instruction.StoreInterface) Error!void { - try requireNoResult(result); - - if (operation.element_index != null) - return Error.UnsupportedOperation; - - const variable = self.lowerer.module.interface_variables.get(operation.variable) orelse return Error.InvalidModule; - - if (variable.direction != .output) - return Error.InvalidModule; - - const source_components = try self.components(operation.value); - const value = self.lowerer.module.values.get(operation.value) orelse return Error.InvalidModule; - - if (value.type != variable.type) - return Error.InvalidModule; - - for (source_components, 0..) |source_component, component_index| { - try self.appendInstruction(block_id, null, .{ - .store_output = .{ - .semantic = try lowerInterfaceSemantic(variable.semantic, @intCast(component_index)), - .source = source_component, - }, - }); - } - } - - fn lowerControlAndTerminators(self: *LoweringState, allocator: std.mem.Allocator) Error!void { - const source_entry = try self.sourceEntryFunction(); - const function = source_entry[1]; - - for (function.blocks.items) |source_block_id| { - const source_block = self.lowerer.module.blocks.get(source_block_id) orelse return Error.InvalidModule; - const target_block_id = try self.mappedBlock(source_block_id); - const structured_control: instruction.StructuredControl = switch (source_block.structured_control) { - .none => .none, - .selection => |selection| .{ .selection = .{ - .merge_block = try self.mappedBlock(selection.merge_block), - } }, - .loop => |loop| .{ .loop = .{ - .merge_block = try self.mappedBlock(loop.merge_block), - .continue_block = try self.mappedBlock(loop.continue_block), - } }, - }; - self.builder.setStructuredControl(target_block_id, structured_control) catch |err| - return mapProgramError(err); - - const source_terminator = source_block.terminator orelse return Error.InvalidModule; - const target_terminator: instruction.Terminator = switch (source_terminator) { - .branch => |edge| .{ .jump = try self.lowerEdge(allocator, edge) }, - .conditional_branch => |branch| conditional: { - switch (try self.predicate(branch.condition)) { - .constant => |condition| { - const edge = if (condition) branch.true_edge else branch.false_edge; - break :conditional .{ .jump = try self.lowerEdge(allocator, edge) }; - }, - .dynamic => |condition| { - const true_edge = try self.lowerEdge(allocator, branch.true_edge); - errdefer allocator.free(true_edge.arguments); - const false_edge = try self.lowerEdge(allocator, branch.false_edge); - break :conditional .{ .conditional_branch = .{ - .predicate = condition, - .true_edge = true_edge, - .false_edge = false_edge, - } }; - }, - } - }, - .return_void => .end_thread, - .return_value => return Error.InvalidEntryPoint, - .discard => return Error.UnsupportedTerminator, - .@"unreachable" => .@"unreachable", - }; - defer freeTerminatorArguments(allocator, target_terminator); - self.builder.setTerminator(target_block_id, target_terminator) catch |err| - return mapProgramError(err); - } - } - - fn lowerEdge(self: *LoweringState, allocator: std.mem.Allocator, edge: shader_ir.module.Edge) Error!instruction.Edge { - const target_source_block = self.lowerer.module.blocks.get(edge.target) orelse return Error.InvalidModule; - if (edge.arguments.len != target_source_block.parameters.items.len) - return Error.InvalidModule; - - var arguments: std.ArrayList(pseudo.EdgeArgument) = .empty; - defer arguments.deinit(allocator); - for (edge.arguments) |argument_id| { - switch (try self.location(argument_id)) { - .components => |bundle| for (bundle) |component| - try arguments.append(allocator, .{ .source = component }), - .predicate => |predicate_value| try arguments.append(allocator, .{ .predicate = predicate_value }), - } - } - - return .{ - .target = try self.mappedBlock(edge.target), - .arguments = try arguments.toOwnedSlice(allocator), - }; - } -}; - -fn freeTerminatorArguments(allocator: std.mem.Allocator, terminator: instruction.Terminator) void { - switch (terminator) { - .jump => |edge| allocator.free(edge.arguments), - .conditional_branch => |branch| { - allocator.free(branch.true_edge.arguments); - allocator.free(branch.false_edge.arguments); - }, - else => {}, - } -} - -pub const Lowerer = struct { - module: *shader_ir.module.Module, - device_info: device.DeviceInfo, - options: Options, - - pub fn init(module: *shader_ir.module.Module, device_info: device.DeviceInfo, options: Options) Lowerer { - return .{ - .module = module, - .device_info = device_info, - .options = options, - }; - } - - pub fn lower(self: *Lowerer, allocator: std.mem.Allocator) Error!program_ir.Program { - // Only supports gen9 for now as it is the only gen I have access to - if (self.device_info.generation != .gen9) - return Error.UnsupportedGeneration; - if (self.module.stage != .vertex) - return Error.UnsupportedStage; - - if (self.options.dispatch_width != .simd8 or !self.device_info.supportsDispatch(self.options.dispatch_width)) - return Error.UnsupportedDispatchWidth; - - shader_ir.validator.validate(self.module) catch |err| return switch (err) { - error.OutOfMemory => Error.OutOfMemory, - error.MissingEntryPoint => Error.MissingEntryPoint, - error.InvalidEntryPoint => Error.InvalidEntryPoint, - else => Error.InvalidModule, - }; - - var transformer_manager = shader_ir.transformer_manager.Manager.init(allocator); - defer transformer_manager.deinit(); - transformer_manager.add(shader_ir.inline_all_functions.transformer) catch return Error.OutOfMemory; - - var transformer_context: shader_ir.transformer_manager.Context = .{ .allocator = allocator }; - _ = transformer_manager.run(self.module, &transformer_context) catch |err| return switch (err) { - error.OutOfMemory => Error.OutOfMemory, - else => Error.SanitizationFailed, - }; - if (!self.module.properties.no_function_calls) - return Error.UnsanitizedModule; - - var program = program_ir.Program.init(allocator, self.module.stage, self.device_info, self.options.dispatch_width); - errdefer program.deinit(); - - const block_map = try allocator.alloc(?ids.BlockId, self.module.blocks.entries.items.len); - defer allocator.free(block_map); - @memset(block_map, null); - - const value_locations = try allocator.alloc(?ValueLocation, self.module.values.entries.items.len); - defer allocator.free(value_locations); - @memset(value_locations, null); - - var state: LoweringState = .{ - .lowerer = self, - .builder = Builder.init(&program), - .storage = program.allocator(), - .block_map = block_map, - .value_locations = value_locations, - }; - - try state.lowerBlocks(); - try state.lowerParameters(); - try state.lowerInstructions(allocator); - try state.lowerControlAndTerminators(allocator); - - program.properties.common_ir_lowered = true; - validator.validate(&program) catch return Error.InvalidLoweredProgram; - - block_arguments.run(allocator, &program) catch |err| return switch (err) { - error.OutOfMemory => Error.OutOfMemory, - else => Error.InvalidLoweredProgram, - }; - validator.validate(&program) catch return Error.InvalidLoweredProgram; - return program; - } -}; - -fn lowerInterfaceSemantic(semantic: shader_ir.module.InterfaceSemantic, component_offset: u8) Error!instruction.InterfaceSemantic { - return switch (semantic) { - .location => |location| location: { - if (location.index != 0) - return Error.UnsupportedOperation; - const component = std.math.add(u8, location.component, component_offset) catch return Error.UnsupportedOperation; - if (component > 3) - return Error.UnsupportedOperation; - break :location .{ - .location = .{ - .location = location.location, - .component = component, - }, - }; - }, - .builtin => |builtin| .{ - .builtin = .{ - .builtin = switch (builtin) { - .position => .position, - .vertex_index => if (component_offset == 0) .vertex_index else return Error.UnsupportedOperation, - .instance_index => if (component_offset == 0) .instance_index else return Error.UnsupportedOperation, - .frag_coord, .frag_depth, .global_invocation_id => return Error.UnsupportedOperation, - }, - .component = component_offset, - }, - }, - }; -} - -fn bitcastImmediate(immediate: operand.Immediate, target_type: operand.DataType) operand.Immediate { - const bits: u32 = switch (immediate) { - .u32 => |value| value, - .i32 => |value| @bitCast(value), - .f32 => |value| @bitCast(value), - }; - return switch (target_type) { - .u32 => .{ .u32 = bits }, - .i32 => .{ .i32 = @bitCast(bits) }, - .f32 => .{ .f32 = @bitCast(bits) }, - else => unreachable, - }; -} - -fn mapProgramError(err: anyerror) Error { - return switch (err) { - Error.OutOfMemory => Error.OutOfMemory, - else => Error.InvalidLoweredProgram, - }; -} - -/// Convenience entry point for callers that do not need to retain a lowerer. -pub inline fn lower(allocator: std.mem.Allocator, module: *shader_ir.module.Module, device_info: device.DeviceInfo, options: Options) Error!program_ir.Program { - var lowerer = Lowerer.init(module, device_info, options); - return lowerer.lower(allocator); -} - -const test_device: device.DeviceInfo = .{ - .generation = .gen9, - .platform = .skylake, - .pci_device_id = 0x1912, - .grf_count = 128, -}; - -fn expectLowered(source: []const u8, expected: []const u8) !void { - var module = try shader_ir.parser.parseString(std.testing.allocator, source); - defer module.deinit(); - - var program = try lower(std.testing.allocator, &module, test_device, .{}); - defer program.deinit(); - try std.testing.expect(program.properties.common_ir_lowered); - try std.testing.expect(!program.properties.instructions_selected); - - const actual = try printer.allocPrint(std.testing.allocator, &program); - defer std.testing.allocator.free(actual); - try std.testing.expectEqualStrings(expected, actual); -} - -fn expectLoweredFragments(source: []const u8, expected: []const []const u8, unexpected: []const []const u8) !void { - var module = try shader_ir.parser.parseString(std.testing.allocator, source); - defer module.deinit(); - - var program = try lower(std.testing.allocator, &module, test_device, .{}); - defer program.deinit(); - try std.testing.expect(program.properties.common_ir_lowered); - try std.testing.expect(!program.properties.instructions_selected); - - const actual = try printer.allocPrint(std.testing.allocator, &program); - defer std.testing.allocator.free(actual); - - for (expected) |fragment| - try std.testing.expect(std.mem.indexOf(u8, actual, fragment) != null); - for (unexpected) |fragment| - try std.testing.expect(std.mem.indexOf(u8, actual, fragment) == null); -} - -fn expectLoweringError(source: []const u8, expected: Error) !void { - var module = try shader_ir.parser.parseString(std.testing.allocator, source); - defer module.deinit(); - - var program = lower(std.testing.allocator, &module, test_device, .{}) catch |actual| { - try std.testing.expectEqual(expected, actual); - return; - }; - defer program.deinit(); - return error.TestExpectedError; -} - -test "[ir] Lower: basic shader" { - const source = - \\shader vertex @main - \\{ - \\ @out_value: u32 = output[location(0), component(0), index(0)] - \\ %one: constant u32 = bits(0x1) - \\ %two: constant u32 = bits(0x2) - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ %sum: u32 = integer_add %one, %two - \\ %condition: bool = cmp_unsigned_less %one, %two - \\ conditional_branch %condition, .left(), .right() - \\ .left(): - \\ branch .merge(%sum) - \\ .right(): - \\ branch .merge(%two) - \\ .merge(%value: u32): - \\ store_interface @out_value, %value - \\ return - \\ } - \\} - ; - - const expected = - \\; Flint program: - \\; .stage: vertex - \\; .generation: gen9 - \\; .platform: skylake - \\; .dispatch_width: simd8 - \\ - \\%value: vgrf u32[8], class(temporary), size(32), alignment(32), spillable - \\%sum: vgrf u32[8], class(temporary), size(32), alignment(32), spillable - \\%condition: vflag - \\ - \\.entry: - \\ [simd8] add %sum:u32, 1:u32, 2:u32 - \\ [simd8] cmp_less_than %condition, 1:u32, 2:u32 - \\ conditional_branch (+%condition), .left, .right - \\ - \\.left: - \\ jump .b4 - \\ - \\.right: - \\ jump .b5 - \\ - \\.merge: - \\ [simd8] store_output location(0), component(0), %value:u32 - \\ end_thread - \\ - \\.b4: - \\ [simd8] parallel_copy [%value:u32 <- %sum:u32] - \\ jump .merge - \\ - \\.b5: - \\ [simd8] parallel_copy [%value:u32 <- 2:u32] - \\ jump .merge - \\ - \\ - ; - - try expectLowered(source, expected); -} - -test "[ir] Lower: control flow" { - const source = - \\shader vertex @main - \\{ - \\ %one: constant u32 = bits(0x1) - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ branch .producer() - \\ .producer(): - \\ %sum: u32 = integer_add %one, %one - \\ branch .merge() - \\ .merge(): - \\ %doubled: u32 = integer_add %sum, %one - \\ return - \\ } - \\} - ; - - const expected = - \\; Flint program: - \\; .stage: vertex - \\; .generation: gen9 - \\; .platform: skylake - \\; .dispatch_width: simd8 - \\ - \\%sum: vgrf u32[8], class(temporary), size(32), alignment(32), spillable - \\%doubled: vgrf u32[8], class(temporary), size(32), alignment(32), spillable - \\ - \\.entry: - \\ jump .producer - \\ - \\.producer: - \\ [simd8] add %sum:u32, 1:u32, 1:u32 - \\ jump .merge - \\ - \\.merge: - \\ [simd8] add %doubled:u32, %sum:u32, 1:u32 - \\ end_thread - \\ - \\ - ; - - try expectLowered(source, expected); -} - -test "[ir] Lower: function call" { - const source = - \\shader vertex @main - \\{ - \\ %one: constant u32 = bits(0x1) - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ %result: u32 = call @identity(%one) - \\ return - \\ } - \\ fn @identity(%value: u32) -> u32 - \\ { - \\ .entry(): - \\ return %value - \\ } - \\} - ; - - const expected = - \\; Flint program: - \\; .stage: vertex - \\; .generation: gen9 - \\; .platform: skylake - \\; .dispatch_width: simd8 - \\ - \\%result: vgrf u32[8], class(temporary), size(32), alignment(32), spillable - \\ - \\.entry: - \\ jump .b2 - \\ - \\.b1: - \\ end_thread - \\ - \\.b2: - \\ jump .b3 - \\ - \\.b3: - \\ [simd8] parallel_copy [%result:u32 <- 1:u32] - \\ jump .b1 - \\ - \\ - ; - - try expectLowered(source, expected); -} - -test "[ir] Lower: unary/binary operations" { - const source = - \\shader vertex @main - \\{ - \\ %u_one: constant u32 = bits(0x1) - \\ %u_two: constant u32 = bits(0x2) - \\ %i_one: constant i32 = bits(0x1) - \\ %i_two: constant i32 = bits(0x2) - \\ %f_one: constant f32 = bits(0x3f800000) - \\ %f_two: constant f32 = bits(0x40000000) - \\ - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ %integer_negated: i32 = negate %i_one - \\ %float_negated: f32 = negate %f_one - \\ %inverted: u32 = bitwise_not %u_one - \\ %integer_difference: i32 = integer_subtract %i_one, %i_two - \\ %float_difference: f32 = float_subtract %f_one, %f_two - \\ %integer_product: u32 = integer_multiply %u_one, %u_two - \\ %float_product: f32 = float_multiply %f_one, %f_two - \\ %shifted_left: u32 = shift_left %u_one, %u_two - \\ %logical_right: u32 = logical_shift_right %u_two, %u_one - \\ %arithmetic_right: i32 = arithmetic_shift_right %i_two, %i_one - \\ %masked: u32 = bitwise_and %u_one, %u_two - \\ %combined: u32 = bitwise_or %u_one, %u_two - \\ %toggled: u32 = bitwise_xor %u_one, %u_two - \\ return - \\ } - \\} - ; - - try expectLoweredFragments(source, &.{ - "[simd8] mov %integer_negated:i32, -1:i32", - "[simd8] mov %float_negated:f32, -1:f32", - "[simd8] bitwise_xor %inverted:u32, 1:u32, 4294967295:u32", - "[simd8] add %integer_difference:i32, 1:i32, -2:i32", - "[simd8] add %float_difference:f32, 1:f32, -2:f32", - "[simd8] multiply %integer_product:u32, 1:u32, 2:u32", - "[simd8] multiply %float_product:f32, 1:f32, 2:f32", - "[simd8] shift_left %shifted_left:u32, 1:u32, 2:u32", - "[simd8] shift_right %logical_right:u32, 2:u32, 1:u32", - "[simd8] shift_right %arithmetic_right:i32, 2:i32, 1:i32", - "[simd8] bitwise_and %masked:u32, 1:u32, 2:u32", - "[simd8] bitwise_or %combined:u32, 1:u32, 2:u32", - "[simd8] bitwise_xor %toggled:u32, 1:u32, 2:u32", - }, &.{}); -} - -test "[ir] Lower: selects and bitcasts" { - const source = - \\shader vertex @main - \\{ - \\ %always: constant bool = true - \\ %one: constant u32 = bits(0x1) - \\ %two: constant u32 = bits(0x2) - \\ %float_one: constant f32 = bits(0x3f800000) - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ %condition: bool = cmp_unsigned_less %one, %two - \\ %dynamic_choice: u32 = select %condition, %one, %two - \\ %inverted_condition: bool = logical_not %condition - \\ %inverted_choice: u32 = select %inverted_condition, %one, %two - \\ %constant_choice: u32 = select %always, %one, %two - \\ %one_bits: u32 = bitcast %float_one - \\ %constant_sum: u32 = integer_add %one_bits, %one - \\ %negative: f32 = negate %float_one - \\ %negative_bits: u32 = bitcast %negative - \\ %register_sum: u32 = integer_add %negative_bits, %one - \\ return - \\ } - \\} - ; - - try expectLoweredFragments(source, &.{ - "[simd8] cmp_less_than %condition, 1:u32, 2:u32", - "[simd8] (-%condition) mov %dynamic_choice:u32, 2:u32", - "[simd8] (+%condition) mov %dynamic_choice:u32, 1:u32", - "[simd8] (+%condition) mov %inverted_choice:u32, 2:u32", - "[simd8] (-%condition) mov %inverted_choice:u32, 1:u32", - "[simd8] mov %constant_choice:u32, 1:u32", - "[simd8] mov %one_bits:u32, 1065353216:u32", - "[simd8] add %constant_sum:u32, %one_bits:u32, 1:u32", - "[simd8] mov %negative:f32, -1:f32", - "[simd8] mov %negative_bits:u32, %negative:u32", - "[simd8] add %register_sum:u32, %negative_bits:u32, 1:u32", - }, &.{}); -} - -test "[ir] Lower: vertex interfaces" { - const source = - \\shader vertex @main - \\{ - \\ @attribute_in: u32 = input[location(2), component(1), index(0)] - \\ @vertex_id_in: u32 = input[builtin(vertex_index)] - \\ @value_out: u32 = output[location(3), component(2), index(0)] - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ %attribute: u32 = load_interface @attribute_in - \\ %vertex_id: u32 = load_interface @vertex_id_in - \\ %value: u32 = integer_add %attribute, %vertex_id - \\ store_interface @value_out, %value - \\ return - \\ } - \\} - ; - - try expectLoweredFragments(source, &.{ - "%attribute: vgrf u32[8], class(varying)", - "%vertex_id: vgrf u32[8], class(varying)", - "[simd8] load_input %attribute:u32, location(2), component(1)", - "[simd8] load_input %vertex_id:u32, builtin(vertex_index), component(0)", - "[simd8] add %value:u32, %attribute:u32, %vertex_id:u32", - "[simd8] store_output location(3), component(2), %value:u32", - }, &.{}); -} - -test "[ir] Lower: vector operations, composites, and interfaces" { - const source = - \\shader vertex @main - \\{ - \\ @attribute_in: vec4[f32] = input[location(0), component(0), index(0)] - \\ @position_out: vec4[f32] = output[builtin(position)] - \\ %one_u32: constant u32 = bits(0x1) - \\ %two_u32: constant u32 = bits(0x2) - \\ %two_f32: constant f32 = bits(0x40000000) - \\ %scale_constant: constant vec4[f32] = [#2, #2, #2, #2] - \\ %zero_constant: constant vec4[f32] = null - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ %attribute: vec4[f32] = load_interface @attribute_in - \\ %scaled: vec4[f32] = float_multiply %attribute, %scale_constant - \\ %with_zero: vec4[f32] = float_add %scaled, %zero_constant - \\ %first: f32 = composite_extract %with_zero[0] - \\ %rebuilt: vec4[f32] = composite_construct %first, %first, %first, %first - \\ %condition: bool = cmp_unsigned_less %one_u32, %two_u32 - \\ %selected: vec4[f32] = select %condition, %with_zero, %rebuilt - \\ %selected_bits: vec4[u32] = bitcast %selected - \\ %restored: vec4[f32] = bitcast %selected_bits - \\ store_interface @position_out, %restored - \\ return - \\ } - \\} - ; - - try expectLoweredFragments(source, &.{ - "%attribute_x: vgrf f32[8], class(varying)", - "%attribute_w: vgrf f32[8], class(varying)", - "[simd8] load_input %attribute_x:f32, location(0), component(0)", - "[simd8] load_input %attribute_w:f32, location(0), component(3)", - "[simd8] multiply %scaled_x:f32, %attribute_x:f32, 2:f32", - "[simd8] multiply %scaled_w:f32, %attribute_w:f32, 2:f32", - "[simd8] add %with_zero_x:f32, %scaled_x:f32, 0:f32", - "[simd8] (+%condition) mov %selected_x:f32, %with_zero_x:f32", - "[simd8] mov %selected_bits_x:u32, %selected_x:u32", - "[simd8] mov %restored_w:f32, %selected_bits_w:f32", - "[simd8] store_output builtin(position), component(0), %restored_x:f32", - "[simd8] store_output builtin(position), component(3), %restored_w:f32", - }, &.{ - "%scale_constant_", - "%zero_constant_", - "%rebuilt_", - }); -} - -test "[ir] Lower: SPIR-V vec4 end-to-end" { - // Assembled from a vertex shader that loads a vec4 input, multiplies it by - // vec4(2.0), and stores the result to Position. - const words = [_]u32{ - 119734787, 65536, 458752, 15, 0, 131089, 1, 196622, - 0, 1, 458767, 0, 1, 1852399981, 0, 2, - 3, 262149, 1, 1852399981, 0, 327685, 2, 1885302377, - 1953067887, 7237481, 393221, 3, 1601467759, 1769172848, 1852795252, 0, - 327685, 4, 1769172848, 1852795252, 0, 262149, 5, 1818321779, - 25701, 262215, 2, 30, 0, 262215, 3, 11, - 0, 131091, 6, 196630, 7, 32, 262167, 8, - 7, 4, 262176, 9, 1, 8, 262176, 10, - 3, 8, 196641, 11, 6, 262187, 7, 12, - 1073741824, 458796, 8, 13, 12, 12, 12, 12, - 262203, 9, 2, 1, 262203, 10, 3, 3, - 327734, 6, 1, 0, 11, 131320, 14, 262205, - 8, 4, 2, 327813, 8, 5, 4, 13, - 196670, 3, 5, 65789, 65592, - }; - - var module = try shader_compiler.spirv.translator.translate(std.testing.allocator, &words, .{ - .entry_point = "main", - .stage = .vertex, - }); - defer module.deinit(); - - var program = try lower(std.testing.allocator, &module, test_device, .{}); - defer program.deinit(); - const text = try printer.allocPrint(std.testing.allocator, &program); - defer std.testing.allocator.free(text); - - for ([_][]const u8{ - "[simd8] load_input %position_x:f32, location(0), component(0)", - "[simd8] load_input %position_w:f32, location(0), component(3)", - "[simd8] multiply %scaled_x:f32, %position_x:f32, 2:f32", - "[simd8] multiply %scaled_w:f32, %position_w:f32, 2:f32", - "[simd8] store_output builtin(position), component(0), %scaled_x:f32", - "[simd8] store_output builtin(position), component(3), %scaled_w:f32", - }) |fragment| - try std.testing.expect(std.mem.indexOf(u8, text, fragment) != null); -} - -test "[ir] Lower: vector block parameter" { - const source = - \\shader vertex @main - \\{ - \\ %one: constant u32 = bits(0x1) - \\ %two: constant u32 = bits(0x2) - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ %pair: vec2[u32] = composite_construct %one, %two - \\ branch .merge(%pair) - \\ .merge(%merged: vec2[u32]): - \\ %first: u32 = composite_extract %merged[0] - \\ return - \\ } - \\} - ; - - try expectLoweredFragments(source, &.{ - "%merged_x: vgrf u32[8]", - "%merged_y: vgrf u32[8]", - "parallel_copy [%merged_x:u32 <- 1:u32, %merged_y:u32 <- 2:u32]", - }, &.{ - ".merge(", - }); -} - -test "[ir] Lower: reject vector interface component overflow" { - try expectLoweringError( - \\shader vertex @main - \\{ - \\ @attribute_in: vec2[f32] = input[location(0), component(3), index(0)] - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ %attribute: vec2[f32] = load_interface @attribute_in - \\ return - \\ } - \\} - , Error.UnsupportedOperation); -} - -test "[ir] Lower: constant conditional branch" { - const source = - \\shader vertex @main - \\{ - \\ %always: constant bool = true - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ conditional_branch %always, .taken(), .untaken() - \\ .taken(): - \\ return - \\ .untaken(): - \\ return - \\ } - \\} - ; - - try expectLoweredFragments(source, &.{ - ".entry:\n jump .taken", - ".taken:\n end_thread", - ".untaken:\n end_thread", - }, &.{ - "conditional_branch", - "vflag", - }); -} - -test "[ir] Lower: boolean block parameter" { - const source = - \\shader vertex @main - \\{ - \\ %one: constant u32 = bits(0x1) - \\ %two: constant u32 = bits(0x2) - \\ %never: constant bool = false - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ %condition: bool = cmp_unsigned_less %one, %two - \\ conditional_branch %condition, .left(), .right() - \\ .left(): - \\ branch .merge(%condition) - \\ .right(): - \\ branch .merge(%never) - \\ .merge(%merged: bool): - \\ conditional_branch %merged, .taken(), .not_taken() - \\ .taken(): - \\ return - \\ .not_taken(): - \\ return - \\ } - \\} - ; - - try expectLoweredFragments(source, &.{ - "%condition: vflag", - "%merged: vflag", - "parallel_copy [%merged <- (+%condition)]", - "parallel_copy [%merged <- false]", - ".merge:\n conditional_branch (+%merged), .taken, .not_taken", - }, &.{}); -} - -test "[ir] Lower: unsupported operations" { - try expectLoweringError( - \\shader vertex @main - \\{ - \\ %one: constant u32 = bits(0x1) - \\ %two: constant u32 = bits(0x2) - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ %quotient: u32 = unsigned_divide %one, %two - \\ return - \\ } - \\} - , Error.UnsupportedOperation); - - try expectLoweringError( - \\shader vertex @main - \\{ - \\ %one: constant u32 = bits(0x1) - \\ %two: constant u32 = bits(0x2) - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ %wide: vec5[u32] = composite_construct %one, %two, %one, %two, %one - \\ return - \\ } - \\} - , Error.UnsupportedType); - - try expectLoweringError( - \\shader vertex @main - \\{ - \\ %one: constant u16 = bits(0x1) - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ %sum: u16 = integer_add %one, %one - \\ return - \\ } - \\} - , Error.UnsupportedType); - - try expectLoweringError( - \\shader vertex @main - \\{ - \\ @storage: u32 = storage_buffer[set(0), binding(0)] - \\ %offset: constant u32 = 0 - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ %value: u32 = load_buffer @storage, %offset - \\ return - \\ } - \\} - , Error.UnsupportedOperation); -} - -test "[ir] Lower: unreachable terminator" { - const source = - \\shader vertex @main - \\{ - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ unreachable - \\ } - \\} - ; - - try expectLoweredFragments(source, &.{ - ".entry:\n unreachable", - }, &.{}); -} - -test "[ir] Lower: unsupported target configuration" { - var module = try shader_ir.parser.parseString(std.testing.allocator, - \\shader vertex @main - \\{ - \\ fn @main() -> void - \\ { - \\ .entry(): - \\ return - \\ } - \\} - ); - defer module.deinit(); - - var gen10 = test_device; - gen10.generation = .gen10; - try std.testing.expectError(Error.UnsupportedGeneration, lower(std.testing.allocator, &module, gen10, .{})); - - module.stage = .fragment; - try std.testing.expectError(Error.UnsupportedStage, lower(std.testing.allocator, &module, test_device, .{})); - module.stage = .vertex; - - try std.testing.expectError(Error.UnsupportedDispatchWidth, lower(std.testing.allocator, &module, test_device, .{ .dispatch_width = .simd16 })); -} +pub const common_ir = @import("common_ir.zig"); diff --git a/src/intel/compiler/lower/vertex_abi.zig b/src/intel/compiler/lower/vertex_abi.zig deleted file mode 100644 index bcdf2b0..0000000 --- a/src/intel/compiler/lower/vertex_abi.zig +++ /dev/null @@ -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); -} diff --git a/src/intel/compiler/targets/gen9/compute/compute.zig b/src/intel/compiler/targets/gen9/compute/compute.zig new file mode 100644 index 0000000..e9c7640 --- /dev/null +++ b/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 })); +} diff --git a/src/intel/compiler/targets/gen9/gen9.zig b/src/intel/compiler/targets/gen9/gen9.zig new file mode 100644 index 0000000..4a19780 --- /dev/null +++ b/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, .{})); +} diff --git a/src/intel/compiler/targets/gen9/validator.zig b/src/intel/compiler/targets/gen9/validator.zig new file mode 100644 index 0000000..17b7ddd --- /dev/null +++ b/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)); +} diff --git a/src/intel/compiler/targets/targets.zig b/src/intel/compiler/targets/targets.zig new file mode 100644 index 0000000..bf35fce --- /dev/null +++ b/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, + }; +} diff --git a/src/software/interpreter/Runtime.zig b/src/software/interpreter/Runtime.zig index 11284b9..e108964 100644 --- a/src/software/interpreter/Runtime.zig +++ b/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 }; } diff --git a/src/software/interpreter/root.zig b/src/software/interpreter/root.zig index 15fd9e4..1162ce5 100644 --- a/src/software/interpreter/root.zig +++ b/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");