blob: 2892d6054a6555e43a8762b68443388851bea6d9 [file] [log] [blame]
James Price73f54952024-05-16 21:44:59 +00001// Copyright 2024 The Dawn & Tint Authors
2//
3// Redistribution and use in source and binary forms, with or without
4// modification, are permitted provided that the following conditions are met:
5//
6// 1. Redistributions of source code must retain the above copyright notice, this
7// list of conditions and the following disclaimer.
8//
9// 2. Redistributions in binary form must reproduce the above copyright notice,
10// this list of conditions and the following disclaimer in the documentation
11// and/or other materials provided with the distribution.
12//
13// 3. Neither the name of the copyright holder nor the names of its
14// contributors may be used to endorse or promote products derived from
15// this software without specific prior written permission.
16//
17// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
18// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
19// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
20// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
21// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
22// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
23// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
24// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
25// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
26// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
27
28#include "src/tint/lang/msl/writer/raise/module_scope_vars.h"
29
30#include <utility>
31
David Neto0be0f1b2024-05-28 16:15:08 +000032#include "src/tint/lang/core/fluent_types.h"
James Price73f54952024-05-16 21:44:59 +000033#include "src/tint/lang/core/ir/transform/helper_test.h"
James Price06f9dbc2024-05-24 16:11:13 +000034#include "src/tint/lang/core/type/sampled_texture.h"
James Price73f54952024-05-16 21:44:59 +000035
36using namespace tint::core::fluent_types; // NOLINT
37using namespace tint::core::number_suffixes; // NOLINT
38
39namespace tint::msl::writer::raise {
40namespace {
41
42using MslWriter_ModuleScopeVarsTest = core::ir::transform::TransformTest;
43
44TEST_F(MslWriter_ModuleScopeVarsTest, NoModuleScopeVars) {
45 auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
46 b.Append(func->Block(), [&] {
47 auto* var = b.Var<function, i32>("v");
48 b.Load(var);
49 b.Return(func);
50 });
51
52 auto* src = R"(
53%foo = @fragment func():void {
54 $B1: {
55 %v:ptr<function, i32, read_write> = var
56 %3:i32 = load %v
57 ret
58 }
59}
60)";
61 EXPECT_EQ(src, str());
62
63 auto* expect = src;
64
65 Run(ModuleScopeVars);
66
67 EXPECT_EQ(expect, str());
68}
69
70TEST_F(MslWriter_ModuleScopeVarsTest, Private) {
71 auto* var_a = b.Var("a", ty.ptr<private_, i32>());
72 auto* var_b = b.Var("b", ty.ptr<private_, i32>());
73 mod.root_block->Append(var_a);
74 mod.root_block->Append(var_b);
75
76 auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
77 b.Append(func->Block(), [&] {
78 auto* load_a = b.Load(var_a);
79 auto* load_b = b.Load(var_b);
80 b.Store(var_a, b.Add<i32>(load_a, load_b));
81 b.Return(func);
82 });
83
84 auto* src = R"(
85$B1: { # root
86 %a:ptr<private, i32, read_write> = var
87 %b:ptr<private, i32, read_write> = var
88}
89
90%foo = @fragment func():void {
91 $B2: {
92 %4:i32 = load %a
93 %5:i32 = load %b
94 %6:i32 = add %4, %5
95 store %a, %6
96 ret
97 }
98}
99)";
100 EXPECT_EQ(src, str());
101
102 auto* expect = R"(
103tint_module_vars_struct = struct @align(1) {
104 a:ptr<private, i32, read_write> @offset(0)
105 b:ptr<private, i32, read_write> @offset(0)
106}
107
108%foo = @fragment func():void {
109 $B1: {
110 %a:ptr<private, i32, read_write> = var
111 %b:ptr<private, i32, read_write> = var
112 %4:tint_module_vars_struct = construct %a, %b
113 %tint_module_vars:tint_module_vars_struct = let %4
114 %6:ptr<private, i32, read_write> = access %tint_module_vars, 0u
115 %7:i32 = load %6
116 %8:ptr<private, i32, read_write> = access %tint_module_vars, 1u
117 %9:i32 = load %8
118 %10:i32 = add %7, %9
119 %11:ptr<private, i32, read_write> = access %tint_module_vars, 0u
120 store %11, %10
121 ret
122 }
123}
124)";
125
126 Run(ModuleScopeVars);
127
128 EXPECT_EQ(expect, str());
129}
130
131TEST_F(MslWriter_ModuleScopeVarsTest, Private_WithInitializers) {
132 auto* var_a = b.Var<private_>("a", 42_i);
133 auto* var_b = b.Var<private_>("b", -1_i);
134 mod.root_block->Append(var_a);
135 mod.root_block->Append(var_b);
136
137 auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
138 b.Append(func->Block(), [&] {
139 auto* load_a = b.Load(var_a);
140 auto* load_b = b.Load(var_b);
141 b.Store(var_a, b.Add<i32>(load_a, load_b));
142 b.Return(func);
143 });
144
145 auto* src = R"(
146$B1: { # root
147 %a:ptr<private, i32, read_write> = var, 42i
148 %b:ptr<private, i32, read_write> = var, -1i
149}
150
151%foo = @fragment func():void {
152 $B2: {
153 %4:i32 = load %a
154 %5:i32 = load %b
155 %6:i32 = add %4, %5
156 store %a, %6
157 ret
158 }
159}
160)";
161 EXPECT_EQ(src, str());
162
163 auto* expect = R"(
164tint_module_vars_struct = struct @align(1) {
165 a:ptr<private, i32, read_write> @offset(0)
166 b:ptr<private, i32, read_write> @offset(0)
167}
168
169%foo = @fragment func():void {
170 $B1: {
171 %a:ptr<private, i32, read_write> = var, 42i
172 %b:ptr<private, i32, read_write> = var, -1i
173 %4:tint_module_vars_struct = construct %a, %b
174 %tint_module_vars:tint_module_vars_struct = let %4
175 %6:ptr<private, i32, read_write> = access %tint_module_vars, 0u
176 %7:i32 = load %6
177 %8:ptr<private, i32, read_write> = access %tint_module_vars, 1u
178 %9:i32 = load %8
179 %10:i32 = add %7, %9
180 %11:ptr<private, i32, read_write> = access %tint_module_vars, 0u
181 store %11, %10
182 ret
183 }
184}
185)";
186
187 Run(ModuleScopeVars);
188
189 EXPECT_EQ(expect, str());
190}
191
192TEST_F(MslWriter_ModuleScopeVarsTest, Storage) {
193 auto* var_a = b.Var("a", ty.ptr<storage, i32, core::Access::kRead>());
194 auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>());
195 var_a->SetBindingPoint(1, 2);
196 var_b->SetBindingPoint(3, 4);
197 mod.root_block->Append(var_a);
198 mod.root_block->Append(var_b);
199
200 auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
201 b.Append(func->Block(), [&] {
202 auto* load_a = b.Load(var_a);
203 auto* load_b = b.Load(var_b);
204 b.Store(var_b, b.Add<i32>(load_a, load_b));
205 b.Return(func);
206 });
207
208 auto* src = R"(
209$B1: { # root
210 %a:ptr<storage, i32, read> = var @binding_point(1, 2)
211 %b:ptr<storage, i32, read_write> = var @binding_point(3, 4)
212}
213
214%foo = @fragment func():void {
215 $B2: {
216 %4:i32 = load %a
217 %5:i32 = load %b
218 %6:i32 = add %4, %5
219 store %b, %6
220 ret
221 }
222}
223)";
224 EXPECT_EQ(src, str());
225
226 auto* expect = R"(
227tint_module_vars_struct = struct @align(1) {
228 a:ptr<storage, i32, read> @offset(0)
229 b:ptr<storage, i32, read_write> @offset(0)
230}
231
232%foo = @fragment func(%a:ptr<storage, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void {
233 $B1: {
234 %4:tint_module_vars_struct = construct %a, %b
235 %tint_module_vars:tint_module_vars_struct = let %4
236 %6:ptr<storage, i32, read> = access %tint_module_vars, 0u
237 %7:i32 = load %6
238 %8:ptr<storage, i32, read_write> = access %tint_module_vars, 1u
239 %9:i32 = load %8
240 %10:i32 = add %7, %9
241 %11:ptr<storage, i32, read_write> = access %tint_module_vars, 1u
242 store %11, %10
243 ret
244 }
245}
246)";
247
248 Run(ModuleScopeVars);
249
250 EXPECT_EQ(expect, str());
251}
252
253TEST_F(MslWriter_ModuleScopeVarsTest, Uniform) {
254 auto* var_a = b.Var("a", ty.ptr<uniform, i32>());
255 auto* var_b = b.Var("b", ty.ptr<uniform, i32>());
256 var_a->SetBindingPoint(1, 2);
257 var_b->SetBindingPoint(3, 4);
258 mod.root_block->Append(var_a);
259 mod.root_block->Append(var_b);
260
261 auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
262 b.Append(func->Block(), [&] {
263 auto* load_a = b.Load(var_a);
264 auto* load_b = b.Load(var_b);
265 b.Add<i32>(load_a, load_b);
266 b.Return(func);
267 });
268
269 auto* src = R"(
270$B1: { # root
271 %a:ptr<uniform, i32, read> = var @binding_point(1, 2)
272 %b:ptr<uniform, i32, read> = var @binding_point(3, 4)
273}
274
275%foo = @fragment func():void {
276 $B2: {
277 %4:i32 = load %a
278 %5:i32 = load %b
279 %6:i32 = add %4, %5
280 ret
281 }
282}
283)";
284 EXPECT_EQ(src, str());
285
286 auto* expect = R"(
287tint_module_vars_struct = struct @align(1) {
288 a:ptr<uniform, i32, read> @offset(0)
289 b:ptr<uniform, i32, read> @offset(0)
290}
291
292%foo = @fragment func(%a:ptr<uniform, i32, read> [@binding_point(1, 2)], %b:ptr<uniform, i32, read> [@binding_point(3, 4)]):void {
293 $B1: {
294 %4:tint_module_vars_struct = construct %a, %b
295 %tint_module_vars:tint_module_vars_struct = let %4
296 %6:ptr<uniform, i32, read> = access %tint_module_vars, 0u
297 %7:i32 = load %6
298 %8:ptr<uniform, i32, read> = access %tint_module_vars, 1u
299 %9:i32 = load %8
300 %10:i32 = add %7, %9
301 ret
302 }
303}
304)";
305
306 Run(ModuleScopeVars);
307
308 EXPECT_EQ(expect, str());
309}
310
James Price06f9dbc2024-05-24 16:11:13 +0000311TEST_F(MslWriter_ModuleScopeVarsTest, HandleTypes) {
312 auto* t = ty.Get<core::type::SampledTexture>(core::type::TextureDimension::k2d, ty.f32());
313 auto* var_t = b.Var("t", ty.ptr<handle>(t));
314 auto* var_s = b.Var("s", ty.ptr<handle>(ty.sampler()));
315 var_t->SetBindingPoint(1, 2);
316 var_s->SetBindingPoint(3, 4);
317 mod.root_block->Append(var_t);
318 mod.root_block->Append(var_s);
319
320 auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
321 b.Append(func->Block(), [&] {
322 auto* load_t = b.Load(var_t);
323 auto* load_s = b.Load(var_s);
324 b.Call<vec4<f32>>(core::BuiltinFn::kTextureSample, load_t, load_s, b.Splat<vec2<f32>>(0_f));
325 b.Return(func);
326 });
327
328 auto* src = R"(
329$B1: { # root
330 %t:ptr<handle, texture_2d<f32>, read> = var @binding_point(1, 2)
331 %s:ptr<handle, sampler, read> = var @binding_point(3, 4)
332}
333
334%foo = @fragment func():void {
335 $B2: {
336 %4:texture_2d<f32> = load %t
337 %5:sampler = load %s
338 %6:vec4<f32> = textureSample %4, %5, vec2<f32>(0.0f)
339 ret
340 }
341}
342)";
343 EXPECT_EQ(src, str());
344
345 auto* expect = R"(
346tint_module_vars_struct = struct @align(1) {
347 t:texture_2d<f32> @offset(0)
348 s:sampler @offset(0)
349}
350
351%foo = @fragment func(%t:texture_2d<f32> [@binding_point(1, 2)], %s:sampler [@binding_point(3, 4)]):void {
352 $B1: {
353 %4:tint_module_vars_struct = construct %t, %s
354 %tint_module_vars:tint_module_vars_struct = let %4
355 %6:texture_2d<f32> = access %tint_module_vars, 0u
356 %7:sampler = access %tint_module_vars, 1u
357 %8:vec4<f32> = textureSample %6, %7, vec2<f32>(0.0f)
358 ret
359 }
360}
361)";
362
363 Run(ModuleScopeVars);
364
365 EXPECT_EQ(expect, str());
366}
367
James Pricec451f622024-05-27 17:28:21 +0000368TEST_F(MslWriter_ModuleScopeVarsTest, Workgroup) {
369 auto* var_a = b.Var("a", ty.ptr<workgroup, i32>());
370 auto* var_b = b.Var("b", ty.ptr<workgroup, i32>());
371 mod.root_block->Append(var_a);
372 mod.root_block->Append(var_b);
373
374 auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kCompute,
375 std::array<uint32_t, 3>{1u, 1u, 1u});
376 b.Append(func->Block(), [&] {
377 auto* load_a = b.Load(var_a);
378 auto* load_b = b.Load(var_b);
379 b.Store(var_a, b.Add<i32>(load_a, load_b));
380 b.Return(func);
381 });
382
383 auto* src = R"(
384$B1: { # root
385 %a:ptr<workgroup, i32, read_write> = var
386 %b:ptr<workgroup, i32, read_write> = var
387}
388
389%foo = @compute @workgroup_size(1, 1, 1) func():void {
390 $B2: {
391 %4:i32 = load %a
392 %5:i32 = load %b
393 %6:i32 = add %4, %5
394 store %a, %6
395 ret
396 }
397}
398)";
399 EXPECT_EQ(src, str());
400
401 auto* expect = R"(
402tint_module_vars_struct = struct @align(1) {
403 a:ptr<workgroup, i32, read_write> @offset(0)
404 b:ptr<workgroup, i32, read_write> @offset(0)
405}
406
407tint_symbol_2 = struct @align(4) {
408 tint_symbol:i32 @offset(0)
409 tint_symbol_1:i32 @offset(4)
410}
411
412%foo = @compute @workgroup_size(1, 1, 1) func(%2:ptr<workgroup, tint_symbol_2, read_write>):void {
413 $B1: {
414 %a:ptr<workgroup, i32, read_write> = access %2, 0u
415 %b:ptr<workgroup, i32, read_write> = access %2, 1u
416 %5:tint_module_vars_struct = construct %a, %b
417 %tint_module_vars:tint_module_vars_struct = let %5
418 %7:ptr<workgroup, i32, read_write> = access %tint_module_vars, 0u
419 %8:i32 = load %7
420 %9:ptr<workgroup, i32, read_write> = access %tint_module_vars, 1u
421 %10:i32 = load %9
422 %11:i32 = add %8, %10
423 %12:ptr<workgroup, i32, read_write> = access %tint_module_vars, 0u
424 store %12, %11
425 ret
426 }
427}
428)";
429
430 Run(ModuleScopeVars);
431
432 EXPECT_EQ(expect, str());
433}
434
James Price73f54952024-05-16 21:44:59 +0000435TEST_F(MslWriter_ModuleScopeVarsTest, MultipleAddressSpaces) {
436 auto* var_a = b.Var("a", ty.ptr<uniform, i32, core::Access::kRead>());
437 auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>());
438 auto* var_c = b.Var("c", ty.ptr<private_, i32, core::Access::kReadWrite>());
439 var_a->SetBindingPoint(1, 2);
440 var_b->SetBindingPoint(3, 4);
441 mod.root_block->Append(var_a);
442 mod.root_block->Append(var_b);
443 mod.root_block->Append(var_c);
444
445 auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
446 b.Append(func->Block(), [&] {
447 auto* load_a = b.Load(var_a);
448 auto* load_b = b.Load(var_b);
449 auto* load_c = b.Load(var_c);
450 b.Store(var_b, b.Add<i32>(load_a, b.Add<i32>(load_b, load_c)));
451 b.Return(func);
452 });
453
454 auto* src = R"(
455$B1: { # root
456 %a:ptr<uniform, i32, read> = var @binding_point(1, 2)
457 %b:ptr<storage, i32, read_write> = var @binding_point(3, 4)
458 %c:ptr<private, i32, read_write> = var
459}
460
461%foo = @fragment func():void {
462 $B2: {
463 %5:i32 = load %a
464 %6:i32 = load %b
465 %7:i32 = load %c
466 %8:i32 = add %6, %7
467 %9:i32 = add %5, %8
468 store %b, %9
469 ret
470 }
471}
472)";
473 EXPECT_EQ(src, str());
474
475 auto* expect = R"(
476tint_module_vars_struct = struct @align(1) {
477 a:ptr<uniform, i32, read> @offset(0)
478 b:ptr<storage, i32, read_write> @offset(0)
479 c:ptr<private, i32, read_write> @offset(0)
480}
481
482%foo = @fragment func(%a:ptr<uniform, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void {
483 $B1: {
484 %c:ptr<private, i32, read_write> = var
485 %5:tint_module_vars_struct = construct %a, %b, %c
486 %tint_module_vars:tint_module_vars_struct = let %5
487 %7:ptr<uniform, i32, read> = access %tint_module_vars, 0u
488 %8:i32 = load %7
489 %9:ptr<storage, i32, read_write> = access %tint_module_vars, 1u
490 %10:i32 = load %9
491 %11:ptr<private, i32, read_write> = access %tint_module_vars, 2u
492 %12:i32 = load %11
493 %13:i32 = add %10, %12
494 %14:i32 = add %8, %13
495 %15:ptr<storage, i32, read_write> = access %tint_module_vars, 1u
496 store %15, %14
497 ret
498 }
499}
500)";
501
502 Run(ModuleScopeVars);
503
504 EXPECT_EQ(expect, str());
505}
506
507TEST_F(MslWriter_ModuleScopeVarsTest, EntryPointHasExistingParameters) {
508 auto* var_a = b.Var("a", ty.ptr<storage, i32, core::Access::kRead>());
509 auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>());
510 var_a->SetBindingPoint(1, 2);
511 var_b->SetBindingPoint(3, 4);
512 mod.root_block->Append(var_a);
513 mod.root_block->Append(var_b);
514
515 auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
516 auto* param = b.FunctionParam<i32>("param");
517 param->SetLocation(
518 core::ir::Location{1_u, core::Interpolation{core::InterpolationType::kFlat}});
519 func->SetParams({param});
520 b.Append(func->Block(), [&] {
521 auto* load_a = b.Load(var_a);
522 auto* load_b = b.Load(var_b);
523 b.Store(var_b, b.Add<i32>(load_a, b.Add<i32>(load_b, param)));
524 b.Return(func);
525 });
526
527 auto* src = R"(
528$B1: { # root
529 %a:ptr<storage, i32, read> = var @binding_point(1, 2)
530 %b:ptr<storage, i32, read_write> = var @binding_point(3, 4)
531}
532
533%foo = @fragment func(%param:i32 [@location(1), @interpolate(flat)]):void {
534 $B2: {
535 %5:i32 = load %a
536 %6:i32 = load %b
537 %7:i32 = add %6, %param
538 %8:i32 = add %5, %7
539 store %b, %8
540 ret
541 }
542}
543)";
544 EXPECT_EQ(src, str());
545
546 auto* expect = R"(
547tint_module_vars_struct = struct @align(1) {
548 a:ptr<storage, i32, read> @offset(0)
549 b:ptr<storage, i32, read_write> @offset(0)
550}
551
552%foo = @fragment func(%param:i32 [@location(1), @interpolate(flat)], %a:ptr<storage, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void {
553 $B1: {
554 %5:tint_module_vars_struct = construct %a, %b
555 %tint_module_vars:tint_module_vars_struct = let %5
556 %7:ptr<storage, i32, read> = access %tint_module_vars, 0u
557 %8:i32 = load %7
558 %9:ptr<storage, i32, read_write> = access %tint_module_vars, 1u
559 %10:i32 = load %9
560 %11:i32 = add %10, %param
561 %12:i32 = add %8, %11
562 %13:ptr<storage, i32, read_write> = access %tint_module_vars, 1u
563 store %13, %12
564 ret
565 }
566}
567)";
568
569 Run(ModuleScopeVars);
570
571 EXPECT_EQ(expect, str());
572}
573
574TEST_F(MslWriter_ModuleScopeVarsTest, CallFunctionThatUsesVars_NoArgs) {
575 auto* var_a = b.Var("a", ty.ptr<storage, i32, core::Access::kRead>());
576 auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>());
577 var_a->SetBindingPoint(1, 2);
578 var_b->SetBindingPoint(3, 4);
579 mod.root_block->Append(var_a);
580 mod.root_block->Append(var_b);
581
582 auto* foo = b.Function("foo", ty.void_());
583 b.Append(foo->Block(), [&] {
584 auto* load_a = b.Load(var_a);
585 auto* load_b = b.Load(var_b);
586 b.Store(var_b, b.Add<i32>(load_a, load_b));
587 b.Return(foo);
588 });
589
590 auto* func = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kFragment);
591 b.Append(func->Block(), [&] {
592 b.Call(foo);
593 b.Return(func);
594 });
595
596 auto* src = R"(
597$B1: { # root
598 %a:ptr<storage, i32, read> = var @binding_point(1, 2)
599 %b:ptr<storage, i32, read_write> = var @binding_point(3, 4)
600}
601
602%foo = func():void {
603 $B2: {
604 %4:i32 = load %a
605 %5:i32 = load %b
606 %6:i32 = add %4, %5
607 store %b, %6
608 ret
609 }
610}
611%main = @fragment func():void {
612 $B3: {
613 %8:void = call %foo
614 ret
615 }
616}
617)";
618 EXPECT_EQ(src, str());
619
620 auto* expect = R"(
621tint_module_vars_struct = struct @align(1) {
622 a:ptr<storage, i32, read> @offset(0)
623 b:ptr<storage, i32, read_write> @offset(0)
624}
625
626%foo = func(%tint_module_vars:tint_module_vars_struct):void {
627 $B1: {
628 %3:ptr<storage, i32, read> = access %tint_module_vars, 0u
629 %4:i32 = load %3
630 %5:ptr<storage, i32, read_write> = access %tint_module_vars, 1u
631 %6:i32 = load %5
632 %7:i32 = add %4, %6
633 %8:ptr<storage, i32, read_write> = access %tint_module_vars, 1u
634 store %8, %7
635 ret
636 }
637}
638%main = @fragment func(%a:ptr<storage, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void {
639 $B2: {
640 %12:tint_module_vars_struct = construct %a, %b
641 %tint_module_vars_1:tint_module_vars_struct = let %12 # %tint_module_vars_1: 'tint_module_vars'
642 %14:void = call %foo, %tint_module_vars_1
643 ret
644 }
645}
646)";
647
648 Run(ModuleScopeVars);
649
650 EXPECT_EQ(expect, str());
651}
652
653TEST_F(MslWriter_ModuleScopeVarsTest, CallFunctionThatUsesVars_WithExistingParameters) {
654 auto* var_a = b.Var("a", ty.ptr<storage, i32, core::Access::kRead>());
655 auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>());
656 var_a->SetBindingPoint(1, 2);
657 var_b->SetBindingPoint(3, 4);
658 mod.root_block->Append(var_a);
659 mod.root_block->Append(var_b);
660
661 auto* foo = b.Function("foo", ty.void_());
662 auto* param = b.FunctionParam<i32>("param");
663 foo->SetParams({param});
664 b.Append(foo->Block(), [&] {
665 auto* load_a = b.Load(var_a);
666 auto* load_b = b.Load(var_b);
667 b.Store(var_b, b.Add<i32>(load_a, b.Add<i32>(load_b, param)));
668 b.Return(foo);
669 });
670
671 auto* func = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kFragment);
672 b.Append(func->Block(), [&] {
673 b.Call(foo, 42_i);
674 b.Return(func);
675 });
676
677 auto* src = R"(
678$B1: { # root
679 %a:ptr<storage, i32, read> = var @binding_point(1, 2)
680 %b:ptr<storage, i32, read_write> = var @binding_point(3, 4)
681}
682
683%foo = func(%param:i32):void {
684 $B2: {
685 %5:i32 = load %a
686 %6:i32 = load %b
687 %7:i32 = add %6, %param
688 %8:i32 = add %5, %7
689 store %b, %8
690 ret
691 }
692}
693%main = @fragment func():void {
694 $B3: {
695 %10:void = call %foo, 42i
696 ret
697 }
698}
699)";
700 EXPECT_EQ(src, str());
701
702 auto* expect = R"(
703tint_module_vars_struct = struct @align(1) {
704 a:ptr<storage, i32, read> @offset(0)
705 b:ptr<storage, i32, read_write> @offset(0)
706}
707
708%foo = func(%param:i32, %tint_module_vars:tint_module_vars_struct):void {
709 $B1: {
710 %4:ptr<storage, i32, read> = access %tint_module_vars, 0u
711 %5:i32 = load %4
712 %6:ptr<storage, i32, read_write> = access %tint_module_vars, 1u
713 %7:i32 = load %6
714 %8:i32 = add %7, %param
715 %9:i32 = add %5, %8
716 %10:ptr<storage, i32, read_write> = access %tint_module_vars, 1u
717 store %10, %9
718 ret
719 }
720}
721%main = @fragment func(%a:ptr<storage, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void {
722 $B2: {
723 %14:tint_module_vars_struct = construct %a, %b
724 %tint_module_vars_1:tint_module_vars_struct = let %14 # %tint_module_vars_1: 'tint_module_vars'
725 %16:void = call %foo, 42i, %tint_module_vars_1
726 ret
727 }
728}
729)";
730
731 Run(ModuleScopeVars);
732
733 EXPECT_EQ(expect, str());
734}
735
James Price06f9dbc2024-05-24 16:11:13 +0000736TEST_F(MslWriter_ModuleScopeVarsTest, CallFunctionThatUsesVars_HandleTypes) {
737 auto* t = ty.Get<core::type::SampledTexture>(core::type::TextureDimension::k2d, ty.f32());
738 auto* var_t = b.Var("t", ty.ptr<handle>(t));
739 auto* var_s = b.Var("s", ty.ptr<handle>(ty.sampler()));
740 var_t->SetBindingPoint(1, 2);
741 var_s->SetBindingPoint(3, 4);
742 mod.root_block->Append(var_t);
743 mod.root_block->Append(var_s);
744
745 auto* foo = b.Function("foo", ty.vec4<f32>());
746 auto* param = b.FunctionParam<i32>("param");
747 foo->SetParams({param});
748 b.Append(foo->Block(), [&] {
749 auto* load_t = b.Load(var_t);
750 auto* load_s = b.Load(var_s);
751 auto* result = b.Call<vec4<f32>>(core::BuiltinFn::kTextureSample, load_t, load_s,
752 b.Splat<vec2<f32>>(0_f));
753 b.Return(foo, result);
754 });
755
756 auto* func = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kFragment);
757 b.Append(func->Block(), [&] {
758 b.Call(foo, 42_i);
759 b.Return(func);
760 });
761
762 auto* src = R"(
763$B1: { # root
764 %t:ptr<handle, texture_2d<f32>, read> = var @binding_point(1, 2)
765 %s:ptr<handle, sampler, read> = var @binding_point(3, 4)
766}
767
768%foo = func(%param:i32):vec4<f32> {
769 $B2: {
770 %5:texture_2d<f32> = load %t
771 %6:sampler = load %s
772 %7:vec4<f32> = textureSample %5, %6, vec2<f32>(0.0f)
773 ret %7
774 }
775}
776%main = @fragment func():void {
777 $B3: {
778 %9:vec4<f32> = call %foo, 42i
779 ret
780 }
781}
782)";
783 EXPECT_EQ(src, str());
784
785 auto* expect = R"(
786tint_module_vars_struct = struct @align(1) {
787 t:texture_2d<f32> @offset(0)
788 s:sampler @offset(0)
789}
790
791%foo = func(%param:i32, %tint_module_vars:tint_module_vars_struct):vec4<f32> {
792 $B1: {
793 %4:texture_2d<f32> = access %tint_module_vars, 0u
794 %5:sampler = access %tint_module_vars, 1u
795 %6:vec4<f32> = textureSample %4, %5, vec2<f32>(0.0f)
796 ret %6
797 }
798}
799%main = @fragment func(%t:texture_2d<f32> [@binding_point(1, 2)], %s:sampler [@binding_point(3, 4)]):void {
800 $B2: {
801 %10:tint_module_vars_struct = construct %t, %s
802 %tint_module_vars_1:tint_module_vars_struct = let %10 # %tint_module_vars_1: 'tint_module_vars'
803 %12:vec4<f32> = call %foo, 42i, %tint_module_vars_1
804 ret
805 }
806}
807)";
808
809 Run(ModuleScopeVars);
810
811 EXPECT_EQ(expect, str());
812}
813
James Price73f54952024-05-16 21:44:59 +0000814TEST_F(MslWriter_ModuleScopeVarsTest, CallFunctionThatUsesVars_OutOfOrder) {
815 auto* var_a = b.Var("a", ty.ptr<storage, i32, core::Access::kRead>());
816 auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>());
817 var_a->SetBindingPoint(1, 2);
818 var_b->SetBindingPoint(3, 4);
819 mod.root_block->Append(var_a);
820 mod.root_block->Append(var_b);
821
822 auto* func = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kFragment);
823
824 auto* foo = b.Function("foo", ty.void_());
825 b.Append(foo->Block(), [&] {
826 auto* load_a = b.Load(var_a);
827 auto* load_b = b.Load(var_b);
828 b.Store(var_b, b.Add<i32>(load_a, load_b));
829 b.Return(foo);
830 });
831
832 b.Append(func->Block(), [&] {
833 b.Call(foo);
834 b.Return(func);
835 });
836
837 auto* src = R"(
838$B1: { # root
839 %a:ptr<storage, i32, read> = var @binding_point(1, 2)
840 %b:ptr<storage, i32, read_write> = var @binding_point(3, 4)
841}
842
843%main = @fragment func():void {
844 $B2: {
845 %4:void = call %foo
846 ret
847 }
848}
849%foo = func():void {
850 $B3: {
851 %6:i32 = load %a
852 %7:i32 = load %b
853 %8:i32 = add %6, %7
854 store %b, %8
855 ret
856 }
857}
858)";
859 EXPECT_EQ(src, str());
860
861 auto* expect = R"(
862tint_module_vars_struct = struct @align(1) {
863 a:ptr<storage, i32, read> @offset(0)
864 b:ptr<storage, i32, read_write> @offset(0)
865}
866
867%main = @fragment func(%a:ptr<storage, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void {
868 $B1: {
869 %4:tint_module_vars_struct = construct %a, %b
870 %tint_module_vars:tint_module_vars_struct = let %4
871 %6:void = call %foo, %tint_module_vars
872 ret
873 }
874}
875%foo = func(%tint_module_vars_1:tint_module_vars_struct):void { # %tint_module_vars_1: 'tint_module_vars'
876 $B2: {
877 %9:ptr<storage, i32, read> = access %tint_module_vars_1, 0u
878 %10:i32 = load %9
879 %11:ptr<storage, i32, read_write> = access %tint_module_vars_1, 1u
880 %12:i32 = load %11
881 %13:i32 = add %10, %12
882 %14:ptr<storage, i32, read_write> = access %tint_module_vars_1, 1u
883 store %14, %13
884 ret
885 }
886}
887)";
888
889 Run(ModuleScopeVars);
890
891 EXPECT_EQ(expect, str());
892}
893
894// Test that we do not add the structure to functions that do not need it.
895TEST_F(MslWriter_ModuleScopeVarsTest, CallFunctionThatDoesNotUseVars) {
896 auto* var_a = b.Var("a", ty.ptr<storage, i32, core::Access::kRead>());
897 auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>());
898 var_a->SetBindingPoint(1, 2);
899 var_b->SetBindingPoint(3, 4);
900 mod.root_block->Append(var_a);
901 mod.root_block->Append(var_b);
902
903 auto* foo = b.Function("foo", ty.i32());
904 b.Append(foo->Block(), [&] { //
905 b.Return(foo, 42_i);
906 });
907
908 auto* func = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kFragment);
909 b.Append(func->Block(), [&] {
910 auto* load_a = b.Load(var_a);
911 auto* load_b = b.Load(var_b);
912 b.Store(var_b, b.Add<i32>(load_a, b.Add<i32>(load_b, b.Call(foo))));
913 b.Return(func);
914 });
915
916 auto* src = R"(
917$B1: { # root
918 %a:ptr<storage, i32, read> = var @binding_point(1, 2)
919 %b:ptr<storage, i32, read_write> = var @binding_point(3, 4)
920}
921
922%foo = func():i32 {
923 $B2: {
924 ret 42i
925 }
926}
927%main = @fragment func():void {
928 $B3: {
929 %5:i32 = load %a
930 %6:i32 = load %b
931 %7:i32 = call %foo
932 %8:i32 = add %6, %7
933 %9:i32 = add %5, %8
934 store %b, %9
935 ret
936 }
937}
938)";
939 EXPECT_EQ(src, str());
940
941 auto* expect = R"(
942tint_module_vars_struct = struct @align(1) {
943 a:ptr<storage, i32, read> @offset(0)
944 b:ptr<storage, i32, read_write> @offset(0)
945}
946
947%foo = func():i32 {
948 $B1: {
949 ret 42i
950 }
951}
952%main = @fragment func(%a:ptr<storage, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void {
953 $B2: {
954 %5:tint_module_vars_struct = construct %a, %b
955 %tint_module_vars:tint_module_vars_struct = let %5
956 %7:ptr<storage, i32, read> = access %tint_module_vars, 0u
957 %8:i32 = load %7
958 %9:ptr<storage, i32, read_write> = access %tint_module_vars, 1u
959 %10:i32 = load %9
960 %11:i32 = call %foo
961 %12:i32 = add %10, %11
962 %13:i32 = add %8, %12
963 %14:ptr<storage, i32, read_write> = access %tint_module_vars, 1u
964 store %14, %13
965 ret
966 }
967}
968)";
969
970 Run(ModuleScopeVars);
971
972 EXPECT_EQ(expect, str());
973}
974
975// Test that we *do* add the structure to functions that only have transitive uses.
976TEST_F(MslWriter_ModuleScopeVarsTest, CallFunctionWithOnlyTransitiveUses) {
977 auto* var_a = b.Var("a", ty.ptr<storage, i32, core::Access::kRead>());
978 auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>());
979 var_a->SetBindingPoint(1, 2);
980 var_b->SetBindingPoint(3, 4);
981 mod.root_block->Append(var_a);
982 mod.root_block->Append(var_b);
983
984 auto* bar = b.Function("bar", ty.i32());
985 b.Append(bar->Block(), [&] { //
986 b.Return(bar, b.Load(var_a));
987 });
988
989 auto* foo = b.Function("foo", ty.i32());
990 b.Append(foo->Block(), [&] { //
991 b.Return(foo, b.Call(bar));
992 });
993
994 auto* func = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kFragment);
995 b.Append(func->Block(), [&] {
996 auto* load_a = b.Load(var_a);
997 auto* load_b = b.Load(var_b);
998 b.Store(var_b, b.Add<i32>(load_a, b.Add<i32>(load_b, b.Call(foo))));
999 b.Return(func);
1000 });
1001
1002 auto* src = R"(
1003$B1: { # root
1004 %a:ptr<storage, i32, read> = var @binding_point(1, 2)
1005 %b:ptr<storage, i32, read_write> = var @binding_point(3, 4)
1006}
1007
1008%bar = func():i32 {
1009 $B2: {
1010 %4:i32 = load %a
1011 ret %4
1012 }
1013}
1014%foo = func():i32 {
1015 $B3: {
1016 %6:i32 = call %bar
1017 ret %6
1018 }
1019}
1020%main = @fragment func():void {
1021 $B4: {
1022 %8:i32 = load %a
1023 %9:i32 = load %b
1024 %10:i32 = call %foo
1025 %11:i32 = add %9, %10
1026 %12:i32 = add %8, %11
1027 store %b, %12
1028 ret
1029 }
1030}
1031)";
1032 EXPECT_EQ(src, str());
1033
1034 auto* expect = R"(
1035tint_module_vars_struct = struct @align(1) {
1036 a:ptr<storage, i32, read> @offset(0)
1037 b:ptr<storage, i32, read_write> @offset(0)
1038}
1039
1040%bar = func(%tint_module_vars:tint_module_vars_struct):i32 {
1041 $B1: {
1042 %3:ptr<storage, i32, read> = access %tint_module_vars, 0u
1043 %4:i32 = load %3
1044 ret %4
1045 }
1046}
1047%foo = func(%tint_module_vars_1:tint_module_vars_struct):i32 { # %tint_module_vars_1: 'tint_module_vars'
1048 $B2: {
1049 %7:i32 = call %bar, %tint_module_vars_1
1050 ret %7
1051 }
1052}
1053%main = @fragment func(%a:ptr<storage, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void {
1054 $B3: {
1055 %11:tint_module_vars_struct = construct %a, %b
1056 %tint_module_vars_2:tint_module_vars_struct = let %11 # %tint_module_vars_2: 'tint_module_vars'
1057 %13:ptr<storage, i32, read> = access %tint_module_vars_2, 0u
1058 %14:i32 = load %13
1059 %15:ptr<storage, i32, read_write> = access %tint_module_vars_2, 1u
1060 %16:i32 = load %15
1061 %17:i32 = call %foo, %tint_module_vars_2
1062 %18:i32 = add %16, %17
1063 %19:i32 = add %14, %18
1064 %20:ptr<storage, i32, read_write> = access %tint_module_vars_2, 1u
1065 store %20, %19
1066 ret
1067 }
1068}
1069)";
1070
1071 Run(ModuleScopeVars);
1072
1073 EXPECT_EQ(expect, str());
1074}
1075
1076// Test that we *do* add the structure to functions that only have transitive uses, where that
1077// function is declared first.
1078TEST_F(MslWriter_ModuleScopeVarsTest, CallFunctionWithOnlyTransitiveUses_OutOfOrder) {
1079 auto* var_a = b.Var("a", ty.ptr<storage, i32, core::Access::kRead>());
1080 auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>());
1081 var_a->SetBindingPoint(1, 2);
1082 var_b->SetBindingPoint(3, 4);
1083 mod.root_block->Append(var_a);
1084 mod.root_block->Append(var_b);
1085
1086 auto* foo = b.Function("foo", ty.i32());
1087
1088 auto* bar = b.Function("bar", ty.i32());
1089 b.Append(bar->Block(), [&] { //
1090 b.Return(bar, b.Load(var_a));
1091 });
1092
1093 b.Append(foo->Block(), [&] { //
1094 b.Return(foo, b.Call(bar));
1095 });
1096
1097 auto* func = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kFragment);
1098 b.Append(func->Block(), [&] {
1099 auto* load_a = b.Load(var_a);
1100 auto* load_b = b.Load(var_b);
1101 b.Store(var_b, b.Add<i32>(load_a, b.Add<i32>(load_b, b.Call(foo))));
1102 b.Return(func);
1103 });
1104
1105 auto* src = R"(
1106$B1: { # root
1107 %a:ptr<storage, i32, read> = var @binding_point(1, 2)
1108 %b:ptr<storage, i32, read_write> = var @binding_point(3, 4)
1109}
1110
1111%foo = func():i32 {
1112 $B2: {
1113 %4:i32 = call %bar
1114 ret %4
1115 }
1116}
1117%bar = func():i32 {
1118 $B3: {
1119 %6:i32 = load %a
1120 ret %6
1121 }
1122}
1123%main = @fragment func():void {
1124 $B4: {
1125 %8:i32 = load %a
1126 %9:i32 = load %b
1127 %10:i32 = call %foo
1128 %11:i32 = add %9, %10
1129 %12:i32 = add %8, %11
1130 store %b, %12
1131 ret
1132 }
1133}
1134)";
1135 EXPECT_EQ(src, str());
1136
1137 auto* expect = R"(
1138tint_module_vars_struct = struct @align(1) {
1139 a:ptr<storage, i32, read> @offset(0)
1140 b:ptr<storage, i32, read_write> @offset(0)
1141}
1142
1143%foo = func(%tint_module_vars:tint_module_vars_struct):i32 {
1144 $B1: {
1145 %3:i32 = call %bar, %tint_module_vars
1146 ret %3
1147 }
1148}
1149%bar = func(%tint_module_vars_1:tint_module_vars_struct):i32 { # %tint_module_vars_1: 'tint_module_vars'
1150 $B2: {
1151 %6:ptr<storage, i32, read> = access %tint_module_vars_1, 0u
1152 %7:i32 = load %6
1153 ret %7
1154 }
1155}
1156%main = @fragment func(%a:ptr<storage, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void {
1157 $B3: {
1158 %11:tint_module_vars_struct = construct %a, %b
1159 %tint_module_vars_2:tint_module_vars_struct = let %11 # %tint_module_vars_2: 'tint_module_vars'
1160 %13:ptr<storage, i32, read> = access %tint_module_vars_2, 0u
1161 %14:i32 = load %13
1162 %15:ptr<storage, i32, read_write> = access %tint_module_vars_2, 1u
1163 %16:i32 = load %15
1164 %17:i32 = call %foo, %tint_module_vars_2
1165 %18:i32 = add %16, %17
1166 %19:i32 = add %14, %18
1167 %20:ptr<storage, i32, read_write> = access %tint_module_vars_2, 1u
1168 store %20, %19
1169 ret
1170 }
1171}
1172)";
1173
1174 Run(ModuleScopeVars);
1175
1176 EXPECT_EQ(expect, str());
1177}
1178
1179TEST_F(MslWriter_ModuleScopeVarsTest, MultipleEntryPoints) {
1180 auto* var_a = b.Var("a", ty.ptr<uniform, i32, core::Access::kRead>());
1181 auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>());
1182 auto* var_c = b.Var("c", ty.ptr<private_, i32, core::Access::kReadWrite>());
1183 var_a->SetBindingPoint(1, 2);
1184 var_b->SetBindingPoint(3, 4);
1185 mod.root_block->Append(var_a);
1186 mod.root_block->Append(var_b);
1187 mod.root_block->Append(var_c);
1188
1189 auto* main_a = b.Function("main_a", ty.void_(), core::ir::Function::PipelineStage::kFragment);
1190 b.Append(main_a->Block(), [&] {
1191 auto* load_a = b.Load(var_a);
1192 auto* load_b = b.Load(var_b);
1193 auto* load_c = b.Load(var_c);
1194 b.Store(var_b, b.Add<i32>(load_a, b.Add<i32>(load_b, load_c)));
1195 b.Return(main_a);
1196 });
1197
1198 auto* main_b = b.Function("main_b", ty.void_(), core::ir::Function::PipelineStage::kFragment);
1199 b.Append(main_b->Block(), [&] {
1200 auto* load_a = b.Load(var_a);
1201 auto* load_b = b.Load(var_b);
1202 auto* load_c = b.Load(var_c);
1203 b.Store(var_b, b.Multiply<i32>(load_a, b.Multiply<i32>(load_b, load_c)));
1204 b.Return(main_b);
1205 });
1206
1207 auto* src = R"(
1208$B1: { # root
1209 %a:ptr<uniform, i32, read> = var @binding_point(1, 2)
1210 %b:ptr<storage, i32, read_write> = var @binding_point(3, 4)
1211 %c:ptr<private, i32, read_write> = var
1212}
1213
1214%main_a = @fragment func():void {
1215 $B2: {
1216 %5:i32 = load %a
1217 %6:i32 = load %b
1218 %7:i32 = load %c
1219 %8:i32 = add %6, %7
1220 %9:i32 = add %5, %8
1221 store %b, %9
1222 ret
1223 }
1224}
1225%main_b = @fragment func():void {
1226 $B3: {
1227 %11:i32 = load %a
1228 %12:i32 = load %b
1229 %13:i32 = load %c
1230 %14:i32 = mul %12, %13
1231 %15:i32 = mul %11, %14
1232 store %b, %15
1233 ret
1234 }
1235}
1236)";
1237 EXPECT_EQ(src, str());
1238
1239 auto* expect = R"(
1240tint_module_vars_struct = struct @align(1) {
1241 a:ptr<uniform, i32, read> @offset(0)
1242 b:ptr<storage, i32, read_write> @offset(0)
1243 c:ptr<private, i32, read_write> @offset(0)
1244}
1245
1246%main_a = @fragment func(%a:ptr<uniform, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void {
1247 $B1: {
1248 %c:ptr<private, i32, read_write> = var
1249 %5:tint_module_vars_struct = construct %a, %b, %c
1250 %tint_module_vars:tint_module_vars_struct = let %5
1251 %7:ptr<uniform, i32, read> = access %tint_module_vars, 0u
1252 %8:i32 = load %7
1253 %9:ptr<storage, i32, read_write> = access %tint_module_vars, 1u
1254 %10:i32 = load %9
1255 %11:ptr<private, i32, read_write> = access %tint_module_vars, 2u
1256 %12:i32 = load %11
1257 %13:i32 = add %10, %12
1258 %14:i32 = add %8, %13
1259 %15:ptr<storage, i32, read_write> = access %tint_module_vars, 1u
1260 store %15, %14
1261 ret
1262 }
1263}
1264%main_b = @fragment func(%a_1:ptr<uniform, i32, read> [@binding_point(1, 2)], %b_1:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void { # %a_1: 'a', %b_1: 'b'
1265 $B2: {
1266 %c_1:ptr<private, i32, read_write> = var # %c_1: 'c'
1267 %20:tint_module_vars_struct = construct %a_1, %b_1, %c_1
1268 %tint_module_vars_1:tint_module_vars_struct = let %20 # %tint_module_vars_1: 'tint_module_vars'
1269 %22:ptr<uniform, i32, read> = access %tint_module_vars_1, 0u
1270 %23:i32 = load %22
1271 %24:ptr<storage, i32, read_write> = access %tint_module_vars_1, 1u
1272 %25:i32 = load %24
1273 %26:ptr<private, i32, read_write> = access %tint_module_vars_1, 2u
1274 %27:i32 = load %26
1275 %28:i32 = mul %25, %27
1276 %29:i32 = mul %23, %28
1277 %30:ptr<storage, i32, read_write> = access %tint_module_vars_1, 1u
1278 store %30, %29
1279 ret
1280 }
1281}
1282)";
1283
1284 Run(ModuleScopeVars);
1285
1286 EXPECT_EQ(expect, str());
1287}
1288
1289TEST_F(MslWriter_ModuleScopeVarsTest, MultipleEntryPoints_DifferentUsageSets) {
1290 auto* var_a = b.Var("a", ty.ptr<uniform, i32, core::Access::kRead>());
1291 auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>());
1292 auto* var_c = b.Var("c", ty.ptr<private_, i32, core::Access::kReadWrite>());
1293 var_a->SetBindingPoint(1, 2);
1294 var_b->SetBindingPoint(3, 4);
1295 mod.root_block->Append(var_a);
1296 mod.root_block->Append(var_b);
1297 mod.root_block->Append(var_c);
1298
1299 auto* main_a = b.Function("main_a", ty.void_(), core::ir::Function::PipelineStage::kFragment);
1300 b.Append(main_a->Block(), [&] {
1301 auto* load_a = b.Load(var_a);
1302 auto* load_b = b.Load(var_b);
1303 b.Store(var_b, b.Add<i32>(load_a, load_b));
1304 b.Return(main_a);
1305 });
1306
1307 auto* main_b = b.Function("main_b", ty.void_(), core::ir::Function::PipelineStage::kFragment);
1308 b.Append(main_b->Block(), [&] {
1309 auto* load_a = b.Load(var_a);
1310 auto* load_c = b.Load(var_c);
1311 b.Store(var_c, b.Multiply<i32>(load_a, load_c));
1312 b.Return(main_b);
1313 });
1314
1315 auto* src = R"(
1316$B1: { # root
1317 %a:ptr<uniform, i32, read> = var @binding_point(1, 2)
1318 %b:ptr<storage, i32, read_write> = var @binding_point(3, 4)
1319 %c:ptr<private, i32, read_write> = var
1320}
1321
1322%main_a = @fragment func():void {
1323 $B2: {
1324 %5:i32 = load %a
1325 %6:i32 = load %b
1326 %7:i32 = add %5, %6
1327 store %b, %7
1328 ret
1329 }
1330}
1331%main_b = @fragment func():void {
1332 $B3: {
1333 %9:i32 = load %a
1334 %10:i32 = load %c
1335 %11:i32 = mul %9, %10
1336 store %c, %11
1337 ret
1338 }
1339}
1340)";
1341 EXPECT_EQ(src, str());
1342
1343 auto* expect = R"(
1344tint_module_vars_struct = struct @align(1) {
1345 a:ptr<uniform, i32, read> @offset(0)
1346 b:ptr<storage, i32, read_write> @offset(0)
1347 c:ptr<private, i32, read_write> @offset(0)
1348}
1349
1350%main_a = @fragment func(%a:ptr<uniform, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void {
1351 $B1: {
1352 %4:tint_module_vars_struct = construct %a, %b, undef
1353 %tint_module_vars:tint_module_vars_struct = let %4
1354 %6:ptr<uniform, i32, read> = access %tint_module_vars, 0u
1355 %7:i32 = load %6
1356 %8:ptr<storage, i32, read_write> = access %tint_module_vars, 1u
1357 %9:i32 = load %8
1358 %10:i32 = add %7, %9
1359 %11:ptr<storage, i32, read_write> = access %tint_module_vars, 1u
1360 store %11, %10
1361 ret
1362 }
1363}
1364%main_b = @fragment func(%a_1:ptr<uniform, i32, read> [@binding_point(1, 2)]):void { # %a_1: 'a'
1365 $B2: {
1366 %c:ptr<private, i32, read_write> = var
1367 %15:tint_module_vars_struct = construct %a_1, undef, %c
1368 %tint_module_vars_1:tint_module_vars_struct = let %15 # %tint_module_vars_1: 'tint_module_vars'
1369 %17:ptr<uniform, i32, read> = access %tint_module_vars_1, 0u
1370 %18:i32 = load %17
1371 %19:ptr<private, i32, read_write> = access %tint_module_vars_1, 2u
1372 %20:i32 = load %19
1373 %21:i32 = mul %18, %20
1374 %22:ptr<private, i32, read_write> = access %tint_module_vars_1, 2u
1375 store %22, %21
1376 ret
1377 }
1378}
1379)";
1380
1381 Run(ModuleScopeVars);
1382
1383 EXPECT_EQ(expect, str());
1384}
1385
1386TEST_F(MslWriter_ModuleScopeVarsTest, MultipleEntryPoints_DifferentUsageSets_CommonHelper) {
1387 auto* var_a = b.Var("a", ty.ptr<uniform, i32, core::Access::kRead>());
1388 auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>());
1389 auto* var_c = b.Var("c", ty.ptr<private_, i32, core::Access::kReadWrite>());
1390 var_a->SetBindingPoint(1, 2);
1391 var_b->SetBindingPoint(3, 4);
1392 mod.root_block->Append(var_a);
1393 mod.root_block->Append(var_b);
1394 mod.root_block->Append(var_c);
1395
1396 auto* foo = b.Function("foo", ty.i32());
1397 b.Append(foo->Block(), [&] { //
1398 b.Return(foo, b.Load(var_a));
1399 });
1400
1401 auto* main_a = b.Function("main_a", ty.void_(), core::ir::Function::PipelineStage::kFragment);
1402 b.Append(main_a->Block(), [&] {
1403 auto* load_b = b.Load(var_b);
1404 b.Store(var_b, b.Add<i32>(b.Call(foo), load_b));
1405 b.Return(main_a);
1406 });
1407
1408 auto* main_b = b.Function("main_b", ty.void_(), core::ir::Function::PipelineStage::kFragment);
1409 b.Append(main_b->Block(), [&] {
1410 auto* load_c = b.Load(var_c);
1411 b.Store(var_c, b.Multiply<i32>(b.Call(foo), load_c));
1412 b.Return(main_b);
1413 });
1414
1415 auto* src = R"(
1416$B1: { # root
1417 %a:ptr<uniform, i32, read> = var @binding_point(1, 2)
1418 %b:ptr<storage, i32, read_write> = var @binding_point(3, 4)
1419 %c:ptr<private, i32, read_write> = var
1420}
1421
1422%foo = func():i32 {
1423 $B2: {
1424 %5:i32 = load %a
1425 ret %5
1426 }
1427}
1428%main_a = @fragment func():void {
1429 $B3: {
1430 %7:i32 = load %b
1431 %8:i32 = call %foo
1432 %9:i32 = add %8, %7
1433 store %b, %9
1434 ret
1435 }
1436}
1437%main_b = @fragment func():void {
1438 $B4: {
1439 %11:i32 = load %c
1440 %12:i32 = call %foo
1441 %13:i32 = mul %12, %11
1442 store %c, %13
1443 ret
1444 }
1445}
1446)";
1447 EXPECT_EQ(src, str());
1448
1449 auto* expect = R"(
1450tint_module_vars_struct = struct @align(1) {
1451 a:ptr<uniform, i32, read> @offset(0)
1452 b:ptr<storage, i32, read_write> @offset(0)
1453 c:ptr<private, i32, read_write> @offset(0)
1454}
1455
1456%foo = func(%tint_module_vars:tint_module_vars_struct):i32 {
1457 $B1: {
1458 %3:ptr<uniform, i32, read> = access %tint_module_vars, 0u
1459 %4:i32 = load %3
1460 ret %4
1461 }
1462}
1463%main_a = @fragment func(%a:ptr<uniform, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void {
1464 $B2: {
1465 %8:tint_module_vars_struct = construct %a, %b, undef
1466 %tint_module_vars_1:tint_module_vars_struct = let %8 # %tint_module_vars_1: 'tint_module_vars'
1467 %10:ptr<storage, i32, read_write> = access %tint_module_vars_1, 1u
1468 %11:i32 = load %10
1469 %12:i32 = call %foo, %tint_module_vars_1
1470 %13:i32 = add %12, %11
1471 %14:ptr<storage, i32, read_write> = access %tint_module_vars_1, 1u
1472 store %14, %13
1473 ret
1474 }
1475}
1476%main_b = @fragment func(%a_1:ptr<uniform, i32, read> [@binding_point(1, 2)]):void { # %a_1: 'a'
1477 $B3: {
1478 %c:ptr<private, i32, read_write> = var
1479 %18:tint_module_vars_struct = construct %a_1, undef, %c
1480 %tint_module_vars_2:tint_module_vars_struct = let %18 # %tint_module_vars_2: 'tint_module_vars'
1481 %20:ptr<private, i32, read_write> = access %tint_module_vars_2, 2u
1482 %21:i32 = load %20
1483 %22:i32 = call %foo, %tint_module_vars_2
1484 %23:i32 = mul %22, %21
1485 %24:ptr<private, i32, read_write> = access %tint_module_vars_2, 2u
1486 store %24, %23
1487 ret
1488 }
1489}
1490)";
1491
1492 Run(ModuleScopeVars);
1493
1494 EXPECT_EQ(expect, str());
1495}
1496
1497TEST_F(MslWriter_ModuleScopeVarsTest, VarsWithNoNames) {
1498 auto* var_a = b.Var(ty.ptr<uniform, i32, core::Access::kRead>());
1499 auto* var_b = b.Var(ty.ptr<storage, i32, core::Access::kReadWrite>());
1500 auto* var_c = b.Var(ty.ptr<private_, i32, core::Access::kReadWrite>());
1501 var_a->SetBindingPoint(1, 2);
1502 var_b->SetBindingPoint(3, 4);
1503 mod.root_block->Append(var_a);
1504 mod.root_block->Append(var_b);
1505 mod.root_block->Append(var_c);
1506
1507 auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
1508 b.Append(func->Block(), [&] {
1509 auto* load_a = b.Load(var_a);
1510 auto* load_b = b.Load(var_b);
1511 auto* load_c = b.Load(var_c);
1512 b.Store(var_b, b.Add<i32>(load_a, b.Add<i32>(load_b, load_c)));
1513 b.Return(func);
1514 });
1515
1516 auto* src = R"(
1517$B1: { # root
1518 %1:ptr<uniform, i32, read> = var @binding_point(1, 2)
1519 %2:ptr<storage, i32, read_write> = var @binding_point(3, 4)
1520 %3:ptr<private, i32, read_write> = var
1521}
1522
1523%foo = @fragment func():void {
1524 $B2: {
1525 %5:i32 = load %1
1526 %6:i32 = load %2
1527 %7:i32 = load %3
1528 %8:i32 = add %6, %7
1529 %9:i32 = add %5, %8
1530 store %2, %9
1531 ret
1532 }
1533}
1534)";
1535 EXPECT_EQ(src, str());
1536
1537 auto* expect = R"(
1538tint_module_vars_struct = struct @align(1) {
1539 tint_symbol:ptr<uniform, i32, read> @offset(0)
1540 tint_symbol_1:ptr<storage, i32, read_write> @offset(0)
1541 tint_symbol_2:ptr<private, i32, read_write> @offset(0)
1542}
1543
1544%foo = @fragment func(%2:ptr<uniform, i32, read> [@binding_point(1, 2)], %3:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void {
1545 $B1: {
1546 %4:ptr<private, i32, read_write> = var
1547 %5:tint_module_vars_struct = construct %2, %3, %4
1548 %tint_module_vars:tint_module_vars_struct = let %5
1549 %7:ptr<uniform, i32, read> = access %tint_module_vars, 0u
1550 %8:i32 = load %7
1551 %9:ptr<storage, i32, read_write> = access %tint_module_vars, 1u
1552 %10:i32 = load %9
1553 %11:ptr<private, i32, read_write> = access %tint_module_vars, 2u
1554 %12:i32 = load %11
1555 %13:i32 = add %10, %12
1556 %14:i32 = add %8, %13
1557 %15:ptr<storage, i32, read_write> = access %tint_module_vars, 1u
1558 store %15, %14
1559 ret
1560 }
1561}
1562)";
1563
1564 Run(ModuleScopeVars);
1565
1566 EXPECT_EQ(expect, str());
1567}
1568
1569} // namespace
1570} // namespace tint::msl::writer::raise