Mirrors/zig
1
0
Fork 0
mirror of https://codeberg.org/ziglang/zig.git synced 2026-04-26 13:01:34 +03:00
Code Issues Packages Projects Releases Wiki Activity

spirv: snake-case the spec

Browse Source
This commit is contained in:
Ali Cheraghi
2025-07-12 16:05:45 +03:30
committed by Alex Rønne Petersen
parent 2f3cd175d3
commit f43f89a705
13 changed files with 13948 additions and 12088 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/codegen/spirv.zig
+258 -292
View File
@@ -15,9 +15,7 @@ const InternPool = @import("../InternPool.zig");
const spec = @import("spirv/spec.zig");
const Opcode = spec.Opcode;
const Word = spec.Word;
const IdRef = spec.IdRef;
const IdResult = spec.IdResult;
const IdResultType = spec.IdResultType;
const Id = spec.Id;
const StorageClass = spec.StorageClass;
const SpvModule = @import("spirv/Module.zig");
@@ -26,7 +24,7 @@ const IdRange = SpvModule.IdRange;
const SpvSection = @import("spirv/Section.zig");
const SpvAssembler = @import("spirv/Assembler.zig");
const InstMap = std.AutoHashMapUnmanaged(Air.Inst.Index, IdRef);
const InstMap = std.AutoHashMapUnmanaged(Air.Inst.Index, Id);
pub fn legalizeFeatures(_: *const std.Target) *const Air.Legalize.Features {
return comptime &.initMany(&.{
@@ -42,10 +40,10 @@ pub fn legalizeFeatures(_: *const std.Target) *const Air.Legalize.Features {
pub const zig_call_abi_ver = 3;
pub const big_int_bits = 32;
const InternMap = std.AutoHashMapUnmanaged(struct { InternPool.Index, NavGen.Repr }, IdResult);
const InternMap = std.AutoHashMapUnmanaged(struct { InternPool.Index, NavGen.Repr }, Id);
const PtrTypeMap = std.AutoHashMapUnmanaged(
struct { InternPool.Index, StorageClass, NavGen.Repr },
struct { ty_id: IdRef, fwd_emitted: bool },
struct { ty_id: Id, fwd_emitted: bool },
);
const ControlFlow = union(enum) {
@@ -55,10 +53,10 @@ const ControlFlow = union(enum) {
/// inside the block must reach the outside.
const Block = union(enum) {
const Incoming = struct {
src_label: IdRef,
src_label: Id,
/// Instruction that returns an u32 value of the
/// `Air.Inst.Index` that control flow should jump to.
next_block: IdRef,
next_block: Id,
};
const SelectionMerge = struct {
@@ -69,7 +67,7 @@ const ControlFlow = union(enum) {
/// The label id of the cond_br's merge block.
/// For the top-most element in the stack, this
/// value is undefined.
merge_block: IdRef,
merge_block: Id,
};
/// For a `selection` type block, we cannot use early exits, and we
@@ -100,7 +98,7 @@ const ControlFlow = union(enum) {
/// of conditions that jump to the loop exit.
merges: std.ArrayListUnmanaged(Incoming) = .empty,
/// The label id of the loop's merge block.
merge_block: IdRef,
merge_block: Id,
},
fn deinit(self: *Structured.Block, a: Allocator) void {
@@ -116,17 +114,17 @@ const ControlFlow = union(enum) {
block_stack: std.ArrayListUnmanaged(*Structured.Block) = .empty,
/// Maps `block` inst indices to the variable that the block's result
/// value must be written to.
block_results: std.AutoHashMapUnmanaged(Air.Inst.Index, IdRef) = .empty,
block_results: std.AutoHashMapUnmanaged(Air.Inst.Index, Id) = .empty,
};
const Unstructured = struct {
const Incoming = struct {
src_label: IdRef,
break_value_id: IdRef,
src_label: Id,
break_value_id: Id,
};
const Block = struct {
label: ?IdRef = null,
label: ?Id = null,
incoming_blocks: std.ArrayListUnmanaged(Incoming) = .empty,
};
@@ -318,7 +316,7 @@ const NavGen = struct {
/// An array of function argument result-ids. Each index corresponds with the
/// function argument of the same index.
args: std.ArrayListUnmanaged(IdRef) = .empty,
args: std.ArrayListUnmanaged(Id) = .empty,
/// A counter to keep track of how many `arg` instructions we've seen yet.
next_arg_index: u32 = 0,
@@ -337,7 +335,7 @@ const NavGen = struct {
control_flow: ControlFlow,
/// The label of the SPIR-V block we are currently generating.
current_block_label: IdRef,
current_block_label: Id,
/// The code (prologue and body) for the function we are currently generating code for.
func: SpvModule.Fn = .{},
@@ -436,17 +434,17 @@ const NavGen = struct {
/// This imports the "default" extended instruction set for the target
/// For OpenCL, OpenCL.std.100. For Vulkan and OpenGL, GLSL.std.450.
fn importExtendedSet(self: *NavGen) !IdResult {
fn importExtendedSet(self: *NavGen) !Id {
const target = self.spv.target;
return switch (target.os.tag) {
.opencl, .amdhsa => try self.spv.importInstructionSet(.@"OpenCL.std"),
.vulkan, .opengl => try self.spv.importInstructionSet(.@"GLSL.std.450"),
.opencl, .amdhsa => try self.spv.importInstructionSet(.open_cl_std),
.vulkan, .opengl => try self.spv.importInstructionSet(.glsl_std_450),
else => unreachable,
};
}
/// Fetch the result-id for a previously generated instruction or constant.
fn resolve(self: *NavGen, inst: Air.Inst.Ref) !IdRef {
fn resolve(self: *NavGen, inst: Air.Inst.Ref) !Id {
const pt = self.pt;
const zcu = pt.zcu;
if (try self.air.value(inst, pt)) |val| {
@@ -468,7 +466,7 @@ const NavGen = struct {
return self.inst_results.get(index).?; // Assertion means instruction does not dominate usage.
}
fn resolveUav(self: *NavGen, val: InternPool.Index) !IdRef {
fn resolveUav(self: *NavGen, val: InternPool.Index) !Id {
// TODO: This cannot be a function at this point, but it should probably be handled anyway.
const zcu = self.pt.zcu;
@@ -476,16 +474,16 @@ const NavGen = struct {
const decl_ptr_ty_id = try self.ptrType(ty, self.spvStorageClass(.generic), .indirect);
const spv_decl_index = blk: {
const entry = try self.object.uav_link.getOrPut(self.object.gpa, .{ val, .Function });
const entry = try self.object.uav_link.getOrPut(self.object.gpa, .{ val, .function });
if (entry.found_existing) {
try self.addFunctionDep(entry.value_ptr.*, .Function);
try self.addFunctionDep(entry.value_ptr.*, .function);
const result_id = self.spv.declPtr(entry.value_ptr.*).result_id;
return try self.castToGeneric(decl_ptr_ty_id, result_id);
}
const spv_decl_index = try self.spv.allocDecl(.invocation_global);
try self.addFunctionDep(spv_decl_index, .Function);
try self.addFunctionDep(spv_decl_index, .function);
entry.value_ptr.* = spv_decl_index;
break :blk spv_decl_index;
};
@@ -536,7 +534,7 @@ const NavGen = struct {
try self.spv.debugNameFmt(initializer_id, "initializer of __anon_{d}", .{@intFromEnum(val)});
const fn_decl_ptr_ty_id = try self.ptrType(ty, .Function, .indirect);
const fn_decl_ptr_ty_id = try self.ptrType(ty, .function, .indirect);
try self.spv.sections.types_globals_constants.emit(self.spv.gpa, .OpExtInst, .{
.id_result_type = fn_decl_ptr_ty_id,
.id_result = result_id,
@@ -552,7 +550,7 @@ const NavGen = struct {
fn addFunctionDep(self: *NavGen, decl_index: SpvModule.Decl.Index, storage_class: StorageClass) !void {
if (self.spv.version.minor < 4) {
// Before version 1.4, the interface’s storage classes are limited to the Input and Output
if (storage_class == .Input or storage_class == .Output) {
if (storage_class == .input or storage_class == .output) {
try self.func.decl_deps.put(self.spv.gpa, decl_index, {});
}
} else {
@@ -560,7 +558,7 @@ const NavGen = struct {
}
}
fn castToGeneric(self: *NavGen, type_id: IdRef, ptr_id: IdRef) !IdRef {
fn castToGeneric(self: *NavGen, type_id: Id, ptr_id: Id) !Id {
if (self.spv.hasFeature(.generic_pointer)) {
const result_id = self.spv.allocId();
try self.func.body.emit(self.spv.gpa, .OpPtrCastToGeneric, .{
@@ -578,7 +576,7 @@ const NavGen = struct {
/// block we are currently generating.
/// Note that there is no such thing as nested blocks like in ZIR or AIR, so we don't need to
/// keep track of the previous block.
fn beginSpvBlock(self: *NavGen, label: IdResult) !void {
fn beginSpvBlock(self: *NavGen, label: Id) !void {
try self.func.body.emit(self.spv.gpa, .OpLabel, .{ .id_result = label });
self.current_block_label = label;
}
@@ -705,7 +703,7 @@ const NavGen = struct {
}
/// Emits a bool constant in a particular representation.
fn constBool(self: *NavGen, value: bool, repr: Repr) !IdRef {
fn constBool(self: *NavGen, value: bool, repr: Repr) !Id {
return switch (repr) {
.indirect => self.constInt(Type.u1, @intFromBool(value)),
.direct => self.spv.constBool(value),
@@ -715,7 +713,7 @@ const NavGen = struct {
/// Emits an integer constant.
/// This function, unlike SpvModule.constInt, takes care to bitcast
/// the value to an unsigned int first for Kernels.
fn constInt(self: *NavGen, ty: Type, value: anytype) !IdRef {
fn constInt(self: *NavGen, ty: Type, value: anytype) !Id {
const zcu = self.pt.zcu;
const scalar_ty = ty.scalarType(zcu);
const int_info = scalar_ty.intInfo(zcu);
@@ -773,7 +771,7 @@ const NavGen = struct {
return self.constructCompositeSplat(ty, result_id);
}
pub fn constructComposite(self: *NavGen, result_ty_id: IdRef, constituents: []const IdRef) !IdRef {
pub fn constructComposite(self: *NavGen, result_ty_id: Id, constituents: []const Id) !Id {
const result_id = self.spv.allocId();
try self.func.body.emit(self.gpa, .OpCompositeConstruct, .{
.id_result_type = result_ty_id,
@@ -785,11 +783,11 @@ const NavGen = struct {
/// Construct a composite at runtime with all lanes set to the same value.
/// ty must be an aggregate type.
fn constructCompositeSplat(self: *NavGen, ty: Type, constituent: IdRef) !IdRef {
fn constructCompositeSplat(self: *NavGen, ty: Type, constituent: Id) !Id {
const zcu = self.pt.zcu;
const n: usize = @intCast(ty.arrayLen(zcu));
const constituents = try self.gpa.alloc(IdRef, n);
const constituents = try self.gpa.alloc(Id, n);
defer self.gpa.free(constituents);
@memset(constituents, constituent);
@@ -803,7 +801,7 @@ const NavGen = struct {
/// is done by emitting a sequence of instructions that initialize the value.
//
/// This function should only be called during function code generation.
fn constant(self: *NavGen, ty: Type, val: Value, repr: Repr) !IdRef {
fn constant(self: *NavGen, ty: Type, val: Value, repr: Repr) !Id {
// Note: Using intern_map can only be used with constants that DO NOT generate any runtime code!!
// Ideally that should be all constants in the future, or it should be cleaned up somehow. For
// now, only use the intern_map on case-by-case basis by breaking to :cache.
@@ -909,7 +907,7 @@ const NavGen = struct {
.payload => |payload| payload,
});
var constituents: [2]IdRef = undefined;
var constituents: [2]Id = undefined;
var types: [2]Type = undefined;
if (eu_layout.error_first) {
constituents[0] = try self.constant(err_ty, err_val, .indirect);
@@ -967,7 +965,7 @@ const NavGen = struct {
inline .array_type, .vector_type => |array_type, tag| {
const elem_ty = Type.fromInterned(array_type.child);
const constituents = try self.gpa.alloc(IdRef, @intCast(ty.arrayLenIncludingSentinel(zcu)));
const constituents = try self.gpa.alloc(Id, @intCast(ty.arrayLenIncludingSentinel(zcu)));
defer self.gpa.free(constituents);
const child_repr: Repr = switch (tag) {
@@ -1015,7 +1013,7 @@ const NavGen = struct {
var types = std.ArrayList(Type).init(self.gpa);
defer types.deinit();
var constituents = std.ArrayList(IdRef).init(self.gpa);
var constituents = std.ArrayList(Id).init(self.gpa);
defer constituents.deinit();
var it = struct_type.iterateRuntimeOrder(ip);
@@ -1064,7 +1062,7 @@ const NavGen = struct {
return cacheable_id;
}
fn constantPtr(self: *NavGen, ptr_val: Value) Error!IdRef {
fn constantPtr(self: *NavGen, ptr_val: Value) Error!Id {
const pt = self.pt;
if (ptr_val.isUndef(pt.zcu)) {
@@ -1080,7 +1078,7 @@ const NavGen = struct {
return self.derivePtr(derivation);
}
fn derivePtr(self: *NavGen, derivation: Value.PointerDeriveStep) Error!IdRef {
fn derivePtr(self: *NavGen, derivation: Value.PointerDeriveStep) Error!Id {
const pt = self.pt;
const zcu = pt.zcu;
switch (derivation) {
@@ -1159,7 +1157,7 @@ const NavGen = struct {
self: *NavGen,
ty: Type,
uav: InternPool.Key.Ptr.BaseAddr.Uav,
) !IdRef {
) !Id {
// TODO: Merge this function with constantDeclRef.
const pt = self.pt;
@@ -1182,7 +1180,7 @@ const NavGen = struct {
// Uav refs are always generic.
assert(ty.ptrAddressSpace(zcu) == .generic);
const decl_ptr_ty_id = try self.ptrType(uav_ty, .Generic, .indirect);
const decl_ptr_ty_id = try self.ptrType(uav_ty, .generic, .indirect);
const ptr_id = try self.resolveUav(uav.val);
if (decl_ptr_ty_id != ty_id) {
@@ -1199,7 +1197,7 @@ const NavGen = struct {
}
}
fn constantNavRef(self: *NavGen, ty: Type, nav_index: InternPool.Nav.Index) !IdRef {
fn constantNavRef(self: *NavGen, ty: Type, nav_index: InternPool.Nav.Index) !Id {
const pt = self.pt;
const zcu = pt.zcu;
const ip = &zcu.intern_pool;
@@ -1240,7 +1238,7 @@ const NavGen = struct {
const decl_ptr_ty_id = try self.ptrType(nav_ty, storage_class, .indirect);
const ptr_id = switch (storage_class) {
.Generic => try self.castToGeneric(decl_ptr_ty_id, decl_id),
.generic => try self.castToGeneric(decl_ptr_ty_id, decl_id),
else => decl_id,
};
@@ -1272,7 +1270,7 @@ const NavGen = struct {
/// The integer type that is returned by this function is the type that is used to perform
/// actual operations (as well as store) a Zig type of a particular number of bits. To create
/// a type with an exact size, use SpvModule.intType.
fn intType(self: *NavGen, signedness: std.builtin.Signedness, bits: u16) !IdRef {
fn intType(self: *NavGen, signedness: std.builtin.Signedness, bits: u16) !Id {
const backing_bits, const big_int = self.backingIntBits(bits);
if (big_int) {
if (backing_bits > 64) {
@@ -1289,12 +1287,12 @@ const NavGen = struct {
};
}
fn arrayType(self: *NavGen, len: u32, child_ty: IdRef) !IdRef {
fn arrayType(self: *NavGen, len: u32, child_ty: Id) !Id {
const len_id = try self.constInt(Type.u32, len);
return self.spv.arrayType(len_id, child_ty);
}
fn ptrType(self: *NavGen, child_ty: Type, storage_class: StorageClass, child_repr: Repr) !IdRef {
fn ptrType(self: *NavGen, child_ty: Type, storage_class: StorageClass, child_repr: Repr) !Id {
const zcu = self.pt.zcu;
const ip = &zcu.intern_pool;
const key = .{ child_ty.toIntern(), storage_class, child_repr };
@@ -1323,7 +1321,7 @@ const NavGen = struct {
.vulkan, .opengl => {
if (child_ty.zigTypeTag(zcu) == .@"struct") {
switch (storage_class) {
.Uniform, .PushConstant => try self.spv.decorate(child_ty_id, .Block),
.uniform, .push_constant => try self.spv.decorate(child_ty_id, .block),
else => {},
}
}
@@ -1331,7 +1329,7 @@ const NavGen = struct {
switch (ip.indexToKey(child_ty.toIntern())) {
.func_type, .opaque_type => {},
else => {
try self.spv.decorate(result_id, .{ .ArrayStride = .{ .array_stride = @intCast(child_ty.abiSize(zcu)) } });
try self.spv.decorate(result_id, .{ .array_stride = .{ .array_stride = @intCast(child_ty.abiSize(zcu)) } });
},
}
},
@@ -1349,9 +1347,9 @@ const NavGen = struct {
return result_id;
}
fn functionType(self: *NavGen, return_ty: Type, param_types: []const Type) !IdRef {
fn functionType(self: *NavGen, return_ty: Type, param_types: []const Type) !Id {
const return_ty_id = try self.resolveFnReturnType(return_ty);
const param_ids = try self.gpa.alloc(IdRef, param_types.len);
const param_ids = try self.gpa.alloc(Id, param_types.len);
defer self.gpa.free(param_ids);
for (param_types, param_ids) |param_ty, *param_id| {
@@ -1379,7 +1377,7 @@ const NavGen = struct {
/// padding: [padding_size]u8,
/// }
/// If any of the fields' size is 0, it will be omitted.
fn resolveUnionType(self: *NavGen, ty: Type) !IdRef {
fn resolveUnionType(self: *NavGen, ty: Type) !Id {
const zcu = self.pt.zcu;
const ip = &zcu.intern_pool;
const union_obj = zcu.typeToUnion(ty).?;
@@ -1394,7 +1392,7 @@ const NavGen = struct {
return try self.resolveType(Type.fromInterned(union_obj.enum_tag_ty), .indirect);
}
var member_types: [4]IdRef = undefined;
var member_types: [4]Id = undefined;
var member_names: [4][]const u8 = undefined;
const u8_ty_id = try self.resolveType(Type.u8, .direct);
@@ -1433,7 +1431,7 @@ const NavGen = struct {
return result_id;
}
fn resolveFnReturnType(self: *NavGen, ret_ty: Type) !IdRef {
fn resolveFnReturnType(self: *NavGen, ret_ty: Type) !Id {
const zcu = self.pt.zcu;
if (!ret_ty.hasRuntimeBitsIgnoreComptime(zcu)) {
// If the return type is an error set or an error union, then we make this
@@ -1450,7 +1448,7 @@ const NavGen = struct {
}
/// Turn a Zig type into a SPIR-V Type, and return a reference to it.
fn resolveType(self: *NavGen, ty: Type, repr: Repr) Error!IdRef {
fn resolveType(self: *NavGen, ty: Type, repr: Repr) Error!Id {
if (self.intern_map.get(.{ ty.toIntern(), repr })) |id| {
return id;
}
@@ -1460,7 +1458,7 @@ const NavGen = struct {
return id;
}
fn resolveTypeInner(self: *NavGen, ty: Type, repr: Repr) Error!IdRef {
fn resolveTypeInner(self: *NavGen, ty: Type, repr: Repr) Error!Id {
const pt = self.pt;
const zcu = pt.zcu;
const ip = &zcu.intern_pool;
@@ -1564,7 +1562,7 @@ const NavGen = struct {
const result_id = try self.arrayType(total_len, elem_ty_id);
switch (self.spv.target.os.tag) {
.vulkan, .opengl => {
try self.spv.decorate(result_id, .{ .ArrayStride = .{
try self.spv.decorate(result_id, .{ .array_stride = .{
.array_stride = @intCast(elem_ty.abiSize(zcu)),
} });
},
@@ -1604,7 +1602,7 @@ const NavGen = struct {
assert(!fn_info.is_var_args);
// Note: Logic is different from functionType().
const param_ty_ids = try self.gpa.alloc(IdRef, fn_info.param_types.len);
const param_ty_ids = try self.gpa.alloc(Id, fn_info.param_types.len);
defer self.gpa.free(param_ty_ids);
var param_index: usize = 0;
for (fn_info.param_types.get(ip)) |param_ty_index| {
@@ -1655,7 +1653,7 @@ const NavGen = struct {
.@"struct" => {
const struct_type = switch (ip.indexToKey(ty.toIntern())) {
.tuple_type => |tuple| {
const member_types = try self.gpa.alloc(IdRef, tuple.values.len);
const member_types = try self.gpa.alloc(Id, tuple.values.len);
defer self.gpa.free(member_types);
var member_index: usize = 0;
@@ -1683,7 +1681,7 @@ const NavGen = struct {
return try self.resolveType(Type.fromInterned(struct_type.backingIntTypeUnordered(ip)), .direct);
}
var member_types = std.ArrayList(IdRef).init(self.gpa);
var member_types = std.ArrayList(Id).init(self.gpa);
defer member_types.deinit();
var member_names = std.ArrayList([]const u8).init(self.gpa);
@@ -1701,7 +1699,7 @@ const NavGen = struct {
switch (self.spv.target.os.tag) {
.vulkan, .opengl => {
try self.spv.decorateMember(result_id, index, .{ .Offset = .{
try self.spv.decorateMember(result_id, index, .{ .offset = .{
.byte_offset = @intCast(ty.structFieldOffset(field_index, zcu)),
} });
},
@@ -1765,7 +1763,7 @@ const NavGen = struct {
const payload_ty_id = try self.resolveType(payload_ty, .indirect);
var member_types: [2]IdRef = undefined;
var member_types: [2]Id = undefined;
var member_names: [2][]const u8 = undefined;
if (eu_layout.error_first) {
// Put the error first
@@ -1809,30 +1807,30 @@ const NavGen = struct {
fn spvStorageClass(self: *NavGen, as: std.builtin.AddressSpace) StorageClass {
return switch (as) {
.generic => if (self.spv.hasFeature(.generic_pointer)) .Generic else .Function,
.generic => if (self.spv.hasFeature(.generic_pointer)) .generic else .function,
.global => switch (self.spv.target.os.tag) {
.opencl, .amdhsa => .CrossWorkgroup,
else => .StorageBuffer,
.opencl, .amdhsa => .cross_workgroup,
else => .storage_buffer,
},
.push_constant => {
return .PushConstant;
return .push_constant;
},
.output => {
return .Output;
return .output;
},
.uniform => {
return .Uniform;
return .uniform;
},
.storage_buffer => {
return .StorageBuffer;
return .storage_buffer;
},
.physical_storage_buffer => {
return .PhysicalStorageBuffer;
return .physical_storage_buffer;
},
.constant => .UniformConstant,
.shared => .Workgroup,
.local => .Function,
.input => .Input,
.constant => .uniform_constant,
.shared => .workgroup,
.local => .function,
.input => .input,
.gs,
.fs,
.ss,
@@ -1980,22 +1978,22 @@ const NavGen = struct {
value: Temporary.Value,
const Value = union(enum) {
singleton: IdResult,
singleton: Id,
exploded_vector: IdRange,
};
fn init(ty: Type, singleton: IdResult) Temporary {
fn init(ty: Type, singleton: Id) Temporary {
return .{ .ty = ty, .value = .{ .singleton = singleton } };
}
fn materialize(self: Temporary, ng: *NavGen) !IdResult {
fn materialize(self: Temporary, ng: *NavGen) !Id {
const zcu = ng.pt.zcu;
switch (self.value) {
.singleton => |id| return id,
.exploded_vector => |range| {
assert(self.ty.isVector(zcu));
assert(self.ty.vectorLen(zcu) == range.len);
const constituents = try ng.gpa.alloc(IdRef, range.len);
const constituents = try ng.gpa.alloc(Id, range.len);
defer ng.gpa.free(constituents);
for (constituents, 0..range.len) |*id, i| {
id.* = range.at(i);
@@ -2170,16 +2168,16 @@ const NavGen = struct {
/// on the operation and input value.
const Value = union(enum) {
/// A single scalar value that is used by a scalar operation.
scalar: IdResult,
scalar: Id,
/// A single scalar that is broadcasted in an unrolled operation.
scalar_broadcast: IdResult,
scalar_broadcast: Id,
/// A vector represented by a consecutive list of IDs that is used in an unrolled operation.
vector_exploded: IdRange,
};
/// Query the value at a particular index of the operation. Note that
/// the index is *not* the component/lane, but the index of the *operation*.
fn at(self: PreparedOperand, i: usize) IdResult {
fn at(self: PreparedOperand, i: usize) Id {
switch (self.value) {
.scalar => |id| {
assert(i == 0);
@@ -2253,9 +2251,9 @@ const NavGen = struct {
for (0..ops) |i| {
try self.func.body.emitRaw(self.spv.gpa, opcode, 3);
self.func.body.writeOperand(spec.IdResultType, op_result_ty_id);
self.func.body.writeOperand(IdResult, results.at(i));
self.func.body.writeOperand(IdResult, op_src.at(i));
self.func.body.writeOperand(spec.Id, op_result_ty_id);
self.func.body.writeOperand(Id, results.at(i));
self.func.body.writeOperand(Id, op_src.at(i));
}
return v.finalize(result_ty, results);
@@ -2388,10 +2386,10 @@ const NavGen = struct {
for (0..ops) |i| {
try self.func.body.emitRaw(self.spv.gpa, opcode, 4);
self.func.body.writeOperand(spec.IdResultType, op_result_ty_id);
self.func.body.writeOperand(IdResult, results.at(i));
self.func.body.writeOperand(IdResult, op_lhs.at(i));
self.func.body.writeOperand(IdResult, op_rhs.at(i));
self.func.body.writeOperand(spec.Id, op_result_ty_id);
self.func.body.writeOperand(Id, results.at(i));
self.func.body.writeOperand(Id, op_lhs.at(i));
self.func.body.writeOperand(Id, op_rhs.at(i));
}
return v.finalize(result_ty, results);
@@ -2442,9 +2440,9 @@ const NavGen = struct {
}) |opcode| {
for (0..ops) |i| {
try self.func.body.emitRaw(self.spv.gpa, opcode, 3);
self.func.body.writeOperand(spec.IdResultType, op_result_ty_id);
self.func.body.writeOperand(IdResult, results.at(i));
self.func.body.writeOperand(IdResult, op_operand.at(i));
self.func.body.writeOperand(spec.Id, op_result_ty_id);
self.func.body.writeOperand(Id, results.at(i));
self.func.body.writeOperand(Id, op_operand.at(i));
}
} else {
const set = try self.importExtendedSet();
@@ -2583,10 +2581,10 @@ const NavGen = struct {
}) |opcode| {
for (0..ops) |i| {
try self.func.body.emitRaw(self.spv.gpa, opcode, 4);
self.func.body.writeOperand(spec.IdResultType, op_result_ty_id);
self.func.body.writeOperand(IdResult, results.at(i));
self.func.body.writeOperand(IdResult, op_lhs.at(i));
self.func.body.writeOperand(IdResult, op_rhs.at(i));
self.func.body.writeOperand(spec.Id, op_result_ty_id);
self.func.body.writeOperand(Id, results.at(i));
self.func.body.writeOperand(Id, op_lhs.at(i));
self.func.body.writeOperand(Id, op_rhs.at(i));
}
} else {
const set = try self.importExtendedSet();
@@ -2702,10 +2700,10 @@ const NavGen = struct {
const op_result = self.spv.allocId();
try self.func.body.emitRaw(self.spv.gpa, opcode, 4);
self.func.body.writeOperand(spec.IdResultType, op_result_ty_id);
self.func.body.writeOperand(IdResult, op_result);
self.func.body.writeOperand(IdResult, lhs_op.at(i));
self.func.body.writeOperand(IdResult, rhs_op.at(i));
self.func.body.writeOperand(spec.Id, op_result_ty_id);
self.func.body.writeOperand(Id, op_result);
self.func.body.writeOperand(Id, lhs_op.at(i));
self.func.body.writeOperand(Id, rhs_op.at(i));
// The above operation returns a struct. We might want to expand
// Temporary to deal with the fact that these are structs eventually,
@@ -2804,8 +2802,8 @@ const NavGen = struct {
const buffer_struct_ty_id = self.spv.allocId();
try self.spv.structType(buffer_struct_ty_id, &.{anyerror_ty_id}, &.{"error_out"});
try self.spv.decorate(buffer_struct_ty_id, .Block);
try self.spv.decorateMember(buffer_struct_ty_id, 0, .{ .Offset = .{ .byte_offset = 0 } });
try self.spv.decorate(buffer_struct_ty_id, .block);
try self.spv.decorateMember(buffer_struct_ty_id, 0, .{ .offset = .{ .byte_offset = 0 } });
const ptr_buffer_struct_ty_id = self.spv.allocId();
try self.spv.sections.types_globals_constants.emit(self.spv.gpa, .OpTypePointer, .{
@@ -2820,15 +2818,15 @@ const NavGen = struct {
.id_result = buffer_struct_id,
.storage_class = self.spvStorageClass(.global),
});
try self.spv.decorate(buffer_struct_id, .{ .DescriptorSet = .{ .descriptor_set = 0 } });
try self.spv.decorate(buffer_struct_id, .{ .Binding = .{ .binding_point = 0 } });
try self.spv.decorate(buffer_struct_id, .{ .descriptor_set = .{ .descriptor_set = 0 } });
try self.spv.decorate(buffer_struct_id, .{ .binding = .{ .binding_point = 0 } });
self.object.error_buffer = spv_err_decl_index;
}
try self.spv.sections.execution_modes.emit(self.spv.gpa, .OpExecutionMode, .{
.entry_point = kernel_id,
.mode = .{ .LocalSize = .{
.mode = .{ .local_size = .{
.x_size = 1,
.y_size = 1,
.z_size = 1,
@@ -2873,7 +2871,7 @@ const NavGen = struct {
.pointer = p_error_id,
.object = error_id,
.memory_access = .{
.Aligned = .{ .literal_integer = @intCast(Type.abiAlignment(.anyerror, zcu).toByteUnits().?) },
.aligned = .{ .literal_integer = @intCast(Type.abiAlignment(.anyerror, zcu).toByteUnits().?) },
},
});
try section.emit(self.spv.gpa, .OpReturn, {});
@@ -2885,8 +2883,8 @@ const NavGen = struct {
defer self.gpa.free(test_name);
const execution_mode: spec.ExecutionModel = switch (target.os.tag) {
.vulkan, .opengl => .GLCompute,
.opencl, .amdhsa => .Kernel,
.vulkan, .opengl => .gl_compute,
.opencl, .amdhsa => .kernel,
else => unreachable,
};
@@ -2983,7 +2981,7 @@ const NavGen = struct {
assert(maybe_init_val == null); // TODO
const storage_class = self.spvStorageClass(nav.getAddrspace());
assert(storage_class != .Generic); // These should be instance globals
assert(storage_class != .generic); // These should be instance globals
const ptr_ty_id = try self.ptrType(ty, storage_class, .indirect);
@@ -2993,38 +2991,8 @@ const NavGen = struct {
.storage_class = storage_class,
});
if (nav.fqn.eqlSlice("position", ip)) {
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("invocation_id", ip)) {
try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .InvocationId } });
} else if (nav.fqn.eqlSlice("frag_coord", ip)) {
try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .FragCoord } });
} else if (nav.fqn.eqlSlice("point_coord", ip)) {
try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .PointCoord } });
} else if (nav.fqn.eqlSlice("front_facing", ip)) {
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("frag_depth", ip)) {
try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .FragDepth } });
} else if (nav.fqn.eqlSlice("num_workgroups", ip)) {
try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .NumWorkgroups } });
} else if (nav.fqn.eqlSlice("workgroup_size", ip)) {
try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .WorkgroupSize } });
} else if (nav.fqn.eqlSlice("workgroup_id", ip)) {
try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .WorkgroupId } });
} else if (nav.fqn.eqlSlice("local_invocation_id", ip)) {
try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .LocalInvocationId } });
} else if (nav.fqn.eqlSlice("global_invocation_id", ip)) {
try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .GlobalInvocationId } });
} else if (nav.fqn.eqlSlice("local_invocation_index", ip)) {
try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .LocalInvocationIndex } });
} else if (nav.fqn.eqlSlice("vertex_index", ip)) {
try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .VertexIndex } });
} else if (nav.fqn.eqlSlice("instance_index", ip)) {
try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .InstanceIndex } });
if (std.meta.stringToEnum(spec.BuiltIn, nav.fqn.toSlice(ip))) |builtin| {
try self.spv.decorate(result_id, .{ .built_in = .{ .built_in = builtin } });
}
try self.spv.debugName(result_id, nav.fqn.toSlice(ip));
@@ -3040,7 +3008,7 @@ const NavGen = struct {
try self.spv.declareDeclDeps(spv_decl_index, &.{});
const ptr_ty_id = try self.ptrType(ty, .Function, .indirect);
const ptr_ty_id = try self.ptrType(ty, .function, .indirect);
if (maybe_init_val) |init_val| {
// TODO: Combine with resolveAnonDecl?
@@ -3109,7 +3077,7 @@ const NavGen = struct {
/// Convert representation from indirect (in memory) to direct (in 'register')
/// This converts the argument type from resolveType(ty, .indirect) to resolveType(ty, .direct).
fn convertToDirect(self: *NavGen, ty: Type, operand_id: IdRef) !IdRef {
fn convertToDirect(self: *NavGen, ty: Type, operand_id: Id) !Id {
const pt = self.pt;
const zcu = pt.zcu;
switch (ty.scalarType(zcu).zigTypeTag(zcu)) {
@@ -3136,7 +3104,7 @@ const NavGen = struct {
/// Convert representation from direct (in 'register) to direct (in memory)
/// This converts the argument type from resolveType(ty, .direct) to resolveType(ty, .indirect).
fn convertToIndirect(self: *NavGen, ty: Type, operand_id: IdRef) !IdRef {
fn convertToIndirect(self: *NavGen, ty: Type, operand_id: Id) !Id {
const zcu = self.pt.zcu;
switch (ty.scalarType(zcu).zigTypeTag(zcu)) {
.bool => {
@@ -3147,7 +3115,7 @@ const NavGen = struct {
}
}
fn extractField(self: *NavGen, result_ty: Type, object: IdRef, field: u32) !IdRef {
fn extractField(self: *NavGen, result_ty: Type, object: Id, field: u32) !Id {
const result_ty_id = try self.resolveType(result_ty, .indirect);
const result_id = self.spv.allocId();
const indexes = [_]u32{field};
@@ -3161,7 +3129,7 @@ const NavGen = struct {
return try self.convertToDirect(result_ty, result_id);
}
fn extractVectorComponent(self: *NavGen, result_ty: Type, vector_id: IdRef, field: u32) !IdRef {
fn extractVectorComponent(self: *NavGen, result_ty: Type, vector_id: Id, field: u32) !Id {
const result_ty_id = try self.resolveType(result_ty, .direct);
const result_id = self.spv.allocId();
const indexes = [_]u32{field};
@@ -3179,14 +3147,14 @@ const NavGen = struct {
is_volatile: bool = false,
};
fn load(self: *NavGen, value_ty: Type, ptr_id: IdRef, options: MemoryOptions) !IdRef {
fn load(self: *NavGen, value_ty: Type, ptr_id: Id, options: MemoryOptions) !Id {
const zcu = self.pt.zcu;
const alignment: u32 = @intCast(value_ty.abiAlignment(zcu).toByteUnits().?);
const indirect_value_ty_id = try self.resolveType(value_ty, .indirect);
const result_id = self.spv.allocId();
const access = spec.MemoryAccess.Extended{
.Volatile = options.is_volatile,
.Aligned = .{ .literal_integer = alignment },
const access: spec.MemoryAccess.Extended = .{
.@"volatile" = options.is_volatile,
.aligned = .{ .literal_integer = alignment },
};
try self.func.body.emit(self.spv.gpa, .OpLoad, .{
.id_result_type = indirect_value_ty_id,
@@ -3197,11 +3165,9 @@ const NavGen = struct {
return try self.convertToDirect(value_ty, result_id);
}
fn store(self: *NavGen, value_ty: Type, ptr_id: IdRef, value_id: IdRef, options: MemoryOptions) !void {
fn store(self: *NavGen, value_ty: Type, ptr_id: Id, value_id: Id, options: MemoryOptions) !void {
const indirect_value_id = try self.convertToIndirect(value_ty, value_id);
const access = spec.MemoryAccess.Extended{
.Volatile = options.is_volatile,
};
const access: spec.MemoryAccess.Extended = .{ .@"volatile" = options.is_volatile };
try self.func.body.emit(self.spv.gpa, .OpStore, .{
.pointer = ptr_id,
.object = indirect_value_id,
@@ -3222,7 +3188,7 @@ const NavGen = struct {
return;
const air_tags = self.air.instructions.items(.tag);
const maybe_result_id: ?IdRef = switch (air_tags[@intFromEnum(inst)]) {
const maybe_result_id: ?Id = switch (air_tags[@intFromEnum(inst)]) {
// zig fmt: off
.add, .add_wrap, .add_optimized => try self.airArithOp(inst, .f_add, .i_add, .i_add),
.sub, .sub_wrap, .sub_optimized => try self.airArithOp(inst, .f_sub, .i_sub, .i_sub),
@@ -3390,7 +3356,7 @@ const NavGen = struct {
try self.inst_results.putNoClobber(self.gpa, inst, result_id);
}
fn airBinOpSimple(self: *NavGen, inst: Air.Inst.Index, op: BinaryOp) !?IdRef {
fn airBinOpSimple(self: *NavGen, inst: Air.Inst.Index, op: BinaryOp) !?Id {
const bin_op = self.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
const lhs = try self.temporary(bin_op.lhs);
const rhs = try self.temporary(bin_op.rhs);
@@ -3399,7 +3365,7 @@ const NavGen = struct {
return try result.materialize(self);
}
fn airShift(self: *NavGen, inst: Air.Inst.Index, unsigned: BinaryOp, signed: BinaryOp) !?IdRef {
fn airShift(self: *NavGen, inst: Air.Inst.Index, unsigned: BinaryOp, signed: BinaryOp) !?Id {
const zcu = self.pt.zcu;
const bin_op = self.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
@@ -3438,7 +3404,7 @@ const NavGen = struct {
const MinMax = enum { min, max };
fn airMinMax(self: *NavGen, inst: Air.Inst.Index, op: MinMax) !?IdRef {
fn airMinMax(self: *NavGen, inst: Air.Inst.Index, op: MinMax) !?Id {
const bin_op = self.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
const lhs = try self.temporary(bin_op.lhs);
@@ -3503,7 +3469,7 @@ const NavGen = struct {
}
}
fn airDivFloor(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airDivFloor(self: *NavGen, inst: Air.Inst.Index) !?Id {
const bin_op = self.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
const lhs = try self.temporary(bin_op.lhs);
@@ -3560,7 +3526,7 @@ const NavGen = struct {
}
}
fn airDivTrunc(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airDivTrunc(self: *NavGen, inst: Air.Inst.Index) !?Id {
const bin_op = self.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
const lhs = try self.temporary(bin_op.lhs);
@@ -3588,7 +3554,7 @@ const NavGen = struct {
}
}
fn airUnOpSimple(self: *NavGen, inst: Air.Inst.Index, op: UnaryOp) !?IdRef {
fn airUnOpSimple(self: *NavGen, inst: Air.Inst.Index, op: UnaryOp) !?Id {
const un_op = self.air.instructions.items(.data)[@intFromEnum(inst)].un_op;
const operand = try self.temporary(un_op);
const result = try self.buildUnary(op, operand);
@@ -3601,7 +3567,7 @@ const NavGen = struct {
comptime fop: BinaryOp,
comptime sop: BinaryOp,
comptime uop: BinaryOp,
) !?IdRef {
) !?Id {
const bin_op = self.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
const lhs = try self.temporary(bin_op.lhs);
@@ -3622,7 +3588,7 @@ const NavGen = struct {
return try result.materialize(self);
}
fn airAbs(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airAbs(self: *NavGen, inst: Air.Inst.Index) !?Id {
const ty_op = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
const operand = try self.temporary(ty_op.operand);
// Note: operand_ty may be signed, while ty is always unsigned!
@@ -3662,7 +3628,7 @@ const NavGen = struct {
comptime add: BinaryOp,
comptime ucmp: CmpPredicate,
comptime scmp: CmpPredicate,
) !?IdRef {
) !?Id {
_ = scmp;
// Note: OpIAddCarry and OpISubBorrow are not really useful here: For unsigned numbers,
// there is in both cases only one extra operation required. For signed operations,
@@ -3725,7 +3691,7 @@ const NavGen = struct {
return try self.constructComposite(result_ty_id, &.{ try result.materialize(self), try ov.materialize(self) });
}
fn airMulOverflow(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airMulOverflow(self: *NavGen, inst: Air.Inst.Index) !?Id {
const pt = self.pt;
const ty_pl = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
@@ -3897,7 +3863,7 @@ const NavGen = struct {
return try self.constructComposite(result_ty_id, &.{ try result.materialize(self), try ov.materialize(self) });
}
fn airShlOverflow(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airShlOverflow(self: *NavGen, inst: Air.Inst.Index) !?Id {
const zcu = self.pt.zcu;
const ty_pl = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
@@ -3938,7 +3904,7 @@ const NavGen = struct {
return try self.constructComposite(result_ty_id, &.{ try result.materialize(self), try ov.materialize(self) });
}
fn airMulAdd(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airMulAdd(self: *NavGen, inst: Air.Inst.Index) !?Id {
const pl_op = self.air.instructions.items(.data)[@intFromEnum(inst)].pl_op;
const extra = self.air.extraData(Air.Bin, pl_op.payload).data;
@@ -3954,7 +3920,7 @@ const NavGen = struct {
return try result.materialize(self);
}
fn airClzCtz(self: *NavGen, inst: Air.Inst.Index, op: UnaryOp) !?IdRef {
fn airClzCtz(self: *NavGen, inst: Air.Inst.Index, op: UnaryOp) !?Id {
if (self.liveness.isUnused(inst)) return null;
const zcu = self.pt.zcu;
@@ -3979,7 +3945,7 @@ const NavGen = struct {
return try result.materialize(self);
}
fn airSelect(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airSelect(self: *NavGen, inst: Air.Inst.Index) !?Id {
const pl_op = self.air.instructions.items(.data)[@intFromEnum(inst)].pl_op;
const extra = self.air.extraData(Air.Bin, pl_op.payload).data;
const pred = try self.temporary(pl_op.operand);
@@ -3990,7 +3956,7 @@ const NavGen = struct {
return try result.materialize(self);
}
fn airSplat(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airSplat(self: *NavGen, inst: Air.Inst.Index) !?Id {
const ty_op = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
const operand_id = try self.resolve(ty_op.operand);
@@ -3999,7 +3965,7 @@ const NavGen = struct {
return try self.constructCompositeSplat(result_ty, operand_id);
}
fn airReduce(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airReduce(self: *NavGen, inst: Air.Inst.Index) !?Id {
const zcu = self.pt.zcu;
const reduce = self.air.instructions.items(.data)[@intFromEnum(inst)].reduce;
const operand = try self.resolve(reduce.operand);
@@ -4062,16 +4028,16 @@ const NavGen = struct {
result_id = self.spv.allocId();
try self.func.body.emitRaw(self.spv.gpa, opcode, 4);
self.func.body.writeOperand(spec.IdResultType, scalar_ty_id);
self.func.body.writeOperand(spec.IdResult, result_id);
self.func.body.writeOperand(spec.IdResultType, lhs);
self.func.body.writeOperand(spec.IdResultType, rhs);
self.func.body.writeOperand(spec.Id, scalar_ty_id);
self.func.body.writeOperand(spec.Id, result_id);
self.func.body.writeOperand(spec.Id, lhs);
self.func.body.writeOperand(spec.Id, rhs);
}
return result_id;
}
fn airShuffleOne(ng: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airShuffleOne(ng: *NavGen, inst: Air.Inst.Index) !?Id {
const pt = ng.pt;
const zcu = pt.zcu;
const gpa = zcu.gpa;
@@ -4082,7 +4048,7 @@ const NavGen = struct {
const elem_ty = result_ty.childType(zcu);
const operand = try ng.resolve(unwrapped.operand);
const constituents = try gpa.alloc(IdRef, mask.len);
const constituents = try gpa.alloc(Id, mask.len);
defer gpa.free(constituents);
for (constituents, mask) |*id, mask_elem| {
@@ -4096,7 +4062,7 @@ const NavGen = struct {
return try ng.constructComposite(result_ty_id, constituents);
}
fn airShuffleTwo(ng: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airShuffleTwo(ng: *NavGen, inst: Air.Inst.Index) !?Id {
const pt = ng.pt;
const zcu = pt.zcu;
const gpa = zcu.gpa;
@@ -4109,7 +4075,7 @@ const NavGen = struct {
const operand_a = try ng.resolve(unwrapped.operand_a);
const operand_b = try ng.resolve(unwrapped.operand_b);
const constituents = try gpa.alloc(IdRef, mask.len);
const constituents = try gpa.alloc(Id, mask.len);
defer gpa.free(constituents);
for (constituents, mask) |*id, mask_elem| {
@@ -4124,8 +4090,8 @@ const NavGen = struct {
return try ng.constructComposite(result_ty_id, constituents);
}
fn indicesToIds(self: *NavGen, indices: []const u32) ![]IdRef {
const ids = try self.gpa.alloc(IdRef, indices.len);
fn indicesToIds(self: *NavGen, indices: []const u32) ![]Id {
const ids = try self.gpa.alloc(Id, indices.len);
errdefer self.gpa.free(ids);
for (indices, ids) |index, *id| {
id.* = try self.constInt(Type.u32, index);
@@ -4136,10 +4102,10 @@ const NavGen = struct {
fn accessChainId(
self: *NavGen,
result_ty_id: IdRef,
base: IdRef,
indices: []const IdRef,
) !IdRef {
result_ty_id: Id,
base: Id,
indices: []const Id,
) !Id {
const result_id = self.spv.allocId();
try self.func.body.emit(self.spv.gpa, .OpInBoundsAccessChain, .{
.id_result_type = result_ty_id,
@@ -4156,10 +4122,10 @@ const NavGen = struct {
/// is the latter and PtrAccessChain is the former.
fn accessChain(
self: *NavGen,
result_ty_id: IdRef,
base: IdRef,
result_ty_id: Id,
base: Id,
indices: []const u32,
) !IdRef {
) !Id {
const ids = try self.indicesToIds(indices);
defer self.gpa.free(ids);
return try self.accessChainId(result_ty_id, base, ids);
@@ -4167,11 +4133,11 @@ const NavGen = struct {
fn ptrAccessChain(
self: *NavGen,
result_ty_id: IdRef,
base: IdRef,
element: IdRef,
result_ty_id: Id,
base: Id,
element: Id,
indices: []const u32,
) !IdRef {
) !Id {
const ids = try self.indicesToIds(indices);
defer self.gpa.free(ids);
@@ -4199,7 +4165,7 @@ const NavGen = struct {
return result_id;
}
fn ptrAdd(self: *NavGen, result_ty: Type, ptr_ty: Type, ptr_id: IdRef, offset_id: IdRef) !IdRef {
fn ptrAdd(self: *NavGen, result_ty: Type, ptr_ty: Type, ptr_id: Id, offset_id: Id) !Id {
const zcu = self.pt.zcu;
const result_ty_id = try self.resolveType(result_ty, .direct);
@@ -4220,7 +4186,7 @@ const NavGen = struct {
}
}
fn airPtrAdd(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airPtrAdd(self: *NavGen, inst: Air.Inst.Index) !?Id {
const ty_pl = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
const bin_op = self.air.extraData(Air.Bin, ty_pl.payload).data;
const ptr_id = try self.resolve(bin_op.lhs);
@@ -4231,7 +4197,7 @@ const NavGen = struct {
return try self.ptrAdd(result_ty, ptr_ty, ptr_id, offset_id);
}
fn airPtrSub(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airPtrSub(self: *NavGen, inst: Air.Inst.Index) !?Id {
const ty_pl = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
const bin_op = self.air.extraData(Air.Bin, ty_pl.payload).data;
const ptr_id = try self.resolve(bin_op.lhs);
@@ -4427,7 +4393,7 @@ const NavGen = struct {
self: *NavGen,
inst: Air.Inst.Index,
comptime op: std.math.CompareOperator,
) !?IdRef {
) !?Id {
const bin_op = self.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
const lhs = try self.temporary(bin_op.lhs);
const rhs = try self.temporary(bin_op.rhs);
@@ -4436,7 +4402,7 @@ const NavGen = struct {
return try result.materialize(self);
}
fn airVectorCmp(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airVectorCmp(self: *NavGen, inst: Air.Inst.Index) !?Id {
const ty_pl = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
const vec_cmp = self.air.extraData(Air.VectorCmp, ty_pl.payload).data;
const lhs = try self.temporary(vec_cmp.lhs);
@@ -4452,8 +4418,8 @@ const NavGen = struct {
self: *NavGen,
dst_ty: Type,
src_ty: Type,
src_id: IdRef,
) !IdRef {
src_id: Id,
) !Id {
const zcu = self.pt.zcu;
const src_ty_id = try self.resolveType(src_ty, .direct);
const dst_ty_id = try self.resolveType(dst_ty, .direct);
@@ -4489,9 +4455,9 @@ const NavGen = struct {
break :blk result_id;
}
const dst_ptr_ty_id = try self.ptrType(dst_ty, .Function, .indirect);
const dst_ptr_ty_id = try self.ptrType(dst_ty, .function, .indirect);
const tmp_id = try self.alloc(src_ty, .{ .storage_class = .Function });
const tmp_id = try self.alloc(src_ty, .{ .storage_class = .function });
try self.store(src_ty, tmp_id, src_id, .{});
const casted_ptr_id = self.spv.allocId();
try self.func.body.emit(self.spv.gpa, .OpBitcast, .{
@@ -4515,7 +4481,7 @@ const NavGen = struct {
return result_id;
}
fn airBitCast(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airBitCast(self: *NavGen, inst: Air.Inst.Index) !?Id {
const ty_op = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
const operand_ty = self.typeOf(ty_op.operand);
const result_ty = self.typeOfIndex(inst);
@@ -4528,7 +4494,7 @@ const NavGen = struct {
return try self.bitCast(result_ty, operand_ty, operand_id);
}
fn airIntCast(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airIntCast(self: *NavGen, inst: Air.Inst.Index) !?Id {
const ty_op = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
const src = try self.temporary(ty_op.operand);
const dst_ty = self.typeOfIndex(inst);
@@ -4554,7 +4520,7 @@ const NavGen = struct {
return try result.materialize(self);
}
fn intFromPtr(self: *NavGen, operand_id: IdRef) !IdRef {
fn intFromPtr(self: *NavGen, operand_id: Id) !Id {
const result_type_id = try self.resolveType(Type.usize, .direct);
const result_id = self.spv.allocId();
try self.func.body.emit(self.spv.gpa, .OpConvertPtrToU, .{
@@ -4565,7 +4531,7 @@ const NavGen = struct {
return result_id;
}
fn airFloatFromInt(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airFloatFromInt(self: *NavGen, inst: Air.Inst.Index) !?Id {
const ty_op = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
const operand_ty = self.typeOf(ty_op.operand);
const operand_id = try self.resolve(ty_op.operand);
@@ -4573,7 +4539,7 @@ const NavGen = struct {
return try self.floatFromInt(result_ty, operand_ty, operand_id);
}
fn floatFromInt(self: *NavGen, result_ty: Type, operand_ty: Type, operand_id: IdRef) !IdRef {
fn floatFromInt(self: *NavGen, result_ty: Type, operand_ty: Type, operand_id: Id) !Id {
const operand_info = self.arithmeticTypeInfo(operand_ty);
const result_id = self.spv.allocId();
const result_ty_id = try self.resolveType(result_ty, .direct);
@@ -4592,14 +4558,14 @@ const NavGen = struct {
return result_id;
}
fn airIntFromFloat(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airIntFromFloat(self: *NavGen, inst: Air.Inst.Index) !?Id {
const ty_op = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
const operand_id = try self.resolve(ty_op.operand);
const result_ty = self.typeOfIndex(inst);
return try self.intFromFloat(result_ty, operand_id);
}
fn intFromFloat(self: *NavGen, result_ty: Type, operand_id: IdRef) !IdRef {
fn intFromFloat(self: *NavGen, result_ty: Type, operand_id: Id) !Id {
const result_info = self.arithmeticTypeInfo(result_ty);
const result_ty_id = try self.resolveType(result_ty, .direct);
const result_id = self.spv.allocId();
@@ -4618,7 +4584,7 @@ const NavGen = struct {
return result_id;
}
fn airFloatCast(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airFloatCast(self: *NavGen, inst: Air.Inst.Index) !?Id {
const ty_op = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
const operand = try self.temporary(ty_op.operand);
const dest_ty = self.typeOfIndex(inst);
@@ -4626,7 +4592,7 @@ const NavGen = struct {
return try result.materialize(self);
}
fn airNot(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airNot(self: *NavGen, inst: Air.Inst.Index) !?Id {
const ty_op = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
const operand = try self.temporary(ty_op.operand);
const result_ty = self.typeOfIndex(inst);
@@ -4645,7 +4611,7 @@ const NavGen = struct {
return try result.materialize(self);
}
fn airArrayToSlice(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airArrayToSlice(self: *NavGen, inst: Air.Inst.Index) !?Id {
const pt = self.pt;
const zcu = pt.zcu;
const ty_op = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
@@ -4670,7 +4636,7 @@ const NavGen = struct {
return try self.constructComposite(slice_ty_id, &.{ elem_ptr_id, len_id });
}
fn airSlice(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airSlice(self: *NavGen, inst: Air.Inst.Index) !?Id {
const ty_pl = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
const bin_op = self.air.extraData(Air.Bin, ty_pl.payload).data;
const ptr_id = try self.resolve(bin_op.lhs);
@@ -4680,7 +4646,7 @@ const NavGen = struct {
return try self.constructComposite(slice_ty_id, &.{ ptr_id, len_id });
}
fn airAggregateInit(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airAggregateInit(self: *NavGen, inst: Air.Inst.Index) !?Id {
const pt = self.pt;
const zcu = pt.zcu;
const ip = &zcu.intern_pool;
@@ -4732,7 +4698,7 @@ const NavGen = struct {
const types = try self.gpa.alloc(Type, elements.len);
defer self.gpa.free(types);
const constituents = try self.gpa.alloc(IdRef, elements.len);
const constituents = try self.gpa.alloc(Id, elements.len);
defer self.gpa.free(constituents);
var index: usize = 0;
@@ -4771,7 +4737,7 @@ const NavGen = struct {
},
.vector => {
const n_elems = result_ty.vectorLen(zcu);
const elem_ids = try self.gpa.alloc(IdRef, n_elems);
const elem_ids = try self.gpa.alloc(Id, n_elems);
defer self.gpa.free(elem_ids);
for (elements, 0..) |element, i| {
@@ -4784,7 +4750,7 @@ const NavGen = struct {
.array => {
const array_info = result_ty.arrayInfo(zcu);
const n_elems: usize = @intCast(result_ty.arrayLenIncludingSentinel(zcu));
const elem_ids = try self.gpa.alloc(IdRef, n_elems);
const elem_ids = try self.gpa.alloc(Id, n_elems);
defer self.gpa.free(elem_ids);
for (elements, 0..) |element, i| {
@@ -4803,7 +4769,7 @@ const NavGen = struct {
}
}
fn sliceOrArrayLen(self: *NavGen, operand_id: IdRef, ty: Type) !IdRef {
fn sliceOrArrayLen(self: *NavGen, operand_id: Id, ty: Type) !Id {
const pt = self.pt;
const zcu = pt.zcu;
switch (ty.ptrSize(zcu)) {
@@ -4819,7 +4785,7 @@ const NavGen = struct {
}
}
fn sliceOrArrayPtr(self: *NavGen, operand_id: IdRef, ty: Type) !IdRef {
fn sliceOrArrayPtr(self: *NavGen, operand_id: Id, ty: Type) !Id {
const zcu = self.pt.zcu;
if (ty.isSlice(zcu)) {
const ptr_ty = ty.slicePtrFieldType(zcu);
@@ -4849,14 +4815,14 @@ const NavGen = struct {
return self.fail("TODO implement airMemcpy for spirv", .{});
}
fn airSliceField(self: *NavGen, inst: Air.Inst.Index, field: u32) !?IdRef {
fn airSliceField(self: *NavGen, inst: Air.Inst.Index, field: u32) !?Id {
const ty_op = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
const field_ty = self.typeOfIndex(inst);
const operand_id = try self.resolve(ty_op.operand);
return try self.extractField(field_ty, operand_id, field);
}
fn airSliceElemPtr(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airSliceElemPtr(self: *NavGen, inst: Air.Inst.Index) !?Id {
const zcu = self.pt.zcu;
const ty_pl = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
const bin_op = self.air.extraData(Air.Bin, ty_pl.payload).data;
@@ -4873,7 +4839,7 @@ const NavGen = struct {
return try self.ptrAccessChain(ptr_ty_id, slice_ptr, index_id, &.{});
}
fn airSliceElemVal(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airSliceElemVal(self: *NavGen, inst: Air.Inst.Index) !?Id {
const zcu = self.pt.zcu;
const bin_op = self.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
const slice_ty = self.typeOf(bin_op.lhs);
@@ -4890,7 +4856,7 @@ const NavGen = struct {
return try self.load(slice_ty.childType(zcu), elem_ptr, .{ .is_volatile = slice_ty.isVolatilePtr(zcu) });
}
fn ptrElemPtr(self: *NavGen, ptr_ty: Type, ptr_id: IdRef, index_id: IdRef) !IdRef {
fn ptrElemPtr(self: *NavGen, ptr_ty: Type, ptr_id: Id, index_id: Id) !Id {
const zcu = self.pt.zcu;
// Construct new pointer type for the resulting pointer
const elem_ty = ptr_ty.elemType2(zcu); // use elemType() so that we get T for *[N]T.
@@ -4905,7 +4871,7 @@ const NavGen = struct {
}
}
fn airPtrElemPtr(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airPtrElemPtr(self: *NavGen, inst: Air.Inst.Index) !?Id {
const pt = self.pt;
const zcu = pt.zcu;
const ty_pl = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
@@ -4923,7 +4889,7 @@ const NavGen = struct {
return try self.ptrElemPtr(src_ptr_ty, ptr_id, index_id);
}
fn airArrayElemVal(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airArrayElemVal(self: *NavGen, inst: Air.Inst.Index) !?Id {
const zcu = self.pt.zcu;
const bin_op = self.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
const array_ty = self.typeOf(bin_op.lhs);
@@ -4938,14 +4904,14 @@ const NavGen = struct {
const is_vector = array_ty.isVector(zcu);
const elem_repr: Repr = if (is_vector) .direct else .indirect;
const ptr_array_ty_id = try self.ptrType(array_ty, .Function, .direct);
const ptr_elem_ty_id = try self.ptrType(elem_ty, .Function, elem_repr);
const ptr_array_ty_id = try self.ptrType(array_ty, .function, .direct);
const ptr_elem_ty_id = try self.ptrType(elem_ty, .function, elem_repr);
const tmp_id = self.spv.allocId();
try self.func.prologue.emit(self.spv.gpa, .OpVariable, .{
.id_result_type = ptr_array_ty_id,
.id_result = tmp_id,
.storage_class = .Function,
.storage_class = .function,
});
try self.func.body.emit(self.spv.gpa, .OpStore, .{
@@ -4973,7 +4939,7 @@ const NavGen = struct {
return try self.convertToDirect(elem_ty, result_id);
}
fn airPtrElemVal(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airPtrElemVal(self: *NavGen, inst: Air.Inst.Index) !?Id {
const zcu = self.pt.zcu;
const bin_op = self.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
const ptr_ty = self.typeOf(bin_op.lhs);
@@ -5029,7 +4995,7 @@ const NavGen = struct {
}
}
fn airGetUnionTag(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airGetUnionTag(self: *NavGen, inst: Air.Inst.Index) !?Id {
const ty_op = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
const un_ty = self.typeOf(ty_op.operand);
@@ -5048,8 +5014,8 @@ const NavGen = struct {
self: *NavGen,
ty: Type,
active_field: u32,
payload: ?IdRef,
) !IdRef {
payload: ?Id,
) !Id {
// To initialize a union, generate a temporary variable with the
// union type, then get the field pointer and pointer-cast it to the
// right type to store it. Finally load the entire union.
@@ -5100,20 +5066,20 @@ const NavGen = struct {
return try self.constInt(tag_ty, tag_int);
}
const tmp_id = try self.alloc(ty, .{ .storage_class = .Function });
const tmp_id = try self.alloc(ty, .{ .storage_class = .function });
if (layout.tag_size != 0) {
const tag_ptr_ty_id = try self.ptrType(tag_ty, .Function, .indirect);
const tag_ptr_ty_id = try self.ptrType(tag_ty, .function, .indirect);
const ptr_id = try self.accessChain(tag_ptr_ty_id, tmp_id, &.{@as(u32, @intCast(layout.tag_index))});
const tag_id = try self.constInt(tag_ty, tag_int);
try self.store(tag_ty, ptr_id, tag_id, .{});
}
if (payload_ty.hasRuntimeBitsIgnoreComptime(zcu)) {
const pl_ptr_ty_id = try self.ptrType(layout.payload_ty, .Function, .indirect);
const pl_ptr_ty_id = try self.ptrType(layout.payload_ty, .function, .indirect);
const pl_ptr_id = try self.accessChain(pl_ptr_ty_id, tmp_id, &.{layout.payload_index});
const active_pl_ptr_id = if (!layout.payload_ty.eql(payload_ty, zcu)) blk: {
const active_pl_ptr_ty_id = try self.ptrType(payload_ty, .Function, .indirect);
const active_pl_ptr_ty_id = try self.ptrType(payload_ty, .function, .indirect);
const active_pl_ptr_id = self.spv.allocId();
try self.func.body.emit(self.spv.gpa, .OpBitcast, .{
.id_result_type = active_pl_ptr_ty_id,
@@ -5134,7 +5100,7 @@ const NavGen = struct {
return try self.load(ty, tmp_id, .{});
}
fn airUnionInit(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airUnionInit(self: *NavGen, inst: Air.Inst.Index) !?Id {
const pt = self.pt;
const zcu = pt.zcu;
const ip = &zcu.intern_pool;
@@ -5151,7 +5117,7 @@ const NavGen = struct {
return try self.unionInit(ty, extra.field_index, payload);
}
fn airStructFieldVal(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airStructFieldVal(self: *NavGen, inst: Air.Inst.Index) !?Id {
const pt = self.pt;
const zcu = pt.zcu;
const ty_pl = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
@@ -5216,13 +5182,13 @@ const NavGen = struct {
const layout = self.unionLayout(object_ty);
assert(layout.has_payload);
const tmp_id = try self.alloc(object_ty, .{ .storage_class = .Function });
const tmp_id = try self.alloc(object_ty, .{ .storage_class = .function });
try self.store(object_ty, tmp_id, object_id, .{});
const pl_ptr_ty_id = try self.ptrType(layout.payload_ty, .Function, .indirect);
const pl_ptr_ty_id = try self.ptrType(layout.payload_ty, .function, .indirect);
const pl_ptr_id = try self.accessChain(pl_ptr_ty_id, tmp_id, &.{layout.payload_index});
const active_pl_ptr_ty_id = try self.ptrType(field_ty, .Function, .indirect);
const active_pl_ptr_ty_id = try self.ptrType(field_ty, .function, .indirect);
const active_pl_ptr_id = self.spv.allocId();
try self.func.body.emit(self.spv.gpa, .OpBitcast, .{
.id_result_type = active_pl_ptr_ty_id,
@@ -5236,7 +5202,7 @@ const NavGen = struct {
}
}
fn airFieldParentPtr(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airFieldParentPtr(self: *NavGen, inst: Air.Inst.Index) !?Id {
const pt = self.pt;
const zcu = pt.zcu;
const ty_pl = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
@@ -5273,9 +5239,9 @@ const NavGen = struct {
self: *NavGen,
result_ptr_ty: Type,
object_ptr_ty: Type,
object_ptr: IdRef,
object_ptr: Id,
field_index: u32,
) !IdRef {
) !Id {
const result_ty_id = try self.resolveType(result_ptr_ty, .direct);
const zcu = self.pt.zcu;
@@ -5318,7 +5284,7 @@ const NavGen = struct {
}
}
fn airStructFieldPtrIndex(self: *NavGen, inst: Air.Inst.Index, field_index: u32) !?IdRef {
fn airStructFieldPtrIndex(self: *NavGen, inst: Air.Inst.Index, field_index: u32) !?Id {
const ty_op = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
const struct_ptr = try self.resolve(ty_op.operand);
const struct_ptr_ty = self.typeOf(ty_op.operand);
@@ -5327,7 +5293,7 @@ const NavGen = struct {
}
const AllocOptions = struct {
initializer: ?IdRef = null,
initializer: ?Id = null,
/// The final storage class of the pointer. This may be either `.Generic` or `.Function`.
/// In either case, the local is allocated in the `.Function` storage class, and optionally
/// cast back to `.Generic`.
@@ -5342,8 +5308,8 @@ const NavGen = struct {
self: *NavGen,
ty: Type,
options: AllocOptions,
) !IdRef {
const ptr_fn_ty_id = try self.ptrType(ty, .Function, .indirect);
) !Id {
const ptr_fn_ty_id = try self.ptrType(ty, .function, .indirect);
// SPIR-V requires that OpVariable declarations for locals go into the first block, so we are just going to
// directly generate them into func.prologue instead of the body.
@@ -5351,7 +5317,7 @@ const NavGen = struct {
try self.func.prologue.emit(self.spv.gpa, .OpVariable, .{
.id_result_type = ptr_fn_ty_id,
.id_result = var_id,
.storage_class = .Function,
.storage_class = .function,
.initializer = options.initializer,
});
@@ -5361,17 +5327,17 @@ const NavGen = struct {
}
switch (options.storage_class) {
.Generic => {
const ptr_gn_ty_id = try self.ptrType(ty, .Generic, .indirect);
.generic => {
const ptr_gn_ty_id = try self.ptrType(ty, .generic, .indirect);
// Convert to a generic pointer
return self.castToGeneric(ptr_gn_ty_id, var_id);
},
.Function => return var_id,
.function => return var_id,
else => unreachable,
}
}
fn airAlloc(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airAlloc(self: *NavGen, inst: Air.Inst.Index) !?Id {
const zcu = self.pt.zcu;
const ptr_ty = self.typeOfIndex(inst);
const child_ty = ptr_ty.childType(zcu);
@@ -5380,7 +5346,7 @@ const NavGen = struct {
});
}
fn airArg(self: *NavGen) IdRef {
fn airArg(self: *NavGen) Id {
defer self.next_arg_index += 1;
return self.args.items[self.next_arg_index];
}
@@ -5389,14 +5355,14 @@ const NavGen = struct {
/// block to jump to. This function emits instructions, so it should be emitted
/// inside the merge block of the block.
/// This function should only be called with structured control flow generation.
fn structuredNextBlock(self: *NavGen, incoming: []const ControlFlow.Structured.Block.Incoming) !IdRef {
fn structuredNextBlock(self: *NavGen, incoming: []const ControlFlow.Structured.Block.Incoming) !Id {
assert(self.control_flow == .structured);
const result_id = self.spv.allocId();
const block_id_ty_id = try self.resolveType(Type.u32, .direct);
try self.func.body.emitRaw(self.spv.gpa, .OpPhi, @intCast(2 + incoming.len * 2)); // result type + result + variable/parent...
self.func.body.writeOperand(spec.IdResultType, block_id_ty_id);
self.func.body.writeOperand(spec.IdRef, result_id);
self.func.body.writeOperand(spec.Id, block_id_ty_id);
self.func.body.writeOperand(spec.Id, result_id);
for (incoming) |incoming_block| {
self.func.body.writeOperand(spec.PairIdRefIdRef, .{ incoming_block.next_block, incoming_block.src_label });
@@ -5408,7 +5374,7 @@ const NavGen = struct {
/// Jumps to the block with the target block-id. This function must only be called when
/// terminating a body, there should be no instructions after it.
/// This function should only be called with structured control flow generation.
fn structuredBreak(self: *NavGen, target_block: IdRef) !void {
fn structuredBreak(self: *NavGen, target_block: Id) !void {
assert(self.control_flow == .structured);
const sblock = self.control_flow.structured.block_stack.getLast();
@@ -5448,12 +5414,12 @@ const NavGen = struct {
/// Using loops; loops can be early exited by jumping to the merge block at
/// any time.
loop: struct {
merge_label: IdRef,
continue_label: IdRef,
merge_label: Id,
continue_label: Id,
},
},
body: []const Air.Inst.Index,
) !IdRef {
) !Id {
assert(self.control_flow == .structured);
var sblock: ControlFlow.Structured.Block = switch (block_merge_type) {
@@ -5533,13 +5499,13 @@ const NavGen = struct {
}
}
fn airBlock(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airBlock(self: *NavGen, inst: Air.Inst.Index) !?Id {
const inst_datas = self.air.instructions.items(.data);
const extra = self.air.extraData(Air.Block, inst_datas[@intFromEnum(inst)].ty_pl.payload);
return self.lowerBlock(inst, @ptrCast(self.air.extra.items[extra.end..][0..extra.data.body_len]));
}
fn lowerBlock(self: *NavGen, inst: Air.Inst.Index, body: []const Air.Inst.Index) !?IdRef {
fn lowerBlock(self: *NavGen, inst: Air.Inst.Index, body: []const Air.Inst.Index) !?Id {
// In AIR, a block doesn't really define an entry point like a block, but
// more like a scope that breaks can jump out of and "return" a value from.
// This cannot be directly modelled in SPIR-V, so in a block instruction,
@@ -5584,8 +5550,8 @@ const NavGen = struct {
// result type + result + variable/parent...
2 + @as(u16, @intCast(block.incoming_blocks.items.len * 2)),
);
self.func.body.writeOperand(spec.IdResultType, result_type_id);
self.func.body.writeOperand(spec.IdRef, result_id);
self.func.body.writeOperand(spec.Id, result_type_id);
self.func.body.writeOperand(spec.Id, result_id);
for (block.incoming_blocks.items) |incoming| {
self.func.body.writeOperand(
@@ -5599,7 +5565,7 @@ const NavGen = struct {
};
const maybe_block_result_var_id = if (have_block_result) blk: {
const block_result_var_id = try self.alloc(ty, .{ .storage_class = .Function });
const block_result_var_id = try self.alloc(ty, .{ .storage_class = .function });
try cf.block_results.putNoClobber(self.gpa, inst, block_result_var_id);
break :blk block_result_var_id;
} else null;
@@ -5823,7 +5789,7 @@ const NavGen = struct {
}
}
fn airLoad(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airLoad(self: *NavGen, inst: Air.Inst.Index) !?Id {
const zcu = self.pt.zcu;
const ty_op = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
const ptr_ty = self.typeOf(ty_op.operand);
@@ -5894,7 +5860,7 @@ const NavGen = struct {
});
}
fn airTry(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airTry(self: *NavGen, inst: Air.Inst.Index) !?Id {
const zcu = self.pt.zcu;
const pl_op = self.air.instructions.items(.data)[@intFromEnum(inst)].pl_op;
const err_union_id = try self.resolve(pl_op.operand);
@@ -5964,7 +5930,7 @@ const NavGen = struct {
return try self.extractField(payload_ty, err_union_id, eu_layout.payloadFieldIndex());
}
fn airErrUnionErr(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airErrUnionErr(self: *NavGen, inst: Air.Inst.Index) !?Id {
const zcu = self.pt.zcu;
const ty_op = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
const operand_id = try self.resolve(ty_op.operand);
@@ -5987,7 +5953,7 @@ const NavGen = struct {
return try self.extractField(Type.anyerror, operand_id, eu_layout.errorFieldIndex());
}
fn airErrUnionPayload(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airErrUnionPayload(self: *NavGen, inst: Air.Inst.Index) !?Id {
const ty_op = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
const operand_id = try self.resolve(ty_op.operand);
const payload_ty = self.typeOfIndex(inst);
@@ -6000,7 +5966,7 @@ const NavGen = struct {
return try self.extractField(payload_ty, operand_id, eu_layout.payloadFieldIndex());
}
fn airWrapErrUnionErr(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airWrapErrUnionErr(self: *NavGen, inst: Air.Inst.Index) !?Id {
const zcu = self.pt.zcu;
const ty_op = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
const err_union_ty = self.typeOfIndex(inst);
@@ -6014,7 +5980,7 @@ const NavGen = struct {
const payload_ty_id = try self.resolveType(payload_ty, .indirect);
var members: [2]IdRef = undefined;
var members: [2]Id = undefined;
members[eu_layout.errorFieldIndex()] = operand_id;
members[eu_layout.payloadFieldIndex()] = try self.spv.constUndef(payload_ty_id);
@@ -6026,7 +5992,7 @@ const NavGen = struct {
return try self.constructComposite(err_union_ty_id, &members);
}
fn airWrapErrUnionPayload(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airWrapErrUnionPayload(self: *NavGen, inst: Air.Inst.Index) !?Id {
const ty_op = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
const err_union_ty = self.typeOfIndex(inst);
const operand_id = try self.resolve(ty_op.operand);
@@ -6037,7 +6003,7 @@ const NavGen = struct {
return try self.constInt(Type.anyerror, 0);
}
var members: [2]IdRef = undefined;
var members: [2]Id = undefined;
members[eu_layout.errorFieldIndex()] = try self.constInt(Type.anyerror, 0);
members[eu_layout.payloadFieldIndex()] = try self.convertToIndirect(payload_ty, operand_id);
@@ -6049,7 +6015,7 @@ const NavGen = struct {
return try self.constructComposite(err_union_ty_id, &members);
}
fn airIsNull(self: *NavGen, inst: Air.Inst.Index, is_pointer: bool, pred: enum { is_null, is_non_null }) !?IdRef {
fn airIsNull(self: *NavGen, inst: Air.Inst.Index, is_pointer: bool, pred: enum { is_null, is_non_null }) !?Id {
const pt = self.pt;
const zcu = pt.zcu;
const un_op = self.air.instructions.items(.data)[@intFromEnum(inst)].un_op;
@@ -6126,7 +6092,7 @@ const NavGen = struct {
};
}
fn airIsErr(self: *NavGen, inst: Air.Inst.Index, pred: enum { is_err, is_non_err }) !?IdRef {
fn airIsErr(self: *NavGen, inst: Air.Inst.Index, pred: enum { is_err, is_non_err }) !?Id {
const zcu = self.pt.zcu;
const un_op = self.air.instructions.items(.data)[@intFromEnum(inst)].un_op;
const operand_id = try self.resolve(un_op);
@@ -6164,7 +6130,7 @@ const NavGen = struct {
return result_id;
}
fn airUnwrapOptional(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airUnwrapOptional(self: *NavGen, inst: Air.Inst.Index) !?Id {
const pt = self.pt;
const zcu = pt.zcu;
const ty_op = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
@@ -6181,7 +6147,7 @@ const NavGen = struct {
return try self.extractField(payload_ty, operand_id, 0);
}
fn airUnwrapOptionalPtr(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airUnwrapOptionalPtr(self: *NavGen, inst: Air.Inst.Index) !?Id {
const pt = self.pt;
const zcu = pt.zcu;
const ty_op = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
@@ -6206,7 +6172,7 @@ const NavGen = struct {
return try self.accessChain(result_ty_id, operand_id, &.{0});
}
fn airWrapOptional(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airWrapOptional(self: *NavGen, inst: Air.Inst.Index) !?Id {
const pt = self.pt;
const zcu = pt.zcu;
const ty_op = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
@@ -6224,7 +6190,7 @@ const NavGen = struct {
}
const payload_id = try self.convertToIndirect(payload_ty, operand_id);
const members = [_]IdRef{ payload_id, try self.constBool(true, .indirect) };
const members = [_]Id{ payload_id, try self.constBool(true, .indirect) };
const optional_ty_id = try self.resolveType(optional_ty, .direct);
return try self.constructComposite(optional_ty_id, &members);
}
@@ -6294,8 +6260,8 @@ const NavGen = struct {
// Emit the instruction before generating the blocks.
try self.func.body.emitRaw(self.spv.gpa, .OpSwitch, 2 + (cond_words + 1) * num_conditions);
self.func.body.writeOperand(IdRef, cond_indirect);
self.func.body.writeOperand(IdRef, default);
self.func.body.writeOperand(Id, cond_indirect);
self.func.body.writeOperand(Id, default);
// Emit each of the cases
{
@@ -6322,7 +6288,7 @@ const NavGen = struct {
else => unreachable,
};
self.func.body.writeOperand(spec.LiteralContextDependentNumber, int_lit);
self.func.body.writeOperand(IdRef, label);
self.func.body.writeOperand(Id, label);
}
}
}
@@ -6399,7 +6365,7 @@ const NavGen = struct {
});
}
fn airDbgInlineBlock(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airDbgInlineBlock(self: *NavGen, inst: Air.Inst.Index) !?Id {
const zcu = self.pt.zcu;
const inst_datas = self.air.instructions.items(.data);
const extra = self.air.extraData(Air.DbgInlineBlock, inst_datas[@intFromEnum(inst)].ty_pl.payload);
@@ -6416,7 +6382,7 @@ const NavGen = struct {
try self.spv.debugName(target_id, name.toSlice(self.air));
}
fn airAssembly(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airAssembly(self: *NavGen, inst: Air.Inst.Index) !?Id {
const zcu = self.pt.zcu;
const ty_pl = self.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
const extra = self.air.extraData(Air.Asm, ty_pl.payload);
@@ -6586,7 +6552,7 @@ const NavGen = struct {
return null;
}
fn airCall(self: *NavGen, inst: Air.Inst.Index, modifier: std.builtin.CallModifier) !?IdRef {
fn airCall(self: *NavGen, inst: Air.Inst.Index, modifier: std.builtin.CallModifier) !?Id {
_ = modifier;
const pt = self.pt;
@@ -6608,7 +6574,7 @@ const NavGen = struct {
const callee_id = try self.resolve(pl_op.operand);
comptime assert(zig_call_abi_ver == 3);
const params = try self.gpa.alloc(spec.IdRef, args.len);
const params = try self.gpa.alloc(spec.Id, args.len);
defer self.gpa.free(params);
var n_params: usize = 0;
for (args) |arg| {
@@ -6637,7 +6603,7 @@ const NavGen = struct {
return result_id;
}
fn builtin3D(self: *NavGen, result_ty: Type, builtin: spec.BuiltIn, dimension: u32, out_of_range_value: anytype) !IdRef {
fn builtin3D(self: *NavGen, result_ty: Type, builtin: spec.BuiltIn, dimension: u32, out_of_range_value: anytype) !Id {
if (dimension >= 3) {
return try self.constInt(result_ty, out_of_range_value);
}
@@ -6645,7 +6611,7 @@ const NavGen = struct {
.len = 3,
.child = result_ty.toIntern(),
});
const ptr_ty_id = try self.ptrType(vec_ty, .Input, .indirect);
const ptr_ty_id = try self.ptrType(vec_ty, .input, .indirect);
const spv_decl_index = try self.spv.builtin(ptr_ty_id, builtin);
try self.func.decl_deps.put(self.spv.gpa, spv_decl_index, {});
const ptr = self.spv.declPtr(spv_decl_index).result_id;
@@ -6653,34 +6619,34 @@ const NavGen = struct {
return try self.extractVectorComponent(result_ty, vec, dimension);
}
fn airWorkItemId(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airWorkItemId(self: *NavGen, inst: Air.Inst.Index) !?Id {
if (self.liveness.isUnused(inst)) return null;
const pl_op = self.air.instructions.items(.data)[@intFromEnum(inst)].pl_op;
const dimension = pl_op.payload;
// TODO: Should we make these builtins return usize?
const result_id = try self.builtin3D(Type.u64, .LocalInvocationId, dimension, 0);
const result_id = try self.builtin3D(Type.u64, .local_invocation_id, dimension, 0);
const tmp = Temporary.init(Type.u64, result_id);
const result = try self.buildConvert(Type.u32, tmp);
return try result.materialize(self);
}
fn airWorkGroupSize(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airWorkGroupSize(self: *NavGen, inst: Air.Inst.Index) !?Id {
if (self.liveness.isUnused(inst)) return null;
const pl_op = self.air.instructions.items(.data)[@intFromEnum(inst)].pl_op;
const dimension = pl_op.payload;
// TODO: Should we make these builtins return usize?
const result_id = try self.builtin3D(Type.u64, .WorkgroupSize, dimension, 0);
const result_id = try self.builtin3D(Type.u64, .workgroup_size, dimension, 0);
const tmp = Temporary.init(Type.u64, result_id);
const result = try self.buildConvert(Type.u32, tmp);
return try result.materialize(self);
}
fn airWorkGroupId(self: *NavGen, inst: Air.Inst.Index) !?IdRef {
fn airWorkGroupId(self: *NavGen, inst: Air.Inst.Index) !?Id {
if (self.liveness.isUnused(inst)) return null;
const pl_op = self.air.instructions.items(.data)[@intFromEnum(inst)].pl_op;
const dimension = pl_op.payload;
// TODO: Should we make these builtins return usize?
const result_id = try self.builtin3D(Type.u64, .WorkgroupId, dimension, 0);
const result_id = try self.builtin3D(Type.u64, .workgroup_id, dimension, 0);
const tmp = Temporary.init(Type.u64, result_id);
const result = try self.buildConvert(Type.u32, tmp);
return try result.materialize(self);
src/codegen/spirv/Assembler.zig
+23 -24
View File
@@ -7,8 +7,7 @@ const assert = std.debug.assert;
const spec = @import("spec.zig");
const Opcode = spec.Opcode;
const Word = spec.Word;
const IdRef = spec.IdRef;
const IdResult = spec.IdResult;
const Id = spec.Id;
const StorageClass = spec.StorageClass;
const SpvModule = @import("Module.zig");
@@ -127,10 +126,10 @@ const AsmValue = union(enum) {
unresolved_forward_reference,
/// This result-value is a normal result produced by a different instruction.
value: IdRef,
value: Id,
/// This result-value represents a type registered into the module's type system.
ty: IdRef,
ty: Id,
/// This is a pre-supplied constant integer value.
constant: u32,
@@ -141,7 +140,7 @@ const AsmValue = union(enum) {
/// Retrieve the result-id of this AsmValue. Asserts that this AsmValue
/// is of a variant that allows the result to be obtained (not an unresolved
/// forward declaration, not in the process of being declared, etc).
pub fn resultId(self: AsmValue) IdRef {
pub fn resultId(self: AsmValue) Id {
return switch (self) {
.just_declared,
.unresolved_forward_reference,
@@ -314,7 +313,7 @@ fn processInstruction(self: *Assembler) !void {
return;
},
else => switch (self.inst.opcode.class()) {
.TypeDeclaration => try self.processTypeInstruction(),
.type_declaration => try self.processTypeInstruction(),
else => (try self.processGenericInstruction()) orelse return,
},
};
@@ -392,7 +391,7 @@ fn processTypeInstruction(self: *Assembler) !AsmValue {
break :blk result_id;
},
.OpTypeStruct => blk: {
const ids = try self.gpa.alloc(IdRef, operands[1..].len);
const ids = try self.gpa.alloc(Id, operands[1..].len);
defer self.gpa.free(ids);
for (operands[1..], ids) |op, *id| id.* = try self.resolveRefId(op.ref_id);
const result_id = self.spv.allocId();
@@ -429,7 +428,7 @@ fn processTypeInstruction(self: *Assembler) !AsmValue {
const param_operands = operands[2..];
const return_type = try self.resolveRefId(operands[1].ref_id);
const param_types = try self.spv.gpa.alloc(IdRef, param_operands.len);
const param_types = try self.spv.gpa.alloc(Id, param_operands.len);
defer self.spv.gpa.free(param_types);
for (param_types, param_operands) |*param, operand| {
param.* = try self.resolveRefId(operand.ref_id);
@@ -457,17 +456,17 @@ fn processGenericInstruction(self: *Assembler) !?AsmValue {
const operands = self.inst.operands.items;
var maybe_spv_decl_index: ?SpvModule.Decl.Index = null;
const section = switch (self.inst.opcode.class()) {
.ConstantCreation => &self.spv.sections.types_globals_constants,
.Annotation => &self.spv.sections.annotations,
.TypeDeclaration => unreachable, // Handled elsewhere.
.constant_creation => &self.spv.sections.types_globals_constants,
.annotation => &self.spv.sections.annotations,
.type_declaration => unreachable, // Handled elsewhere.
else => switch (self.inst.opcode) {
.OpEntryPoint => unreachable,
.OpExecutionMode, .OpExecutionModeId => &self.spv.sections.execution_modes,
.OpVariable => section: {
const storage_class: spec.StorageClass = @enumFromInt(operands[2].value);
if (storage_class == .Function) break :section &self.func.prologue;
if (storage_class == .function) break :section &self.func.prologue;
maybe_spv_decl_index = try self.spv.allocDecl(.global);
if (self.spv.version.minor < 4 and storage_class != .Input and storage_class != .Output) {
if (self.spv.version.minor < 4 and storage_class != .input and storage_class != .output) {
// Before version 1.4, the interface’s storage classes are limited to the Input and Output
break :section &self.spv.sections.types_globals_constants;
}
@@ -481,7 +480,7 @@ fn processGenericInstruction(self: *Assembler) !?AsmValue {
},
};
var maybe_result_id: ?IdResult = null;
var maybe_result_id: ?Id = null;
const first_word = section.instructions.items.len;
// At this point we're not quite sure how many operands this instruction is going to have,
// so insert 0 and patch up the actual opcode word later.
@@ -504,12 +503,12 @@ fn processGenericInstruction(self: *Assembler) !?AsmValue {
else
self.spv.allocId();
try section.ensureUnusedCapacity(self.spv.gpa, 1);
section.writeOperand(IdResult, maybe_result_id.?);
section.writeOperand(Id, maybe_result_id.?);
},
.ref_id => |index| {
const result = try self.resolveRef(index);
try section.ensureUnusedCapacity(self.spv.gpa, 1);
section.writeOperand(spec.IdRef, result.resultId());
section.writeOperand(spec.Id, result.resultId());
},
.string => |offset| {
const text = std.mem.sliceTo(self.inst.string_bytes.items[offset..], 0);
@@ -558,7 +557,7 @@ fn resolveRef(self: *Assembler, ref: AsmValue.Ref) !AsmValue {
}
}
fn resolveRefId(self: *Assembler, ref: AsmValue.Ref) !IdRef {
fn resolveRefId(self: *Assembler, ref: AsmValue.Ref) !Id {
const value = try self.resolveRef(ref);
return value.resultId();
}
@@ -600,7 +599,7 @@ fn parseInstruction(self: *Assembler) !void {
const expected_operands = inst.operands;
// This is a loop because the result-id is not always the first operand.
const requires_lhs_result = for (expected_operands) |op| {
if (op.kind == .IdResult) break true;
if (op.kind == .id_result) break true;
} else false;
if (requires_lhs_result and maybe_lhs_result == null) {
@@ -614,7 +613,7 @@ fn parseInstruction(self: *Assembler) !void {
}
for (expected_operands) |operand| {
if (operand.kind == .IdResult) {
if (operand.kind == .id_result) {
try self.inst.operands.append(self.gpa, .{ .result_id = maybe_lhs_result.? });
continue;
}
@@ -646,11 +645,11 @@ fn parseOperand(self: *Assembler, kind: spec.OperandKind) Error!void {
.value_enum => try self.parseValueEnum(kind),
.id => try self.parseRefId(),
else => switch (kind) {
.LiteralInteger => try self.parseLiteralInteger(),
.LiteralString => try self.parseString(),
.LiteralContextDependentNumber => try self.parseContextDependentNumber(),
.LiteralExtInstInteger => try self.parseLiteralExtInstInteger(),
.PairIdRefIdRef => try self.parsePhiSource(),
.literal_integer => try self.parseLiteralInteger(),
.literal_string => try self.parseString(),
.literal_context_dependent_number => try self.parseContextDependentNumber(),
.literal_ext_inst_integer => try self.parseLiteralExtInstInteger(),
.pair_id_ref_id_ref => try self.parsePhiSource(),
else => return self.todo("parse operand of type {s}", .{@tagName(kind)}),
},
}
src/codegen/spirv/Module.zig
+68 -70
View File
@@ -15,9 +15,7 @@ const Wyhash = std.hash.Wyhash;
const spec = @import("spec.zig");
const Word = spec.Word;
const IdRef = spec.IdRef;
const IdResult = spec.IdResult;
const IdResultType = spec.IdResultType;
const Id = spec.Id;
const Section = @import("Section.zig");
@@ -82,7 +80,7 @@ pub const Decl = struct {
/// - For `func`, this is the result-id of the associated OpFunction instruction.
/// - For `global`, this is the result-id of the associated OpVariable instruction.
/// - For `invocation_global`, this is the result-id of the associated InvocationGlobal instruction.
result_id: IdRef,
result_id: Id,
/// The offset of the first dependency of this decl in the `decl_deps` array.
begin_dep: u32,
/// The past-end offset of the dependencies of this decl in the `decl_deps` array.
@@ -150,7 +148,7 @@ sections: struct {
next_result_id: Word,
/// Cache for results of OpString instructions.
strings: std.StringArrayHashMapUnmanaged(IdRef) = .empty,
strings: std.StringArrayHashMapUnmanaged(Id) = .empty,
/// Some types shouldn't be emitted more than one time, but cannot be caught by
/// the `intern_map` during codegen. Sometimes, IDs are compared to check if
@@ -161,20 +159,20 @@ strings: std.StringArrayHashMapUnmanaged(IdRef) = .empty,
/// Additionally, this is used for other values which can be cached, for example,
/// built-in variables.
cache: struct {
bool_type: ?IdRef = null,
void_type: ?IdRef = null,
int_types: std.AutoHashMapUnmanaged(std.builtin.Type.Int, IdRef) = .empty,
float_types: std.AutoHashMapUnmanaged(std.builtin.Type.Float, IdRef) = .empty,
vector_types: std.AutoHashMapUnmanaged(struct { IdRef, u32 }, IdRef) = .empty,
array_types: std.AutoHashMapUnmanaged(struct { IdRef, IdRef }, IdRef) = .empty,
bool_type: ?Id = null,
void_type: ?Id = null,
int_types: std.AutoHashMapUnmanaged(std.builtin.Type.Int, Id) = .empty,
float_types: std.AutoHashMapUnmanaged(std.builtin.Type.Float, Id) = .empty,
vector_types: std.AutoHashMapUnmanaged(struct { Id, u32 }, Id) = .empty,
array_types: std.AutoHashMapUnmanaged(struct { Id, Id }, Id) = .empty,
capabilities: std.AutoHashMapUnmanaged(spec.Capability, void) = .empty,
extensions: std.StringHashMapUnmanaged(void) = .empty,
extended_instruction_set: std.AutoHashMapUnmanaged(spec.InstructionSet, IdRef) = .empty,
decorations: std.AutoHashMapUnmanaged(struct { IdRef, spec.Decoration }, void) = .empty,
builtins: std.AutoHashMapUnmanaged(struct { IdRef, spec.BuiltIn }, Decl.Index) = .empty,
extended_instruction_set: std.AutoHashMapUnmanaged(spec.InstructionSet, Id) = .empty,
decorations: std.AutoHashMapUnmanaged(struct { Id, spec.Decoration }, void) = .empty,
builtins: std.AutoHashMapUnmanaged(struct { Id, spec.BuiltIn }, Decl.Index) = .empty,
bool_const: [2]?IdRef = .{ null, null },
bool_const: [2]?Id = .{ null, null },
} = .{},
/// Set of Decls, referred to by Decl.Index.
@@ -185,7 +183,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.AutoArrayHashMapUnmanaged(IdRef, EntryPoint) = .empty,
entry_points: std.AutoArrayHashMapUnmanaged(Id, EntryPoint) = .empty,
pub fn init(gpa: Allocator, target: *const std.Target) Module {
const version_minor: u8 = blk: {
@@ -245,7 +243,7 @@ pub const IdRange = struct {
base: u32,
len: u32,
pub fn at(range: IdRange, i: usize) IdResult {
pub fn at(range: IdRange, i: usize) Id {
assert(i < range.len);
return @enumFromInt(range.base + i);
}
@@ -259,7 +257,7 @@ pub fn allocIds(self: *Module, n: u32) IdRange {
};
}
pub fn allocId(self: *Module) IdResult {
pub fn allocId(self: *Module) Id {
return self.allocIds(1).at(0);
}
@@ -275,7 +273,7 @@ fn addEntryPointDeps(
self: *Module,
decl_index: Decl.Index,
seen: *std.DynamicBitSetUnmanaged,
interface: *std.ArrayList(IdRef),
interface: *std.ArrayList(Id),
) !void {
const decl = self.declPtr(decl_index);
const deps = self.decl_deps.items[decl.begin_dep..decl.end_dep];
@@ -299,7 +297,7 @@ fn entryPoints(self: *Module) !Section {
var entry_points = Section{};
errdefer entry_points.deinit(self.gpa);
var interface = std.ArrayList(IdRef).init(self.gpa);
var interface = std.ArrayList(Id).init(self.gpa);
defer interface.deinit();
var seen = try std.DynamicBitSetUnmanaged.initEmpty(self.gpa, self.decls.items.len);
@@ -317,12 +315,12 @@ fn entryPoints(self: *Module) !Section {
.interface = interface.items,
});
if (entry_point.exec_mode == null and entry_point.exec_model == .Fragment) {
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,
.mode = if (tag == .vulkan) .origin_upper_left else .origin_lower_left,
});
},
.opencl => {},
@@ -338,59 +336,59 @@ pub fn finalize(self: *Module, a: Allocator) ![]Word {
// Emit capabilities and extensions
switch (self.target.os.tag) {
.opengl => {
try self.addCapability(.Shader);
try self.addCapability(.Matrix);
try self.addCapability(.shader);
try self.addCapability(.matrix);
},
.vulkan => {
try self.addCapability(.Shader);
try self.addCapability(.Matrix);
try self.addCapability(.shader);
try self.addCapability(.matrix);
if (self.target.cpu.arch == .spirv64) {
try self.addExtension("SPV_KHR_physical_storage_buffer");
try self.addCapability(.PhysicalStorageBufferAddresses);
try self.addCapability(.physical_storage_buffer_addresses);
}
},
.opencl, .amdhsa => {
try self.addCapability(.Kernel);
try self.addCapability(.Addresses);
try self.addCapability(.kernel);
try self.addCapability(.addresses);
},
else => unreachable,
}
if (self.target.cpu.arch == .spirv64) try self.addCapability(.Int64);
if (self.target.cpu.has(.spirv, .int64)) try self.addCapability(.Int64);
if (self.target.cpu.has(.spirv, .float16)) try self.addCapability(.Float16);
if (self.target.cpu.has(.spirv, .float64)) try self.addCapability(.Float64);
if (self.target.cpu.has(.spirv, .generic_pointer)) try self.addCapability(.GenericPointer);
if (self.target.cpu.has(.spirv, .vector16)) try self.addCapability(.Vector16);
if (self.target.cpu.arch == .spirv64) try self.addCapability(.int64);
if (self.target.cpu.has(.spirv, .int64)) try self.addCapability(.int64);
if (self.target.cpu.has(.spirv, .float16)) try self.addCapability(.float16);
if (self.target.cpu.has(.spirv, .float64)) try self.addCapability(.float64);
if (self.target.cpu.has(.spirv, .generic_pointer)) try self.addCapability(.generic_pointer);
if (self.target.cpu.has(.spirv, .vector16)) try self.addCapability(.vector16);
if (self.target.cpu.has(.spirv, .storage_push_constant16)) {
try self.addExtension("SPV_KHR_16bit_storage");
try self.addCapability(.StoragePushConstant16);
try self.addCapability(.storage_push_constant16);
}
if (self.target.cpu.has(.spirv, .arbitrary_precision_integers)) {
try self.addExtension("SPV_INTEL_arbitrary_precision_integers");
try self.addCapability(.ArbitraryPrecisionIntegersINTEL);
try self.addCapability(.arbitrary_precision_integers_intel);
}
if (self.target.cpu.has(.spirv, .variable_pointers)) {
try self.addExtension("SPV_KHR_variable_pointers");
try self.addCapability(.VariablePointersStorageBuffer);
try self.addCapability(.VariablePointers);
try self.addCapability(.variable_pointers_storage_buffer);
try self.addCapability(.variable_pointers);
}
// These are well supported
try self.addCapability(.Int8);
try self.addCapability(.Int16);
try self.addCapability(.int8);
try self.addCapability(.int16);
// Emit memory model
const addressing_model: spec.AddressingModel = switch (self.target.os.tag) {
.opengl => .Logical,
.vulkan => if (self.target.cpu.arch == .spirv32) .Logical else .PhysicalStorageBuffer64,
.opencl => if (self.target.cpu.arch == .spirv32) .Physical32 else .Physical64,
.amdhsa => .Physical64,
.opengl => .logical,
.vulkan => if (self.target.cpu.arch == .spirv32) .logical else .physical_storage_buffer64,
.opencl => if (self.target.cpu.arch == .spirv32) .physical32 else .physical64,
.amdhsa => .physical64,
else => unreachable,
};
try self.sections.memory_model.emit(self.gpa, .OpMemoryModel, .{
.addressing_model = addressing_model,
.memory_model = switch (self.target.os.tag) {
.opencl => .OpenCL,
.vulkan, .opengl => .GLSL450,
.opencl => .open_cl,
.vulkan, .opengl => .glsl450,
else => unreachable,
},
});
@@ -411,7 +409,7 @@ pub fn finalize(self: *Module, a: Allocator) ![]Word {
var source = Section{};
defer source.deinit(self.gpa);
try self.sections.debug_strings.emit(self.gpa, .OpSource, .{
.source_language = .Zig,
.source_language = .zig,
.version = 0,
// We cannot emit these because the Khronos translator does not parse this instruction
// correctly.
@@ -473,7 +471,7 @@ pub fn addExtension(self: *Module, ext: []const u8) !void {
}
/// Imports or returns the existing id of an extended instruction set
pub fn importInstructionSet(self: *Module, set: spec.InstructionSet) !IdRef {
pub fn importInstructionSet(self: *Module, set: spec.InstructionSet) !Id {
assert(set != .core);
const gop = try self.cache.extended_instruction_set.getOrPut(self.gpa, set);
@@ -490,7 +488,7 @@ pub fn importInstructionSet(self: *Module, set: spec.InstructionSet) !IdRef {
}
/// Fetch the result-id of an instruction corresponding to a string.
pub fn resolveString(self: *Module, string: []const u8) !IdRef {
pub fn resolveString(self: *Module, string: []const u8) !Id {
if (self.strings.get(string)) |id| {
return id;
}
@@ -506,7 +504,7 @@ pub fn resolveString(self: *Module, string: []const u8) !IdRef {
return id;
}
pub fn structType(self: *Module, result_id: IdResult, types: []const IdRef, maybe_names: ?[]const []const u8) !void {
pub fn structType(self: *Module, result_id: Id, types: []const Id, maybe_names: ?[]const []const u8) !void {
try self.sections.types_globals_constants.emit(self.gpa, .OpTypeStruct, .{
.id_result = result_id,
.id_ref = types,
@@ -520,7 +518,7 @@ pub fn structType(self: *Module, result_id: IdResult, types: []const IdRef, mayb
}
}
pub fn boolType(self: *Module) !IdRef {
pub fn boolType(self: *Module) !Id {
if (self.cache.bool_type) |id| return id;
const result_id = self.allocId();
@@ -531,7 +529,7 @@ pub fn boolType(self: *Module) !IdRef {
return result_id;
}
pub fn voidType(self: *Module) !IdRef {
pub fn voidType(self: *Module) !Id {
if (self.cache.void_type) |id| return id;
const result_id = self.allocId();
@@ -543,7 +541,7 @@ pub fn voidType(self: *Module) !IdRef {
return result_id;
}
pub fn intType(self: *Module, signedness: std.builtin.Signedness, bits: u16) !IdRef {
pub fn intType(self: *Module, signedness: std.builtin.Signedness, bits: u16) !Id {
assert(bits > 0);
const entry = try self.cache.int_types.getOrPut(self.gpa, .{ .signedness = signedness, .bits = bits });
if (!entry.found_existing) {
@@ -566,7 +564,7 @@ pub fn intType(self: *Module, signedness: std.builtin.Signedness, bits: u16) !Id
return entry.value_ptr.*;
}
pub fn floatType(self: *Module, bits: u16) !IdRef {
pub fn floatType(self: *Module, bits: u16) !Id {
assert(bits > 0);
const entry = try self.cache.float_types.getOrPut(self.gpa, .{ .bits = bits });
if (!entry.found_existing) {
@@ -581,7 +579,7 @@ pub fn floatType(self: *Module, bits: u16) !IdRef {
return entry.value_ptr.*;
}
pub fn vectorType(self: *Module, len: u32, child_ty_id: IdRef) !IdRef {
pub fn vectorType(self: *Module, len: u32, child_ty_id: Id) !Id {
const entry = try self.cache.vector_types.getOrPut(self.gpa, .{ child_ty_id, len });
if (!entry.found_existing) {
const result_id = self.allocId();
@@ -595,7 +593,7 @@ pub fn vectorType(self: *Module, len: u32, child_ty_id: IdRef) !IdRef {
return entry.value_ptr.*;
}
pub fn arrayType(self: *Module, len_id: IdRef, child_ty_id: IdRef) !IdRef {
pub fn arrayType(self: *Module, len_id: Id, child_ty_id: Id) !Id {
const entry = try self.cache.array_types.getOrPut(self.gpa, .{ child_ty_id, len_id });
if (!entry.found_existing) {
const result_id = self.allocId();
@@ -609,7 +607,7 @@ pub fn arrayType(self: *Module, len_id: IdRef, child_ty_id: IdRef) !IdRef {
return entry.value_ptr.*;
}
pub fn functionType(self: *Module, return_ty_id: IdRef, param_type_ids: []const IdRef) !IdRef {
pub fn functionType(self: *Module, return_ty_id: Id, param_type_ids: []const Id) !Id {
const result_id = self.allocId();
try self.sections.types_globals_constants.emit(self.gpa, .OpTypeFunction, .{
.id_result = result_id,
@@ -619,7 +617,7 @@ pub fn functionType(self: *Module, return_ty_id: IdRef, param_type_ids: []const
return result_id;
}
pub fn constant(self: *Module, result_ty_id: IdRef, value: spec.LiteralContextDependentNumber) !IdRef {
pub fn constant(self: *Module, result_ty_id: Id, value: spec.LiteralContextDependentNumber) !Id {
const result_id = self.allocId();
const section = &self.sections.types_globals_constants;
try section.emit(self.gpa, .OpConstant, .{
@@ -630,7 +628,7 @@ pub fn constant(self: *Module, result_ty_id: IdRef, value: spec.LiteralContextDe
return result_id;
}
pub fn constBool(self: *Module, value: bool) !IdRef {
pub fn constBool(self: *Module, value: bool) !Id {
if (self.cache.bool_const[@intFromBool(value)]) |b| return b;
const result_ty_id = try self.boolType();
@@ -653,7 +651,7 @@ pub fn constBool(self: *Module, value: bool) !IdRef {
/// Return a pointer to a builtin variable. `result_ty_id` must be a **pointer**
/// with storage class `.Input`.
pub fn builtin(self: *Module, result_ty_id: IdRef, spirv_builtin: spec.BuiltIn) !Decl.Index {
pub fn builtin(self: *Module, result_ty_id: Id, spirv_builtin: spec.BuiltIn) !Decl.Index {
const entry = try self.cache.builtins.getOrPut(self.gpa, .{ result_ty_id, spirv_builtin });
if (!entry.found_existing) {
const decl_index = try self.allocDecl(.global);
@@ -662,15 +660,15 @@ pub fn builtin(self: *Module, result_ty_id: IdRef, spirv_builtin: spec.BuiltIn)
try self.sections.types_globals_constants.emit(self.gpa, .OpVariable, .{
.id_result_type = result_ty_id,
.id_result = result_id,
.storage_class = .Input,
.storage_class = .input,
});
try self.decorate(result_id, .{ .BuiltIn = .{ .built_in = spirv_builtin } });
try self.decorate(result_id, .{ .built_in = .{ .built_in = spirv_builtin } });
try self.declareDeclDeps(decl_index, &.{});
}
return entry.value_ptr.*;
}
pub fn constUndef(self: *Module, ty_id: IdRef) !IdRef {
pub fn constUndef(self: *Module, ty_id: Id) !Id {
const result_id = self.allocId();
try self.sections.types_globals_constants.emit(self.gpa, .OpUndef, .{
.id_result_type = ty_id,
@@ -679,7 +677,7 @@ pub fn constUndef(self: *Module, ty_id: IdRef) !IdRef {
return result_id;
}
pub fn constNull(self: *Module, ty_id: IdRef) !IdRef {
pub fn constNull(self: *Module, ty_id: Id) !Id {
const result_id = self.allocId();
try self.sections.types_globals_constants.emit(self.gpa, .OpConstantNull, .{
.id_result_type = ty_id,
@@ -691,7 +689,7 @@ pub fn constNull(self: *Module, ty_id: IdRef) !IdRef {
/// Decorate a result-id.
pub fn decorate(
self: *Module,
target: IdRef,
target: Id,
decoration: spec.Decoration.Extended,
) !void {
const entry = try self.cache.decorations.getOrPut(self.gpa, .{ target, decoration });
@@ -707,7 +705,7 @@ pub fn decorate(
/// We really don't have to and shouldn't need to cache this.
pub fn decorateMember(
self: *Module,
structure_type: IdRef,
structure_type: Id,
member: u32,
decoration: spec.Decoration.Extended,
) !void {
@@ -762,20 +760,20 @@ pub fn declareEntryPoint(
if (!gop.found_existing) gop.value_ptr.exec_mode = exec_mode;
}
pub fn debugName(self: *Module, target: IdResult, name: []const u8) !void {
pub fn debugName(self: *Module, target: Id, name: []const u8) !void {
try self.sections.debug_names.emit(self.gpa, .OpName, .{
.target = target,
.name = name,
});
}
pub fn debugNameFmt(self: *Module, target: IdResult, comptime fmt: []const u8, args: anytype) !void {
pub fn debugNameFmt(self: *Module, target: Id, comptime fmt: []const u8, args: anytype) !void {
const name = try std.fmt.allocPrint(self.gpa, fmt, args);
defer self.gpa.free(name);
try self.debugName(target, name);
}
pub fn memberDebugName(self: *Module, target: IdResult, member: u32, name: []const u8) !void {
pub fn memberDebugName(self: *Module, target: Id, member: u32, name: []const u8) !void {
try self.sections.debug_names.emit(self.gpa, .OpMemberName, .{
.type = target,
.member = member,
src/codegen/spirv/Section.zig
+5 -5
View File
@@ -79,7 +79,7 @@ pub fn emit(
pub fn emitBranch(
section: *Section,
allocator: Allocator,
target_label: spec.IdRef,
target_label: spec.Id,
) !void {
try section.emit(allocator, .OpBranch, .{
.target_label = target_label,
@@ -94,8 +94,8 @@ pub fn emitSpecConstantOp(
) !void {
const word_count = operandsSize(opcode.Operands(), operands);
try section.emitRaw(allocator, .OpSpecConstantOp, 1 + word_count);
section.writeOperand(spec.IdRef, operands.id_result_type);
section.writeOperand(spec.IdRef, operands.id_result);
section.writeOperand(spec.Id, operands.id_result_type);
section.writeOperand(spec.Id, operands.id_result);
section.writeOperand(Opcode, opcode);
const fields = @typeInfo(opcode.Operands()).@"struct".fields;
@@ -134,7 +134,7 @@ fn writeOperands(section: *Section, comptime Operands: type, operands: Operands)
pub fn writeOperand(section: *Section, comptime Operand: type, operand: Operand) void {
switch (Operand) {
spec.IdResult => section.writeWord(@intFromEnum(operand)),
spec.Id => section.writeWord(@intFromEnum(operand)),
spec.LiteralInteger => section.writeWord(operand),
@@ -266,7 +266,7 @@ fn operandsSize(comptime Operands: type, operands: Operands) usize {
fn operandSize(comptime Operand: type, operand: Operand) usize {
return switch (Operand) {
spec.IdResult,
spec.Id,
spec.LiteralInteger,
spec.LiteralExtInstInteger,
=> 1,
src/codegen/spirv/extinst.zig.grammar.json
+9 -11
View File
@@ -1,13 +1,11 @@
{
"version": 0,
"revision": 0,
"instructions": [
{
"opname": "InvocationGlobal",
"opcode": 0,
"operands": [
{ "kind": "IdRef", "name": "initializer function" }
]
}
]
"version": 0,
"revision": 0,
"instructions": [
{
"opname": "InvocationGlobal",
"opcode": 0,
"operands": [{ "kind": "IdRef", "name": "initializer function" }]
}
]
}
src/codegen/spirv/spec.zig
+13379 -11494
View File
File diff suppressed because it is too large Load Diff
src/link/SpirV.zig
+5 -5
View File
@@ -42,7 +42,7 @@ const Value = @import("../Value.zig");
const SpvModule = @import("../codegen/spirv/Module.zig");
const Section = @import("../codegen/spirv/Section.zig");
const spec = @import("../codegen/spirv/spec.zig");
const IdResult = spec.IdResult;
const Id = spec.Id;
const Word = spec.Word;
const BinaryModule = @import("SpirV/BinaryModule.zig");
@@ -144,15 +144,15 @@ pub fn updateExports(
const cc = Type.fromInterned(nav_ty).fnCallingConvention(zcu);
const exec_model: spec.ExecutionModel = switch (target.os.tag) {
.vulkan, .opengl => switch (cc) {
.spirv_vertex => .Vertex,
.spirv_fragment => .Fragment,
.spirv_kernel => .GLCompute,
.spirv_vertex => .vertex,
.spirv_fragment => .fragment,
.spirv_kernel => .gl_compute,
// TODO: We should integrate with the Linkage capability and export this function
.spirv_device => return,
else => unreachable,
},
.opencl => switch (cc) {
.spirv_kernel => .Kernel,
.spirv_kernel => .kernel,
// TODO: We should integrate with the Linkage capability and export this function
.spirv_device => return,
else => unreachable,
src/link/SpirV/BinaryModule.zig
+17 -17
View File
@@ -7,7 +7,7 @@ const spec = @import("../../codegen/spirv/spec.zig");
const Opcode = spec.Opcode;
const Word = spec.Word;
const InstructionSet = spec.InstructionSet;
const ResultId = spec.IdResult;
const ResultId = spec.Id;
const BinaryModule = @This();
@@ -254,8 +254,8 @@ pub const Parser = struct {
// with ALL operations that return an int or float.
const spec_operands = inst_spec.operands;
if (spec_operands.len >= 2 and
spec_operands[0].kind == .IdResultType and
spec_operands[1].kind == .IdResult)
spec_operands[0].kind == .id_result_type and
spec_operands[1].kind == .id_result)
{
if (operands.len < 2) return error.InvalidOperands;
if (binary.arith_type_width.get(@enumFromInt(operands[0]))) |width| {
@@ -288,8 +288,8 @@ pub const Parser = struct {
var offset: usize = 0;
switch (inst.opcode) {
.OpSpecConstantOp => {
assert(operands[0].kind == .IdResultType);
assert(operands[1].kind == .IdResult);
assert(operands[0].kind == .id_result_type);
assert(operands[1].kind == .id_result);
offset = try self.parseOperandsResultIds(binary, inst, operands[0..2], offset, offsets);
if (offset >= inst.operands.len) return error.InvalidPhysicalFormat;
@@ -297,13 +297,13 @@ pub const Parser = struct {
const spec_index = self.opcode_table.get(mapSetAndOpcode(.core, spec_opcode)) orelse
return error.InvalidPhysicalFormat;
const spec_operands = InstructionSet.core.instructions()[spec_index].operands;
assert(spec_operands[0].kind == .IdResultType);
assert(spec_operands[1].kind == .IdResult);
assert(spec_operands[0].kind == .id_result_type);
assert(spec_operands[1].kind == .id_result);
offset = try self.parseOperandsResultIds(binary, inst, spec_operands[2..], offset + 1, offsets);
},
.OpExtInst => {
assert(operands[0].kind == .IdResultType);
assert(operands[1].kind == .IdResult);
assert(operands[0].kind == .id_result_type);
assert(operands[1].kind == .id_result);
offset = try self.parseOperandsResultIds(binary, inst, operands[0..2], offset, offsets);
if (offset + 1 >= inst.operands.len) return error.InvalidPhysicalFormat;
@@ -405,8 +405,8 @@ pub const Parser = struct {
offset += 1;
},
else => switch (kind) {
.LiteralInteger, .LiteralFloat => offset += 1,
.LiteralString => while (true) {
.literal_integer, .literal_float => offset += 1,
.literal_string => while (true) {
if (offset >= inst.operands.len) return error.InvalidPhysicalFormat;
const word = inst.operands[offset];
offset += 1;
@@ -419,7 +419,7 @@ pub const Parser = struct {
break;
}
},
.LiteralContextDependentNumber => {
.literal_context_dependent_number => {
assert(inst.opcode == .OpConstant or inst.opcode == .OpSpecConstantOp);
const bit_width = binary.arith_type_width.get(@enumFromInt(inst.operands[0])) orelse {
log.err("invalid LiteralContextDependentNumber type {}", .{inst.operands[0]});
@@ -431,9 +431,9 @@ pub const Parser = struct {
else => unreachable,
};
},
.LiteralExtInstInteger => unreachable,
.LiteralSpecConstantOpInteger => unreachable,
.PairLiteralIntegerIdRef => { // Switch case
.literal_ext_inst_integer => unreachable,
.literal_spec_constant_op_integer => unreachable,
.pair_literal_integer_id_ref => { // Switch case
assert(inst.opcode == .OpSwitch);
const bit_width = binary.arith_type_width.get(@enumFromInt(inst.operands[0])) orelse {
log.err("invalid OpSwitch type {}", .{inst.operands[0]});
@@ -447,11 +447,11 @@ pub const Parser = struct {
try offsets.append(@intCast(offset));
offset += 1;
},
.PairIdRefLiteralInteger => {
.pair_id_ref_literal_integer => {
try offsets.append(@intCast(offset));
offset += 2;
},
.PairIdRefIdRef => {
.pair_id_ref_id_ref => {
try offsets.append(@intCast(offset));
try offsets.append(@intCast(offset + 1));
offset += 2;
src/link/SpirV/deduplicate.zig
+11 -11
View File
@@ -7,7 +7,7 @@ const BinaryModule = @import("BinaryModule.zig");
const Section = @import("../../codegen/spirv/Section.zig");
const spec = @import("../../codegen/spirv/spec.zig");
const Opcode = spec.Opcode;
const ResultId = spec.IdResult;
const ResultId = spec.Id;
const Word = spec.Word;
fn canDeduplicate(opcode: Opcode) bool {
@@ -20,9 +20,9 @@ fn canDeduplicate(opcode: Opcode) bool {
// Debug decoration-style instructions
.OpName, .OpMemberName => true,
else => switch (opcode.class()) {
.TypeDeclaration,
.ConstantCreation,
.Annotation,
.type_declaration,
.constant_creation,
.annotation,
=> true,
else => false,
},
@@ -86,8 +86,8 @@ const ModuleInfo = struct {
if (!canDeduplicate(inst.opcode)) continue;
const result_id_index: u16 = switch (inst.opcode.class()) {
.TypeDeclaration, .Annotation, .Debug => 0,
.ConstantCreation => 1,
.type_declaration, .annotation, .debug => 0,
.constant_creation => 1,
else => unreachable,
};
@@ -101,13 +101,13 @@ const ModuleInfo = struct {
};
switch (inst.opcode.class()) {
.Annotation, .Debug => {
.annotation, .debug => {
try decorations.append(arena, .{
.target_id = result_id,
.entity = entity,
});
},
.TypeDeclaration, .ConstantCreation => {
.type_declaration, .constant_creation => {
const entry = try entities.getOrPut(result_id);
if (entry.found_existing) {
log.err("type or constant {f} has duplicate definition", .{result_id});
@@ -469,7 +469,7 @@ pub fn run(parser: *BinaryModule.Parser, binary: *BinaryModule, progress: std.Pr
const inst_spec = parser.getInstSpec(inst.opcode).?;
const maybe_result_id_offset: ?u16 = for (0..2) |i| {
if (inst_spec.operands.len > i and inst_spec.operands[i].kind == .IdResult) {
if (inst_spec.operands.len > i and inst_spec.operands[i].kind == .id_result) {
break @intCast(i);
}
} else null;
@@ -488,7 +488,7 @@ pub fn run(parser: *BinaryModule.Parser, binary: *BinaryModule, progress: std.Pr
}
switch (inst.opcode.class()) {
.Annotation, .Debug => {
.annotation, .debug => {
// For decoration-style instructions, only emit them
// if the target is not removed.
const target: ResultId = @enumFromInt(inst.operands[0]);
@@ -515,7 +515,7 @@ pub fn run(parser: *BinaryModule.Parser, binary: *BinaryModule, progress: std.Pr
// Debug and Annotation instructions don't need the forward pointer, and it
// messes up the logical layout of the module.
switch (inst.opcode.class()) {
.TypeDeclaration, .ConstantCreation, .Memory => {},
.type_declaration, .constant_creation, .memory => {},
else => continue,
}
src/link/SpirV/lower_invocation_globals.zig
+2 -2
View File
@@ -6,7 +6,7 @@ const log = std.log.scoped(.spirv_link);
const BinaryModule = @import("BinaryModule.zig");
const Section = @import("../../codegen/spirv/Section.zig");
const spec = @import("../../codegen/spirv/spec.zig");
const ResultId = spec.IdResult;
const ResultId = spec.Id;
const Word = spec.Word;
/// This structure contains all the stuff that we need to parse from the module in
@@ -626,7 +626,7 @@ const ModuleBuilder = struct {
try self.section.emit(self.arena, .OpVariable, .{
.id_result_type = global_info.ty,
.id_result = id,
.storage_class = .Function,
.storage_class = .function,
.initializer = null,
});
}
src/link/SpirV/prune_unused.zig
+10 -10
View File
@@ -15,14 +15,14 @@ const BinaryModule = @import("BinaryModule.zig");
const Section = @import("../../codegen/spirv/Section.zig");
const spec = @import("../../codegen/spirv/spec.zig");
const Opcode = spec.Opcode;
const ResultId = spec.IdResult;
const ResultId = spec.Id;
const Word = spec.Word;
/// Return whether a particular opcode's instruction can be pruned.
/// These are idempotent instructions at globals scope and instructions
/// within functions that do not have any side effects.
/// The opcodes that return true here do not necessarily need to
/// have an .IdResult. If they don't, then they are regarded
/// have an .Id. If they don't, then they are regarded
/// as 'decoration'-style instructions that don't keep their
/// operands alive, but will be emitted if they are.
fn canPrune(op: Opcode) bool {
@@ -34,12 +34,12 @@ fn canPrune(op: Opcode) bool {
// instruction has any non-trivial side effects (like OpLoad
// with the Volatile memory semantics).
return switch (op.class()) {
.TypeDeclaration,
.Conversion,
.Arithmetic,
.RelationalAndLogical,
.Bit,
.Annotation,
.type_declaration,
.conversion,
.arithmetic,
.relational_and_logical,
.bit,
.annotation,
=> true,
else => switch (op) {
.OpFunction,
@@ -111,7 +111,7 @@ const ModuleInfo = struct {
// Result-id can only be the first or second operand
const maybe_result_id: ?ResultId = for (0..2) |i| {
if (inst_spec.operands.len > i and inst_spec.operands[i].kind == .IdResult) {
if (inst_spec.operands.len > i and inst_spec.operands[i].kind == .id_result) {
break @enumFromInt(inst.operands[i]);
}
} else null;
@@ -305,7 +305,7 @@ pub fn run(parser: *BinaryModule.Parser, binary: *BinaryModule, progress: std.Pr
// Result-id can only be the first or second operand
const result_id: ResultId = for (0..2) |i| {
if (inst_spec.operands.len > i and inst_spec.operands[i].kind == .IdResult) {
if (inst_spec.operands.len > i and inst_spec.operands[i].kind == .id_result) {
break @enumFromInt(inst.operands[i]);
}
} else {
tools/gen_spirv_spec.zig
+157 -147
View File
@@ -59,26 +59,28 @@ const set_names = std.StaticStringMap([]const u8).initComptime(.{
.{ "nonsemantic.debugprintf", "NonSemantic.DebugPrintf" },
.{ "spv-amd-shader-explicit-vertex-parameter", "SPV_AMD_shader_explicit_vertex_parameter" },
.{ "nonsemantic.debugbreak", "NonSemantic.DebugBreak" },
.{ "tosa.001000.1", "SPV_EXT_INST_TYPE_TOSA_001000_1" },
.{ "zig", "zig" },
});
pub fn main() !void {
var arena = std.heap.ArenaAllocator.init(std.heap.page_allocator);
defer arena.deinit();
const a = arena.allocator();
var arena = std.heap.ArenaAllocator.init(std.heap.smp_allocator);
const allocator = arena.allocator();
const args = try std.process.argsAlloc(a);
pub fn main() !void {
defer arena.deinit();
const args = try std.process.argsAlloc(allocator);
if (args.len != 3) {
usageAndExit(args[0], 1);
}
const json_path = try std.fs.path.join(a, &.{ args[1], "include/spirv/unified1/" });
const json_path = try std.fs.path.join(allocator, &.{ args[1], "include/spirv/unified1/" });
const dir = try std.fs.cwd().openDir(json_path, .{ .iterate = true });
const core_spec = try readRegistry(CoreRegistry, a, dir, "spirv.core.grammar.json");
const core_spec = try readRegistry(CoreRegistry, dir, "spirv.core.grammar.json");
std.sort.block(Instruction, core_spec.instructions, CmpInst{}, CmpInst.lt);
var exts = std.ArrayList(Extension).init(a);
var exts = std.ArrayList(Extension).init(allocator);
var it = dir.iterate();
while (try it.next()) |entry| {
@@ -86,18 +88,43 @@ pub fn main() !void {
continue;
}
try readExtRegistry(&exts, a, dir, entry.name);
try readExtRegistry(&exts, dir, entry.name);
}
try readExtRegistry(&exts, a, std.fs.cwd(), args[2]);
try readExtRegistry(&exts, std.fs.cwd(), args[2]);
var buffer: [4000]u8 = undefined;
var w = std.fs.File.stdout().writerStreaming(&buffer);
try render(&w, a, core_spec, exts.items);
try w.flush();
const output_buf = try allocator.alloc(u8, 1024 * 1024);
var fbs = std.io.fixedBufferStream(output_buf);
var adapter = fbs.writer().adaptToNewApi();
const w = &adapter.new_interface;
try render(w, core_spec, exts.items);
var output: [:0]u8 = @ptrCast(fbs.getWritten());
output[output.len] = 0;
var tree = try std.zig.Ast.parse(allocator, output, .zig);
var color: std.zig.Color = .on;
if (tree.errors.len != 0) {
try std.zig.printAstErrorsToStderr(allocator, tree, "", color);
return;
}
var zir = try std.zig.AstGen.generate(allocator, tree);
if (zir.hasCompileErrors()) {
var wip_errors: std.zig.ErrorBundle.Wip = undefined;
try wip_errors.init(allocator);
defer wip_errors.deinit();
try wip_errors.addZirErrorMessages(zir, tree, output, "");
var error_bundle = try wip_errors.toOwnedBundle("");
defer error_bundle.deinit(allocator);
error_bundle.renderToStdErr(color.renderOptions());
}
const formatted_output = try tree.render(allocator);
_ = try std.fs.File.stdout().write(formatted_output);
}
fn readExtRegistry(exts: *std.ArrayList(Extension), a: Allocator, dir: std.fs.Dir, sub_path: []const u8) !void {
fn readExtRegistry(exts: *std.ArrayList(Extension), dir: std.fs.Dir, sub_path: []const u8) !void {
const filename = std.fs.path.basename(sub_path);
if (!std.mem.startsWith(u8, filename, "extinst.")) {
return;
@@ -105,22 +132,22 @@ fn readExtRegistry(exts: *std.ArrayList(Extension), a: Allocator, dir: std.fs.Di
std.debug.assert(std.mem.endsWith(u8, filename, ".grammar.json"));
const name = filename["extinst.".len .. filename.len - ".grammar.json".len];
const spec = try readRegistry(ExtensionRegistry, a, dir, sub_path);
const spec = try readRegistry(ExtensionRegistry, dir, sub_path);
std.sort.block(Instruction, spec.instructions, CmpInst{}, CmpInst.lt);
try exts.append(.{ .name = set_names.get(name).?, .spec = spec });
}
fn readRegistry(comptime RegistryType: type, a: Allocator, dir: std.fs.Dir, path: []const u8) !RegistryType {
const spec = try dir.readFileAlloc(a, path, std.math.maxInt(usize));
fn readRegistry(comptime RegistryType: type, dir: std.fs.Dir, path: []const u8) !RegistryType {
const spec = try dir.readFileAlloc(allocator, path, std.math.maxInt(usize));
// Required for json parsing.
@setEvalBranchQuota(10000);
var scanner = std.json.Scanner.initCompleteInput(a, spec);
var scanner = std.json.Scanner.initCompleteInput(allocator, spec);
var diagnostics = std.json.Diagnostics{};
scanner.enableDiagnostics(&diagnostics);
const parsed = std.json.parseFromTokenSource(RegistryType, a, &scanner, .{}) catch |err| {
const parsed = std.json.parseFromTokenSource(RegistryType, allocator, &scanner, .{}) catch |err| {
std.debug.print("{s}:{}:{}:\n", .{ path, diagnostics.getLine(), diagnostics.getColumn() });
return err;
};
@@ -129,11 +156,8 @@ fn readRegistry(comptime RegistryType: type, a: Allocator, dir: std.fs.Dir, path
/// Returns a set with types that require an extra struct for the `Instruction` interface
/// to the spir-v spec, or whether the original type can be used.
fn extendedStructs(
a: Allocator,
kinds: []const OperandKind,
) !ExtendedStructSet {
var map = ExtendedStructSet.init(a);
fn extendedStructs(kinds: []const OperandKind) !ExtendedStructSet {
var map = ExtendedStructSet.init(allocator);
try map.ensureTotalCapacity(@as(u32, @intCast(kinds.len)));
for (kinds) |kind| {
@@ -167,7 +191,7 @@ fn tagPriorityScore(tag: []const u8) usize {
}
}
fn render(writer: *std.io.Writer, a: Allocator, registry: CoreRegistry, extensions: []const Extension) !void {
fn render(writer: *std.io.Writer, registry: CoreRegistry, extensions: []const Extension) !void {
try writer.writeAll(
\\//! This file is auto-generated by tools/gen_spirv_spec.zig.
\\
@@ -185,22 +209,17 @@ fn render(writer: *std.io.Writer, a: Allocator, registry: CoreRegistry, extensio
\\};
\\
\\pub const Word = u32;
\\pub const IdResult = enum(Word) {
\\pub const Id = enum(Word) {
\\ none,
\\ _,
\\
\\ pub fn format(self: IdResult, writer: *std.io.Writer) std.io.Writer.Error!void {
\\ pub fn format(self: Id, writer: *std.io.Writer) std.io.Writer.Error!void {
\\ switch (self) {
\\ .none => try writer.writeAll("(none)"),
\\ else => try writer.print("%{d}", .{@intFromEnum(self)}),
\\ }
\\ }
\\};
\\pub const IdResultType = IdResult;
\\pub const IdRef = IdResult;
\\
\\pub const IdMemorySemantics = IdRef;
\\pub const IdScope = IdRef;
\\
\\pub const LiteralInteger = Word;
\\pub const LiteralFloat = Word;
@@ -215,9 +234,9 @@ fn render(writer: *std.io.Writer, a: Allocator, registry: CoreRegistry, extensio
\\};
\\pub const LiteralExtInstInteger = struct{ inst: Word };
\\pub const LiteralSpecConstantOpInteger = struct { opcode: Opcode };
\\pub const PairLiteralIntegerIdRef = struct { value: LiteralInteger, label: IdRef };
\\pub const PairIdRefLiteralInteger = struct { target: IdRef, member: LiteralInteger };
\\pub const PairIdRefIdRef = [2]IdRef;
\\pub const PairLiteralIntegerIdRef = struct { value: LiteralInteger, label: Id };
\\pub const PairIdRefLiteralInteger = struct { target: Id, member: LiteralInteger };
\\pub const PairIdRefIdRef = [2]Id;
\\
\\pub const Quantifier = enum {
\\ required,
@@ -255,7 +274,7 @@ fn render(writer: *std.io.Writer, a: Allocator, registry: CoreRegistry, extensio
);
try writer.print(
\\pub const version = Version{{ .major = {}, .minor = {}, .patch = {} }};
\\pub const version: Version = .{{ .major = {}, .minor = {}, .patch = {} }};
\\pub const magic_number: Word = {s};
\\
\\
@@ -266,7 +285,7 @@ fn render(writer: *std.io.Writer, a: Allocator, registry: CoreRegistry, extensio
// Merge the operand kinds from all extensions together.
// var all_operand_kinds = std.ArrayList(OperandKind).init(a);
// try all_operand_kinds.appendSlice(registry.operand_kinds);
var all_operand_kinds = OperandKindMap.init(a);
var all_operand_kinds = OperandKindMap.init(allocator);
for (registry.operand_kinds) |kind| {
try all_operand_kinds.putNoClobber(.{ "core", kind.kind }, kind);
}
@@ -279,35 +298,33 @@ fn render(writer: *std.io.Writer, a: Allocator, registry: CoreRegistry, extensio
try all_operand_kinds.ensureUnusedCapacity(ext.spec.operand_kinds.len);
for (ext.spec.operand_kinds) |kind| {
var new_kind = kind;
new_kind.kind = try std.mem.join(a, ".", &.{ ext.name, kind.kind });
new_kind.kind = try std.mem.join(allocator, ".", &.{ ext.name, kind.kind });
try all_operand_kinds.putNoClobber(.{ ext.name, kind.kind }, new_kind);
}
}
const extended_structs = try extendedStructs(a, all_operand_kinds.values());
const extended_structs = try extendedStructs(all_operand_kinds.values());
// Note: extensions don't seem to have class.
try renderClass(writer, a, registry.instructions);
try renderClass(writer, registry.instructions);
try renderOperandKind(writer, all_operand_kinds.values());
try renderOpcodes(writer, a, registry.instructions, extended_structs);
try renderOperandKinds(writer, a, all_operand_kinds.values(), extended_structs);
try renderInstructionSet(writer, a, registry, extensions, all_operand_kinds);
try renderOpcodes(writer, registry.instructions, extended_structs);
try renderOperandKinds(writer, all_operand_kinds.values(), extended_structs);
try renderInstructionSet(writer, registry, extensions, all_operand_kinds);
}
fn renderInstructionSet(
writer: anytype,
a: Allocator,
core: CoreRegistry,
extensions: []const Extension,
all_operand_kinds: OperandKindMap,
) !void {
_ = a;
try writer.writeAll(
\\pub const InstructionSet = enum {
\\ core,
);
for (extensions) |ext| {
try writer.print("{p},\n", .{std.zig.fmtId(ext.name)});
try writer.print("{f},\n", .{formatId(ext.name)});
}
try writer.writeAll(
@@ -340,14 +357,14 @@ fn renderInstructionsCase(
// but there aren't so many total aliases and that would add more overhead in total. We will
// just filter those out when needed.
try writer.print(".{p_} => &[_]Instruction{{\n", .{std.zig.fmtId(set_name)});
try writer.print(".{f} => &.{{\n", .{formatId(set_name)});
for (instructions) |inst| {
try writer.print(
\\.{{
\\ .name = "{s}",
\\ .opcode = {},
\\ .operands = &[_]Operand{{
\\ .operands = &.{{
\\
, .{ inst.opname, inst.opcode });
@@ -362,7 +379,7 @@ fn renderInstructionsCase(
const kind = all_operand_kinds.get(.{ set_name, operand.kind }) orelse
all_operand_kinds.get(.{ "core", operand.kind }).?;
try writer.print(".{{.kind = .{p_}, .quantifier = .{s}}},\n", .{ std.zig.fmtId(kind.kind), quantifier });
try writer.print(".{{.kind = .{f}, .quantifier = .{s}}},\n", .{ formatId(kind.kind), quantifier });
}
try writer.writeAll(
@@ -378,54 +395,69 @@ fn renderInstructionsCase(
);
}
fn renderClass(writer: anytype, a: Allocator, instructions: []const Instruction) !void {
var class_map = std.StringArrayHashMap(void).init(a);
fn renderClass(writer: anytype, instructions: []const Instruction) !void {
var class_map = std.StringArrayHashMap(void).init(allocator);
for (instructions) |inst| {
if (std.mem.eql(u8, inst.class.?, "@exclude")) {
continue;
}
if (std.mem.eql(u8, inst.class.?, "@exclude")) continue;
try class_map.put(inst.class.?, {});
}
try writer.writeAll("pub const Class = enum {\n");
for (class_map.keys()) |class| {
try renderInstructionClass(writer, class);
try writer.writeAll(",\n");
try writer.print("{f},\n", .{formatId(class)});
}
try writer.writeAll("};\n\n");
}
fn renderInstructionClass(writer: anytype, class: []const u8) !void {
// Just assume that these wont clobber zig builtin types.
var prev_was_sep = true;
for (class) |c| {
switch (c) {
'-', '_' => prev_was_sep = true,
else => if (prev_was_sep) {
try writer.writeByte(std.ascii.toUpper(c));
prev_was_sep = false;
} else {
try writer.writeByte(std.ascii.toLower(c));
},
const Formatter = struct {
data: []const u8,
fn format(f: Formatter, writer: *std.io.Writer) std.io.Writer.Error!void {
var id_buf: [128]u8 = undefined;
var fbs = std.io.fixedBufferStream(&id_buf);
const fw = fbs.writer();
for (f.data, 0..) |c, i| {
switch (c) {
'-', '_', '.', '~', ' ' => fw.writeByte('_') catch return error.WriteFailed,
'a'...'z', '0'...'9' => fw.writeByte(c) catch return error.WriteFailed,
'A'...'Z' => {
if ((i > 0 and std.ascii.isLower(f.data[i - 1])) or
(i > 0 and std.ascii.isUpper(f.data[i - 1]) and
i + 1 < f.data.len and std.ascii.isLower(f.data[i + 1])))
{
_ = fw.write(&.{ '_', std.ascii.toLower(c) }) catch return error.WriteFailed;
} else {
fw.writeByte(std.ascii.toLower(c)) catch return error.WriteFailed;
}
},
else => unreachable,
}
}
// make sure that this won't clobber with zig keywords
try writer.print("{f}", .{std.zig.fmtId(fbs.getWritten())});
}
};
fn formatId(identifier: []const u8) std.fmt.Alt(Formatter, Formatter.format) {
return .{ .data = .{ .data = identifier } };
}
fn renderOperandKind(writer: anytype, operands: []const OperandKind) !void {
try writer.writeAll(
\\pub const OperandKind = enum {
\\ Opcode,
\\ opcode,
\\
);
for (operands) |operand| {
try writer.print("{p},\n", .{std.zig.fmtId(operand.kind)});
try writer.print("{f},\n", .{formatId(operand.kind)});
}
try writer.writeAll(
\\
\\pub fn category(self: OperandKind) OperandCategory {
\\ return switch (self) {
\\ .Opcode => .literal,
\\ .opcode => .literal,
\\
);
for (operands) |operand| {
@@ -436,26 +468,26 @@ fn renderOperandKind(writer: anytype, operands: []const OperandKind) !void {
.Literal => "literal",
.Composite => "composite",
};
try writer.print(".{p_} => .{s},\n", .{ std.zig.fmtId(operand.kind), cat });
try writer.print(".{f} => .{s},\n", .{ formatId(operand.kind), cat });
}
try writer.writeAll(
\\ };
\\}
\\pub fn enumerants(self: OperandKind) []const Enumerant {
\\ return switch (self) {
\\ .Opcode => unreachable,
\\ .opcode => unreachable,
\\
);
for (operands) |operand| {
switch (operand.category) {
.BitEnum, .ValueEnum => {},
else => {
try writer.print(".{p_} => unreachable,\n", .{std.zig.fmtId(operand.kind)});
try writer.print(".{f} => unreachable,\n", .{formatId(operand.kind)});
continue;
},
}
try writer.print(".{p_} => &[_]Enumerant{{", .{std.zig.fmtId(operand.kind)});
try writer.print(".{f} => &.{{", .{formatId(operand.kind)});
for (operand.enumerants.?) |enumerant| {
if (enumerant.value == .bitflag and std.mem.eql(u8, enumerant.enumerant, "None")) {
continue;
@@ -474,32 +506,30 @@ fn renderEnumerant(writer: anytype, enumerant: Enumerant) !void {
.bitflag => |flag| try writer.writeAll(flag),
.int => |int| try writer.print("{}", .{int}),
}
try writer.writeAll(", .parameters = &[_]OperandKind{");
try writer.writeAll(", .parameters = &.{");
for (enumerant.parameters, 0..) |param, i| {
if (i != 0)
try writer.writeAll(", ");
// Note, param.quantifier will always be one.
try writer.print(".{p_}", .{std.zig.fmtId(param.kind)});
try writer.print(".{f}", .{formatId(param.kind)});
}
try writer.writeAll("}}");
}
fn renderOpcodes(
writer: anytype,
a: Allocator,
instructions: []const Instruction,
extended_structs: ExtendedStructSet,
) !void {
var inst_map = std.AutoArrayHashMap(u32, usize).init(a);
var inst_map = std.AutoArrayHashMap(u32, usize).init(allocator);
try inst_map.ensureTotalCapacity(instructions.len);
var aliases = std.ArrayList(struct { inst: usize, alias: usize }).init(a);
var aliases = std.ArrayList(struct { inst: usize, alias: usize }).init(allocator);
try aliases.ensureTotalCapacity(instructions.len);
for (instructions, 0..) |inst, i| {
if (std.mem.eql(u8, inst.class.?, "@exclude")) {
continue;
}
if (std.mem.eql(u8, inst.class.?, "@exclude")) continue;
const result = inst_map.getOrPutAssumeCapacity(inst.opcode);
if (!result.found_existing) {
result.value_ptr.* = i;
@@ -525,7 +555,7 @@ fn renderOpcodes(
try writer.writeAll("pub const Opcode = enum(u16) {\n");
for (instructions_indices) |i| {
const inst = instructions[i];
try writer.print("{p} = {},\n", .{ std.zig.fmtId(inst.opname), inst.opcode });
try writer.print("{f} = {},\n", .{ std.zig.fmtId(inst.opname), inst.opcode });
}
try writer.writeAll(
@@ -533,9 +563,9 @@ fn renderOpcodes(
);
for (aliases.items) |alias| {
try writer.print("pub const {} = Opcode.{p_};\n", .{
std.zig.fmtId(instructions[alias.inst].opname),
std.zig.fmtId(instructions[alias.alias].opname),
try writer.print("pub const {f} = Opcode.{f};\n", .{
formatId(instructions[alias.inst].opname),
formatId(instructions[alias.alias].opname),
});
}
@@ -548,7 +578,7 @@ fn renderOpcodes(
for (instructions_indices) |i| {
const inst = instructions[i];
try renderOperand(writer, .instruction, inst.opname, inst.operands, extended_structs);
try renderOperand(writer, .instruction, inst.opname, inst.operands, extended_structs, false);
}
try writer.writeAll(
@@ -561,9 +591,7 @@ fn renderOpcodes(
for (instructions_indices) |i| {
const inst = instructions[i];
try writer.print(".{p_} => .", .{std.zig.fmtId(inst.opname)});
try renderInstructionClass(writer, inst.class.?);
try writer.writeAll(",\n");
try writer.print(".{f} => .{f},\n", .{ std.zig.fmtId(inst.opname), formatId(inst.class.?) });
}
try writer.writeAll(
@@ -576,14 +604,13 @@ fn renderOpcodes(
fn renderOperandKinds(
writer: anytype,
a: Allocator,
kinds: []const OperandKind,
extended_structs: ExtendedStructSet,
) !void {
for (kinds) |kind| {
switch (kind.category) {
.ValueEnum => try renderValueEnum(writer, a, kind, extended_structs),
.BitEnum => try renderBitEnum(writer, a, kind, extended_structs),
.ValueEnum => try renderValueEnum(writer, kind, extended_structs),
.BitEnum => try renderBitEnum(writer, kind, extended_structs),
else => {},
}
}
@@ -591,20 +618,18 @@ fn renderOperandKinds(
fn renderValueEnum(
writer: anytype,
a: Allocator,
enumeration: OperandKind,
extended_structs: ExtendedStructSet,
) !void {
const enumerants = enumeration.enumerants orelse return error.InvalidRegistry;
var enum_map = std.AutoArrayHashMap(u32, usize).init(a);
var enum_map = std.AutoArrayHashMap(u32, usize).init(allocator);
try enum_map.ensureTotalCapacity(enumerants.len);
var aliases = std.ArrayList(struct { enumerant: usize, alias: usize }).init(a);
var aliases = std.ArrayList(struct { enumerant: usize, alias: usize }).init(allocator);
try aliases.ensureTotalCapacity(enumerants.len);
for (enumerants, 0..) |enumerant, i| {
try writer.context.flush();
const value: u31 = switch (enumerant.value) {
.int => |value| value,
// Some extensions declare ints as string
@@ -632,25 +657,25 @@ fn renderValueEnum(
const enum_indices = enum_map.values();
try writer.print("pub const {} = enum(u32) {{\n", .{std.zig.fmtId(enumeration.kind)});
try writer.print("pub const {f} = enum(u32) {{\n", .{std.zig.fmtId(enumeration.kind)});
for (enum_indices) |i| {
const enumerant = enumerants[i];
// if (enumerant.value != .int) return error.InvalidRegistry;
switch (enumerant.value) {
.int => |value| try writer.print("{p} = {},\n", .{ std.zig.fmtId(enumerant.enumerant), value }),
.bitflag => |value| try writer.print("{p} = {s},\n", .{ std.zig.fmtId(enumerant.enumerant), value }),
.int => |value| try writer.print("{f} = {},\n", .{ formatId(enumerant.enumerant), value }),
.bitflag => |value| try writer.print("{f} = {s},\n", .{ formatId(enumerant.enumerant), value }),
}
}
try writer.writeByte('\n');
for (aliases.items) |alias| {
try writer.print("pub const {} = {}.{p_};\n", .{
std.zig.fmtId(enumerants[alias.enumerant].enumerant),
try writer.print("pub const {f} = {f}.{f};\n", .{
formatId(enumerants[alias.enumerant].enumerant),
std.zig.fmtId(enumeration.kind),
std.zig.fmtId(enumerants[alias.alias].enumerant),
formatId(enumerants[alias.alias].enumerant),
});
}
@@ -659,11 +684,11 @@ fn renderValueEnum(
return;
}
try writer.print("\npub const Extended = union({}) {{\n", .{std.zig.fmtId(enumeration.kind)});
try writer.print("\npub const Extended = union({f}) {{\n", .{std.zig.fmtId(enumeration.kind)});
for (enum_indices) |i| {
const enumerant = enumerants[i];
try renderOperand(writer, .@"union", enumerant.enumerant, enumerant.parameters, extended_structs);
try renderOperand(writer, .@"union", enumerant.enumerant, enumerant.parameters, extended_structs, true);
}
try writer.writeAll("};\n};\n");
@@ -671,16 +696,15 @@ fn renderValueEnum(
fn renderBitEnum(
writer: anytype,
a: Allocator,
enumeration: OperandKind,
extended_structs: ExtendedStructSet,
) !void {
try writer.print("pub const {} = packed struct {{\n", .{std.zig.fmtId(enumeration.kind)});
try writer.print("pub const {f} = packed struct {{\n", .{std.zig.fmtId(enumeration.kind)});
var flags_by_bitpos = [_]?usize{null} ** 32;
const enumerants = enumeration.enumerants orelse return error.InvalidRegistry;
var aliases = std.ArrayList(struct { flag: usize, alias: u5 }).init(a);
var aliases = std.ArrayList(struct { flag: usize, alias: u5 }).init(allocator);
try aliases.ensureTotalCapacity(enumerants.len);
for (enumerants, 0..) |enumerant, i| {
@@ -715,7 +739,7 @@ fn renderBitEnum(
for (flags_by_bitpos, 0..) |maybe_flag_index, bitpos| {
if (maybe_flag_index) |flag_index| {
try writer.print("{p_}", .{std.zig.fmtId(enumerants[flag_index].enumerant)});
try writer.print("{f}", .{formatId(enumerants[flag_index].enumerant)});
} else {
try writer.print("_reserved_bit_{}", .{bitpos});
}
@@ -726,10 +750,10 @@ fn renderBitEnum(
try writer.writeByte('\n');
for (aliases.items) |alias| {
try writer.print("pub const {}: {} = .{{.{p_} = true}};\n", .{
std.zig.fmtId(enumerants[alias.flag].enumerant),
try writer.print("pub const {f}: {f} = .{{.{f} = true}};\n", .{
formatId(enumerants[alias.flag].enumerant),
std.zig.fmtId(enumeration.kind),
std.zig.fmtId(enumerants[flags_by_bitpos[alias.alias].?].enumerant),
formatId(enumerants[flags_by_bitpos[alias.alias].?].enumerant),
});
}
@@ -747,7 +771,7 @@ fn renderBitEnum(
};
const enumerant = enumerants[flag_index];
try renderOperand(writer, .mask, enumerant.enumerant, enumerant.parameters, extended_structs);
try renderOperand(writer, .mask, enumerant.enumerant, enumerant.parameters, extended_structs, true);
}
try writer.writeAll("};\n};\n");
@@ -763,11 +787,18 @@ fn renderOperand(
field_name: []const u8,
parameters: []const Operand,
extended_structs: ExtendedStructSet,
snake_case: bool,
) !void {
if (kind == .instruction) {
try writer.writeByte('.');
}
try writer.print("{}", .{std.zig.fmtId(field_name)});
if (snake_case) {
try writer.print("{f}", .{formatId(field_name)});
} else {
try writer.print("{f}", .{std.zig.fmtId(field_name)});
}
if (parameters.len == 0) {
switch (kind) {
.@"union" => try writer.writeAll(",\n"),
@@ -787,7 +818,7 @@ fn renderOperand(
try writer.writeByte('?');
}
try writer.writeAll("struct{");
try writer.writeAll("struct {");
for (parameters, 0..) |param, j| {
if (j != 0) {
@@ -804,7 +835,11 @@ fn renderOperand(
}
}
try writer.print("{}", .{std.zig.fmtId(param.kind)});
if (std.mem.startsWith(u8, param.kind, "Id")) {
_ = try writer.write("Id");
} else {
try writer.print("{f}", .{std.zig.fmtId(param.kind)});
}
if (extended_structs.contains(param.kind)) {
try writer.writeAll(".Extended");
@@ -830,49 +865,24 @@ fn renderOperand(
fn renderFieldName(writer: anytype, operands: []const Operand, field_index: usize) !void {
const operand = operands[field_index];
// Should be enough for all names - adjust as needed.
var name_backing_buffer: [64]u8 = undefined;
var name_buffer = std.ArrayListUnmanaged(u8).initBuffer(&name_backing_buffer);
derive_from_kind: {
// Operand names are often in the json encoded as "'Name'" (with two sets of quotes).
// Additionally, some operands have ~ in them at the end (D~ref~).
const name = std.mem.trim(u8, operand.name, "'~");
if (name.len == 0) {
break :derive_from_kind;
}
if (name.len == 0) break :derive_from_kind;
// Some names have weird characters in them (like newlines) - skip any such ones.
// Use the same loop to transform to snake-case.
for (name) |c| {
switch (c) {
'a'...'z', '0'...'9' => name_buffer.appendAssumeCapacity(c),
'A'...'Z' => name_buffer.appendAssumeCapacity(std.ascii.toLower(c)),
' ', '~' => name_buffer.appendAssumeCapacity('_'),
'a'...'z', '0'...'9', 'A'...'Z', ' ', '~' => continue,
else => break :derive_from_kind,
}
}
// Assume there are no duplicate 'name' fields.
try writer.print("{p_}", .{std.zig.fmtId(name_buffer.items)});
try writer.print("{f}", .{formatId(name)});
return;
}
// Translate to snake case.
name_buffer.items.len = 0;
for (operand.kind, 0..) |c, i| {
switch (c) {
'a'...'z', '0'...'9' => name_buffer.appendAssumeCapacity(c),
'A'...'Z' => if (i > 0 and std.ascii.isLower(operand.kind[i - 1])) {
name_buffer.appendSliceAssumeCapacity(&[_]u8{ '_', std.ascii.toLower(c) });
} else {
name_buffer.appendAssumeCapacity(std.ascii.toLower(c));
},
else => unreachable, // Assume that the name is valid C-syntax (and contains no underscores).
}
}
try writer.print("{p_}", .{std.zig.fmtId(name_buffer.items)});
try writer.print("{f}", .{formatId(operand.kind)});
// For fields derived from type name, there could be any amount.
// Simply check against all other fields, and if another similar one exists, add a number.
tools/spirv/grammar.zig
+4
View File
@@ -37,9 +37,11 @@ pub const InstructionPrintingClass = struct {
pub const Instruction = struct {
opname: []const u8,
class: ?[]const u8 = null, // Note: Only available in the core registry.
aliases: [][]const u8 = &[_][]const u8{},
opcode: u32,
operands: []Operand = &[_]Operand{},
capabilities: [][]const u8 = &[_][]const u8{},
provisional: bool = false,
// DebugModuleINTEL has this...
capability: ?[]const u8 = null,
extensions: [][]const u8 = &[_][]const u8{},
@@ -81,6 +83,7 @@ pub const OperandKind = struct {
pub const Enumerant = struct {
enumerant: []const u8,
aliases: [][]const u8 = &[_][]const u8{},
value: union(enum) {
bitflag: []const u8, // Hexadecimal representation of the value
int: u31,
@@ -100,6 +103,7 @@ pub const Enumerant = struct {
pub const jsonStringify = @compileError("not supported");
},
capabilities: [][]const u8 = &[_][]const u8{},
provisional: bool = false,
/// Valid for .ValueEnum and .BitEnum
extensions: [][]const u8 = &[_][]const u8{},
/// `quantifier` will always be `null`.