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 <bclayton@google.com> Commit-Queue: Dan Sinclair <dsinclair@chromium.org> Kokoro: Kokoro <noreply+kokoro@google.com>
This commit is contained in:
dan sinclair
Dawn LUCI CQ
parent
f6a9404978
commit
e1854b2d72
|
|
@ -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<sem::Vector>(builder_.create<sem::U32>(), 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<sem::GlobalVariable>(wgsize[i].overridable_const);
|
||||
if (!sem_const->Declaration()->Is<ast::Override>()) {
|
||||
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<ast::Override>(); 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<ast::Override>()) {
|
||||
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<uint32_t>(result);
|
||||
|
||||
|
|
@ -1293,15 +1206,9 @@ uint32_t Builder::GetGLSLstd450Import() {
|
|||
|
||||
uint32_t Builder::GenerateConstructorExpression(const ast::Variable* var,
|
||||
const ast::Expression* expr) {
|
||||
if (Is<ast::Override>(var)) {
|
||||
if (auto* literal = expr->As<ast::LiteralExpression>()) {
|
||||
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<sem::Call>(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<ast::Override>()) {
|
||||
auto constant_id = global_var->OverrideId().value;
|
||||
if (result_type->Is<sem::I32>()) {
|
||||
return GenerateConstantIfNeeded(ScalarConstant::I32(0).AsSpecOp(constant_id));
|
||||
}
|
||||
if (result_type->Is<sem::U32>()) {
|
||||
return GenerateConstantIfNeeded(ScalarConstant::U32(0).AsSpecOp(constant_id));
|
||||
}
|
||||
if (result_type->Is<sem::F32>()) {
|
||||
return GenerateConstantIfNeeded(ScalarConstant::F32(0).AsSpecOp(constant_id));
|
||||
}
|
||||
if (result_type->Is<sem::F16>()) {
|
||||
return GenerateConstantIfNeeded(ScalarConstant::F16(0).AsSpecOp(constant_id));
|
||||
}
|
||||
if (result_type->Is<sem::Bool>()) {
|
||||
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<sem::GlobalVariable>(var);
|
||||
if (global && global->Declaration()->Is<ast::Override>()) {
|
||||
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<uint32_t>(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;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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<uint32_t> merge_stack_;
|
||||
std::vector<uint32_t> continue_stack_;
|
||||
std::unordered_set<uint32_t> capability_set_;
|
||||
bool has_overridable_workgroup_size_ = false;
|
||||
bool zero_initialize_workgroup_memory_ = false;
|
||||
|
||||
struct ContinuingInfo {
|
||||
|
|
|
|||
|
|
@ -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{
|
||||
|
|
|
|||
|
|
@ -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<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_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<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_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;
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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) {
|
||||
|
|
|
|||
Loading…
Reference in New Issue