master
1const std = @import("std");
2const Allocator = std.mem.Allocator;
3const Target = std.Target;
4const Signedness = std.builtin.Signedness;
5const assert = std.debug.assert;
6const log = std.log.scoped(.codegen);
7
8const Zcu = @import("../../Zcu.zig");
9const Type = @import("../../Type.zig");
10const Value = @import("../../Value.zig");
11const Air = @import("../../Air.zig");
12const InternPool = @import("../../InternPool.zig");
13const Section = @import("Section.zig");
14const Assembler = @import("Assembler.zig");
15
16const spec = @import("spec.zig");
17const Opcode = spec.Opcode;
18const Word = spec.Word;
19const Id = spec.Id;
20const IdRange = spec.IdRange;
21const StorageClass = spec.StorageClass;
22
23const Module = @import("Module.zig");
24const Decl = Module.Decl;
25const Repr = Module.Repr;
26const InternMap = Module.InternMap;
27const PtrTypeMap = Module.PtrTypeMap;
28
29const CodeGen = @This();
30
31pub fn legalizeFeatures(_: *const std.Target) *const Air.Legalize.Features {
32 return comptime &.initMany(&.{
33 .expand_intcast_safe,
34 .expand_int_from_float_safe,
35 .expand_int_from_float_optimized_safe,
36 .expand_add_safe,
37 .expand_sub_safe,
38 .expand_mul_safe,
39 });
40}
41
42pub const zig_call_abi_ver = 3;
43
44const ControlFlow = union(enum) {
45 const Structured = struct {
46 /// This type indicates the way that a block is terminated. The
47 /// state of a particular block is used to track how a jump from
48 /// inside the block must reach the outside.
49 const Block = union(enum) {
50 const Incoming = struct {
51 src_label: Id,
52 /// Instruction that returns an u32 value of the
53 /// `Air.Inst.Index` that control flow should jump to.
54 next_block: Id,
55 };
56
57 const SelectionMerge = struct {
58 /// Incoming block from the `then` label.
59 /// Note that hte incoming block from the `else` label is
60 /// either given by the next element in the stack.
61 incoming: Incoming,
62 /// The label id of the cond_br's merge block.
63 /// For the top-most element in the stack, this
64 /// value is undefined.
65 merge_block: Id,
66 };
67
68 /// For a `selection` type block, we cannot use early exits, and we
69 /// must generate a 'merge ladder' of OpSelection instructions. To that end,
70 /// we keep a stack of the merges that still must be closed at the end of
71 /// a block.
72 ///
73 /// This entire structure basically just resembles a tree like
74 /// a x
75 /// \ /
76 /// b o merge
77 /// \ /
78 /// c o merge
79 /// \ /
80 /// o merge
81 /// /
82 /// o jump to next block
83 selection: struct {
84 /// In order to know which merges we still need to do, we need to keep
85 /// a stack of those.
86 merge_stack: std.ArrayList(SelectionMerge) = .empty,
87 },
88 /// For a `loop` type block, we can early-exit the block by
89 /// jumping to the loop exit node, and we don't need to generate
90 /// an entire stack of merges.
91 loop: struct {
92 /// The next block to jump to can be determined from any number
93 /// of conditions that jump to the loop exit.
94 merges: std.ArrayList(Incoming) = .empty,
95 /// The label id of the loop's merge block.
96 merge_block: Id,
97 },
98
99 fn deinit(block: *Structured.Block, gpa: Allocator) void {
100 switch (block.*) {
101 .selection => |*merge| merge.merge_stack.deinit(gpa),
102 .loop => |*merge| merge.merges.deinit(gpa),
103 }
104 block.* = undefined;
105 }
106 };
107 /// This determines how exits from the current block must be handled.
108 block_stack: std.ArrayList(*Structured.Block) = .empty,
109 block_results: std.AutoHashMapUnmanaged(Air.Inst.Index, Id) = .empty,
110 };
111
112 const Unstructured = struct {
113 const Incoming = struct {
114 src_label: Id,
115 break_value_id: Id,
116 };
117
118 const Block = struct {
119 label: ?Id = null,
120 incoming_blocks: std.ArrayList(Incoming) = .empty,
121 };
122
123 /// We need to keep track of result ids for block labels, as well as the 'incoming'
124 /// blocks for a block.
125 blocks: std.AutoHashMapUnmanaged(Air.Inst.Index, *Block) = .empty,
126 };
127
128 structured: Structured,
129 unstructured: Unstructured,
130
131 pub fn deinit(cg: *ControlFlow, gpa: Allocator) void {
132 switch (cg.*) {
133 .structured => |*cf| {
134 cf.block_stack.deinit(gpa);
135 cf.block_results.deinit(gpa);
136 },
137 .unstructured => |*cf| {
138 cf.blocks.deinit(gpa);
139 },
140 }
141 cg.* = undefined;
142 }
143};
144
145pt: Zcu.PerThread,
146air: Air,
147liveness: Air.Liveness,
148owner_nav: InternPool.Nav.Index,
149module: *Module,
150control_flow: ControlFlow,
151base_line: u32,
152block_label: Id = .none,
153next_arg_index: u32 = 0,
154args: std.ArrayList(Id) = .empty,
155inst_results: std.AutoHashMapUnmanaged(Air.Inst.Index, Id) = .empty,
156id_scratch: std.ArrayList(Id) = .empty,
157prologue: Section = .{},
158body: Section = .{},
159error_msg: ?*Zcu.ErrorMsg = null,
160
161pub fn deinit(cg: *CodeGen) void {
162 const gpa = cg.module.gpa;
163 cg.control_flow.deinit(gpa);
164 cg.args.deinit(gpa);
165 cg.inst_results.deinit(gpa);
166 cg.id_scratch.deinit(gpa);
167 cg.prologue.deinit(gpa);
168 cg.body.deinit(gpa);
169}
170
171const Error = error{ CodegenFail, OutOfMemory };
172
173pub fn genNav(cg: *CodeGen, do_codegen: bool) Error!void {
174 const gpa = cg.module.gpa;
175 const zcu = cg.module.zcu;
176 const ip = &zcu.intern_pool;
177 const target = zcu.getTarget();
178
179 const nav = ip.getNav(cg.owner_nav);
180 const val = zcu.navValue(cg.owner_nav);
181 const ty = val.typeOf(zcu);
182
183 if (!do_codegen and !ty.hasRuntimeBits(zcu)) return;
184
185 const spv_decl_index = try cg.module.resolveNav(ip, cg.owner_nav);
186 const decl = cg.module.declPtr(spv_decl_index);
187 const result_id = decl.result_id;
188 decl.begin_dep = cg.module.decl_deps.items.len;
189
190 switch (decl.kind) {
191 .func => {
192 const fn_info = zcu.typeToFunc(ty).?;
193 const return_ty_id = try cg.resolveFnReturnType(.fromInterned(fn_info.return_type));
194 const is_test = zcu.test_functions.contains(cg.owner_nav);
195
196 const func_result_id = if (is_test) cg.module.allocId() else result_id;
197 const prototype_ty_id = try cg.resolveType(ty, .direct);
198 try cg.prologue.emit(gpa, .OpFunction, .{
199 .id_result_type = return_ty_id,
200 .id_result = func_result_id,
201 .function_type = prototype_ty_id,
202 // Note: the backend will never be asked to generate an inline function
203 // (this is handled in sema), so we don't need to set function_control here.
204 .function_control = .{},
205 });
206
207 comptime assert(zig_call_abi_ver == 3);
208 try cg.args.ensureUnusedCapacity(gpa, fn_info.param_types.len);
209 for (fn_info.param_types.get(ip)) |param_ty_index| {
210 const param_ty: Type = .fromInterned(param_ty_index);
211 if (!param_ty.hasRuntimeBitsIgnoreComptime(zcu)) continue;
212
213 const param_type_id = try cg.resolveType(param_ty, .direct);
214 const arg_result_id = cg.module.allocId();
215 try cg.prologue.emit(gpa, .OpFunctionParameter, .{
216 .id_result_type = param_type_id,
217 .id_result = arg_result_id,
218 });
219 cg.args.appendAssumeCapacity(arg_result_id);
220 }
221
222 // TODO: This could probably be done in a better way...
223 const root_block_id = cg.module.allocId();
224
225 // The root block of a function declaration should appear before OpVariable instructions,
226 // so it is generated into the function's prologue.
227 try cg.prologue.emit(gpa, .OpLabel, .{
228 .id_result = root_block_id,
229 });
230 cg.block_label = root_block_id;
231
232 const main_body = cg.air.getMainBody();
233 switch (cg.control_flow) {
234 .structured => {
235 _ = try cg.genStructuredBody(.selection, main_body);
236 // We always expect paths to here to end, but we still need the block
237 // to act as a dummy merge block.
238 try cg.body.emit(gpa, .OpUnreachable, {});
239 },
240 .unstructured => {
241 try cg.genBody(main_body);
242 },
243 }
244 try cg.body.emit(gpa, .OpFunctionEnd, {});
245 // Append the actual code into the functions section.
246 try cg.module.sections.functions.append(gpa, cg.prologue);
247 try cg.module.sections.functions.append(gpa, cg.body);
248
249 // Temporarily generate a test kernel declaration if this is a test function.
250 if (is_test) {
251 try cg.generateTestEntryPoint(nav.fqn.toSlice(ip), spv_decl_index, func_result_id);
252 }
253
254 try cg.module.debugName(func_result_id, nav.fqn.toSlice(ip));
255 },
256 .global => {
257 assert(ip.indexToKey(val.toIntern()) == .@"extern");
258
259 const storage_class = cg.module.storageClass(nav.getAddrspace());
260 assert(storage_class != .generic); // These should be instance globals
261
262 const ty_id = try cg.resolveType(ty, .indirect);
263 const ptr_ty_id = try cg.module.ptrType(ty_id, storage_class);
264
265 try cg.module.sections.globals.emit(gpa, .OpVariable, .{
266 .id_result_type = ptr_ty_id,
267 .id_result = result_id,
268 .storage_class = storage_class,
269 });
270
271 switch (target.os.tag) {
272 .vulkan, .opengl => {
273 if (ty.zigTypeTag(zcu) == .@"struct") {
274 switch (storage_class) {
275 .uniform, .push_constant => try cg.module.decorate(ty_id, .block),
276 else => {},
277 }
278 }
279
280 switch (ip.indexToKey(ty.toIntern())) {
281 .func_type, .opaque_type => {},
282 else => {
283 try cg.module.decorate(ptr_ty_id, .{
284 .array_stride = .{ .array_stride = @intCast(ty.abiSize(zcu)) },
285 });
286 },
287 }
288 },
289 else => {},
290 }
291
292 if (std.meta.stringToEnum(spec.BuiltIn, nav.fqn.toSlice(ip))) |builtin| {
293 try cg.module.decorate(result_id, .{ .built_in = .{ .built_in = builtin } });
294 }
295
296 try cg.module.debugName(result_id, nav.fqn.toSlice(ip));
297 },
298 .invocation_global => {
299 const maybe_init_val: ?Value = switch (ip.indexToKey(val.toIntern())) {
300 .func => unreachable,
301 .variable => |variable| .fromInterned(variable.init),
302 .@"extern" => null,
303 else => val,
304 };
305
306 const ty_id = try cg.resolveType(ty, .indirect);
307 const ptr_ty_id = try cg.module.ptrType(ty_id, .function);
308
309 if (maybe_init_val) |init_val| {
310 // TODO: Combine with resolveAnonDecl?
311 const void_ty_id = try cg.resolveType(.void, .direct);
312 const initializer_proto_ty_id = try cg.module.functionType(void_ty_id, &.{});
313
314 const initializer_id = cg.module.allocId();
315 try cg.prologue.emit(gpa, .OpFunction, .{
316 .id_result_type = try cg.resolveType(.void, .direct),
317 .id_result = initializer_id,
318 .function_control = .{},
319 .function_type = initializer_proto_ty_id,
320 });
321
322 const root_block_id = cg.module.allocId();
323 try cg.prologue.emit(gpa, .OpLabel, .{
324 .id_result = root_block_id,
325 });
326 cg.block_label = root_block_id;
327
328 const val_id = try cg.constant(ty, init_val, .indirect);
329 try cg.body.emit(gpa, .OpStore, .{
330 .pointer = result_id,
331 .object = val_id,
332 });
333
334 try cg.body.emit(gpa, .OpReturn, {});
335 try cg.body.emit(gpa, .OpFunctionEnd, {});
336 try cg.module.sections.functions.append(gpa, cg.prologue);
337 try cg.module.sections.functions.append(gpa, cg.body);
338
339 try cg.module.debugNameFmt(initializer_id, "initializer of {f}", .{nav.fqn.fmt(ip)});
340
341 try cg.module.sections.globals.emit(gpa, .OpExtInst, .{
342 .id_result_type = ptr_ty_id,
343 .id_result = result_id,
344 .set = try cg.module.importInstructionSet(.zig),
345 .instruction = .{ .inst = @intFromEnum(spec.Zig.InvocationGlobal) },
346 .id_ref_4 = &.{initializer_id},
347 });
348 } else {
349 try cg.module.sections.globals.emit(gpa, .OpExtInst, .{
350 .id_result_type = ptr_ty_id,
351 .id_result = result_id,
352 .set = try cg.module.importInstructionSet(.zig),
353 .instruction = .{ .inst = @intFromEnum(spec.Zig.InvocationGlobal) },
354 .id_ref_4 = &.{},
355 });
356 }
357 },
358 }
359
360 cg.module.declPtr(spv_decl_index).end_dep = cg.module.decl_deps.items.len;
361}
362
363pub fn fail(cg: *CodeGen, comptime format: []const u8, args: anytype) Error {
364 @branchHint(.cold);
365 const zcu = cg.module.zcu;
366 const src_loc = zcu.navSrcLoc(cg.owner_nav);
367 assert(cg.error_msg == null);
368 cg.error_msg = try Zcu.ErrorMsg.create(zcu.gpa, src_loc, format, args);
369 return error.CodegenFail;
370}
371
372pub fn todo(cg: *CodeGen, comptime format: []const u8, args: anytype) Error {
373 return cg.fail("TODO (SPIR-V): " ++ format, args);
374}
375
376/// This imports the "default" extended instruction set for the target
377/// For OpenCL, OpenCL.std.100. For Vulkan and OpenGL, GLSL.std.450.
378fn importExtendedSet(cg: *CodeGen) !Id {
379 const target = cg.module.zcu.getTarget();
380 return switch (target.os.tag) {
381 .opencl, .amdhsa => try cg.module.importInstructionSet(.@"OpenCL.std"),
382 .vulkan, .opengl => try cg.module.importInstructionSet(.@"GLSL.std.450"),
383 else => unreachable,
384 };
385}
386
387/// Fetch the result-id for a previously generated instruction or constant.
388fn resolve(cg: *CodeGen, inst: Air.Inst.Ref) !Id {
389 const pt = cg.pt;
390 const zcu = cg.module.zcu;
391 const ip = &zcu.intern_pool;
392 if (try cg.air.value(inst, pt)) |val| {
393 const ty = cg.typeOf(inst);
394 if (ty.zigTypeTag(zcu) == .@"fn") {
395 const fn_nav = switch (zcu.intern_pool.indexToKey(val.ip_index)) {
396 .@"extern" => |@"extern"| @"extern".owner_nav,
397 .func => |func| func.owner_nav,
398 else => unreachable,
399 };
400 const spv_decl_index = try cg.module.resolveNav(ip, fn_nav);
401 try cg.module.decl_deps.append(cg.module.gpa, spv_decl_index);
402 return cg.module.declPtr(spv_decl_index).result_id;
403 }
404
405 return try cg.constant(ty, val, .direct);
406 }
407 const index = inst.toIndex().?;
408 return cg.inst_results.get(index).?; // Assertion means instruction does not dominate usage.
409}
410
411fn resolveUav(cg: *CodeGen, val: InternPool.Index) !Id {
412 const gpa = cg.module.gpa;
413
414 // TODO: This cannot be a function at this point, but it should probably be handled anyway.
415
416 const zcu = cg.module.zcu;
417 const ty: Type = .fromInterned(zcu.intern_pool.typeOf(val));
418 const ty_id = try cg.resolveType(ty, .indirect);
419
420 const spv_decl_index = blk: {
421 const entry = try cg.module.uav_link.getOrPut(gpa, .{ val, .function });
422 if (entry.found_existing) {
423 try cg.addFunctionDep(entry.value_ptr.*, .function);
424 return cg.module.declPtr(entry.value_ptr.*).result_id;
425 }
426
427 const spv_decl_index = try cg.module.allocDecl(.invocation_global);
428 try cg.addFunctionDep(spv_decl_index, .function);
429 entry.value_ptr.* = spv_decl_index;
430 break :blk spv_decl_index;
431 };
432
433 // TODO: At some point we will be able to generate this all constant here, but then all of
434 // constant() will need to be implemented such that it doesn't generate any at-runtime code.
435 // NOTE: Because this is a global, we really only want to initialize it once. Therefore the
436 // constant lowering of this value will need to be deferred to an initializer similar to
437 // other globals.
438
439 const result_id = cg.module.declPtr(spv_decl_index).result_id;
440
441 {
442 // Save the current state so that we can temporarily generate into a different function.
443 // TODO: This should probably be made a little more robust.
444 const func_prologue = cg.prologue;
445 const func_body = cg.body;
446 const block_label = cg.block_label;
447 defer {
448 cg.prologue = func_prologue;
449 cg.body = func_body;
450 cg.block_label = block_label;
451 }
452
453 cg.prologue = .{};
454 cg.body = .{};
455 defer {
456 cg.prologue.deinit(gpa);
457 cg.body.deinit(gpa);
458 }
459
460 const void_ty_id = try cg.resolveType(.void, .direct);
461 const initializer_proto_ty_id = try cg.module.functionType(void_ty_id, &.{});
462
463 const initializer_id = cg.module.allocId();
464 try cg.prologue.emit(gpa, .OpFunction, .{
465 .id_result_type = try cg.resolveType(.void, .direct),
466 .id_result = initializer_id,
467 .function_control = .{},
468 .function_type = initializer_proto_ty_id,
469 });
470 const root_block_id = cg.module.allocId();
471 try cg.prologue.emit(gpa, .OpLabel, .{
472 .id_result = root_block_id,
473 });
474 cg.block_label = root_block_id;
475
476 const val_id = try cg.constant(ty, .fromInterned(val), .indirect);
477 try cg.body.emit(gpa, .OpStore, .{
478 .pointer = result_id,
479 .object = val_id,
480 });
481
482 try cg.body.emit(gpa, .OpReturn, {});
483 try cg.body.emit(gpa, .OpFunctionEnd, {});
484
485 try cg.module.sections.functions.append(gpa, cg.prologue);
486 try cg.module.sections.functions.append(gpa, cg.body);
487
488 try cg.module.debugNameFmt(initializer_id, "initializer of __anon_{d}", .{@intFromEnum(val)});
489
490 const fn_decl_ptr_ty_id = try cg.module.ptrType(ty_id, .function);
491 try cg.module.sections.globals.emit(gpa, .OpExtInst, .{
492 .id_result_type = fn_decl_ptr_ty_id,
493 .id_result = result_id,
494 .set = try cg.module.importInstructionSet(.zig),
495 .instruction = .{ .inst = @intFromEnum(spec.Zig.InvocationGlobal) },
496 .id_ref_4 = &.{initializer_id},
497 });
498 }
499
500 return result_id;
501}
502
503fn addFunctionDep(cg: *CodeGen, decl_index: Module.Decl.Index, storage_class: StorageClass) !void {
504 const gpa = cg.module.gpa;
505 const target = cg.module.zcu.getTarget();
506 if (target.cpu.has(.spirv, .v1_4)) {
507 try cg.module.decl_deps.append(gpa, decl_index);
508 } else {
509 // Before version 1.4, the interface’s storage classes are limited to the Input and Output
510 if (storage_class == .input or storage_class == .output) {
511 try cg.module.decl_deps.append(gpa, decl_index);
512 }
513 }
514}
515
516/// Start a new SPIR-V block, Emits the label of the new block, and stores which
517/// block we are currently generating.
518/// Note that there is no such thing as nested blocks like in ZIR or AIR, so we don't need to
519/// keep track of the previous block.
520fn beginSpvBlock(cg: *CodeGen, label: Id) !void {
521 try cg.body.emit(cg.module.gpa, .OpLabel, .{ .id_result = label });
522 cg.block_label = label;
523}
524
525/// Return the amount of bits in the largest supported integer type. This is either 32 (always supported), or 64 (if
526/// the Int64 capability is enabled).
527/// Note: The extension SPV_INTEL_arbitrary_precision_integers allows any integer size (at least up to 32 bits).
528/// In theory that could also be used, but since the spec says that it only guarantees support up to 32-bit ints there
529/// is no way of knowing whether those are actually supported.
530/// TODO: Maybe this should be cached?
531fn largestSupportedIntBits(cg: *CodeGen) u16 {
532 const target = cg.module.zcu.getTarget();
533 if (target.cpu.has(.spirv, .int64) or target.cpu.arch == .spirv64) {
534 return 64;
535 }
536 return 32;
537}
538
539const ArithmeticTypeInfo = struct {
540 const Class = enum {
541 bool,
542 /// A regular, **native**, integer.
543 /// This is only returned when the backend supports this int as a native type (when
544 /// the relevant capability is enabled).
545 integer,
546 /// A regular float. These are all required to be natively supported. Floating points
547 /// for which the relevant capability is not enabled are not emulated.
548 float,
549 /// An integer of a 'strange' size (which' bit size is not the same as its backing
550 /// type. **Note**: this may **also** include power-of-2 integers for which the
551 /// relevant capability is not enabled), but still within the limits of the largest
552 /// natively supported integer type.
553 strange_integer,
554 /// An integer with more bits than the largest natively supported integer type.
555 composite_integer,
556 };
557
558 /// A classification of the inner type.
559 /// These scenarios will all have to be handled slightly different.
560 class: Class,
561 /// The number of bits in the inner type.
562 /// This is the actual number of bits of the type, not the size of the backing integer.
563 bits: u16,
564 /// The number of bits required to store the type.
565 /// For `integer` and `float`, this is equal to `bits`.
566 /// For `strange_integer` and `bool` this is the size of the backing integer.
567 /// For `composite_integer` this is the elements count.
568 backing_bits: u16,
569 /// Null if this type is a scalar, or the length of the vector otherwise.
570 vector_len: ?u32,
571 /// Whether the inner type is signed. Only relevant for integers.
572 signedness: std.builtin.Signedness,
573};
574
575fn arithmeticTypeInfo(cg: *CodeGen, ty: Type) ArithmeticTypeInfo {
576 const zcu = cg.module.zcu;
577 const target = cg.module.zcu.getTarget();
578 var scalar_ty = ty.scalarType(zcu);
579 if (scalar_ty.zigTypeTag(zcu) == .@"enum") {
580 scalar_ty = scalar_ty.intTagType(zcu);
581 }
582 const vector_len = if (ty.isVector(zcu)) ty.vectorLen(zcu) else null;
583 return switch (scalar_ty.zigTypeTag(zcu)) {
584 .bool => .{
585 .bits = 1, // Doesn't matter for this class.
586 .backing_bits = cg.module.backingIntBits(1).@"0",
587 .vector_len = vector_len,
588 .signedness = .unsigned, // Technically, but doesn't matter for this class.
589 .class = .bool,
590 },
591 .float => .{
592 .bits = scalar_ty.floatBits(target),
593 .backing_bits = scalar_ty.floatBits(target), // TODO: F80?
594 .vector_len = vector_len,
595 .signedness = .signed, // Technically, but doesn't matter for this class.
596 .class = .float,
597 },
598 .int => blk: {
599 const int_info = scalar_ty.intInfo(zcu);
600 // TODO: Maybe it's useful to also return this value.
601 const backing_bits, const big_int = cg.module.backingIntBits(int_info.bits);
602 break :blk .{
603 .bits = int_info.bits,
604 .backing_bits = backing_bits,
605 .vector_len = vector_len,
606 .signedness = int_info.signedness,
607 .class = class: {
608 if (big_int) break :class .composite_integer;
609 break :class if (backing_bits == int_info.bits) .integer else .strange_integer;
610 },
611 };
612 },
613 .@"enum" => unreachable,
614 .vector => unreachable,
615 else => unreachable, // Unhandled arithmetic type
616 };
617}
618
619/// Checks whether the type can be directly translated to SPIR-V vectors
620fn isSpvVector(cg: *CodeGen, ty: Type) bool {
621 const zcu = cg.module.zcu;
622 const target = cg.module.zcu.getTarget();
623 if (ty.zigTypeTag(zcu) != .vector) return false;
624
625 // TODO: This check must be expanded for types that can be represented
626 // as integers (enums / packed structs?) and types that are represented
627 // by multiple SPIR-V values.
628 const scalar_ty = ty.scalarType(zcu);
629 switch (scalar_ty.zigTypeTag(zcu)) {
630 .bool,
631 .int,
632 .float,
633 => {},
634 else => return false,
635 }
636
637 const elem_ty = ty.childType(zcu);
638 const len = ty.vectorLen(zcu);
639
640 if (elem_ty.isNumeric(zcu) or elem_ty.toIntern() == .bool_type) {
641 if (len > 1 and len <= 4) return true;
642 if (target.cpu.has(.spirv, .vector16)) return (len == 8 or len == 16);
643 }
644
645 return false;
646}
647
648/// Emits a bool constant in a particular representation.
649fn constBool(cg: *CodeGen, value: bool, repr: Repr) !Id {
650 return switch (repr) {
651 .indirect => cg.constInt(.u1, @intFromBool(value)),
652 .direct => cg.module.constBool(value),
653 };
654}
655
656/// Emits an integer constant.
657/// This function, unlike Module.constInt, takes care to bitcast
658/// the value to an unsigned int first for Kernels.
659fn constInt(cg: *CodeGen, ty: Type, value: anytype) !Id {
660 const zcu = cg.module.zcu;
661 const target = cg.module.zcu.getTarget();
662 const scalar_ty = ty.scalarType(zcu);
663 const int_info = scalar_ty.intInfo(zcu);
664 // Use backing bits so that negatives are sign extended
665 const backing_bits, const big_int = cg.module.backingIntBits(int_info.bits);
666 assert(backing_bits != 0); // u0 is comptime
667
668 const result_ty_id = try cg.resolveType(scalar_ty, .indirect);
669 const signedness: Signedness = switch (@typeInfo(@TypeOf(value))) {
670 .int => |int| int.signedness,
671 .comptime_int => if (value < 0) .signed else .unsigned,
672 else => unreachable,
673 };
674 if (@sizeOf(@TypeOf(value)) >= 4 and big_int) {
675 const value64: u64 = switch (signedness) {
676 .signed => @bitCast(@as(i64, @intCast(value))),
677 .unsigned => @as(u64, @intCast(value)),
678 };
679 assert(backing_bits == 64);
680 return cg.constructComposite(result_ty_id, &.{
681 try cg.constInt(.u32, @as(u32, @truncate(value64))),
682 try cg.constInt(.u32, @as(u32, @truncate(value64 << 32))),
683 });
684 }
685
686 const final_value: spec.LiteralContextDependentNumber = switch (target.os.tag) {
687 .opencl, .amdhsa => blk: {
688 const value64: u64 = switch (signedness) {
689 .signed => @bitCast(@as(i64, @intCast(value))),
690 .unsigned => @as(u64, @intCast(value)),
691 };
692
693 // Manually truncate the value to the right amount of bits.
694 const truncated_value = if (backing_bits == 64)
695 value64
696 else
697 value64 & (@as(u64, 1) << @intCast(backing_bits)) - 1;
698
699 break :blk switch (backing_bits) {
700 1...32 => .{ .uint32 = @truncate(truncated_value) },
701 33...64 => .{ .uint64 = truncated_value },
702 else => unreachable,
703 };
704 },
705 else => switch (backing_bits) {
706 1...32 => if (signedness == .signed) .{ .int32 = @intCast(value) } else .{ .uint32 = @intCast(value) },
707 33...64 => if (signedness == .signed) .{ .int64 = value } else .{ .uint64 = value },
708 else => unreachable,
709 },
710 };
711
712 const result_id = try cg.module.constant(result_ty_id, final_value);
713
714 if (!ty.isVector(zcu)) return result_id;
715 return cg.constructCompositeSplat(ty, result_id);
716}
717
718pub fn constructComposite(cg: *CodeGen, result_ty_id: Id, constituents: []const Id) !Id {
719 const gpa = cg.module.gpa;
720 const result_id = cg.module.allocId();
721 try cg.body.emit(gpa, .OpCompositeConstruct, .{
722 .id_result_type = result_ty_id,
723 .id_result = result_id,
724 .constituents = constituents,
725 });
726 return result_id;
727}
728
729/// Construct a composite at runtime with all lanes set to the same value.
730/// ty must be an aggregate type.
731fn constructCompositeSplat(cg: *CodeGen, ty: Type, constituent: Id) !Id {
732 const gpa = cg.module.gpa;
733 const zcu = cg.module.zcu;
734 const n: usize = @intCast(ty.arrayLen(zcu));
735
736 const scratch_top = cg.id_scratch.items.len;
737 defer cg.id_scratch.shrinkRetainingCapacity(scratch_top);
738
739 const constituents = try cg.id_scratch.addManyAsSlice(gpa, n);
740 @memset(constituents, constituent);
741
742 const result_ty_id = try cg.resolveType(ty, .direct);
743 return cg.constructComposite(result_ty_id, constituents);
744}
745
746/// This function generates a load for a constant in direct (ie, non-memory) representation.
747/// When the constant is simple, it can be generated directly using OpConstant instructions.
748/// When the constant is more complicated however, it needs to be constructed using multiple values. This
749/// is done by emitting a sequence of instructions that initialize the value.
750//
751/// This function should only be called during function code generation.
752fn constant(cg: *CodeGen, ty: Type, val: Value, repr: Repr) Error!Id {
753 const gpa = cg.module.gpa;
754
755 // Note: Using intern_map can only be used with constants that DO NOT generate any runtime code!!
756 // Ideally that should be all constants in the future, or it should be cleaned up somehow. For
757 // now, only use the intern_map on case-by-case basis by breaking to :cache.
758 if (cg.module.intern_map.get(.{ val.toIntern(), repr })) |id| {
759 return id;
760 }
761
762 const pt = cg.pt;
763 const zcu = cg.module.zcu;
764 const target = cg.module.zcu.getTarget();
765 const result_ty_id = try cg.resolveType(ty, repr);
766 const ip = &zcu.intern_pool;
767
768 log.debug("lowering constant: ty = {f}, val = {f}, key = {s}", .{ ty.fmt(pt), val.fmtValue(pt), @tagName(ip.indexToKey(val.toIntern())) });
769 if (val.isUndef(zcu)) {
770 return cg.module.constUndef(result_ty_id);
771 }
772
773 const cacheable_id = cache: {
774 switch (ip.indexToKey(val.toIntern())) {
775 .int_type,
776 .ptr_type,
777 .array_type,
778 .vector_type,
779 .opt_type,
780 .anyframe_type,
781 .error_union_type,
782 .simple_type,
783 .struct_type,
784 .tuple_type,
785 .union_type,
786 .opaque_type,
787 .enum_type,
788 .func_type,
789 .error_set_type,
790 .inferred_error_set_type,
791 => unreachable, // types, not values
792
793 .undef => unreachable, // handled above
794
795 .variable,
796 .@"extern",
797 .func,
798 .enum_literal,
799 .empty_enum_value,
800 => unreachable, // non-runtime values
801
802 .simple_value => |simple_value| switch (simple_value) {
803 .undefined,
804 .void,
805 .null,
806 .empty_tuple,
807 .@"unreachable",
808 => unreachable, // non-runtime values
809
810 .false, .true => break :cache try cg.constBool(val.toBool(), repr),
811 },
812 .int => {
813 if (ty.isSignedInt(zcu)) {
814 break :cache try cg.constInt(ty, val.toSignedInt(zcu));
815 } else {
816 break :cache try cg.constInt(ty, val.toUnsignedInt(zcu));
817 }
818 },
819 .float => {
820 const lit: spec.LiteralContextDependentNumber = switch (ty.floatBits(target)) {
821 16 => .{ .uint32 = @as(u16, @bitCast(val.toFloat(f16, zcu))) },
822 32 => .{ .float32 = val.toFloat(f32, zcu) },
823 64 => .{ .float64 = val.toFloat(f64, zcu) },
824 80, 128 => unreachable, // TODO
825 else => unreachable,
826 };
827 break :cache try cg.module.constant(result_ty_id, lit);
828 },
829 .err => |err| {
830 const value = try pt.getErrorValue(err.name);
831 break :cache try cg.constInt(ty, value);
832 },
833 .error_union => |error_union| {
834 // TODO: Error unions may be constructed with constant instructions if the payload type
835 // allows it. For now, just generate it here regardless.
836 const err_ty = ty.errorUnionSet(zcu);
837 const payload_ty = ty.errorUnionPayload(zcu);
838 const err_val_id = switch (error_union.val) {
839 .err_name => |err_name| try cg.constInt(
840 err_ty,
841 try pt.getErrorValue(err_name),
842 ),
843 .payload => try cg.constInt(err_ty, 0),
844 };
845 const eu_layout = cg.errorUnionLayout(payload_ty);
846 if (!eu_layout.payload_has_bits) {
847 // We use the error type directly as the type.
848 break :cache err_val_id;
849 }
850
851 const payload_val_id = switch (error_union.val) {
852 .err_name => try cg.constant(payload_ty, .undef, .indirect),
853 .payload => |p| try cg.constant(payload_ty, .fromInterned(p), .indirect),
854 };
855
856 var constituents: [2]Id = undefined;
857 var types: [2]Type = undefined;
858 if (eu_layout.error_first) {
859 constituents[0] = err_val_id;
860 constituents[1] = payload_val_id;
861 types = .{ err_ty, payload_ty };
862 } else {
863 constituents[0] = payload_val_id;
864 constituents[1] = err_val_id;
865 types = .{ payload_ty, err_ty };
866 }
867
868 const comp_ty_id = try cg.resolveType(ty, .direct);
869 return try cg.constructComposite(comp_ty_id, &constituents);
870 },
871 .enum_tag => {
872 const int_val = try val.intFromEnum(ty, pt);
873 const int_ty = ty.intTagType(zcu);
874 break :cache try cg.constant(int_ty, int_val, repr);
875 },
876 .ptr => return cg.constantPtr(val),
877 .slice => |slice| {
878 const ptr_id = try cg.constantPtr(.fromInterned(slice.ptr));
879 const len_id = try cg.constant(.usize, .fromInterned(slice.len), .indirect);
880 const comp_ty_id = try cg.resolveType(ty, .direct);
881 return try cg.constructComposite(comp_ty_id, &.{ ptr_id, len_id });
882 },
883 .opt => {
884 const payload_ty = ty.optionalChild(zcu);
885 const maybe_payload_val = val.optionalValue(zcu);
886
887 if (!payload_ty.hasRuntimeBits(zcu)) {
888 break :cache try cg.constBool(maybe_payload_val != null, .indirect);
889 } else if (ty.optionalReprIsPayload(zcu)) {
890 // Optional representation is a nullable pointer or slice.
891 if (maybe_payload_val) |payload_val| {
892 return try cg.constant(payload_ty, payload_val, .indirect);
893 } else {
894 break :cache try cg.module.constNull(result_ty_id);
895 }
896 }
897
898 // Optional representation is a structure.
899 // { Payload, Bool }
900
901 const has_pl_id = try cg.constBool(maybe_payload_val != null, .indirect);
902 const payload_id = if (maybe_payload_val) |payload_val|
903 try cg.constant(payload_ty, payload_val, .indirect)
904 else
905 try cg.module.constUndef(try cg.resolveType(payload_ty, .indirect));
906
907 const comp_ty_id = try cg.resolveType(ty, .direct);
908 return try cg.constructComposite(comp_ty_id, &.{ payload_id, has_pl_id });
909 },
910 .aggregate => |aggregate| switch (ip.indexToKey(ty.ip_index)) {
911 inline .array_type, .vector_type => |array_type, tag| {
912 const elem_ty: Type = .fromInterned(array_type.child);
913
914 const scratch_top = cg.id_scratch.items.len;
915 defer cg.id_scratch.shrinkRetainingCapacity(scratch_top);
916 const constituents = try cg.id_scratch.addManyAsSlice(gpa, @intCast(ty.arrayLenIncludingSentinel(zcu)));
917
918 const child_repr: Repr = switch (tag) {
919 .array_type => .indirect,
920 .vector_type => .direct,
921 else => unreachable,
922 };
923
924 switch (aggregate.storage) {
925 .bytes => |bytes| {
926 // TODO: This is really space inefficient, perhaps there is a better
927 // way to do it?
928 for (constituents, bytes.toSlice(constituents.len, ip)) |*constituent, byte| {
929 constituent.* = try cg.constInt(elem_ty, byte);
930 }
931 },
932 .elems => |elems| {
933 for (constituents, elems) |*constituent, elem| {
934 constituent.* = try cg.constant(elem_ty, .fromInterned(elem), child_repr);
935 }
936 },
937 .repeated_elem => |elem| {
938 @memset(constituents, try cg.constant(elem_ty, .fromInterned(elem), child_repr));
939 },
940 }
941
942 const comp_ty_id = try cg.resolveType(ty, .direct);
943 return cg.constructComposite(comp_ty_id, constituents);
944 },
945 .struct_type => {
946 const struct_type = zcu.typeToStruct(ty).?;
947
948 if (struct_type.layout == .@"packed") {
949 // TODO: composite int
950 // TODO: endianness
951 const bits: u16 = @intCast(ty.bitSize(zcu));
952 const bytes = std.mem.alignForward(u16, cg.module.backingIntBits(bits).@"0", 8) / 8;
953 var limbs: [8]u8 = undefined;
954 @memset(&limbs, 0);
955 val.writeToPackedMemory(ty, pt, limbs[0..bytes], 0) catch unreachable;
956 const backing_ty: Type = .fromInterned(struct_type.backingIntTypeUnordered(ip));
957 return try cg.constInt(backing_ty, @as(u64, @bitCast(limbs)));
958 }
959
960 var types = std.array_list.Managed(Type).init(gpa);
961 defer types.deinit();
962
963 var constituents = std.array_list.Managed(Id).init(gpa);
964 defer constituents.deinit();
965
966 var it = struct_type.iterateRuntimeOrder(ip);
967 while (it.next()) |field_index| {
968 const field_ty: Type = .fromInterned(struct_type.field_types.get(ip)[field_index]);
969 if (!field_ty.hasRuntimeBitsIgnoreComptime(zcu)) {
970 // This is a zero-bit field - we only needed it for the alignment.
971 continue;
972 }
973
974 // TODO: Padding?
975 const field_val = try val.fieldValue(pt, field_index);
976 const field_id = try cg.constant(field_ty, field_val, .indirect);
977
978 try types.append(field_ty);
979 try constituents.append(field_id);
980 }
981
982 const comp_ty_id = try cg.resolveType(ty, .direct);
983 return try cg.constructComposite(comp_ty_id, constituents.items);
984 },
985 .tuple_type => return cg.todo("implement tuple types", .{}),
986 else => unreachable,
987 },
988 .un => |un| {
989 if (un.tag == .none) {
990 assert(ty.containerLayout(zcu) == .@"packed"); // TODO
991 const int_ty = try pt.intType(.unsigned, @intCast(ty.bitSize(zcu)));
992 return try cg.constInt(int_ty, Value.toUnsignedInt(.fromInterned(un.val), zcu));
993 }
994 const active_field = ty.unionTagFieldIndex(.fromInterned(un.tag), zcu).?;
995 const union_obj = zcu.typeToUnion(ty).?;
996 const field_ty: Type = .fromInterned(union_obj.field_types.get(ip)[active_field]);
997 const payload = if (field_ty.hasRuntimeBitsIgnoreComptime(zcu))
998 try cg.constant(field_ty, .fromInterned(un.val), .direct)
999 else
1000 null;
1001 return try cg.unionInit(ty, active_field, payload);
1002 },
1003 .memoized_call => unreachable,
1004 }
1005 };
1006
1007 try cg.module.intern_map.putNoClobber(gpa, .{ val.toIntern(), repr }, cacheable_id);
1008
1009 return cacheable_id;
1010}
1011
1012fn constantPtr(cg: *CodeGen, ptr_val: Value) !Id {
1013 const pt = cg.pt;
1014 const zcu = cg.module.zcu;
1015 const gpa = cg.module.gpa;
1016
1017 if (ptr_val.isUndef(zcu)) {
1018 const result_ty = ptr_val.typeOf(zcu);
1019 const result_ty_id = try cg.resolveType(result_ty, .direct);
1020 return cg.module.constUndef(result_ty_id);
1021 }
1022
1023 var arena = std.heap.ArenaAllocator.init(gpa);
1024 defer arena.deinit();
1025
1026 const derivation = try ptr_val.pointerDerivation(arena.allocator(), pt);
1027 return cg.derivePtr(derivation);
1028}
1029
1030fn derivePtr(cg: *CodeGen, derivation: Value.PointerDeriveStep) !Id {
1031 const gpa = cg.module.gpa;
1032 const pt = cg.pt;
1033 const zcu = cg.module.zcu;
1034 const target = zcu.getTarget();
1035 switch (derivation) {
1036 .comptime_alloc_ptr, .comptime_field_ptr => unreachable,
1037 .int => |int| {
1038 if (target.os.tag != .opencl) {
1039 if (int.ptr_ty.ptrAddressSpace(zcu) != .physical_storage_buffer) {
1040 return cg.fail(
1041 "cannot cast integer to pointer with address space '{s}'",
1042 .{@tagName(int.ptr_ty.ptrAddressSpace(zcu))},
1043 );
1044 }
1045 }
1046 const result_ty_id = try cg.resolveType(int.ptr_ty, .direct);
1047 // TODO: This can probably be an OpSpecConstantOp Bitcast, but
1048 // that is not implemented by Mesa yet. Therefore, just generate it
1049 // as a runtime operation.
1050 const result_ptr_id = cg.module.allocId();
1051 const value_id = try cg.constInt(.usize, int.addr);
1052 try cg.body.emit(gpa, .OpConvertUToPtr, .{
1053 .id_result_type = result_ty_id,
1054 .id_result = result_ptr_id,
1055 .integer_value = value_id,
1056 });
1057 return result_ptr_id;
1058 },
1059 .nav_ptr => |nav| {
1060 const result_ptr_ty = try pt.navPtrType(nav);
1061 return cg.constantNavRef(result_ptr_ty, nav);
1062 },
1063 .uav_ptr => |uav| {
1064 const result_ptr_ty: Type = .fromInterned(uav.orig_ty);
1065 return cg.constantUavRef(result_ptr_ty, uav);
1066 },
1067 .eu_payload_ptr => @panic("TODO"),
1068 .opt_payload_ptr => @panic("TODO"),
1069 .field_ptr => |field| {
1070 const parent_ptr_id = try cg.derivePtr(field.parent.*);
1071 const parent_ptr_ty = try field.parent.ptrType(pt);
1072 return cg.structFieldPtr(field.result_ptr_ty, parent_ptr_ty, parent_ptr_id, field.field_idx);
1073 },
1074 .elem_ptr => |elem| {
1075 const parent_ptr_id = try cg.derivePtr(elem.parent.*);
1076 const parent_ptr_ty = try elem.parent.ptrType(pt);
1077 const index_id = try cg.constInt(.usize, elem.elem_idx);
1078 return cg.ptrElemPtr(parent_ptr_ty, parent_ptr_id, index_id);
1079 },
1080 .offset_and_cast => |oac| {
1081 const parent_ptr_id = try cg.derivePtr(oac.parent.*);
1082 const parent_ptr_ty = try oac.parent.ptrType(pt);
1083 const result_ty_id = try cg.resolveType(oac.new_ptr_ty, .direct);
1084 const child_size = oac.new_ptr_ty.childType(zcu).abiSize(zcu);
1085
1086 if (parent_ptr_ty.childType(zcu).isVector(zcu) and oac.byte_offset % child_size == 0) {
1087 // Vector element ptr accesses are derived as offset_and_cast.
1088 // We can just use OpAccessChain.
1089 return cg.accessChain(
1090 result_ty_id,
1091 parent_ptr_id,
1092 &.{@intCast(@divExact(oac.byte_offset, child_size))},
1093 );
1094 }
1095
1096 if (oac.byte_offset == 0) {
1097 // Allow changing the pointer type child only to restructure arrays.
1098 // e.g. [3][2]T to T is fine, as is [2]T -> [2][1]T.
1099 const result_ptr_id = cg.module.allocId();
1100 try cg.body.emit(gpa, .OpBitcast, .{
1101 .id_result_type = result_ty_id,
1102 .id_result = result_ptr_id,
1103 .operand = parent_ptr_id,
1104 });
1105 return result_ptr_id;
1106 }
1107
1108 return cg.fail("cannot perform pointer cast: '{f}' to '{f}'", .{
1109 parent_ptr_ty.fmt(pt),
1110 oac.new_ptr_ty.fmt(pt),
1111 });
1112 },
1113 }
1114}
1115
1116fn constantUavRef(
1117 cg: *CodeGen,
1118 ty: Type,
1119 uav: InternPool.Key.Ptr.BaseAddr.Uav,
1120) !Id {
1121 // TODO: Merge this function with constantDeclRef.
1122
1123 const zcu = cg.module.zcu;
1124 const ip = &zcu.intern_pool;
1125 const ty_id = try cg.resolveType(ty, .direct);
1126 const uav_ty: Type = .fromInterned(ip.typeOf(uav.val));
1127
1128 switch (ip.indexToKey(uav.val)) {
1129 .func => unreachable, // TODO
1130 .@"extern" => assert(!ip.isFunctionType(uav_ty.toIntern())),
1131 else => {},
1132 }
1133
1134 // const is_fn_body = decl_ty.zigTypeTag(zcu) == .@"fn";
1135 if (!uav_ty.isFnOrHasRuntimeBitsIgnoreComptime(zcu)) {
1136 // Pointer to nothing - return undefined
1137 return cg.module.constUndef(ty_id);
1138 }
1139
1140 // Uav refs are always generic.
1141 assert(ty.ptrAddressSpace(zcu) == .generic);
1142 const uav_ty_id = try cg.resolveType(uav_ty, .indirect);
1143 const decl_ptr_ty_id = try cg.module.ptrType(uav_ty_id, .function);
1144 const ptr_id = try cg.resolveUav(uav.val);
1145
1146 if (decl_ptr_ty_id != ty_id) {
1147 // Differing pointer types, insert a cast.
1148 const casted_ptr_id = cg.module.allocId();
1149 try cg.body.emit(cg.module.gpa, .OpBitcast, .{
1150 .id_result_type = ty_id,
1151 .id_result = casted_ptr_id,
1152 .operand = ptr_id,
1153 });
1154 return casted_ptr_id;
1155 } else {
1156 return ptr_id;
1157 }
1158}
1159
1160fn constantNavRef(cg: *CodeGen, ty: Type, nav_index: InternPool.Nav.Index) !Id {
1161 const zcu = cg.module.zcu;
1162 const ip = &zcu.intern_pool;
1163 const ty_id = try cg.resolveType(ty, .direct);
1164 const nav = ip.getNav(nav_index);
1165 const nav_ty: Type = .fromInterned(nav.typeOf(ip));
1166
1167 switch (nav.status) {
1168 .unresolved => unreachable,
1169 .type_resolved => {}, // this is not a function or extern
1170 .fully_resolved => |r| switch (ip.indexToKey(r.val)) {
1171 .func => {
1172 // TODO: Properly lower function pointers. For now we are going to hack around it and
1173 // just generate an empty pointer. Function pointers are represented by a pointer to usize.
1174 return try cg.module.constUndef(ty_id);
1175 },
1176 .@"extern" => if (ip.isFunctionType(nav_ty.toIntern())) @panic("TODO"),
1177 else => {},
1178 },
1179 }
1180
1181 if (!nav_ty.isFnOrHasRuntimeBitsIgnoreComptime(zcu)) {
1182 // Pointer to nothing - return undefined.
1183 return cg.module.constUndef(ty_id);
1184 }
1185
1186 const spv_decl_index = try cg.module.resolveNav(ip, nav_index);
1187 const spv_decl = cg.module.declPtr(spv_decl_index);
1188 const spv_decl_result_id = spv_decl.result_id;
1189 assert(spv_decl.kind != .func);
1190
1191 const storage_class = cg.module.storageClass(nav.getAddrspace());
1192 try cg.addFunctionDep(spv_decl_index, storage_class);
1193
1194 const nav_ty_id = try cg.resolveType(nav_ty, .indirect);
1195 const decl_ptr_ty_id = try cg.module.ptrType(nav_ty_id, storage_class);
1196
1197 if (decl_ptr_ty_id != ty_id) {
1198 // Differing pointer types, insert a cast.
1199 const casted_ptr_id = cg.module.allocId();
1200 try cg.body.emit(cg.module.gpa, .OpBitcast, .{
1201 .id_result_type = ty_id,
1202 .id_result = casted_ptr_id,
1203 .operand = spv_decl_result_id,
1204 });
1205 return casted_ptr_id;
1206 }
1207
1208 return spv_decl_result_id;
1209}
1210
1211// Turn a Zig type's name into a cache reference.
1212fn resolveTypeName(cg: *CodeGen, ty: Type) ![]const u8 {
1213 const gpa = cg.module.gpa;
1214 var aw: std.Io.Writer.Allocating = .init(gpa);
1215 defer aw.deinit();
1216 ty.print(&aw.writer, cg.pt, null) catch |err| switch (err) {
1217 error.WriteFailed => return error.OutOfMemory,
1218 };
1219 return try aw.toOwnedSlice();
1220}
1221
1222/// Generate a union type. Union types are always generated with the
1223/// most aligned field active. If the tag alignment is greater
1224/// than that of the payload, a regular union (non-packed, with both tag and
1225/// payload), will be generated as follows:
1226/// struct {
1227/// tag: TagType,
1228/// payload: MostAlignedFieldType,
1229/// payload_padding: [payload_size - @sizeOf(MostAlignedFieldType)]u8,
1230/// padding: [padding_size]u8,
1231/// }
1232/// If the payload alignment is greater than that of the tag:
1233/// struct {
1234/// payload: MostAlignedFieldType,
1235/// payload_padding: [payload_size - @sizeOf(MostAlignedFieldType)]u8,
1236/// tag: TagType,
1237/// padding: [padding_size]u8,
1238/// }
1239/// If any of the fields' size is 0, it will be omitted.
1240fn resolveUnionType(cg: *CodeGen, ty: Type) !Id {
1241 const gpa = cg.module.gpa;
1242 const zcu = cg.module.zcu;
1243 const ip = &zcu.intern_pool;
1244 const union_obj = zcu.typeToUnion(ty).?;
1245
1246 if (union_obj.flagsUnordered(ip).layout == .@"packed") {
1247 return try cg.module.intType(.unsigned, @intCast(ty.bitSize(zcu)));
1248 }
1249
1250 const layout = cg.unionLayout(ty);
1251 if (!layout.has_payload) {
1252 // No payload, so represent this as just the tag type.
1253 return try cg.resolveType(.fromInterned(union_obj.enum_tag_ty), .indirect);
1254 }
1255
1256 var member_types: [4]Id = undefined;
1257 var member_names: [4][]const u8 = undefined;
1258
1259 const u8_ty_id = try cg.resolveType(.u8, .direct);
1260
1261 if (layout.tag_size != 0) {
1262 const tag_ty_id = try cg.resolveType(.fromInterned(union_obj.enum_tag_ty), .indirect);
1263 member_types[layout.tag_index] = tag_ty_id;
1264 member_names[layout.tag_index] = "(tag)";
1265 }
1266
1267 if (layout.payload_size != 0) {
1268 const payload_ty_id = try cg.resolveType(layout.payload_ty, .indirect);
1269 member_types[layout.payload_index] = payload_ty_id;
1270 member_names[layout.payload_index] = "(payload)";
1271 }
1272
1273 if (layout.payload_padding_size != 0) {
1274 const len_id = try cg.constInt(.u32, layout.payload_padding_size);
1275 const payload_padding_ty_id = try cg.module.arrayType(len_id, u8_ty_id);
1276 member_types[layout.payload_padding_index] = payload_padding_ty_id;
1277 member_names[layout.payload_padding_index] = "(payload padding)";
1278 }
1279
1280 if (layout.padding_size != 0) {
1281 const len_id = try cg.constInt(.u32, layout.padding_size);
1282 const padding_ty_id = try cg.module.arrayType(len_id, u8_ty_id);
1283 member_types[layout.padding_index] = padding_ty_id;
1284 member_names[layout.padding_index] = "(padding)";
1285 }
1286
1287 const result_id = try cg.module.structType(
1288 member_types[0..layout.total_fields],
1289 member_names[0..layout.total_fields],
1290 null,
1291 .none,
1292 );
1293
1294 const type_name = try cg.resolveTypeName(ty);
1295 defer gpa.free(type_name);
1296 try cg.module.debugName(result_id, type_name);
1297
1298 return result_id;
1299}
1300
1301fn resolveFnReturnType(cg: *CodeGen, ret_ty: Type) !Id {
1302 const zcu = cg.module.zcu;
1303 if (!ret_ty.hasRuntimeBitsIgnoreComptime(zcu)) {
1304 // If the return type is an error set or an error union, then we make this
1305 // anyerror return type instead, so that it can be coerced into a function
1306 // pointer type which has anyerror as the return type.
1307 if (ret_ty.isError(zcu)) {
1308 return cg.resolveType(.anyerror, .direct);
1309 } else {
1310 return cg.resolveType(.void, .direct);
1311 }
1312 }
1313
1314 return try cg.resolveType(ret_ty, .direct);
1315}
1316
1317fn resolveType(cg: *CodeGen, ty: Type, repr: Repr) Error!Id {
1318 const gpa = cg.module.gpa;
1319 const pt = cg.pt;
1320 const zcu = cg.module.zcu;
1321 const ip = &zcu.intern_pool;
1322 const target = cg.module.zcu.getTarget();
1323
1324 log.debug("resolveType: ty = {f}", .{ty.fmt(pt)});
1325
1326 switch (ty.zigTypeTag(zcu)) {
1327 .noreturn => {
1328 assert(repr == .direct);
1329 return try cg.module.voidType();
1330 },
1331 .void => switch (repr) {
1332 .direct => return try cg.module.voidType(),
1333 .indirect => {
1334 if (target.os.tag != .opencl) return cg.fail("cannot generate opaque type", .{});
1335 return try cg.module.opaqueType("void");
1336 },
1337 },
1338 .bool => switch (repr) {
1339 .direct => return try cg.module.boolType(),
1340 .indirect => return try cg.resolveType(.u1, .indirect),
1341 },
1342 .int => {
1343 const int_info = ty.intInfo(zcu);
1344 if (int_info.bits == 0) {
1345 assert(repr == .indirect);
1346 if (target.os.tag != .opencl) return cg.fail("cannot generate opaque type", .{});
1347 return try cg.module.opaqueType("u0");
1348 }
1349 return try cg.module.intType(int_info.signedness, int_info.bits);
1350 },
1351 .@"enum" => return try cg.resolveType(ty.intTagType(zcu), repr),
1352 .float => {
1353 const bits = ty.floatBits(target);
1354 const supported = switch (bits) {
1355 16 => target.cpu.has(.spirv, .float16),
1356 32 => true,
1357 64 => target.cpu.has(.spirv, .float64),
1358 else => false,
1359 };
1360
1361 if (!supported) {
1362 return cg.fail(
1363 "floating point width of {} bits is not supported for the current SPIR-V feature set",
1364 .{bits},
1365 );
1366 }
1367
1368 return try cg.module.floatType(bits);
1369 },
1370 .array => {
1371 const elem_ty = ty.childType(zcu);
1372 const elem_ty_id = try cg.resolveType(elem_ty, .indirect);
1373 const total_len = std.math.cast(u32, ty.arrayLenIncludingSentinel(zcu)) orelse {
1374 return cg.fail("array type of {} elements is too large", .{ty.arrayLenIncludingSentinel(zcu)});
1375 };
1376
1377 if (!elem_ty.hasRuntimeBitsIgnoreComptime(zcu)) {
1378 assert(repr == .indirect);
1379 if (target.os.tag != .opencl) return cg.fail("cannot generate opaque type", .{});
1380 return try cg.module.opaqueType("zero-sized-array");
1381 } else if (total_len == 0) {
1382 // The size of the array would be 0, but that is not allowed in SPIR-V.
1383 // This path can be reached for example when there is a slicing of a pointer
1384 // that produces a zero-length array. In all cases where this type can be generated,
1385 // this should be an indirect path.
1386 assert(repr == .indirect);
1387 // In this case, we have an array of a non-zero sized type. In this case,
1388 // generate an array of 1 element instead, so that ptr_elem_ptr instructions
1389 // can be lowered to ptrAccessChain instead of manually performing the math.
1390 const len_id = try cg.constInt(.u32, 1);
1391 return try cg.module.arrayType(len_id, elem_ty_id);
1392 } else {
1393 const total_len_id = try cg.constInt(.u32, total_len);
1394 const result_id = try cg.module.arrayType(total_len_id, elem_ty_id);
1395 switch (target.os.tag) {
1396 .vulkan, .opengl => {
1397 try cg.module.decorate(result_id, .{
1398 .array_stride = .{
1399 .array_stride = @intCast(elem_ty.abiSize(zcu)),
1400 },
1401 });
1402 },
1403 else => {},
1404 }
1405 return result_id;
1406 }
1407 },
1408 .vector => {
1409 const elem_ty = ty.childType(zcu);
1410 const elem_ty_id = try cg.resolveType(elem_ty, repr);
1411 const len = ty.vectorLen(zcu);
1412 if (cg.isSpvVector(ty)) return try cg.module.vectorType(len, elem_ty_id);
1413 const len_id = try cg.constInt(.u32, len);
1414 return try cg.module.arrayType(len_id, elem_ty_id);
1415 },
1416 .@"fn" => switch (repr) {
1417 .direct => {
1418 const fn_info = zcu.typeToFunc(ty).?;
1419
1420 comptime assert(zig_call_abi_ver == 3);
1421 assert(!fn_info.is_var_args);
1422 switch (fn_info.cc) {
1423 .auto,
1424 .spirv_kernel,
1425 .spirv_fragment,
1426 .spirv_vertex,
1427 .spirv_device,
1428 => {},
1429 else => unreachable,
1430 }
1431
1432 const return_ty_id = try cg.resolveFnReturnType(.fromInterned(fn_info.return_type));
1433
1434 const scratch_top = cg.id_scratch.items.len;
1435 defer cg.id_scratch.shrinkRetainingCapacity(scratch_top);
1436 const param_ty_ids = try cg.id_scratch.addManyAsSlice(gpa, fn_info.param_types.len);
1437
1438 var param_index: usize = 0;
1439 for (fn_info.param_types.get(ip)) |param_ty_index| {
1440 const param_ty: Type = .fromInterned(param_ty_index);
1441 if (!param_ty.hasRuntimeBitsIgnoreComptime(zcu)) continue;
1442
1443 param_ty_ids[param_index] = try cg.resolveType(param_ty, .direct);
1444 param_index += 1;
1445 }
1446
1447 return try cg.module.functionType(return_ty_id, param_ty_ids[0..param_index]);
1448 },
1449 .indirect => {
1450 // TODO: Represent function pointers properly.
1451 // For now, just use an usize type.
1452 return try cg.resolveType(.usize, .indirect);
1453 },
1454 },
1455 .pointer => {
1456 const ptr_info = ty.ptrInfo(zcu);
1457
1458 const child_ty: Type = .fromInterned(ptr_info.child);
1459 const child_ty_id = try cg.resolveType(child_ty, .indirect);
1460 const storage_class = cg.module.storageClass(ptr_info.flags.address_space);
1461 const ptr_ty_id = try cg.module.ptrType(child_ty_id, storage_class);
1462
1463 if (ptr_info.flags.size != .slice) {
1464 return ptr_ty_id;
1465 }
1466
1467 const size_ty_id = try cg.resolveType(.usize, .direct);
1468 return try cg.module.structType(
1469 &.{ ptr_ty_id, size_ty_id },
1470 &.{ "ptr", "len" },
1471 null,
1472 .none,
1473 );
1474 },
1475 .@"struct" => {
1476 const struct_type = switch (ip.indexToKey(ty.toIntern())) {
1477 .tuple_type => |tuple| {
1478 const scratch_top = cg.id_scratch.items.len;
1479 defer cg.id_scratch.shrinkRetainingCapacity(scratch_top);
1480 const member_types = try cg.id_scratch.addManyAsSlice(gpa, tuple.values.len);
1481
1482 var member_index: usize = 0;
1483 for (tuple.types.get(ip), tuple.values.get(ip)) |field_ty, field_val| {
1484 if (field_val != .none or !Type.fromInterned(field_ty).hasRuntimeBits(zcu)) continue;
1485
1486 member_types[member_index] = try cg.resolveType(.fromInterned(field_ty), .indirect);
1487 member_index += 1;
1488 }
1489
1490 const result_id = try cg.module.structType(
1491 member_types[0..member_index],
1492 null,
1493 null,
1494 .none,
1495 );
1496 const type_name = try cg.resolveTypeName(ty);
1497 defer gpa.free(type_name);
1498 try cg.module.debugName(result_id, type_name);
1499 return result_id;
1500 },
1501 .struct_type => ip.loadStructType(ty.toIntern()),
1502 else => unreachable,
1503 };
1504
1505 if (struct_type.layout == .@"packed") {
1506 return try cg.resolveType(.fromInterned(struct_type.backingIntTypeUnordered(ip)), .direct);
1507 }
1508
1509 var member_types = std.array_list.Managed(Id).init(gpa);
1510 defer member_types.deinit();
1511
1512 var member_names = std.array_list.Managed([]const u8).init(gpa);
1513 defer member_names.deinit();
1514
1515 var member_offsets = std.array_list.Managed(u32).init(gpa);
1516 defer member_offsets.deinit();
1517
1518 var it = struct_type.iterateRuntimeOrder(ip);
1519 while (it.next()) |field_index| {
1520 const field_ty: Type = .fromInterned(struct_type.field_types.get(ip)[field_index]);
1521 if (!field_ty.hasRuntimeBitsIgnoreComptime(zcu)) continue;
1522
1523 const field_name = struct_type.fieldName(ip, field_index);
1524 try member_types.append(try cg.resolveType(field_ty, .indirect));
1525 try member_names.append(field_name.toSlice(ip));
1526 try member_offsets.append(@intCast(ty.structFieldOffset(field_index, zcu)));
1527 }
1528
1529 const result_id = try cg.module.structType(
1530 member_types.items,
1531 member_names.items,
1532 member_offsets.items,
1533 ty.toIntern(),
1534 );
1535
1536 const type_name = try cg.resolveTypeName(ty);
1537 defer gpa.free(type_name);
1538 try cg.module.debugName(result_id, type_name);
1539
1540 return result_id;
1541 },
1542 .optional => {
1543 const payload_ty = ty.optionalChild(zcu);
1544 if (!payload_ty.hasRuntimeBitsIgnoreComptime(zcu)) {
1545 // Just use a bool.
1546 // Note: Always generate the bool with indirect format, to save on some sanity
1547 // Perform the conversion to a direct bool when the field is extracted.
1548 return try cg.resolveType(.bool, .indirect);
1549 }
1550
1551 const payload_ty_id = try cg.resolveType(payload_ty, .indirect);
1552 if (ty.optionalReprIsPayload(zcu)) {
1553 // Optional is actually a pointer or a slice.
1554 return payload_ty_id;
1555 }
1556
1557 const bool_ty_id = try cg.resolveType(.bool, .indirect);
1558
1559 return try cg.module.structType(
1560 &.{ payload_ty_id, bool_ty_id },
1561 &.{ "payload", "valid" },
1562 null,
1563 .none,
1564 );
1565 },
1566 .@"union" => return try cg.resolveUnionType(ty),
1567 .error_set => {
1568 const err_int_ty = try pt.errorIntType();
1569 return try cg.resolveType(err_int_ty, repr);
1570 },
1571 .error_union => {
1572 const payload_ty = ty.errorUnionPayload(zcu);
1573 const err_ty = ty.errorUnionSet(zcu);
1574 const error_ty_id = try cg.resolveType(err_ty, .indirect);
1575
1576 const eu_layout = cg.errorUnionLayout(payload_ty);
1577 if (!eu_layout.payload_has_bits) {
1578 return error_ty_id;
1579 }
1580
1581 const payload_ty_id = try cg.resolveType(payload_ty, .indirect);
1582
1583 var member_types: [2]Id = undefined;
1584 var member_names: [2][]const u8 = undefined;
1585 if (eu_layout.error_first) {
1586 // Put the error first
1587 member_types = .{ error_ty_id, payload_ty_id };
1588 member_names = .{ "error", "payload" };
1589 // TODO: ABI padding?
1590 } else {
1591 // Put the payload first.
1592 member_types = .{ payload_ty_id, error_ty_id };
1593 member_names = .{ "payload", "error" };
1594 // TODO: ABI padding?
1595 }
1596
1597 return try cg.module.structType(&member_types, &member_names, null, .none);
1598 },
1599 .@"opaque" => {
1600 if (target.os.tag != .opencl) return cg.fail("cannot generate opaque type", .{});
1601 const type_name = try cg.resolveTypeName(ty);
1602 defer gpa.free(type_name);
1603 return try cg.module.opaqueType(type_name);
1604 },
1605
1606 .null,
1607 .undefined,
1608 .enum_literal,
1609 .comptime_float,
1610 .comptime_int,
1611 .type,
1612 => unreachable, // Must be comptime.
1613
1614 .frame, .@"anyframe" => unreachable, // TODO
1615 }
1616}
1617
1618const ErrorUnionLayout = struct {
1619 payload_has_bits: bool,
1620 error_first: bool,
1621
1622 fn errorFieldIndex(cg: @This()) u32 {
1623 assert(cg.payload_has_bits);
1624 return if (cg.error_first) 0 else 1;
1625 }
1626
1627 fn payloadFieldIndex(cg: @This()) u32 {
1628 assert(cg.payload_has_bits);
1629 return if (cg.error_first) 1 else 0;
1630 }
1631};
1632
1633fn errorUnionLayout(cg: *CodeGen, payload_ty: Type) ErrorUnionLayout {
1634 const zcu = cg.module.zcu;
1635
1636 const error_align = Type.abiAlignment(.anyerror, zcu);
1637 const payload_align = payload_ty.abiAlignment(zcu);
1638
1639 const error_first = error_align.compare(.gt, payload_align);
1640 return .{
1641 .payload_has_bits = payload_ty.hasRuntimeBitsIgnoreComptime(zcu),
1642 .error_first = error_first,
1643 };
1644}
1645
1646const UnionLayout = struct {
1647 /// If false, this union is represented
1648 /// by only an integer of the tag type.
1649 has_payload: bool,
1650 tag_size: u32,
1651 tag_index: u32,
1652 /// Note: This is the size of the payload type itcg, NOT the size of the ENTIRE payload.
1653 /// Use `has_payload` instead!!
1654 payload_ty: Type,
1655 payload_size: u32,
1656 payload_index: u32,
1657 payload_padding_size: u32,
1658 payload_padding_index: u32,
1659 padding_size: u32,
1660 padding_index: u32,
1661 total_fields: u32,
1662};
1663
1664fn unionLayout(cg: *CodeGen, ty: Type) UnionLayout {
1665 const zcu = cg.module.zcu;
1666 const ip = &zcu.intern_pool;
1667 const layout = ty.unionGetLayout(zcu);
1668 const union_obj = zcu.typeToUnion(ty).?;
1669
1670 var union_layout: UnionLayout = .{
1671 .has_payload = layout.payload_size != 0,
1672 .tag_size = @intCast(layout.tag_size),
1673 .tag_index = undefined,
1674 .payload_ty = undefined,
1675 .payload_size = undefined,
1676 .payload_index = undefined,
1677 .payload_padding_size = undefined,
1678 .payload_padding_index = undefined,
1679 .padding_size = @intCast(layout.padding),
1680 .padding_index = undefined,
1681 .total_fields = undefined,
1682 };
1683
1684 if (union_layout.has_payload) {
1685 const most_aligned_field = layout.most_aligned_field;
1686 const most_aligned_field_ty: Type = .fromInterned(union_obj.field_types.get(ip)[most_aligned_field]);
1687 union_layout.payload_ty = most_aligned_field_ty;
1688 union_layout.payload_size = @intCast(most_aligned_field_ty.abiSize(zcu));
1689 } else {
1690 union_layout.payload_size = 0;
1691 }
1692
1693 union_layout.payload_padding_size = @intCast(layout.payload_size - union_layout.payload_size);
1694
1695 const tag_first = layout.tag_align.compare(.gte, layout.payload_align);
1696 var field_index: u32 = 0;
1697
1698 if (union_layout.tag_size != 0 and tag_first) {
1699 union_layout.tag_index = field_index;
1700 field_index += 1;
1701 }
1702
1703 if (union_layout.payload_size != 0) {
1704 union_layout.payload_index = field_index;
1705 field_index += 1;
1706 }
1707
1708 if (union_layout.payload_padding_size != 0) {
1709 union_layout.payload_padding_index = field_index;
1710 field_index += 1;
1711 }
1712
1713 if (union_layout.tag_size != 0 and !tag_first) {
1714 union_layout.tag_index = field_index;
1715 field_index += 1;
1716 }
1717
1718 if (union_layout.padding_size != 0) {
1719 union_layout.padding_index = field_index;
1720 field_index += 1;
1721 }
1722
1723 union_layout.total_fields = field_index;
1724
1725 return union_layout;
1726}
1727
1728/// This structure represents a "temporary" value: Something we are currently
1729/// operating on. It typically lives no longer than the function that
1730/// implements a particular AIR operation. These are used to easier
1731/// implement vectorizable operations (see Vectorization and the build*
1732/// functions), and typically are only used for vectors of primitive types.
1733const Temporary = struct {
1734 /// The type of the temporary. This is here mainly
1735 /// for easier bookkeeping. Because we will never really
1736 /// store Temporaries, they only cause extra stack space,
1737 /// therefore no real storage is wasted.
1738 ty: Type,
1739 /// The value that this temporary holds. This is not necessarily
1740 /// a value that is actually usable, or a single value: It is virtual
1741 /// until materialize() is called, at which point is turned into
1742 /// the usual SPIR-V representation of `cg.ty`.
1743 value: Temporary.Value,
1744
1745 const Value = union(enum) {
1746 singleton: Id,
1747 exploded_vector: IdRange,
1748 };
1749
1750 fn init(ty: Type, singleton: Id) Temporary {
1751 return .{ .ty = ty, .value = .{ .singleton = singleton } };
1752 }
1753
1754 fn materialize(temp: Temporary, cg: *CodeGen) !Id {
1755 const gpa = cg.module.gpa;
1756 const zcu = cg.module.zcu;
1757 switch (temp.value) {
1758 .singleton => |id| return id,
1759 .exploded_vector => |range| {
1760 assert(temp.ty.isVector(zcu));
1761 assert(temp.ty.vectorLen(zcu) == range.len);
1762
1763 const scratch_top = cg.id_scratch.items.len;
1764 defer cg.id_scratch.shrinkRetainingCapacity(scratch_top);
1765 const constituents = try cg.id_scratch.addManyAsSlice(gpa, range.len);
1766 for (constituents, 0..range.len) |*id, i| {
1767 id.* = range.at(i);
1768 }
1769
1770 const result_ty_id = try cg.resolveType(temp.ty, .direct);
1771 return cg.constructComposite(result_ty_id, constituents);
1772 },
1773 }
1774 }
1775
1776 fn vectorization(temp: Temporary, cg: *CodeGen) Vectorization {
1777 return .fromType(temp.ty, cg);
1778 }
1779
1780 fn pun(temp: Temporary, new_ty: Type) Temporary {
1781 return .{
1782 .ty = new_ty,
1783 .value = temp.value,
1784 };
1785 }
1786
1787 /// 'Explode' a temporary into separate elements. This turns a vector
1788 /// into a bag of elements.
1789 fn explode(temp: Temporary, cg: *CodeGen) !IdRange {
1790 const zcu = cg.module.zcu;
1791
1792 // If the value is a scalar, then this is a no-op.
1793 if (!temp.ty.isVector(zcu)) {
1794 return switch (temp.value) {
1795 .singleton => |id| .{ .base = @intFromEnum(id), .len = 1 },
1796 .exploded_vector => |range| range,
1797 };
1798 }
1799
1800 const ty_id = try cg.resolveType(temp.ty.scalarType(zcu), .direct);
1801 const n = temp.ty.vectorLen(zcu);
1802 const results = cg.module.allocIds(n);
1803
1804 const id = switch (temp.value) {
1805 .singleton => |id| id,
1806 .exploded_vector => |range| return range,
1807 };
1808
1809 for (0..n) |i| {
1810 const indexes = [_]u32{@intCast(i)};
1811 try cg.body.emit(cg.module.gpa, .OpCompositeExtract, .{
1812 .id_result_type = ty_id,
1813 .id_result = results.at(i),
1814 .composite = id,
1815 .indexes = &indexes,
1816 });
1817 }
1818
1819 return results;
1820 }
1821};
1822
1823/// Initialize a `Temporary` from an AIR value.
1824fn temporary(cg: *CodeGen, inst: Air.Inst.Ref) !Temporary {
1825 return .{
1826 .ty = cg.typeOf(inst),
1827 .value = .{ .singleton = try cg.resolve(inst) },
1828 };
1829}
1830
1831/// This union describes how a particular operation should be vectorized.
1832/// That depends on the operation and number of components of the inputs.
1833const Vectorization = union(enum) {
1834 /// This is an operation between scalars.
1835 scalar,
1836 /// This operation is unrolled into separate operations.
1837 /// Inputs may still be SPIR-V vectors, for example,
1838 /// when the operation can't be vectorized in SPIR-V.
1839 /// Value is number of components.
1840 unrolled: u32,
1841
1842 /// Derive a vectorization from a particular type
1843 fn fromType(ty: Type, cg: *CodeGen) Vectorization {
1844 const zcu = cg.module.zcu;
1845 if (!ty.isVector(zcu)) return .scalar;
1846 return .{ .unrolled = ty.vectorLen(zcu) };
1847 }
1848
1849 /// Given two vectorization methods, compute a "unification": a fallback
1850 /// that works for both, according to the following rules:
1851 /// - Scalars may broadcast
1852 /// - SPIR-V vectorized operations will unroll
1853 /// - Prefer scalar > unrolled
1854 fn unify(a: Vectorization, b: Vectorization) Vectorization {
1855 if (a == .scalar and b == .scalar) return .scalar;
1856 if (a == .unrolled or b == .unrolled) {
1857 if (a == .unrolled and b == .unrolled) assert(a.components() == b.components());
1858 if (a == .unrolled) return .{ .unrolled = a.components() };
1859 return .{ .unrolled = b.components() };
1860 }
1861 unreachable;
1862 }
1863
1864 /// Query the number of components that inputs of this operation have.
1865 /// Note: for broadcasting scalars, this returns the number of elements
1866 /// that the broadcasted vector would have.
1867 fn components(vec: Vectorization) u32 {
1868 return switch (vec) {
1869 .scalar => 1,
1870 .unrolled => |n| n,
1871 };
1872 }
1873
1874 /// Turns `ty` into the result-type of the entire operation.
1875 /// `ty` may be a scalar or vector, it doesn't matter.
1876 fn resultType(vec: Vectorization, cg: *CodeGen, ty: Type) !Type {
1877 const pt = cg.pt;
1878 const zcu = cg.module.zcu;
1879 const scalar_ty = ty.scalarType(zcu);
1880 return switch (vec) {
1881 .scalar => scalar_ty,
1882 .unrolled => |n| try pt.vectorType(.{ .len = n, .child = scalar_ty.toIntern() }),
1883 };
1884 }
1885
1886 /// Before a temporary can be used, some setup may need to be one. This function implements
1887 /// this setup, and returns a new type that holds the relevant information on how to access
1888 /// elements of the input.
1889 fn prepare(vec: Vectorization, cg: *CodeGen, tmp: Temporary) !PreparedOperand {
1890 const zcu = cg.module.zcu;
1891 const is_vector = tmp.ty.isVector(zcu);
1892 const value: PreparedOperand.Value = switch (tmp.value) {
1893 .singleton => |id| switch (vec) {
1894 .scalar => blk: {
1895 assert(!is_vector);
1896 break :blk .{ .scalar = id };
1897 },
1898 .unrolled => blk: {
1899 if (is_vector) break :blk .{ .vector_exploded = try tmp.explode(cg) };
1900 break :blk .{ .scalar_broadcast = id };
1901 },
1902 },
1903 .exploded_vector => |range| switch (vec) {
1904 .scalar => unreachable,
1905 .unrolled => |n| blk: {
1906 assert(range.len == n);
1907 break :blk .{ .vector_exploded = range };
1908 },
1909 },
1910 };
1911
1912 return .{
1913 .ty = tmp.ty,
1914 .value = value,
1915 };
1916 }
1917
1918 /// Finalize the results of an operation back into a temporary. `results` is
1919 /// a list of result-ids of the operation.
1920 fn finalize(vec: Vectorization, ty: Type, results: IdRange) Temporary {
1921 assert(vec.components() == results.len);
1922 return .{
1923 .ty = ty,
1924 .value = switch (vec) {
1925 .scalar => .{ .singleton = results.at(0) },
1926 .unrolled => .{ .exploded_vector = results },
1927 },
1928 };
1929 }
1930
1931 /// This struct represents an operand that has gone through some setup, and is
1932 /// ready to be used as part of an operation.
1933 const PreparedOperand = struct {
1934 ty: Type,
1935 value: PreparedOperand.Value,
1936
1937 /// The types of value that a prepared operand can hold internally. Depends
1938 /// on the operation and input value.
1939 const Value = union(enum) {
1940 /// A single scalar value that is used by a scalar operation.
1941 scalar: Id,
1942 /// A single scalar that is broadcasted in an unrolled operation.
1943 scalar_broadcast: Id,
1944 /// A vector represented by a consecutive list of IDs that is used in an unrolled operation.
1945 vector_exploded: IdRange,
1946 };
1947
1948 /// Query the value at a particular index of the operation. Note that
1949 /// the index is *not* the component/lane, but the index of the *operation*.
1950 fn at(op: PreparedOperand, i: usize) Id {
1951 switch (op.value) {
1952 .scalar => |id| {
1953 assert(i == 0);
1954 return id;
1955 },
1956 .scalar_broadcast => |id| return id,
1957 .vector_exploded => |range| return range.at(i),
1958 }
1959 }
1960 };
1961};
1962
1963/// A utility function to compute the vectorization style of
1964/// a list of values. These values may be any of the following:
1965/// - A `Vectorization` instance
1966/// - A Type, in which case the vectorization is computed via `Vectorization.fromType`.
1967/// - A Temporary, in which case the vectorization is computed via `Temporary.vectorization`.
1968fn vectorization(cg: *CodeGen, args: anytype) Vectorization {
1969 var v: Vectorization = undefined;
1970 assert(args.len >= 1);
1971 inline for (args, 0..) |arg, i| {
1972 const iv: Vectorization = switch (@TypeOf(arg)) {
1973 Vectorization => arg,
1974 Type => Vectorization.fromType(arg, cg),
1975 Temporary => arg.vectorization(cg),
1976 else => @compileError("invalid type"),
1977 };
1978 if (i == 0) {
1979 v = iv;
1980 } else {
1981 v = v.unify(iv);
1982 }
1983 }
1984 return v;
1985}
1986
1987/// This function builds an OpSConvert of OpUConvert depending on the
1988/// signedness of the types.
1989fn buildConvert(cg: *CodeGen, dst_ty: Type, src: Temporary) !Temporary {
1990 const zcu = cg.module.zcu;
1991
1992 const dst_ty_id = try cg.resolveType(dst_ty.scalarType(zcu), .direct);
1993 const src_ty_id = try cg.resolveType(src.ty.scalarType(zcu), .direct);
1994
1995 const v = cg.vectorization(.{ dst_ty, src });
1996 const result_ty = try v.resultType(cg, dst_ty);
1997
1998 // We can directly compare integers, because those type-IDs are cached.
1999 if (dst_ty_id == src_ty_id) {
2000 // Nothing to do, type-pun to the right value.
2001 // Note, Caller guarantees that the types fit (or caller will normalize after),
2002 // so we don't have to normalize here.
2003 // Note, dst_ty may be a scalar type even if we expect a vector, so we have to
2004 // convert to the right type here.
2005 return src.pun(result_ty);
2006 }
2007
2008 const ops = v.components();
2009 const results = cg.module.allocIds(ops);
2010
2011 const op_result_ty = dst_ty.scalarType(zcu);
2012 const op_result_ty_id = try cg.resolveType(op_result_ty, .direct);
2013
2014 const opcode: Opcode = blk: {
2015 if (dst_ty.scalarType(zcu).isAnyFloat()) break :blk .OpFConvert;
2016 if (dst_ty.scalarType(zcu).isSignedInt(zcu)) break :blk .OpSConvert;
2017 break :blk .OpUConvert;
2018 };
2019
2020 const op_src = try v.prepare(cg, src);
2021
2022 for (0..ops) |i| {
2023 try cg.body.emitRaw(cg.module.gpa, opcode, 3);
2024 cg.body.writeOperand(Id, op_result_ty_id);
2025 cg.body.writeOperand(Id, results.at(i));
2026 cg.body.writeOperand(Id, op_src.at(i));
2027 }
2028
2029 return v.finalize(result_ty, results);
2030}
2031
2032fn buildFma(cg: *CodeGen, a: Temporary, b: Temporary, c: Temporary) !Temporary {
2033 const zcu = cg.module.zcu;
2034 const target = cg.module.zcu.getTarget();
2035
2036 const v = cg.vectorization(.{ a, b, c });
2037 const ops = v.components();
2038 const results = cg.module.allocIds(ops);
2039
2040 const op_result_ty = a.ty.scalarType(zcu);
2041 const op_result_ty_id = try cg.resolveType(op_result_ty, .direct);
2042 const result_ty = try v.resultType(cg, a.ty);
2043
2044 const op_a = try v.prepare(cg, a);
2045 const op_b = try v.prepare(cg, b);
2046 const op_c = try v.prepare(cg, c);
2047
2048 const set = try cg.importExtendedSet();
2049 const opcode: u32 = switch (target.os.tag) {
2050 .opencl => @intFromEnum(spec.OpenClOpcode.fma),
2051 // NOTE: Vulkan's FMA instruction does *NOT* produce the right values!
2052 // its precision guarantees do NOT match zigs and it does NOT match OpenCLs!
2053 // it needs to be emulated!
2054 .vulkan, .opengl => @intFromEnum(spec.GlslOpcode.Fma),
2055 else => unreachable,
2056 };
2057
2058 for (0..ops) |i| {
2059 try cg.body.emit(cg.module.gpa, .OpExtInst, .{
2060 .id_result_type = op_result_ty_id,
2061 .id_result = results.at(i),
2062 .set = set,
2063 .instruction = .{ .inst = opcode },
2064 .id_ref_4 = &.{ op_a.at(i), op_b.at(i), op_c.at(i) },
2065 });
2066 }
2067
2068 return v.finalize(result_ty, results);
2069}
2070
2071fn buildSelect(cg: *CodeGen, condition: Temporary, lhs: Temporary, rhs: Temporary) !Temporary {
2072 const zcu = cg.module.zcu;
2073
2074 const v = cg.vectorization(.{ condition, lhs, rhs });
2075 const ops = v.components();
2076 const results = cg.module.allocIds(ops);
2077
2078 const op_result_ty = lhs.ty.scalarType(zcu);
2079 const op_result_ty_id = try cg.resolveType(op_result_ty, .direct);
2080 const result_ty = try v.resultType(cg, lhs.ty);
2081
2082 assert(condition.ty.scalarType(zcu).zigTypeTag(zcu) == .bool);
2083
2084 const cond = try v.prepare(cg, condition);
2085 const object_1 = try v.prepare(cg, lhs);
2086 const object_2 = try v.prepare(cg, rhs);
2087
2088 for (0..ops) |i| {
2089 try cg.body.emit(cg.module.gpa, .OpSelect, .{
2090 .id_result_type = op_result_ty_id,
2091 .id_result = results.at(i),
2092 .condition = cond.at(i),
2093 .object_1 = object_1.at(i),
2094 .object_2 = object_2.at(i),
2095 });
2096 }
2097
2098 return v.finalize(result_ty, results);
2099}
2100
2101fn buildCmp(cg: *CodeGen, opcode: Opcode, lhs: Temporary, rhs: Temporary) !Temporary {
2102 const v = cg.vectorization(.{ lhs, rhs });
2103 const ops = v.components();
2104 const results = cg.module.allocIds(ops);
2105
2106 const op_result_ty: Type = .bool;
2107 const op_result_ty_id = try cg.resolveType(op_result_ty, .direct);
2108 const result_ty = try v.resultType(cg, Type.bool);
2109
2110 const op_lhs = try v.prepare(cg, lhs);
2111 const op_rhs = try v.prepare(cg, rhs);
2112
2113 for (0..ops) |i| {
2114 try cg.body.emitRaw(cg.module.gpa, opcode, 4);
2115 cg.body.writeOperand(Id, op_result_ty_id);
2116 cg.body.writeOperand(Id, results.at(i));
2117 cg.body.writeOperand(Id, op_lhs.at(i));
2118 cg.body.writeOperand(Id, op_rhs.at(i));
2119 }
2120
2121 return v.finalize(result_ty, results);
2122}
2123
2124const UnaryOp = enum {
2125 l_not,
2126 bit_not,
2127 i_neg,
2128 f_neg,
2129 i_abs,
2130 f_abs,
2131 clz,
2132 ctz,
2133 floor,
2134 ceil,
2135 trunc,
2136 round,
2137 sqrt,
2138 sin,
2139 cos,
2140 tan,
2141 exp,
2142 exp2,
2143 log,
2144 log2,
2145 log10,
2146
2147 pub fn extInstOpcode(op: UnaryOp, target: *const std.Target) ?u32 {
2148 return switch (target.os.tag) {
2149 .opencl => @intFromEnum(@as(spec.OpenClOpcode, switch (op) {
2150 .i_abs => .s_abs,
2151 .f_abs => .fabs,
2152 .clz => .clz,
2153 .ctz => .ctz,
2154 .floor => .floor,
2155 .ceil => .ceil,
2156 .trunc => .trunc,
2157 .round => .round,
2158 .sqrt => .sqrt,
2159 .sin => .sin,
2160 .cos => .cos,
2161 .tan => .tan,
2162 .exp => .exp,
2163 .exp2 => .exp2,
2164 .log => .log,
2165 .log2 => .log2,
2166 .log10 => .log10,
2167 else => return null,
2168 })),
2169 // Note: We'll need to check these for floating point accuracy
2170 // Vulkan does not put tight requirements on these, for correction
2171 // we might want to emulate them at some point.
2172 .vulkan, .opengl => @intFromEnum(@as(spec.GlslOpcode, switch (op) {
2173 .i_abs => .SAbs,
2174 .f_abs => .FAbs,
2175 .floor => .Floor,
2176 .ceil => .Ceil,
2177 .trunc => .Trunc,
2178 .round => .Round,
2179 .sin => .Sin,
2180 .cos => .Cos,
2181 .tan => .Tan,
2182 .sqrt => .Sqrt,
2183 .exp => .Exp,
2184 .exp2 => .Exp2,
2185 .log => .Log,
2186 .log2 => .Log2,
2187 else => return null,
2188 })),
2189 else => unreachable,
2190 };
2191 }
2192};
2193
2194fn buildUnary(cg: *CodeGen, op: UnaryOp, operand: Temporary) !Temporary {
2195 const zcu = cg.module.zcu;
2196 const target = cg.module.zcu.getTarget();
2197 const v = cg.vectorization(.{operand});
2198 const ops = v.components();
2199 const results = cg.module.allocIds(ops);
2200 const op_result_ty = operand.ty.scalarType(zcu);
2201 const op_result_ty_id = try cg.resolveType(op_result_ty, .direct);
2202 const result_ty = try v.resultType(cg, operand.ty);
2203 const op_operand = try v.prepare(cg, operand);
2204
2205 if (op.extInstOpcode(target)) |opcode| {
2206 const set = try cg.importExtendedSet();
2207 for (0..ops) |i| {
2208 try cg.body.emit(cg.module.gpa, .OpExtInst, .{
2209 .id_result_type = op_result_ty_id,
2210 .id_result = results.at(i),
2211 .set = set,
2212 .instruction = .{ .inst = opcode },
2213 .id_ref_4 = &.{op_operand.at(i)},
2214 });
2215 }
2216 } else {
2217 const opcode: Opcode = switch (op) {
2218 .l_not => .OpLogicalNot,
2219 .bit_not => .OpNot,
2220 .i_neg => .OpSNegate,
2221 .f_neg => .OpFNegate,
2222 else => return cg.todo(
2223 "implement unary operation '{s}' for {s} os",
2224 .{ @tagName(op), @tagName(target.os.tag) },
2225 ),
2226 };
2227 for (0..ops) |i| {
2228 try cg.body.emitRaw(cg.module.gpa, opcode, 3);
2229 cg.body.writeOperand(Id, op_result_ty_id);
2230 cg.body.writeOperand(Id, results.at(i));
2231 cg.body.writeOperand(Id, op_operand.at(i));
2232 }
2233 }
2234
2235 return v.finalize(result_ty, results);
2236}
2237
2238fn buildBinary(cg: *CodeGen, opcode: Opcode, lhs: Temporary, rhs: Temporary) !Temporary {
2239 const zcu = cg.module.zcu;
2240
2241 const v = cg.vectorization(.{ lhs, rhs });
2242 const ops = v.components();
2243 const results = cg.module.allocIds(ops);
2244
2245 const op_result_ty = lhs.ty.scalarType(zcu);
2246 const op_result_ty_id = try cg.resolveType(op_result_ty, .direct);
2247 const result_ty = try v.resultType(cg, lhs.ty);
2248
2249 const op_lhs = try v.prepare(cg, lhs);
2250 const op_rhs = try v.prepare(cg, rhs);
2251
2252 for (0..ops) |i| {
2253 try cg.body.emitRaw(cg.module.gpa, opcode, 4);
2254 cg.body.writeOperand(Id, op_result_ty_id);
2255 cg.body.writeOperand(Id, results.at(i));
2256 cg.body.writeOperand(Id, op_lhs.at(i));
2257 cg.body.writeOperand(Id, op_rhs.at(i));
2258 }
2259
2260 return v.finalize(result_ty, results);
2261}
2262
2263/// This function builds an extended multiplication, either OpSMulExtended or OpUMulExtended on Vulkan,
2264/// or OpIMul and s_mul_hi or u_mul_hi on OpenCL.
2265fn buildWideMul(
2266 cg: *CodeGen,
2267 signedness: std.builtin.Signedness,
2268 lhs: Temporary,
2269 rhs: Temporary,
2270) !struct { Temporary, Temporary } {
2271 const pt = cg.pt;
2272 const zcu = cg.module.zcu;
2273 const target = cg.module.zcu.getTarget();
2274 const ip = &zcu.intern_pool;
2275
2276 const v = lhs.vectorization(cg).unify(rhs.vectorization(cg));
2277 const ops = v.components();
2278
2279 const arith_op_ty = lhs.ty.scalarType(zcu);
2280 const arith_op_ty_id = try cg.resolveType(arith_op_ty, .direct);
2281
2282 const lhs_op = try v.prepare(cg, lhs);
2283 const rhs_op = try v.prepare(cg, rhs);
2284
2285 const value_results = cg.module.allocIds(ops);
2286 const overflow_results = cg.module.allocIds(ops);
2287
2288 switch (target.os.tag) {
2289 .opencl => {
2290 // Currently, SPIRV-LLVM-Translator based backends cannot deal with OpSMulExtended and
2291 // OpUMulExtended. For these we will use the OpenCL s_mul_hi to compute the high-order bits
2292 // instead.
2293 const set = try cg.importExtendedSet();
2294 const overflow_inst: spec.OpenClOpcode = switch (signedness) {
2295 .signed => .s_mul_hi,
2296 .unsigned => .u_mul_hi,
2297 };
2298
2299 for (0..ops) |i| {
2300 try cg.body.emit(cg.module.gpa, .OpIMul, .{
2301 .id_result_type = arith_op_ty_id,
2302 .id_result = value_results.at(i),
2303 .operand_1 = lhs_op.at(i),
2304 .operand_2 = rhs_op.at(i),
2305 });
2306
2307 try cg.body.emit(cg.module.gpa, .OpExtInst, .{
2308 .id_result_type = arith_op_ty_id,
2309 .id_result = overflow_results.at(i),
2310 .set = set,
2311 .instruction = .{ .inst = @intFromEnum(overflow_inst) },
2312 .id_ref_4 = &.{ lhs_op.at(i), rhs_op.at(i) },
2313 });
2314 }
2315 },
2316 .vulkan, .opengl => {
2317 // Operations return a struct{T, T}
2318 // where T is maybe vectorized.
2319 const op_result_ty: Type = .fromInterned(try ip.getTupleType(zcu.gpa, pt.tid, .{
2320 .types = &.{ arith_op_ty.toIntern(), arith_op_ty.toIntern() },
2321 .values = &.{ .none, .none },
2322 }));
2323 const op_result_ty_id = try cg.resolveType(op_result_ty, .direct);
2324
2325 const opcode: Opcode = switch (signedness) {
2326 .signed => .OpSMulExtended,
2327 .unsigned => .OpUMulExtended,
2328 };
2329
2330 for (0..ops) |i| {
2331 const op_result = cg.module.allocId();
2332
2333 try cg.body.emitRaw(cg.module.gpa, opcode, 4);
2334 cg.body.writeOperand(Id, op_result_ty_id);
2335 cg.body.writeOperand(Id, op_result);
2336 cg.body.writeOperand(Id, lhs_op.at(i));
2337 cg.body.writeOperand(Id, rhs_op.at(i));
2338
2339 // The above operation returns a struct. We might want to expand
2340 // Temporary to deal with the fact that these are structs eventually,
2341 // but for now, take the struct apart and return two separate vectors.
2342
2343 try cg.body.emit(cg.module.gpa, .OpCompositeExtract, .{
2344 .id_result_type = arith_op_ty_id,
2345 .id_result = value_results.at(i),
2346 .composite = op_result,
2347 .indexes = &.{0},
2348 });
2349
2350 try cg.body.emit(cg.module.gpa, .OpCompositeExtract, .{
2351 .id_result_type = arith_op_ty_id,
2352 .id_result = overflow_results.at(i),
2353 .composite = op_result,
2354 .indexes = &.{1},
2355 });
2356 }
2357 },
2358 else => unreachable,
2359 }
2360
2361 const result_ty = try v.resultType(cg, lhs.ty);
2362 return .{
2363 v.finalize(result_ty, value_results),
2364 v.finalize(result_ty, overflow_results),
2365 };
2366}
2367
2368/// The SPIR-V backend is not yet advanced enough to support the std testing infrastructure.
2369/// In order to be able to run tests, we "temporarily" lower test kernels into separate entry-
2370/// points. The test executor will then be able to invoke these to run the tests.
2371/// Note that tests are lowered according to std.builtin.TestFn, which is `fn () anyerror!void`.
2372/// (anyerror!void has the same layout as anyerror).
2373/// Each test declaration generates a function like.
2374/// %anyerror = OpTypeInt 0 16
2375/// %p_invocation_globals_struct_ty = ...
2376/// %p_anyerror = OpTypePointer CrossWorkgroup %anyerror
2377/// %K = OpTypeFunction %void %p_invocation_globals_struct_ty %p_anyerror
2378///
2379/// %test = OpFunction %void %K
2380/// %p_invocation_globals = OpFunctionParameter p_invocation_globals_struct_ty
2381/// %p_err = OpFunctionParameter %p_anyerror
2382/// %lbl = OpLabel
2383/// %result = OpFunctionCall %anyerror %func %p_invocation_globals
2384/// OpStore %p_err %result
2385/// OpFunctionEnd
2386/// TODO is to also write out the error as a function call parameter, and to somehow fetch
2387/// the name of an error in the text executor.
2388fn generateTestEntryPoint(
2389 cg: *CodeGen,
2390 name: []const u8,
2391 spv_decl_index: Module.Decl.Index,
2392 test_id: Id,
2393) !void {
2394 const gpa = cg.module.gpa;
2395 const zcu = cg.module.zcu;
2396 const target = cg.module.zcu.getTarget();
2397
2398 const anyerror_ty_id = try cg.resolveType(.anyerror, .direct);
2399 const ptr_anyerror_ty = try cg.pt.ptrType(.{
2400 .child = .anyerror_type,
2401 .flags = .{ .address_space = .global },
2402 });
2403 const ptr_anyerror_ty_id = try cg.resolveType(ptr_anyerror_ty, .direct);
2404
2405 const kernel_id = cg.module.declPtr(spv_decl_index).result_id;
2406
2407 const section = &cg.module.sections.functions;
2408
2409 const p_error_id = cg.module.allocId();
2410 switch (target.os.tag) {
2411 .opencl, .amdhsa => {
2412 const void_ty_id = try cg.resolveType(.void, .direct);
2413 const kernel_proto_ty_id = try cg.module.functionType(void_ty_id, &.{ptr_anyerror_ty_id});
2414
2415 try section.emit(gpa, .OpFunction, .{
2416 .id_result_type = try cg.resolveType(.void, .direct),
2417 .id_result = kernel_id,
2418 .function_control = .{},
2419 .function_type = kernel_proto_ty_id,
2420 });
2421
2422 try section.emit(gpa, .OpFunctionParameter, .{
2423 .id_result_type = ptr_anyerror_ty_id,
2424 .id_result = p_error_id,
2425 });
2426
2427 try section.emit(gpa, .OpLabel, .{
2428 .id_result = cg.module.allocId(),
2429 });
2430 },
2431 .vulkan, .opengl => {
2432 if (cg.module.error_buffer == null) {
2433 const spv_err_decl_index = try cg.module.allocDecl(.global);
2434 const err_buf_result_id = cg.module.declPtr(spv_err_decl_index).result_id;
2435
2436 const buffer_struct_ty_id = try cg.module.structType(
2437 &.{anyerror_ty_id},
2438 &.{"error_out"},
2439 null,
2440 .none,
2441 );
2442 try cg.module.decorate(buffer_struct_ty_id, .block);
2443 try cg.module.decorateMember(buffer_struct_ty_id, 0, .{ .offset = .{ .byte_offset = 0 } });
2444
2445 const ptr_buffer_struct_ty_id = cg.module.allocId();
2446 try cg.module.sections.globals.emit(gpa, .OpTypePointer, .{
2447 .id_result = ptr_buffer_struct_ty_id,
2448 .storage_class = cg.module.storageClass(.global),
2449 .type = buffer_struct_ty_id,
2450 });
2451
2452 try cg.module.sections.globals.emit(gpa, .OpVariable, .{
2453 .id_result_type = ptr_buffer_struct_ty_id,
2454 .id_result = err_buf_result_id,
2455 .storage_class = cg.module.storageClass(.global),
2456 });
2457 try cg.module.decorate(err_buf_result_id, .{ .descriptor_set = .{ .descriptor_set = 0 } });
2458 try cg.module.decorate(err_buf_result_id, .{ .binding = .{ .binding_point = 0 } });
2459
2460 cg.module.error_buffer = spv_err_decl_index;
2461 }
2462
2463 try cg.module.sections.execution_modes.emit(gpa, .OpExecutionMode, .{
2464 .entry_point = kernel_id,
2465 .mode = .{ .local_size = .{
2466 .x_size = 1,
2467 .y_size = 1,
2468 .z_size = 1,
2469 } },
2470 });
2471
2472 const void_ty_id = try cg.resolveType(.void, .direct);
2473 const kernel_proto_ty_id = try cg.module.functionType(void_ty_id, &.{});
2474 try section.emit(gpa, .OpFunction, .{
2475 .id_result_type = try cg.resolveType(.void, .direct),
2476 .id_result = kernel_id,
2477 .function_control = .{},
2478 .function_type = kernel_proto_ty_id,
2479 });
2480 try section.emit(gpa, .OpLabel, .{
2481 .id_result = cg.module.allocId(),
2482 });
2483
2484 const spv_err_decl_index = cg.module.error_buffer.?;
2485 const buffer_id = cg.module.declPtr(spv_err_decl_index).result_id;
2486 try cg.module.decl_deps.append(gpa, spv_err_decl_index);
2487
2488 const zero_id = try cg.constInt(.u32, 0);
2489 try section.emit(gpa, .OpInBoundsAccessChain, .{
2490 .id_result_type = ptr_anyerror_ty_id,
2491 .id_result = p_error_id,
2492 .base = buffer_id,
2493 .indexes = &.{zero_id},
2494 });
2495 },
2496 else => unreachable,
2497 }
2498
2499 const error_id = cg.module.allocId();
2500 try section.emit(gpa, .OpFunctionCall, .{
2501 .id_result_type = anyerror_ty_id,
2502 .id_result = error_id,
2503 .function = test_id,
2504 });
2505 // Note: Convert to direct not required.
2506 try section.emit(gpa, .OpStore, .{
2507 .pointer = p_error_id,
2508 .object = error_id,
2509 .memory_access = .{
2510 .aligned = .{ .literal_integer = @intCast(Type.abiAlignment(.anyerror, zcu).toByteUnits().?) },
2511 },
2512 });
2513 try section.emit(gpa, .OpReturn, {});
2514 try section.emit(gpa, .OpFunctionEnd, {});
2515
2516 // Just generate a quick other name because the intel runtime crashes when the entry-
2517 // point name is the same as a different OpName.
2518 const test_name = try std.fmt.allocPrint(cg.module.arena, "test {s}", .{name});
2519
2520 const execution_mode: spec.ExecutionModel = switch (target.os.tag) {
2521 .vulkan, .opengl => .gl_compute,
2522 .opencl, .amdhsa => .kernel,
2523 else => unreachable,
2524 };
2525
2526 try cg.module.declareEntryPoint(spv_decl_index, test_name, execution_mode, null);
2527}
2528
2529fn intFromBool(cg: *CodeGen, value: Temporary, result_ty: Type) !Temporary {
2530 const zero_id = try cg.constInt(result_ty, 0);
2531 const one_id = try cg.constInt(result_ty, 1);
2532
2533 return try cg.buildSelect(
2534 value,
2535 Temporary.init(result_ty, one_id),
2536 Temporary.init(result_ty, zero_id),
2537 );
2538}
2539
2540/// Convert representation from indirect (in memory) to direct (in 'register')
2541/// This converts the argument type from resolveType(ty, .indirect) to resolveType(ty, .direct).
2542fn convertToDirect(cg: *CodeGen, ty: Type, operand_id: Id) !Id {
2543 const pt = cg.pt;
2544 const zcu = cg.module.zcu;
2545 switch (ty.scalarType(zcu).zigTypeTag(zcu)) {
2546 .bool => {
2547 const false_id = try cg.constBool(false, .indirect);
2548 const operand_ty = blk: {
2549 if (!ty.isVector(zcu)) break :blk Type.u1;
2550 break :blk try pt.vectorType(.{
2551 .len = ty.vectorLen(zcu),
2552 .child = .u1_type,
2553 });
2554 };
2555
2556 const result = try cg.buildCmp(
2557 .OpINotEqual,
2558 Temporary.init(operand_ty, operand_id),
2559 Temporary.init(.u1, false_id),
2560 );
2561 return try result.materialize(cg);
2562 },
2563 else => return operand_id,
2564 }
2565}
2566
2567/// Convert representation from direct (in 'register) to direct (in memory)
2568/// This converts the argument type from resolveType(ty, .direct) to resolveType(ty, .indirect).
2569fn convertToIndirect(cg: *CodeGen, ty: Type, operand_id: Id) !Id {
2570 const zcu = cg.module.zcu;
2571 switch (ty.scalarType(zcu).zigTypeTag(zcu)) {
2572 .bool => {
2573 const result = try cg.intFromBool(.init(ty, operand_id), .u1);
2574 return try result.materialize(cg);
2575 },
2576 else => return operand_id,
2577 }
2578}
2579
2580fn extractField(cg: *CodeGen, result_ty: Type, object: Id, field: u32) !Id {
2581 const result_ty_id = try cg.resolveType(result_ty, .indirect);
2582 const result_id = cg.module.allocId();
2583 const indexes = [_]u32{field};
2584 try cg.body.emit(cg.module.gpa, .OpCompositeExtract, .{
2585 .id_result_type = result_ty_id,
2586 .id_result = result_id,
2587 .composite = object,
2588 .indexes = &indexes,
2589 });
2590 // Convert bools; direct structs have their field types as indirect values.
2591 return try cg.convertToDirect(result_ty, result_id);
2592}
2593
2594fn extractVectorComponent(cg: *CodeGen, result_ty: Type, vector_id: Id, field: u32) !Id {
2595 const result_ty_id = try cg.resolveType(result_ty, .direct);
2596 const result_id = cg.module.allocId();
2597 const indexes = [_]u32{field};
2598 try cg.body.emit(cg.module.gpa, .OpCompositeExtract, .{
2599 .id_result_type = result_ty_id,
2600 .id_result = result_id,
2601 .composite = vector_id,
2602 .indexes = &indexes,
2603 });
2604 // Vector components are already stored in direct representation.
2605 return result_id;
2606}
2607
2608const MemoryOptions = struct {
2609 is_volatile: bool = false,
2610};
2611
2612fn load(cg: *CodeGen, value_ty: Type, ptr_id: Id, options: MemoryOptions) !Id {
2613 const zcu = cg.module.zcu;
2614 const alignment: u32 = @intCast(value_ty.abiAlignment(zcu).toByteUnits().?);
2615 const indirect_value_ty_id = try cg.resolveType(value_ty, .indirect);
2616 const result_id = cg.module.allocId();
2617 const access: spec.MemoryAccess.Extended = .{
2618 .@"volatile" = options.is_volatile,
2619 .aligned = .{ .literal_integer = alignment },
2620 };
2621 try cg.body.emit(cg.module.gpa, .OpLoad, .{
2622 .id_result_type = indirect_value_ty_id,
2623 .id_result = result_id,
2624 .pointer = ptr_id,
2625 .memory_access = access,
2626 });
2627 return try cg.convertToDirect(value_ty, result_id);
2628}
2629
2630fn store(cg: *CodeGen, value_ty: Type, ptr_id: Id, value_id: Id, options: MemoryOptions) !void {
2631 const indirect_value_id = try cg.convertToIndirect(value_ty, value_id);
2632 const access: spec.MemoryAccess.Extended = .{ .@"volatile" = options.is_volatile };
2633 try cg.body.emit(cg.module.gpa, .OpStore, .{
2634 .pointer = ptr_id,
2635 .object = indirect_value_id,
2636 .memory_access = access,
2637 });
2638}
2639
2640fn genBody(cg: *CodeGen, body: []const Air.Inst.Index) !void {
2641 for (body) |inst| {
2642 try cg.genInst(inst);
2643 }
2644}
2645
2646fn genInst(cg: *CodeGen, inst: Air.Inst.Index) Error!void {
2647 const gpa = cg.module.gpa;
2648 const zcu = cg.module.zcu;
2649 const ip = &zcu.intern_pool;
2650 if (cg.liveness.isUnused(inst) and !cg.air.mustLower(inst, ip))
2651 return;
2652
2653 const air_tags = cg.air.instructions.items(.tag);
2654 const maybe_result_id: ?Id = switch (air_tags[@intFromEnum(inst)]) {
2655 // zig fmt: off
2656 .add, .add_wrap, .add_optimized => try cg.airArithOp(inst, .OpFAdd, .OpIAdd, .OpIAdd),
2657 .sub, .sub_wrap, .sub_optimized => try cg.airArithOp(inst, .OpFSub, .OpISub, .OpISub),
2658 .mul, .mul_wrap, .mul_optimized => try cg.airArithOp(inst, .OpFMul, .OpIMul, .OpIMul),
2659
2660 .sqrt => try cg.airUnOpSimple(inst, .sqrt),
2661 .sin => try cg.airUnOpSimple(inst, .sin),
2662 .cos => try cg.airUnOpSimple(inst, .cos),
2663 .tan => try cg.airUnOpSimple(inst, .tan),
2664 .exp => try cg.airUnOpSimple(inst, .exp),
2665 .exp2 => try cg.airUnOpSimple(inst, .exp2),
2666 .log => try cg.airUnOpSimple(inst, .log),
2667 .log2 => try cg.airUnOpSimple(inst, .log2),
2668 .log10 => try cg.airUnOpSimple(inst, .log10),
2669 .abs => try cg.airAbs(inst),
2670 .floor => try cg.airUnOpSimple(inst, .floor),
2671 .ceil => try cg.airUnOpSimple(inst, .ceil),
2672 .round => try cg.airUnOpSimple(inst, .round),
2673 .trunc_float => try cg.airUnOpSimple(inst, .trunc),
2674 .neg, .neg_optimized => try cg.airUnOpSimple(inst, .f_neg),
2675
2676 .div_float, .div_float_optimized => try cg.airArithOp(inst, .OpFDiv, .OpSDiv, .OpUDiv),
2677 .div_floor, .div_floor_optimized => try cg.airDivFloor(inst),
2678 .div_trunc, .div_trunc_optimized => try cg.airDivTrunc(inst),
2679
2680 .rem, .rem_optimized => try cg.airArithOp(inst, .OpFRem, .OpSRem, .OpUMod),
2681 .mod, .mod_optimized => try cg.airArithOp(inst, .OpFMod, .OpSMod, .OpUMod),
2682
2683 .add_with_overflow => try cg.airAddSubOverflow(inst, .OpIAdd, .OpULessThan, .OpSLessThan),
2684 .sub_with_overflow => try cg.airAddSubOverflow(inst, .OpISub, .OpUGreaterThan, .OpSGreaterThan),
2685 .mul_with_overflow => try cg.airMulOverflow(inst),
2686 .shl_with_overflow => try cg.airShlOverflow(inst),
2687
2688 .mul_add => try cg.airMulAdd(inst),
2689
2690 .ctz => try cg.airClzCtz(inst, .ctz),
2691 .clz => try cg.airClzCtz(inst, .clz),
2692
2693 .select => try cg.airSelect(inst),
2694
2695 .splat => try cg.airSplat(inst),
2696 .reduce, .reduce_optimized => try cg.airReduce(inst),
2697 .shuffle_one => try cg.airShuffleOne(inst),
2698 .shuffle_two => try cg.airShuffleTwo(inst),
2699
2700 .ptr_add => try cg.airPtrAdd(inst),
2701 .ptr_sub => try cg.airPtrSub(inst),
2702
2703 .bit_and => try cg.airBinOpSimple(inst, .OpBitwiseAnd),
2704 .bit_or => try cg.airBinOpSimple(inst, .OpBitwiseOr),
2705 .xor => try cg.airBinOpSimple(inst, .OpBitwiseXor),
2706 .bool_and => try cg.airBinOpSimple(inst, .OpLogicalAnd),
2707 .bool_or => try cg.airBinOpSimple(inst, .OpLogicalOr),
2708
2709 .shl, .shl_exact => try cg.airShift(inst, .OpShiftLeftLogical, .OpShiftLeftLogical),
2710 .shr, .shr_exact => try cg.airShift(inst, .OpShiftRightLogical, .OpShiftRightArithmetic),
2711
2712 .min => try cg.airMinMax(inst, .min),
2713 .max => try cg.airMinMax(inst, .max),
2714
2715 .bitcast => try cg.airBitCast(inst),
2716 .intcast, .trunc => try cg.airIntCast(inst),
2717 .float_from_int => try cg.airFloatFromInt(inst),
2718 .int_from_float => try cg.airIntFromFloat(inst),
2719 .fpext, .fptrunc => try cg.airFloatCast(inst),
2720 .not => try cg.airNot(inst),
2721
2722 .array_to_slice => try cg.airArrayToSlice(inst),
2723 .slice => try cg.airSlice(inst),
2724 .aggregate_init => try cg.airAggregateInit(inst),
2725 .memcpy => return cg.airMemcpy(inst),
2726 .memmove => return cg.airMemmove(inst),
2727
2728 .slice_ptr => try cg.airSliceField(inst, 0),
2729 .slice_len => try cg.airSliceField(inst, 1),
2730 .slice_elem_ptr => try cg.airSliceElemPtr(inst),
2731 .slice_elem_val => try cg.airSliceElemVal(inst),
2732 .ptr_elem_ptr => try cg.airPtrElemPtr(inst),
2733 .ptr_elem_val => try cg.airPtrElemVal(inst),
2734 .array_elem_val => try cg.airArrayElemVal(inst),
2735
2736 .set_union_tag => return cg.airSetUnionTag(inst),
2737 .get_union_tag => try cg.airGetUnionTag(inst),
2738 .union_init => try cg.airUnionInit(inst),
2739
2740 .struct_field_val => try cg.airStructFieldVal(inst),
2741 .field_parent_ptr => try cg.airFieldParentPtr(inst),
2742
2743 .struct_field_ptr_index_0 => try cg.airStructFieldPtrIndex(inst, 0),
2744 .struct_field_ptr_index_1 => try cg.airStructFieldPtrIndex(inst, 1),
2745 .struct_field_ptr_index_2 => try cg.airStructFieldPtrIndex(inst, 2),
2746 .struct_field_ptr_index_3 => try cg.airStructFieldPtrIndex(inst, 3),
2747
2748 .cmp_eq => try cg.airCmp(inst, .eq),
2749 .cmp_neq => try cg.airCmp(inst, .neq),
2750 .cmp_gt => try cg.airCmp(inst, .gt),
2751 .cmp_gte => try cg.airCmp(inst, .gte),
2752 .cmp_lt => try cg.airCmp(inst, .lt),
2753 .cmp_lte => try cg.airCmp(inst, .lte),
2754 .cmp_vector => try cg.airVectorCmp(inst),
2755
2756 .arg => cg.airArg(),
2757 .alloc => try cg.airAlloc(inst),
2758 // TODO: We probably need to have a special implementation of this for the C abi.
2759 .ret_ptr => try cg.airAlloc(inst),
2760 .block => try cg.airBlock(inst),
2761
2762 .load => try cg.airLoad(inst),
2763 .store, .store_safe => return cg.airStore(inst),
2764
2765 .br => return cg.airBr(inst),
2766 // For now just ignore this instruction. This effectively falls back on the old implementation,
2767 // this doesn't change anything for us.
2768 .repeat => return,
2769 .breakpoint => return,
2770 .cond_br => return cg.airCondBr(inst),
2771 .loop => return cg.airLoop(inst),
2772 .ret => return cg.airRet(inst),
2773 .ret_safe => return cg.airRet(inst), // TODO
2774 .ret_load => return cg.airRetLoad(inst),
2775 .@"try" => try cg.airTry(inst),
2776 .switch_br => return cg.airSwitchBr(inst),
2777 .unreach, .trap => return cg.airUnreach(),
2778
2779 .dbg_empty_stmt => return,
2780 .dbg_stmt => return cg.airDbgStmt(inst),
2781 .dbg_inline_block => try cg.airDbgInlineBlock(inst),
2782 .dbg_var_ptr, .dbg_var_val, .dbg_arg_inline => return cg.airDbgVar(inst),
2783
2784 .unwrap_errunion_err => try cg.airErrUnionErr(inst),
2785 .unwrap_errunion_payload => try cg.airErrUnionPayload(inst),
2786 .wrap_errunion_err => try cg.airWrapErrUnionErr(inst),
2787 .wrap_errunion_payload => try cg.airWrapErrUnionPayload(inst),
2788
2789 .is_null => try cg.airIsNull(inst, false, .is_null),
2790 .is_non_null => try cg.airIsNull(inst, false, .is_non_null),
2791 .is_null_ptr => try cg.airIsNull(inst, true, .is_null),
2792 .is_non_null_ptr => try cg.airIsNull(inst, true, .is_non_null),
2793 .is_err => try cg.airIsErr(inst, .is_err),
2794 .is_non_err => try cg.airIsErr(inst, .is_non_err),
2795
2796 .optional_payload => try cg.airUnwrapOptional(inst),
2797 .optional_payload_ptr => try cg.airUnwrapOptionalPtr(inst),
2798 .wrap_optional => try cg.airWrapOptional(inst),
2799
2800 .assembly => try cg.airAssembly(inst),
2801
2802 .call => try cg.airCall(inst, .auto),
2803 .call_always_tail => try cg.airCall(inst, .always_tail),
2804 .call_never_tail => try cg.airCall(inst, .never_tail),
2805 .call_never_inline => try cg.airCall(inst, .never_inline),
2806
2807 .work_item_id => try cg.airWorkItemId(inst),
2808 .work_group_size => try cg.airWorkGroupSize(inst),
2809 .work_group_id => try cg.airWorkGroupId(inst),
2810
2811 // zig fmt: on
2812
2813 else => |tag| return cg.todo("implement AIR tag {s}", .{@tagName(tag)}),
2814 };
2815
2816 const result_id = maybe_result_id orelse return;
2817 try cg.inst_results.putNoClobber(gpa, inst, result_id);
2818}
2819
2820fn airBinOpSimple(cg: *CodeGen, inst: Air.Inst.Index, op: Opcode) !?Id {
2821 const bin_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
2822 const lhs = try cg.temporary(bin_op.lhs);
2823 const rhs = try cg.temporary(bin_op.rhs);
2824
2825 const result = try cg.buildBinary(op, lhs, rhs);
2826 return try result.materialize(cg);
2827}
2828
2829fn airShift(cg: *CodeGen, inst: Air.Inst.Index, unsigned: Opcode, signed: Opcode) !?Id {
2830 const zcu = cg.module.zcu;
2831 const bin_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
2832
2833 if (cg.typeOf(bin_op.lhs).isVector(zcu) and !cg.typeOf(bin_op.rhs).isVector(zcu)) {
2834 return cg.fail("vector shift with scalar rhs", .{});
2835 }
2836
2837 const base = try cg.temporary(bin_op.lhs);
2838 const shift = try cg.temporary(bin_op.rhs);
2839
2840 const result_ty = cg.typeOfIndex(inst);
2841
2842 const info = cg.arithmeticTypeInfo(result_ty);
2843 switch (info.class) {
2844 .composite_integer => return cg.todo("shift ops for composite integers", .{}),
2845 .integer, .strange_integer => {},
2846 .float, .bool => unreachable,
2847 }
2848
2849 // Sometimes Zig doesn't make both of the arguments the same types here. SPIR-V expects that,
2850 // so just manually upcast it if required.
2851
2852 // Note: The sign may differ here between the shift and the base type, in case
2853 // of an arithmetic right shift. SPIR-V still expects the same type,
2854 // so in that case we have to cast convert to signed.
2855 const casted_shift = try cg.buildConvert(base.ty.scalarType(zcu), shift);
2856
2857 const shifted = switch (info.signedness) {
2858 .unsigned => try cg.buildBinary(unsigned, base, casted_shift),
2859 .signed => try cg.buildBinary(signed, base, casted_shift),
2860 };
2861
2862 const result = try cg.normalize(shifted, info);
2863 return try result.materialize(cg);
2864}
2865
2866const MinMax = enum {
2867 min,
2868 max,
2869
2870 pub fn extInstOpcode(
2871 op: MinMax,
2872 target: *const std.Target,
2873 info: ArithmeticTypeInfo,
2874 ) u32 {
2875 return switch (target.os.tag) {
2876 .opencl => @intFromEnum(@as(spec.OpenClOpcode, switch (info.class) {
2877 .float => switch (op) {
2878 .min => .fmin,
2879 .max => .fmax,
2880 },
2881 .integer, .strange_integer, .composite_integer => switch (info.signedness) {
2882 .signed => switch (op) {
2883 .min => .s_min,
2884 .max => .s_max,
2885 },
2886 .unsigned => switch (op) {
2887 .min => .u_min,
2888 .max => .u_max,
2889 },
2890 },
2891 .bool => unreachable,
2892 })),
2893 .vulkan, .opengl => @intFromEnum(@as(spec.GlslOpcode, switch (info.class) {
2894 .float => switch (op) {
2895 .min => .FMin,
2896 .max => .FMax,
2897 },
2898 .integer, .strange_integer, .composite_integer => switch (info.signedness) {
2899 .signed => switch (op) {
2900 .min => .SMin,
2901 .max => .SMax,
2902 },
2903 .unsigned => switch (op) {
2904 .min => .UMin,
2905 .max => .UMax,
2906 },
2907 },
2908 .bool => unreachable,
2909 })),
2910 else => unreachable,
2911 };
2912 }
2913};
2914
2915fn airMinMax(cg: *CodeGen, inst: Air.Inst.Index, op: MinMax) !?Id {
2916 const bin_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
2917
2918 const lhs = try cg.temporary(bin_op.lhs);
2919 const rhs = try cg.temporary(bin_op.rhs);
2920
2921 const result = try cg.minMax(lhs, rhs, op);
2922 return try result.materialize(cg);
2923}
2924
2925fn minMax(cg: *CodeGen, lhs: Temporary, rhs: Temporary, op: MinMax) !Temporary {
2926 const zcu = cg.module.zcu;
2927 const target = zcu.getTarget();
2928 const info = cg.arithmeticTypeInfo(lhs.ty);
2929
2930 const v = cg.vectorization(.{ lhs, rhs });
2931 const ops = v.components();
2932 const results = cg.module.allocIds(ops);
2933
2934 const op_result_ty = lhs.ty.scalarType(zcu);
2935 const op_result_ty_id = try cg.resolveType(op_result_ty, .direct);
2936 const result_ty = try v.resultType(cg, lhs.ty);
2937
2938 const op_lhs = try v.prepare(cg, lhs);
2939 const op_rhs = try v.prepare(cg, rhs);
2940
2941 const set = try cg.importExtendedSet();
2942 const opcode = op.extInstOpcode(target, info);
2943 for (0..ops) |i| {
2944 try cg.body.emit(cg.module.gpa, .OpExtInst, .{
2945 .id_result_type = op_result_ty_id,
2946 .id_result = results.at(i),
2947 .set = set,
2948 .instruction = .{ .inst = opcode },
2949 .id_ref_4 = &.{ op_lhs.at(i), op_rhs.at(i) },
2950 });
2951 }
2952
2953 return v.finalize(result_ty, results);
2954}
2955
2956/// This function normalizes values to a canonical representation
2957/// after some arithmetic operation. This mostly consists of wrapping
2958/// behavior for strange integers:
2959/// - Unsigned integers are bitwise masked with a mask that only passes
2960/// the valid bits through.
2961/// - Signed integers are also sign extended if they are negative.
2962/// All other values are returned unmodified (this makes strange integer
2963/// wrapping easier to use in generic operations).
2964fn normalize(cg: *CodeGen, value: Temporary, info: ArithmeticTypeInfo) !Temporary {
2965 const zcu = cg.module.zcu;
2966 const ty = value.ty;
2967 switch (info.class) {
2968 .composite_integer, .integer, .bool, .float => return value,
2969 .strange_integer => switch (info.signedness) {
2970 .unsigned => {
2971 const mask_value = @as(u64, std.math.maxInt(u64)) >> @as(u6, @intCast(64 - info.bits));
2972 const mask_id = try cg.constInt(ty.scalarType(zcu), mask_value);
2973 return try cg.buildBinary(.OpBitwiseAnd, value, Temporary.init(ty.scalarType(zcu), mask_id));
2974 },
2975 .signed => {
2976 // Shift left and right so that we can copy the sight bit that way.
2977 const shift_amt_id = try cg.constInt(ty.scalarType(zcu), info.backing_bits - info.bits);
2978 const shift_amt: Temporary = .init(ty.scalarType(zcu), shift_amt_id);
2979 const left = try cg.buildBinary(.OpShiftLeftLogical, value, shift_amt);
2980 return try cg.buildBinary(.OpShiftRightArithmetic, left, shift_amt);
2981 },
2982 },
2983 }
2984}
2985
2986fn airDivFloor(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
2987 const bin_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
2988
2989 const lhs = try cg.temporary(bin_op.lhs);
2990 const rhs = try cg.temporary(bin_op.rhs);
2991
2992 const info = cg.arithmeticTypeInfo(lhs.ty);
2993 switch (info.class) {
2994 .composite_integer => unreachable, // TODO
2995 .integer, .strange_integer => {
2996 switch (info.signedness) {
2997 .unsigned => {
2998 const result = try cg.buildBinary(.OpUDiv, lhs, rhs);
2999 return try result.materialize(cg);
3000 },
3001 .signed => {},
3002 }
3003
3004 // For signed integers:
3005 // (a / b) - (a % b != 0 && a < 0 != b < 0);
3006 // There shouldn't be any overflow issues.
3007
3008 const div = try cg.buildBinary(.OpSDiv, lhs, rhs);
3009 const rem = try cg.buildBinary(.OpSRem, lhs, rhs);
3010 const zero: Temporary = .init(lhs.ty, try cg.constInt(lhs.ty, 0));
3011 const rem_non_zero = try cg.buildCmp(.OpINotEqual, rem, zero);
3012 const lhs_rhs_xor = try cg.buildBinary(.OpBitwiseXor, lhs, rhs);
3013 const signs_differ = try cg.buildCmp(.OpSLessThan, lhs_rhs_xor, zero);
3014 const adjust = try cg.buildBinary(.OpLogicalAnd, rem_non_zero, signs_differ);
3015 const result = try cg.buildBinary(.OpISub, div, try cg.intFromBool(adjust, div.ty));
3016 return try result.materialize(cg);
3017 },
3018 .float => {
3019 const div = try cg.buildBinary(.OpFDiv, lhs, rhs);
3020 const result = try cg.buildUnary(.floor, div);
3021 return try result.materialize(cg);
3022 },
3023 .bool => unreachable,
3024 }
3025}
3026
3027fn airDivTrunc(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
3028 const bin_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
3029 const lhs = try cg.temporary(bin_op.lhs);
3030 const rhs = try cg.temporary(bin_op.rhs);
3031 const info = cg.arithmeticTypeInfo(lhs.ty);
3032 switch (info.class) {
3033 .composite_integer => unreachable, // TODO
3034 .integer, .strange_integer => switch (info.signedness) {
3035 .unsigned => {
3036 const result = try cg.buildBinary(.OpUDiv, lhs, rhs);
3037 return try result.materialize(cg);
3038 },
3039 .signed => {
3040 const result = try cg.buildBinary(.OpSDiv, lhs, rhs);
3041 return try result.materialize(cg);
3042 },
3043 },
3044 .float => {
3045 const div = try cg.buildBinary(.OpFDiv, lhs, rhs);
3046 const result = try cg.buildUnary(.trunc, div);
3047 return try result.materialize(cg);
3048 },
3049 .bool => unreachable,
3050 }
3051}
3052
3053fn airUnOpSimple(cg: *CodeGen, inst: Air.Inst.Index, op: UnaryOp) !?Id {
3054 const un_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].un_op;
3055 const operand = try cg.temporary(un_op);
3056 const result = try cg.buildUnary(op, operand);
3057 return try result.materialize(cg);
3058}
3059
3060fn airArithOp(
3061 cg: *CodeGen,
3062 inst: Air.Inst.Index,
3063 comptime fop: Opcode,
3064 comptime sop: Opcode,
3065 comptime uop: Opcode,
3066) !?Id {
3067 const bin_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
3068 const lhs = try cg.temporary(bin_op.lhs);
3069 const rhs = try cg.temporary(bin_op.rhs);
3070 const info = cg.arithmeticTypeInfo(lhs.ty);
3071 const result = switch (info.class) {
3072 .composite_integer => unreachable, // TODO
3073 .integer, .strange_integer => switch (info.signedness) {
3074 .signed => try cg.buildBinary(sop, lhs, rhs),
3075 .unsigned => try cg.buildBinary(uop, lhs, rhs),
3076 },
3077 .float => try cg.buildBinary(fop, lhs, rhs),
3078 .bool => unreachable,
3079 };
3080 return try result.materialize(cg);
3081}
3082
3083fn airAbs(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
3084 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
3085 const operand = try cg.temporary(ty_op.operand);
3086 // Note: operand_ty may be signed, while ty is always unsigned!
3087 const result_ty = cg.typeOfIndex(inst);
3088 const result = try cg.abs(result_ty, operand);
3089 return try result.materialize(cg);
3090}
3091
3092fn abs(cg: *CodeGen, result_ty: Type, value: Temporary) !Temporary {
3093 const zcu = cg.module.zcu;
3094 const target = cg.module.zcu.getTarget();
3095 const operand_info = cg.arithmeticTypeInfo(value.ty);
3096 switch (operand_info.class) {
3097 .float => return try cg.buildUnary(.f_abs, value),
3098 .integer, .strange_integer => {
3099 const abs_value = try cg.buildUnary(.i_abs, value);
3100 switch (target.os.tag) {
3101 .vulkan, .opengl => {
3102 if (value.ty.intInfo(zcu).signedness == .signed) {
3103 return cg.todo("perform bitcast after @abs", .{});
3104 }
3105 },
3106 else => {},
3107 }
3108 return try cg.normalize(abs_value, cg.arithmeticTypeInfo(result_ty));
3109 },
3110 .composite_integer => unreachable, // TODO
3111 .bool => unreachable,
3112 }
3113}
3114
3115fn airAddSubOverflow(
3116 cg: *CodeGen,
3117 inst: Air.Inst.Index,
3118 comptime add: Opcode,
3119 u_opcode: Opcode,
3120 s_opcode: Opcode,
3121) !?Id {
3122 // Note: OpIAddCarry and OpISubBorrow are not really useful here: For unsigned numbers,
3123 // there is in both cases only one extra operation required. For signed operations,
3124 // the overflow bit is set then going from 0x80.. to 0x00.., but this doesn't actually
3125 // normally set a carry bit. So the SPIR-V overflow operations are not particularly
3126 // useful here.
3127
3128 _ = s_opcode;
3129
3130 const ty_pl = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
3131 const extra = cg.air.extraData(Air.Bin, ty_pl.payload).data;
3132 const lhs = try cg.temporary(extra.lhs);
3133 const rhs = try cg.temporary(extra.rhs);
3134 const result_ty = cg.typeOfIndex(inst);
3135
3136 const info = cg.arithmeticTypeInfo(lhs.ty);
3137 switch (info.class) {
3138 .composite_integer => unreachable, // TODO
3139 .strange_integer, .integer => {},
3140 .float, .bool => unreachable,
3141 }
3142
3143 const sum = try cg.buildBinary(add, lhs, rhs);
3144 const result = try cg.normalize(sum, info);
3145 const overflowed = switch (info.signedness) {
3146 // Overflow happened if the result is smaller than either of the operands. It doesn't matter which.
3147 // For subtraction the conditions need to be swapped.
3148 .unsigned => try cg.buildCmp(u_opcode, result, lhs),
3149 // For signed operations, we check the signs of the operands and the result.
3150 .signed => blk: {
3151 // Signed overflow detection using the sign bits of the operands and the result.
3152 // For addition (a + b), overflow occurs if the operands have the same sign
3153 // and the result's sign is different from the operands' sign.
3154 // (sign(a) == sign(b)) && (sign(a) != sign(result))
3155 // For subtraction (a - b), overflow occurs if the operands have different signs
3156 // and the result's sign is different from the minuend's (a's) sign.
3157 // (sign(a) != sign(b)) && (sign(a) != sign(result))
3158 const zero: Temporary = .init(rhs.ty, try cg.constInt(rhs.ty, 0));
3159 const lhs_is_neg = try cg.buildCmp(.OpSLessThan, lhs, zero);
3160 const rhs_is_neg = try cg.buildCmp(.OpSLessThan, rhs, zero);
3161 const result_is_neg = try cg.buildCmp(.OpSLessThan, result, zero);
3162 const signs_match = try cg.buildCmp(.OpLogicalEqual, lhs_is_neg, rhs_is_neg);
3163 const result_sign_differs = try cg.buildCmp(.OpLogicalNotEqual, lhs_is_neg, result_is_neg);
3164 const overflow_condition = switch (add) {
3165 .OpIAdd => signs_match,
3166 .OpISub => try cg.buildUnary(.l_not, signs_match),
3167 else => unreachable,
3168 };
3169 break :blk try cg.buildCmp(.OpLogicalAnd, overflow_condition, result_sign_differs);
3170 },
3171 };
3172
3173 const ov = try cg.intFromBool(overflowed, .u1);
3174 const result_ty_id = try cg.resolveType(result_ty, .direct);
3175 return try cg.constructComposite(result_ty_id, &.{ try result.materialize(cg), try ov.materialize(cg) });
3176}
3177
3178fn airMulOverflow(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
3179 const pt = cg.pt;
3180 const ty_pl = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
3181 const extra = cg.air.extraData(Air.Bin, ty_pl.payload).data;
3182 const lhs = try cg.temporary(extra.lhs);
3183 const rhs = try cg.temporary(extra.rhs);
3184 const result_ty = cg.typeOfIndex(inst);
3185
3186 const info = cg.arithmeticTypeInfo(lhs.ty);
3187 switch (info.class) {
3188 .composite_integer => unreachable, // TODO
3189 .strange_integer, .integer => {},
3190 .float, .bool => unreachable,
3191 }
3192
3193 // There are 3 cases which we have to deal with:
3194 // - If info.bits < 32 / 2, we will upcast to 32 and check the higher bits
3195 // - If info.bits > 32 / 2, we have to use extended multiplication
3196 // - Additionally, if info.bits != 32, we'll have to check the high bits
3197 // of the result too.
3198
3199 const largest_int_bits = cg.largestSupportedIntBits();
3200 // If non-null, the number of bits that the multiplication should be performed in. If
3201 // null, we have to use wide multiplication.
3202 const maybe_op_ty_bits: ?u16 = switch (info.bits) {
3203 0 => unreachable,
3204 1...16 => 32,
3205 17...32 => if (largest_int_bits > 32) 64 else null, // Upcast if we can.
3206 33...64 => null, // Always use wide multiplication.
3207 else => unreachable, // TODO: Composite integers
3208 };
3209
3210 const result, const overflowed = switch (info.signedness) {
3211 .unsigned => blk: {
3212 if (maybe_op_ty_bits) |op_ty_bits| {
3213 const op_ty = try pt.intType(.unsigned, op_ty_bits);
3214 const casted_lhs = try cg.buildConvert(op_ty, lhs);
3215 const casted_rhs = try cg.buildConvert(op_ty, rhs);
3216 const full_result = try cg.buildBinary(.OpIMul, casted_lhs, casted_rhs);
3217 const low_bits = try cg.buildConvert(lhs.ty, full_result);
3218 const result = try cg.normalize(low_bits, info);
3219 // Shift the result bits away to get the overflow bits.
3220 const shift: Temporary = .init(full_result.ty, try cg.constInt(full_result.ty, info.bits));
3221 const overflow = try cg.buildBinary(.OpShiftRightLogical, full_result, shift);
3222 // Directly check if its zero in the op_ty without converting first.
3223 const zero: Temporary = .init(full_result.ty, try cg.constInt(full_result.ty, 0));
3224 const overflowed = try cg.buildCmp(.OpINotEqual, zero, overflow);
3225 break :blk .{ result, overflowed };
3226 }
3227
3228 const low_bits, const high_bits = try cg.buildWideMul(.unsigned, lhs, rhs);
3229
3230 // Truncate the result, if required.
3231 const result = try cg.normalize(low_bits, info);
3232
3233 // Overflow happened if the high-bits of the result are non-zero OR if the
3234 // high bits of the low word of the result (those outside the range of the
3235 // int) are nonzero.
3236 const zero: Temporary = .init(lhs.ty, try cg.constInt(lhs.ty, 0));
3237 const high_overflowed = try cg.buildCmp(.OpINotEqual, zero, high_bits);
3238
3239 // If no overflow bits in low_bits, no extra work needs to be done.
3240 if (info.backing_bits == info.bits) break :blk .{ result, high_overflowed };
3241
3242 // Shift the result bits away to get the overflow bits.
3243 const shift: Temporary = .init(lhs.ty, try cg.constInt(lhs.ty, info.bits));
3244 const low_overflow = try cg.buildBinary(.OpShiftRightLogical, low_bits, shift);
3245 const low_overflowed = try cg.buildCmp(.OpINotEqual, zero, low_overflow);
3246
3247 const overflowed = try cg.buildCmp(.OpLogicalOr, low_overflowed, high_overflowed);
3248
3249 break :blk .{ result, overflowed };
3250 },
3251 .signed => blk: {
3252 // - lhs >= 0, rhxs >= 0: expect positive; overflow should be 0
3253 // - lhs == 0 : expect positive; overflow should be 0
3254 // - rhs == 0: expect positive; overflow should be 0
3255 // - lhs > 0, rhs < 0: expect negative; overflow should be -1
3256 // - lhs < 0, rhs > 0: expect negative; overflow should be -1
3257 // - lhs <= 0, rhs <= 0: expect positive; overflow should be 0
3258 // ------
3259 // overflow should be -1 when
3260 // (lhs > 0 && rhs < 0) || (lhs < 0 && rhs > 0)
3261
3262 const zero: Temporary = .init(lhs.ty, try cg.constInt(lhs.ty, 0));
3263 const lhs_negative = try cg.buildCmp(.OpSLessThan, lhs, zero);
3264 const rhs_negative = try cg.buildCmp(.OpSLessThan, rhs, zero);
3265 const lhs_positive = try cg.buildCmp(.OpSGreaterThan, lhs, zero);
3266 const rhs_positive = try cg.buildCmp(.OpSGreaterThan, rhs, zero);
3267
3268 // Set to `true` if we expect -1.
3269 const expected_overflow_bit = try cg.buildBinary(
3270 .OpLogicalOr,
3271 try cg.buildCmp(.OpLogicalAnd, lhs_positive, rhs_negative),
3272 try cg.buildCmp(.OpLogicalAnd, lhs_negative, rhs_positive),
3273 );
3274
3275 if (maybe_op_ty_bits) |op_ty_bits| {
3276 const op_ty = try pt.intType(.signed, op_ty_bits);
3277 // Assume normalized; sign bit is set. We want a sign extend.
3278 const casted_lhs = try cg.buildConvert(op_ty, lhs);
3279 const casted_rhs = try cg.buildConvert(op_ty, rhs);
3280
3281 const full_result = try cg.buildBinary(.OpIMul, casted_lhs, casted_rhs);
3282
3283 // Truncate to the result type.
3284 const low_bits = try cg.buildConvert(lhs.ty, full_result);
3285 const result = try cg.normalize(low_bits, info);
3286
3287 // Now, we need to check the overflow bits AND the sign
3288 // bit for the expected overflow bits.
3289 // To do that, shift out everything bit the sign bit and
3290 // then check what remains.
3291 const shift: Temporary = .init(full_result.ty, try cg.constInt(full_result.ty, info.bits - 1));
3292 // Use SRA so that any sign bits are duplicated. Now we can just check if ALL bits are set
3293 // for negative cases.
3294 const overflow = try cg.buildBinary(.OpShiftRightArithmetic, full_result, shift);
3295
3296 const long_all_set: Temporary = .init(full_result.ty, try cg.constInt(full_result.ty, -1));
3297 const long_zero: Temporary = .init(full_result.ty, try cg.constInt(full_result.ty, 0));
3298 const mask = try cg.buildSelect(expected_overflow_bit, long_all_set, long_zero);
3299
3300 const overflowed = try cg.buildCmp(.OpINotEqual, mask, overflow);
3301
3302 break :blk .{ result, overflowed };
3303 }
3304
3305 const low_bits, const high_bits = try cg.buildWideMul(.signed, lhs, rhs);
3306
3307 // Truncate result if required.
3308 const result = try cg.normalize(low_bits, info);
3309
3310 const all_set: Temporary = .init(lhs.ty, try cg.constInt(lhs.ty, -1));
3311 const mask = try cg.buildSelect(expected_overflow_bit, all_set, zero);
3312
3313 // Like with unsigned, overflow happened if high_bits are not the ones we expect,
3314 // and we also need to check some ones from the low bits.
3315
3316 const high_overflowed = try cg.buildCmp(.OpINotEqual, mask, high_bits);
3317
3318 // If no overflow bits in low_bits, no extra work needs to be done.
3319 // Careful, we still have to check the sign bit, so this branch
3320 // only goes for i33 and such.
3321 if (info.backing_bits == info.bits + 1) break :blk .{ result, high_overflowed };
3322
3323 // Shift the result bits away to get the overflow bits.
3324 const shift: Temporary = .init(lhs.ty, try cg.constInt(lhs.ty, info.bits - 1));
3325 // Use SRA so that any sign bits are duplicated. Now we can just check if ALL bits are set
3326 // for negative cases.
3327 const low_overflow = try cg.buildBinary(.OpShiftRightArithmetic, low_bits, shift);
3328 const low_overflowed = try cg.buildCmp(.OpINotEqual, mask, low_overflow);
3329
3330 const overflowed = try cg.buildCmp(.OpLogicalOr, low_overflowed, high_overflowed);
3331
3332 break :blk .{ result, overflowed };
3333 },
3334 };
3335
3336 const ov = try cg.intFromBool(overflowed, .u1);
3337
3338 const result_ty_id = try cg.resolveType(result_ty, .direct);
3339 return try cg.constructComposite(result_ty_id, &.{ try result.materialize(cg), try ov.materialize(cg) });
3340}
3341
3342fn airShlOverflow(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
3343 const zcu = cg.module.zcu;
3344
3345 const ty_pl = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
3346 const extra = cg.air.extraData(Air.Bin, ty_pl.payload).data;
3347
3348 if (cg.typeOf(extra.lhs).isVector(zcu) and !cg.typeOf(extra.rhs).isVector(zcu)) {
3349 return cg.fail("vector shift with scalar rhs", .{});
3350 }
3351
3352 const base = try cg.temporary(extra.lhs);
3353 const shift = try cg.temporary(extra.rhs);
3354
3355 const result_ty = cg.typeOfIndex(inst);
3356
3357 const info = cg.arithmeticTypeInfo(base.ty);
3358 switch (info.class) {
3359 .composite_integer => unreachable, // TODO
3360 .integer, .strange_integer => {},
3361 .float, .bool => unreachable,
3362 }
3363
3364 // Sometimes Zig doesn't make both of the arguments the same types here. SPIR-V expects that,
3365 // so just manually upcast it if required.
3366 const casted_shift = try cg.buildConvert(base.ty.scalarType(zcu), shift);
3367
3368 const left = try cg.buildBinary(.OpShiftLeftLogical, base, casted_shift);
3369 const result = try cg.normalize(left, info);
3370
3371 const right = switch (info.signedness) {
3372 .unsigned => try cg.buildBinary(.OpShiftRightLogical, result, casted_shift),
3373 .signed => try cg.buildBinary(.OpShiftRightArithmetic, result, casted_shift),
3374 };
3375
3376 const overflowed = try cg.buildCmp(.OpINotEqual, base, right);
3377 const ov = try cg.intFromBool(overflowed, .u1);
3378
3379 const result_ty_id = try cg.resolveType(result_ty, .direct);
3380 return try cg.constructComposite(result_ty_id, &.{ try result.materialize(cg), try ov.materialize(cg) });
3381}
3382
3383fn airMulAdd(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
3384 const pl_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].pl_op;
3385 const extra = cg.air.extraData(Air.Bin, pl_op.payload).data;
3386
3387 const a = try cg.temporary(extra.lhs);
3388 const b = try cg.temporary(extra.rhs);
3389 const c = try cg.temporary(pl_op.operand);
3390
3391 const result_ty = cg.typeOfIndex(inst);
3392 const info = cg.arithmeticTypeInfo(result_ty);
3393 assert(info.class == .float); // .mul_add is only emitted for floats
3394
3395 const result = try cg.buildFma(a, b, c);
3396 return try result.materialize(cg);
3397}
3398
3399fn airClzCtz(cg: *CodeGen, inst: Air.Inst.Index, op: UnaryOp) !?Id {
3400 if (cg.liveness.isUnused(inst)) return null;
3401
3402 const zcu = cg.module.zcu;
3403 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
3404 const operand = try cg.temporary(ty_op.operand);
3405
3406 const scalar_result_ty = cg.typeOfIndex(inst).scalarType(zcu);
3407
3408 const info = cg.arithmeticTypeInfo(operand.ty);
3409 switch (info.class) {
3410 .composite_integer => unreachable, // TODO
3411 .integer, .strange_integer => {},
3412 .float, .bool => unreachable,
3413 }
3414
3415 const count = try cg.buildUnary(op, operand);
3416
3417 // Result of OpenCL ctz/clz returns operand.ty, and we want result_ty.
3418 // result_ty is always large enough to hold the result, so we might have to down
3419 // cast it.
3420 const result = try cg.buildConvert(scalar_result_ty, count);
3421 return try result.materialize(cg);
3422}
3423
3424fn airSelect(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
3425 const pl_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].pl_op;
3426 const extra = cg.air.extraData(Air.Bin, pl_op.payload).data;
3427 const pred = try cg.temporary(pl_op.operand);
3428 const a = try cg.temporary(extra.lhs);
3429 const b = try cg.temporary(extra.rhs);
3430
3431 const result = try cg.buildSelect(pred, a, b);
3432 return try result.materialize(cg);
3433}
3434
3435fn airSplat(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
3436 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
3437
3438 const operand_id = try cg.resolve(ty_op.operand);
3439 const result_ty = cg.typeOfIndex(inst);
3440
3441 return try cg.constructCompositeSplat(result_ty, operand_id);
3442}
3443
3444fn airReduce(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
3445 const zcu = cg.module.zcu;
3446 const reduce = cg.air.instructions.items(.data)[@intFromEnum(inst)].reduce;
3447 const operand = try cg.resolve(reduce.operand);
3448 const operand_ty = cg.typeOf(reduce.operand);
3449 const scalar_ty = operand_ty.scalarType(zcu);
3450 const scalar_ty_id = try cg.resolveType(scalar_ty, .direct);
3451 const info = cg.arithmeticTypeInfo(operand_ty);
3452 const len = operand_ty.vectorLen(zcu);
3453 const first = try cg.extractVectorComponent(scalar_ty, operand, 0);
3454
3455 switch (reduce.operation) {
3456 .Min, .Max => |op| {
3457 var result: Temporary = .init(scalar_ty, first);
3458 const cmp_op: MinMax = switch (op) {
3459 .Max => .max,
3460 .Min => .min,
3461 else => unreachable,
3462 };
3463 for (1..len) |i| {
3464 const lhs = result;
3465 const rhs_id = try cg.extractVectorComponent(scalar_ty, operand, @intCast(i));
3466 const rhs: Temporary = .init(scalar_ty, rhs_id);
3467
3468 result = try cg.minMax(lhs, rhs, cmp_op);
3469 }
3470
3471 return try result.materialize(cg);
3472 },
3473 else => {},
3474 }
3475
3476 var result_id = first;
3477
3478 const opcode: Opcode = switch (info.class) {
3479 .bool => switch (reduce.operation) {
3480 .And => .OpLogicalAnd,
3481 .Or => .OpLogicalOr,
3482 .Xor => .OpLogicalNotEqual,
3483 else => unreachable,
3484 },
3485 .strange_integer, .integer => switch (reduce.operation) {
3486 .And => .OpBitwiseAnd,
3487 .Or => .OpBitwiseOr,
3488 .Xor => .OpBitwiseXor,
3489 .Add => .OpIAdd,
3490 .Mul => .OpIMul,
3491 else => unreachable,
3492 },
3493 .float => switch (reduce.operation) {
3494 .Add => .OpFAdd,
3495 .Mul => .OpFMul,
3496 else => unreachable,
3497 },
3498 .composite_integer => unreachable, // TODO
3499 };
3500
3501 for (1..len) |i| {
3502 const lhs = result_id;
3503 const rhs = try cg.extractVectorComponent(scalar_ty, operand, @intCast(i));
3504 result_id = cg.module.allocId();
3505
3506 try cg.body.emitRaw(cg.module.gpa, opcode, 4);
3507 cg.body.writeOperand(Id, scalar_ty_id);
3508 cg.body.writeOperand(Id, result_id);
3509 cg.body.writeOperand(Id, lhs);
3510 cg.body.writeOperand(Id, rhs);
3511 }
3512
3513 return result_id;
3514}
3515
3516fn airShuffleOne(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
3517 const zcu = cg.module.zcu;
3518 const gpa = zcu.gpa;
3519
3520 const unwrapped = cg.air.unwrapShuffleOne(zcu, inst);
3521 const mask = unwrapped.mask;
3522 const result_ty = unwrapped.result_ty;
3523 const elem_ty = result_ty.childType(zcu);
3524 const operand = try cg.resolve(unwrapped.operand);
3525
3526 const scratch_top = cg.id_scratch.items.len;
3527 defer cg.id_scratch.shrinkRetainingCapacity(scratch_top);
3528 const constituents = try cg.id_scratch.addManyAsSlice(gpa, mask.len);
3529
3530 for (constituents, mask) |*id, mask_elem| {
3531 id.* = switch (mask_elem.unwrap()) {
3532 .elem => |idx| try cg.extractVectorComponent(elem_ty, operand, idx),
3533 .value => |val| try cg.constant(elem_ty, .fromInterned(val), .direct),
3534 };
3535 }
3536
3537 const result_ty_id = try cg.resolveType(result_ty, .direct);
3538 return try cg.constructComposite(result_ty_id, constituents);
3539}
3540
3541fn airShuffleTwo(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
3542 const zcu = cg.module.zcu;
3543 const gpa = zcu.gpa;
3544
3545 const unwrapped = cg.air.unwrapShuffleTwo(zcu, inst);
3546 const mask = unwrapped.mask;
3547 const result_ty = unwrapped.result_ty;
3548 const elem_ty = result_ty.childType(zcu);
3549 const elem_ty_id = try cg.resolveType(elem_ty, .direct);
3550 const operand_a = try cg.resolve(unwrapped.operand_a);
3551 const operand_b = try cg.resolve(unwrapped.operand_b);
3552
3553 const scratch_top = cg.id_scratch.items.len;
3554 defer cg.id_scratch.shrinkRetainingCapacity(scratch_top);
3555 const constituents = try cg.id_scratch.addManyAsSlice(gpa, mask.len);
3556
3557 for (constituents, mask) |*id, mask_elem| {
3558 id.* = switch (mask_elem.unwrap()) {
3559 .a_elem => |idx| try cg.extractVectorComponent(elem_ty, operand_a, idx),
3560 .b_elem => |idx| try cg.extractVectorComponent(elem_ty, operand_b, idx),
3561 .undef => try cg.module.constUndef(elem_ty_id),
3562 };
3563 }
3564
3565 const result_ty_id = try cg.resolveType(result_ty, .direct);
3566 return try cg.constructComposite(result_ty_id, constituents);
3567}
3568
3569fn accessChainId(
3570 cg: *CodeGen,
3571 result_ty_id: Id,
3572 base: Id,
3573 indices: []const Id,
3574) !Id {
3575 const result_id = cg.module.allocId();
3576 try cg.body.emit(cg.module.gpa, .OpInBoundsAccessChain, .{
3577 .id_result_type = result_ty_id,
3578 .id_result = result_id,
3579 .base = base,
3580 .indexes = indices,
3581 });
3582 return result_id;
3583}
3584
3585/// AccessChain is essentially PtrAccessChain with 0 as initial argument. The effective
3586/// difference lies in whether the resulting type of the first dereference will be the
3587/// same as that of the base pointer, or that of a dereferenced base pointer. AccessChain
3588/// is the latter and PtrAccessChain is the former.
3589fn accessChain(
3590 cg: *CodeGen,
3591 result_ty_id: Id,
3592 base: Id,
3593 indices: []const u32,
3594) !Id {
3595 const gpa = cg.module.gpa;
3596 const scratch_top = cg.id_scratch.items.len;
3597 defer cg.id_scratch.shrinkRetainingCapacity(scratch_top);
3598 const ids = try cg.id_scratch.addManyAsSlice(gpa, indices.len);
3599 for (indices, ids) |index, *id| {
3600 id.* = try cg.constInt(.u32, index);
3601 }
3602 return try cg.accessChainId(result_ty_id, base, ids);
3603}
3604
3605fn ptrAccessChain(
3606 cg: *CodeGen,
3607 result_ty_id: Id,
3608 base: Id,
3609 element: Id,
3610 indices: []const u32,
3611) !Id {
3612 const gpa = cg.module.gpa;
3613 const target = cg.module.zcu.getTarget();
3614
3615 const scratch_top = cg.id_scratch.items.len;
3616 defer cg.id_scratch.shrinkRetainingCapacity(scratch_top);
3617 const ids = try cg.id_scratch.addManyAsSlice(gpa, indices.len);
3618 for (indices, ids) |index, *id| {
3619 id.* = try cg.constInt(.u32, index);
3620 }
3621
3622 const result_id = cg.module.allocId();
3623 switch (target.os.tag) {
3624 .opencl, .amdhsa => {
3625 try cg.body.emit(gpa, .OpInBoundsPtrAccessChain, .{
3626 .id_result_type = result_ty_id,
3627 .id_result = result_id,
3628 .base = base,
3629 .element = element,
3630 .indexes = ids,
3631 });
3632 },
3633 .vulkan, .opengl => {
3634 try cg.body.emit(gpa, .OpPtrAccessChain, .{
3635 .id_result_type = result_ty_id,
3636 .id_result = result_id,
3637 .base = base,
3638 .element = element,
3639 .indexes = ids,
3640 });
3641 },
3642 else => unreachable,
3643 }
3644 return result_id;
3645}
3646
3647fn ptrAdd(cg: *CodeGen, result_ty: Type, ptr_ty: Type, ptr_id: Id, offset_id: Id) !Id {
3648 const zcu = cg.module.zcu;
3649 const result_ty_id = try cg.resolveType(result_ty, .direct);
3650
3651 switch (ptr_ty.ptrSize(zcu)) {
3652 .one => {
3653 // Pointer to array
3654 // TODO: Is this correct?
3655 return try cg.accessChainId(result_ty_id, ptr_id, &.{offset_id});
3656 },
3657 .c, .many => {
3658 return try cg.ptrAccessChain(result_ty_id, ptr_id, offset_id, &.{});
3659 },
3660 .slice => {
3661 // TODO: This is probably incorrect. A slice should be returned here, though this is what llvm does.
3662 const slice_ptr_id = try cg.extractField(result_ty, ptr_id, 0);
3663 return try cg.ptrAccessChain(result_ty_id, slice_ptr_id, offset_id, &.{});
3664 },
3665 }
3666}
3667
3668fn airPtrAdd(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
3669 const ty_pl = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
3670 const bin_op = cg.air.extraData(Air.Bin, ty_pl.payload).data;
3671 const ptr_id = try cg.resolve(bin_op.lhs);
3672 const offset_id = try cg.resolve(bin_op.rhs);
3673 const ptr_ty = cg.typeOf(bin_op.lhs);
3674 const result_ty = cg.typeOfIndex(inst);
3675
3676 return try cg.ptrAdd(result_ty, ptr_ty, ptr_id, offset_id);
3677}
3678
3679fn airPtrSub(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
3680 const ty_pl = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
3681 const bin_op = cg.air.extraData(Air.Bin, ty_pl.payload).data;
3682 const ptr_id = try cg.resolve(bin_op.lhs);
3683 const ptr_ty = cg.typeOf(bin_op.lhs);
3684 const offset_id = try cg.resolve(bin_op.rhs);
3685 const offset_ty = cg.typeOf(bin_op.rhs);
3686 const offset_ty_id = try cg.resolveType(offset_ty, .direct);
3687 const result_ty = cg.typeOfIndex(inst);
3688
3689 const negative_offset_id = cg.module.allocId();
3690 try cg.body.emit(cg.module.gpa, .OpSNegate, .{
3691 .id_result_type = offset_ty_id,
3692 .id_result = negative_offset_id,
3693 .operand = offset_id,
3694 });
3695 return try cg.ptrAdd(result_ty, ptr_ty, ptr_id, negative_offset_id);
3696}
3697
3698fn cmp(
3699 cg: *CodeGen,
3700 op: std.math.CompareOperator,
3701 lhs: Temporary,
3702 rhs: Temporary,
3703) !Temporary {
3704 const gpa = cg.module.gpa;
3705 const pt = cg.pt;
3706 const zcu = cg.module.zcu;
3707 const ip = &zcu.intern_pool;
3708 const scalar_ty = lhs.ty.scalarType(zcu);
3709 const is_vector = lhs.ty.isVector(zcu);
3710
3711 switch (scalar_ty.zigTypeTag(zcu)) {
3712 .int, .bool, .float => {},
3713 .@"enum" => {
3714 assert(!is_vector);
3715 const ty = lhs.ty.intTagType(zcu);
3716 return try cg.cmp(op, lhs.pun(ty), rhs.pun(ty));
3717 },
3718 .@"struct" => {
3719 const struct_ty = zcu.typeToPackedStruct(scalar_ty).?;
3720 const ty: Type = .fromInterned(struct_ty.backingIntTypeUnordered(ip));
3721 return try cg.cmp(op, lhs.pun(ty), rhs.pun(ty));
3722 },
3723 .error_set => {
3724 assert(!is_vector);
3725 const err_int_ty = try pt.errorIntType();
3726 return try cg.cmp(op, lhs.pun(err_int_ty), rhs.pun(err_int_ty));
3727 },
3728 .pointer => {
3729 assert(!is_vector);
3730 // Note that while SPIR-V offers OpPtrEqual and OpPtrNotEqual, they are
3731 // currently not implemented in the SPIR-V LLVM translator. Thus, we emit these using
3732 // OpConvertPtrToU...
3733
3734 const usize_ty_id = try cg.resolveType(.usize, .direct);
3735
3736 const lhs_int_id = cg.module.allocId();
3737 try cg.body.emit(gpa, .OpConvertPtrToU, .{
3738 .id_result_type = usize_ty_id,
3739 .id_result = lhs_int_id,
3740 .pointer = try lhs.materialize(cg),
3741 });
3742
3743 const rhs_int_id = cg.module.allocId();
3744 try cg.body.emit(gpa, .OpConvertPtrToU, .{
3745 .id_result_type = usize_ty_id,
3746 .id_result = rhs_int_id,
3747 .pointer = try rhs.materialize(cg),
3748 });
3749
3750 const lhs_int: Temporary = .init(.usize, lhs_int_id);
3751 const rhs_int: Temporary = .init(.usize, rhs_int_id);
3752 return try cg.cmp(op, lhs_int, rhs_int);
3753 },
3754 .optional => {
3755 assert(!is_vector);
3756
3757 const ty = lhs.ty;
3758
3759 const payload_ty = ty.optionalChild(zcu);
3760 if (ty.optionalReprIsPayload(zcu)) {
3761 assert(payload_ty.hasRuntimeBitsIgnoreComptime(zcu));
3762 assert(!payload_ty.isSlice(zcu));
3763
3764 return try cg.cmp(op, lhs.pun(payload_ty), rhs.pun(payload_ty));
3765 }
3766
3767 const lhs_id = try lhs.materialize(cg);
3768 const rhs_id = try rhs.materialize(cg);
3769
3770 const lhs_valid_id = if (payload_ty.hasRuntimeBitsIgnoreComptime(zcu))
3771 try cg.extractField(.bool, lhs_id, 1)
3772 else
3773 try cg.convertToDirect(.bool, lhs_id);
3774
3775 const rhs_valid_id = if (payload_ty.hasRuntimeBitsIgnoreComptime(zcu))
3776 try cg.extractField(.bool, rhs_id, 1)
3777 else
3778 try cg.convertToDirect(.bool, rhs_id);
3779
3780 const lhs_valid: Temporary = .init(.bool, lhs_valid_id);
3781 const rhs_valid: Temporary = .init(.bool, rhs_valid_id);
3782
3783 if (!payload_ty.hasRuntimeBitsIgnoreComptime(zcu)) {
3784 return try cg.cmp(op, lhs_valid, rhs_valid);
3785 }
3786
3787 // a = lhs_valid
3788 // b = rhs_valid
3789 // c = lhs_pl == rhs_pl
3790 //
3791 // For op == .eq we have:
3792 // a == b && a -> c
3793 // = a == b && (!a || c)
3794 //
3795 // For op == .neq we have
3796 // a == b && a -> c
3797 // = !(a == b && a -> c)
3798 // = a != b || !(a -> c
3799 // = a != b || !(!a || c)
3800 // = a != b || a && !c
3801
3802 const lhs_pl_id = try cg.extractField(payload_ty, lhs_id, 0);
3803 const rhs_pl_id = try cg.extractField(payload_ty, rhs_id, 0);
3804
3805 const lhs_pl: Temporary = .init(payload_ty, lhs_pl_id);
3806 const rhs_pl: Temporary = .init(payload_ty, rhs_pl_id);
3807
3808 return switch (op) {
3809 .eq => try cg.buildBinary(
3810 .OpLogicalAnd,
3811 try cg.cmp(.eq, lhs_valid, rhs_valid),
3812 try cg.buildBinary(
3813 .OpLogicalOr,
3814 try cg.buildUnary(.l_not, lhs_valid),
3815 try cg.cmp(.eq, lhs_pl, rhs_pl),
3816 ),
3817 ),
3818 .neq => try cg.buildBinary(
3819 .OpLogicalOr,
3820 try cg.cmp(.neq, lhs_valid, rhs_valid),
3821 try cg.buildBinary(
3822 .OpLogicalAnd,
3823 lhs_valid,
3824 try cg.cmp(.neq, lhs_pl, rhs_pl),
3825 ),
3826 ),
3827 else => unreachable,
3828 };
3829 },
3830 else => |ty| return cg.todo("implement cmp operation for '{s}' type", .{@tagName(ty)}),
3831 }
3832
3833 const info = cg.arithmeticTypeInfo(scalar_ty);
3834 const pred: Opcode = switch (info.class) {
3835 .composite_integer => unreachable, // TODO
3836 .float => switch (op) {
3837 .eq => .OpFOrdEqual,
3838 .neq => .OpFUnordNotEqual,
3839 .lt => .OpFOrdLessThan,
3840 .lte => .OpFOrdLessThanEqual,
3841 .gt => .OpFOrdGreaterThan,
3842 .gte => .OpFOrdGreaterThanEqual,
3843 },
3844 .bool => switch (op) {
3845 .eq => .OpLogicalEqual,
3846 .neq => .OpLogicalNotEqual,
3847 else => unreachable,
3848 },
3849 .integer, .strange_integer => switch (info.signedness) {
3850 .signed => switch (op) {
3851 .eq => .OpIEqual,
3852 .neq => .OpINotEqual,
3853 .lt => .OpSLessThan,
3854 .lte => .OpSLessThanEqual,
3855 .gt => .OpSGreaterThan,
3856 .gte => .OpSGreaterThanEqual,
3857 },
3858 .unsigned => switch (op) {
3859 .eq => .OpIEqual,
3860 .neq => .OpINotEqual,
3861 .lt => .OpULessThan,
3862 .lte => .OpULessThanEqual,
3863 .gt => .OpUGreaterThan,
3864 .gte => .OpUGreaterThanEqual,
3865 },
3866 },
3867 };
3868
3869 return try cg.buildCmp(pred, lhs, rhs);
3870}
3871
3872fn airCmp(
3873 cg: *CodeGen,
3874 inst: Air.Inst.Index,
3875 comptime op: std.math.CompareOperator,
3876) !?Id {
3877 const bin_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
3878 const lhs = try cg.temporary(bin_op.lhs);
3879 const rhs = try cg.temporary(bin_op.rhs);
3880
3881 const result = try cg.cmp(op, lhs, rhs);
3882 return try result.materialize(cg);
3883}
3884
3885fn airVectorCmp(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
3886 const ty_pl = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
3887 const vec_cmp = cg.air.extraData(Air.VectorCmp, ty_pl.payload).data;
3888 const lhs = try cg.temporary(vec_cmp.lhs);
3889 const rhs = try cg.temporary(vec_cmp.rhs);
3890 const op = vec_cmp.compareOperator();
3891
3892 const result = try cg.cmp(op, lhs, rhs);
3893 return try result.materialize(cg);
3894}
3895
3896/// Bitcast one type to another. Note: both types, input, output are expected in **direct** representation.
3897fn bitCast(
3898 cg: *CodeGen,
3899 dst_ty: Type,
3900 src_ty: Type,
3901 src_id: Id,
3902) !Id {
3903 const gpa = cg.module.gpa;
3904 const zcu = cg.module.zcu;
3905 const target = zcu.getTarget();
3906 const src_ty_id = try cg.resolveType(src_ty, .direct);
3907 const dst_ty_id = try cg.resolveType(dst_ty, .direct);
3908
3909 const result_id = blk: {
3910 if (src_ty_id == dst_ty_id) break :blk src_id;
3911
3912 // TODO: Some more cases are missing here
3913 // See fn bitCast in llvm.zig
3914
3915 if (src_ty.zigTypeTag(zcu) == .int and dst_ty.isPtrAtRuntime(zcu)) {
3916 if (target.os.tag != .opencl) {
3917 if (dst_ty.ptrAddressSpace(zcu) != .physical_storage_buffer) {
3918 return cg.fail(
3919 "cannot cast integer to pointer with address space '{s}'",
3920 .{@tagName(dst_ty.ptrAddressSpace(zcu))},
3921 );
3922 }
3923 }
3924
3925 const result_id = cg.module.allocId();
3926 try cg.body.emit(gpa, .OpConvertUToPtr, .{
3927 .id_result_type = dst_ty_id,
3928 .id_result = result_id,
3929 .integer_value = src_id,
3930 });
3931 break :blk result_id;
3932 }
3933
3934 // We can only use OpBitcast for specific conversions: between numerical types, and
3935 // between pointers. If the resolved spir-v types fall into this category then emit OpBitcast,
3936 // otherwise use a temporary and perform a pointer cast.
3937 const can_bitcast = (src_ty.isNumeric(zcu) and dst_ty.isNumeric(zcu)) or (src_ty.isPtrAtRuntime(zcu) and dst_ty.isPtrAtRuntime(zcu));
3938 if (can_bitcast) {
3939 const result_id = cg.module.allocId();
3940 try cg.body.emit(gpa, .OpBitcast, .{
3941 .id_result_type = dst_ty_id,
3942 .id_result = result_id,
3943 .operand = src_id,
3944 });
3945
3946 break :blk result_id;
3947 }
3948
3949 const dst_ptr_ty_id = try cg.module.ptrType(dst_ty_id, .function);
3950
3951 const src_ty_indirect_id = try cg.resolveType(src_ty, .indirect);
3952 const tmp_id = try cg.alloc(src_ty_indirect_id, null);
3953 try cg.store(src_ty, tmp_id, src_id, .{});
3954 const casted_ptr_id = cg.module.allocId();
3955 try cg.body.emit(gpa, .OpBitcast, .{
3956 .id_result_type = dst_ptr_ty_id,
3957 .id_result = casted_ptr_id,
3958 .operand = tmp_id,
3959 });
3960 break :blk try cg.load(dst_ty, casted_ptr_id, .{});
3961 };
3962
3963 // Because strange integers use sign-extended representation, we may need to normalize
3964 // the result here.
3965 // TODO: This detail could cause stuff like @as(*const i1, @ptrCast(&@as(u1, 1))) to break
3966 // should we change the representation of strange integers?
3967 if (dst_ty.zigTypeTag(zcu) == .int) {
3968 const info = cg.arithmeticTypeInfo(dst_ty);
3969 const result = try cg.normalize(Temporary.init(dst_ty, result_id), info);
3970 return try result.materialize(cg);
3971 }
3972
3973 return result_id;
3974}
3975
3976fn airBitCast(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
3977 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
3978 const operand_ty = cg.typeOf(ty_op.operand);
3979 const result_ty = cg.typeOfIndex(inst);
3980 if (operand_ty.toIntern() == .bool_type) {
3981 const operand = try cg.temporary(ty_op.operand);
3982 const result = try cg.intFromBool(operand, .u1);
3983 return try result.materialize(cg);
3984 }
3985 const operand_id = try cg.resolve(ty_op.operand);
3986 return try cg.bitCast(result_ty, operand_ty, operand_id);
3987}
3988
3989fn airIntCast(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
3990 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
3991 const src = try cg.temporary(ty_op.operand);
3992 const dst_ty = cg.typeOfIndex(inst);
3993
3994 const src_info = cg.arithmeticTypeInfo(src.ty);
3995 const dst_info = cg.arithmeticTypeInfo(dst_ty);
3996
3997 if (src_info.backing_bits == dst_info.backing_bits) {
3998 return try src.materialize(cg);
3999 }
4000
4001 const converted = try cg.buildConvert(dst_ty, src);
4002
4003 // Make sure to normalize the result if shrinking.
4004 // Because strange ints are sign extended in their backing
4005 // type, we don't need to normalize when growing the type. The
4006 // representation is already the same.
4007 const result = if (dst_info.bits < src_info.bits)
4008 try cg.normalize(converted, dst_info)
4009 else
4010 converted;
4011
4012 return try result.materialize(cg);
4013}
4014
4015fn intFromPtr(cg: *CodeGen, operand_id: Id) !Id {
4016 const result_type_id = try cg.resolveType(.usize, .direct);
4017 const result_id = cg.module.allocId();
4018 try cg.body.emit(cg.module.gpa, .OpConvertPtrToU, .{
4019 .id_result_type = result_type_id,
4020 .id_result = result_id,
4021 .pointer = operand_id,
4022 });
4023 return result_id;
4024}
4025
4026fn airFloatFromInt(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
4027 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
4028 const operand_ty = cg.typeOf(ty_op.operand);
4029 const operand_id = try cg.resolve(ty_op.operand);
4030 const result_ty = cg.typeOfIndex(inst);
4031 return try cg.floatFromInt(result_ty, operand_ty, operand_id);
4032}
4033
4034fn floatFromInt(cg: *CodeGen, result_ty: Type, operand_ty: Type, operand_id: Id) !Id {
4035 const gpa = cg.module.gpa;
4036 const operand_info = cg.arithmeticTypeInfo(operand_ty);
4037 const result_id = cg.module.allocId();
4038 const result_ty_id = try cg.resolveType(result_ty, .direct);
4039 switch (operand_info.signedness) {
4040 .signed => try cg.body.emit(gpa, .OpConvertSToF, .{
4041 .id_result_type = result_ty_id,
4042 .id_result = result_id,
4043 .signed_value = operand_id,
4044 }),
4045 .unsigned => try cg.body.emit(gpa, .OpConvertUToF, .{
4046 .id_result_type = result_ty_id,
4047 .id_result = result_id,
4048 .unsigned_value = operand_id,
4049 }),
4050 }
4051 return result_id;
4052}
4053
4054fn airIntFromFloat(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
4055 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
4056 const operand_id = try cg.resolve(ty_op.operand);
4057 const result_ty = cg.typeOfIndex(inst);
4058 return try cg.intFromFloat(result_ty, operand_id);
4059}
4060
4061fn intFromFloat(cg: *CodeGen, result_ty: Type, operand_id: Id) !Id {
4062 const gpa = cg.module.gpa;
4063 const result_info = cg.arithmeticTypeInfo(result_ty);
4064 const result_ty_id = try cg.resolveType(result_ty, .direct);
4065 const result_id = cg.module.allocId();
4066 switch (result_info.signedness) {
4067 .signed => try cg.body.emit(gpa, .OpConvertFToS, .{
4068 .id_result_type = result_ty_id,
4069 .id_result = result_id,
4070 .float_value = operand_id,
4071 }),
4072 .unsigned => try cg.body.emit(gpa, .OpConvertFToU, .{
4073 .id_result_type = result_ty_id,
4074 .id_result = result_id,
4075 .float_value = operand_id,
4076 }),
4077 }
4078 return result_id;
4079}
4080
4081fn airFloatCast(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
4082 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
4083 const operand = try cg.temporary(ty_op.operand);
4084 const dest_ty = cg.typeOfIndex(inst);
4085 const result = try cg.buildConvert(dest_ty, operand);
4086 return try result.materialize(cg);
4087}
4088
4089fn airNot(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
4090 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
4091 const operand = try cg.temporary(ty_op.operand);
4092 const result_ty = cg.typeOfIndex(inst);
4093 const info = cg.arithmeticTypeInfo(result_ty);
4094
4095 const result = switch (info.class) {
4096 .bool => try cg.buildUnary(.l_not, operand),
4097 .float => unreachable,
4098 .composite_integer => unreachable, // TODO
4099 .strange_integer, .integer => blk: {
4100 const complement = try cg.buildUnary(.bit_not, operand);
4101 break :blk try cg.normalize(complement, info);
4102 },
4103 };
4104
4105 return try result.materialize(cg);
4106}
4107
4108fn airArrayToSlice(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
4109 const zcu = cg.module.zcu;
4110 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
4111 const array_ptr_ty = cg.typeOf(ty_op.operand);
4112 const array_ty = array_ptr_ty.childType(zcu);
4113 const slice_ty = cg.typeOfIndex(inst);
4114 const elem_ptr_ty = slice_ty.slicePtrFieldType(zcu);
4115
4116 const elem_ptr_ty_id = try cg.resolveType(elem_ptr_ty, .direct);
4117
4118 const array_ptr_id = try cg.resolve(ty_op.operand);
4119 const len_id = try cg.constInt(.usize, array_ty.arrayLen(zcu));
4120
4121 const elem_ptr_id = if (!array_ty.hasRuntimeBitsIgnoreComptime(zcu))
4122 // Note: The pointer is something like *opaque{}, so we need to bitcast it to the element type.
4123 try cg.bitCast(elem_ptr_ty, array_ptr_ty, array_ptr_id)
4124 else
4125 // Convert the pointer-to-array to a pointer to the first element.
4126 try cg.accessChain(elem_ptr_ty_id, array_ptr_id, &.{0});
4127
4128 const slice_ty_id = try cg.resolveType(slice_ty, .direct);
4129 return try cg.constructComposite(slice_ty_id, &.{ elem_ptr_id, len_id });
4130}
4131
4132fn airSlice(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
4133 const ty_pl = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
4134 const bin_op = cg.air.extraData(Air.Bin, ty_pl.payload).data;
4135 const ptr_id = try cg.resolve(bin_op.lhs);
4136 const len_id = try cg.resolve(bin_op.rhs);
4137 const slice_ty = cg.typeOfIndex(inst);
4138 const slice_ty_id = try cg.resolveType(slice_ty, .direct);
4139 return try cg.constructComposite(slice_ty_id, &.{ ptr_id, len_id });
4140}
4141
4142fn airAggregateInit(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
4143 const gpa = cg.module.gpa;
4144 const pt = cg.pt;
4145 const zcu = cg.module.zcu;
4146 const ip = &zcu.intern_pool;
4147 const target = cg.module.zcu.getTarget();
4148 const ty_pl = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
4149 const result_ty = cg.typeOfIndex(inst);
4150 const len: usize = @intCast(result_ty.arrayLen(zcu));
4151 const elements: []const Air.Inst.Ref = @ptrCast(cg.air.extra.items[ty_pl.payload..][0..len]);
4152
4153 switch (result_ty.zigTypeTag(zcu)) {
4154 .@"struct" => {
4155 if (zcu.typeToPackedStruct(result_ty)) |struct_type| {
4156 comptime assert(Type.packed_struct_layout_version == 2);
4157 const backing_int_ty: Type = .fromInterned(struct_type.backingIntTypeUnordered(ip));
4158 var running_int_id = try cg.constInt(backing_int_ty, 0);
4159 var running_bits: u16 = 0;
4160 for (struct_type.field_types.get(ip), elements) |field_ty_ip, element| {
4161 const field_ty: Type = .fromInterned(field_ty_ip);
4162 if (!field_ty.hasRuntimeBitsIgnoreComptime(zcu)) continue;
4163 const field_id = try cg.resolve(element);
4164 const ty_bit_size: u16 = @intCast(field_ty.bitSize(zcu));
4165 const field_int_ty = try cg.pt.intType(.unsigned, ty_bit_size);
4166 const field_int_id = blk: {
4167 if (field_ty.isPtrAtRuntime(zcu)) {
4168 assert(target.cpu.arch == .spirv64 and
4169 field_ty.ptrAddressSpace(zcu) == .storage_buffer);
4170 break :blk try cg.intFromPtr(field_id);
4171 }
4172 break :blk try cg.bitCast(field_int_ty, field_ty, field_id);
4173 };
4174 const shift_rhs = try cg.constInt(backing_int_ty, running_bits);
4175 const extended_int_conv = try cg.buildConvert(backing_int_ty, .{
4176 .ty = field_int_ty,
4177 .value = .{ .singleton = field_int_id },
4178 });
4179 const shifted = try cg.buildBinary(.OpShiftLeftLogical, extended_int_conv, .{
4180 .ty = backing_int_ty,
4181 .value = .{ .singleton = shift_rhs },
4182 });
4183 const running_int_tmp = try cg.buildBinary(
4184 .OpBitwiseOr,
4185 .{ .ty = backing_int_ty, .value = .{ .singleton = running_int_id } },
4186 shifted,
4187 );
4188 running_int_id = try running_int_tmp.materialize(cg);
4189 running_bits += ty_bit_size;
4190 }
4191 return running_int_id;
4192 }
4193
4194 const scratch_top = cg.id_scratch.items.len;
4195 defer cg.id_scratch.shrinkRetainingCapacity(scratch_top);
4196 const constituents = try cg.id_scratch.addManyAsSlice(gpa, elements.len);
4197
4198 const types = try gpa.alloc(Type, elements.len);
4199 defer gpa.free(types);
4200
4201 var index: usize = 0;
4202
4203 switch (ip.indexToKey(result_ty.toIntern())) {
4204 .tuple_type => |tuple| {
4205 for (tuple.types.get(ip), elements, 0..) |field_ty, element, i| {
4206 if ((try result_ty.structFieldValueComptime(pt, i)) != null) continue;
4207 assert(Type.fromInterned(field_ty).hasRuntimeBits(zcu));
4208
4209 const id = try cg.resolve(element);
4210 types[index] = .fromInterned(field_ty);
4211 constituents[index] = try cg.convertToIndirect(.fromInterned(field_ty), id);
4212 index += 1;
4213 }
4214 },
4215 .struct_type => {
4216 const struct_type = ip.loadStructType(result_ty.toIntern());
4217 var it = struct_type.iterateRuntimeOrder(ip);
4218 for (elements, 0..) |element, i| {
4219 const field_index = it.next().?;
4220 if ((try result_ty.structFieldValueComptime(pt, i)) != null) continue;
4221 const field_ty: Type = .fromInterned(struct_type.field_types.get(ip)[field_index]);
4222 assert(field_ty.hasRuntimeBitsIgnoreComptime(zcu));
4223
4224 const id = try cg.resolve(element);
4225 types[index] = field_ty;
4226 constituents[index] = try cg.convertToIndirect(field_ty, id);
4227 index += 1;
4228 }
4229 },
4230 else => unreachable,
4231 }
4232
4233 const result_ty_id = try cg.resolveType(result_ty, .direct);
4234 return try cg.constructComposite(result_ty_id, constituents[0..index]);
4235 },
4236 .vector => {
4237 const n_elems = result_ty.vectorLen(zcu);
4238 const scratch_top = cg.id_scratch.items.len;
4239 defer cg.id_scratch.shrinkRetainingCapacity(scratch_top);
4240 const elem_ids = try cg.id_scratch.addManyAsSlice(gpa, n_elems);
4241
4242 for (elements, 0..) |element, i| {
4243 elem_ids[i] = try cg.resolve(element);
4244 }
4245
4246 const result_ty_id = try cg.resolveType(result_ty, .direct);
4247 return try cg.constructComposite(result_ty_id, elem_ids);
4248 },
4249 .array => {
4250 const array_info = result_ty.arrayInfo(zcu);
4251 const n_elems: usize = @intCast(result_ty.arrayLenIncludingSentinel(zcu));
4252 const scratch_top = cg.id_scratch.items.len;
4253 defer cg.id_scratch.shrinkRetainingCapacity(scratch_top);
4254 const elem_ids = try cg.id_scratch.addManyAsSlice(gpa, n_elems);
4255
4256 for (elements, 0..) |element, i| {
4257 const id = try cg.resolve(element);
4258 elem_ids[i] = try cg.convertToIndirect(array_info.elem_type, id);
4259 }
4260
4261 if (array_info.sentinel) |sentinel_val| {
4262 elem_ids[n_elems - 1] = try cg.constant(array_info.elem_type, sentinel_val, .indirect);
4263 }
4264
4265 const result_ty_id = try cg.resolveType(result_ty, .direct);
4266 return try cg.constructComposite(result_ty_id, elem_ids);
4267 },
4268 else => unreachable,
4269 }
4270}
4271
4272fn sliceOrArrayLen(cg: *CodeGen, operand_id: Id, ty: Type) !Id {
4273 const zcu = cg.module.zcu;
4274 switch (ty.ptrSize(zcu)) {
4275 .slice => return cg.extractField(.usize, operand_id, 1),
4276 .one => {
4277 const array_ty = ty.childType(zcu);
4278 const elem_ty = array_ty.childType(zcu);
4279 const abi_size = elem_ty.abiSize(zcu);
4280 const size = array_ty.arrayLenIncludingSentinel(zcu) * abi_size;
4281 return try cg.constInt(.usize, size);
4282 },
4283 .many, .c => unreachable,
4284 }
4285}
4286
4287fn sliceOrArrayPtr(cg: *CodeGen, operand_id: Id, ty: Type) !Id {
4288 const zcu = cg.module.zcu;
4289 if (ty.isSlice(zcu)) {
4290 const ptr_ty = ty.slicePtrFieldType(zcu);
4291 return cg.extractField(ptr_ty, operand_id, 0);
4292 }
4293 return operand_id;
4294}
4295
4296fn airMemcpy(cg: *CodeGen, inst: Air.Inst.Index) !void {
4297 const bin_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
4298 const dest_slice = try cg.resolve(bin_op.lhs);
4299 const src_slice = try cg.resolve(bin_op.rhs);
4300 const dest_ty = cg.typeOf(bin_op.lhs);
4301 const src_ty = cg.typeOf(bin_op.rhs);
4302 const dest_ptr = try cg.sliceOrArrayPtr(dest_slice, dest_ty);
4303 const src_ptr = try cg.sliceOrArrayPtr(src_slice, src_ty);
4304 const len = try cg.sliceOrArrayLen(dest_slice, dest_ty);
4305 try cg.body.emit(cg.module.gpa, .OpCopyMemorySized, .{
4306 .target = dest_ptr,
4307 .source = src_ptr,
4308 .size = len,
4309 });
4310}
4311
4312fn airMemmove(cg: *CodeGen, inst: Air.Inst.Index) !void {
4313 _ = inst;
4314 return cg.fail("TODO implement airMemcpy for spirv", .{});
4315}
4316
4317fn airSliceField(cg: *CodeGen, inst: Air.Inst.Index, field: u32) !?Id {
4318 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
4319 const field_ty = cg.typeOfIndex(inst);
4320 const operand_id = try cg.resolve(ty_op.operand);
4321 return try cg.extractField(field_ty, operand_id, field);
4322}
4323
4324fn airSliceElemPtr(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
4325 const zcu = cg.module.zcu;
4326 const ty_pl = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
4327 const bin_op = cg.air.extraData(Air.Bin, ty_pl.payload).data;
4328 const slice_ty = cg.typeOf(bin_op.lhs);
4329 if (!slice_ty.isVolatilePtr(zcu) and cg.liveness.isUnused(inst)) return null;
4330
4331 const slice_id = try cg.resolve(bin_op.lhs);
4332 const index_id = try cg.resolve(bin_op.rhs);
4333
4334 const ptr_ty = cg.typeOfIndex(inst);
4335 const ptr_ty_id = try cg.resolveType(ptr_ty, .direct);
4336
4337 const slice_ptr = try cg.extractField(ptr_ty, slice_id, 0);
4338 return try cg.ptrAccessChain(ptr_ty_id, slice_ptr, index_id, &.{});
4339}
4340
4341fn airSliceElemVal(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
4342 const zcu = cg.module.zcu;
4343 const bin_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
4344 const slice_ty = cg.typeOf(bin_op.lhs);
4345 if (!slice_ty.isVolatilePtr(zcu) and cg.liveness.isUnused(inst)) return null;
4346
4347 const slice_id = try cg.resolve(bin_op.lhs);
4348 const index_id = try cg.resolve(bin_op.rhs);
4349
4350 const ptr_ty = slice_ty.slicePtrFieldType(zcu);
4351 const ptr_ty_id = try cg.resolveType(ptr_ty, .direct);
4352
4353 const slice_ptr = try cg.extractField(ptr_ty, slice_id, 0);
4354 const elem_ptr = try cg.ptrAccessChain(ptr_ty_id, slice_ptr, index_id, &.{});
4355 return try cg.load(slice_ty.childType(zcu), elem_ptr, .{ .is_volatile = slice_ty.isVolatilePtr(zcu) });
4356}
4357
4358fn ptrElemPtr(cg: *CodeGen, ptr_ty: Type, ptr_id: Id, index_id: Id) !Id {
4359 const zcu = cg.module.zcu;
4360 // Construct new pointer type for the resulting pointer
4361 const elem_ty = ptr_ty.elemType2(zcu); // use elemType() so that we get T for *[N]T.
4362 const elem_ty_id = try cg.resolveType(elem_ty, .indirect);
4363 const elem_ptr_ty_id = try cg.module.ptrType(elem_ty_id, cg.module.storageClass(ptr_ty.ptrAddressSpace(zcu)));
4364 if (ptr_ty.isSinglePointer(zcu)) {
4365 // Pointer-to-array. In this case, the resulting pointer is not of the same type
4366 // as the ptr_ty (we want a *T, not a *[N]T), and hence we need to use accessChain.
4367 return try cg.accessChainId(elem_ptr_ty_id, ptr_id, &.{index_id});
4368 } else {
4369 // Resulting pointer type is the same as the ptr_ty, so use ptrAccessChain
4370 return try cg.ptrAccessChain(elem_ptr_ty_id, ptr_id, index_id, &.{});
4371 }
4372}
4373
4374fn airPtrElemPtr(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
4375 const zcu = cg.module.zcu;
4376 const ty_pl = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
4377 const bin_op = cg.air.extraData(Air.Bin, ty_pl.payload).data;
4378 const src_ptr_ty = cg.typeOf(bin_op.lhs);
4379 const elem_ty = src_ptr_ty.childType(zcu);
4380 const ptr_id = try cg.resolve(bin_op.lhs);
4381
4382 if (!elem_ty.hasRuntimeBitsIgnoreComptime(zcu)) {
4383 const dst_ptr_ty = cg.typeOfIndex(inst);
4384 return try cg.bitCast(dst_ptr_ty, src_ptr_ty, ptr_id);
4385 }
4386
4387 const index_id = try cg.resolve(bin_op.rhs);
4388 return try cg.ptrElemPtr(src_ptr_ty, ptr_id, index_id);
4389}
4390
4391fn airArrayElemVal(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
4392 const gpa = cg.module.gpa;
4393 const zcu = cg.module.zcu;
4394 const bin_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
4395 const array_ty = cg.typeOf(bin_op.lhs);
4396 const elem_ty = array_ty.childType(zcu);
4397 const array_id = try cg.resolve(bin_op.lhs);
4398 const index_id = try cg.resolve(bin_op.rhs);
4399
4400 // SPIR-V doesn't have an array indexing function for some damn reason.
4401 // For now, just generate a temporary and use that.
4402 // TODO: This backend probably also should use isByRef from llvm...
4403
4404 const is_vector = array_ty.isVector(zcu);
4405 const elem_repr: Repr = if (is_vector) .direct else .indirect;
4406 const array_ty_id = try cg.resolveType(array_ty, .direct);
4407 const elem_ty_id = try cg.resolveType(elem_ty, elem_repr);
4408 const ptr_array_ty_id = try cg.module.ptrType(array_ty_id, .function);
4409 const ptr_elem_ty_id = try cg.module.ptrType(elem_ty_id, .function);
4410
4411 const tmp_id = cg.module.allocId();
4412 try cg.prologue.emit(gpa, .OpVariable, .{
4413 .id_result_type = ptr_array_ty_id,
4414 .id_result = tmp_id,
4415 .storage_class = .function,
4416 });
4417
4418 try cg.body.emit(gpa, .OpStore, .{
4419 .pointer = tmp_id,
4420 .object = array_id,
4421 });
4422
4423 const elem_ptr_id = try cg.accessChainId(ptr_elem_ty_id, tmp_id, &.{index_id});
4424
4425 const result_id = cg.module.allocId();
4426 try cg.body.emit(gpa, .OpLoad, .{
4427 .id_result_type = try cg.resolveType(elem_ty, elem_repr),
4428 .id_result = result_id,
4429 .pointer = elem_ptr_id,
4430 });
4431
4432 if (is_vector) {
4433 // Result is already in direct representation
4434 return result_id;
4435 }
4436
4437 // This is an array type; the elements are stored in indirect representation.
4438 // We have to convert the type to direct.
4439
4440 return try cg.convertToDirect(elem_ty, result_id);
4441}
4442
4443fn airPtrElemVal(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
4444 const zcu = cg.module.zcu;
4445 const bin_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
4446 const ptr_ty = cg.typeOf(bin_op.lhs);
4447 const elem_ty = cg.typeOfIndex(inst);
4448 const ptr_id = try cg.resolve(bin_op.lhs);
4449 const index_id = try cg.resolve(bin_op.rhs);
4450 const elem_ptr_id = try cg.ptrElemPtr(ptr_ty, ptr_id, index_id);
4451 return try cg.load(elem_ty, elem_ptr_id, .{ .is_volatile = ptr_ty.isVolatilePtr(zcu) });
4452}
4453
4454fn airSetUnionTag(cg: *CodeGen, inst: Air.Inst.Index) !void {
4455 const zcu = cg.module.zcu;
4456 const bin_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
4457 const un_ptr_ty = cg.typeOf(bin_op.lhs);
4458 const un_ty = un_ptr_ty.childType(zcu);
4459 const layout = cg.unionLayout(un_ty);
4460
4461 if (layout.tag_size == 0) return;
4462
4463 const tag_ty = un_ty.unionTagTypeSafety(zcu).?;
4464 const tag_ty_id = try cg.resolveType(tag_ty, .indirect);
4465 const tag_ptr_ty_id = try cg.module.ptrType(tag_ty_id, cg.module.storageClass(un_ptr_ty.ptrAddressSpace(zcu)));
4466
4467 const union_ptr_id = try cg.resolve(bin_op.lhs);
4468 const new_tag_id = try cg.resolve(bin_op.rhs);
4469
4470 if (!layout.has_payload) {
4471 try cg.store(tag_ty, union_ptr_id, new_tag_id, .{ .is_volatile = un_ptr_ty.isVolatilePtr(zcu) });
4472 } else {
4473 const ptr_id = try cg.accessChain(tag_ptr_ty_id, union_ptr_id, &.{layout.tag_index});
4474 try cg.store(tag_ty, ptr_id, new_tag_id, .{ .is_volatile = un_ptr_ty.isVolatilePtr(zcu) });
4475 }
4476}
4477
4478fn airGetUnionTag(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
4479 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
4480 const un_ty = cg.typeOf(ty_op.operand);
4481
4482 const zcu = cg.module.zcu;
4483 const layout = cg.unionLayout(un_ty);
4484 if (layout.tag_size == 0) return null;
4485
4486 const union_handle = try cg.resolve(ty_op.operand);
4487 if (!layout.has_payload) return union_handle;
4488
4489 const tag_ty = un_ty.unionTagTypeSafety(zcu).?;
4490 return try cg.extractField(tag_ty, union_handle, layout.tag_index);
4491}
4492
4493fn unionInit(
4494 cg: *CodeGen,
4495 ty: Type,
4496 active_field: u32,
4497 payload: ?Id,
4498) !Id {
4499 // To initialize a union, generate a temporary variable with the
4500 // union type, then get the field pointer and pointer-cast it to the
4501 // right type to store it. Finally load the entire union.
4502
4503 // Note: The result here is not cached, because it generates runtime code.
4504
4505 const pt = cg.pt;
4506 const zcu = cg.module.zcu;
4507 const ip = &zcu.intern_pool;
4508 const union_ty = zcu.typeToUnion(ty).?;
4509 const tag_ty: Type = .fromInterned(union_ty.enum_tag_ty);
4510
4511 const layout = cg.unionLayout(ty);
4512 const payload_ty: Type = .fromInterned(union_ty.field_types.get(ip)[active_field]);
4513
4514 if (union_ty.flagsUnordered(ip).layout == .@"packed") {
4515 if (!payload_ty.hasRuntimeBitsIgnoreComptime(zcu)) {
4516 const int_ty = try pt.intType(.unsigned, @intCast(ty.bitSize(zcu)));
4517 return cg.constInt(int_ty, 0);
4518 }
4519
4520 assert(payload != null);
4521 if (payload_ty.isInt(zcu)) {
4522 if (ty.bitSize(zcu) == payload_ty.bitSize(zcu)) {
4523 return cg.bitCast(ty, payload_ty, payload.?);
4524 }
4525
4526 const trunc = try cg.buildConvert(ty, .{ .ty = payload_ty, .value = .{ .singleton = payload.? } });
4527 return try trunc.materialize(cg);
4528 }
4529
4530 const payload_int_ty = try pt.intType(.unsigned, @intCast(payload_ty.bitSize(zcu)));
4531 const payload_int = if (payload_ty.ip_index == .bool_type)
4532 try cg.convertToIndirect(payload_ty, payload.?)
4533 else
4534 try cg.bitCast(payload_int_ty, payload_ty, payload.?);
4535 const trunc = try cg.buildConvert(ty, .{ .ty = payload_int_ty, .value = .{ .singleton = payload_int } });
4536 return try trunc.materialize(cg);
4537 }
4538
4539 const tag_int = if (layout.tag_size != 0) blk: {
4540 const tag_val = try pt.enumValueFieldIndex(tag_ty, active_field);
4541 const tag_int_val = try tag_val.intFromEnum(tag_ty, pt);
4542 break :blk tag_int_val.toUnsignedInt(zcu);
4543 } else 0;
4544
4545 if (!layout.has_payload) {
4546 return try cg.constInt(tag_ty, tag_int);
4547 }
4548
4549 const ty_id = try cg.resolveType(ty, .indirect);
4550 const tmp_id = try cg.alloc(ty_id, null);
4551
4552 if (layout.tag_size != 0) {
4553 const tag_ty_id = try cg.resolveType(tag_ty, .indirect);
4554 const tag_ptr_ty_id = try cg.module.ptrType(tag_ty_id, .function);
4555 const ptr_id = try cg.accessChain(tag_ptr_ty_id, tmp_id, &.{@as(u32, @intCast(layout.tag_index))});
4556 const tag_id = try cg.constInt(tag_ty, tag_int);
4557 try cg.store(tag_ty, ptr_id, tag_id, .{});
4558 }
4559
4560 if (payload_ty.hasRuntimeBitsIgnoreComptime(zcu)) {
4561 const layout_payload_ty_id = try cg.resolveType(layout.payload_ty, .indirect);
4562 const pl_ptr_ty_id = try cg.module.ptrType(layout_payload_ty_id, .function);
4563 const pl_ptr_id = try cg.accessChain(pl_ptr_ty_id, tmp_id, &.{layout.payload_index});
4564 const active_pl_ptr_id = if (!layout.payload_ty.eql(payload_ty, zcu)) blk: {
4565 const payload_ty_id = try cg.resolveType(payload_ty, .indirect);
4566 const active_pl_ptr_ty_id = try cg.module.ptrType(payload_ty_id, .function);
4567 const active_pl_ptr_id = cg.module.allocId();
4568 try cg.body.emit(cg.module.gpa, .OpBitcast, .{
4569 .id_result_type = active_pl_ptr_ty_id,
4570 .id_result = active_pl_ptr_id,
4571 .operand = pl_ptr_id,
4572 });
4573 break :blk active_pl_ptr_id;
4574 } else pl_ptr_id;
4575
4576 try cg.store(payload_ty, active_pl_ptr_id, payload.?, .{});
4577 } else {
4578 assert(payload == null);
4579 }
4580
4581 // Just leave the padding fields uninitialized...
4582 // TODO: Or should we initialize them with undef explicitly?
4583
4584 return try cg.load(ty, tmp_id, .{});
4585}
4586
4587fn airUnionInit(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
4588 const zcu = cg.module.zcu;
4589 const ip = &zcu.intern_pool;
4590 const ty_pl = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
4591 const extra = cg.air.extraData(Air.UnionInit, ty_pl.payload).data;
4592 const ty = cg.typeOfIndex(inst);
4593
4594 const union_obj = zcu.typeToUnion(ty).?;
4595 const field_ty: Type = .fromInterned(union_obj.field_types.get(ip)[extra.field_index]);
4596 const payload = if (field_ty.hasRuntimeBitsIgnoreComptime(zcu))
4597 try cg.resolve(extra.init)
4598 else
4599 null;
4600 return try cg.unionInit(ty, extra.field_index, payload);
4601}
4602
4603fn airStructFieldVal(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
4604 const pt = cg.pt;
4605 const zcu = cg.module.zcu;
4606 const ty_pl = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
4607 const struct_field = cg.air.extraData(Air.StructField, ty_pl.payload).data;
4608
4609 const object_ty = cg.typeOf(struct_field.struct_operand);
4610 const object_id = try cg.resolve(struct_field.struct_operand);
4611 const field_index = struct_field.field_index;
4612 const field_ty = object_ty.fieldType(field_index, zcu);
4613
4614 if (!field_ty.hasRuntimeBitsIgnoreComptime(zcu)) return null;
4615
4616 switch (object_ty.zigTypeTag(zcu)) {
4617 .@"struct" => switch (object_ty.containerLayout(zcu)) {
4618 .@"packed" => {
4619 const struct_ty = zcu.typeToPackedStruct(object_ty).?;
4620 const struct_backing_int_bits = cg.module.backingIntBits(@intCast(object_ty.bitSize(zcu))).@"0";
4621 const bit_offset = zcu.structPackedFieldBitOffset(struct_ty, field_index);
4622 // We use the same int type the packed struct is backed by, because even though it would
4623 // be valid SPIR-V to use an smaller type like u16, some implementations like PoCL will complain.
4624 const bit_offset_id = try cg.constInt(object_ty, bit_offset);
4625 const signedness = if (field_ty.isInt(zcu)) field_ty.intInfo(zcu).signedness else .unsigned;
4626 const field_bit_size: u16 = @intCast(field_ty.bitSize(zcu));
4627 const field_int_ty = try pt.intType(signedness, field_bit_size);
4628 const shift_lhs: Temporary = .{ .ty = object_ty, .value = .{ .singleton = object_id } };
4629 const shift = try cg.buildBinary(.OpShiftRightLogical, shift_lhs, .{ .ty = object_ty, .value = .{ .singleton = bit_offset_id } });
4630 const mask_id = try cg.constInt(object_ty, (@as(u64, 1) << @as(u6, @intCast(field_bit_size))) - 1);
4631 const masked = try cg.buildBinary(.OpBitwiseAnd, shift, .{ .ty = object_ty, .value = .{ .singleton = mask_id } });
4632 const result_id = blk: {
4633 if (cg.module.backingIntBits(field_bit_size).@"0" == struct_backing_int_bits)
4634 break :blk try cg.bitCast(field_int_ty, object_ty, try masked.materialize(cg));
4635 const trunc = try cg.buildConvert(field_int_ty, masked);
4636 break :blk try trunc.materialize(cg);
4637 };
4638 if (field_ty.ip_index == .bool_type) return try cg.convertToDirect(.bool, result_id);
4639 if (field_ty.isInt(zcu)) return result_id;
4640 return try cg.bitCast(field_ty, field_int_ty, result_id);
4641 },
4642 else => return try cg.extractField(field_ty, object_id, field_index),
4643 },
4644 .@"union" => switch (object_ty.containerLayout(zcu)) {
4645 .@"packed" => {
4646 const backing_int_ty = try pt.intType(.unsigned, @intCast(object_ty.bitSize(zcu)));
4647 const signedness = if (field_ty.isInt(zcu)) field_ty.intInfo(zcu).signedness else .unsigned;
4648 const field_bit_size: u16 = @intCast(field_ty.bitSize(zcu));
4649 const int_ty = try pt.intType(signedness, field_bit_size);
4650 const mask_id = try cg.constInt(backing_int_ty, (@as(u64, 1) << @as(u6, @intCast(field_bit_size))) - 1);
4651 const masked = try cg.buildBinary(
4652 .OpBitwiseAnd,
4653 .{ .ty = backing_int_ty, .value = .{ .singleton = object_id } },
4654 .{ .ty = backing_int_ty, .value = .{ .singleton = mask_id } },
4655 );
4656 const result_id = blk: {
4657 if (cg.module.backingIntBits(field_bit_size).@"0" == cg.module.backingIntBits(@intCast(backing_int_ty.bitSize(zcu))).@"0")
4658 break :blk try cg.bitCast(int_ty, backing_int_ty, try masked.materialize(cg));
4659 const trunc = try cg.buildConvert(int_ty, masked);
4660 break :blk try trunc.materialize(cg);
4661 };
4662 if (field_ty.ip_index == .bool_type) return try cg.convertToDirect(.bool, result_id);
4663 if (field_ty.isInt(zcu)) return result_id;
4664 return try cg.bitCast(field_ty, int_ty, result_id);
4665 },
4666 else => {
4667 // Store, ptr-elem-ptr, pointer-cast, load
4668 const layout = cg.unionLayout(object_ty);
4669 assert(layout.has_payload);
4670
4671 const object_ty_id = try cg.resolveType(object_ty, .indirect);
4672 const tmp_id = try cg.alloc(object_ty_id, null);
4673 try cg.store(object_ty, tmp_id, object_id, .{});
4674
4675 const layout_payload_ty_id = try cg.resolveType(layout.payload_ty, .indirect);
4676 const pl_ptr_ty_id = try cg.module.ptrType(layout_payload_ty_id, .function);
4677 const pl_ptr_id = try cg.accessChain(pl_ptr_ty_id, tmp_id, &.{layout.payload_index});
4678
4679 const field_ty_id = try cg.resolveType(field_ty, .indirect);
4680 const active_pl_ptr_ty_id = try cg.module.ptrType(field_ty_id, .function);
4681 const active_pl_ptr_id = cg.module.allocId();
4682 try cg.body.emit(cg.module.gpa, .OpBitcast, .{
4683 .id_result_type = active_pl_ptr_ty_id,
4684 .id_result = active_pl_ptr_id,
4685 .operand = pl_ptr_id,
4686 });
4687 return try cg.load(field_ty, active_pl_ptr_id, .{});
4688 },
4689 },
4690 else => unreachable,
4691 }
4692}
4693
4694fn airFieldParentPtr(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
4695 const zcu = cg.module.zcu;
4696 const target = zcu.getTarget();
4697 const ty_pl = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
4698 const extra = cg.air.extraData(Air.FieldParentPtr, ty_pl.payload).data;
4699
4700 const parent_ptr_ty = ty_pl.ty.toType();
4701 const parent_ty = parent_ptr_ty.childType(zcu);
4702 const result_ty_id = try cg.resolveType(parent_ptr_ty, .indirect);
4703
4704 const field_ptr = try cg.resolve(extra.field_ptr);
4705 const field_ptr_ty = cg.typeOf(extra.field_ptr);
4706 const field_ptr_int = try cg.intFromPtr(field_ptr);
4707 const field_offset = parent_ty.structFieldOffset(extra.field_index, zcu);
4708
4709 const base_ptr_int = base_ptr_int: {
4710 if (field_offset == 0) break :base_ptr_int field_ptr_int;
4711
4712 const field_offset_id = try cg.constInt(.usize, field_offset);
4713 const field_ptr_tmp: Temporary = .init(.usize, field_ptr_int);
4714 const field_offset_tmp: Temporary = .init(.usize, field_offset_id);
4715 const result = try cg.buildBinary(.OpISub, field_ptr_tmp, field_offset_tmp);
4716 break :base_ptr_int try result.materialize(cg);
4717 };
4718
4719 if (target.os.tag != .opencl) {
4720 if (field_ptr_ty.ptrAddressSpace(zcu) != .physical_storage_buffer) {
4721 return cg.fail(
4722 "cannot cast integer to pointer with address space '{s}'",
4723 .{@tagName(field_ptr_ty.ptrAddressSpace(zcu))},
4724 );
4725 }
4726 }
4727
4728 const base_ptr = cg.module.allocId();
4729 try cg.body.emit(cg.module.gpa, .OpConvertUToPtr, .{
4730 .id_result_type = result_ty_id,
4731 .id_result = base_ptr,
4732 .integer_value = base_ptr_int,
4733 });
4734
4735 return base_ptr;
4736}
4737
4738fn structFieldPtr(
4739 cg: *CodeGen,
4740 result_ptr_ty: Type,
4741 object_ptr_ty: Type,
4742 object_ptr: Id,
4743 field_index: u32,
4744) !Id {
4745 const result_ty_id = try cg.resolveType(result_ptr_ty, .direct);
4746
4747 const zcu = cg.module.zcu;
4748 const object_ty = object_ptr_ty.childType(zcu);
4749 switch (object_ty.zigTypeTag(zcu)) {
4750 .pointer => {
4751 assert(object_ty.isSlice(zcu));
4752 return cg.accessChain(result_ty_id, object_ptr, &.{field_index});
4753 },
4754 .@"struct" => switch (object_ty.containerLayout(zcu)) {
4755 .@"packed" => return cg.todo("implement field access for packed structs", .{}),
4756 else => {
4757 return try cg.accessChain(result_ty_id, object_ptr, &.{field_index});
4758 },
4759 },
4760 .@"union" => {
4761 const layout = cg.unionLayout(object_ty);
4762 if (!layout.has_payload) {
4763 // Asked to get a pointer to a zero-sized field. Just lower this
4764 // to undefined, there is no reason to make it be a valid pointer.
4765 return try cg.module.constUndef(result_ty_id);
4766 }
4767
4768 const storage_class = cg.module.storageClass(object_ptr_ty.ptrAddressSpace(zcu));
4769 const layout_payload_ty_id = try cg.resolveType(layout.payload_ty, .indirect);
4770 const pl_ptr_ty_id = try cg.module.ptrType(layout_payload_ty_id, storage_class);
4771 const pl_ptr_id = blk: {
4772 if (object_ty.containerLayout(zcu) == .@"packed") break :blk object_ptr;
4773 break :blk try cg.accessChain(pl_ptr_ty_id, object_ptr, &.{layout.payload_index});
4774 };
4775
4776 const active_pl_ptr_id = cg.module.allocId();
4777 try cg.body.emit(cg.module.gpa, .OpBitcast, .{
4778 .id_result_type = result_ty_id,
4779 .id_result = active_pl_ptr_id,
4780 .operand = pl_ptr_id,
4781 });
4782 return active_pl_ptr_id;
4783 },
4784 else => unreachable,
4785 }
4786}
4787
4788fn airStructFieldPtrIndex(cg: *CodeGen, inst: Air.Inst.Index, field_index: u32) !?Id {
4789 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
4790 const struct_ptr = try cg.resolve(ty_op.operand);
4791 const struct_ptr_ty = cg.typeOf(ty_op.operand);
4792 const result_ptr_ty = cg.typeOfIndex(inst);
4793 return try cg.structFieldPtr(result_ptr_ty, struct_ptr_ty, struct_ptr, field_index);
4794}
4795
4796fn alloc(cg: *CodeGen, ty_id: Id, initializer: ?Id) !Id {
4797 const ptr_ty_id = try cg.module.ptrType(ty_id, .function);
4798 const result_id = cg.module.allocId();
4799 try cg.prologue.emit(cg.module.gpa, .OpVariable, .{
4800 .id_result_type = ptr_ty_id,
4801 .id_result = result_id,
4802 .storage_class = .function,
4803 .initializer = initializer,
4804 });
4805 return result_id;
4806}
4807
4808fn airAlloc(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
4809 const zcu = cg.module.zcu;
4810 const target = zcu.getTarget();
4811 const ptr_ty = cg.typeOfIndex(inst);
4812 const child_ty = ptr_ty.childType(zcu);
4813 const child_ty_id = try cg.resolveType(child_ty, .indirect);
4814 const ptr_align = ptr_ty.ptrAlignment(zcu);
4815 const result_id = try cg.alloc(child_ty_id, null);
4816 if (ptr_align != child_ty.abiAlignment(zcu)) {
4817 if (target.os.tag != .opencl) return cg.fail("cannot apply alignment to variables", .{});
4818 try cg.module.decorate(result_id, .{
4819 .alignment = .{ .alignment = @intCast(ptr_align.toByteUnits().?) },
4820 });
4821 }
4822 return result_id;
4823}
4824
4825fn airArg(cg: *CodeGen) Id {
4826 defer cg.next_arg_index += 1;
4827 return cg.args.items[cg.next_arg_index];
4828}
4829
4830/// Given a slice of incoming block connections, returns the block-id of the next
4831/// block to jump to. This function emits instructions, so it should be emitted
4832/// inside the merge block of the block.
4833/// This function should only be called with structured control flow generation.
4834fn structuredNextBlock(cg: *CodeGen, incoming: []const ControlFlow.Structured.Block.Incoming) !Id {
4835 assert(cg.control_flow == .structured);
4836
4837 const result_id = cg.module.allocId();
4838 const block_id_ty_id = try cg.resolveType(.u32, .direct);
4839 try cg.body.emitRaw(cg.module.gpa, .OpPhi, @intCast(2 + incoming.len * 2)); // result type + result + variable/parent...
4840 cg.body.writeOperand(Id, block_id_ty_id);
4841 cg.body.writeOperand(Id, result_id);
4842
4843 for (incoming) |incoming_block| {
4844 cg.body.writeOperand(spec.PairIdRefIdRef, .{ incoming_block.next_block, incoming_block.src_label });
4845 }
4846
4847 return result_id;
4848}
4849
4850/// Jumps to the block with the target block-id. This function must only be called when
4851/// terminating a body, there should be no instructions after it.
4852/// This function should only be called with structured control flow generation.
4853fn structuredBreak(cg: *CodeGen, target_block: Id) !void {
4854 assert(cg.control_flow == .structured);
4855
4856 const gpa = cg.module.gpa;
4857 const sblock = cg.control_flow.structured.block_stack.getLast();
4858 const merge_block = switch (sblock.*) {
4859 .selection => |*merge| blk: {
4860 const merge_label = cg.module.allocId();
4861 try merge.merge_stack.append(gpa, .{
4862 .incoming = .{
4863 .src_label = cg.block_label,
4864 .next_block = target_block,
4865 },
4866 .merge_block = merge_label,
4867 });
4868 break :blk merge_label;
4869 },
4870 // Loop blocks do not end in a break. Not through a direct break,
4871 // and also not through another instruction like cond_br or unreachable (these
4872 // situations are replaced by `cond_br` in sema, or there is a `block` instruction
4873 // placed around them).
4874 .loop => unreachable,
4875 };
4876
4877 try cg.body.emit(gpa, .OpBranch, .{ .target_label = merge_block });
4878}
4879
4880/// Generate a body in a way that exits the body using only structured constructs.
4881/// Returns the block-id of the next block to jump to. After this function, a jump
4882/// should still be emitted to the block that should follow this structured body.
4883/// This function should only be called with structured control flow generation.
4884fn genStructuredBody(
4885 cg: *CodeGen,
4886 /// This parameter defines the method that this structured body is exited with.
4887 block_merge_type: union(enum) {
4888 /// Using selection; early exits from this body are surrounded with
4889 /// if() statements.
4890 selection,
4891 /// Using loops; loops can be early exited by jumping to the merge block at
4892 /// any time.
4893 loop: struct {
4894 merge_label: Id,
4895 continue_label: Id,
4896 },
4897 },
4898 body: []const Air.Inst.Index,
4899) !Id {
4900 assert(cg.control_flow == .structured);
4901
4902 const gpa = cg.module.gpa;
4903
4904 var sblock: ControlFlow.Structured.Block = switch (block_merge_type) {
4905 .loop => |merge| .{ .loop = .{
4906 .merge_block = merge.merge_label,
4907 } },
4908 .selection => .{ .selection = .{} },
4909 };
4910 defer sblock.deinit(gpa);
4911
4912 {
4913 try cg.control_flow.structured.block_stack.append(gpa, &sblock);
4914 defer _ = cg.control_flow.structured.block_stack.pop();
4915
4916 try cg.genBody(body);
4917 }
4918
4919 switch (sblock) {
4920 .selection => |merge| {
4921 // Now generate the merge block for all merges that
4922 // still need to be performed.
4923 const merge_stack = merge.merge_stack.items;
4924
4925 // If no merges on the stack, this block didn't generate any jumps (all paths
4926 // ended with a return or an unreachable). In that case, we don't need to do
4927 // any merging.
4928 if (merge_stack.len == 0) {
4929 // We still need to return a value of a next block to jump to.
4930 // For example, if we have code like
4931 // if (x) {
4932 // if (y) return else return;
4933 // } else {}
4934 // then we still need the outer to have an OpSelectionMerge and consequently
4935 // a phi node. In that case we can just return bogus, since we know that its
4936 // path will never be taken.
4937
4938 // Make sure that we are still in a block when exiting the function.
4939 // TODO: Can we get rid of that?
4940 try cg.beginSpvBlock(cg.module.allocId());
4941 const block_id_ty_id = try cg.resolveType(.u32, .direct);
4942 return try cg.module.constUndef(block_id_ty_id);
4943 }
4944
4945 // The top-most merge actually only has a single source, the
4946 // final jump of the block, or the merge block of a sub-block, cond_br,
4947 // or loop. Therefore we just need to generate a block with a jump to the
4948 // next merge block.
4949 try cg.beginSpvBlock(merge_stack[merge_stack.len - 1].merge_block);
4950
4951 // Now generate a merge ladder for the remaining merges in the stack.
4952 var incoming: ControlFlow.Structured.Block.Incoming = .{
4953 .src_label = cg.block_label,
4954 .next_block = merge_stack[merge_stack.len - 1].incoming.next_block,
4955 };
4956 var i = merge_stack.len - 1;
4957 while (i > 0) {
4958 i -= 1;
4959 const step = merge_stack[i];
4960
4961 try cg.body.emit(gpa, .OpBranch, .{ .target_label = step.merge_block });
4962 try cg.beginSpvBlock(step.merge_block);
4963 const next_block = try cg.structuredNextBlock(&.{ incoming, step.incoming });
4964 incoming = .{
4965 .src_label = step.merge_block,
4966 .next_block = next_block,
4967 };
4968 }
4969
4970 return incoming.next_block;
4971 },
4972 .loop => |merge| {
4973 // Close the loop by jumping to the continue label
4974
4975 try cg.body.emit(gpa, .OpBranch, .{ .target_label = block_merge_type.loop.continue_label });
4976 // For blocks we must simple merge all the incoming blocks to get the next block.
4977 try cg.beginSpvBlock(merge.merge_block);
4978 return try cg.structuredNextBlock(merge.merges.items);
4979 },
4980 }
4981}
4982
4983fn airBlock(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
4984 const inst_datas = cg.air.instructions.items(.data);
4985 const extra = cg.air.extraData(Air.Block, inst_datas[@intFromEnum(inst)].ty_pl.payload);
4986 return cg.lowerBlock(inst, @ptrCast(cg.air.extra.items[extra.end..][0..extra.data.body_len]));
4987}
4988
4989fn lowerBlock(cg: *CodeGen, inst: Air.Inst.Index, body: []const Air.Inst.Index) !?Id {
4990 // In AIR, a block doesn't really define an entry point like a block, but
4991 // more like a scope that breaks can jump out of and "return" a value from.
4992 // This cannot be directly modelled in SPIR-V, so in a block instruction,
4993 // we're going to split up the current block by first generating the code
4994 // of the block, then a label, and then generate the rest of the current
4995 // ir.Block in a different SPIR-V block.
4996
4997 const gpa = cg.module.gpa;
4998 const zcu = cg.module.zcu;
4999 const ty = cg.typeOfIndex(inst);
5000 const have_block_result = ty.isFnOrHasRuntimeBitsIgnoreComptime(zcu);
5001
5002 const cf = switch (cg.control_flow) {
5003 .structured => |*cf| cf,
5004 .unstructured => |*cf| {
5005 var block: ControlFlow.Unstructured.Block = .{};
5006 defer block.incoming_blocks.deinit(gpa);
5007
5008 // 4 chosen as arbitrary initial capacity.
5009 try block.incoming_blocks.ensureUnusedCapacity(gpa, 4);
5010
5011 try cf.blocks.putNoClobber(gpa, inst, &block);
5012 defer assert(cf.blocks.remove(inst));
5013
5014 try cg.genBody(body);
5015
5016 // Only begin a new block if there were actually any breaks towards it.
5017 if (block.label) |label| {
5018 try cg.beginSpvBlock(label);
5019 }
5020
5021 if (!have_block_result)
5022 return null;
5023
5024 assert(block.label != null);
5025 const result_id = cg.module.allocId();
5026 const result_type_id = try cg.resolveType(ty, .direct);
5027
5028 try cg.body.emitRaw(
5029 gpa,
5030 .OpPhi,
5031 // result type + result + variable/parent...
5032 2 + @as(u16, @intCast(block.incoming_blocks.items.len * 2)),
5033 );
5034 cg.body.writeOperand(Id, result_type_id);
5035 cg.body.writeOperand(Id, result_id);
5036
5037 for (block.incoming_blocks.items) |incoming| {
5038 cg.body.writeOperand(
5039 spec.PairIdRefIdRef,
5040 .{ incoming.break_value_id, incoming.src_label },
5041 );
5042 }
5043
5044 return result_id;
5045 },
5046 };
5047
5048 const maybe_block_result_var_id = if (have_block_result) blk: {
5049 const ty_id = try cg.resolveType(ty, .indirect);
5050 const block_result_var_id = try cg.alloc(ty_id, null);
5051 try cf.block_results.putNoClobber(gpa, inst, block_result_var_id);
5052 break :blk block_result_var_id;
5053 } else null;
5054 defer if (have_block_result) assert(cf.block_results.remove(inst));
5055
5056 const next_block = try cg.genStructuredBody(.selection, body);
5057
5058 // When encountering a block instruction, we are always at least in the function's scope,
5059 // so there always has to be another entry.
5060 assert(cf.block_stack.items.len > 0);
5061
5062 // Check if the target of the branch was this current block.
5063 const this_block = try cg.constInt(.u32, @intFromEnum(inst));
5064 const jump_to_this_block_id = cg.module.allocId();
5065 const bool_ty_id = try cg.resolveType(.bool, .direct);
5066 try cg.body.emit(gpa, .OpIEqual, .{
5067 .id_result_type = bool_ty_id,
5068 .id_result = jump_to_this_block_id,
5069 .operand_1 = next_block,
5070 .operand_2 = this_block,
5071 });
5072
5073 const sblock = cf.block_stack.getLast();
5074
5075 if (ty.isNoReturn(zcu)) {
5076 // If this block is noreturn, this instruction is the last of a block,
5077 // and we must simply jump to the block's merge unconditionally.
5078 try cg.structuredBreak(next_block);
5079 } else {
5080 switch (sblock.*) {
5081 .selection => |*merge| {
5082 // To jump out of a selection block, push a new entry onto its merge stack and
5083 // generate a conditional branch to there and to the instructions following this block.
5084 const merge_label = cg.module.allocId();
5085 const then_label = cg.module.allocId();
5086 try cg.body.emit(gpa, .OpSelectionMerge, .{
5087 .merge_block = merge_label,
5088 .selection_control = .{},
5089 });
5090 try cg.body.emit(gpa, .OpBranchConditional, .{
5091 .condition = jump_to_this_block_id,
5092 .true_label = then_label,
5093 .false_label = merge_label,
5094 });
5095 try merge.merge_stack.append(gpa, .{
5096 .incoming = .{
5097 .src_label = cg.block_label,
5098 .next_block = next_block,
5099 },
5100 .merge_block = merge_label,
5101 });
5102
5103 try cg.beginSpvBlock(then_label);
5104 },
5105 .loop => |*merge| {
5106 // To jump out of a loop block, generate a conditional that exits the block
5107 // to the loop merge if the target ID is not the one of this block.
5108 const continue_label = cg.module.allocId();
5109 try cg.body.emit(gpa, .OpBranchConditional, .{
5110 .condition = jump_to_this_block_id,
5111 .true_label = continue_label,
5112 .false_label = merge.merge_block,
5113 });
5114 try merge.merges.append(gpa, .{
5115 .src_label = cg.block_label,
5116 .next_block = next_block,
5117 });
5118 try cg.beginSpvBlock(continue_label);
5119 },
5120 }
5121 }
5122
5123 if (maybe_block_result_var_id) |block_result_var_id| {
5124 return try cg.load(ty, block_result_var_id, .{});
5125 }
5126
5127 return null;
5128}
5129
5130fn airBr(cg: *CodeGen, inst: Air.Inst.Index) !void {
5131 const gpa = cg.module.gpa;
5132 const zcu = cg.module.zcu;
5133 const br = cg.air.instructions.items(.data)[@intFromEnum(inst)].br;
5134 const operand_ty = cg.typeOf(br.operand);
5135
5136 switch (cg.control_flow) {
5137 .structured => |*cf| {
5138 if (operand_ty.isFnOrHasRuntimeBitsIgnoreComptime(zcu)) {
5139 const operand_id = try cg.resolve(br.operand);
5140 const block_result_var_id = cf.block_results.get(br.block_inst).?;
5141 try cg.store(operand_ty, block_result_var_id, operand_id, .{});
5142 }
5143
5144 const next_block = try cg.constInt(.u32, @intFromEnum(br.block_inst));
5145 try cg.structuredBreak(next_block);
5146 },
5147 .unstructured => |cf| {
5148 const block = cf.blocks.get(br.block_inst).?;
5149 if (operand_ty.isFnOrHasRuntimeBitsIgnoreComptime(zcu)) {
5150 const operand_id = try cg.resolve(br.operand);
5151 // block_label should not be undefined here, lest there
5152 // is a br or br_void in the function's body.
5153 try block.incoming_blocks.append(gpa, .{
5154 .src_label = cg.block_label,
5155 .break_value_id = operand_id,
5156 });
5157 }
5158
5159 if (block.label == null) {
5160 block.label = cg.module.allocId();
5161 }
5162
5163 try cg.body.emit(gpa, .OpBranch, .{ .target_label = block.label.? });
5164 },
5165 }
5166}
5167
5168fn airCondBr(cg: *CodeGen, inst: Air.Inst.Index) !void {
5169 const gpa = cg.module.gpa;
5170 const pl_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].pl_op;
5171 const cond_br = cg.air.extraData(Air.CondBr, pl_op.payload);
5172 const then_body: []const Air.Inst.Index = @ptrCast(cg.air.extra.items[cond_br.end..][0..cond_br.data.then_body_len]);
5173 const else_body: []const Air.Inst.Index = @ptrCast(cg.air.extra.items[cond_br.end + then_body.len ..][0..cond_br.data.else_body_len]);
5174 const condition_id = try cg.resolve(pl_op.operand);
5175
5176 const then_label = cg.module.allocId();
5177 const else_label = cg.module.allocId();
5178
5179 switch (cg.control_flow) {
5180 .structured => {
5181 const merge_label = cg.module.allocId();
5182
5183 try cg.body.emit(gpa, .OpSelectionMerge, .{
5184 .merge_block = merge_label,
5185 .selection_control = .{},
5186 });
5187 try cg.body.emit(gpa, .OpBranchConditional, .{
5188 .condition = condition_id,
5189 .true_label = then_label,
5190 .false_label = else_label,
5191 });
5192
5193 try cg.beginSpvBlock(then_label);
5194 const then_next = try cg.genStructuredBody(.selection, then_body);
5195 const then_incoming: ControlFlow.Structured.Block.Incoming = .{
5196 .src_label = cg.block_label,
5197 .next_block = then_next,
5198 };
5199
5200 try cg.body.emit(gpa, .OpBranch, .{ .target_label = merge_label });
5201
5202 try cg.beginSpvBlock(else_label);
5203 const else_next = try cg.genStructuredBody(.selection, else_body);
5204 const else_incoming: ControlFlow.Structured.Block.Incoming = .{
5205 .src_label = cg.block_label,
5206 .next_block = else_next,
5207 };
5208
5209 try cg.body.emit(gpa, .OpBranch, .{ .target_label = merge_label });
5210
5211 try cg.beginSpvBlock(merge_label);
5212 const next_block = try cg.structuredNextBlock(&.{ then_incoming, else_incoming });
5213
5214 try cg.structuredBreak(next_block);
5215 },
5216 .unstructured => {
5217 try cg.body.emit(gpa, .OpBranchConditional, .{
5218 .condition = condition_id,
5219 .true_label = then_label,
5220 .false_label = else_label,
5221 });
5222
5223 try cg.beginSpvBlock(then_label);
5224 try cg.genBody(then_body);
5225 try cg.beginSpvBlock(else_label);
5226 try cg.genBody(else_body);
5227 },
5228 }
5229}
5230
5231fn airLoop(cg: *CodeGen, inst: Air.Inst.Index) !void {
5232 const gpa = cg.module.gpa;
5233 const ty_pl = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
5234 const loop = cg.air.extraData(Air.Block, ty_pl.payload);
5235 const body: []const Air.Inst.Index = @ptrCast(cg.air.extra.items[loop.end..][0..loop.data.body_len]);
5236
5237 const body_label = cg.module.allocId();
5238
5239 switch (cg.control_flow) {
5240 .structured => {
5241 const header_label = cg.module.allocId();
5242 const merge_label = cg.module.allocId();
5243 const continue_label = cg.module.allocId();
5244
5245 // The back-edge must point to the loop header, so generate a separate block for the
5246 // loop header so that we don't accidentally include some instructions from there
5247 // in the loop.
5248
5249 try cg.body.emit(gpa, .OpBranch, .{ .target_label = header_label });
5250 try cg.beginSpvBlock(header_label);
5251
5252 // Emit loop header and jump to loop body
5253 try cg.body.emit(gpa, .OpLoopMerge, .{
5254 .merge_block = merge_label,
5255 .continue_target = continue_label,
5256 .loop_control = .{},
5257 });
5258
5259 try cg.body.emit(gpa, .OpBranch, .{ .target_label = body_label });
5260
5261 try cg.beginSpvBlock(body_label);
5262
5263 const next_block = try cg.genStructuredBody(.{ .loop = .{
5264 .merge_label = merge_label,
5265 .continue_label = continue_label,
5266 } }, body);
5267 try cg.structuredBreak(next_block);
5268
5269 try cg.beginSpvBlock(continue_label);
5270
5271 try cg.body.emit(gpa, .OpBranch, .{ .target_label = header_label });
5272 },
5273 .unstructured => {
5274 try cg.body.emit(gpa, .OpBranch, .{ .target_label = body_label });
5275 try cg.beginSpvBlock(body_label);
5276 try cg.genBody(body);
5277
5278 try cg.body.emit(gpa, .OpBranch, .{ .target_label = body_label });
5279 },
5280 }
5281}
5282
5283fn airLoad(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
5284 const zcu = cg.module.zcu;
5285 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
5286 const ptr_ty = cg.typeOf(ty_op.operand);
5287 const elem_ty = cg.typeOfIndex(inst);
5288 const operand = try cg.resolve(ty_op.operand);
5289 if (!ptr_ty.isVolatilePtr(zcu) and cg.liveness.isUnused(inst)) return null;
5290
5291 return try cg.load(elem_ty, operand, .{ .is_volatile = ptr_ty.isVolatilePtr(zcu) });
5292}
5293
5294fn airStore(cg: *CodeGen, inst: Air.Inst.Index) !void {
5295 const zcu = cg.module.zcu;
5296 const bin_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].bin_op;
5297 const ptr_ty = cg.typeOf(bin_op.lhs);
5298 const elem_ty = ptr_ty.childType(zcu);
5299 const ptr = try cg.resolve(bin_op.lhs);
5300 const value = try cg.resolve(bin_op.rhs);
5301
5302 try cg.store(elem_ty, ptr, value, .{ .is_volatile = ptr_ty.isVolatilePtr(zcu) });
5303}
5304
5305fn airRet(cg: *CodeGen, inst: Air.Inst.Index) !void {
5306 const gpa = cg.module.gpa;
5307 const zcu = cg.module.zcu;
5308 const operand = cg.air.instructions.items(.data)[@intFromEnum(inst)].un_op;
5309 const ret_ty = cg.typeOf(operand);
5310 if (!ret_ty.hasRuntimeBitsIgnoreComptime(zcu)) {
5311 const fn_info = zcu.typeToFunc(zcu.navValue(cg.owner_nav).typeOf(zcu)).?;
5312 if (Type.fromInterned(fn_info.return_type).isError(zcu)) {
5313 // Functions with an empty error set are emitted with an error code
5314 // return type and return zero so they can be function pointers coerced
5315 // to functions that return anyerror.
5316 const no_err_id = try cg.constInt(.anyerror, 0);
5317 return try cg.body.emit(gpa, .OpReturnValue, .{ .value = no_err_id });
5318 } else {
5319 return try cg.body.emit(gpa, .OpReturn, {});
5320 }
5321 }
5322
5323 const operand_id = try cg.resolve(operand);
5324 try cg.body.emit(gpa, .OpReturnValue, .{ .value = operand_id });
5325}
5326
5327fn airRetLoad(cg: *CodeGen, inst: Air.Inst.Index) !void {
5328 const gpa = cg.module.gpa;
5329 const zcu = cg.module.zcu;
5330 const un_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].un_op;
5331 const ptr_ty = cg.typeOf(un_op);
5332 const ret_ty = ptr_ty.childType(zcu);
5333
5334 if (!ret_ty.hasRuntimeBitsIgnoreComptime(zcu)) {
5335 const fn_info = zcu.typeToFunc(zcu.navValue(cg.owner_nav).typeOf(zcu)).?;
5336 if (Type.fromInterned(fn_info.return_type).isError(zcu)) {
5337 // Functions with an empty error set are emitted with an error code
5338 // return type and return zero so they can be function pointers coerced
5339 // to functions that return anyerror.
5340 const no_err_id = try cg.constInt(.anyerror, 0);
5341 return try cg.body.emit(gpa, .OpReturnValue, .{ .value = no_err_id });
5342 } else {
5343 return try cg.body.emit(gpa, .OpReturn, {});
5344 }
5345 }
5346
5347 const ptr = try cg.resolve(un_op);
5348 const value = try cg.load(ret_ty, ptr, .{ .is_volatile = ptr_ty.isVolatilePtr(zcu) });
5349 try cg.body.emit(gpa, .OpReturnValue, .{
5350 .value = value,
5351 });
5352}
5353
5354fn airTry(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
5355 const gpa = cg.module.gpa;
5356 const zcu = cg.module.zcu;
5357 const pl_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].pl_op;
5358 const err_union_id = try cg.resolve(pl_op.operand);
5359 const extra = cg.air.extraData(Air.Try, pl_op.payload);
5360 const body: []const Air.Inst.Index = @ptrCast(cg.air.extra.items[extra.end..][0..extra.data.body_len]);
5361
5362 const err_union_ty = cg.typeOf(pl_op.operand);
5363 const payload_ty = cg.typeOfIndex(inst);
5364
5365 const bool_ty_id = try cg.resolveType(.bool, .direct);
5366
5367 const eu_layout = cg.errorUnionLayout(payload_ty);
5368
5369 if (!err_union_ty.errorUnionSet(zcu).errorSetIsEmpty(zcu)) {
5370 const err_id = if (eu_layout.payload_has_bits)
5371 try cg.extractField(.anyerror, err_union_id, eu_layout.errorFieldIndex())
5372 else
5373 err_union_id;
5374
5375 const zero_id = try cg.constInt(.anyerror, 0);
5376 const is_err_id = cg.module.allocId();
5377 try cg.body.emit(gpa, .OpINotEqual, .{
5378 .id_result_type = bool_ty_id,
5379 .id_result = is_err_id,
5380 .operand_1 = err_id,
5381 .operand_2 = zero_id,
5382 });
5383
5384 // When there is an error, we must evaluate `body`. Otherwise we must continue
5385 // with the current body.
5386 // Just generate a new block here, then generate a new block inline for the remainder of the body.
5387
5388 const err_block = cg.module.allocId();
5389 const ok_block = cg.module.allocId();
5390
5391 switch (cg.control_flow) {
5392 .structured => {
5393 // According to AIR documentation, this block is guaranteed
5394 // to not break and end in a return instruction. Thus,
5395 // for structured control flow, we can just naively use
5396 // the ok block as the merge block here.
5397 try cg.body.emit(gpa, .OpSelectionMerge, .{
5398 .merge_block = ok_block,
5399 .selection_control = .{},
5400 });
5401 },
5402 .unstructured => {},
5403 }
5404
5405 try cg.body.emit(gpa, .OpBranchConditional, .{
5406 .condition = is_err_id,
5407 .true_label = err_block,
5408 .false_label = ok_block,
5409 });
5410
5411 try cg.beginSpvBlock(err_block);
5412 try cg.genBody(body);
5413
5414 try cg.beginSpvBlock(ok_block);
5415 }
5416
5417 if (!eu_layout.payload_has_bits) {
5418 return null;
5419 }
5420
5421 // Now just extract the payload, if required.
5422 return try cg.extractField(payload_ty, err_union_id, eu_layout.payloadFieldIndex());
5423}
5424
5425fn airErrUnionErr(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
5426 const zcu = cg.module.zcu;
5427 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
5428 const operand_id = try cg.resolve(ty_op.operand);
5429 const err_union_ty = cg.typeOf(ty_op.operand);
5430 const err_ty_id = try cg.resolveType(.anyerror, .direct);
5431
5432 if (err_union_ty.errorUnionSet(zcu).errorSetIsEmpty(zcu)) {
5433 // No error possible, so just return undefined.
5434 return try cg.module.constUndef(err_ty_id);
5435 }
5436
5437 const payload_ty = err_union_ty.errorUnionPayload(zcu);
5438 const eu_layout = cg.errorUnionLayout(payload_ty);
5439
5440 if (!eu_layout.payload_has_bits) {
5441 // If no payload, error union is represented by error set.
5442 return operand_id;
5443 }
5444
5445 return try cg.extractField(.anyerror, operand_id, eu_layout.errorFieldIndex());
5446}
5447
5448fn airErrUnionPayload(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
5449 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
5450 const operand_id = try cg.resolve(ty_op.operand);
5451 const payload_ty = cg.typeOfIndex(inst);
5452 const eu_layout = cg.errorUnionLayout(payload_ty);
5453
5454 if (!eu_layout.payload_has_bits) {
5455 return null; // No error possible.
5456 }
5457
5458 return try cg.extractField(payload_ty, operand_id, eu_layout.payloadFieldIndex());
5459}
5460
5461fn airWrapErrUnionErr(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
5462 const zcu = cg.module.zcu;
5463 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
5464 const err_union_ty = cg.typeOfIndex(inst);
5465 const payload_ty = err_union_ty.errorUnionPayload(zcu);
5466 const operand_id = try cg.resolve(ty_op.operand);
5467 const eu_layout = cg.errorUnionLayout(payload_ty);
5468
5469 if (!eu_layout.payload_has_bits) {
5470 return operand_id;
5471 }
5472
5473 const payload_ty_id = try cg.resolveType(payload_ty, .indirect);
5474
5475 var members: [2]Id = undefined;
5476 members[eu_layout.errorFieldIndex()] = operand_id;
5477 members[eu_layout.payloadFieldIndex()] = try cg.module.constUndef(payload_ty_id);
5478
5479 var types: [2]Type = undefined;
5480 types[eu_layout.errorFieldIndex()] = .anyerror;
5481 types[eu_layout.payloadFieldIndex()] = payload_ty;
5482
5483 const err_union_ty_id = try cg.resolveType(err_union_ty, .direct);
5484 return try cg.constructComposite(err_union_ty_id, &members);
5485}
5486
5487fn airWrapErrUnionPayload(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
5488 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
5489 const err_union_ty = cg.typeOfIndex(inst);
5490 const operand_id = try cg.resolve(ty_op.operand);
5491 const payload_ty = cg.typeOf(ty_op.operand);
5492 const eu_layout = cg.errorUnionLayout(payload_ty);
5493
5494 if (!eu_layout.payload_has_bits) {
5495 return try cg.constInt(.anyerror, 0);
5496 }
5497
5498 var members: [2]Id = undefined;
5499 members[eu_layout.errorFieldIndex()] = try cg.constInt(.anyerror, 0);
5500 members[eu_layout.payloadFieldIndex()] = try cg.convertToIndirect(payload_ty, operand_id);
5501
5502 var types: [2]Type = undefined;
5503 types[eu_layout.errorFieldIndex()] = .anyerror;
5504 types[eu_layout.payloadFieldIndex()] = payload_ty;
5505
5506 const err_union_ty_id = try cg.resolveType(err_union_ty, .direct);
5507 return try cg.constructComposite(err_union_ty_id, &members);
5508}
5509
5510fn airIsNull(cg: *CodeGen, inst: Air.Inst.Index, is_pointer: bool, pred: enum { is_null, is_non_null }) !?Id {
5511 const zcu = cg.module.zcu;
5512 const un_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].un_op;
5513 const operand_id = try cg.resolve(un_op);
5514 const operand_ty = cg.typeOf(un_op);
5515 const optional_ty = if (is_pointer) operand_ty.childType(zcu) else operand_ty;
5516 const payload_ty = optional_ty.optionalChild(zcu);
5517
5518 const bool_ty_id = try cg.resolveType(.bool, .direct);
5519
5520 if (optional_ty.optionalReprIsPayload(zcu)) {
5521 // Pointer payload represents nullability: pointer or slice.
5522 const loaded_id = if (is_pointer)
5523 try cg.load(optional_ty, operand_id, .{})
5524 else
5525 operand_id;
5526
5527 const ptr_ty = if (payload_ty.isSlice(zcu))
5528 payload_ty.slicePtrFieldType(zcu)
5529 else
5530 payload_ty;
5531
5532 const ptr_id = if (payload_ty.isSlice(zcu))
5533 try cg.extractField(ptr_ty, loaded_id, 0)
5534 else
5535 loaded_id;
5536
5537 const ptr_ty_id = try cg.resolveType(ptr_ty, .direct);
5538 const null_id = try cg.module.constNull(ptr_ty_id);
5539 const null_tmp: Temporary = .init(ptr_ty, null_id);
5540 const ptr: Temporary = .init(ptr_ty, ptr_id);
5541
5542 const op: std.math.CompareOperator = switch (pred) {
5543 .is_null => .eq,
5544 .is_non_null => .neq,
5545 };
5546 const result = try cg.cmp(op, ptr, null_tmp);
5547 return try result.materialize(cg);
5548 }
5549
5550 const is_non_null_id = blk: {
5551 if (is_pointer) {
5552 if (payload_ty.hasRuntimeBitsIgnoreComptime(zcu)) {
5553 const storage_class = cg.module.storageClass(operand_ty.ptrAddressSpace(zcu));
5554 const bool_indirect_ty_id = try cg.resolveType(.bool, .indirect);
5555 const bool_ptr_ty_id = try cg.module.ptrType(bool_indirect_ty_id, storage_class);
5556 const tag_ptr_id = try cg.accessChain(bool_ptr_ty_id, operand_id, &.{1});
5557 break :blk try cg.load(.bool, tag_ptr_id, .{});
5558 }
5559
5560 break :blk try cg.load(.bool, operand_id, .{});
5561 }
5562
5563 break :blk if (payload_ty.hasRuntimeBitsIgnoreComptime(zcu))
5564 try cg.extractField(.bool, operand_id, 1)
5565 else
5566 // Optional representation is bool indicating whether the optional is set
5567 // Optionals with no payload are represented as an (indirect) bool, so convert
5568 // it back to the direct bool here.
5569 try cg.convertToDirect(.bool, operand_id);
5570 };
5571
5572 return switch (pred) {
5573 .is_null => blk: {
5574 // Invert condition
5575 const result_id = cg.module.allocId();
5576 try cg.body.emit(cg.module.gpa, .OpLogicalNot, .{
5577 .id_result_type = bool_ty_id,
5578 .id_result = result_id,
5579 .operand = is_non_null_id,
5580 });
5581 break :blk result_id;
5582 },
5583 .is_non_null => is_non_null_id,
5584 };
5585}
5586
5587fn airIsErr(cg: *CodeGen, inst: Air.Inst.Index, pred: enum { is_err, is_non_err }) !?Id {
5588 const zcu = cg.module.zcu;
5589 const un_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].un_op;
5590 const operand_id = try cg.resolve(un_op);
5591 const err_union_ty = cg.typeOf(un_op);
5592
5593 if (err_union_ty.errorUnionSet(zcu).errorSetIsEmpty(zcu)) {
5594 return try cg.constBool(pred == .is_non_err, .direct);
5595 }
5596
5597 const payload_ty = err_union_ty.errorUnionPayload(zcu);
5598 const eu_layout = cg.errorUnionLayout(payload_ty);
5599 const bool_ty_id = try cg.resolveType(.bool, .direct);
5600
5601 const error_id = if (!eu_layout.payload_has_bits)
5602 operand_id
5603 else
5604 try cg.extractField(.anyerror, operand_id, eu_layout.errorFieldIndex());
5605
5606 const result_id = cg.module.allocId();
5607 switch (pred) {
5608 inline else => |pred_ct| try cg.body.emit(
5609 cg.module.gpa,
5610 switch (pred_ct) {
5611 .is_err => .OpINotEqual,
5612 .is_non_err => .OpIEqual,
5613 },
5614 .{
5615 .id_result_type = bool_ty_id,
5616 .id_result = result_id,
5617 .operand_1 = error_id,
5618 .operand_2 = try cg.constInt(.anyerror, 0),
5619 },
5620 ),
5621 }
5622 return result_id;
5623}
5624
5625fn airUnwrapOptional(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
5626 const zcu = cg.module.zcu;
5627 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
5628 const operand_id = try cg.resolve(ty_op.operand);
5629 const optional_ty = cg.typeOf(ty_op.operand);
5630 const payload_ty = cg.typeOfIndex(inst);
5631
5632 if (!payload_ty.hasRuntimeBitsIgnoreComptime(zcu)) return null;
5633
5634 if (optional_ty.optionalReprIsPayload(zcu)) {
5635 return operand_id;
5636 }
5637
5638 return try cg.extractField(payload_ty, operand_id, 0);
5639}
5640
5641fn airUnwrapOptionalPtr(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
5642 const zcu = cg.module.zcu;
5643 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
5644 const operand_id = try cg.resolve(ty_op.operand);
5645 const operand_ty = cg.typeOf(ty_op.operand);
5646 const optional_ty = operand_ty.childType(zcu);
5647 const payload_ty = optional_ty.optionalChild(zcu);
5648 const result_ty = cg.typeOfIndex(inst);
5649 const result_ty_id = try cg.resolveType(result_ty, .direct);
5650
5651 if (!payload_ty.hasRuntimeBitsIgnoreComptime(zcu)) {
5652 // There is no payload, but we still need to return a valid pointer.
5653 // We can just return anything here, so just return a pointer to the operand.
5654 return try cg.bitCast(result_ty, operand_ty, operand_id);
5655 }
5656
5657 if (optional_ty.optionalReprIsPayload(zcu)) {
5658 // They are the same value.
5659 return try cg.bitCast(result_ty, operand_ty, operand_id);
5660 }
5661
5662 return try cg.accessChain(result_ty_id, operand_id, &.{0});
5663}
5664
5665fn airWrapOptional(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
5666 const zcu = cg.module.zcu;
5667 const ty_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_op;
5668 const payload_ty = cg.typeOf(ty_op.operand);
5669
5670 if (!payload_ty.hasRuntimeBitsIgnoreComptime(zcu)) {
5671 return try cg.constBool(true, .indirect);
5672 }
5673
5674 const operand_id = try cg.resolve(ty_op.operand);
5675
5676 const optional_ty = cg.typeOfIndex(inst);
5677 if (optional_ty.optionalReprIsPayload(zcu)) {
5678 return operand_id;
5679 }
5680
5681 const payload_id = try cg.convertToIndirect(payload_ty, operand_id);
5682 const members = [_]Id{ payload_id, try cg.constBool(true, .indirect) };
5683 const optional_ty_id = try cg.resolveType(optional_ty, .direct);
5684 return try cg.constructComposite(optional_ty_id, &members);
5685}
5686
5687fn airSwitchBr(cg: *CodeGen, inst: Air.Inst.Index) !void {
5688 const gpa = cg.module.gpa;
5689 const pt = cg.pt;
5690 const zcu = cg.module.zcu;
5691 const target = cg.module.zcu.getTarget();
5692 const switch_br = cg.air.unwrapSwitch(inst);
5693 const cond_ty = cg.typeOf(switch_br.operand);
5694 const cond = try cg.resolve(switch_br.operand);
5695 var cond_indirect = try cg.convertToIndirect(cond_ty, cond);
5696
5697 const cond_words: u32 = switch (cond_ty.zigTypeTag(zcu)) {
5698 .bool, .error_set => 1,
5699 .int => blk: {
5700 const bits = cond_ty.intInfo(zcu).bits;
5701 const backing_bits, const big_int = cg.module.backingIntBits(bits);
5702 if (big_int) return cg.todo("implement composite int switch", .{});
5703 break :blk if (backing_bits <= 32) 1 else 2;
5704 },
5705 .@"enum" => blk: {
5706 const int_ty = cond_ty.intTagType(zcu);
5707 const int_info = int_ty.intInfo(zcu);
5708 const backing_bits, const big_int = cg.module.backingIntBits(int_info.bits);
5709 if (big_int) return cg.todo("implement composite int switch", .{});
5710 break :blk if (backing_bits <= 32) 1 else 2;
5711 },
5712 .pointer => blk: {
5713 cond_indirect = try cg.intFromPtr(cond_indirect);
5714 break :blk target.ptrBitWidth() / 32;
5715 },
5716 // TODO: Figure out which types apply here, and work around them as we can only do integers.
5717 else => return cg.todo("implement switch for type {s}", .{@tagName(cond_ty.zigTypeTag(zcu))}),
5718 };
5719
5720 const num_cases = switch_br.cases_len;
5721
5722 // Compute the total number of arms that we need.
5723 // Zig switches are grouped by condition, so we need to loop through all of them
5724 const num_conditions = blk: {
5725 var num_conditions: u32 = 0;
5726 var it = switch_br.iterateCases();
5727 while (it.next()) |case| {
5728 if (case.ranges.len > 0) return cg.todo("switch with ranges", .{});
5729 num_conditions += @intCast(case.items.len);
5730 }
5731 break :blk num_conditions;
5732 };
5733
5734 // First, pre-allocate the labels for the cases.
5735 const case_labels = cg.module.allocIds(num_cases);
5736 // We always need the default case - if zig has none, we will generate unreachable there.
5737 const default = cg.module.allocId();
5738
5739 const merge_label = switch (cg.control_flow) {
5740 .structured => cg.module.allocId(),
5741 .unstructured => null,
5742 };
5743
5744 if (cg.control_flow == .structured) {
5745 try cg.body.emit(gpa, .OpSelectionMerge, .{
5746 .merge_block = merge_label.?,
5747 .selection_control = .{},
5748 });
5749 }
5750
5751 // Emit the instruction before generating the blocks.
5752 try cg.body.emitRaw(gpa, .OpSwitch, 2 + (cond_words + 1) * num_conditions);
5753 cg.body.writeOperand(Id, cond_indirect);
5754 cg.body.writeOperand(Id, default);
5755
5756 // Emit each of the cases
5757 {
5758 var it = switch_br.iterateCases();
5759 while (it.next()) |case| {
5760 // SPIR-V needs a literal here, which' width depends on the case condition.
5761 const label = case_labels.at(case.idx);
5762
5763 for (case.items) |item| {
5764 const value = (try cg.air.value(item, pt)) orelse unreachable;
5765 const int_val: u64 = switch (cond_ty.zigTypeTag(zcu)) {
5766 .bool, .int => if (cond_ty.isSignedInt(zcu)) @bitCast(value.toSignedInt(zcu)) else value.toUnsignedInt(zcu),
5767 .@"enum" => blk: {
5768 // TODO: figure out of cond_ty is correct (something with enum literals)
5769 break :blk (try value.intFromEnum(cond_ty, pt)).toUnsignedInt(zcu); // TODO: composite integer constants
5770 },
5771 .error_set => value.getErrorInt(zcu),
5772 .pointer => value.toUnsignedInt(zcu),
5773 else => unreachable,
5774 };
5775 const int_lit: spec.LiteralContextDependentNumber = switch (cond_words) {
5776 1 => .{ .uint32 = @intCast(int_val) },
5777 2 => .{ .uint64 = int_val },
5778 else => unreachable,
5779 };
5780 cg.body.writeOperand(spec.LiteralContextDependentNumber, int_lit);
5781 cg.body.writeOperand(Id, label);
5782 }
5783 }
5784 }
5785
5786 var incoming_structured_blocks: std.ArrayList(ControlFlow.Structured.Block.Incoming) = .empty;
5787 defer incoming_structured_blocks.deinit(gpa);
5788
5789 if (cg.control_flow == .structured) {
5790 try incoming_structured_blocks.ensureUnusedCapacity(gpa, num_cases + 1);
5791 }
5792
5793 // Now, finally, we can start emitting each of the cases.
5794 var it = switch_br.iterateCases();
5795 while (it.next()) |case| {
5796 const label = case_labels.at(case.idx);
5797
5798 try cg.beginSpvBlock(label);
5799
5800 switch (cg.control_flow) {
5801 .structured => {
5802 const next_block = try cg.genStructuredBody(.selection, case.body);
5803 incoming_structured_blocks.appendAssumeCapacity(.{
5804 .src_label = cg.block_label,
5805 .next_block = next_block,
5806 });
5807
5808 try cg.body.emit(gpa, .OpBranch, .{ .target_label = merge_label.? });
5809 },
5810 .unstructured => {
5811 try cg.genBody(case.body);
5812 },
5813 }
5814 }
5815
5816 const else_body = it.elseBody();
5817 try cg.beginSpvBlock(default);
5818 if (else_body.len != 0) {
5819 switch (cg.control_flow) {
5820 .structured => {
5821 const next_block = try cg.genStructuredBody(.selection, else_body);
5822 incoming_structured_blocks.appendAssumeCapacity(.{
5823 .src_label = cg.block_label,
5824 .next_block = next_block,
5825 });
5826
5827 try cg.body.emit(gpa, .OpBranch, .{ .target_label = merge_label.? });
5828 },
5829 .unstructured => {
5830 try cg.genBody(else_body);
5831 },
5832 }
5833 } else {
5834 try cg.body.emit(gpa, .OpUnreachable, {});
5835 }
5836
5837 if (cg.control_flow == .structured) {
5838 try cg.beginSpvBlock(merge_label.?);
5839 const next_block = try cg.structuredNextBlock(incoming_structured_blocks.items);
5840 try cg.structuredBreak(next_block);
5841 }
5842}
5843
5844fn airUnreach(cg: *CodeGen) !void {
5845 try cg.body.emit(cg.module.gpa, .OpUnreachable, {});
5846}
5847
5848fn airDbgStmt(cg: *CodeGen, inst: Air.Inst.Index) !void {
5849 const zcu = cg.module.zcu;
5850 const dbg_stmt = cg.air.instructions.items(.data)[@intFromEnum(inst)].dbg_stmt;
5851 const path = zcu.navFileScope(cg.owner_nav).sub_file_path;
5852
5853 if (zcu.comp.config.root_strip) return;
5854
5855 try cg.body.emit(cg.module.gpa, .OpLine, .{
5856 .file = try cg.module.debugString(path),
5857 .line = cg.base_line + dbg_stmt.line + 1,
5858 .column = dbg_stmt.column + 1,
5859 });
5860}
5861
5862fn airDbgInlineBlock(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
5863 const zcu = cg.module.zcu;
5864 const inst_datas = cg.air.instructions.items(.data);
5865 const extra = cg.air.extraData(Air.DbgInlineBlock, inst_datas[@intFromEnum(inst)].ty_pl.payload);
5866 const old_base_line = cg.base_line;
5867 defer cg.base_line = old_base_line;
5868 cg.base_line = zcu.navSrcLine(zcu.funcInfo(extra.data.func).owner_nav);
5869 return cg.lowerBlock(inst, @ptrCast(cg.air.extra.items[extra.end..][0..extra.data.body_len]));
5870}
5871
5872fn airDbgVar(cg: *CodeGen, inst: Air.Inst.Index) !void {
5873 const pl_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].pl_op;
5874 const target_id = try cg.resolve(pl_op.operand);
5875 const name: Air.NullTerminatedString = @enumFromInt(pl_op.payload);
5876 try cg.module.debugName(target_id, name.toSlice(cg.air));
5877}
5878
5879fn airAssembly(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
5880 const gpa = cg.module.gpa;
5881 const zcu = cg.module.zcu;
5882 const ty_pl = cg.air.instructions.items(.data)[@intFromEnum(inst)].ty_pl;
5883 const extra = cg.air.extraData(Air.Asm, ty_pl.payload);
5884
5885 const is_volatile = extra.data.flags.is_volatile;
5886 const outputs_len = extra.data.flags.outputs_len;
5887
5888 if (!is_volatile and cg.liveness.isUnused(inst)) return null;
5889
5890 var extra_i: usize = extra.end;
5891 const outputs: []const Air.Inst.Ref = @ptrCast(cg.air.extra.items[extra_i..][0..outputs_len]);
5892 extra_i += outputs.len;
5893 const inputs: []const Air.Inst.Ref = @ptrCast(cg.air.extra.items[extra_i..][0..extra.data.inputs_len]);
5894 extra_i += inputs.len;
5895
5896 if (outputs.len > 1) {
5897 return cg.todo("implement inline asm with more than 1 output", .{});
5898 }
5899
5900 var ass: Assembler = .{ .cg = cg };
5901 defer ass.deinit();
5902
5903 var output_extra_i = extra_i;
5904 for (outputs) |output| {
5905 if (output != .none) {
5906 return cg.todo("implement inline asm with non-returned output", .{});
5907 }
5908 const extra_bytes = std.mem.sliceAsBytes(cg.air.extra.items[extra_i..]);
5909 const constraint = std.mem.sliceTo(std.mem.sliceAsBytes(cg.air.extra.items[extra_i..]), 0);
5910 const name = std.mem.sliceTo(extra_bytes[constraint.len + 1 ..], 0);
5911 extra_i += (constraint.len + name.len + (2 + 3)) / 4;
5912 // TODO: Record output and use it somewhere.
5913 }
5914
5915 for (inputs) |input| {
5916 const extra_bytes = std.mem.sliceAsBytes(cg.air.extra.items[extra_i..]);
5917 const constraint = std.mem.sliceTo(extra_bytes, 0);
5918 const name = std.mem.sliceTo(extra_bytes[constraint.len + 1 ..], 0);
5919 // This equation accounts for the fact that even if we have exactly 4 bytes
5920 // for the string, we still use the next u32 for the null terminator.
5921 extra_i += (constraint.len + name.len + (2 + 3)) / 4;
5922
5923 const input_ty = cg.typeOf(input);
5924
5925 if (std.mem.eql(u8, constraint, "c")) {
5926 // constant
5927 const val = (try cg.air.value(input, cg.pt)) orelse {
5928 return cg.fail("assembly inputs with 'c' constraint have to be compile-time known", .{});
5929 };
5930
5931 // TODO: This entire function should be handled a bit better...
5932 const ip = &zcu.intern_pool;
5933 switch (ip.indexToKey(val.toIntern())) {
5934 .int_type,
5935 .ptr_type,
5936 .array_type,
5937 .vector_type,
5938 .opt_type,
5939 .anyframe_type,
5940 .error_union_type,
5941 .simple_type,
5942 .struct_type,
5943 .union_type,
5944 .opaque_type,
5945 .enum_type,
5946 .func_type,
5947 .error_set_type,
5948 .inferred_error_set_type,
5949 => unreachable, // types, not values
5950
5951 .undef => return cg.fail("assembly input with 'c' constraint cannot be undefined", .{}),
5952
5953 .int => try ass.value_map.put(gpa, name, .{ .constant = @intCast(val.toUnsignedInt(zcu)) }),
5954 .enum_literal => |str| try ass.value_map.put(gpa, name, .{ .string = str.toSlice(ip) }),
5955
5956 else => unreachable, // TODO
5957 }
5958 } else if (std.mem.eql(u8, constraint, "t")) {
5959 // type
5960 if (input_ty.zigTypeTag(zcu) == .type) {
5961 // This assembly input is a type instead of a value.
5962 // That's fine for now, just make sure to resolve it as such.
5963 const val = (try cg.air.value(input, cg.pt)).?;
5964 const ty_id = try cg.resolveType(val.toType(), .direct);
5965 try ass.value_map.put(gpa, name, .{ .ty = ty_id });
5966 } else {
5967 const ty_id = try cg.resolveType(input_ty, .direct);
5968 try ass.value_map.put(gpa, name, .{ .ty = ty_id });
5969 }
5970 } else {
5971 if (input_ty.zigTypeTag(zcu) == .type) {
5972 return cg.fail("use the 't' constraint to supply types to SPIR-V inline assembly", .{});
5973 }
5974
5975 const val_id = try cg.resolve(input);
5976 try ass.value_map.put(gpa, name, .{ .value = val_id });
5977 }
5978 }
5979
5980 // TODO: do something with clobbers
5981 _ = extra.data.clobbers;
5982
5983 const asm_source = std.mem.sliceAsBytes(cg.air.extra.items[extra_i..])[0..extra.data.source_len];
5984
5985 ass.assemble(asm_source) catch |err| switch (err) {
5986 error.AssembleFail => {
5987 // TODO: For now the compiler only supports a single error message per decl,
5988 // so to translate the possible multiple errors from the assembler, emit
5989 // them as notes here.
5990 // TODO: Translate proper error locations.
5991 assert(ass.errors.items.len != 0);
5992 assert(cg.error_msg == null);
5993 const src_loc = zcu.navSrcLoc(cg.owner_nav);
5994 cg.error_msg = try Zcu.ErrorMsg.create(zcu.gpa, src_loc, "failed to assemble SPIR-V inline assembly", .{});
5995 const notes = try zcu.gpa.alloc(Zcu.ErrorMsg, ass.errors.items.len);
5996
5997 // Sub-scope to prevent `return error.CodegenFail` from running the errdefers.
5998 {
5999 errdefer zcu.gpa.free(notes);
6000 var i: usize = 0;
6001 errdefer for (notes[0..i]) |*note| {
6002 note.deinit(zcu.gpa);
6003 };
6004
6005 while (i < ass.errors.items.len) : (i += 1) {
6006 notes[i] = try Zcu.ErrorMsg.init(zcu.gpa, src_loc, "{s}", .{ass.errors.items[i].msg});
6007 }
6008 }
6009 cg.error_msg.?.notes = notes;
6010 return error.CodegenFail;
6011 },
6012 else => |others| return others,
6013 };
6014
6015 for (outputs) |output| {
6016 _ = output;
6017 const extra_bytes = std.mem.sliceAsBytes(cg.air.extra.items[output_extra_i..]);
6018 const constraint = std.mem.sliceTo(std.mem.sliceAsBytes(cg.air.extra.items[output_extra_i..]), 0);
6019 const name = std.mem.sliceTo(extra_bytes[constraint.len + 1 ..], 0);
6020 output_extra_i += (constraint.len + name.len + (2 + 3)) / 4;
6021
6022 const result = ass.value_map.get(name) orelse return {
6023 return cg.fail("invalid asm output '{s}'", .{name});
6024 };
6025
6026 switch (result) {
6027 .just_declared, .unresolved_forward_reference => unreachable,
6028 .ty => return cg.fail("cannot return spir-v type as value from assembly", .{}),
6029 .value => |ref| return ref,
6030 .constant, .string => return cg.fail("cannot return constant from assembly", .{}),
6031 }
6032
6033 // TODO: Multiple results
6034 // TODO: Check that the output type from assembly is the same as the type actually expected by Zig.
6035 }
6036
6037 return null;
6038}
6039
6040fn airCall(cg: *CodeGen, inst: Air.Inst.Index, modifier: std.builtin.CallModifier) !?Id {
6041 _ = modifier;
6042
6043 const gpa = cg.module.gpa;
6044 const zcu = cg.module.zcu;
6045 const pl_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].pl_op;
6046 const extra = cg.air.extraData(Air.Call, pl_op.payload);
6047 const args: []const Air.Inst.Ref = @ptrCast(cg.air.extra.items[extra.end..][0..extra.data.args_len]);
6048 const callee_ty = cg.typeOf(pl_op.operand);
6049 const zig_fn_ty = switch (callee_ty.zigTypeTag(zcu)) {
6050 .@"fn" => callee_ty,
6051 .pointer => return cg.fail("cannot call function pointers", .{}),
6052 else => unreachable,
6053 };
6054 const fn_info = zcu.typeToFunc(zig_fn_ty).?;
6055 const return_type = fn_info.return_type;
6056
6057 const result_type_id = try cg.resolveFnReturnType(.fromInterned(return_type));
6058 const result_id = cg.module.allocId();
6059 const callee_id = try cg.resolve(pl_op.operand);
6060
6061 comptime assert(zig_call_abi_ver == 3);
6062
6063 const scratch_top = cg.id_scratch.items.len;
6064 defer cg.id_scratch.shrinkRetainingCapacity(scratch_top);
6065 const params = try cg.id_scratch.addManyAsSlice(gpa, args.len);
6066
6067 var n_params: usize = 0;
6068 for (args) |arg| {
6069 // Note: resolve() might emit instructions, so we need to call it
6070 // before starting to emit OpFunctionCall instructions. Hence the
6071 // temporary params buffer.
6072 const arg_ty = cg.typeOf(arg);
6073 if (!arg_ty.hasRuntimeBitsIgnoreComptime(zcu)) continue;
6074 const arg_id = try cg.resolve(arg);
6075
6076 params[n_params] = arg_id;
6077 n_params += 1;
6078 }
6079
6080 try cg.body.emit(gpa, .OpFunctionCall, .{
6081 .id_result_type = result_type_id,
6082 .id_result = result_id,
6083 .function = callee_id,
6084 .id_ref_3 = params[0..n_params],
6085 });
6086
6087 if (cg.liveness.isUnused(inst) or !Type.fromInterned(return_type).hasRuntimeBitsIgnoreComptime(zcu)) {
6088 return null;
6089 }
6090
6091 return result_id;
6092}
6093
6094fn builtin3D(
6095 cg: *CodeGen,
6096 result_ty: Type,
6097 builtin: spec.BuiltIn,
6098 dimension: u32,
6099 out_of_range_value: anytype,
6100) !Id {
6101 const gpa = cg.module.gpa;
6102 if (dimension >= 3) return try cg.constInt(result_ty, out_of_range_value);
6103 const u32_ty_id = try cg.module.intType(.unsigned, 32);
6104 const vec_ty_id = try cg.module.vectorType(3, u32_ty_id);
6105 const ptr_ty_id = try cg.module.ptrType(vec_ty_id, .input);
6106 const spv_decl_index = try cg.module.builtin(ptr_ty_id, builtin, .input);
6107 try cg.module.decl_deps.append(gpa, spv_decl_index);
6108 const ptr_id = cg.module.declPtr(spv_decl_index).result_id;
6109 const vec_id = cg.module.allocId();
6110 try cg.body.emit(gpa, .OpLoad, .{
6111 .id_result_type = vec_ty_id,
6112 .id_result = vec_id,
6113 .pointer = ptr_id,
6114 });
6115 return try cg.extractVectorComponent(result_ty, vec_id, dimension);
6116}
6117
6118fn airWorkItemId(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
6119 if (cg.liveness.isUnused(inst)) return null;
6120 const pl_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].pl_op;
6121 const dimension = pl_op.payload;
6122 return try cg.builtin3D(.u32, .local_invocation_id, dimension, 0);
6123}
6124
6125// TODO: this must be an OpConstant/OpSpec but even then the driver crashes.
6126fn airWorkGroupSize(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
6127 if (cg.liveness.isUnused(inst)) return null;
6128 const pl_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].pl_op;
6129 const dimension = pl_op.payload;
6130 return try cg.builtin3D(.u32, .workgroup_size, dimension, 0);
6131}
6132
6133fn airWorkGroupId(cg: *CodeGen, inst: Air.Inst.Index) !?Id {
6134 if (cg.liveness.isUnused(inst)) return null;
6135 const pl_op = cg.air.instructions.items(.data)[@intFromEnum(inst)].pl_op;
6136 const dimension = pl_op.payload;
6137 return try cg.builtin3D(.u32, .workgroup_id, dimension, 0);
6138}
6139
6140fn typeOf(cg: *CodeGen, inst: Air.Inst.Ref) Type {
6141 const zcu = cg.module.zcu;
6142 return cg.air.typeOf(inst, &zcu.intern_pool);
6143}
6144
6145fn typeOfIndex(cg: *CodeGen, inst: Air.Inst.Index) Type {
6146 const zcu = cg.module.zcu;
6147 return cg.air.typeOfIndex(inst, &zcu.intern_pool);
6148}