[wgsl] Disallow subgroup matrix in the private address space
The private address space is always considered to be non-uniform,
which means subgroup matrices in var<private> would be unusable
without the diagnostic to disable the uniformity analysis.
Fixed: 410549252
Change-Id: I7b2a7f110a4b346ece518c9af001d5dac0a30a15
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/236818
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/wgsl/inspector/inspector_test.cc b/src/tint/lang/wgsl/inspector/inspector_test.cc
index f891332..a1d3f9d 100644
--- a/src/tint/lang/wgsl/inspector/inspector_test.cc
+++ b/src/tint/lang/wgsl/inspector/inspector_test.cc
@@ -3968,8 +3968,9 @@
auto* src = R"(
enable chromium_experimental_subgroup_matrix;
-var<private> sm: subgroup_matrix_result<f32, 8, 8>;
-@compute @workgroup_size(1) fn foo() { _ = sm; }
+@compute @workgroup_size(1) fn foo() {
+ _ = subgroup_matrix_result<f32, 8, 8>();
+}
)";
Inspector& inspector = Initialize(src);
@@ -3984,8 +3985,7 @@
auto* src = R"(
enable chromium_experimental_subgroup_matrix;
-var<private> sm: subgroup_matrix_result<f32, 8, 8>;
-fn foo() { _ = sm; }
+fn foo() { _ = subgroup_matrix_result<f32, 8, 8>(); }
@compute @workgroup_size(1) fn main() { foo(); }
)";
Inspector& inspector = Initialize(src);
diff --git a/src/tint/lang/wgsl/resolver/resolver.cc b/src/tint/lang/wgsl/resolver/resolver.cc
index 06c6317..b8ae4ee 100644
--- a/src/tint/lang/wgsl/resolver/resolver.cc
+++ b/src/tint/lang/wgsl/resolver/resolver.cc
@@ -4930,11 +4930,10 @@
const_cast<core::type::Type*>(arr->ElemType()), usage);
}
- // Subgroup matrix types can only be declared in the `function` and `private` address space, or
- // in value declarations (the `undefined` address space).
+ // Subgroup matrix types can only be declared in the `function` address space, or in value
+ // declarations (the `undefined` address space).
if (ty->Is<core::type::SubgroupMatrix>() && address_space != core::AddressSpace::kUndefined &&
- address_space != core::AddressSpace::kFunction &&
- address_space != core::AddressSpace::kPrivate) {
+ address_space != core::AddressSpace::kFunction) {
AddError(usage) << "subgroup matrix types cannot be declared in the "
<< style::Enum(address_space) << " address space";
return false;
diff --git a/src/tint/lang/wgsl/resolver/subgroup_matrix_test.cc b/src/tint/lang/wgsl/resolver/subgroup_matrix_test.cc
index ccedf7a..79a9dd1 100644
--- a/src/tint/lang/wgsl/resolver/subgroup_matrix_test.cc
+++ b/src/tint/lang/wgsl/resolver/subgroup_matrix_test.cc
@@ -62,12 +62,12 @@
StringStream kind;
kind << "subgroup_matrix_" << ToString(params.kind);
- auto* var = GlobalVar("m", private_,
- ty(kind.str(), params.el_ast(*this), u32(params.cols), u32(params.rows)));
+ auto* alias =
+ Alias("m", ty(kind.str(), params.el_ast(*this), u32(params.cols), u32(params.rows)));
EXPECT_TRUE(r()->Resolve()) << r()->error();
- auto* m = TypeOf(var)->UnwrapRef()->As<core::type::SubgroupMatrix>();
+ auto* m = TypeOf(alias)->UnwrapRef()->As<core::type::SubgroupMatrix>();
ASSERT_NE(m, nullptr);
EXPECT_EQ(m->Kind(), params.kind);
EXPECT_EQ(m->Type(), params.el_sem(*this));
@@ -90,11 +90,11 @@
TEST_F(ResolverSubgroupMatrixTest, SignedColumnCount) {
Enable(wgsl::Extension::kChromiumExperimentalSubgroupMatrix);
- auto* var = GlobalVar("left", private_, ty("subgroup_matrix_result", ty.f32(), 4_i, 2_u));
+ auto* alias = Alias("left", ty("subgroup_matrix_result", ty.f32(), 4_i, 2_u));
EXPECT_TRUE(r()->Resolve()) << r()->error();
- auto* m = TypeOf(var)->UnwrapRef()->As<core::type::SubgroupMatrix>();
+ auto* m = TypeOf(alias)->UnwrapRef()->As<core::type::SubgroupMatrix>();
ASSERT_NE(m, nullptr);
EXPECT_EQ(m->Columns(), 4u);
EXPECT_EQ(m->Rows(), 2u);
@@ -102,18 +102,18 @@
TEST_F(ResolverSubgroupMatrixTest, SignedRowCount) {
Enable(wgsl::Extension::kChromiumExperimentalSubgroupMatrix);
- auto* var = GlobalVar("left", private_, ty("subgroup_matrix_result", ty.f32(), 4_u, 2_i));
+ auto* alias = Alias("left", ty("subgroup_matrix_result", ty.f32(), 4_u, 2_i));
EXPECT_TRUE(r()->Resolve()) << r()->error();
- auto* m = TypeOf(var)->UnwrapRef()->As<core::type::SubgroupMatrix>();
+ auto* m = TypeOf(alias)->UnwrapRef()->As<core::type::SubgroupMatrix>();
ASSERT_NE(m, nullptr);
EXPECT_EQ(m->Columns(), 4u);
EXPECT_EQ(m->Rows(), 2u);
}
TEST_F(ResolverSubgroupMatrixTest, DeclareTypeWithoutExtension) {
- GlobalVar("left", private_, ty("subgroup_matrix_result", ty.f32(), 8_a, 8_a));
+ Alias("left", ty("subgroup_matrix_result", ty.f32(), 8_a, 8_a));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(
@@ -123,7 +123,7 @@
TEST_F(ResolverSubgroupMatrixTest, MissingTemplateArgs) {
Enable(wgsl::Extension::kChromiumExperimentalSubgroupMatrix);
- GlobalVar("left", private_, ty("subgroup_matrix_result"));
+ Alias("left", ty("subgroup_matrix_result"));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(), R"(error: expected '<' for 'subgroup_matrix_result')");
@@ -131,7 +131,7 @@
TEST_F(ResolverSubgroupMatrixTest, MissingColsAndRows) {
Enable(wgsl::Extension::kChromiumExperimentalSubgroupMatrix);
- GlobalVar("left", private_, ty("subgroup_matrix_result", ty.f32()));
+ Alias("left", ty("subgroup_matrix_result", ty.f32()));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(), R"(error: 'subgroup_matrix_result' requires 3 template arguments)");
@@ -139,7 +139,7 @@
TEST_F(ResolverSubgroupMatrixTest, MissingRows) {
Enable(wgsl::Extension::kChromiumExperimentalSubgroupMatrix);
- GlobalVar("left", private_, ty("subgroup_matrix_result", ty.f32(), 8_a));
+ Alias("left", ty("subgroup_matrix_result", ty.f32(), 8_a));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(), R"(error: 'subgroup_matrix_result' requires 3 template arguments)");
@@ -147,7 +147,7 @@
TEST_F(ResolverSubgroupMatrixTest, MissingType) {
Enable(wgsl::Extension::kChromiumExperimentalSubgroupMatrix);
- GlobalVar("left", private_, ty("subgroup_matrix_result", 8_a, 8_a));
+ Alias("left", ty("subgroup_matrix_result", 8_a, 8_a));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(), R"(error: 'subgroup_matrix_result' requires 3 template arguments)");
@@ -155,7 +155,7 @@
TEST_F(ResolverSubgroupMatrixTest, BadType) {
Enable(wgsl::Extension::kChromiumExperimentalSubgroupMatrix);
- GlobalVar("left", private_, ty("subgroup_matrix_result", ty.bool_(), 8_a, 8_a));
+ Alias("left", ty("subgroup_matrix_result", ty.bool_(), 8_a, 8_a));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(), R"(error: subgroup_matrix element type must be f32, f16, i32, or u32)");
@@ -176,7 +176,7 @@
TEST_F(ResolverSubgroupMatrixTest, ZeroColumnCount) {
Enable(wgsl::Extension::kChromiumExperimentalSubgroupMatrix);
- GlobalVar("left", private_, ty("subgroup_matrix_result", ty.f32(), 0_a, 8_a));
+ Alias("left", ty("subgroup_matrix_result", ty.f32(), 0_a, 8_a));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
@@ -185,7 +185,7 @@
TEST_F(ResolverSubgroupMatrixTest, NegativeColumnCount) {
Enable(wgsl::Extension::kChromiumExperimentalSubgroupMatrix);
- GlobalVar("left", private_, ty("subgroup_matrix_result", ty.f32(), -1_i, 8_a));
+ Alias("left", ty("subgroup_matrix_result", ty.f32(), -1_i, 8_a));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
@@ -207,7 +207,7 @@
TEST_F(ResolverSubgroupMatrixTest, ZeroRowCount) {
Enable(wgsl::Extension::kChromiumExperimentalSubgroupMatrix);
- GlobalVar("left", private_, ty("subgroup_matrix_result", ty.f32(), 8_a, 0_a));
+ Alias("left", ty("subgroup_matrix_result", ty.f32(), 8_a, 0_a));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
@@ -216,7 +216,7 @@
TEST_F(ResolverSubgroupMatrixTest, NegativeRowCount) {
Enable(wgsl::Extension::kChromiumExperimentalSubgroupMatrix);
- GlobalVar("left", private_, ty("subgroup_matrix_result", ty.f32(), 8_a, -1_i));
+ Alias("left", ty("subgroup_matrix_result", ty.f32(), 8_a, -1_i));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
@@ -553,11 +553,15 @@
EXPECT_TRUE(r()->Resolve()) << r()->error();
}
-TEST_F(ResolverSubgroupMatrixTest, PrivateVar_Valid) {
+TEST_F(ResolverSubgroupMatrixTest, PrivateVar_Invalid) {
Enable(wgsl::Extension::kChromiumExperimentalSubgroupMatrix);
GlobalVar("result", private_, ty("subgroup_matrix_result", ty.f32(), 8_a, 8_a));
- EXPECT_TRUE(r()->Resolve()) << r()->error();
+ EXPECT_FALSE(r()->Resolve());
+ EXPECT_THAT(
+ r()->error(),
+ testing::HasSubstr(
+ R"(error: subgroup matrix types cannot be declared in the 'private' address space)"));
}
TEST_F(ResolverSubgroupMatrixTest, StorageVar_Invalid) {
@@ -595,9 +599,13 @@
R"(error: subgroup matrix types cannot be declared in the 'workgroup' address space)"));
}
-TEST_F(ResolverSubgroupMatrixTest, PrivateVar_ArrayElement_Valid) {
+TEST_F(ResolverSubgroupMatrixTest, FunctionVar_ArrayElement_Valid) {
Enable(wgsl::Extension::kChromiumExperimentalSubgroupMatrix);
- GlobalVar("result", private_, ty.array(ty("subgroup_matrix_result", ty.f32(), 8_a, 8_a), 8_a));
+ auto matrix_type = ty("subgroup_matrix_result", ty.f32(), 8_a, 8_a);
+ Func("foo", Empty, ty.void_(),
+ Vector{
+ Decl(Var("result", function, ty.array(matrix_type, 8_a))),
+ });
EXPECT_TRUE(r()->Resolve()) << r()->error();
}
@@ -613,13 +621,16 @@
R"(error: subgroup matrix types cannot be declared in the 'workgroup' address space)"));
}
-TEST_F(ResolverSubgroupMatrixTest, PrivateVar_StructMember_Valid) {
+TEST_F(ResolverSubgroupMatrixTest, FunctionVar_StructMember_Valid) {
Enable(wgsl::Extension::kChromiumExperimentalSubgroupMatrix);
auto* s = Structure("S", Vector{
Member("m", ty("subgroup_matrix_result", ty.f32(), 8_a, 8_a)),
});
- GlobalVar("result", private_, ty.Of(s));
+ Func("foo", Empty, ty.void_(),
+ Vector{
+ Decl(Var("result", function, ty.Of(s))),
+ });
EXPECT_TRUE(r()->Resolve()) << r()->error();
}
@@ -715,52 +726,6 @@
EXPECT_TRUE(r()->Resolve()) << r()->error();
}
-TEST_F(ResolverSubgroupMatrixTest, FragmentShader_ReferenceModuleScope) {
- Enable(wgsl::Extension::kChromiumExperimentalSubgroupMatrix);
- GlobalVar("result", private_, ty("subgroup_matrix_result", ty.f32(), 8_u, 8_u));
- Func("foo", Empty, ty.void_(),
- Vector{
- Assign(Phony(), AddressOf(Ident(Source({12, 34}), "result"))),
- },
- Vector{Stage(ast::PipelineStage::kFragment)});
-
- EXPECT_FALSE(r()->Resolve());
- EXPECT_EQ(r()->error(),
- R"(12:34 error: subgroup matrix type cannot be used in fragment pipeline stage)");
-}
-
-TEST_F(ResolverSubgroupMatrixTest, FragmentShader_ReferenceModuleScopeInArray) {
- Enable(wgsl::Extension::kChromiumExperimentalSubgroupMatrix);
- GlobalVar("result", private_, ty.array(ty("subgroup_matrix_result", ty.f32(), 8_u, 8_u), 4_a));
- Func("foo", Empty, ty.void_(),
- Vector{
- Assign(Phony(), AddressOf(Ident(Source({12, 34}), "result"))),
- },
- Vector{Stage(ast::PipelineStage::kFragment)});
-
- EXPECT_FALSE(r()->Resolve());
- EXPECT_EQ(r()->error(),
- R"(12:34 error: subgroup matrix type cannot be used in fragment pipeline stage)");
-}
-
-TEST_F(ResolverSubgroupMatrixTest, FragmentShader_ReferenceModuleScopeInStruct) {
- Enable(wgsl::Extension::kChromiumExperimentalSubgroupMatrix);
-
- auto* s = Structure("S", Vector{
- Member("m", ty("subgroup_matrix_result", ty.f32(), 8_a, 8_a)),
- });
- GlobalVar("result", private_, ty.Of(s));
- Func("foo", Empty, ty.void_(),
- Vector{
- Assign(Phony(), AddressOf(Ident(Source({12, 34}), "result"))),
- },
- Vector{Stage(ast::PipelineStage::kFragment)});
-
- EXPECT_FALSE(r()->Resolve());
- EXPECT_EQ(r()->error(),
- R"(12:34 error: subgroup matrix type cannot be used in fragment pipeline stage)");
-}
-
TEST_F(ResolverSubgroupMatrixTest, FragmentShader_FunctionVar) {
Enable(wgsl::Extension::kChromiumExperimentalSubgroupMatrix);
Func(
