encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

validation: Allow unsigned workgroup_size component 

Bug: tint:923
Change-Id: I7bd7d22279d9c6ce4c3225bdfd8693261b9084f9
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/58121
Commit-Queue: Sarah Mashayekhi <sarahmashay@google.com>
Auto-Submit: Sarah Mashayekhi <sarahmashay@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
This commit is contained in:
Sarah 2021-07-14 23:33:24 +00:00 committed by Tint LUCI CQ
parent 3242d3e8c9
commit 58bbe32cf4
2 changed files with 164 additions and 35 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

139
src/resolver/function_validation_test.cc
View File

@ -357,6 +357,108 @@ TEST_F(ResolverFunctionValidationTest, FunctionParamsConst) {
"declared here:");
}
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_GoodType_ConstU32) {
// let x = 4u;
// let x = 8u;
// [[stage(compute), workgroup_size(x, y, 16u]
// fn main() {}
GlobalConst("x", ty.u32(), Expr(4u));
GlobalConst("y", ty.u32(), Expr(8u));
Func("main", {}, ty.void_(), {},
{Stage(ast::PipelineStage::kCompute),
WorkgroupSize(Expr("x"), Expr("y"), Expr(16u))});
EXPECT_TRUE(r()->Resolve()) << r()->error();
}
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_GoodType_U32) {
// [[stage(compute), workgroup_size(1u, 2u, 3u)]
// fn main() {}
Func("main", {}, ty.void_(), {},
{Stage(ast::PipelineStage::kCompute),
WorkgroupSize(Source{{12, 34}}, Expr(1u), Expr(2u), Expr(3u))});
EXPECT_TRUE(r()->Resolve()) << r()->error();
}
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_MismatchTypeU32) {
// [[stage(compute), workgroup_size(1u, 2u, 3)]
// fn main() {}
Func("main", {}, ty.void_(), {},
{Stage(ast::PipelineStage::kCompute),
WorkgroupSize(Expr(1u), Expr(2u), Expr(Source{{12, 34}}, 3))});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
"12:34 error: workgroup_size parameters must be of the same "
"type, either i32 or u32");
}
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_MismatchTypeI32) {
// [[stage(compute), workgroup_size(1, 2u, 3)]
// fn main() {}
Func("main", {}, ty.void_(), {},
{Stage(ast::PipelineStage::kCompute),
WorkgroupSize(Expr(1), Expr(Source{{12, 34}}, 2u), Expr(3))});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
"12:34 error: workgroup_size parameters must be of the same "
"type, either i32 or u32");
}
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_TypeMismatch) {
// let x = 64u;
// [[stage(compute), workgroup_size(1, x)]
// fn main() {}
GlobalConst("x", ty.u32(), Expr(64u));
Func("main", {}, ty.void_(), {},
{Stage(ast::PipelineStage::kCompute),
WorkgroupSize(Expr(1), Expr(Source{Source::Location{12, 34}}, "x"))});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
"12:34 error: workgroup_size parameters must be of the same "
"type, either i32 or u32");
}
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_TypeMismatch2) {
// let x = 64u;
// let y = 32;
// [[stage(compute), workgroup_size(x, y)]
// fn main() {}
GlobalConst("x", ty.u32(), Expr(64u));
GlobalConst("y", ty.i32(), Expr(32));
Func("main", {}, ty.void_(), {},
{Stage(ast::PipelineStage::kCompute),
WorkgroupSize(Expr("x"), Expr(Source{Source::Location{12, 34}}, "y"))});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
"12:34 error: workgroup_size parameters must be of the same "
"type, either i32 or u32");
}
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Mismatch_ConstU32) {
// let x = 4u;
// let x = 8u;
// [[stage(compute), workgroup_size(x, y, 16]
// fn main() {}
GlobalConst("x", ty.u32(), Expr(4u));
GlobalConst("y", ty.u32(), Expr(8u));
Func("main", {}, ty.void_(), {},
{Stage(ast::PipelineStage::kCompute),
WorkgroupSize(Expr("x"), Expr("y"),
Expr(Source{Source::Location{12, 34}}, 16))});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
"12:34 error: workgroup_size parameters must be of the same "
"type, either i32 or u32");
}
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Literal_BadType) {
// [[stage(compute), workgroup_size(64.0)]
// fn main() {}
@ -368,8 +470,8 @@ TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Literal_BadType) {
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
"12:34 error: workgroup_size parameter must be a literal i32 or an "
"i32 module-scope constant");
"12:34 error: workgroup_size parameter must be either literal or "
"module-scope constant of type i32 or u32");
}
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Literal_Negative) {
@ -382,9 +484,8 @@ TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Literal_Negative) {
Source{Source::Location{12, 34}}, Literal(-2)))});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(
r()->error(),
"12:34 error: workgroup_size parameter must be a positive i32 value");
EXPECT_EQ(r()->error(),
"12:34 error: workgroup_size parameter must be at least 1");
}
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Literal_Zero) {
@ -397,9 +498,8 @@ TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Literal_Zero) {
Source{Source::Location{12, 34}}, Literal(0)))});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(
r()->error(),
"12:34 error: workgroup_size parameter must be a positive i32 value");
EXPECT_EQ(r()->error(),
"12:34 error: workgroup_size parameter must be at least 1");
}
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_BadType) {
@ -413,8 +513,8 @@ TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_BadType) {
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
"12:34 error: workgroup_size parameter must be a literal i32 or an "
"i32 module-scope constant");
"12:34 error: workgroup_size parameter must be either literal or "
"module-scope constant of type i32 or u32");
}
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_Negative) {
@ -427,9 +527,8 @@ TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_Negative) {
WorkgroupSize(Expr(Source{Source::Location{12, 34}}, "x"))});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(
r()->error(),
"12:34 error: workgroup_size parameter must be a positive i32 value");
EXPECT_EQ(r()->error(),
"12:34 error: workgroup_size parameter must be at least 1");
}
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_Zero) {
@ -442,9 +541,8 @@ TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_Zero) {
WorkgroupSize(Expr(Source{Source::Location{12, 34}}, "x"))});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(
r()->error(),
"12:34 error: workgroup_size parameter must be a positive i32 value");
EXPECT_EQ(r()->error(),
"12:34 error: workgroup_size parameter must be at least 1");
}
TEST_F(ResolverFunctionValidationTest,
@ -459,9 +557,8 @@ TEST_F(ResolverFunctionValidationTest,
WorkgroupSize(Expr(Source{Source::Location{12, 34}}, "x"))});
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(
r()->error(),
"12:34 error: workgroup_size parameter must be a positive i32 value");
EXPECT_EQ(r()->error(),
"12:34 error: workgroup_size parameter must be at least 1");
}
TEST_F(ResolverFunctionValidationTest, WorkgroupSize_NonConst) {
@ -475,8 +572,8 @@ TEST_F(ResolverFunctionValidationTest, WorkgroupSize_NonConst) {
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
"12:34 error: workgroup_size parameter must be a literal i32 or an "
"i32 module-scope constant");
"12:34 error: workgroup_size parameter must be either literal or "
"module-scope constant of type i32 or u32");
}
TEST_F(ResolverFunctionValidationTest, ReturnIsConstructible_NonPlain) {

60
src/resolver/resolver.cc
View File

@ -1608,6 +1608,8 @@ bool Resolver::Function(ast::Function* func) {
if (auto* workgroup =
ast::GetDecoration<ast::WorkgroupDecoration>(func->decorations())) {
auto values = workgroup->values();
auto is_i32 = false;
auto is_less_than_one = true;
for (int i = 0; i < 3; i++) {
// Each argument to this decoration can either be a literal, an
// identifier for a module-scope constants, or nullptr if not specified.
@ -1619,7 +1621,7 @@ bool Resolver::Function(ast::Function* func) {
Mark(values[i]);
int32_t value = 0;
uint32_t value = 0;
if (auto* ident = values[i]->As<ast::IdentifierExpression>()) {
// We have an identifier of a module-scope constant.
if (!Identifier(ident)) {
@ -1628,10 +1630,10 @@ bool Resolver::Function(ast::Function* func) {
VariableInfo* var;
if (!variable_stack_.get(ident->symbol(), &var) ||
!(var->declaration->is_const() && var->type->Is<sem::I32>())) {
!(var->declaration->is_const() && var->type->is_integer_scalar())) {
AddError(
"workgroup_size parameter must be a literal i32 or an i32 "
"module-scope constant",
"workgroup_size parameter must be either literal or module-scope "
"constant of type i32 or u32",
values[i]->source());
return false;
}
@ -1646,12 +1648,28 @@ bool Resolver::Function(ast::Function* func) {
if (constructor) {
// Resolve the constructor expression to use as the default value.
auto val = ConstantValueOf(constructor);
if (!val.IsValid() || !val.Type()->Is<sem::I32>()) {
if (!val.IsValid() || !val.Type()->is_integer_scalar()) {
TINT_ICE(Resolver, diagnostics_)
<< "failed to resolve workgroup_size constant value";
return false;
}
value = val.Elements()[0].i32;
if (i == 0) {
is_i32 = val.Type()->Is<sem::I32>();
} else {
if (is_i32 != val.Type()->Is<sem::I32>()) {
AddError(
"workgroup_size parameters must be of the same type, "
"either i32 or u32",
values[i]->source());
return false;
}
}
is_less_than_one =
is_i32 ? val.Elements()[0].i32 < 1 : val.Elements()[0].u32 < 1;
value = is_i32 ? static_cast<uint32_t>(val.Elements()[0].i32)
: val.Elements()[0].u32;
} else {
// No constructor means this value must be overriden by the user.
info->workgroup_size[i].value = 0;
@ -1661,22 +1679,36 @@ bool Resolver::Function(ast::Function* func) {
values[i]->As<ast::ScalarConstructorExpression>()) {
// We have a literal.
Mark(scalar->literal());
auto* i32_literal = scalar->literal()->As<ast::IntLiteral>();
if (!i32_literal) {
auto* literal = scalar->literal()->As<ast::IntLiteral>();
if (!literal) {
AddError(
"workgroup_size parameter must be a literal i32 or an i32 "
"module-scope constant",
"workgroup_size parameter must be either literal or module-scope "
"constant of type i32 or u32",
values[i]->source());
return false;
}
value = i32_literal->value_as_i32();
if (i == 0) {
is_i32 = literal->Is<ast::SintLiteral>();
} else {
if (literal->Is<ast::SintLiteral>() != is_i32) {
AddError(
"workgroup_size parameters must be of the same type, "
"either i32 or u32",
values[i]->source());
return false;
}
}
is_less_than_one =
is_i32 ? literal->value_as_i32() < 1 : literal->value_as_u32() < 1;
value = is_i32 ? static_cast<uint32_t>(literal->value_as_i32())
: literal->value_as_u32();
}
// Validate and set the default value for this dimension.
if (value < 1) {
AddError("workgroup_size parameter must be a positive i32 value",
if (is_less_than_one) {
AddError("workgroup_size parameter must be at least 1",
values[i]->source());
return false;
}