diff options
Diffstat (limited to 'src')
| -rw-r--r-- | src/arch/spirv/CodeGen.zig | 410 | ||||
| -rw-r--r-- | src/arch/spirv/Module.zig | 117 |
2 files changed, 246 insertions, 281 deletions
diff --git a/src/arch/spirv/CodeGen.zig b/src/arch/spirv/CodeGen.zig index bde3ef33e8..c641c43cf8 100644 --- a/src/arch/spirv/CodeGen.zig +++ b/src/arch/spirv/CodeGen.zig @@ -40,7 +40,6 @@ pub fn legalizeFeatures(_: *const std.Target) *const Air.Legalize.Features { } pub const zig_call_abi_ver = 3; -pub const big_int_bits = 32; const ControlFlow = union(enum) { const Structured = struct { @@ -183,6 +182,7 @@ pub fn genNav(cg: *CodeGen, do_codegen: bool) Error!void { const gpa = cg.module.gpa; const zcu = cg.module.zcu; const ip = &zcu.intern_pool; + const target = zcu.getTarget(); const nav = ip.getNav(cg.owner_nav); const val = zcu.navValue(cg.owner_nav); @@ -251,19 +251,19 @@ pub fn genNav(cg: *CodeGen, do_codegen: bool) Error!void { // Append the actual code into the functions section. try cg.module.sections.functions.append(cg.module.gpa, cg.prologue); try cg.module.sections.functions.append(cg.module.gpa, cg.body); - try cg.module.declareDeclDeps(spv_decl_index, cg.decl_deps.keys()); - - try cg.module.debugName(func_result_id, nav.fqn.toSlice(ip)); // Temporarily generate a test kernel declaration if this is a test function. if (is_test) { try cg.generateTestEntryPoint(nav.fqn.toSlice(ip), spv_decl_index, func_result_id); } + + try cg.module.declareDeclDeps(spv_decl_index, cg.decl_deps.keys()); + try cg.module.debugName(func_result_id, nav.fqn.toSlice(ip)); }, .global => { const maybe_init_val: ?Value = switch (ip.indexToKey(val.toIntern())) { .func => unreachable, - .variable => |variable| Value.fromInterned(variable.init), + .variable => |variable| .fromInterned(variable.init), .@"extern" => null, else => val, }; @@ -272,7 +272,8 @@ pub fn genNav(cg: *CodeGen, do_codegen: bool) Error!void { const storage_class = cg.module.storageClass(nav.getAddrspace()); assert(storage_class != .generic); // These should be instance globals - const ptr_ty_id = try cg.ptrType(ty, storage_class, .indirect); + const ty_id = try cg.resolveType(ty, .indirect); + const ptr_ty_id = try cg.module.ptrType(ty_id, storage_class); try cg.module.sections.globals.emit(cg.module.gpa, .OpVariable, .{ .id_result_type = ptr_ty_id, @@ -280,6 +281,27 @@ pub fn genNav(cg: *CodeGen, do_codegen: bool) Error!void { .storage_class = storage_class, }); + switch (target.os.tag) { + .vulkan, .opengl => { + if (ty.zigTypeTag(zcu) == .@"struct") { + switch (storage_class) { + .uniform, .push_constant => try cg.module.decorate(ty_id, .block), + else => {}, + } + } + + switch (ip.indexToKey(ty.toIntern())) { + .func_type, .opaque_type => {}, + else => { + try cg.module.decorate(ptr_ty_id, .{ + .array_stride = .{ .array_stride = @intCast(ty.abiSize(zcu)) }, + }); + }, + } + }, + else => {}, + } + if (std.meta.stringToEnum(spec.BuiltIn, nav.fqn.toSlice(ip))) |builtin| { try cg.module.decorate(result_id, .{ .built_in = .{ .built_in = builtin } }); } @@ -290,18 +312,20 @@ pub fn genNav(cg: *CodeGen, do_codegen: bool) Error!void { .invocation_global => { const maybe_init_val: ?Value = switch (ip.indexToKey(val.toIntern())) { .func => unreachable, - .variable => |variable| Value.fromInterned(variable.init), + .variable => |variable| .fromInterned(variable.init), .@"extern" => null, else => val, }; try cg.module.declareDeclDeps(spv_decl_index, &.{}); - const ptr_ty_id = try cg.ptrType(ty, .function, .indirect); + const ty_id = try cg.resolveType(ty, .indirect); + const ptr_ty_id = try cg.module.ptrType(ty_id, .function); if (maybe_init_val) |init_val| { // TODO: Combine with resolveAnonDecl? - const initializer_proto_ty_id = try cg.functionType(.void, &.{}); + const void_ty_id = try cg.resolveType(.void, .direct); + const initializer_proto_ty_id = try cg.module.functionType(void_ty_id, &.{}); const initializer_id = cg.module.allocId(); try cg.prologue.emit(cg.module.gpa, .OpFunction, .{ @@ -406,7 +430,8 @@ fn resolveUav(cg: *CodeGen, val: InternPool.Index) !Id { const zcu = cg.module.zcu; const ty: Type = .fromInterned(zcu.intern_pool.typeOf(val)); - const decl_ptr_ty_id = try cg.ptrType(ty, cg.module.storageClass(.generic), .indirect); + const ty_id = try cg.resolveType(ty, .indirect); + const decl_ptr_ty_id = try cg.module.ptrType(ty_id, cg.module.storageClass(.generic)); const spv_decl_index = blk: { const entry = try cg.module.uav_link.getOrPut(cg.module.gpa, .{ val, .function }); @@ -454,7 +479,8 @@ fn resolveUav(cg: *CodeGen, val: InternPool.Index) !Id { cg.decl_deps.deinit(gpa); } - const initializer_proto_ty_id = try cg.functionType(.void, &.{}); + const void_ty_id = try cg.resolveType(.void, .direct); + const initializer_proto_ty_id = try cg.module.functionType(void_ty_id, &.{}); const initializer_id = cg.module.allocId(); try cg.prologue.emit(cg.module.gpa, .OpFunction, .{ @@ -469,7 +495,7 @@ fn resolveUav(cg: *CodeGen, val: InternPool.Index) !Id { }); cg.block_label = root_block_id; - const val_id = try cg.constant(ty, Value.fromInterned(val), .indirect); + const val_id = try cg.constant(ty, .fromInterned(val), .indirect); try cg.body.emit(cg.module.gpa, .OpStore, .{ .pointer = result_id, .object = val_id, @@ -484,7 +510,7 @@ fn resolveUav(cg: *CodeGen, val: InternPool.Index) !Id { try cg.module.debugNameFmt(initializer_id, "initializer of __anon_{d}", .{@intFromEnum(val)}); - const fn_decl_ptr_ty_id = try cg.ptrType(ty, .function, .indirect); + const fn_decl_ptr_ty_id = try cg.module.ptrType(ty_id, .function); try cg.module.sections.globals.emit(cg.module.gpa, .OpExtInst, .{ .id_result_type = fn_decl_ptr_ty_id, .id_result = result_id, @@ -533,44 +559,6 @@ fn beginSpvBlock(cg: *CodeGen, label: Id) !void { cg.block_label = label; } -/// SPIR-V requires enabling specific integer sizes through capabilities, and so if they are not enabled, we need -/// to emulate them in other instructions/types. This function returns, given an integer bit width (signed or unsigned, sign -/// included), the width of the underlying type which represents it, given the enabled features for the current target. -/// If the result is `null`, the largest type the target platform supports natively is not able to perform computations using -/// that size. In this case, multiple elements of the largest type should be used. -/// The backing type will be chosen as the smallest supported integer larger or equal to it in number of bits. -/// The result is valid to be used with OpTypeInt. -/// TODO: Should the result of this function be cached? -fn backingIntBits(cg: *CodeGen, bits: u16) struct { u16, bool } { - const target = cg.module.zcu.getTarget(); - - // The backend will never be asked to compiler a 0-bit integer, so we won't have to handle those in this function. - assert(bits != 0); - - if (target.cpu.has(.spirv, .arbitrary_precision_integers) and bits <= 32) { - return .{ bits, false }; - } - - // We require Int8 and Int16 capabilities and benefit Int64 when available. - // 32-bit integers are always supported (see spec, 2.16.1, Data rules). - const ints = [_]struct { bits: u16, enabled: bool }{ - .{ .bits = 8, .enabled = true }, - .{ .bits = 16, .enabled = true }, - .{ .bits = 32, .enabled = true }, - .{ - .bits = 64, - .enabled = target.cpu.has(.spirv, .int64) or target.cpu.arch == .spirv64, - }, - }; - - for (ints) |int| { - if (bits <= int.bits and int.enabled) return .{ int.bits, false }; - } - - // Big int - return .{ std.mem.alignForward(u16, bits, big_int_bits), true }; -} - /// Return the amount of bits in the largest supported integer type. This is either 32 (always supported), or 64 (if /// the Int64 capability is enabled). /// Note: The extension SPV_INTEL_arbitrary_precision_integers allows any integer size (at least up to 32 bits). @@ -632,7 +620,7 @@ fn arithmeticTypeInfo(cg: *CodeGen, ty: Type) ArithmeticTypeInfo { return switch (scalar_ty.zigTypeTag(zcu)) { .bool => .{ .bits = 1, // Doesn't matter for this class. - .backing_bits = cg.backingIntBits(1).@"0", + .backing_bits = cg.module.backingIntBits(1).@"0", .vector_len = vector_len, .signedness = .unsigned, // Technically, but doesn't matter for this class. .class = .bool, @@ -647,7 +635,7 @@ fn arithmeticTypeInfo(cg: *CodeGen, ty: Type) ArithmeticTypeInfo { .int => blk: { const int_info = scalar_ty.intInfo(zcu); // TODO: Maybe it's useful to also return this value. - const backing_bits, const big_int = cg.backingIntBits(int_info.bits); + const backing_bits, const big_int = cg.module.backingIntBits(int_info.bits); break :blk .{ .bits = int_info.bits, .backing_bits = backing_bits, @@ -711,7 +699,7 @@ fn constInt(cg: *CodeGen, ty: Type, value: anytype) !Id { const scalar_ty = ty.scalarType(zcu); const int_info = scalar_ty.intInfo(zcu); // Use backing bits so that negatives are sign extended - const backing_bits, const big_int = cg.backingIntBits(int_info.bits); + const backing_bits, const big_int = cg.module.backingIntBits(int_info.bits); assert(backing_bits != 0); // u0 is comptime const result_ty_id = try cg.resolveType(scalar_ty, .indirect); @@ -922,8 +910,8 @@ fn constant(cg: *CodeGen, ty: Type, val: Value, repr: Repr) Error!Id { }, .ptr => return cg.constantPtr(val), .slice => |slice| { - const ptr_id = try cg.constantPtr(Value.fromInterned(slice.ptr)); - const len_id = try cg.constant(.usize, Value.fromInterned(slice.len), .indirect); + const ptr_id = try cg.constantPtr(.fromInterned(slice.ptr)); + const len_id = try cg.constant(.usize, .fromInterned(slice.len), .indirect); const comp_ty_id = try cg.resolveType(ty, .direct); return try cg.constructComposite(comp_ty_id, &.{ ptr_id, len_id }); }, @@ -977,11 +965,11 @@ fn constant(cg: *CodeGen, ty: Type, val: Value, repr: Repr) Error!Id { }, .elems => |elems| { for (constituents, elems) |*constituent, elem| { - constituent.* = try cg.constant(elem_ty, Value.fromInterned(elem), child_repr); + constituent.* = try cg.constant(elem_ty, .fromInterned(elem), child_repr); } }, .repeated_elem => |elem| { - @memset(constituents, try cg.constant(elem_ty, Value.fromInterned(elem), child_repr)); + @memset(constituents, try cg.constant(elem_ty, .fromInterned(elem), child_repr)); }, } @@ -995,7 +983,7 @@ fn constant(cg: *CodeGen, ty: Type, val: Value, repr: Repr) Error!Id { // TODO: composite int // TODO: endianness const bits: u16 = @intCast(ty.bitSize(zcu)); - const bytes = std.mem.alignForward(u16, cg.backingIntBits(bits).@"0", 8) / 8; + const bytes = std.mem.alignForward(u16, cg.module.backingIntBits(bits).@"0", 8) / 8; var limbs: [8]u8 = undefined; @memset(&limbs, 0); val.writeToPackedMemory(ty, pt, limbs[0..bytes], 0) catch unreachable; @@ -1035,13 +1023,13 @@ fn constant(cg: *CodeGen, ty: Type, val: Value, repr: Repr) Error!Id { if (un.tag == .none) { assert(ty.containerLayout(zcu) == .@"packed"); // TODO const int_ty = try pt.intType(.unsigned, @intCast(ty.bitSize(zcu))); - return try cg.constant(int_ty, Value.fromInterned(un.val), .direct); + return try cg.constInt(int_ty, Value.toUnsignedInt(.fromInterned(un.val), zcu)); } - const active_field = ty.unionTagFieldIndex(Value.fromInterned(un.tag), zcu).?; + const active_field = ty.unionTagFieldIndex(.fromInterned(un.tag), zcu).?; const union_obj = zcu.typeToUnion(ty).?; const field_ty: Type = .fromInterned(union_obj.field_types.get(ip)[active_field]); const payload = if (field_ty.hasRuntimeBitsIgnoreComptime(zcu)) - try cg.constant(field_ty, Value.fromInterned(un.val), .direct) + try cg.constant(field_ty, .fromInterned(un.val), .direct) else null; return try cg.unionInit(ty, active_field, payload); @@ -1084,10 +1072,11 @@ fn derivePtr(cg: *CodeGen, derivation: Value.PointerDeriveStep) !Id { // that is not implemented by Mesa yet. Therefore, just generate it // as a runtime operation. const result_ptr_id = cg.module.allocId(); + const value_id = try cg.constInt(.usize, int.addr); try cg.body.emit(cg.module.gpa, .OpConvertUToPtr, .{ .id_result_type = result_ty_id, .id_result = result_ptr_id, - .integer_value = try cg.constant(.usize, try pt.intValue(.usize, int.addr), .direct), + .integer_value = value_id, }); return result_ptr_id; }, @@ -1174,7 +1163,8 @@ fn constantUavRef( // Uav refs are always generic. assert(ty.ptrAddressSpace(zcu) == .generic); - const decl_ptr_ty_id = try cg.ptrType(uav_ty, .generic, .indirect); + const uav_ty_id = try cg.resolveType(uav_ty, .indirect); + const decl_ptr_ty_id = try cg.module.ptrType(uav_ty_id, .generic); const ptr_id = try cg.resolveUav(uav.val); if (decl_ptr_ty_id != ty_id) { @@ -1228,7 +1218,8 @@ fn constantNavRef(cg: *CodeGen, ty: Type, nav_index: InternPool.Nav.Index) !Id { const storage_class = cg.module.storageClass(nav.getAddrspace()); try cg.addFunctionDep(spv_decl_index, storage_class); - const decl_ptr_ty_id = try cg.ptrType(nav_ty, storage_class, .indirect); + const nav_ty_id = try cg.resolveType(nav_ty, .indirect); + const decl_ptr_ty_id = try cg.module.ptrType(nav_ty_id, storage_class); const ptr_id = switch (storage_class) { .generic => try cg.castToGeneric(decl_ptr_ty_id, decl_id), @@ -1260,104 +1251,6 @@ fn resolveTypeName(cg: *CodeGen, ty: Type) ![]const u8 { return try aw.toOwnedSlice(); } -/// Create an integer type suitable for storing at least 'bits' bits. -/// 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 Module.intType. -fn intType(cg: *CodeGen, signedness: std.builtin.Signedness, bits: u16) !Id { - const target = cg.module.zcu.getTarget(); - - const backing_bits, const big_int = cg.backingIntBits(bits); - if (big_int) { - if (backing_bits > 64) { - return cg.fail("composite integers larger than 64bit aren't supported", .{}); - } - const int_ty = try cg.resolveType(.u32, .direct); - return cg.arrayType(backing_bits / big_int_bits, int_ty); - } - - return switch (target.os.tag) { - // Kernel only supports unsigned ints. - .opencl, .amdhsa => return cg.module.intType(.unsigned, backing_bits), - else => cg.module.intType(signedness, backing_bits), - }; -} - -fn arrayType(cg: *CodeGen, len: u32, child_ty: Id) !Id { - const len_id = try cg.constInt(.u32, len); - return cg.module.arrayType(len_id, child_ty); -} - -fn ptrType(cg: *CodeGen, child_ty: Type, storage_class: StorageClass, child_repr: Repr) !Id { - const gpa = cg.module.gpa; - const zcu = cg.module.zcu; - const ip = &zcu.intern_pool; - const target = cg.module.zcu.getTarget(); - - const child_ty_id = try cg.resolveType(child_ty, child_repr); - const key = .{ child_ty_id, storage_class }; - const entry = try cg.module.ptr_types.getOrPut(gpa, key); - if (entry.found_existing) { - const fwd_id = entry.value_ptr.ty_id; - if (!entry.value_ptr.fwd_emitted) { - try cg.module.sections.globals.emit(cg.module.gpa, .OpTypeForwardPointer, .{ - .pointer_type = fwd_id, - .storage_class = storage_class, - }); - entry.value_ptr.fwd_emitted = true; - } - return fwd_id; - } - - const result_id = cg.module.allocId(); - entry.value_ptr.* = .{ - .ty_id = result_id, - .fwd_emitted = false, - }; - - switch (target.os.tag) { - .vulkan, .opengl => { - if (child_ty.zigTypeTag(zcu) == .@"struct") { - switch (storage_class) { - .uniform, .push_constant => try cg.module.decorate(child_ty_id, .block), - else => {}, - } - } - - switch (ip.indexToKey(child_ty.toIntern())) { - .func_type, .opaque_type => {}, - else => { - try cg.module.decorate(result_id, .{ .array_stride = .{ .array_stride = @intCast(child_ty.abiSize(zcu)) } }); - }, - } - }, - else => {}, - } - - try cg.module.sections.globals.emit(cg.module.gpa, .OpTypePointer, .{ - .id_result = result_id, - .storage_class = storage_class, - .type = child_ty_id, - }); - - cg.module.ptr_types.getPtr(key).?.fwd_emitted = true; - - return result_id; -} - -fn functionType(cg: *CodeGen, return_ty: Type, param_types: []const Type) !Id { - const gpa = cg.module.gpa; - const return_ty_id = try cg.resolveFnReturnType(return_ty); - const param_ids = try gpa.alloc(Id, param_types.len); - defer gpa.free(param_ids); - - for (param_types, param_ids) |param_ty, *param_id| { - param_id.* = try cg.resolveType(param_ty, .direct); - } - - return cg.module.functionType(return_ty_id, param_ids); -} - /// Generate a union type. Union types are always generated with the /// most aligned field active. If the tag alignment is greater /// than that of the payload, a regular union (non-packed, with both tag and @@ -1383,7 +1276,7 @@ fn resolveUnionType(cg: *CodeGen, ty: Type) !Id { const union_obj = zcu.typeToUnion(ty).?; if (union_obj.flagsUnordered(ip).layout == .@"packed") { - return try cg.intType(.unsigned, @intCast(ty.bitSize(zcu))); + return try cg.module.intType(.unsigned, @intCast(ty.bitSize(zcu))); } const layout = cg.unionLayout(ty); @@ -1410,13 +1303,15 @@ fn resolveUnionType(cg: *CodeGen, ty: Type) !Id { } if (layout.payload_padding_size != 0) { - const payload_padding_ty_id = try cg.arrayType(@intCast(layout.payload_padding_size), u8_ty_id); + const len_id = try cg.constInt(.u32, layout.payload_padding_size); + const payload_padding_ty_id = try cg.module.arrayType(len_id, u8_ty_id); member_types[layout.payload_padding_index] = payload_padding_ty_id; member_names[layout.payload_padding_index] = "(payload padding)"; } if (layout.padding_size != 0) { - const padding_ty_id = try cg.arrayType(@intCast(layout.padding_size), u8_ty_id); + const len_id = try cg.constInt(.u32, layout.padding_size); + const padding_ty_id = try cg.module.arrayType(len_id, u8_ty_id); member_types[layout.padding_index] = padding_ty_id; member_names[layout.padding_index] = "(padding)"; } @@ -1479,7 +1374,7 @@ fn resolveType(cg: *CodeGen, ty: Type, repr: Repr) Error!Id { assert(repr == .indirect); return try cg.module.opaqueType("u0"); } - return try cg.intType(int_info.signedness, int_info.bits); + return try cg.module.intType(int_info.signedness, int_info.bits); }, .@"enum" => return try cg.resolveType(ty.intTagType(zcu), repr), .float => { @@ -1519,9 +1414,11 @@ fn resolveType(cg: *CodeGen, ty: Type, repr: Repr) Error!Id { // In this case, we have an array of a non-zero sized type. In this case, // generate an array of 1 element instead, so that ptr_elem_ptr instructions // can be lowered to ptrAccessChain instead of manually performing the math. - return try cg.arrayType(1, elem_ty_id); + const len_id = try cg.constInt(.u32, 1); + return try cg.module.arrayType(len_id, elem_ty_id); } else { - const result_id = try cg.arrayType(total_len, elem_ty_id); + const total_len_id = try cg.constInt(.u32, total_len); + const result_id = try cg.module.arrayType(total_len_id, elem_ty_id); switch (target.os.tag) { .vulkan, .opengl => { try cg.module.decorate(result_id, .{ @@ -1540,7 +1437,8 @@ fn resolveType(cg: *CodeGen, ty: Type, repr: Repr) Error!Id { const elem_ty_id = try cg.resolveType(elem_ty, repr); const len = ty.vectorLen(zcu); if (cg.isSpvVector(ty)) return try cg.module.vectorType(len, elem_ty_id); - return try cg.arrayType(len, elem_ty_id); + const len_id = try cg.constInt(.u32, len); + return try cg.module.arrayType(len_id, elem_ty_id); }, .@"fn" => switch (repr) { .direct => { @@ -1582,8 +1480,9 @@ fn resolveType(cg: *CodeGen, ty: Type, repr: Repr) Error!Id { const ptr_info = ty.ptrInfo(zcu); const child_ty: Type = .fromInterned(ptr_info.child); + const child_ty_id = try cg.resolveType(child_ty, .indirect); const storage_class = cg.module.storageClass(ptr_info.flags.address_space); - const ptr_ty_id = try cg.ptrType(child_ty, storage_class, .indirect); + const ptr_ty_id = try cg.module.ptrType(child_ty_id, storage_class); if (ptr_info.flags.size != .slice) { return ptr_ty_id; @@ -2142,7 +2041,7 @@ fn buildConvert(cg: *CodeGen, dst_ty: Type, src: Temporary) !Temporary { for (0..ops) |i| { try cg.body.emitRaw(cg.module.gpa, opcode, 3); - cg.body.writeOperand(spec.Id, op_result_ty_id); + cg.body.writeOperand(Id, op_result_ty_id); cg.body.writeOperand(Id, results.at(i)); cg.body.writeOperand(Id, op_src.at(i)); } @@ -2277,7 +2176,7 @@ fn buildCmp(cg: *CodeGen, pred: CmpPredicate, lhs: Temporary, rhs: Temporary) !T for (0..ops) |i| { try cg.body.emitRaw(cg.module.gpa, opcode, 4); - cg.body.writeOperand(spec.Id, op_result_ty_id); + cg.body.writeOperand(Id, op_result_ty_id); cg.body.writeOperand(Id, results.at(i)); cg.body.writeOperand(Id, op_lhs.at(i)); cg.body.writeOperand(Id, op_rhs.at(i)); @@ -2331,7 +2230,7 @@ fn buildUnary(cg: *CodeGen, op: UnaryOp, operand: Temporary) !Temporary { }) |opcode| { for (0..ops) |i| { try cg.body.emitRaw(cg.module.gpa, opcode, 3); - cg.body.writeOperand(spec.Id, op_result_ty_id); + cg.body.writeOperand(Id, op_result_ty_id); cg.body.writeOperand(Id, results.at(i)); cg.body.writeOperand(Id, op_operand.at(i)); } @@ -2472,7 +2371,7 @@ fn buildBinary(cg: *CodeGen, op: BinaryOp, lhs: Temporary, rhs: Temporary) !Temp }) |opcode| { for (0..ops) |i| { try cg.body.emitRaw(cg.module.gpa, opcode, 4); - cg.body.writeOperand(spec.Id, op_result_ty_id); + cg.body.writeOperand(Id, op_result_ty_id); cg.body.writeOperand(Id, results.at(i)); cg.body.writeOperand(Id, op_lhs.at(i)); cg.body.writeOperand(Id, op_rhs.at(i)); @@ -2591,7 +2490,7 @@ fn buildWideMul( const op_result = cg.module.allocId(); try cg.body.emitRaw(cg.module.gpa, opcode, 4); - cg.body.writeOperand(spec.Id, op_result_ty_id); + cg.body.writeOperand(Id, op_result_ty_id); cg.body.writeOperand(Id, op_result); cg.body.writeOperand(Id, lhs_op.at(i)); cg.body.writeOperand(Id, rhs_op.at(i)); @@ -2664,30 +2563,27 @@ fn generateTestEntryPoint( const kernel_id = cg.module.declPtr(spv_decl_index).result_id; - var decl_deps = std.ArrayList(Module.Decl.Index).init(gpa); - defer decl_deps.deinit(); - try decl_deps.append(spv_decl_index); - const section = &cg.module.sections.functions; const p_error_id = cg.module.allocId(); switch (target.os.tag) { .opencl, .amdhsa => { - const kernel_proto_ty_id = try cg.functionType(.void, &.{ptr_anyerror_ty}); + const void_ty_id = try cg.resolveType(.void, .direct); + const kernel_proto_ty_id = try cg.module.functionType(void_ty_id, &.{ptr_anyerror_ty_id}); - try section.emit(cg.module.gpa, .OpFunction, .{ + try section.emit(gpa, .OpFunction, .{ .id_result_type = try cg.resolveType(.void, .direct), .id_result = kernel_id, .function_control = .{}, .function_type = kernel_proto_ty_id, }); - try section.emit(cg.module.gpa, .OpFunctionParameter, .{ + try section.emit(gpa, .OpFunctionParameter, .{ .id_result_type = ptr_anyerror_ty_id, .id_result = p_error_id, }); - try section.emit(cg.module.gpa, .OpLabel, .{ + try section.emit(gpa, .OpLabel, .{ .id_result = cg.module.allocId(), }); }, @@ -2706,14 +2602,14 @@ fn generateTestEntryPoint( try cg.module.decorateMember(buffer_struct_ty_id, 0, .{ .offset = .{ .byte_offset = 0 } }); const ptr_buffer_struct_ty_id = cg.module.allocId(); - try cg.module.sections.globals.emit(cg.module.gpa, .OpTypePointer, .{ + try cg.module.sections.globals.emit(gpa, .OpTypePointer, .{ .id_result = ptr_buffer_struct_ty_id, .storage_class = cg.module.storageClass(.global), .type = buffer_struct_ty_id, }); const buffer_struct_id = cg.module.declPtr(spv_err_decl_index).result_id; - try cg.module.sections.globals.emit(cg.module.gpa, .OpVariable, .{ + try cg.module.sections.globals.emit(gpa, .OpVariable, .{ .id_result_type = ptr_buffer_struct_ty_id, .id_result = buffer_struct_id, .storage_class = cg.module.storageClass(.global), @@ -2724,7 +2620,7 @@ fn generateTestEntryPoint( cg.module.error_buffer = spv_err_decl_index; } - try cg.module.sections.execution_modes.emit(cg.module.gpa, .OpExecutionMode, .{ + try cg.module.sections.execution_modes.emit(gpa, .OpExecutionMode, .{ .entry_point = kernel_id, .mode = .{ .local_size = .{ .x_size = 1, @@ -2733,23 +2629,24 @@ fn generateTestEntryPoint( } }, }); - const kernel_proto_ty_id = try cg.functionType(.void, &.{}); - try section.emit(cg.module.gpa, .OpFunction, .{ + const void_ty_id = try cg.resolveType(.void, .direct); + const kernel_proto_ty_id = try cg.module.functionType(void_ty_id, &.{}); + try section.emit(gpa, .OpFunction, .{ .id_result_type = try cg.resolveType(.void, .direct), .id_result = kernel_id, .function_control = .{}, .function_type = kernel_proto_ty_id, }); - try section.emit(cg.module.gpa, .OpLabel, .{ + try section.emit(gpa, .OpLabel, .{ .id_result = cg.module.allocId(), }); const spv_err_decl_index = cg.module.error_buffer.?; const buffer_id = cg.module.declPtr(spv_err_decl_index).result_id; - try decl_deps.append(spv_err_decl_index); + try cg.decl_deps.put(gpa, spv_err_decl_index, {}); const zero_id = try cg.constInt(.u32, 0); - try section.emit(cg.module.gpa, .OpInBoundsAccessChain, .{ + try section.emit(gpa, .OpInBoundsAccessChain, .{ .id_result_type = ptr_anyerror_ty_id, .id_result = p_error_id, .base = buffer_id, @@ -2760,25 +2657,25 @@ fn generateTestEntryPoint( } const error_id = cg.module.allocId(); - try section.emit(cg.module.gpa, .OpFunctionCall, .{ + try section.emit(gpa, .OpFunctionCall, .{ .id_result_type = anyerror_ty_id, .id_result = error_id, .function = test_id, }); // Note: Convert to direct not required. - try section.emit(cg.module.gpa, .OpStore, .{ + try section.emit(gpa, .OpStore, .{ .pointer = p_error_id, .object = error_id, .memory_access = .{ .aligned = .{ .literal_integer = @intCast(Type.abiAlignment(.anyerror, zcu).toByteUnits().?) }, }, }); - try section.emit(cg.module.gpa, .OpReturn, {}); - try section.emit(cg.module.gpa, .OpFunctionEnd, {}); + try section.emit(gpa, .OpReturn, {}); + try section.emit(gpa, .OpFunctionEnd, {}); // Just generate a quick other name because the intel runtime crashes when the entry- // point name is the same as a different OpName. - const test_name = try std.fmt.allocPrint(gpa, "test {s}", .{name}); + const test_name = try std.fmt.allocPrint(cg.module.arena, "test {s}", .{name}); const execution_mode: spec.ExecutionModel = switch (target.os.tag) { .vulkan, .opengl => .gl_compute, @@ -2786,7 +2683,6 @@ fn generateTestEntryPoint( else => unreachable, }; - try cg.module.declareDeclDeps(spv_decl_index, decl_deps.items); try cg.module.declareEntryPoint(spv_decl_index, test_name, execution_mode, null); } @@ -3760,10 +3656,10 @@ fn airReduce(cg: *CodeGen, inst: Air.Inst.Index) !?Id { result_id = cg.module.allocId(); try cg.body.emitRaw(cg.module.gpa, opcode, 4); - cg.body.writeOperand(spec.Id, scalar_ty_id); - cg.body.writeOperand(spec.Id, result_id); - cg.body.writeOperand(spec.Id, lhs); - cg.body.writeOperand(spec.Id, rhs); + cg.body.writeOperand(Id, scalar_ty_id); + cg.body.writeOperand(Id, result_id); + cg.body.writeOperand(Id, lhs); + cg.body.writeOperand(Id, rhs); } return result_id; @@ -4189,7 +4085,7 @@ fn bitCast( break :blk result_id; } - const dst_ptr_ty_id = try cg.ptrType(dst_ty, .function, .indirect); + const dst_ptr_ty_id = try cg.module.ptrType(dst_ty_id, .function); const tmp_id = try cg.alloc(src_ty, .{ .storage_class = .function }); try cg.store(src_ty, tmp_id, src_id, .{}); @@ -4594,7 +4490,8 @@ fn ptrElemPtr(cg: *CodeGen, ptr_ty: Type, ptr_id: Id, index_id: Id) !Id { const zcu = cg.module.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. - const elem_ptr_ty_id = try cg.ptrType(elem_ty, cg.module.storageClass(ptr_ty.ptrAddressSpace(zcu)), .indirect); + const elem_ty_id = try cg.resolveType(elem_ty, .indirect); + const elem_ptr_ty_id = try cg.module.ptrType(elem_ty_id, cg.module.storageClass(ptr_ty.ptrAddressSpace(zcu))); if (ptr_ty.isSinglePointer(zcu)) { // Pointer-to-array. In this case, the resulting pointer is not of the same type // as the ptr_ty (we want a *T, not a *[N]T), and hence we need to use accessChain. @@ -4637,8 +4534,10 @@ fn airArrayElemVal(cg: *CodeGen, inst: Air.Inst.Index) !?Id { const is_vector = array_ty.isVector(zcu); const elem_repr: Repr = if (is_vector) .direct else .indirect; - const ptr_array_ty_id = try cg.ptrType(array_ty, .function, .direct); - const ptr_elem_ty_id = try cg.ptrType(elem_ty, .function, elem_repr); + const array_ty_id = try cg.resolveType(array_ty, .direct); + const elem_ty_id = try cg.resolveType(elem_ty, elem_repr); + const ptr_array_ty_id = try cg.module.ptrType(array_ty_id, .function); + const ptr_elem_ty_id = try cg.module.ptrType(elem_ty_id, .function); const tmp_id = cg.module.allocId(); try cg.prologue.emit(cg.module.gpa, .OpVariable, .{ @@ -4692,8 +4591,9 @@ fn airVectorStoreElem(cg: *CodeGen, inst: Air.Inst.Index) !void { const vector_ty = vector_ptr_ty.childType(zcu); const scalar_ty = vector_ty.scalarType(zcu); + const scalar_ty_id = try cg.resolveType(scalar_ty, .indirect); const storage_class = cg.module.storageClass(vector_ptr_ty.ptrAddressSpace(zcu)); - const scalar_ptr_ty_id = try cg.ptrType(scalar_ty, storage_class, .indirect); + const scalar_ptr_ty_id = try cg.module.ptrType(scalar_ty_id, storage_class); const vector_ptr = try cg.resolve(data.vector_ptr); const index = try cg.resolve(extra.lhs); @@ -4715,7 +4615,8 @@ fn airSetUnionTag(cg: *CodeGen, inst: Air.Inst.Index) !void { if (layout.tag_size == 0) return; const tag_ty = un_ty.unionTagTypeSafety(zcu).?; - const tag_ptr_ty_id = try cg.ptrType(tag_ty, cg.module.storageClass(un_ptr_ty.ptrAddressSpace(zcu)), .indirect); + const tag_ty_id = try cg.resolveType(tag_ty, .indirect); + const tag_ptr_ty_id = try cg.module.ptrType(tag_ty_id, cg.module.storageClass(un_ptr_ty.ptrAddressSpace(zcu))); const union_ptr_id = try cg.resolve(bin_op.lhs); const new_tag_id = try cg.resolve(bin_op.rhs); @@ -4802,17 +4703,20 @@ fn unionInit( const tmp_id = try cg.alloc(ty, .{ .storage_class = .function }); if (layout.tag_size != 0) { - const tag_ptr_ty_id = try cg.ptrType(tag_ty, .function, .indirect); + const tag_ty_id = try cg.resolveType(tag_ty, .indirect); + const tag_ptr_ty_id = try cg.module.ptrType(tag_ty_id, .function); const ptr_id = try cg.accessChain(tag_ptr_ty_id, tmp_id, &.{@as(u32, @intCast(layout.tag_index))}); const tag_id = try cg.constInt(tag_ty, tag_int); try cg.store(tag_ty, ptr_id, tag_id, .{}); } if (payload_ty.hasRuntimeBitsIgnoreComptime(zcu)) { - const pl_ptr_ty_id = try cg.ptrType(layout.payload_ty, .function, .indirect); + const layout_payload_ty_id = try cg.resolveType(layout.payload_ty, .indirect); + const pl_ptr_ty_id = try cg.module.ptrType(layout_payload_ty_id, .function); const pl_ptr_id = try cg.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 cg.ptrType(payload_ty, .function, .indirect); + const payload_ty_id = try cg.resolveType(payload_ty, .indirect); + const active_pl_ptr_ty_id = try cg.module.ptrType(payload_ty_id, .function); const active_pl_ptr_id = cg.module.allocId(); try cg.body.emit(cg.module.gpa, .OpBitcast, .{ .id_result_type = active_pl_ptr_ty_id, @@ -4876,7 +4780,7 @@ fn airStructFieldVal(cg: *CodeGen, inst: Air.Inst.Index) !?Id { const mask_id = try cg.constInt(object_ty, (@as(u64, 1) << @as(u6, @intCast(field_bit_size))) - 1); const masked = try cg.buildBinary(.bit_and, shift, .{ .ty = object_ty, .value = .{ .singleton = mask_id } }); const result_id = blk: { - if (cg.backingIntBits(field_bit_size).@"0" == cg.backingIntBits(@intCast(object_ty.bitSize(zcu))).@"0") + if (cg.module.backingIntBits(field_bit_size).@"0" == cg.module.backingIntBits(@intCast(object_ty.bitSize(zcu))).@"0") break :blk try cg.bitCast(field_int_ty, object_ty, try masked.materialize(cg)); const trunc = try cg.buildConvert(field_int_ty, masked); break :blk try trunc.materialize(cg); @@ -4900,7 +4804,7 @@ fn airStructFieldVal(cg: *CodeGen, inst: Air.Inst.Index) !?Id { .{ .ty = backing_int_ty, .value = .{ .singleton = mask_id } }, ); const result_id = blk: { - if (cg.backingIntBits(field_bit_size).@"0" == cg.backingIntBits(@intCast(backing_int_ty.bitSize(zcu))).@"0") + if (cg.module.backingIntBits(field_bit_size).@"0" == cg.module.backingIntBits(@intCast(backing_int_ty.bitSize(zcu))).@"0") break :blk try cg.bitCast(int_ty, backing_int_ty, try masked.materialize(cg)); const trunc = try cg.buildConvert(int_ty, masked); break :blk try trunc.materialize(cg); @@ -4917,10 +4821,12 @@ fn airStructFieldVal(cg: *CodeGen, inst: Air.Inst.Index) !?Id { const tmp_id = try cg.alloc(object_ty, .{ .storage_class = .function }); try cg.store(object_ty, tmp_id, object_id, .{}); - const pl_ptr_ty_id = try cg.ptrType(layout.payload_ty, .function, .indirect); + const layout_payload_ty_id = try cg.resolveType(layout.payload_ty, .indirect); + const pl_ptr_ty_id = try cg.module.ptrType(layout_payload_ty_id, .function); const pl_ptr_id = try cg.accessChain(pl_ptr_ty_id, tmp_id, &.{layout.payload_index}); - const active_pl_ptr_ty_id = try cg.ptrType(field_ty, .function, .indirect); + const field_ty_id = try cg.resolveType(field_ty, .indirect); + const active_pl_ptr_ty_id = try cg.module.ptrType(field_ty_id, .function); const active_pl_ptr_id = cg.module.allocId(); try cg.body.emit(cg.module.gpa, .OpBitcast, .{ .id_result_type = active_pl_ptr_ty_id, @@ -4997,7 +4903,8 @@ fn structFieldPtr( } const storage_class = cg.module.storageClass(object_ptr_ty.ptrAddressSpace(zcu)); - const pl_ptr_ty_id = try cg.ptrType(layout.payload_ty, storage_class, .indirect); + const layout_payload_ty_id = try cg.resolveType(layout.payload_ty, .indirect); + const pl_ptr_ty_id = try cg.module.ptrType(layout_payload_ty_id, storage_class); const pl_ptr_id = blk: { if (object_ty.containerLayout(zcu) == .@"packed") break :blk object_ptr; break :blk try cg.accessChain(pl_ptr_ty_id, object_ptr, &.{layout.payload_index}); @@ -5041,7 +4948,8 @@ fn alloc( options: AllocOptions, ) !Id { const target = cg.module.zcu.getTarget(); - const ptr_fn_ty_id = try cg.ptrType(ty, .function, .indirect); + const ty_id = try cg.resolveType(ty, .indirect); + const ptr_fn_ty_id = try cg.module.ptrType(ty_id, .function); // 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. @@ -5060,7 +4968,7 @@ fn alloc( switch (options.storage_class) { .generic => { - const ptr_gn_ty_id = try cg.ptrType(ty, .generic, .indirect); + const ptr_gn_ty_id = try cg.module.ptrType(ty_id, .generic); // Convert to a generic pointer return cg.castToGeneric(ptr_gn_ty_id, var_id); }, @@ -5093,8 +5001,8 @@ fn structuredNextBlock(cg: *CodeGen, incoming: []const ControlFlow.Structured.Bl const result_id = cg.module.allocId(); const block_id_ty_id = try cg.resolveType(.u32, .direct); try cg.body.emitRaw(cg.module.gpa, .OpPhi, @intCast(2 + incoming.len * 2)); // result type + result + variable/parent... - cg.body.writeOperand(spec.Id, block_id_ty_id); - cg.body.writeOperand(spec.Id, result_id); + cg.body.writeOperand(Id, block_id_ty_id); + cg.body.writeOperand(Id, result_id); for (incoming) |incoming_block| { cg.body.writeOperand(spec.PairIdRefIdRef, .{ incoming_block.next_block, incoming_block.src_label }); @@ -5285,8 +5193,8 @@ fn lowerBlock(cg: *CodeGen, inst: Air.Inst.Index, body: []const Air.Inst.Index) // result type + result + variable/parent... 2 + @as(u16, @intCast(block.incoming_blocks.items.len * 2)), ); - cg.body.writeOperand(spec.Id, result_type_id); - cg.body.writeOperand(spec.Id, result_id); + cg.body.writeOperand(Id, result_type_id); + cg.body.writeOperand(Id, result_id); for (block.incoming_blocks.items) |incoming| { cg.body.writeOperand( @@ -5793,7 +5701,8 @@ fn airIsNull(cg: *CodeGen, inst: Air.Inst.Index, is_pointer: bool, pred: enum { if (is_pointer) { if (payload_ty.hasRuntimeBitsIgnoreComptime(zcu)) { const storage_class = cg.module.storageClass(operand_ty.ptrAddressSpace(zcu)); - const bool_ptr_ty_id = try cg.ptrType(.bool, storage_class, .indirect); + const bool_indirect_ty_id = try cg.resolveType(.bool, .indirect); + const bool_ptr_ty_id = try cg.module.ptrType(bool_indirect_ty_id, storage_class); const tag_ptr_id = try cg.accessChain(bool_ptr_ty_id, operand_id, &.{1}); break :blk try cg.load(.bool, tag_ptr_id, .{}); } @@ -5939,14 +5848,14 @@ fn airSwitchBr(cg: *CodeGen, inst: Air.Inst.Index) !void { .bool, .error_set => 1, .int => blk: { const bits = cond_ty.intInfo(zcu).bits; - const backing_bits, const big_int = cg.backingIntBits(bits); + const backing_bits, const big_int = cg.module.backingIntBits(bits); if (big_int) return cg.todo("implement composite int switch", .{}); break :blk if (backing_bits <= 32) 1 else 2; }, .@"enum" => blk: { const int_ty = cond_ty.intTagType(zcu); const int_info = int_ty.intInfo(zcu); - const backing_bits, const big_int = cg.backingIntBits(int_info.bits); + const backing_bits, const big_int = cg.module.backingIntBits(int_info.bits); if (big_int) return cg.todo("implement composite int switch", .{}); break :blk if (backing_bits <= 32) 1 else 2; }, @@ -6298,7 +6207,7 @@ fn airCall(cg: *CodeGen, inst: Air.Inst.Index, modifier: std.builtin.CallModifie const callee_id = try cg.resolve(pl_op.operand); comptime assert(zig_call_abi_ver == 3); - const params = try gpa.alloc(spec.Id, args.len); + const params = try gpa.alloc(Id, args.len); defer gpa.free(params); var n_params: usize = 0; for (args) |arg| { @@ -6327,50 +6236,49 @@ fn airCall(cg: *CodeGen, inst: Air.Inst.Index, modifier: std.builtin.CallModifie return result_id; } -fn builtin3D(cg: *CodeGen, result_ty: Type, builtin: spec.BuiltIn, dimension: u32, out_of_range_value: anytype) !Id { - if (dimension >= 3) { - return try cg.constInt(result_ty, out_of_range_value); - } - const vec_ty = try cg.pt.vectorType(.{ - .len = 3, - .child = result_ty.toIntern(), - }); - const ptr_ty_id = try cg.ptrType(vec_ty, .input, .indirect); - const spv_decl_index = try cg.module.builtin(ptr_ty_id, builtin); +fn builtin3D( + cg: *CodeGen, + result_ty: Type, + builtin: spec.BuiltIn, + dimension: u32, + out_of_range_value: anytype, +) !Id { + if (dimension >= 3) return try cg.constInt(result_ty, out_of_range_value); + const u32_ty_id = try cg.module.intType(.unsigned, 32); + const vec_ty_id = try cg.module.vectorType(3, u32_ty_id); + const ptr_ty_id = try cg.module.ptrType(vec_ty_id, .input); + const spv_decl_index = try cg.module.builtin(ptr_ty_id, builtin, .input); try cg.decl_deps.put(cg.module.gpa, spv_decl_index, {}); - const ptr = cg.module.declPtr(spv_decl_index).result_id; - const vec = try cg.load(vec_ty, ptr, .{}); - return try cg.extractVectorComponent(result_ty, vec, dimension); + const ptr_id = cg.module.declPtr(spv_decl_index).result_id; + const vec_id = cg.module.allocId(); + try cg.body.emit(cg.module.gpa, .OpLoad, .{ + .id_result_type = vec_ty_id, + .id_result = vec_id, + .pointer = ptr_id, + }); + return try cg.extractVectorComponent(result_ty, vec_id, dimension); } fn airWorkItemId(cg: *CodeGen, inst: Air.Inst.Index) !?Id { if (cg.liveness.isUnused(inst)) return null; const pl_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].pl_op; const dimension = pl_op.payload; - const result_id = try cg.builtin3D(.u32, .local_invocation_id, dimension, 0); - const tmp: Temporary = .init(.u32, result_id); - const result = try cg.buildConvert(.u32, tmp); - return try result.materialize(cg); + return try cg.builtin3D(.u32, .local_invocation_id, dimension, 0); } +// TODO: this must be an OpConstant/OpSpec but even then the driver crashes. fn airWorkGroupSize(cg: *CodeGen, inst: Air.Inst.Index) !?Id { if (cg.liveness.isUnused(inst)) return null; const pl_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].pl_op; const dimension = pl_op.payload; - const result_id = try cg.builtin3D(.u32, .workgroup_size, dimension, 0); - const tmp: Temporary = .init(.u32, result_id); - const result = try cg.buildConvert(.u32, tmp); - return try result.materialize(cg); + return try cg.builtin3D(.u32, .workgroup_id, dimension, 0); } fn airWorkGroupId(cg: *CodeGen, inst: Air.Inst.Index) !?Id { if (cg.liveness.isUnused(inst)) return null; const pl_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].pl_op; const dimension = pl_op.payload; - const result_id = try cg.builtin3D(.u32, .workgroup_id, dimension, 0); - const tmp: Temporary = .init(.u32, result_id); - const result = try cg.buildConvert(.u32, tmp); - return try result.materialize(cg); + return try cg.builtin3D(.u32, .workgroup_id, dimension, 0); } fn typeOf(cg: *CodeGen, inst: Air.Inst.Ref) Type { diff --git a/src/arch/spirv/Module.zig b/src/arch/spirv/Module.zig index 68207fead8..2e60cbfd7c 100644 --- a/src/arch/spirv/Module.zig +++ b/src/arch/spirv/Module.zig @@ -35,10 +35,7 @@ entry_points: std.AutoArrayHashMapUnmanaged(Id, EntryPoint) = .empty, /// - It caches pointers by child-type. This is required because sometimes we rely on /// ID-equality for pointers, and pointers constructed via `ptrType()` aren't interned /// via the usual `intern_map` mechanism. -ptr_types: std.AutoHashMapUnmanaged( - struct { Id, spec.StorageClass }, - struct { ty_id: Id, fwd_emitted: bool }, -) = .{}, +ptr_types: std.AutoHashMapUnmanaged(struct { Id, spec.StorageClass }, Id) = .{}, /// For test declarations compiled for Vulkan target, we have to add a buffer. /// We only need to generate this once, this holds the link information related to that. error_buffer: ?Decl.Index = null, @@ -68,7 +65,7 @@ cache: struct { extensions: std.StringHashMapUnmanaged(void) = .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, + builtins: std.AutoHashMapUnmanaged(struct { spec.BuiltIn, spec.StorageClass }, Decl.Index) = .empty, strings: std.StringArrayHashMapUnmanaged(Id) = .empty, bool_const: [2]?Id = .{ null, null }, @@ -88,6 +85,8 @@ sections: struct { functions: Section = .{}, } = .{}, +pub const big_int_bits = 32; + /// Data can be lowered into in two basic representations: indirect, which is when /// a type is stored in memory, and direct, which is how a type is stored when its /// a direct SPIR-V value. @@ -241,10 +240,6 @@ pub fn deinit(module: *Module) void { module.decls.deinit(module.gpa); module.decl_deps.deinit(module.gpa); - - for (module.entry_points.values()) |ep| { - module.gpa.free(ep.name); - } module.entry_points.deinit(module.gpa); module.* = undefined; @@ -546,24 +541,68 @@ pub fn opaqueType(module: *Module, name: []const u8) !Id { return result_id; } +pub fn backingIntBits(module: *Module, bits: u16) struct { u16, bool } { + assert(bits != 0); + const target = module.zcu.getTarget(); + + if (target.cpu.has(.spirv, .arbitrary_precision_integers) and bits <= 32) { + return .{ bits, false }; + } + + // We require Int8 and Int16 capabilities and benefit Int64 when available. + // 32-bit integers are always supported (see spec, 2.16.1, Data rules). + const ints = [_]struct { bits: u16, enabled: bool }{ + .{ .bits = 8, .enabled = true }, + .{ .bits = 16, .enabled = true }, + .{ .bits = 32, .enabled = true }, + .{ + .bits = 64, + .enabled = target.cpu.has(.spirv, .int64) or target.cpu.arch == .spirv64, + }, + }; + + for (ints) |int| { + if (bits <= int.bits and int.enabled) return .{ int.bits, false }; + } + + // Big int + return .{ std.mem.alignForward(u16, bits, big_int_bits), true }; +} + pub fn intType(module: *Module, signedness: std.builtin.Signedness, bits: u16) !Id { assert(bits > 0); - const entry = try module.cache.int_types.getOrPut(module.gpa, .{ .signedness = signedness, .bits = bits }); + + const target = module.zcu.getTarget(); + const actual_signedness = switch (target.os.tag) { + // Kernel only supports unsigned ints. + .opencl, .amdhsa => .unsigned, + else => signedness, + }; + const backing_bits, const big_int = module.backingIntBits(bits); + if (big_int) { + // TODO: support composite integers larger than 64 bit + assert(backing_bits <= 64); + const u32_ty = try module.intType(.unsigned, 32); + const len_id = try module.constant(u32_ty, .{ .uint32 = backing_bits / big_int_bits }); + return module.arrayType(len_id, u32_ty); + } + + const entry = try module.cache.int_types.getOrPut(module.gpa, .{ .signedness = actual_signedness, .bits = backing_bits }); if (!entry.found_existing) { const result_id = module.allocId(); entry.value_ptr.* = result_id; try module.sections.globals.emit(module.gpa, .OpTypeInt, .{ .id_result = result_id, - .width = bits, - .signedness = switch (signedness) { + .width = backing_bits, + .signedness = switch (actual_signedness) { .signed => 1, .unsigned => 0, }, }); - switch (signedness) { - .signed => try module.debugNameFmt(result_id, "i{}", .{bits}), - .unsigned => try module.debugNameFmt(result_id, "u{}", .{bits}), + switch (actual_signedness) { + .signed => try module.debugNameFmt(result_id, "i{}", .{backing_bits}), + .unsigned => try module.debugNameFmt(result_id, "u{}", .{backing_bits}), } } return entry.value_ptr.*; @@ -612,6 +651,21 @@ pub fn arrayType(module: *Module, len_id: Id, child_ty_id: Id) !Id { return entry.value_ptr.*; } +pub fn ptrType(module: *Module, child_ty_id: Id, storage_class: spec.StorageClass) !Id { + const key = .{ child_ty_id, storage_class }; + const gop = try module.ptr_types.getOrPut(module.gpa, key); + if (!gop.found_existing) { + gop.value_ptr.* = module.allocId(); + try module.sections.globals.emit(module.gpa, .OpTypePointer, .{ + .id_result = gop.value_ptr.*, + .storage_class = storage_class, + .type = child_ty_id, + }); + return gop.value_ptr.*; + } + return gop.value_ptr.*; +} + pub fn structType( module: *Module, types: []const Id, @@ -683,16 +737,16 @@ pub fn functionType(module: *Module, return_ty_id: Id, param_type_ids: []const I } pub fn constant(module: *Module, ty_id: Id, value: spec.LiteralContextDependentNumber) !Id { - const entry = try module.cache.constants.getOrPut(module.gpa, .{ .ty = ty_id, .value = value }); - if (!entry.found_existing) { - entry.value_ptr.* = module.allocId(); + const gop = try module.cache.constants.getOrPut(module.gpa, .{ .ty = ty_id, .value = value }); + if (!gop.found_existing) { + gop.value_ptr.* = module.allocId(); try module.sections.globals.emit(module.gpa, .OpConstant, .{ .id_result_type = ty_id, - .id_result = entry.value_ptr.*, + .id_result = gop.value_ptr.*, .value = value, }); } - return entry.value_ptr.*; + return gop.value_ptr.*; } pub fn constBool(module: *Module, value: bool) !Id { @@ -716,23 +770,26 @@ pub fn constBool(module: *Module, value: bool) !Id { return result_id; } -/// Return a pointer to a builtin variable. `result_ty_id` must be a **pointer** -/// with storage class `.Input`. -pub fn builtin(module: *Module, result_ty_id: Id, spirv_builtin: spec.BuiltIn) !Decl.Index { - const entry = try module.cache.builtins.getOrPut(module.gpa, .{ result_ty_id, spirv_builtin }); - if (!entry.found_existing) { +pub fn builtin( + module: *Module, + result_ty_id: Id, + spirv_builtin: spec.BuiltIn, + storage_class: spec.StorageClass, +) !Decl.Index { + const gop = try module.cache.builtins.getOrPut(module.gpa, .{ spirv_builtin, storage_class }); + if (!gop.found_existing) { const decl_index = try module.allocDecl(.global); const result_id = module.declPtr(decl_index).result_id; - entry.value_ptr.* = decl_index; + gop.value_ptr.* = decl_index; try module.sections.globals.emit(module.gpa, .OpVariable, .{ .id_result_type = result_ty_id, .id_result = result_id, - .storage_class = .input, + .storage_class = storage_class, }); try module.decorate(result_id, .{ .built_in = .{ .built_in = spirv_builtin } }); try module.declareDeclDeps(decl_index, &.{}); } - return entry.value_ptr.*; + return gop.value_ptr.*; } pub fn constUndef(module: *Module, ty_id: Id) !Id { @@ -759,8 +816,8 @@ pub fn decorate( target: Id, decoration: spec.Decoration.Extended, ) !void { - const entry = try module.cache.decorations.getOrPut(module.gpa, .{ target, decoration }); - if (!entry.found_existing) { + const gop = try module.cache.decorations.getOrPut(module.gpa, .{ target, decoration }); + if (!gop.found_existing) { try module.sections.annotations.emit(module.gpa, .OpDecorate, .{ .target = target, .decoration = decoration, |