From 937b84cbc366145ebf5f4914c3df562baeb97fdd Mon Sep 17 00:00:00 2001 From: Kbz-8 Date: Wed, 5 Aug 2026 17:20:49 +0200 Subject: [PATCH] [IR/Flint] adding shared SPIR-V IR and Gen9 vertex lowering --- build.zig | 19 +- src/compiler/README.md | 35 +- src/compiler/spirv/SourceModule.zig | 54 ++ src/compiler/spirv/root.zig | 14 +- src/compiler/spirv/spirv.zig | 1 + src/compiler/spirv/translator.zig | 274 +++++++++- src/intel/FlintPhysicalDevice.zig | 3 + src/intel/FlintPipeline.zig | 177 ++++++- src/intel/FlintShaderModule.zig | 11 +- src/intel/compiler/compiler.zig | 1 + src/intel/compiler/device.zig | 57 +++ src/intel/compiler/ir/Builder.zig | 22 +- src/intel/compiler/ir/program.zig | 3 +- src/intel/compiler/ir/validator.zig | 62 ++- src/intel/compiler/lower/lower.zig | 649 +++++++++++++++++------- src/intel/compiler/lower/vertex_abi.zig | 485 ++++++++++++++++++ src/phi/PhiShaderModule.zig | 5 +- src/software/SoftShaderModule.zig | 16 +- src/vulkan/ShaderModule.zig | 34 +- 19 files changed, 1672 insertions(+), 250 deletions(-) create mode 100644 src/compiler/spirv/SourceModule.zig create mode 100644 src/intel/compiler/lower/vertex_abi.zig diff --git a/build.zig b/build.zig index 06a3f75..8d628d5 100644 --- a/build.zig +++ b/build.zig @@ -17,6 +17,7 @@ const ImplementationDesc = struct { *std.Build.Module, *std.Build.Module, *std.Build.Module, + *std.Build.Module, std.Build.ResolvedTarget, std.builtin.OptimizeMode, bool, @@ -125,6 +126,7 @@ pub fn build(b: *std.Build) !void { base_mod.addImport("vulkan", vulkan); base_mod.addImport("zmath", zmath); base_mod.addImport("drm", drm); + base_mod.addImport("shader_ir", ir_mod); const base_c_includes = b.addTranslateC(.{ .root_source_file = b.path("src/vulkan/c_includes.h"), @@ -174,7 +176,7 @@ pub fn build(b: *std.Build) !void { for (implementations[0..impl_index], implementation_modules[0..impl_index]) |child_impl, child_mod| lib_mod.addImport(child_impl.name, child_mod); } else if (impl.custom) |func| { - func(b, options, lib, lib_mod, base_mod, vulkan, base_c_mod, target, optimize, use_llvm) catch continue; + func(b, options, lib, lib_mod, base_mod, vulkan, base_c_mod, ir_mod, target, optimize, use_llvm) catch continue; } const icd_file = b.addWriteFile( @@ -433,6 +435,7 @@ fn customSoft( _: *std.Build.Module, _: *std.Build.Module, base_c_mod: *std.Build.Module, + _: *std.Build.Module, target: std.Build.ResolvedTarget, optimize: std.builtin.OptimizeMode, use_llvm: bool, @@ -462,23 +465,20 @@ fn customSoft( // Flint specialized functions fn customFlint( - b: *std.Build, + _: *std.Build, _: *Step.Options, _: *Step.Compile, lib_mod: *std.Build.Module, _: *std.Build.Module, _: *std.Build.Module, base_c_mod: *std.Build.Module, - target: std.Build.ResolvedTarget, - optimize: std.builtin.OptimizeMode, + shader_ir_mod: *std.Build.Module, + _: std.Build.ResolvedTarget, + _: std.builtin.OptimizeMode, _: bool, ) !void { lib_mod.addImport("intel_c", base_c_mod); - lib_mod.addImport("shader_ir", b.createModule(.{ - .root_source_file = b.path("src/compiler/root.zig"), - .target = target, - .optimize = optimize, - })); + lib_mod.addImport("shader_ir", shader_ir_mod); } // Phi specialized functions @@ -491,6 +491,7 @@ fn customPhi( _: *std.Build.Module, _: *std.Build.Module, base_c_mod: *std.Build.Module, + _: *std.Build.Module, target: std.Build.ResolvedTarget, optimize: std.builtin.OptimizeMode, use_llvm: bool, diff --git a/src/compiler/README.md b/src/compiler/README.md index 7382e1c..459a425 100644 --- a/src/compiler/README.md +++ b/src/compiler/README.md @@ -573,14 +573,39 @@ final large-shader implementation. ## SPIR-V frontend -The compiler currently provides a word parser and an initial translator in -`spirv/`. The parser validates the header, word counts, truncation, and literal -strings. The translator selects one entry point and lowers a defined subset: +The compiler currently provides a word parser, an owned `SourceModule`, and an +initial translator in `spirv/`. The parser validates the header, word counts, +truncation, and literal strings. `SourceModule` copies and retains validated +SPIR-V so API objects can instantiate multiple entry points without borrowing +application memory. + +Use `translator.instantiate` when retaining a source module: + +```zig +var source = try ir.spirv.SourceModule.init(allocator, words); +defer source.deinit(allocator); + +var module = try ir.spirv.translator.instantiate(allocator, &source, .{ + .entry_point = "main", + .stage = .compute, + .specializations = &.{.{ + .constant_id = 7, // SPIR-V SpecId + .data = std.mem.asBytes(&workgroup_width), + }}, +}); +defer module.deinit(); +``` + +`translator.translate` remains a convenience wrapper for borrowed words. Each +translation selects one entry point and returns an independent mutable IR +module. The translator lowers a defined subset: - Vertex, fragment, and compute stages. - Basic scalar, vector, array, structure, pointer, and function types. -- Ordinary and composite constants; unapplied specialization constants are - refused. +- Ordinary constants plus scalar boolean, integer, and floating-point + specialization constants selected through `SpecId`. Missing overrides use the + SPIR-V defaults, and specialization composites are rebuilt from their + specialized elements. `OpSpecConstantOp` is not evaluated yet. - Functions, blocks, branches, structured merge marks, and returns. - `OpPhi` into block parameters and edge arguments. - The arithmetic, comparison, select, bitcast, and composite operations named diff --git a/src/compiler/spirv/SourceModule.zig b/src/compiler/spirv/SourceModule.zig new file mode 100644 index 0000000..75f03d3 --- /dev/null +++ b/src/compiler/spirv/SourceModule.zig @@ -0,0 +1,54 @@ +const std = @import("std"); +const Parser = @import("Parser.zig"); + +const Self = @This(); + +words: []u32, +parsed: Parser, + +pub const Error = std.mem.Allocator.Error || Parser.Error; + +pub fn init(allocator: std.mem.Allocator, words: []const u32) Error!Self { + const owned_words = try allocator.dupe(u32, words); + errdefer allocator.free(owned_words); + + return .{ + .words = owned_words, + .parsed = try Parser.init(owned_words), + }; +} + +pub fn deinit(self: *Self, allocator: std.mem.Allocator) void { + allocator.free(self.words); + self.* = undefined; +} + +pub fn code(self: *const Self) []const u32 { + return self.words; +} + +pub fn parser(self: *const Self) Parser { + return self.parsed; +} + +test "SPIR-V: source module owns and validates its words" { + var words = [_]u32{ + 0x07230203, + 0x00010000, + 0, + 1, + 0, + }; + + var source = try Self.init(std.testing.allocator, &words); + defer source.deinit(std.testing.allocator); + + words[0] = 0; + try std.testing.expectEqual(@as(u32, 0x07230203), source.code()[0]); + try std.testing.expectEqual(@as(u8, 1), source.parser().header.major()); +} + +test "SPIR-V: source module rejects malformed input" { + const malformed = [_]u32{ 0, 0, 0, 0, 0 }; + try std.testing.expectError(error.InvalidMagic, Self.init(std.testing.allocator, &malformed)); +} diff --git a/src/compiler/spirv/root.zig b/src/compiler/spirv/root.zig index 49ae358..aff0d92 100644 --- a/src/compiler/spirv/root.zig +++ b/src/compiler/spirv/root.zig @@ -3,19 +3,21 @@ //! This namespace contains the SPIR-V parser and translator used to import shader //! modules into the compiler IR. //! -//! `Parser` validates the SPIR-V header and iterates over binary instructions. -//! `spec` exposes a minimalistic SPIR-V header translation. +//! `Parser` validates borrowed SPIR-V words, while `SourceModule` owns and +//! structurally validates words that need to outlive an API call. `spec` exposes a +//! minimalistic SPIR-V header translation. //! -//! The main entry point is `translator.translate`, which finds the requested entry -//! point, maps its execution model to an IR shader stage, lowers supported types, -//! constants, interfaces, instructions, and structured control flow, then validates -//! the generated IR module. +//! Use `translator.instantiate` to lower one entry point from a retained source +//! module. `translator.translate` remains as a convenience wrapper for borrowed +//! words. pub const Parser = @import("Parser.zig"); +pub const SourceModule = @import("SourceModule.zig"); pub const translator = @import("translator.zig"); pub const spec = @import("spirv.zig"); test { _ = Parser; + _ = SourceModule; _ = translator; } diff --git a/src/compiler/spirv/spirv.zig b/src/compiler/spirv/spirv.zig index 455fb29..606c836 100644 --- a/src/compiler/spirv/spirv.zig +++ b/src/compiler/spirv/spirv.zig @@ -161,6 +161,7 @@ pub const ExecutionMode = enum(u32) { }; pub const Decoration = enum(u32) { + spec_id = 1, built_in = 11, location = 30, component = 31, diff --git a/src/compiler/spirv/translator.zig b/src/compiler/spirv/translator.zig index 77441b9..b7ffc1e 100644 --- a/src/compiler/spirv/translator.zig +++ b/src/compiler/spirv/translator.zig @@ -1,10 +1,19 @@ const std = @import("std"); +const builtin_info = @import("builtin"); const Parser = @import("Parser.zig"); +const SourceModule = @import("SourceModule.zig"); const spirv = @import("spirv.zig"); const ir = @import("../ir/ir.zig"); +pub const SpecializationValue = struct { + constant_id: u32, + data: []const u8, +}; + pub const Options = struct { entry_point: []const u8, + stage: ?ir.module.Stage = null, + specializations: []const SpecializationValue = &.{}, }; pub const TranslationError = error{ @@ -24,6 +33,8 @@ pub const TranslationError = error{ UnsupportedType, UnsupportedConstant, SpecializationConstantsNotApplied, + InvalidSpecialization, + DuplicateSpecializationConstant, UnsupportedOpcode, }; @@ -34,6 +45,7 @@ const EntryPoint = struct { }; const Decorations = struct { + spec_id: ?u32 = null, location: ?u32 = null, component: u8 = 0, index: u8 = 0, @@ -57,6 +69,7 @@ const Context = struct { variable_defs: []?Parser.Instruction, names: []?[]const u8, decorations: []Decorations, + specializations: []const SpecializationValue, types: []?ir.id.TypeId, values: []?ir.id.ValueId, @@ -82,6 +95,17 @@ const Context = struct { return self.names[index]; } + fn specializationData(self: *const Context, result_id: u32) TranslationError!?[]const u8 { + const index = try self.idIndex(result_id); + const spec_id = self.decorations[index].spec_id orelse return null; + + for (self.specializations) |specialization| { + if (specialization.constant_id == spec_id) + return specialization.data; + } + return null; + } + fn translateType(self: *Context, spv_id: u32) anyerror!ir.id.TypeId { const index = try self.idIndex(spv_id); if (self.types[index]) |translated| @@ -232,7 +256,7 @@ const Context = struct { try expectOperandCount(operands, 2); break :blk try self.builder.internConstant(try self.translateType(operands[0]), .null); }, - .constant_composite => blk: { + .constant_composite, .spec_constant_composite => blk: { if (operands.len < 2) return error.InvalidInstruction; @@ -252,12 +276,45 @@ const Context = struct { .{ .composite = elements }, ); }, - .spec_constant_true, - .spec_constant_false, - .spec_constant, - .spec_constant_composite, - .spec_constant_op, - => return error.SpecializationConstantsNotApplied, + .spec_constant_true, .spec_constant_false => blk: { + try expectOperandCount(operands, 2); + const ty = try self.translateType(operands[0]); + const type_data = self.module.types.get(ty) orelse return error.InvalidId; + if (type_data.* != .boolean) + return error.UnsupportedConstant; + + const value = if (try self.specializationData(spv_id)) |data| + try specializationBoolean(data) + else + instruction.opcode == .spec_constant_true; + break :blk try self.builder.internConstant(ty, .{ .boolean = value }); + }, + .spec_constant => blk: { + if (operands.len < 3 or operands.len > 4) + return error.InvalidInstruction; + + const ty = try self.translateType(operands[0]); + const type_data = self.module.types.get(ty) orelse return error.InvalidId; + const default_bits = try literalBits(operands[2..]); + const override = try self.specializationData(spv_id); + + break :blk switch (type_data.*) { + .integer => |integer| try self.builder.internConstant(ty, .{ + .integer_bits = if (override) |data| + try specializationBits(data, integer.bits) + else + default_bits, + }), + .floating => |floating| try self.builder.internConstant(ty, .{ + .float_bits = if (override) |data| + try specializationBits(data, floating.bits) + else + default_bits, + }), + else => return error.UnsupportedConstant, + }; + }, + .spec_constant_op => return error.SpecializationConstantsNotApplied, else => return error.MissingDefinition, }; @@ -290,9 +347,12 @@ const Context = struct { } }; -pub fn translate(allocator: std.mem.Allocator, words: []const u32, options: Options) !ir.module.Module { - const parser = try Parser.init(words); - const entry_point = try findEntryPoint(parser, options.entry_point); +/// Translates one entry point from a retained SPIR-V source into an independent +/// common IR module. The returned module does not borrow from `source`. +pub fn instantiate(allocator: std.mem.Allocator, source: *const SourceModule, options: Options) !ir.module.Module { + try validateSpecializations(options.specializations); + const parser = source.parser(); + const entry_point = try findEntryPoint(parser, options.entry_point, options.stage); const stage = try translateStage(entry_point.model); var module = ir.module.Module.init(allocator, stage); @@ -314,6 +374,7 @@ pub fn translate(allocator: std.mem.Allocator, words: []const u32, options: Opti .variable_defs = try allocOptional(Parser.Instruction, scratch, bound), .names = try allocOptional([]const u8, scratch, bound), .decorations = try scratch.alloc(Decorations, bound), + .specializations = options.specializations, .types = try allocOptional(ir.id.TypeId, scratch, bound), .values = try allocOptional(ir.id.ValueId, scratch, bound), .blocks = try allocOptional(ir.id.BlockId, scratch, bound), @@ -334,6 +395,13 @@ pub fn translate(allocator: std.mem.Allocator, words: []const u32, options: Opti return module; } +/// Convenience wrapper for callers that do not retain a source module. +pub fn translate(allocator: std.mem.Allocator, words: []const u32, options: Options) !ir.module.Module { + var source = try SourceModule.init(allocator, words); + defer source.deinit(allocator); + return instantiate(allocator, &source, options); +} + fn collectDeclarations(context: *Context) !void { var iterator = context.parser.iterator(); while (try iterator.next()) |instruction| { @@ -378,6 +446,12 @@ fn collectDecoration(context: *Context, operands: []const u32) !void { const index = try context.idIndex(operands[0]); const decoration: spirv.Decoration = @enumFromInt(operands[1]); switch (decoration) { + .spec_id => { + try expectOperandCount(operands, 3); + if (context.decorations[index].spec_id != null) + return error.InvalidInstruction; + context.decorations[index].spec_id = operands[2]; + }, .built_in => { try expectOperandCount(operands, 3); context.decorations[index].builtin = operands[2]; @@ -455,7 +529,7 @@ fn translateInterfaces(context: *Context, interface_ids: []const u32) !void { } } -fn findEntryPoint(parser: Parser, requested_name: []const u8) !EntryPoint { +fn findEntryPoint(parser: Parser, requested_name: []const u8, requested_stage: ?ir.module.Stage) !EntryPoint { var found: ?EntryPoint = null; var iterator = parser.iterator(); while (try iterator.next()) |instruction| { @@ -473,11 +547,21 @@ fn findEntryPoint(parser: Parser, requested_name: []const u8) !EntryPoint { if (!try Parser.literalStringEquals(instruction.operands[2 .. 2 + string_words], requested_name)) continue; + const model: spirv.ExecutionModel = @enumFromInt(instruction.operands[0]); + if (requested_stage) |stage| { + const candidate_stage = translateStage(model) catch |err| switch (err) { + error.UnsupportedExecutionModel => continue, + else => return err, + }; + if (candidate_stage != stage) + continue; + } + if (found != null) return error.AmbiguousEntryPoint; found = .{ - .model = @enumFromInt(instruction.operands[0]), + .model = model, .function_id = instruction.operands[1], .interface_ids = instruction.operands[2 + string_words ..], }; @@ -1012,6 +1096,35 @@ fn translateCompareOpcode(opcode: spirv.Opcode) ir.instruction.CompareOpcode { }; } +fn validateSpecializations(specializations: []const SpecializationValue) TranslationError!void { + for (specializations, 0..) |specialization, index| { + for (specializations[0..index]) |previous| { + if (previous.constant_id == specialization.constant_id) + return error.DuplicateSpecializationConstant; + } + } +} + +fn specializationBoolean(data: []const u8) TranslationError!bool { + if (data.len != @sizeOf(u32)) + return error.InvalidSpecialization; + return std.mem.readInt(u32, data[0..4], builtin_info.target.cpu.arch.endian()) != 0; +} + +fn specializationBits(data: []const u8, bit_width: u16) TranslationError!u64 { + const expected_size: usize = (@as(usize, bit_width) + 7) / 8; + if (data.len != expected_size) + return error.InvalidSpecialization; + + return switch (expected_size) { + 1 => data[0], + 2 => std.mem.readInt(u16, data[0..2], builtin_info.target.cpu.arch.endian()), + 4 => std.mem.readInt(u32, data[0..4], builtin_info.target.cpu.arch.endian()), + 8 => std.mem.readInt(u64, data[0..8], builtin_info.target.cpu.arch.endian()), + else => error.InvalidSpecialization, + }; +} + fn literalBits(words: []const u32) TranslationError!u64 { return switch (words.len) { 1 => words[0], @@ -1230,6 +1343,120 @@ test "SPIR-V: fragment execution modes and translated properties" { try std.testing.expect(entry.terminator.? == .return_void); } +test "SPIR-V: retained source instantiates independent entry points" { + const assembly = + \\OpCapability Shader + \\OpMemoryModel Logical GLSL450 + \\OpEntryPoint Vertex %vertex_main "main" + \\OpEntryPoint GLCompute %compute_main "main" + \\OpExecutionMode %compute_main LocalSize 2 1 1 + \\%void = OpTypeVoid + \\%fn_void = OpTypeFunction %void + \\%vertex_main = OpFunction %void None %fn_void + \\ %vertex_entry = OpLabel + \\ OpReturn + \\OpFunctionEnd + \\%compute_main = OpFunction %void None %fn_void + \\ %compute_entry = OpLabel + \\ OpReturn + \\OpFunctionEnd + ; + const words = try assembleSpirv(std.testing.allocator, assembly); + defer std.testing.allocator.free(words); + + var source = try SourceModule.init(std.testing.allocator, words); + defer source.deinit(std.testing.allocator); + + var vertex_module = try instantiate(std.testing.allocator, &source, .{ + .entry_point = "main", + .stage = .vertex, + }); + defer vertex_module.deinit(); + var compute_module = try instantiate(std.testing.allocator, &source, .{ + .entry_point = "main", + .stage = .compute, + }); + defer compute_module.deinit(); + + try std.testing.expectEqual(ir.module.Stage.vertex, vertex_module.stage); + try std.testing.expectEqual(ir.module.Stage.compute, compute_module.stage); + try std.testing.expectEqual(@as(?[3]u32, .{ 2, 1, 1 }), compute_module.execution_modes.workgroup_size); + try std.testing.expect(vertex_module.entry_point != null); + try std.testing.expect(compute_module.entry_point != null); +} + +test "SPIR-V: scalar specialization constants and defaults" { + const assembly = + \\OpCapability Shader + \\OpMemoryModel Logical GLSL450 + \\OpEntryPoint GLCompute %main "main" + \\OpExecutionMode %main LocalSize 1 1 1 + \\OpName %number "number" + \\OpName %enabled "enabled" + \\OpName %pair "pair" + \\OpDecorate %number SpecId 7 + \\OpDecorate %enabled SpecId 8 + \\%void = OpTypeVoid + \\%bool = OpTypeBool + \\%u32 = OpTypeInt 32 0 + \\%vec2_u32 = OpTypeVector %u32 2 + \\%fn_void = OpTypeFunction %void + \\%number = OpSpecConstant %u32 3 + \\%enabled = OpSpecConstantFalse %bool + \\%pair = OpSpecConstantComposite %vec2_u32 %number %number + \\%main = OpFunction %void None %fn_void + \\ %entry = OpLabel + \\ %sum = OpIAdd %u32 %number %number + \\ %selected = OpSelect %u32 %enabled %sum %number + \\ %first = OpCompositeExtract %u32 %pair 0 + \\ OpReturn + \\OpFunctionEnd + ; + const words = try assembleSpirv(std.testing.allocator, assembly); + defer std.testing.allocator.free(words); + + var source = try SourceModule.init(std.testing.allocator, words); + defer source.deinit(std.testing.allocator); + + var defaults = try instantiate(std.testing.allocator, &source, .{ + .entry_point = "main", + .stage = .compute, + }); + defer defaults.deinit(); + try expectNamedIntegerConstant(&defaults, "number", 3); + try expectNamedBooleanConstant(&defaults, "enabled", false); + + const number_override: u32 = 42; + const enabled_override: u32 = 1; + const specializations = [_]SpecializationValue{ + .{ .constant_id = 7, .data = std.mem.asBytes(&number_override) }, + .{ .constant_id = 8, .data = std.mem.asBytes(&enabled_override) }, + }; + var specialized = try instantiate(std.testing.allocator, &source, .{ + .entry_point = "main", + .stage = .compute, + .specializations = &specializations, + }); + defer specialized.deinit(); + try expectNamedIntegerConstant(&specialized, "number", 42); + try expectNamedBooleanConstant(&specialized, "enabled", true); + + const invalid_size: u16 = 9; + try std.testing.expectError(error.InvalidSpecialization, instantiate(std.testing.allocator, &source, .{ + .entry_point = "main", + .stage = .compute, + .specializations = &.{.{ .constant_id = 7, .data = std.mem.asBytes(&invalid_size) }}, + })); + try std.testing.expectError(error.DuplicateSpecializationConstant, instantiate(std.testing.allocator, &source, .{ + .entry_point = "main", + .stage = .compute, + .specializations = &.{ + .{ .constant_id = 7, .data = std.mem.asBytes(&number_override) }, + .{ .constant_id = 7, .data = std.mem.asBytes(&number_override) }, + }, + })); +} + test "SPIR-V: entry point lookup errors" { const single_entry_assembly = \\OpCapability Shader @@ -1475,6 +1702,29 @@ test "SPIR-V: preserves location components and builtin interfaces" { try std.testing.expectEqual(ir.module.Builtin.position, position.semantic.builtin); } +fn expectNamedIntegerConstant(module: *const ir.module.Module, name: []const u8, expected: u64) !void { + const value = findNamedConstant(module, name) orelse return error.MissingNamedConstant; + try std.testing.expect(value == .integer_bits); + try std.testing.expectEqual(expected, value.integer_bits); +} + +fn expectNamedBooleanConstant(module: *const ir.module.Module, name: []const u8, expected: bool) !void { + const value = findNamedConstant(module, name) orelse return error.MissingNamedConstant; + try std.testing.expect(value == .boolean); + try std.testing.expectEqual(expected, value.boolean); +} + +fn findNamedConstant(module: *const ir.module.Module, name: []const u8) ?ir.constant.ConstantValue { + for (module.values.entries.items) |entry| { + const value = entry orelse continue; + const value_name = value.name orelse continue; + if (!std.mem.eql(u8, value_name, name) or value.definition != .constant) + continue; + return module.constants.get(value.definition.constant).?.value; + } + return null; +} + fn assembleSpirv(allocator: std.mem.Allocator, assembly: []const u8) ![]u32 { var io_backend: std.Io.Threaded = .init(allocator, .{}); defer io_backend.deinit(); diff --git a/src/intel/FlintPhysicalDevice.zig b/src/intel/FlintPhysicalDevice.zig index 3866663..b10a052 100644 --- a/src/intel/FlintPhysicalDevice.zig +++ b/src/intel/FlintPhysicalDevice.zig @@ -7,6 +7,7 @@ const lib = @import("lib.zig"); const pci_ids = @import("pci_ids.zig").map; const FlintDevice = @import("FlintDevice.zig"); +const compiler_device = @import("compiler/device.zig"); const VkError = base.VkError; const SurfaceKHR = base.SurfaceKHR; @@ -29,6 +30,7 @@ pub const extensions = [_]vk.ExtensionProperties{ interface: Interface, kmd_type: lib.KmdType, +compiler_info: ?compiler_device.DeviceInfo, node_path: [base.drm.max_node_name:0]u8, pub fn create(allocator: std.mem.Allocator, instance: *base.Instance, drm_device: *const base.drm.Device, kmd_type: lib.KmdType) VkError!*Self { @@ -224,6 +226,7 @@ pub fn create(allocator: std.mem.Allocator, instance: *base.Instance, drm_device self.* = .{ .interface = interface, .kmd_type = kmd_type, + .compiler_info = compiler_device.DeviceInfo.fromPciDeviceId(interface.props.device_id), .node_path = @splat(0), }; const node_path = drm_device.nodePath(); diff --git a/src/intel/FlintPipeline.zig b/src/intel/FlintPipeline.zig index 80a819c..bee7071 100644 --- a/src/intel/FlintPipeline.zig +++ b/src/intel/FlintPipeline.zig @@ -1,47 +1,210 @@ const std = @import("std"); const vk = @import("vulkan"); const base = @import("base"); +const shader_ir = @import("shader_ir"); +const compiler = @import("compiler/compiler.zig"); +const FlintPhysicalDevice = @import("FlintPhysicalDevice.zig"); const VkError = base.VkError; const Self = @This(); pub const Interface = base.Pipeline; +const PipelineKind = enum { + graphics, + compute, +}; + +const CommonStage = struct { + stage: shader_ir.ir.module.Stage, + module: base.ShaderModule.IrModule, + program: ?compiler.Program, + + fn deinit(self: *CommonStage) void { + if (self.program) |*program| + program.deinit(); + self.module.deinit(); + self.* = undefined; + } +}; + interface: Interface, -host_allocator: base.VulkanAllocator, +artifact_allocator: base.VulkanAllocator, +stages: []CommonStage, pub fn createCompute(device: *base.Device, allocator: std.mem.Allocator, cache: ?*base.PipelineCache, info: *const vk.ComputePipelineCreateInfo) VkError!*Self { const self = allocator.create(Self) catch return VkError.OutOfHostMemory; - errdefer allocator.destroy(self); + var initialized = false; + errdefer if (initialized) self.interface.destroy(allocator) else allocator.destroy(self); var interface = try Interface.initCompute(device, allocator, cache, info); interface.vtable = &.{ .destroy = destroy }; + self.* = .{ .interface = interface, - .host_allocator = base.VulkanAllocator.from(allocator).clone(), + .artifact_allocator = base.VulkanAllocator.from(allocator).clone(), + .stages = &.{}, }; - errdefer self.interface.layout.unref(allocator); + initialized = true; + self.stages = try compileStages(self.artifact_allocator.allocator(), &.{info.stage}, .compute, compilerDeviceInfo(device)); return self; } pub fn createGraphics(device: *base.Device, allocator: std.mem.Allocator, cache: ?*base.PipelineCache, info: *const vk.GraphicsPipelineCreateInfo) VkError!*Self { const self = allocator.create(Self) catch return VkError.OutOfHostMemory; - errdefer allocator.destroy(self); + var initialized = false; + errdefer if (initialized) self.interface.destroy(allocator) else allocator.destroy(self); var interface = try Interface.initGraphics(device, allocator, cache, info); interface.vtable = &.{ .destroy = destroy }; self.* = .{ .interface = interface, - .host_allocator = base.VulkanAllocator.from(allocator).clone(), + .artifact_allocator = base.VulkanAllocator.from(allocator).clone(), + .stages = &.{}, }; - errdefer self.interface.layout.unref(allocator); + initialized = true; + const stage_infos = if (info.p_stages) |stages| + stages[0..info.stage_count] + else + return VkError.ValidationFailed; + self.stages = try compileStages(self.artifact_allocator.allocator(), stage_infos, .graphics, compilerDeviceInfo(device)); return self; } +fn compileStages(allocator: std.mem.Allocator, infos: []const vk.PipelineShaderStageCreateInfo, pipeline_kind: PipelineKind, device_info: ?compiler.device.DeviceInfo) VkError![]CommonStage { + if (infos.len == 0) + return VkError.ValidationFailed; + + const stages = allocator.alloc(CommonStage, infos.len) catch return VkError.OutOfHostMemory; + var initialized: usize = 0; + errdefer { + for (stages[0..initialized]) |*stage| + stage.deinit(); + allocator.free(stages); + } + + for (infos, stages) |*info, *stage| { + stage.* = try compileStage(allocator, info, pipeline_kind, device_info); + initialized += 1; + } + return stages; +} + +fn compileStage(allocator: std.mem.Allocator, info: *const vk.PipelineShaderStageCreateInfo, pipeline_kind: PipelineKind, device_info: ?compiler.device.DeviceInfo) VkError!CommonStage { + const specializations = try specializationValues(allocator, info.p_specialization_info); + defer if (specializations.len != 0) allocator.free(specializations); + + const expected_stage = commonStage(info.stage) orelse return VkError.ValidationFailed; + switch (pipeline_kind) { + .compute => if (expected_stage != .compute) return VkError.ValidationFailed, + .graphics => if (expected_stage == .compute) return VkError.ValidationFailed, + } + + const shader_module = base.NonDispatchable(base.ShaderModule).fromHandleObject(info.module) catch |err| return err; + var module = shader_module.instantiateIr(allocator, .{ + .entry_point = std.mem.span(info.p_name), + .stage = expected_stage, + .specializations = specializations, + }) catch |err| { + std.log.scoped(.FlintPipeline).err("common shader translation failed: {s}", .{@errorName(err)}); + return switch (err) { + error.OutOfMemory => VkError.OutOfHostMemory, + else => VkError.ValidationFailed, + }; + }; + errdefer module.deinit(); + + std.debug.assert(module.stage == expected_stage); + + var program = try lowerToFlint(allocator, &module, device_info); + errdefer if (program) |*value| value.deinit(); + + return .{ + .stage = expected_stage, + .module = module, + .program = 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, + error.UnsupportedGeneration, + error.UnsupportedStage, + error.UnsupportedDispatchWidth, + error.UnsupportedType, + error.UnsupportedOperation, + error.UnsupportedTerminator, + => null, + else => { + std.log.scoped(.FlintPipeline).err("Flint shader lowering failed: {s}", .{@errorName(err)}); + return VkError.ValidationFailed; + }, + }; +} + +fn compilerDeviceInfo(device: *const base.Device) ?compiler.device.DeviceInfo { + const physical_device: *const FlintPhysicalDevice = @alignCast(@fieldParentPtr("interface", device.physical_device)); + return physical_device.compiler_info; +} + +fn specializationValues(allocator: std.mem.Allocator, info: ?*const vk.SpecializationInfo) VkError![]shader_ir.spirv.translator.SpecializationValue { + const specialization = info orelse return &.{}; + if (specialization.map_entry_count == 0) + return &.{}; + + const entries = specialization.p_map_entries orelse return VkError.ValidationFailed; + const data: []const u8 = if (specialization.data_size == 0) + &.{} + else + @as([*]const u8, @ptrCast(@alignCast(specialization.p_data)))[0..specialization.data_size]; + + const values = allocator.alloc(shader_ir.spirv.translator.SpecializationValue, specialization.map_entry_count) catch + return VkError.OutOfHostMemory; + errdefer allocator.free(values); + + for (entries[0..specialization.map_entry_count], values) |entry, *value| { + const offset: usize = entry.offset; + const end = std.math.add(usize, offset, entry.size) catch return VkError.ValidationFailed; + if (end > data.len) + return VkError.ValidationFailed; + value.* = .{ + .constant_id = entry.constant_id, + .data = data[offset..end], + }; + } + return values; +} + +fn commonStage(stage: vk.ShaderStageFlags) ?shader_ir.ir.module.Stage { + const bits: u32 = @bitCast(stage); + const vertex_bits: u32 = @bitCast(vk.ShaderStageFlags{ .vertex_bit = true }); + const fragment_bits: u32 = @bitCast(vk.ShaderStageFlags{ .fragment_bit = true }); + const compute_bits: u32 = @bitCast(vk.ShaderStageFlags{ .compute_bit = true }); + + return if (bits == vertex_bits) + .vertex + else if (bits == fragment_bits) + .fragment + else if (bits == compute_bits) + .compute + else + null; +} + +fn deinitStages(allocator: std.mem.Allocator, stages: []CommonStage) void { + for (stages) |*stage| + stage.deinit(); + if (stages.len != 0) + allocator.free(stages); +} + pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void { const self: *Self = @alignCast(@fieldParentPtr("interface", interface)); + deinitStages(self.artifact_allocator.allocator(), self.stages); allocator.destroy(self); } diff --git a/src/intel/FlintShaderModule.zig b/src/intel/FlintShaderModule.zig index 970e099..199d166 100644 --- a/src/intel/FlintShaderModule.zig +++ b/src/intel/FlintShaderModule.zig @@ -8,7 +8,6 @@ const Self = @This(); pub const Interface = base.ShaderModule; interface: Interface, -code: []u32, ref_count: std.atomic.Value(usize), pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const vk.ShaderModuleCreateInfo) VkError!*Self { @@ -16,14 +15,11 @@ pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const v errdefer allocator.destroy(self); var interface = try Interface.init(device, allocator, info); + errdefer interface.deinit(); interface.vtable = &.{ .destroy = destroy }; - if (info.code_size % @sizeOf(u32) != 0) return VkError.ValidationFailed; - const code = allocator.dupe(u32, info.p_code[0 .. info.code_size / @sizeOf(u32)]) catch return VkError.OutOfHostMemory; - errdefer allocator.free(code); self.* = .{ .interface = interface, - .code = code, .ref_count = std.atomic.Value(usize).init(1), }; return self; @@ -35,7 +31,7 @@ pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void { } pub fn drop(self: *Self, allocator: std.mem.Allocator) void { - allocator.free(self.code); + self.interface.deinit(); allocator.destroy(self); } @@ -44,7 +40,6 @@ pub fn ref(self: *Self) void { } pub fn unref(self: *Self, allocator: std.mem.Allocator) void { - if (self.ref_count.fetchSub(1, .release) == 1) { + if (self.ref_count.fetchSub(1, .acq_rel) == 1) self.drop(allocator); - } } diff --git a/src/intel/compiler/compiler.zig b/src/intel/compiler/compiler.zig index b05f481..0c34a25 100644 --- a/src/intel/compiler/compiler.zig +++ b/src/intel/compiler/compiler.zig @@ -183,4 +183,5 @@ test "[ir] ID stability after removal" { test { _ = lower; + _ = lower.vertex_abi; } diff --git a/src/intel/compiler/device.zig b/src/intel/compiler/device.zig index 2129ea0..1cbe902 100644 --- a/src/intel/compiler/device.zig +++ b/src/intel/compiler/device.zig @@ -38,6 +38,50 @@ pub const DeviceInfo = struct { .simd32 => self.supports_simd32, }; } + + pub fn fromPciDeviceId(raw_pci_device_id: u32) ?DeviceInfo { + if (raw_pci_device_id > 0xffff) + return null; + const pci_device_id: u16 = @intCast(raw_pci_device_id); + + const platform: Platform = switch (pci_device_id & 0xff00) { + 0x1900 => .skylake, + 0x5900 => .kabylake, + 0x3e00 => switch (pci_device_id) { + 0x3ea0, 0x3ea1, 0x3ea2, 0x3ea3, 0x3ea4 => .whiskey_lake, + else => .coffee_lake, + }, + 0x9b00 => .comet_lake, + 0x8a00 => .ice_lake, + 0x4500 => .elkhart_lake, + 0x4e00 => .jasper_lake, + else => switch (pci_device_id) { + 0x0a84, 0x1a84, 0x1a85, 0x5a84, 0x5a85 => .broxton, + 0x3184, 0x3185 => .gemini_lake, + 0x87c0, 0x87ca => .kabylake, + else => return null, + }, + }; + + const generation: Generation = switch (platform) { + .skylake, + .broxton, + .kabylake, + .gemini_lake, + .coffee_lake, + .whiskey_lake, + .comet_lake, + => .gen9, + .ice_lake, .elkhart_lake, .jasper_lake => .gen11, + }; + + return .{ + .generation = generation, + .platform = platform, + .pci_device_id = pci_device_id, + .grf_count = 128, + }; + } }; pub const DispatchWidth = enum(u8) { @@ -54,3 +98,16 @@ pub const ExecutionSize = enum(u8) { simd16 = 16, simd32 = 32, }; + +test "compiler device: classify supported Intel PCI IDs" { + const std = @import("std"); + + try std.testing.expectEqual(Platform.skylake, DeviceInfo.fromPciDeviceId(0x1912).?.platform); + try std.testing.expectEqual(Platform.broxton, DeviceInfo.fromPciDeviceId(0x5a84).?.platform); + try std.testing.expectEqual(Platform.kabylake, DeviceInfo.fromPciDeviceId(0x5916).?.platform); + try std.testing.expectEqual(Platform.whiskey_lake, DeviceInfo.fromPciDeviceId(0x3ea0).?.platform); + try std.testing.expectEqual(Platform.comet_lake, DeviceInfo.fromPciDeviceId(0x9bc5).?.platform); + try std.testing.expectEqual(Generation.gen11, DeviceInfo.fromPciDeviceId(0x8a52).?.generation); + try std.testing.expectEqual(Generation.gen11, DeviceInfo.fromPciDeviceId(0x4e55).?.generation); + try std.testing.expectEqual(@as(?DeviceInfo, null), DeviceInfo.fromPciDeviceId(0x46a6)); +} diff --git a/src/intel/compiler/ir/Builder.zig b/src/intel/compiler/ir/Builder.zig index 9f98759..87b394f 100644 --- a/src/intel/compiler/ir/Builder.zig +++ b/src/intel/compiler/ir/Builder.zig @@ -10,6 +10,7 @@ const Self = @This(); pub const Error = std.mem.Allocator.Error || error{ InvalidBlock, + InvalidInstruction, InvalidInsertionIndex, TerminatorAlreadySet, }; @@ -55,13 +56,7 @@ pub fn edge(self: *Self, target: ids.BlockId, arguments: []const pseudo.EdgeArgu }; } -pub fn appendInstruction( - self: *Self, - block_id: ids.BlockId, - execution_size: device.ExecutionSize, - predicate: ?operand.Predicate, - operation: instruction.Operation, -) Error!ids.InstructionId { +pub fn appendInstruction(self: *Self, block_id: ids.BlockId, execution_size: device.ExecutionSize, predicate: ?operand.Predicate, operation: instruction.Operation) Error!ids.InstructionId { const block = self.program.blocks.get(block_id) orelse return Error.InvalidBlock; return self.insertInstruction(block_id, block.instructions.items.len, execution_size, predicate, operation); } @@ -91,6 +86,12 @@ pub fn insertInstruction( return instruction_id; } +pub fn replaceOperation(self: *Self, instruction_id: ids.InstructionId, operation: instruction.Operation) Error!void { + const inst = self.program.instructions.getMut(instruction_id) orelse return Error.InvalidInstruction; + const owned_operation = try instruction.cloneOperation(self.program.allocator(), operation); + inst.operation = owned_operation; +} + pub fn setStructuredControl(self: *Self, block_id: ids.BlockId, control: instruction.StructuredControl) Error!void { const block = self.program.blocks.getMut(block_id) orelse return Error.InvalidBlock; block.structured_control = control; @@ -162,6 +163,13 @@ test "[ir] Builder: construction and ordered insertion" { try std.testing.expectEqual(entry, program.instructions.get(first).?.parent_block); try std.testing.expectEqual(entry, program.instructions.get(second).?.parent_block); + try builder.replaceOperation(first, moveImmediate(register_id, 3)); + const replaced = program.instructions.get(first).?; + try std.testing.expectEqual(entry, replaced.parent_block); + try std.testing.expectEqual(device.ExecutionSize.simd8, replaced.execution_size); + try std.testing.expectEqual(@as(u32, 3), replaced.operation.move.source.register.immediate.u32); + try std.testing.expectError(Error.InvalidInstruction, builder.replaceOperation(ids.InstructionId.fromIndex(999), moveImmediate(register_id, 4))); + try builder.setStructuredControl(entry, .{ .selection = .{ .merge_block = exit } }); try builder.setTerminator(entry, .{ .jump = try builder.edge(exit, &.{}) }); try builder.setTerminator(exit, .end_thread); diff --git a/src/intel/compiler/ir/program.zig b/src/intel/compiler/ir/program.zig index c2d3716..f5f69c7 100644 --- a/src/intel/compiler/ir/program.zig +++ b/src/intel/compiler/ir/program.zig @@ -8,6 +8,7 @@ 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, @@ -24,7 +25,7 @@ pub const Properties = packed struct { flags_allocated: bool = false, branches_resolved: bool = false, - _padding: u20 = 0, + _padding: u19 = 0, }; pub const VertexPayload = struct { diff --git a/src/intel/compiler/ir/validator.zig b/src/intel/compiler/ir/validator.zig index a4dfc58..8cd3e19 100644 --- a/src/intel/compiler/ir/validator.zig +++ b/src/intel/compiler/ir/validator.zig @@ -31,6 +31,10 @@ pub const Error = error{ DuplicateParallelCopyDestination, PredicatedParallelCopy, UnloweredParallelCopy, + UnloweredStageIo, + UnloweredMessage, + InvalidInterfaceSemantic, + InvalidMessage, EntryBlockHasParameters, DuplicateBlockParameter, EdgeArgumentCountMismatch, @@ -125,8 +129,18 @@ fn validateInstruction(program: *const program_ir.Program, inst: instruction.Ins try validateFlag(program, predicate.flag); switch (inst.operation) { - .load_input => |op| try validateDestination(program, op.destination), - .store_output => |op| try validateSource(program, op.source), + .load_input => |op| { + if (program.properties.stage_io_lowered) + return Error.UnloweredStageIo; + try validateDestination(program, op.destination); + try validateInterfaceSemantic(op.semantic, .input); + }, + .store_output => |op| { + if (program.properties.stage_io_lowered) + return Error.UnloweredStageIo; + try validateSource(program, op.source); + try validateInterfaceSemantic(op.semantic, .output); + }, .move => |op| { try validateDestination(program, op.destination); try validateSource(program, op.source); @@ -142,9 +156,17 @@ fn validateInstruction(program: *const program_ir.Program, inst: instruction.Ins 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) @@ -156,6 +178,29 @@ 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 validateParallelCopy(program: *const program_ir.Program, copy: pseudo.ParallelCopy) Error!void { if (copy.register_copies.len == 0 and copy.flag_copies.len == 0) return Error.EmptyParallelCopy; @@ -280,7 +325,18 @@ fn validateSpan(program: *const program_ir.Program, span: operand.RegisterSpan) if (span.register_count == 0) return Error.InvalidRegisterSpan; switch (span.base) { - .virtual, .physical_grf => try validateRegisterRef(program, 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, } } diff --git a/src/intel/compiler/lower/lower.zig b/src/intel/compiler/lower/lower.zig index db0fbbb..0bf5424 100644 --- a/src/intel/compiler/lower/lower.zig +++ b/src/intel/compiler/lower/lower.zig @@ -1,5 +1,6 @@ const std = @import("std"); -const shader_ir = @import("shader_ir").ir; +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"); @@ -11,6 +12,7 @@ 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, @@ -33,21 +35,26 @@ pub const Error = std.mem.Allocator.Error || error{ const PredicateValue = pseudo.PredicateValue; +const LoweredType = struct { + element_type: operand.DataType, + component_count: usize, +}; + const ValueLocation = union(enum) { - source: operand.Source, + 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 lowerType(self: *const LoweringState, type_id: shader_ir.id.TypeId) Error!operand.DataType { + 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.*) { - .void => Error.UnsupportedType, .integer => |integer| if (integer.bits == 32) switch (integer.signedness) { .unsigned => .u32, @@ -60,6 +67,24 @@ const LoweringState = struct { }; } + 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; @@ -97,13 +122,28 @@ const LoweringState = struct { }; } - fn addRegisterLocation(self: *LoweringState, value_id: shader_ir.id.ValueId, class: operand.RegisterClass) Error!operand.Source { + 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 data_type = try self.lowerType(value.type); - const register_id = try self.addRegister(data_type, class, value.name); - const register_source = self.registerSource(register_id, data_type); - try self.putLocation(value_id, .{ .source = register_source }); - return register_source; + 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 { @@ -117,7 +157,6 @@ const LoweringState = struct { 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; @@ -125,7 +164,7 @@ const LoweringState = struct { .boolean => |boolean| .{ .predicate = .{ .constant = boolean } }, else => return Error.UnsupportedType, } else .{ - .source = try self.constantSource(value.type, constant.value), + .components = try self.constantComponents(value.type, constant.value), }; self.value_locations[value_id.index()] = result; return result; @@ -140,23 +179,28 @@ const LoweringState = struct { } } - fn source(self: *LoweringState, value_id: shader_ir.id.ValueId) Error!operand.Source { + fn components(self: *LoweringState, value_id: shader_ir.id.ValueId) Error![]const operand.Source { return switch (try self.location(value_id)) { - .source => |value| value, + .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)) { - .source => Error.UnsupportedType, + .components => Error.UnsupportedType, .predicate => |value| value, }; } - fn destination(self: *LoweringState, value_id: shader_ir.id.ValueId) Error!operand.Destination { - const source_value = try self.source(value_id); - + fn destinationFromSource(source_value: operand.Source) Error!operand.Destination { if (source_value.negate or source_value.absolute) return Error.InvalidLoweredProgram; @@ -170,8 +214,50 @@ const LoweringState = struct { }; } - fn constantSource(self: *const LoweringState, type_id: shader_ir.id.TypeId, value: shader_ir.constant.ConstantValue) Error!operand.Source { - const data_type = try self.lowerType(type_id); + 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) }, @@ -254,13 +340,15 @@ const LoweringState = struct { self.builder.addBlockParameter(target_block_id, .{ .flag = flag_id }) catch |err| return mapProgramError(err); } else { - const parameter_source = try self.addRegisterLocation(parameter_id, .temporary); - 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); + 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); + } } } } @@ -315,10 +403,11 @@ const LoweringState = struct { .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(source_instruction.result, value_id), + .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, .composite_extract => return Error.UnsupportedOperation, + .composite_construct => |operation| try self.lowerCompositeConstruct(source_instruction.result, operation), + .composite_extract => |operation| try self.lowerCompositeExtract(source_instruction.result, operation), .call => return Error.UnsanitizedModule, } } @@ -334,85 +423,81 @@ const LoweringState = struct { 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); - switch (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 }); - }, - .negate => { - const source_value = try self.source(operation.operand); - if (source_value.type != .i32 and source_value.type != .f32) - return Error.UnsupportedOperation; - _ = try self.addRegisterLocation(result_id, .temporary); - var negated = source_value; - negated.negate = !negated.negate; - try self.appendMove(block_id, null, try self.destination(result_id), negated); - }, - .bitwise_not => { - const source_value = try self.source(operation.operand); - if (source_value.type != .u32 and source_value.type != .i32) - return Error.UnsupportedOperation; - _ = try self.addRegisterLocation(result_id, .temporary); - const all_ones: operand.Immediate = switch (source_value.type) { - .u32 => .{ .u32 = std.math.maxInt(u32) }, - .i32 => .{ .i32 = -1 }, - else => unreachable, - }; - try self.appendInstruction(block_id, null, .{ - .binary = .{ - .opcode = .bitwise_xor, - .destination = try self.destination(result_id), - .lhs = source_value, - .rhs = .{ - .register = .{ .immediate = all_ones }, - .type = source_value.type, - .region = operand.Region.broadcast(), + 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 = try self.source(operation.lhs); - var rhs = try self.source(operation.rhs); - _ = try self.addRegisterLocation(result_id, .temporary); - const destination_value = try self.destination(result_id); - - if (lhs.type != destination_value.type or rhs.type != destination_value.type) + 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 (lhs.type == .u32 or lhs.type == .i32) .add else return Error.UnsupportedOperation, - .float_add => if (lhs.type == .f32) .add else return Error.UnsupportedOperation, - - .integer_subtract => if (lhs.type == .u32 or lhs.type == .i32) subtract: { - rhs.negate = !rhs.negate; - break :subtract .add; - } else return Error.UnsupportedOperation, - - .float_subtract => if (lhs.type == .f32) subtract: { - rhs.negate = !rhs.negate; - break :subtract .add; - } else return Error.UnsupportedOperation, - - .integer_multiply => if (lhs.type == .u32 or lhs.type == .i32) .multiply else return Error.UnsupportedOperation, - .float_multiply => if (lhs.type == .f32) .multiply else return Error.UnsupportedOperation, - .shift_left => if (lhs.type == .u32 or lhs.type == .i32) .shift_left else return Error.UnsupportedOperation, - .logical_shift_right => if (lhs.type == .u32) .shift_right else return Error.UnsupportedOperation, - .arithmetic_shift_right => if (lhs.type == .i32) .shift_right else return Error.UnsupportedOperation, - .bitwise_and => if (lhs.type == .u32 or lhs.type == .i32) .bitwise_and else return Error.UnsupportedOperation, - .bitwise_or => if (lhs.type == .u32 or lhs.type == .i32) .bitwise_or else return Error.UnsupportedOperation, - .bitwise_xor => if (lhs.type == .u32 or lhs.type == .i32) .bitwise_xor else return Error.UnsupportedOperation, + .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, @@ -424,14 +509,21 @@ const LoweringState = struct { => return Error.UnsupportedOperation, }; - try self.appendInstruction(block_id, null, .{ - .binary = .{ - .opcode = opcode, - .destination = destination_value, - .lhs = lhs, - .rhs = rhs, - }, - }); + 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 { @@ -441,8 +533,12 @@ const LoweringState = struct { if (!try self.isBoolean(result_value.type)) return Error.InvalidModule; - const lhs = try self.source(operation.lhs); - const rhs = try self.source(operation.rhs); + 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; @@ -477,44 +573,89 @@ const LoweringState = struct { 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_value = try self.source(operation.true_value); - const false_value = try self.source(operation.false_value); - _ = try self.addRegisterLocation(result_id, .temporary); - const destination_value = try self.destination(result_id); - - if (true_value.type != destination_value.type or false_value.type != destination_value.type) + 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; - switch (try self.predicate(operation.condition)) { - .constant => |condition| try self.appendMove( - block_id, - null, - destination_value, - if (condition) true_value else false_value, - ), - .dynamic => |condition| { - try self.appendMove(block_id, .{ - .flag = condition.flag, - .inverse = !condition.inverse, - }, destination_value, false_value); - try self.appendMove(block_id, condition, destination_value, true_value); - }, + 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, result: ?shader_ir.id.ValueId, source_id: shader_ir.id.ValueId) Error!void { + 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); - var source_value = try self.source(source_id); + 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; - source_value.register = switch (source_value.register) { - .immediate => |immediate| .{ .immediate = bitcastImmediate(immediate, target_type) }, - else => source_value.register, - }; + 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); + } + } - source_value.type = target_type; - try self.putLocation(result_id, .{ .source = source_value }); + 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 { @@ -533,13 +674,15 @@ const LoweringState = struct { if (result_value.type != variable.type) return Error.InvalidModule; - _ = try self.addRegisterLocation(result_id, .varying); - try self.appendInstruction(block_id, null, .{ - .load_input = .{ - .destination = try self.destination(result_id), - .semantic = try lowerInterfaceSemantic(variable.semantic), - }, - }); + 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 { @@ -553,18 +696,20 @@ const LoweringState = struct { if (variable.direction != .output) return Error.InvalidModule; - const source_value = try self.source(operation.value); + 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; - try self.appendInstruction(block_id, null, .{ - .store_output = .{ - .semantic = try lowerInterfaceSemantic(variable.semantic), - .source = source_value, - }, - }); + 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 { @@ -624,18 +769,19 @@ const LoweringState = struct { if (edge.arguments.len != target_source_block.parameters.items.len) return Error.InvalidModule; - const arguments = try allocator.alloc(pseudo.EdgeArgument, edge.arguments.len); - errdefer allocator.free(arguments); - for (edge.arguments, arguments) |argument_id, *argument| { - argument.* = switch (try self.location(argument_id)) { - .source => |source_value| .{ .source = source_value }, - .predicate => |predicate_value| .{ .predicate = predicate_value }, - }; + 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 = arguments, + .arguments = try arguments.toOwnedSlice(allocator), }; } }; @@ -668,6 +814,8 @@ pub const Lowerer = struct { // 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; @@ -705,6 +853,7 @@ pub const Lowerer = struct { var state: LoweringState = .{ .lowerer = self, .builder = Builder.init(&program), + .storage = program.allocator(), .block_map = block_map, .value_locations = value_locations, }; @@ -714,7 +863,7 @@ pub const Lowerer = struct { try state.lowerInstructions(allocator); try state.lowerControlAndTerminators(allocator); - program.properties.instructions_selected = true; + program.properties.common_ir_lowered = true; validator.validate(&program) catch return Error.InvalidLoweredProgram; block_arguments.run(allocator, &program) catch |err| return switch (err) { @@ -726,25 +875,30 @@ pub const Lowerer = struct { } }; -fn lowerInterfaceSemantic(semantic: shader_ir.module.InterfaceSemantic) Error!instruction.InterfaceSemantic { +fn lowerInterfaceSemantic(semantic: shader_ir.module.InterfaceSemantic, component_offset: u8) Error!instruction.InterfaceSemantic { return switch (semantic) { - .location => |location| if (location.index == 0) - .{ + .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 = location.component, + .component = component, }, - } - else - Error.UnsupportedOperation, + }; + }, .builtin => |builtin| .{ .builtin = .{ .builtin = switch (builtin) { .position => .position, - .vertex_index => .vertex_index, - .instance_index => .instance_index, + .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, }, }, }; @@ -790,6 +944,8 @@ fn expectLowered(source: []const u8, expected: []const u8) !void { 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); @@ -802,6 +958,8 @@ fn expectLoweredFragments(source: []const u8, expected: []const []const u8, unex 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); @@ -1063,13 +1221,12 @@ test "[ir] Lower: selects and bitcasts" { "[simd8] (+%condition) mov %inverted_choice:u32, 2:u32", "[simd8] (-%condition) mov %inverted_choice:u32, 1:u32", "[simd8] mov %constant_choice:u32, 1:u32", - "[simd8] add %constant_sum:u32, 1065353216: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] add %register_sum:u32, %negative:u32, 1:u32", - }, &.{ - "%one_bits: vgrf", - "%negative_bits: vgrf", - }); + "[simd8] mov %negative_bits:u32, %negative:u32", + "[simd8] add %register_sum:u32, %negative_bits:u32, 1:u32", + }, &.{}); } test "[ir] Lower: vertex interfaces" { @@ -1101,6 +1258,139 @@ test "[ir] Lower: vertex interfaces" { }, &.{}); } +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 @@ -1186,11 +1476,11 @@ test "[ir] Lower: unsupported operations" { \\ fn @main() -> void \\ { \\ .entry(): - \\ %pair: vec2[u32] = composite_construct %one, %two + \\ %wide: vec5[u32] = composite_construct %one, %two, %one, %two, %one \\ return \\ } \\} - , Error.UnsupportedOperation); + , Error.UnsupportedType); try expectLoweringError( \\shader vertex @main @@ -1239,5 +1529,10 @@ test "[ir] Lower: unsupported target configuration" { 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 })); } diff --git a/src/intel/compiler/lower/vertex_abi.zig b/src/intel/compiler/lower/vertex_abi.zig new file mode 100644 index 0000000..6b7cb89 --- /dev/null +++ b/src/intel/compiler/lower/vertex_abi.zig @@ -0,0 +1,485 @@ +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/phi/PhiShaderModule.zig b/src/phi/PhiShaderModule.zig index c91d96a..199d166 100644 --- a/src/phi/PhiShaderModule.zig +++ b/src/phi/PhiShaderModule.zig @@ -15,6 +15,7 @@ pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const v errdefer allocator.destroy(self); var interface = try Interface.init(device, allocator, info); + errdefer interface.deinit(); interface.vtable = &.{ .destroy = destroy }; self.* = .{ @@ -30,6 +31,7 @@ pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void { } pub fn drop(self: *Self, allocator: std.mem.Allocator) void { + self.interface.deinit(); allocator.destroy(self); } @@ -38,7 +40,6 @@ pub fn ref(self: *Self) void { } pub fn unref(self: *Self, allocator: std.mem.Allocator) void { - if (self.ref_count.fetchSub(1, .release) == 1) { + if (self.ref_count.fetchSub(1, .acq_rel) == 1) self.drop(allocator); - } } diff --git a/src/software/SoftShaderModule.zig b/src/software/SoftShaderModule.zig index bb17187..c5c02e4 100644 --- a/src/software/SoftShaderModule.zig +++ b/src/software/SoftShaderModule.zig @@ -5,15 +5,14 @@ const spv = @import("spv"); const VkError = base.VkError; - const Self = @This(); pub const Interface = base.ShaderModule; interface: Interface, module: spv.Module, -/// Pipelines need SPIR-V module reference so shader module may not -/// be destroy on call to `vkDestroyShaderModule` +/// Pipelines need a SPIR-V module reference so the shader module may outlive +/// the application's `vkDestroyShaderModule` call. ref_count: std.atomic.Value(usize), pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const vk.ShaderModuleCreateInfo) VkError!*Self { @@ -21,6 +20,7 @@ pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const v errdefer allocator.destroy(self); var interface = try Interface.init(device, allocator, info); + errdefer interface.deinit(); const device_allocator = device.device_allocator.allocator(); @@ -28,11 +28,9 @@ pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const v .destroy = destroy, }; - const code = info.p_code[0..@divExact(info.code_size, 4)]; - self.* = .{ .interface = interface, - .module = spv.Module.init(device_allocator, code, .{ + .module = spv.Module.init(device_allocator, interface.code(), .{ .use_simd_vectors_specializations = base.config.soft_shaders_simd, }) catch |err| switch (err) { spv.Module.ModuleError.OutOfMemory => return VkError.OutOfHostMemory, @@ -59,9 +57,8 @@ pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void { pub fn drop(self: *Self, allocator: std.mem.Allocator) void { const device_allocator = self.interface.owner.device_allocator.allocator(); - self.module.deinit(device_allocator); - + self.interface.deinit(); allocator.destroy(self); } @@ -70,7 +67,6 @@ pub fn ref(self: *Self) void { } pub fn unref(self: *Self, allocator: std.mem.Allocator) void { - if (self.ref_count.fetchSub(1, .release) == 1) { + if (self.ref_count.fetchSub(1, .acq_rel) == 1) self.drop(allocator); - } } diff --git a/src/vulkan/ShaderModule.zig b/src/vulkan/ShaderModule.zig index 0f71c63..58fce27 100644 --- a/src/vulkan/ShaderModule.zig +++ b/src/vulkan/ShaderModule.zig @@ -1,5 +1,6 @@ const std = @import("std"); const vk = @import("vulkan"); +const shader_ir = @import("shader_ir"); const VkError = @import("error_set.zig").VkError; @@ -7,8 +8,11 @@ const Device = @import("Device.zig"); const Self = @This(); pub const ObjectType: vk.ObjectType = .shader_module; +pub const IrModule = shader_ir.ir.module.Module; +pub const InstantiateOptions = shader_ir.spirv.translator.Options; owner: *Device, +source: shader_ir.spirv.SourceModule, vtable: *const VTable, @@ -16,16 +20,40 @@ pub const VTable = struct { destroy: *const fn (*Self, std.mem.Allocator) void, }; -pub fn init(device: *Device, allocator: std.mem.Allocator, info: *const vk.ShaderModuleCreateInfo) VkError!Self { - _ = allocator; - _ = info; +pub fn init(device: *Device, _: std.mem.Allocator, info: *const vk.ShaderModuleCreateInfo) VkError!Self { + if (info.code_size % @sizeOf(u32) != 0) + return VkError.ValidationFailed; + + const source_allocator = device.device_allocator.allocator(); + const words = info.p_code[0 .. info.code_size / @sizeOf(u32)]; + var source = shader_ir.spirv.SourceModule.init(source_allocator, words) catch |err| return switch (err) { + error.OutOfMemory => VkError.OutOfHostMemory, + else => VkError.ValidationFailed, + }; + errdefer source.deinit(source_allocator); + return .{ .owner = device, + .source = source, // SAFETY: the backend assigns the vtable before returning the shader module. .vtable = undefined, }; } +pub fn deinit(self: *Self) void { + self.source.deinit(self.owner.device_allocator.allocator()); +} + +pub fn code(self: *const Self) []const u32 { + return self.source.code(); +} + +/// Instantiates one entry point as a fresh, backend-agnostic IR module. +/// The returned module does not borrow from this shader module. +pub fn instantiateIr(self: *const Self, allocator: std.mem.Allocator, options: InstantiateOptions) !IrModule { + return shader_ir.spirv.translator.instantiate(allocator, &self.source, options); +} + pub inline fn destroy(self: *Self, allocator: std.mem.Allocator) void { self.vtable.destroy(self, allocator); }