[spirv-reader][ir] Support `OpSpecConstantComposite`.
The `OpSpecConstantComposite` is the only way to build a composite
object from spec constants. In WGSL there is no equivalent, so if the
`OpSpecConstantComposite` has an attached `SpecId`, we error. In the
case there is no `SpecId` associated, we can turn the
`OpSpecConstantComposite` into a `Construct`. This construct needs to be
emitted in the block of usage, as it can not appear in the root block.
Bug: 398008657
Change-Id: I0ab4e55957aab63597d4e4eba87382738162422f
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/245994
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index 54be1de..b7c6cb2 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -309,6 +309,35 @@
CreateOverride(inst, Value(inst.result_id()), std::nullopt);
break;
}
+ case spv::Op::OpSpecConstantComposite: {
+ auto spec_id = GetSpecId(inst);
+ if (spec_id.has_value()) {
+ TINT_ICE()
+ << "OpSpecConstantCompositeOp not supported when set with a SpecId";
+ }
+
+ auto* cnst = SpvConstant(inst.result_id());
+ if (cnst != nullptr) {
+ // The spec constant is made of literals, so it's return as a constant from
+ // SPIR-V Tools Opt. We can just ignore it and let the normal constant
+ // handling take over.
+ break;
+ }
+
+ Vector<uint32_t, 4> args;
+ args.Reserve(inst.NumInOperands());
+
+ for (uint32_t i = 0; i < inst.NumInOperands(); ++i) {
+ uint32_t id = inst.GetSingleWordInOperand(i);
+ args.Push(id);
+ }
+
+ spec_composites_.insert({inst.result_id(), SpecComposite{
+ .type = Type(inst.type_id()),
+ .args = args,
+ }});
+ break;
+ }
default:
break;
}
@@ -882,6 +911,11 @@
return false;
}
+ // Get the spirv constant for the given `id`. `nullptr` if no constant exists.
+ const spvtools::opt::analysis::Constant* SpvConstant(uint32_t id) {
+ return spirv_context_->get_constant_mgr()->FindDeclaredConstant(id);
+ }
+
/// Attempts to retrieve the current Tint IR value for `id`. This ignores scoping for the
/// variable, if it exists it's returned (or if it's constant it's created). The value will not
/// propagate up through control instructions.
@@ -894,12 +928,27 @@
return *v;
}
- if (auto* c = spirv_context_->get_constant_mgr()->FindDeclaredConstant(id)) {
+ if (auto* c = SpvConstant(id)) {
auto* val = b_.Constant(Constant(c));
values_.Add(id, val);
return val;
}
+ // If this was a spec composite, then it currently isn't in scope, so we construct
+ // a new copy and assign the constant ID to the new construct in this scope.
+ auto iter = spec_composites_.find(id);
+ if (iter != spec_composites_.end()) {
+ Vector<core::ir::Value*, 4> args;
+ for (auto arg : iter->second.args) {
+ args.Push(Value(arg));
+ }
+
+ auto* construct = b_.Construct(iter->second.type, args);
+ current_block_->Append(construct);
+ values_.Replace(id, construct->Result());
+ return construct->Result();
+ }
+
TINT_UNREACHABLE() << "missing value for result ID " << id;
}
@@ -3328,6 +3377,17 @@
// Map of certain instructions back to their originating spirv block
std::unordered_map<core::ir::Instruction*, uint32_t> inst_to_spirv_block_;
+
+ // Structure hold spec composite information
+ struct SpecComposite {
+ // The composite type
+ const core::type::Type* type;
+ // The composite arguments
+ Vector<uint32_t, 4> args;
+ };
+
+ // The set of SPIR-V IDs which map to `OpSpecConstantComposite` information
+ std::unordered_map<uint32_t, SpecComposite> spec_composites_;
};
} // namespace
diff --git a/src/tint/lang/spirv/reader/parser/var_test.cc b/src/tint/lang/spirv/reader/parser/var_test.cc
index c7f333c..d1e845c 100644
--- a/src/tint/lang/spirv/reader/parser/var_test.cc
+++ b/src/tint/lang/spirv/reader/parser/var_test.cc
@@ -1069,5 +1069,264 @@
)");
}
+// In the case of all literals, SPIR-V opt treats the `OpSpecConstantComposite` as an
+// `OpConstantComposite` so it appears in the constant manager already. This then needs no handling
+// on our side.
+TEST_F(SpirvParserTest, Var_OpSpecConstantComposite_vec2_literals) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %1 "myconst"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %v2i = OpTypeVector %i32 2
+ %one = OpConstant %i32 1
+ %two = OpConstant %i32 2
+ %1 = OpSpecConstantComposite %v2i %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %2 = OpIAdd %v2i %1 %1
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = spirv.add<i32> vec2<i32>(1i, 2i), vec2<i32>(1i, 2i)
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantComposite_vec2_SpecConstants) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %1 "myconst"
+ OpDecorate %one SpecId 1
+ OpDecorate %two SpecId 2
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %v2i = OpTypeVector %i32 2
+ %one = OpSpecConstant %i32 1
+ %two = OpSpecConstant %i32 2
+ %1 = OpSpecConstantComposite %v2i %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %2 = OpIAdd %v2i %1 %1
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:i32 = override 1i @id(1)
+ %2:i32 = override 2i @id(2)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:vec2<i32> = construct %1, %2
+ %5:vec2<i32> = spirv.add<i32> %4, %4
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantComposite_vec4_Mixed) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %1 "myconst"
+ OpDecorate %one SpecId 1
+ OpDecorate %three SpecId 3
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %v4i = OpTypeVector %i32 4
+ %one = OpSpecConstant %i32 1
+ %two = OpConstant %i32 2
+ %three = OpSpecConstant %i32 3
+ %four = OpConstant %i32 4
+ %1 = OpSpecConstantComposite %v4i %one %two %three %four
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %2 = OpIAdd %v4i %1 %1
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:i32 = override 1i @id(1)
+ %2:i32 = override 3i @id(3)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:vec4<i32> = construct %1, 2i, %2, 4i
+ %5:vec4<i32> = spirv.add<i32> %4, %4
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantComposite_mat3x4_Mixed) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %1 "myconst"
+ OpDecorate %one SpecId 1
+ OpDecorate %three SpecId 3
+ %void = OpTypeVoid
+ %f32 = OpTypeFloat 32
+ %v3f = OpTypeVector %f32 3
+ %mat4x3 = OpTypeMatrix %v3f 4
+ %one = OpSpecConstant %f32 1
+ %two = OpConstant %f32 2
+ %three = OpSpecConstant %f32 3
+ %four = OpConstant %f32 4
+ %1 = OpSpecConstantComposite %v3f %one %two %three
+ %2 = OpSpecConstantComposite %mat4x3 %1 %1 %1 %1
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %3 = OpMatrixTimesScalar %mat4x3 %2 %four
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:f32 = override 1.0f @id(1)
+ %2:f32 = override 3.0f @id(3)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:vec3<f32> = construct %1, 2.0f, %2
+ %5:mat4x3<f32> = construct %4, %4, %4, %4
+ %6:mat4x3<f32> = mul %5, 4.0f
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantComposite_array_Mixed) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %1 "myconst"
+ OpDecorate %one SpecId 1
+ OpDecorate %three SpecId 3
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %uint_2 = OpConstant %uint 2
+ %uint_3 = OpConstant %uint 3
+ %f32 = OpTypeFloat 32
+ %ary = OpTypeArray %f32 %uint_3
+ %ptr = OpTypePointer Function %f32
+ %fn_ptr = OpTypePointer Function %ary
+ %one = OpSpecConstant %f32 1
+ %two = OpConstant %f32 2
+ %three = OpSpecConstant %f32 3
+ %1 = OpSpecConstantComposite %ary %one %two %three
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %indexable = OpVariable %fn_ptr Function
+ OpStore %indexable %1
+ %20 = OpAccessChain %ptr %indexable %uint_2
+ %21 = OpLoad %f32 %20
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:f32 = override 1.0f @id(1)
+ %2:f32 = override 3.0f @id(3)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:ptr<function, array<f32, 3>, read_write> = var undef
+ %5:array<f32, 3> = construct %1, 2.0f, %2
+ store %4, %5
+ %6:ptr<function, f32, read_write> = access %4, 2u
+ %7:f32 = load %6
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantComposite_struct_Mixed) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpExtension "SPV_KHR_storage_buffer_storage_class"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %1 "main"
+ OpExecutionMode %1 LocalSize 1 1 1
+ OpDecorate %str Block
+ OpMemberDecorate %str 0 Offset 0
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %int = OpTypeInt 32 1
+ %f32 = OpTypeFloat 32
+ %str = OpTypeStruct %uint %f32
+ %ptr_str = OpTypePointer Function %str
+ %ptr_f32 = OpTypePointer Function %f32
+ %int_1 = OpConstant %int 1
+ %one = OpSpecConstant %uint 1
+ %two = OpConstant %f32 2
+ %5 = OpTypeFunction %void
+ %2 = OpSpecConstantComposite %str %one %two
+ %1 = OpFunction %void None %5
+ %11 = OpLabel
+ %b = OpVariable %ptr_str Function
+ OpStore %b %2
+ %24 = OpAccessChain %ptr_f32 %b %int_1
+ %17 = OpLoad %f32 %24
+ %25 = OpFAdd %f32 %17 %17
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+tint_symbol_2 = struct @align(4) {
+ tint_symbol:u32 @offset(0)
+ tint_symbol_1:f32 @offset(4)
+}
+
+$B1: { # root
+ %1:u32 = override 1u
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %3:ptr<function, tint_symbol_2, read_write> = var undef
+ %4:tint_symbol_2 = construct %1, 2.0f
+ store %3, %4
+ %5:ptr<function, f32, read_write> = access %3, 1i
+ %6:f32 = load %5
+ %7:f32 = add %6, %6
+ ret
+ }
+}
+)");
+}
+
} // namespace
} // namespace tint::spirv::reader
