lib/accy/src/kernel/compile/artifact.zig
daab053ee43316e1809a84551d573ddd1e5bf3d2
1 const std = @import("std");
2 const gpu = @import("gpu");
3 const choir = @import("choir");
4 const choir_abi = @import("choir_abi");
5 const accy_root = @import("../../root.zig");
6 const artifact_product = @import("../../artifact/root.zig");
7 const preparation = @import("../../preparation/root.zig");
8 const target = @import("../../target/root.zig");
9 const kernel_model = @import("../model/root.zig");
10 const builder = kernel_model.core.builder;
11 const plan_mod = kernel_model.plan;
12 const program_mod = @import("../program/root.zig");
13
14 const DType = choir_abi.DType;
15
16 pub const Options = struct {
17 format: ?gpu.ArtifactFormat = null,
18 kernel_plan: plan_mod.Options = .{},
19 disable_cpu_vectorization: bool = false,
20 };
21
22 const CompilePlan = struct {
23 entry_name: []const u8,
24 argument_count: u32,
25 payload: gpu.CompilePayload,
26 required_features: choir_abi.Features = .{},
27 required_subgroup: choir_abi.SubgroupRequirements = .{},
28 push_constants: choir_abi.PushConstants = .{},
29 runtime_scalar_argument_count: u32 = 0,
30 static_arguments: []choir_abi.ScalarArgument = &.{},
31
32 fn deinit(self: *CompilePlan, allocator: std.mem.Allocator) void {
33 if (self.static_arguments.len != 0) allocator.free(self.static_arguments);
34 deinitCompilePayload(allocator, self.payload);
35 self.* = undefined;
36 }
37 };
38
39 pub fn createJob(
40 allocator: std.mem.Allocator,
41 handle: gpu.BackendHandle,
42 program: *program_mod.Program,
43 options: Options,
44 ) !*artifact_product.ArtifactJob {
45 var artifact_plan = try createPlan(allocator, handle, program, options);
46 var plan_owned = true;
47 errdefer if (plan_owned) artifact_plan.deinit();
48
49 const artifact_job = try artifact_product.ArtifactJob.init(
50 allocator,
51 artifact_plan,
52 );
53 plan_owned = false;
54 return artifact_job;
55 }
56
57 pub fn createPlan(
58 allocator: std.mem.Allocator,
59 handle: gpu.BackendHandle,
60 program: *program_mod.Program,
61 options: Options,
62 ) !artifact_product.BackendArtifactPlan {
63 var authored_plan = try program.createCheckedPlan(allocator, options.kernel_plan);
64 defer authored_plan.deinit();
65
66 const caps = try handle.queryCapabilities();
67 const backend_kind = handle.backendKind() orelse caps.identity.backend;
68 const format = options.format orelse artifact_product.defaultArtifactFormat(backend_kind) orelse {
69 return error.UnsupportedOperation;
70 };
71 if (!caps.supportsArtifactFormat(format)) return error.UnsupportedArtifactFormat;
72 const profile = try preparation.BackendTargetProfile.init(caps, backend_kind, format);
73
74 const launch_resources = try createLaunchResourcePlan(caps, format, &authored_plan);
75 var artifact_plan = artifact_product.BackendArtifactPlan.init(allocator, profile);
76 errdefer artifact_plan.deinit();
77
78 var compile_plan = try compilePlan(allocator, format, &authored_plan, program, launch_resources, options.disable_cpu_vectorization);
79 defer compile_plan.deinit(allocator);
80 const required_dtypes = requiredDTypesForParams(authored_plan.params);
81 var compile = try artifact_product.PlannedKernelCompileContract.init(
82 allocator,
83 .choir_kernel,
84 .authored,
85 format,
86 compile_plan.entry_name,
87 compile_plan.argument_count,
88 required_dtypes,
89 compile_plan.required_features,
90 compile_plan.required_subgroup,
91 null,
92 compile_plan.payload,
93 );
94 var compile_owned = true;
95 errdefer if (compile_owned) compile.deinit(allocator);
96
97 const artifact = try createKernelArtifact(
98 handle,
99 format,
100 &compile_plan,
101 authored_plan.diagnostic_id,
102 required_dtypes,
103 compile_plan.required_features,
104 );
105
106 compile_owned = false;
107 try artifact_plan.addStandaloneKernel(artifact, launch_resources, compile, .{
108 .runtime_scalar_argument_count = compile_plan.runtime_scalar_argument_count,
109 .static_arguments = compile_plan.static_arguments,
110 });
111 return artifact_plan;
112 }
113
114 fn requiredDTypesForParams(params: []const builder.Param) gpu.DTypeSet {
115 var dtypes: gpu.DTypeSet = .{};
116 for (params) |param| {
117 switch (param) {
118 .scalar => |dtype| dtypes.insert(dtype),
119 .buffer => |buffer| dtypes.insert(buffer.dtype),
120 }
121 }
122 return dtypes;
123 }
124
125 fn createLaunchResourcePlan(
126 caps: gpu.BackendCapabilities,
127 format: gpu.ArtifactFormat,
128 authored_plan: *const plan_mod.Plan,
129 ) gpu.BackendError!artifact_product.LaunchResourcePlan {
130 const geometry = launchGeometry(authored_plan);
131 try caps.validateLaunchGeometry(geometry);
132 const element_count = try launchElementCount(geometry);
133
134 var resources: artifact_product.LaunchResourcePlan = .{
135 .format = format,
136 .element_count = element_count,
137 .geometry = geometry,
138 .fixed_threadgroup = true,
139 .candidate_count = 1,
140 };
141 resources.candidates[0] = .{
142 .geometry = geometry,
143 };
144 return resources;
145 }
146
147 fn launchGeometry(authored_plan: *const plan_mod.Plan) choir_abi.LaunchGeometry {
148 return .{
149 .grid = authored_plan.launch.grid,
150 .threadgroup = authored_plan.launch.block,
151 };
152 }
153
154 fn launchElementCount(geometry: choir_abi.LaunchGeometry) gpu.BackendError!u64 {
155 return try geometry.threadCount();
156 }
157
158 fn compilePlan(
159 allocator: std.mem.Allocator,
160 format: gpu.ArtifactFormat,
161 authored_plan: *const plan_mod.Plan,
162 program: *const program_mod.Program,
163 launch_resources: artifact_product.LaunchResourcePlan,
164 disable_cpu_vectorization: bool,
165 ) gpu.BackendError!CompilePlan {
166 const module = program.kernelModule();
167 const required_features = choir.backends.gpu.featureRequirementsForModule(module);
168 const required_subgroup = choir.backends.gpu.subgroupRequirementsForModule(module);
169 const static_arguments = try compileStaticArguments(allocator, format, launch_resources);
170 errdefer if (static_arguments.len != 0) allocator.free(static_arguments);
171 const compilation = try target.compileKernelForArtifactFormat(
172 allocator,
173 format,
174 authored_plan.entry_name,
175 module,
176 compileOptions(format, launch_resources, disable_cpu_vectorization),
177 );
178 errdefer deinitCompilePayload(allocator, compilation.payload);
179 return .{
180 .entry_name = authored_plan.entry_name,
181 .argument_count = try compileArgumentCount(format, authored_plan.argument_count),
182 .payload = compilation.payload,
183 .required_features = required_features,
184 .required_subgroup = required_subgroup,
185 .push_constants = compilation.push_constants,
186 .runtime_scalar_argument_count = try runtimeScalarArgumentCount(authored_plan.params),
187 .static_arguments = static_arguments,
188 };
189 }
190
191 fn compileOptions(
192 format: gpu.ArtifactFormat,
193 launch_resources: artifact_product.LaunchResourcePlan,
194 disable_cpu_vectorization: bool,
195 ) target.CompileOptions {
196 var options = target.compileOptionsForArtifactFormat(format, launch_resources.element_count);
197 if (disable_cpu_vectorization) options.cpu_vector_width = null;
198 return options;
199 }
200
201 fn runtimeScalarArgumentCount(params: []const builder.Param) gpu.BackendError!u32 {
202 var count: u32 = 0;
203 for (params) |param| {
204 switch (param) {
205 .scalar => count = std.math.add(u32, count, 1) catch return error.InvalidArtifact,
206 .buffer => {},
207 }
208 }
209 return count;
210 }
211
212 fn compileArgumentCount(format: gpu.ArtifactFormat, argument_count: u32) gpu.BackendError!u32 {
213 if (!gpu.artifactFormatUsesHostLoopLaunch(format)) return argument_count;
214 return choir_abi.kernelArgumentCount(argument_count);
215 }
216
217 fn compileStaticArguments(
218 allocator: std.mem.Allocator,
219 format: gpu.ArtifactFormat,
220 launch_resources: artifact_product.LaunchResourcePlan,
221 ) gpu.BackendError![]choir_abi.ScalarArgument {
222 if (!gpu.artifactFormatUsesHostLoopLaunch(format)) return &.{};
223 return choir_abi.launchShapeArguments(allocator, launch_resources.element_count, launch_resources.geometry);
224 }
225
226 fn createKernelArtifact(
227 handle: gpu.BackendHandle,
228 format: gpu.ArtifactFormat,
229 compile_plan: *const CompilePlan,
230 diagnostic_id: []const u8,
231 required_dtypes: gpu.DTypeSet,
232 required_features: choir_abi.Features,
233 ) gpu.BackendError!gpu.KernelArtifact {
234 return try handle.createArtifact(.{
235 .kernel_name = compile_plan.entry_name,
236 .requested_format = format,
237 .argument_count = compile_plan.argument_count,
238 .scalar_argument_count = compile_plan.runtime_scalar_argument_count + @as(u32, @intCast(compile_plan.static_arguments.len)),
239 .required_dtypes = required_dtypes,
240 .required_features = required_features,
241 .required_subgroup = compile_plan.required_subgroup,
242 .push_constants = compile_plan.push_constants,
243 .diagnostic_id = diagnostic_id,
244 .payload = compile_plan.payload,
245 });
246 }
247
248 fn deinitCompilePayload(allocator: std.mem.Allocator, payload: gpu.CompilePayload) void {
249 switch (payload) {
250 .bytes => |bytes| allocator.free(@constCast(bytes)),
251 .words_u32 => |words| allocator.free(@constCast(words)),
252 .text => |text| allocator.free(@constCast(text)),
253 .none => {},
254 }
255 }
256
257 const testing = std.testing;
258
259 const TestBackendState = struct {
260 allocator: std.mem.Allocator,
261 kind: gpu.BackendKind,
262 format: gpu.ArtifactFormat,
263
264 fn init(allocator: std.mem.Allocator, kind: gpu.BackendKind) TestBackendState {
265 return .{
266 .allocator = allocator,
267 .kind = kind,
268 .format = artifact_product.defaultArtifactFormat(kind).?,
269 };
270 }
271
272 fn handle(self: *TestBackendState) gpu.BackendHandle {
273 return .{
274 .ptr = self,
275 .vtable = &test_backend_vtable,
276 .kind = self.kind,
277 };
278 }
279 };
280
281 fn testQueryCapabilities(ptr: *anyopaque) gpu.BackendError!gpu.BackendCapabilities {
282 const state: *TestBackendState = @ptrCast(@alignCast(ptr));
283 const subgroup_size: u32 = switch (state.kind) {
284 .vulkan, .webgpu, .cpu => 0,
285 else => 32,
286 };
287 const dtypes = switch (state.kind) {
288 .cpu => gpu.DTypeSet.init(&.{ .i1, .i32, .u32, .i64, .u64, .f32, .f64 }),
289 .webgpu => gpu.DTypeSet.init(&.{ .i1, .i32, .u32, .f32 }),
290 else => gpu.DTypeSet.init(&.{ .i1, .i32, .u32, .f16, .f32 }),
291 };
292 return .{
293 .identity = .{
294 .backend = state.kind,
295 .family = gpu.familyForBackendKind(state.kind),
296 .name = "kernel-artifact-test",
297 },
298 .subgroup = if (subgroup_size == 0) .{} else .{
299 .supported = true,
300 .size_min = subgroup_size,
301 .size_max = subgroup_size,
302 .shuffle = true,
303 .ballot = true,
304 .vote = true,
305 .arithmetic = true,
306 .scan = state.kind == .cuda or state.kind == .metal,
307 },
308 .threadgroup = .{
309 .max_threads = 1024,
310 .max_blocks = .{ 65_535, 65_535, 65_535 },
311 .max_threads_per_dim = .{ 1024, 1024, 64 },
312 .max_grid_per_dim = .{ 65_535, 65_535, 65_535 },
313 },
314 .dtypes = dtypes,
315 .artifact_formats = gpu.ArtifactFormatSet.init(&.{state.format}),
316 .features = .{
317 .atomic_i32 = state.kind != .webgpu,
318 .atomic_u32 = state.kind != .webgpu,
319 .atomic_index = state.kind != .webgpu,
320 .atomic_f32_add_device = state.kind == .cuda or state.kind == .metal,
321 .atomic_f32_add_shared = state.kind == .cuda,
322 .dynamic_shared_memory = state.kind == .cuda,
323 },
324 };
325 }
326
327 test "kernel artifact plan emits native cpu object for authored global x kernel" {
328 const allocator = testing.allocator;
329
330 var builder_state = try program_mod.Builder.init(allocator, program_mod.Builder.Limits.testing, "authored_cpu_copy_f32", &.{
331 builder.dynamicBuffer(.f32),
332 builder.dynamicBuffer(.f32),
333 });
334 errdefer builder_state.deinit();
335
336 const axis = try builder_state.axis("i", 4);
337 try builder_state.bind(axis, .thread_x);
338 const src = builder_state.argument(0);
339 const dst = builder_state.argument(1);
340 const index = try builder_state.globalId(.x);
341 const value = try builder_state.load(src, index);
342 try builder_state.store(value, dst, index);
343 try builder_state.return_();
344
345 var program = try builder_state.finish();
346 defer program.deinit();
347
348 var state = TestBackendState.init(allocator, .cpu);
349 state.format = .cpu_object;
350 var artifact_plan = createPlan(allocator, state.handle(), &program, .{ .format = .cpu_object }) catch |err| switch (err) {
351 error.UnsupportedOperation => return error.SkipZigTest,
352 else => return err,
353 };
354 defer artifact_plan.deinit();
355
356 try testing.expectEqual(gpu.BackendKind.cpu, artifact_plan.backend_kind);
357 try testing.expectEqual(gpu.ArtifactFormat.cpu_object, artifact_plan.format);
358 try testing.expectEqual(@as(usize, 1), artifact_plan.kernelCount());
359
360 const planned = artifact_plan.kernels.items[0];
361 try testing.expectEqual(gpu.ArtifactFormat.cpu_object, planned.compile.format);
362 try testing.expectEqual(artifact_product.PlannedKernelCompilePayload.bytes, planned.compile.payload);
363 try testing.expect(planned.compile.payload_byte_count > 0);
364 try testing.expectEqual(@as(u32, 9), planned.compile.argument_count);
365 try testing.expectEqual(@as(u32, 9), planned.artifact.argument_count);
366 try testing.expectEqual(@as(usize, 7), planned.static_arguments.len);
367 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 4 }, planned.static_arguments[0]);
368 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[1]);
369 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[2]);
370 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[3]);
371 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 4 }, planned.static_arguments[4]);
372 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[5]);
373 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[6]);
374 switch (planned.artifact.payload) {
375 .bytes => |bytes| {
376 try testing.expect(bytes.len > 4);
377 try testing.expectEqual(@as(u8, 0x7f), bytes[0]);
378 try testing.expectEqual(@as(u8, 'E'), bytes[1]);
379 try testing.expectEqual(@as(u8, 'L'), bytes[2]);
380 try testing.expectEqual(@as(u8, 'F'), bytes[3]);
381 },
382 else => return error.ExpectedCpuObjectBytes,
383 }
384 }
385
386 test "kernel artifact plan emits webassembly module for authored global x kernel" {
387 const allocator = testing.allocator;
388
389 var builder_state = try program_mod.Builder.init(allocator, program_mod.Builder.Limits.testing, "authored_wasm_copy_f32", &.{
390 builder.dynamicBuffer(.f32),
391 builder.dynamicBuffer(.f32),
392 });
393 errdefer builder_state.deinit();
394
395 const axis = try builder_state.axis("i", 4);
396 try builder_state.bind(axis, .thread_x);
397 const src = builder_state.argument(0);
398 const dst = builder_state.argument(1);
399 const index = try builder_state.globalId(.x);
400 const value = try builder_state.load(src, index);
401 try builder_state.store(value, dst, index);
402 try builder_state.return_();
403
404 var program = try builder_state.finish();
405 defer program.deinit();
406
407 var state = gpu.wasm.State.init(allocator);
408 defer state.deinit();
409 var artifact_plan = try createPlan(allocator, state.handle(), &program, .{});
410 defer artifact_plan.deinit();
411
412 try testing.expectEqual(gpu.BackendKind.wasm, artifact_plan.backend_kind);
413 try testing.expectEqual(gpu.ArtifactFormat.webassembly_module, artifact_plan.format);
414 try testing.expectEqual(@as(usize, 1), artifact_plan.kernelCount());
415
416 const planned = artifact_plan.kernels.items[0];
417 try testing.expectEqual(gpu.ArtifactFormat.webassembly_module, planned.compile.format);
418 try testing.expectEqual(artifact_product.PlannedKernelCompilePayload.bytes, planned.compile.payload);
419 try testing.expectEqual(@as(u32, 9), planned.compile.argument_count);
420 try testing.expectEqual(@as(u32, 9), planned.artifact.argument_count);
421 try testing.expectEqual(@as(usize, 7), planned.static_arguments.len);
422 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 4 }, planned.static_arguments[0]);
423 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[1]);
424 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[2]);
425 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[3]);
426 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 4 }, planned.static_arguments[4]);
427 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[5]);
428 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[6]);
429 switch (planned.artifact.payload) {
430 .bytes => |bytes| try testing.expectEqualSlices(u8, &.{ 0x00, 0x61, 0x73, 0x6d, 0x01, 0x00, 0x00, 0x00 }, bytes[0..8]),
431 else => return error.ExpectedWebAssemblyModuleBytes,
432 }
433 }
434
435 test "kernel artifact plan runs native cpu machine code for authored global x kernel" {
436 try runAuthoredCpuCopyKernel(testing.allocator, .cpu_machine_code);
437 }
438
439 test "kernel artifact plan runs native cpu object for authored global x kernel" {
440 try runAuthoredCpuCopyKernel(testing.allocator, .cpu_object);
441 }
442
443 test "kernel artifact plan runs native cpu machine code for automatic vectorized scalar add kernel" {
444 try runAutomaticScalarAddKernel(
445 testing.allocator,
446 .cpu_machine_code,
447 f32,
448 .f32,
449 "automatic_cpu_vectorized_scalar_add_f32",
450 4,
451 [_]f32{ 1.25, -2.5, 3.75, 8.0 },
452 [_]f32{ 10.0, 4.0, -0.75, -9.0 },
453 );
454 }
455
456 test "kernel artifact plan runs native cpu object for automatic vectorized scalar add kernel" {
457 try runAutomaticScalarAddKernel(
458 testing.allocator,
459 .cpu_object,
460 f32,
461 .f32,
462 "automatic_cpu_vectorized_scalar_add_f32",
463 4,
464 [_]f32{ 1.25, -2.5, 3.75, 8.0 },
465 [_]f32{ 10.0, 4.0, -0.75, -9.0 },
466 );
467 }
468
469 test "kernel artifact plan runs native cpu machine code for automatic scalar affine tail kernel" {
470 try runAutomaticScalarAffineKernel(
471 testing.allocator,
472 .cpu_machine_code,
473 "automatic_cpu_vectorized_scalar_affine_tail_f32",
474 [_]f32{ 1.0, -2.0, 3.5, 8.0, -4.0, 0.5 },
475 [_]f32{ 10.0, 4.0, -0.5, -9.0, 2.0, 6.0 },
476 2.0,
477 );
478 }
479
480 test "kernel artifact plan runs native cpu object for automatic scalar affine tail kernel" {
481 try runAutomaticScalarAffineKernel(
482 testing.allocator,
483 .cpu_object,
484 "automatic_cpu_vectorized_scalar_affine_tail_f32",
485 [_]f32{ 1.0, -2.0, 3.5, 8.0, -4.0, 0.5 },
486 [_]f32{ 10.0, 4.0, -0.5, -9.0, 2.0, 6.0 },
487 2.0,
488 );
489 }
490
491 test "kernel artifact plan runs native cpu machine code for automatic f64 scalar add kernel" {
492 try runAutomaticScalarAddKernel(
493 testing.allocator,
494 .cpu_machine_code,
495 f64,
496 .f64,
497 "automatic_cpu_scalar_add_f64",
498 4,
499 [_]f64{ 1.25, -2.5, 3.75, 8.0 },
500 [_]f64{ 10.0, 4.0, -0.75, -9.0 },
501 );
502 }
503
504 test "kernel artifact plan runs native cpu object for automatic f64 scalar add kernel" {
505 try runAutomaticScalarAddKernel(
506 testing.allocator,
507 .cpu_object,
508 f64,
509 .f64,
510 "automatic_cpu_scalar_add_f64",
511 4,
512 [_]f64{ 1.25, -2.5, 3.75, 8.0 },
513 [_]f64{ 10.0, 4.0, -0.75, -9.0 },
514 );
515 }
516
517 test "kernel artifact plan runs native cpu machine code for authored global y kernel" {
518 try runAuthoredCpuRowKernel(testing.allocator, .cpu_machine_code);
519 }
520
521 test "kernel artifact plan runs native cpu object for authored global y kernel" {
522 try runAuthoredCpuRowKernel(testing.allocator, .cpu_object);
523 }
524
525 test "kernel artifact plan runs native cpu machine code for authored thread block kernel" {
526 try runAuthoredCpuThreadBlockKernel(testing.allocator, .cpu_machine_code);
527 }
528
529 test "kernel artifact plan runs native cpu object for authored thread block kernel" {
530 try runAuthoredCpuThreadBlockKernel(testing.allocator, .cpu_object);
531 }
532
533 test "kernel artifact plan runs native cpu machine code for authored vector add kernel" {
534 try runAuthoredCpuVectorBinaryKernel(
535 testing.allocator,
536 .cpu_machine_code,
537 .add,
538 f32,
539 .f32,
540 "authored_cpu_vec4_add_f32",
541 8,
542 4,
543 2,
544 [_]f32{ 1.0, -2.5, 3.25, 10.0, 0.5, 8.0, -3.0, 2.25 },
545 [_]f32{ 4.0, 2.0, -1.25, -11.5, 6.5, -1.0, 7.0, -2.25 },
546 [_]f32{ 5.0, -0.5, 2.0, -1.5, 7.0, 7.0, 4.0, 0.0 },
547 );
548 }
549
550 test "kernel artifact plan runs native cpu object for authored vector add kernel" {
551 try runAuthoredCpuVectorBinaryKernel(
552 testing.allocator,
553 .cpu_object,
554 .add,
555 f32,
556 .f32,
557 "authored_cpu_vec4_add_f32",
558 8,
559 4,
560 2,
561 [_]f32{ 1.0, -2.5, 3.25, 10.0, 0.5, 8.0, -3.0, 2.25 },
562 [_]f32{ 4.0, 2.0, -1.25, -11.5, 6.5, -1.0, 7.0, -2.25 },
563 [_]f32{ 5.0, -0.5, 2.0, -1.5, 7.0, 7.0, 4.0, 0.0 },
564 );
565 }
566
567 test "kernel artifact plan runs native cpu machine code for authored f32 vector min kernel" {
568 try runAuthoredCpuVectorBinaryKernel(
569 testing.allocator,
570 .cpu_machine_code,
571 .min,
572 f32,
573 .f32,
574 "authored_cpu_vec4_min_f32",
575 8,
576 4,
577 2,
578 [_]f32{ std.math.nan(f32), 5.0, -8.0, -5.0, 1.0, -2.0, 3.0, 4.0 },
579 [_]f32{ 5.0, std.math.nan(f32), -1.0, -3.0, 2.0, -3.0, 2.0, 10.0 },
580 [_]f32{ 5.0, 5.0, -8.0, -5.0, 1.0, -3.0, 2.0, 4.0 },
581 );
582 }
583
584 test "kernel artifact plan runs native cpu object for authored f32 vector min kernel" {
585 try runAuthoredCpuVectorBinaryKernel(
586 testing.allocator,
587 .cpu_object,
588 .min,
589 f32,
590 .f32,
591 "authored_cpu_vec4_min_f32",
592 8,
593 4,
594 2,
595 [_]f32{ std.math.nan(f32), 5.0, -8.0, -5.0, 1.0, -2.0, 3.0, 4.0 },
596 [_]f32{ 5.0, std.math.nan(f32), -1.0, -3.0, 2.0, -3.0, 2.0, 10.0 },
597 [_]f32{ 5.0, 5.0, -8.0, -5.0, 1.0, -3.0, 2.0, 4.0 },
598 );
599 }
600
601 test "kernel artifact plan runs native cpu machine code for authored u32 vector max kernel" {
602 try runAuthoredCpuVectorBinaryKernel(
603 testing.allocator,
604 .cpu_machine_code,
605 .max,
606 u32,
607 .u32,
608 "authored_cpu_vec4_max_u32",
609 8,
610 4,
611 2,
612 [_]u32{ 0x8000_0000, 1, 0xffff_ffff, 7, 3, 4, 0, 10 },
613 [_]u32{ 1, 0x8000_0000, 2, 0xffff_fffe, 9, 2, 0xffff_ffff, 5 },
614 [_]u32{ 0x8000_0000, 0x8000_0000, 0xffff_ffff, 0xffff_fffe, 9, 4, 0xffff_ffff, 10 },
615 );
616 }
617
618 test "kernel artifact plan runs native cpu object for authored u32 vector max kernel" {
619 try runAuthoredCpuVectorBinaryKernel(
620 testing.allocator,
621 .cpu_object,
622 .max,
623 u32,
624 .u32,
625 "authored_cpu_vec4_max_u32",
626 8,
627 4,
628 2,
629 [_]u32{ 0x8000_0000, 1, 0xffff_ffff, 7, 3, 4, 0, 10 },
630 [_]u32{ 1, 0x8000_0000, 2, 0xffff_fffe, 9, 2, 0xffff_ffff, 5 },
631 [_]u32{ 0x8000_0000, 0x8000_0000, 0xffff_ffff, 0xffff_fffe, 9, 4, 0xffff_ffff, 10 },
632 );
633 }
634
635 test "kernel artifact plan runs native cpu machine code for authored u32 vector xor kernel" {
636 try runAuthoredCpuVectorBinaryKernel(
637 testing.allocator,
638 .cpu_machine_code,
639 .bxor,
640 u32,
641 .u32,
642 "authored_cpu_vec4_xor_u32",
643 8,
644 4,
645 2,
646 [_]u32{ 0, 1, 0xaaaa_aaaa, 0xffff_0000, 0x8000_0000, 0xffff_ffff, 7, 0x1234_5678 },
647 [_]u32{ 0xffff_ffff, 1, 0x5555_5555, 0x00ff_ff00, 0x7fff_ffff, 0, 3, 0x8765_4321 },
648 [_]u32{ 0xffff_ffff, 0, 0xffff_ffff, 0xff00_ff00, 0xffff_ffff, 0xffff_ffff, 4, 0x9551_1559 },
649 );
650 }
651
652 test "kernel artifact plan runs native cpu object for authored u32 vector xor kernel" {
653 try runAuthoredCpuVectorBinaryKernel(
654 testing.allocator,
655 .cpu_object,
656 .bxor,
657 u32,
658 .u32,
659 "authored_cpu_vec4_xor_u32",
660 8,
661 4,
662 2,
663 [_]u32{ 0, 1, 0xaaaa_aaaa, 0xffff_0000, 0x8000_0000, 0xffff_ffff, 7, 0x1234_5678 },
664 [_]u32{ 0xffff_ffff, 1, 0x5555_5555, 0x00ff_ff00, 0x7fff_ffff, 0, 3, 0x8765_4321 },
665 [_]u32{ 0xffff_ffff, 0, 0xffff_ffff, 0xff00_ff00, 0xffff_ffff, 0xffff_ffff, 4, 0x9551_1559 },
666 );
667 }
668
669 test "kernel artifact plan runs native cpu machine code for authored u32 vector select true kernel" {
670 try runAuthoredCpuVectorBinaryKernel(
671 testing.allocator,
672 .cpu_machine_code,
673 .select_true,
674 u32,
675 .u32,
676 "authored_cpu_vec4_select_true_u32",
677 8,
678 4,
679 2,
680 [_]u32{ 1, 2, 3, 4, 0x8000_0000, 0xffff_ffff, 7, 8 },
681 [_]u32{ 10, 20, 30, 40, 50, 60, 70, 80 },
682 [_]u32{ 1, 2, 3, 4, 0x8000_0000, 0xffff_ffff, 7, 8 },
683 );
684 }
685
686 test "kernel artifact plan runs native cpu object for authored u32 vector select true kernel" {
687 try runAuthoredCpuVectorBinaryKernel(
688 testing.allocator,
689 .cpu_object,
690 .select_true,
691 u32,
692 .u32,
693 "authored_cpu_vec4_select_true_u32",
694 8,
695 4,
696 2,
697 [_]u32{ 1, 2, 3, 4, 0x8000_0000, 0xffff_ffff, 7, 8 },
698 [_]u32{ 10, 20, 30, 40, 50, 60, 70, 80 },
699 [_]u32{ 1, 2, 3, 4, 0x8000_0000, 0xffff_ffff, 7, 8 },
700 );
701 }
702
703 test "kernel artifact plan runs native cpu machine code for authored u32 vector select false kernel" {
704 try runAuthoredCpuVectorBinaryKernel(
705 testing.allocator,
706 .cpu_machine_code,
707 .select_false,
708 u32,
709 .u32,
710 "authored_cpu_vec4_select_false_u32",
711 8,
712 4,
713 2,
714 [_]u32{ 1, 2, 3, 4, 0x8000_0000, 0xffff_ffff, 7, 8 },
715 [_]u32{ 10, 20, 30, 40, 50, 60, 70, 80 },
716 [_]u32{ 10, 20, 30, 40, 50, 60, 70, 80 },
717 );
718 }
719
720 test "kernel artifact plan runs native cpu object for authored u32 vector select false kernel" {
721 try runAuthoredCpuVectorBinaryKernel(
722 testing.allocator,
723 .cpu_object,
724 .select_false,
725 u32,
726 .u32,
727 "authored_cpu_vec4_select_false_u32",
728 8,
729 4,
730 2,
731 [_]u32{ 1, 2, 3, 4, 0x8000_0000, 0xffff_ffff, 7, 8 },
732 [_]u32{ 10, 20, 30, 40, 50, 60, 70, 80 },
733 [_]u32{ 10, 20, 30, 40, 50, 60, 70, 80 },
734 );
735 }
736
737 test "kernel artifact plan runs native cpu machine code for authored f64 vector add kernel" {
738 try runAuthoredCpuVectorBinaryKernel(
739 testing.allocator,
740 .cpu_machine_code,
741 .add,
742 f64,
743 .f64,
744 "authored_cpu_vec2_add_f64",
745 4,
746 2,
747 2,
748 [_]f64{ 1.0, -2.5, 3.25, 10.0 },
749 [_]f64{ 4.0, 2.0, -1.25, -11.5 },
750 [_]f64{ 5.0, -0.5, 2.0, -1.5 },
751 );
752 }
753
754 test "kernel artifact plan runs native cpu object for authored f64 vector add kernel" {
755 try runAuthoredCpuVectorBinaryKernel(
756 testing.allocator,
757 .cpu_object,
758 .add,
759 f64,
760 .f64,
761 "authored_cpu_vec2_add_f64",
762 4,
763 2,
764 2,
765 [_]f64{ 1.0, -2.5, 3.25, 10.0 },
766 [_]f64{ 4.0, 2.0, -1.25, -11.5 },
767 [_]f64{ 5.0, -0.5, 2.0, -1.5 },
768 );
769 }
770
771 test "kernel artifact plan runs native cpu machine code for authored f64 vector mul kernel" {
772 try runAuthoredCpuVectorBinaryKernel(
773 testing.allocator,
774 .cpu_machine_code,
775 .mul,
776 f64,
777 .f64,
778 "authored_cpu_vec2_mul_f64",
779 4,
780 2,
781 2,
782 [_]f64{ 1.5, -2.0, 3.25, -4.0 },
783 [_]f64{ 2.0, 4.0, -2.0, -0.25 },
784 [_]f64{ 3.0, -8.0, -6.5, 1.0 },
785 );
786 }
787
788 test "kernel artifact plan runs native cpu object for authored f64 vector mul kernel" {
789 try runAuthoredCpuVectorBinaryKernel(
790 testing.allocator,
791 .cpu_object,
792 .mul,
793 f64,
794 .f64,
795 "authored_cpu_vec2_mul_f64",
796 4,
797 2,
798 2,
799 [_]f64{ 1.5, -2.0, 3.25, -4.0 },
800 [_]f64{ 2.0, 4.0, -2.0, -0.25 },
801 [_]f64{ 3.0, -8.0, -6.5, 1.0 },
802 );
803 }
804
805 test "kernel artifact plan runs native cpu machine code for authored f32 vector mul kernel" {
806 try runAuthoredCpuVectorBinaryKernel(
807 testing.allocator,
808 .cpu_machine_code,
809 .mul,
810 f32,
811 .f32,
812 "authored_cpu_vec4_mul_f32",
813 8,
814 4,
815 2,
816 [_]f32{ 1.5, -2.0, 3.0, -4.0, 0.5, 8.0, -3.0, 2.25 },
817 [_]f32{ 2.0, 4.0, -1.5, -0.25, 6.0, -1.0, 7.0, -4.0 },
818 [_]f32{ 3.0, -8.0, -4.5, 1.0, 3.0, -8.0, -21.0, -9.0 },
819 );
820 }
821
822 test "kernel artifact plan runs native cpu object for authored f32 vector mul kernel" {
823 try runAuthoredCpuVectorBinaryKernel(
824 testing.allocator,
825 .cpu_object,
826 .mul,
827 f32,
828 .f32,
829 "authored_cpu_vec4_mul_f32",
830 8,
831 4,
832 2,
833 [_]f32{ 1.5, -2.0, 3.0, -4.0, 0.5, 8.0, -3.0, 2.25 },
834 [_]f32{ 2.0, 4.0, -1.5, -0.25, 6.0, -1.0, 7.0, -4.0 },
835 [_]f32{ 3.0, -8.0, -4.5, 1.0, 3.0, -8.0, -21.0, -9.0 },
836 );
837 }
838
839 test "kernel artifact plan runs native cpu machine code for authored f32 vector div kernel" {
840 try runAuthoredCpuVectorBinaryKernel(
841 testing.allocator,
842 .cpu_machine_code,
843 .div,
844 f32,
845 .f32,
846 "authored_cpu_vec4_div_f32",
847 8,
848 4,
849 2,
850 [_]f32{ 8.0, -9.0, 7.5, -12.0, 4.0, -6.0, 0.5, -2.25 },
851 [_]f32{ 2.0, 3.0, 2.5, -4.0, -2.0, -1.5, 0.25, 0.75 },
852 [_]f32{ 4.0, -3.0, 3.0, 3.0, -2.0, 4.0, 2.0, -3.0 },
853 );
854 }
855
856 test "kernel artifact plan runs native cpu object for authored f32 vector div kernel" {
857 try runAuthoredCpuVectorBinaryKernel(
858 testing.allocator,
859 .cpu_object,
860 .div,
861 f32,
862 .f32,
863 "authored_cpu_vec4_div_f32",
864 8,
865 4,
866 2,
867 [_]f32{ 8.0, -9.0, 7.5, -12.0, 4.0, -6.0, 0.5, -2.25 },
868 [_]f32{ 2.0, 3.0, 2.5, -4.0, -2.0, -1.5, 0.25, 0.75 },
869 [_]f32{ 4.0, -3.0, 3.0, 3.0, -2.0, 4.0, 2.0, -3.0 },
870 );
871 }
872
873 test "kernel artifact plan runs native cpu object for authored f32 floor kernel" {
874 try runAuthoredCpuFloatUnaryKernel(
875 testing.allocator,
876 .cpu_object,
877 .floor,
878 "authored_cpu_floor_f32",
879 [_]f32{ 1.75, -1.25, 0.0, 8.0 },
880 [_]f32{ 1.0, -2.0, 0.0, 8.0 },
881 );
882 }
883
884 test "kernel artifact plan runs native cpu machine code for authored f64 vector div kernel" {
885 try runAuthoredCpuVectorBinaryKernel(
886 testing.allocator,
887 .cpu_machine_code,
888 .div,
889 f64,
890 .f64,
891 "authored_cpu_vec2_div_f64",
892 4,
893 2,
894 2,
895 [_]f64{ 8.0, -9.0, 7.5, -12.0 },
896 [_]f64{ 2.0, 3.0, 2.5, -4.0 },
897 [_]f64{ 4.0, -3.0, 3.0, 3.0 },
898 );
899 }
900
901 test "kernel artifact plan runs native cpu object for authored f64 vector div kernel" {
902 try runAuthoredCpuVectorBinaryKernel(
903 testing.allocator,
904 .cpu_object,
905 .div,
906 f64,
907 .f64,
908 "authored_cpu_vec2_div_f64",
909 4,
910 2,
911 2,
912 [_]f64{ 8.0, -9.0, 7.5, -12.0 },
913 [_]f64{ 2.0, 3.0, 2.5, -4.0 },
914 [_]f64{ 4.0, -3.0, 3.0, 3.0 },
915 );
916 }
917
918 test "kernel artifact plan runs native cpu machine code for authored u32 vector add kernel" {
919 try runAuthoredCpuVectorBinaryKernel(
920 testing.allocator,
921 .cpu_machine_code,
922 .add,
923 u32,
924 .u32,
925 "authored_cpu_vec4_add_u32",
926 8,
927 4,
928 2,
929 [_]u32{ 1, 0x8000_0000, 0xffff_ffff, 7, 3, 4, 5, 6 },
930 [_]u32{ 4, 0x8000_0000, 2, 0xffff_fffe, 7, 8, 9, 10 },
931 [_]u32{ 5, 0, 1, 5, 10, 12, 14, 16 },
932 );
933 }
934
935 test "kernel artifact plan runs native cpu object for authored u32 vector add kernel" {
936 try runAuthoredCpuVectorBinaryKernel(
937 testing.allocator,
938 .cpu_object,
939 .add,
940 u32,
941 .u32,
942 "authored_cpu_vec4_add_u32",
943 8,
944 4,
945 2,
946 [_]u32{ 1, 0x8000_0000, 0xffff_ffff, 7, 3, 4, 5, 6 },
947 [_]u32{ 4, 0x8000_0000, 2, 0xffff_fffe, 7, 8, 9, 10 },
948 [_]u32{ 5, 0, 1, 5, 10, 12, 14, 16 },
949 );
950 }
951
952 test "kernel artifact plan runs native cpu machine code for authored u32 vector mul kernel" {
953 try runAuthoredCpuVectorBinaryKernel(
954 testing.allocator,
955 .cpu_machine_code,
956 .mul,
957 u32,
958 .u32,
959 "authored_cpu_vec4_mul_u32",
960 8,
961 4,
962 2,
963 [_]u32{ 1, 0x8000_0000, 0xffff_ffff, 7, 3, 4, 0x0001_0000, 0x0000_ffff },
964 [_]u32{ 4, 2, 2, 0xffff_ffff, 7, 0x4000_0000, 0x0001_0000, 0x0001_0001 },
965 [_]u32{ 4, 0, 0xffff_fffe, 0xffff_fff9, 21, 0, 0, 0xffff_ffff },
966 );
967 }
968
969 test "kernel artifact plan runs native cpu object for authored u32 vector mul kernel" {
970 try runAuthoredCpuVectorBinaryKernel(
971 testing.allocator,
972 .cpu_object,
973 .mul,
974 u32,
975 .u32,
976 "authored_cpu_vec4_mul_u32",
977 8,
978 4,
979 2,
980 [_]u32{ 1, 0x8000_0000, 0xffff_ffff, 7, 3, 4, 0x0001_0000, 0x0000_ffff },
981 [_]u32{ 4, 2, 2, 0xffff_ffff, 7, 0x4000_0000, 0x0001_0000, 0x0001_0001 },
982 [_]u32{ 4, 0, 0xffff_fffe, 0xffff_fff9, 21, 0, 0, 0xffff_ffff },
983 );
984 }
985
986 test "kernel artifact plan runs native cpu machine code for authored u32 vector umulhi kernel" {
987 try runAuthoredCpuVectorBinaryKernel(
988 testing.allocator,
989 .cpu_machine_code,
990 .umulhi,
991 u32,
992 .u32,
993 "authored_cpu_vec4_umulhi_u32",
994 8,
995 4,
996 2,
997 [_]u32{ 0xffff_ffff, 0x8000_0000, 0x0001_0000, 123456789, 0, 0xffff_ffff, 0x4000_0000, 0xffff_0000 },
998 [_]u32{ 2, 4, 0x0001_0000, 987654321, 9, 0xffff_ffff, 8, 0x0001_0000 },
999 [_]u32{ 1, 2, 1, 28389652, 0, 0xffff_fffe, 2, 65535 },
1000 );
1001 }
1002
1003 test "kernel artifact plan runs native cpu object for authored u32 vector umulhi kernel" {
1004 try runAuthoredCpuVectorBinaryKernel(
1005 testing.allocator,
1006 .cpu_object,
1007 .umulhi,
1008 u32,
1009 .u32,
1010 "authored_cpu_vec4_umulhi_u32",
1011 8,
1012 4,
1013 2,
1014 [_]u32{ 0xffff_ffff, 0x8000_0000, 0x0001_0000, 123456789, 0, 0xffff_ffff, 0x4000_0000, 0xffff_0000 },
1015 [_]u32{ 2, 4, 0x0001_0000, 987654321, 9, 0xffff_ffff, 8, 0x0001_0000 },
1016 [_]u32{ 1, 2, 1, 28389652, 0, 0xffff_fffe, 2, 65535 },
1017 );
1018 }
1019
1020 test "kernel artifact plan runs native cpu machine code for authored u32 vector popcount kernel" {
1021 try runAuthoredCpuVectorBinaryKernel(
1022 testing.allocator,
1023 .cpu_machine_code,
1024 .popcount,
1025 u32,
1026 .u32,
1027 "authored_cpu_vec4_popcount_u32",
1028 8,
1029 4,
1030 2,
1031 [_]u32{ 0, 1, 0xffff_ffff, 0xf0f0_00ff, 0x8000_0000, 0x5555_5555, 0x1234_5678, 0xffff_0000 },
1032 [_]u32{ 0, 0, 0, 0, 0, 0, 0, 0 },
1033 [_]u32{ 0, 1, 32, 16, 1, 16, 13, 16 },
1034 );
1035 }
1036
1037 test "kernel artifact plan runs native cpu object for authored u32 vector popcount kernel" {
1038 try runAuthoredCpuVectorBinaryKernel(
1039 testing.allocator,
1040 .cpu_object,
1041 .popcount,
1042 u32,
1043 .u32,
1044 "authored_cpu_vec4_popcount_u32",
1045 8,
1046 4,
1047 2,
1048 [_]u32{ 0, 1, 0xffff_ffff, 0xf0f0_00ff, 0x8000_0000, 0x5555_5555, 0x1234_5678, 0xffff_0000 },
1049 [_]u32{ 0, 0, 0, 0, 0, 0, 0, 0 },
1050 [_]u32{ 0, 1, 32, 16, 1, 16, 13, 16 },
1051 );
1052 }
1053
1054 test "kernel artifact plan runs native cpu machine code for authored u32 vector div kernel" {
1055 try runAuthoredCpuVectorBinaryKernel(
1056 testing.allocator,
1057 .cpu_machine_code,
1058 .div,
1059 u32,
1060 .u32,
1061 "authored_cpu_vec4_div_u32",
1062 8,
1063 4,
1064 2,
1065 [_]u32{ 0x8000_0000, 0xffff_ffff, 21, 100, 7, 0, 0x7fff_ffff, 0x4000_0000 },
1066 [_]u32{ 2, 2, 5, 4, 3, 1, 3, 0x10 },
1067 [_]u32{ 0x4000_0000, 0x7fff_ffff, 4, 25, 2, 0, 0x2aaa_aaaa, 0x0400_0000 },
1068 );
1069 }
1070
1071 test "kernel artifact plan runs native cpu object for authored u32 vector div kernel" {
1072 try runAuthoredCpuVectorBinaryKernel(
1073 testing.allocator,
1074 .cpu_object,
1075 .div,
1076 u32,
1077 .u32,
1078 "authored_cpu_vec4_div_u32",
1079 8,
1080 4,
1081 2,
1082 [_]u32{ 0x8000_0000, 0xffff_ffff, 21, 100, 7, 0, 0x7fff_ffff, 0x4000_0000 },
1083 [_]u32{ 2, 2, 5, 4, 3, 1, 3, 0x10 },
1084 [_]u32{ 0x4000_0000, 0x7fff_ffff, 4, 25, 2, 0, 0x2aaa_aaaa, 0x0400_0000 },
1085 );
1086 }
1087
1088 test "kernel artifact plan runs native cpu machine code for authored u32 vector shl kernel" {
1089 try runAuthoredCpuVectorBinaryKernel(
1090 testing.allocator,
1091 .cpu_machine_code,
1092 .shl,
1093 u32,
1094 .u32,
1095 "authored_cpu_vec4_shl_u32",
1096 8,
1097 4,
1098 2,
1099 [_]u32{ 1, 0x8000_0000, 0x0000_ffff, 0xffff_ffff, 3, 4, 5, 6 },
1100 [_]u32{ 0, 1, 4, 8, 2, 3, 4, 5 },
1101 [_]u32{ 1, 0, 0x000f_fff0, 0xffff_ff00, 12, 32, 80, 192 },
1102 );
1103 }
1104
1105 test "kernel artifact plan runs native cpu object for authored u32 vector shl kernel" {
1106 try runAuthoredCpuVectorBinaryKernel(
1107 testing.allocator,
1108 .cpu_object,
1109 .shl,
1110 u32,
1111 .u32,
1112 "authored_cpu_vec4_shl_u32",
1113 8,
1114 4,
1115 2,
1116 [_]u32{ 1, 0x8000_0000, 0x0000_ffff, 0xffff_ffff, 3, 4, 5, 6 },
1117 [_]u32{ 0, 1, 4, 8, 2, 3, 4, 5 },
1118 [_]u32{ 1, 0, 0x000f_fff0, 0xffff_ff00, 12, 32, 80, 192 },
1119 );
1120 }
1121
1122 test "kernel artifact plan runs native cpu machine code for authored i32 vector add kernel" {
1123 try runAuthoredCpuVectorBinaryKernel(
1124 testing.allocator,
1125 .cpu_machine_code,
1126 .add,
1127 i32,
1128 .i32,
1129 "authored_cpu_vec4_add_i32",
1130 8,
1131 4,
1132 2,
1133 [_]i32{ 1, -2, 0x7fff_ffff, -2147483648, -5, 6, 123, -456 },
1134 [_]i32{ -3, -4, 1, -1, 9, -10, -200, 300 },
1135 [_]i32{ -2, -6, -2147483648, 2147483647, 4, -4, -77, -156 },
1136 );
1137 }
1138
1139 test "kernel artifact plan runs native cpu object for authored i32 vector add kernel" {
1140 try runAuthoredCpuVectorBinaryKernel(
1141 testing.allocator,
1142 .cpu_object,
1143 .add,
1144 i32,
1145 .i32,
1146 "authored_cpu_vec4_add_i32",
1147 8,
1148 4,
1149 2,
1150 [_]i32{ 1, -2, 0x7fff_ffff, -2147483648, -5, 6, 123, -456 },
1151 [_]i32{ -3, -4, 1, -1, 9, -10, -200, 300 },
1152 [_]i32{ -2, -6, -2147483648, 2147483647, 4, -4, -77, -156 },
1153 );
1154 }
1155
1156 test "kernel artifact plan runs native cpu machine code for authored i32 vector mul kernel" {
1157 try runAuthoredCpuVectorBinaryKernel(
1158 testing.allocator,
1159 .cpu_machine_code,
1160 .mul,
1161 i32,
1162 .i32,
1163 "authored_cpu_vec4_mul_i32",
1164 8,
1165 4,
1166 2,
1167 [_]i32{ 2, -3, 0x4000_0000, std.math.minInt(i32), -1, 65536, -65536, 12345 },
1168 [_]i32{ -4, -5, 4, 2, -1, 65536, 65536, -2 },
1169 [_]i32{ -8, 15, 0, 0, 1, 0, 0, -24690 },
1170 );
1171 }
1172
1173 test "kernel artifact plan runs native cpu object for authored i32 vector mul kernel" {
1174 try runAuthoredCpuVectorBinaryKernel(
1175 testing.allocator,
1176 .cpu_object,
1177 .mul,
1178 i32,
1179 .i32,
1180 "authored_cpu_vec4_mul_i32",
1181 8,
1182 4,
1183 2,
1184 [_]i32{ 2, -3, 0x4000_0000, std.math.minInt(i32), -1, 65536, -65536, 12345 },
1185 [_]i32{ -4, -5, 4, 2, -1, 65536, 65536, -2 },
1186 [_]i32{ -8, 15, 0, 0, 1, 0, 0, -24690 },
1187 );
1188 }
1189
1190 test "kernel artifact plan runs native cpu machine code for authored i32 vector div kernel" {
1191 try runAuthoredCpuVectorBinaryKernel(
1192 testing.allocator,
1193 .cpu_machine_code,
1194 .div,
1195 i32,
1196 .i32,
1197 "authored_cpu_vec4_div_i32",
1198 8,
1199 4,
1200 2,
1201 [_]i32{ 21, -21, -18, 100, -100, 7, 2_147_483_646, -1024 },
1202 [_]i32{ 5, 5, 5, -4, 9, -2, 2, 8 },
1203 [_]i32{ 4, -4, -3, -25, -11, -3, 1_073_741_823, -128 },
1204 );
1205 }
1206
1207 test "kernel artifact plan runs native cpu object for authored i32 vector div kernel" {
1208 try runAuthoredCpuVectorBinaryKernel(
1209 testing.allocator,
1210 .cpu_object,
1211 .div,
1212 i32,
1213 .i32,
1214 "authored_cpu_vec4_div_i32",
1215 8,
1216 4,
1217 2,
1218 [_]i32{ 21, -21, -18, 100, -100, 7, 2_147_483_646, -1024 },
1219 [_]i32{ 5, 5, 5, -4, 9, -2, 2, 8 },
1220 [_]i32{ 4, -4, -3, -25, -11, -3, 1_073_741_823, -128 },
1221 );
1222 }
1223
1224 test "kernel artifact plan runs native cpu machine code for authored i32 vector shr kernel" {
1225 try runAuthoredCpuVectorBinaryKernel(
1226 testing.allocator,
1227 .cpu_machine_code,
1228 .shr,
1229 i32,
1230 .i32,
1231 "authored_cpu_vec4_shr_i32",
1232 8,
1233 4,
1234 2,
1235 [_]i32{ -16, -1, 1024, std.math.minInt(i32), 7, -128, 123456, -123456 },
1236 [_]i32{ 2, 1, 5, 31, 0, 3, 4, 4 },
1237 [_]i32{ -4, -1, 32, -1, 7, -16, 7716, -7716 },
1238 );
1239 }
1240
1241 test "kernel artifact plan runs native cpu object for authored i32 vector shr kernel" {
1242 try runAuthoredCpuVectorBinaryKernel(
1243 testing.allocator,
1244 .cpu_object,
1245 .shr,
1246 i32,
1247 .i32,
1248 "authored_cpu_vec4_shr_i32",
1249 8,
1250 4,
1251 2,
1252 [_]i32{ -16, -1, 1024, std.math.minInt(i32), 7, -128, 123456, -123456 },
1253 [_]i32{ 2, 1, 5, 31, 0, 3, 4, 4 },
1254 [_]i32{ -4, -1, 32, -1, 7, -16, 7716, -7716 },
1255 );
1256 }
1257
1258 test "kernel artifact plan runs native cpu machine code for authored i32 vector sub kernel" {
1259 try runAuthoredCpuVectorBinaryKernel(
1260 testing.allocator,
1261 .cpu_machine_code,
1262 .sub,
1263 i32,
1264 .i32,
1265 "authored_cpu_vec4_sub_i32",
1266 8,
1267 4,
1268 2,
1269 [_]i32{ 1, -2, -2147483648, 2147483647, 9, -10, 123, -456 },
1270 [_]i32{ 3, -4, 1, -1, -5, 6, -200, 300 },
1271 [_]i32{ -2, 2, 2147483647, -2147483648, 14, -16, 323, -756 },
1272 );
1273 }
1274
1275 test "kernel artifact plan runs native cpu object for authored i32 vector sub kernel" {
1276 try runAuthoredCpuVectorBinaryKernel(
1277 testing.allocator,
1278 .cpu_object,
1279 .sub,
1280 i32,
1281 .i32,
1282 "authored_cpu_vec4_sub_i32",
1283 8,
1284 4,
1285 2,
1286 [_]i32{ 1, -2, -2147483648, 2147483647, 9, -10, 123, -456 },
1287 [_]i32{ 3, -4, 1, -1, -5, 6, -200, 300 },
1288 [_]i32{ -2, 2, 2147483647, -2147483648, 14, -16, 323, -756 },
1289 );
1290 }
1291
1292 test "kernel artifact plan runs native cpu machine code for authored i64 vector add kernel" {
1293 try runAuthoredCpuVectorBinaryKernel(
1294 testing.allocator,
1295 .cpu_machine_code,
1296 .add,
1297 i64,
1298 .i64,
1299 "authored_cpu_vec2_add_i64",
1300 4,
1301 2,
1302 2,
1303 [_]i64{ 1, -2, std.math.maxInt(i64), std.math.minInt(i64) },
1304 [_]i64{ -3, -4, 1, -1 },
1305 [_]i64{ -2, -6, std.math.minInt(i64), std.math.maxInt(i64) },
1306 );
1307 }
1308
1309 test "kernel artifact plan runs native cpu object for authored i64 vector add kernel" {
1310 try runAuthoredCpuVectorBinaryKernel(
1311 testing.allocator,
1312 .cpu_object,
1313 .add,
1314 i64,
1315 .i64,
1316 "authored_cpu_vec2_add_i64",
1317 4,
1318 2,
1319 2,
1320 [_]i64{ 1, -2, std.math.maxInt(i64), std.math.minInt(i64) },
1321 [_]i64{ -3, -4, 1, -1 },
1322 [_]i64{ -2, -6, std.math.minInt(i64), std.math.maxInt(i64) },
1323 );
1324 }
1325
1326 test "kernel artifact plan runs native cpu machine code for authored i64 vector mul kernel" {
1327 try runAuthoredCpuVectorBinaryKernel(
1328 testing.allocator,
1329 .cpu_machine_code,
1330 .mul,
1331 i64,
1332 .i64,
1333 "authored_cpu_vec2_mul_i64",
1334 4,
1335 2,
1336 2,
1337 [_]i64{ 0x0000_0001_0000_0000, -3, std.math.maxInt(i64), std.math.minInt(i64) },
1338 [_]i64{ 3, 7, 2, -1 },
1339 [_]i64{ 0x0000_0003_0000_0000, -21, -2, std.math.minInt(i64) },
1340 );
1341 }
1342
1343 test "kernel artifact plan runs native cpu object for authored i64 vector mul kernel" {
1344 try runAuthoredCpuVectorBinaryKernel(
1345 testing.allocator,
1346 .cpu_object,
1347 .mul,
1348 i64,
1349 .i64,
1350 "authored_cpu_vec2_mul_i64",
1351 4,
1352 2,
1353 2,
1354 [_]i64{ 0x0000_0001_0000_0000, -3, std.math.maxInt(i64), std.math.minInt(i64) },
1355 [_]i64{ 3, 7, 2, -1 },
1356 [_]i64{ 0x0000_0003_0000_0000, -21, -2, std.math.minInt(i64) },
1357 );
1358 }
1359
1360 test "kernel artifact plan runs native cpu machine code for authored i64 vector div kernel" {
1361 try runAuthoredCpuVectorBinaryKernel(
1362 testing.allocator,
1363 .cpu_machine_code,
1364 .div,
1365 i64,
1366 .i64,
1367 "authored_cpu_vec2_div_i64",
1368 4,
1369 2,
1370 2,
1371 [_]i64{ 21, -21, 9_000_000_000, -9_000_000_000 },
1372 [_]i64{ 5, 5, 3, 4 },
1373 [_]i64{ 4, -4, 3_000_000_000, -2_250_000_000 },
1374 );
1375 }
1376
1377 test "kernel artifact plan runs native cpu object for authored i64 vector div kernel" {
1378 try runAuthoredCpuVectorBinaryKernel(
1379 testing.allocator,
1380 .cpu_object,
1381 .div,
1382 i64,
1383 .i64,
1384 "authored_cpu_vec2_div_i64",
1385 4,
1386 2,
1387 2,
1388 [_]i64{ 21, -21, 9_000_000_000, -9_000_000_000 },
1389 [_]i64{ 5, 5, 3, 4 },
1390 [_]i64{ 4, -4, 3_000_000_000, -2_250_000_000 },
1391 );
1392 }
1393
1394 test "kernel artifact plan runs native cpu machine code for authored i64 vector max kernel" {
1395 try runAuthoredCpuVectorBinaryKernel(
1396 testing.allocator,
1397 .cpu_machine_code,
1398 .max,
1399 i64,
1400 .i64,
1401 "authored_cpu_vec2_max_i64",
1402 4,
1403 2,
1404 2,
1405 [_]i64{ std.math.minInt(i64), -1, 5, -3 },
1406 [_]i64{ 1, std.math.maxInt(i64), -7, -2 },
1407 [_]i64{ 1, std.math.maxInt(i64), 5, -2 },
1408 );
1409 }
1410
1411 test "kernel artifact plan runs native cpu object for authored i64 vector max kernel" {
1412 try runAuthoredCpuVectorBinaryKernel(
1413 testing.allocator,
1414 .cpu_object,
1415 .max,
1416 i64,
1417 .i64,
1418 "authored_cpu_vec2_max_i64",
1419 4,
1420 2,
1421 2,
1422 [_]i64{ std.math.minInt(i64), -1, 5, -3 },
1423 [_]i64{ 1, std.math.maxInt(i64), -7, -2 },
1424 [_]i64{ 1, std.math.maxInt(i64), 5, -2 },
1425 );
1426 }
1427
1428 test "kernel artifact plan runs native cpu machine code for authored u64 vector sub kernel" {
1429 try runAuthoredCpuVectorBinaryKernel(
1430 testing.allocator,
1431 .cpu_machine_code,
1432 .sub,
1433 u64,
1434 .u64,
1435 "authored_cpu_vec2_sub_u64",
1436 4,
1437 2,
1438 2,
1439 [_]u64{ 0, 0x8000_0000_0000_0000, 7, 5 },
1440 [_]u64{ 1, 1, 10, 5 },
1441 [_]u64{ 0xffff_ffff_ffff_ffff, 0x7fff_ffff_ffff_ffff, 0xffff_ffff_ffff_fffd, 0 },
1442 );
1443 }
1444
1445 test "kernel artifact plan runs native cpu object for authored u64 vector sub kernel" {
1446 try runAuthoredCpuVectorBinaryKernel(
1447 testing.allocator,
1448 .cpu_object,
1449 .sub,
1450 u64,
1451 .u64,
1452 "authored_cpu_vec2_sub_u64",
1453 4,
1454 2,
1455 2,
1456 [_]u64{ 0, 0x8000_0000_0000_0000, 7, 5 },
1457 [_]u64{ 1, 1, 10, 5 },
1458 [_]u64{ 0xffff_ffff_ffff_ffff, 0x7fff_ffff_ffff_ffff, 0xffff_ffff_ffff_fffd, 0 },
1459 );
1460 }
1461
1462 test "kernel artifact plan runs native cpu machine code for authored u64 vector div kernel" {
1463 try runAuthoredCpuVectorBinaryKernel(
1464 testing.allocator,
1465 .cpu_machine_code,
1466 .div,
1467 u64,
1468 .u64,
1469 "authored_cpu_vec2_div_u64",
1470 4,
1471 2,
1472 2,
1473 [_]u64{ 0x8000_0000_0000_0000, 0xffff_ffff_ffff_ffff, 21, 100 },
1474 [_]u64{ 2, 2, 5, 4 },
1475 [_]u64{ 0x4000_0000_0000_0000, 0x7fff_ffff_ffff_ffff, 4, 25 },
1476 );
1477 }
1478
1479 test "kernel artifact plan runs native cpu object for authored u64 vector div kernel" {
1480 try runAuthoredCpuVectorBinaryKernel(
1481 testing.allocator,
1482 .cpu_object,
1483 .div,
1484 u64,
1485 .u64,
1486 "authored_cpu_vec2_div_u64",
1487 4,
1488 2,
1489 2,
1490 [_]u64{ 0x8000_0000_0000_0000, 0xffff_ffff_ffff_ffff, 21, 100 },
1491 [_]u64{ 2, 2, 5, 4 },
1492 [_]u64{ 0x4000_0000_0000_0000, 0x7fff_ffff_ffff_ffff, 4, 25 },
1493 );
1494 }
1495
1496 test "kernel artifact plan runs native cpu machine code for authored u64 vector ushr kernel" {
1497 try runAuthoredCpuVectorBinaryKernel(
1498 testing.allocator,
1499 .cpu_machine_code,
1500 .ushr,
1501 u64,
1502 .u64,
1503 "authored_cpu_vec2_ushr_u64",
1504 4,
1505 2,
1506 2,
1507 [_]u64{ 0xffff_ffff_ffff_ffff, 0x8000_0000_0000_0000, 16, 0x7fff_ffff_ffff_ffff },
1508 [_]u64{ 1, 63, 4, 60 },
1509 [_]u64{ 0x7fff_ffff_ffff_ffff, 1, 1, 7 },
1510 );
1511 }
1512
1513 test "kernel artifact plan runs native cpu object for authored u64 vector ushr kernel" {
1514 try runAuthoredCpuVectorBinaryKernel(
1515 testing.allocator,
1516 .cpu_object,
1517 .ushr,
1518 u64,
1519 .u64,
1520 "authored_cpu_vec2_ushr_u64",
1521 4,
1522 2,
1523 2,
1524 [_]u64{ 0xffff_ffff_ffff_ffff, 0x8000_0000_0000_0000, 16, 0x7fff_ffff_ffff_ffff },
1525 [_]u64{ 1, 63, 4, 60 },
1526 [_]u64{ 0x7fff_ffff_ffff_ffff, 1, 1, 7 },
1527 );
1528 }
1529
1530 test "kernel artifact plan runs native cpu machine code for authored u64 vector mul kernel" {
1531 try runAuthoredCpuVectorBinaryKernel(
1532 testing.allocator,
1533 .cpu_machine_code,
1534 .mul,
1535 u64,
1536 .u64,
1537 "authored_cpu_vec2_mul_u64",
1538 4,
1539 2,
1540 2,
1541 [_]u64{ 0x0000_0001_0000_0000, 0xffff_ffff_ffff_ffff, 0x8000_0000_0000_0000, 0x0000_0001_0000_0001 },
1542 [_]u64{ 3, 2, 2, 0x0000_0001_0000_0001 },
1543 [_]u64{ 0x0000_0003_0000_0000, 0xffff_ffff_ffff_fffe, 0, 0x0000_0002_0000_0001 },
1544 );
1545 }
1546
1547 test "kernel artifact plan runs native cpu object for authored u64 vector mul kernel" {
1548 try runAuthoredCpuVectorBinaryKernel(
1549 testing.allocator,
1550 .cpu_object,
1551 .mul,
1552 u64,
1553 .u64,
1554 "authored_cpu_vec2_mul_u64",
1555 4,
1556 2,
1557 2,
1558 [_]u64{ 0x0000_0001_0000_0000, 0xffff_ffff_ffff_ffff, 0x8000_0000_0000_0000, 0x0000_0001_0000_0001 },
1559 [_]u64{ 3, 2, 2, 0x0000_0001_0000_0001 },
1560 [_]u64{ 0x0000_0003_0000_0000, 0xffff_ffff_ffff_fffe, 0, 0x0000_0002_0000_0001 },
1561 );
1562 }
1563
1564 test "kernel artifact plan runs native cpu machine code for authored u64 vector min kernel" {
1565 try runAuthoredCpuVectorBinaryKernel(
1566 testing.allocator,
1567 .cpu_machine_code,
1568 .min,
1569 u64,
1570 .u64,
1571 "authored_cpu_vec2_min_u64",
1572 4,
1573 2,
1574 2,
1575 [_]u64{ 0xffff_ffff_ffff_ffff, 1, 0x8000_0000_0000_0000, 7 },
1576 [_]u64{ 1, 0xffff_ffff_ffff_ffff, 2, 0x8000_0000_0000_0000 },
1577 [_]u64{ 1, 1, 2, 7 },
1578 );
1579 }
1580
1581 test "kernel artifact plan runs native cpu object for authored u64 vector min kernel" {
1582 try runAuthoredCpuVectorBinaryKernel(
1583 testing.allocator,
1584 .cpu_object,
1585 .min,
1586 u64,
1587 .u64,
1588 "authored_cpu_vec2_min_u64",
1589 4,
1590 2,
1591 2,
1592 [_]u64{ 0xffff_ffff_ffff_ffff, 1, 0x8000_0000_0000_0000, 7 },
1593 [_]u64{ 1, 0xffff_ffff_ffff_ffff, 2, 0x8000_0000_0000_0000 },
1594 [_]u64{ 1, 1, 2, 7 },
1595 );
1596 }
1597
1598 test "kernel artifact plan runs native cpu machine code for authored u64 vector or kernel" {
1599 try runAuthoredCpuVectorBinaryKernel(
1600 testing.allocator,
1601 .cpu_machine_code,
1602 .bor,
1603 u64,
1604 .u64,
1605 "authored_cpu_vec2_or_u64",
1606 4,
1607 2,
1608 2,
1609 [_]u64{ 0, 1, 0x8000_0000_0000_0000, 0x00ff_00ff_00ff_00ff },
1610 [_]u64{ 0xffff_ffff_ffff_ffff, 2, 0x7fff_ffff_ffff_ffff, 0xff00_ff00_ff00_ff00 },
1611 [_]u64{ 0xffff_ffff_ffff_ffff, 3, 0xffff_ffff_ffff_ffff, 0xffff_ffff_ffff_ffff },
1612 );
1613 }
1614
1615 test "kernel artifact plan runs native cpu object for authored u64 vector or kernel" {
1616 try runAuthoredCpuVectorBinaryKernel(
1617 testing.allocator,
1618 .cpu_object,
1619 .bor,
1620 u64,
1621 .u64,
1622 "authored_cpu_vec2_or_u64",
1623 4,
1624 2,
1625 2,
1626 [_]u64{ 0, 1, 0x8000_0000_0000_0000, 0x00ff_00ff_00ff_00ff },
1627 [_]u64{ 0xffff_ffff_ffff_ffff, 2, 0x7fff_ffff_ffff_ffff, 0xff00_ff00_ff00_ff00 },
1628 [_]u64{ 0xffff_ffff_ffff_ffff, 3, 0xffff_ffff_ffff_ffff, 0xffff_ffff_ffff_ffff },
1629 );
1630 }
1631
1632 test "kernel artifact plan runs native cpu machine code for authored i32 vector compare kernel" {
1633 try runAuthoredCpuVectorCompareKernel(
1634 testing.allocator,
1635 .cpu_machine_code,
1636 .slt,
1637 i32,
1638 .i32,
1639 "authored_cpu_vec4_cmp_i32",
1640 8,
1641 4,
1642 2,
1643 [_]i32{ -3, -2, 0, 7, std.math.minInt(i32), std.math.maxInt(i32), 9, -10 },
1644 [_]i32{ -2, -2, -1, 8, 0, -1, 9, -11 },
1645 [_]i32{ 1, 0, 0, 1, 1, 0, 0, 0 },
1646 );
1647 }
1648
1649 test "kernel artifact plan runs native cpu object for authored i32 vector compare kernel" {
1650 try runAuthoredCpuVectorCompareKernel(
1651 testing.allocator,
1652 .cpu_object,
1653 .slt,
1654 i32,
1655 .i32,
1656 "authored_cpu_vec4_cmp_i32",
1657 8,
1658 4,
1659 2,
1660 [_]i32{ -3, -2, 0, 7, std.math.minInt(i32), std.math.maxInt(i32), 9, -10 },
1661 [_]i32{ -2, -2, -1, 8, 0, -1, 9, -11 },
1662 [_]i32{ 1, 0, 0, 1, 1, 0, 0, 0 },
1663 );
1664 }
1665
1666 test "kernel artifact plan runs native cpu machine code for authored u32 vector compare kernel" {
1667 try runAuthoredCpuVectorCompareKernel(
1668 testing.allocator,
1669 .cpu_machine_code,
1670 .uge,
1671 u32,
1672 .u32,
1673 "authored_cpu_vec4_cmp_u32",
1674 8,
1675 4,
1676 2,
1677 [_]u32{ 0, 1, 0x8000_0000, 0xffff_ffff, 7, 9, 10, 3 },
1678 [_]u32{ 1, 1, 0x7fff_ffff, 0xffff_ffff, 8, 8, 10, 4 },
1679 [_]u32{ 0, 1, 1, 1, 0, 1, 1, 0 },
1680 );
1681 }
1682
1683 test "kernel artifact plan runs native cpu object for authored u32 vector compare kernel" {
1684 try runAuthoredCpuVectorCompareKernel(
1685 testing.allocator,
1686 .cpu_object,
1687 .uge,
1688 u32,
1689 .u32,
1690 "authored_cpu_vec4_cmp_u32",
1691 8,
1692 4,
1693 2,
1694 [_]u32{ 0, 1, 0x8000_0000, 0xffff_ffff, 7, 9, 10, 3 },
1695 [_]u32{ 1, 1, 0x7fff_ffff, 0xffff_ffff, 8, 8, 10, 4 },
1696 [_]u32{ 0, 1, 1, 1, 0, 1, 1, 0 },
1697 );
1698 }
1699
1700 test "kernel artifact plan runs native cpu machine code for authored u64 vector compare kernel" {
1701 try runAuthoredCpuVectorCompareKernel(
1702 testing.allocator,
1703 .cpu_machine_code,
1704 .ne,
1705 u64,
1706 .u64,
1707 "authored_cpu_vec2_cmp_u64",
1708 4,
1709 2,
1710 2,
1711 [_]u64{ 0, 1, 0xffff_ffff_ffff_ffff, 0x8000_0000_0000_0000 },
1712 [_]u64{ 0, 2, 0xffff_ffff_ffff_ffff, 0x7fff_ffff_ffff_ffff },
1713 [_]u64{ 0, 1, 0, 1 },
1714 );
1715 }
1716
1717 test "kernel artifact plan runs native cpu object for authored u64 vector compare kernel" {
1718 try runAuthoredCpuVectorCompareKernel(
1719 testing.allocator,
1720 .cpu_object,
1721 .ne,
1722 u64,
1723 .u64,
1724 "authored_cpu_vec2_cmp_u64",
1725 4,
1726 2,
1727 2,
1728 [_]u64{ 0, 1, 0xffff_ffff_ffff_ffff, 0x8000_0000_0000_0000 },
1729 [_]u64{ 0, 2, 0xffff_ffff_ffff_ffff, 0x7fff_ffff_ffff_ffff },
1730 [_]u64{ 0, 1, 0, 1 },
1731 );
1732 }
1733
1734 test "kernel artifact plan runs native cpu machine code for authored i64 ordered vector compare kernel" {
1735 try runAuthoredCpuVectorCompareKernel(
1736 testing.allocator,
1737 .cpu_machine_code,
1738 .slt,
1739 i64,
1740 .i64,
1741 "authored_cpu_vec2_order_cmp_i64",
1742 4,
1743 2,
1744 2,
1745 [_]i64{ std.math.minInt(i64), -3, 5, std.math.maxInt(i64) },
1746 [_]i64{ 0, -2, 7, -1 },
1747 [_]i64{ 1, 1, 1, 0 },
1748 );
1749 }
1750
1751 test "kernel artifact plan runs native cpu object for authored i64 ordered vector compare kernel" {
1752 try runAuthoredCpuVectorCompareKernel(
1753 testing.allocator,
1754 .cpu_object,
1755 .slt,
1756 i64,
1757 .i64,
1758 "authored_cpu_vec2_order_cmp_i64",
1759 4,
1760 2,
1761 2,
1762 [_]i64{ std.math.minInt(i64), -3, 5, std.math.maxInt(i64) },
1763 [_]i64{ 0, -2, 7, -1 },
1764 [_]i64{ 1, 1, 1, 0 },
1765 );
1766 }
1767
1768 test "kernel artifact plan runs native cpu machine code for authored u64 ordered vector compare kernel" {
1769 try runAuthoredCpuVectorCompareKernel(
1770 testing.allocator,
1771 .cpu_machine_code,
1772 .uge,
1773 u64,
1774 .u64,
1775 "authored_cpu_vec2_order_cmp_u64",
1776 4,
1777 2,
1778 2,
1779 [_]u64{ 0, 1, 0x8000_0000_0000_0000, 0xffff_ffff_ffff_fffe },
1780 [_]u64{ 1, 1, 0x7fff_ffff_ffff_ffff, 0xffff_ffff_ffff_ffff },
1781 [_]u64{ 0, 1, 1, 0 },
1782 );
1783 }
1784
1785 test "kernel artifact plan runs native cpu object for authored u64 ordered vector compare kernel" {
1786 try runAuthoredCpuVectorCompareKernel(
1787 testing.allocator,
1788 .cpu_object,
1789 .uge,
1790 u64,
1791 .u64,
1792 "authored_cpu_vec2_order_cmp_u64",
1793 4,
1794 2,
1795 2,
1796 [_]u64{ 0, 1, 0x8000_0000_0000_0000, 0xffff_ffff_ffff_fffe },
1797 [_]u64{ 1, 1, 0x7fff_ffff_ffff_ffff, 0xffff_ffff_ffff_ffff },
1798 [_]u64{ 0, 1, 1, 0 },
1799 );
1800 }
1801
1802 test "kernel artifact plan runs native cpu machine code for authored u32 vector scalar-bias kernel" {
1803 try runAuthoredCpuVectorBiasAddKernel(
1804 testing.allocator,
1805 .cpu_machine_code,
1806 u32,
1807 .u32,
1808 "authored_cpu_vec4_bias_u32",
1809 8,
1810 4,
1811 2,
1812 .{ .u32 = 0x8000_0000 },
1813 [_]u32{ 0, 1, 0x7fff_ffff, 0xffff_ffff, 3, 4, 5, 6 },
1814 [_]u32{ 0x8000_0000, 0x8000_0001, 0xffff_ffff, 0x7fff_ffff, 0x8000_0003, 0x8000_0004, 0x8000_0005, 0x8000_0006 },
1815 );
1816 }
1817
1818 test "kernel artifact plan runs native cpu object for authored u32 vector scalar-bias kernel" {
1819 try runAuthoredCpuVectorBiasAddKernel(
1820 testing.allocator,
1821 .cpu_object,
1822 u32,
1823 .u32,
1824 "authored_cpu_vec4_bias_u32",
1825 8,
1826 4,
1827 2,
1828 .{ .u32 = 0x8000_0000 },
1829 [_]u32{ 0, 1, 0x7fff_ffff, 0xffff_ffff, 3, 4, 5, 6 },
1830 [_]u32{ 0x8000_0000, 0x8000_0001, 0xffff_ffff, 0x7fff_ffff, 0x8000_0003, 0x8000_0004, 0x8000_0005, 0x8000_0006 },
1831 );
1832 }
1833
1834 test "kernel artifact plan runs native cpu machine code for authored u32 vector neg kernel" {
1835 try runAuthoredCpuVectorUnaryU32Kernel(
1836 testing.allocator,
1837 .cpu_machine_code,
1838 .neg,
1839 "authored_cpu_vec4_neg_u32",
1840 [_]u32{ 0, 1, 0x8000_0000, 0xffff_ffff, 7, 8, 123, 0xffff_ff00 },
1841 [_]u32{ 0, 0xffff_ffff, 0x8000_0000, 1, 0xffff_fff9, 0xffff_fff8, 0xffff_ff85, 0x100 },
1842 );
1843 }
1844
1845 test "kernel artifact plan runs native cpu object for authored u32 vector neg kernel" {
1846 try runAuthoredCpuVectorUnaryU32Kernel(
1847 testing.allocator,
1848 .cpu_object,
1849 .neg,
1850 "authored_cpu_vec4_neg_u32",
1851 [_]u32{ 0, 1, 0x8000_0000, 0xffff_ffff, 7, 8, 123, 0xffff_ff00 },
1852 [_]u32{ 0, 0xffff_ffff, 0x8000_0000, 1, 0xffff_fff9, 0xffff_fff8, 0xffff_ff85, 0x100 },
1853 );
1854 }
1855
1856 test "kernel artifact plan runs native cpu machine code for authored u32 vector not kernel" {
1857 try runAuthoredCpuVectorUnaryU32Kernel(
1858 testing.allocator,
1859 .cpu_machine_code,
1860 .bnot,
1861 "authored_cpu_vec4_not_u32",
1862 [_]u32{ 0, 1, 0x8000_0000, 0xffff_ffff, 0x00ff_00ff, 0xff00_ff00, 0x1234_5678, 0x8765_4321 },
1863 [_]u32{ 0xffff_ffff, 0xffff_fffe, 0x7fff_ffff, 0, 0xff00_ff00, 0x00ff_00ff, 0xedcb_a987, 0x789a_bcde },
1864 );
1865 }
1866
1867 test "kernel artifact plan runs native cpu object for authored u32 vector not kernel" {
1868 try runAuthoredCpuVectorUnaryU32Kernel(
1869 testing.allocator,
1870 .cpu_object,
1871 .bnot,
1872 "authored_cpu_vec4_not_u32",
1873 [_]u32{ 0, 1, 0x8000_0000, 0xffff_ffff, 0x00ff_00ff, 0xff00_ff00, 0x1234_5678, 0x8765_4321 },
1874 [_]u32{ 0xffff_ffff, 0xffff_fffe, 0x7fff_ffff, 0, 0xff00_ff00, 0x00ff_00ff, 0xedcb_a987, 0x789a_bcde },
1875 );
1876 }
1877
1878 test "kernel artifact plan runs native cpu machine code for authored u32 vector shuffle kernel" {
1879 try runAuthoredCpuVectorUnaryU32Kernel(
1880 testing.allocator,
1881 .cpu_machine_code,
1882 .shuffle_reverse,
1883 "authored_cpu_vec4_shuffle_u32",
1884 [_]u32{ 1, 2, 3, 4, 0x8000_0000, 0xffff_ffff, 7, 8 },
1885 [_]u32{ 4, 3, 2, 1, 8, 7, 0xffff_ffff, 0x8000_0000 },
1886 );
1887 }
1888
1889 test "kernel artifact plan runs native cpu object for authored u32 vector shuffle kernel" {
1890 try runAuthoredCpuVectorUnaryU32Kernel(
1891 testing.allocator,
1892 .cpu_object,
1893 .shuffle_reverse,
1894 "authored_cpu_vec4_shuffle_u32",
1895 [_]u32{ 1, 2, 3, 4, 0x8000_0000, 0xffff_ffff, 7, 8 },
1896 [_]u32{ 4, 3, 2, 1, 8, 7, 0xffff_ffff, 0x8000_0000 },
1897 );
1898 }
1899
1900 test "kernel artifact plan runs native cpu machine code for authored u32 vector constant shift kernel" {
1901 try runAuthoredCpuVectorUnaryU32Kernel(
1902 testing.allocator,
1903 .cpu_machine_code,
1904 .shl_const_3,
1905 "authored_cpu_vec4_shl_const_u32",
1906 [_]u32{ 1, 2, 0x1000_0000, 0xffff_ffff, 5, 6, 7, 8 },
1907 [_]u32{ 8, 16, 0x8000_0000, 0xffff_fff8, 40, 48, 56, 64 },
1908 );
1909 }
1910
1911 test "kernel artifact plan runs native cpu object for authored u32 vector constant shift kernel" {
1912 try runAuthoredCpuVectorUnaryU32Kernel(
1913 testing.allocator,
1914 .cpu_object,
1915 .shl_const_3,
1916 "authored_cpu_vec4_shl_const_u32",
1917 [_]u32{ 1, 2, 0x1000_0000, 0xffff_ffff, 5, 6, 7, 8 },
1918 [_]u32{ 8, 16, 0x8000_0000, 0xffff_fff8, 40, 48, 56, 64 },
1919 );
1920 }
1921
1922 test "kernel artifact plan runs native cpu machine code for mixed buffer scalar kernel" {
1923 try runAuthoredCpuIntegerScalarKernel(
1924 testing.allocator,
1925 .cpu_machine_code,
1926 i32,
1927 .i32,
1928 "authored_cpu_mixed_scalar_i32",
1929 .{ .i32 = 7 },
1930 [_]i32{ 1, -2, 30, 400 },
1931 [_]i32{ 8, 5, 37, 407 },
1932 );
1933 }
1934
1935 test "kernel artifact plan runs native cpu object for mixed buffer scalar kernel" {
1936 try runAuthoredCpuIntegerScalarKernel(
1937 testing.allocator,
1938 .cpu_object,
1939 i32,
1940 .i32,
1941 "authored_cpu_mixed_scalar_i32",
1942 .{ .i32 = 7 },
1943 [_]i32{ 1, -2, 30, 400 },
1944 [_]i32{ 8, 5, 37, 407 },
1945 );
1946 }
1947
1948 test "kernel artifact plan runs native cpu machine code for u32 unsigned max scalar kernel" {
1949 try runAuthoredCpuIntegerScalarOpKernel(
1950 testing.allocator,
1951 .cpu_machine_code,
1952 .max,
1953 u32,
1954 .u32,
1955 "authored_cpu_unsigned_max_u32",
1956 .{ .u32 = 0x8000_0000 },
1957 [_]u32{ 0, 1, 0x7fff_ffff, 0xffff_ffff },
1958 [_]u32{ 0x8000_0000, 0x8000_0000, 0x8000_0000, 0xffff_ffff },
1959 );
1960 }
1961
1962 test "kernel artifact plan runs native cpu object for u32 unsigned max scalar kernel" {
1963 try runAuthoredCpuIntegerScalarOpKernel(
1964 testing.allocator,
1965 .cpu_object,
1966 .max,
1967 u32,
1968 .u32,
1969 "authored_cpu_unsigned_max_u32",
1970 .{ .u32 = 0x8000_0000 },
1971 [_]u32{ 0, 1, 0x7fff_ffff, 0xffff_ffff },
1972 [_]u32{ 0x8000_0000, 0x8000_0000, 0x8000_0000, 0xffff_ffff },
1973 );
1974 }
1975
1976 test "kernel artifact plan runs native cpu machine code for u32 unsigned div scalar kernel" {
1977 try runAuthoredCpuIntegerScalarOpKernel(
1978 testing.allocator,
1979 .cpu_machine_code,
1980 .div,
1981 u32,
1982 .u32,
1983 "authored_cpu_unsigned_div_u32",
1984 .{ .u32 = 2 },
1985 [_]u32{ 0x8000_0000, 0xffff_ffff, 4, 5 },
1986 [_]u32{ 0x4000_0000, 0x7fff_ffff, 2, 2 },
1987 );
1988 }
1989
1990 test "kernel artifact plan runs native cpu object for u32 unsigned div scalar kernel" {
1991 try runAuthoredCpuIntegerScalarOpKernel(
1992 testing.allocator,
1993 .cpu_object,
1994 .div,
1995 u32,
1996 .u32,
1997 "authored_cpu_unsigned_div_u32",
1998 .{ .u32 = 2 },
1999 [_]u32{ 0x8000_0000, 0xffff_ffff, 4, 5 },
2000 [_]u32{ 0x4000_0000, 0x7fff_ffff, 2, 2 },
2001 );
2002 }
2003
2004 test "kernel artifact plan runs native cpu machine code for u32 unsigned umulhi scalar kernel" {
2005 try runAuthoredCpuIntegerScalarOpKernel(
2006 testing.allocator,
2007 .cpu_machine_code,
2008 .umulhi,
2009 u32,
2010 .u32,
2011 "authored_cpu_unsigned_umulhi_u32",
2012 .{ .u32 = 2 },
2013 [_]u32{ 0x7fff_ffff, 0xffff_ffff, 0x8000_0000, 123456789 },
2014 [_]u32{ 0, 1, 1, 0 },
2015 );
2016 }
2017
2018 test "kernel artifact plan runs native cpu object for u32 unsigned umulhi scalar kernel" {
2019 try runAuthoredCpuIntegerScalarOpKernel(
2020 testing.allocator,
2021 .cpu_object,
2022 .umulhi,
2023 u32,
2024 .u32,
2025 "authored_cpu_unsigned_umulhi_u32",
2026 .{ .u32 = 2 },
2027 [_]u32{ 0x7fff_ffff, 0xffff_ffff, 0x8000_0000, 123456789 },
2028 [_]u32{ 0, 1, 1, 0 },
2029 );
2030 }
2031
2032 test "kernel artifact plan runs native cpu machine code for u32 popcount scalar kernel" {
2033 try runAuthoredCpuIntegerScalarOpKernel(
2034 testing.allocator,
2035 .cpu_machine_code,
2036 .popcount,
2037 u32,
2038 .u32,
2039 "authored_cpu_popcount_u32",
2040 .{ .u32 = 0 },
2041 [_]u32{ 0, 1, 0x5555_5555, 0xffff_ffff },
2042 [_]u32{ 0, 1, 16, 32 },
2043 );
2044 }
2045
2046 test "kernel artifact plan runs native cpu object for u32 popcount scalar kernel" {
2047 try runAuthoredCpuIntegerScalarOpKernel(
2048 testing.allocator,
2049 .cpu_object,
2050 .popcount,
2051 u32,
2052 .u32,
2053 "authored_cpu_popcount_u32",
2054 .{ .u32 = 0 },
2055 [_]u32{ 0, 1, 0x5555_5555, 0xffff_ffff },
2056 [_]u32{ 0, 1, 16, 32 },
2057 );
2058 }
2059
2060 test "kernel artifact plan runs native cpu machine code for u32 to f64 cast kernel" {
2061 try runAuthoredCpuCastKernel(
2062 testing.allocator,
2063 .cpu_machine_code,
2064 u32,
2065 .u32,
2066 f64,
2067 .f64,
2068 "authored_cpu_cast_u32_f64",
2069 [_]u32{ 0, 1, 0x8000_0000, 0xffff_ffff },
2070 [_]f64{ 0.0, 1.0, 2147483648.0, 4294967295.0 },
2071 );
2072 }
2073
2074 test "kernel artifact plan runs native cpu object for u32 to f64 cast kernel" {
2075 try runAuthoredCpuCastKernel(
2076 testing.allocator,
2077 .cpu_object,
2078 u32,
2079 .u32,
2080 f64,
2081 .f64,
2082 "authored_cpu_cast_u32_f64",
2083 [_]u32{ 0, 1, 0x8000_0000, 0xffff_ffff },
2084 [_]f64{ 0.0, 1.0, 2147483648.0, 4294967295.0 },
2085 );
2086 }
2087
2088 test "kernel artifact plan runs native cpu machine code for f64 to u32 cast kernel" {
2089 try runAuthoredCpuCastKernel(
2090 testing.allocator,
2091 .cpu_machine_code,
2092 f64,
2093 .f64,
2094 u32,
2095 .u32,
2096 "authored_cpu_cast_f64_u32",
2097 [_]f64{ 0.0, 1.0, 2147483648.0, 4294967295.0 },
2098 [_]u32{ 0, 1, 0x8000_0000, 0xffff_ffff },
2099 );
2100 }
2101
2102 test "kernel artifact plan runs native cpu object for f64 to u32 cast kernel" {
2103 try runAuthoredCpuCastKernel(
2104 testing.allocator,
2105 .cpu_object,
2106 f64,
2107 .f64,
2108 u32,
2109 .u32,
2110 "authored_cpu_cast_f64_u32",
2111 [_]f64{ 0.0, 1.0, 2147483648.0, 4294967295.0 },
2112 [_]u32{ 0, 1, 0x8000_0000, 0xffff_ffff },
2113 );
2114 }
2115
2116 test "kernel artifact plan runs native cpu machine code for i64 scalar kernel" {
2117 try runAuthoredCpuIntegerScalarKernel(
2118 testing.allocator,
2119 .cpu_machine_code,
2120 i64,
2121 .i64,
2122 "authored_cpu_mixed_scalar_i64",
2123 .{ .i64 = 5_000_000_000 },
2124 [_]i64{ 1, -2, 30, 400 },
2125 [_]i64{ 5_000_000_001, 4_999_999_998, 5_000_000_030, 5_000_000_400 },
2126 );
2127 }
2128
2129 test "kernel artifact plan runs native cpu object for i64 scalar kernel" {
2130 try runAuthoredCpuIntegerScalarKernel(
2131 testing.allocator,
2132 .cpu_object,
2133 i64,
2134 .i64,
2135 "authored_cpu_mixed_scalar_i64",
2136 .{ .i64 = 5_000_000_000 },
2137 [_]i64{ 1, -2, 30, 400 },
2138 [_]i64{ 5_000_000_001, 4_999_999_998, 5_000_000_030, 5_000_000_400 },
2139 );
2140 }
2141
2142 test "kernel artifact plan runs native cpu machine code for u64 scalar kernel" {
2143 try runAuthoredCpuIntegerScalarKernel(
2144 testing.allocator,
2145 .cpu_machine_code,
2146 u64,
2147 .u64,
2148 "authored_cpu_mixed_scalar_u64",
2149 .{ .u64 = 0x8000_0000_0000_0000 },
2150 [_]u64{ 0, 1, 0x7fff_ffff_ffff_ffff, 0xffff_ffff_ffff_ffff },
2151 [_]u64{ 0x8000_0000_0000_0000, 0x8000_0000_0000_0001, 0xffff_ffff_ffff_ffff, 0x7fff_ffff_ffff_ffff },
2152 );
2153 }
2154
2155 test "kernel artifact plan runs native cpu object for u64 scalar kernel" {
2156 try runAuthoredCpuIntegerScalarKernel(
2157 testing.allocator,
2158 .cpu_object,
2159 u64,
2160 .u64,
2161 "authored_cpu_mixed_scalar_u64",
2162 .{ .u64 = 0x8000_0000_0000_0000 },
2163 [_]u64{ 0, 1, 0x7fff_ffff_ffff_ffff, 0xffff_ffff_ffff_ffff },
2164 [_]u64{ 0x8000_0000_0000_0000, 0x8000_0000_0000_0001, 0xffff_ffff_ffff_ffff, 0x7fff_ffff_ffff_ffff },
2165 );
2166 }
2167
2168 test "kernel artifact plan runs native cpu machine code for f32 scalar kernel" {
2169 try runAuthoredCpuFloatScalarKernel(
2170 testing.allocator,
2171 .cpu_machine_code,
2172 f32,
2173 .f32,
2174 "authored_cpu_scalar_f32",
2175 1.5,
2176 [_]f32{ 1.0, -2.0, 0.5, 8.0 },
2177 [_]f32{ 1.5, -3.0, 0.75, 12.0 },
2178 );
2179 }
2180
2181 test "kernel artifact plan runs native cpu object for f32 scalar kernel" {
2182 try runAuthoredCpuFloatScalarKernel(
2183 testing.allocator,
2184 .cpu_object,
2185 f32,
2186 .f32,
2187 "authored_cpu_scalar_f32",
2188 1.5,
2189 [_]f32{ 1.0, -2.0, 0.5, 8.0 },
2190 [_]f32{ 1.5, -3.0, 0.75, 12.0 },
2191 );
2192 }
2193
2194 test "kernel artifact plan runs native cpu machine code for f64 scalar kernel" {
2195 try runAuthoredCpuFloatScalarKernel(
2196 testing.allocator,
2197 .cpu_machine_code,
2198 f64,
2199 .f64,
2200 "authored_cpu_scalar_f64",
2201 -2.5,
2202 [_]f64{ 1.0, -2.0, 0.5, 8.0 },
2203 [_]f64{ -2.5, 5.0, -1.25, -20.0 },
2204 );
2205 }
2206
2207 test "kernel artifact plan runs native cpu object for f64 scalar kernel" {
2208 try runAuthoredCpuFloatScalarKernel(
2209 testing.allocator,
2210 .cpu_object,
2211 f64,
2212 .f64,
2213 "authored_cpu_scalar_f64",
2214 -2.5,
2215 [_]f64{ 1.0, -2.0, 0.5, 8.0 },
2216 [_]f64{ -2.5, 5.0, -1.25, -20.0 },
2217 );
2218 }
2219
2220 fn runAuthoredCpuCopyKernel(
2221 allocator: std.mem.Allocator,
2222 format: gpu.ArtifactFormat,
2223 ) !void {
2224 var builder_state = try program_mod.Builder.init(allocator, program_mod.Builder.Limits.testing, "authored_cpu_run_copy_f32", &.{
2225 builder.dynamicBuffer(.f32),
2226 builder.dynamicBuffer(.f32),
2227 });
2228 errdefer builder_state.deinit();
2229
2230 const axis = try builder_state.axis("i", 4);
2231 try builder_state.bind(axis, .thread_x);
2232 const src = builder_state.argument(0);
2233 const dst = builder_state.argument(1);
2234 const index = try builder_state.globalId(.x);
2235 const value = try builder_state.load(src, index);
2236 try builder_state.store(value, dst, index);
2237 try builder_state.return_();
2238
2239 var program = try builder_state.finish();
2240 defer program.deinit();
2241
2242 var state = gpu.cpu.State.init(allocator);
2243 defer state.deinit();
2244 const handle = state.handle();
2245
2246 var artifact_plan = createPlan(allocator, handle, &program, .{ .format = format }) catch |err| switch (err) {
2247 error.UnsupportedOperation => return error.SkipZigTest,
2248 else => return err,
2249 };
2250 defer artifact_plan.deinit();
2251
2252 try testing.expectEqual(format, artifact_plan.format);
2253 const planned = artifact_plan.kernels.items[0];
2254 try testing.expectEqual(format, planned.compile.format);
2255 try testing.expectEqual(@as(u32, 9), planned.artifact.argument_count);
2256
2257 const loaded = try handle.loadArtifact(&planned.artifact);
2258 defer handle.destroyObject(loaded.id);
2259
2260 const src_buffer = try handle.allocateBuffer(.{
2261 .byte_size = 4 * @sizeOf(f32),
2262 .alignment = @alignOf(f32),
2263 .dtype = .f32,
2264 .element_count = 4,
2265 });
2266 defer handle.destroyObject(src_buffer.id);
2267
2268 const dst_buffer = try handle.allocateBuffer(.{
2269 .byte_size = 4 * @sizeOf(f32),
2270 .alignment = @alignOf(f32),
2271 .dtype = .f32,
2272 .element_count = 4,
2273 });
2274 defer handle.destroyObject(dst_buffer.id);
2275
2276 const src_values = [_]f32{ 1.25, -2.5, 3.75, 8.0 };
2277 try handle.writeBuffer(.{
2278 .handle = src_buffer,
2279 .bytes = std.mem.sliceAsBytes(src_values[0..]),
2280 });
2281
2282 const bindings = [_]gpu.BufferBinding{
2283 .{
2284 .handle = src_buffer,
2285 .access = .read_only,
2286 .ownership = .backend,
2287 .byte_size = src_buffer.byte_size,
2288 },
2289 .{
2290 .handle = dst_buffer,
2291 .access = .write_only,
2292 .ownership = .backend,
2293 .byte_size = dst_buffer.byte_size,
2294 },
2295 };
2296 try handle.launch(.{
2297 .artifact = &planned.artifact,
2298 .loaded_artifact = loaded,
2299 .buffers = bindings[0..],
2300 .scalar_arguments = planned.static_arguments,
2301 .geometry = planned.launch_resources.geometry,
2302 });
2303
2304 var dst_values = @as([4]f32, @splat(0.0));
2305 try handle.readBuffer(.{
2306 .handle = dst_buffer,
2307 .bytes = std.mem.sliceAsBytes(dst_values[0..]),
2308 });
2309 try testing.expectEqualSlices(f32, src_values[0..], dst_values[0..]);
2310 }
2311
2312 fn runAutomaticScalarAddKernel(
2313 allocator: std.mem.Allocator,
2314 format: gpu.ArtifactFormat,
2315 comptime Element: type,
2316 comptime dtype: DType,
2317 entry_name: []const u8,
2318 comptime extent: usize,
2319 lhs_values: [extent]Element,
2320 rhs_values: [extent]Element,
2321 ) !void {
2322 var builder_state = try program_mod.Builder.init(allocator, program_mod.Builder.Limits.testing, entry_name, &.{
2323 builder.dynamicBuffer(dtype),
2324 builder.dynamicBuffer(dtype),
2325 builder.dynamicBuffer(dtype),
2326 });
2327 errdefer builder_state.deinit();
2328
2329 const axis = try builder_state.axis("i", @intCast(extent));
2330 try builder_state.bind(axis, .thread_x);
2331 const dst = builder_state.argument(0);
2332 const lhs = builder_state.argument(1);
2333 const rhs = builder_state.argument(2);
2334 const index = try builder_state.globalId(.x);
2335 const lhs_value = try builder_state.load(lhs, index);
2336 const rhs_value = try builder_state.load(rhs, index);
2337 const sum = try builder_state.add(lhs_value, rhs_value);
2338 try builder_state.store(sum, dst, index);
2339 try builder_state.return_();
2340
2341 var program = try builder_state.finish();
2342 defer program.deinit();
2343
2344 var expected_values = @as([extent]Element, @splat(@as(Element, 0)));
2345 try program.runCpu(allocator, &.{
2346 accy_root.kernel.argumentBuffer(Element, expected_values[0..]),
2347 accy_root.kernel.argumentBuffer(Element, @constCast(lhs_values[0..])),
2348 accy_root.kernel.argumentBuffer(Element, @constCast(rhs_values[0..])),
2349 });
2350
2351 var state = gpu.cpu.State.init(allocator);
2352 defer state.deinit();
2353 const handle = state.handle();
2354
2355 var artifact_plan = createPlan(allocator, handle, &program, .{ .format = format }) catch |err| switch (err) {
2356 error.UnsupportedOperation => return error.SkipZigTest,
2357 else => return err,
2358 };
2359 defer artifact_plan.deinit();
2360
2361 const planned = artifact_plan.kernels.items[0];
2362 try testing.expectEqual(format, planned.artifact.format);
2363 try testing.expectEqual(@as(u32, 10), planned.artifact.argument_count);
2364 try testing.expectEqual(@as(usize, 7), planned.static_arguments.len);
2365 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = @intCast(extent) }, planned.static_arguments[0]);
2366
2367 const loaded = try handle.loadArtifact(&planned.artifact);
2368 defer handle.destroyObject(loaded.id);
2369
2370 const dst_buffer = try handle.allocateBuffer(.{
2371 .byte_size = extent * @sizeOf(Element),
2372 .alignment = @alignOf(Element),
2373 .dtype = dtype,
2374 .element_count = extent,
2375 });
2376 defer handle.destroyObject(dst_buffer.id);
2377
2378 const lhs_buffer = try handle.allocateBuffer(.{
2379 .byte_size = extent * @sizeOf(Element),
2380 .alignment = @alignOf(Element),
2381 .dtype = dtype,
2382 .element_count = extent,
2383 });
2384 defer handle.destroyObject(lhs_buffer.id);
2385
2386 const rhs_buffer = try handle.allocateBuffer(.{
2387 .byte_size = extent * @sizeOf(Element),
2388 .alignment = @alignOf(Element),
2389 .dtype = dtype,
2390 .element_count = extent,
2391 });
2392 defer handle.destroyObject(rhs_buffer.id);
2393
2394 try handle.writeBuffer(.{
2395 .handle = lhs_buffer,
2396 .bytes = std.mem.sliceAsBytes(lhs_values[0..]),
2397 });
2398 try handle.writeBuffer(.{
2399 .handle = rhs_buffer,
2400 .bytes = std.mem.sliceAsBytes(rhs_values[0..]),
2401 });
2402
2403 const bindings = [_]gpu.BufferBinding{
2404 .{
2405 .handle = dst_buffer,
2406 .access = .write_only,
2407 .ownership = .backend,
2408 .byte_size = dst_buffer.byte_size,
2409 },
2410 .{
2411 .handle = lhs_buffer,
2412 .access = .read_only,
2413 .ownership = .backend,
2414 .byte_size = lhs_buffer.byte_size,
2415 },
2416 .{
2417 .handle = rhs_buffer,
2418 .access = .read_only,
2419 .ownership = .backend,
2420 .byte_size = rhs_buffer.byte_size,
2421 },
2422 };
2423 try handle.launch(.{
2424 .artifact = &planned.artifact,
2425 .loaded_artifact = loaded,
2426 .buffers = bindings[0..],
2427 .scalar_arguments = planned.static_arguments,
2428 .geometry = planned.launch_resources.geometry,
2429 });
2430
2431 var actual_values = @as([extent]Element, @splat(@as(Element, 0)));
2432 try handle.readBuffer(.{
2433 .handle = dst_buffer,
2434 .bytes = std.mem.sliceAsBytes(actual_values[0..]),
2435 });
2436 try testing.expectEqualSlices(Element, expected_values[0..], actual_values[0..]);
2437 }
2438
2439 fn runAutomaticScalarAffineKernel(
2440 allocator: std.mem.Allocator,
2441 format: gpu.ArtifactFormat,
2442 entry_name: []const u8,
2443 lhs_values: [6]f32,
2444 rhs_values: [6]f32,
2445 scale: f32,
2446 ) !void {
2447 const extent = 6;
2448 var builder_state = try program_mod.Builder.init(allocator, program_mod.Builder.Limits.testing, entry_name, &.{
2449 builder.dynamicBuffer(.f32),
2450 builder.scalar(.f32),
2451 builder.dynamicBuffer(.f32),
2452 builder.dynamicBuffer(.f32),
2453 });
2454 errdefer builder_state.deinit();
2455
2456 const axis = try builder_state.axis("i", extent);
2457 try builder_state.bind(axis, .thread_x);
2458 const dst = builder_state.argument(0);
2459 const scale_arg = builder_state.argument(1);
2460 const lhs = builder_state.argument(2);
2461 const rhs = builder_state.argument(3);
2462 const index = try builder_state.globalId(.x);
2463 const lhs_value = try builder_state.load(lhs, index);
2464 const rhs_value = try builder_state.load(rhs, index);
2465 const scaled = try builder_state.mul(lhs_value, scale_arg);
2466 const value = try builder_state.add(scaled, rhs_value);
2467 try builder_state.store(value, dst, index);
2468 try builder_state.return_();
2469
2470 var program = try builder_state.finish();
2471 defer program.deinit();
2472
2473 var expected_values = @as([extent]f32, @splat(@as(f32, 0)));
2474 try program.runCpu(allocator, &.{
2475 accy_root.kernel.argumentBuffer(f32, expected_values[0..]),
2476 accy_root.kernel.argumentF32(scale),
2477 accy_root.kernel.argumentBuffer(f32, @constCast(lhs_values[0..])),
2478 accy_root.kernel.argumentBuffer(f32, @constCast(rhs_values[0..])),
2479 });
2480
2481 var state = gpu.cpu.State.init(allocator);
2482 defer state.deinit();
2483 const handle = state.handle();
2484
2485 var artifact_plan = createPlan(allocator, handle, &program, .{ .format = format }) catch |err| switch (err) {
2486 error.UnsupportedOperation => return error.SkipZigTest,
2487 else => return err,
2488 };
2489 defer artifact_plan.deinit();
2490
2491 const planned = artifact_plan.kernels.items[0];
2492 try testing.expectEqual(format, planned.artifact.format);
2493 try testing.expectEqual(@as(u32, 11), planned.artifact.argument_count);
2494 try testing.expectEqual(@as(u32, 1), planned.runtime_scalar_argument_count);
2495 try testing.expectEqual(@as(usize, 7), planned.static_arguments.len);
2496 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = extent }, planned.static_arguments[0]);
2497
2498 const loaded = try handle.loadArtifact(&planned.artifact);
2499 defer handle.destroyObject(loaded.id);
2500
2501 const dst_buffer = try handle.allocateBuffer(.{
2502 .byte_size = extent * @sizeOf(f32),
2503 .alignment = @alignOf(f32),
2504 .dtype = .f32,
2505 .element_count = extent,
2506 });
2507 defer handle.destroyObject(dst_buffer.id);
2508
2509 const lhs_buffer = try handle.allocateBuffer(.{
2510 .byte_size = extent * @sizeOf(f32),
2511 .alignment = @alignOf(f32),
2512 .dtype = .f32,
2513 .element_count = extent,
2514 });
2515 defer handle.destroyObject(lhs_buffer.id);
2516
2517 const rhs_buffer = try handle.allocateBuffer(.{
2518 .byte_size = extent * @sizeOf(f32),
2519 .alignment = @alignOf(f32),
2520 .dtype = .f32,
2521 .element_count = extent,
2522 });
2523 defer handle.destroyObject(rhs_buffer.id);
2524
2525 try handle.writeBuffer(.{
2526 .handle = lhs_buffer,
2527 .bytes = std.mem.sliceAsBytes(lhs_values[0..]),
2528 });
2529 try handle.writeBuffer(.{
2530 .handle = rhs_buffer,
2531 .bytes = std.mem.sliceAsBytes(rhs_values[0..]),
2532 });
2533
2534 const bindings = [_]gpu.BufferBinding{
2535 .{
2536 .handle = dst_buffer,
2537 .access = .write_only,
2538 .ownership = .backend,
2539 .byte_size = dst_buffer.byte_size,
2540 },
2541 .{
2542 .handle = lhs_buffer,
2543 .access = .read_only,
2544 .ownership = .backend,
2545 .byte_size = lhs_buffer.byte_size,
2546 },
2547 .{
2548 .handle = rhs_buffer,
2549 .access = .read_only,
2550 .ownership = .backend,
2551 .byte_size = rhs_buffer.byte_size,
2552 },
2553 };
2554 var scalar_arguments: [8]choir_abi.ScalarArgument = undefined;
2555 scalar_arguments[0] = .{ .f32 = scale };
2556 @memcpy(scalar_arguments[1..][0..planned.static_arguments.len], planned.static_arguments);
2557 try handle.launch(.{
2558 .artifact = &planned.artifact,
2559 .loaded_artifact = loaded,
2560 .buffers = bindings[0..],
2561 .scalar_arguments = scalar_arguments[0 .. 1 + planned.static_arguments.len],
2562 .geometry = planned.launch_resources.geometry,
2563 });
2564
2565 var actual_values = @as([extent]f32, @splat(@as(f32, 0)));
2566 try handle.readBuffer(.{
2567 .handle = dst_buffer,
2568 .bytes = std.mem.sliceAsBytes(actual_values[0..]),
2569 });
2570 try testing.expectEqualSlices(f32, expected_values[0..], actual_values[0..]);
2571 }
2572
2573 const CpuVectorBinaryOp = enum {
2574 add,
2575 sub,
2576 mul,
2577 umulhi,
2578 div,
2579 min,
2580 max,
2581 band,
2582 bor,
2583 bxor,
2584 popcount,
2585 shl,
2586 shr,
2587 ushr,
2588 select_true,
2589 select_false,
2590 };
2591
2592 fn CpuVectorBinaryContextType(comptime dtype: DType) type {
2593 return struct {
2594 dst: program_mod.BufferView(dtype),
2595 lhs_mem: program_mod.BufferView(dtype),
2596 rhs_mem: program_mod.BufferView(dtype),
2597 op: CpuVectorBinaryOp,
2598
2599 const Self = @This();
2600
2601 fn run(
2602 k: *program_mod.Builder,
2603 index: program_mod.VectorIndex1D,
2604 ctx: Self,
2605 ) !void {
2606 const lhs = try ctx.lhs_mem.loadVector(k, index, index.width);
2607 const rhs = try ctx.rhs_mem.loadVector(k, index, index.width);
2608 const output = switch (ctx.op) {
2609 .add => try k.add(lhs, rhs),
2610 .sub => try k.sub(lhs, rhs),
2611 .mul => try k.mul(lhs, rhs),
2612 .umulhi => try k.umulhi(lhs, rhs),
2613 .div => try k.div(lhs, rhs),
2614 .min => try k.min(lhs, rhs),
2615 .max => try k.max(lhs, rhs),
2616 .band => try k.and_(lhs, rhs),
2617 .bor => try k.or_(lhs, rhs),
2618 .bxor => try k.xor(lhs, rhs),
2619 .popcount => try k.popcount(lhs),
2620 .shl => try k.shl(lhs, rhs),
2621 .shr => try k.shr(lhs, rhs),
2622 .ushr => try k.ushr(lhs, rhs),
2623 .select_true => value: {
2624 const condition = try k.constantBool(true);
2625 break :value try k.select(condition, lhs, rhs);
2626 },
2627 .select_false => value: {
2628 const condition = try k.constantBool(false);
2629 break :value try k.select(condition, lhs, rhs);
2630 },
2631 };
2632 try ctx.dst.storeVector(k, output, index);
2633 }
2634 };
2635 }
2636
2637 fn runAuthoredCpuVectorBinaryKernel(
2638 allocator: std.mem.Allocator,
2639 format: gpu.ArtifactFormat,
2640 op_kind: CpuVectorBinaryOp,
2641 comptime Element: type,
2642 comptime dtype: DType,
2643 entry_name: []const u8,
2644 comptime extent: usize,
2645 comptime width: u32,
2646 comptime threads_per_block: u32,
2647 lhs_values: [extent]Element,
2648 rhs_values: [extent]Element,
2649 expected_values: [extent]Element,
2650 ) !void {
2651 var builder_state = try program_mod.Builder.init(allocator, program_mod.Builder.Limits.testing, entry_name, &.{
2652 builder.dynamicBuffer(dtype),
2653 builder.dynamicBuffer(dtype),
2654 builder.dynamicBuffer(dtype),
2655 });
2656 errdefer builder_state.deinit();
2657
2658 const BinaryContext: type = CpuVectorBinaryContextType(dtype);
2659 const context: BinaryContext = .{
2660 .dst = builder_state.bufferArgument(dtype, 0),
2661 .lhs_mem = builder_state.bufferArgument(dtype, 1),
2662 .rhs_mem = builder_state.bufferArgument(dtype, 2),
2663 .op = op_kind,
2664 };
2665 _ = try builder_state.forEachVector1D(
2666 "i",
2667 extent,
2668 width,
2669 threads_per_block,
2670 context,
2671 BinaryContext.run,
2672 );
2673 try builder_state.return_();
2674
2675 var program = try builder_state.finish();
2676 defer program.deinit();
2677
2678 try runAuthoredCpuVectorThreeBufferProgram(
2679 allocator,
2680 format,
2681 &program,
2682 Element,
2683 dtype,
2684 extent,
2685 width,
2686 threads_per_block,
2687 lhs_values,
2688 rhs_values,
2689 expected_values,
2690 );
2691 }
2692
2693 fn CpuVectorCompareContextType(comptime dtype: DType) type {
2694 return struct {
2695 dst: program_mod.BufferView(dtype),
2696 lhs_mem: program_mod.BufferView(dtype),
2697 rhs_mem: program_mod.BufferView(dtype),
2698 predicate: builder.Compare,
2699
2700 const Self = @This();
2701
2702 fn run(
2703 k: *program_mod.Builder,
2704 index: program_mod.VectorIndex1D,
2705 ctx: Self,
2706 ) !void {
2707 const lhs = try ctx.lhs_mem.loadVector(k, index, index.width);
2708 const rhs = try ctx.rhs_mem.loadVector(k, index, index.width);
2709 const output = try k.compare(ctx.predicate, lhs, rhs);
2710 try ctx.dst.storeVector(k, output, index);
2711 }
2712 };
2713 }
2714
2715 fn runAuthoredCpuVectorCompareKernel(
2716 allocator: std.mem.Allocator,
2717 format: gpu.ArtifactFormat,
2718 predicate: builder.Compare,
2719 comptime Element: type,
2720 comptime dtype: DType,
2721 entry_name: []const u8,
2722 comptime extent: usize,
2723 comptime width: u32,
2724 comptime threads_per_block: u32,
2725 lhs_values: [extent]Element,
2726 rhs_values: [extent]Element,
2727 expected_values: [extent]Element,
2728 ) !void {
2729 var builder_state = try program_mod.Builder.init(allocator, program_mod.Builder.Limits.testing, entry_name, &.{
2730 builder.dynamicBuffer(dtype),
2731 builder.dynamicBuffer(dtype),
2732 builder.dynamicBuffer(dtype),
2733 });
2734 errdefer builder_state.deinit();
2735
2736 const CompareContext: type = CpuVectorCompareContextType(dtype);
2737 const context: CompareContext = .{
2738 .dst = builder_state.bufferArgument(dtype, 0),
2739 .lhs_mem = builder_state.bufferArgument(dtype, 1),
2740 .rhs_mem = builder_state.bufferArgument(dtype, 2),
2741 .predicate = predicate,
2742 };
2743 _ = try builder_state.forEachVector1D(
2744 "i",
2745 extent,
2746 width,
2747 threads_per_block,
2748 context,
2749 CompareContext.run,
2750 );
2751 try builder_state.return_();
2752
2753 var program = try builder_state.finish();
2754 defer program.deinit();
2755
2756 try runAuthoredCpuVectorThreeBufferProgram(
2757 allocator,
2758 format,
2759 &program,
2760 Element,
2761 dtype,
2762 extent,
2763 width,
2764 threads_per_block,
2765 lhs_values,
2766 rhs_values,
2767 expected_values,
2768 );
2769 }
2770
2771 fn runAuthoredCpuVectorThreeBufferProgram(
2772 allocator: std.mem.Allocator,
2773 format: gpu.ArtifactFormat,
2774 program: *program_mod.Program,
2775 comptime Element: type,
2776 comptime dtype: DType,
2777 comptime extent: usize,
2778 comptime width: u32,
2779 comptime threads_per_block: u32,
2780 lhs_values: [extent]Element,
2781 rhs_values: [extent]Element,
2782 expected_values: [extent]Element,
2783 ) !void {
2784 var schedule_snapshot = try program.scheduleSnapshot(allocator);
2785 defer schedule_snapshot.deinit(allocator);
2786 try testing.expectEqual(@as(usize, 1), schedule_snapshot.allAxes().len);
2787 try testing.expectEqual(@as(?u32, width), schedule_snapshot.allAxes()[0].vector_width);
2788 const launch = try schedule_snapshot.launch();
2789 try testing.expectEqual(@as(u32, threads_per_block), launch.block[0]);
2790
2791 var state = gpu.cpu.State.init(allocator);
2792 defer state.deinit();
2793 const handle = state.handle();
2794
2795 var artifact_plan = createPlan(allocator, handle, program, .{ .format = format }) catch |err| switch (err) {
2796 error.UnsupportedOperation => return error.SkipZigTest,
2797 else => return err,
2798 };
2799 defer artifact_plan.deinit();
2800
2801 const planned = artifact_plan.kernels.items[0];
2802 const packet_count: u32 = @intCast(extent / width);
2803 try testing.expectEqual(format, planned.artifact.format);
2804 try testing.expectEqual(@as(u32, 10), planned.artifact.argument_count);
2805 try testing.expectEqual(@as(usize, 7), planned.static_arguments.len);
2806 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = packet_count }, planned.static_arguments[0]);
2807 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[1]);
2808 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[2]);
2809 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[3]);
2810 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = threads_per_block }, planned.static_arguments[4]);
2811 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[5]);
2812 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[6]);
2813
2814 const loaded = try handle.loadArtifact(&planned.artifact);
2815 defer handle.destroyObject(loaded.id);
2816
2817 const dst_buffer = try handle.allocateBuffer(.{
2818 .byte_size = extent * @sizeOf(Element),
2819 .alignment = @alignOf(Element),
2820 .dtype = dtype,
2821 .element_count = extent,
2822 });
2823 defer handle.destroyObject(dst_buffer.id);
2824
2825 const lhs_buffer = try handle.allocateBuffer(.{
2826 .byte_size = extent * @sizeOf(Element),
2827 .alignment = @alignOf(Element),
2828 .dtype = dtype,
2829 .element_count = extent,
2830 });
2831 defer handle.destroyObject(lhs_buffer.id);
2832
2833 const rhs_buffer = try handle.allocateBuffer(.{
2834 .byte_size = extent * @sizeOf(Element),
2835 .alignment = @alignOf(Element),
2836 .dtype = dtype,
2837 .element_count = extent,
2838 });
2839 defer handle.destroyObject(rhs_buffer.id);
2840
2841 try handle.writeBuffer(.{
2842 .handle = lhs_buffer,
2843 .bytes = std.mem.sliceAsBytes(lhs_values[0..]),
2844 });
2845 try handle.writeBuffer(.{
2846 .handle = rhs_buffer,
2847 .bytes = std.mem.sliceAsBytes(rhs_values[0..]),
2848 });
2849
2850 const bindings = [_]gpu.BufferBinding{
2851 .{
2852 .handle = dst_buffer,
2853 .access = .write_only,
2854 .ownership = .backend,
2855 .byte_size = dst_buffer.byte_size,
2856 },
2857 .{
2858 .handle = lhs_buffer,
2859 .access = .read_only,
2860 .ownership = .backend,
2861 .byte_size = lhs_buffer.byte_size,
2862 },
2863 .{
2864 .handle = rhs_buffer,
2865 .access = .read_only,
2866 .ownership = .backend,
2867 .byte_size = rhs_buffer.byte_size,
2868 },
2869 };
2870 try handle.launch(.{
2871 .artifact = &planned.artifact,
2872 .loaded_artifact = loaded,
2873 .buffers = bindings[0..],
2874 .scalar_arguments = planned.static_arguments,
2875 .geometry = planned.launch_resources.geometry,
2876 });
2877
2878 var dst_values = @as([extent]Element, @splat(@as(Element, 0)));
2879 try handle.readBuffer(.{
2880 .handle = dst_buffer,
2881 .bytes = std.mem.sliceAsBytes(dst_values[0..]),
2882 });
2883 try testing.expectEqualSlices(Element, expected_values[0..], dst_values[0..]);
2884 }
2885
2886 fn CpuVectorBiasContextType(comptime dtype: DType) type {
2887 return struct {
2888 dst: program_mod.BufferView(dtype),
2889 bias: builder.Value,
2890 src_mem: program_mod.BufferView(dtype),
2891
2892 const Self = @This();
2893
2894 fn run(
2895 k: *program_mod.Builder,
2896 index: program_mod.VectorIndex1D,
2897 ctx: Self,
2898 ) !void {
2899 const src = try ctx.src_mem.loadVector(k, index, index.width);
2900 const bias = try k.splatVector(ctx.bias, index.width);
2901 const sum = try k.add(src, bias);
2902 try ctx.dst.storeVector(k, sum, index);
2903 }
2904 };
2905 }
2906
2907 fn runAuthoredCpuVectorBiasAddKernel(
2908 allocator: std.mem.Allocator,
2909 format: gpu.ArtifactFormat,
2910 comptime Element: type,
2911 comptime dtype: DType,
2912 entry_name: []const u8,
2913 comptime extent: usize,
2914 comptime width: u32,
2915 comptime threads_per_block: u32,
2916 scalar_argument: choir_abi.ScalarArgument,
2917 src_values: [extent]Element,
2918 expected_values: [extent]Element,
2919 ) !void {
2920 var builder_state = try program_mod.Builder.init(allocator, program_mod.Builder.Limits.testing, entry_name, &.{
2921 builder.dynamicBuffer(dtype),
2922 builder.scalar(dtype),
2923 builder.dynamicBuffer(dtype),
2924 });
2925 errdefer builder_state.deinit();
2926
2927 const BiasContext: type = CpuVectorBiasContextType(dtype);
2928 const context: BiasContext = .{
2929 .dst = builder_state.bufferArgument(dtype, 0),
2930 .bias = builder_state.argument(1),
2931 .src_mem = builder_state.bufferArgument(dtype, 2),
2932 };
2933 _ = try builder_state.forEachVector1D(
2934 "i",
2935 extent,
2936 width,
2937 threads_per_block,
2938 context,
2939 BiasContext.run,
2940 );
2941 try builder_state.return_();
2942
2943 var program = try builder_state.finish();
2944 defer program.deinit();
2945
2946 var schedule_snapshot = try program.scheduleSnapshot(allocator);
2947 defer schedule_snapshot.deinit(allocator);
2948 try testing.expectEqual(@as(usize, 1), schedule_snapshot.allAxes().len);
2949 try testing.expectEqual(@as(?u32, width), schedule_snapshot.allAxes()[0].vector_width);
2950 const launch = try schedule_snapshot.launch();
2951 try testing.expectEqual(@as(u32, threads_per_block), launch.block[0]);
2952
2953 var state = gpu.cpu.State.init(allocator);
2954 defer state.deinit();
2955 const handle = state.handle();
2956
2957 var artifact_plan = createPlan(allocator, handle, &program, .{ .format = format }) catch |err| switch (err) {
2958 error.UnsupportedOperation => return error.SkipZigTest,
2959 else => return err,
2960 };
2961 defer artifact_plan.deinit();
2962
2963 const planned = artifact_plan.kernels.items[0];
2964 const packet_count: u32 = @intCast(extent / width);
2965 try testing.expectEqual(format, planned.artifact.format);
2966 try testing.expectEqual(@as(u32, 10), planned.artifact.argument_count);
2967 try testing.expectEqual(@as(u32, 1), planned.runtime_scalar_argument_count);
2968 try testing.expectEqual(@as(usize, 7), planned.static_arguments.len);
2969 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = packet_count }, planned.static_arguments[0]);
2970 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[1]);
2971 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[2]);
2972 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[3]);
2973 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = threads_per_block }, planned.static_arguments[4]);
2974 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[5]);
2975 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[6]);
2976
2977 const loaded = try handle.loadArtifact(&planned.artifact);
2978 defer handle.destroyObject(loaded.id);
2979
2980 const dst_buffer = try handle.allocateBuffer(.{
2981 .byte_size = extent * @sizeOf(Element),
2982 .alignment = @alignOf(Element),
2983 .dtype = dtype,
2984 .element_count = extent,
2985 });
2986 defer handle.destroyObject(dst_buffer.id);
2987
2988 const src_buffer = try handle.allocateBuffer(.{
2989 .byte_size = extent * @sizeOf(Element),
2990 .alignment = @alignOf(Element),
2991 .dtype = dtype,
2992 .element_count = extent,
2993 });
2994 defer handle.destroyObject(src_buffer.id);
2995
2996 try handle.writeBuffer(.{
2997 .handle = src_buffer,
2998 .bytes = std.mem.sliceAsBytes(src_values[0..]),
2999 });
3000
3001 const bindings = [_]gpu.BufferBinding{
3002 .{
3003 .handle = dst_buffer,
3004 .access = .write_only,
3005 .ownership = .backend,
3006 .byte_size = dst_buffer.byte_size,
3007 },
3008 .{
3009 .handle = src_buffer,
3010 .access = .read_only,
3011 .ownership = .backend,
3012 .byte_size = src_buffer.byte_size,
3013 },
3014 };
3015 var scalar_arguments: [8]choir_abi.ScalarArgument = undefined;
3016 scalar_arguments[0] = scalar_argument;
3017 @memcpy(scalar_arguments[1..][0..planned.static_arguments.len], planned.static_arguments);
3018 try handle.launch(.{
3019 .artifact = &planned.artifact,
3020 .loaded_artifact = loaded,
3021 .buffers = bindings[0..],
3022 .scalar_arguments = scalar_arguments[0 .. 1 + planned.static_arguments.len],
3023 .geometry = planned.launch_resources.geometry,
3024 });
3025
3026 var dst_values = @as([extent]Element, @splat(@as(Element, 0)));
3027 try handle.readBuffer(.{
3028 .handle = dst_buffer,
3029 .bytes = std.mem.sliceAsBytes(dst_values[0..]),
3030 });
3031 try testing.expectEqualSlices(Element, expected_values[0..], dst_values[0..]);
3032 }
3033
3034 const CpuFloatUnaryOp = enum {
3035 floor,
3036 };
3037
3038 const CpuFloatUnaryContext = struct {
3039 dst: program_mod.BufferView(.f32),
3040 src_mem: program_mod.BufferView(.f32),
3041 op: CpuFloatUnaryOp,
3042
3043 fn run(
3044 k: *program_mod.Builder,
3045 index: program_mod.Index1D,
3046 ctx: CpuFloatUnaryContext,
3047 ) !void {
3048 const src = try ctx.src_mem.load(k, index);
3049 const output = switch (ctx.op) {
3050 .floor => try src.floor(k),
3051 };
3052 try ctx.dst.store(k, output, index);
3053 }
3054 };
3055
3056 fn runAuthoredCpuFloatUnaryKernel(
3057 allocator: std.mem.Allocator,
3058 format: gpu.ArtifactFormat,
3059 op_kind: CpuFloatUnaryOp,
3060 entry_name: []const u8,
3061 src_values: [4]f32,
3062 expected_values: [4]f32,
3063 ) !void {
3064 const extent: usize = 4;
3065 const threads_per_block: u32 = 4;
3066
3067 var builder_state = try program_mod.Builder.init(allocator, program_mod.Builder.Limits.testing, entry_name, &.{
3068 builder.dynamicBuffer(.f32),
3069 builder.dynamicBuffer(.f32),
3070 });
3071 errdefer builder_state.deinit();
3072
3073 const context: CpuFloatUnaryContext = .{
3074 .dst = builder_state.bufferArgument(.f32, 0),
3075 .src_mem = builder_state.bufferArgument(.f32, 1),
3076 .op = op_kind,
3077 };
3078 _ = try builder_state.forEach1D(
3079 "i",
3080 extent,
3081 threads_per_block,
3082 context,
3083 CpuFloatUnaryContext.run,
3084 );
3085 try builder_state.return_();
3086
3087 var program = try builder_state.finish();
3088 defer program.deinit();
3089
3090 var state = gpu.cpu.State.init(allocator);
3091 defer state.deinit();
3092 const handle = state.handle();
3093
3094 var artifact_plan = createPlan(allocator, handle, &program, .{ .format = format }) catch |err| switch (err) {
3095 error.UnsupportedOperation => return error.SkipZigTest,
3096 else => return err,
3097 };
3098 defer artifact_plan.deinit();
3099
3100 const planned = artifact_plan.kernels.items[0];
3101 try testing.expectEqual(format, planned.artifact.format);
3102 try testing.expectEqual(@as(u32, 9), planned.artifact.argument_count);
3103 try testing.expectEqual(@as(u32, 0), planned.runtime_scalar_argument_count);
3104 try testing.expectEqual(@as(usize, 7), planned.static_arguments.len);
3105
3106 const loaded = try handle.loadArtifact(&planned.artifact);
3107 defer handle.destroyObject(loaded.id);
3108
3109 const dst_buffer = try handle.allocateBuffer(.{
3110 .byte_size = extent * @sizeOf(f32),
3111 .alignment = @alignOf(f32),
3112 .dtype = .f32,
3113 .element_count = extent,
3114 });
3115 defer handle.destroyObject(dst_buffer.id);
3116
3117 const src_buffer = try handle.allocateBuffer(.{
3118 .byte_size = extent * @sizeOf(f32),
3119 .alignment = @alignOf(f32),
3120 .dtype = .f32,
3121 .element_count = extent,
3122 });
3123 defer handle.destroyObject(src_buffer.id);
3124
3125 try handle.writeBuffer(.{
3126 .handle = src_buffer,
3127 .bytes = std.mem.sliceAsBytes(src_values[0..]),
3128 });
3129
3130 const bindings = [_]gpu.BufferBinding{
3131 .{
3132 .handle = dst_buffer,
3133 .access = .write_only,
3134 .ownership = .backend,
3135 .byte_size = dst_buffer.byte_size,
3136 },
3137 .{
3138 .handle = src_buffer,
3139 .access = .read_only,
3140 .ownership = .backend,
3141 .byte_size = src_buffer.byte_size,
3142 },
3143 };
3144 try handle.launch(.{
3145 .artifact = &planned.artifact,
3146 .loaded_artifact = loaded,
3147 .buffers = bindings[0..],
3148 .scalar_arguments = planned.static_arguments,
3149 .geometry = planned.launch_resources.geometry,
3150 });
3151
3152 var dst_values = @as([extent]f32, @splat(0.0));
3153 try handle.readBuffer(.{
3154 .handle = dst_buffer,
3155 .bytes = std.mem.sliceAsBytes(dst_values[0..]),
3156 });
3157 try testing.expectEqualSlices(f32, expected_values[0..], dst_values[0..]);
3158 }
3159
3160 const CpuVectorUnaryOp = enum {
3161 neg,
3162 bnot,
3163 shuffle_reverse,
3164 shl_const_3,
3165 };
3166
3167 const CpuVectorUnaryU32Context = struct {
3168 dst: program_mod.BufferView(.u32),
3169 src_mem: program_mod.BufferView(.u32),
3170 op: CpuVectorUnaryOp,
3171
3172 fn run(
3173 k: *program_mod.Builder,
3174 index: program_mod.VectorIndex1D,
3175 ctx: CpuVectorUnaryU32Context,
3176 ) !void {
3177 const src = try ctx.src_mem.loadVector(k, index, index.width);
3178 const output = switch (ctx.op) {
3179 .neg => try k.neg(src),
3180 .bnot => try k.not(src),
3181 .shuffle_reverse => value: {
3182 const indices = [_]i64{ 3, 2, 1, 0 };
3183 break :value try k.shuffleVector(src, indices[0..]);
3184 },
3185 .shl_const_3 => value: {
3186 const shift_scalar = try k.constantInt(.u32, 3);
3187 const shift = try k.splatVector(shift_scalar, index.width);
3188 break :value try k.shl(src, shift);
3189 },
3190 };
3191 try ctx.dst.storeVector(k, output, index);
3192 }
3193 };
3194
3195 fn runAuthoredCpuVectorUnaryU32Kernel(
3196 allocator: std.mem.Allocator,
3197 format: gpu.ArtifactFormat,
3198 op_kind: CpuVectorUnaryOp,
3199 entry_name: []const u8,
3200 src_values: [8]u32,
3201 expected_values: [8]u32,
3202 ) !void {
3203 const extent: usize = 8;
3204 const width: u32 = 4;
3205 const threads_per_block: u32 = 2;
3206
3207 var builder_state = try program_mod.Builder.init(allocator, program_mod.Builder.Limits.testing, entry_name, &.{
3208 builder.dynamicBuffer(.u32),
3209 builder.dynamicBuffer(.u32),
3210 });
3211 errdefer builder_state.deinit();
3212
3213 const context: CpuVectorUnaryU32Context = .{
3214 .dst = builder_state.bufferArgument(.u32, 0),
3215 .src_mem = builder_state.bufferArgument(.u32, 1),
3216 .op = op_kind,
3217 };
3218 _ = try builder_state.forEachVector1D(
3219 "i",
3220 extent,
3221 width,
3222 threads_per_block,
3223 context,
3224 CpuVectorUnaryU32Context.run,
3225 );
3226 try builder_state.return_();
3227
3228 var program = try builder_state.finish();
3229 defer program.deinit();
3230
3231 var schedule_snapshot = try program.scheduleSnapshot(allocator);
3232 defer schedule_snapshot.deinit(allocator);
3233 try testing.expectEqual(@as(usize, 1), schedule_snapshot.allAxes().len);
3234 try testing.expectEqual(@as(?u32, width), schedule_snapshot.allAxes()[0].vector_width);
3235 const launch = try schedule_snapshot.launch();
3236 try testing.expectEqual(@as(u32, threads_per_block), launch.block[0]);
3237
3238 var state = gpu.cpu.State.init(allocator);
3239 defer state.deinit();
3240 const handle = state.handle();
3241
3242 var artifact_plan = createPlan(allocator, handle, &program, .{ .format = format }) catch |err| switch (err) {
3243 error.UnsupportedOperation => return error.SkipZigTest,
3244 else => return err,
3245 };
3246 defer artifact_plan.deinit();
3247
3248 const planned = artifact_plan.kernels.items[0];
3249 const packet_count: u32 = @intCast(extent / width);
3250 try testing.expectEqual(format, planned.artifact.format);
3251 try testing.expectEqual(@as(u32, 9), planned.artifact.argument_count);
3252 try testing.expectEqual(@as(u32, 0), planned.runtime_scalar_argument_count);
3253 try testing.expectEqual(@as(usize, 7), planned.static_arguments.len);
3254 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = packet_count }, planned.static_arguments[0]);
3255 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[1]);
3256 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[2]);
3257 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[3]);
3258 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = threads_per_block }, planned.static_arguments[4]);
3259 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[5]);
3260 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[6]);
3261
3262 const loaded = try handle.loadArtifact(&planned.artifact);
3263 defer handle.destroyObject(loaded.id);
3264
3265 const dst_buffer = try handle.allocateBuffer(.{
3266 .byte_size = extent * @sizeOf(u32),
3267 .alignment = @alignOf(u32),
3268 .dtype = .u32,
3269 .element_count = extent,
3270 });
3271 defer handle.destroyObject(dst_buffer.id);
3272
3273 const src_buffer = try handle.allocateBuffer(.{
3274 .byte_size = extent * @sizeOf(u32),
3275 .alignment = @alignOf(u32),
3276 .dtype = .u32,
3277 .element_count = extent,
3278 });
3279 defer handle.destroyObject(src_buffer.id);
3280
3281 try handle.writeBuffer(.{
3282 .handle = src_buffer,
3283 .bytes = std.mem.sliceAsBytes(src_values[0..]),
3284 });
3285
3286 const bindings = [_]gpu.BufferBinding{
3287 .{
3288 .handle = dst_buffer,
3289 .access = .write_only,
3290 .ownership = .backend,
3291 .byte_size = dst_buffer.byte_size,
3292 },
3293 .{
3294 .handle = src_buffer,
3295 .access = .read_only,
3296 .ownership = .backend,
3297 .byte_size = src_buffer.byte_size,
3298 },
3299 };
3300 try handle.launch(.{
3301 .artifact = &planned.artifact,
3302 .loaded_artifact = loaded,
3303 .buffers = bindings[0..],
3304 .scalar_arguments = planned.static_arguments,
3305 .geometry = planned.launch_resources.geometry,
3306 });
3307
3308 var dst_values = @as([extent]u32, @splat(0));
3309 try handle.readBuffer(.{
3310 .handle = dst_buffer,
3311 .bytes = std.mem.sliceAsBytes(dst_values[0..]),
3312 });
3313 try testing.expectEqualSlices(u32, expected_values[0..], dst_values[0..]);
3314 }
3315
3316 const CpuIntegerScalarOp = enum {
3317 add,
3318 max,
3319 div,
3320 umulhi,
3321 popcount,
3322 };
3323
3324 fn runAuthoredCpuIntegerScalarKernel(
3325 allocator: std.mem.Allocator,
3326 format: gpu.ArtifactFormat,
3327 comptime Element: type,
3328 comptime dtype: DType,
3329 entry_name: []const u8,
3330 scalar_argument: choir_abi.ScalarArgument,
3331 src_values: [4]Element,
3332 expected_values: [4]Element,
3333 ) !void {
3334 try runAuthoredCpuIntegerScalarOpKernel(
3335 allocator,
3336 format,
3337 .add,
3338 Element,
3339 dtype,
3340 entry_name,
3341 scalar_argument,
3342 src_values,
3343 expected_values,
3344 );
3345 }
3346
3347 fn runAuthoredCpuIntegerScalarOpKernel(
3348 allocator: std.mem.Allocator,
3349 format: gpu.ArtifactFormat,
3350 op: CpuIntegerScalarOp,
3351 comptime Element: type,
3352 comptime dtype: DType,
3353 entry_name: []const u8,
3354 scalar_argument: choir_abi.ScalarArgument,
3355 src_values: [4]Element,
3356 expected_values: [4]Element,
3357 ) !void {
3358 var builder_state = try program_mod.Builder.init(allocator, program_mod.Builder.Limits.testing, entry_name, &.{
3359 builder.dynamicBuffer(dtype),
3360 builder.scalar(dtype),
3361 builder.dynamicBuffer(dtype),
3362 });
3363 errdefer builder_state.deinit();
3364
3365 const axis = try builder_state.axis("i", 4);
3366 try builder_state.bind(axis, .thread_x);
3367 const dst = builder_state.argument(0);
3368 const bias = builder_state.argument(1);
3369 const src = builder_state.argument(2);
3370 const index = try builder_state.globalId(.x);
3371 const value = try builder_state.load(src, index);
3372 const adjusted = switch (op) {
3373 .add => try builder_state.add(value, bias),
3374 .max => try builder_state.max(value, bias),
3375 .div => try builder_state.div(value, bias),
3376 .umulhi => try builder_state.umulhi(value, bias),
3377 .popcount => try builder_state.popcount(value),
3378 };
3379 try builder_state.store(adjusted, dst, index);
3380 try builder_state.return_();
3381
3382 var program = try builder_state.finish();
3383 defer program.deinit();
3384
3385 var state = gpu.cpu.State.init(allocator);
3386 defer state.deinit();
3387 const handle = state.handle();
3388
3389 var artifact_plan = createPlan(allocator, handle, &program, .{ .format = format }) catch |err| switch (err) {
3390 error.UnsupportedOperation => return error.SkipZigTest,
3391 else => return err,
3392 };
3393 defer artifact_plan.deinit();
3394
3395 const planned = artifact_plan.kernels.items[0];
3396 try testing.expectEqual(format, planned.artifact.format);
3397 try testing.expectEqual(@as(u32, 10), planned.artifact.argument_count);
3398 try testing.expectEqual(@as(u32, 1), planned.runtime_scalar_argument_count);
3399 try testing.expectEqual(@as(usize, 7), planned.static_arguments.len);
3400 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 4 }, planned.static_arguments[0]);
3401 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[1]);
3402 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[2]);
3403 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[3]);
3404 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 4 }, planned.static_arguments[4]);
3405 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[5]);
3406 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[6]);
3407
3408 const loaded = try handle.loadArtifact(&planned.artifact);
3409 defer handle.destroyObject(loaded.id);
3410
3411 const dst_buffer = try handle.allocateBuffer(.{
3412 .byte_size = 4 * @sizeOf(Element),
3413 .alignment = @alignOf(Element),
3414 .dtype = dtype,
3415 .element_count = 4,
3416 });
3417 defer handle.destroyObject(dst_buffer.id);
3418
3419 const src_buffer = try handle.allocateBuffer(.{
3420 .byte_size = 4 * @sizeOf(Element),
3421 .alignment = @alignOf(Element),
3422 .dtype = dtype,
3423 .element_count = 4,
3424 });
3425 defer handle.destroyObject(src_buffer.id);
3426
3427 try handle.writeBuffer(.{
3428 .handle = src_buffer,
3429 .bytes = std.mem.sliceAsBytes(src_values[0..]),
3430 });
3431
3432 const bindings = [_]gpu.BufferBinding{
3433 .{
3434 .handle = dst_buffer,
3435 .access = .write_only,
3436 .ownership = .backend,
3437 .byte_size = dst_buffer.byte_size,
3438 },
3439 .{
3440 .handle = src_buffer,
3441 .access = .read_only,
3442 .ownership = .backend,
3443 .byte_size = src_buffer.byte_size,
3444 },
3445 };
3446 var scalar_arguments: [8]choir_abi.ScalarArgument = undefined;
3447 scalar_arguments[0] = scalar_argument;
3448 @memcpy(scalar_arguments[1..][0..planned.static_arguments.len], planned.static_arguments);
3449 try handle.launch(.{
3450 .artifact = &planned.artifact,
3451 .loaded_artifact = loaded,
3452 .buffers = bindings[0..],
3453 .scalar_arguments = scalar_arguments[0 .. 1 + planned.static_arguments.len],
3454 .geometry = planned.launch_resources.geometry,
3455 });
3456
3457 var dst_values = @as([4]Element, @splat(@as(Element, 0)));
3458 try handle.readBuffer(.{
3459 .handle = dst_buffer,
3460 .bytes = std.mem.sliceAsBytes(dst_values[0..]),
3461 });
3462 try testing.expectEqualSlices(Element, expected_values[0..], dst_values[0..]);
3463 }
3464
3465 fn runAuthoredCpuCastKernel(
3466 allocator: std.mem.Allocator,
3467 format: gpu.ArtifactFormat,
3468 comptime SrcElement: type,
3469 comptime src_dtype: DType,
3470 comptime DstElement: type,
3471 comptime dst_dtype: DType,
3472 entry_name: []const u8,
3473 src_values: [4]SrcElement,
3474 expected_values: [4]DstElement,
3475 ) !void {
3476 var builder_state = try program_mod.Builder.init(allocator, program_mod.Builder.Limits.testing, entry_name, &.{
3477 builder.dynamicBuffer(dst_dtype),
3478 builder.dynamicBuffer(src_dtype),
3479 });
3480 errdefer builder_state.deinit();
3481
3482 const axis = try builder_state.axis("i", 4);
3483 try builder_state.bind(axis, .thread_x);
3484 const dst = builder_state.argument(0);
3485 const src = builder_state.argument(1);
3486 const index = try builder_state.globalId(.x);
3487 const value = try builder_state.load(src, index);
3488 const casted = try builder_state.cast(value, dst_dtype);
3489 try builder_state.store(casted, dst, index);
3490 try builder_state.return_();
3491
3492 var program = try builder_state.finish();
3493 defer program.deinit();
3494
3495 var state = gpu.cpu.State.init(allocator);
3496 defer state.deinit();
3497 const handle = state.handle();
3498
3499 var artifact_plan = createPlan(allocator, handle, &program, .{ .format = format }) catch |err| switch (err) {
3500 error.UnsupportedOperation => return error.SkipZigTest,
3501 else => return err,
3502 };
3503 defer artifact_plan.deinit();
3504
3505 const planned = artifact_plan.kernels.items[0];
3506 try testing.expectEqual(format, planned.artifact.format);
3507 try testing.expectEqual(@as(u32, 9), planned.artifact.argument_count);
3508 try testing.expectEqual(@as(u32, 0), planned.runtime_scalar_argument_count);
3509 try testing.expectEqual(@as(usize, 7), planned.static_arguments.len);
3510 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 4 }, planned.static_arguments[0]);
3511 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[1]);
3512 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[2]);
3513 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[3]);
3514 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 4 }, planned.static_arguments[4]);
3515 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[5]);
3516 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[6]);
3517
3518 const loaded = try handle.loadArtifact(&planned.artifact);
3519 defer handle.destroyObject(loaded.id);
3520
3521 const dst_buffer = try handle.allocateBuffer(.{
3522 .byte_size = 4 * @sizeOf(DstElement),
3523 .alignment = @alignOf(DstElement),
3524 .dtype = dst_dtype,
3525 .element_count = 4,
3526 });
3527 defer handle.destroyObject(dst_buffer.id);
3528
3529 const src_buffer = try handle.allocateBuffer(.{
3530 .byte_size = 4 * @sizeOf(SrcElement),
3531 .alignment = @alignOf(SrcElement),
3532 .dtype = src_dtype,
3533 .element_count = 4,
3534 });
3535 defer handle.destroyObject(src_buffer.id);
3536
3537 try handle.writeBuffer(.{
3538 .handle = src_buffer,
3539 .bytes = std.mem.sliceAsBytes(src_values[0..]),
3540 });
3541
3542 const bindings = [_]gpu.BufferBinding{
3543 .{
3544 .handle = dst_buffer,
3545 .access = .write_only,
3546 .ownership = .backend,
3547 .byte_size = dst_buffer.byte_size,
3548 },
3549 .{
3550 .handle = src_buffer,
3551 .access = .read_only,
3552 .ownership = .backend,
3553 .byte_size = src_buffer.byte_size,
3554 },
3555 };
3556 try handle.launch(.{
3557 .artifact = &planned.artifact,
3558 .loaded_artifact = loaded,
3559 .buffers = bindings[0..],
3560 .scalar_arguments = planned.static_arguments,
3561 .geometry = planned.launch_resources.geometry,
3562 });
3563
3564 var dst_values = @as([4]DstElement, @splat(@as(DstElement, 0)));
3565 try handle.readBuffer(.{
3566 .handle = dst_buffer,
3567 .bytes = std.mem.sliceAsBytes(dst_values[0..]),
3568 });
3569 try testing.expectEqualSlices(DstElement, expected_values[0..], dst_values[0..]);
3570 }
3571
3572 fn runAuthoredCpuFloatScalarKernel(
3573 allocator: std.mem.Allocator,
3574 format: gpu.ArtifactFormat,
3575 comptime Element: type,
3576 comptime dtype: DType,
3577 entry_name: []const u8,
3578 scale: Element,
3579 src_values: [4]Element,
3580 expected_values: [4]Element,
3581 ) !void {
3582 var builder_state = try program_mod.Builder.init(allocator, program_mod.Builder.Limits.testing, entry_name, &.{
3583 builder.dynamicBuffer(dtype),
3584 builder.scalar(dtype),
3585 builder.dynamicBuffer(dtype),
3586 });
3587 errdefer builder_state.deinit();
3588
3589 const axis = try builder_state.axis("i", 4);
3590 try builder_state.bind(axis, .thread_x);
3591 const dst = builder_state.argument(0);
3592 const factor = builder_state.argument(1);
3593 const src = builder_state.argument(2);
3594 const index = try builder_state.globalId(.x);
3595 const value = try builder_state.load(src, index);
3596 const adjusted = try builder_state.mul(value, factor);
3597 try builder_state.store(adjusted, dst, index);
3598 try builder_state.return_();
3599
3600 var program = try builder_state.finish();
3601 defer program.deinit();
3602
3603 var state = gpu.cpu.State.init(allocator);
3604 defer state.deinit();
3605 const handle = state.handle();
3606
3607 var artifact_plan = createPlan(allocator, handle, &program, .{ .format = format }) catch |err| switch (err) {
3608 error.UnsupportedOperation => return error.SkipZigTest,
3609 else => return err,
3610 };
3611 defer artifact_plan.deinit();
3612
3613 const planned = artifact_plan.kernels.items[0];
3614 try testing.expectEqual(format, planned.artifact.format);
3615 try testing.expectEqual(@as(u32, 10), planned.artifact.argument_count);
3616 try testing.expectEqual(@as(u32, 1), planned.runtime_scalar_argument_count);
3617 try testing.expectEqual(@as(usize, 7), planned.static_arguments.len);
3618 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 4 }, planned.static_arguments[0]);
3619 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[1]);
3620 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[2]);
3621 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[3]);
3622 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 4 }, planned.static_arguments[4]);
3623 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[5]);
3624 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[6]);
3625
3626 const loaded = try handle.loadArtifact(&planned.artifact);
3627 defer handle.destroyObject(loaded.id);
3628
3629 const dst_buffer = try handle.allocateBuffer(.{
3630 .byte_size = 4 * @sizeOf(Element),
3631 .alignment = @alignOf(Element),
3632 .dtype = dtype,
3633 .element_count = 4,
3634 });
3635 defer handle.destroyObject(dst_buffer.id);
3636
3637 const src_buffer = try handle.allocateBuffer(.{
3638 .byte_size = 4 * @sizeOf(Element),
3639 .alignment = @alignOf(Element),
3640 .dtype = dtype,
3641 .element_count = 4,
3642 });
3643 defer handle.destroyObject(src_buffer.id);
3644
3645 try handle.writeBuffer(.{
3646 .handle = src_buffer,
3647 .bytes = std.mem.sliceAsBytes(src_values[0..]),
3648 });
3649
3650 const bindings = [_]gpu.BufferBinding{
3651 .{
3652 .handle = dst_buffer,
3653 .access = .write_only,
3654 .ownership = .backend,
3655 .byte_size = dst_buffer.byte_size,
3656 },
3657 .{
3658 .handle = src_buffer,
3659 .access = .read_only,
3660 .ownership = .backend,
3661 .byte_size = src_buffer.byte_size,
3662 },
3663 };
3664 var scalar_arguments: [8]choir_abi.ScalarArgument = undefined;
3665 scalar_arguments[0] = switch (dtype) {
3666 .f32 => .{ .f32 = @as(f32, @floatCast(scale)) },
3667 .f64 => .{ .f64 = @as(f64, @floatCast(scale)) },
3668 else => unreachable,
3669 };
3670 @memcpy(scalar_arguments[1..][0..planned.static_arguments.len], planned.static_arguments);
3671 try handle.launch(.{
3672 .artifact = &planned.artifact,
3673 .loaded_artifact = loaded,
3674 .buffers = bindings[0..],
3675 .scalar_arguments = scalar_arguments[0 .. 1 + planned.static_arguments.len],
3676 .geometry = planned.launch_resources.geometry,
3677 });
3678
3679 var dst_values = @as([4]Element, @splat(@as(Element, 0)));
3680 try handle.readBuffer(.{
3681 .handle = dst_buffer,
3682 .bytes = std.mem.sliceAsBytes(dst_values[0..]),
3683 });
3684 try testing.expectEqualSlices(Element, expected_values[0..], dst_values[0..]);
3685 }
3686
3687 fn runAuthoredCpuRowKernel(
3688 allocator: std.mem.Allocator,
3689 format: gpu.ArtifactFormat,
3690 ) !void {
3691 var builder_state = try program_mod.Builder.init(allocator, program_mod.Builder.Limits.testing, "authored_cpu_run_row_f32", &.{
3692 builder.dynamicBuffer(.f32),
3693 builder.dynamicBuffer(.f32),
3694 });
3695 errdefer builder_state.deinit();
3696
3697 const col_axis = try builder_state.axis("col", 3);
3698 const row_axis = try builder_state.axis("row", 2);
3699 try builder_state.bind(col_axis, .thread_x);
3700 try builder_state.bind(row_axis, .thread_y);
3701
3702 const rows = builder_state.argument(0);
3703 const dst = builder_state.argument(1);
3704 const col = try builder_state.globalId(.x);
3705 const row = try builder_state.globalId(.y);
3706 const width = try builder_state.constantIndex(3);
3707 const row_offset = try builder_state.mul(row, width);
3708 const index = try builder_state.add(row_offset, col);
3709 const value = try builder_state.load(rows, row);
3710 try builder_state.store(value, dst, index);
3711 try builder_state.return_();
3712
3713 var program = try builder_state.finish();
3714 defer program.deinit();
3715
3716 var state = gpu.cpu.State.init(allocator);
3717 defer state.deinit();
3718 const handle = state.handle();
3719
3720 var artifact_plan = createPlan(allocator, handle, &program, .{ .format = format }) catch |err| switch (err) {
3721 error.UnsupportedOperation => return error.SkipZigTest,
3722 else => return err,
3723 };
3724 defer artifact_plan.deinit();
3725
3726 const planned = artifact_plan.kernels.items[0];
3727 try testing.expectEqual(format, planned.artifact.format);
3728 try testing.expectEqual(@as(u32, 9), planned.artifact.argument_count);
3729 try testing.expectEqual(@as(usize, 7), planned.static_arguments.len);
3730 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 6 }, planned.static_arguments[0]);
3731 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[1]);
3732 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[2]);
3733 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[3]);
3734 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 3 }, planned.static_arguments[4]);
3735 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 2 }, planned.static_arguments[5]);
3736 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[6]);
3737
3738 const loaded = try handle.loadArtifact(&planned.artifact);
3739 defer handle.destroyObject(loaded.id);
3740
3741 const row_buffer = try handle.allocateBuffer(.{
3742 .byte_size = 2 * @sizeOf(f32),
3743 .alignment = @alignOf(f32),
3744 .dtype = .f32,
3745 .element_count = 2,
3746 });
3747 defer handle.destroyObject(row_buffer.id);
3748
3749 const dst_buffer = try handle.allocateBuffer(.{
3750 .byte_size = 6 * @sizeOf(f32),
3751 .alignment = @alignOf(f32),
3752 .dtype = .f32,
3753 .element_count = 6,
3754 });
3755 defer handle.destroyObject(dst_buffer.id);
3756
3757 const row_values = [_]f32{ 10.0, 20.0 };
3758 try handle.writeBuffer(.{
3759 .handle = row_buffer,
3760 .bytes = std.mem.sliceAsBytes(row_values[0..]),
3761 });
3762
3763 const bindings = [_]gpu.BufferBinding{
3764 .{
3765 .handle = row_buffer,
3766 .access = .read_only,
3767 .ownership = .backend,
3768 .byte_size = row_buffer.byte_size,
3769 },
3770 .{
3771 .handle = dst_buffer,
3772 .access = .write_only,
3773 .ownership = .backend,
3774 .byte_size = dst_buffer.byte_size,
3775 },
3776 };
3777 try handle.launch(.{
3778 .artifact = &planned.artifact,
3779 .loaded_artifact = loaded,
3780 .buffers = bindings[0..],
3781 .scalar_arguments = planned.static_arguments,
3782 .geometry = planned.launch_resources.geometry,
3783 });
3784
3785 var dst_values = @as([6]f32, @splat(0.0));
3786 try handle.readBuffer(.{
3787 .handle = dst_buffer,
3788 .bytes = std.mem.sliceAsBytes(dst_values[0..]),
3789 });
3790 try testing.expectEqualSlices(f32, &[_]f32{ 10.0, 10.0, 10.0, 20.0, 20.0, 20.0 }, dst_values[0..]);
3791 }
3792
3793 fn runAuthoredCpuThreadBlockKernel(
3794 allocator: std.mem.Allocator,
3795 format: gpu.ArtifactFormat,
3796 ) !void {
3797 var builder_state = try program_mod.Builder.init(allocator, program_mod.Builder.Limits.testing, "authored_cpu_thread_block_f32", &.{
3798 builder.dynamicBuffer(.f32),
3799 builder.dynamicBuffer(.f32),
3800 });
3801 errdefer builder_state.deinit();
3802
3803 const block_axis = try builder_state.axis("block", 2);
3804 const thread_axis = try builder_state.axis("thread", 4);
3805 try builder_state.bind(block_axis, .block_x);
3806 try builder_state.bind(thread_axis, .thread_x);
3807
3808 const src = builder_state.argument(0);
3809 const dst = builder_state.argument(1);
3810 const thread = try builder_state.threadId(.x);
3811 const block = try builder_state.blockId(.x);
3812 const block_dim = try builder_state.blockDim(.x);
3813 const grid_dim = try builder_state.gridDim(.x);
3814 const block_offset = try builder_state.mul(block, block_dim);
3815 const base_index = try builder_state.add(block_offset, thread);
3816 const zero = try builder_state.sub(grid_dim, grid_dim);
3817 const index = try builder_state.add(base_index, zero);
3818 const value = try builder_state.load(src, index);
3819 try builder_state.store(value, dst, index);
3820 try builder_state.return_();
3821
3822 var program = try builder_state.finish();
3823 defer program.deinit();
3824
3825 var state = gpu.cpu.State.init(allocator);
3826 defer state.deinit();
3827 const handle = state.handle();
3828
3829 var artifact_plan = createPlan(allocator, handle, &program, .{ .format = format }) catch |err| switch (err) {
3830 error.UnsupportedOperation => return error.SkipZigTest,
3831 else => return err,
3832 };
3833 defer artifact_plan.deinit();
3834
3835 const planned = artifact_plan.kernels.items[0];
3836 try testing.expectEqual(format, planned.artifact.format);
3837 try testing.expectEqual(@as(u32, 9), planned.artifact.argument_count);
3838 try testing.expectEqual(@as(usize, 7), planned.static_arguments.len);
3839 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 8 }, planned.static_arguments[0]);
3840 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 2 }, planned.static_arguments[1]);
3841 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[2]);
3842 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[3]);
3843 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 4 }, planned.static_arguments[4]);
3844 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[5]);
3845 try testing.expectEqual(choir_abi.ScalarArgument{ .u32 = 1 }, planned.static_arguments[6]);
3846
3847 const loaded = try handle.loadArtifact(&planned.artifact);
3848 defer handle.destroyObject(loaded.id);
3849
3850 const src_buffer = try handle.allocateBuffer(.{
3851 .byte_size = 8 * @sizeOf(f32),
3852 .alignment = @alignOf(f32),
3853 .dtype = .f32,
3854 .element_count = 8,
3855 });
3856 defer handle.destroyObject(src_buffer.id);
3857
3858 const dst_buffer = try handle.allocateBuffer(.{
3859 .byte_size = 8 * @sizeOf(f32),
3860 .alignment = @alignOf(f32),
3861 .dtype = .f32,
3862 .element_count = 8,
3863 });
3864 defer handle.destroyObject(dst_buffer.id);
3865
3866 const src_values = [_]f32{ 1.0, 2.0, 3.0, 4.0, -1.0, -2.0, -3.0, -4.0 };
3867 try handle.writeBuffer(.{
3868 .handle = src_buffer,
3869 .bytes = std.mem.sliceAsBytes(src_values[0..]),
3870 });
3871
3872 const bindings = [_]gpu.BufferBinding{
3873 .{
3874 .handle = src_buffer,
3875 .access = .read_only,
3876 .ownership = .backend,
3877 .byte_size = src_buffer.byte_size,
3878 },
3879 .{
3880 .handle = dst_buffer,
3881 .access = .write_only,
3882 .ownership = .backend,
3883 .byte_size = dst_buffer.byte_size,
3884 },
3885 };
3886 try handle.launch(.{
3887 .artifact = &planned.artifact,
3888 .loaded_artifact = loaded,
3889 .buffers = bindings[0..],
3890 .scalar_arguments = planned.static_arguments,
3891 .geometry = planned.launch_resources.geometry,
3892 });
3893
3894 var dst_values = @as([8]f32, @splat(0.0));
3895 try handle.readBuffer(.{
3896 .handle = dst_buffer,
3897 .bytes = std.mem.sliceAsBytes(dst_values[0..]),
3898 });
3899 try testing.expectEqualSlices(f32, src_values[0..], dst_values[0..]);
3900 }
3901
3902 fn testCreateArtifact(
3903 ptr: *anyopaque,
3904 request: gpu.CompileRequest,
3905 ) gpu.BackendError!gpu.KernelArtifact {
3906 const state: *TestBackendState = @ptrCast(@alignCast(ptr));
3907 if (request.requested_format != state.format) return error.UnsupportedOperation;
3908 var artifact = gpu.KernelArtifact.init(state.allocator, .{
3909 .backend = state.kind,
3910 .format = state.format,
3911 .entry_name = request.kernel_name,
3912 .argument_count = request.argument_count,
3913 .scalar_argument_count = request.scalar_argument_count,
3914 .diagnostic_id = request.diagnostic_id,
3915 }) catch return error.OutOfMemory;
3916 errdefer artifact.deinit();
3917 switch (request.payload) {
3918 .text => |text| try artifact.setOwnedText(text),
3919 .bytes => |bytes| try artifact.setOwnedBytes(bytes),
3920 .words_u32 => |words| try artifact.setOwnedWords(words),
3921 .none => return error.InvalidArtifact,
3922 }
3923 return artifact;
3924 }
3925
3926 const test_backend_vtable = gpu.BackendVTable{
3927 .query_capabilities = testQueryCapabilities,
3928 .create_artifact = testCreateArtifact,
3929 };
3930
3931 test "kernel artifact plan creates wgsl for authored webgpu kernels" {
3932 const allocator = testing.allocator;
3933
3934 var builder_state = try program_mod.Builder.init(allocator, program_mod.Builder.Limits.testing, "authored_webgpu_add_f32", &.{
3935 builder.dynamicBuffer(.f32),
3936 builder.dynamicBuffer(.f32),
3937 builder.dynamicBuffer(.f32),
3938 });
3939 errdefer builder_state.deinit();
3940
3941 const axis = try builder_state.axis("i", 64);
3942 try builder_state.bind(axis, .thread_x);
3943 const dst = builder_state.argument(0);
3944 const lhs = builder_state.argument(1);
3945 const rhs = builder_state.argument(2);
3946 const index = try builder_state.globalId(.x);
3947 const value = try builder_state.add(try builder_state.load(lhs, index), try builder_state.load(rhs, index));
3948 try builder_state.store(value, dst, index);
3949 try builder_state.return_();
3950
3951 var program = try builder_state.finish();
3952 defer program.deinit();
3953
3954 var state = TestBackendState.init(allocator, .webgpu);
3955 var artifact_plan = try createPlan(allocator, state.handle(), &program, .{});
3956 defer artifact_plan.deinit();
3957
3958 try testing.expectEqual(gpu.BackendKind.webgpu, artifact_plan.backend_kind);
3959 try testing.expectEqual(gpu.ArtifactFormat.webgpu_wgsl, artifact_plan.format);
3960 try testing.expectEqual(@as(usize, 1), artifact_plan.kernelCount());
3961 const artifact = artifact_plan.kernels.items[0].artifact;
3962 try testing.expectEqual(gpu.BackendKind.webgpu, artifact.backend);
3963 try testing.expectEqual(gpu.ArtifactFormat.webgpu_wgsl, artifact.format);
3964 const text = artifact.payload.text;
3965 try testing.expect(std.mem.indexOf(u8, text, "@compute @workgroup_size") != null);
3966 try testing.expect(std.mem.indexOf(u8, text, "@group(0) @binding(0) var<storage, read_write> arg0: array<f32>;") != null);
3967 try testing.expect(std.mem.indexOf(u8, text, "fn authored_webgpu_add_f32") != null);
3968 }
3969
3970 test "kernel artifact plan infers subgroup requirements from authored warp collectives" {
3971 const allocator = testing.allocator;
3972
3973 var builder_state = try program_mod.Builder.init(allocator, program_mod.Builder.Limits.testing, "authored_warp_scan_i32", &.{
3974 builder.dynamicBuffer(.i32),
3975 builder.dynamicBuffer(.i32),
3976 });
3977 errdefer builder_state.deinit();
3978
3979 const axis = try builder_state.axis("i", 32);
3980 try builder_state.bind(axis, .thread_x);
3981 const src = builder_state.argument(0);
3982 const dst = builder_state.argument(1);
3983 const index = try builder_state.globalId(.x);
3984 const value = try builder_state.load(src, index);
3985 const reduced = try builder_state.warpReduce(.add, value);
3986 const scanned = try builder_state.warpScan(.add, .inclusive, reduced);
3987 try builder_state.store(scanned, dst, index);
3988 try builder_state.return_();
3989
3990 var program = try builder_state.finish();
3991 defer program.deinit();
3992 const body_fingerprint = try program.bodyFingerprint(allocator);
3993
3994 var state = TestBackendState.init(allocator, .cuda);
3995 var artifact_plan = try createPlan(allocator, state.handle(), &program, .{});
3996 defer artifact_plan.deinit();
3997
3998 try testing.expectEqual(body_fingerprint, try program.bodyFingerprint(allocator));
3999 try testing.expectEqual(@as(usize, 1), artifact_plan.kernelCount());
4000 const compile = artifact_plan.kernels.items[0].compile;
4001 try testing.expect(compile.required_subgroup.supported);
4002 try testing.expect(compile.required_subgroup.arithmetic);
4003 try testing.expect(compile.required_subgroup.scan);
4004 }
4005
4006 test "kernel artifact plan rejects authored warp collectives without subgroup support" {
4007 const allocator = testing.allocator;
4008
4009 var builder_state = try program_mod.Builder.init(allocator, program_mod.Builder.Limits.testing, "authored_warp_reduce_i32", &.{
4010 builder.dynamicBuffer(.i32),
4011 builder.dynamicBuffer(.i32),
4012 });
4013 errdefer builder_state.deinit();
4014
4015 const axis = try builder_state.axis("i", 32);
4016 try builder_state.bind(axis, .thread_x);
4017 const src = builder_state.argument(0);
4018 const dst = builder_state.argument(1);
4019 const index = try builder_state.globalId(.x);
4020 const value = try builder_state.load(src, index);
4021 const reduced = try builder_state.warpReduce(.add, value);
4022 try builder_state.store(reduced, dst, index);
4023 try builder_state.return_();
4024
4025 var program = try builder_state.finish();
4026 defer program.deinit();
4027
4028 var state = TestBackendState.init(allocator, .vulkan);
4029 try testing.expectError(
4030 error.CapabilityMismatch,
4031 createPlan(allocator, state.handle(), &program, .{}),
4032 );
4033 }
4034
4035 test "kernel artifact plan infers atomics from authored atomic operations" {
4036 const allocator = testing.allocator;
4037
4038 var builder_state = try program_mod.Builder.init(allocator, program_mod.Builder.Limits.testing, "authored_atomic_add_i32", &.{
4039 builder.dynamicBuffer(.i32),
4040 builder.dynamicBuffer(.i32),
4041 });
4042 errdefer builder_state.deinit();
4043
4044 const axis = try builder_state.axis("i", 32);
4045 try builder_state.bind(axis, .thread_x);
4046 const acc = builder_state.argument(0);
4047 const dst = builder_state.argument(1);
4048 const index = try builder_state.globalId(.x);
4049 const one = try builder_state.constantInt(.i32, 1);
4050 const old = try builder_state.atomicRmw(.add, one, acc, index);
4051 try builder_state.store(old, dst, index);
4052 try builder_state.return_();
4053
4054 var program = try builder_state.finish();
4055 defer program.deinit();
4056
4057 var state = TestBackendState.init(allocator, .vulkan);
4058 var artifact_plan = try createPlan(allocator, state.handle(), &program, .{});
4059 defer artifact_plan.deinit();
4060
4061 try testing.expectEqual(@as(usize, 1), artifact_plan.kernelCount());
4062 const compile = artifact_plan.kernels.items[0].compile;
4063 try testing.expect(compile.required_features.atomic_i32);
4064 try testing.expect(std.meta.eql(choir_abi.SubgroupRequirements{}, compile.required_subgroup));
4065 }
4066
4067 test "kernel artifact plan rejects unsupported authored f32 atomic operations" {
4068 const allocator = testing.allocator;
4069
4070 var builder_state = try program_mod.Builder.init(allocator, program_mod.Builder.Limits.testing, "authored_atomic_add_f32_unsupported", &.{
4071 builder.dynamicBuffer(.f32),
4072 builder.dynamicBuffer(.f32),
4073 });
4074 errdefer builder_state.deinit();
4075
4076 const axis = try builder_state.axis("i", 32);
4077 try builder_state.bind(axis, .thread_x);
4078 const acc = builder_state.argument(0);
4079 const dst = builder_state.argument(1);
4080 const index = try builder_state.globalId(.x);
4081 const one = try builder_state.constantFloat(.f32, 1.0);
4082 const old = try builder_state.atomicRmw(.add, one, acc, index);
4083 try builder_state.store(old, dst, index);
4084 try builder_state.return_();
4085
4086 var program = try builder_state.finish();
4087 defer program.deinit();
4088
4089 var state = TestBackendState.init(allocator, .vulkan);
4090 try testing.expectError(
4091 error.CapabilityMismatch,
4092 createPlan(allocator, state.handle(), &program, .{}),
4093 );
4094 }