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