From 0901328f12e7ea3d05dc1d5b4a588e595c4bc0bc Mon Sep 17 00:00:00 2001 From: Ali Cheraghi Date: Wed, 7 May 2025 15:03:42 +0330 Subject: spirv: write error value in an storage buffer --- lib/std/Target.zig | 2 +- lib/std/Target/spirv.zig | 8 +++++++- lib/std/builtin.zig | 1 + 3 files changed, 9 insertions(+), 2 deletions(-) (limited to 'lib') diff --git a/lib/std/Target.zig b/lib/std/Target.zig index 9148fd5fdc..bf5a6369b5 100644 --- a/lib/std/Target.zig +++ b/lib/std/Target.zig @@ -2014,7 +2014,7 @@ pub const Cpu = struct { .global, .local, .shared => is_gpu, .constant => is_gpu and (context == null or context == .constant), .param => is_nvptx, - .input, .output, .uniform, .push_constant, .storage_buffer => is_spirv, + .input, .output, .uniform, .push_constant, .storage_buffer, .physical_storage_buffer => is_spirv, }; } }; diff --git a/lib/std/Target/spirv.zig b/lib/std/Target/spirv.zig index a2575b2fe8..90abacdd08 100644 --- a/lib/std/Target/spirv.zig +++ b/lib/std/Target/spirv.zig @@ -21,6 +21,7 @@ pub const Feature = enum { generic_pointer, vector16, shader, + variable_pointers, physical_storage_buffer, }; @@ -129,6 +130,11 @@ pub const all_features = blk: { .description = "Enable SPV_KHR_physical_storage_buffer extension and the PhysicalStorageBufferAddresses capability", .dependencies = featureSet(&[_]Feature{.v1_0}), }; + result[@intFromEnum(Feature.variable_pointers)] = .{ + .llvm_name = null, + .description = "Enable SPV_KHR_variable_pointers extension and the (VariablePointers, VariablePointersStorageBuffer) capabilities", + .dependencies = featureSet(&[_]Feature{.v1_0}), + }; const ti = @typeInfo(Feature); for (&result, 0..) |*elem, i| { elem.index = i; @@ -147,7 +153,7 @@ pub const cpu = struct { pub const vulkan_v1_2: CpuModel = .{ .name = "vulkan_v1_2", .llvm_name = null, - .features = featureSet(&[_]Feature{ .v1_5, .shader, .physical_storage_buffer }), + .features = featureSet(&[_]Feature{ .v1_5, .shader }), }; pub const opencl_v2: CpuModel = .{ diff --git a/lib/std/builtin.zig b/lib/std/builtin.zig index 852b94c324..1683cc500b 100644 --- a/lib/std/builtin.zig +++ b/lib/std/builtin.zig @@ -531,6 +531,7 @@ pub const AddressSpace = enum(u5) { uniform, push_constant, storage_buffer, + physical_storage_buffer, // AVR address spaces. flash, -- cgit v1.2.3 From 8fa54eb7987bdb8138c625f03aa9fb91239dba48 Mon Sep 17 00:00:00 2001 From: Ali Cheraghi Date: Sun, 11 May 2025 15:45:44 +0330 Subject: spirv: error when execution mode is set more than once --- lib/std/gpu.zig | 228 ++++++++++++++++------------------------ src/codegen/spirv.zig | 8 +- src/codegen/spirv/Assembler.zig | 22 +++- src/codegen/spirv/Module.zig | 46 +++++--- src/link/SpirV.zig | 5 +- 5 files changed, 143 insertions(+), 166 deletions(-) (limited to 'lib') diff --git a/lib/std/gpu.zig b/lib/std/gpu.zig index d02b2424d4..d72d298b32 100644 --- a/lib/std/gpu.zig +++ b/lib/std/gpu.zig @@ -1,81 +1,24 @@ const std = @import("std.zig"); -/// Will make `ptr` contain the location of the current invocation within the -/// global workgroup. Each component is equal to the index of the local workgroup -/// multiplied by the size of the local workgroup plus `localInvocationId`. -/// `ptr` must be a reference to variable or struct field. -pub fn globalInvocationId(comptime ptr: *addrspace(.input) @Vector(3, u32)) void { - asm volatile ( - \\OpDecorate %ptr BuiltIn GlobalInvocationId - : - : [ptr] "" (ptr), - ); -} - -/// Will make that variable contain the location of the current cluster -/// culling, task, mesh, or compute shader invocation within the local -/// workgroup. Each component ranges from zero through to the size of the -/// workgroup in that dimension minus one. -/// `ptr` must be a reference to variable or struct field. -pub fn localInvocationId(comptime ptr: *addrspace(.input) @Vector(3, u32)) void { - asm volatile ( - \\OpDecorate %ptr BuiltIn LocalInvocationId - : - : [ptr] "" (ptr), - ); -} - -/// Output vertex position from a `Vertex` entrypoint -/// `ptr` must be a reference to variable or struct field. -pub fn position(comptime ptr: *addrspace(.output) @Vector(4, f32)) void { - asm volatile ( - \\OpDecorate %ptr BuiltIn Position - : - : [ptr] "" (ptr), - ); -} - -/// Will make `ptr` contain the index of the vertex that is -/// being processed by the current vertex shader invocation. -/// `ptr` must be a reference to variable or struct field. -pub fn vertexIndex(comptime ptr: *addrspace(.input) u32) void { - asm volatile ( - \\OpDecorate %ptr BuiltIn VertexIndex - : - : [ptr] "" (ptr), - ); -} - -/// Will make `ptr` contain the index of the instance that is -/// being processed by the current vertex shader invocation. -/// `ptr` must be a reference to variable or struct field. -pub fn instanceIndex(comptime ptr: *addrspace(.input) u32) void { - asm volatile ( - \\OpDecorate %ptr BuiltIn InstanceIndex - : - : [ptr] "" (ptr), - ); -} - -/// Output fragment depth from a `Fragment` entrypoint -/// `ptr` must be a reference to variable or struct field. -pub fn fragmentCoord(comptime ptr: *addrspace(.input) @Vector(4, f32)) void { - asm volatile ( - \\OpDecorate %ptr BuiltIn FragCoord - : - : [ptr] "" (ptr), - ); -} - -/// Output fragment depth from a `Fragment` entrypoint -/// `ptr` must be a reference to variable or struct field. -pub fn fragmentDepth(comptime ptr: *addrspace(.output) f32) void { - asm volatile ( - \\OpDecorate %ptr BuiltIn FragDepth - : - : [ptr] "" (ptr), - ); -} +pub const position_in = @extern(*addrspace(.input) @Vector(4, f32), .{ .name = "position" }); +pub const position_out = @extern(*addrspace(.output) @Vector(4, f32), .{ .name = "position" }); +pub const point_size_in = @extern(*addrspace(.input) f32, .{ .name = "point_size" }); +pub const point_size_out = @extern(*addrspace(.output) f32, .{ .name = "point_size" }); +pub extern const invocation_id: u32 addrspace(.input); +pub extern const frag_coord: @Vector(4, f32) addrspace(.input); +pub extern const point_coord: @Vector(2, f32) addrspace(.input); +// TODO: direct/indirect values +// pub extern const front_facing: bool addrspace(.input); +// TODO: runtime array +// pub extern const sample_mask; +pub extern var frag_depth: f32 addrspace(.output); +pub extern const num_workgroups: @Vector(3, u32) addrspace(.input); +pub extern const workgroup_size: @Vector(3, u32) addrspace(.input); +pub extern const workgroup_id: @Vector(3, u32) addrspace(.input); +pub extern const local_invocation_id: @Vector(3, u32) addrspace(.input); +pub extern const global_invocation_id: @Vector(3, u32) addrspace(.input); +pub extern const vertex_index: u32 addrspace(.input); +pub extern const instance_index: u32 addrspace(.input); /// Forms the main linkage for `input` and `output` address spaces. /// `ptr` must be a reference to variable or struct field. @@ -101,74 +44,85 @@ pub fn binding(comptime ptr: anytype, comptime set: u32, comptime bind: u32) voi ); } -pub const Origin = enum(u32) { - /// Increase toward the right and downward - upper_left = 7, - /// Increase toward the right and upward - lower_left = 8, -}; - -/// The coordinates appear to originate in the specified `origin`. -/// Only valid with the `Fragment` calling convention. -pub fn fragmentOrigin(comptime entry_point: anytype, comptime origin: Origin) void { - asm volatile ( - \\OpExecutionMode %entry_point $origin - : - : [entry_point] "" (entry_point), - [origin] "c" (@intFromEnum(origin)), - ); -} - -pub const DepthMode = enum(u32) { - /// Declares that this entry point dynamically writes the - /// `fragmentDepth` built in-decorated variable. - replacing = 12, +pub const ExecutionMode = union(Tag) { + /// Sets origin of the framebuffer to the upper-left corner + origin_upper_left, + /// Sets origin of the framebuffer to the lower-left corner + origin_lower_left, + /// Indicates that the fragment shader writes to `frag_depth`, + /// replacing the fixed-function depth value. + depth_replacing, /// Indicates that per-fragment tests may assume that - /// any `fragmentDepth` built in-decorated value written by the shader is + /// any `frag_depth` built in-decorated value written by the shader is /// greater-than-or-equal to the fragment’s interpolated depth value - greater = 14, + depth_greater, /// Indicates that per-fragment tests may assume that - /// any `fragmentDepth` built in-decorated value written by the shader is + /// any `frag_depth` built in-decorated value written by the shader is /// less-than-or-equal to the fragment’s interpolated depth value - less = 15, + depth_less, /// Indicates that per-fragment tests may assume that - /// any `fragmentDepth` built in-decorated value written by the shader is + /// any `frag_depth` built in-decorated value written by the shader is /// the same as the fragment’s interpolated depth value - unchanged = 16, -}; + depth_unchanged, + /// Indicates the workgroup size in the x, y, and z dimensions. + local_size: LocalSize, -/// Only valid with the `Fragment` calling convention. -pub fn depthMode(comptime entry_point: anytype, comptime mode: DepthMode) void { - asm volatile ( - \\OpExecutionMode %entry_point $mode - : - : [entry_point] "" (entry_point), - [mode] "c" (mode), - ); -} + pub const Tag = enum(u32) { + origin_upper_left = 7, + origin_lower_left = 8, + depth_replacing = 12, + depth_greater = 14, + depth_less = 15, + depth_unchanged = 16, + local_size = 17, + }; -/// Indicates the workgroup size in the `x`, `y`, and `z` dimensions. -/// Only valid with the `GLCompute` or `Kernel` calling conventions. -pub fn workgroupSize(comptime entry_point: anytype, comptime size: @Vector(3, u32)) void { - asm volatile ( - \\OpExecutionMode %entry_point LocalSize %x %y %z - : - : [entry_point] "" (entry_point), - [x] "c" (size[0]), - [y] "c" (size[1]), - [z] "c" (size[2]), - ); -} + pub const LocalSize = struct { x: u32, y: u32, z: u32 }; +}; -/// A hint to the client, which indicates the workgroup size in the `x`, `y`, and `z` dimensions. -/// Only valid with the `GLCompute` or `Kernel` calling conventions. -pub fn workgroupSizeHint(comptime entry_point: anytype, comptime size: @Vector(3, u32)) void { - asm volatile ( - \\OpExecutionMode %entry_point LocalSizeHint %x %y %z - : - : [entry_point] "" (entry_point), - [x] "c" (size[0]), - [y] "c" (size[1]), - [z] "c" (size[2]), - ); +/// Declare the mode entry point executes in. +pub fn executionMode(comptime entry_point: anytype, comptime mode: ExecutionMode) void { + const cc = @typeInfo(@TypeOf(entry_point)).@"fn".calling_convention; + switch (mode) { + .origin_upper_left, + .origin_lower_left, + .depth_replacing, + .depth_greater, + .depth_less, + .depth_unchanged, + => { + if (cc != .spirv_fragment) { + @compileError( + \\invalid execution mode ' + ++ @tagName(mode) ++ + \\' for function with ' + ++ @tagName(cc) ++ + \\' calling convention + ); + } + asm volatile ( + \\OpExecutionMode %entry_point $mode + : + : [entry_point] "" (entry_point), + [mode] "c" (@intFromEnum(mode)), + ); + }, + .local_size => |size| { + if (cc != .spirv_kernel) { + @compileError( + \\invalid execution mode 'local_size' for function with ' + ++ @tagName(cc) ++ + \\' calling convention + ); + } + asm volatile ( + \\OpExecutionMode %entry_point LocalSize $x $y $z + : + : [entry_point] "" (entry_point), + [x] "c" (size.x), + [y] "c" (size.y), + [z] "c" (size.z), + ); + }, + } } diff --git a/src/codegen/spirv.zig b/src/codegen/spirv.zig index 99f948e789..b2ab76e2c7 100644 --- a/src/codegen/spirv.zig +++ b/src/codegen/spirv.zig @@ -2870,7 +2870,7 @@ const NavGen = struct { }; try self.spv.declareDeclDeps(spv_decl_index, decl_deps.items); - try self.spv.declareEntryPoint(spv_decl_index, test_name, execution_mode); + try self.spv.declareEntryPoint(spv_decl_index, test_name, execution_mode, null); } fn genNav(self: *NavGen, do_codegen: bool) !void { @@ -2976,10 +2976,6 @@ const NavGen = struct { try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .Position } }); } else if (nav.fqn.eqlSlice("point_size", ip)) { try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .PointSize } }); - } else if (nav.fqn.eqlSlice("vertex_id", ip)) { - try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .VertexId } }); - } else if (nav.fqn.eqlSlice("instance_id", ip)) { - try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .InstanceId } }); } else if (nav.fqn.eqlSlice("invocation_id", ip)) { try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .InvocationId } }); } else if (nav.fqn.eqlSlice("frag_coord", ip)) { @@ -2990,8 +2986,6 @@ const NavGen = struct { try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .FrontFacing } }); } else if (nav.fqn.eqlSlice("sample_mask", ip)) { try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .SampleMask } }); - } else if (nav.fqn.eqlSlice("sample_mask", ip)) { - try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .SampleMask } }); } else if (nav.fqn.eqlSlice("frag_depth", ip)) { try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .FragDepth } }); } else if (nav.fqn.eqlSlice("num_workgroups", ip)) { diff --git a/src/codegen/spirv/Assembler.zig b/src/codegen/spirv/Assembler.zig index e4ad326006..2cf336b9c4 100644 --- a/src/codegen/spirv/Assembler.zig +++ b/src/codegen/spirv/Assembler.zig @@ -296,12 +296,26 @@ fn processInstruction(self: *Assembler) !void { }; break :blk .{ .value = try self.spv.importInstructionSet(set_tag) }; }, + .OpExecutionMode, .OpExecutionModeId => { + assert(try self.processGenericInstruction() == null); + const entry_point_id = try self.resolveRefId(self.inst.operands.items[0].ref_id); + const exec_mode: spec.ExecutionMode = @enumFromInt(self.inst.operands.items[1].value); + const gop = try self.spv.entry_points.getOrPut(self.gpa, entry_point_id); + if (!gop.found_existing) { + gop.value_ptr.* = .{}; + } else if (gop.value_ptr.exec_mode != null) { + return self.fail( + self.currentToken().start, + "cannot set execution mode more than once to any entry point", + .{}, + ); + } + gop.value_ptr.exec_mode = exec_mode; + return; + }, else => switch (self.inst.opcode.class()) { .TypeDeclaration => try self.processTypeInstruction(), - else => if (try self.processGenericInstruction()) |result| - result - else - return, + else => (try self.processGenericInstruction()) orelse return, }, }; diff --git a/src/codegen/spirv/Module.zig b/src/codegen/spirv/Module.zig index 920215bee1..691749bf1d 100644 --- a/src/codegen/spirv/Module.zig +++ b/src/codegen/spirv/Module.zig @@ -92,11 +92,12 @@ pub const Decl = struct { /// This models a kernel entry point. pub const EntryPoint = struct { /// The declaration that should be exported. - decl_index: Decl.Index, + decl_index: ?Decl.Index = null, /// The name of the kernel to be exported. - name: []const u8, + name: ?[]const u8 = null, /// Calling Convention - execution_model: spec.ExecutionModel, + exec_model: ?spec.ExecutionModel = null, + exec_mode: ?spec.ExecutionMode = null, }; /// A general-purpose allocator which may be used to allocate resources for this module @@ -184,7 +185,7 @@ decls: std.ArrayListUnmanaged(Decl) = .empty, decl_deps: std.ArrayListUnmanaged(Decl.Index) = .empty, /// The list of entry points that should be exported from this module. -entry_points: std.ArrayListUnmanaged(EntryPoint) = .empty, +entry_points: std.AutoArrayHashMapUnmanaged(IdRef, EntryPoint) = .empty, pub fn init(gpa: Allocator, target: std.Target) Module { const version_minor: u8 = blk: { @@ -304,19 +305,30 @@ fn entryPoints(self: *Module) !Section { var seen = try std.DynamicBitSetUnmanaged.initEmpty(self.gpa, self.decls.items.len); defer seen.deinit(self.gpa); - for (self.entry_points.items) |entry_point| { + for (self.entry_points.keys(), self.entry_points.values()) |entry_point_id, entry_point| { interface.items.len = 0; seen.setRangeValue(.{ .start = 0, .end = self.decls.items.len }, false); - try self.addEntryPointDeps(entry_point.decl_index, &seen, &interface); - - const entry_point_id = self.declPtr(entry_point.decl_index).result_id; + try self.addEntryPointDeps(entry_point.decl_index.?, &seen, &interface); try entry_points.emit(self.gpa, .OpEntryPoint, .{ - .execution_model = entry_point.execution_model, + .execution_model = entry_point.exec_model.?, .entry_point = entry_point_id, - .name = entry_point.name, + .name = entry_point.name.?, .interface = interface.items, }); + + if (entry_point.exec_mode == null and entry_point.exec_model == .Fragment) { + switch (self.target.os.tag) { + .vulkan, .opengl => |tag| { + try self.sections.execution_modes.emit(self.gpa, .OpExecutionMode, .{ + .entry_point = entry_point_id, + .mode = if (tag == .vulkan) .OriginUpperLeft else .OriginLowerLeft, + }); + }, + .opencl => {}, + else => unreachable, + } + } } return entry_points; @@ -749,13 +761,15 @@ pub fn declareEntryPoint( self: *Module, decl_index: Decl.Index, name: []const u8, - execution_model: spec.ExecutionModel, + exec_model: spec.ExecutionModel, + exec_mode: ?spec.ExecutionMode, ) !void { - try self.entry_points.append(self.gpa, .{ - .decl_index = decl_index, - .name = try self.arena.allocator().dupe(u8, name), - .execution_model = execution_model, - }); + const gop = try self.entry_points.getOrPut(self.gpa, self.declPtr(decl_index).result_id); + gop.value_ptr.decl_index = decl_index; + gop.value_ptr.name = try self.arena.allocator().dupe(u8, name); + gop.value_ptr.exec_model = exec_model; + // Might've been set by assembler + if (!gop.found_existing) gop.value_ptr.exec_mode = exec_mode; } pub fn debugName(self: *Module, target: IdResult, name: []const u8) !void { diff --git a/src/link/SpirV.zig b/src/link/SpirV.zig index f5e569ce69..f3c2922725 100644 --- a/src/link/SpirV.zig +++ b/src/link/SpirV.zig @@ -162,7 +162,7 @@ pub fn updateExports( if (ip.isFunctionType(nav_ty)) { const spv_decl_index = try self.object.resolveNav(zcu, nav_index); const cc = Type.fromInterned(nav_ty).fnCallingConvention(zcu); - const execution_model: spec.ExecutionModel = switch (target.os.tag) { + const exec_model: spec.ExecutionModel = switch (target.os.tag) { .vulkan, .opengl => switch (cc) { .spirv_vertex => .Vertex, .spirv_fragment => .Fragment, @@ -185,7 +185,8 @@ pub fn updateExports( try self.object.spv.declareEntryPoint( spv_decl_index, exp.opts.name.toSlice(ip), - execution_model, + exec_model, + null, ); } } -- cgit v1.2.3 From 4bf1e4d198abd2018bf23f9067617800a2bc0554 Mon Sep 17 00:00:00 2001 From: Ali Cheraghi Date: Wed, 21 May 2025 15:26:18 +0330 Subject: target: auto-generated spirv features --- lib/std/Target/spirv.zig | 205 +++++++++++++++++++++++++----------------- tools/update_cpu_features.zig | 122 +++++++++++++++++++++++++ 2 files changed, 247 insertions(+), 80 deletions(-) (limited to 'lib') diff --git a/lib/std/Target/spirv.zig b/lib/std/Target/spirv.zig index 90abacdd08..229d77a6d6 100644 --- a/lib/std/Target/spirv.zig +++ b/lib/std/Target/spirv.zig @@ -1,8 +1,21 @@ +//! This file is auto-generated by tools/update_cpu_features.zig. + const std = @import("../std.zig"); const CpuFeature = std.Target.Cpu.Feature; const CpuModel = std.Target.Cpu.Model; pub const Feature = enum { + addresses, + arbitrary_precision_integers, + float16, + float64, + generic_pointer, + int64, + kernel, + matrix, + physical_storage_buffer, + shader, + storage_push_constant16, v1_0, v1_1, v1_2, @@ -10,19 +23,8 @@ pub const Feature = enum { v1_4, v1_5, v1_6, - int64, - float16, - float64, - matrix, - storage_push_constant16, - arbitrary_precision_integers, - kernel, - addresses, - generic_pointer, - vector16, - shader, variable_pointers, - physical_storage_buffer, + vector16, }; pub const featureSet = CpuFeature.FeatureSetFns(Feature).featureSet; @@ -35,105 +37,143 @@ pub const all_features = blk: { const len = @typeInfo(Feature).@"enum".fields.len; std.debug.assert(len <= CpuFeature.Set.needed_bit_count); var result: [len]CpuFeature = undefined; - result[@intFromEnum(Feature.v1_0)] = .{ - .llvm_name = null, - .description = "Enable version 1.0", - .dependencies = featureSet(&[_]Feature{}), - }; - result[@intFromEnum(Feature.v1_1)] = .{ - .llvm_name = null, - .description = "Enable version 1.1", - .dependencies = featureSet(&[_]Feature{.v1_0}), - }; - result[@intFromEnum(Feature.v1_2)] = .{ + result[@intFromEnum(Feature.addresses)] = .{ .llvm_name = null, - .description = "Enable version 1.2", - .dependencies = featureSet(&[_]Feature{.v1_1}), + .description = "Enable Addresses capability", + .dependencies = featureSet(&[_]Feature{ + .v1_0, + }), }; - result[@intFromEnum(Feature.v1_3)] = .{ + result[@intFromEnum(Feature.arbitrary_precision_integers)] = .{ .llvm_name = null, - .description = "Enable version 1.3", - .dependencies = featureSet(&[_]Feature{.v1_2}), + .description = "Enable SPV_INTEL_arbitrary_precision_integers extension and the ArbitraryPrecisionIntegersINTEL capability", + .dependencies = featureSet(&[_]Feature{ + .v1_5, + }), }; - result[@intFromEnum(Feature.v1_4)] = .{ + result[@intFromEnum(Feature.float16)] = .{ .llvm_name = null, - .description = "Enable version 1.4", - .dependencies = featureSet(&[_]Feature{.v1_3}), + .description = "Enable Float16 capability", + .dependencies = featureSet(&[_]Feature{ + .v1_0, + }), }; - result[@intFromEnum(Feature.v1_5)] = .{ + result[@intFromEnum(Feature.float64)] = .{ .llvm_name = null, - .description = "Enable version 1.5", - .dependencies = featureSet(&[_]Feature{.v1_4}), + .description = "Enable Float64 capability", + .dependencies = featureSet(&[_]Feature{ + .v1_0, + }), }; - result[@intFromEnum(Feature.v1_6)] = .{ + result[@intFromEnum(Feature.generic_pointer)] = .{ .llvm_name = null, - .description = "Enable version 1.6", - .dependencies = featureSet(&[_]Feature{.v1_5}), + .description = "Enable GenericPointer capability", + .dependencies = featureSet(&[_]Feature{ + .addresses, + }), }; result[@intFromEnum(Feature.int64)] = .{ .llvm_name = null, .description = "Enable Int64 capability", - .dependencies = featureSet(&[_]Feature{.v1_0}), - }; - result[@intFromEnum(Feature.float16)] = .{ - .llvm_name = null, - .description = "Enable Float16 capability", - .dependencies = featureSet(&[_]Feature{.v1_0}), + .dependencies = featureSet(&[_]Feature{ + .v1_0, + }), }; - result[@intFromEnum(Feature.float64)] = .{ + result[@intFromEnum(Feature.kernel)] = .{ .llvm_name = null, - .description = "Enable Float64 capability", - .dependencies = featureSet(&[_]Feature{.v1_0}), + .description = "Enable Kernel capability", + .dependencies = featureSet(&[_]Feature{ + .v1_0, + }), }; result[@intFromEnum(Feature.matrix)] = .{ .llvm_name = null, .description = "Enable Matrix capability", - .dependencies = featureSet(&[_]Feature{.v1_0}), + .dependencies = featureSet(&[_]Feature{ + .v1_0, + }), + }; + result[@intFromEnum(Feature.physical_storage_buffer)] = .{ + .llvm_name = null, + .description = "Enable SPV_KHR_variable_pointers extension and the (VariablePointers, VariablePointersStorageBuffer) capabilities", + .dependencies = featureSet(&[_]Feature{ + .v1_0, + }), + }; + result[@intFromEnum(Feature.shader)] = .{ + .llvm_name = null, + .description = "Enable Shader capability", + .dependencies = featureSet(&[_]Feature{ + .matrix, + }), }; result[@intFromEnum(Feature.storage_push_constant16)] = .{ .llvm_name = null, .description = "Enable SPV_KHR_16bit_storage extension and the StoragePushConstant16 capability", - .dependencies = featureSet(&[_]Feature{.v1_3}), + .dependencies = featureSet(&[_]Feature{ + .v1_3, + }), }; - result[@intFromEnum(Feature.arbitrary_precision_integers)] = .{ + result[@intFromEnum(Feature.v1_0)] = .{ .llvm_name = null, - .description = "Enable SPV_INTEL_arbitrary_precision_integers extension and the ArbitraryPrecisionIntegersINTEL capability", - .dependencies = featureSet(&[_]Feature{.v1_5}), + .description = "Enable version 1.0", + .dependencies = featureSet(&[_]Feature{}), }; - result[@intFromEnum(Feature.kernel)] = .{ + result[@intFromEnum(Feature.v1_1)] = .{ .llvm_name = null, - .description = "Enable Kernel capability", - .dependencies = featureSet(&[_]Feature{.v1_0}), + .description = "Enable version 1.1", + .dependencies = featureSet(&[_]Feature{ + .v1_0, + }), }; - result[@intFromEnum(Feature.addresses)] = .{ + result[@intFromEnum(Feature.v1_2)] = .{ .llvm_name = null, - .description = "Enable Addresses capability", - .dependencies = featureSet(&[_]Feature{.v1_0}), + .description = "Enable version 1.2", + .dependencies = featureSet(&[_]Feature{ + .v1_1, + }), }; - result[@intFromEnum(Feature.generic_pointer)] = .{ + result[@intFromEnum(Feature.v1_3)] = .{ .llvm_name = null, - .description = "Enable GenericPointer capability", - .dependencies = featureSet(&[_]Feature{ .v1_0, .addresses }), + .description = "Enable version 1.3", + .dependencies = featureSet(&[_]Feature{ + .v1_2, + }), }; - result[@intFromEnum(Feature.vector16)] = .{ + result[@intFromEnum(Feature.v1_4)] = .{ .llvm_name = null, - .description = "Enable Vector16 capability", - .dependencies = featureSet(&[_]Feature{ .v1_0, .kernel }), + .description = "Enable version 1.4", + .dependencies = featureSet(&[_]Feature{ + .v1_3, + }), }; - result[@intFromEnum(Feature.shader)] = .{ + result[@intFromEnum(Feature.v1_5)] = .{ .llvm_name = null, - .description = "Enable Shader capability", - .dependencies = featureSet(&[_]Feature{ .v1_0, .matrix }), + .description = "Enable version 1.5", + .dependencies = featureSet(&[_]Feature{ + .v1_4, + }), }; - result[@intFromEnum(Feature.physical_storage_buffer)] = .{ + result[@intFromEnum(Feature.v1_6)] = .{ .llvm_name = null, - .description = "Enable SPV_KHR_physical_storage_buffer extension and the PhysicalStorageBufferAddresses capability", - .dependencies = featureSet(&[_]Feature{.v1_0}), + .description = "Enable version 1.6", + .dependencies = featureSet(&[_]Feature{ + .v1_5, + }), }; result[@intFromEnum(Feature.variable_pointers)] = .{ .llvm_name = null, - .description = "Enable SPV_KHR_variable_pointers extension and the (VariablePointers, VariablePointersStorageBuffer) capabilities", - .dependencies = featureSet(&[_]Feature{.v1_0}), + .description = "Enable SPV_KHR_physical_storage_buffer extension and the PhysicalStorageBufferAddresses capability", + .dependencies = featureSet(&[_]Feature{ + .v1_0, + }), + }; + result[@intFromEnum(Feature.vector16)] = .{ + .llvm_name = null, + .description = "Enable Vector16 capability", + .dependencies = featureSet(&[_]Feature{ + .kernel, + }), }; const ti = @typeInfo(Feature); for (&result, 0..) |*elem, i| { @@ -147,18 +187,23 @@ pub const cpu = struct { pub const generic: CpuModel = .{ .name = "generic", .llvm_name = "generic", - .features = featureSet(&[_]Feature{.v1_0}), - }; - - pub const vulkan_v1_2: CpuModel = .{ - .name = "vulkan_v1_2", - .llvm_name = null, - .features = featureSet(&[_]Feature{ .v1_5, .shader }), + .features = featureSet(&[_]Feature{}), }; - pub const opencl_v2: CpuModel = .{ .name = "opencl_v2", .llvm_name = null, - .features = featureSet(&[_]Feature{ .v1_2, .kernel, .addresses, .generic_pointer }), + .features = featureSet(&[_]Feature{ + .generic_pointer, + .kernel, + .v1_2, + }), + }; + pub const vulkan_v1_2: CpuModel = .{ + .name = "vulkan_v1_2", + .llvm_name = null, + .features = featureSet(&[_]Feature{ + .shader, + .v1_5, + }), }; }; diff --git a/tools/update_cpu_features.zig b/tools/update_cpu_features.zig index 9ac7b7ef2a..f033195aac 100644 --- a/tools/update_cpu_features.zig +++ b/tools/update_cpu_features.zig @@ -1047,6 +1047,128 @@ const targets = [_]ArchTarget{ }, }, }, + .{ + .zig_name = "spirv", + .llvm = .{ + .name = "SPIRV", + .td_name = "SPIRV", + }, + .branch_quota = 2000, + .extra_features = &.{ + .{ + .zig_name = "v1_0", + .desc = "Enable version 1.0", + .deps = &.{}, + }, + .{ + .zig_name = "v1_1", + .desc = "Enable version 1.1", + .deps = &.{"v1_0"}, + }, + .{ + .zig_name = "v1_2", + .desc = "Enable version 1.2", + .deps = &.{"v1_1"}, + }, + .{ + .zig_name = "v1_3", + .desc = "Enable version 1.3", + .deps = &.{"v1_2"}, + }, + .{ + .zig_name = "v1_4", + .desc = "Enable version 1.4", + .deps = &.{"v1_3"}, + }, + .{ + .zig_name = "v1_5", + .desc = "Enable version 1.5", + .deps = &.{"v1_4"}, + }, + .{ + .zig_name = "v1_6", + .desc = "Enable version 1.6", + .deps = &.{"v1_5"}, + }, + .{ + .zig_name = "int64", + .desc = "Enable Int64 capability", + .deps = &.{"v1_0"}, + }, + .{ + .zig_name = "float16", + .desc = "Enable Float16 capability", + .deps = &.{"v1_0"}, + }, + .{ + .zig_name = "float64", + .desc = "Enable Float64 capability", + .deps = &.{"v1_0"}, + }, + .{ + .zig_name = "matrix", + .desc = "Enable Matrix capability", + .deps = &.{"v1_0"}, + }, + .{ + .zig_name = "storage_push_constant16", + .desc = "Enable SPV_KHR_16bit_storage extension and the StoragePushConstant16 capability", + .deps = &.{"v1_3"}, + }, + .{ + .zig_name = "arbitrary_precision_integers", + .desc = "Enable SPV_INTEL_arbitrary_precision_integers extension and the ArbitraryPrecisionIntegersINTEL capability", + .deps = &.{"v1_5"}, + }, + .{ + .zig_name = "kernel", + .desc = "Enable Kernel capability", + .deps = &.{"v1_0"}, + }, + .{ + .zig_name = "addresses", + .desc = "Enable Addresses capability", + .deps = &.{"v1_0"}, + }, + .{ + .zig_name = "generic_pointer", + .desc = "Enable GenericPointer capability", + .deps = &.{ "v1_0", "addresses" }, + }, + .{ + .zig_name = "vector16", + .desc = "Enable Vector16 capability", + .deps = &.{ "v1_0", "kernel" }, + }, + .{ + .zig_name = "shader", + .desc = "Enable Shader capability", + .deps = &.{ "v1_0", "matrix" }, + }, + .{ + .zig_name = "variable_pointers", + .desc = "Enable SPV_KHR_physical_storage_buffer extension and the PhysicalStorageBufferAddresses capability", + .deps = &.{"v1_0"}, + }, + .{ + .zig_name = "physical_storage_buffer", + .desc = "Enable SPV_KHR_variable_pointers extension and the (VariablePointers, VariablePointersStorageBuffer) capabilities", + .deps = &.{"v1_0"}, + }, + }, + .extra_cpus = &.{ + .{ + .llvm_name = null, + .zig_name = "vulkan_v1_2", + .features = &.{ "v1_5", "shader" }, + }, + .{ + .llvm_name = null, + .zig_name = "opencl_v2", + .features = &.{ "v1_2", "kernel", "addresses", "generic_pointer" }, + }, + }, + }, .{ .zig_name = "riscv", .llvm = .{ -- cgit v1.2.3