Consistent formatting for Dawn/Tint.
This CL updates the clang format files to have a single shared format
between Dawn and Tint. The major changes are tabs are 4 spaces, lines
are 100 columns and namespaces are not indented.
Bug: dawn:1339
Change-Id: I4208742c95643998d9fd14e77a9cc558071ded39
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/87603
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/tint/transform/decompose_strided_matrix_test.cc b/src/tint/transform/decompose_strided_matrix_test.cc
index 4246c6e..8784839 100644
--- a/src/tint/transform/decompose_strided_matrix_test.cc
+++ b/src/tint/transform/decompose_strided_matrix_test.cc
@@ -31,64 +31,61 @@
using f32 = ProgramBuilder::f32;
TEST_F(DecomposeStridedMatrixTest, ShouldRunEmptyModule) {
- auto* src = R"()";
+ auto* src = R"()";
- EXPECT_FALSE(ShouldRun<DecomposeStridedMatrix>(src));
+ EXPECT_FALSE(ShouldRun<DecomposeStridedMatrix>(src));
}
TEST_F(DecomposeStridedMatrixTest, ShouldRunNonStridedMatrox) {
- auto* src = R"(
+ auto* src = R"(
var<private> m : mat3x2<f32>;
)";
- EXPECT_FALSE(ShouldRun<DecomposeStridedMatrix>(src));
+ EXPECT_FALSE(ShouldRun<DecomposeStridedMatrix>(src));
}
TEST_F(DecomposeStridedMatrixTest, Empty) {
- auto* src = R"()";
- auto* expect = src;
+ auto* src = R"()";
+ auto* expect = src;
- auto got = Run<DecomposeStridedMatrix>(src);
+ auto got = Run<DecomposeStridedMatrix>(src);
- EXPECT_EQ(expect, str(got));
+ EXPECT_EQ(expect, str(got));
}
TEST_F(DecomposeStridedMatrixTest, ReadUniformMatrix) {
- // struct S {
- // @offset(16) @stride(32)
- // @internal(ignore_stride_attribute)
- // m : mat2x2<f32>,
- // };
- // @group(0) @binding(0) var<uniform> s : S;
- //
- // @stage(compute) @workgroup_size(1)
- // fn f() {
- // let x : mat2x2<f32> = s.m;
- // }
- ProgramBuilder b;
- auto* S = b.Structure(
- "S",
- {
- b.Member(
- "m", b.ty.mat2x2<f32>(),
- {
- b.create<ast::StructMemberOffsetAttribute>(16),
- b.create<ast::StrideAttribute>(32),
- b.Disable(ast::DisabledValidation::kIgnoreStrideAttribute),
- }),
- });
- b.Global("s", b.ty.Of(S), ast::StorageClass::kUniform,
- b.GroupAndBinding(0, 0));
- b.Func("f", {}, b.ty.void_(),
- {
- b.Decl(b.Let("x", b.ty.mat2x2<f32>(), b.MemberAccessor("s", "m"))),
- },
- {
- b.Stage(ast::PipelineStage::kCompute),
- b.WorkgroupSize(1),
- });
+ // struct S {
+ // @offset(16) @stride(32)
+ // @internal(ignore_stride_attribute)
+ // m : mat2x2<f32>,
+ // };
+ // @group(0) @binding(0) var<uniform> s : S;
+ //
+ // @stage(compute) @workgroup_size(1)
+ // fn f() {
+ // let x : mat2x2<f32> = s.m;
+ // }
+ ProgramBuilder b;
+ auto* S = b.Structure(
+ "S", {
+ b.Member("m", b.ty.mat2x2<f32>(),
+ {
+ b.create<ast::StructMemberOffsetAttribute>(16),
+ b.create<ast::StrideAttribute>(32),
+ b.Disable(ast::DisabledValidation::kIgnoreStrideAttribute),
+ }),
+ });
+ b.Global("s", b.ty.Of(S), ast::StorageClass::kUniform, b.GroupAndBinding(0, 0));
+ b.Func("f", {}, b.ty.void_(),
+ {
+ b.Decl(b.Let("x", b.ty.mat2x2<f32>(), b.MemberAccessor("s", "m"))),
+ },
+ {
+ b.Stage(ast::PipelineStage::kCompute),
+ b.WorkgroupSize(1),
+ });
- auto* expect = R"(
+ auto* expect = R"(
struct S {
@size(16)
padding : u32,
@@ -107,49 +104,44 @@
}
)";
- auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
- Program(std::move(b)));
+ auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
- EXPECT_EQ(expect, str(got));
+ EXPECT_EQ(expect, str(got));
}
TEST_F(DecomposeStridedMatrixTest, ReadUniformColumn) {
- // struct S {
- // @offset(16) @stride(32)
- // @internal(ignore_stride_attribute)
- // m : mat2x2<f32>,
- // };
- // @group(0) @binding(0) var<uniform> s : S;
- //
- // @stage(compute) @workgroup_size(1)
- // fn f() {
- // let x : vec2<f32> = s.m[1];
- // }
- ProgramBuilder b;
- auto* S = b.Structure(
- "S",
- {
- b.Member(
- "m", b.ty.mat2x2<f32>(),
- {
- b.create<ast::StructMemberOffsetAttribute>(16),
- b.create<ast::StrideAttribute>(32),
- b.Disable(ast::DisabledValidation::kIgnoreStrideAttribute),
- }),
- });
- b.Global("s", b.ty.Of(S), ast::StorageClass::kUniform,
- b.GroupAndBinding(0, 0));
- b.Func("f", {}, b.ty.void_(),
- {
- b.Decl(b.Let("x", b.ty.vec2<f32>(),
- b.IndexAccessor(b.MemberAccessor("s", "m"), 1))),
- },
- {
- b.Stage(ast::PipelineStage::kCompute),
- b.WorkgroupSize(1),
- });
+ // struct S {
+ // @offset(16) @stride(32)
+ // @internal(ignore_stride_attribute)
+ // m : mat2x2<f32>,
+ // };
+ // @group(0) @binding(0) var<uniform> s : S;
+ //
+ // @stage(compute) @workgroup_size(1)
+ // fn f() {
+ // let x : vec2<f32> = s.m[1];
+ // }
+ ProgramBuilder b;
+ auto* S = b.Structure(
+ "S", {
+ b.Member("m", b.ty.mat2x2<f32>(),
+ {
+ b.create<ast::StructMemberOffsetAttribute>(16),
+ b.create<ast::StrideAttribute>(32),
+ b.Disable(ast::DisabledValidation::kIgnoreStrideAttribute),
+ }),
+ });
+ b.Global("s", b.ty.Of(S), ast::StorageClass::kUniform, b.GroupAndBinding(0, 0));
+ b.Func("f", {}, b.ty.void_(),
+ {
+ b.Decl(b.Let("x", b.ty.vec2<f32>(), b.IndexAccessor(b.MemberAccessor("s", "m"), 1))),
+ },
+ {
+ b.Stage(ast::PipelineStage::kCompute),
+ b.WorkgroupSize(1),
+ });
- auto* expect = R"(
+ auto* expect = R"(
struct S {
@size(16)
padding : u32,
@@ -164,48 +156,44 @@
}
)";
- auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
- Program(std::move(b)));
+ auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
- EXPECT_EQ(expect, str(got));
+ EXPECT_EQ(expect, str(got));
}
TEST_F(DecomposeStridedMatrixTest, ReadUniformMatrix_DefaultStride) {
- // struct S {
- // @offset(16) @stride(8)
- // @internal(ignore_stride_attribute)
- // m : mat2x2<f32>,
- // };
- // @group(0) @binding(0) var<uniform> s : S;
- //
- // @stage(compute) @workgroup_size(1)
- // fn f() {
- // let x : mat2x2<f32> = s.m;
- // }
- ProgramBuilder b;
- auto* S = b.Structure(
- "S",
- {
- b.Member(
- "m", b.ty.mat2x2<f32>(),
- {
- b.create<ast::StructMemberOffsetAttribute>(16),
- b.create<ast::StrideAttribute>(8),
- b.Disable(ast::DisabledValidation::kIgnoreStrideAttribute),
- }),
- });
- b.Global("s", b.ty.Of(S), ast::StorageClass::kUniform,
- b.GroupAndBinding(0, 0));
- b.Func("f", {}, b.ty.void_(),
- {
- b.Decl(b.Let("x", b.ty.mat2x2<f32>(), b.MemberAccessor("s", "m"))),
- },
- {
- b.Stage(ast::PipelineStage::kCompute),
- b.WorkgroupSize(1),
- });
+ // struct S {
+ // @offset(16) @stride(8)
+ // @internal(ignore_stride_attribute)
+ // m : mat2x2<f32>,
+ // };
+ // @group(0) @binding(0) var<uniform> s : S;
+ //
+ // @stage(compute) @workgroup_size(1)
+ // fn f() {
+ // let x : mat2x2<f32> = s.m;
+ // }
+ ProgramBuilder b;
+ auto* S = b.Structure(
+ "S", {
+ b.Member("m", b.ty.mat2x2<f32>(),
+ {
+ b.create<ast::StructMemberOffsetAttribute>(16),
+ b.create<ast::StrideAttribute>(8),
+ b.Disable(ast::DisabledValidation::kIgnoreStrideAttribute),
+ }),
+ });
+ b.Global("s", b.ty.Of(S), ast::StorageClass::kUniform, b.GroupAndBinding(0, 0));
+ b.Func("f", {}, b.ty.void_(),
+ {
+ b.Decl(b.Let("x", b.ty.mat2x2<f32>(), b.MemberAccessor("s", "m"))),
+ },
+ {
+ b.Stage(ast::PipelineStage::kCompute),
+ b.WorkgroupSize(1),
+ });
- auto* expect = R"(
+ auto* expect = R"(
struct S {
@size(16)
padding : u32,
@@ -221,48 +209,45 @@
}
)";
- auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
- Program(std::move(b)));
+ auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
- EXPECT_EQ(expect, str(got));
+ EXPECT_EQ(expect, str(got));
}
TEST_F(DecomposeStridedMatrixTest, ReadStorageMatrix) {
- // struct S {
- // @offset(8) @stride(32)
- // @internal(ignore_stride_attribute)
- // m : mat2x2<f32>,
- // };
- // @group(0) @binding(0) var<storage, read_write> s : S;
- //
- // @stage(compute) @workgroup_size(1)
- // fn f() {
- // let x : mat2x2<f32> = s.m;
- // }
- ProgramBuilder b;
- auto* S = b.Structure(
- "S",
- {
- b.Member(
- "m", b.ty.mat2x2<f32>(),
- {
- b.create<ast::StructMemberOffsetAttribute>(8),
- b.create<ast::StrideAttribute>(32),
- b.Disable(ast::DisabledValidation::kIgnoreStrideAttribute),
- }),
- });
- b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage,
- ast::Access::kReadWrite, b.GroupAndBinding(0, 0));
- b.Func("f", {}, b.ty.void_(),
- {
- b.Decl(b.Let("x", b.ty.mat2x2<f32>(), b.MemberAccessor("s", "m"))),
- },
- {
- b.Stage(ast::PipelineStage::kCompute),
- b.WorkgroupSize(1),
- });
+ // struct S {
+ // @offset(8) @stride(32)
+ // @internal(ignore_stride_attribute)
+ // m : mat2x2<f32>,
+ // };
+ // @group(0) @binding(0) var<storage, read_write> s : S;
+ //
+ // @stage(compute) @workgroup_size(1)
+ // fn f() {
+ // let x : mat2x2<f32> = s.m;
+ // }
+ ProgramBuilder b;
+ auto* S = b.Structure(
+ "S", {
+ b.Member("m", b.ty.mat2x2<f32>(),
+ {
+ b.create<ast::StructMemberOffsetAttribute>(8),
+ b.create<ast::StrideAttribute>(32),
+ b.Disable(ast::DisabledValidation::kIgnoreStrideAttribute),
+ }),
+ });
+ b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage, ast::Access::kReadWrite,
+ b.GroupAndBinding(0, 0));
+ b.Func("f", {}, b.ty.void_(),
+ {
+ b.Decl(b.Let("x", b.ty.mat2x2<f32>(), b.MemberAccessor("s", "m"))),
+ },
+ {
+ b.Stage(ast::PipelineStage::kCompute),
+ b.WorkgroupSize(1),
+ });
- auto* expect = R"(
+ auto* expect = R"(
struct S {
@size(8)
padding : u32,
@@ -281,49 +266,45 @@
}
)";
- auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
- Program(std::move(b)));
+ auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
- EXPECT_EQ(expect, str(got));
+ EXPECT_EQ(expect, str(got));
}
TEST_F(DecomposeStridedMatrixTest, ReadStorageColumn) {
- // struct S {
- // @offset(16) @stride(32)
- // @internal(ignore_stride_attribute)
- // m : mat2x2<f32>,
- // };
- // @group(0) @binding(0) var<storage, read_write> s : S;
- //
- // @stage(compute) @workgroup_size(1)
- // fn f() {
- // let x : vec2<f32> = s.m[1];
- // }
- ProgramBuilder b;
- auto* S = b.Structure(
- "S",
- {
- b.Member(
- "m", b.ty.mat2x2<f32>(),
- {
- b.create<ast::StructMemberOffsetAttribute>(16),
- b.create<ast::StrideAttribute>(32),
- b.Disable(ast::DisabledValidation::kIgnoreStrideAttribute),
- }),
- });
- b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage,
- ast::Access::kReadWrite, b.GroupAndBinding(0, 0));
- b.Func("f", {}, b.ty.void_(),
- {
- b.Decl(b.Let("x", b.ty.vec2<f32>(),
- b.IndexAccessor(b.MemberAccessor("s", "m"), 1))),
- },
- {
- b.Stage(ast::PipelineStage::kCompute),
- b.WorkgroupSize(1),
- });
+ // struct S {
+ // @offset(16) @stride(32)
+ // @internal(ignore_stride_attribute)
+ // m : mat2x2<f32>,
+ // };
+ // @group(0) @binding(0) var<storage, read_write> s : S;
+ //
+ // @stage(compute) @workgroup_size(1)
+ // fn f() {
+ // let x : vec2<f32> = s.m[1];
+ // }
+ ProgramBuilder b;
+ auto* S = b.Structure(
+ "S", {
+ b.Member("m", b.ty.mat2x2<f32>(),
+ {
+ b.create<ast::StructMemberOffsetAttribute>(16),
+ b.create<ast::StrideAttribute>(32),
+ b.Disable(ast::DisabledValidation::kIgnoreStrideAttribute),
+ }),
+ });
+ b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage, ast::Access::kReadWrite,
+ b.GroupAndBinding(0, 0));
+ b.Func("f", {}, b.ty.void_(),
+ {
+ b.Decl(b.Let("x", b.ty.vec2<f32>(), b.IndexAccessor(b.MemberAccessor("s", "m"), 1))),
+ },
+ {
+ b.Stage(ast::PipelineStage::kCompute),
+ b.WorkgroupSize(1),
+ });
- auto* expect = R"(
+ auto* expect = R"(
struct S {
@size(16)
padding : u32,
@@ -338,50 +319,46 @@
}
)";
- auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
- Program(std::move(b)));
+ auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
- EXPECT_EQ(expect, str(got));
+ EXPECT_EQ(expect, str(got));
}
TEST_F(DecomposeStridedMatrixTest, WriteStorageMatrix) {
- // struct S {
- // @offset(8) @stride(32)
- // @internal(ignore_stride_attribute)
- // m : mat2x2<f32>,
- // };
- // @group(0) @binding(0) var<storage, read_write> s : S;
- //
- // @stage(compute) @workgroup_size(1)
- // fn f() {
- // s.m = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
- // }
- ProgramBuilder b;
- auto* S = b.Structure(
- "S",
- {
- b.Member(
- "m", b.ty.mat2x2<f32>(),
- {
- b.create<ast::StructMemberOffsetAttribute>(8),
- b.create<ast::StrideAttribute>(32),
- b.Disable(ast::DisabledValidation::kIgnoreStrideAttribute),
- }),
- });
- b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage,
- ast::Access::kReadWrite, b.GroupAndBinding(0, 0));
- b.Func("f", {}, b.ty.void_(),
- {
- b.Assign(b.MemberAccessor("s", "m"),
- b.mat2x2<f32>(b.vec2<f32>(1.0f, 2.0f),
- b.vec2<f32>(3.0f, 4.0f))),
- },
- {
- b.Stage(ast::PipelineStage::kCompute),
- b.WorkgroupSize(1),
- });
+ // struct S {
+ // @offset(8) @stride(32)
+ // @internal(ignore_stride_attribute)
+ // m : mat2x2<f32>,
+ // };
+ // @group(0) @binding(0) var<storage, read_write> s : S;
+ //
+ // @stage(compute) @workgroup_size(1)
+ // fn f() {
+ // s.m = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
+ // }
+ ProgramBuilder b;
+ auto* S = b.Structure(
+ "S", {
+ b.Member("m", b.ty.mat2x2<f32>(),
+ {
+ b.create<ast::StructMemberOffsetAttribute>(8),
+ b.create<ast::StrideAttribute>(32),
+ b.Disable(ast::DisabledValidation::kIgnoreStrideAttribute),
+ }),
+ });
+ b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage, ast::Access::kReadWrite,
+ b.GroupAndBinding(0, 0));
+ b.Func("f", {}, b.ty.void_(),
+ {
+ b.Assign(b.MemberAccessor("s", "m"),
+ b.mat2x2<f32>(b.vec2<f32>(1.0f, 2.0f), b.vec2<f32>(3.0f, 4.0f))),
+ },
+ {
+ b.Stage(ast::PipelineStage::kCompute),
+ b.WorkgroupSize(1),
+ });
- auto* expect = R"(
+ auto* expect = R"(
struct S {
@size(8)
padding : u32,
@@ -400,49 +377,45 @@
}
)";
- auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
- Program(std::move(b)));
+ auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
- EXPECT_EQ(expect, str(got));
+ EXPECT_EQ(expect, str(got));
}
TEST_F(DecomposeStridedMatrixTest, WriteStorageColumn) {
- // struct S {
- // @offset(8) @stride(32)
- // @internal(ignore_stride_attribute)
- // m : mat2x2<f32>,
- // };
- // @group(0) @binding(0) var<storage, read_write> s : S;
- //
- // @stage(compute) @workgroup_size(1)
- // fn f() {
- // s.m[1] = vec2<f32>(1.0, 2.0);
- // }
- ProgramBuilder b;
- auto* S = b.Structure(
- "S",
- {
- b.Member(
- "m", b.ty.mat2x2<f32>(),
- {
- b.create<ast::StructMemberOffsetAttribute>(8),
- b.create<ast::StrideAttribute>(32),
- b.Disable(ast::DisabledValidation::kIgnoreStrideAttribute),
- }),
- });
- b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage,
- ast::Access::kReadWrite, b.GroupAndBinding(0, 0));
- b.Func("f", {}, b.ty.void_(),
- {
- b.Assign(b.IndexAccessor(b.MemberAccessor("s", "m"), 1),
- b.vec2<f32>(1.0f, 2.0f)),
- },
- {
- b.Stage(ast::PipelineStage::kCompute),
- b.WorkgroupSize(1),
- });
+ // struct S {
+ // @offset(8) @stride(32)
+ // @internal(ignore_stride_attribute)
+ // m : mat2x2<f32>,
+ // };
+ // @group(0) @binding(0) var<storage, read_write> s : S;
+ //
+ // @stage(compute) @workgroup_size(1)
+ // fn f() {
+ // s.m[1] = vec2<f32>(1.0, 2.0);
+ // }
+ ProgramBuilder b;
+ auto* S = b.Structure(
+ "S", {
+ b.Member("m", b.ty.mat2x2<f32>(),
+ {
+ b.create<ast::StructMemberOffsetAttribute>(8),
+ b.create<ast::StrideAttribute>(32),
+ b.Disable(ast::DisabledValidation::kIgnoreStrideAttribute),
+ }),
+ });
+ b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage, ast::Access::kReadWrite,
+ b.GroupAndBinding(0, 0));
+ b.Func("f", {}, b.ty.void_(),
+ {
+ b.Assign(b.IndexAccessor(b.MemberAccessor("s", "m"), 1), b.vec2<f32>(1.0f, 2.0f)),
+ },
+ {
+ b.Stage(ast::PipelineStage::kCompute),
+ b.WorkgroupSize(1),
+ });
- auto* expect = R"(
+ auto* expect = R"(
struct S {
@size(8)
padding : u32,
@@ -457,63 +430,58 @@
}
)";
- auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
- Program(std::move(b)));
+ auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
- EXPECT_EQ(expect, str(got));
+ EXPECT_EQ(expect, str(got));
}
TEST_F(DecomposeStridedMatrixTest, ReadWriteViaPointerLets) {
- // struct S {
- // @offset(8) @stride(32)
- // @internal(ignore_stride_attribute)
- // m : mat2x2<f32>,
- // };
- // @group(0) @binding(0) var<storage, read_write> s : S;
- //
- // @stage(compute) @workgroup_size(1)
- // fn f() {
- // let a = &s.m;
- // let b = &*&*(a);
- // let x = *b;
- // let y = (*b)[1];
- // let z = x[1];
- // (*b) = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
- // (*b)[1] = vec2<f32>(5.0, 6.0);
- // }
- ProgramBuilder b;
- auto* S = b.Structure(
- "S",
- {
- b.Member(
- "m", b.ty.mat2x2<f32>(),
- {
- b.create<ast::StructMemberOffsetAttribute>(8),
- b.create<ast::StrideAttribute>(32),
- b.Disable(ast::DisabledValidation::kIgnoreStrideAttribute),
- }),
- });
- b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage,
- ast::Access::kReadWrite, b.GroupAndBinding(0, 0));
- b.Func(
- "f", {}, b.ty.void_(),
- {
- b.Decl(b.Let("a", nullptr, b.AddressOf(b.MemberAccessor("s", "m")))),
- b.Decl(b.Let("b", nullptr,
- b.AddressOf(b.Deref(b.AddressOf(b.Deref("a")))))),
- b.Decl(b.Let("x", nullptr, b.Deref("b"))),
- b.Decl(b.Let("y", nullptr, b.IndexAccessor(b.Deref("b"), 1))),
- b.Decl(b.Let("z", nullptr, b.IndexAccessor("x", 1))),
- b.Assign(b.Deref("b"), b.mat2x2<f32>(b.vec2<f32>(1.0f, 2.0f),
- b.vec2<f32>(3.0f, 4.0f))),
- b.Assign(b.IndexAccessor(b.Deref("b"), 1), b.vec2<f32>(5.0f, 6.0f)),
- },
- {
- b.Stage(ast::PipelineStage::kCompute),
- b.WorkgroupSize(1),
- });
+ // struct S {
+ // @offset(8) @stride(32)
+ // @internal(ignore_stride_attribute)
+ // m : mat2x2<f32>,
+ // };
+ // @group(0) @binding(0) var<storage, read_write> s : S;
+ //
+ // @stage(compute) @workgroup_size(1)
+ // fn f() {
+ // let a = &s.m;
+ // let b = &*&*(a);
+ // let x = *b;
+ // let y = (*b)[1];
+ // let z = x[1];
+ // (*b) = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
+ // (*b)[1] = vec2<f32>(5.0, 6.0);
+ // }
+ ProgramBuilder b;
+ auto* S = b.Structure(
+ "S", {
+ b.Member("m", b.ty.mat2x2<f32>(),
+ {
+ b.create<ast::StructMemberOffsetAttribute>(8),
+ b.create<ast::StrideAttribute>(32),
+ b.Disable(ast::DisabledValidation::kIgnoreStrideAttribute),
+ }),
+ });
+ b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage, ast::Access::kReadWrite,
+ b.GroupAndBinding(0, 0));
+ b.Func(
+ "f", {}, b.ty.void_(),
+ {
+ b.Decl(b.Let("a", nullptr, b.AddressOf(b.MemberAccessor("s", "m")))),
+ b.Decl(b.Let("b", nullptr, b.AddressOf(b.Deref(b.AddressOf(b.Deref("a")))))),
+ b.Decl(b.Let("x", nullptr, b.Deref("b"))),
+ b.Decl(b.Let("y", nullptr, b.IndexAccessor(b.Deref("b"), 1))),
+ b.Decl(b.Let("z", nullptr, b.IndexAccessor("x", 1))),
+ b.Assign(b.Deref("b"), b.mat2x2<f32>(b.vec2<f32>(1.0f, 2.0f), b.vec2<f32>(3.0f, 4.0f))),
+ b.Assign(b.IndexAccessor(b.Deref("b"), 1), b.vec2<f32>(5.0f, 6.0f)),
+ },
+ {
+ b.Stage(ast::PipelineStage::kCompute),
+ b.WorkgroupSize(1),
+ });
- auto* expect = R"(
+ auto* expect = R"(
struct S {
@size(8)
padding : u32,
@@ -540,47 +508,44 @@
}
)";
- auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
- Program(std::move(b)));
+ auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
- EXPECT_EQ(expect, str(got));
+ EXPECT_EQ(expect, str(got));
}
TEST_F(DecomposeStridedMatrixTest, ReadPrivateMatrix) {
- // struct S {
- // @offset(8) @stride(32)
- // @internal(ignore_stride_attribute)
- // m : mat2x2<f32>,
- // };
- // var<private> s : S;
- //
- // @stage(compute) @workgroup_size(1)
- // fn f() {
- // let x : mat2x2<f32> = s.m;
- // }
- ProgramBuilder b;
- auto* S = b.Structure(
- "S",
- {
- b.Member(
- "m", b.ty.mat2x2<f32>(),
- {
- b.create<ast::StructMemberOffsetAttribute>(8),
- b.create<ast::StrideAttribute>(32),
- b.Disable(ast::DisabledValidation::kIgnoreStrideAttribute),
- }),
- });
- b.Global("s", b.ty.Of(S), ast::StorageClass::kPrivate);
- b.Func("f", {}, b.ty.void_(),
- {
- b.Decl(b.Let("x", b.ty.mat2x2<f32>(), b.MemberAccessor("s", "m"))),
- },
- {
- b.Stage(ast::PipelineStage::kCompute),
- b.WorkgroupSize(1),
- });
+ // struct S {
+ // @offset(8) @stride(32)
+ // @internal(ignore_stride_attribute)
+ // m : mat2x2<f32>,
+ // };
+ // var<private> s : S;
+ //
+ // @stage(compute) @workgroup_size(1)
+ // fn f() {
+ // let x : mat2x2<f32> = s.m;
+ // }
+ ProgramBuilder b;
+ auto* S = b.Structure(
+ "S", {
+ b.Member("m", b.ty.mat2x2<f32>(),
+ {
+ b.create<ast::StructMemberOffsetAttribute>(8),
+ b.create<ast::StrideAttribute>(32),
+ b.Disable(ast::DisabledValidation::kIgnoreStrideAttribute),
+ }),
+ });
+ b.Global("s", b.ty.Of(S), ast::StorageClass::kPrivate);
+ b.Func("f", {}, b.ty.void_(),
+ {
+ b.Decl(b.Let("x", b.ty.mat2x2<f32>(), b.MemberAccessor("s", "m"))),
+ },
+ {
+ b.Stage(ast::PipelineStage::kCompute),
+ b.WorkgroupSize(1),
+ });
- auto* expect = R"(
+ auto* expect = R"(
struct S {
@size(8)
padding : u32,
@@ -596,49 +561,45 @@
}
)";
- auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
- Program(std::move(b)));
+ auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
- EXPECT_EQ(expect, str(got));
+ EXPECT_EQ(expect, str(got));
}
TEST_F(DecomposeStridedMatrixTest, WritePrivateMatrix) {
- // struct S {
- // @offset(8) @stride(32)
- // @internal(ignore_stride_attribute)
- // m : mat2x2<f32>,
- // };
- // var<private> s : S;
- //
- // @stage(compute) @workgroup_size(1)
- // fn f() {
- // s.m = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
- // }
- ProgramBuilder b;
- auto* S = b.Structure(
- "S",
- {
- b.Member(
- "m", b.ty.mat2x2<f32>(),
- {
- b.create<ast::StructMemberOffsetAttribute>(8),
- b.create<ast::StrideAttribute>(32),
- b.Disable(ast::DisabledValidation::kIgnoreStrideAttribute),
- }),
- });
- b.Global("s", b.ty.Of(S), ast::StorageClass::kPrivate);
- b.Func("f", {}, b.ty.void_(),
- {
- b.Assign(b.MemberAccessor("s", "m"),
- b.mat2x2<f32>(b.vec2<f32>(1.0f, 2.0f),
- b.vec2<f32>(3.0f, 4.0f))),
- },
- {
- b.Stage(ast::PipelineStage::kCompute),
- b.WorkgroupSize(1),
- });
+ // struct S {
+ // @offset(8) @stride(32)
+ // @internal(ignore_stride_attribute)
+ // m : mat2x2<f32>,
+ // };
+ // var<private> s : S;
+ //
+ // @stage(compute) @workgroup_size(1)
+ // fn f() {
+ // s.m = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
+ // }
+ ProgramBuilder b;
+ auto* S = b.Structure(
+ "S", {
+ b.Member("m", b.ty.mat2x2<f32>(),
+ {
+ b.create<ast::StructMemberOffsetAttribute>(8),
+ b.create<ast::StrideAttribute>(32),
+ b.Disable(ast::DisabledValidation::kIgnoreStrideAttribute),
+ }),
+ });
+ b.Global("s", b.ty.Of(S), ast::StorageClass::kPrivate);
+ b.Func("f", {}, b.ty.void_(),
+ {
+ b.Assign(b.MemberAccessor("s", "m"),
+ b.mat2x2<f32>(b.vec2<f32>(1.0f, 2.0f), b.vec2<f32>(3.0f, 4.0f))),
+ },
+ {
+ b.Stage(ast::PipelineStage::kCompute),
+ b.WorkgroupSize(1),
+ });
- auto* expect = R"(
+ auto* expect = R"(
struct S {
@size(8)
padding : u32,
@@ -654,10 +615,9 @@
}
)";
- auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
- Program(std::move(b)));
+ auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
- EXPECT_EQ(expect, str(got));
+ EXPECT_EQ(expect, str(got));
}
} // namespace
