lib/choir/src/backends/gpu/spirv/emitter/gpu.zig
daab053ee43316e1809a84551d573ddd1e5bf3d2
1 const std = @import("std");
2 const choir = @import("../../../../root.zig");
3
4 const ir = choir.ir;
5 const dialects = choir.dialects;
6 const spirv_target = @import("../root.zig");
7 const gpu_target = @import("../../../../dialects/gpu/root.zig");
8 const binary = @import("module.zig");
9 const scalar = @import("scalar.zig");
10 const spec = @import("spec.zig");
11 const spirv_ops = @import("ops.zig");
12
13 const ArithDialect = dialects.arith.ArithDialect;
14 const GpuDialect = gpu_target.GpuDialect;
15 const SpirvDialect = spirv_target.SpirvDialect;
16 const GpuDimension = gpu_target.Dimension;
17 const GpuScope = gpu_target.Scope;
18 const WarpOpKind = gpu_target.WarpOpKind;
19 const SpirvOp = spirv_ops.SpirvOp;
20 const SpvCapability = spec.Capability;
21 const SpvDecoration = spec.Decoration;
22 const SpvStorageClass = spec.StorageClass;
23 const SpirvVersion = binary.Version;
24 const SpvScope = spec.Scope;
25 const SpvMemorySemanticsMask = spec.MemorySemanticsMask;
26
27 const SpvGroupOperation = struct {
28 const Reduce: u32 = 0;
29 const InclusiveScan: u32 = 1;
30 const ExclusiveScan: u32 = 2;
31 const ClusteredReduce: u32 = 3;
32 };
33
34 pub const SpvBuiltIn = struct {
35 pub const Position: u32 = 0;
36 pub const FragCoord: u32 = 15;
37 pub const FrontFacing: u32 = 17;
38 pub const NumWorkgroups: u32 = 24;
39 pub const WorkgroupSize: u32 = 25;
40 pub const WorkgroupId: u32 = 26;
41 pub const LocalInvocationId: u32 = 27;
42 pub const GlobalInvocationId: u32 = 28;
43 pub const SubgroupId: u32 = 40;
44 pub const SubgroupLocalInvocationId: u32 = 41;
45 pub const VertexIndex: u32 = 42;
46 pub const InstanceIndex: u32 = 43;
47 };
48
49 pub const BuiltinKind = enum {
50 num_workgroups,
51 workgroup_size,
52 workgroup_id,
53 local_invocation_id,
54 global_invocation_id,
55 subgroup_id,
56 subgroup_local_invocation_id,
57 position,
58 frag_coord,
59 vertex_index,
60 instance_index,
61 front_facing,
62 };
63
64 pub const AllAnyKind = enum { all, any };
65
66 const GroupCapability = enum { vote, ballot, shuffle, shuffle_relative, arithmetic };
67
68 pub fn emitIndex(self: anytype, op: *ir.Operation, builtin: BuiltinKind) !void {
69 const dim = getGpuDimension(op) orelse return error.MissingAttribute;
70 const u32_type = try self.getScalarType(.u32);
71
72 const composite_id = if (builtin == .workgroup_size)
73 try self.getWorkgroupSizeConstant()
74 else blk: {
75 const builtin_id = try getBuiltinVar(self, builtin);
76 const vec_type = try self.getVectorType(.u32, 3);
77 const load_id = self.builder.newId();
78 try self.builder.emit(&self.builder.functions, SpirvOp.Load, &.{
79 vec_type,
80 load_id,
81 builtin_id,
82 });
83 break :blk load_id;
84 };
85
86 const result_id = self.builder.newId();
87 const component_index: u32 = switch (dim) {
88 .x => 0,
89 .y => 1,
90 .z => 2,
91 };
92
93 try self.builder.emit(&self.builder.functions, SpirvOp.CompositeExtract, &.{
94 u32_type,
95 result_id,
96 composite_id,
97 component_index,
98 });
99
100 const result = op.getResult(0) orelse return error.UnsupportedOperation;
101 try self.bindValue(result, result_id);
102 }
103
104 pub fn emitBarrier(self: anytype, op: *ir.Operation) !void {
105 const barrier = GpuDialect.BarrierOp{ .op = op };
106 const scope = barrier.getScope() orelse return error.MissingAttribute;
107 const exec_scope = spvScopeForGpuScope(scope);
108 const mem_scope = exec_scope;
109 const semantics = spvMemorySemanticsForScope(scope);
110 const u32_type = try self.getScalarType(.u32);
111 const exec_id = try self.getIntConstant(u32_type, .u32, @intCast(exec_scope));
112 const mem_id = try self.getIntConstant(u32_type, .u32, @intCast(mem_scope));
113 const sem_id = try self.getIntConstant(u32_type, .u32, @intCast(semantics));
114
115 try self.builder.emit(&self.builder.functions, SpirvOp.ControlBarrier, &.{
116 exec_id,
117 mem_id,
118 sem_id,
119 });
120 }
121
122 pub fn emitSubgroupIndex(self: anytype, op: *ir.Operation, builtin: BuiltinKind) !void {
123 const result = op.getResult(0) orelse return error.UnsupportedOperation;
124 const builtin_id = try getBuiltinVar(self, builtin);
125 const u32_type = try self.getScalarType(.u32);
126
127 const load_id = self.builder.newId();
128 try self.builder.emit(&self.builder.functions, SpirvOp.Load, &.{
129 u32_type,
130 load_id,
131 builtin_id,
132 });
133
134 try self.bindValue(result, load_id);
135 }
136
137 pub fn emitSyncWarp(self: anytype, op: *ir.Operation) !void {
138 const sync = GpuDialect.SyncWarpOp{ .op = op };
139 try requireFullMask(self, sync.getMask());
140 const u32_type = try self.getScalarType(.u32);
141 const exec_id = try self.getIntConstant(u32_type, .u32, SpvScope.Subgroup);
142 const mem_id = exec_id;
143 const semantics = SpvMemorySemanticsMask.AcquireRelease | SpvMemorySemanticsMask.SubgroupMemory;
144 const sem_id = try self.getIntConstant(u32_type, .u32, semantics);
145
146 try self.builder.emit(&self.builder.functions, SpirvOp.ControlBarrier, &.{
147 exec_id,
148 mem_id,
149 sem_id,
150 });
151 }
152
153 pub fn emitActiveMask(self: anytype, op: *ir.Operation) !void {
154 const result = op.getResult(0) orelse return error.UnsupportedOperation;
155 const true_id = try self.getBoolConstant(true);
156 const mask_id = try emitGroupBallotMask(self, true_id);
157 try bindMaskValue(self, result, mask_id);
158 }
159
160 pub fn emitAllAny(self: anytype, op: *ir.Operation, kind: AllAnyKind) !void {
161 try requireFullMask(self, op.operands.items[0].value);
162 try requireGroupNonUniform(self, .vote);
163 const result = op.getResult(0) orelse return error.UnsupportedOperation;
164 const predicate = op.operands.items[1].value;
165 const pred_id = try self.getValue(predicate);
166 const result_type_id = try self.getTypeForValue(result);
167
168 const u32_type = try self.getScalarType(.u32);
169 const scope_id = try self.getIntConstant(u32_type, .u32, SpvScope.Subgroup);
170 const opcode = switch (kind) {
171 .all => SpirvOp.GroupNonUniformAll,
172 .any => SpirvOp.GroupNonUniformAny,
173 };
174
175 const result_id = self.builder.newId();
176 try self.builder.emit(&self.builder.functions, opcode, &.{
177 result_type_id,
178 result_id,
179 scope_id,
180 pred_id,
181 });
182 try self.bindValue(result, result_id);
183 }
184
185 pub fn emitBallot(self: anytype, op: *ir.Operation) !void {
186 const ballot = GpuDialect.BallotSyncOp{ .op = op };
187 try requireFullMask(self, ballot.getMask());
188 const pred_id = try self.getValue(ballot.getPredicate());
189 const result = ballot.getResult();
190 const mask_id = try emitGroupBallotMask(self, pred_id);
191 try bindMaskValue(self, result, mask_id);
192 }
193
194 pub fn emitShuffle(self: anytype, op: *ir.Operation) !void {
195 const shfl = GpuDialect.ShflSyncOp{ .op = op };
196 const mode = shfl.getMode() orelse return error.MissingAttribute;
197 try requireFullMask(self, shfl.getMask());
198 try requireGroupNonUniform(self, switch (mode) {
199 .sync, .xor => .shuffle,
200 .down, .up => .shuffle_relative,
201 });
202 const src_id = try self.getValue(shfl.getSrc());
203 const lane_id = try self.getValue(shfl.getLaneOrDelta());
204 const result = shfl.getResult();
205 const result_type_id = try self.getTypeForValue(result);
206
207 const u32_type = try self.getScalarType(.u32);
208 const scope_id = try self.getIntConstant(u32_type, .u32, SpvScope.Subgroup);
209 const opcode = switch (mode) {
210 .sync => SpirvOp.GroupNonUniformShuffle,
211 .down => SpirvOp.GroupNonUniformShuffleDown,
212 .up => SpirvOp.GroupNonUniformShuffleUp,
213 .xor => SpirvOp.GroupNonUniformShuffleXor,
214 };
215
216 const result_id = self.builder.newId();
217 try self.builder.emit(&self.builder.functions, opcode, &.{
218 result_type_id,
219 result_id,
220 scope_id,
221 src_id,
222 lane_id,
223 });
224 try self.bindValue(result, result_id);
225 }
226
227 pub fn emitWarpReduce(self: anytype, op: *ir.Operation) !void {
228 const reduce = GpuDialect.WarpReduceOp{ .op = op };
229 try requireFullMask(self, reduce.getMask());
230 try requireGroupNonUniform(self, .arithmetic);
231 const op_kind = reduce.getOpKind() orelse return error.MissingAttribute;
232 const value = reduce.getValue();
233 const value_id = try self.getValue(value);
234 const result = reduce.getResult();
235 const kind = scalar.kindFromType(value.type) orelse return error.UnsupportedType;
236 const opcode = groupNonUniformOpcode(op_kind, kind) orelse return error.UnsupportedType;
237
238 const u32_type = try self.getScalarType(.u32);
239 const scope_id = try self.getIntConstant(u32_type, .u32, SpvScope.Subgroup);
240 const result_type_id = try self.getTypeForValue(result);
241 const result_id = self.builder.newId();
242 try self.builder.emit(&self.builder.functions, opcode, &.{
243 result_type_id,
244 result_id,
245 scope_id,
246 SpvGroupOperation.Reduce,
247 value_id,
248 });
249 try self.bindValue(result, result_id);
250 }
251
252 pub fn emitWarpScan(self: anytype, op: *ir.Operation) !void {
253 const scan = GpuDialect.WarpScanOp{ .op = op };
254 try requireFullMask(self, scan.getMask());
255 try requireGroupNonUniform(self, .arithmetic);
256 const op_kind = scan.getOpKind() orelse return error.MissingAttribute;
257 const value = scan.getValue();
258 const value_id = try self.getValue(value);
259 const result = scan.getResult();
260 const kind = scalar.kindFromType(value.type) orelse return error.UnsupportedType;
261 const opcode = groupNonUniformOpcode(op_kind, kind) orelse return error.UnsupportedType;
262
263 const u32_type = try self.getScalarType(.u32);
264 const scope_id = try self.getIntConstant(u32_type, .u32, SpvScope.Subgroup);
265 const group_op = if (scan.isInclusive())
266 SpvGroupOperation.InclusiveScan
267 else
268 SpvGroupOperation.ExclusiveScan;
269 const result_type_id = try self.getTypeForValue(result);
270 const result_id = self.builder.newId();
271 try self.builder.emit(&self.builder.functions, opcode, &.{
272 result_type_id,
273 result_id,
274 scope_id,
275 group_op,
276 value_id,
277 });
278 try self.bindValue(result, result_id);
279 }
280
281 fn requireGroupNonUniform(self: anytype, capability: GroupCapability) !void {
282 self.builder.requireVersion(SpirvVersion.v13);
283 try self.requireCapability(SpvCapability.GroupNonUniform);
284 switch (capability) {
285 .vote => try self.requireCapability(SpvCapability.GroupNonUniformVote),
286 .ballot => try self.requireCapability(SpvCapability.GroupNonUniformBallot),
287 .shuffle => try self.requireCapability(SpvCapability.GroupNonUniformShuffle),
288 .shuffle_relative => try self.requireCapability(SpvCapability.GroupNonUniformShuffleRelative),
289 .arithmetic => try self.requireCapability(SpvCapability.GroupNonUniformArithmetic),
290 }
291 }
292
293 fn requireFullMask(self: anytype, mask: *ir.Value) !void {
294 if (isFullMaskValue(self, mask)) return;
295 return error.UnsupportedMask;
296 }
297
298 fn isFullMaskValue(self: anytype, mask: *ir.Value) bool {
299 _ = self;
300 const kind = scalar.kindFromType(mask.type) orelse return false;
301 if (!scalar.isInt(kind)) return false;
302
303 const defining = mask.getDefiningOp() orelse return false;
304 const def_op: *ir.Operation = @ptrCast(@alignCast(defining));
305 if (std.mem.eql(u8, def_op.name.name, ArithDialect.ConstantOp.operation_name)) {
306 const constant = ArithDialect.ConstantOp{ .op = def_op };
307 if (constant.getIntValue()) |int_value| {
308 return scalar.isFullMaskConstant(kind, int_value);
309 }
310 }
311 if (std.mem.eql(u8, def_op.name.name, SpirvDialect.ConstantOp.operation_name)) {
312 const constant = SpirvDialect.ConstantOp{ .op = def_op };
313 if (constant.getIntValue()) |int_value| {
314 return scalar.isFullMaskConstant(kind, int_value);
315 }
316 }
317 return false;
318 }
319
320 fn emitGroupBallotMask(self: anytype, pred_id: u32) !u32 {
321 try requireGroupNonUniform(self, .ballot);
322 const vec_type = try self.getVectorType(.u32, 4);
323 const u32_type = try self.getScalarType(.u32);
324 const scope_id = try self.getIntConstant(u32_type, .u32, SpvScope.Subgroup);
325
326 const ballot_id = self.builder.newId();
327 try self.builder.emit(&self.builder.functions, SpirvOp.GroupNonUniformBallot, &.{
328 vec_type,
329 ballot_id,
330 scope_id,
331 pred_id,
332 });
333
334 const mask_id = self.builder.newId();
335 try self.builder.emit(&self.builder.functions, SpirvOp.CompositeExtract, &.{
336 u32_type,
337 mask_id,
338 ballot_id,
339 0,
340 });
341 return mask_id;
342 }
343
344 fn bindMaskValue(self: anytype, result: *ir.Value, mask_id: u32) !void {
345 const result_kind = scalar.kindFromType(result.type) orelse return error.UnsupportedType;
346 if (result_kind == .u32) {
347 try self.bindValue(result, mask_id);
348 return;
349 }
350 if (result_kind == .i32) {
351 const result_type_id = try self.getScalarType(.i32);
352 const cast_id = self.builder.newId();
353 try self.builder.emit(&self.builder.functions, SpirvOp.Bitcast, &.{
354 result_type_id,
355 cast_id,
356 mask_id,
357 });
358 try self.bindValue(result, cast_id);
359 return;
360 }
361 return error.UnsupportedType;
362 }
363
364 pub fn getBuiltinVar(self: anytype, builtin: BuiltinKind) !u32 {
365 if (self.builtin_vars.get(builtin)) |id| {
366 try self.addInterfaceVar(id);
367 return id;
368 }
369
370 const value_type = try builtinType(self, builtin);
371 const storage_class: u32 = if (builtin == .position)
372 SpvStorageClass.Output
373 else
374 SpvStorageClass.Input;
375 const ptr_type = try self.getPointerType(storage_class, value_type);
376
377 const var_id = self.builder.newId();
378 try self.builder.emit(&self.builder.globals, SpirvOp.Variable, &.{
379 ptr_type,
380 var_id,
381 storage_class,
382 });
383
384 const builtin_value: u32 = switch (builtin) {
385 .num_workgroups => SpvBuiltIn.NumWorkgroups,
386 .workgroup_size => return error.UnsupportedOperation,
387 .workgroup_id => SpvBuiltIn.WorkgroupId,
388 .local_invocation_id => SpvBuiltIn.LocalInvocationId,
389 .global_invocation_id => SpvBuiltIn.GlobalInvocationId,
390 .subgroup_id => SpvBuiltIn.SubgroupId,
391 .subgroup_local_invocation_id => SpvBuiltIn.SubgroupLocalInvocationId,
392 .position => SpvBuiltIn.Position,
393 .frag_coord => SpvBuiltIn.FragCoord,
394 .vertex_index => SpvBuiltIn.VertexIndex,
395 .instance_index => SpvBuiltIn.InstanceIndex,
396 .front_facing => SpvBuiltIn.FrontFacing,
397 };
398
399 try self.builder.emit(&self.builder.annotations, SpirvOp.Decorate, &.{
400 var_id,
401 SpvDecoration.BuiltIn,
402 builtin_value,
403 });
404
405 try self.builtin_vars.put(self.allocator, builtin, var_id);
406 try self.addInterfaceVar(var_id);
407 return var_id;
408 }
409
410 /// The type a builtin variable holds.
411 pub fn builtinType(self: anytype, builtin: BuiltinKind) !u32 {
412 return switch (builtin) {
413 .num_workgroups,
414 .workgroup_size,
415 .workgroup_id,
416 .local_invocation_id,
417 .global_invocation_id,
418 => self.getVectorType(.u32, 3),
419 .subgroup_id,
420 .subgroup_local_invocation_id,
421 .vertex_index,
422 .instance_index,
423 => self.getScalarType(.u32),
424 .position, .frag_coord => self.getVectorType(.f32, 4),
425 .front_facing => self.getScalarType(.bool),
426 };
427 }
428
429 fn spvScopeForGpuScope(scope: GpuScope) u32 {
430 return switch (scope) {
431 .thread => SpvScope.Invocation,
432 .warp => SpvScope.Subgroup,
433 .block => SpvScope.Workgroup,
434 .cluster => SpvScope.Device,
435 .device => SpvScope.Device,
436 .system => SpvScope.CrossDevice,
437 };
438 }
439
440 fn spvMemorySemanticsForScope(scope: GpuScope) u32 {
441 var mask: u32 = SpvMemorySemanticsMask.AcquireRelease;
442 switch (scope) {
443 .thread => {},
444 .warp => mask |= SpvMemorySemanticsMask.SubgroupMemory,
445 .block => mask |= SpvMemorySemanticsMask.WorkgroupMemory,
446 .cluster, .device, .system => mask |= SpvMemorySemanticsMask.CrossWorkgroupMemory,
447 }
448 return mask;
449 }
450
451 fn groupNonUniformOpcode(op_kind: WarpOpKind, scalar_kind: scalar.Kind) ?u16 {
452 return switch (op_kind) {
453 .add => if (scalar.isFloat(scalar_kind))
454 SpirvOp.GroupNonUniformFAdd
455 else if (scalar.isInt(scalar_kind))
456 SpirvOp.GroupNonUniformIAdd
457 else
458 null,
459 .min => if (scalar.isFloat(scalar_kind))
460 SpirvOp.GroupNonUniformFMin
461 else if (scalar.isSignedInt(scalar_kind))
462 SpirvOp.GroupNonUniformSMin
463 else if (scalar.isUnsignedInt(scalar_kind))
464 SpirvOp.GroupNonUniformUMin
465 else
466 null,
467 .max => if (scalar.isFloat(scalar_kind))
468 SpirvOp.GroupNonUniformFMax
469 else if (scalar.isSignedInt(scalar_kind))
470 SpirvOp.GroupNonUniformSMax
471 else if (scalar.isUnsignedInt(scalar_kind))
472 SpirvOp.GroupNonUniformUMax
473 else
474 null,
475 .and_ => if (scalar_kind == .bool)
476 SpirvOp.GroupNonUniformLogicalAnd
477 else if (scalar.isInt(scalar_kind))
478 SpirvOp.GroupNonUniformBitwiseAnd
479 else
480 null,
481 .or_ => if (scalar_kind == .bool)
482 SpirvOp.GroupNonUniformLogicalOr
483 else if (scalar.isInt(scalar_kind))
484 SpirvOp.GroupNonUniformBitwiseOr
485 else
486 null,
487 .xor => if (scalar_kind == .bool)
488 SpirvOp.GroupNonUniformLogicalXor
489 else if (scalar.isInt(scalar_kind))
490 SpirvOp.GroupNonUniformBitwiseXor
491 else
492 null,
493 };
494 }
495
496 fn getGpuDimension(op: *ir.Operation) ?GpuDimension {
497 if (std.mem.eql(u8, op.name.name, GpuDialect.ThreadIdxOp.operation_name)) {
498 const thread = GpuDialect.ThreadIdxOp{ .op = op };
499 return thread.getDimension();
500 }
501 if (std.mem.eql(u8, op.name.name, GpuDialect.BlockIdxOp.operation_name)) {
502 const block = GpuDialect.BlockIdxOp{ .op = op };
503 return block.getDimension();
504 }
505 if (std.mem.eql(u8, op.name.name, GpuDialect.BlockDimOp.operation_name)) {
506 const dim = GpuDialect.BlockDimOp{ .op = op };
507 return dim.getDimension();
508 }
509 if (std.mem.eql(u8, op.name.name, GpuDialect.GridDimOp.operation_name)) {
510 const dim = GpuDialect.GridDimOp{ .op = op };
511 return dim.getDimension();
512 }
513 if (std.mem.eql(u8, op.name.name, GpuDialect.GlobalIdxOp.operation_name)) {
514 const global = GpuDialect.GlobalIdxOp{ .op = op };
515 return global.getDimension();
516 }
517 if (std.mem.eql(u8, op.name.name, SpirvDialect.LocalInvocationIdOp.operation_name)) {
518 const local = SpirvDialect.LocalInvocationIdOp{ .op = op };
519 return local.getDimension();
520 }
521 if (std.mem.eql(u8, op.name.name, SpirvDialect.WorkgroupIdOp.operation_name)) {
522 const wg = SpirvDialect.WorkgroupIdOp{ .op = op };
523 return wg.getDimension();
524 }
525 if (std.mem.eql(u8, op.name.name, SpirvDialect.WorkgroupSizeOp.operation_name)) {
526 const wg_size = SpirvDialect.WorkgroupSizeOp{ .op = op };
527 return wg_size.getDimension();
528 }
529 if (std.mem.eql(u8, op.name.name, SpirvDialect.NumWorkgroupsOp.operation_name)) {
530 const num = SpirvDialect.NumWorkgroupsOp{ .op = op };
531 return num.getDimension();
532 }
533 if (std.mem.eql(u8, op.name.name, SpirvDialect.GlobalInvocationIdOp.operation_name)) {
534 const global = SpirvDialect.GlobalInvocationIdOp{ .op = op };
535 return global.getDimension();
536 }
537 const dialect_attr = op.getAttrAs(ir.Attribute.DialectAttr, "dim") orelse return null;
538 return GpuDimension.fromString(dialect_attr.payload);
539 }
540
541 test "spirv gpu maps barrier scopes to memory semantics" {
542 try std.testing.expectEqual(SpvScope.Invocation, spvScopeForGpuScope(.thread));
543 try std.testing.expectEqual(SpvScope.Subgroup, spvScopeForGpuScope(.warp));
544 try std.testing.expectEqual(SpvScope.Workgroup, spvScopeForGpuScope(.block));
545 try std.testing.expectEqual(SpvScope.Device, spvScopeForGpuScope(.device));
546 try std.testing.expectEqual(SpvScope.CrossDevice, spvScopeForGpuScope(.system));
547 try std.testing.expectEqual(SpvMemorySemanticsMask.AcquireRelease | SpvMemorySemanticsMask.SubgroupMemory, spvMemorySemanticsForScope(.warp));
548 try std.testing.expectEqual(SpvMemorySemanticsMask.AcquireRelease | SpvMemorySemanticsMask.WorkgroupMemory, spvMemorySemanticsForScope(.block));
549 try std.testing.expectEqual(SpvMemorySemanticsMask.AcquireRelease | SpvMemorySemanticsMask.CrossWorkgroupMemory, spvMemorySemanticsForScope(.device));
550 }
551
552 test "spirv gpu maps warp reductions by scalar kind" {
553 try std.testing.expectEqual(SpirvOp.GroupNonUniformFAdd, groupNonUniformOpcode(.add, .f32).?);
554 try std.testing.expectEqual(SpirvOp.GroupNonUniformIAdd, groupNonUniformOpcode(.add, .i32).?);
555 try std.testing.expectEqual(SpirvOp.GroupNonUniformSMin, groupNonUniformOpcode(.min, .i32).?);
556 try std.testing.expectEqual(SpirvOp.GroupNonUniformUMax, groupNonUniformOpcode(.max, .u32).?);
557 try std.testing.expectEqual(SpirvOp.GroupNonUniformLogicalAnd, groupNonUniformOpcode(.and_, .bool).?);
558 try std.testing.expectEqual(@as(?u16, null), groupNonUniformOpcode(.and_, .f32));
559 }