• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
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