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>

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
 

src/tint/resolver/function_validation_test.cc

@@ -1071,7 +1071,7 @@ INSTANTIATE_TEST_SUITE_P(ResolverTest,
                                          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},

src/tint/resolver/validator.cc

@@ -853,8 +853,7 @@ bool Validator::Parameter(const ast::Function* func, const sem::Variable* var) c
     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::kPrivate) &&
-              address_space == ast::AddressSpace::kWorkgroup) &&
             IsValidationEnabled(decl->attributes, ast::DisabledValidation::kIgnoreAddressSpace)) {
             std::stringstream ss;
             ss << "function parameter of pointer type cannot be in '" << address_space

src/tint/transform/module_scope_var_to_entry_point_param.cc

@@ -441,10 +441,12 @@ struct ModuleScopeVarToEntryPointParam::State {
                     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 =
+                auto* param = ctx.dst->Param(
-                    ctx.dst->Disable(ast::DisabledValidation::kEntryPointParameter);
+                    workgroup_param(), param_type,
-                auto* param = ctx.dst->Param(workgroup_param(), param_type,
+                    utils::Vector{
-                                             utils::Vector{disable_validation});
+                        ctx.dst->Disable(ast::DisabledValidation::kEntryPointParameter),
+                        ctx.dst->Disable(ast::DisabledValidation::kIgnoreAddressSpace),
+                    });
                 ctx.InsertFront(func_ast->params, param);
             }
 

src/tint/transform/module_scope_var_to_entry_point_param_test.cc

@@ -325,7 +325,9 @@ fn main() {
     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 @@ fn main() {
     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 @@ struct tint_symbol_2 {
 }
 
 @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 @@ struct tint_symbol_2 {
 }
 
 @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 @@ struct tint_symbol_3 {
 }
 
 @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 @@ struct tint_symbol_3 {
 }
 
 @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);