tint/resolver: Forbid workgroup pointer parameters
Fixed: tint:1721
Change-Id: I80a2cfc753f928facc165e1c8cd8c5af6c497550
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/108701
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Auto-Submit: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/docs/tint/origin-trial-changes.md b/docs/tint/origin-trial-changes.md
index 1126461..514308f 100644
--- a/docs/tint/origin-trial-changes.md
+++ b/docs/tint/origin-trial-changes.md
@@ -9,6 +9,7 @@
instead (`@vertex`, `@fragment`, or `@compute`). [tint:1503](crbug.com/tint/1503)
* Module-scope `let` is now an error. Use module-scope `const` instead. [tint:1580](crbug.com/tint/1584)
* Reserved words are now an error instead of a deprecation. [tint:1463](crbug.com/tint/1463)
+* You may no longer use pointer parameters in `workgroup` address space. [tint:1721](crbug.com/tint/1721)
### New features
diff --git a/src/tint/resolver/function_validation_test.cc b/src/tint/resolver/function_validation_test.cc
index a4ff9d8..808f556 100644
--- a/src/tint/resolver/function_validation_test.cc
+++ b/src/tint/resolver/function_validation_test.cc
@@ -1071,7 +1071,7 @@
TestParams{ast::AddressSpace::kIn, false},
TestParams{ast::AddressSpace::kOut, false},
TestParams{ast::AddressSpace::kUniform, false},
- TestParams{ast::AddressSpace::kWorkgroup, true},
+ TestParams{ast::AddressSpace::kWorkgroup, false},
TestParams{ast::AddressSpace::kHandle, false},
TestParams{ast::AddressSpace::kStorage, false},
TestParams{ast::AddressSpace::kPrivate, true},
diff --git a/src/tint/resolver/validator.cc b/src/tint/resolver/validator.cc
index 5b38e68..a6d98e2 100644
--- a/src/tint/resolver/validator.cc
+++ b/src/tint/resolver/validator.cc
@@ -853,8 +853,7 @@
if (auto* ref = var->Type()->As<sem::Pointer>()) {
auto address_space = ref->AddressSpace();
if (!(address_space == ast::AddressSpace::kFunction ||
- address_space == ast::AddressSpace::kPrivate ||
- address_space == ast::AddressSpace::kWorkgroup) &&
+ address_space == ast::AddressSpace::kPrivate) &&
IsValidationEnabled(decl->attributes, ast::DisabledValidation::kIgnoreAddressSpace)) {
std::stringstream ss;
ss << "function parameter of pointer type cannot be in '" << address_space
diff --git a/src/tint/transform/module_scope_var_to_entry_point_param.cc b/src/tint/transform/module_scope_var_to_entry_point_param.cc
index 16a622e..599591a 100644
--- a/src/tint/transform/module_scope_var_to_entry_point_param.cc
+++ b/src/tint/transform/module_scope_var_to_entry_point_param.cc
@@ -441,10 +441,12 @@
ctx.dst->Structure(ctx.dst->Sym(), std::move(workgroup_parameter_members));
auto* param_type =
ctx.dst->ty.pointer(ctx.dst->ty.Of(str), ast::AddressSpace::kWorkgroup);
- auto* disable_validation =
- ctx.dst->Disable(ast::DisabledValidation::kEntryPointParameter);
- auto* param = ctx.dst->Param(workgroup_param(), param_type,
- utils::Vector{disable_validation});
+ auto* param = ctx.dst->Param(
+ workgroup_param(), param_type,
+ utils::Vector{
+ ctx.dst->Disable(ast::DisabledValidation::kEntryPointParameter),
+ ctx.dst->Disable(ast::DisabledValidation::kIgnoreAddressSpace),
+ });
ctx.InsertFront(func_ast->params, param);
}
diff --git a/src/tint/transform/module_scope_var_to_entry_point_param_test.cc b/src/tint/transform/module_scope_var_to_entry_point_param_test.cc
index 232737b..8b9ecba 100644
--- a/src/tint/transform/module_scope_var_to_entry_point_param_test.cc
+++ b/src/tint/transform/module_scope_var_to_entry_point_param_test.cc
@@ -325,7 +325,9 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(ModuleScopeVarToEntryPointParamTest, FoldAddressOfDeref) {
+// TODO(crbug.com/tint/1758): Requires support for workgroup pointer parameters, which is
+// unsupported until WGSL 1.1
+TEST_F(ModuleScopeVarToEntryPointParamTest, DISABLED_FoldAddressOfDeref) {
auto* src = R"(
var<workgroup> v : f32;
@@ -364,7 +366,9 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(ModuleScopeVarToEntryPointParamTest, FoldAddressOfDeref_OutOfOrder) {
+// TODO(crbug.com/tint/1758): Requires support for workgroup pointer parameters, which is
+// unsupported until WGSL 1.1
+TEST_F(ModuleScopeVarToEntryPointParamTest, DISABLED_FoldAddressOfDeref_OutOfOrder) {
auto* src = R"(
@compute @workgroup_size(1)
fn main() {
@@ -998,7 +1002,7 @@
}
@compute @workgroup_size(1)
-fn main(@internal(disable_validation__entry_point_parameter) tint_symbol_1 : ptr<workgroup, tint_symbol_2>) {
+fn main(@internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr<workgroup, tint_symbol_2>) {
let tint_symbol : ptr<workgroup, mat2x2<f32>> = &((*(tint_symbol_1)).m);
let x = *(tint_symbol);
}
@@ -1039,7 +1043,7 @@
}
@compute @workgroup_size(1)
-fn main(@internal(disable_validation__entry_point_parameter) tint_symbol_1 : ptr<workgroup, tint_symbol_2>) {
+fn main(@internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr<workgroup, tint_symbol_2>) {
let tint_symbol : ptr<workgroup, array<S2, 4u>> = &((*(tint_symbol_1)).m);
let x = *(tint_symbol);
}
@@ -1080,7 +1084,7 @@
}
@compute @workgroup_size(1)
-fn main(@internal(disable_validation__entry_point_parameter) tint_symbol_1 : ptr<workgroup, tint_symbol_3>) {
+fn main(@internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr<workgroup, tint_symbol_3>) {
let tint_symbol : ptr<workgroup, S> = &((*(tint_symbol_1)).a);
let tint_symbol_2 : ptr<workgroup, S> = &((*(tint_symbol_1)).b);
let x = *(tint_symbol);
@@ -1122,7 +1126,7 @@
}
@compute @workgroup_size(1)
-fn main(@internal(disable_validation__entry_point_parameter) tint_symbol_1 : ptr<workgroup, tint_symbol_3>) {
+fn main(@internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr<workgroup, tint_symbol_3>) {
let tint_symbol : ptr<workgroup, S> = &((*(tint_symbol_1)).a);
let tint_symbol_2 : ptr<workgroup, S> = &((*(tint_symbol_1)).b);
let x = *(tint_symbol);
