|  | // Copyright 2022 The Tint Authors. | 
|  | // | 
|  | // Licensed under the Apache License, Version 2.0 (the "License"); | 
|  | // you may not use this file except in compliance with the License. | 
|  | // You may obtain a copy of the License at | 
|  | // | 
|  | //     http://www.apache.org/licenses/LICENSE-2.0 | 
|  | // | 
|  | // Unless required by applicable law or agreed to in writing, software | 
|  | // distributed under the License is distributed on an "AS IS" BASIS, | 
|  | // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | 
|  | // See the License for the specific language governing permissions and | 
|  | // limitations under the License. | 
|  |  | 
|  | #include "src/tint/transform/decompose_strided_array.h" | 
|  |  | 
|  | #include <memory> | 
|  | #include <utility> | 
|  | #include <vector> | 
|  |  | 
|  | #include "src/tint/program_builder.h" | 
|  | #include "src/tint/transform/simplify_pointers.h" | 
|  | #include "src/tint/transform/test_helper.h" | 
|  | #include "src/tint/transform/unshadow.h" | 
|  |  | 
|  | using namespace tint::number_suffixes;  // NOLINT | 
|  |  | 
|  | namespace tint::transform { | 
|  | namespace { | 
|  |  | 
|  | using DecomposeStridedArrayTest = TransformTest; | 
|  |  | 
|  | TEST_F(DecomposeStridedArrayTest, ShouldRunEmptyModule) { | 
|  | ProgramBuilder b; | 
|  | EXPECT_FALSE(ShouldRun<DecomposeStridedArray>(Program(std::move(b)))); | 
|  | } | 
|  |  | 
|  | TEST_F(DecomposeStridedArrayTest, ShouldRunNonStridedArray) { | 
|  | // var<private> arr : array<f32, 4u> | 
|  |  | 
|  | ProgramBuilder b; | 
|  | b.GlobalVar("arr", b.ty.array<f32, 4u>(), ast::AddressSpace::kPrivate); | 
|  | EXPECT_FALSE(ShouldRun<DecomposeStridedArray>(Program(std::move(b)))); | 
|  | } | 
|  |  | 
|  | TEST_F(DecomposeStridedArrayTest, ShouldRunDefaultStridedArray) { | 
|  | // var<private> arr : @stride(4) array<f32, 4u> | 
|  |  | 
|  | ProgramBuilder b; | 
|  | b.GlobalVar("arr", b.ty.array<f32, 4u>(4), ast::AddressSpace::kPrivate); | 
|  | EXPECT_TRUE(ShouldRun<DecomposeStridedArray>(Program(std::move(b)))); | 
|  | } | 
|  |  | 
|  | TEST_F(DecomposeStridedArrayTest, ShouldRunExplicitStridedArray) { | 
|  | // var<private> arr : @stride(16) array<f32, 4u> | 
|  |  | 
|  | ProgramBuilder b; | 
|  | b.GlobalVar("arr", b.ty.array<f32, 4u>(16), ast::AddressSpace::kPrivate); | 
|  | EXPECT_TRUE(ShouldRun<DecomposeStridedArray>(Program(std::move(b)))); | 
|  | } | 
|  |  | 
|  | TEST_F(DecomposeStridedArrayTest, Empty) { | 
|  | auto* src = R"()"; | 
|  | auto* expect = src; | 
|  |  | 
|  | auto got = Run<DecomposeStridedArray>(src); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(DecomposeStridedArrayTest, PrivateDefaultStridedArray) { | 
|  | // var<private> arr : @stride(4) array<f32, 4u> | 
|  | // | 
|  | // @compute @workgroup_size(1) | 
|  | // fn f() { | 
|  | //   let a : @stride(4) array<f32, 4u> = a; | 
|  | //   let b : f32 = arr[1]; | 
|  | // } | 
|  |  | 
|  | ProgramBuilder b; | 
|  | b.GlobalVar("arr", b.ty.array<f32, 4u>(4), ast::AddressSpace::kPrivate); | 
|  | b.Func("f", utils::Empty, b.ty.void_(), | 
|  | utils::Vector{ | 
|  | b.Decl(b.Let("a", b.ty.array<f32, 4u>(4), b.Expr("arr"))), | 
|  | b.Decl(b.Let("b", b.ty.f32(), b.IndexAccessor("arr", 1_i))), | 
|  | }, | 
|  | utils::Vector{ | 
|  | b.Stage(ast::PipelineStage::kCompute), | 
|  | b.WorkgroupSize(1_i), | 
|  | }); | 
|  |  | 
|  | auto* expect = R"( | 
|  | var<private> arr : array<f32, 4u>; | 
|  |  | 
|  | @compute @workgroup_size(1i) | 
|  | fn f() { | 
|  | let a : array<f32, 4u> = arr; | 
|  | let b : f32 = arr[1i]; | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedArray>(Program(std::move(b))); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(DecomposeStridedArrayTest, PrivateStridedArray) { | 
|  | // var<private> arr : @stride(32) array<f32, 4u> | 
|  | // | 
|  | // @compute @workgroup_size(1) | 
|  | // fn f() { | 
|  | //   let a : @stride(32) array<f32, 4u> = a; | 
|  | //   let b : f32 = arr[1]; | 
|  | // } | 
|  |  | 
|  | ProgramBuilder b; | 
|  | b.GlobalVar("arr", b.ty.array<f32, 4u>(32), ast::AddressSpace::kPrivate); | 
|  | b.Func("f", utils::Empty, b.ty.void_(), | 
|  | utils::Vector{ | 
|  | b.Decl(b.Let("a", b.ty.array<f32, 4u>(32), b.Expr("arr"))), | 
|  | b.Decl(b.Let("b", b.ty.f32(), b.IndexAccessor("arr", 1_i))), | 
|  | }, | 
|  | utils::Vector{ | 
|  | b.Stage(ast::PipelineStage::kCompute), | 
|  | b.WorkgroupSize(1_i), | 
|  | }); | 
|  |  | 
|  | auto* expect = R"( | 
|  | struct strided_arr { | 
|  | @size(32) | 
|  | el : f32, | 
|  | } | 
|  |  | 
|  | var<private> arr : array<strided_arr, 4u>; | 
|  |  | 
|  | @compute @workgroup_size(1i) | 
|  | fn f() { | 
|  | let a : array<strided_arr, 4u> = arr; | 
|  | let b : f32 = arr[1i].el; | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedArray>(Program(std::move(b))); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(DecomposeStridedArrayTest, ReadUniformStridedArray) { | 
|  | // struct S { | 
|  | //   a : @stride(32) array<f32, 4u>, | 
|  | // }; | 
|  | // @group(0) @binding(0) var<uniform> s : S; | 
|  | // | 
|  | // @compute @workgroup_size(1) | 
|  | // fn f() { | 
|  | //   let a : @stride(32) array<f32, 4u> = s.a; | 
|  | //   let b : f32 = s.a[1]; | 
|  | // } | 
|  | ProgramBuilder b; | 
|  | auto* S = b.Structure("S", utils::Vector{b.Member("a", b.ty.array<f32, 4u>(32))}); | 
|  | b.GlobalVar("s", b.ty.Of(S), ast::AddressSpace::kUniform, b.Group(0_a), b.Binding(0_a)); | 
|  | b.Func("f", utils::Empty, b.ty.void_(), | 
|  | utils::Vector{ | 
|  | b.Decl(b.Let("a", b.ty.array<f32, 4u>(32), b.MemberAccessor("s", "a"))), | 
|  | b.Decl(b.Let("b", b.ty.f32(), b.IndexAccessor(b.MemberAccessor("s", "a"), 1_i))), | 
|  | }, | 
|  | utils::Vector{ | 
|  | b.Stage(ast::PipelineStage::kCompute), | 
|  | b.WorkgroupSize(1_i), | 
|  | }); | 
|  |  | 
|  | auto* expect = R"( | 
|  | struct strided_arr { | 
|  | @size(32) | 
|  | el : f32, | 
|  | } | 
|  |  | 
|  | struct S { | 
|  | a : array<strided_arr, 4u>, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<uniform> s : S; | 
|  |  | 
|  | @compute @workgroup_size(1i) | 
|  | fn f() { | 
|  | let a : array<strided_arr, 4u> = s.a; | 
|  | let b : f32 = s.a[1i].el; | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedArray>(Program(std::move(b))); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(DecomposeStridedArrayTest, ReadUniformDefaultStridedArray) { | 
|  | // struct S { | 
|  | //   a : @stride(16) array<vec4<f32>, 4u>, | 
|  | // }; | 
|  | // @group(0) @binding(0) var<uniform> s : S; | 
|  | // | 
|  | // @compute @workgroup_size(1) | 
|  | // fn f() { | 
|  | //   let a : @stride(16) array<vec4<f32>, 4u> = s.a; | 
|  | //   let b : f32 = s.a[1][2]; | 
|  | // } | 
|  | ProgramBuilder b; | 
|  | auto* S = b.Structure("S", utils::Vector{b.Member("a", b.ty.array(b.ty.vec4<f32>(), 4_u, 16))}); | 
|  | b.GlobalVar("s", b.ty.Of(S), ast::AddressSpace::kUniform, b.Group(0_a), b.Binding(0_a)); | 
|  | b.Func( | 
|  | "f", utils::Empty, b.ty.void_(), | 
|  | utils::Vector{ | 
|  | b.Decl(b.Let("a", b.ty.array(b.ty.vec4<f32>(), 4_u, 16), b.MemberAccessor("s", "a"))), | 
|  | b.Decl(b.Let("b", b.ty.f32(), | 
|  | b.IndexAccessor(b.IndexAccessor(b.MemberAccessor("s", "a"), 1_i), 2_i))), | 
|  | }, | 
|  | utils::Vector{ | 
|  | b.Stage(ast::PipelineStage::kCompute), | 
|  | b.WorkgroupSize(1_i), | 
|  | }); | 
|  |  | 
|  | auto* expect = | 
|  | R"( | 
|  | struct S { | 
|  | a : array<vec4<f32>, 4u>, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<uniform> s : S; | 
|  |  | 
|  | @compute @workgroup_size(1i) | 
|  | fn f() { | 
|  | let a : array<vec4<f32>, 4u> = s.a; | 
|  | let b : f32 = s.a[1i][2i]; | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedArray>(Program(std::move(b))); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(DecomposeStridedArrayTest, ReadStorageStridedArray) { | 
|  | // struct S { | 
|  | //   a : @stride(32) array<f32, 4u>, | 
|  | // }; | 
|  | // @group(0) @binding(0) var<storage> s : S; | 
|  | // | 
|  | // @compute @workgroup_size(1) | 
|  | // fn f() { | 
|  | //   let a : @stride(32) array<f32, 4u> = s.a; | 
|  | //   let b : f32 = s.a[1]; | 
|  | // } | 
|  | ProgramBuilder b; | 
|  | auto* S = b.Structure("S", utils::Vector{b.Member("a", b.ty.array<f32, 4u>(32))}); | 
|  | b.GlobalVar("s", b.ty.Of(S), ast::AddressSpace::kStorage, b.Group(0_a), b.Binding(0_a)); | 
|  | b.Func("f", utils::Empty, b.ty.void_(), | 
|  | utils::Vector{ | 
|  | b.Decl(b.Let("a", b.ty.array<f32, 4u>(32), b.MemberAccessor("s", "a"))), | 
|  | b.Decl(b.Let("b", b.ty.f32(), b.IndexAccessor(b.MemberAccessor("s", "a"), 1_i))), | 
|  | }, | 
|  | utils::Vector{ | 
|  | b.Stage(ast::PipelineStage::kCompute), | 
|  | b.WorkgroupSize(1_i), | 
|  | }); | 
|  |  | 
|  | auto* expect = R"( | 
|  | struct strided_arr { | 
|  | @size(32) | 
|  | el : f32, | 
|  | } | 
|  |  | 
|  | struct S { | 
|  | a : array<strided_arr, 4u>, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage> s : S; | 
|  |  | 
|  | @compute @workgroup_size(1i) | 
|  | fn f() { | 
|  | let a : array<strided_arr, 4u> = s.a; | 
|  | let b : f32 = s.a[1i].el; | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedArray>(Program(std::move(b))); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(DecomposeStridedArrayTest, ReadStorageDefaultStridedArray) { | 
|  | // struct S { | 
|  | //   a : @stride(4) array<f32, 4u>, | 
|  | // }; | 
|  | // @group(0) @binding(0) var<storage> s : S; | 
|  | // | 
|  | // @compute @workgroup_size(1) | 
|  | // fn f() { | 
|  | //   let a : @stride(4) array<f32, 4u> = s.a; | 
|  | //   let b : f32 = s.a[1]; | 
|  | // } | 
|  | ProgramBuilder b; | 
|  | auto* S = b.Structure("S", utils::Vector{b.Member("a", b.ty.array<f32, 4u>(4))}); | 
|  | b.GlobalVar("s", b.ty.Of(S), ast::AddressSpace::kStorage, b.Group(0_a), b.Binding(0_a)); | 
|  | b.Func("f", utils::Empty, b.ty.void_(), | 
|  | utils::Vector{ | 
|  | b.Decl(b.Let("a", b.ty.array<f32, 4u>(4), b.MemberAccessor("s", "a"))), | 
|  | b.Decl(b.Let("b", b.ty.f32(), b.IndexAccessor(b.MemberAccessor("s", "a"), 1_i))), | 
|  | }, | 
|  | utils::Vector{ | 
|  | b.Stage(ast::PipelineStage::kCompute), | 
|  | b.WorkgroupSize(1_i), | 
|  | }); | 
|  |  | 
|  | auto* expect = R"( | 
|  | struct S { | 
|  | a : array<f32, 4u>, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage> s : S; | 
|  |  | 
|  | @compute @workgroup_size(1i) | 
|  | fn f() { | 
|  | let a : array<f32, 4u> = s.a; | 
|  | let b : f32 = s.a[1i]; | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedArray>(Program(std::move(b))); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(DecomposeStridedArrayTest, WriteStorageStridedArray) { | 
|  | // struct S { | 
|  | //   a : @stride(32) array<f32, 4u>, | 
|  | // }; | 
|  | // @group(0) @binding(0) var<storage, read_write> s : S; | 
|  | // | 
|  | // @compute @workgroup_size(1) | 
|  | // fn f() { | 
|  | //   s.a = @stride(32) array<f32, 4u>(); | 
|  | //   s.a = @stride(32) array<f32, 4u>(1.0, 2.0, 3.0, 4.0); | 
|  | //   s.a[1i] = 5.0; | 
|  | // } | 
|  | ProgramBuilder b; | 
|  | auto* S = b.Structure("S", utils::Vector{b.Member("a", b.ty.array<f32, 4u>(32))}); | 
|  | b.GlobalVar("s", b.ty.Of(S), ast::AddressSpace::kStorage, ast::Access::kReadWrite, b.Group(0_a), | 
|  | b.Binding(0_a)); | 
|  | b.Func("f", utils::Empty, b.ty.void_(), | 
|  | utils::Vector{ | 
|  | b.Assign(b.MemberAccessor("s", "a"), b.Construct(b.ty.array<f32, 4u>(32))), | 
|  | b.Assign(b.MemberAccessor("s", "a"), | 
|  | b.Construct(b.ty.array<f32, 4u>(32), 1_f, 2_f, 3_f, 4_f)), | 
|  | b.Assign(b.IndexAccessor(b.MemberAccessor("s", "a"), 1_i), 5_f), | 
|  | }, | 
|  | utils::Vector{ | 
|  | b.Stage(ast::PipelineStage::kCompute), | 
|  | b.WorkgroupSize(1_i), | 
|  | }); | 
|  |  | 
|  | auto* expect = | 
|  | R"( | 
|  | struct strided_arr { | 
|  | @size(32) | 
|  | el : f32, | 
|  | } | 
|  |  | 
|  | struct S { | 
|  | a : array<strided_arr, 4u>, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> s : S; | 
|  |  | 
|  | @compute @workgroup_size(1i) | 
|  | fn f() { | 
|  | s.a = array<strided_arr, 4u>(); | 
|  | s.a = array<strided_arr, 4u>(strided_arr(1.0f), strided_arr(2.0f), strided_arr(3.0f), strided_arr(4.0f)); | 
|  | s.a[1i].el = 5.0f; | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedArray>(Program(std::move(b))); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(DecomposeStridedArrayTest, WriteStorageDefaultStridedArray) { | 
|  | // struct S { | 
|  | //   a : @stride(4) array<f32, 4u>, | 
|  | // }; | 
|  | // @group(0) @binding(0) var<storage, read_write> s : S; | 
|  | // | 
|  | // @compute @workgroup_size(1) | 
|  | // fn f() { | 
|  | //   s.a = @stride(4) array<f32, 4u>(); | 
|  | //   s.a = @stride(4) array<f32, 4u>(1.0, 2.0, 3.0, 4.0); | 
|  | //   s.a[1] = 5.0; | 
|  | // } | 
|  | ProgramBuilder b; | 
|  | auto* S = b.Structure("S", utils::Vector{b.Member("a", b.ty.array<f32, 4u>(4))}); | 
|  | b.GlobalVar("s", b.ty.Of(S), ast::AddressSpace::kStorage, ast::Access::kReadWrite, b.Group(0_a), | 
|  | b.Binding(0_a)); | 
|  | b.Func("f", utils::Empty, b.ty.void_(), | 
|  | utils::Vector{ | 
|  | b.Assign(b.MemberAccessor("s", "a"), b.Construct(b.ty.array<f32, 4u>(4))), | 
|  | b.Assign(b.MemberAccessor("s", "a"), | 
|  | b.Construct(b.ty.array<f32, 4u>(4), 1_f, 2_f, 3_f, 4_f)), | 
|  | b.Assign(b.IndexAccessor(b.MemberAccessor("s", "a"), 1_i), 5_f), | 
|  | }, | 
|  | utils::Vector{ | 
|  | b.Stage(ast::PipelineStage::kCompute), | 
|  | b.WorkgroupSize(1_i), | 
|  | }); | 
|  |  | 
|  | auto* expect = | 
|  | R"( | 
|  | struct S { | 
|  | a : array<f32, 4u>, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> s : S; | 
|  |  | 
|  | @compute @workgroup_size(1i) | 
|  | fn f() { | 
|  | s.a = array<f32, 4u>(); | 
|  | s.a = array<f32, 4u>(1.0f, 2.0f, 3.0f, 4.0f); | 
|  | s.a[1i] = 5.0f; | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedArray>(Program(std::move(b))); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(DecomposeStridedArrayTest, ReadWriteViaPointerLets) { | 
|  | // struct S { | 
|  | //   a : @stride(32) array<f32, 4u>, | 
|  | // }; | 
|  | // @group(0) @binding(0) var<storage, read_write> s : S; | 
|  | // | 
|  | // @compute @workgroup_size(1) | 
|  | // fn f() { | 
|  | //   let a = &s.a; | 
|  | //   let b = &*&*(a); | 
|  | //   let c = *b; | 
|  | //   let d = (*b)[1]; | 
|  | //   (*b) = @stride(32) array<f32, 4u>(1.0, 2.0, 3.0, 4.0); | 
|  | //   (*b)[1] = 5.0; | 
|  | // } | 
|  | ProgramBuilder b; | 
|  | auto* S = b.Structure("S", utils::Vector{b.Member("a", b.ty.array<f32, 4u>(32))}); | 
|  | b.GlobalVar("s", b.ty.Of(S), ast::AddressSpace::kStorage, ast::Access::kReadWrite, b.Group(0_a), | 
|  | b.Binding(0_a)); | 
|  | b.Func("f", utils::Empty, b.ty.void_(), | 
|  | utils::Vector{ | 
|  | b.Decl(b.Let("a", b.AddressOf(b.MemberAccessor("s", "a")))), | 
|  | b.Decl(b.Let("b", b.AddressOf(b.Deref(b.AddressOf(b.Deref("a")))))), | 
|  | b.Decl(b.Let("c", b.Deref("b"))), | 
|  | b.Decl(b.Let("d", b.IndexAccessor(b.Deref("b"), 1_i))), | 
|  | b.Assign(b.Deref("b"), b.Construct(b.ty.array<f32, 4u>(32), 1_f, 2_f, 3_f, 4_f)), | 
|  | b.Assign(b.IndexAccessor(b.Deref("b"), 1_i), 5_f), | 
|  | }, | 
|  | utils::Vector{ | 
|  | b.Stage(ast::PipelineStage::kCompute), | 
|  | b.WorkgroupSize(1_i), | 
|  | }); | 
|  |  | 
|  | auto* expect = | 
|  | R"( | 
|  | struct strided_arr { | 
|  | @size(32) | 
|  | el : f32, | 
|  | } | 
|  |  | 
|  | struct S { | 
|  | a : array<strided_arr, 4u>, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> s : S; | 
|  |  | 
|  | @compute @workgroup_size(1i) | 
|  | fn f() { | 
|  | let c = s.a; | 
|  | let d = s.a[1i].el; | 
|  | s.a = array<strided_arr, 4u>(strided_arr(1.0f), strided_arr(2.0f), strided_arr(3.0f), strided_arr(4.0f)); | 
|  | s.a[1i].el = 5.0f; | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedArray>(Program(std::move(b))); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(DecomposeStridedArrayTest, PrivateAliasedStridedArray) { | 
|  | // type ARR = @stride(32) array<f32, 4u>; | 
|  | // struct S { | 
|  | //   a : ARR, | 
|  | // }; | 
|  | // @group(0) @binding(0) var<storage, read_write> s : S; | 
|  | // | 
|  | // @compute @workgroup_size(1) | 
|  | // fn f() { | 
|  | //   let a : ARR = s.a; | 
|  | //   let b : f32 = s.a[1]; | 
|  | //   s.a = ARR(); | 
|  | //   s.a = ARR(1.0, 2.0, 3.0, 4.0); | 
|  | //   s.a[1] = 5.0; | 
|  | // } | 
|  | ProgramBuilder b; | 
|  | b.Alias("ARR", b.ty.array<f32, 4u>(32)); | 
|  | auto* S = b.Structure("S", utils::Vector{b.Member("a", b.ty.type_name("ARR"))}); | 
|  | b.GlobalVar("s", b.ty.Of(S), ast::AddressSpace::kStorage, ast::Access::kReadWrite, b.Group(0_a), | 
|  | b.Binding(0_a)); | 
|  | b.Func("f", utils::Empty, b.ty.void_(), | 
|  | utils::Vector{ | 
|  | b.Decl(b.Let("a", b.ty.type_name("ARR"), b.MemberAccessor("s", "a"))), | 
|  | b.Decl(b.Let("b", b.ty.f32(), b.IndexAccessor(b.MemberAccessor("s", "a"), 1_i))), | 
|  | b.Assign(b.MemberAccessor("s", "a"), b.Construct(b.ty.type_name("ARR"))), | 
|  | b.Assign(b.MemberAccessor("s", "a"), | 
|  | b.Construct(b.ty.type_name("ARR"), 1_f, 2_f, 3_f, 4_f)), | 
|  | b.Assign(b.IndexAccessor(b.MemberAccessor("s", "a"), 1_i), 5_f), | 
|  | }, | 
|  | utils::Vector{ | 
|  | b.Stage(ast::PipelineStage::kCompute), | 
|  | b.WorkgroupSize(1_i), | 
|  | }); | 
|  |  | 
|  | auto* expect = R"( | 
|  | struct strided_arr { | 
|  | @size(32) | 
|  | el : f32, | 
|  | } | 
|  |  | 
|  | type ARR = array<strided_arr, 4u>; | 
|  |  | 
|  | struct S { | 
|  | a : ARR, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> s : S; | 
|  |  | 
|  | @compute @workgroup_size(1i) | 
|  | fn f() { | 
|  | let a : ARR = s.a; | 
|  | let b : f32 = s.a[1i].el; | 
|  | s.a = ARR(); | 
|  | s.a = ARR(strided_arr(1.0f), strided_arr(2.0f), strided_arr(3.0f), strided_arr(4.0f)); | 
|  | s.a[1i].el = 5.0f; | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedArray>(Program(std::move(b))); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(DecomposeStridedArrayTest, PrivateNestedStridedArray) { | 
|  | // type ARR_A = @stride(8) array<f32, 2u>; | 
|  | // type ARR_B = @stride(128) array<@stride(16) array<ARR_A, 3u>, 4u>; | 
|  | // struct S { | 
|  | //   a : ARR_B, | 
|  | // }; | 
|  | // @group(0) @binding(0) var<storage, read_write> s : S; | 
|  | // | 
|  | // @compute @workgroup_size(1) | 
|  | // fn f() { | 
|  | //   let a : ARR_B = s.a; | 
|  | //   let b : array<@stride(8) array<f32, 2u>, 3u> = s.a[3]; | 
|  | //   let c = s.a[3][2]; | 
|  | //   let d = s.a[3][2][1]; | 
|  | //   s.a = ARR_B(); | 
|  | //   s.a[3][2][1] = 5.0; | 
|  | // } | 
|  |  | 
|  | ProgramBuilder b; | 
|  | b.Alias("ARR_A", b.ty.array<f32, 2>(8)); | 
|  | b.Alias("ARR_B", | 
|  | b.ty.array(                                        // | 
|  | b.ty.array(b.ty.type_name("ARR_A"), 3_u, 16),  // | 
|  | 4_u, 128)); | 
|  | auto* S = b.Structure("S", utils::Vector{b.Member("a", b.ty.type_name("ARR_B"))}); | 
|  | b.GlobalVar("s", b.ty.Of(S), ast::AddressSpace::kStorage, ast::Access::kReadWrite, b.Group(0_a), | 
|  | b.Binding(0_a)); | 
|  | b.Func("f", utils::Empty, b.ty.void_(), | 
|  | utils::Vector{ | 
|  | b.Decl(b.Let("a", b.ty.type_name("ARR_B"), b.MemberAccessor("s", "a"))), | 
|  | b.Decl(b.Let("b", b.ty.array(b.ty.type_name("ARR_A"), 3_u, 16), | 
|  | b.IndexAccessor(                 // | 
|  | b.MemberAccessor("s", "a"),  // | 
|  | 3_i))), | 
|  | b.Decl(b.Let("c", b.ty.type_name("ARR_A"), | 
|  | b.IndexAccessor(                     // | 
|  | b.IndexAccessor(                 // | 
|  | b.MemberAccessor("s", "a"),  // | 
|  | 3_i), | 
|  | 2_i))), | 
|  | b.Decl(b.Let("d", b.ty.f32(), | 
|  | b.IndexAccessor(                         // | 
|  | b.IndexAccessor(                     // | 
|  | b.IndexAccessor(                 // | 
|  | b.MemberAccessor("s", "a"),  // | 
|  | 3_i), | 
|  | 2_i), | 
|  | 1_i))), | 
|  | b.Assign(b.MemberAccessor("s", "a"), b.Construct(b.ty.type_name("ARR_B"))), | 
|  | b.Assign(b.IndexAccessor(                         // | 
|  | b.IndexAccessor(                     // | 
|  | b.IndexAccessor(                 // | 
|  | b.MemberAccessor("s", "a"),  // | 
|  | 3_i), | 
|  | 2_i), | 
|  | 1_i), | 
|  | 5_f), | 
|  | }, | 
|  | utils::Vector{ | 
|  | b.Stage(ast::PipelineStage::kCompute), | 
|  | b.WorkgroupSize(1_i), | 
|  | }); | 
|  |  | 
|  | auto* expect = | 
|  | R"( | 
|  | struct strided_arr { | 
|  | @size(8) | 
|  | el : f32, | 
|  | } | 
|  |  | 
|  | type ARR_A = array<strided_arr, 2u>; | 
|  |  | 
|  | struct strided_arr_1 { | 
|  | @size(128) | 
|  | el : array<ARR_A, 3u>, | 
|  | } | 
|  |  | 
|  | type ARR_B = array<strided_arr_1, 4u>; | 
|  |  | 
|  | struct S { | 
|  | a : ARR_B, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> s : S; | 
|  |  | 
|  | @compute @workgroup_size(1i) | 
|  | fn f() { | 
|  | let a : ARR_B = s.a; | 
|  | let b : array<ARR_A, 3u> = s.a[3i].el; | 
|  | let c : ARR_A = s.a[3i].el[2i]; | 
|  | let d : f32 = s.a[3i].el[2i][1i].el; | 
|  | s.a = ARR_B(); | 
|  | s.a[3i].el[2i][1i].el = 5.0f; | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedArray>(Program(std::move(b))); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  | }  // namespace | 
|  | }  // namespace tint::transform |