[spirv-reader][ir] Support names.
Add support for `OpName` and `OpMemberName` decorations.
Bug: 391482385, 391482466
Change-Id: I5640494ebd30100d0f6141b99c1df1ca7ca9666c
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/226157
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.bazel b/src/tint/lang/spirv/reader/parser/BUILD.bazel
index 165aba1..17609ee 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.bazel
+++ b/src/tint/lang/spirv/reader/parser/BUILD.bazel
@@ -91,6 +91,7 @@
"import_test.cc",
"memory_test.cc",
"misc_test.cc",
+ "name_test.cc",
"struct_test.cc",
"var_test.cc",
],
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.cmake b/src/tint/lang/spirv/reader/parser/BUILD.cmake
index c9d1414..9fedf1b 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.cmake
+++ b/src/tint/lang/spirv/reader/parser/BUILD.cmake
@@ -100,6 +100,7 @@
lang/spirv/reader/parser/import_test.cc
lang/spirv/reader/parser/memory_test.cc
lang/spirv/reader/parser/misc_test.cc
+ lang/spirv/reader/parser/name_test.cc
lang/spirv/reader/parser/struct_test.cc
lang/spirv/reader/parser/var_test.cc
)
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.gn b/src/tint/lang/spirv/reader/parser/BUILD.gn
index 55175db..937008a 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.gn
+++ b/src/tint/lang/spirv/reader/parser/BUILD.gn
@@ -99,6 +99,7 @@
"import_test.cc",
"memory_test.cc",
"misc_test.cc",
+ "name_test.cc",
"struct_test.cc",
"var_test.cc",
]
diff --git a/src/tint/lang/spirv/reader/parser/name_test.cc b/src/tint/lang/spirv/reader/parser/name_test.cc
new file mode 100644
index 0000000..9bc1f4b3
--- /dev/null
+++ b/src/tint/lang/spirv/reader/parser/name_test.cc
@@ -0,0 +1,314 @@
+// Copyright 2025 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, Name_Set) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %1 "my_name"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %ep_type = OpTypeFunction %void
+ %one = OpConstant %i32 1
+
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ %1 = OpCopyObject %i32 %one
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %my_name:i32 = let 1i
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Name_IgnoreEmptyName) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %1 ""
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %ep_type = OpTypeFunction %void
+ %one = OpConstant %i32 1
+
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ %1 = OpCopyObject %i32 %one
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = let 1i
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Name_Duplicate) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %1 "vanilla"
+ OpName %2 "vanilla"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %ep_type = OpTypeFunction %void
+ %one = OpConstant %i32 1
+ %two = OpConstant %i32 2
+
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ %1 = OpCopyObject %i32 %one
+ %2 = OpCopyObject %i32 %two
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %vanilla:i32 = let 1i
+ %vanilla_1:i32 = let 2i # %vanilla_1: 'vanilla'
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Name_MemberName) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %S "Desert"
+ OpMemberName %S 0 "strawberry"
+ OpMemberName %S 1 "vanilla"
+ OpMemberName %S 2 "chocolate"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %S = OpTypeStruct %i32 %i32 %i32
+ %one = OpConstant %i32 1
+ %d = OpConstantComposite %S %one %one %one
+ %ep_type = OpTypeFunction %void
+
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ %2 = OpCopyObject %S %d
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+Desert = struct @align(4) {
+ strawberry:i32 @offset(0)
+ vanilla:i32 @offset(4)
+ chocolate:i32 @offset(8)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:Desert = let Desert(1i)
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Name_IgnoreEmptyMemberName) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %S "Desert"
+ OpMemberName %S 0 "strawberry"
+ OpMemberName %S 1 ""
+ OpMemberName %S 2 "chocolate"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %S = OpTypeStruct %i32 %i32 %i32
+ %one = OpConstant %i32 1
+ %d = OpConstantComposite %S %one %one %one
+ %ep_type = OpTypeFunction %void
+
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ %2 = OpCopyObject %S %d
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+Desert = struct @align(4) {
+ strawberry:i32 @offset(0)
+ tint_symbol:i32 @offset(4)
+ chocolate:i32 @offset(8)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:Desert = let Desert(1i)
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Name_StructSkipFirstMember) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %S "Desert"
+ OpMemberName %S 1 "strawberry"
+ OpMemberName %S 2 "chocolate"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %S = OpTypeStruct %i32 %i32 %i32
+ %one = OpConstant %i32 1
+ %d = OpConstantComposite %S %one %one %one
+ %ep_type = OpTypeFunction %void
+
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ %2 = OpCopyObject %S %d
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+Desert = struct @align(4) {
+ tint_symbol:i32 @offset(0)
+ strawberry:i32 @offset(4)
+ chocolate:i32 @offset(8)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:Desert = let Desert(1i)
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Name_StructSkipMember) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %S "Desert"
+ OpMemberName %S 0 "strawberry"
+ OpMemberName %S 2 "chocolate"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %S = OpTypeStruct %i32 %i32 %i32
+ %one = OpConstant %i32 1
+ %d = OpConstantComposite %S %one %one %one
+ %ep_type = OpTypeFunction %void
+
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ %2 = OpCopyObject %S %d
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+Desert = struct @align(4) {
+ strawberry:i32 @offset(0)
+ tint_symbol:i32 @offset(4)
+ chocolate:i32 @offset(8)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:Desert = let Desert(1i)
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Name_StructSkipLast) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %S "Desert"
+ OpMemberName %S 0 "strawberry"
+ OpMemberName %S 1 "chocolate"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %S = OpTypeStruct %i32 %i32 %i32
+ %one = OpConstant %i32 1
+ %d = OpConstantComposite %S %one %one %one
+ %ep_type = OpTypeFunction %void
+
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ %2 = OpCopyObject %S %d
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+Desert = struct @align(4) {
+ strawberry:i32 @offset(0)
+ chocolate:i32 @offset(4)
+ tint_symbol:i32 @offset(8)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:Desert = let Desert(1i)
+ ret
+ }
+}
+)");
+}
+
+} // 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 3323527..6e81010 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -30,6 +30,7 @@
#include <algorithm>
#include <memory>
#include <string>
+#include <unordered_map>
#include <unordered_set>
#include <utility>
#include <vector>
@@ -112,15 +113,48 @@
EmitModuleScopeVariables();
}
+ RegisterNames();
+
EmitFunctions();
EmitEntryPointAttributes();
// TODO(crbug.com/tint/1907): Handle annotation instructions.
- // TODO(crbug.com/tint/1907): Handle names.
return std::move(ir_);
}
+ void RegisterNames() {
+ // Register names from OpName
+ for (const auto& inst : spirv_context_->debugs2()) {
+ switch (inst.opcode()) {
+ case spv::Op::OpName: {
+ const auto name = inst.GetInOperand(1).AsString();
+ if (!name.empty()) {
+ id_to_name_[inst.GetSingleWordInOperand(0)] = name;
+ }
+ break;
+ }
+ case spv::Op::OpMemberName: {
+ const auto name = inst.GetInOperand(2).AsString();
+ if (!name.empty()) {
+ uint32_t struct_id = inst.GetSingleWordInOperand(0);
+ uint32_t member_idx = inst.GetSingleWordInOperand(1);
+ auto iter = struct_to_member_names_.insert({struct_id, {}});
+ auto& members = (*(iter.first)).second;
+
+ if (members.size() < (member_idx + 1)) {
+ members.resize(member_idx + 1);
+ }
+ members[member_idx] = name;
+ }
+ break;
+ }
+ default:
+ break;
+ }
+ }
+ }
+
/// @param sc a SPIR-V storage class
/// @returns the Tint address space for a SPIR-V storage class
core::AddressSpace AddressSpace(spv::StorageClass sc) {
@@ -282,6 +316,15 @@
TINT_ICE() << "empty structures are not supported";
}
+ auto* type_mgr = spirv_context_->get_type_mgr();
+ auto struct_id = type_mgr->GetId(struct_ty);
+
+ std::vector<std::string>* member_names = nullptr;
+ auto struct_to_member_iter = struct_to_member_names_.find(struct_id);
+ if (struct_to_member_iter != struct_to_member_names_.end()) {
+ member_names = &((*struct_to_member_iter).second);
+ }
+
// Build a list of struct members.
uint32_t current_size = 0u;
Vector<core::type::StructMember*, 4> members;
@@ -335,15 +378,28 @@
}
}
- // TODO(crbug.com/tint/1907): Use OpMemberName to name it.
- members.Push(ty_.Get<core::type::StructMember>(ir_.symbols.New(), member_ty, i, offset,
- align, member_ty->Size(),
- std::move(attributes)));
+ Symbol name;
+ if (member_names && member_names->size() > i) {
+ auto n = (*member_names)[i];
+ name = ir_.symbols.New(n);
+ } else {
+ name = ir_.symbols.New();
+ }
+
+ members.Push(ty_.Get<core::type::StructMember>(
+ name, member_ty, i, offset, align, member_ty->Size(), std::move(attributes)));
current_size = offset + member_ty->Size();
}
- // TODO(crbug.com/tint/1907): Use OpName to name it.
- return ty_.Struct(ir_.symbols.New(), std::move(members));
+
+ Symbol name;
+ auto iter = id_to_name_.find(struct_id);
+ if (iter != id_to_name_.end()) {
+ name = ir_.symbols.New(iter->second);
+ } else {
+ name = ir_.symbols.New();
+ }
+ return ty_.Struct(name, std::move(members));
}
/// @param id a SPIR-V result ID for a function declaration instruction
@@ -437,6 +493,11 @@
current_block_->Append(inst);
TINT_ASSERT(inst->Results().Length() == 1u);
AddValue(result_id, inst->Result(0));
+
+ auto iter = id_to_name_.find(result_id);
+ if (iter != id_to_name_.end()) {
+ ir_.SetName(inst, iter->second);
+ }
}
/// Emit an instruction to the current block.
@@ -1191,6 +1252,11 @@
// The set of IDs of imports that are ignored. For example, any "NonSemanticInfo." import is
// ignored.
std::unordered_set<uint32_t> ignored_imports_;
+
+ // Map of SPIR-V IDs to string names
+ std::unordered_map<uint32_t, std::string> id_to_name_;
+ // Map of SPIR-V Struct IDs to a list of member string names
+ std::unordered_map<uint32_t, std::vector<std::string>> struct_to_member_names_;
};
} // namespace
diff --git a/src/tint/lang/spirv/reader/reader_test.cc b/src/tint/lang/spirv/reader/reader_test.cc
index eb351b2..3b475a3 100644
--- a/src/tint/lang/spirv/reader/reader_test.cc
+++ b/src/tint/lang/spirv/reader/reader_test.cc
@@ -317,15 +317,15 @@
)");
ASSERT_EQ(got, Success);
EXPECT_EQ(got, R"(
-tint_symbol_2 = struct @align(16) {
- tint_symbol:vec4<f32> @offset(0)
- tint_symbol_1:array<f32, 1> @offset(16)
+VertexOutputs = struct @align(16) {
+ position:vec4<f32> @offset(0)
+ clipDistance:array<f32, 1> @offset(16)
}
-tint_symbol_6 = struct @align(16) {
- tint_symbol_3:vec4<f32> @offset(0), @builtin(position)
- tint_symbol_4:array<f32, 1> @offset(16), @builtin(clip_distances)
- tint_symbol_5:f32 @offset(20), @builtin(__point_size)
+tint_symbol_3 = struct @align(16) {
+ tint_symbol:vec4<f32> @offset(0), @builtin(position)
+ tint_symbol_1:array<f32, 1> @offset(16), @builtin(clip_distances)
+ tint_symbol_2:f32 @offset(20), @builtin(__point_size)
}
$B1: { # root
@@ -334,14 +334,14 @@
%3:ptr<private, f32, read_write> = var
}
-%4 = func():tint_symbol_2 {
+%4 = func():VertexOutputs {
$B2: {
- ret tint_symbol_2(vec4<f32>(0.0f), array<f32, 1>(0.0f))
+ ret VertexOutputs(vec4<f32>(0.0f), array<f32, 1>(0.0f))
}
}
%main_inner = func():void {
$B3: {
- %6:tint_symbol_2 = call %4
+ %6:VertexOutputs = call %4
%7:vec4<f32> = access %6, 0u
store %1, %7
%8:array<f32, 1> = access %6, 1u
@@ -350,13 +350,13 @@
ret
}
}
-%main = @vertex func():tint_symbol_6 {
+%main = @vertex func():tint_symbol_3 {
$B4: {
%10:void = call %main_inner
%11:vec4<f32> = load %1
%12:array<f32, 1> = load %2
%13:f32 = load %3
- %14:tint_symbol_6 = construct %11, %12, %13
+ %14:tint_symbol_3 = construct %11, %12, %13
ret %14
}
}
@@ -568,14 +568,14 @@
ASSERT_EQ(got, Success);
EXPECT_EQ(got, R"(
-tint_symbol_2 = struct @align(16) {
- tint_symbol:vec4<f32> @offset(0)
- tint_symbol_1:vec4<f32> @offset(16)
+FragOutput = struct @align(16) {
+ color:vec4<f32> @offset(0)
+ blend:vec4<f32> @offset(16)
}
-tint_symbol_5 = struct @align(16) {
- tint_symbol_3:vec4<f32> @offset(0), @location(0), @blend_src(0)
- tint_symbol_4:vec4<f32> @offset(16), @location(0), @blend_src(1)
+tint_symbol_2 = struct @align(16) {
+ tint_symbol:vec4<f32> @offset(0), @location(0), @blend_src(0)
+ tint_symbol_1:vec4<f32> @offset(16), @location(0), @blend_src(1)
}
$B1: { # root
@@ -583,20 +583,20 @@
%2:ptr<private, vec4<f32>, read_write> = var
}
-%3 = func():tint_symbol_2 {
+%3 = func():FragOutput {
$B2: {
- %4:ptr<function, tint_symbol_2, read_write> = var, tint_symbol_2(vec4<f32>(0.0f))
- %5:ptr<function, vec4<f32>, read_write> = access %4, 0u
+ %output:ptr<function, FragOutput, read_write> = var, FragOutput(vec4<f32>(0.0f))
+ %5:ptr<function, vec4<f32>, read_write> = access %output, 0u
store %5, vec4<f32>(0.5f, 0.5f, 0.5f, 1.0f)
- %6:ptr<function, vec4<f32>, read_write> = access %4, 1u
+ %6:ptr<function, vec4<f32>, read_write> = access %output, 1u
store %6, vec4<f32>(0.5f, 0.5f, 0.5f, 1.0f)
- %7:tint_symbol_2 = load %4
+ %7:FragOutput = load %output
ret %7
}
}
%frag_main_inner = func():void {
$B3: {
- %9:tint_symbol_2 = call %3
+ %9:FragOutput = call %3
%10:vec4<f32> = access %9, 0u
store %1, %10
%11:vec4<f32> = access %9, 1u
@@ -604,12 +604,12 @@
ret
}
}
-%frag_main = @fragment func():tint_symbol_5 {
+%frag_main = @fragment func():tint_symbol_2 {
$B4: {
%13:void = call %frag_main_inner
%14:vec4<f32> = load %1
%15:vec4<f32> = load %2
- %16:tint_symbol_5 = construct %14, %15
+ %16:tint_symbol_2 = construct %14, %15
ret %16
}
}
