Import Tint changes from Dawn
Changes:
- 1a2543e754a31ed889566b9bb93d6bce83ea9eb7 [tint][cmd] Don't vary HLSL output based on validation mode by Ben Clayton <bclayton@google.com>
- 0e15f4ab994611a50eb63ff35c8d2a430f5f358a [spirv-reader] Handle OpCompositeExtract by James Price <jrprice@google.com>
- 383c99157767e8a9de5cef59fc783a20b6235302 [spirv-reader] Handle OpCompositeConstruct by James Price <jrprice@google.com>
- a158d5164ebc9406e3e331052f6ddf1dfad6c76c spirv-reader: Fix invalid memory semantics in test by James Price <jrprice@google.com>
GitOrigin-RevId: 1a2543e754a31ed889566b9bb93d6bce83ea9eb7
Change-Id: Ie963cf5d8dc9f396dbd503317082f14f330e59a5
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/170520
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
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.
