// Copyright 2021 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/localize_struct_array_assignment.h" #include "src/tint/transform/simplify_pointers.h" #include "src/tint/transform/unshadow.h" #include "src/tint/transform/test_helper.h" namespace tint::transform { namespace { using LocalizeStructArrayAssignmentTest = TransformTest; TEST_F(LocalizeStructArrayAssignmentTest, EmptyModule) { auto* src = R"()"; EXPECT_FALSE(ShouldRun(src)); } TEST_F(LocalizeStructArrayAssignmentTest, StructArray) { auto* src = R"( struct Uniforms { i : u32, }; struct InnerS { v : i32, }; struct OuterS { a1 : array, }; @group(1) @binding(4) var uniforms : Uniforms; @compute @workgroup_size(1) fn main() { var v : InnerS; var s1 : OuterS; s1.a1[uniforms.i] = v; } )"; auto* expect = R"( struct Uniforms { i : u32, } struct InnerS { v : i32, } struct OuterS { a1 : array, } @group(1) @binding(4) var uniforms : Uniforms; @compute @workgroup_size(1) fn main() { var v : InnerS; var s1 : OuterS; { let tint_symbol = &(s1.a1); var tint_symbol_1 = *(tint_symbol); tint_symbol_1[uniforms.i] = v; *(tint_symbol) = tint_symbol_1; } } )"; auto got = Run(src); EXPECT_EQ(expect, str(got)); } TEST_F(LocalizeStructArrayAssignmentTest, StructArray_OutOfOrder) { auto* src = R"( @compute @workgroup_size(1) fn main() { var v : InnerS; var s1 : OuterS; s1.a1[uniforms.i] = v; } @group(1) @binding(4) var uniforms : Uniforms; struct OuterS { a1 : array, }; struct InnerS { v : i32, }; struct Uniforms { i : u32, }; )"; auto* expect = R"( @compute @workgroup_size(1) fn main() { var v : InnerS; var s1 : OuterS; { let tint_symbol = &(s1.a1); var tint_symbol_1 = *(tint_symbol); tint_symbol_1[uniforms.i] = v; *(tint_symbol) = tint_symbol_1; } } @group(1) @binding(4) var uniforms : Uniforms; struct OuterS { a1 : array, } struct InnerS { v : i32, } struct Uniforms { i : u32, } )"; auto got = Run(src); EXPECT_EQ(expect, str(got)); } TEST_F(LocalizeStructArrayAssignmentTest, StructStructArray) { auto* src = R"( struct Uniforms { i : u32, }; struct InnerS { v : i32, }; struct S1 { a : array, }; struct OuterS { s2 : S1, }; @group(1) @binding(4) var uniforms : Uniforms; @compute @workgroup_size(1) fn main() { var v : InnerS; var s1 : OuterS; s1.s2.a[uniforms.i] = v; } )"; auto* expect = R"( struct Uniforms { i : u32, } struct InnerS { v : i32, } struct S1 { a : array, } struct OuterS { s2 : S1, } @group(1) @binding(4) var uniforms : Uniforms; @compute @workgroup_size(1) fn main() { var v : InnerS; var s1 : OuterS; { let tint_symbol = &(s1.s2.a); var tint_symbol_1 = *(tint_symbol); tint_symbol_1[uniforms.i] = v; *(tint_symbol) = tint_symbol_1; } } )"; auto got = Run(src); EXPECT_EQ(expect, str(got)); } TEST_F(LocalizeStructArrayAssignmentTest, StructStructArray_OutOfOrder) { auto* src = R"( @compute @workgroup_size(1) fn main() { var v : InnerS; var s1 : OuterS; s1.s2.a[uniforms.i] = v; } @group(1) @binding(4) var uniforms : Uniforms; struct OuterS { s2 : S1, }; struct S1 { a : array, }; struct InnerS { v : i32, }; struct Uniforms { i : u32, }; )"; auto* expect = R"( @compute @workgroup_size(1) fn main() { var v : InnerS; var s1 : OuterS; { let tint_symbol = &(s1.s2.a); var tint_symbol_1 = *(tint_symbol); tint_symbol_1[uniforms.i] = v; *(tint_symbol) = tint_symbol_1; } } @group(1) @binding(4) var uniforms : Uniforms; struct OuterS { s2 : S1, } struct S1 { a : array, } struct InnerS { v : i32, } struct Uniforms { i : u32, } )"; auto got = Run(src); EXPECT_EQ(expect, str(got)); } TEST_F(LocalizeStructArrayAssignmentTest, StructArrayArray) { auto* src = R"( struct Uniforms { i : u32, j : u32, }; struct InnerS { v : i32, }; struct OuterS { a1 : array, 8>, }; @group(1) @binding(4) var uniforms : Uniforms; @compute @workgroup_size(1) fn main() { var v : InnerS; var s1 : OuterS; s1.a1[uniforms.i][uniforms.j] = v; } )"; auto* expect = R"( struct Uniforms { i : u32, j : u32, } struct InnerS { v : i32, } struct OuterS { a1 : array, 8>, } @group(1) @binding(4) var uniforms : Uniforms; @compute @workgroup_size(1) fn main() { var v : InnerS; var s1 : OuterS; { let tint_symbol = &(s1.a1); var tint_symbol_1 = *(tint_symbol); tint_symbol_1[uniforms.i][uniforms.j] = v; *(tint_symbol) = tint_symbol_1; } } )"; auto got = Run(src); EXPECT_EQ(expect, str(got)); } TEST_F(LocalizeStructArrayAssignmentTest, StructArrayStruct) { auto* src = R"( struct Uniforms { i : u32, }; struct InnerS { v : i32, }; struct S1 { s2 : InnerS, }; struct OuterS { a1 : array, }; @group(1) @binding(4) var uniforms : Uniforms; @compute @workgroup_size(1) fn main() { var v : InnerS; var s1 : OuterS; s1.a1[uniforms.i].s2 = v; } )"; auto* expect = R"( struct Uniforms { i : u32, } struct InnerS { v : i32, } struct S1 { s2 : InnerS, } struct OuterS { a1 : array, } @group(1) @binding(4) var uniforms : Uniforms; @compute @workgroup_size(1) fn main() { var v : InnerS; var s1 : OuterS; { let tint_symbol = &(s1.a1); var tint_symbol_1 = *(tint_symbol); tint_symbol_1[uniforms.i].s2 = v; *(tint_symbol) = tint_symbol_1; } } )"; auto got = Run(src); EXPECT_EQ(expect, str(got)); } TEST_F(LocalizeStructArrayAssignmentTest, StructArrayStructArray) { auto* src = R"( struct Uniforms { i : u32, j : u32, }; struct InnerS { v : i32, }; struct S1 { a2 : array, }; struct OuterS { a1 : array, }; @group(1) @binding(4) var uniforms : Uniforms; @compute @workgroup_size(1) fn main() { var v : InnerS; var s : OuterS; s.a1[uniforms.i].a2[uniforms.j] = v; } )"; auto* expect = R"( struct Uniforms { i : u32, j : u32, } struct InnerS { v : i32, } struct S1 { a2 : array, } struct OuterS { a1 : array, } @group(1) @binding(4) var uniforms : Uniforms; @compute @workgroup_size(1) fn main() { var v : InnerS; var s : OuterS; { let tint_symbol = &(s.a1); var tint_symbol_1 = *(tint_symbol); let tint_symbol_2 = &(tint_symbol_1[uniforms.i].a2); var tint_symbol_3 = *(tint_symbol_2); tint_symbol_3[uniforms.j] = v; *(tint_symbol_2) = tint_symbol_3; *(tint_symbol) = tint_symbol_1; } } )"; auto got = Run(src); EXPECT_EQ(expect, str(got)); } TEST_F(LocalizeStructArrayAssignmentTest, IndexingWithSideEffectFunc) { auto* src = R"( struct Uniforms { i : u32, j : u32, }; struct InnerS { v : i32, }; struct S1 { a2 : array, }; struct OuterS { a1 : array, }; var nextIndex : u32; fn getNextIndex() -> u32 { nextIndex = nextIndex + 1u; return nextIndex; } @group(1) @binding(4) var uniforms : Uniforms; @compute @workgroup_size(1) fn main() { var v : InnerS; var s : OuterS; s.a1[getNextIndex()].a2[uniforms.j] = v; } )"; auto* expect = R"( struct Uniforms { i : u32, j : u32, } struct InnerS { v : i32, } struct S1 { a2 : array, } struct OuterS { a1 : array, } var nextIndex : u32; fn getNextIndex() -> u32 { nextIndex = (nextIndex + 1u); return nextIndex; } @group(1) @binding(4) var uniforms : Uniforms; @compute @workgroup_size(1) fn main() { var v : InnerS; var s : OuterS; { let tint_symbol = &(s.a1); var tint_symbol_1 = *(tint_symbol); let tint_symbol_2 = &(tint_symbol_1[getNextIndex()].a2); var tint_symbol_3 = *(tint_symbol_2); tint_symbol_3[uniforms.j] = v; *(tint_symbol_2) = tint_symbol_3; *(tint_symbol) = tint_symbol_1; } } )"; auto got = Run(src); EXPECT_EQ(expect, str(got)); } TEST_F(LocalizeStructArrayAssignmentTest, IndexingWithSideEffectFunc_OutOfOrder) { auto* src = R"( @compute @workgroup_size(1) fn main() { var v : InnerS; var s : OuterS; s.a1[getNextIndex()].a2[uniforms.j] = v; } @group(1) @binding(4) var uniforms : Uniforms; struct Uniforms { i : u32, j : u32, }; var nextIndex : u32; fn getNextIndex() -> u32 { nextIndex = nextIndex + 1u; return nextIndex; } struct OuterS { a1 : array, }; struct S1 { a2 : array, }; struct InnerS { v : i32, }; )"; auto* expect = R"( @compute @workgroup_size(1) fn main() { var v : InnerS; var s : OuterS; { let tint_symbol = &(s.a1); var tint_symbol_1 = *(tint_symbol); let tint_symbol_2 = &(tint_symbol_1[getNextIndex()].a2); var tint_symbol_3 = *(tint_symbol_2); tint_symbol_3[uniforms.j] = v; *(tint_symbol_2) = tint_symbol_3; *(tint_symbol) = tint_symbol_1; } } @group(1) @binding(4) var uniforms : Uniforms; struct Uniforms { i : u32, j : u32, } var nextIndex : u32; fn getNextIndex() -> u32 { nextIndex = (nextIndex + 1u); return nextIndex; } struct OuterS { a1 : array, } struct S1 { a2 : array, } struct InnerS { v : i32, } )"; auto got = Run(src); EXPECT_EQ(expect, str(got)); } TEST_F(LocalizeStructArrayAssignmentTest, ViaPointerArg) { auto* src = R"( struct Uniforms { i : u32, }; struct InnerS { v : i32, }; struct OuterS { a1 : array, }; @group(1) @binding(4) var uniforms : Uniforms; fn f(p : ptr) { var v : InnerS; (*p).a1[uniforms.i] = v; } @compute @workgroup_size(1) fn main() { var s1 : OuterS; f(&s1); } )"; auto* expect = R"( struct Uniforms { i : u32, } struct InnerS { v : i32, } struct OuterS { a1 : array, } @group(1) @binding(4) var uniforms : Uniforms; fn f(p : ptr) { var v : InnerS; { let tint_symbol = &((*(p)).a1); var tint_symbol_1 = *(tint_symbol); tint_symbol_1[uniforms.i] = v; *(tint_symbol) = tint_symbol_1; } } @compute @workgroup_size(1) fn main() { var s1 : OuterS; f(&(s1)); } )"; auto got = Run(src); EXPECT_EQ(expect, str(got)); } TEST_F(LocalizeStructArrayAssignmentTest, ViaPointerArg_OutOfOrder) { auto* src = R"( @compute @workgroup_size(1) fn main() { var s1 : OuterS; f(&s1); } fn f(p : ptr) { var v : InnerS; (*p).a1[uniforms.i] = v; } struct InnerS { v : i32, }; struct OuterS { a1 : array, }; @group(1) @binding(4) var uniforms : Uniforms; struct Uniforms { i : u32, }; )"; auto* expect = R"( @compute @workgroup_size(1) fn main() { var s1 : OuterS; f(&(s1)); } fn f(p : ptr) { var v : InnerS; { let tint_symbol = &((*(p)).a1); var tint_symbol_1 = *(tint_symbol); tint_symbol_1[uniforms.i] = v; *(tint_symbol) = tint_symbol_1; } } struct InnerS { v : i32, } struct OuterS { a1 : array, } @group(1) @binding(4) var uniforms : Uniforms; struct Uniforms { i : u32, } )"; auto got = Run(src); EXPECT_EQ(expect, str(got)); } TEST_F(LocalizeStructArrayAssignmentTest, ViaPointerVar) { auto* src = R"( struct Uniforms { i : u32, }; struct InnerS { v : i32, }; struct OuterS { a1 : array, }; @group(1) @binding(4) var uniforms : Uniforms; fn f(p : ptr, v : InnerS) { *(p) = v; } @compute @workgroup_size(1) fn main() { var v : InnerS; var s1 : OuterS; let p = &(s1.a1[uniforms.i]); *(p) = v; } )"; auto* expect = R"( struct Uniforms { i : u32, } struct InnerS { v : i32, } struct OuterS { a1 : array, } @group(1) @binding(4) var uniforms : Uniforms; fn f(p : ptr, v : InnerS) { *(p) = v; } @compute @workgroup_size(1) fn main() { var v : InnerS; var s1 : OuterS; let p_save = uniforms.i; { let tint_symbol = &(s1.a1); var tint_symbol_1 = *(tint_symbol); tint_symbol_1[p_save] = v; *(tint_symbol) = tint_symbol_1; } } )"; auto got = Run(src); EXPECT_EQ(expect, str(got)); } TEST_F(LocalizeStructArrayAssignmentTest, VectorAssignment) { auto* src = R"( struct Uniforms { i : u32, } struct OuterS { a1 : array, } @group(1) @binding(4) var uniforms : Uniforms; fn f(i : u32) -> u32 { return (i + 1u); } @compute @workgroup_size(1) fn main() { var s1 : OuterS; var v : vec3; v[s1.a1[uniforms.i]] = 1.0; v[f(s1.a1[uniforms.i])] = 1.0; } )"; // Transform does nothing here as we're not actually assigning to the array in // the struct. EXPECT_FALSE(ShouldRun(src)); } } // namespace } // namespace tint::transform