diff --git a/docs/tint/origin-trial-changes.md b/docs/tint/origin-trial-changes.md index 1126461ed3..514308feed 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 a4ff9d8e6d..808f556777 100644 --- a/src/tint/resolver/function_validation_test.cc +++ b/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}, diff --git a/src/tint/resolver/validator.cc b/src/tint/resolver/validator.cc index 5b38e684c2..a6d98e2e06 100644 --- a/src/tint/resolver/validator.cc +++ b/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()) { 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 16a622e891..599591ab1b 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 @@ 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 = - 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 232737b5e5..8b9ecbad28 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 @@ 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 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) { +fn main(@internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr) { let tint_symbol : ptr> = &((*(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) { +fn main(@internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr) { let tint_symbol : ptr> = &((*(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) { +fn main(@internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr) { let tint_symbol : ptr = &((*(tint_symbol_1)).a); let tint_symbol_2 : ptr = &((*(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) { +fn main(@internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr) { let tint_symbol : ptr = &((*(tint_symbol_1)).a); let tint_symbol_2 : ptr = &((*(tint_symbol_1)).b); let x = *(tint_symbol);