lib/choir/src/backends/gpu/metal/test.zig
daab053ee43316e1809a84551d573ddd1e5bf3d2
1 const std = @import("std");
2 const fixture = @import("../fixture/root.zig");
3 const metal = @import("root.zig");
4
5 const expectContains = fixture.expectContains;
6
7 test "metal declaration coverage" {
8 std.testing.refAllDecls(metal.msl);
9 std.testing.refAllDecls(metal);
10 }
11
12 fn emit(source: []const u8, entry_name: []const u8) ![]u8 {
13 return fixture.emitText(source, entry_name, null, metal.msl.emitMsl);
14 }
15
16 test "msl emitter lowers a parsed elementwise kernel" {
17 const text = try emit(@embedFile("fixture/elementwise.txt"), "test_elementwise");
18 defer std.testing.allocator.free(text);
19 try expectContains(text, "kernel void test_elementwise");
20 try expectContains(text, "device float* arg0 [[buffer(0)]]");
21 try expectContains(text, "thread_position_in_grid");
22 try expectContains(text, "if (");
23 try expectContains(text, "floor(");
24 try expectContains(text, "0x80000000u");
25 try expectContains(text, "trunc(");
26 try expectContains(text, "atan2(");
27 try expectContains(text, "arg0[");
28 }
29
30 test "msl emitter lowers a parsed while loop with carried exits" {
31 const text = try emit(@embedFile("fixture/loop.txt"), "test_loop");
32 defer std.testing.allocator.free(text);
33 try expectContains(text, "while (true) {");
34 try expectContains(text, "if (!(");
35 try expectContains(text, "break;");
36 try std.testing.expect(std.mem.count(u8, text, "int l") >= 4);
37 }
38
39 test "msl emitter lowers parsed shared memory barriers and atomics" {
40 const text = try emit(@embedFile("fixture/shared.txt"), "test_shared");
41 defer std.testing.allocator.free(text);
42 try expectContains(text, "threadgroup int shared");
43 try expectContains(text, "threadgroup_barrier(mem_flags::mem_threadgroup)");
44 try expectContains(text, "threadgroup atomic_int*");
45 }
46
47 test "msl emitter lowers parsed device compare and swap" {
48 const text = try emit(@embedFile("fixture/exchange.txt"), "test_exchange");
49 defer std.testing.allocator.free(text);
50 try expectContains(text, "atomic_compare_exchange_weak_explicit");
51 try expectContains(text, "device atomic_int*");
52 try expectContains(text, "do {");
53 try expectContains(text, "while (!");
54 }
55
56 test "msl emitter lowers parsed simdgroup ids scan and reduce" {
57 const text = try emit(@embedFile("fixture/simdgroup.txt"), "test_simdgroup");
58 defer std.testing.allocator.free(text);
59 try expectContains(text, "thread_index_in_simdgroup");
60 try expectContains(text, "simdgroup_index_in_threadgroup");
61 try expectContains(text, "simd_prefix_inclusive_sum");
62 try expectContains(text, "simd_sum");
63 }
64
65 test "msl emitter rejects a parsed shared f32 atomic add" {
66 try std.testing.expectError(
67 error.UnsupportedOperation,
68 emit(@embedFile("fixture/rejection.txt"), "test_rejection"),
69 );
70 }
71
72 const choir = @import("../../../root.zig");
73 const registration = @import("../registration.zig");
74 const stages = fixture.stages;
75
76 fn emitStages(source: []const u8, case: stages.Case) ![]u8 {
77 const allocator = std.testing.allocator;
78 var ctx = try choir.ir.Context.init(allocator, choir.ir.Context.Limits.testing);
79 defer ctx.deinit(allocator);
80 try registration.prepareCompilationDialects(&ctx);
81 const module = try fixture.lower(&ctx, source, null);
82 defer module.erase();
83 return stages.emitMsl(allocator, module, case);
84 }
85
86 /// `case`'s source with its one `needle` replaced. The caller owns it.
87 fn variant(case: stages.Case, needle: []const u8, replacement: []const u8) ![]u8 {
88 try std.testing.expectEqual(@as(usize, 1), std.mem.count(u8, case.source, needle));
89 return std.mem.replaceOwned(u8, std.testing.allocator, case.source, needle, replacement);
90 }
91
92 test "msl emitter writes each stage module as its golden vertex and fragment source" {
93 const goldens = [_][]const u8{
94 @embedFile("fixture/textured.metal"),
95 @embedFile("fixture/indexed.metal"),
96 @embedFile("fixture/pushed.metal"),
97 };
98 for (stages.cases, goldens) |case, golden| {
99 const text = try emitStages(case.source, case);
100 defer std.testing.allocator.free(text);
101 try std.testing.expectEqualStrings(golden, text);
102 }
103 }
104
105 test "msl stage emission keeps static local arrays in thread address space" {
106 const local = stages.alloca_case;
107 const text = try emitStages(local.source, local);
108 defer std.testing.allocator.free(text);
109 try expectContains(text, "thread float v");
110 try expectContains(text, "thread bool v");
111 try expectContains(text, "[4];");
112 try expectContains(text, "[1];");
113 try std.testing.expectEqual(@as(usize, 0), std.mem.count(u8, text, "threadgroup float"));
114 }
115
116 test "msl stage emission refuses dynamic local arrays" {
117 const local = stages.alloca_case;
118 const source = try variant(
119 local,
120 "%16 = memref.alloca() : !memref<4,arith.f32,local>",
121 "%16 = memref.alloca(%12) : !memref<?,arith.f32,local>",
122 );
123 defer std.testing.allocator.free(source);
124 try std.testing.expectError(error.UnsupportedStageMemory, emitStages(source, local));
125 }
126
127 test "msl stage calls define helpers once and refuse recursion" {
128 const calls = stages.calls_case;
129 const text = try emitStages(calls.source, calls);
130 defer std.testing.allocator.free(text);
131 try std.testing.expectEqual(@as(usize, 1), std.mem.count(u8, text, "float invert("));
132 try std.testing.expectEqual(@as(usize, 1), std.mem.count(u8, text, "float nested_invert("));
133 try expectContains(text, "= invert(");
134 try expectContains(text, "= nested_invert(");
135 try expectContains(text, "thread float v");
136 const recursive = try variant(calls, "%9 = func.call(%8) {callee = @invert}", "%9 = func.call(%8) {callee = @nested_invert}");
137 defer std.testing.allocator.free(recursive);
138 try std.testing.expectError(error.UnsupportedOperation, emitStages(recursive, calls));
139 }
140
141 test "msl stage emission rejects ops outside their stage and texture groups past zero" {
142 const indexed = stages.cases[1];
143 const textured = stages.cases[0];
144 const implicit_in_vertex = try variant(
145 indexed,
146 "gpu.sample_lod(%4, %2, %3, %5)",
147 "gpu.sample(%4, %2, %3)",
148 );
149 defer std.testing.allocator.free(implicit_in_vertex);
150 const implicit_result = emitStages(implicit_in_vertex, indexed);
151 try std.testing.expectError(error.UnsupportedOperation, implicit_result);
152 const second_group = try variant(
153 textured,
154 "{binding = 0:i64, group = 0:i64}",
155 "{binding = 0:i64, group = 1:i64}",
156 );
157 defer std.testing.allocator.free(second_group);
158 try std.testing.expectError(error.CapabilityMismatch, emitStages(second_group, textured));
159 const split_slot = try variant(
160 textured,
161 "%12, %13, %14, %15 = gpu.stage_input() {location = 1:i64}",
162 "%12, %13, %14, %15 = gpu.stage_input() {location = 0:i64}",
163 );
164 defer std.testing.allocator.free(split_slot);
165 try std.testing.expectError(error.InvalidArtifact, emitStages(split_slot, textured));
166 }
167
168 test "msl block reads refuse uniforms past group zero or binding seven" {
169 const pushed = stages.cases[2];
170 const needle = "{binding = 1:i64, group = 0:i64, offset = 12:i64}";
171 const refusals = [_][]const u8{
172 "{binding = 1:i64, group = 1:i64, offset = 12:i64}",
173 "{binding = 8:i64, group = 0:i64, offset = 12:i64}",
174 "{binding = 1:i64, group = 0:i64, offset = 16384:i64}",
175 };
176 for (refusals) |replacement| {
177 const source = try variant(pushed, needle, replacement);
178 defer std.testing.allocator.free(source);
179 try std.testing.expectError(error.CapabilityMismatch, emitStages(source, pushed));
180 }
181 }
182
183 test "msl kernel emission rejects stage ops" {
184 const textured = stages.cases[0];
185 const kernel = try variant(
186 textured,
187 "stage = #attr<gpu.stage>(\"vertex\")",
188 "kernel = #attr<func.kernel>",
189 );
190 defer std.testing.allocator.free(kernel);
191 try std.testing.expectError(error.UnsupportedOperation, emit(kernel, textured.vertex));
192 }