From e1854b2d721dc07224291dcfa2c842ff749212ef Mon Sep 17 00:00:00 2001 From: dan sinclair Date: Fri, 9 Sep 2022 16:18:48 +0000 Subject: [PATCH] Remove SPIR-V override generation This CL removes the override emission from the SPIR-V backend. The override should be removed by the substitute_override transform before making it to the backend. Bug: tint:1155 Change-Id: Id00a58d497988908e15e3746ea05b57838acc8ba Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/101740 Reviewed-by: Ben Clayton Commit-Queue: Dan Sinclair Kokoro: Kokoro --- src/tint/writer/spirv/builder.cc | 169 ++---------------- src/tint/writer/spirv/builder.h | 4 +- .../spirv/builder_function_attribute_test.cc | 57 ------ .../spirv/builder_global_variable_test.cc | 140 --------------- src/tint/writer/spirv/builder_literal_test.cc | 34 ++-- src/tint/writer/spirv/scalar_constant.h | 17 +- src/tint/writer/spirv/scalar_constant_test.cc | 10 -- 7 files changed, 38 insertions(+), 393 deletions(-) diff --git a/src/tint/writer/spirv/builder.cc b/src/tint/writer/spirv/builder.cc index 3ee241da5a..f51d808a97 100644 --- a/src/tint/writer/spirv/builder.cc +++ b/src/tint/writer/spirv/builder.cc @@ -506,60 +506,13 @@ bool Builder::GenerateExecutionModes(const ast::Function* func, uint32_t id) { } else if (func->PipelineStage() == ast::PipelineStage::kCompute) { auto& wgsize = func_sem->WorkgroupSize(); - // Check if the workgroup_size uses pipeline-overridable constants. - if (wgsize[0].overridable_const || wgsize[1].overridable_const || - wgsize[2].overridable_const) { - if (has_overridable_workgroup_size_) { - // Only one stage can have a pipeline-overridable workgroup size. - // TODO(crbug.com/tint/810): Use LocalSizeId to handle this scenario. - TINT_ICE(Writer, builder_.Diagnostics()) - << "multiple stages using pipeline-overridable workgroup sizes"; - } - has_overridable_workgroup_size_ = true; - - auto* vec3_u32 = builder_.create(builder_.create(), 3u); - uint32_t vec3_u32_type_id = GenerateTypeIfNeeded(vec3_u32); - if (vec3_u32_type_id == 0) { - return 0; - } - - OperandList wgsize_ops; - auto wgsize_result = result_op(); - wgsize_ops.push_back(Operand(vec3_u32_type_id)); - wgsize_ops.push_back(wgsize_result); - - // Generate OpConstant instructions for each dimension. - for (size_t i = 0; i < 3; i++) { - auto constant = ScalarConstant::U32(wgsize[i].value); - if (wgsize[i].overridable_const) { - // Make the constant specializable. - auto* sem_const = - builder_.Sem().Get(wgsize[i].overridable_const); - if (!sem_const->Declaration()->Is()) { - TINT_ICE(Writer, builder_.Diagnostics()) - << "expected a pipeline-overridable constant"; - } - constant.is_spec_op = true; - constant.constant_id = sem_const->OverrideId().value; - } - - auto result = GenerateConstantIfNeeded(constant); - wgsize_ops.push_back(Operand(result)); - } - - // Generate the WorkgroupSize builtin. - push_type(spv::Op::OpSpecConstantComposite, wgsize_ops); - push_annot(spv::Op::OpDecorate, {wgsize_result, U32Operand(SpvDecorationBuiltIn), - U32Operand(SpvBuiltInWorkgroupSize)}); - } else { - // Not overridable, so just use OpExecutionMode LocalSize. - uint32_t x = wgsize[0].value; - uint32_t y = wgsize[1].value; - uint32_t z = wgsize[2].value; - push_execution_mode(spv::Op::OpExecutionMode, - {Operand(id), U32Operand(SpvExecutionModeLocalSize), Operand(x), - Operand(y), Operand(z)}); - } + // SubstituteOverride replaced all overrides with constants. + uint32_t x = wgsize[0].value; + uint32_t y = wgsize[1].value; + uint32_t z = wgsize[2].value; + push_execution_mode(spv::Op::OpExecutionMode, + {Operand(id), U32Operand(SpvExecutionModeLocalSize), Operand(x), + Operand(y), Operand(z)}); } for (auto builtin : func_sem->TransitivelyReferencedBuiltinVariables()) { @@ -585,7 +538,7 @@ uint32_t Builder::GenerateExpression(const ast::Expression* expr) { [&](const ast::BitcastExpression* b) { return GenerateBitcastExpression(b); }, [&](const ast::CallExpression* c) { return GenerateCallExpression(c); }, [&](const ast::IdentifierExpression* i) { return GenerateIdentifierExpression(i); }, - [&](const ast::LiteralExpression* l) { return GenerateLiteralIfNeeded(nullptr, l); }, + [&](const ast::LiteralExpression* l) { return GenerateLiteralIfNeeded(l); }, [&](const ast::MemberAccessorExpression* m) { return GenerateAccessorExpression(m); }, [&](const ast::UnaryOpExpression* u) { return GenerateUnaryOpExpression(u); }, [&](Default) { @@ -778,46 +731,6 @@ bool Builder::GenerateGlobalVariable(const ast::Variable* v) { } } - if (auto* override = v->As(); override && !override->constructor) { - // SPIR-V requires specialization constants to have initializers. - init_id = Switch( - type, // - [&](const sem::F32*) { - ast::FloatLiteralExpression l(ProgramID{}, ast::NodeID{}, Source{}, 0, - ast::FloatLiteralExpression::Suffix::kF); - return GenerateLiteralIfNeeded(override, &l); - }, - [&](const sem::U32*) { - ast::IntLiteralExpression l(ProgramID{}, ast::NodeID{}, Source{}, 0, - ast::IntLiteralExpression::Suffix::kU); - return GenerateLiteralIfNeeded(override, &l); - }, - [&](const sem::I32*) { - ast::IntLiteralExpression l(ProgramID{}, ast::NodeID{}, Source{}, 0, - ast::IntLiteralExpression::Suffix::kI); - return GenerateLiteralIfNeeded(override, &l); - }, - [&](const sem::Bool*) { - ast::BoolLiteralExpression l(ProgramID{}, ast::NodeID{}, Source{}, false); - return GenerateLiteralIfNeeded(override, &l); - }, - [&](Default) { - error_ = "invalid type for pipeline constant ID, must be scalar"; - return 0; - }); - if (init_id == 0) { - return 0; - } - } - - if (v->Is()) { - push_debug(spv::Op::OpName, - {Operand(init_id), Operand(builder_.Symbols().NameFor(v->symbol))}); - - RegisterVariable(sem, init_id); - return true; - } - auto result = result_op(); auto var_id = std::get(result); @@ -1293,15 +1206,9 @@ uint32_t Builder::GetGLSLstd450Import() { uint32_t Builder::GenerateConstructorExpression(const ast::Variable* var, const ast::Expression* expr) { - if (Is(var)) { - if (auto* literal = expr->As()) { - return GenerateLiteralIfNeeded(var, literal); - } - } else { - if (auto* sem = builder_.Sem().Get(expr)) { - if (auto constant = sem->ConstantValue()) { - return GenerateConstantIfNeeded(constant); - } + if (auto* sem = builder_.Sem().Get(expr)) { + if (auto constant = sem->ConstantValue()) { + return GenerateConstantIfNeeded(constant); } } if (auto* call = builder_.Sem().Get(expr)) { @@ -1346,24 +1253,6 @@ uint32_t Builder::GenerateTypeConstructorOrConversion(const sem::Call* call, // Generate the zero initializer if there are no values provided. if (args.IsEmpty()) { - if (global_var && global_var->Declaration()->Is()) { - auto constant_id = global_var->OverrideId().value; - if (result_type->Is()) { - return GenerateConstantIfNeeded(ScalarConstant::I32(0).AsSpecOp(constant_id)); - } - if (result_type->Is()) { - return GenerateConstantIfNeeded(ScalarConstant::U32(0).AsSpecOp(constant_id)); - } - if (result_type->Is()) { - return GenerateConstantIfNeeded(ScalarConstant::F32(0).AsSpecOp(constant_id)); - } - if (result_type->Is()) { - return GenerateConstantIfNeeded(ScalarConstant::F16(0).AsSpecOp(constant_id)); - } - if (result_type->Is()) { - return GenerateConstantIfNeeded(ScalarConstant::Bool(false).AsSpecOp(constant_id)); - } - } return GenerateConstantNullIfNeeded(result_type->UnwrapRef()); } @@ -1679,16 +1568,8 @@ uint32_t Builder::GenerateCastOrCopyOrPassthrough(const sem::Type* to_type, return result_id; } -uint32_t Builder::GenerateLiteralIfNeeded(const ast::Variable* var, - const ast::LiteralExpression* lit) { +uint32_t Builder::GenerateLiteralIfNeeded(const ast::LiteralExpression* lit) { ScalarConstant constant; - - auto* global = builder_.Sem().Get(var); - if (global && global->Declaration()->Is()) { - constant.is_spec_op = true; - constant.constant_id = global->OverrideId().value; - } - Switch( lit, [&](const ast::BoolLiteralExpression* l) { @@ -1837,42 +1718,30 @@ uint32_t Builder::GenerateConstantIfNeeded(const ScalarConstant& constant) { auto result = result_op(); auto result_id = std::get(result); - if (constant.is_spec_op) { - push_annot(spv::Op::OpDecorate, {Operand(result_id), U32Operand(SpvDecorationSpecId), - Operand(constant.constant_id)}); - } - switch (constant.kind) { case ScalarConstant::Kind::kU32: { - push_type(constant.is_spec_op ? spv::Op::OpSpecConstant : spv::Op::OpConstant, - {Operand(type_id), result, Operand(constant.value.u32)}); + push_type(spv::Op::OpConstant, {Operand(type_id), result, Operand(constant.value.u32)}); break; } case ScalarConstant::Kind::kI32: { - push_type(constant.is_spec_op ? spv::Op::OpSpecConstant : spv::Op::OpConstant, + push_type(spv::Op::OpConstant, {Operand(type_id), result, U32Operand(constant.value.i32)}); break; } case ScalarConstant::Kind::kF32: { - push_type(constant.is_spec_op ? spv::Op::OpSpecConstant : spv::Op::OpConstant, - {Operand(type_id), result, Operand(constant.value.f32)}); + push_type(spv::Op::OpConstant, {Operand(type_id), result, Operand(constant.value.f32)}); break; } case ScalarConstant::Kind::kF16: { - push_type( - constant.is_spec_op ? spv::Op::OpSpecConstant : spv::Op::OpConstant, - {Operand(type_id), result, U32Operand(constant.value.f16.bits_representation)}); + push_type(spv::Op::OpConstant, {Operand(type_id), result, + U32Operand(constant.value.f16.bits_representation)}); break; } case ScalarConstant::Kind::kBool: { if (constant.value.b) { - push_type( - constant.is_spec_op ? spv::Op::OpSpecConstantTrue : spv::Op::OpConstantTrue, - {Operand(type_id), result}); + push_type(spv::Op::OpConstantTrue, {Operand(type_id), result}); } else { - push_type( - constant.is_spec_op ? spv::Op::OpSpecConstantFalse : spv::Op::OpConstantFalse, - {Operand(type_id), result}); + push_type(spv::Op::OpConstantFalse, {Operand(type_id), result}); } break; } diff --git a/src/tint/writer/spirv/builder.h b/src/tint/writer/spirv/builder.h index a0eff107bc..ba88443480 100644 --- a/src/tint/writer/spirv/builder.h +++ b/src/tint/writer/spirv/builder.h @@ -333,10 +333,9 @@ class Builder { /// @returns the ID of the expression or 0 on failure. uint32_t GenerateConstructorExpression(const ast::Variable* var, const ast::Expression* expr); /// Generates a literal constant if needed - /// @param var the variable generated for, nullptr if no variable associated. /// @param lit the literal to generate /// @returns the ID on success or 0 on failure - uint32_t GenerateLiteralIfNeeded(const ast::Variable* var, const ast::LiteralExpression* lit); + uint32_t GenerateLiteralIfNeeded(const ast::LiteralExpression* lit); /// Generates a binary expression /// @param expr the expression to generate /// @returns the expression ID on success or 0 otherwise @@ -625,7 +624,6 @@ class Builder { std::vector merge_stack_; std::vector continue_stack_; std::unordered_set capability_set_; - bool has_overridable_workgroup_size_ = false; bool zero_initialize_workgroup_memory_ = false; struct ContinuingInfo { diff --git a/src/tint/writer/spirv/builder_function_attribute_test.cc b/src/tint/writer/spirv/builder_function_attribute_test.cc index 8825c793d5..9e6c338814 100644 --- a/src/tint/writer/spirv/builder_function_attribute_test.cc +++ b/src/tint/writer/spirv/builder_function_attribute_test.cc @@ -149,63 +149,6 @@ TEST_F(BuilderTest, Decoration_ExecutionMode_WorkgroupSize_Const) { )"); } -TEST_F(BuilderTest, Decoration_ExecutionMode_WorkgroupSize_OverridableConst) { - Override("width", ty.i32(), Construct(ty.i32(), 2_i), Id(7_u)); - Override("height", ty.i32(), Construct(ty.i32(), 3_i), Id(8_u)); - Override("depth", ty.i32(), Construct(ty.i32(), 4_i), Id(9_u)); - auto* func = Func("main", utils::Empty, ty.void_(), utils::Empty, - utils::Vector{ - WorkgroupSize("width", "height", "depth"), - Stage(ast::PipelineStage::kCompute), - }); - - spirv::Builder& b = Build(); - - ASSERT_TRUE(b.GenerateExecutionModes(func, 3)) << b.error(); - EXPECT_EQ(DumpInstructions(b.execution_modes()), ""); - EXPECT_EQ(DumpInstructions(b.types()), - R"(%2 = OpTypeInt 32 0 -%1 = OpTypeVector %2 3 -%4 = OpSpecConstant %2 2 -%5 = OpSpecConstant %2 3 -%6 = OpSpecConstant %2 4 -%3 = OpSpecConstantComposite %1 %4 %5 %6 -)"); - EXPECT_EQ(DumpInstructions(b.annots()), - R"(OpDecorate %4 SpecId 7 -OpDecorate %5 SpecId 8 -OpDecorate %6 SpecId 9 -OpDecorate %3 BuiltIn WorkgroupSize -)"); -} - -TEST_F(BuilderTest, Decoration_ExecutionMode_WorkgroupSize_LiteralAndConst) { - Override("height", ty.i32(), Construct(ty.i32(), 2_i), Id(7_u)); - GlobalConst("depth", ty.i32(), Construct(ty.i32(), 3_i)); - auto* func = Func("main", utils::Empty, ty.void_(), utils::Empty, - utils::Vector{ - WorkgroupSize(4_i, "height", "depth"), - Stage(ast::PipelineStage::kCompute), - }); - - spirv::Builder& b = Build(); - - ASSERT_TRUE(b.GenerateExecutionModes(func, 3)) << b.error(); - EXPECT_EQ(DumpInstructions(b.execution_modes()), ""); - EXPECT_EQ(DumpInstructions(b.types()), - R"(%2 = OpTypeInt 32 0 -%1 = OpTypeVector %2 3 -%4 = OpConstant %2 4 -%5 = OpSpecConstant %2 2 -%6 = OpConstant %2 3 -%3 = OpSpecConstantComposite %1 %4 %5 %6 -)"); - EXPECT_EQ(DumpInstructions(b.annots()), - R"(OpDecorate %5 SpecId 7 -OpDecorate %3 BuiltIn WorkgroupSize -)"); -} - TEST_F(BuilderTest, Decoration_ExecutionMode_MultipleFragment) { auto* func1 = Func("main1", utils::Empty, ty.void_(), utils::Empty, utils::Vector{ diff --git a/src/tint/writer/spirv/builder_global_variable_test.cc b/src/tint/writer/spirv/builder_global_variable_test.cc index dfb38e5f0a..2c53023070 100644 --- a/src/tint/writer/spirv/builder_global_variable_test.cc +++ b/src/tint/writer/spirv/builder_global_variable_test.cc @@ -249,146 +249,6 @@ OpDecorate %1 DescriptorSet 3 )"); } -TEST_F(BuilderTest, GlobalVar_Override_Bool) { - auto* v = Override("var", ty.bool_(), Expr(true), Id(1200_a)); - - spirv::Builder& b = Build(); - - EXPECT_TRUE(b.GenerateGlobalVariable(v)) << b.error(); - EXPECT_EQ(DumpInstructions(b.debug()), R"(OpName %2 "var" -)"); - EXPECT_EQ(DumpInstructions(b.annots()), R"(OpDecorate %2 SpecId 1200 -)"); - EXPECT_EQ(DumpInstructions(b.types()), R"(%1 = OpTypeBool -%2 = OpSpecConstantTrue %1 -)"); -} - -TEST_F(BuilderTest, GlobalVar_Override_Bool_ZeroValue) { - auto* v = Override("var", ty.bool_(), Construct(), Id(1200_a)); - - spirv::Builder& b = Build(); - - EXPECT_TRUE(b.GenerateGlobalVariable(v)) << b.error(); - EXPECT_EQ(DumpInstructions(b.debug()), R"(OpName %2 "var" -)"); - EXPECT_EQ(DumpInstructions(b.annots()), R"(OpDecorate %2 SpecId 1200 -)"); - EXPECT_EQ(DumpInstructions(b.types()), R"(%1 = OpTypeBool -%2 = OpSpecConstantFalse %1 -)"); -} - -TEST_F(BuilderTest, GlobalVar_Override_Bool_NoConstructor) { - auto* v = Override("var", ty.bool_(), Id(1200_a)); - - spirv::Builder& b = Build(); - - EXPECT_TRUE(b.GenerateGlobalVariable(v)) << b.error(); - EXPECT_EQ(DumpInstructions(b.debug()), R"(OpName %2 "var" -)"); - EXPECT_EQ(DumpInstructions(b.annots()), R"(OpDecorate %2 SpecId 1200 -)"); - EXPECT_EQ(DumpInstructions(b.types()), R"(%1 = OpTypeBool -%2 = OpSpecConstantFalse %1 -)"); -} - -TEST_F(BuilderTest, GlobalVar_Override_Scalar) { - auto* v = Override("var", ty.f32(), Expr(2_f), Id(0_a)); - - spirv::Builder& b = Build(); - - EXPECT_TRUE(b.GenerateGlobalVariable(v)) << b.error(); - EXPECT_EQ(DumpInstructions(b.debug()), R"(OpName %2 "var" -)"); - EXPECT_EQ(DumpInstructions(b.annots()), R"(OpDecorate %2 SpecId 0 -)"); - EXPECT_EQ(DumpInstructions(b.types()), R"(%1 = OpTypeFloat 32 -%2 = OpSpecConstant %1 2 -)"); -} - -TEST_F(BuilderTest, GlobalVar_Override_Scalar_ZeroValue) { - auto* v = Override("var", ty.f32(), Construct(), Id(0_a)); - - spirv::Builder& b = Build(); - - EXPECT_TRUE(b.GenerateGlobalVariable(v)) << b.error(); - EXPECT_EQ(DumpInstructions(b.debug()), R"(OpName %2 "var" -)"); - EXPECT_EQ(DumpInstructions(b.annots()), R"(OpDecorate %2 SpecId 0 -)"); - EXPECT_EQ(DumpInstructions(b.types()), R"(%1 = OpTypeFloat 32 -%2 = OpSpecConstant %1 0 -)"); -} - -TEST_F(BuilderTest, GlobalVar_Override_Scalar_F32_NoConstructor) { - auto* v = Override("var", ty.f32(), Id(0_a)); - - spirv::Builder& b = Build(); - - EXPECT_TRUE(b.GenerateGlobalVariable(v)) << b.error(); - EXPECT_EQ(DumpInstructions(b.debug()), R"(OpName %2 "var" -)"); - EXPECT_EQ(DumpInstructions(b.annots()), R"(OpDecorate %2 SpecId 0 -)"); - EXPECT_EQ(DumpInstructions(b.types()), R"(%1 = OpTypeFloat 32 -%2 = OpSpecConstant %1 0 -)"); -} - -TEST_F(BuilderTest, GlobalVar_Override_Scalar_I32_NoConstructor) { - auto* v = Override("var", ty.i32(), Id(0_a)); - - spirv::Builder& b = Build(); - - EXPECT_TRUE(b.GenerateGlobalVariable(v)) << b.error(); - EXPECT_EQ(DumpInstructions(b.debug()), R"(OpName %2 "var" -)"); - EXPECT_EQ(DumpInstructions(b.annots()), R"(OpDecorate %2 SpecId 0 -)"); - EXPECT_EQ(DumpInstructions(b.types()), R"(%1 = OpTypeInt 32 1 -%2 = OpSpecConstant %1 0 -)"); -} - -TEST_F(BuilderTest, GlobalVar_Override_Scalar_U32_NoConstructor) { - auto* v = Override("var", ty.u32(), Id(0_a)); - - spirv::Builder& b = Build(); - - EXPECT_TRUE(b.GenerateGlobalVariable(v)) << b.error(); - EXPECT_EQ(DumpInstructions(b.debug()), R"(OpName %2 "var" -)"); - EXPECT_EQ(DumpInstructions(b.annots()), R"(OpDecorate %2 SpecId 0 -)"); - EXPECT_EQ(DumpInstructions(b.types()), R"(%1 = OpTypeInt 32 0 -%2 = OpSpecConstant %1 0 -)"); -} - -TEST_F(BuilderTest, GlobalVar_Override_NoId) { - auto* var_a = Override("a", ty.bool_(), Expr(true), Id(0_a)); - auto* var_b = Override("b", ty.bool_(), Expr(false)); - - spirv::Builder& b = Build(); - - EXPECT_TRUE(b.GenerateGlobalVariable(var_a)) << b.error(); - EXPECT_TRUE(b.GenerateGlobalVariable(var_b)) << b.error(); - EXPECT_EQ(DumpInstructions(b.debug()), R"(OpName %2 "a" -OpName %3 "b" -)"); - EXPECT_EQ(DumpInstructions(b.annots()), R"(OpDecorate %2 SpecId 0 -OpDecorate %3 SpecId 1 -)"); - EXPECT_EQ(DumpInstructions(b.types()), R"(%1 = OpTypeBool -%2 = OpSpecConstantTrue %1 -%3 = OpSpecConstantFalse %1 -)"); -} - struct BuiltinData { ast::BuiltinValue builtin; ast::StorageClass storage; diff --git a/src/tint/writer/spirv/builder_literal_test.cc b/src/tint/writer/spirv/builder_literal_test.cc index 374c80be12..0d53237ecc 100644 --- a/src/tint/writer/spirv/builder_literal_test.cc +++ b/src/tint/writer/spirv/builder_literal_test.cc @@ -27,7 +27,7 @@ TEST_F(BuilderTest, Literal_Bool_True) { spirv::Builder& b = Build(); - auto id = b.GenerateLiteralIfNeeded(nullptr, b_true); + auto id = b.GenerateLiteralIfNeeded(b_true); ASSERT_FALSE(b.has_error()) << b.error(); EXPECT_EQ(2u, id); @@ -42,7 +42,7 @@ TEST_F(BuilderTest, Literal_Bool_False) { spirv::Builder& b = Build(); - auto id = b.GenerateLiteralIfNeeded(nullptr, b_false); + auto id = b.GenerateLiteralIfNeeded(b_false); ASSERT_FALSE(b.has_error()) << b.error(); EXPECT_EQ(2u, id); @@ -58,11 +58,11 @@ TEST_F(BuilderTest, Literal_Bool_Dedup) { spirv::Builder& b = Build(); - ASSERT_NE(b.GenerateLiteralIfNeeded(nullptr, b_true), 0u); + ASSERT_NE(b.GenerateLiteralIfNeeded(b_true), 0u); ASSERT_FALSE(b.has_error()) << b.error(); - ASSERT_NE(b.GenerateLiteralIfNeeded(nullptr, b_false), 0u); + ASSERT_NE(b.GenerateLiteralIfNeeded(b_false), 0u); ASSERT_FALSE(b.has_error()) << b.error(); - ASSERT_NE(b.GenerateLiteralIfNeeded(nullptr, b_true), 0u); + ASSERT_NE(b.GenerateLiteralIfNeeded(b_true), 0u); ASSERT_FALSE(b.has_error()) << b.error(); EXPECT_EQ(DumpInstructions(b.types()), R"(%1 = OpTypeBool @@ -76,7 +76,7 @@ TEST_F(BuilderTest, Literal_I32) { WrapInFunction(i); spirv::Builder& b = Build(); - auto id = b.GenerateLiteralIfNeeded(nullptr, i); + auto id = b.GenerateLiteralIfNeeded(i); ASSERT_FALSE(b.has_error()) << b.error(); EXPECT_EQ(2u, id); @@ -92,8 +92,8 @@ TEST_F(BuilderTest, Literal_I32_Dedup) { spirv::Builder& b = Build(); - ASSERT_NE(b.GenerateLiteralIfNeeded(nullptr, i1), 0u); - ASSERT_NE(b.GenerateLiteralIfNeeded(nullptr, i2), 0u); + ASSERT_NE(b.GenerateLiteralIfNeeded(i1), 0u); + ASSERT_NE(b.GenerateLiteralIfNeeded(i2), 0u); ASSERT_FALSE(b.has_error()) << b.error(); EXPECT_EQ(DumpInstructions(b.types()), R"(%1 = OpTypeInt 32 1 @@ -107,7 +107,7 @@ TEST_F(BuilderTest, Literal_U32) { spirv::Builder& b = Build(); - auto id = b.GenerateLiteralIfNeeded(nullptr, i); + auto id = b.GenerateLiteralIfNeeded(i); ASSERT_FALSE(b.has_error()) << b.error(); EXPECT_EQ(2u, id); @@ -123,8 +123,8 @@ TEST_F(BuilderTest, Literal_U32_Dedup) { spirv::Builder& b = Build(); - ASSERT_NE(b.GenerateLiteralIfNeeded(nullptr, i1), 0u); - ASSERT_NE(b.GenerateLiteralIfNeeded(nullptr, i2), 0u); + ASSERT_NE(b.GenerateLiteralIfNeeded(i1), 0u); + ASSERT_NE(b.GenerateLiteralIfNeeded(i2), 0u); ASSERT_FALSE(b.has_error()) << b.error(); EXPECT_EQ(DumpInstructions(b.types()), R"(%1 = OpTypeInt 32 0 @@ -138,7 +138,7 @@ TEST_F(BuilderTest, Literal_F32) { spirv::Builder& b = Build(); - auto id = b.GenerateLiteralIfNeeded(nullptr, i); + auto id = b.GenerateLiteralIfNeeded(i); ASSERT_FALSE(b.has_error()) << b.error(); EXPECT_EQ(2u, id); @@ -154,8 +154,8 @@ TEST_F(BuilderTest, Literal_F32_Dedup) { spirv::Builder& b = Build(); - ASSERT_NE(b.GenerateLiteralIfNeeded(nullptr, i1), 0u); - ASSERT_NE(b.GenerateLiteralIfNeeded(nullptr, i2), 0u); + ASSERT_NE(b.GenerateLiteralIfNeeded(i1), 0u); + ASSERT_NE(b.GenerateLiteralIfNeeded(i2), 0u); ASSERT_FALSE(b.has_error()) << b.error(); EXPECT_EQ(DumpInstructions(b.types()), R"(%1 = OpTypeFloat 32 @@ -171,7 +171,7 @@ TEST_F(BuilderTest, Literal_F16) { spirv::Builder& b = Build(); - auto id = b.GenerateLiteralIfNeeded(nullptr, i); + auto id = b.GenerateLiteralIfNeeded(i); ASSERT_FALSE(b.has_error()) << b.error(); EXPECT_EQ(2u, id); @@ -189,8 +189,8 @@ TEST_F(BuilderTest, Literal_F16_Dedup) { spirv::Builder& b = Build(); - ASSERT_NE(b.GenerateLiteralIfNeeded(nullptr, i1), 0u); - ASSERT_NE(b.GenerateLiteralIfNeeded(nullptr, i2), 0u); + ASSERT_NE(b.GenerateLiteralIfNeeded(i1), 0u); + ASSERT_NE(b.GenerateLiteralIfNeeded(i2), 0u); ASSERT_FALSE(b.has_error()) << b.error(); EXPECT_EQ(DumpInstructions(b.types()), R"(%1 = OpTypeFloat 16 diff --git a/src/tint/writer/spirv/scalar_constant.h b/src/tint/writer/spirv/scalar_constant.h index 0629d0cdc4..dd4dfd073c 100644 --- a/src/tint/writer/spirv/scalar_constant.h +++ b/src/tint/writer/spirv/scalar_constant.h @@ -111,8 +111,7 @@ struct ScalarConstant { /// @param rhs the ScalarConstant to compare against /// @returns true if this ScalarConstant is equal to `rhs` inline bool operator==(const ScalarConstant& rhs) const { - return value.u64 == rhs.value.u64 && kind == rhs.kind && is_spec_op == rhs.is_spec_op && - constant_id == rhs.constant_id; + return value.u64 == rhs.value.u64 && kind == rhs.kind; } /// Inequality operator @@ -120,24 +119,10 @@ struct ScalarConstant { /// @returns true if this ScalarConstant is not equal to `rhs` inline bool operator!=(const ScalarConstant& rhs) const { return !(*this == rhs); } - /// @returns this ScalarConstant as a specialization op with the given - /// specialization constant identifier - /// @param id the constant identifier - ScalarConstant AsSpecOp(uint32_t id) const { - auto ret = *this; - ret.is_spec_op = true; - ret.constant_id = id; - return ret; - } - /// The constant value Value value; /// The constant value kind Kind kind = Kind::kBool; - /// True if the constant is a specialization op - bool is_spec_op = false; - /// The identifier if a specialization op - uint32_t constant_id = 0; }; } // namespace tint::writer::spirv diff --git a/src/tint/writer/spirv/scalar_constant_test.cc b/src/tint/writer/spirv/scalar_constant_test.cc index b00f82ab0b..d38a050455 100644 --- a/src/tint/writer/spirv/scalar_constant_test.cc +++ b/src/tint/writer/spirv/scalar_constant_test.cc @@ -34,16 +34,6 @@ TEST_F(SpirvScalarConstantTest, Equality) { EXPECT_NE(a, b); b.value.b = true; EXPECT_EQ(a, b); - - a.is_spec_op = true; - EXPECT_NE(a, b); - b.is_spec_op = true; - EXPECT_EQ(a, b); - - a.constant_id = 3; - EXPECT_NE(a, b); - b.constant_id = 3; - EXPECT_EQ(a, b); } TEST_F(SpirvScalarConstantTest, U32) {