diff --git a/src/tint/cmd/tint/main.cc b/src/tint/cmd/tint/main.cc
index ff85ff7..474088e 100644
--- a/src/tint/cmd/tint/main.cc
+++ b/src/tint/cmd/tint/main.cc
@@ -145,6 +145,8 @@
 #if TINT_BUILD_HLSL_WRITER
 constexpr uint32_t kMinShaderModelForDXC = 60u;
 constexpr uint32_t kMaxSupportedShaderModelForDXC = 66u;
+constexpr uint32_t kMinShaderModelForDP4aInHLSL = 64u;
+constexpr uint32_t kMinShaderModelForPackUnpack4x8InHLSL = 66u;
 #endif  // TINT_BUILD_HLSL_WRITER
 
 struct Options {
@@ -917,17 +919,9 @@
         tint::cmd::GenerateExternalTextureBindings(program);
     gen_options.root_constant_binding_point = options.hlsl_root_constant_binding_point;
     gen_options.pixel_local_options = options.pixel_local_options;
-    if (must_validate_dxc) {
-        constexpr uint32_t kMinShaderModelForDP4aInHLSL = 64u;
-        constexpr uint32_t kMinShaderModelForPackUnPack4x8InHLSL = 66u;
-        gen_options.polyfill_dot_4x8_packed =
-            options.hlsl_shader_model < kMinShaderModelForDP4aInHLSL;
-        gen_options.polyfill_pack_unpack_4x8 =
-            options.hlsl_shader_model < kMinShaderModelForPackUnPack4x8InHLSL;
-    } else {
-        gen_options.polyfill_dot_4x8_packed = true;
-        gen_options.polyfill_pack_unpack_4x8 = true;
-    }
+    gen_options.polyfill_dot_4x8_packed = options.hlsl_shader_model < kMinShaderModelForDP4aInHLSL;
+    gen_options.polyfill_pack_unpack_4x8 =
+        options.hlsl_shader_model < kMinShaderModelForPackUnpack4x8InHLSL;
     auto result = tint::hlsl::writer::Generate(program, gen_options);
     if (result != tint::Success) {
         tint::cmd::PrintWGSL(std::cerr, program);
diff --git a/src/tint/lang/spirv/reader/ast_parser/parser_test.cc b/src/tint/lang/spirv/reader/ast_parser/parser_test.cc
index 404e925..5805440 100644
--- a/src/tint/lang/spirv/reader/ast_parser/parser_test.cc
+++ b/src/tint/lang/spirv/reader/ast_parser/parser_test.cc
@@ -103,7 +103,7 @@
 %_ptr_Input_vec3u = OpTypePointer Input %vec3u
      %uint_0 = OpConstant %uint 0
      %uint_2 = OpConstant %uint 2
-     %uint_8 = OpConstant %uint 8
+   %uint_264 = OpConstant %uint 264
        %wgid = OpVariable %_ptr_Input_vec3u Input
        %void = OpTypeVoid
        %bool = OpTypeBool
@@ -116,7 +116,7 @@
                OpSelectionMerge %merge None
                OpBranchConditional %condition %true_branch %merge
 %true_branch = OpLabel
-               OpControlBarrier %uint_2 %uint_2 %uint_8
+               OpControlBarrier %uint_2 %uint_2 %uint_264
                OpBranch %merge
       %merge = OpLabel
                OpReturn
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.bazel b/src/tint/lang/spirv/reader/parser/BUILD.bazel
index 2bfaad6..22ae352 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.bazel
+++ b/src/tint/lang/spirv/reader/parser/BUILD.bazel
@@ -79,6 +79,7 @@
   name = "test",
   alwayslink = True,
   srcs = [
+    "composite_test.cc",
     "constant_test.cc",
     "function_test.cc",
     "helper_test.h",
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.cmake b/src/tint/lang/spirv/reader/parser/BUILD.cmake
index c8abe87..52c7c0c 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.cmake
+++ b/src/tint/lang/spirv/reader/parser/BUILD.cmake
@@ -85,6 +85,7 @@
 # Condition: TINT_BUILD_SPV_READER
 ################################################################################
 tint_add_target(tint_lang_spirv_reader_parser_test test
+  lang/spirv/reader/parser/composite_test.cc
   lang/spirv/reader/parser/constant_test.cc
   lang/spirv/reader/parser/function_test.cc
   lang/spirv/reader/parser/helper_test.h
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.gn b/src/tint/lang/spirv/reader/parser/BUILD.gn
index 26c2954..f6c3fba 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.gn
+++ b/src/tint/lang/spirv/reader/parser/BUILD.gn
@@ -86,6 +86,7 @@
   if (tint_build_spv_reader) {
     tint_unittests_source_set("unittests") {
       sources = [
+        "composite_test.cc",
         "constant_test.cc",
         "function_test.cc",
         "helper_test.h",
diff --git a/src/tint/lang/spirv/reader/parser/composite_test.cc b/src/tint/lang/spirv/reader/parser/composite_test.cc
new file mode 100644
index 0000000..6a0ce9f
--- /dev/null
+++ b/src/tint/lang/spirv/reader/parser/composite_test.cc
@@ -0,0 +1,459 @@
+// 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/parser/helper_test.h"
+
+namespace tint::spirv::reader {
+namespace {
+
+TEST_F(SpirvParserTest, CompositeConstruct_Vector) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+        %u32 = OpTypeInt 32 0
+      %vec4u = OpTypeVector %u32 4
+    %ep_type = OpTypeFunction %void
+    %fn_type = OpTypeFunction %vec4u
+      %u32_1 = OpConstant %u32 1
+      %u32_2 = OpConstant %u32 2
+      %u32_3 = OpConstant %u32 3
+      %u32_4 = OpConstant %u32 4
+       %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+               OpReturn
+               OpFunctionEnd
+
+        %foo = OpFunction %vec4u None %fn_type
+  %foo_start = OpLabel
+        %vec = OpCompositeConstruct %vec4u %u32_1 %u32_2 %u32_3 %u32_4
+               OpReturnValue %vec
+               OpFunctionEnd
+)",
+              R"(
+  %b2 = block {
+    %3:vec4<u32> = construct 1u, 2u, 3u, 4u
+    ret %3
+  }
+)");
+}
+
+TEST_F(SpirvParserTest, CompositeConstruct_Matrix) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+        %f32 = OpTypeFloat 32
+      %vec3f = OpTypeVector %f32 3
+     %mat4x3 = OpTypeMatrix %vec3f 4
+    %ep_type = OpTypeFunction %void
+    %fn_type = OpTypeFunction %mat4x3
+      %f32_1 = OpConstant %f32 1.0
+      %f32_2 = OpConstant %f32 2.0
+      %f32_3 = OpConstant %f32 3.0
+       %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+               OpReturn
+               OpFunctionEnd
+
+        %foo = OpFunction %mat4x3 None %fn_type
+  %foo_start = OpLabel
+      %col_0 = OpCompositeConstruct %vec3f %f32_1 %f32_2 %f32_3
+      %col_1 = OpCompositeConstruct %vec3f %f32_2 %f32_3 %f32_1
+      %col_2 = OpCompositeConstruct %vec3f %f32_3 %f32_2 %f32_1
+      %col_3 = OpCompositeConstruct %vec3f %f32_3 %f32_3 %f32_3
+        %mat = OpCompositeConstruct %mat4x3 %col_0 %col_1 %col_2 %col_3
+               OpReturnValue %mat
+               OpFunctionEnd
+)",
+              R"(
+  %b2 = block {
+    %3:vec3<f32> = construct 1.0f, 2.0f, 3.0f
+    %4:vec3<f32> = construct 2.0f, 3.0f, 1.0f
+    %5:vec3<f32> = construct 3.0f, 2.0f, 1.0f
+    %6:vec3<f32> = construct 3.0f, 3.0f, 3.0f
+    %7:mat4x3<f32> = construct %3, %4, %5, %6
+    ret %7
+  }
+)");
+}
+
+TEST_F(SpirvParserTest, CompositeConstruct_Array) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+        %u32 = OpTypeInt 32 0
+    %ep_type = OpTypeFunction %void
+      %u32_1 = OpConstant %u32 1
+      %u32_2 = OpConstant %u32 2
+      %u32_3 = OpConstant %u32 3
+      %u32_4 = OpConstant %u32 4
+     %arr_ty = OpTypeArray %u32 %u32_4
+    %fn_type = OpTypeFunction %arr_ty
+       %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+               OpReturn
+               OpFunctionEnd
+
+        %foo = OpFunction %arr_ty None %fn_type
+  %foo_start = OpLabel
+        %arr = OpCompositeConstruct %arr_ty %u32_1 %u32_2 %u32_3 %u32_4
+               OpReturnValue %arr
+               OpFunctionEnd
+)",
+              R"(
+  %b2 = block {
+    %3:array<u32, 4> = construct 1u, 2u, 3u, 4u
+    ret %3
+  }
+)");
+}
+
+TEST_F(SpirvParserTest, CompositeConstruct_ArrayOfVec) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+        %u32 = OpTypeInt 32 0
+      %vec4u = OpTypeVector %u32 4
+    %ep_type = OpTypeFunction %void
+      %u32_1 = OpConstant %u32 1
+      %u32_2 = OpConstant %u32 2
+      %u32_3 = OpConstant %u32 3
+      %u32_4 = OpConstant %u32 4
+     %arr_ty = OpTypeArray %vec4u %u32_4
+    %fn_type = OpTypeFunction %arr_ty
+       %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+               OpReturn
+               OpFunctionEnd
+
+       %foo = OpFunction %arr_ty None %fn_type
+ %foo_start = OpLabel
+       %el_0 = OpCompositeConstruct %vec4u %u32_1 %u32_2 %u32_3 %u32_4
+       %el_1 = OpCompositeConstruct %vec4u %u32_2 %u32_3 %u32_4 %u32_1
+       %el_2 = OpCompositeConstruct %vec4u %u32_3 %u32_4 %u32_1 %u32_2
+       %el_3 = OpCompositeConstruct %vec4u %u32_4 %u32_1 %u32_2 %u32_3
+        %arr = OpCompositeConstruct %arr_ty %el_0 %el_1 %el_2 %el_3
+               OpReturnValue %arr
+               OpFunctionEnd
+)",
+              R"(
+  %b2 = block {
+    %3:vec4<u32> = construct 1u, 2u, 3u, 4u
+    %4:vec4<u32> = construct 2u, 3u, 4u, 1u
+    %5:vec4<u32> = construct 3u, 4u, 1u, 2u
+    %6:vec4<u32> = construct 4u, 1u, 2u, 3u
+    %7:array<vec4<u32>, 4> = construct %3, %4, %5, %6
+    ret %7
+  }
+)");
+}
+
+TEST_F(SpirvParserTest, CompositeConstruct_Struct) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+        %u32 = OpTypeInt 32 0
+      %vec4u = OpTypeVector %u32 4
+    %ep_type = OpTypeFunction %void
+      %u32_1 = OpConstant %u32 1
+      %u32_2 = OpConstant %u32 2
+      %u32_3 = OpConstant %u32 3
+      %u32_4 = OpConstant %u32 4
+     %arr_ty = OpTypeArray %vec4u %u32_2
+     %str_ty = OpTypeStruct %u32 %vec4u %arr_ty
+    %fn_type = OpTypeFunction %str_ty
+       %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+               OpReturn
+               OpFunctionEnd
+
+        %foo = OpFunction %str_ty None %fn_type
+  %foo_start = OpLabel
+       %el_0 = OpCompositeConstruct %vec4u %u32_1 %u32_2 %u32_3 %u32_4
+       %el_1 = OpCompositeConstruct %vec4u %u32_2 %u32_3 %u32_4 %u32_1
+       %el_2 = OpCompositeConstruct %vec4u %u32_3 %u32_4 %u32_1 %u32_2
+        %arr = OpCompositeConstruct %arr_ty %el_1 %el_2
+        %str = OpCompositeConstruct %str_ty %u32_4 %el_0 %arr
+               OpReturnValue %str
+               OpFunctionEnd
+)",
+              R"(
+tint_symbol_3 = struct @align(16) {
+  tint_symbol:u32 @offset(0)
+  tint_symbol_1:vec4<u32> @offset(16)
+  tint_symbol_2:array<vec4<u32>, 2> @offset(32)
+}
+
+%main = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+  %b1 = block {
+    ret
+  }
+}
+%2 = func():tint_symbol_3 -> %b2 {
+  %b2 = block {
+    %3:vec4<u32> = construct 1u, 2u, 3u, 4u
+    %4:vec4<u32> = construct 2u, 3u, 4u, 1u
+    %5:vec4<u32> = construct 3u, 4u, 1u, 2u
+    %6:array<vec4<u32>, 2> = construct %4, %5
+    %7:tint_symbol_3 = construct 4u, %3, %6
+    ret %7
+  }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, CompositeExtract_Vector) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+        %u32 = OpTypeInt 32 0
+      %vec4u = OpTypeVector %u32 4
+    %ep_type = OpTypeFunction %void
+    %fn_type = OpTypeFunction %u32 %vec4u
+       %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+               OpReturn
+               OpFunctionEnd
+
+        %foo = OpFunction %u32 None %fn_type
+        %vec = OpFunctionParameter %vec4u
+  %foo_start = OpLabel
+    %extract = OpCompositeExtract %u32 %vec 2
+               OpReturnValue %extract
+               OpFunctionEnd
+)",
+              R"(
+%2 = func(%3:vec4<u32>):u32 -> %b2 {
+  %b2 = block {
+    %4:u32 = access %3, 2u
+    ret %4
+  }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, CompositeExtract_MatrixColumn) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+        %u32 = OpTypeInt 32 0
+        %f32 = OpTypeFloat 32
+      %vec3f = OpTypeVector %f32 3
+     %mat4x3 = OpTypeMatrix %vec3f 4
+    %ep_type = OpTypeFunction %void
+    %fn_type = OpTypeFunction %vec3f %mat4x3
+       %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+               OpReturn
+               OpFunctionEnd
+
+        %foo = OpFunction %vec3f None %fn_type
+        %mat = OpFunctionParameter %mat4x3
+  %foo_start = OpLabel
+    %extract = OpCompositeExtract %vec3f %mat 2
+               OpReturnValue %extract
+               OpFunctionEnd
+)",
+              R"(
+%2 = func(%3:mat4x3<f32>):vec3<f32> -> %b2 {
+  %b2 = block {
+    %4:vec3<f32> = access %3, 2u
+    ret %4
+  }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, CompositeExtract_MatrixElement) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+        %u32 = OpTypeInt 32 0
+        %f32 = OpTypeFloat 32
+      %vec3f = OpTypeVector %f32 3
+     %mat4x3 = OpTypeMatrix %vec3f 4
+    %ep_type = OpTypeFunction %void
+    %fn_type = OpTypeFunction %f32 %mat4x3
+       %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+               OpReturn
+               OpFunctionEnd
+
+        %foo = OpFunction %f32 None %fn_type
+        %mat = OpFunctionParameter %mat4x3
+  %foo_start = OpLabel
+    %extract = OpCompositeExtract %f32 %mat 2 1
+               OpReturnValue %extract
+               OpFunctionEnd
+)",
+              R"(
+%2 = func(%3:mat4x3<f32>):f32 -> %b2 {
+  %b2 = block {
+    %4:f32 = access %3, 2u, 1u
+    ret %4
+  }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, CompositeExtract_Array) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+        %u32 = OpTypeInt 32 0
+      %u32_4 = OpConstant %u32 4
+     %arr_ty = OpTypeArray %u32 %u32_4
+    %ep_type = OpTypeFunction %void
+    %fn_type = OpTypeFunction %u32 %arr_ty
+       %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+               OpReturn
+               OpFunctionEnd
+
+        %foo = OpFunction %u32 None %fn_type
+        %arr = OpFunctionParameter %arr_ty
+  %foo_start = OpLabel
+    %extract = OpCompositeExtract %u32 %arr 2
+               OpReturnValue %extract
+               OpFunctionEnd
+)",
+              R"(
+%2 = func(%3:array<u32, 4>):u32 -> %b2 {
+  %b2 = block {
+    %4:u32 = access %3, 2u
+    ret %4
+  }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, CompositeExtract_ArrayOfVec) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+        %u32 = OpTypeInt 32 0
+      %vec3u = OpTypeVector %u32 3
+      %u32_4 = OpConstant %u32 4
+     %arr_ty = OpTypeArray %vec3u %u32_4
+    %ep_type = OpTypeFunction %void
+    %fn_type = OpTypeFunction %u32 %arr_ty
+       %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+               OpReturn
+               OpFunctionEnd
+
+        %foo = OpFunction %u32 None %fn_type
+        %arr = OpFunctionParameter %arr_ty
+  %foo_start = OpLabel
+    %extract = OpCompositeExtract %u32 %arr 1 2
+               OpReturnValue %extract
+               OpFunctionEnd
+)",
+              R"(
+%2 = func(%3:array<vec3<u32>, 4>):u32 -> %b2 {
+  %b2 = block {
+    %4:u32 = access %3, 1u, 2u
+    ret %4
+  }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, CompositeExtract_Struct) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+        %u32 = OpTypeInt 32 0
+     %str_ty = OpTypeStruct %u32 %u32
+    %ep_type = OpTypeFunction %void
+    %fn_type = OpTypeFunction %u32 %str_ty
+       %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+               OpReturn
+               OpFunctionEnd
+
+        %foo = OpFunction %u32 None %fn_type
+        %str = OpFunctionParameter %str_ty
+  %foo_start = OpLabel
+    %extract = OpCompositeExtract %u32 %str 1
+               OpReturnValue %extract
+               OpFunctionEnd
+)",
+              R"(
+tint_symbol_2 = struct @align(4) {
+  tint_symbol:u32 @offset(0)
+  tint_symbol_1:u32 @offset(4)
+}
+
+%main = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+  %b1 = block {
+    ret
+  }
+}
+%2 = func(%3:tint_symbol_2):u32 -> %b2 {
+  %b2 = block {
+    %4:u32 = access %3, 1u
+    ret %4
+  }
+}
+)");
+}
+
+}  // namespace
+}  // namespace tint::spirv::reader
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index 56a3d47..6fe4eb9 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -420,6 +420,12 @@
                 case spv::Op::OpInBoundsAccessChain:
                     EmitAccess(inst);
                     break;
+                case spv::Op::OpCompositeConstruct:
+                    EmitConstruct(inst);
+                    break;
+                case spv::Op::OpCompositeExtract:
+                    EmitCompositeExtract(inst);
+                    break;
                 case spv::Op::OpFunctionCall:
                     EmitFunctionCall(inst);
                     break;
@@ -457,6 +463,27 @@
         Emit(access, inst.result_id());
     }
 
+    /// @param inst the SPIR-V instruction for OpCompositeExtract
+    void EmitCompositeExtract(const spvtools::opt::Instruction& inst) {
+        Vector<core::ir::Value*, 4> indices;
+        for (uint32_t i = 3; i < inst.NumOperandWords(); i++) {
+            indices.Push(b_.Constant(u32(inst.GetSingleWordOperand(i))));
+        }
+        auto* object = Value(inst.GetSingleWordOperand(2));
+        auto* access = b_.Access(Type(inst.type_id()), object, std::move(indices));
+        Emit(access, inst.result_id());
+    }
+
+    /// @param inst the SPIR-V instruction for OpCompositeConstruct
+    void EmitConstruct(const spvtools::opt::Instruction& inst) {
+        Vector<core::ir::Value*, 4> values;
+        for (uint32_t i = 2; i < inst.NumOperandWords(); i++) {
+            values.Push(Value(inst.GetSingleWordOperand(i)));
+        }
+        auto* construct = b_.Construct(Type(inst.type_id()), std::move(values));
+        Emit(construct, inst.result_id());
+    }
+
     /// @param inst the SPIR-V instruction for OpFunctionCall
     void EmitFunctionCall(const spvtools::opt::Instruction& inst) {
         // TODO(crbug.com/tint/1907): Capture result.