[tint][ir][val] Add basic const checks for @workgroup_size
This checks that the value for params are in range if they are
constant values, and rejects non-constant value params if
kAllowsOverrides is not set.
It does not do the look-up to determine if more complex cases are
override-expression yet, since there is still implementation side work
needed for overrides.
Issue: 376624999
Change-Id: I370c982d503e32e6db03f5b0ef91dee5bfe2eb75
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/214095
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Commit-Queue: Ryan Harrison <rharrison@chromium.org>
Auto-Submit: Ryan Harrison <rharrison@chromium.org>
diff --git a/src/tint/lang/core/ir/validator.cc b/src/tint/lang/core/ir/validator.cc
index 2221117..e467b07 100644
--- a/src/tint/lang/core/ir/validator.cc
+++ b/src/tint/lang/core/ir/validator.cc
@@ -2015,7 +2015,33 @@
return;
}
- // TODO(376624999): Implement enforcing rules around override and constant expressions
+ if (auto* c = size->As<ir::Constant>()) {
+ if (c->Value()->ValueAs<int64_t>() <= 0) {
+ AddError(func) << "@workgroup_size params must be greater than 0";
+ return;
+ }
+ continue;
+ }
+
+ if (!capabilities_.Contains(Capability::kAllowOverrides)) {
+ AddError(func) << "@workgroup_size param is not a constant value, and IR capability "
+ "'kAllowOverrides' is not set";
+ return;
+ }
+
+ if (auto* r = size->As<ir::InstructionResult>()) {
+ if (r->Instruction() && r->Instruction()->Is<core::ir::Override>()) {
+ continue;
+ }
+
+ // TODO(376624999): Finish implementing checking that this is a override/constant
+ // expression, i.e. calculated from only appropriate values/operations, once override
+ // implementation is complete
+ // for each value/operation used to calculate param:
+ // if not constant expression && not override expression:
+ // fail
+ // pass
+ }
}
}
diff --git a/src/tint/lang/core/ir/validator_test.cc b/src/tint/lang/core/ir/validator_test.cc
index b2fccba..01b25b8 100644
--- a/src/tint/lang/core/ir/validator_test.cc
+++ b/src/tint/lang/core/ir/validator_test.cc
@@ -1054,6 +1054,52 @@
)");
}
+TEST_F(IR_ValidatorTest, Function_WorkgroupSize_ParamsTooSmall) {
+ auto* f = ComputeEntryPoint();
+ f->SetWorkgroupSize({b.Constant(-1_i), b.Constant(2_i), b.Constant(3_i)});
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:1 error: @workgroup_size params must be greater than 0
+%f = @compute @workgroup_size(-1i, 2i, 3i) func():void {
+^^
+
+note: # Disassembly
+%f = @compute @workgroup_size(-1i, 2i, 3i) func():void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Function_WorkgroupSize_OverrideWithoutAllowOverrides) {
+ auto* o = b.Override(ty.u32());
+ auto* f = ComputeEntryPoint();
+ f->SetWorkgroupSize({o->Result(0), o->Result(0), o->Result(0)});
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(
+ res.Failure().reason.Str(),
+ R"(:1:1 error: @workgroup_size param is not a constant value, and IR capability 'kAllowOverrides' is not set
+%f = @compute @workgroup_size(%2, %2, %2) func():void {
+^^
+
+note: # Disassembly
+%f = @compute @workgroup_size(%2, %2, %2) func():void {
+ $B1: {
+ unreachable
+ }
+}
+)");
+}
+
TEST_F(IR_ValidatorTest, Function_Vertex_BasicPosition) {
auto* f = b.Function("my_func", ty.vec4<f32>(), Function::PipelineStage::kVertex);
f->SetReturnBuiltin(BuiltinValue::kPosition);
