tint: Simplify workgroup size resolving
A `@workgroup_size()` value must be a constant or override expression.
There's nothing specific here about literals or variable expressions.
Remove the semantic tracking of override variables, as these can be override expressions.
The backends will require the `SubstituteOverride` transform to be run, so gut the workgroup_size override handling from the backends.
Bug: tint:1633
Change-Id: Ib3ff843fc64a3595d49223c661b4d58130c0ab30
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/100142
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
diff --git a/src/tint/inspector/inspector.cc b/src/tint/inspector/inspector.cc
index 087e786..d32ef6f 100644
--- a/src/tint/inspector/inspector.cc
+++ b/src/tint/inspector/inspector.cc
@@ -148,9 +148,9 @@
entry_point.stage = PipelineStage::kCompute;
auto wgsize = sem->WorkgroupSize();
- if (!wgsize[0].overridable_const && !wgsize[1].overridable_const &&
- !wgsize[2].overridable_const) {
- entry_point.workgroup_size = {wgsize[0].value, wgsize[1].value, wgsize[2].value};
+ if (wgsize[0].has_value() && wgsize[1].has_value() && wgsize[2].has_value()) {
+ entry_point.workgroup_size = {wgsize[0].value(), wgsize[1].value(),
+ wgsize[2].value()};
}
break;
}
@@ -849,19 +849,18 @@
auto* t = c->args[static_cast<size_t>(texture_index)];
auto* s = c->args[static_cast<size_t>(sampler_index)];
- GetOriginatingResources(
- std::array<const ast::Expression*, 2>{t, s},
- [&](std::array<const sem::GlobalVariable*, 2> globals) {
- auto texture_binding_point = globals[0]->BindingPoint();
- auto sampler_binding_point = globals[1]->BindingPoint();
+ GetOriginatingResources(std::array<const ast::Expression*, 2>{t, s},
+ [&](std::array<const sem::GlobalVariable*, 2> globals) {
+ auto texture_binding_point = globals[0]->BindingPoint();
+ auto sampler_binding_point = globals[1]->BindingPoint();
- for (auto* entry_point : entry_points) {
- const auto& ep_name =
- program_->Symbols().NameFor(entry_point->Declaration()->symbol);
- (*sampler_targets_)[ep_name].Add(
- {sampler_binding_point, texture_binding_point});
- }
- });
+ for (auto* entry_point : entry_points) {
+ const auto& ep_name = program_->Symbols().NameFor(
+ entry_point->Declaration()->symbol);
+ (*sampler_targets_)[ep_name].Add(
+ {sampler_binding_point, texture_binding_point});
+ }
+ });
}
}
diff --git a/src/tint/resolver/function_validation_test.cc b/src/tint/resolver/function_validation_test.cc
index 22f2620..cf87296 100644
--- a/src/tint/resolver/function_validation_test.cc
+++ b/src/tint/resolver/function_validation_test.cc
@@ -468,7 +468,7 @@
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_GoodType_ConstU32) {
// const x = 4u;
- // const x = 8u;
+ // const y = 8u;
// @compute @workgroup_size(x, y, 16u)
// fn main() {}
auto* x = GlobalConst("x", ty.u32(), Expr(4_u));
@@ -489,10 +489,29 @@
ASSERT_NE(sem_x, nullptr);
ASSERT_NE(sem_y, nullptr);
+ EXPECT_EQ(sem_func->WorkgroupSize(), (sem::WorkgroupSize{4u, 8u, 16u}));
+
EXPECT_TRUE(sem_func->DirectlyReferencedGlobals().Contains(sem_x));
EXPECT_TRUE(sem_func->DirectlyReferencedGlobals().Contains(sem_y));
}
+TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Cast) {
+ // @compute @workgroup_size(i32(5))
+ // fn main() {}
+ auto* func = Func("main", utils::Empty, ty.void_(), utils::Empty,
+ utils::Vector{
+ Stage(ast::PipelineStage::kCompute),
+ WorkgroupSize(Construct(Source{{12, 34}}, ty.i32(), 5_a)),
+ });
+
+ ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+ auto* sem_func = Sem().Get(func);
+
+ ASSERT_NE(sem_func, nullptr);
+ EXPECT_EQ(sem_func->WorkgroupSize(), (sem::WorkgroupSize{5u, 1u, 1u}));
+}
+
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_GoodType_I32) {
// @compute @workgroup_size(1i, 2i, 3i)
// fn main() {}
@@ -651,9 +670,10 @@
});
EXPECT_FALSE(r()->Resolve());
- EXPECT_EQ(r()->error(),
- "12:34 error: workgroup_size argument must be either a literal, constant, or "
- "overridable of type abstract-integer, i32 or u32");
+ EXPECT_EQ(
+ r()->error(),
+ "12:34 error: workgroup_size argument must be a constant or override expression of type "
+ "abstract-integer, i32 or u32");
}
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Literal_Negative) {
@@ -696,9 +716,10 @@
});
EXPECT_FALSE(r()->Resolve());
- EXPECT_EQ(r()->error(),
- "12:34 error: workgroup_size argument must be either a literal, constant, or "
- "overridable of type abstract-integer, i32 or u32");
+ EXPECT_EQ(
+ r()->error(),
+ "12:34 error: workgroup_size argument must be a constant or override expression of type "
+ "abstract-integer, i32 or u32");
}
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_Negative) {
@@ -759,8 +780,8 @@
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
- "12:34 error: workgroup_size argument must be either a literal, constant, or "
- "overridable of type abstract-integer, i32 or u32");
+ "12:34 error: workgroup_size argument must be a constant or override expression of "
+ "type abstract-integer, i32 or u32");
}
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_InvalidExpr_x) {
@@ -774,8 +795,8 @@
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
- "12:34 error: workgroup_size argument must be either a literal, constant, or "
- "overridable of type abstract-integer, i32 or u32");
+ "12:34 error: workgroup_size argument must be a constant or override expression of "
+ "type abstract-integer, i32 or u32");
}
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_InvalidExpr_y) {
@@ -789,8 +810,8 @@
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
- "12:34 error: workgroup_size argument must be either a literal, constant, or "
- "overridable of type abstract-integer, i32 or u32");
+ "12:34 error: workgroup_size argument must be a constant or override expression of "
+ "type abstract-integer, i32 or u32");
}
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_InvalidExpr_z) {
@@ -804,8 +825,8 @@
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
- "12:34 error: workgroup_size argument must be either a literal, constant, or "
- "overridable of type abstract-integer, i32 or u32");
+ "12:34 error: workgroup_size argument must be a constant or override expression of "
+ "type abstract-integer, i32 or u32");
}
TEST_F(ResolverFunctionValidationTest, ReturnIsConstructible_NonPlain) {
diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc
index d57be4c..c1920d8 100644
--- a/src/tint/resolver/resolver.cc
+++ b/src/tint/resolver/resolver.cc
@@ -1050,8 +1050,7 @@
// Set work-group size defaults.
sem::WorkgroupSize ws;
for (size_t i = 0; i < 3; i++) {
- ws[i].value = 1;
- ws[i].overridable_const = nullptr;
+ ws[i] = 1;
}
auto* attr = ast::GetAttribute<ast::WorkgroupAttribute>(func->attributes);
@@ -1064,7 +1063,7 @@
utils::Vector<const sem::Type*, 3> arg_tys;
constexpr const char* kErrBadExpr =
- "workgroup_size argument must be either a literal, constant, or overridable of type "
+ "workgroup_size argument must be a constant or override expression of type "
"abstract-integer, i32 or u32";
for (size_t i = 0; i < 3; i++) {
@@ -1084,6 +1083,12 @@
return false;
}
+ if (expr->Stage() != sem::EvaluationStage::kConstant &&
+ expr->Stage() != sem::EvaluationStage::kOverride) {
+ AddError(kErrBadExpr, value->source);
+ return false;
+ }
+
args.Push(expr);
arg_tys.Push(ty);
}
@@ -1105,47 +1110,15 @@
if (!materialized) {
return false;
}
-
- const sem::Constant* value = nullptr;
-
- if (auto* user = args[i]->As<sem::VariableUser>()) {
- // We have an variable of a module-scope constant.
- auto* decl = user->Variable()->Declaration();
- if (!decl->IsAnyOf<ast::Const, ast::Override>()) {
- AddError(kErrBadExpr, values[i]->source);
+ if (auto* value = materialized->ConstantValue()) {
+ if (value->As<AInt>() < 1) {
+ AddError("workgroup_size argument must be at least 1", values[i]->source);
return false;
}
- // Capture the constant if it is pipeline-overridable.
- if (decl->Is<ast::Override>()) {
- ws[i].overridable_const = decl;
- }
-
- if (decl->constructor) {
- value = sem_.Get(decl->constructor)->ConstantValue();
- } else {
- // No constructor means this value must be overriden by the user.
- ws[i].value = 0;
- continue;
- }
- } else if (values[i]->Is<ast::LiteralExpression>() || args[i]->ConstantValue()) {
- value = materialized->ConstantValue();
+ ws[i] = value->As<uint32_t>();
} else {
- AddError(kErrBadExpr, values[i]->source);
- return false;
+ ws[i] = std::nullopt;
}
-
- if (!value) {
- TINT_ICE(Resolver, diagnostics_)
- << "could not resolve constant workgroup_size constant value";
- continue;
- }
- // validator_.Validate and set the default value for this dimension.
- if (value->As<AInt>() < 1) {
- AddError("workgroup_size argument must be at least 1", values[i]->source);
- return false;
- }
-
- ws[i].value = value->As<uint32_t>();
}
current_function_->SetWorkgroupSize(std::move(ws));
diff --git a/src/tint/resolver/resolver_test.cc b/src/tint/resolver/resolver_test.cc
index 1a3c623..79cc139 100644
--- a/src/tint/resolver/resolver_test.cc
+++ b/src/tint/resolver/resolver_test.cc
@@ -993,12 +993,9 @@
auto* func_sem = Sem().Get(func);
ASSERT_NE(func_sem, nullptr);
- EXPECT_EQ(func_sem->WorkgroupSize()[0].value, 1u);
- EXPECT_EQ(func_sem->WorkgroupSize()[1].value, 1u);
- EXPECT_EQ(func_sem->WorkgroupSize()[2].value, 1u);
- EXPECT_EQ(func_sem->WorkgroupSize()[0].overridable_const, nullptr);
- EXPECT_EQ(func_sem->WorkgroupSize()[1].overridable_const, nullptr);
- EXPECT_EQ(func_sem->WorkgroupSize()[2].overridable_const, nullptr);
+ EXPECT_EQ(func_sem->WorkgroupSize()[0], 1u);
+ EXPECT_EQ(func_sem->WorkgroupSize()[1], 1u);
+ EXPECT_EQ(func_sem->WorkgroupSize()[2], 1u);
}
TEST_F(ResolverTest, Function_WorkgroupSize_Literals) {
@@ -1015,12 +1012,9 @@
auto* func_sem = Sem().Get(func);
ASSERT_NE(func_sem, nullptr);
- EXPECT_EQ(func_sem->WorkgroupSize()[0].value, 8u);
- EXPECT_EQ(func_sem->WorkgroupSize()[1].value, 2u);
- EXPECT_EQ(func_sem->WorkgroupSize()[2].value, 3u);
- EXPECT_EQ(func_sem->WorkgroupSize()[0].overridable_const, nullptr);
- EXPECT_EQ(func_sem->WorkgroupSize()[1].overridable_const, nullptr);
- EXPECT_EQ(func_sem->WorkgroupSize()[2].overridable_const, nullptr);
+ EXPECT_EQ(func_sem->WorkgroupSize()[0], 8u);
+ EXPECT_EQ(func_sem->WorkgroupSize()[1], 2u);
+ EXPECT_EQ(func_sem->WorkgroupSize()[2], 3u);
}
TEST_F(ResolverTest, Function_WorkgroupSize_ViaConst) {
@@ -1043,12 +1037,9 @@
auto* func_sem = Sem().Get(func);
ASSERT_NE(func_sem, nullptr);
- EXPECT_EQ(func_sem->WorkgroupSize()[0].value, 16u);
- EXPECT_EQ(func_sem->WorkgroupSize()[1].value, 8u);
- EXPECT_EQ(func_sem->WorkgroupSize()[2].value, 2u);
- EXPECT_EQ(func_sem->WorkgroupSize()[0].overridable_const, nullptr);
- EXPECT_EQ(func_sem->WorkgroupSize()[1].overridable_const, nullptr);
- EXPECT_EQ(func_sem->WorkgroupSize()[2].overridable_const, nullptr);
+ EXPECT_EQ(func_sem->WorkgroupSize()[0], 16u);
+ EXPECT_EQ(func_sem->WorkgroupSize()[1], 8u);
+ EXPECT_EQ(func_sem->WorkgroupSize()[2], 2u);
}
TEST_F(ResolverTest, Function_WorkgroupSize_ViaConst_NestedInitializer) {
@@ -1071,12 +1062,9 @@
auto* func_sem = Sem().Get(func);
ASSERT_NE(func_sem, nullptr);
- EXPECT_EQ(func_sem->WorkgroupSize()[0].value, 8u);
- EXPECT_EQ(func_sem->WorkgroupSize()[1].value, 4u);
- EXPECT_EQ(func_sem->WorkgroupSize()[2].value, 1u);
- EXPECT_EQ(func_sem->WorkgroupSize()[0].overridable_const, nullptr);
- EXPECT_EQ(func_sem->WorkgroupSize()[1].overridable_const, nullptr);
- EXPECT_EQ(func_sem->WorkgroupSize()[2].overridable_const, nullptr);
+ EXPECT_EQ(func_sem->WorkgroupSize()[0], 8u);
+ EXPECT_EQ(func_sem->WorkgroupSize()[1], 4u);
+ EXPECT_EQ(func_sem->WorkgroupSize()[2], 1u);
}
TEST_F(ResolverTest, Function_WorkgroupSize_OverridableConsts) {
@@ -1085,9 +1073,9 @@
// @id(2) override depth = 2i;
// @compute @workgroup_size(width, height, depth)
// fn main() {}
- auto* width = Override("width", ty.i32(), Expr(16_i), Id(0_a));
- auto* height = Override("height", ty.i32(), Expr(8_i), Id(1_a));
- auto* depth = Override("depth", ty.i32(), Expr(2_i), Id(2_a));
+ Override("width", ty.i32(), Expr(16_i), Id(0_a));
+ Override("height", ty.i32(), Expr(8_i), Id(1_a));
+ Override("depth", ty.i32(), Expr(2_i), Id(2_a));
auto* func = Func("main", utils::Empty, ty.void_(), utils::Empty,
utils::Vector{
Stage(ast::PipelineStage::kCompute),
@@ -1099,12 +1087,9 @@
auto* func_sem = Sem().Get(func);
ASSERT_NE(func_sem, nullptr);
- EXPECT_EQ(func_sem->WorkgroupSize()[0].value, 16u);
- EXPECT_EQ(func_sem->WorkgroupSize()[1].value, 8u);
- EXPECT_EQ(func_sem->WorkgroupSize()[2].value, 2u);
- EXPECT_EQ(func_sem->WorkgroupSize()[0].overridable_const, width);
- EXPECT_EQ(func_sem->WorkgroupSize()[1].overridable_const, height);
- EXPECT_EQ(func_sem->WorkgroupSize()[2].overridable_const, depth);
+ EXPECT_EQ(func_sem->WorkgroupSize()[0], std::nullopt);
+ EXPECT_EQ(func_sem->WorkgroupSize()[1], std::nullopt);
+ EXPECT_EQ(func_sem->WorkgroupSize()[2], std::nullopt);
}
TEST_F(ResolverTest, Function_WorkgroupSize_OverridableConsts_NoInit) {
@@ -1113,9 +1098,9 @@
// @id(2) override depth : i32;
// @compute @workgroup_size(width, height, depth)
// fn main() {}
- auto* width = Override("width", ty.i32(), Id(0_a));
- auto* height = Override("height", ty.i32(), Id(1_a));
- auto* depth = Override("depth", ty.i32(), Id(2_a));
+ Override("width", ty.i32(), Id(0_a));
+ Override("height", ty.i32(), Id(1_a));
+ Override("depth", ty.i32(), Id(2_a));
auto* func = Func("main", utils::Empty, ty.void_(), utils::Empty,
utils::Vector{
Stage(ast::PipelineStage::kCompute),
@@ -1127,12 +1112,9 @@
auto* func_sem = Sem().Get(func);
ASSERT_NE(func_sem, nullptr);
- EXPECT_EQ(func_sem->WorkgroupSize()[0].value, 0u);
- EXPECT_EQ(func_sem->WorkgroupSize()[1].value, 0u);
- EXPECT_EQ(func_sem->WorkgroupSize()[2].value, 0u);
- EXPECT_EQ(func_sem->WorkgroupSize()[0].overridable_const, width);
- EXPECT_EQ(func_sem->WorkgroupSize()[1].overridable_const, height);
- EXPECT_EQ(func_sem->WorkgroupSize()[2].overridable_const, depth);
+ EXPECT_EQ(func_sem->WorkgroupSize()[0], std::nullopt);
+ EXPECT_EQ(func_sem->WorkgroupSize()[1], std::nullopt);
+ EXPECT_EQ(func_sem->WorkgroupSize()[2], std::nullopt);
}
TEST_F(ResolverTest, Function_WorkgroupSize_Mixed) {
@@ -1140,7 +1122,7 @@
// const depth = 3i;
// @compute @workgroup_size(8, height, depth)
// fn main() {}
- auto* height = Override("height", ty.i32(), Expr(2_i), Id(0_a));
+ Override("height", ty.i32(), Expr(2_i), Id(0_a));
GlobalConst("depth", ty.i32(), Expr(3_i));
auto* func = Func("main", utils::Empty, ty.void_(), utils::Empty,
utils::Vector{
@@ -1153,12 +1135,9 @@
auto* func_sem = Sem().Get(func);
ASSERT_NE(func_sem, nullptr);
- EXPECT_EQ(func_sem->WorkgroupSize()[0].value, 8u);
- EXPECT_EQ(func_sem->WorkgroupSize()[1].value, 2u);
- EXPECT_EQ(func_sem->WorkgroupSize()[2].value, 3u);
- EXPECT_EQ(func_sem->WorkgroupSize()[0].overridable_const, nullptr);
- EXPECT_EQ(func_sem->WorkgroupSize()[1].overridable_const, height);
- EXPECT_EQ(func_sem->WorkgroupSize()[2].overridable_const, nullptr);
+ EXPECT_EQ(func_sem->WorkgroupSize()[0], 8u);
+ EXPECT_EQ(func_sem->WorkgroupSize()[1], std::nullopt);
+ EXPECT_EQ(func_sem->WorkgroupSize()[2], 3u);
}
TEST_F(ResolverTest, Expr_MemberAccessor_Struct) {
diff --git a/src/tint/sem/function.cc b/src/tint/sem/function.cc
index 6562526..fc7809b 100644
--- a/src/tint/sem/function.cc
+++ b/src/tint/sem/function.cc
@@ -44,7 +44,7 @@
utils::VectorRef<Parameter*> parameters)
: Base(return_type, SetOwner(std::move(parameters), this), EvaluationStage::kRuntime),
declaration_(declaration),
- workgroup_size_{WorkgroupDimension{1}, WorkgroupDimension{1}, WorkgroupDimension{1}},
+ workgroup_size_{1, 1, 1},
return_location_(return_location) {}
Function::~Function() = default;
diff --git a/src/tint/sem/function.h b/src/tint/sem/function.h
index 3f7256a..50d853c 100644
--- a/src/tint/sem/function.h
+++ b/src/tint/sem/function.h
@@ -39,18 +39,10 @@
namespace tint::sem {
-/// WorkgroupDimension describes the size of a single dimension of an entry
-/// point's workgroup size.
-struct WorkgroupDimension {
- /// The size of this dimension.
- uint32_t value;
- /// A pipeline-overridable constant that overrides the size, or nullptr if
- /// this dimension is not overridable.
- const ast::Variable* overridable_const = nullptr;
-};
-
/// WorkgroupSize is a three-dimensional array of WorkgroupDimensions.
-using WorkgroupSize = std::array<WorkgroupDimension, 3>;
+/// Each dimension is a std::optional as a workgroup size can be a constant or override expression.
+/// Override expressions are not known at compilation time, so these will be std::nullopt.
+using WorkgroupSize = std::array<std::optional<uint32_t>, 3>;
/// Function holds the semantic information for function nodes.
class Function final : public Castable<Function, CallTarget> {
diff --git a/src/tint/writer/glsl/generator_impl.cc b/src/tint/writer/glsl/generator_impl.cc
index e332495..21840bc 100644
--- a/src/tint/writer/glsl/generator_impl.cc
+++ b/src/tint/writer/glsl/generator_impl.cc
@@ -102,7 +102,6 @@
namespace {
const char kTempNamePrefix[] = "tint_tmp";
-const char kSpecConstantPrefix[] = "WGSL_SPEC_CONSTANT_";
bool last_is_break_or_fallthrough(const ast::BlockStatement* stmts) {
return IsAnyOf<ast::BreakStatement, ast::FallthroughStatement>(stmts->Last());
@@ -1886,8 +1885,9 @@
[&](const ast::Let* let) { return EmitProgramConstVariable(let); },
[&](const ast::Override*) {
// Override is removed with SubstituteOverride
- TINT_ICE(Writer, diagnostics_)
- << "Override should have been removed by the substitute_override transform.";
+ diagnostics_.add_error(diag::System::Writer,
+ "override expressions should have been removed with the "
+ "SubstituteOverride transform");
return false;
},
[&](const ast::Const*) {
@@ -2104,16 +2104,14 @@
}
out << "local_size_" << (i == 0 ? "x" : i == 1 ? "y" : "z") << " = ";
- if (wgsize[i].overridable_const) {
- auto* global = builder_.Sem().Get<sem::GlobalVariable>(wgsize[i].overridable_const);
- if (!global->Declaration()->Is<ast::Override>()) {
- TINT_ICE(Writer, builder_.Diagnostics())
- << "expected a pipeline-overridable constant";
- }
- out << kSpecConstantPrefix << global->OverrideId().value;
- } else {
- out << std::to_string(wgsize[i].value);
+ if (!wgsize[i].has_value()) {
+ diagnostics_.add_error(
+ diag::System::Writer,
+ "override expressions should have been removed with the SubstituteOverride "
+ "transform");
+ return false;
}
+ out << std::to_string(wgsize[i].value());
}
out << ") in;";
}
diff --git a/src/tint/writer/glsl/generator_impl_function_test.cc b/src/tint/writer/glsl/generator_impl_function_test.cc
index eb0a9a8..6afac16 100644
--- a/src/tint/writer/glsl/generator_impl_function_test.cc
+++ b/src/tint/writer/glsl/generator_impl_function_test.cc
@@ -783,6 +783,25 @@
)");
}
+TEST_F(GlslGeneratorImplTest_Function,
+ Emit_Attribute_EntryPoint_Compute_WithWorkgroup_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));
+ Func("main", utils::Empty, ty.void_(), {},
+ utils::Vector{
+ Stage(ast::PipelineStage::kCompute),
+ WorkgroupSize("width", "height", "depth"),
+ });
+
+ GeneratorImpl& gen = Build();
+
+ EXPECT_FALSE(gen.Generate()) << gen.error();
+ EXPECT_EQ(
+ gen.error(),
+ R"(error: override expressions should have been removed with the SubstituteOverride transform)");
+}
+
TEST_F(GlslGeneratorImplTest_Function, Emit_Function_WithArrayParams) {
Func("my_func", utils::Vector{Param("a", ty.array<f32, 5>())}, ty.void_(),
utils::Vector{
diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc
index 3186955..516acdd 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -81,7 +81,6 @@
namespace {
const char kTempNamePrefix[] = "tint_tmp";
-const char kSpecConstantPrefix[] = "WGSL_SPEC_CONSTANT_";
const char* image_format_to_rwtexture_type(ast::TexelFormat image_format) {
switch (image_format) {
@@ -2842,8 +2841,9 @@
},
[&](const ast::Override*) {
// Override is removed with SubstituteOverride
- TINT_ICE(Writer, diagnostics_)
- << "Override should have been removed by the substitute_override transform.";
+ diagnostics_.add_error(diag::System::Writer,
+ "override expressions should have been removed with the "
+ "SubstituteOverride transform");
return false;
},
[&](const ast::Const*) {
@@ -3044,18 +3044,14 @@
if (i > 0) {
out << ", ";
}
-
- if (wgsize[i].overridable_const) {
- auto* global =
- builder_.Sem().Get<sem::GlobalVariable>(wgsize[i].overridable_const);
- if (!global->Declaration()->Is<ast::Override>()) {
- TINT_ICE(Writer, diagnostics_)
- << "expected a pipeline-overridable constant";
- }
- out << kSpecConstantPrefix << global->OverrideId().value;
- } else {
- out << std::to_string(wgsize[i].value);
+ if (!wgsize[i].has_value()) {
+ diagnostics_.add_error(
+ diag::System::Writer,
+ "override expressions should have been removed with the SubstituteOverride "
+ "transform");
+ return false;
}
+ out << std::to_string(wgsize[i].value());
}
out << ")]" << std::endl;
}
diff --git a/src/tint/writer/hlsl/generator_impl_function_test.cc b/src/tint/writer/hlsl/generator_impl_function_test.cc
index 5c167fa..322b560 100644
--- a/src/tint/writer/hlsl/generator_impl_function_test.cc
+++ b/src/tint/writer/hlsl/generator_impl_function_test.cc
@@ -712,6 +712,25 @@
)");
}
+TEST_F(HlslGeneratorImplTest_Function,
+ Emit_Attribute_EntryPoint_Compute_WithWorkgroup_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));
+ Func("main", utils::Empty, ty.void_(), utils::Empty,
+ utils::Vector{
+ Stage(ast::PipelineStage::kCompute),
+ WorkgroupSize("width", "height", "depth"),
+ });
+
+ GeneratorImpl& gen = Build();
+
+ EXPECT_FALSE(gen.Generate()) << gen.error();
+ EXPECT_EQ(
+ gen.error(),
+ R"(error: override expressions should have been removed with the SubstituteOverride transform)");
+}
+
TEST_F(HlslGeneratorImplTest_Function, Emit_Function_WithArrayParams) {
Func("my_func",
utils::Vector{
diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc
index 5ff5751..d1bab39 100644
--- a/src/tint/writer/msl/generator_impl.cc
+++ b/src/tint/writer/msl/generator_impl.cc
@@ -273,8 +273,9 @@
},
[&](const ast::Override*) {
// Override is removed with SubstituteOverride
- TINT_ICE(Writer, diagnostics_)
- << "Override should have been removed by the substitute_override transform.";
+ diagnostics_.add_error(diag::System::Writer,
+ "override expressions should have been removed with the "
+ "SubstituteOverride transform.");
return false;
},
[&](const ast::Function* func) {
diff --git a/src/tint/writer/spirv/builder.cc b/src/tint/writer/spirv/builder.cc
index f51d808..05f3716 100644
--- a/src/tint/writer/spirv/builder.cc
+++ b/src/tint/writer/spirv/builder.cc
@@ -506,13 +506,17 @@
} else if (func->PipelineStage() == ast::PipelineStage::kCompute) {
auto& wgsize = func_sem->WorkgroupSize();
- // 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)});
+ // Check if the workgroup_size uses pipeline-overridable constants.
+ if (!wgsize[0].has_value() || !wgsize[1].has_value() || !wgsize[2].has_value()) {
+ error_ =
+ "override expressions should have been removed with the SubstituteOverride "
+ "transform";
+ return false;
+ }
+ push_execution_mode(
+ spv::Op::OpExecutionMode,
+ {Operand(id), U32Operand(SpvExecutionModeLocalSize), //
+ Operand(wgsize[0].value()), Operand(wgsize[1].value()), Operand(wgsize[2].value())});
}
for (auto builtin : func_sem->TransitivelyReferencedBuiltinVariables()) {
diff --git a/src/tint/writer/spirv/builder_function_attribute_test.cc b/src/tint/writer/spirv/builder_function_attribute_test.cc
index 9e6c338..60bf062 100644
--- a/src/tint/writer/spirv/builder_function_attribute_test.cc
+++ b/src/tint/writer/spirv/builder_function_attribute_test.cc
@@ -149,6 +149,41 @@
)");
}
+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();
+
+ EXPECT_FALSE(b.GenerateExecutionModes(func, 3)) << b.error();
+ EXPECT_EQ(
+ b.error(),
+ R"(override expressions should have been removed with the SubstituteOverride transform)");
+}
+
+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();
+
+ EXPECT_FALSE(b.GenerateExecutionModes(func, 3)) << b.error();
+ EXPECT_EQ(
+ b.error(),
+ R"(override expressions should have been removed with the SubstituteOverride transform)");
+}
+
TEST_F(BuilderTest, Decoration_ExecutionMode_MultipleFragment) {
auto* func1 = Func("main1", utils::Empty, ty.void_(), utils::Empty,
utils::Vector{
