[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