1 // Copyright 2021 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/transform/decompose_strided_matrix.h"
16
17 #include <memory>
18 #include <utility>
19 #include <vector>
20
21 #include "src/ast/disable_validation_decoration.h"
22 #include "src/program_builder.h"
23 #include "src/transform/simplify_pointers.h"
24 #include "src/transform/test_helper.h"
25 #include "src/transform/unshadow.h"
26
27 namespace tint {
28 namespace transform {
29 namespace {
30
31 using DecomposeStridedMatrixTest = TransformTest;
32 using f32 = ProgramBuilder::f32;
33
TEST_F(DecomposeStridedMatrixTest,Empty)34 TEST_F(DecomposeStridedMatrixTest, Empty) {
35 auto* src = R"()";
36 auto* expect = src;
37
38 auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(src);
39
40 EXPECT_EQ(expect, str(got));
41 }
42
TEST_F(DecomposeStridedMatrixTest,MissingDependencySimplify)43 TEST_F(DecomposeStridedMatrixTest, MissingDependencySimplify) {
44 auto* src = R"()";
45 auto* expect =
46 R"(error: tint::transform::DecomposeStridedMatrix depends on tint::transform::SimplifyPointers but the dependency was not run)";
47
48 auto got = Run<DecomposeStridedMatrix>(src);
49
50 EXPECT_EQ(expect, str(got));
51 }
52
TEST_F(DecomposeStridedMatrixTest,ReadUniformMatrix)53 TEST_F(DecomposeStridedMatrixTest, ReadUniformMatrix) {
54 // [[block]]
55 // struct S {
56 // [[offset(16), stride(32)]]
57 // [[internal(ignore_stride_decoration)]]
58 // m : mat2x2<f32>;
59 // };
60 // [[group(0), binding(0)]] var<uniform> s : S;
61 //
62 // [[stage(compute), workgroup_size(1)]]
63 // fn f() {
64 // let x : mat2x2<f32> = s.m;
65 // }
66 ProgramBuilder b;
67 auto* S = b.Structure(
68 "S",
69 {
70 b.Member(
71 "m", b.ty.mat2x2<f32>(),
72 {
73 b.create<ast::StructMemberOffsetDecoration>(16),
74 b.create<ast::StrideDecoration>(32),
75 b.Disable(ast::DisabledValidation::kIgnoreStrideDecoration),
76 }),
77 },
78 {
79 b.StructBlock(),
80 });
81 b.Global("s", b.ty.Of(S), ast::StorageClass::kUniform,
82 b.GroupAndBinding(0, 0));
83 b.Func(
84 "f", {}, b.ty.void_(),
85 {
86 b.Decl(b.Const("x", b.ty.mat2x2<f32>(), b.MemberAccessor("s", "m"))),
87 },
88 {
89 b.Stage(ast::PipelineStage::kCompute),
90 b.WorkgroupSize(1),
91 });
92
93 auto* expect = R"(
94 [[block]]
95 struct S {
96 [[size(16)]]
97 padding : u32;
98 m : [[stride(32)]] array<vec2<f32>, 2u>;
99 };
100
101 [[group(0), binding(0)]] var<uniform> s : S;
102
103 fn arr_to_mat2x2_stride_32(arr : [[stride(32)]] array<vec2<f32>, 2u>) -> mat2x2<f32> {
104 return mat2x2<f32>(arr[0u], arr[1u]);
105 }
106
107 [[stage(compute), workgroup_size(1)]]
108 fn f() {
109 let x : mat2x2<f32> = arr_to_mat2x2_stride_32(s.m);
110 }
111 )";
112
113 auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
114 Program(std::move(b)));
115
116 EXPECT_EQ(expect, str(got));
117 }
118
TEST_F(DecomposeStridedMatrixTest,ReadUniformColumn)119 TEST_F(DecomposeStridedMatrixTest, ReadUniformColumn) {
120 // [[block]]
121 // struct S {
122 // [[offset(16), stride(32)]]
123 // [[internal(ignore_stride_decoration)]]
124 // m : mat2x2<f32>;
125 // };
126 // [[group(0), binding(0)]] var<uniform> s : S;
127 //
128 // [[stage(compute), workgroup_size(1)]]
129 // fn f() {
130 // let x : vec2<f32> = s.m[1];
131 // }
132 ProgramBuilder b;
133 auto* S = b.Structure(
134 "S",
135 {
136 b.Member(
137 "m", b.ty.mat2x2<f32>(),
138 {
139 b.create<ast::StructMemberOffsetDecoration>(16),
140 b.create<ast::StrideDecoration>(32),
141 b.Disable(ast::DisabledValidation::kIgnoreStrideDecoration),
142 }),
143 },
144 {
145 b.StructBlock(),
146 });
147 b.Global("s", b.ty.Of(S), ast::StorageClass::kUniform,
148 b.GroupAndBinding(0, 0));
149 b.Func("f", {}, b.ty.void_(),
150 {
151 b.Decl(b.Const("x", b.ty.vec2<f32>(),
152 b.IndexAccessor(b.MemberAccessor("s", "m"), 1))),
153 },
154 {
155 b.Stage(ast::PipelineStage::kCompute),
156 b.WorkgroupSize(1),
157 });
158
159 auto* expect = R"(
160 [[block]]
161 struct S {
162 [[size(16)]]
163 padding : u32;
164 m : [[stride(32)]] array<vec2<f32>, 2u>;
165 };
166
167 [[group(0), binding(0)]] var<uniform> s : S;
168
169 [[stage(compute), workgroup_size(1)]]
170 fn f() {
171 let x : vec2<f32> = s.m[1];
172 }
173 )";
174
175 auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
176 Program(std::move(b)));
177
178 EXPECT_EQ(expect, str(got));
179 }
180
TEST_F(DecomposeStridedMatrixTest,ReadUniformMatrix_DefaultStride)181 TEST_F(DecomposeStridedMatrixTest, ReadUniformMatrix_DefaultStride) {
182 // [[block]]
183 // struct S {
184 // [[offset(16), stride(8)]]
185 // [[internal(ignore_stride_decoration)]]
186 // m : mat2x2<f32>;
187 // };
188 // [[group(0), binding(0)]] var<uniform> s : S;
189 //
190 // [[stage(compute), workgroup_size(1)]]
191 // fn f() {
192 // let x : mat2x2<f32> = s.m;
193 // }
194 ProgramBuilder b;
195 auto* S = b.Structure(
196 "S",
197 {
198 b.Member(
199 "m", b.ty.mat2x2<f32>(),
200 {
201 b.create<ast::StructMemberOffsetDecoration>(16),
202 b.create<ast::StrideDecoration>(8),
203 b.Disable(ast::DisabledValidation::kIgnoreStrideDecoration),
204 }),
205 },
206 {
207 b.StructBlock(),
208 });
209 b.Global("s", b.ty.Of(S), ast::StorageClass::kUniform,
210 b.GroupAndBinding(0, 0));
211 b.Func(
212 "f", {}, b.ty.void_(),
213 {
214 b.Decl(b.Const("x", b.ty.mat2x2<f32>(), b.MemberAccessor("s", "m"))),
215 },
216 {
217 b.Stage(ast::PipelineStage::kCompute),
218 b.WorkgroupSize(1),
219 });
220
221 auto* expect = R"(
222 [[block]]
223 struct S {
224 [[size(16)]]
225 padding : u32;
226 [[stride(8), internal(disable_validation__ignore_stride)]]
227 m : mat2x2<f32>;
228 };
229
230 [[group(0), binding(0)]] var<uniform> s : S;
231
232 [[stage(compute), workgroup_size(1)]]
233 fn f() {
234 let x : mat2x2<f32> = s.m;
235 }
236 )";
237
238 auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
239 Program(std::move(b)));
240
241 EXPECT_EQ(expect, str(got));
242 }
243
TEST_F(DecomposeStridedMatrixTest,ReadStorageMatrix)244 TEST_F(DecomposeStridedMatrixTest, ReadStorageMatrix) {
245 // [[block]]
246 // struct S {
247 // [[offset(8), stride(32)]]
248 // [[internal(ignore_stride_decoration)]]
249 // m : mat2x2<f32>;
250 // };
251 // [[group(0), binding(0)]] var<storage, read_write> s : S;
252 //
253 // [[stage(compute), workgroup_size(1)]]
254 // fn f() {
255 // let x : mat2x2<f32> = s.m;
256 // }
257 ProgramBuilder b;
258 auto* S = b.Structure(
259 "S",
260 {
261 b.Member(
262 "m", b.ty.mat2x2<f32>(),
263 {
264 b.create<ast::StructMemberOffsetDecoration>(8),
265 b.create<ast::StrideDecoration>(32),
266 b.Disable(ast::DisabledValidation::kIgnoreStrideDecoration),
267 }),
268 },
269 {
270 b.StructBlock(),
271 });
272 b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage,
273 ast::Access::kReadWrite, b.GroupAndBinding(0, 0));
274 b.Func(
275 "f", {}, b.ty.void_(),
276 {
277 b.Decl(b.Const("x", b.ty.mat2x2<f32>(), b.MemberAccessor("s", "m"))),
278 },
279 {
280 b.Stage(ast::PipelineStage::kCompute),
281 b.WorkgroupSize(1),
282 });
283
284 auto* expect = R"(
285 [[block]]
286 struct S {
287 [[size(8)]]
288 padding : u32;
289 m : [[stride(32)]] array<vec2<f32>, 2u>;
290 };
291
292 [[group(0), binding(0)]] var<storage, read_write> s : S;
293
294 fn arr_to_mat2x2_stride_32(arr : [[stride(32)]] array<vec2<f32>, 2u>) -> mat2x2<f32> {
295 return mat2x2<f32>(arr[0u], arr[1u]);
296 }
297
298 [[stage(compute), workgroup_size(1)]]
299 fn f() {
300 let x : mat2x2<f32> = arr_to_mat2x2_stride_32(s.m);
301 }
302 )";
303
304 auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
305 Program(std::move(b)));
306
307 EXPECT_EQ(expect, str(got));
308 }
309
TEST_F(DecomposeStridedMatrixTest,ReadStorageColumn)310 TEST_F(DecomposeStridedMatrixTest, ReadStorageColumn) {
311 // [[block]]
312 // struct S {
313 // [[offset(16), stride(32)]]
314 // [[internal(ignore_stride_decoration)]]
315 // m : mat2x2<f32>;
316 // };
317 // [[group(0), binding(0)]] var<storage, read_write> s : S;
318 //
319 // [[stage(compute), workgroup_size(1)]]
320 // fn f() {
321 // let x : vec2<f32> = s.m[1];
322 // }
323 ProgramBuilder b;
324 auto* S = b.Structure(
325 "S",
326 {
327 b.Member(
328 "m", b.ty.mat2x2<f32>(),
329 {
330 b.create<ast::StructMemberOffsetDecoration>(16),
331 b.create<ast::StrideDecoration>(32),
332 b.Disable(ast::DisabledValidation::kIgnoreStrideDecoration),
333 }),
334 },
335 {
336 b.StructBlock(),
337 });
338 b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage,
339 ast::Access::kReadWrite, b.GroupAndBinding(0, 0));
340 b.Func("f", {}, b.ty.void_(),
341 {
342 b.Decl(b.Const("x", b.ty.vec2<f32>(),
343 b.IndexAccessor(b.MemberAccessor("s", "m"), 1))),
344 },
345 {
346 b.Stage(ast::PipelineStage::kCompute),
347 b.WorkgroupSize(1),
348 });
349
350 auto* expect = R"(
351 [[block]]
352 struct S {
353 [[size(16)]]
354 padding : u32;
355 m : [[stride(32)]] array<vec2<f32>, 2u>;
356 };
357
358 [[group(0), binding(0)]] var<storage, read_write> s : S;
359
360 [[stage(compute), workgroup_size(1)]]
361 fn f() {
362 let x : vec2<f32> = s.m[1];
363 }
364 )";
365
366 auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
367 Program(std::move(b)));
368
369 EXPECT_EQ(expect, str(got));
370 }
371
TEST_F(DecomposeStridedMatrixTest,WriteStorageMatrix)372 TEST_F(DecomposeStridedMatrixTest, WriteStorageMatrix) {
373 // [[block]]
374 // struct S {
375 // [[offset(8), stride(32)]]
376 // [[internal(ignore_stride_decoration)]]
377 // m : mat2x2<f32>;
378 // };
379 // [[group(0), binding(0)]] var<storage, read_write> s : S;
380 //
381 // [[stage(compute), workgroup_size(1)]]
382 // fn f() {
383 // s.m = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
384 // }
385 ProgramBuilder b;
386 auto* S = b.Structure(
387 "S",
388 {
389 b.Member(
390 "m", b.ty.mat2x2<f32>(),
391 {
392 b.create<ast::StructMemberOffsetDecoration>(8),
393 b.create<ast::StrideDecoration>(32),
394 b.Disable(ast::DisabledValidation::kIgnoreStrideDecoration),
395 }),
396 },
397 {
398 b.StructBlock(),
399 });
400 b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage,
401 ast::Access::kReadWrite, b.GroupAndBinding(0, 0));
402 b.Func("f", {}, b.ty.void_(),
403 {
404 b.Assign(b.MemberAccessor("s", "m"),
405 b.mat2x2<f32>(b.vec2<f32>(1.0f, 2.0f),
406 b.vec2<f32>(3.0f, 4.0f))),
407 },
408 {
409 b.Stage(ast::PipelineStage::kCompute),
410 b.WorkgroupSize(1),
411 });
412
413 auto* expect = R"(
414 [[block]]
415 struct S {
416 [[size(8)]]
417 padding : u32;
418 m : [[stride(32)]] array<vec2<f32>, 2u>;
419 };
420
421 [[group(0), binding(0)]] var<storage, read_write> s : S;
422
423 fn mat2x2_stride_32_to_arr(mat : mat2x2<f32>) -> [[stride(32)]] array<vec2<f32>, 2u> {
424 return [[stride(32)]] array<vec2<f32>, 2u>(mat[0u], mat[1u]);
425 }
426
427 [[stage(compute), workgroup_size(1)]]
428 fn f() {
429 s.m = mat2x2_stride_32_to_arr(mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0)));
430 }
431 )";
432
433 auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
434 Program(std::move(b)));
435
436 EXPECT_EQ(expect, str(got));
437 }
438
TEST_F(DecomposeStridedMatrixTest,WriteStorageColumn)439 TEST_F(DecomposeStridedMatrixTest, WriteStorageColumn) {
440 // [[block]]
441 // struct S {
442 // [[offset(8), stride(32)]]
443 // [[internal(ignore_stride_decoration)]]
444 // m : mat2x2<f32>;
445 // };
446 // [[group(0), binding(0)]] var<storage, read_write> s : S;
447 //
448 // [[stage(compute), workgroup_size(1)]]
449 // fn f() {
450 // s.m[1] = vec2<f32>(1.0, 2.0);
451 // }
452 ProgramBuilder b;
453 auto* S = b.Structure(
454 "S",
455 {
456 b.Member(
457 "m", b.ty.mat2x2<f32>(),
458 {
459 b.create<ast::StructMemberOffsetDecoration>(8),
460 b.create<ast::StrideDecoration>(32),
461 b.Disable(ast::DisabledValidation::kIgnoreStrideDecoration),
462 }),
463 },
464 {
465 b.StructBlock(),
466 });
467 b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage,
468 ast::Access::kReadWrite, b.GroupAndBinding(0, 0));
469 b.Func("f", {}, b.ty.void_(),
470 {
471 b.Assign(b.IndexAccessor(b.MemberAccessor("s", "m"), 1),
472 b.vec2<f32>(1.0f, 2.0f)),
473 },
474 {
475 b.Stage(ast::PipelineStage::kCompute),
476 b.WorkgroupSize(1),
477 });
478
479 auto* expect = R"(
480 [[block]]
481 struct S {
482 [[size(8)]]
483 padding : u32;
484 m : [[stride(32)]] array<vec2<f32>, 2u>;
485 };
486
487 [[group(0), binding(0)]] var<storage, read_write> s : S;
488
489 [[stage(compute), workgroup_size(1)]]
490 fn f() {
491 s.m[1] = vec2<f32>(1.0, 2.0);
492 }
493 )";
494
495 auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
496 Program(std::move(b)));
497
498 EXPECT_EQ(expect, str(got));
499 }
500
TEST_F(DecomposeStridedMatrixTest,ReadWriteViaPointerLets)501 TEST_F(DecomposeStridedMatrixTest, ReadWriteViaPointerLets) {
502 // [[block]]
503 // struct S {
504 // [[offset(8), stride(32)]]
505 // [[internal(ignore_stride_decoration)]]
506 // m : mat2x2<f32>;
507 // };
508 // [[group(0), binding(0)]] var<storage, read_write> s : S;
509 //
510 // [[stage(compute), workgroup_size(1)]]
511 // fn f() {
512 // let a = &s.m;
513 // let b = &*&*(a);
514 // let x = *b;
515 // let y = (*b)[1];
516 // let z = x[1];
517 // (*b) = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
518 // (*b)[1] = vec2<f32>(5.0, 6.0);
519 // }
520 ProgramBuilder b;
521 auto* S = b.Structure(
522 "S",
523 {
524 b.Member(
525 "m", b.ty.mat2x2<f32>(),
526 {
527 b.create<ast::StructMemberOffsetDecoration>(8),
528 b.create<ast::StrideDecoration>(32),
529 b.Disable(ast::DisabledValidation::kIgnoreStrideDecoration),
530 }),
531 },
532 {
533 b.StructBlock(),
534 });
535 b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage,
536 ast::Access::kReadWrite, b.GroupAndBinding(0, 0));
537 b.Func(
538 "f", {}, b.ty.void_(),
539 {
540 b.Decl(
541 b.Const("a", nullptr, b.AddressOf(b.MemberAccessor("s", "m")))),
542 b.Decl(b.Const("b", nullptr,
543 b.AddressOf(b.Deref(b.AddressOf(b.Deref("a")))))),
544 b.Decl(b.Const("x", nullptr, b.Deref("b"))),
545 b.Decl(b.Const("y", nullptr, b.IndexAccessor(b.Deref("b"), 1))),
546 b.Decl(b.Const("z", nullptr, b.IndexAccessor("x", 1))),
547 b.Assign(b.Deref("b"), b.mat2x2<f32>(b.vec2<f32>(1.0f, 2.0f),
548 b.vec2<f32>(3.0f, 4.0f))),
549 b.Assign(b.IndexAccessor(b.Deref("b"), 1), b.vec2<f32>(5.0f, 6.0f)),
550 },
551 {
552 b.Stage(ast::PipelineStage::kCompute),
553 b.WorkgroupSize(1),
554 });
555
556 auto* expect = R"(
557 [[block]]
558 struct S {
559 [[size(8)]]
560 padding : u32;
561 m : [[stride(32)]] array<vec2<f32>, 2u>;
562 };
563
564 [[group(0), binding(0)]] var<storage, read_write> s : S;
565
566 fn arr_to_mat2x2_stride_32(arr : [[stride(32)]] array<vec2<f32>, 2u>) -> mat2x2<f32> {
567 return mat2x2<f32>(arr[0u], arr[1u]);
568 }
569
570 fn mat2x2_stride_32_to_arr(mat : mat2x2<f32>) -> [[stride(32)]] array<vec2<f32>, 2u> {
571 return [[stride(32)]] array<vec2<f32>, 2u>(mat[0u], mat[1u]);
572 }
573
574 [[stage(compute), workgroup_size(1)]]
575 fn f() {
576 let x = arr_to_mat2x2_stride_32(s.m);
577 let y = s.m[1];
578 let z = x[1];
579 s.m = mat2x2_stride_32_to_arr(mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0)));
580 s.m[1] = vec2<f32>(5.0, 6.0);
581 }
582 )";
583
584 auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
585 Program(std::move(b)));
586
587 EXPECT_EQ(expect, str(got));
588 }
589
TEST_F(DecomposeStridedMatrixTest,ReadPrivateMatrix)590 TEST_F(DecomposeStridedMatrixTest, ReadPrivateMatrix) {
591 // struct S {
592 // [[offset(8), stride(32)]]
593 // [[internal(ignore_stride_decoration)]]
594 // m : mat2x2<f32>;
595 // };
596 // var<private> s : S;
597 //
598 // [[stage(compute), workgroup_size(1)]]
599 // fn f() {
600 // let x : mat2x2<f32> = s.m;
601 // }
602 ProgramBuilder b;
603 auto* S = b.Structure(
604 "S",
605 {
606 b.Member(
607 "m", b.ty.mat2x2<f32>(),
608 {
609 b.create<ast::StructMemberOffsetDecoration>(8),
610 b.create<ast::StrideDecoration>(32),
611 b.Disable(ast::DisabledValidation::kIgnoreStrideDecoration),
612 }),
613 });
614 b.Global("s", b.ty.Of(S), ast::StorageClass::kPrivate);
615 b.Func(
616 "f", {}, b.ty.void_(),
617 {
618 b.Decl(b.Const("x", b.ty.mat2x2<f32>(), b.MemberAccessor("s", "m"))),
619 },
620 {
621 b.Stage(ast::PipelineStage::kCompute),
622 b.WorkgroupSize(1),
623 });
624
625 auto* expect = R"(
626 struct S {
627 [[size(8)]]
628 padding : u32;
629 [[stride(32), internal(disable_validation__ignore_stride)]]
630 m : mat2x2<f32>;
631 };
632
633 var<private> s : S;
634
635 [[stage(compute), workgroup_size(1)]]
636 fn f() {
637 let x : mat2x2<f32> = s.m;
638 }
639 )";
640
641 auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
642 Program(std::move(b)));
643
644 EXPECT_EQ(expect, str(got));
645 }
646
TEST_F(DecomposeStridedMatrixTest,WritePrivateMatrix)647 TEST_F(DecomposeStridedMatrixTest, WritePrivateMatrix) {
648 // struct S {
649 // [[offset(8), stride(32)]]
650 // [[internal(ignore_stride_decoration)]]
651 // m : mat2x2<f32>;
652 // };
653 // var<private> s : S;
654 //
655 // [[stage(compute), workgroup_size(1)]]
656 // fn f() {
657 // s.m = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
658 // }
659 ProgramBuilder b;
660 auto* S = b.Structure(
661 "S",
662 {
663 b.Member(
664 "m", b.ty.mat2x2<f32>(),
665 {
666 b.create<ast::StructMemberOffsetDecoration>(8),
667 b.create<ast::StrideDecoration>(32),
668 b.Disable(ast::DisabledValidation::kIgnoreStrideDecoration),
669 }),
670 });
671 b.Global("s", b.ty.Of(S), ast::StorageClass::kPrivate);
672 b.Func("f", {}, b.ty.void_(),
673 {
674 b.Assign(b.MemberAccessor("s", "m"),
675 b.mat2x2<f32>(b.vec2<f32>(1.0f, 2.0f),
676 b.vec2<f32>(3.0f, 4.0f))),
677 },
678 {
679 b.Stage(ast::PipelineStage::kCompute),
680 b.WorkgroupSize(1),
681 });
682
683 auto* expect = R"(
684 struct S {
685 [[size(8)]]
686 padding : u32;
687 [[stride(32), internal(disable_validation__ignore_stride)]]
688 m : mat2x2<f32>;
689 };
690
691 var<private> s : S;
692
693 [[stage(compute), workgroup_size(1)]]
694 fn f() {
695 s.m = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
696 }
697 )";
698
699 auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
700 Program(std::move(b)));
701
702 EXPECT_EQ(expect, str(got));
703 }
704
705 } // namespace
706 } // namespace transform
707 } // namespace tint
708