[spirv-reader] Support row-major matrices
Apply the `@row_major` attribute when a `RowMajor` decoration is seen,
and then use the `TransposeRowMajor` transform to transpose the
matrices.
Add unit tests that cover the full SPIR-V -> WGSL translation.
Bug: 364267168
Change-Id: I1f7d0746a9ab23c425a41a0d7c9224f6a7d41280
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/207375
Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/spirv/reader/ast_parser/BUILD.bazel b/src/tint/lang/spirv/reader/ast_parser/BUILD.bazel
index 5ebf2be..174fbd3 100644
--- a/src/tint/lang/spirv/reader/ast_parser/BUILD.bazel
+++ b/src/tint/lang/spirv/reader/ast_parser/BUILD.bazel
@@ -138,6 +138,7 @@
"named_types_test.cc",
"namer_test.cc",
"parser_test.cc",
+ "row_major_matrix_test.cc",
"spirv_tools_helpers_test.cc",
"spirv_tools_helpers_test.h",
"type_test.cc",
@@ -155,6 +156,7 @@
"//src/tint/lang/wgsl/features",
"//src/tint/lang/wgsl/program",
"//src/tint/lang/wgsl/sem",
+ "//src/tint/lang/wgsl/writer/ir_to_program",
"//src/tint/utils/containers",
"//src/tint/utils/diagnostic",
"//src/tint/utils/generator",
@@ -186,6 +188,7 @@
"//conditions:default": [],
}) + select({
":tint_build_wgsl_writer": [
+ "//src/tint/lang/wgsl/writer",
"//src/tint/lang/wgsl/writer/ast_printer",
],
"//conditions:default": [],
diff --git a/src/tint/lang/spirv/reader/ast_parser/BUILD.cmake b/src/tint/lang/spirv/reader/ast_parser/BUILD.cmake
index 76e2429..6708a76 100644
--- a/src/tint/lang/spirv/reader/ast_parser/BUILD.cmake
+++ b/src/tint/lang/spirv/reader/ast_parser/BUILD.cmake
@@ -146,6 +146,7 @@
lang/spirv/reader/ast_parser/named_types_test.cc
lang/spirv/reader/ast_parser/namer_test.cc
lang/spirv/reader/ast_parser/parser_test.cc
+ lang/spirv/reader/ast_parser/row_major_matrix_test.cc
lang/spirv/reader/ast_parser/spirv_tools_helpers_test.cc
lang/spirv/reader/ast_parser/spirv_tools_helpers_test.h
lang/spirv/reader/ast_parser/type_test.cc
@@ -164,6 +165,7 @@
tint_lang_wgsl_features
tint_lang_wgsl_program
tint_lang_wgsl_sem
+ tint_lang_wgsl_writer_ir_to_program
tint_utils_containers
tint_utils_diagnostic
tint_utils_generator
@@ -202,6 +204,7 @@
if(TINT_BUILD_WGSL_WRITER)
tint_target_add_dependencies(tint_lang_spirv_reader_ast_parser_test test
+ tint_lang_wgsl_writer
tint_lang_wgsl_writer_ast_printer
)
endif(TINT_BUILD_WGSL_WRITER)
diff --git a/src/tint/lang/spirv/reader/ast_parser/BUILD.gn b/src/tint/lang/spirv/reader/ast_parser/BUILD.gn
index 7a5c65c..4125ab5 100644
--- a/src/tint/lang/spirv/reader/ast_parser/BUILD.gn
+++ b/src/tint/lang/spirv/reader/ast_parser/BUILD.gn
@@ -146,6 +146,7 @@
"named_types_test.cc",
"namer_test.cc",
"parser_test.cc",
+ "row_major_matrix_test.cc",
"spirv_tools_helpers_test.cc",
"spirv_tools_helpers_test.h",
"type_test.cc",
@@ -165,6 +166,7 @@
"${tint_src_dir}/lang/wgsl/features",
"${tint_src_dir}/lang/wgsl/program",
"${tint_src_dir}/lang/wgsl/sem",
+ "${tint_src_dir}/lang/wgsl/writer/ir_to_program",
"${tint_src_dir}/utils/containers",
"${tint_src_dir}/utils/diagnostic",
"${tint_src_dir}/utils/generator",
@@ -200,7 +202,10 @@
}
if (tint_build_wgsl_writer) {
- deps += [ "${tint_src_dir}/lang/wgsl/writer/ast_printer" ]
+ deps += [
+ "${tint_src_dir}/lang/wgsl/writer",
+ "${tint_src_dir}/lang/wgsl/writer/ast_printer",
+ ]
}
public_configs = [ "${tint_spirv_tools_dir}/:spvtools_internal_config" ]
}
diff --git a/src/tint/lang/spirv/reader/ast_parser/ast_parser.cc b/src/tint/lang/spirv/reader/ast_parser/ast_parser.cc
index b79fb8c..a14ccaa 100644
--- a/src/tint/lang/spirv/reader/ast_parser/ast_parser.cc
+++ b/src/tint/lang/spirv/reader/ast_parser/ast_parser.cc
@@ -42,6 +42,7 @@
#include "src/tint/lang/wgsl/ast/disable_validation_attribute.h"
#include "src/tint/lang/wgsl/ast/id_attribute.h"
#include "src/tint/lang/wgsl/ast/interpolate_attribute.h"
+#include "src/tint/lang/wgsl/ast/row_major_attribute.h"
#include "src/tint/lang/wgsl/ast/unary_op_expression.h"
#include "src/tint/lang/wgsl/resolver/resolve.h"
#include "src/tint/utils/containers/unique_vector.h"
@@ -506,9 +507,12 @@
case spv::Decoration::RelaxedPrecision: // WGSL doesn't support relaxed precision.
break;
case spv::Decoration::RowMajor:
- Fail() << "WGSL does not support row-major matrices: can't "
- "translate member "
- << member_index << " of " << ShowType(struct_type_id);
+ if (!member_ty->Is<Matrix>()) {
+ Fail() << "row-major matrix layout not currently supported on type "
+ << member_ty->String();
+ break;
+ }
+ out.Add(create<ast::RowMajorAttribute>(Source{}));
break;
case spv::Decoration::MatrixStride: {
if (decoration.size() != 2) {
diff --git a/src/tint/lang/spirv/reader/ast_parser/module_var_test.cc b/src/tint/lang/spirv/reader/ast_parser/module_var_test.cc
index f92a464..77f0d06 100644
--- a/src/tint/lang/spirv/reader/ast_parser/module_var_test.cc
+++ b/src/tint/lang/spirv/reader/ast_parser/module_var_test.cc
@@ -1341,6 +1341,38 @@
)")) << module_str;
}
+TEST_F(SpvModuleScopeVarParserTest, MatrixStrideDecoration_Natural_RowMajor) {
+ auto p = parser(test::Assemble(Preamble() + FragMain() + R"(
+ OpName %myvar "myvar"
+ OpDecorate %myvar DescriptorSet 0
+ OpDecorate %myvar Binding 0
+ OpDecorate %s Block
+ OpMemberDecorate %s 0 MatrixStride 8
+ OpMemberDecorate %s 0 Offset 0
+ OpMemberDecorate %s 0 RowMajor
+ %void = OpTypeVoid
+ %voidfn = OpTypeFunction %void
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+ %m2v3float = OpTypeMatrix %v3float 2
+
+ %s = OpTypeStruct %m2v3float
+ %ptr_sb_s = OpTypePointer StorageBuffer %s
+ %myvar = OpVariable %ptr_sb_s StorageBuffer
+ )" + MainBody()));
+ ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
+ EXPECT_TRUE(p->error().empty());
+ const auto module_str = test::ToString(p->program());
+ EXPECT_THAT(module_str, HasSubstr(R"(struct S {
+ /* @offset(0) */
+ @stride(8) @internal(disable_validation__ignore_stride) @row_major
+ field0 : mat2x3f,
+}
+
+@group(0) @binding(0) var<storage, read_write> myvar : S;
+)")) << module_str;
+}
+
TEST_F(SpvModuleScopeVarParserTest, MatrixStrideDecoration) {
auto p = parser(test::Assemble(Preamble() + FragMain() + R"(
OpName %myvar "myvar"
@@ -1373,7 +1405,7 @@
)")) << module_str;
}
-TEST_F(SpvModuleScopeVarParserTest, RowMajorDecoration_IsError) {
+TEST_F(SpvModuleScopeVarParserTest, RowMajorDecoration) {
auto p = parser(test::Assemble(Preamble() + FragMain() + R"(
OpName %myvar "myvar"
OpDecorate %s Block
@@ -1389,11 +1421,17 @@
%ptr_sb_s = OpTypePointer StorageBuffer %s
%myvar = OpVariable %ptr_sb_s StorageBuffer
)" + MainBody()));
- EXPECT_FALSE(p->BuildAndParseInternalModuleExceptFunctions());
- EXPECT_THAT(
- p->error(),
- Eq(R"(WGSL does not support row-major matrices: can't translate member 0 of %3 = OpTypeStruct %8)"))
- << p->error();
+ ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
+ EXPECT_TRUE(p->error().empty());
+ const auto module_str = test::ToString(p->program());
+ EXPECT_THAT(module_str, HasSubstr(R"(struct S {
+ /* @offset(0) */
+ @row_major
+ field0 : mat3x2f,
+}
+
+var<storage, read_write> myvar : S;
+)")) << module_str;
}
TEST_F(SpvModuleScopeVarParserTest, StorageBuffer_NonWritable_Var) {
diff --git a/src/tint/lang/spirv/reader/ast_parser/parse.cc b/src/tint/lang/spirv/reader/ast_parser/parse.cc
index 48f9443..9106cbb 100644
--- a/src/tint/lang/spirv/reader/ast_parser/parse.cc
+++ b/src/tint/lang/spirv/reader/ast_parser/parse.cc
@@ -34,6 +34,7 @@
#include "src/tint/lang/spirv/reader/ast_lower/decompose_strided_matrix.h"
#include "src/tint/lang/spirv/reader/ast_lower/fold_trivial_lets.h"
#include "src/tint/lang/spirv/reader/ast_lower/pass_workgroup_id_as_argument.h"
+#include "src/tint/lang/spirv/reader/ast_lower/transpose_row_major.h"
#include "src/tint/lang/spirv/reader/ast_parser/ast_parser.h"
#include "src/tint/lang/wgsl/ast/transform/manager.h"
#include "src/tint/lang/wgsl/ast/transform/remove_unreachable_statements.h"
@@ -115,6 +116,7 @@
manager.Add<ast::transform::SimplifyPointers>();
manager.Add<FoldTrivialLets>();
manager.Add<PassWorkgroupIdAsArgument>();
+ manager.Add<TransposeRowMajor>();
manager.Add<DecomposeStridedMatrix>();
manager.Add<DecomposeStridedArray>();
manager.Add<ast::transform::RemoveUnreachableStatements>();
diff --git a/src/tint/lang/spirv/reader/ast_parser/row_major_matrix_test.cc b/src/tint/lang/spirv/reader/ast_parser/row_major_matrix_test.cc
new file mode 100644
index 0000000..ddb2ec5
--- /dev/null
+++ b/src/tint/lang/spirv/reader/ast_parser/row_major_matrix_test.cc
@@ -0,0 +1,374 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/spirv/reader/ast_parser/parse.h"
+
+#include "gmock/gmock.h"
+#include "gtest/gtest.h"
+#include "src/tint/lang/spirv/reader/ast_parser/spirv_tools_helpers_test.h"
+#include "src/tint/lang/wgsl/writer/writer.h"
+
+namespace tint::spirv::reader::ast_parser {
+namespace {
+
+class SpirvReaderRowMajorMatrixTest : public testing::Test {
+ public:
+ std::string Run(std::string decorations, std::string types, std::string body) {
+ auto spv = test::Assemble(R"(
+ OpCapability Shader
+ OpExtension "SPV_KHR_storage_buffer_storage_class"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %foo "foo"
+ OpExecutionMode %foo LocalSize 1 1 1
+ OpDecorate %buffer DescriptorSet 0
+ OpDecorate %buffer Binding 0
+ OpDecorate %S Block
+
+ )" + decorations + R"(
+
+ %u32 = OpTypeInt 32 0
+ %u32_0 = OpConstant %u32 0
+ %u32_1 = OpConstant %u32 1
+ %u32_2 = OpConstant %u32 2
+
+ %f32 = OpTypeFloat 32
+ %f32_2 = OpConstant %f32 2.0
+
+ %vec3f = OpTypeVector %f32 3
+ %mat2x3f = OpTypeMatrix %vec3f 2
+
+ )" + types + R"(
+
+ %_ptr_Storage_f32 = OpTypePointer StorageBuffer %f32
+ %_ptr_Storage_vec3f = OpTypePointer StorageBuffer %vec3f
+%_ptr_Storage_mat2x3f = OpTypePointer StorageBuffer %mat2x3f
+%_ptr_Storage_S = OpTypePointer StorageBuffer %S
+ %buffer = OpVariable %_ptr_Storage_S StorageBuffer
+
+ %void = OpTypeVoid
+ %func_type = OpTypeFunction %void
+ %foo = OpFunction %void None %func_type
+ %func_start = OpLabel
+
+ )" + body + R"(
+
+ OpReturn
+ OpFunctionEnd
+)");
+ auto program = Parse(spv, {});
+ auto errs = program.Diagnostics().Str();
+ EXPECT_TRUE(program.IsValid()) << errs;
+ EXPECT_EQ(program.Diagnostics().Count(), 0u) << errs;
+ auto result = wgsl::writer::Generate(program, {});
+ EXPECT_EQ(result, Success);
+ return "\n" + result->wgsl;
+ }
+};
+
+TEST_F(SpirvReaderRowMajorMatrixTest, LoadMatrix_DefaultStride) {
+ auto result = Run(R"(
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 0 RowMajor
+ OpMemberDecorate %S 0 MatrixStride 8
+ OpMemberDecorate %S 1 Offset 64)",
+ R"(
+ %S = OpTypeStruct %mat2x3f %vec3f)",
+ R"(
+ %m_ptr = OpAccessChain %_ptr_Storage_mat2x3f %buffer %u32_0
+ %m = OpLoad %mat2x3f %m_ptr)");
+
+ EXPECT_EQ(result, R"(
+struct S {
+ /* @offset(0) */
+ field0 : mat3x2<f32>,
+ @size(32)
+ padding_0 : u32,
+ /* @offset(64) */
+ field1 : vec3f,
+}
+
+@group(0) @binding(0) var<storage, read_write> x_2 : S;
+
+fn foo_1() {
+ let x_20 = transpose(x_2.field0);
+ return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn foo() {
+ foo_1();
+}
+)");
+}
+
+TEST_F(SpirvReaderRowMajorMatrixTest, LoadMatrix_CustomStride) {
+ auto result = Run(R"(
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 0 RowMajor
+ OpMemberDecorate %S 0 MatrixStride 32
+ OpMemberDecorate %S 1 Offset 128)",
+ R"(
+ %S = OpTypeStruct %mat2x3f %vec3f)",
+ R"(
+ %m_ptr = OpAccessChain %_ptr_Storage_mat2x3f %buffer %u32_0
+ %m = OpLoad %mat2x3f %m_ptr)");
+
+ EXPECT_EQ(result, R"(
+struct strided_arr {
+ @size(32)
+ el : vec2<f32>,
+}
+
+struct S {
+ /* @offset(0) */
+ field0 : array<strided_arr, 3u>,
+ @size(32)
+ padding_0 : u32,
+ /* @offset(128) */
+ field1 : vec3f,
+}
+
+@group(0) @binding(0) var<storage, read_write> x_2 : S;
+
+fn arr_to_mat3x2_stride_32(arr : array<strided_arr, 3u>) -> mat3x2<f32> {
+ return mat3x2<f32>(arr[0u].el, arr[1u].el, arr[2u].el);
+}
+
+fn foo_1() {
+ let x_20 = transpose(arr_to_mat3x2_stride_32(x_2.field0));
+ return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn foo() {
+ foo_1();
+}
+)");
+}
+
+TEST_F(SpirvReaderRowMajorMatrixTest, LoadColumn) {
+ auto result = Run(R"(
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 0 RowMajor
+ OpMemberDecorate %S 0 MatrixStride 8
+ OpMemberDecorate %S 1 Offset 64)",
+ R"(
+ %S = OpTypeStruct %mat2x3f %vec3f)",
+ R"(
+ %c_ptr = OpAccessChain %_ptr_Storage_vec3f %buffer %u32_0 %u32_1
+ %c = OpLoad %vec3f %c_ptr)");
+
+ EXPECT_EQ(result, R"(
+fn tint_load_row_major_column(tint_from : ptr<storage, mat3x2<f32>, read_write>, tint_idx : u32) -> vec3<f32> {
+ return vec3<f32>(tint_from[0][tint_idx], tint_from[1][tint_idx], tint_from[2][tint_idx]);
+}
+
+struct S {
+ /* @offset(0) */
+ field0 : mat3x2<f32>,
+ @size(32)
+ padding_0 : u32,
+ /* @offset(64) */
+ field1 : vec3f,
+}
+
+@group(0) @binding(0) var<storage, read_write> x_2 : S;
+
+fn foo_1() {
+ let x_20 = tint_load_row_major_column(&(x_2.field0), u32(1u));
+ return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn foo() {
+ foo_1();
+}
+)");
+}
+
+TEST_F(SpirvReaderRowMajorMatrixTest, LoadElement) {
+ auto result = Run(R"(
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 0 RowMajor
+ OpMemberDecorate %S 0 MatrixStride 8
+ OpMemberDecorate %S 1 Offset 64)",
+ R"(
+ %S = OpTypeStruct %mat2x3f %vec3f)",
+ R"(
+ %e_ptr = OpAccessChain %_ptr_Storage_f32 %buffer %u32_0 %u32_1 %u32_2
+ %e = OpLoad %f32 %e_ptr)");
+
+ EXPECT_EQ(result, R"(
+struct S {
+ /* @offset(0) */
+ field0 : mat3x2<f32>,
+ @size(32)
+ padding_0 : u32,
+ /* @offset(64) */
+ field1 : vec3f,
+}
+
+@group(0) @binding(0) var<storage, read_write> x_2 : S;
+
+fn foo_1() {
+ let x_20 = x_2.field0[2u][1u];
+ return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn foo() {
+ foo_1();
+}
+)");
+}
+
+TEST_F(SpirvReaderRowMajorMatrixTest, StoreMatrix) {
+ auto result = Run(R"(
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 0 RowMajor
+ OpMemberDecorate %S 0 MatrixStride 8
+ OpMemberDecorate %S 1 Offset 64)",
+ R"(
+ %S = OpTypeStruct %mat2x3f %vec3f)",
+ R"(
+ %m_ptr = OpAccessChain %_ptr_Storage_mat2x3f %buffer %u32_0
+ %m = OpLoad %mat2x3f %m_ptr
+ %mul = OpMatrixTimesScalar %mat2x3f %m %f32_2
+ OpStore %m_ptr %mul)");
+
+ EXPECT_EQ(result, R"(
+struct S {
+ /* @offset(0) */
+ field0 : mat3x2<f32>,
+ @size(32)
+ padding_0 : u32,
+ /* @offset(64) */
+ field1 : vec3f,
+}
+
+@group(0) @binding(0) var<storage, read_write> x_2 : S;
+
+fn foo_1() {
+ x_2.field0 = transpose((transpose(x_2.field0) * 2.0f));
+ return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn foo() {
+ foo_1();
+}
+)");
+}
+
+TEST_F(SpirvReaderRowMajorMatrixTest, StoreColumn) {
+ auto result = Run(R"(
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 0 RowMajor
+ OpMemberDecorate %S 0 MatrixStride 8
+ OpMemberDecorate %S 1 Offset 64)",
+ R"(
+ %S = OpTypeStruct %mat2x3f %vec3f)",
+ R"(
+ %c_ptr = OpAccessChain %_ptr_Storage_vec3f %buffer %u32_0 %u32_1
+ %c = OpLoad %vec3f %c_ptr
+ %mul = OpVectorTimesScalar %vec3f %c %f32_2
+ OpStore %c_ptr %mul)");
+
+ EXPECT_EQ(result, R"(
+fn tint_load_row_major_column(tint_from : ptr<storage, mat3x2<f32>, read_write>, tint_idx : u32) -> vec3<f32> {
+ return vec3<f32>(tint_from[0][tint_idx], tint_from[1][tint_idx], tint_from[2][tint_idx]);
+}
+
+fn tint_store_row_major_column(tint_to : ptr<storage, mat3x2<f32>, read_write>, tint_idx : u32, tint_col : vec3<f32>) {
+ tint_to[0][tint_idx] = tint_col[0];
+ tint_to[1][tint_idx] = tint_col[1];
+ tint_to[2][tint_idx] = tint_col[2];
+}
+
+struct S {
+ /* @offset(0) */
+ field0 : mat3x2<f32>,
+ @size(32)
+ padding_0 : u32,
+ /* @offset(64) */
+ field1 : vec3f,
+}
+
+@group(0) @binding(0) var<storage, read_write> x_2 : S;
+
+fn foo_1() {
+ tint_store_row_major_column(&(x_2.field0), u32(1u), (tint_load_row_major_column(&(x_2.field0), u32(1u)) * 2.0f));
+ return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn foo() {
+ foo_1();
+}
+)");
+}
+
+TEST_F(SpirvReaderRowMajorMatrixTest, StoreElement) {
+ auto result = Run(R"(
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 0 RowMajor
+ OpMemberDecorate %S 0 MatrixStride 8
+ OpMemberDecorate %S 1 Offset 64)",
+ R"(
+ %S = OpTypeStruct %mat2x3f %vec3f)",
+ R"(
+ %e_ptr = OpAccessChain %_ptr_Storage_f32 %buffer %u32_0 %u32_1 %u32_2
+ %e = OpLoad %f32 %e_ptr
+ %mul = OpFMul %f32 %e %f32_2
+ OpStore %e_ptr %mul)");
+
+ EXPECT_EQ(result, R"(
+struct S {
+ /* @offset(0) */
+ field0 : mat3x2<f32>,
+ @size(32)
+ padding_0 : u32,
+ /* @offset(64) */
+ field1 : vec3f,
+}
+
+@group(0) @binding(0) var<storage, read_write> x_2 : S;
+
+fn foo_1() {
+ x_2.field0[2u][1u] = (x_2.field0[2u][1u] * 2.0f);
+ return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn foo() {
+ foo_1();
+}
+)");
+}
+
+} // namespace
+} // namespace tint::spirv::reader::ast_parser
