lib/gpu/src/metal.zig
daab053ee43316e1809a84551d573ddd1e5bf3d2
1 const std = @import("std");
2 const choir_abi = @import("choir_abi");
3 const pretty = @import("pretty");
4 const build_options = @import("build_options");
5 const sys = @import("sys");
6
7 const backend = @import("root.zig");
8 const runtime_root = @import("runtime/root.zig");
9
10 const runtime_mod = runtime_root.metal.runtime;
11 const raster_mod = runtime_mod.raster;
12 const foundation = sys.apple.foundation;
13 const metal = sys.apple.metal;
14
15 const Allocator = std.mem.Allocator;
16 const Runtime = runtime_mod.Runtime;
17 const BackendObjectId = backend.BackendObjectId;
18
19 const native_available = metal.available;
20
21 /// Releases a native object. On a host without Metal no native object exists.
22 fn releaseNative(object: anytype) void {
23 if (comptime native_available) sys.apple.objc.release(object);
24 }
25
26 /// Vertex buffers a pipeline reads, at Metal buffer indices 0 to 7, below the uniform blocks.
27 const max_vertex_buffers = 8;
28 /// Vertex attributes a pipeline reads, at `[[attribute(0)]]` to `[[attribute(30)]]`.
29 const max_vertex_attributes = 16;
30 const max_attribute_location = 31;
31 /// Resource bindings a pipeline declares, each a uniform block or a sampled texture of group 0.
32 const max_bindings = choir_abi.metal.max_uniform_bindings;
33 /// Bytes between two draws' push constants in a bundle's push buffer, which meets the constant
34 /// buffer offset alignment of every Metal GPU family.
35 const push_stride = 256;
36 const max_push_constant_bytes = choir_abi.stage.push_words * 4;
37 /// Events a pass may wait for.
38 const max_wait_events = 8;
39
40 comptime {
41 std.debug.assert(max_vertex_buffers <= choir_abi.metal.uniform_buffer_base);
42 std.debug.assert(max_push_constant_bytes <= push_stride);
43 }
44
45 const RuntimeStorage = union(enum) {
46 none,
47 borrowed: *Runtime,
48 owned: Runtime,
49
50 fn ptr(self: *RuntimeStorage) ?*Runtime {
51 return switch (self.*) {
52 .none => null,
53 .borrowed => |runtime| runtime,
54 .owned => |*runtime| runtime,
55 };
56 }
57
58 fn deinit(self: *RuntimeStorage) void {
59 switch (self.*) {
60 .none, .borrowed => {},
61 .owned => |*runtime| runtime.deinit(),
62 }
63 self.* = .none;
64 }
65 };
66
67 pub const State = struct {
68 allocator: Allocator,
69 runtime: RuntimeStorage = .none,
70 next_id: BackendObjectId = 1,
71 objects: std.AutoHashMapUnmanaged(BackendObjectId, Object) = .{},
72
73 fn init(allocator: Allocator) State {
74 return .{
75 .allocator = allocator,
76 };
77 }
78
79 pub fn initDevice(allocator: Allocator) backend.BackendError!State {
80 const runtime = Runtime.init(allocator) catch |err| return mapRuntimeError(err);
81 return .{
82 .allocator = allocator,
83 .runtime = .{ .owned = runtime },
84 };
85 }
86
87 fn initWithRuntime(allocator: Allocator, runtime: ?*Runtime) State {
88 return .{
89 .allocator = allocator,
90 .runtime = if (runtime) |rt| .{ .borrowed = rt } else .none,
91 };
92 }
93
94 pub fn deinit(self: *State) void {
95 var it = self.objects.iterator();
96 while (it.next()) |entry| {
97 deinitObject(self, entry.value_ptr);
98 }
99 self.objects.deinit(self.allocator);
100 self.objects = .{};
101 self.next_id = 1;
102 self.runtime.deinit();
103 }
104
105 pub fn handle(self: *State) backend.BackendHandle {
106 return .{
107 .ptr = self,
108 .vtable = &vtable,
109 .kind = .metal,
110 };
111 }
112
113 fn putObject(self: *State, object: Object) backend.BackendError!BackendObjectId {
114 const id = self.next_id;
115 if (id == std.math.maxInt(BackendObjectId)) return error.OutOfMemory;
116 self.next_id += 1;
117 self.objects.put(self.allocator, id, object) catch return error.OutOfMemory;
118 return id;
119 }
120
121 fn getLoaded(self: *State, loaded: backend.LoadedArtifact) backend.BackendError!*runtime_mod.LoadedKernel {
122 if (loaded.backend != .metal or !isLoadableArtifactFormat(loaded.format)) return error.InvalidArtifact;
123 const object = self.objects.getPtr(loaded.id) orelse return error.InvalidArtifact;
124 return switch (object.*) {
125 .loaded_artifact => |*artifact| if (artifact.format == loaded.format) &artifact.kernel else error.InvalidArtifact,
126 else => error.InvalidArtifact,
127 };
128 }
129
130 fn getBuffer(self: *State, buffer_handle: backend.BufferHandle) backend.BackendError!*runtime_mod.DeviceBuffer {
131 if (buffer_handle.backend != .metal) return error.InvalidBuffer;
132 const object = self.objects.getPtr(buffer_handle.id) orelse return error.InvalidBuffer;
133 return switch (object.*) {
134 .buffer => |*buffer| buffer,
135 else => error.InvalidBuffer,
136 };
137 }
138
139 fn getStream(self: *State, stream_handle: backend.StreamHandle) backend.BackendError!*MetalStream {
140 if (stream_handle.backend != .metal) return error.InvalidStream;
141 const object = self.objects.getPtr(stream_handle.id) orelse return error.InvalidStream;
142 return switch (object.*) {
143 .stream => |*stream| stream,
144 else => error.InvalidStream,
145 };
146 }
147
148 fn getEvent(self: *State, event_handle: backend.EventHandle) backend.BackendError!*MetalEvent {
149 if (event_handle.backend != .metal) return error.InvalidEvent;
150 const object = self.objects.getPtr(event_handle.id) orelse return error.InvalidEvent;
151 return switch (object.*) {
152 .event => |*event| event,
153 else => error.InvalidEvent,
154 };
155 }
156
157 fn getTexture(self: *State, texture_handle: backend.TextureHandle) backend.BackendError!*MetalTexture {
158 if (texture_handle.backend != .metal) return error.InvalidTexture;
159 const object = self.objects.getPtr(texture_handle.id) orelse return error.InvalidTexture;
160 return switch (object.*) {
161 .texture => |*texture| texture,
162 else => error.InvalidTexture,
163 };
164 }
165
166 fn getRenderPipeline(self: *State, loaded: backend.LoadedRenderArtifact) backend.BackendError!*MetalPipeline {
167 if (loaded.backend != .metal) return error.InvalidRenderArtifact;
168 const object = self.objects.getPtr(loaded.id) orelse return error.InvalidRenderArtifact;
169 return switch (object.*) {
170 .render_pipeline => |*pipeline| pipeline,
171 else => error.InvalidRenderArtifact,
172 };
173 }
174
175 fn getRenderBindings(self: *State, bindings: backend.RenderBindings) backend.BackendError!*MetalBindings {
176 if (bindings.backend != .metal) return error.RenderArgumentMismatch;
177 const object = self.objects.getPtr(bindings.id) orelse return error.RenderArgumentMismatch;
178 return switch (object.*) {
179 .render_bindings => |*set| set,
180 else => error.RenderArgumentMismatch,
181 };
182 }
183
184 fn getRenderBundle(self: *State, bundle: backend.RenderBundle) backend.BackendError!*MetalBundle {
185 if (bundle.backend != .metal) return error.RenderArgumentMismatch;
186 const object = self.objects.getPtr(bundle.id) orelse return error.RenderArgumentMismatch;
187 return switch (object.*) {
188 .render_bundle => |*recorded| recorded,
189 else => error.RenderArgumentMismatch,
190 };
191 }
192
193 fn getSurface(self: *State, surface: backend.SurfaceHandle) backend.BackendError!*MetalSurface {
194 if (surface.backend != .metal) return error.InvalidSurface;
195 const object = self.objects.getPtr(surface.id) orelse return error.InvalidSurface;
196 return switch (object.*) {
197 .surface => |*stored| stored,
198 else => error.InvalidSurface,
199 };
200 }
201
202 fn getSurfaceById(self: *State, id: BackendObjectId) ?*MetalSurface {
203 const object = self.objects.getPtr(id) orelse return null;
204 return switch (object.*) {
205 .surface => |*stored| stored,
206 else => null,
207 };
208 }
209
210 fn getFrame(self: *State, frame: backend.SurfaceFrame) backend.BackendError!*MetalFrame {
211 if (frame.backend != .metal) return error.InvalidSurfaceFrame;
212 const object = self.objects.getPtr(frame.id) orelse return error.InvalidSurfaceFrame;
213 return switch (object.*) {
214 .surface_frame => |*stored| stored,
215 else => error.InvalidSurfaceFrame,
216 };
217 }
218
219 fn nativeStream(self: *State, stream: ?backend.StreamHandle) backend.BackendError!?*runtime_mod.Stream {
220 const stream_handle = stream orelse return null;
221 return (try self.getStream(stream_handle)).native orelse error.InvalidStream;
222 }
223
224 fn nativeWaits(
225 self: *State,
226 events: []const backend.EventHandle,
227 storage: *[max_wait_events]*runtime_mod.Event,
228 ) backend.BackendError![]const *runtime_mod.Event {
229 if (events.len > max_wait_events) return error.LaunchArgumentMismatch;
230 for (events, storage[0..events.len]) |event_handle, *slot| {
231 const event = try self.getEvent(event_handle);
232 if (!event.recorded) return error.InvalidEvent;
233 slot.* = event.native orelse return error.InvalidEvent;
234 }
235 return storage[0..events.len];
236 }
237
238 fn nativeSignal(self: *State, event: ?backend.EventHandle) backend.BackendError!?*MetalEvent {
239 const event_handle = event orelse return null;
240 const signal = try self.getEvent(event_handle);
241 if (signal.native == null) return error.InvalidEvent;
242 return signal;
243 }
244
245 /// Encodes `pass` with no checks and no state beyond the first draw's pipeline, bindings and
246 /// vertex buffers, then drops the command buffer uncommitted: the least a pass of these draws
247 /// can cost to record. Every draw must be non-indexed.
248 pub fn recordFloor(self: *State, pass: backend.RenderPass) backend.BackendError!void {
249 if (comptime !native_available) return error.RuntimeUnavailable;
250 const rt = self.runtime.ptr() orelse return error.RuntimeUnavailable;
251 if (pass.draws.len == 0) return;
252 var pool = foundation.AutoreleasePool.init() orelse return error.RuntimeUnavailable;
253 defer pool.deinit();
254 const targets = try passTargets(self, pass);
255 const stream = rt.defaultStream() catch |err| return mapRenderError(err);
256 const command_buffer = metal.commandBuffer(stream.command_queue) orelse return error.RenderFailed;
257 const encoder = try beginPass(command_buffer, targets);
258 const first = pass.draws[0];
259 const pipeline = try self.getRenderPipeline(first.pipeline);
260 var sink = EncoderSink{ .encoder = encoder };
261 sink.setPipeline(pipeline);
262 if (first.bindings) |bindings| sink.setBindings(try self.getRenderBindings(bindings));
263 for (first.vertex_buffers, pipeline.layouts[0..first.vertex_buffers.len]) |range, layout| {
264 metal.setVertexBuffer(encoder, (try self.getBuffer(range.buffer)).buffer, range.offset, layout.binding);
265 }
266 for (pass.draws) |draw| {
267 if (draw.push_constants.len != 0) {
268 metal.setVertexBytes(encoder, draw.push_constants, choir_abi.metal.push_buffer);
269 metal.setFragmentBytes(encoder, draw.push_constants, choir_abi.metal.push_buffer);
270 }
271 metal.drawPrimitives(encoder, pipeline.primitive, draw.range.first_vertex, draw.range.vertex_count, draw.range.instance_count, draw.range.first_instance);
272 }
273 metal.endRenderEncoding(encoder);
274 }
275
276 /// Waits for the last frame `surface` presented and copies its texels, rows from the top, in
277 /// the surface's format: the texels the display was handed, read back in the command buffer
278 /// that presented them.
279 pub fn readPresented(self: *State, surface: backend.SurfaceHandle, bytes: []u8) backend.BackendError!void {
280 const stored = try self.getSurface(surface);
281 stored.surface.readPresented(bytes) catch |err| return mapSurfaceError(err);
282 }
283
284 /// Compiles `source` into a library with fast math off and drops it, so a caller can see that
285 /// the Metal compiler takes a module whatever its entry points.
286 pub fn compileMsl(self: *State, source: []const u8) backend.BackendError!void {
287 if (comptime !native_available) return error.RuntimeUnavailable;
288 const rt = self.runtime.ptr() orelse return error.RuntimeUnavailable;
289 var pool = foundation.AutoreleasePool.init() orelse return error.RuntimeUnavailable;
290 defer pool.deinit();
291 const library = rt.newLibrary(.{ .msl = source }) catch |err| return mapRuntimeError(err);
292 releaseNative(library);
293 }
294
295 /// Whether this macOS lets an indirect render command set its own depth-stencil state and
296 /// depth bias. Bundles inherit both from the encoder either way.
297 pub fn indirectCommandsCarryDepthState(self: *State) backend.BackendError!bool {
298 if (comptime !native_available) return error.RuntimeUnavailable;
299 const rt = self.runtime.ptr() orelse return error.RuntimeUnavailable;
300 var pool = foundation.AutoreleasePool.init() orelse return error.RuntimeUnavailable;
301 defer pool.deinit();
302 const commands = raster_mod.createIndirectCommands(rt, 1) catch |err| return mapRuntimeError(err);
303 defer releaseNative(commands);
304 return metal.indirectCommandsCarryDepthState(metal.indirectRenderCommand(commands, 0));
305 }
306
307 /// The device's name, copied into `buffer` and cut to fit.
308 pub fn deviceName(self: *State, buffer: []u8) []const u8 {
309 if (comptime !native_available) return "";
310 const rt = self.runtime.ptr() orelse return "";
311 var pool = foundation.AutoreleasePool.init() orelse return "";
312 defer pool.deinit();
313 const name = metal.deviceName(rt.device) orelse return "";
314 const bytes = std.mem.span(foundation.utf8(name) orelse return "");
315 const length = @min(bytes.len, buffer.len);
316 @memcpy(buffer[0..length], bytes[0..length]);
317 return buffer[0..length];
318 }
319
320 /// The message of the last failed library compile or pipeline build.
321 pub fn compileLog(self: *State) []const u8 {
322 const rt = self.runtime.ptr() orelse return "";
323 return rt.compile_log.text();
324 }
325
326 /// How the last library compiled from MSL treated floating point.
327 pub fn mathSetting(self: *State) ?metal.MathSetting {
328 const rt = self.runtime.ptr() orelse return null;
329 return rt.math_setting;
330 }
331 };
332
333 const MetalStream = struct {
334 native: ?*runtime_mod.Stream = null,
335 };
336
337 const MetalEvent = struct {
338 native: ?*runtime_mod.Event = null,
339 recorded: bool = false,
340 };
341
342 const MetalLoadedArtifact = struct {
343 kernel: runtime_mod.LoadedKernel,
344 format: backend.ArtifactFormat,
345 };
346
347 /// A texture this backend allocated, or the texture of a drawable a surface lent it until the
348 /// frame presents.
349 const MetalTexture = struct {
350 handle: backend.TextureHandle,
351 native: raster_mod.Texture,
352 lent: bool = false,
353 };
354
355 const VertexLayout = struct {
356 binding: u32,
357 stride: u32,
358 per_instance: bool,
359 };
360
361 /// A render pipeline and the facts its draws are checked against, with the depth test and bias
362 /// that Metal keeps as encoder state rather than pipeline state.
363 const MetalPipeline = struct {
364 state: metal.RenderPipelineState,
365 depth_state: ?metal.DepthStencilState,
366 bias: backend.RenderDepthBias,
367 primitive: metal.PrimitiveType,
368 loaded: backend.LoadedRenderArtifact,
369 layouts: [max_vertex_buffers]VertexLayout,
370 bindings: [max_bindings]backend.RenderBindingDesc,
371 /// Whether an indirect command buffer may set the pipeline: only when it samples no texture.
372 indirect: bool,
373 };
374
375 const MetalResource = union(enum) {
376 /// A uniform block at buffer index `uniform_buffer_base + binding` of both stages.
377 uniform: metal.Buffer,
378 /// A texture and its sampler at index `binding` of both stages.
379 texture: struct { texture: metal.Texture, sampler: metal.SamplerState },
380 };
381
382 const MetalBinding = struct {
383 binding: u32,
384 resource: MetalResource,
385 };
386
387 const MetalBindings = struct {
388 pipeline_id: BackendObjectId,
389 entries: [max_bindings]MetalBinding,
390 count: u32,
391
392 fn slice(self: *const MetalBindings) []const MetalBinding {
393 return self.entries[0..self.count];
394 }
395 };
396
397 /// Draws of a bundle that share their encoder state: the depth test, the bias and the textures,
398 /// which indirect commands inherit from the encoder that executes them.
399 const BundleRun = struct {
400 start: u32,
401 count: u32,
402 depth_state: ?metal.DepthStencilState,
403 bias: backend.RenderDepthBias,
404 bindings: MetalBindings,
405 };
406
407 /// A pass recorded as an indirect command buffer. Each command sets its own pipeline, vertex
408 /// buffers, uniform blocks and push constants, the last from a buffer the bundle filled when it
409 /// recorded. Runs set the rest on the encoder before executing their commands.
410 /// A recorded pass. When every pipeline it draws with may join an indirect command buffer, its
411 /// draws live in `commands`, run by run. Otherwise `replay` keeps the pass and each submission
412 /// encodes it again.
413 const MetalBundle = struct {
414 replay: ?ReplayPass = null,
415 commands: ?metal.IndirectCommandBuffer,
416 push: ?metal.Buffer,
417 targets: Targets,
418 runs: []BundleRun,
419 /// Every buffer a command reads, made resident before the commands execute.
420 resources: []metal.Buffer,
421 draw_count: u32,
422 };
423
424 /// A copy of a pass that owns its draws, their vertex buffer ranges and their push constants.
425 const ReplayPass = struct {
426 pass: backend.RenderPass,
427 draws: []backend.RenderDraw,
428 ranges: []backend.RenderBufferRange,
429 bytes: []u8,
430
431 fn init(allocator: Allocator, pass: backend.RenderPass) error{OutOfMemory}!ReplayPass {
432 var range_count: usize = 0;
433 var byte_count: usize = 0;
434 for (pass.draws) |draw| {
435 range_count += draw.vertex_buffers.len;
436 byte_count += draw.push_constants.len;
437 }
438 const draws = try allocator.dupe(backend.RenderDraw, pass.draws);
439 errdefer allocator.free(draws);
440 const ranges = try allocator.alloc(backend.RenderBufferRange, range_count);
441 errdefer allocator.free(ranges);
442 const bytes = try allocator.alloc(u8, byte_count);
443 var range_next: usize = 0;
444 var byte_next: usize = 0;
445 for (draws) |*draw| {
446 const vertex_buffers = ranges[range_next..][0..draw.vertex_buffers.len];
447 @memcpy(vertex_buffers, draw.vertex_buffers);
448 draw.vertex_buffers = vertex_buffers;
449 range_next += vertex_buffers.len;
450 const push_constants = bytes[byte_next..][0..draw.push_constants.len];
451 @memcpy(push_constants, draw.push_constants);
452 draw.push_constants = push_constants;
453 byte_next += push_constants.len;
454 }
455 var copied = pass;
456 copied.draws = draws;
457 copied.diagnostic_id = null;
458 return .{ .pass = copied, .draws = draws, .ranges = ranges, .bytes = bytes };
459 }
460
461 fn deinit(self: *ReplayPass, allocator: Allocator) void {
462 allocator.free(self.draws);
463 allocator.free(self.ranges);
464 allocator.free(self.bytes);
465 }
466 };
467
468 const MetalSurface = struct {
469 surface: raster_mod.Surface,
470 handle: backend.SurfaceHandle,
471 acquired_frame: ?BackendObjectId = null,
472 acquired_texture: ?BackendObjectId = null,
473 };
474
475 const MetalFrame = struct {
476 surface_id: BackendObjectId,
477 texture_id: BackendObjectId,
478 frame: raster_mod.Frame,
479 generation: u64,
480 presented: bool = false,
481 written: bool = false,
482 };
483
484 const Object = union(enum) {
485 loaded_artifact: MetalLoadedArtifact,
486 buffer: runtime_mod.DeviceBuffer,
487 stream: MetalStream,
488 event: MetalEvent,
489 texture: MetalTexture,
490 render_pipeline: MetalPipeline,
491 render_bindings: MetalBindings,
492 render_bundle: MetalBundle,
493 surface: MetalSurface,
494 surface_frame: MetalFrame,
495 };
496
497 fn isLoadableArtifactFormat(format: backend.ArtifactFormat) bool {
498 return switch (format) {
499 .metal_msl, .metal_metallib => true,
500 else => false,
501 };
502 }
503
504 pub fn staticCapabilities(has_runtime: bool) backend.BackendCapabilities {
505 return .{
506 .identity = .{
507 .backend = .metal,
508 .family = .apple_metal,
509 .name = "apple-metal",
510 },
511 .memory = .{
512 .min_buffer_alignment = 256,
513 .unified_memory = true,
514 .host_visible_device_memory = true,
515 },
516 .subgroup = .{
517 .supported = true,
518 .size_min = 32,
519 .size_max = 32,
520 .shuffle = true,
521 .ballot = true,
522 .vote = true,
523 .arithmetic = true,
524 .scan = true,
525 },
526 .threadgroup = .{
527 .max_threads = 1024,
528 .max_blocks = .{ 65_535, 65_535, 65_535 },
529 .max_threads_per_dim = .{ 1024, 1024, 64 },
530 .max_grid_per_dim = .{ 65_535, 65_535, 65_535 },
531 },
532 .dtypes = backend.DTypeSet.init(&.{ .i1, .i32, .u32, .f16, .f32 }),
533 .layouts = .{
534 .row_major = true,
535 .compact_strides = true,
536 .broadcast_strides = true,
537 .tiled = true,
538 .opaque_backend_layouts = true,
539 },
540 .runtime = .{
541 .driver_loaded = has_runtime,
542 .device_context = has_runtime,
543 .streams = true,
544 .events = true,
545 },
546 .features = .{
547 .atomic_i32 = true,
548 .atomic_u32 = true,
549 .atomic_index = true,
550 .atomic_f32_add_device = true,
551 .async_copy = true,
552 },
553 .artifact_formats = backend.ArtifactFormatSet.init(&.{ .metal_msl, .metal_metallib }),
554 .textures = .{
555 .supported = has_runtime,
556 .formats = backend.TextureFormatSet.init(&(color_formats ++ [_]backend.TextureFormat{.depth32_float})),
557 .usages = .{
558 .copy_src = true,
559 .copy_dst = true,
560 .sampled = true,
561 .color_attachment = true,
562 .depth_attachment = true,
563 },
564 .max_extent = .{ .width = 16_384, .height = 16_384, .depth = 1 },
565 .max_sample_count = 1,
566 },
567 .raster = .{
568 .supported = has_runtime,
569 .artifact_formats = backend.RenderArtifactFormatSet.init(&.{ .metal_msl, .metal_metallib }),
570 .target_formats = backend.TextureFormatSet.init(&color_formats),
571 .depth_formats = backend.TextureFormatSet.init(&.{.depth32_float}),
572 .blend_modes = backend.RenderBlendModeSet.init(&.{ .replace, .alpha_premultiplied, .alpha_straight, .additive }),
573 .topologies = backend.RenderPrimitiveTopologySet.init(&.{ .triangle_list, .triangle_strip, .line_list, .line_strip }),
574 .vertex_formats = backend.RenderVertexFormatSet.init(&.{ .float32, .float32x2, .float32x3, .float32x4, .uint32, .uint32x2, .uint32x4 }),
575 .binding_kinds = backend.RenderBindingKindSet.init(&.{ .uniform_buffer, .sampled_texture }),
576 .index_formats = backend.RenderIndexFormatSet.init(&.{ .none, .u16, .u32 }),
577 .max_vertex_buffers = max_vertex_buffers,
578 .max_vertex_attributes = max_vertex_attributes,
579 .max_bindings = max_bindings,
580 .instancing = true,
581 .max_push_constant_bytes = max_push_constant_bytes,
582 .depth_bias = true,
583 .depth_bias_clamp = true,
584 },
585 .surfaces = .{
586 .supported = has_runtime,
587 .platforms = backend.SurfacePlatformSet.init(&.{.cocoa}),
588 .formats = backend.TextureFormatSet.init(&.{ .bgra8_unorm, .bgra8_srgb }),
589 .color_spaces = backend.ColorSpaceSet.init(&.{.srgb}),
590 .present_modes = backend.PresentModeSet.init(&.{ .fifo, .immediate }),
591 .usages = .{
592 .copy_src = true,
593 .copy_dst = true,
594 .color_attachment = true,
595 .present = true,
596 },
597 .max_extent = .{ .width = 16_384, .height = 16_384 },
598 .max_frames_in_flight = 3,
599 },
600 };
601 }
602
603 const color_formats = [_]backend.TextureFormat{ .rgba8_unorm, .bgra8_unorm, .rgba8_srgb, .bgra8_srgb };
604
605 fn pixelFormat(format: backend.TextureFormat) metal.PixelFormat {
606 return switch (format) {
607 .rgba8_unorm => .rgba8_unorm,
608 .bgra8_unorm => .bgra8_unorm,
609 .rgba8_srgb => .rgba8_unorm_srgb,
610 .bgra8_srgb => .bgra8_unorm_srgb,
611 .depth32_float => .depth32_float,
612 };
613 }
614
615 fn queryCapabilities(ptr: *anyopaque) backend.BackendError!backend.BackendCapabilities {
616 const state: *State = @ptrCast(@alignCast(ptr));
617 return staticCapabilities(state.runtime.ptr() != null);
618 }
619
620 fn createArtifact(ptr: *anyopaque, request: backend.CompileRequest) backend.BackendError!backend.KernelArtifact {
621 if (request.requested_format != .metal_msl) return error.UnsupportedOperation;
622 const state: *State = @ptrCast(@alignCast(ptr));
623 return switch (request.payload) {
624 .text => |source| createMslArtifact(state, request, source),
625 .bytes => |source| createMslArtifact(state, request, source),
626 .none => error.UnsupportedOperation,
627 .words_u32 => error.UnsupportedArtifactFormat,
628 };
629 }
630
631 fn createMslArtifact(
632 state: *State,
633 request: backend.CompileRequest,
634 source: []const u8,
635 ) backend.BackendError!backend.KernelArtifact {
636 if (request.kernel_name.len == 0) return error.InvalidArtifact;
637 if (source.len == 0) return error.InvalidArtifact;
638
639 var artifact = backend.KernelArtifact.init(state.allocator, .{
640 .backend = .metal,
641 .format = .metal_msl,
642 .entry_name = request.kernel_name,
643 .argument_count = request.argument_count,
644 .scalar_argument_count = request.scalar_argument_count,
645 .diagnostic_id = request.diagnostic_id,
646 }) catch return error.OutOfMemory;
647 errdefer artifact.deinit();
648 try artifact.setOwnedText(source);
649 return artifact;
650 }
651
652 fn loadArtifact(ptr: *anyopaque, artifact: *const backend.KernelArtifact) backend.BackendError!backend.LoadedArtifact {
653 if (artifact.backend != .metal) return error.CapabilityMismatch;
654 if (artifact.entry_name.len == 0) return error.InvalidArtifact;
655 if (!isLoadableArtifactFormat(artifact.format)) return error.UnsupportedArtifactFormat;
656
657 const state: *State = @ptrCast(@alignCast(ptr));
658 const rt = state.runtime.ptr() orelse return error.RuntimeUnavailable;
659 const kernel = switch (artifact.format) {
660 .metal_msl => switch (artifact.payload) {
661 .text => |source| if (source.len == 0)
662 return error.InvalidArtifact
663 else
664 rt.loadMsl(
665 source,
666 artifact.entry_name,
667 artifact.argument_count,
668 ) catch |err| return mapRuntimeError(err),
669 else => return error.InvalidArtifact,
670 },
671 .metal_metallib => switch (artifact.payload) {
672 .bytes => |bytes| if (bytes.len == 0)
673 return error.InvalidArtifact
674 else
675 rt.loadMetallib(
676 bytes,
677 artifact.entry_name,
678 artifact.argument_count,
679 ) catch |err| return mapRuntimeError(err),
680 else => return error.InvalidArtifact,
681 },
682 else => unreachable,
683 };
684 errdefer {
685 var owned = kernel;
686 owned.deinit();
687 }
688
689 const id = try state.putObject(.{ .loaded_artifact = .{
690 .kernel = kernel,
691 .format = artifact.format,
692 } });
693 return .{
694 .id = id,
695 .backend = .metal,
696 .format = artifact.format,
697 };
698 }
699
700 fn allocateBuffer(ptr: *anyopaque, request: backend.BufferAllocation) backend.BackendError!backend.BufferHandle {
701 if (request.byte_size == 0) return error.InvalidBuffer;
702 const state: *State = @ptrCast(@alignCast(ptr));
703 const rt = state.runtime.ptr() orelse return error.RuntimeUnavailable;
704
705 const buffer = rt.allocBuffer(request.byte_size) catch |err| return mapRuntimeError(err);
706 errdefer {
707 var owned = buffer;
708 owned.deinit();
709 }
710
711 const id = try state.putObject(.{ .buffer = buffer });
712 return .{
713 .id = id,
714 .backend = .metal,
715 .byte_size = request.byte_size,
716 .ownership = .backend,
717 };
718 }
719
720 fn createStream(ptr: *anyopaque, _: backend.StreamAllocation) backend.BackendError!backend.StreamHandle {
721 const state: *State = @ptrCast(@alignCast(ptr));
722 const rt = state.runtime.ptr() orelse return error.RuntimeUnavailable;
723 const stream = rt.createStream() catch |err| return mapRuntimeError(err);
724 errdefer rt.destroyStream(stream);
725
726 const id = try state.putObject(.{ .stream = .{ .native = stream } });
727 return .{
728 .id = id,
729 .backend = .metal,
730 };
731 }
732
733 fn createEvent(ptr: *anyopaque, _: backend.EventAllocation) backend.BackendError!backend.EventHandle {
734 const state: *State = @ptrCast(@alignCast(ptr));
735 const rt = state.runtime.ptr() orelse return error.RuntimeUnavailable;
736 const event = rt.createEvent() catch |err| return mapRuntimeError(err);
737 errdefer rt.destroyEvent(event);
738
739 const id = try state.putObject(.{ .event = .{ .native = event } });
740 return .{
741 .id = id,
742 .backend = .metal,
743 };
744 }
745
746 fn writeBuffer(ptr: *anyopaque, request: backend.BufferWriteRequest) backend.BackendError!void {
747 if (request.handle.backend != .metal) return error.InvalidBuffer;
748 const state: *State = @ptrCast(@alignCast(ptr));
749 _ = state.runtime.ptr() orelse return error.RuntimeUnavailable;
750 const buffer = try state.getBuffer(request.handle);
751 if (request.bytes.len > buffer.byte_size) return error.InvalidBuffer;
752 buffer.write(request.bytes) catch |err| return mapRuntimeError(err);
753 }
754
755 fn readBuffer(ptr: *anyopaque, request: backend.BufferReadRequest) backend.BackendError!void {
756 if (request.handle.backend != .metal) return error.InvalidBuffer;
757 const state: *State = @ptrCast(@alignCast(ptr));
758 _ = state.runtime.ptr() orelse return error.RuntimeUnavailable;
759 const buffer = try state.getBuffer(request.handle);
760 if (request.bytes.len < buffer.byte_size) return error.ReadBufferDestinationTooSmall;
761 if (request.bytes.len > buffer.byte_size) return error.InvalidBuffer;
762 buffer.read(request.bytes) catch |err| return mapRuntimeError(err);
763 }
764
765 fn launch(ptr: *anyopaque, request: backend.LaunchRequest) backend.BackendError!void {
766 if (request.artifact.backend != .metal) return error.CapabilityMismatch;
767 if (!isLoadableArtifactFormat(request.artifact.format)) return error.UnsupportedArtifactFormat;
768
769 const state: *State = @ptrCast(@alignCast(ptr));
770 const rt = state.runtime.ptr() orelse return error.RuntimeUnavailable;
771 const loaded_handle = request.loaded_artifact orelse return error.InvalidArtifact;
772 const loaded = try state.getLoaded(loaded_handle);
773 const argument_count = request.buffers.len + request.scalar_arguments.len;
774 if (argument_count != @as(usize, @intCast(loaded.argument_count))) {
775 return error.LaunchArgumentMismatch;
776 }
777 if (argument_count > runtime_mod.max_kernel_arguments) return error.LaunchArgumentMismatch;
778
779 var buffers: [runtime_mod.max_kernel_arguments]*runtime_mod.DeviceBuffer = undefined;
780 for (request.buffers, 0..) |binding, i| {
781 if (binding.handle.backend != .metal or binding.ownership != .backend) {
782 return error.InvalidBuffer;
783 }
784 const buffer = try state.getBuffer(binding.handle);
785 if (binding.byte_size > buffer.byte_size or binding.handle.byte_size != buffer.byte_size) {
786 return error.InvalidBuffer;
787 }
788 buffers[i] = buffer;
789 }
790 var scalars: [runtime_mod.max_kernel_arguments]runtime_mod.ScalarArgument = undefined;
791 for (request.scalar_arguments, 0..) |arg, i| {
792 scalars[i] = lowerScalarArgument(arg);
793 }
794
795 const stream = if (request.stream) |stream_handle|
796 try state.getStream(stream_handle)
797 else
798 null;
799 if (request.wait_events.len > runtime_mod.max_kernel_arguments) return error.LaunchArgumentMismatch;
800 var wait_events: [runtime_mod.max_kernel_arguments]*runtime_mod.Event = undefined;
801 for (request.wait_events, 0..) |event_handle, i| {
802 const event = try state.getEvent(event_handle);
803 if (!event.recorded) return error.InvalidEvent;
804 wait_events[i] = event.native orelse return error.InvalidEvent;
805 }
806 const signal_event = if (request.signal_event) |event_handle|
807 try state.getEvent(event_handle)
808 else
809 null;
810
811 rt.launch(
812 loaded,
813 buffers[0..request.buffers.len],
814 scalars[0..request.scalar_arguments.len],
815 if (stream) |value| value.native orelse return error.InvalidStream else null,
816 wait_events[0..request.wait_events.len],
817 if (signal_event) |event| event.native orelse return error.InvalidEvent else null,
818 request.geometry.grid,
819 request.geometry.threadgroup,
820 ) catch |err| return mapRuntimeError(err);
821
822 if (signal_event) |event| event.recorded = true;
823 }
824
825 fn lowerScalarArgument(arg: choir_abi.ScalarArgument) runtime_mod.ScalarArgument {
826 return switch (arg) {
827 .i32 => |value| .{ .tag = .i32, .bits = @as(u32, @bitCast(value)) },
828 .u32 => |value| .{ .tag = .u32, .bits = value },
829 .i64 => |value| .{ .tag = .i64, .bits = @bitCast(value) },
830 .u64 => |value| .{ .tag = .u64, .bits = value },
831 .f32 => |value| .{ .tag = .f32, .bits = @as(u32, @bitCast(value)) },
832 .f64 => |value| .{ .tag = .f64, .bits = @bitCast(value) },
833 };
834 }
835
836 fn synchronize(ptr: *anyopaque, request: backend.SyncRequest) backend.BackendError!void {
837 const state: *State = @ptrCast(@alignCast(ptr));
838 switch (request.scope) {
839 .default_stream => {
840 const rt = state.runtime.ptr() orelse return error.RuntimeUnavailable;
841 const stream = rt.defaultStream() catch |err| return mapRuntimeError(err);
842 stream.synchronize() catch |err| return mapRuntimeError(err);
843 },
844 .device => {
845 const rt = state.runtime.ptr() orelse return error.RuntimeUnavailable;
846 return rt.synchronize() catch |err| mapRuntimeError(err);
847 },
848 .stream => {
849 _ = state.runtime.ptr() orelse return error.RuntimeUnavailable;
850 const stream = try state.getStream(request.stream.?);
851 if (stream.native) |native_stream| {
852 native_stream.synchronize() catch |err| return mapRuntimeError(err);
853 }
854 },
855 .event => {
856 _ = state.runtime.ptr() orelse return error.RuntimeUnavailable;
857 const event = try state.getEvent(request.event.?);
858 if (!event.recorded) return error.InvalidEvent;
859 if (event.native) |native_event| {
860 native_event.synchronize() catch |err| return mapRuntimeError(err);
861 }
862 },
863 }
864 }
865
866 fn queryEvent(ptr: *anyopaque, request: backend.EventQueryRequest) backend.BackendError!bool {
867 const state: *State = @ptrCast(@alignCast(ptr));
868 _ = state.runtime.ptr() orelse return error.RuntimeUnavailable;
869 const event = try state.getEvent(request.event);
870 if (event.native) |native_event| {
871 return native_event.query() catch |err| return mapRuntimeError(err);
872 }
873 return event.recorded;
874 }
875
876 fn recordEvent(ptr: *anyopaque, request: backend.EventRecordRequest) backend.BackendError!void {
877 const state: *State = @ptrCast(@alignCast(ptr));
878 _ = state.runtime.ptr() orelse return error.RuntimeUnavailable;
879 const stream = try state.getStream(request.stream);
880 const event = try state.getEvent(request.event);
881 if (stream.native) |native_stream| {
882 const native_event = event.native orelse return error.InvalidEvent;
883 native_event.record(native_stream) catch |err| return mapRuntimeError(err);
884 }
885 event.recorded = true;
886 }
887
888 fn destroyObject(ptr: *anyopaque, id: BackendObjectId) void {
889 const state: *State = @ptrCast(@alignCast(ptr));
890 if (state.objects.fetchRemove(id)) |entry| {
891 var object = entry.value;
892 releaseLinked(state, &object);
893 deinitObject(state, &object);
894 }
895 }
896
897 /// Drops the objects that live only through `object`: a surface's unpresented frame and the
898 /// texture that frame lent.
899 fn releaseLinked(state: *State, object: *Object) void {
900 switch (object.*) {
901 .surface => |*surface| {
902 if (surface.acquired_texture) |texture_id| _ = state.objects.remove(texture_id);
903 if (surface.acquired_frame) |frame_id| if (state.objects.fetchRemove(frame_id)) |frame_entry| {
904 var frame_object = frame_entry.value;
905 deinitObject(state, &frame_object);
906 };
907 },
908 .surface_frame => |*frame| if (!frame.presented) {
909 _ = state.objects.remove(frame.texture_id);
910 if (state.getSurfaceById(frame.surface_id)) |surface| {
911 surface.acquired_frame = null;
912 surface.acquired_texture = null;
913 }
914 },
915 else => {},
916 }
917 }
918
919 fn deinitHandle(ptr: *anyopaque, allocator: Allocator) void {
920 _ = allocator;
921 const state: *State = @ptrCast(@alignCast(ptr));
922 state.deinit();
923 }
924
925 fn deinitObject(state: *State, object: *Object) void {
926 switch (object.*) {
927 .loaded_artifact => |*artifact| artifact.kernel.deinit(),
928 .buffer => |*buffer| buffer.deinit(),
929 .stream => |stream| if (stream.native) |native_stream| {
930 if (state.runtime.ptr()) |rt| rt.destroyStream(native_stream);
931 },
932 .event => |event| if (event.native) |native_event| {
933 if (state.runtime.ptr()) |rt| rt.destroyEvent(native_event);
934 },
935 .texture => |*texture| if (!texture.lent) texture.native.deinit(),
936 .render_pipeline => |pipeline| {
937 releaseNative(pipeline.state);
938 if (pipeline.depth_state) |depth_state| releaseNative(depth_state);
939 },
940 .render_bindings => {},
941 .render_bundle => |*bundle| {
942 if (bundle.replay) |*replay| replay.deinit(state.allocator);
943 if (bundle.commands) |commands| releaseNative(commands);
944 if (bundle.push) |push| releaseNative(push);
945 state.allocator.free(bundle.runs);
946 state.allocator.free(bundle.resources);
947 },
948 .surface => |*surface| surface.surface.deinit(),
949 .surface_frame => |*frame| if (!frame.presented) frame.frame.release(),
950 }
951 object.* = undefined;
952 }
953
954 fn allocateTexture(ptr: *anyopaque, request: backend.TextureAllocation) backend.BackendError!backend.TextureHandle {
955 const state: *State = @ptrCast(@alignCast(ptr));
956 const rt = state.runtime.ptr() orelse return error.RuntimeUnavailable;
957 if (request.sample_count != 1 or request.extent.depth != 1) return error.CapabilityMismatch;
958 if (request.usage.present or request.usage.storage) return error.CapabilityMismatch;
959 const depth = request.format.isDepth();
960 if (depth and request.usage.color_attachment) return error.CapabilityMismatch;
961 if (!depth and request.usage.depth_attachment) return error.CapabilityMismatch;
962 try state.objects.ensureUnusedCapacity(state.allocator, 1);
963 var usage: metal.TextureUsage = 0;
964 if (request.usage.sampled) usage |= metal.texture_usage_shader_read;
965 if (request.usage.color_attachment or request.usage.depth_attachment) usage |= metal.texture_usage_render_target;
966 var native = raster_mod.Texture.create(
967 rt,
968 pixelFormat(request.format),
969 request.extent.width,
970 request.extent.height,
971 request.format.texelBytes(),
972 usage,
973 ) catch |err| return mapRuntimeError(err);
974 errdefer native.deinit();
975 var texture = backend.TextureHandle{
976 .id = 0,
977 .backend = .metal,
978 .extent = request.extent,
979 .format = request.format,
980 .usage = request.usage,
981 .sample_count = 1,
982 .ownership = .backend,
983 };
984 texture.id = try state.putObject(.{ .texture = .{ .handle = texture, .native = native } });
985 state.objects.getPtr(texture.id).?.texture.handle.id = texture.id;
986 return texture;
987 }
988
989 fn destroyTexture(ptr: *anyopaque, handle: backend.TextureHandle) backend.BackendError!void {
990 const state: *State = @ptrCast(@alignCast(ptr));
991 const texture = try state.getTexture(handle);
992 if (texture.lent) return error.SurfaceAlreadyAcquired;
993 const entry = state.objects.fetchRemove(handle.id).?;
994 var object = entry.value;
995 deinitObject(state, &object);
996 }
997
998 fn writeTexture(ptr: *anyopaque, request: backend.TextureWriteRequest) backend.BackendError!void {
999 const state: *State = @ptrCast(@alignCast(ptr));
1000 const rt = state.runtime.ptr() orelse return error.RuntimeUnavailable;
1001 const texture = try state.getTexture(request.texture);
1002 const native = &texture.native;
1003 raster_mod.transferTexture(rt, native.texture, native.size, native.texel_bytes, .{ .write = request.bytes }) catch |err| {
1004 return mapRuntimeError(err);
1005 };
1006 }
1007
1008 fn readTexture(ptr: *anyopaque, request: backend.TextureReadRequest) backend.BackendError!void {
1009 const state: *State = @ptrCast(@alignCast(ptr));
1010 const rt = state.runtime.ptr() orelse return error.RuntimeUnavailable;
1011 const texture = try state.getTexture(request.texture);
1012 const native = &texture.native;
1013 raster_mod.transferTexture(rt, native.texture, native.size, native.texel_bytes, .{ .read = request.bytes }) catch |err| {
1014 return mapRuntimeError(err);
1015 };
1016 }
1017
1018 fn createRenderArtifact(ptr: *anyopaque, desc: backend.RenderPipelineDesc) backend.BackendError!backend.RenderArtifact {
1019 const state: *State = @ptrCast(@alignCast(ptr));
1020 var artifact = backend.RenderArtifact.init(state.allocator, .{
1021 .backend = .metal,
1022 .pipeline = desc,
1023 }) catch return error.OutOfMemory;
1024 errdefer artifact.deinit();
1025 switch (desc.format) {
1026 .metal_msl => switch (desc.payload) {
1027 .text => |text| if (text.len == 0) return error.InvalidRenderArtifact else try artifact.setOwnedText(text),
1028 else => return error.InvalidRenderArtifact,
1029 },
1030 .metal_metallib => switch (desc.payload) {
1031 .bytes => |bytes| if (bytes.len == 0) return error.InvalidRenderArtifact else try artifact.setOwnedBytes(bytes),
1032 else => return error.InvalidRenderArtifact,
1033 },
1034 else => return error.UnsupportedArtifactFormat,
1035 }
1036 return artifact;
1037 }
1038
1039 /// Builds the pipeline an artifact describes. Vertex layout `binding` reads Metal buffer index
1040 /// `binding`, below 8, and attribute `location` is `[[attribute(location)]]`. A uniform block of
1041 /// binding `b` sits at buffer index 16 + b, below 24, and a sampled texture at texture and
1042 /// sampler index `b`, below 16, all in group 0.
1043 fn loadRenderArtifact(ptr: *anyopaque, artifact: *const backend.RenderArtifact) backend.BackendError!backend.LoadedRenderArtifact {
1044 const state: *State = @ptrCast(@alignCast(ptr));
1045 if (artifact.backend != .metal) return error.CapabilityMismatch;
1046 const library: runtime_mod.LibrarySource = switch (artifact.format) {
1047 .metal_msl => switch (artifact.payload) {
1048 .text => |text| .{ .msl = text },
1049 else => return error.InvalidRenderArtifact,
1050 },
1051 .metal_metallib => switch (artifact.payload) {
1052 .bytes => |bytes| .{ .metallib = bytes },
1053 else => return error.InvalidRenderArtifact,
1054 },
1055 else => return error.UnsupportedArtifactFormat,
1056 };
1057 const rt = state.runtime.ptr() orelse return error.RuntimeUnavailable;
1058 if (artifact.vertex_layouts.len > max_vertex_buffers) return error.CapabilityMismatch;
1059 if (artifact.bindings.len > max_bindings) return error.CapabilityMismatch;
1060 if (artifact.push_constant_bytes > max_push_constant_bytes) return error.CapabilityMismatch;
1061 try state.objects.ensureUnusedCapacity(state.allocator, 1);
1062
1063 var object: MetalPipeline = .{
1064 .state = undefined,
1065 .depth_state = null,
1066 .bias = if (artifact.depth) |depth| depth.bias else .{},
1067 .primitive = metalPrimitive(artifact.topology),
1068 .loaded = undefined,
1069 .layouts = undefined,
1070 .bindings = undefined,
1071 .indirect = for (artifact.bindings) |binding| {
1072 if (binding.kind == .sampled_texture) break false;
1073 } else true,
1074 };
1075 var layouts: [max_vertex_buffers]raster_mod.VertexLayout = undefined;
1076 var attributes: [max_vertex_attributes]raster_mod.VertexAttribute = undefined;
1077 var attribute_count: usize = 0;
1078 var used_buffers: u32 = 0;
1079 for (artifact.vertex_layouts, 0..) |layout, index| {
1080 if (layout.binding >= max_vertex_buffers) return error.CapabilityMismatch;
1081 const bit = @as(u32, 1) << @intCast(layout.binding);
1082 if (used_buffers & bit != 0) return error.InvalidRenderArtifact;
1083 used_buffers |= bit;
1084 const per_instance = layout.step_mode == .instance;
1085 layouts[index] = .{ .buffer_index = layout.binding, .stride = layout.stride, .per_instance = per_instance };
1086 object.layouts[index] = .{ .binding = layout.binding, .stride = layout.stride, .per_instance = per_instance };
1087 const start: usize = layout.attribute_start;
1088 const count: usize = layout.attribute_count;
1089 if (start > artifact.vertex_attributes.len or count > artifact.vertex_attributes.len - start) return error.InvalidRenderArtifact;
1090 for (artifact.vertex_attributes[start..][0..count]) |attribute| {
1091 if (attribute_count == max_vertex_attributes) return error.CapabilityMismatch;
1092 if (attribute.location >= max_attribute_location) return error.CapabilityMismatch;
1093 attributes[attribute_count] = .{
1094 .location = attribute.location,
1095 .format = vertexFormat(attribute.format),
1096 .offset = attribute.offset,
1097 .buffer_index = layout.binding,
1098 };
1099 attribute_count += 1;
1100 }
1101 }
1102 var used_bindings: u32 = 0;
1103 for (artifact.bindings, 0..) |binding, index| {
1104 if (binding.group != 0) return error.CapabilityMismatch;
1105 const limit: u32 = switch (binding.kind) {
1106 .uniform_buffer => choir_abi.metal.max_uniform_bindings,
1107 .sampled_texture => choir_abi.metal.max_texture_bindings,
1108 .storage_buffer, .storage_texture => return error.CapabilityMismatch,
1109 };
1110 if (binding.binding >= limit) return error.CapabilityMismatch;
1111 const bit = @as(u32, 1) << @intCast(binding.binding);
1112 if (used_bindings & bit != 0) return error.InvalidRenderArtifact;
1113 used_bindings |= bit;
1114 object.bindings[index] = binding;
1115 }
1116 if (artifact.depth) |depth| {
1117 if (depth.format != .depth32_float) return error.CapabilityMismatch;
1118 }
1119
1120 object.state = raster_mod.createPipelineState(rt, .{
1121 .library = library,
1122 .vertex_entry = artifact.vertex_entry_name,
1123 .fragment_entry = artifact.fragment_entry_name,
1124 .color_format = pixelFormat(artifact.target_format),
1125 .blend = metalBlend(artifact.blend_mode),
1126 .depth_format = if (artifact.depth != null) .depth32_float else null,
1127 .layouts = layouts[0..artifact.vertex_layouts.len],
1128 .attributes = attributes[0..attribute_count],
1129 .indirect = object.indirect,
1130 }) catch |err| return mapRuntimeError(err);
1131 errdefer releaseNative(object.state);
1132 if (artifact.depth) |depth| {
1133 object.depth_state = raster_mod.createDepthState(rt, metalCompare(depth.compare), depth.write) catch |err| {
1134 return mapRuntimeError(err);
1135 };
1136 }
1137 errdefer if (object.depth_state) |depth_state| releaseNative(depth_state);
1138
1139 const id = try state.putObject(.{ .render_pipeline = object });
1140 const loaded = backend.LoadedRenderArtifact.describing(artifact, id);
1141 state.objects.getPtr(id).?.render_pipeline.loaded = loaded;
1142 return loaded;
1143 }
1144
1145 fn createRenderBindings(ptr: *anyopaque, request: backend.RenderBindingsRequest) backend.BackendError!backend.RenderBindings {
1146 const state: *State = @ptrCast(@alignCast(ptr));
1147 const rt = state.runtime.ptr() orelse return error.RuntimeUnavailable;
1148 try state.objects.ensureUnusedCapacity(state.allocator, 1);
1149 const pipeline = try state.getRenderPipeline(request.pipeline);
1150 const count: usize = pipeline.loaded.binding_count;
1151 if (request.resources.len != count) return error.RenderArgumentMismatch;
1152 var set = MetalBindings{ .pipeline_id = request.pipeline.id, .entries = undefined, .count = @intCast(count) };
1153 for (request.resources, pipeline.bindings[0..count], set.entries[0..count]) |resource, binding, *entry| {
1154 if (std.meta.activeTag(resource) != binding.kind) return error.RenderArgumentMismatch;
1155 entry.* = .{ .binding = binding.binding, .resource = switch (resource) {
1156 .uniform_buffer => |buffer_handle| .{ .uniform = (try state.getBuffer(buffer_handle)).buffer },
1157 .sampled_texture => |sampled| blk: {
1158 const texture = try state.getTexture(sampled.texture);
1159 if (!texture.handle.usage.sampled) return error.InvalidTexture;
1160 const sampler = rt.sampler(
1161 switch (sampled.sampler.filter) {
1162 .nearest => .nearest,
1163 .linear => .linear,
1164 },
1165 switch (sampled.sampler.address) {
1166 .clamp_to_edge => .clamp_to_edge,
1167 .repeat => .repeat,
1168 },
1169 ) catch |err| return mapRuntimeError(err);
1170 break :blk .{ .texture = .{ .texture = texture.native.texture, .sampler = sampler } };
1171 },
1172 .storage_buffer, .storage_texture => return error.CapabilityMismatch,
1173 } };
1174 }
1175 const id = try state.putObject(.{ .render_bindings = set });
1176 return .{ .id = id, .backend = .metal, .pipeline = request.pipeline.id };
1177 }
1178
1179 fn render(ptr: *anyopaque, request: backend.RenderRequest) backend.BackendError!void {
1180 if (comptime !native_available) return error.RuntimeUnavailable;
1181 const state: *State = @ptrCast(@alignCast(ptr));
1182 const rt = state.runtime.ptr() orelse return error.RuntimeUnavailable;
1183 const stream = try state.nativeStream(request.stream);
1184 var wait_storage: [max_wait_events]*runtime_mod.Event = undefined;
1185 const waits = try state.nativeWaits(request.wait_events, &wait_storage);
1186 const signal = try state.nativeSignal(request.signal_event);
1187 const targets = try passTargets(state, request.pass);
1188
1189 var pool = foundation.AutoreleasePool.init() orelse return error.RuntimeUnavailable;
1190 defer pool.deinit();
1191 var submission = raster_mod.Submission.begin(rt, stream, waits) catch |err| return mapRenderError(err);
1192 const encoder = try beginPass(submission.command_buffer, targets);
1193 var sink = EncoderSink{ .encoder = encoder };
1194 recordDraws(state, request.pass, &sink) catch |err| {
1195 metal.endRenderEncoding(encoder);
1196 return err;
1197 };
1198 metal.endRenderEncoding(encoder);
1199 submission.finish(if (signal) |event| event.native.? else null) catch |err| return mapRenderError(err);
1200 if (signal) |event| event.recorded = true;
1201 }
1202
1203 fn recordRenderBundle(ptr: *anyopaque, pass: backend.RenderPass) backend.BackendError!backend.RenderBundle {
1204 if (comptime !native_available) return error.RuntimeUnavailable;
1205 const state: *State = @ptrCast(@alignCast(ptr));
1206 const rt = state.runtime.ptr() orelse return error.RuntimeUnavailable;
1207 try state.objects.ensureUnusedCapacity(state.allocator, 1);
1208 const draw_count: u32 = @intCast(pass.draws.len);
1209 const targets = try passTargets(state, pass);
1210 const indirect = for (pass.draws) |draw| {
1211 if (!(try state.getRenderPipeline(draw.pipeline)).indirect) break false;
1212 } else true;
1213 if (!indirect) {
1214 var check = CheckSink{};
1215 try recordDraws(state, pass, &check);
1216 var replay = ReplayPass.init(state.allocator, pass) catch return error.OutOfMemory;
1217 errdefer replay.deinit(state.allocator);
1218 const id = try state.putObject(.{ .render_bundle = .{
1219 .replay = replay,
1220 .commands = null,
1221 .push = null,
1222 .targets = targets,
1223 .runs = &.{},
1224 .resources = &.{},
1225 .draw_count = draw_count,
1226 } });
1227 return .{ .id = id, .backend = .metal, .draw_count = draw_count };
1228 }
1229
1230 var pool = foundation.AutoreleasePool.init() orelse return error.RuntimeUnavailable;
1231 defer pool.deinit();
1232 const commands = if (draw_count == 0) null else raster_mod.createIndirectCommands(rt, draw_count) catch |err| {
1233 return mapRenderError(err);
1234 };
1235 errdefer if (commands) |owned| releaseNative(owned);
1236 var push_bytes: usize = 0;
1237 for (pass.draws) |draw| {
1238 if (draw.push_constants.len != 0) push_bytes = @as(usize, draw_count) * push_stride;
1239 }
1240 const push = if (push_bytes == 0) null else raster_mod.createSharedBuffer(rt, push_bytes) catch |err| {
1241 return mapRenderError(err);
1242 };
1243 errdefer if (push) |owned| releaseNative(owned);
1244
1245 var sink = BundleSink{
1246 .allocator = state.allocator,
1247 .commands = commands,
1248 .push = push,
1249 .push_contents = if (push) |buffer| @ptrCast(metal.bufferContents(buffer) orelse return error.RenderFailed) else null,
1250 .runs = std.ArrayList(BundleRun).initCapacity(state.allocator, pass.draws.len) catch return error.OutOfMemory,
1251 };
1252 defer sink.deinit();
1253 if (push) |buffer| sink.addResource(buffer) catch return error.OutOfMemory;
1254 try recordDraws(state, pass, &sink);
1255
1256 const runs = sink.runs.toOwnedSlice(state.allocator) catch return error.OutOfMemory;
1257 errdefer state.allocator.free(runs);
1258 const resources = state.allocator.dupe(metal.Buffer, sink.resources.keys()) catch return error.OutOfMemory;
1259 errdefer state.allocator.free(resources);
1260 const id = try state.putObject(.{ .render_bundle = .{
1261 .commands = commands,
1262 .push = push,
1263 .targets = targets,
1264 .runs = runs,
1265 .resources = resources,
1266 .draw_count = draw_count,
1267 } });
1268 return .{ .id = id, .backend = .metal, .draw_count = draw_count };
1269 }
1270
1271 fn submitRenderBundle(ptr: *anyopaque, request: backend.RenderBundleSubmit) backend.BackendError!void {
1272 if (comptime !native_available) return error.RuntimeUnavailable;
1273 const state: *State = @ptrCast(@alignCast(ptr));
1274 const rt = state.runtime.ptr() orelse return error.RuntimeUnavailable;
1275 const bundle = try state.getRenderBundle(request.bundle);
1276 const stream = try state.nativeStream(request.stream);
1277 var wait_storage: [max_wait_events]*runtime_mod.Event = undefined;
1278 const waits = try state.nativeWaits(request.wait_events, &wait_storage);
1279 const signal = try state.nativeSignal(request.signal_event);
1280
1281 var pool = foundation.AutoreleasePool.init() orelse return error.RuntimeUnavailable;
1282 defer pool.deinit();
1283 var submission = raster_mod.Submission.begin(rt, stream, waits) catch |err| return mapRenderError(err);
1284 const encoder = try beginPass(submission.command_buffer, bundle.targets);
1285 if (bundle.replay) |replay| {
1286 var sink = EncoderSink{ .encoder = encoder };
1287 recordDraws(state, replay.pass, &sink) catch |err| {
1288 metal.endRenderEncoding(encoder);
1289 return err;
1290 };
1291 }
1292 metal.useResources(encoder, bundle.resources);
1293 for (bundle.runs) |run| {
1294 setEncoderState(encoder, run.depth_state, run.bias);
1295 for (run.bindings.slice()) |entry| switch (entry.resource) {
1296 .uniform => {},
1297 .texture => |bound| {
1298 metal.setVertexTexture(encoder, bound.texture, bound.sampler, entry.binding);
1299 metal.setFragmentTexture(encoder, bound.texture, bound.sampler, entry.binding);
1300 },
1301 };
1302 metal.executeCommands(encoder, bundle.commands.?, .{ .location = run.start, .length = run.count });
1303 }
1304 metal.endRenderEncoding(encoder);
1305 submission.finish(if (signal) |event| event.native.? else null) catch |err| return mapRenderError(err);
1306 if (signal) |event| event.recorded = true;
1307 }
1308
1309 /// The attachments, viewport and scissor of a pass, in Metal's terms.
1310 const Targets = struct {
1311 color: metal.Texture,
1312 color_load: metal.LoadAction,
1313 clear: metal.ClearColor,
1314 depth: ?metal.Texture,
1315 depth_load: metal.LoadAction,
1316 clear_depth: f64,
1317 viewport: metal.Viewport,
1318 scissor: metal.ScissorRect,
1319 };
1320
1321 fn passTargets(state: *State, pass: backend.RenderPass) backend.BackendError!Targets {
1322 const color = try state.getTexture(pass.color.view.texture);
1323 const depth = if (pass.depth) |attachment| try state.getTexture(attachment.view.texture) else null;
1324 const clear: backend.SurfaceClearColor = switch (pass.color.load) {
1325 .load => .{},
1326 .clear => |value| value,
1327 };
1328 return .{
1329 .color = color.native.texture,
1330 .color_load = switch (pass.color.load) {
1331 .load => .load,
1332 .clear => .clear,
1333 },
1334 .clear = .{ .red = clear.r, .green = clear.g, .blue = clear.b, .alpha = clear.a },
1335 .depth = if (depth) |target| target.native.texture else null,
1336 .depth_load = if (pass.depth) |attachment| switch (attachment.load) {
1337 .load => .load,
1338 .clear => .clear,
1339 } else .dont_care,
1340 .clear_depth = if (pass.depth) |attachment| switch (attachment.load) {
1341 .load => 1,
1342 .clear => |value| value,
1343 } else 1,
1344 .viewport = .{
1345 .origin_x = pass.viewport.x,
1346 .origin_y = pass.viewport.y,
1347 .width = pass.viewport.width,
1348 .height = pass.viewport.height,
1349 .znear = pass.viewport.min_depth,
1350 .zfar = pass.viewport.max_depth,
1351 },
1352 .scissor = .{
1353 .x = pass.scissor.x,
1354 .y = pass.scissor.y,
1355 .width = pass.scissor.width,
1356 .height = pass.scissor.height,
1357 },
1358 };
1359 }
1360
1361 /// Opens a render encoder on `targets`. A triangle is front-facing when counter-clockwise on
1362 /// screen, as the contract has it; Metal's default is clockwise. The caller holds an autorelease
1363 /// pool.
1364 fn beginPass(command_buffer: metal.CommandBuffer, targets: Targets) backend.BackendError!metal.RenderCommandEncoder {
1365 const descriptor = metal.renderPassDescriptor() orelse return error.RenderFailed;
1366 metal.setColorAttachment(descriptor, targets.color, targets.color_load, targets.clear);
1367 if (targets.depth) |depth| metal.setDepthAttachment(descriptor, depth, targets.depth_load, targets.clear_depth);
1368 const encoder = metal.renderCommandEncoder(command_buffer, descriptor) orelse return error.RenderFailed;
1369 metal.setViewport(encoder, targets.viewport);
1370 metal.setScissorRect(encoder, targets.scissor);
1371 metal.setFrontFacingWinding(encoder, .counter_clockwise);
1372 return encoder;
1373 }
1374
1375 fn setEncoderState(encoder: metal.RenderCommandEncoder, depth_state: ?metal.DepthStencilState, bias: backend.RenderDepthBias) void {
1376 metal.setDepthStencilState(encoder, depth_state);
1377 metal.setDepthBias(encoder, bias.constant, bias.slope, bias.clamp);
1378 }
1379
1380 /// The draw a pass issues once its state is bound.
1381 const DrawCall = struct {
1382 draw: backend.RenderDraw,
1383 primitive: metal.PrimitiveType,
1384 index: ?struct { buffer: metal.Buffer, offset: u64, format: metal.IndexType },
1385 };
1386
1387 /// Encodes a pass straight into a render encoder, binding only what changed.
1388 const EncoderSink = struct {
1389 encoder: metal.RenderCommandEncoder,
1390
1391 fn setPipeline(self: *EncoderSink, pipeline: *const MetalPipeline) void {
1392 metal.setRenderPipelineState(self.encoder, pipeline.state);
1393 setEncoderState(self.encoder, pipeline.depth_state, pipeline.bias);
1394 }
1395
1396 fn setBindings(self: *EncoderSink, bindings: *const MetalBindings) void {
1397 for (bindings.slice()) |entry| switch (entry.resource) {
1398 .uniform => |buffer| {
1399 const index = choir_abi.metal.uniform_buffer_base + entry.binding;
1400 metal.setVertexBuffer(self.encoder, buffer, 0, index);
1401 metal.setFragmentBuffer(self.encoder, buffer, 0, index);
1402 },
1403 .texture => |bound| {
1404 metal.setVertexTexture(self.encoder, bound.texture, bound.sampler, entry.binding);
1405 metal.setFragmentTexture(self.encoder, bound.texture, bound.sampler, entry.binding);
1406 },
1407 };
1408 }
1409
1410 fn setVertexBuffer(self: *EncoderSink, binding: u32, buffer: metal.Buffer, offset: u64) error{OutOfMemory}!void {
1411 metal.setVertexBuffer(self.encoder, buffer, offset, binding);
1412 }
1413
1414 fn issue(self: *EncoderSink, call: DrawCall) error{OutOfMemory}!void {
1415 const draw = call.draw;
1416 if (draw.push_constants.len != 0) {
1417 metal.setVertexBytes(self.encoder, draw.push_constants, choir_abi.metal.push_buffer);
1418 metal.setFragmentBytes(self.encoder, draw.push_constants, choir_abi.metal.push_buffer);
1419 }
1420 const range = draw.range;
1421 if (call.index) |index| {
1422 metal.drawIndexedPrimitives(self.encoder, call.primitive, range.index_count, index.format, index.buffer, index.offset, range.instance_count, range.base_vertex, range.first_instance);
1423 } else {
1424 metal.drawPrimitives(self.encoder, call.primitive, range.first_vertex, range.vertex_count, range.instance_count, range.first_instance);
1425 }
1426 }
1427 };
1428
1429 /// Takes a pass's draws and encodes nothing: `recordDraws` through it only checks the pass.
1430 const CheckSink = struct {
1431 fn setPipeline(_: *CheckSink, _: *const MetalPipeline) void {}
1432 fn setBindings(_: *CheckSink, _: *const MetalBindings) void {}
1433 fn setVertexBuffer(_: *CheckSink, _: u32, _: metal.Buffer, _: u64) error{OutOfMemory}!void {}
1434 fn issue(_: *CheckSink, _: DrawCall) error{OutOfMemory}!void {}
1435 };
1436
1437 /// Encodes a pass into an indirect command buffer. A command inherits nothing from the command
1438 /// before it, so each one sets the pipeline, vertex buffers and uniform blocks its draw reads.
1439 const BundleSink = struct {
1440 allocator: Allocator,
1441 commands: ?metal.IndirectCommandBuffer,
1442 push: ?metal.Buffer,
1443 push_contents: ?[*]u8,
1444 runs: std.ArrayList(BundleRun),
1445 resources: std.AutoArrayHashMapUnmanaged(metal.Buffer, void) = .empty,
1446 pipeline: ?*const MetalPipeline = null,
1447 bindings: ?*const MetalBindings = null,
1448 vertex_buffers: [max_vertex_buffers]?struct { buffer: metal.Buffer, offset: u64 } = @splat(null),
1449 next: u32 = 0,
1450 run_open: bool = false,
1451
1452 fn deinit(self: *BundleSink) void {
1453 self.runs.deinit(self.allocator);
1454 self.resources.deinit(self.allocator);
1455 }
1456
1457 fn addResource(self: *BundleSink, buffer: metal.Buffer) error{OutOfMemory}!void {
1458 try self.resources.put(self.allocator, buffer, {});
1459 }
1460
1461 fn setPipeline(self: *BundleSink, pipeline: *const MetalPipeline) void {
1462 self.pipeline = pipeline;
1463 self.run_open = false;
1464 }
1465
1466 fn setBindings(self: *BundleSink, bindings: *const MetalBindings) void {
1467 self.bindings = bindings;
1468 self.run_open = false;
1469 }
1470
1471 fn setVertexBuffer(self: *BundleSink, binding: u32, buffer: metal.Buffer, offset: u64) error{OutOfMemory}!void {
1472 self.vertex_buffers[binding] = .{ .buffer = buffer, .offset = offset };
1473 try self.addResource(buffer);
1474 }
1475
1476 fn issue(self: *BundleSink, call: DrawCall) error{OutOfMemory}!void {
1477 const pipeline = self.pipeline.?;
1478 const draw = call.draw;
1479 if (!self.run_open) {
1480 self.runs.appendAssumeCapacity(.{
1481 .start = self.next,
1482 .count = 0,
1483 .depth_state = pipeline.depth_state,
1484 .bias = pipeline.bias,
1485 .bindings = if (draw.bindings != null) self.bindings.?.* else .{ .pipeline_id = 0, .entries = undefined, .count = 0 },
1486 });
1487 self.run_open = true;
1488 }
1489 self.runs.items[self.runs.items.len - 1].count += 1;
1490 const command = metal.indirectRenderCommand(self.commands.?, self.next);
1491 metal.setIndirectPipelineState(command, pipeline.state);
1492 for (pipeline.layouts[0..draw.vertex_buffers.len]) |layout| {
1493 const bound = self.vertex_buffers[layout.binding].?;
1494 metal.setIndirectVertexBuffer(command, bound.buffer, bound.offset, layout.binding);
1495 }
1496 if (draw.bindings != null) for (self.bindings.?.slice()) |entry| switch (entry.resource) {
1497 .uniform => |buffer| {
1498 const index = choir_abi.metal.uniform_buffer_base + entry.binding;
1499 metal.setIndirectVertexBuffer(command, buffer, 0, index);
1500 metal.setIndirectFragmentBuffer(command, buffer, 0, index);
1501 try self.addResource(buffer);
1502 },
1503 .texture => {},
1504 };
1505 if (draw.push_constants.len != 0) {
1506 const offset = @as(u64, self.next) * push_stride;
1507 @memcpy(self.push_contents.?[@intCast(offset)..][0..draw.push_constants.len], draw.push_constants);
1508 metal.setIndirectVertexBuffer(command, self.push.?, offset, choir_abi.metal.push_buffer);
1509 metal.setIndirectFragmentBuffer(command, self.push.?, offset, choir_abi.metal.push_buffer);
1510 }
1511 const range = draw.range;
1512 if (call.index) |index| {
1513 try self.addResource(index.buffer);
1514 metal.indirectDrawIndexedPrimitives(command, call.primitive, range.index_count, index.format, index.buffer, index.offset, range.instance_count, range.base_vertex, range.first_instance);
1515 } else {
1516 metal.indirectDrawPrimitives(command, call.primitive, range.first_vertex, range.vertex_count, range.instance_count, range.first_instance);
1517 }
1518 self.next += 1;
1519 }
1520 };
1521
1522 const VertexSlot = struct {
1523 id: BackendObjectId,
1524 offset: u64,
1525 size: u64,
1526 };
1527
1528 /// Checks every draw of `pass` and hands `sink` what changed before each one, as the Vulkan
1529 /// backend's pass recording does: each range a draw reads is checked against its buffer, except
1530 /// the vertices an indexed draw reaches through its indices, which only the indices name.
1531 fn recordDraws(state: *State, pass: backend.RenderPass, sink: anytype) backend.BackendError!void {
1532 var pipeline: ?*MetalPipeline = null;
1533 var bindings_id: ?BackendObjectId = null;
1534 var vertex_slots: [max_vertex_buffers]?VertexSlot = @splat(null);
1535 for (pass.draws) |draw| {
1536 const current = if (pipeline != null and pipeline.?.loaded.id == draw.pipeline.id) pipeline.? else blk: {
1537 const next = try state.getRenderPipeline(draw.pipeline);
1538 if (next.loaded.target_format != pass.color.view.format) return error.RenderArgumentMismatch;
1539 if ((next.loaded.depth != null) != (pass.depth != null)) return error.RenderArgumentMismatch;
1540 sink.setPipeline(next);
1541 bindings_id = null;
1542 pipeline = next;
1543 break :blk next;
1544 };
1545 if (draw.vertex_buffers.len != current.loaded.vertex_buffer_count) return error.RenderArgumentMismatch;
1546 if ((draw.bindings != null) != (current.loaded.binding_count != 0)) return error.RenderArgumentMismatch;
1547 if (draw.bindings) |bindings| if (bindings_id != bindings.id) {
1548 const set = try state.getRenderBindings(bindings);
1549 if (set.pipeline_id != current.loaded.id) return error.RenderArgumentMismatch;
1550 sink.setBindings(set);
1551 bindings_id = bindings.id;
1552 };
1553 if (draw.push_constants.len != current.loaded.push_constant_bytes) return error.RenderArgumentMismatch;
1554 const indexed = draw.range.index_format != .none;
1555 for (draw.vertex_buffers, current.layouts[0..draw.vertex_buffers.len]) |range, layout| {
1556 const slot = &vertex_slots[layout.binding];
1557 const bound = if (slot.*) |existing| existing.id == range.buffer.id and existing.offset == range.offset else false;
1558 if (!bound) {
1559 const buffer = try state.getBuffer(range.buffer);
1560 if (range.offset >= buffer.byte_size) return error.RenderArgumentMismatch;
1561 sink.setVertexBuffer(layout.binding, buffer.buffer, range.offset) catch return error.OutOfMemory;
1562 slot.* = .{ .id = range.buffer.id, .offset = range.offset, .size = buffer.byte_size };
1563 }
1564 const elements: u64 = if (layout.per_instance)
1565 @as(u64, draw.range.first_instance) + draw.range.instance_count
1566 else if (indexed)
1567 0
1568 else
1569 @as(u64, draw.range.first_vertex) + draw.range.vertex_count;
1570 if (elements * layout.stride > slot.*.?.size - slot.*.?.offset) return error.RenderArgumentMismatch;
1571 }
1572 if (!indexed) {
1573 if (draw.index_buffer != null) return error.RenderArgumentMismatch;
1574 sink.issue(.{ .draw = draw, .primitive = current.primitive, .index = null }) catch return error.OutOfMemory;
1575 continue;
1576 }
1577 const range = draw.index_buffer orelse return error.RenderArgumentMismatch;
1578 const index_bytes: u64 = if (draw.range.index_format == .u16) 2 else 4;
1579 const buffer = try state.getBuffer(range.buffer);
1580 if (range.offset >= buffer.byte_size or range.offset % index_bytes != 0) return error.RenderArgumentMismatch;
1581 const indices = @as(u64, draw.range.first_index) + draw.range.index_count;
1582 if (indices * index_bytes > buffer.byte_size - range.offset) return error.RenderArgumentMismatch;
1583 sink.issue(.{ .draw = draw, .primitive = current.primitive, .index = .{
1584 .buffer = buffer.buffer,
1585 .offset = range.offset + @as(u64, draw.range.first_index) * index_bytes,
1586 .format = if (index_bytes == 2) .uint16 else .uint32,
1587 } }) catch return error.OutOfMemory;
1588 }
1589 }
1590
1591 /// Makes a `CAMetalLayer` over the layer of the Cocoa window `request` names, as
1592 /// `Window.getNativeSurface` returns it. The layer draws at `request.extent`, in pixels. FIFO
1593 /// presentation syncs to the display; immediate does not.
1594 fn createSurface(ptr: *anyopaque, request: backend.SurfaceCreationRequest) backend.BackendError!backend.SurfaceHandle {
1595 const state: *State = @ptrCast(@alignCast(ptr));
1596 const rt = state.runtime.ptr() orelse return error.RuntimeUnavailable;
1597 const cocoa = switch (request.platform) {
1598 .cocoa => |native| native,
1599 else => return error.CapabilityMismatch,
1600 };
1601 if (cocoa.layer == 0) return error.InvalidSurface;
1602 const format: metal.PixelFormat = switch (request.format) {
1603 .bgra8_unorm, .bgra8_srgb => pixelFormat(request.format),
1604 else => return error.CapabilityMismatch,
1605 };
1606 const display_sync = switch (request.present_mode) {
1607 .fifo => true,
1608 .immediate => false,
1609 .mailbox => return error.CapabilityMismatch,
1610 };
1611 try state.objects.ensureUnusedCapacity(state.allocator, 1);
1612 var surface = raster_mod.Surface.create(
1613 rt,
1614 @ptrFromInt(cocoa.layer),
1615 format,
1616 request.extent.width,
1617 request.extent.height,
1618 std.math.clamp(request.max_frames_in_flight, 2, 3),
1619 display_sync,
1620 ) catch |err| return mapSurfaceError(err);
1621 errdefer surface.deinit();
1622 var handle = backend.SurfaceHandle{
1623 .id = 0,
1624 .backend = .metal,
1625 .platform = .cocoa,
1626 .extent = request.extent,
1627 .format = request.format,
1628 .color_space = request.color_space,
1629 .present_mode = request.present_mode,
1630 };
1631 handle.id = try state.putObject(.{ .surface = .{ .surface = surface, .handle = handle } });
1632 state.objects.getPtr(handle.id).?.surface.handle.id = handle.id;
1633 return handle;
1634 }
1635
1636 fn destroySurface(ptr: *anyopaque, handle: backend.SurfaceHandle) backend.BackendError!void {
1637 const state: *State = @ptrCast(@alignCast(ptr));
1638 const surface = try state.getSurface(handle);
1639 if (surface.handle.generation != handle.generation) return error.SurfaceFrameExpired;
1640 if (surface.acquired_frame != null) return error.SurfaceAlreadyAcquired;
1641 const entry = state.objects.fetchRemove(handle.id).?;
1642 var object = entry.value;
1643 deinitObject(state, &object);
1644 }
1645
1646 fn acquireSurfaceFrame(ptr: *anyopaque, request: backend.SurfaceFrameAcquireRequest) backend.BackendError!backend.SurfaceFrame {
1647 const state: *State = @ptrCast(@alignCast(ptr));
1648 try state.objects.ensureUnusedCapacity(state.allocator, 2);
1649 const surface = try state.getSurface(request.surface);
1650 if (surface.handle.generation != request.surface.generation) return error.SurfaceFrameExpired;
1651 if (surface.acquired_frame != null) return error.SurfaceAlreadyAcquired;
1652 const texture_id = state.next_id;
1653 if (texture_id >= std.math.maxInt(BackendObjectId) - 1) return error.OutOfMemory;
1654 const frame_id = texture_id + 1;
1655
1656 var native = surface.surface.acquire() catch |err| return mapSurfaceError(err);
1657 errdefer native.release();
1658 state.next_id = frame_id + 1;
1659 const texture = backend.TextureHandle{
1660 .id = texture_id,
1661 .backend = .metal,
1662 .extent = .{ .width = surface.handle.extent.width, .height = surface.handle.extent.height, .depth = 1 },
1663 .format = surface.handle.format,
1664 .usage = .{ .copy_src = true, .copy_dst = true, .color_attachment = true, .present = true },
1665 .sample_count = 1,
1666 .ownership = .acquired_surface,
1667 };
1668 surface.acquired_frame = frame_id;
1669 surface.acquired_texture = texture_id;
1670 state.objects.putAssumeCapacityNoClobber(texture_id, .{ .texture = .{
1671 .handle = texture,
1672 .native = .{ .texture = native.texture, .size = .{ .width = texture.extent.width, .height = texture.extent.height, .depth = 1 }, .texel_bytes = 4 },
1673 .lent = true,
1674 } });
1675 state.objects.putAssumeCapacityNoClobber(frame_id, .{ .surface_frame = .{
1676 .surface_id = request.surface.id,
1677 .texture_id = texture_id,
1678 .frame = native,
1679 .generation = surface.handle.generation,
1680 } });
1681 return .{
1682 .id = frame_id,
1683 .backend = .metal,
1684 .surface = surface.handle,
1685 .texture = texture,
1686 .view = .{ .texture = texture, .format = texture.format },
1687 .generation = surface.handle.generation,
1688 };
1689 }
1690
1691 /// Checks that `request_frame` is the frame `request_surface` lent and has not presented.
1692 fn liveFrame(
1693 state: *State,
1694 request_surface: backend.SurfaceHandle,
1695 request_frame: backend.SurfaceFrame,
1696 ) backend.BackendError!struct { *MetalSurface, *MetalFrame } {
1697 const surface = try state.getSurface(request_surface);
1698 const frame = try state.getFrame(request_frame);
1699 if (surface.handle.generation != request_surface.generation) return error.SurfaceFrameExpired;
1700 if (frame.generation != request_frame.generation) return error.SurfaceFrameExpired;
1701 if (frame.presented) return error.SurfaceFrameExpired;
1702 if (frame.surface_id != request_surface.id) return error.InvalidSurfaceFrame;
1703 if (frame.texture_id != request_frame.texture.id) return error.InvalidSurfaceFrame;
1704 if (surface.acquired_frame == null or surface.acquired_frame.? != request_frame.id) return error.InvalidSurfaceFrame;
1705 return .{ surface, frame };
1706 }
1707
1708 /// Presents the frame after the work already queued on the default stream, and copies what it
1709 /// presents into the surface's readback buffer in the same command buffer.
1710 /// Presents a frame. Presenting consumes the drawable whether or not it succeeds, so the frame and
1711 /// its texture leave the surface before the native call, and a failed present releases nothing twice.
1712 fn presentSurfaceFrame(ptr: *anyopaque, request: backend.PresentRequest) backend.BackendError!void {
1713 const state: *State = @ptrCast(@alignCast(ptr));
1714 const rt = state.runtime.ptr() orelse return error.RuntimeUnavailable;
1715 if (request.signal_event != null) return error.UnsupportedOperation;
1716 const surface, const frame = try liveFrame(state, request.surface, request.frame);
1717 var wait_storage: [max_wait_events]*runtime_mod.Event = undefined;
1718 const waits = try state.nativeWaits(request.wait_events, &wait_storage);
1719 frame.presented = true;
1720 surface.acquired_frame = null;
1721 surface.acquired_texture = null;
1722 _ = state.objects.remove(frame.texture_id);
1723 surface.surface.present(rt, &frame.frame, waits) catch |err| return mapSurfaceError(err);
1724 }
1725
1726 fn writeSurfaceFrame(ptr: *anyopaque, request: backend.SurfaceFrameWriteRequest) backend.BackendError!void {
1727 if (comptime !native_available) return error.RuntimeUnavailable;
1728 const state: *State = @ptrCast(@alignCast(ptr));
1729 const rt = state.runtime.ptr() orelse return error.RuntimeUnavailable;
1730 _, const frame = try liveFrame(state, request.surface, request.frame);
1731 if (frame.written) return error.SurfaceFrameExpired;
1732 var wait_storage: [max_wait_events]*runtime_mod.Event = undefined;
1733 const waits = try state.nativeWaits(request.wait_events, &wait_storage);
1734 const signal = try state.nativeSignal(request.signal_event);
1735 const extent = request.frame.texture.extent;
1736 const size = metal.Size{ .width = extent.width, .height = extent.height, .depth = 1 };
1737
1738 var pool = foundation.AutoreleasePool.init() orelse return error.RuntimeUnavailable;
1739 defer pool.deinit();
1740 var submission = raster_mod.Submission.begin(rt, null, waits) catch |err| return mapRenderError(err);
1741 for (request.operations) |op| switch (op) {
1742 .clear => |color| {
1743 const encoder = try beginPass(submission.command_buffer, .{
1744 .color = frame.frame.texture,
1745 .color_load = .clear,
1746 .clear = .{ .red = color.r, .green = color.g, .blue = color.b, .alpha = color.a },
1747 .depth = null,
1748 .depth_load = .dont_care,
1749 .clear_depth = 1,
1750 .viewport = .{ .origin_x = 0, .origin_y = 0, .width = @floatFromInt(extent.width), .height = @floatFromInt(extent.height), .znear = 0, .zfar = 1 },
1751 .scissor = .{ .x = 0, .y = 0, .width = extent.width, .height = extent.height },
1752 });
1753 metal.endRenderEncoding(encoder);
1754 },
1755 .copy_buffer => |buffer_handle| {
1756 const buffer = try state.getBuffer(buffer_handle);
1757 if (buffer.byte_size < @as(usize, extent.width) * extent.height * 4) return error.InvalidBuffer;
1758 const blit = metal.blitCommandEncoder(submission.command_buffer) orelse return error.RenderFailed;
1759 metal.copyBufferToTexture(blit, buffer.buffer, @as(u64, extent.width) * 4, frame.frame.texture, size);
1760 metal.endBlitEncoding(blit);
1761 },
1762 };
1763 submission.finish(if (signal) |event| event.native.? else null) catch |err| return mapRenderError(err);
1764 if (signal) |event| event.recorded = true;
1765 frame.written = true;
1766 }
1767
1768 fn metalPrimitive(topology: backend.RenderPrimitiveTopology) metal.PrimitiveType {
1769 return switch (topology) {
1770 .triangle_list => .triangle,
1771 .triangle_strip => .triangle_strip,
1772 .line_list => .line,
1773 .line_strip => .line_strip,
1774 };
1775 }
1776
1777 fn metalCompare(compare: backend.RenderCompare) metal.CompareFunction {
1778 return switch (compare) {
1779 .never => .never,
1780 .less => .less,
1781 .equal => .equal,
1782 .less_equal => .less_equal,
1783 .greater => .greater,
1784 .not_equal => .not_equal,
1785 .greater_equal => .greater_equal,
1786 .always => .always,
1787 };
1788 }
1789
1790 /// The factors of each blend mode's equations, as `RenderBlendMode` states them.
1791 fn metalBlend(mode: backend.RenderBlendMode) ?metal.Blend {
1792 return switch (mode) {
1793 .replace => null,
1794 .alpha_premultiplied => .{
1795 .source_rgb = .one,
1796 .destination_rgb = .one_minus_source_alpha,
1797 .source_alpha = .one,
1798 .destination_alpha = .one_minus_source_alpha,
1799 },
1800 .alpha_straight => .{
1801 .source_rgb = .source_alpha,
1802 .destination_rgb = .one_minus_source_alpha,
1803 .source_alpha = .one,
1804 .destination_alpha = .one_minus_source_alpha,
1805 },
1806 .additive => .{
1807 .source_rgb = .one,
1808 .destination_rgb = .one,
1809 .source_alpha = .one,
1810 .destination_alpha = .one,
1811 },
1812 };
1813 }
1814
1815 fn vertexFormat(format: backend.RenderVertexFormat) metal.VertexFormat {
1816 return switch (format) {
1817 .float32 => .float,
1818 .float32x2 => .float2,
1819 .float32x3 => .float3,
1820 .float32x4 => .float4,
1821 .uint32 => .uint,
1822 .uint32x2 => .uint2,
1823 .uint32x4 => .uint4,
1824 };
1825 }
1826
1827 /// Maps a runtime failure met while recording or submitting a pass.
1828 fn mapRenderError(err: runtime_mod.Error) backend.BackendError {
1829 return switch (err) {
1830 error.LaunchFailed => error.RenderFailed,
1831 else => mapRuntimeError(err),
1832 };
1833 }
1834
1835 fn mapSurfaceError(err: runtime_mod.Error) backend.BackendError {
1836 return switch (err) {
1837 error.SurfaceUnavailable => error.InvalidSurface,
1838 error.FrameUnavailable => error.SurfaceFrameExpired,
1839 error.LaunchFailed => error.RenderFailed,
1840 else => mapRuntimeError(err),
1841 };
1842 }
1843
1844 fn mapRuntimeError(err: runtime_mod.Error) backend.BackendError {
1845 return switch (err) {
1846 error.UnsupportedPlatform,
1847 error.RuntimeUnavailable,
1848 error.DeviceUnavailable,
1849 => error.RuntimeUnavailable,
1850 error.CompilationFailed => error.CompilationFailed,
1851 error.LaunchFailed,
1852 => error.LaunchFailed,
1853 error.ResultMismatch => error.ResultMismatch,
1854 error.DeviceLost => error.DeviceLost,
1855 error.InvalidArtifact => error.InvalidArtifact,
1856 error.InvalidBuffer => error.InvalidBuffer,
1857 error.InvalidStream => error.InvalidStream,
1858 error.InvalidEvent => error.InvalidEvent,
1859 error.OutOfMemory => error.OutOfMemory,
1860 error.SurfaceUnavailable => error.InvalidSurface,
1861 error.FrameUnavailable => error.SurfaceFrameExpired,
1862 };
1863 }
1864
1865 const vtable = backend.BackendVTable{
1866 .query_capabilities = queryCapabilities,
1867 .create_artifact = createArtifact,
1868 .load_artifact = loadArtifact,
1869 .allocate_buffer = allocateBuffer,
1870 .create_stream = createStream,
1871 .create_event = createEvent,
1872 .write_buffer = writeBuffer,
1873 .read_buffer = readBuffer,
1874 .launch = launch,
1875 .create_render_artifact = createRenderArtifact,
1876 .load_render_artifact = loadRenderArtifact,
1877 .allocate_texture = allocateTexture,
1878 .destroy_texture = destroyTexture,
1879 .write_texture = writeTexture,
1880 .read_texture = readTexture,
1881 .create_render_bindings = createRenderBindings,
1882 .render = render,
1883 .record_render_bundle = recordRenderBundle,
1884 .submit_render_bundle = submitRenderBundle,
1885 .create_surface = createSurface,
1886 .destroy_surface = destroySurface,
1887 .acquire_surface_frame = acquireSurfaceFrame,
1888 .present_surface_frame = presentSurfaceFrame,
1889 .write_surface_frame = writeSurfaceFrame,
1890 .synchronize = synchronize,
1891 .query_event = queryEvent,
1892 .record_event = recordEvent,
1893 .destroy_object = destroyObject,
1894 .deinit = deinitHandle,
1895 };
1896
1897 fn fakeRuntime(handle: *anyopaque) Runtime {
1898 return .{
1899 .allocator = std.testing.allocator,
1900 .device = @ptrCast(handle),
1901 };
1902 }
1903
1904 test "metal reports its render, texture and surface capabilities with a runtime" {
1905 var fake: u8 = 0;
1906 var runtime = fakeRuntime(@ptrCast(&fake));
1907 var state = State.initWithRuntime(std.testing.allocator, &runtime);
1908 defer state.deinit();
1909 const caps = try state.handle().queryCapabilities();
1910 try std.testing.expect(caps.raster.supported);
1911 try std.testing.expect(caps.raster.supportsArtifactFormat(.metal_msl));
1912 try std.testing.expect(caps.raster.supportsArtifactFormat(.metal_metallib));
1913 try std.testing.expect(!caps.raster.supportsArtifactFormat(.vulkan_spirv));
1914 try std.testing.expectEqual(@as(u32, 128), caps.raster.max_push_constant_bytes);
1915 try std.testing.expect(caps.raster.depth_bias and caps.raster.depth_bias_clamp);
1916 try std.testing.expect(caps.textures.supportsFormat(.depth32_float));
1917 try std.testing.expect(caps.surfaces.supportsPlatform(.{ .cocoa = .{ .layer = 1 } }));
1918 try std.testing.expect(caps.surfaces.supportsFormat(.bgra8_unorm));
1919 try std.testing.expect(!caps.surfaces.supportsFormat(.rgba8_unorm));
1920 try std.testing.expect(!caps.surfaces.supportsPresentMode(.mailbox));
1921
1922 var bare = State.init(std.testing.allocator);
1923 defer bare.deinit();
1924 const bare_caps = try bare.handle().queryCapabilities();
1925 try std.testing.expect(!bare_caps.raster.supported);
1926 try std.testing.expect(!bare_caps.surfaces.supported);
1927 }
1928
1929 test "metal render pipelines refuse layouts and bindings past Metal's index plan before compiling" {
1930 var fake: u8 = 0;
1931 var runtime = fakeRuntime(@ptrCast(&fake));
1932 var state = State.initWithRuntime(std.testing.allocator, &runtime);
1933 defer state.deinit();
1934 const handle = state.handle();
1935 const attributes = [_]backend.RenderVertexAttribute{.{ .location = 0, .format = .float32x4, .offset = 0 }};
1936 const cases = [_]struct { desc: backend.RenderPipelineDesc, expected: backend.BackendError }{
1937 .{ .desc = .{
1938 .format = .metal_msl,
1939 .vertex_entry_name = "v",
1940 .fragment_entry_name = "f",
1941 .target_format = .rgba8_unorm,
1942 .vertex_layouts = &.{.{ .binding = max_vertex_buffers, .stride = 16, .attribute_start = 0, .attribute_count = 1 }},
1943 .vertex_attributes = &attributes,
1944 .push_extent = 0,
1945 .payload = .{ .text = "// msl" },
1946 }, .expected = error.CapabilityMismatch },
1947 .{ .desc = .{
1948 .format = .metal_msl,
1949 .vertex_entry_name = "v",
1950 .fragment_entry_name = "f",
1951 .target_format = .rgba8_unorm,
1952 .bindings = &.{.{ .binding = choir_abi.metal.max_uniform_bindings, .kind = .uniform_buffer }},
1953 .push_extent = 0,
1954 .payload = .{ .text = "// msl" },
1955 }, .expected = error.CapabilityMismatch },
1956 .{ .desc = .{
1957 .format = .metal_msl,
1958 .vertex_entry_name = "v",
1959 .fragment_entry_name = "f",
1960 .target_format = .rgba8_unorm,
1961 .bindings = &.{ .{ .binding = 1, .kind = .uniform_buffer }, .{ .binding = 1, .kind = .sampled_texture } },
1962 .push_extent = 0,
1963 .payload = .{ .text = "// msl" },
1964 }, .expected = error.InvalidRenderArtifact },
1965 };
1966 for (cases) |case| {
1967 var artifact = try handle.createRenderArtifact(case.desc);
1968 defer artifact.deinit();
1969 try std.testing.expectError(case.expected, handle.loadRenderArtifact(&artifact));
1970 }
1971 try std.testing.expectError(error.CapabilityMismatch, handle.createRenderArtifact(.{
1972 .format = .vulkan_spirv,
1973 .vertex_entry_name = "v",
1974 .fragment_entry_name = "f",
1975 .target_format = .rgba8_unorm,
1976 .push_extent = 0,
1977 .payload = .{ .words_u32 = &.{0x07230203} },
1978 }));
1979 }
1980
1981 test "metal blend modes follow the contract's equations" {
1982 try std.testing.expect(metalBlend(.replace) == null);
1983 const premultiplied = metalBlend(.alpha_premultiplied).?;
1984 try std.testing.expectEqual(metal.BlendFactor.one, premultiplied.source_rgb);
1985 try std.testing.expectEqual(metal.BlendFactor.one_minus_source_alpha, premultiplied.destination_rgb);
1986 const straight = metalBlend(.alpha_straight).?;
1987 try std.testing.expectEqual(metal.BlendFactor.source_alpha, straight.source_rgb);
1988 try std.testing.expectEqual(metal.BlendFactor.one, straight.source_alpha);
1989 const additive = metalBlend(.additive).?;
1990 try std.testing.expectEqual(metal.BlendFactor.one, additive.destination_alpha);
1991 }
1992
1993 /// Fills the top half of the target, in the contract's clip space where y = -1 is the top edge,
1994 /// with the color its push constants carry.
1995 const live_top_half_msl =
1996 \\#include <metal_stdlib>
1997 \\using namespace metal;
1998 \\
1999 \\struct live_out { float4 position [[position]]; };
2000 \\
2001 \\vertex live_out live_vertex(uint id [[vertex_id]]) {
2002 \\ const float2 corners[6] = {
2003 \\ float2(-1, -1), float2(1, -1), float2(-1, 0),
2004 \\ float2(1, -1), float2(1, 0), float2(-1, 0),
2005 \\ };
2006 \\ live_out out = {};
2007 \\ out.position = float4(corners[id], 0.5, 1);
2008 \\ out.position.y = -out.position.y;
2009 \\ return out;
2010 \\}
2011 \\
2012 \\fragment float4 live_fragment(live_out in [[stage_in]], constant uint* push [[buffer(30)]]) {
2013 \\ return float4(as_type<float>(push[0]), as_type<float>(push[1]), as_type<float>(push[2]), 1);
2014 \\}
2015 ;
2016
2017 test "metal live pass and bundle draw the top half with push constants and read it back" {
2018 var rt = try initRuntimeOrSkip(std.testing.allocator);
2019 defer rt.deinit();
2020 var state = State.initWithRuntime(std.testing.allocator, &rt);
2021 defer state.deinit();
2022 const handle = state.handle();
2023
2024 const side = 8;
2025 const color = try handle.allocateTexture(.{
2026 .extent = .{ .width = side, .height = side },
2027 .format = .rgba8_unorm,
2028 .usage = .{ .color_attachment = true, .copy_src = true },
2029 });
2030 defer handle.destroyTexture(color) catch {};
2031 var artifact = try handle.createRenderArtifact(.{
2032 .format = .metal_msl,
2033 .vertex_entry_name = "live_vertex",
2034 .fragment_entry_name = "live_fragment",
2035 .target_format = .rgba8_unorm,
2036 .push_constant_bytes = 16,
2037 .push_extent = 12,
2038 .payload = .{ .text = live_top_half_msl },
2039 });
2040 defer artifact.deinit();
2041 const pipeline = handle.loadRenderArtifact(&artifact) catch |err| {
2042 pretty.diagnostic.writeStderrText("metal pipeline: {s}\n", .{state.compileLog()});
2043 return err;
2044 };
2045 defer handle.destroyObject(pipeline.id);
2046 const push = [4]f32{ 1, 0, 0, 0 };
2047 const pass = backend.RenderPass{
2048 .color = .{ .view = .{ .texture = color, .format = .rgba8_unorm }, .load = .{ .clear = .{ .r = 0, .g = 0, .b = 1, .a = 1 } } },
2049 .viewport = .{ .width = side, .height = side },
2050 .scissor = .{ .width = side, .height = side },
2051 .draws = &.{.{ .pipeline = pipeline, .range = .{ .vertex_count = 6 }, .push_constants = std.mem.asBytes(&push) }},
2052 };
2053 try handle.render(.{ .pass = pass });
2054 try handle.synchronize(.{ .scope = .default_stream });
2055 var texels: [side * side * 4]u8 = undefined;
2056 try handle.readTexture(.{ .texture = color, .bytes = &texels });
2057 try expectTopHalf(&texels, side);
2058
2059 const bundle = try handle.recordRenderBundle(pass);
2060 defer handle.destroyObject(bundle.id);
2061 @memset(&texels, 0);
2062 try handle.submitRenderBundle(.{ .bundle = bundle });
2063 try handle.submitRenderBundle(.{ .bundle = bundle });
2064 try handle.synchronize(.{ .scope = .default_stream });
2065 try handle.readTexture(.{ .texture = color, .bytes = &texels });
2066 try expectTopHalf(&texels, side);
2067 }
2068
2069 /// Two triangles in contract clip space, y = -1 at the top, flipped for Metal as Choir's emitter
2070 /// flips them. The first covers the top-left half, counter-clockwise on screen and so front-facing
2071 /// by the contract; the second covers the bottom-right half, clockwise and so back-facing.
2072 const live_facing_msl =
2073 \\#include <metal_stdlib>
2074 \\using namespace metal;
2075 \\struct live_out { float4 position [[position]]; };
2076 \\vertex live_out facing_vertex(uint id [[vertex_id]]) {
2077 \\ const float2 corners[6] = {
2078 \\ float2(-1, -1), float2(-1, 1), float2(1, -1),
2079 \\ float2(-1, 1), float2(1, -1), float2(1, 1),
2080 \\ };
2081 \\ live_out out;
2082 \\ out.position = float4(corners[id].x, -corners[id].y, 0, 1);
2083 \\ return out;
2084 \\}
2085 \\fragment float4 facing_fragment(bool front [[front_facing]]) {
2086 \\ return front ? float4(1, 0, 0, 1) : float4(0, 1, 0, 1);
2087 \\}
2088 ;
2089
2090 test "metal live front faces follow the contract's counter-clockwise winding" {
2091 var rt = try initRuntimeOrSkip(std.testing.allocator);
2092 defer rt.deinit();
2093 var state = State.initWithRuntime(std.testing.allocator, &rt);
2094 defer state.deinit();
2095 const handle = state.handle();
2096
2097 const side = 8;
2098 const color = try handle.allocateTexture(.{
2099 .extent = .{ .width = side, .height = side },
2100 .format = .rgba8_unorm,
2101 .usage = .{ .color_attachment = true, .copy_src = true },
2102 });
2103 defer handle.destroyTexture(color) catch {};
2104 var artifact = try handle.createRenderArtifact(.{
2105 .format = .metal_msl,
2106 .vertex_entry_name = "facing_vertex",
2107 .fragment_entry_name = "facing_fragment",
2108 .target_format = .rgba8_unorm,
2109 .push_extent = 0,
2110 .payload = .{ .text = live_facing_msl },
2111 });
2112 defer artifact.deinit();
2113 const pipeline = handle.loadRenderArtifact(&artifact) catch |err| {
2114 pretty.diagnostic.writeStderrText("metal pipeline: {s}\n", .{state.compileLog()});
2115 return err;
2116 };
2117 defer handle.destroyObject(pipeline.id);
2118 const pass = backend.RenderPass{
2119 .color = .{ .view = .{ .texture = color, .format = .rgba8_unorm }, .load = .{ .clear = .{} } },
2120 .viewport = .{ .width = side, .height = side },
2121 .scissor = .{ .width = side, .height = side },
2122 .draws = &.{.{ .pipeline = pipeline, .range = .{ .vertex_count = 6 } }},
2123 };
2124 try handle.render(.{ .pass = pass });
2125 try handle.synchronize(.{ .scope = .default_stream });
2126 var texels: [side * side * 4]u8 = undefined;
2127 try handle.readTexture(.{ .texture = color, .bytes = &texels });
2128 try std.testing.expectEqualSlices(u8, &.{ 255, 0, 0, 255 }, texels[0..4]);
2129 try std.testing.expectEqualSlices(u8, &.{ 0, 255, 0, 255 }, texels[(side * side - 1) * 4 ..][0..4]);
2130
2131 const bundle = try handle.recordRenderBundle(pass);
2132 defer handle.destroyObject(bundle.id);
2133 @memset(&texels, 0);
2134 try handle.submitRenderBundle(.{ .bundle = bundle });
2135 try handle.synchronize(.{ .scope = .default_stream });
2136 try handle.readTexture(.{ .texture = color, .bytes = &texels });
2137 try std.testing.expectEqualSlices(u8, &.{ 255, 0, 0, 255 }, texels[0..4]);
2138 try std.testing.expectEqualSlices(u8, &.{ 0, 255, 0, 255 }, texels[(side * side - 1) * 4 ..][0..4]);
2139 }
2140
2141 fn expectTopHalf(texels: []const u8, side: usize) !void {
2142 for (0..side) |y| {
2143 const expected: [4]u8 = if (y < side / 2) .{ 255, 0, 0, 255 } else .{ 0, 0, 255, 255 };
2144 for (0..side) |x| {
2145 try std.testing.expectEqualSlices(u8, &expected, texels[(y * side + x) * 4 ..][0..4]);
2146 }
2147 }
2148 }
2149
2150 test "metal contract reports portable capability shape" {
2151 var state = State.init(std.testing.allocator);
2152 defer state.deinit();
2153 const handle = state.handle();
2154
2155 const caps = try handle.queryCapabilities();
2156 try std.testing.expectEqual(backend.BackendKind.metal, caps.identity.backend);
2157 try std.testing.expectEqual(backend.DeviceFamily.apple_metal, caps.identity.family);
2158 try std.testing.expect(caps.supportsDType(.i1));
2159 try std.testing.expect(caps.supportsDType(.f32));
2160 try std.testing.expect(caps.supportsDType(.f16));
2161 try std.testing.expect(caps.supportsArtifactFormat(.metal_msl));
2162 try std.testing.expect(caps.supportsArtifactFormat(.metal_metallib));
2163 try std.testing.expect(!caps.runtime.driver_loaded);
2164 }
2165
2166 test "metal contract reports runtime-backed capability shape" {
2167 var fake: u8 = 0;
2168 var runtime = fakeRuntime(@ptrCast(&fake));
2169 var state = State.initWithRuntime(std.testing.allocator, &runtime);
2170 defer state.deinit();
2171 const handle = state.handle();
2172
2173 const caps = try handle.queryCapabilities();
2174 try std.testing.expect(caps.runtime.driver_loaded);
2175 try std.testing.expect(caps.runtime.device_context);
2176 try std.testing.expect(caps.runtime.streams);
2177 try std.testing.expect(caps.runtime.events);
2178 try std.testing.expect(caps.supportsArtifactFormat(.metal_msl));
2179 }
2180
2181 test "metal contract supports synchronous stream event tokens" {
2182 var fake: u8 = 0;
2183 var runtime = fakeRuntime(@ptrCast(&fake));
2184 var state = State.initWithRuntime(std.testing.allocator, &runtime);
2185 defer state.deinit();
2186 const handle = state.handle();
2187
2188 const stream = backend.StreamHandle{
2189 .id = try state.putObject(.{ .stream = .{} }),
2190 .backend = .metal,
2191 };
2192 const event = backend.EventHandle{
2193 .id = try state.putObject(.{ .event = .{} }),
2194 .backend = .metal,
2195 };
2196
2197 try std.testing.expect(!try handle.queryEvent(.{ .event = event }));
2198 try std.testing.expectError(error.InvalidEvent, handle.synchronize(.{
2199 .scope = .event,
2200 .event = event,
2201 }));
2202
2203 try handle.recordEvent(.{
2204 .stream = stream,
2205 .event = event,
2206 });
2207 try std.testing.expect(try handle.queryEvent(.{ .event = event }));
2208 try handle.synchronize(.{
2209 .scope = .stream,
2210 .stream = stream,
2211 });
2212 try handle.synchronize(.{
2213 .scope = .event,
2214 .event = event,
2215 });
2216
2217 handle.destroyObject(event.id);
2218 try std.testing.expectError(error.InvalidEvent, handle.queryEvent(.{ .event = event }));
2219 handle.destroyObject(stream.id);
2220 try std.testing.expectError(error.InvalidStream, handle.synchronize(.{
2221 .scope = .stream,
2222 .stream = stream,
2223 }));
2224 }
2225
2226 test "metal contract rejects unrecorded wait events before native launch" {
2227 var fake: u8 = 0;
2228 var runtime = fakeRuntime(@ptrCast(&fake));
2229 var state = State.initWithRuntime(std.testing.allocator, &runtime);
2230 defer state.deinit();
2231 const handle = state.handle();
2232
2233 var artifact = try backend.KernelArtifact.init(std.testing.allocator, .{
2234 .backend = .metal,
2235 .format = .metal_msl,
2236 .entry_name = "main0",
2237 .argument_count = 0,
2238 });
2239 defer artifact.deinit();
2240 artifact.setBorrowedText("kernel void main0() {}");
2241
2242 const loaded_id = try state.putObject(.{ .loaded_artifact = .{
2243 .kernel = .{
2244 .pipeline = @ptrCast(&fake),
2245 .argument_count = 0,
2246 },
2247 .format = .metal_msl,
2248 } });
2249 defer _ = state.objects.fetchRemove(loaded_id);
2250
2251 const event = backend.EventHandle{
2252 .id = try state.putObject(.{ .event = .{} }),
2253 .backend = .metal,
2254 };
2255 defer handle.destroyObject(event.id);
2256
2257 try std.testing.expectError(error.InvalidEvent, handle.launch(.{
2258 .artifact = &artifact,
2259 .loaded_artifact = .{
2260 .id = loaded_id,
2261 .backend = .metal,
2262 .format = .metal_msl,
2263 },
2264 .buffers = &.{},
2265 .wait_events = &.{event},
2266 .geometry = .{ .grid = .{ 1, 1, 1 }, .threadgroup = .{ 1, 1, 1 } },
2267 }));
2268 }
2269
2270 test "metal contract validates artifacts before runtime access" {
2271 var state = State.init(std.testing.allocator);
2272 defer state.deinit();
2273 const handle = state.handle();
2274
2275 var msl = try backend.KernelArtifact.init(std.testing.allocator, .{
2276 .backend = .metal,
2277 .format = .metal_msl,
2278 .entry_name = "main0",
2279 .argument_count = 1,
2280 });
2281 defer msl.deinit();
2282 msl.setBorrowedText("kernel void main0() {}");
2283 try std.testing.expectError(error.RuntimeUnavailable, handle.loadArtifact(&msl));
2284
2285 var metallib = try backend.KernelArtifact.init(std.testing.allocator, .{
2286 .backend = .metal,
2287 .format = .metal_metallib,
2288 .entry_name = "main0",
2289 .argument_count = 1,
2290 });
2291 defer metallib.deinit();
2292 metallib.setBorrowedBytes(&.{ 0xca, 0xfe, 0xba, 0xbe });
2293 try std.testing.expectError(error.RuntimeUnavailable, handle.loadArtifact(&metallib));
2294
2295 var bad = try backend.KernelArtifact.init(std.testing.allocator, .{
2296 .backend = .metal,
2297 .format = .cuda_ptx,
2298 .entry_name = "main0",
2299 .argument_count = 1,
2300 });
2301 defer bad.deinit();
2302 bad.setBorrowedText("// ptx");
2303 try std.testing.expectError(error.UnsupportedArtifactFormat, handle.loadArtifact(&bad));
2304 }
2305
2306 test "metal contract rejects loaded artifact format forgery" {
2307 var fake: u8 = 0;
2308 var runtime = fakeRuntime(@ptrCast(&fake));
2309 var state = State.initWithRuntime(std.testing.allocator, &runtime);
2310 defer state.deinit();
2311 const handle = state.handle();
2312
2313 var artifact = try backend.KernelArtifact.init(std.testing.allocator, .{
2314 .backend = .metal,
2315 .format = .metal_metallib,
2316 .entry_name = "main0",
2317 .argument_count = 0,
2318 });
2319 defer artifact.deinit();
2320 artifact.setBorrowedBytes(&.{ 0xca, 0xfe, 0xba, 0xbe });
2321
2322 const loaded_id = try state.putObject(.{ .loaded_artifact = .{
2323 .kernel = .{
2324 .pipeline = @ptrCast(&fake),
2325 .argument_count = 0,
2326 },
2327 .format = .metal_msl,
2328 } });
2329 defer _ = state.objects.fetchRemove(loaded_id);
2330
2331 try std.testing.expectError(error.InvalidArtifact, handle.launch(.{
2332 .artifact = &artifact,
2333 .loaded_artifact = .{
2334 .id = loaded_id,
2335 .backend = .metal,
2336 .format = .metal_metallib,
2337 },
2338 .buffers = &.{},
2339 .geometry = .{ .grid = .{ 1, 1, 1 }, .threadgroup = .{ 1, 1, 1 } },
2340 }));
2341 }
2342
2343 test "metal contract creates MSL artifacts before runtime access" {
2344 var state = State.init(std.testing.allocator);
2345 defer state.deinit();
2346 const handle = state.handle();
2347
2348 var artifact = try handle.createArtifact(.{
2349 .kernel_name = "direct_add_f32",
2350 .requested_format = .metal_msl,
2351 .argument_count = 4,
2352 .required_dtypes = backend.DTypeSet.init(&.{.f32}),
2353 .payload = .{ .text = "kernel void direct_add_f32() {}" },
2354 });
2355 defer artifact.deinit();
2356
2357 try std.testing.expectEqual(backend.BackendKind.metal, artifact.backend);
2358 try std.testing.expectEqual(backend.ArtifactFormat.metal_msl, artifact.format);
2359 try std.testing.expectEqualStrings("direct_add_f32", artifact.entry_name);
2360 try std.testing.expectEqual(@as(u32, 4), artifact.argument_count);
2361 try std.testing.expect(std.mem.indexOf(u8, artifact.payload.text, "direct_add_f32") != null);
2362 }
2363
2364 test "metal contract validates launch argument count before runtime access" {
2365 var state = State.init(std.testing.allocator);
2366 defer state.deinit();
2367 const handle = state.handle();
2368
2369 var artifact = try backend.KernelArtifact.init(std.testing.allocator, .{
2370 .backend = .metal,
2371 .format = .metal_msl,
2372 .entry_name = "direct_add_f32",
2373 .argument_count = 4,
2374 });
2375 defer artifact.deinit();
2376 artifact.setBorrowedText("kernel void direct_add_f32() {}");
2377
2378 const out = backend.BufferHandle{ .id = 1, .backend = .metal, .byte_size = 1024, .ownership = .backend };
2379 const lhs = backend.BufferHandle{ .id = 2, .backend = .metal, .byte_size = 1024, .ownership = .backend };
2380 const rhs = backend.BufferHandle{ .id = 3, .backend = .metal, .byte_size = 1024, .ownership = .backend };
2381 const n = backend.BufferHandle{ .id = 4, .backend = .metal, .byte_size = @sizeOf(u32), .ownership = .backend };
2382
2383 const short_bindings = [_]backend.BufferBinding{
2384 .{ .handle = out, .access = .write_only, .ownership = .backend, .byte_size = 1024 },
2385 .{ .handle = lhs, .access = .read_only, .ownership = .backend, .byte_size = 1024 },
2386 .{ .handle = rhs, .access = .read_only, .ownership = .backend, .byte_size = 1024 },
2387 };
2388 try std.testing.expectError(error.LaunchArgumentMismatch, handle.launch(.{
2389 .artifact = &artifact,
2390 .loaded_artifact = .{ .id = 99, .backend = .metal, .format = .metal_msl },
2391 .buffers = &short_bindings,
2392 .geometry = .{ .grid = .{ 4, 1, 1 }, .threadgroup = .{ 64, 1, 1 } },
2393 }));
2394
2395 const bindings = [_]backend.BufferBinding{
2396 .{ .handle = out, .access = .write_only, .ownership = .backend, .byte_size = 1024 },
2397 .{ .handle = lhs, .access = .read_only, .ownership = .backend, .byte_size = 1024 },
2398 .{ .handle = rhs, .access = .read_only, .ownership = .backend, .byte_size = 1024 },
2399 };
2400 try std.testing.expectError(error.RuntimeUnavailable, handle.launch(.{
2401 .artifact = &artifact,
2402 .loaded_artifact = .{ .id = 99, .backend = .metal, .format = .metal_msl },
2403 .buffers = &bindings,
2404 .scalar_arguments = &.{.{ .u32 = 256 }},
2405 .geometry = .{ .grid = .{ 4, 1, 1 }, .threadgroup = .{ 64, 1, 1 } },
2406 }));
2407
2408 const too_many = [_]backend.BufferBinding{
2409 .{ .handle = out, .access = .write_only, .ownership = .backend, .byte_size = 1024 },
2410 .{ .handle = lhs, .access = .read_only, .ownership = .backend, .byte_size = 1024 },
2411 .{ .handle = rhs, .access = .read_only, .ownership = .backend, .byte_size = 1024 },
2412 .{ .handle = n, .access = .read_only, .ownership = .backend, .byte_size = @sizeOf(u32) },
2413 };
2414 try std.testing.expectError(error.LaunchArgumentMismatch, handle.launch(.{
2415 .artifact = &artifact,
2416 .loaded_artifact = .{ .id = 99, .backend = .metal, .format = .metal_msl },
2417 .buffers = &too_many,
2418 .scalar_arguments = &.{.{ .u32 = 256 }},
2419 .geometry = .{ .grid = .{ 4, 1, 1 }, .threadgroup = .{ 64, 1, 1 } },
2420 }));
2421 }
2422
2423 test "metal scalar argument lowering preserves tags and bit patterns" {
2424 try std.testing.expectEqual(runtime_mod.ScalarArgument{ .tag = .i32, .bits = @as(u32, @bitCast(@as(i32, -7))) }, lowerScalarArgument(.{ .i32 = -7 }));
2425 try std.testing.expectEqual(runtime_mod.ScalarArgument{ .tag = .u32, .bits = 7 }, lowerScalarArgument(.{ .u32 = 7 }));
2426 try std.testing.expectEqual(runtime_mod.ScalarArgument{ .tag = .i64, .bits = @bitCast(@as(i64, -9)) }, lowerScalarArgument(.{ .i64 = -9 }));
2427 try std.testing.expectEqual(runtime_mod.ScalarArgument{ .tag = .u64, .bits = 9 }, lowerScalarArgument(.{ .u64 = 9 }));
2428 try std.testing.expectEqual(runtime_mod.ScalarArgument{ .tag = .f32, .bits = @as(u32, @bitCast(@as(f32, -1.25))) }, lowerScalarArgument(.{ .f32 = -1.25 }));
2429 try std.testing.expectEqual(runtime_mod.ScalarArgument{ .tag = .f64, .bits = @bitCast(@as(f64, 2.5)) }, lowerScalarArgument(.{ .f64 = 2.5 }));
2430 }
2431
2432 test "metal live validation row launches add_f32 through backend handle" {
2433 var rt = try initRuntimeOrSkip(std.testing.allocator);
2434 defer rt.deinit();
2435
2436 var state = State.initWithRuntime(std.testing.allocator, &rt);
2437 defer state.deinit();
2438 const handle = state.handle();
2439
2440 const element_count: usize = 257;
2441 const threads_per_group: u32 = 64;
2442 const group_count: u32 = @intCast((element_count + threads_per_group - 1) / threads_per_group);
2443
2444 var lhs: [element_count]f32 = undefined;
2445 var rhs: [element_count]f32 = undefined;
2446 var out: [element_count]f32 = @as([element_count]f32, @splat(0));
2447 for (&lhs, &rhs, 0..) |*left, *right, i| {
2448 left.* = @as(f32, @floatFromInt(i)) * 0.25;
2449 right.* = @as(f32, @floatFromInt(i % 17)) * -0.5;
2450 }
2451
2452 var artifact = try backend.KernelArtifact.init(std.testing.allocator, .{
2453 .backend = .metal,
2454 .format = .metal_msl,
2455 .entry_name = "add_f32",
2456 .argument_count = 4,
2457 });
2458 defer artifact.deinit();
2459 artifact.setBorrowedText(
2460 \\#include <metal_stdlib>
2461 \\using namespace metal;
2462 \\
2463 \\kernel void add_f32(
2464 \\ device float* out [[buffer(0)]],
2465 \\ device float* lhs [[buffer(1)]],
2466 \\ device float* rhs [[buffer(2)]],
2467 \\ constant uint* n [[buffer(3)]],
2468 \\ uint id [[thread_position_in_grid]]
2469 \\) {
2470 \\ if (id < *n) {
2471 \\ out[id] = lhs[id] + rhs[id];
2472 \\ }
2473 \\}
2474 );
2475
2476 const loaded = try handle.loadArtifact(&artifact);
2477 const out_buffer = try handle.allocateBuffer(.{
2478 .byte_size = @sizeOf(@TypeOf(out)),
2479 .alignment = 256,
2480 .dtype = .f32,
2481 .element_count = element_count,
2482 });
2483 const lhs_buffer = try handle.allocateBuffer(.{
2484 .byte_size = @sizeOf(@TypeOf(lhs)),
2485 .alignment = 256,
2486 .dtype = .f32,
2487 .element_count = element_count,
2488 });
2489 const rhs_buffer = try handle.allocateBuffer(.{
2490 .byte_size = @sizeOf(@TypeOf(rhs)),
2491 .alignment = 256,
2492 .dtype = .f32,
2493 .element_count = element_count,
2494 });
2495
2496 try handle.writeBuffer(.{
2497 .handle = out_buffer,
2498 .bytes = std.mem.sliceAsBytes(out[0..]),
2499 });
2500 try handle.writeBuffer(.{
2501 .handle = lhs_buffer,
2502 .bytes = std.mem.sliceAsBytes(lhs[0..]),
2503 });
2504 try handle.writeBuffer(.{
2505 .handle = rhs_buffer,
2506 .bytes = std.mem.sliceAsBytes(rhs[0..]),
2507 });
2508
2509 const stream = try handle.createStream(.{});
2510 const event = try handle.createEvent(.{});
2511 const bindings = [_]backend.BufferBinding{
2512 .{ .handle = out_buffer, .access = .write_only, .ownership = out_buffer.ownership, .byte_size = out_buffer.byte_size },
2513 .{ .handle = lhs_buffer, .access = .read_only, .ownership = lhs_buffer.ownership, .byte_size = lhs_buffer.byte_size },
2514 .{ .handle = rhs_buffer, .access = .read_only, .ownership = rhs_buffer.ownership, .byte_size = rhs_buffer.byte_size },
2515 };
2516
2517 try handle.launch(.{
2518 .artifact = &artifact,
2519 .loaded_artifact = loaded,
2520 .buffers = &bindings,
2521 .scalar_arguments = &.{.{ .u32 = @intCast(element_count) }},
2522 .geometry = .{
2523 .grid = .{ group_count, 1, 1 },
2524 .threadgroup = .{ threads_per_group, 1, 1 },
2525 },
2526 .stream = stream,
2527 .signal_event = event,
2528 });
2529
2530 try handle.synchronize(.{ .scope = .event, .event = event });
2531 try std.testing.expect(try handle.queryEvent(.{ .event = event }));
2532 try handle.readBuffer(.{
2533 .handle = out_buffer,
2534 .bytes = std.mem.sliceAsBytes(out[0..]),
2535 });
2536
2537 for (lhs, rhs, out) |left, right, observed| {
2538 try std.testing.expectApproxEqAbs(left + right, observed, 1e-6);
2539 }
2540 }
2541
2542 test "metal live loads precompiled metallib artifact through backend handle" {
2543 var rt = try initRuntimeOrSkip(std.testing.allocator);
2544 defer rt.deinit();
2545
2546 const source =
2547 \\#include <metal_stdlib>
2548 \\using namespace metal;
2549 \\
2550 \\kernel void scale_f32(
2551 \\ device const float* input [[buffer(0)]],
2552 \\ device float* output [[buffer(1)]],
2553 \\ constant float& scale [[buffer(2)]],
2554 \\ uint id [[thread_position_in_grid]]
2555 \\) {
2556 \\ if (id < 64) {
2557 \\ output[id] = input[id] * scale;
2558 \\ }
2559 \\}
2560 ;
2561 const metallib = try compileMetallibOrSkip(std.testing.allocator, source);
2562 defer std.testing.allocator.free(metallib);
2563
2564 var state = State.initWithRuntime(std.testing.allocator, &rt);
2565 defer state.deinit();
2566 const handle = state.handle();
2567
2568 const element_count: usize = 64;
2569 var input: [element_count]f32 = undefined;
2570 var output: [element_count]f32 = @as([element_count]f32, @splat(0));
2571 for (&input, 0..) |*value, i| {
2572 value.* = @as(f32, @floatFromInt(i)) * 1.5 - 7.0;
2573 }
2574
2575 var artifact = try backend.KernelArtifact.init(std.testing.allocator, .{
2576 .backend = .metal,
2577 .format = .metal_metallib,
2578 .entry_name = "scale_f32",
2579 .argument_count = 3,
2580 });
2581 defer artifact.deinit();
2582 artifact.setBorrowedBytes(metallib);
2583
2584 const loaded = try handle.loadArtifact(&artifact);
2585 try std.testing.expectEqual(backend.ArtifactFormat.metal_metallib, loaded.format);
2586
2587 const input_buffer = try handle.allocateBuffer(.{
2588 .byte_size = @sizeOf(@TypeOf(input)),
2589 .alignment = 256,
2590 .dtype = .f32,
2591 .element_count = element_count,
2592 });
2593 const output_buffer = try handle.allocateBuffer(.{
2594 .byte_size = @sizeOf(@TypeOf(output)),
2595 .alignment = 256,
2596 .dtype = .f32,
2597 .element_count = element_count,
2598 });
2599 try handle.writeBuffer(.{
2600 .handle = input_buffer,
2601 .bytes = std.mem.sliceAsBytes(input[0..]),
2602 });
2603 try handle.writeBuffer(.{
2604 .handle = output_buffer,
2605 .bytes = std.mem.sliceAsBytes(output[0..]),
2606 });
2607
2608 const bindings = [_]backend.BufferBinding{
2609 .{ .handle = input_buffer, .access = .read_only, .ownership = input_buffer.ownership, .byte_size = input_buffer.byte_size },
2610 .{ .handle = output_buffer, .access = .write_only, .ownership = output_buffer.ownership, .byte_size = output_buffer.byte_size },
2611 };
2612 const scale: f32 = -2.25;
2613 try handle.launch(.{
2614 .artifact = &artifact,
2615 .loaded_artifact = loaded,
2616 .buffers = &bindings,
2617 .scalar_arguments = &.{.{ .f32 = scale }},
2618 .geometry = .{
2619 .grid = .{ 1, 1, 1 },
2620 .threadgroup = .{ element_count, 1, 1 },
2621 },
2622 });
2623 try handle.readBuffer(.{
2624 .handle = output_buffer,
2625 .bytes = std.mem.sliceAsBytes(output[0..]),
2626 });
2627
2628 for (input, output) |source_value, observed| {
2629 try std.testing.expectApproxEqAbs(source_value * scale, observed, 1e-6);
2630 }
2631 }
2632
2633 test "metal live stream event record synchronizes native completion" {
2634 var rt = try initRuntimeOrSkip(std.testing.allocator);
2635 defer rt.deinit();
2636
2637 var state = State.initWithRuntime(std.testing.allocator, &rt);
2638 defer state.deinit();
2639 const handle = state.handle();
2640
2641 const stream = try handle.createStream(.{});
2642 const event = try handle.createEvent(.{});
2643
2644 try std.testing.expect(!try handle.queryEvent(.{ .event = event }));
2645 try handle.recordEvent(.{
2646 .stream = stream,
2647 .event = event,
2648 });
2649 try handle.synchronize(.{
2650 .scope = .event,
2651 .event = event,
2652 });
2653 try std.testing.expect(try handle.queryEvent(.{ .event = event }));
2654 try handle.synchronize(.{
2655 .scope = .stream,
2656 .stream = stream,
2657 });
2658 }
2659
2660 fn compileMetallibOrSkip(allocator: Allocator, source: []const u8) ![]u8 {
2661 const dir = try metallibScratchDir(allocator);
2662 defer removeMetallibScratchDir(allocator, dir);
2663
2664 const source_path = try sys.path.join(allocator, &.{ dir, "kernel.metal" });
2665 defer allocator.free(source_path);
2666 const air_path = try sys.path.join(allocator, &.{ dir, "kernel.air" });
2667 defer allocator.free(air_path);
2668 const metallib_path = try sys.path.join(allocator, &.{ dir, "kernel.metallib" });
2669 defer allocator.free(metallib_path);
2670
2671 try sys.fs.writeFile(source_path, source);
2672
2673 const metal_result = try runMetalToolOrSkip(allocator, &.{ "xcrun", "-sdk", "macosx", "metal", "-c", source_path, "-o", air_path });
2674 defer allocator.free(metal_result.stdout);
2675 defer allocator.free(metal_result.stderr);
2676 if (sys.process.exitCode(metal_result.term) != 0) {
2677 pretty.diagnostic.writeStderrText(
2678 "xcrun metal failed\nstdout:\n{s}\nstderr:\n{s}\n",
2679 .{ metal_result.stdout, metal_result.stderr },
2680 );
2681 return error.MetalToolRejected;
2682 }
2683
2684 const metallib_result = try runMetalToolOrSkip(allocator, &.{ "xcrun", "-sdk", "macosx", "metallib", air_path, "-o", metallib_path });
2685 defer allocator.free(metallib_result.stdout);
2686 defer allocator.free(metallib_result.stderr);
2687 if (sys.process.exitCode(metallib_result.term) != 0) {
2688 pretty.diagnostic.writeStderrText(
2689 "xcrun metallib failed\nstdout:\n{s}\nstderr:\n{s}\n",
2690 .{ metallib_result.stdout, metallib_result.stderr },
2691 );
2692 return error.MetalToolRejected;
2693 }
2694
2695 return try sys.fs.readFileAlloc(allocator, metallib_path, 4 * 1024 * 1024);
2696 }
2697
2698 /// Runs an Xcode tool, skipping the test when the tool is absent. Xcode ships `xcrun metal` as a
2699 /// shim until its Metal Toolchain component is downloaded, which counts as absent.
2700 fn runMetalToolOrSkip(allocator: Allocator, argv: []const []const u8) !sys.process.RunResult {
2701 var io_state = std.Io.Threaded.init(allocator, .{ .environ = sys.env.current() });
2702 defer io_state.deinit();
2703 const result = sys.process.run(allocator, io_state.io(), .{
2704 .argv = argv,
2705 .expand_arg0 = .expand,
2706 .stdout_limit = .limited(64 * 1024),
2707 .stderr_limit = .limited(64 * 1024),
2708 }) catch |err| switch (err) {
2709 error.OutOfMemory => return error.OutOfMemory,
2710 else => return error.SkipZigTest,
2711 };
2712 if (std.mem.indexOf(u8, result.stderr, "missing Metal Toolchain") != null) {
2713 allocator.free(result.stdout);
2714 allocator.free(result.stderr);
2715 return error.SkipZigTest;
2716 }
2717 return result;
2718 }
2719
2720 fn metallibScratchDir(allocator: Allocator) ![]u8 {
2721 const base = sys.env.get("TMPDIR") orelse "/tmp";
2722 const pid = sys.process.currentProcessId() catch 0;
2723 const dir = try std.fmt.allocPrint(allocator, "{s}/gpu-metal-metallib-{d}", .{ base, pid });
2724 errdefer allocator.free(dir);
2725 try sys.fs.createDirPath(dir);
2726 return dir;
2727 }
2728
2729 fn removeMetallibScratchDir(allocator: Allocator, dir: []u8) void {
2730 sys.fs.deleteTree(dir) catch {};
2731 allocator.free(dir);
2732 }
2733
2734 fn initRuntimeOrSkip(allocator: Allocator) (runtime_mod.Error || error{SkipZigTest})!Runtime {
2735 if (!build_options.metal_tests) return error.SkipZigTest;
2736 switch (runtime_mod.detectAvailability()) {
2737 .available => {},
2738 .unsupported_platform,
2739 .runtime_missing,
2740 .device_missing,
2741 => return error.SkipZigTest,
2742 }
2743 return Runtime.init(allocator) catch |err| switch (err) {
2744 error.UnsupportedPlatform,
2745 error.RuntimeUnavailable,
2746 error.DeviceUnavailable,
2747 => return error.SkipZigTest,
2748 else => return err,
2749 };
2750 }