1 // Copyright 2020 The Tint Authors.
2 //
3 // Licensed under the Apache License, Version 2.0 (the "License");
4 // you may not use this file except in compliance with the License.
5 // You may obtain a copy of the License at
6 //
7 // http://www.apache.org/licenses/LICENSE-2.0
8 //
9 // Unless required by applicable law or agreed to in writing, software
10 // distributed under the License is distributed on an "AS IS" BASIS,
11 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12 // See the License for the specific language governing permissions and
13 // limitations under the License.
14
15 #include "src/ast/stage_decoration.h"
16 #include "src/ast/struct_block_decoration.h"
17 #include "src/ast/variable_decl_statement.h"
18 #include "src/writer/msl/test_helper.h"
19
20 namespace tint {
21 namespace writer {
22 namespace msl {
23 namespace {
24
25 using MslGeneratorImplTest = TestHelper;
26
TEST_F(MslGeneratorImplTest,Emit_Function)27 TEST_F(MslGeneratorImplTest, Emit_Function) {
28 Func("my_func", ast::VariableList{}, ty.void_(),
29 ast::StatementList{
30 Return(),
31 },
32 {});
33
34 GeneratorImpl& gen = Build();
35
36 gen.increment_indent();
37
38 ASSERT_TRUE(gen.Generate()) << gen.error();
39 EXPECT_EQ(gen.result(), R"( #include <metal_stdlib>
40
41 using namespace metal;
42 void my_func() {
43 return;
44 }
45
46 )");
47 }
48
TEST_F(MslGeneratorImplTest,Emit_Function_WithParams)49 TEST_F(MslGeneratorImplTest, Emit_Function_WithParams) {
50 ast::VariableList params;
51 params.push_back(Param("a", ty.f32()));
52 params.push_back(Param("b", ty.i32()));
53
54 Func("my_func", params, ty.void_(),
55 ast::StatementList{
56 Return(),
57 },
58 {});
59
60 GeneratorImpl& gen = Build();
61
62 gen.increment_indent();
63
64 ASSERT_TRUE(gen.Generate()) << gen.error();
65 EXPECT_EQ(gen.result(), R"( #include <metal_stdlib>
66
67 using namespace metal;
68 void my_func(float a, int b) {
69 return;
70 }
71
72 )");
73 }
74
TEST_F(MslGeneratorImplTest,Emit_Decoration_EntryPoint_NoReturn_Void)75 TEST_F(MslGeneratorImplTest, Emit_Decoration_EntryPoint_NoReturn_Void) {
76 Func("main", ast::VariableList{}, ty.void_(),
77 ast::StatementList{/* no explicit return */},
78 {Stage(ast::PipelineStage::kFragment)});
79
80 GeneratorImpl& gen = Build();
81
82 ASSERT_TRUE(gen.Generate()) << gen.error();
83 EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
84
85 using namespace metal;
86 fragment void main() {
87 return;
88 }
89
90 )");
91 }
92
TEST_F(MslGeneratorImplTest,Emit_Decoration_EntryPoint_WithInOutVars)93 TEST_F(MslGeneratorImplTest, Emit_Decoration_EntryPoint_WithInOutVars) {
94 // fn frag_main([[location(0)]] foo : f32) -> [[location(1)]] f32 {
95 // return foo;
96 // }
97 auto* foo_in = Param("foo", ty.f32(), {Location(0)});
98 Func("frag_main", ast::VariableList{foo_in}, ty.f32(), {Return("foo")},
99 {Stage(ast::PipelineStage::kFragment)}, {Location(1)});
100
101 GeneratorImpl& gen = SanitizeAndBuild();
102
103 ASSERT_TRUE(gen.Generate()) << gen.error();
104 EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
105
106 using namespace metal;
107 struct tint_symbol_1 {
108 float foo [[user(locn0)]];
109 };
110 struct tint_symbol_2 {
111 float value [[color(1)]];
112 };
113
114 float frag_main_inner(float foo) {
115 return foo;
116 }
117
118 fragment tint_symbol_2 frag_main(tint_symbol_1 tint_symbol [[stage_in]]) {
119 float const inner_result = frag_main_inner(tint_symbol.foo);
120 tint_symbol_2 wrapper_result = {};
121 wrapper_result.value = inner_result;
122 return wrapper_result;
123 }
124
125 )");
126 }
127
TEST_F(MslGeneratorImplTest,Emit_Decoration_EntryPoint_WithInOut_Builtins)128 TEST_F(MslGeneratorImplTest, Emit_Decoration_EntryPoint_WithInOut_Builtins) {
129 // fn frag_main([[position(0)]] coord : vec4<f32>) -> [[frag_depth]] f32 {
130 // return coord.x;
131 // }
132 auto* coord_in =
133 Param("coord", ty.vec4<f32>(), {Builtin(ast::Builtin::kPosition)});
134 Func("frag_main", ast::VariableList{coord_in}, ty.f32(),
135 {Return(MemberAccessor("coord", "x"))},
136 {Stage(ast::PipelineStage::kFragment)},
137 {Builtin(ast::Builtin::kFragDepth)});
138
139 GeneratorImpl& gen = SanitizeAndBuild();
140
141 ASSERT_TRUE(gen.Generate()) << gen.error();
142 EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
143
144 using namespace metal;
145 struct tint_symbol {
146 float value [[depth(any)]];
147 };
148
149 float frag_main_inner(float4 coord) {
150 return coord[0];
151 }
152
153 fragment tint_symbol frag_main(float4 coord [[position]]) {
154 float const inner_result = frag_main_inner(coord);
155 tint_symbol wrapper_result = {};
156 wrapper_result.value = inner_result;
157 return wrapper_result;
158 }
159
160 )");
161 }
162
TEST_F(MslGeneratorImplTest,Emit_Decoration_EntryPoint_SharedStruct_DifferentStages)163 TEST_F(MslGeneratorImplTest,
164 Emit_Decoration_EntryPoint_SharedStruct_DifferentStages) {
165 // struct Interface {
166 // [[location(1)]] col1 : f32;
167 // [[location(2)]] col2 : f32;
168 // [[builtin(position)]] pos : vec4<f32>;
169 // };
170 // fn vert_main() -> Interface {
171 // return Interface(0.4, 0.6, vec4<f32>());
172 // }
173 // fn frag_main(colors : Interface) {
174 // const r = colors.col1;
175 // const g = colors.col2;
176 // }
177 auto* interface_struct = Structure(
178 "Interface",
179 {
180 Member("col1", ty.f32(), {Location(1)}),
181 Member("col2", ty.f32(), {Location(2)}),
182 Member("pos", ty.vec4<f32>(), {Builtin(ast::Builtin::kPosition)}),
183 });
184
185 Func("vert_main", {}, ty.Of(interface_struct),
186 {Return(Construct(ty.Of(interface_struct), Expr(0.5f), Expr(0.25f),
187 Construct(ty.vec4<f32>())))},
188 {Stage(ast::PipelineStage::kVertex)});
189
190 Func("frag_main", {Param("colors", ty.Of(interface_struct))}, ty.void_(),
191 {
192 WrapInStatement(
193 Const("r", ty.f32(), MemberAccessor("colors", "col1"))),
194 WrapInStatement(
195 Const("g", ty.f32(), MemberAccessor("colors", "col2"))),
196 },
197 {Stage(ast::PipelineStage::kFragment)});
198
199 GeneratorImpl& gen = SanitizeAndBuild();
200
201 ASSERT_TRUE(gen.Generate()) << gen.error();
202 EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
203
204 using namespace metal;
205 struct Interface {
206 float col1;
207 float col2;
208 float4 pos;
209 };
210 struct tint_symbol {
211 float col1 [[user(locn1)]];
212 float col2 [[user(locn2)]];
213 float4 pos [[position]];
214 };
215 struct tint_symbol_2 {
216 float col1 [[user(locn1)]];
217 float col2 [[user(locn2)]];
218 };
219
220 Interface vert_main_inner() {
221 Interface const tint_symbol_3 = {.col1=0.5f, .col2=0.25f, .pos=float4()};
222 return tint_symbol_3;
223 }
224
225 vertex tint_symbol vert_main() {
226 Interface const inner_result = vert_main_inner();
227 tint_symbol wrapper_result = {};
228 wrapper_result.col1 = inner_result.col1;
229 wrapper_result.col2 = inner_result.col2;
230 wrapper_result.pos = inner_result.pos;
231 return wrapper_result;
232 }
233
234 void frag_main_inner(Interface colors) {
235 float const r = colors.col1;
236 float const g = colors.col2;
237 }
238
239 fragment void frag_main(float4 pos [[position]], tint_symbol_2 tint_symbol_1 [[stage_in]]) {
240 Interface const tint_symbol_4 = {.col1=tint_symbol_1.col1, .col2=tint_symbol_1.col2, .pos=pos};
241 frag_main_inner(tint_symbol_4);
242 return;
243 }
244
245 )");
246 }
247
TEST_F(MslGeneratorImplTest,Emit_Decoration_EntryPoint_SharedStruct_HelperFunction)248 TEST_F(MslGeneratorImplTest,
249 Emit_Decoration_EntryPoint_SharedStruct_HelperFunction) {
250 // struct VertexOutput {
251 // [[builtin(position)]] pos : vec4<f32>;
252 // };
253 // fn foo(x : f32) -> VertexOutput {
254 // return VertexOutput(vec4<f32>(x, x, x, 1.0));
255 // }
256 // fn vert_main1() -> VertexOutput {
257 // return foo(0.5);
258 // }
259 // fn vert_main2() -> VertexOutput {
260 // return foo(0.25);
261 // }
262 auto* vertex_output_struct = Structure(
263 "VertexOutput",
264 {Member("pos", ty.vec4<f32>(), {Builtin(ast::Builtin::kPosition)})});
265
266 Func("foo", {Param("x", ty.f32())}, ty.Of(vertex_output_struct),
267 {Return(Construct(ty.Of(vertex_output_struct),
268 Construct(ty.vec4<f32>(), "x", "x", "x", Expr(1.f))))},
269 {});
270
271 Func("vert_main1", {}, ty.Of(vertex_output_struct),
272 {Return(Expr(Call("foo", Expr(0.5f))))},
273 {Stage(ast::PipelineStage::kVertex)});
274
275 Func("vert_main2", {}, ty.Of(vertex_output_struct),
276 {Return(Expr(Call("foo", Expr(0.25f))))},
277 {Stage(ast::PipelineStage::kVertex)});
278
279 GeneratorImpl& gen = SanitizeAndBuild();
280
281 ASSERT_TRUE(gen.Generate()) << gen.error();
282 EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
283
284 using namespace metal;
285 struct VertexOutput {
286 float4 pos;
287 };
288 struct tint_symbol {
289 float4 pos [[position]];
290 };
291 struct tint_symbol_1 {
292 float4 pos [[position]];
293 };
294
295 VertexOutput foo(float x) {
296 VertexOutput const tint_symbol_2 = {.pos=float4(x, x, x, 1.0f)};
297 return tint_symbol_2;
298 }
299
300 VertexOutput vert_main1_inner() {
301 return foo(0.5f);
302 }
303
304 vertex tint_symbol vert_main1() {
305 VertexOutput const inner_result = vert_main1_inner();
306 tint_symbol wrapper_result = {};
307 wrapper_result.pos = inner_result.pos;
308 return wrapper_result;
309 }
310
311 VertexOutput vert_main2_inner() {
312 return foo(0.25f);
313 }
314
315 vertex tint_symbol_1 vert_main2() {
316 VertexOutput const inner_result_1 = vert_main2_inner();
317 tint_symbol_1 wrapper_result_1 = {};
318 wrapper_result_1.pos = inner_result_1.pos;
319 return wrapper_result_1;
320 }
321
322 )");
323 }
324
TEST_F(MslGeneratorImplTest,Emit_FunctionDecoration_EntryPoint_With_RW_StorageBuffer)325 TEST_F(MslGeneratorImplTest,
326 Emit_FunctionDecoration_EntryPoint_With_RW_StorageBuffer) {
327 auto* s = Structure("Data",
328 {
329 Member("a", ty.i32()),
330 Member("b", ty.f32()),
331 },
332 {create<ast::StructBlockDecoration>()});
333
334 Global("coord", ty.Of(s), ast::StorageClass::kStorage,
335 ast::Access::kReadWrite,
336 ast::DecorationList{
337 create<ast::BindingDecoration>(0),
338 create<ast::GroupDecoration>(0),
339 });
340
341 auto* var = Var("v", ty.f32(), ast::StorageClass::kNone,
342 MemberAccessor("coord", "b"));
343
344 Func("frag_main", ast::VariableList{}, ty.void_(),
345 ast::StatementList{
346 Decl(var),
347 Return(),
348 },
349 {
350 Stage(ast::PipelineStage::kFragment),
351 });
352
353 GeneratorImpl& gen = SanitizeAndBuild();
354
355 ASSERT_TRUE(gen.Generate()) << gen.error();
356 EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
357
358 using namespace metal;
359 struct Data {
360 /* 0x0000 */ int a;
361 /* 0x0004 */ float b;
362 };
363
364 fragment void frag_main(device Data* tint_symbol [[buffer(0)]]) {
365 float v = (*(tint_symbol)).b;
366 return;
367 }
368
369 )");
370 }
371
TEST_F(MslGeneratorImplTest,Emit_FunctionDecoration_EntryPoint_With_RO_StorageBuffer)372 TEST_F(MslGeneratorImplTest,
373 Emit_FunctionDecoration_EntryPoint_With_RO_StorageBuffer) {
374 auto* s = Structure("Data",
375 {
376 Member("a", ty.i32()),
377 Member("b", ty.f32()),
378 },
379 {create<ast::StructBlockDecoration>()});
380
381 Global("coord", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kRead,
382 ast::DecorationList{
383 create<ast::BindingDecoration>(0),
384 create<ast::GroupDecoration>(0),
385 });
386
387 auto* var = Var("v", ty.f32(), ast::StorageClass::kNone,
388 MemberAccessor("coord", "b"));
389
390 Func("frag_main", ast::VariableList{}, ty.void_(),
391 ast::StatementList{
392 Decl(var),
393 Return(),
394 },
395 {
396 Stage(ast::PipelineStage::kFragment),
397 });
398
399 GeneratorImpl& gen = SanitizeAndBuild();
400
401 ASSERT_TRUE(gen.Generate()) << gen.error();
402 EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
403
404 using namespace metal;
405 struct Data {
406 /* 0x0000 */ int a;
407 /* 0x0004 */ float b;
408 };
409
410 fragment void frag_main(const device Data* tint_symbol [[buffer(0)]]) {
411 float v = (*(tint_symbol)).b;
412 return;
413 }
414
415 )");
416 }
417
TEST_F(MslGeneratorImplTest,Emit_Decoration_Called_By_EntryPoint_With_Uniform)418 TEST_F(MslGeneratorImplTest,
419 Emit_Decoration_Called_By_EntryPoint_With_Uniform) {
420 auto* ubo_ty = Structure("UBO", {Member("coord", ty.vec4<f32>())},
421 {create<ast::StructBlockDecoration>()});
422 auto* ubo = Global("ubo", ty.Of(ubo_ty), ast::StorageClass::kUniform,
423 ast::DecorationList{
424 create<ast::BindingDecoration>(0),
425 create<ast::GroupDecoration>(0),
426 });
427
428 Func("sub_func",
429 {
430 Param("param", ty.f32()),
431 },
432 ty.f32(),
433 {
434 Return(MemberAccessor(MemberAccessor(ubo, "coord"), "x")),
435 });
436
437 auto* var =
438 Var("v", ty.f32(), ast::StorageClass::kNone, Call("sub_func", 1.0f));
439
440 Func("frag_main", {}, ty.void_(),
441 {
442 Decl(var),
443 Return(),
444 },
445 {
446 Stage(ast::PipelineStage::kFragment),
447 });
448
449 GeneratorImpl& gen = SanitizeAndBuild();
450
451 ASSERT_TRUE(gen.Generate()) << gen.error();
452 EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
453
454 using namespace metal;
455 struct UBO {
456 /* 0x0000 */ float4 coord;
457 };
458
459 float sub_func(float param, const constant UBO* const tint_symbol) {
460 return (*(tint_symbol)).coord[0];
461 }
462
463 fragment void frag_main(const constant UBO* tint_symbol_1 [[buffer(0)]]) {
464 float v = sub_func(1.0f, tint_symbol_1);
465 return;
466 }
467
468 )");
469 }
470
TEST_F(MslGeneratorImplTest,Emit_FunctionDecoration_Called_By_EntryPoint_With_RW_StorageBuffer)471 TEST_F(MslGeneratorImplTest,
472 Emit_FunctionDecoration_Called_By_EntryPoint_With_RW_StorageBuffer) {
473 auto* s = Structure("Data",
474 {
475 Member("a", ty.i32()),
476 Member("b", ty.f32()),
477 },
478 {create<ast::StructBlockDecoration>()});
479
480 Global("coord", ty.Of(s), ast::StorageClass::kStorage,
481 ast::Access::kReadWrite,
482 ast::DecorationList{
483 create<ast::BindingDecoration>(0),
484 create<ast::GroupDecoration>(0),
485 });
486
487 ast::VariableList params;
488 params.push_back(Param("param", ty.f32()));
489
490 auto body = ast::StatementList{Return(MemberAccessor("coord", "b"))};
491
492 Func("sub_func", params, ty.f32(), body, {});
493
494 auto* var =
495 Var("v", ty.f32(), ast::StorageClass::kNone, Call("sub_func", 1.0f));
496
497 Func("frag_main", ast::VariableList{}, ty.void_(),
498 ast::StatementList{
499 Decl(var),
500 Return(),
501 },
502 {
503 Stage(ast::PipelineStage::kFragment),
504 });
505
506 GeneratorImpl& gen = SanitizeAndBuild();
507
508 ASSERT_TRUE(gen.Generate()) << gen.error();
509 EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
510
511 using namespace metal;
512 struct Data {
513 /* 0x0000 */ int a;
514 /* 0x0004 */ float b;
515 };
516
517 float sub_func(float param, device Data* const tint_symbol) {
518 return (*(tint_symbol)).b;
519 }
520
521 fragment void frag_main(device Data* tint_symbol_1 [[buffer(0)]]) {
522 float v = sub_func(1.0f, tint_symbol_1);
523 return;
524 }
525
526 )");
527 }
528
TEST_F(MslGeneratorImplTest,Emit_FunctionDecoration_Called_By_EntryPoint_With_RO_StorageBuffer)529 TEST_F(MslGeneratorImplTest,
530 Emit_FunctionDecoration_Called_By_EntryPoint_With_RO_StorageBuffer) {
531 auto* s = Structure("Data",
532 {
533 Member("a", ty.i32()),
534 Member("b", ty.f32()),
535 },
536 {create<ast::StructBlockDecoration>()});
537
538 Global("coord", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kRead,
539 ast::DecorationList{
540 create<ast::BindingDecoration>(0),
541 create<ast::GroupDecoration>(0),
542 });
543
544 ast::VariableList params;
545 params.push_back(Param("param", ty.f32()));
546
547 auto body = ast::StatementList{Return(MemberAccessor("coord", "b"))};
548
549 Func("sub_func", params, ty.f32(), body, {});
550
551 auto* var =
552 Var("v", ty.f32(), ast::StorageClass::kNone, Call("sub_func", 1.0f));
553
554 Func("frag_main", ast::VariableList{}, ty.void_(),
555 ast::StatementList{
556 Decl(var),
557 Return(),
558 },
559 {
560 Stage(ast::PipelineStage::kFragment),
561 });
562
563 GeneratorImpl& gen = SanitizeAndBuild();
564
565 ASSERT_TRUE(gen.Generate()) << gen.error();
566 EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
567
568 using namespace metal;
569 struct Data {
570 /* 0x0000 */ int a;
571 /* 0x0004 */ float b;
572 };
573
574 float sub_func(float param, const device Data* const tint_symbol) {
575 return (*(tint_symbol)).b;
576 }
577
578 fragment void frag_main(const device Data* tint_symbol_1 [[buffer(0)]]) {
579 float v = sub_func(1.0f, tint_symbol_1);
580 return;
581 }
582
583 )");
584 }
585
TEST_F(MslGeneratorImplTest,Emit_Function_WithArrayParams)586 TEST_F(MslGeneratorImplTest, Emit_Function_WithArrayParams) {
587 ast::VariableList params;
588 params.push_back(Param("a", ty.array<f32, 5>()));
589
590 Func("my_func", params, ty.void_(),
591 {
592 Return(),
593 });
594
595 GeneratorImpl& gen = SanitizeAndBuild();
596
597 gen.increment_indent();
598
599 ASSERT_TRUE(gen.Generate()) << gen.error();
600 EXPECT_EQ(gen.result(), R"( #include <metal_stdlib>
601
602 using namespace metal;
603 struct tint_array_wrapper {
604 float arr[5];
605 };
606
607 void my_func(tint_array_wrapper a) {
608 return;
609 }
610
611 )");
612 }
613
TEST_F(MslGeneratorImplTest,Emit_Function_WithArrayReturn)614 TEST_F(MslGeneratorImplTest, Emit_Function_WithArrayReturn) {
615 Func("my_func", {}, ty.array<f32, 5>(),
616 {
617 Return(Construct(ty.array<f32, 5>())),
618 });
619
620 GeneratorImpl& gen = SanitizeAndBuild();
621
622 gen.increment_indent();
623
624 ASSERT_TRUE(gen.Generate()) << gen.error();
625 EXPECT_EQ(gen.result(), R"( #include <metal_stdlib>
626
627 using namespace metal;
628 struct tint_array_wrapper {
629 float arr[5];
630 };
631
632 tint_array_wrapper my_func() {
633 tint_array_wrapper const tint_symbol = {.arr={}};
634 return tint_symbol;
635 }
636
637 )");
638 }
639
640 // https://crbug.com/tint/297
TEST_F(MslGeneratorImplTest,Emit_Function_Multiple_EntryPoint_With_Same_ModuleVar)641 TEST_F(MslGeneratorImplTest,
642 Emit_Function_Multiple_EntryPoint_With_Same_ModuleVar) {
643 // [[block]] struct Data {
644 // d : f32;
645 // };
646 // [[binding(0), group(0)]] var<storage> data : Data;
647 //
648 // [[stage(compute), workgroup_size(1)]]
649 // fn a() {
650 // return;
651 // }
652 //
653 // [[stage(compute), workgroup_size(1)]]
654 // fn b() {
655 // return;
656 // }
657
658 auto* s = Structure("Data", {Member("d", ty.f32())},
659 {create<ast::StructBlockDecoration>()});
660
661 Global("data", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kReadWrite,
662 ast::DecorationList{
663 create<ast::BindingDecoration>(0),
664 create<ast::GroupDecoration>(0),
665 });
666
667 {
668 auto* var = Var("v", ty.f32(), ast::StorageClass::kNone,
669 MemberAccessor("data", "d"));
670
671 Func("a", ast::VariableList{}, ty.void_(),
672 ast::StatementList{
673 Decl(var),
674 Return(),
675 },
676 {
677 Stage(ast::PipelineStage::kCompute),
678 WorkgroupSize(1),
679 });
680 }
681
682 {
683 auto* var = Var("v", ty.f32(), ast::StorageClass::kNone,
684 MemberAccessor("data", "d"));
685
686 Func("b", ast::VariableList{}, ty.void_(),
687 ast::StatementList{Decl(var), Return()},
688 {
689 Stage(ast::PipelineStage::kCompute),
690 WorkgroupSize(1),
691 });
692 }
693
694 GeneratorImpl& gen = SanitizeAndBuild();
695
696 ASSERT_TRUE(gen.Generate()) << gen.error();
697 EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
698
699 using namespace metal;
700 struct Data {
701 /* 0x0000 */ float d;
702 };
703
704 kernel void a(device Data* tint_symbol [[buffer(0)]]) {
705 float v = (*(tint_symbol)).d;
706 return;
707 }
708
709 kernel void b(device Data* tint_symbol_1 [[buffer(0)]]) {
710 float v = (*(tint_symbol_1)).d;
711 return;
712 }
713
714 )");
715 }
716
717 } // namespace
718 } // namespace msl
719 } // namespace writer
720 } // namespace tint
721