|  | // 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/preserve_padding.h" | 
|  |  | 
|  | #include <utility> | 
|  |  | 
|  | #include "src/tint/transform/test_helper.h" | 
|  |  | 
|  | namespace tint::transform { | 
|  | namespace { | 
|  |  | 
|  | using PreservePaddingTest = TransformTest; | 
|  |  | 
|  | TEST_F(PreservePaddingTest, ShouldRun_EmptyModule) { | 
|  | auto* src = R"()"; | 
|  |  | 
|  | EXPECT_FALSE(ShouldRun<PreservePadding>(src)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, ShouldRun_NonStructVec3) { | 
|  | auto* src = R"( | 
|  | @group(0) @binding(0) var<storage, read_write> v : vec3<u32>; | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | v = vec3<u32>(); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | EXPECT_FALSE(ShouldRun<PreservePadding>(src)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, ShouldRun_StructWithoutPadding) { | 
|  | auto* src = R"( | 
|  | struct S { | 
|  | a : u32, | 
|  | b : u32, | 
|  | c : u32, | 
|  | d : u32, | 
|  | e : vec3<u32>, | 
|  | f : u32, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> v : S; | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | v = S(); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | EXPECT_FALSE(ShouldRun<PreservePadding>(src)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, ShouldRun_ArrayWithoutPadding) { | 
|  | auto* src = R"( | 
|  | @group(0) @binding(0) var<storage, read_write> v : array<vec4<u32>, 4>; | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | v = array<vec4<u32>, 4>(); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | EXPECT_FALSE(ShouldRun<PreservePadding>(src)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, EmptyModule) { | 
|  | auto* src = R"()"; | 
|  |  | 
|  | auto* expect = src; | 
|  |  | 
|  | auto got = Run<PreservePadding>(src); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, StructTrailingPadding) { | 
|  | auto* src = R"( | 
|  | struct S { | 
|  | a : u32, | 
|  | b : u32, | 
|  | c : u32, | 
|  | d : u32, | 
|  | e : vec3<u32>, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> v : S; | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | v = S(); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto* expect = R"( | 
|  | enable chromium_experimental_full_ptr_parameters; | 
|  |  | 
|  | struct S { | 
|  | a : u32, | 
|  | b : u32, | 
|  | c : u32, | 
|  | d : u32, | 
|  | e : vec3<u32>, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> v : S; | 
|  |  | 
|  | fn assign_and_preserve_padding(dest : ptr<storage, S, read_write>, value : S) { | 
|  | (*(dest)).a = value.a; | 
|  | (*(dest)).b = value.b; | 
|  | (*(dest)).c = value.c; | 
|  | (*(dest)).d = value.d; | 
|  | (*(dest)).e = value.e; | 
|  | } | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | assign_and_preserve_padding(&(v), S()); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto got = Run<PreservePadding>(src); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, StructInternalPadding) { | 
|  | auto* src = R"( | 
|  | struct S { | 
|  | a : u32, | 
|  | b : vec4<u32>, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> v : S; | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | v = S(); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto* expect = R"( | 
|  | enable chromium_experimental_full_ptr_parameters; | 
|  |  | 
|  | struct S { | 
|  | a : u32, | 
|  | b : vec4<u32>, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> v : S; | 
|  |  | 
|  | fn assign_and_preserve_padding(dest : ptr<storage, S, read_write>, value : S) { | 
|  | (*(dest)).a = value.a; | 
|  | (*(dest)).b = value.b; | 
|  | } | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | assign_and_preserve_padding(&(v), S()); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto got = Run<PreservePadding>(src); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, StructExplicitSize_TrailingPadding) { | 
|  | auto* src = R"( | 
|  | struct S { | 
|  | @size(16) a : u32, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> v : S; | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | v = S(); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto* expect = R"( | 
|  | enable chromium_experimental_full_ptr_parameters; | 
|  |  | 
|  | struct S { | 
|  | @size(16) | 
|  | a : u32, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> v : S; | 
|  |  | 
|  | fn assign_and_preserve_padding(dest : ptr<storage, S, read_write>, value : S) { | 
|  | (*(dest)).a = value.a; | 
|  | } | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | assign_and_preserve_padding(&(v), S()); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto got = Run<PreservePadding>(src); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, StructExplicitSize_InternalPadding) { | 
|  | auto* src = R"( | 
|  | struct S { | 
|  | @size(16) a : u32, | 
|  | b : u32, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> v : S; | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | v = S(); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto* expect = R"( | 
|  | enable chromium_experimental_full_ptr_parameters; | 
|  |  | 
|  | struct S { | 
|  | @size(16) | 
|  | a : u32, | 
|  | b : u32, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> v : S; | 
|  |  | 
|  | fn assign_and_preserve_padding(dest : ptr<storage, S, read_write>, value : S) { | 
|  | (*(dest)).a = value.a; | 
|  | (*(dest)).b = value.b; | 
|  | } | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | assign_and_preserve_padding(&(v), S()); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto got = Run<PreservePadding>(src); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, NestedStructs) { | 
|  | auto* src = R"( | 
|  | struct S1 { | 
|  | a1 : u32, | 
|  | b1 : vec3<u32>, | 
|  | c1 : u32, | 
|  | } | 
|  |  | 
|  | struct S2 { | 
|  | a2 : u32, | 
|  | b2 : S1, | 
|  | c2 : S1, | 
|  | } | 
|  |  | 
|  | struct S3 { | 
|  | a3 : S1, | 
|  | b3 : S2, | 
|  | c3 : S2, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> v : S3; | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | v = S3(); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto* expect = R"( | 
|  | enable chromium_experimental_full_ptr_parameters; | 
|  |  | 
|  | struct S1 { | 
|  | a1 : u32, | 
|  | b1 : vec3<u32>, | 
|  | c1 : u32, | 
|  | } | 
|  |  | 
|  | struct S2 { | 
|  | a2 : u32, | 
|  | b2 : S1, | 
|  | c2 : S1, | 
|  | } | 
|  |  | 
|  | struct S3 { | 
|  | a3 : S1, | 
|  | b3 : S2, | 
|  | c3 : S2, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> v : S3; | 
|  |  | 
|  | fn assign_and_preserve_padding_1(dest : ptr<storage, S1, read_write>, value : S1) { | 
|  | (*(dest)).a1 = value.a1; | 
|  | (*(dest)).b1 = value.b1; | 
|  | (*(dest)).c1 = value.c1; | 
|  | } | 
|  |  | 
|  | fn assign_and_preserve_padding_2(dest : ptr<storage, S2, read_write>, value : S2) { | 
|  | (*(dest)).a2 = value.a2; | 
|  | assign_and_preserve_padding_1(&((*(dest)).b2), value.b2); | 
|  | assign_and_preserve_padding_1(&((*(dest)).c2), value.c2); | 
|  | } | 
|  |  | 
|  | fn assign_and_preserve_padding(dest : ptr<storage, S3, read_write>, value : S3) { | 
|  | assign_and_preserve_padding_1(&((*(dest)).a3), value.a3); | 
|  | assign_and_preserve_padding_2(&((*(dest)).b3), value.b3); | 
|  | assign_and_preserve_padding_2(&((*(dest)).c3), value.c3); | 
|  | } | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | assign_and_preserve_padding(&(v), S3()); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto got = Run<PreservePadding>(src); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, ArrayOfVec3) { | 
|  | auto* src = R"( | 
|  | @group(0) @binding(0) var<storage, read_write> v : array<vec3<u32>, 4>; | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | v = array<vec3<u32>, 4>(); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto* expect = R"( | 
|  | enable chromium_experimental_full_ptr_parameters; | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> v : array<vec3<u32>, 4>; | 
|  |  | 
|  | fn assign_and_preserve_padding(dest : ptr<storage, array<vec3<u32>, 4u>, read_write>, value : array<vec3<u32>, 4u>) { | 
|  | for(var i = 0u; (i < 4u); i = (i + 1u)) { | 
|  | (*(dest))[i] = value[i]; | 
|  | } | 
|  | } | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | assign_and_preserve_padding(&(v), array<vec3<u32>, 4>()); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto got = Run<PreservePadding>(src); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, ArrayOfArray) { | 
|  | auto* src = R"( | 
|  | type Array = array<array<vec3<u32>, 4>, 3>; | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> v : Array; | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | v = Array(); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto* expect = R"( | 
|  | enable chromium_experimental_full_ptr_parameters; | 
|  |  | 
|  | alias Array = array<array<vec3<u32>, 4>, 3>; | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> v : Array; | 
|  |  | 
|  | fn assign_and_preserve_padding_1(dest : ptr<storage, array<vec3<u32>, 4u>, read_write>, value : array<vec3<u32>, 4u>) { | 
|  | for(var i = 0u; (i < 4u); i = (i + 1u)) { | 
|  | (*(dest))[i] = value[i]; | 
|  | } | 
|  | } | 
|  |  | 
|  | fn assign_and_preserve_padding(dest : ptr<storage, array<array<vec3<u32>, 4u>, 3u>, read_write>, value : array<array<vec3<u32>, 4u>, 3u>) { | 
|  | for(var i = 0u; (i < 3u); i = (i + 1u)) { | 
|  | assign_and_preserve_padding_1(&((*(dest))[i]), value[i]); | 
|  | } | 
|  | } | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | assign_and_preserve_padding(&(v), Array()); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto got = Run<PreservePadding>(src); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, ArrayOfStructOfArray) { | 
|  | auto* src = R"( | 
|  | struct S { | 
|  | a : u32, | 
|  | b : array<vec3<u32>, 4>, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> v : array<S, 3>; | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | v = array<S, 3>(); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto* expect = R"( | 
|  | enable chromium_experimental_full_ptr_parameters; | 
|  |  | 
|  | struct S { | 
|  | a : u32, | 
|  | b : array<vec3<u32>, 4>, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> v : array<S, 3>; | 
|  |  | 
|  | fn assign_and_preserve_padding_2(dest : ptr<storage, array<vec3<u32>, 4u>, read_write>, value : array<vec3<u32>, 4u>) { | 
|  | for(var i = 0u; (i < 4u); i = (i + 1u)) { | 
|  | (*(dest))[i] = value[i]; | 
|  | } | 
|  | } | 
|  |  | 
|  | fn assign_and_preserve_padding_1(dest : ptr<storage, S, read_write>, value : S) { | 
|  | (*(dest)).a = value.a; | 
|  | assign_and_preserve_padding_2(&((*(dest)).b), value.b); | 
|  | } | 
|  |  | 
|  | fn assign_and_preserve_padding(dest : ptr<storage, array<S, 3u>, read_write>, value : array<S, 3u>) { | 
|  | for(var i = 0u; (i < 3u); i = (i + 1u)) { | 
|  | assign_and_preserve_padding_1(&((*(dest))[i]), value[i]); | 
|  | } | 
|  | } | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | assign_and_preserve_padding(&(v), array<S, 3>()); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto got = Run<PreservePadding>(src); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, NoModify_Vec3) { | 
|  | auto* src = R"( | 
|  | @group(0) @binding(0) var<storage, read_write> v : vec3<u32>; | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | v = vec3<u32>(); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto* expect = src; | 
|  |  | 
|  | auto got = Run<PreservePadding>(src); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, AvoidDuplicateEnables) { | 
|  | auto* src = R"( | 
|  | enable chromium_experimental_full_ptr_parameters; | 
|  |  | 
|  | struct S { | 
|  | @size(16) a : u32, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> v : S; | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | v = S(); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto* expect = R"( | 
|  | enable chromium_experimental_full_ptr_parameters; | 
|  |  | 
|  | struct S { | 
|  | @size(16) | 
|  | a : u32, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> v : S; | 
|  |  | 
|  | fn assign_and_preserve_padding(dest : ptr<storage, S, read_write>, value : S) { | 
|  | (*(dest)).a = value.a; | 
|  | } | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | assign_and_preserve_padding(&(v), S()); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto got = Run<PreservePadding>(src); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, NoModify_Mat3x3) { | 
|  | auto* src = R"( | 
|  | @group(0) @binding(0) var<storage, read_write> v : mat3x3<f32>; | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | v = mat3x3<f32>(); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto* expect = src; | 
|  |  | 
|  | auto got = Run<PreservePadding>(src); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, NoModify_StructNoPadding) { | 
|  | auto* src = R"( | 
|  | struct S { | 
|  | a : u32, | 
|  | b : u32, | 
|  | c : u32, | 
|  | d : u32, | 
|  | e : vec4<u32>, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> v : S; | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | v = S(); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto* expect = src; | 
|  |  | 
|  | auto got = Run<PreservePadding>(src); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, NoModify_ArrayNoPadding) { | 
|  | auto* src = R"( | 
|  | @group(0) @binding(0) var<storage, read_write> v : array<vec4<u32>, 4>; | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | v = array<vec4<u32>, 4>(); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto* expect = src; | 
|  |  | 
|  | auto got = Run<PreservePadding>(src); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, NoModify_ArrayOfStructNoPadding) { | 
|  | auto* src = R"( | 
|  | struct S { | 
|  | a : u32, | 
|  | b : u32, | 
|  | c : u32, | 
|  | d : u32, | 
|  | e : vec4<u32>, | 
|  | } | 
|  |  | 
|  | @group(0) @binding(0) var<storage, read_write> v : array<S, 4>; | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | v = array<S, 4>(); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto* expect = src; | 
|  |  | 
|  | auto got = Run<PreservePadding>(src); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, NoModify_Workgroup) { | 
|  | auto* src = R"( | 
|  | struct S { | 
|  | a : u32, | 
|  | b : vec3<u32>, | 
|  | } | 
|  |  | 
|  | var<workgroup> v : S; | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | v = S(); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto* expect = src; | 
|  |  | 
|  | auto got = Run<PreservePadding>(src); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, NoModify_Private) { | 
|  | auto* src = R"( | 
|  | struct S { | 
|  | a : u32, | 
|  | b : vec3<u32>, | 
|  | } | 
|  |  | 
|  | var<private> v : S; | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | v = S(); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto* expect = src; | 
|  |  | 
|  | auto got = Run<PreservePadding>(src); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | TEST_F(PreservePaddingTest, NoModify_Function) { | 
|  | auto* src = R"( | 
|  | struct S { | 
|  | a : u32, | 
|  | b : vec3<u32>, | 
|  | } | 
|  |  | 
|  | @compute @workgroup_size(1) | 
|  | fn foo() { | 
|  | var<function> v : S; | 
|  | v = S(); | 
|  | } | 
|  | )"; | 
|  |  | 
|  | auto* expect = src; | 
|  |  | 
|  | auto got = Run<PreservePadding>(src); | 
|  |  | 
|  | EXPECT_EQ(expect, str(got)); | 
|  | } | 
|  |  | 
|  | }  // namespace | 
|  | }  // namespace tint::transform |